c++ - 用于将 RGB 图像转换为灰度图像的共享内存 Cuda

Question

我是 Cuda 编程的新手，我有一个将 RGB 图像转换为灰度图像的代码。读取像素的RGB值并将其转换为GreyScale的算法已经提供给我们。并行化代码使我的速度提高了大约 40-50 倍。我想进一步优化它以实现大约 100 倍的加速。为此，我想使用共享内存访问，因为它的数量级比全局内存访问快。我浏览了不同的在线资源，对共享内存访问有了基本的了解。但是在我的代码中，我无法理解如何实现共享内存，读取 RGB 值并转换为灰度的代码

    for ( int y = 0; y < height; y++ ) {
      for ( int x = 0; x < width; x++ ) {
        float grayPix = 0.0f;
        float r = static_cast< float >(inputImage[(y * width) + x]);
        float g = static_cast< float >(inputImage[(width * height) + (y * width) + x]);
        float b = static_cast< float >(inputImage[(2 * width * height) + (y * width) + x]);

        grayPix = ((0.3f * r) + (0.59f * g) + (0.11f * b));
        grayPix = (grayPix * 0.6f) + 0.5f;

        darkGrayImage[(y * width) + x] = static_cast< unsigned char >(grayPix);
        }
     }

输入图像一个 char*，我们使用 CImg 库来处理图像

CImg< unsigned char > inputImage = CImg< unsigned char >(argv[1]);

用户在运行代码时将图像路径作为参数传递

这是我的 Cuda 实现

unsigned int y = (blockIdx.x * blockDim.x) + threadIdx.x;
unsigned int x = (blockIdx.y * blockDim.y) + threadIdx.y;
float grayPix = 0.0f;
float r = static_cast< float >(inputImage[(y * height) + x]);
float g = static_cast< float >(inputImage[(width * height) + (y * height) + x]);
float b = static_cast< float >(inputImage[(2 * width * height) + (y * height) + x]);    
grayPix = ((0.3f * r) + (0.59f * g) + (0.11f * b));
grayPix = (grayPix * 0.6f) + 0.5f;

darkGrayImage[(y * height) + x] = static_cast< unsigned char >(grayPix);

Grid和 block 以及调用代码

    dim3 gridSize(width/16,height/16);
    dim3 blockSize(16,16);
    greyScale<<< gridSize, blockSize >>>(width,height,d_in, d_out);

其中 width 和 height 是输入图像的宽度和高度。我尝试使用 (32,32) 的 block 大小，但它减慢了代码速度而不是加快速度

现在我想添加共享内存，但问题是对输入变量 InputImage 的访问是非线性的，所以我应该向共享内存添加什么值我试过类似的东西

 unsigned int y = (blockIdx.x * blockDim.x) + threadIdx.x;
 unsigned int x = (blockIdx.y * blockDim.y) + threadIdx.y;
 extern __shared__ int s[];
 s[x]=inputImage[x];
 __syncthreads();

然后在实现中用 s 替换 inputImage 但这只是给出了错误的输出(全黑图像)你能帮我在这里了解我如何实现共享内存，如果它可能和有用的话，有没有办法让我以更合并的方式进行访问？

如有帮助将不胜感激

Answer 1

由于以下几个原因，这无法工作:

 unsigned int x = (blockIdx.y * blockDim.y) + threadIdx.y;
 extern __shared__ int s[];
 s[x]=inputImage[x];

一个原因是我们不能使用全局索引 (x) 作为共享内存索引，除非数据集足够小以适合共享内存。对于尺寸相当大的图像，您无法将整个图像放入共享内存的单个实例中。此外，您仅使用二维数据集的一维索引 (x)，因此这可能没有意义。

这表明人们普遍缺乏对如何在程序中使用共享内存的理解。然而，我们并没有试图解决这个问题，而是观察到对于正确编写的 RGB-> 灰度代码，共享内存的使用不太可能提供任何好处。

当存在数据重用时，共享内存带宽优势(这就是您所说的“速度更快”时指的是什么)很有值(value)。 RGB-> 灰度代码不应要求任何数据重用。您从全局内存中加载每个 R、G、B 数量恰好一次，并将计算出的灰度数量恰好存储在全局内存中一次。将其中一些数据临时移动到共享内存不会加快任何速度。您仍然需要执行全局加载和全局存储，对于正确编写的代码，这应该是所有必要的。

但是在您的问题中，您已经提出了一个可能的改进路径:合并访问。如果您要分析发布的代码，您会发现完全未合并的访问模式。为了实现良好的合并，我们希望复合索引计算具有 threadIdx.x 变量不乘以任何值的属性:

unsigned int y = (blockIdx.x * blockDim.x) + threadIdx.x;
unsigned int x = (blockIdx.y * blockDim.y) + threadIdx.y;
float grayPix = 0.0f;
float r = static_cast< float >(inputImage[(y * height) + x]);
                                           ^
                                           |
                                           y depends on threadIdx.x

但在您的情况下，您的索引计算是将 threadIdx.x 乘以 height。这将导致非合并访问。 warp 中的相邻线程将具有不同的 threadIdx.x，我们希望 warp 中相邻线程的索引计算结果是内存中的相邻位置，以实现良好的合并访问。如果将 threadIdx.x 乘以任何值，则无法实现此目的。

这个问题的解决方案非常简单。您应该只使用几乎与您显示的非 CUDA 代码完全相同的内核代码，并为 x 和 y 定义适当的代码:

    unsigned int x = (blockIdx.x * blockDim.x) + threadIdx.x;
    unsigned int y = (blockIdx.y * blockDim.y) + threadIdx.y;
    if ((x < width) && (y < height)){ 
      float grayPix = 0.0f;
      float r = static_cast< float >(inputImage[(y * width) + x]);
      float g = static_cast< float >(inputImage[(width * height) + (y * width) + x]);
      float b = static_cast< float >(inputImage[(2 * width * height) + (y * width) + x]);

      grayPix = ((0.3f * r) + (0.59f * g) + (0.11f * b));
      grayPix = (grayPix * 0.6f) + 0.5f;

      darkGrayImage[(y * width) + x] = static_cast< unsigned char >(grayPix);
      }

当然，这不是完整的代码。你没有显示完整的代码，所以如果你回答“我试过了，但它不起作用”，我不太可能帮到你很多，因为我不知道你实际运行的是什么代码.但是:

1. 共享内存不是该算法的正确方法
2. 由于我指出的原因，您发布的代码中无疑存在合并问题
3. 合并修复应遵循我概述的路径
4. 您的性能应该会随着合并修复而提高。

请注意，“它不起作用”的响应意味着您实际上是在寻求调试帮助，而不是概念性解释，在这种情况下，您是 supposed to提供一个MCVE .您展示的不是 MCVE。最好您的 MCVE 不应依赖于 CImg 之类的外部库，这意味着您需要努力创建一个独立测试，但要证明您遇到的问题。

此外，我建议您在使用 CUDA 代码时遇到问题时，使用 proper CUDA error checking以及使用 cuda-memcheck 运行您的代码。

(适当的 CUDA 错误检查会发现您尝试使用共享内存时出现问题，例如，由于共享内存中的越界索引。)