CUDA 2d 卷积边界不正确答案

【问题标题】：CUDA 2d convolution boundary incorrectCUDA 2d 卷积边界不正确
【发布时间】：2014-04-01 04:49:31
【问题描述】：

我用幼稚的方式实现了一个 CUDA 2D 卷积码，无法正确获取边界值。错误发生在顶部和左侧边界，宽度为过滤器的一半。例如，如果我的过滤器是 7x7，则错误位于顶部 3 像素和左侧 3 像素（与 C 结果相比）。有人可以帮我解决这个错误吗？非常感谢您的帮助！

附上我的cuda代码和c代码：

#define ISIZE 32//input image size ISIZE*ISIZE
#define MASK_RADIUS 3
#define MASK_WIDTH (2 * MASK_RADIUS + 1)
const int FILTER_SIZE = MASK_WIDTH * MASK_WIDTH * sizeof(float);
__device__ __constant__ float d_filter[FILTER_SIZE];

__global__ void convolution2D_cuda(float* d_Result, float* d_Data, int dataH, int dataW)
{
// global mem address for this thread
const int gLoc = threadIdx.x + blockIdx.x * blockDim.x +
                 (threadIdx.y + blockIdx.y * blockDim.y) * dataW; 

float sum = 0;
float value = 0;

for(int i = -MASK_RADIUS; i <= MASK_RADIUS; i++) //row wise
{ 
    for (int j = -MASK_RADIUS; j <= MASK_RADIUS; j++) //col wise
    { 
        // check row 
        if ( (blockIdx.x == 0) && ((threadIdx.x + j) < 0) ) //left apron
            value = 0;
        else if ( blockIdx.x == (gridDim.x -1) && (threadIdx.x + j) > (blockDim.x-1) ) //right apron
            value = 0;          
        else {
            // check col
            if ( blockIdx.y == 0 && (threadIdx.y + i) < 0) //top apron
                value = 0;
            else if ( blockIdx.y == (gridDim.y-1) && (threadIdx.y + i) > (blockDim.y-1) ) //bottom apron
                value = 0;
            else // load data
                value = d_Data[gLoc + i * dataW + j];
        }
        //2d array case: non-separable filter
        sum += value * d_filter[ (MASK_RADIUS - i) * MASK_WIDTH + (MASK_RADIUS - j) ];
    }
}
d_Result[gLoc] = sum;
}

//c code
void convolution2D_cpu(float* result, float* input, float* filter, int dataW, int dataH, int k_Width, int k_Height, int radiusY, int radiusX)
{
      int y, x, ky, kx;
      for (y = 0; y < dataH; y++) { //row
         for (x = 0; x < dataW; x++) {
         result[y*dataW + x] = 0;
         float sum=0;
         for(ky = -radiusY; ky <= radiusY; ky++) {
                 for(kx = -radiusX; kx <= radiusX; kx++) {
                int dy = y + ky;
            int dx = x + kx;
            if (dy >= 0 && dy < dataH) //left & upper borders
                if (dx >= 0 && dx < dataW) //right & lower borders
            sum += input[dy*dataW + dx] * filter[(radiusY-ky)*k_Width + (radiusX - kx)];            
                }
            }
        result[y*dataW+x] = sum;
        }
    }
}



Part of the main() code is :

    dim3 blocks(16, 16);
dim3 grids(width/16, height/16);


    checkCudaErrors( cudaMalloc( (void **)&d_data, data_size ));
checkCudaErrors( cudaMalloc( (void **)&d_result, data_size ));
    checkCudaErrors( cudaMemcpy(d_data, indata, data_size, cudaMemcpyHostToDevice) );
    checkCudaErrors( cudaThreadSynchronize() );

convolution2D_cuda<<<grids, blocks>>>(d_result, d_data, width, height);
    checkCudaErrors( cudaThreadSynchronize() );

    checkCudaErrors( cudaMemcpy(output, d_result, data_size, cudaMemcpyDeviceToHost) );
    checkCudaErrors( cudaThreadSynchronize() );

    //check with result of CPU
    convolution2D_cpu(c_result, indata, filter, width, height, len, len, MASK_RADIUS, MASK_RADIUS);

【问题讨论】：

您使用的块和网格尺寸是多少？能否也给出调用内核的那部分代码？
只是一个简单的问题，在 CPU 部分，您有两个具有不同边界的不同循环（radiusY，radiusX），但在 GPU 部分，您的循环只有一个边界（@987654324 @) 这可能是错误，不是吗？由于主要代码不完整。我无法对此下结论。但它可能。
是的，radiusX 和 radiusY 设置为等于 MASK_RADIUS 以进行测试。
您的数组d_filter[FILTER_SIZE] 的大小不正确。那应该是d_filter[FILTER_SIZE/sizeof(float)]
你的网格有你图片的大小吗？

标签： cuda

【解决方案1】：

我要解开这个谜团。错误发生在线程索引计算上。 threadIdx 是 uint，nvcc 认为 (threadIdx.x + j) 为 unsigned int。前任。如果 j 为 -1，则解释为 4294967295 (ffffffff)，边界索引不正确。

【讨论】：