主页 / coding / 问题 / 76934852
Accepted
user366312
Asked: 2023-08-19 19:26:31 +0800 CST

为什么他们只填充一个共享内存？

  • 772

上面的链接说：

__global__ void coalescedMultiply(float *a, float *c, int M)
{
  __shared__ float aTile[TILE_DIM][TILE_DIM],
                   transposedTile[TILE_DIM][TILE_DIM];
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  float sum = 0.0f;
  aTile[threadIdx.y][threadIdx.x] = > a[row*TILE_DIM+threadIdx.x];
  transposedTile[threadIdx.x][threadIdx.y] =
      a[(blockIdx.x*blockDim.x + threadIdx.y)*TILE_DIM +
      threadIdx.x];  
  __syncthreads();
  for (int i = 0; i < TILE_DIM; i++) {
      sum += aTile[threadIdx.y][i]* transposedTile[i][threadIdx.x];
  }
  c[row*M+col] = sum;
}

(...)
这些多方银行冲突的代价非常高昂。简单的补救措施是填充共享内存阵列，使其具有额外的列，如以下代码行所示。

__shared__ float transposedTile[TILE_DIM][TILE_DIM+1];

我的问题是，

为什么它们只填充右侧 ( transposedTile) 共享内存，而不是同时填充 (transposedTile和aTile)？

cuda

1 个回答

  1. Best Answer

    因为根据代码中的访问模式，只有对该图块的访问才会出现库冲突。因此，只需填充该图块即可避免银行冲突。

    简而言之，当 warp 中的 2 个或更多线程尝试在同一事务中访问同一存储体中的值时（基本上就像在同一指令问题中所说的那样），共享访问中的存储体冲突就会发生。为简单起见，当我们有一个 32x32 线程块数组和 4 字节数量的 32x32 平铺排列时，库和列是同义词。因此，前面的语句可以修改为如果两个或多个线程在同一个周期内访问tile中的同一列，就会出现bank冲突。

    此规则的例外是广播例外。如果两个或多个线程由于访问同一位置而访问同一存储体（或本例中的列），则该位置将被广播到该周期的所有必要线程，并且广播规则并不意味着存储体冲突。

    因此，这意味着我们必须查看每个图块的所有访问模式（加载和存储），并查看原始代码中的哪些模式会发生银行冲突。我们的规则是，如果线程束中的线程之间的列相同，但行不同，则会发生组冲突。我们还应该能够观察或快速实现 32x32 线程块的一般“布局”：该线程块中有 32 个扭曲。在第一个 warp 中，所有线程都等于threadIdx.y0，并且threadIdx.x在整个 warp 上从 0 到 31 变化。第二个扭曲的唯一区别是其threadIdx.y值全部为 1。第三个扭曲的值为 2，依此类推。

    现在让我们看看每个访问（每个图块一次加载和一次存储），并查看哪些访问在原始代码中具有银行冲突模式：

    aTile[threadIdx.y][threadIdx.x] = ...

    这是 的商店操作aTile。 threadIdx.y位于行索引中，threadIdx.x位于列索引中。我们知道，对于 32x32 线程块中的给定扭曲，threadIdx.y在扭曲上是恒定的，并且threadIdx.x在扭曲上以线性方式变化：0,1,2,3,...31。因此，这个共享加载事务将有线程在同一行共享内存中请求数据，并且不会发生存储体冲突。每个线程都从一个单独的列读取，这就像在这种情况下说一个单独的存储体，这是共享内存的典型最佳模式：没有存储体冲突。

    transposedTile[threadIdx.x][threadIdx.y] = ...

    这是 的共享商店交易transposedTile。我们看到，与之前的情况相比，索引是相反的。 threadIdx.y，它在整个扭曲上是恒定的，现在位于列索引中，并且在整个threadIdx.x扭曲上变化，位于行索引中。这意味着我们在经线中选择单个列，不同的线程访问该列的不同成员。这是列式访问（而不是前一种情况中的行式访问），并且当我们进行平铺排列使得每列对应于共享内存组时，这通常是共享内存中最糟糕的可能访问模式。所有线程都访问同一个存储体，并且在扭曲范围内，我们将出现 32 路存储体冲突，这意味着访问时间将比最佳情况长 32 倍。

    ... * transposedTile[i][threadIdx.x];

    这是 的从共享加载操作transposedTile。i在整个扭曲上是恒定的（在任何给定的循环迭代中），并且threadIdx.x在整个扭曲上变化。因此，我们选择特定的行，并询问该行中的元素。这是共享访问的最佳情况。

    sum += aTile[threadIdx.y][i] ...

    考虑到整个扭曲，在任何给定的循环迭代中，行和列都是固定/恒定的。这是广播案例。在任何给定的循环迭代和指令发出周期中，warp 中的每个线程都要求共享内存中的相同位置。这也是典型的好。

    因此，4 种可能的访问模式中只有 1 种可能会发生存储体冲突。此访问模式属于transposedTile，因此只transposedTile需要填充，以避免此代码中的库冲突。

    • 1

相关问题