用于重叠数据的内核设计，启动单独的 warp

Posted 2023-03-23

【中文标题】用于重叠数据的内核设计，启动单独的 warp

【英文标题】：Kernel design for overlapping data, launch of a seperate warp

【发布时间】：2014-01-13 09:42:57

【问题描述】：

我有一个关于我试图根据我在网上找到的paper 实施的 CFD 应用程序的问题。这可能是一个初学者的问题，但它是这样的：

情况如下： 2D 域被分解为瓦片。这些瓦片中的每一个都由相关内核块处理。正在执行的计算非常适合并行执行，因为它们只考虑了少数邻居（这是一个浅水应用程序）。瓷砖确实重叠。每个图块在它应该计算结果的域的每一侧都有 2 个额外的单元格。

您在左侧看到 1 个区块，在右侧看到 4 个区块，并带有随之而来的重叠。灰色是计算所需的“幽灵细胞”。浅绿色是每个块实际写回全局内存的域。不用说整个域将有超过 4 个图块。

每个线程的想法如下：

(1) copy data from global memory to shared memory
__synchthreads();
(2) perform some calculations
__synchthreads();
(3) perform some more calculations
(4) write back to globabl memory

对于绿色区域的单元格，内核是直截了当的，您根据您的 threadId 复制数据，并使用共享内存中的邻居进行计算。然而，由于数据依赖的性质，这还不够：

(1) 必须在所有单元格（灰色和绿色）上运行。没有依赖。

(2) 必须在所有绿色单元格以及灰色单元格的内部行/列上运行。取决于相邻数据 N、S、E 和 W。

(3) 必须在所有绿色单元格上运行。取决于步骤 (2) 中关于邻居 N、S、E 和 W 的数据。

所以我的问题是： 如果没有非常混乱的代码，如何做到这一点？

我能想到的只是大量的“if”语句来决定一个线程是否应该执行其中一些步骤两次，具体取决于他的 threadId。

我也考虑过使用重叠块（而不仅仅是重叠数据），但这会导致另一个问题：__synchthreads() 调用必须在代码的条件部分中。 将内核拆开并让步骤 (2)/(3) 在不同的内核中运行也不是一个真正的选择，因为它们会产生中间结果，由于它们的数量/大小而不能全部写回内存。

作者自己写道（Brodtkorb et al. 2010, Efficient Shallow Water Simulations on GPUs： 实施、可视化、验证和确认）：

启动内核时，我们首先从全局内存读取到片上共享内存。除了我们块的内部单元格之外，我们还需要使用来自每个方向的两个相邻单元格的数据来完成数据 模具的依赖关系。将数据读入共享内存后，我们继续计算一维 分别在 x 和 y 方向上的通量。使用图 1 中所示的步骤，通过存储计算通量 共享内存中多个线程使用的所有值。我们还共同执行计算 一个块以避免重复计算。然而，因为我们计算每个单元的净贡献，我们有 执行比线程数更多的重建和通量计算，使我们的内核复杂化。 这是 在我们的代码中通过指定一个执行额外计算的warp来解决；一种产生了 比在多个 warp 之间划分额外的计算具有更好的性能。

那么，他指定一个经线来进行这些计算是什么意思？如何做到这一点？

【参考方案1】：

那么，他指定一个经线来进行这些计算是什么意思？如何做到这一点？

你可以这样做：

  // work that is done by all threads in a block
  __syncthreads();  // may or may not be needed
  if (threadIdx.x < 32) 
    // work that is done only by the designated single warp
    

尽管这非常简单，但就您问题中的最后一个问题以及突出显示的段落而言，我认为这很可能是他们所指的。我认为它符合我在这里阅读的内容。此外，我不知道有任何其他方法可以将工作限制为单个扭曲，除非使用条件。他们可能还选择了单个 warp 来利用 warp 同步行为，这绕过了您前面提到的条件代码问题中的__syncthreads();。

所以我的问题是：如何在没有非常混乱的代码的情况下做到这一点？

我能想到的只是大量的“if”语句来决定一个线程是否应该执行其中一些步骤两次，具体取决于他的 threadId。

实际上，我认为任何普通的“if”语句序列，无论多么混乱，都无法解决您描述的问题。

您已经提到的解决第 2 步和第 3 步之间依赖关系的典型方法是将工作分成两个（或更多）内核。您表示这“不是一个真正的选择”，但据我所知，您正在寻找的是全局同步。除了内核启动/退出点之外，这样的概念在 CUDA 中没有明确定义。 CUDA 不保证网格中块之间的执行顺序。如果您在第 3 步中的块计算依赖于第 2 步中的 相邻块 计算，那么在我看来，您肯定需要全局同步，如果您不实现它，您的代码将会变得丑陋随着内核启动。在我看来，使用全局信号量或全局块计数器等替代方法是脆弱的，并且难以应用于广泛数据依赖的一般情况（其中每个块都依赖于上一步的邻居计算）。

如果相邻计算仅依赖于一组较薄的相邻单元（“晕”）中的数据，而不是整个相邻块，并且这些单元可以独立计算，那么您的块可能是扩展为包括相邻单元格（即重叠），有效地计算相邻块之间的光环区域两次，但您已经表明您已经考虑并放弃了这个想法。但是，我个人希望在接受完全基于 __syncthreads(); 的难度而拒绝该代码的想法之前详细考虑代码 根据我的经验，由于条件代码执行避风港而说他们不能使用 __syncthreads(); 的人即使在条件代码中，也没有在详细代码级别准确考虑所有选项以使__syncthreads(); 工作。

【讨论】：

感谢您花时间回答。关于全局同步：由于我的解释，这可能是一个误解。这些块不需要在相关步骤中进行通信。他们使用重叠（光环）来完全避免彼此交谈。所以这不是问题。我没有像你说的那样放弃这个，这实际上正是应该发生的事情。但是我的意思是通过 192 个线程（绿色单元的数量）来做到这一点。要处理的灰色单元格的数量是（取决于是第 1 步还是第 2 步）128 或 48。所以我需要将 threadIds 巧妙地映射到 memoryIds。但是怎么做？

还有一件事，要明确一点：我确实丢弃的（至少现在）是为每个单元格（包括光环）使用一个线程的选项，因此使用 320 个线程。那么这些块将重叠（不仅仅是内存）。您可能是对的，有一种方法可以通过改进算法将__synchthreads(); 放在这种情况下的条件语句中。但是我还没有看到。

我重新考虑并重写了内核以使用更多线程，以不同的方式处理__synchthreads();。谢谢你的回答。