Skip to content

cudev: fix CUDA texture pitch alignment for createContinuous GpuMat #4068

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
Shubh3155 wants to merge 1 commit into
base: 4.x
Choose a base branch
from
+17 −6
Open

cudev: fix CUDA texture pitch alignment for createContinuous GpuMat #4068

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
23 changes: 17 additions & 6 deletions modules/cudev/include/opencv2/cudev/ptr2d/texture.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -176,17 +176,28 @@ namespace cv { namespace cudev {
texRes.res.pitch2D.height = rows;
texRes.res.pitch2D.width = cols;
// temporary fix for single row/columns until TexturePtr is reworked
Copy link

Copilot AI Jan 23, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The comment "temporary fix for single row/columns until TexturePtr is reworked" is now outdated. The fix now also handles pitch alignment issues, not just single row/column cases. Consider updating the comment to reflect the expanded scope, such as: "temporary fix for single row/columns and misaligned pitch until TexturePtr is reworked".

Suggested change
// temporary fix for single row/columns until TexturePtr is reworked
// temporary fix for single row/columns and misaligned pitch until TexturePtr is reworked

Copilot uses AI. Check for mistakes.
if (rows == 1 || cols == 1) {
int currentDevice = 0;
CV_CUDEV_SAFE_CALL(cudaGetDevice(&currentDevice));

cudaDeviceProp prop;
CV_CUDEV_SAFE_CALL(cudaGetDeviceProperties(&prop, currentDevice));

Comment on lines +182 to +184
Copy link

Copilot AI Jan 23, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Performance consideration: cudaGetDevice and cudaGetDeviceProperties are called every time this function is invoked. For frequently created textures, this could add overhead. Consider caching the device properties or checking if there's a pattern in the codebase for reusing device properties. However, this is a minor concern since texture creation is typically not done in a tight loop, and the fix is necessary for correctness.

Suggested change
cudaDeviceProp prop;
CV_CUDEV_SAFE_CALL(cudaGetDeviceProperties(&prop, currentDevice));
// Cache device properties per device to avoid repeated queries.
static int cachedDevice = -1;
static cudaDeviceProp cachedProp;
static bool propInitialized = false;
if (!propInitialized || currentDevice != cachedDevice)
{
CV_CUDEV_SAFE_CALL(cudaGetDeviceProperties(&cachedProp, currentDevice));
cachedDevice = currentDevice;
propInitialized = true;
}
const cudaDeviceProp& prop = cachedProp;

Copilot uses AI. Check for mistakes.
if (rows == 1 || cols == 1 ||
(step % prop.texturePitchAlignment) != 0)
{
size_t dStep = 0;
CV_CUDEV_SAFE_CALL(cudaMallocPitch(&internalSrc, &dStep, cols * sizeof(T1), rows));
CV_CUDEV_SAFE_CALL(cudaMemcpy2D(internalSrc, dStep, data, step, cols * sizeof(T1), rows, cudaMemcpyDeviceToDevice));
CV_CUDEV_SAFE_CALL(cudaMemcpy2D(internalSrc, dStep, data, step,
cols * sizeof(T1), rows,
cudaMemcpyDeviceToDevice));
texRes.res.pitch2D.devPtr = internalSrc;
texRes.res.pitch2D.pitchInBytes = dStep;
}
else {
texRes.res.pitch2D.devPtr = data;
texRes.res.pitch2D.pitchInBytes = step;
}
Comment on lines +195 to 199
Copy link

Copilot AI Jan 23, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Inconsistent indentation: the closing brace and else statement use incorrect indentation. They should align with the opening brace on line 187, using 12 spaces instead of 11.

Suggested change
}
else {
texRes.res.pitch2D.devPtr = data;
texRes.res.pitch2D.pitchInBytes = step;
}
}
else {
texRes.res.pitch2D.devPtr = data;
texRes.res.pitch2D.pitchInBytes = step;
}

Copilot uses AI. Check for mistakes.
else {
texRes.res.pitch2D.devPtr = data;
texRes.res.pitch2D.pitchInBytes = step;
}

texRes.res.pitch2D.desc = cudaCreateChannelDesc<T1>();
createTextureObject(texRes, normalizedCoords, filterMode, addressMode, readMode);
}
Expand Down
Loading