CUDA 中 __syncthreads() 的替代方案答案

【问题标题】：Alternatives to __syncthreads() in CUDACUDA 中 __syncthreads() 的替代方案
【发布时间】：2019-01-13 20:01:13
【问题描述】：

我正在使用 CUDA 并行化 C 代码。我发现我们可以按照以下模式进行计算：

因此，我们在第一步只能计算一个标记为“1”的元素，只有在第一个元素的计算完成后，我们才能开始计算接下来的两个标记为“2”的对角线元素，因为我们有数据依赖性。其他元素以此类推...

我们解决这个问题的方法是为每一行分配一个线程，每个线程在每一步结束时执行__syncthreads()，以实现上述同步。

但是，__syncthreads() 需要很多时间。有没有其他解决方案可以解决这个问题。

编辑 1： 计算每个元素X的依赖模式如下：

这里，元素X 需要红色和绿色元素的值。它仅依赖于红色元素（在上一次迭代中计算）。

提前致谢。

【问题讨论】：

您能否更具体地了解依赖模式？ nth 步骤上的每个元素是否都取决于步骤 n-1 上的每个元素，还是存在某种局部性？
@SteelRaven 我已经编辑了我的问题。是的！有某种地方性。
@AnastasiyaAsadullayeva 此矩阵中每个元素 X 的计算遵循结构，如编辑 1 中所示。此计算在顺序版本中逐行发生。我认为上面的平行对角线结构最好地解决了这种计算模式。但是，我对__syncthreads() 的副作用并不满意。
这是一个已解决的问题（谷歌“A FAST ITERATIVE METHOD FOR EIKONAL EQUATIONS” by Whitaker and Jeong）

标签： cuda gpu

【解决方案1】：

您可以利用 warp 中的所有线程同时执行且不需要显式同步的事实。为此，让我们让每个 warp 评估一个独立的大块结果，并仅在大块处理之间使用同步。让我们加载大方块的边界数据并在独立的warp中处理它们。这里我尝试绘制新的图案，每个方块都围绕着单个warp的数据，颜色开关表示需要同步，L表示warp需要加载该元素，E表示warp将负责元素评估。当然，chunk 必须和 warp 一样大，而不是 image 那样小。

代码可能如下所示：

volatile shared sharedChunks[warpsInBlock][33][33];
int warpId = threadIdx.x / 32;
int inWarpId = threadIdx.x % 32;
while(not done){
  sharedChunks[warpId][0][inWarpId + 1] =
    data[mapToCorrectHorisontalLoadId(threadIdx, iteration)];
  sharedChunks[warpId][inWarpId + 1][0] =
    data[mapToCorrectVerticalLoadId(threadIdx, iteration)];
  // Filling upper left triangle of array
  for(int i = 0; i < 32; ++i){
    if(inWarpId <= i){
      sharedChunks[warpId][i - inWarpId + 1][inWarpId + 1] =
        sharedChunks[warpId][i - inWarpId][inWarpId + 1] +
        sharedChunks[warpId][i - inWarpId + 1][inWarpId];
    }
  }
  // Filling lower right triangle of array
  for(int i = 1; i < 32; ++i){
    if(inWarpId >= i)
      sharedChunks[warpId][i + inWarpId + 1][31 - inWarpId + 1] =
        sharedChunks[warpId][i + inWarpId][31 - inWarpId + 1] +
        sharedChunks[warpId][i + inWarpId + 1][31 - inWarpId];
  }
  for(int i = 0; i < 32; ++i){
    data[backwardMapping(threadIdx, iteration, i)] =
      sharedChunks[warpId][i + 1][inWarpId + 1];
  }
  __syncthreads();
}

在这里要全面评估 32 条对角线上的元素，您只需进行两次同步而不是 32 次。

但是这个解决方案有它自己的缺点：

您必须弄清楚从全局内存数据索引到曲内索引的映射。如果加载或存储的合并不够充分，您可能无法获得任何加速。
正如您所见，在块评估期间，一些线程在很长一段时间内都没有做任何事情。大约有一半的时间不会被使用。
高度依赖可用于块的共享内存。更多同步的成本可以降低共享内存使用量。例如，您可以使用大小为 (warpsInBlock*2)x17x17 的数组，这里每个 warp 将评估两个较小的块，共享内存使用量大约会低两倍，但同步数会高出两倍。

也许还有什么我忘记了。

您可能会尝试这样做，但真正的加速将在很大程度上取决于许多因素。

【讨论】：

【解决方案2】：

首先，很抱歉复活一个超级老线程。

您可以重新排序项目（通过巧妙的寻址，而不是在内存中）以获得如下访问模式：

... 等等。您只需要在每个这样的行之后使用__syncthreads()，这样会快得多。

【讨论】：