用简单的矩阵乘法核控制散度

Question

给定以下简单矩阵乘法内核

`__global__ void MatrixMulKernel(float* M, float* N, float* P, int 
Width)
{
  int Row = blockIdx.y*blockDim.y+threadIdx.y;
  int Col = blockIdx.x*blockDim.x+threadIdx.x;
  if ((Row < Width) && (Col < Width)) {
      float Pvalue = 0;
      for (int k = 0; k < Width; ++k) 
         {
            Pvalue += M[Row*Width+k]*N[k*Width+Col];
         }
  P[Row*Width+Col] = Pvalue;
   }
 }`

如果我们在 1000X1000 的矩阵上启动块大小为 16X16 的内核，有多少个 warp 会有控制发散？

答案：500

解释：水平方向有63个方块。 每行中x维度的8个线程将在无效范围内。每两行形成一个warp。因此，有 1000/2=500 个扭曲将在水平方向上跨越有效和无效范围。至于底部块中的扭曲，有效范围内有 8 个扭曲，无效范围内有 8 个扭曲。这些 warp 中的线程要么完全在有效范围内，要么完全在无效范围内。

问题：我想了解为什么在这种情况下 x 维度中的 8 个线程会在无效范围内？

Answer 1

每个块包含一个 16x16 的元素数组。要覆盖 1000x1000 元素的矩阵，我需要一个方形线程块数组，其尺寸为 1000/16 = 62.5 块在水平方向和 62.5 块在垂直方向。

但我无法启动 62.5x62.5 块，因此为了获得完整覆盖，我必须启动 63x63 块，承认这将在 "invalid range" 中创建额外的线程（即映射到1000x1000 矩阵之外的元素位置）。

当我在水平方向上启动63个块时，我在水平方向上得到63x16 = 1008个线程。但我只需要 1000 个，所以 "invalid range".

中有 8 个线程（每行）

用简单的矩阵乘法核控制散度

Control Divergence with simple matrix multiplication kernel

parallel-processing

cuda

computer-science

gpu

gpu-warp