CUDA中的卷积，带滤波器的数组答案

【问题标题】：Convolution, array with filter, in CUDACUDA中的卷积，带滤波器的数组
【发布时间】：2011-05-02 10:43:04
【问题描述】：

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

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

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);
}

当我运行代码时，如果我运行卷积部分，我会得到一个内核错误。有什么见解吗？或建议？

【问题讨论】：

取决于您拥有的 GPU 以及您尝试运行共享内存分配的线程数可能会太大。您可能需要重新考虑您的实现，因为您将无法以如此大的分配运行许多线程。
CUDA SDK 有几个卷积示例。您可能想与之进行比较，看看您的实现有何不同。 CUFFT 库也是另一种可能性。

标签： c cuda

【解决方案1】：

查看 sobelFilter SDK 示例。

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

共享内存的微妙之处在于，如果您读取相邻的字节，您会遇到 4-way bank 冲突。解决此问题的一种方法（如 sobelFilter 示例所示）是将循环展开 4 倍并访问每四个字节。

【讨论】：