CUDA处理16位图像输出异常求助：输出为输入1/4叠加

你遇到的输出图像只有1/4内容且叠加的问题，核心原因是核函数中的坐标判断逻辑错误，导致大量线程错误地写入了重复的像素位置。下面一步步分析并给出修正方案：

问题根源分析 #

你的核函数中使用了错误的线程有效性判断条件：

if(x < width * height)

这里的x是线程对应的横向像素坐标（范围是0到图像宽度-1），而width * height是图像的总像素数。这个条件会让所有x小于总像素数的线程执行，包括那些x >= width或者y >= height的越界线程。

举个例子：当x=640（刚好等于你的图像宽度），y=0时，这个线程会通过判断，但计算出的索引index = y*width +x = 0*640+640=640，这实际上对应y=1, x=0的像素位置。多个越界线程会重复写入同一个像素，最终导致图像出现叠加、只显示1/4内容的异常。

修正方案 #

1. 修正核函数的判断逻辑 #

将核函数中的判断条件改为检查x是否在图像宽度范围内，且y是否在图像高度范围内，这样只有有效像素位置的线程才会执行：

__global__ void convolve(uint16_t *source, int width, int height, uint16_t *destination)
{
    // Calculate our pixel's location
    int x = (blockIdx.x * blockDim.x) + threadIdx.x;
    int y = (blockIdx.y * blockDim.y) + threadIdx.y;
    // Only execute for valid pixels
    if(x < width && y < height)
    {
        int index = (y * width) + x;
        destination[index] = source[index];
    }
}

2. 优化线程块的计算（可选） #

你的线程块计算dim3 blocks (w/16 + 1, h/16 + 1)是为了处理非16倍数的图像尺寸，这个逻辑是对的，但如果图像尺寸刚好是16的倍数（比如你设置的640x480，640/16=40，480/16=30），会多创建一个空的线程块，不过因为我们已经添加了正确的越界判断，这部分不会影响结果，只是会有少量空闲线程，不影响性能。

3. 其他注意事项 #

• 你使用的cudaThreadSynchronize()在CUDA 9及以后版本已经被弃用，建议替换为cudaDeviceSynchronize()，功能完全一致。
• 确保cudaMemcpy的字节数计算正确：你使用的640*480*sizeof(uint16_t)是正确的，因为CV_16UC1格式每个像素占2字节。

验证修正效果 #

修改核函数后，重新编译运行，应该就能得到和输入图像完全一致的输出（因为你的核函数目前只是做一对一复制），后续添加滤镜逻辑也不会再出现坐标错误的问题。

内容的提问来源于stack exchange，提问作者M. P.