CUDA：每个线程计算的最佳像素数（灰度）

Question

我正在开发一个程序来转换灰度图像。我正在使用 CImg 库。我必须为每个像素读取 3 个值 R-G-B，计算相应的灰度值并将灰度像素存储在输出图像上。我正在使用 NVIDIA GTX 480。关于卡的一些细节：

• 微架构：费米
• 计算能力（版本）：2.0
• 每个 SM 的核心数（经线大小）：32
• 流式多处理器：15
• 每个多处理器的最大驻留扭曲数：48
• 每个多处理器的最大共享内存量：48KB
• 每个多处理器的最大驻留线程数：1536
• 每个多处理器的 32 位寄存器数：32K

我正在使用具有 256 个线程块的方形网格。 该程序可以有不同尺寸的输入图像（例如 512x512 像素、10000x10000 像素）。我观察到增加分配给每个线程的像素数会提高性能，因此它比每个线程计算一个像素要好。问题是，如何确定静态分配给每个线程的像素数？用所有可能的数字计算测试？我知道在 GTX 480 上，1536 是每个多处理器的最大驻留线程数。我需要考虑这个数字吗？以下，是内核执行的代码。

for(i = ((gridDim.x + blockIdx.x) * blockDim.x) + threadIdx.x; i < width * height; i += (gridDim.x * blockDim.x)) {
    float grayPix = 0.0f;
    float r = static_cast< float >(inputImage[i]);
    float g = static_cast< float >(inputImage[(width * height) + i]);
    float b = static_cast< float >(inputImage[(2 * width * height) + i]);

    grayPix = ((0.3f * r) + (0.59f * g) + (0.11f * b));
    grayPix = (grayPix * 0.6f) + 0.5f;
    darkGrayImage[i] = static_cast< unsigned char >(grayPix);
}

Answer 1

问题是，如何确定静态分配给每个线程的像素数？用所有可能的数字计算测试？

虽然您没有显示任何代码，但您提到了一个观察到的特征：

我观察到增加分配给每个线程的像素数会提高性能，

对于这些类型的工作负载，这实际上是一个相当普遍的观察结果，而且这种情况在 Fermi 上可能比在新架构上更明显。在矩阵转置期间会发生类似的观察。如果您编写一个“简单”的矩阵转置，每个线程转置一个元素，并将其与here 讨论的每个线程转置多个元素的矩阵转置进行比较，您会发现，特别是在 Fermi 上，每个线程的多个元素转置可以在设备上实现大约可用的内存带宽，而每线程一个元素的转置则不能。这最终与机器隐藏延迟的能力有关，以及您的代码公开足够工作以允许机器隐藏延迟的能力。理解底层行为有些复杂，但幸运的是，优化目标相当简单。

GPU 在等待之前发出的操作完成时，通过切换到大量可用工作来隐藏延迟。因此，如果我有很多内存流量，那么对内存的各个请求都会有很长的延迟。如果机器在等待内存流量返回数据时我还有其他工作可以做（即使该工作会产生更多的内存流量），那么机器可以使用该工作来保持自己忙碌并隐藏延迟。

让机器完成大量工作的方法首先是确保我们已启用最大数量的经线，以适应机器的瞬时容量。这个数字计算起来相当简单，它是 GPU 上的 SM 数量和每个 SM 上可以驻留的最大扭曲数的乘积。我们希望启动一个满足或超过这个数字的内核，但是超过这个数字的额外扭曲/块不一定帮助我们隐藏延迟。

一旦我们遇到了上述数字，我们希望将尽可能多的“工作”打包到每个线程中。实际上，对于您描述的问题和矩阵转置情况，将尽可能多的工作打包到每个线程中意味着每个线程处理多个元素。

所以步骤相当简单：

1. 启动机器可以立即处理的尽可能多的经纱
2. 如果可能，将所有剩余的工作放在线程代码中。

让我们举一个简单的例子。假设我的 GPU 有 2 个 SM，每个 SM 可以处理 4 个 warp（128 个线程）。请注意，这不是核心数，而是 deviceQuery 输出所指示的“每个多处理器的最大驻留扭曲数”。

然后我的目标是创建一个包含 8 个线程的网格，即总共 256 个线程（在至少 2 个线程块中，因此它们可以分配给 2 个 SM 中的每一个）并通过每个线程处理多个元素来使这些线程执行整个问题线。因此，如果我的整体问题空间总共有 1024x1024 个元素，那么理想情况下我希望每个线程处理 1024*1024/256 个元素。

请注意，此方法为我们提供了优化方向。我们不一定要完全实现这个目标才能使机器饱和。例如，可能只需要每个线程处理 8 个元素，以便让机器完全隐藏延迟，并且通常会出现另一个限制因素，如下所述。

采用这种方法将倾向于消除 延迟 作为内核性能的限制因素。使用分析器，您可以通过多种方式评估延迟在多大程度上是一个限制因素，但一个相当简单的方法是捕获sm_efficiencymetric，并可能在您概述的两种情况下比较该指标（每个线程一个元素，每个线程多个元素）。我怀疑您会发现，对于您的代码，sm_efficiency 指标表明每个线程情况下的多个元素的效率更高，这表明在这种情况下延迟不是限制因素。

一旦您将延迟作为限制因素排除，您将倾向于遇到其他两个机器性能限制因素之一：计算吞吐量和内存吞吐量（带宽）。在矩阵转置的情况下，一旦我们充分处理了延迟问题，内核就会倾向于以受内存带宽限制的速度运行。