使用共享内存计算 cuda 中像素的大小答案

【问题标题】：calculating the magnitude of a pixel in cuda with shared memory使用共享内存计算 cuda 中像素的大小
【发布时间】：2012-10-08 18:06:31
【问题描述】：

我正在做一个学校项目，我们在 cuda 中运行 sift 算法。我有一点要根据其邻居（A，B，C，D）的值来计算图像的每个像素（X）的幅度值：

   A
B  X  C
   D

我设法通过使用全局内存来实现它，因为我可以轻松地从输入数组中获取我想要的值。

但是现在我想通过首先将输入数组放入共享内存来实现它，但是我很难让线程将正确的像素放在共享内存上。我必须考虑图像边界上的填充。

我知道我需要比我想要放入其中的图像部分更多的共享内存，以便包含填充，但我不知道我的线程块是否应该包含比共享内存空间更多或更少的线程以及如何指定要阅读的内容。如果有人可以就如何思考这个问题给我一个大致的想法，我可以从那里得到它......

谢谢！

【问题讨论】：

对于图像处理应用程序，共享内存的使用有点复杂。我建议您使用 CUDA 纹理。纹理读取被缓存，处理边界情况非常容易。纹理最适合邻域操作。
您的图形设备具备哪种计算能力？

标签： cuda shared-memory sift

【解决方案1】：

我提供了一个通过灰度图像并应用 sobel 过滤器的代码：（Sobel 是一个类似于你的邻居（A，B，C，D）函数的过滤器）

#define QUANTUM_TYPE short
__global__ void sobel_gpu(QUANTUM_TYPE *img_out, QUANTUM_TYPE *img_in, int WIDTH, int HEIGHT){
    int x,y;
    x=blockDim.x*blockIdx.x+threadIdx.x;
    y=blockDim.y*blockIdx.y+threadIdx.y;
    QUANTUM_TYPE LUp,LCnt,LDw,RUp,RCnt,RDw;
    int pixel;

    if(x<WIDTH && y<HEIGHT){
        LUp = (x-1>=0 && y-1>=0)? img_in[(x-1)+(y-1)*WIDTH]:0;
        LCnt= (x-1>=0)? img_in[(x-1)+y*WIDTH]:0;
        LDw = (x-1>=0 && y+1<HEIGHT)? img_in[(x-1)+(y+1)*WIDTH]:0;
        RUp = (x+1<WIDTH && y-1>=0)? img_in[(x+1)+(y-1)*WIDTH]:0;
        RCnt= (x+1<WIDTH)? img_in[(x+1)+y*WIDTH]:0;
        RDw = (x+1<WIDTH && y+1<HEIGHT)? img_in[(x+1)+(y+1)*WIDTH]:0;
        pixel = -1*LUp  + 1*RUp +
                -2*LCnt + 2*RCnt +
                -1*LDw  + 1*RDw;
        pixel=(pixel<0)?0:pixel;
        pixel=(pixel>MAXRGB)?MAXRGB:pixel;
        img_out[x+y*WIDTH]=pixel;
    }
}

代码适用于全局内存并安全地处理边界。我的完整代码读取 BMP 图像并对其应用过滤器并将生成的 BMP 存储回磁盘。它在here 上提供（CPU 和 GPU 实现集成，适用于 Linux 和 Windows）。

您可以通过一些工作将其转换为共享内存样式。首先，您应该决定给每个块分配多少任务。然后将任务分解为多个共享内存接收器/排水管。 CUDA SDK 中的矩阵乘法示例为您提供了一个完美的想法。

【讨论】：