-2

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

`__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 的内核，有多少扭曲将具有控制散度？

答案：500

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

问题：我试图理解为什么在这种情况下，x 维度中的 8 个线程将在无效范围内？

4

1 回答 1

3

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

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

当我在水平方向上启动 63 个块时，我在水平方向上得到 63x16 = 1008 个线程。但我只需要 1000 个，所以 8 个线程（每行）在“无效范围”内。

于 2017-06-02T20:39:10.330 回答