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

Posted 2023-04-15

【中文标题】CUDA：每个线程计算的最佳像素数（灰度）

【英文标题】：CUDA: Best number of pixel computed per thread (grayscale)

【发布时间】：2016-03-22 13:10:52

【问题描述】：

我正在开发一个程序来转换灰度图像。我正在使用 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);

【问题讨论】：

性能取决于许多因素：寄存器使用、内存合并，当然还有块和网格大小。您可以通过在“NVIDIA CUDA 占用计算器”中输入您的数字来获得一些信息 - 这是一个 XLS (Excel) 文件，可在 developer.download.nvidia.com/compute/cuda/…

找到

@Marco：甚至不再需要使用占用电子表格。运行时 API 功能 cudaOccupancyMaxActiveBlocksPerMultiprocessor 将为您完成所有艰苦的工作

@talonmies 对，那是在... CUDA 6 左右添加的？但是，我认为电子表格仍然比编写自定义代码、查阅 API 文档和在一些修改-编译-运行-重复循环中尝试不同的参数化更方便。我只是想提一下（虽然我自己并没有广泛使用它，也不能说它对有针对性的优化有多大帮助）

@Marco：我已经尝试过“NVIDIA CUDA 占用计算器”。但我唯一能看到的是变化 1）每个线程的寄存器计数，2）每个块的共享内存和 3）块大小的影响。没有关于用于获得性能的最佳块数量的建议。我错了吗？

我对它也不是很熟悉，所以我认为@talonmies 可能会在这里给出更集中的建议。但是例如，当您使用默认设置打开 XLS，然后将“每个块的线程数”更改为 128，您会看到占用率（在上图中，在“每个多处理器的占用率”字段中）减少。最大值似乎已达到，例如用于 512 或 672 个线程（100% 或 98% 占用率）

【参考方案1】：

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

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

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

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

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

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

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

所以步骤相当简单：

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

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

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

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

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

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

【讨论】：

感谢您快速详细的回复！因此，为了检查我是否正确理解，使用 10200x6500 像素（66300000 像素）的图像、15 SM 和 48 每个多处理器的最大驻留扭曲数。我当时有 15x48x32= 23040 个线程可以驻留。 66300000/23040=2877.60，这意味着每个线程 2878 像素。对吗？

是的，但是您从这种情况中消除延迟的点很可能远小于每个线程 2878 像素。此外，在您的实际示例中，您可能希望选择每个块 512 个线程的线程块大小，目标是每个 SM 驻留 3 个线程块，以实现 1536 个线程（48 个线程）驻留目标。所以你会启动一个包含 45 个线程块的内核，每个线程块有 512 个线程。

最后，由于您没有显示任何代码，请注意，以这种方式在线程级别公开工作取决于编译器能够移动独立操作。

我觉得这个答案写得特别好！

我尝试了你的方法，使用 'dim3 gridSize(45);昏暗3块大小（512）；' .但是使用大小为的方形网格：'unsigned int grid_size = static_cast(ceil(sqrt(ceil(width * height / PX_TH) / (float)256)))' 其中 PX_TH 是要分配给的像素数每个线程和 256 个块，我获得了更多的性能。在内核中，我读取每个像素的 R-G-B 值，计算灰度值并存储新像素。最后，我增加要计算的像素的索引。例如。 65536000 px，第一种方法需要 0.004402s，第二种方法每个线程 11 px 需要 0.003255s。有什么想法吗？