如何为邻居访问优化 OpenCL 代码？

Posted 2023-04-15

【中文标题】如何为邻居访问优化 OpenCL 代码？

【英文标题】：How to optimize OpenCL code for neighbors accessing?

【发布时间】：2013-06-29 12:34:14

【问题描述】：

编辑：建议的解决方案结果添加在问题的末尾。

我开始使用 OpenCL 进行编程，并且我已经为我的问题创建了一个幼稚的实现。

理论是：我有一个元素的 3D 网格，其中每个元素都有一堆信息（大约 200 个字节）。每一步，每个元素都会访问它的邻居信息并积累这些信息以准备更新自己。之后有一个步骤，每个元素使用之前收集的信息进行自我更新。这个过程是迭代执行的。

我的 OpenCL 实现是：我创建一个 1 维 OpenCL 缓冲区，用代表元素的结构填充它，这些元素有一个“int neighbors 6 ”，我将邻居的索引存储在缓冲区中。我启动了一个内核，它咨询邻居并将它们的信息累积到此步骤中未咨询的元素变量中，然后我启动另一个使用此变量更新元素的内核。这些内核仅使用 __global 变量。

示例代码：

typedef struct
  float4 var1;
  float4 var2;
  float4 nextStepVar1;
  int neighbors[8];
  int var3;
  int nextStepVar2;
  bool var4;
 Element;

__kernel void step1(__global Element *elements, int nelements)
  int id = get_global_id(0);
  if (id >= nelements)
    return;
  
  Element elem = elements[id];

  for (int i=0; i < 6; ++i)
    if (elem.neighbors[i] != -1)
      //Gather information of the neighbor and accumulate it in elem.nextStepVars
    
  
  elements[id] = elem;


__kernel void step2(__global Element *elements, int nelements)
  int id = get_global_id(0);
  if (id >= nelements)
    return;
  
  Element elem = elements[id];

  //update elem variables by using elem.nextStepVariables
  //restart elem.nextStepVariables

现在，我的 OpenCL 实现与我的 C++ 实现所用的时间基本相同。

所以，问题是：您（专家：P）将如何解决这个问题？ 我已阅读有关 3D 图像的信息，以存储信息并通过将 NDRange 更改为 3D 来更改邻域访问模式。另外，我读过__local memory，首先加载工作组中的所有邻居，与屏障同步然后使用它们，从而减少对内存的访问。

您能否给我一些提示来优化我所描述的流程，如果可能的话，给我一些 sn-ps？

编辑：Huseyin Tugrul 提出的第三和第五个优化已经在代码中。正如here 所提到的，要使结构正常运行，它们需要满足一些限制，因此值得理解以避免头痛。

编辑 1：应用Huseyin Tugrul 提出的第七个优化，性能从 7 fps 提高到 60 fps。在更一般的实验中，性能增益约为 x8。

编辑 2：应用 Huseyin Tugrul 提出的第一个优化，性能提高了大约 x1.2 。我认为真正的收益更高，但因为另一个尚未解决的瓶颈而隐藏起来。

编辑 3：应用 Huseyin Tugrul 提出的第 8 和第 9 优化并没有改变性能，因为缺少利用这些优化的重要代码，但值得在其他内核中尝试。

编辑 4：如 here 所述，将不变参数（例如 n_elements 或 workgroupsize）作为 #DEFINE 而不是内核参数传递给内核，提高了 x1.33 左右的性能。正如文档中所解释的，这是因为编译器在编译时知道变量时可以进行积极的优化。

编辑 5：应用 Huseyin Tugrul 提出的第二个优化，但使用每个邻居 1 位并使用按位运算来检查邻居是否存在（因此，如果邻居 & 1 != 0，存在顶部邻居，如果邻居 & 2 != 0，则存在机器人邻居，如果邻居 & 4 ！= 0，存在右邻居等），性能提高了 x1.11 倍。我认为这主要是因为数据传输减少，因为数据移动一直是我的瓶颈。很快我将尝试摆脱用于向我的结构添加填充的虚拟变量。

编辑 6：通过消除我正在使用的结构，并为每个属性创建单独的缓冲区，我消除了填充变量，节省了空间，并且能够优化全局内存访问和本地内存分配。性能提高了 x1.25 倍，非常好。值得这样做，尽管程序复杂且不可读。

【问题讨论】：

这周有点忙，把共享数据移到本地内存，没时间测试性能。这些天我会试着去做。之后，我想尝试使用您在第四次优化中指出的技术来减少 IF 条件。我会保持更新:)

现在是 72 FPS 吗？从 8 倍到 9.2 倍？

好吧，由于 OpenGL swapBuffers，我有 60 FPS 的限制，但随着问题大小的增加，FPS 下降了，所以我可以测量加速（它总是给我 60 FPS），但没有 Vsync我猜大概是 72 FPS，是的。

所以你在做同步计算+绘图。如果后期绘制成为瓶颈，可以在绘制前计算4-5次。无论如何，人眼无法在 60fps 和 100fps 之间感知（也许专业 fps 游戏玩家可以）

现在，有 133x133x133（230 万）个元素，除了边界之外，所有元素都有 6 个邻居，我得到 20 FPS，这非常好。我不知道我是否会获得更多的表现，但我会继续努力。

【参考方案1】：

根据您的第 1 步和第 2 步，您并没有让您的 gpu 核心努力工作。你的内核的复杂度是多少？你的gpu使用率是多少？您是否检查过加力燃烧器等监控程序？中端桌面游戏卡可以获得 10k 个线程，每个线程进行 10k 次迭代。

由于您只与邻居一起工作，因此数据大小/计算大小太大，您的内核可能会受到 vram bandiwdth 的限制。您的主系统内存可能与您的 pci-e 带宽一样快，这可能是问题所在。

1) 使用专用缓存 可以让您将线程的实际网格单元放入最快的私有寄存器中。然后邻居进入 __local 数组，所以比较/计算只在芯片中完成。

将当前单元格加载到 __private

将邻居加载到 __local

开始循环本地数组

让下一个邻居从 __local 进入 __private

计算

结束循环

（如果它有很多邻居，“将邻居加载到__local”之后的行可以在另一个循环中，通过补丁从主内存中获取）

你的显卡是什么？很好，它是 GTX660。每个计算单元应该有 64kB 的可控缓存。 CPU 只有 1kB 的寄存器，不能用于数组操作。

2) 更短的索引 可以使用单个字节作为存储的邻居的索引，而不是 int。从“id”提取中节省宝贵的 L1 缓存空间很重要，这样其他线程可以更多地访问 L1 缓存！

例子：

 0=neighbour from left
 1=neighbour from right
 2=neighbour from up
 3=neighbour from down
 4=neighbour from front
 5=neighbour from back
 6=neighbour from upper left
 ...
 ...

所以您可以只从单个字节而不是 4 字节 int 派生邻居索引，这减少了至少对邻居访问的主内存访问。您的内核将使用其计算能力而不是内存能力从上表导出邻居索引，因为您将使用核心寄存器（__privates）进行此操作。如果您的总网格大小是恒定的，这非常简单，例如只需添加 1 个实际单元格 id、将 256 加到 id 或将 256*256 加到 id 左右。

3) 最佳对象大小 可能使您的结构/单元对象大小成为 4 个字节的倍数。如果您的对象总大小约为 200 字节，您可以用一些空字节填充它或增加它，使之正好为 200 字节、220 字节或 256 字节。

4) 无分支代码（编辑： 取决于！）使用较少的 if 语句。使用 if 语句会使计算速度变慢。您可以使用另一种方法，而不是检查 -1 作为邻居索引的结尾。因为轻量级的核心没有重量级的能力。您可以使用表面缓冲单元来包装表面，因此计算单元将始终具有 6 个邻居，因此您可以摆脱 if (elem.neighbors[i] != -1) 。值得一试，尤其是 GPU。

仅计算所有邻居比执行 if 语句更快。当它不是一个有效的邻居时，只需将结果变化乘以零。我们怎么知道它不是一个有效的邻居？通过使用每个单元格 6 个元素的字节数组（平行于邻居 id 数组）（无效=0，有效=1 -->将结果乘以这个）

if 语句在一个循环中，该循环计数六次。如果循环中的工作负载相对容易，则循环展开可以提供类似的加速。

但是，如果同一个 warp 中的所有线程都进入同一个 if-or-else 分支，它们不会失去性能。所以这取决于你的代码是否有分歧。

5) 数据元素重新排序 您可以将 int[8] 元素移动到 struct 的最上侧，这样内存访问可能会变得更加高效，因此可以在一次读取中读取到下侧的较小尺寸的元素-操作。

6) 工作组大小尝试不同的本地工作组大小可以提供 2-3 倍的性能。从 16 到 512 会给出不同的结果。例如，AMD GPU 喜欢 64 的整数倍，而 NVIDIA GPU 喜欢 32 的整数倍。INTEL 在 8 上做得很好，因为它可以将多个计算单元融合在一起以在同一个工作组上工作。

7) 变量分离（仅当你不能摆脱 if 语句时）比较元素与结构的分离。这样你就不需要从主内存中加载整个结构来比较一个 int 或一个 boolean。当比较需要时，然后从主内存加载结构（如果你已经有本地内存优化，那么你应该把这个操作放在它之前，这样加载到本地内存只对选定的邻居进行）

这种优化使最佳情况（没有邻居或只有一个邻居）更快。不影响最坏情况（最大邻居情况）。

8a) 魔法 使用移位而不是除以 2 的幂。对模数执行类似操作。将“f”放在浮动文字的末尾（1.0f 而不是 1.0）以避免从 double 到 float 的自动转换。

8b) Magic-2 -cl-mad-enable 编译器选项可以提高乘加运算速度。

9) 延迟隐藏 执行配置优化。您需要隐藏内存访问延迟并注意占用。

 Get maximum cycles of latency for instructions and global memory access.
 Then divide memory latency by instruction latency.
 Now you have the ratio of: arithmetic instruction number per memory access to hide latency. 
 If you have to use N instructions to hide mem latency and you have only M  instructions in your code, then you will need N/M warps(wavefronts?) to hide latency because a thread in gpu can do arithmetics while other thread getting things from mem.

10) 混合类型计算 优化内存访问后，在适用的地方交换或移动一些指令以获得更好的占用率，使用半类型来帮助精度不重要的浮点运算。 p>

11) 再次隐藏延迟 尝试仅使用算术运算的内核代码（注释掉所有内存访问并使用 0 或您喜欢的某个值启动它们）然后尝试仅使用内存访问指令的内核代码（评论出计算/如果）

将内核时间与原始内核时间进行比较。哪个对原始时间的影响更大？专注于此..

12) Lane & Bank Conflicts 纠正任何 LDS-lane 冲突和全局内存 bank 冲突，因为相同的地址访问可以以串行方式完成，减慢进程（较新的卡具有广播能力来减少这种情况）

13) 使用寄存器尝试将任何独立的本地变量替换为私有变量，因为您的 GPU 可以使用寄存器提供接近 10TB/s 的吞吐量。

14) 不使用寄存器 不要使用太多寄存器，否则它们会溢出到全局内存并减慢进程。

15) 职业的简约方法查看本地/私人使用情况以了解职业。如果您使用更多的本地和私有，那么可以在同一计算单元中使用更少的线程并导致更少的占用。更少的资源使用导致更高的占用机会（如果你有足够的总线程数）

16) Gather Scatter 当邻居是来自随机内存地址的不同粒子（如 nbody NNS）时，它可能难以应用，但 gather read 优化可以提供在优化之前（需要本地内存优化才能工作）速度提高了 2x-3x，因此它从内存中按顺序读取而不是随机读取，并根据需要在本地内存中重新排序以在（分散）到线程之间共享。

17) 分而治之 以防缓冲区太大，在主机和设备之间复制，让gpu等待空闲，然后将其一分为二，分别发送，尽快开始计算一个到达，最后并发回结果。即使是进程级别的并行性也可以通过这种方式将 gpu 推向极限。此外，GPU 的 L2 缓存可能不足以存储全部数据。缓存平铺计算，但隐式完成，而不是直接使用本地内存。

18) 来自内存限定符的带宽。当内核需要一些额外的“读取”带宽时，您可以在一些尺寸较小且仅用于读取的参数上使用“__constant”（而不是 __global）关键字。如果这些参数太大，那么您仍然可以从 '__read_only' 限定符（在 '__global' 限定符之后）获得良好的流式传输。类似的 '__write_only' 增加了吞吐量，但这些主要提供了特定于硬件的性能。如果是amd的HD5000系列，恒是不错的。也许 GTX660 的缓存速度更快，因此 __read_only 可能会变得更可用（或者 Nvidia 使用缓存作为 __constant？）。

拥有相同缓冲区的三个部分，其中一个作为 __global __read_only，一个作为 __constant，一个作为 __global（如果构建它们不会比读取带来的好处更多）。

刚刚使用 AMD APP SDK 示例测试了我的卡，LDS 带宽显示为 2TB/s，而常数为 5TB/s（相同的索引而不是线性/随机），主内存为 120 GB/s。

也不要忘记在可能的情况下将 restrict 添加到内核参数。这让编译器可以对它们进行更多优化（如果你没有给它们起别名的话）。

19) 现代硬件超越函数比旧的 bit hack（如 Quake-3 快速反平方根）版本更快

20) 现在有 Opencl 2.0 可以在内核中生成内核，因此您可以进一步提高二维网格点的分辨率并在需要时将其卸载到工作组（例如增加涡度细节在流体边缘动态）

分析器可以帮助解决所有这些问题，但任何 FPS 指标都可以做到，如果每一步只进行一次优化。

即使基准测试不适用于依赖于架构的代码路径，您也可以尝试在计算空间中每行包含 192 个点的倍数，因为您的 gpu 具有该数量的多个内核，并且如果它使 gpu 更多，则可以进行基准测试占用并且每秒有更多的千兆浮点操作。

在所有这些选项之后，肯定还有一些优化空间，但是如果它损坏了您的卡或对您的项目的生产时间可行，请注意。例如：

21) 查找表当内存带宽余量增加 10% 但没有计算能力余量时，将这些工作项中的 10% 卸载到 LUT 版本，以便获得表中预先计算的值。我没有尝试，但这样的事情应该可以工作：

8 个计算组 2 个 LUT 组 8 个计算组 2 个 LUT 组

因此它们被均匀地分布到“运行中的线程”中，并利用延迟隐藏的东西。我不确定这是否是一种更好的科学研究方式。

21) Z 顺序模式对于旅行的邻居提高缓存命中率。缓存命中率为其他作业节省了一些全局内存带宽，从而提高了整体性能。但这取决于缓存的大小、数据布局和其他一些我不记得的东西。

22) 异步邻居遍历

iteration-1：加载邻居 2 + 计算邻居 1 + 存储邻居 0 iteration-2：加载邻居 3 + 计算邻居 2 + 存储邻居 1 iteration-3：加载邻居 4 + 计算邻居 3 + 存储邻居 2

因此每个循环体都没有任何依赖链，并且在 GPU 处理元素上完全流水线化，并且 OpenCL 具有使用工作组的所有内核异步加载/存储全局变量的特殊指令。检查这个：

https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/async_work_group_copy.html

也许您甚至可以将计算部分一分为二，让一部分使用超越函数，另一部分使用加法/乘法，这样加法/乘法运算就不会等待缓慢的 sqrt。如果至少有几个邻居要遍历，这应该会在其他迭代之后隐藏一些延迟。