c++ - CUDA估算2D网格数据的每个 block 的线程数和 block 号

Question

首先，我已经仔细阅读了关于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

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

我的数据 不是像我所看到的大多数例子中那样经常是正方形格子(矩阵)，它是两个 vector 产生一个矩阵，行与列不相等:

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];
}

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

我必须将每个块的线程计算为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)

我试图了解如何推断/计算网格大小，每块线程数和块数。

让我们假设我的权重 vector 为800，输入 vector 为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:我的问题不是关于优化代码，而是了解如何在设备上分配线程和网格数据。

Answer 1

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

减少是问题的一种，它占用较大的输入数据集大小，而产生较小的输出数据集大小。例如，拍摄图像并找到最大像素值将是一种减少。对于此讨论，我们将忽略减少量。

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

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

按照上述线程策略，我们将需要网格中的线程总数等于我需要创建的输出点总数。对于2D问题，通常以二维方式考虑这些问题很方便，并且CUDA为此提供2D(或3D)线程块组织和2D(或3D)网格组织。

CUDA线程块尺寸的选择通常有些随意。一般而言，我们通常希望针对每个块范围内的128-512个线程中的线程块(出于其他方面介绍的原因)，并且我们希望线程块的整数倍是32(扭曲大小)的整数倍，以提高线程块的效率。 segmentation 为扭曲，这是CUDA执行的实际单位。在当前支持的GPU上，每个块的线程块限制为1024个线程(总计-即尺寸的乘积)。但是，对于许多问题，此范围内的线程块选择(例如256个线程与512个线程)通常对性能的影响相对较小。为了使某些功能正常运行，我们目前不详细介绍。 (当您回来进行优化时，可以重新选择该选项。)

到目前为止，我们已经了解到，对于这种问题类型，我们需要总数为64的线程来覆盖我们的问题空间，并且我们将有一些任意的线程块维选择。因此，让我们选择(32,16)(x，y)开始，总共有512个线程。没有规则说明adblocks必须是“square”，或者网格必须是“square”，或者甚至在线程块尺寸和问题尺寸(或网格尺寸)之间应该存在任何比例的奇偶校验。

现在考虑到线程块选择为(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内核代码的简便方法，但这不是唯一的方法。另一种可能的方法是使用某种形式的网格跨越循环，以便较小的网格可以覆盖较大的问题空间。这些策略的讨论不在此答案的范围内，并且在解决其他方法之前应了解此答案中讨论的基本方法。