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

首先我要说的是，我已经仔细阅读了所有类似的问题：

1. 确定每个块的线程和每个网格的块 https://stackoverflow.com/questions/4391162/cuda-determining-threads-per-block-blocks-per-grid
2. 每个 SM 的线程数、每个块的线程数 https://stackoverflow.com/questions/17816136/cuda-what-is-the-threads-per-multiprocessor-and-threads-per-block-distinction
3. CUDA 块和线程 https://stackoverflow.com/questions/9342599/cuda-blocks-and-threads
4. 扭曲和最佳块数 https://stackoverflow.com/questions/32855684/cuda-warps-and-optimal-number-of-threads-per-block

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

My data is not与我见过的大多数例子一样，它是一个方格（矩阵），它是两个向量生成一个矩阵，行与列不相等：

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 计算中有意义。 我想保持这一点的一致性，并且据我所知，使用 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 个项目的输入向量。

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

我知道每个块的最大线程数是 1024，但因为它是 2D 网格，所以更有可能是：

dim3 threadPerBlock(X,Y);

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

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

最后，由于我的经纱尺寸是 32，

4. 无论所有其他参数如何，是否最小网格尺寸需要至少为 32，或 32 的倍数？我需要at least每个块 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：我的问题不是关于优化代码，而是了解如何在设备上分发线程和网格数据。

对计算问题进行分类的一种方法是讨论转变 and 减少.

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

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

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

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

CUDA 线程块尺寸的选择通常有些随意。一般来说，我们通常希望线程块的目标是每块 128 - 512 个线程范围内（原因已在其他地方介绍），并且我们希望线程块是 32（扭曲大小）的整数倍，以提高线程块获取时的效率。细分为 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 内核代码的一种简单方法，但这不是唯一的方法。另一种可能的方法是使用某种形式的网格跨步循环，以便较小的网格可以覆盖较大的问题空间。对这些策略的讨论超出了本答案的范围，在解决其他方法之前应先了解本答案中讨论的基本方法。