CUDA 中的卷积、带滤波器的数组

Question

我正在尝试使用共享内存在 GPU 上对 256x256 数据数组与 3x3 滤波器进行卷积。我知道我要将数组分成块，然后在每个块中应用过滤器。这最终意味着块沿边缘重叠，并且需要在没有数据的边缘周围进行一些填充，以便过滤器正常工作。

int 网格 = (256/(16+3-1))*(256/(16+3-1)) 其中 256 是我的数组的长度或宽度，16 是我的共享内存中的块的长度或宽度，3 是我的过滤器的长度或宽度，我减一以使其均匀。

int thread = (16+3-1)*(16+3-1)

现在我调用我的内核<<>>(output, input, 256) 输入和输出是大小为 256*256 的数组

__global__ void kernel(float *input, float *output, int size)
{
    __shared__ float tile[16+3-1][16+3-1];
    blockIdx.x = bIdx;
    blockIdy.y = bIdy;
    threadIdx.x = tIdx;
    threadIdy.y = tIdy

    //i is for input
    unsigned int iX = bIdx * 3 + tIdx;
    unsigned int iY = bIdy * 3 + tIdy;

    if (tIdx == 0 || tIdx == width || tIdy == 0 || tIdy == height)
    {
        //this will pad the outside edges
        block[tIdy][tIdx] = 0;
    }
    else 
    {
        //This will fill in the block with real data
        unsigned int iin = iY * size + iX;
        block[tIdy][tIdx] = idata[iin];
    }

    __syncthreads();

    //I believe is above is correct; below, where I do the convolution, I feel is wrong
    float result = 0;
    for(int fX=-N/2; fX<=N/2; fX++){
        for(int fY=-N/2; fY<=N/2; fY++){
            if(iY+fX>=0 && iY+fX<size && iX+fY>=0 && iX+fY<size)
                result+=tile[tIdx+fX][tIdy+fY];
        }
    }
    output[iY*size+iX] = result/(3*3);
}

当我运行代码时，如果运行卷积部分，则会出现内核错误。有什么见解吗？或者建议？

Answer 1

查看 sobelFilter SDK 示例。

它使用纹理来处理边缘情况，稍微过度获取块（但纹理缓存使之更加高效），并使用共享内存进行处理。

关于共享内存的微妙之处在于，如果读取相邻字节，则会出现 4 路存储体冲突。解决这个问题的一种方法（如 sobelFilter 示例所示）是将循环展开 4 倍并访问每四个字节。