11

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

我开始使用 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，首先加载工作组中的所有邻域，与屏障同步然后使用它们，从而减少对内存的访问。

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

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

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

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

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

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

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

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

4

1 回答 1

19

根据您的第 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）以避免从双精度自动转换为浮点。

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

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

将内核时间与原始内核时间进行比较。哪个对原始时间的影响更大？专心那个。。

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

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。

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

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

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

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

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

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

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

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

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

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

22) 异步邻居遍历

  • 迭代 1：加载邻居 2 + 计算邻居 1 + 存储邻居 0
  • 迭代 2：加载邻居 3 + 计算邻居 2 + 存储邻居 1
  • 迭代 3：加载邻居 4 + 计算邻居 3 + 存储邻居 2

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

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

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

于 2013-07-02T09:30:54.333 回答