cuda - CUDA 内核中的 2D 图像索引错误

Question

我正在使用 CUDA 对图像进行线性过滤。我使用 2D 线程块和 2D 网格来使问题变得自然。这是我的索引方式:( 高度 和 宽度 是图像尺寸)

dim3 BlockDim(16,16);

dim3 GridDim;
GridDim.x = (width + 15) / 16;
GridDim.y = (height + 15) / 16;

在内核中，我按如下方式访问位置:

unsigned int xIndex = blockIdx.x*16+ threadIdx.x;
unsigned int yIndex = blockIdx.y*16+ threadIdx.y;
unsigned int tid = yIndex * width + xIndex;

我想返回四个边界(稍后我会满足它们)。我这样做:

if(yIndex>=height-N || xIndex>=width-N || yIndex<N || xIndex<N)
  return;

其中 N 是我不想计算的每个边界处的像素数。

问题:

该代码在所有标准图像尺寸上运行良好。但是对于某些随机图像大小，它显示对角线。例如，在我的情况下，500x333 图像(即使没有尺寸是 16 的倍数)显示正确的输出，而 450x365 显示输出中的对角线。即使我只返回网格的额外线程，问题仍然存在，没有其他类似的东西:

if(yIndex>=height || xIndex>=width)
return;

代码保持不变，一些输入运行良好，而另一些则没有。任何人都可以发现错误吗？我在这里附上了输入和输出样本: IMAGES谢谢!

更新:

内核代码(简化为返回输入图像，但出现同样的问题)

__global__ void filter_8u_c1_kernel(unsigned char* in, unsigned char* out, int width, int height, float* filter, int fSize)
{
    unsigned int xIndex = blockIdx.x*BLOCK_SIZE + threadIdx.x;
    unsigned int yIndex = blockIdx.y*BLOCK_SIZE + threadIdx.y;
    unsigned int tid = yIndex * width + xIndex;

    unsigned int N = filterSize/2;

    if(yIndex>=height-N || xIndex>=width-N || yIndex<N || xIndex<N)
        return;

       /*Filter code removed, still gives the same problem*/

    out[tid] = in[tid];
}

更新 2:

我也删除了 返回 通过反转 的声明如果 健康)状况。但问题仍然存在。

if(yIndex<=height-N && xIndex<=width-N && yIndex>N && xIndex>N){

  /*Kernel Code*/

}

Answer 1

您还没有很好地描述很多事情，但是根据您发布的信息，我构建了一个合理的重现案例，其参数与您说失败的案例相匹配(450 x 364 与filterSize=5):

#include <stdio.h>
#include <assert.h>

template<int filterSize>
__global__ void filter_8u_c1_kernel(unsigned char* in, unsigned char* out, int width, int height, float* filter, int fSize)
{
    unsigned int xIndex = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int yIndex = blockIdx.y*blockDim.y + threadIdx.y;
    unsigned int tid = yIndex * width + xIndex;

    unsigned int N = filterSize/2;

    if(yIndex>=height-N || xIndex>=width-N || yIndex<N || xIndex<N)
        return;

    out[tid] = in[tid];
}

int main(void)
{
    const int width = 450, height = 365, filterSize=5;
    const size_t isize = sizeof(unsigned char) * size_t(width * height);
    unsigned char * _in, * _out, * out;

    assert( cudaMalloc((void **)&_in, isize) == cudaSuccess ); 
    assert( cudaMalloc((void **)&_out, isize) == cudaSuccess ); 
    assert( cudaMemset(_in, 'Z', isize) == cudaSuccess );
    assert( cudaMemset(_out, 'A', isize) == cudaSuccess );

    const dim3 BlockDim(16,16);
    dim3 GridDim;
    GridDim.x = (width + BlockDim.x - 1) / BlockDim.x;
    GridDim.y = (height + BlockDim.y - 1) / BlockDim.y;

    filter_8u_c1_kernel<filterSize><<<GridDim,BlockDim>>>(_in,_out,width,height,0,0);
    assert( cudaPeekAtLastError() == cudaSuccess );

    out = (unsigned char *)malloc(isize);
    assert( cudaMemcpy(out, _out, isize, cudaMemcpyDeviceToHost) == cudaSuccess);

    for(int i=0; i<width; i++) {
        fprintf(stdout, "%d: ", i);
        for(int j=0; j<height; j++) {
            unsigned int idx = i + j*width;
            fprintf(stdout, "%c", out[idx]);
        }
        fprintf(stdout, "\n");
    }

    return cudaThreadExit();
}

运行时它完全符合我的预期，除了第一行和最后两行以及中间所有行中的第一和最后两个条目外，到处都用输入覆盖输出内存。这是在具有计算 1.2 GPU 的 OS X 10.6.5 上使用 CUDA 3.2 运行的。因此，无论您的代码发生了什么，在我的重现案例中都不会发生，这要么意味着我误解了您所写的内容，要么是您未描述的其他原因导致了问题。