CUDA 估计每个块的线程数和 2D 网格数据的块数

Posted 2023-04-15

【中文标题】CUDA 估计每个块的线程数和 2D 网格数据的块数

【英文标题】：CUDA estimating threads per blocks and block numbers for 2D grid data

【发布时间】：2016-01-19 15:34:27

【问题描述】：

首先让我说我已经仔细阅读了关于 SO 的所有类似问题：

Determining threads per block and block per grid Threads per SM, threads per block CUDA Blocks and Threads Warps and optimal number of blocks

我的目的是为我正在开发的前馈神经网络库尝试动态计算（而不是硬编码值）。

我的数据不是像我见过的大多数示例中经常出现的方格（矩阵），而是两个向量产生一个矩阵，行与列不相等：

float x[6] 1.f, 1.f, 0.f, 1.f, 1.f, 0.f; 
thrust::device_vector<float> in_vec( x, x+6 );
float y[9] 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f, 1.f;
thrust::device_vector<float> w_vec( y, y+9 );
thrust::device_vector<float> o_wec(9);
thrust::device_vector<float> mtx_vec( 9 * 6 );

float * i_ptr = thrust::raw_pointer_cast( in_vec.data() );
float * w_ptr = thrust::raw_pointer_cast( w_vec.data() );
float * out_ptr = thrust::raw_pointer_cast( mtx_vec.data() );

dim3 threadsPerBlock(9,6);
dim3 numBlocks(1,1);
prop_mtx<<<numBlocks,threadsPerBlock>>>( w_ptr, i_ptr, out_ptr, 6 );

和内核：

__global__ void prop_mtx( float * w, float * i, float * o, int s ) 

    int x = blockIdx.x * blockDim.x + threadIdx.x; 
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    o[y + x * s] = w[x] * i[y];

我之所以采用这种方法，是因为它在 ANN 计算中有意义，当涉及到向量/矩阵计算时。 我想保持一致，AFAIK 使用 2D 网格进行权重 * 输入计算是合理的。

我必须将每个块的线程数计算为 2D，并且网格中的线程数不相等。

我正在使用 GTX 660，它有：

  CUDA Capability Major/Minor version number:    3.0
  Total amount of global memory:                 2047 MBytes 
  ( 5) Multiprocessors, (192) CUDA Cores/MP:     960 CUDA Cores
  Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  2048
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)

我正在尝试了解如何推断/计算网格大小、每个块的线程数和块数。

假设我有一个包含 800 个项目的权重向量和一个包含 6500 个项目的输入向量。

这是否意味着我真正需要的是 800,6500 的 2D 网格？据我了解，还有什么会提供不正确的结果？

我知道我每个块的最大线程数是 1024，但由于它是一个 2D 网格，它更有可能是：

dim3 threadPerBlock(X,Y);

由于我的网格不是方阵，我需要以不同的方式计算每个块的 X、Y 线程吗？

或者我需要先推断出需要的块数？

最后，由于我的线程经纱大小是 32，

无论其他所有参数如何，最小网格尺寸是否需要至少为 32 或 32 的倍数？我需要至少每个块有 32 个线程，还是需要最小数量为 32 的网格大小？

任何伪代码，或我应该如何做的解释，将不胜感激。

我尝试的是通过将我的数据除以 32 环绕大小来计算我的 2D 网格大小。 然后我考虑使用可用的 SM 计算网格线程。例如

800 weights / 5 SM, = 160 x's per SM
6500 inputs  / 5 SM, = 1300 y's per SM

但我不知道从那里开始做什么。 最后，我考虑先找到输入权重比：

6500/800 = 8.125

暗示对 X 使用 32 的最小网格大小， Y 必须乘以 8.125 * 32 因此，我的 threadsPerBlock 将是：

dim3 threadsPerBlock(32,260);

那当然是每块8320个线程，远远超过每块1024个。

所以这是我的问题：如何不超过每个块的 1024 个线程，同时保持数据的正确网格大小？

PS：我的问题不是优化代码，而是了解如何在设备上分配线程和网格数据。

【问题讨论】：

***.com/questions/9985912/…

@talonmies 虽然您的回复非常有帮助，但它并没有回答我所有的问题：我如何推断线程数（总数？）以便网格与数据对齐，或者我不需要对齐它吗？答案之一如下：gridSize = (N + blockSize - 1) / blockSize;我是否只需要计算每个块的线程数而不考虑网格 X、Y？

是的，选择任意块大小，例如 32x32。然后将您的总 x-grid-width (800) 除以 x-block-dimension (32) 并在网格的 x 方向上启动那么多块（加一个）。然后将您的总 y-grid-width (6500) 除以 y-block-dimension (32) 并在网格的 y 方向上启动那么多块（加一个）。您的 GPU 中的 SM 数量不考虑在内。我会假设就像在你的小例子中你需要 9x6 线程一样，在更大的情况下你总共需要 800x6500 线程。这种方法在很多地方都有介绍。

您的内核还需要一个“线程检查”来防止所需的 800x6500 区域之外的线程执行任何操作。没有“32 最小网格大小”。建议将您的 threadblock size 设为 32 的整数倍。32x32 满足这一要求。

“线程检查”就是我所说的。如果您想查看其他示例，请尝试在本页右上角的搜索框中输入“user:1695960 线程检查”。

【参考方案1】：

对计算问题进行分类的一种方法是讨论转换和归约。

归约是一类问题，它采用较大的输入数据集大小，并产生较小的输出数据集大小。例如，拍摄图像并找到最大像素值将是一种缩减。对于本次讨论，我们将忽略归约。

转换是一种计算类别，其中输出数据集大小（元素数量）与输入数据集大小“大”或“大致相同”。例如，拍摄图像并生成模糊图像将是一种转换。

对于转换，编写 cuda 内核（线程代码）的一种常用方法（“线程策略”）是让一个唯一线程负责输出数组中的每个点。因此，我必须拥有的最小线程总数等于我的输出数组的大小。线程代码只是输入数据所需的一组计算，以产生一个输出数据点。粗略地说，您的问题和简化的内核符合这个定义；这是一个转变。

按照上面的线程策略，我们需要网格中的线程总数等于我需要创建的输出点的总数。对于 2D 问题，通常可以方便地考虑这些二维问题，而 CUDA 为此提供了 2D（或 3D）线程块组织和 2D（或 3D）网格组织。

CUDA 线程块尺寸的选择通常有些随意。一般来说，我们通常希望每个块范围内的线程块在 128 - 512 个线程范围内（原因将在其他地方介绍），并且我们希望线程块是 32 的整数倍（warp 大小）以提高线程块获取时的效率细分为warp，它们是CUDA执行的实际单元。在当前支持的 GPU 上，线程块被限制为每个块 1024 个线程（总计 - 即维度的乘积）。然而，对于许多问题，这个范围内的线程块选择（例如 256 线程与 512 线程）通常对性能的影响相对较小。为了让某些东西发挥作用，我们此时不会过多关注细节。 （当你回来进行优化时，你可以重新考虑这个选择。）

到目前为止，我们已经了解到，对于这种问题类型，我们需要总线程数来覆盖我们的问题空间，并且我们将有一个有点随意的线程块维度选择。所以让我们选择 (32,16) (x,y) 开始，总共 512 个线程。没有规则规定 adblocks 必须是“方形”，或者网格需要是“方形”，或者线程块尺寸和问题大小（或网格尺寸）之间甚至应该有任何类型的比率奇偶校验。

现在我们已经考虑到 (32,16) 的线程块选择，我们必须问自己“我需要多少这些？”。这个问题是 2D 的，因此我们选择了 2D 线程块以简化线程代码中的索引生成。让我们也选择一个 2D 网格 - 这对于 2D 问题是有意义的，同样对于索引生成的 2D 简单性。所以我们可以独立考虑这两个维度。

那么，我在 x 方向上需要多少块？我至少需要（我在 x 中的问题大小）/（我在 x 中的线程块大小）。由于我们在这里处理所有整数，这就引出了一个问题“如果我的问题大小不能被我的线程块大小整除怎么办？”规范的解决方案是启动足够多的线程来覆盖空间，或者足够的块来覆盖空间。但在非均分情况下，这将导致“额外线程”。我们将很快讨论和处理这些问题。因此，如果我有一个像这样用于线程块尺寸的 dim3 变量：

    #define BX 32
    #define BY 16   
    ...
    dim3 block(BX,BY);

那么我可以像这样构造我的 dim3 网格变量：

    #define DX 800
    #define DY 6500
    ...
    dim3 grid((DX+block.x-1)/block.x, (DY+block.y-1)/block.y);

如果你通过这个算法，你会发现这会导致我们在 x 和 y 方向启动 足够的块，这样我们将至少有足够的线程来覆盖我们的问题空间(DX,DY)，每个输出点一个线程。

希望 Y 维度与 x 维度分开处理是很清楚的。

上述计算通常会导致在我的网格中生成“太多”线程。在我需要处理的问题空间（DX、DY）的末尾之外，我将有一些“额外线程”。我们希望这些线程“什么都不做”。处理这个问题的规范方法是将问题空间维度传递给我的内核，在我的内核中创建一个适当的全局唯一线程索引，然后将该索引与我的问题空间中的最大索引进行比较。如果超过它，我们只需让该线程跳过所有剩余的线程代码。

以您的内核为例，它可能如下所示：

__global__ void prop_mtx( float * w, float * i, float * o, int s, const size_t d_size_x, const size_t d_size_y ) 

    int x = blockIdx.x * blockDim.x + threadIdx.x; 
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if ((x < d_size_x) && (y < d_size_y))  // thread check
      o[y + x * s] = w[x] * i[y];

请注意，这样的线程检查将创建“不参与”后续代码的线程（在某些块中）。这里要注意的一点是__syncthreads() 的使用取决于参与的块中的所有线程。因此，在这种情况下，我们不应该直接使用__syncthreads()。相反，我们必须适当地调节线程块行为：

__global__ void prop_mtx( float * w, float * i, float * o, int s, const size_t d_size_x, const size_t d_size_y ) 

    int x = blockIdx.x * blockDim.x + threadIdx.x; 
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if ((x < d_size_x) && (y < d_size_y))  // thread check
      
         o[y + x * s] = w[x] * i[y];
         // and other code not dependent on __syncthreads()
       
     // now it is safe to use since all threads are participating
     __syncthreads();
    if ((x < d_size_x) && (y < d_size_y))  // thread check
      
          // rest of kernel code
       

请注意，可以让较少数量的线程为较大数量的输出数据点执行必要的计算。线程和输出数据之间的 1:1 对应是一种简单的思考和编写 cuda 内核代码的方法，但这不是唯一的方法。另一种可能的方法是使用某种形式的网格跨步循环，以便较小的网格可以覆盖较大的问题空间。对这些策略的讨论超出了本答案的范围，在处理其他方法之前，应了解本答案中讨论的基本方法。

【讨论】：

感谢您精彩而详细的回答！我很困惑，因为我假设块网格线程应该以某种方式与网格数据大小相关。我猜不是，尽管仍然不明白 (32,16) 或 (32,32) 甚至 (32,8) 块网格将如何影响性能。 PS：如果我使用threadId = x + (gridDim.x * gridDim.y * y);其中 x = blockIdx.x * blockDim.x + threadIdx.x；和 y 同样，它与您发布的代码相同吗？无论如何，你回答了我所有的问题，所以再次感谢！

优化线程块大小以提高性能确实与我的回答不同。正如我所指出的，有一系列线程块选择（通常）可以提供大致相同的性能，如果不研究具体情况，很难确定哪个是最好的。

当您问“它与您发布的代码是否相同”时，我很高兴，因为我发布的代码是您的代码。代码“我已发布”中显示的计算创建了一个 x 和 y 线程索引，即每个线程的 2D 索引。您现在建议的为每个线程创建一个数字索引。它们将用于不同的索引目的。从数字上讲，您应该能够确定哪个是合适的。询问一维索引是否与二维索引相同对我来说没有意义。

你说得对，这不是我在网上找到的，然后愚蠢地复制的。非常感谢您的帮助！