通知 OpenCL 内核许多内存对象的正确方法？

Posted 2023-02-16

【中文标题】通知 OpenCL 内核许多内存对象的正确方法？

【英文标题】：Proper way to inform OpenCL kernels of many memory objects?

【发布时间】：2012-06-16 11:20:22

【问题描述】：

在我的 OpenCL 程序中，我将最终获得 60 多个全局内存缓冲区，每个内核都需要能够访问这些缓冲区。让每个内核知道每个缓冲区的位置的推荐方法是什么？

缓冲区本身在应用程序的整个生命周期中都是稳定的——也就是说，我们将在应用程序启动时分配缓冲区，调用多个内核，然后只在应用程序结束时释放缓冲区。然而，它们的内容可能会随着内核读取/写入它们而改变。

在 CUDA 中，我这样做的方式是在我的 CUDA 代码中创建 60 多个程序范围的全局变量。然后，我会在主机上将分配的设备缓冲区的地址写入这些全局变量中。然后内核将简单地使用这些全局变量来找到它需要使用的缓冲区。

在 OpenCL 中执行此操作的最佳方法是什么？似乎 CL 的全局变量与 CUDA 的有点不同，但我无法找到关于我的 CUDA 方法是否有效的明确答案，如果可以，如何将缓冲区指针转换为全局变量。如果这不起作用，那么最好的方法是什么？

【参考方案1】：

60 个全局变量肯定不少！您确定没有办法稍微重构您的算法以使用更小的数据块吗？请记住，每个内核都应该是最小的工作单元，而不是巨大的！

但是，有一种可能的解决方案。假设您的 60 个数组的大小已知，您可以将它们全部存储到一个大缓冲区中，然后使用偏移量来访问该大数组的各个部分。这是一个非常简单的三个数组示例：

A is 100 elements
B is 200 elements
C is 100 elements

big_array = A[0:100] B[0:200] C[0:100]
offsets = [0, 100, 300]

然后，您只需将 big_array 和偏移量传递给您的内核，您就可以访问每个数组。例如：

A[50] = big_array[offsets[0] + 50]
B[20] = big_array[offsets[1] + 20]
C[0] = big_array[offsets[2] + 0]

我不确定这会如何影响您特定设备上的缓存，但我最初的猜测是“不太好”。这种数组访问也有点讨厌。我不确定它是否有效，但您可以使用一些代码启动每个内核，这些代码提取每个偏移量并将其添加到原始指针的副本中。

在主机端，为了让您的数组更易于访问，您可以使用 clCreateSubBuffer:http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateSubBuffer.html 这也允许您在没有偏移数组的情况下传递对特定数组的引用。

我认为这个解决方案不会比传递 60 个内核参数更好，但取决于您的 OpenCL 实现的 clSetKernelArgs，它可能会更快。它肯定会减少参数列表的长度。

【讨论】：

60 个参数是因为此代码是由我参与的研究项目的特殊代码合成器生成的。不幸的是，我无法控制那部分。我确实最终使用了您概述的缓冲区打包方法。希望这是比 60 个参数更好的方法。感谢您的帮助！

【参考方案2】：

你需要做两件事。首先，使用每个全局内存缓冲区的每个内核都应该为 each 声明一个参数，如下所示：

kernel void awesome_parallel_stuff(global float* buf1, ..., global float* buf60)

以便列出该内核的每个已使用缓冲区。然后，在 host 端，您需要创建每个缓冲区并使用 clSetKernelArg 将给定的内存缓冲区附加到给定的内核参数，然后再调用 clEnqueueNDRangeKernel 来开始聚会。

请注意，如果内核将在每次内核执行时继续使用相同的缓冲区，则您只需设置内核参数一次次。我看到的一个常见错误是在完全没有必要的情况下重复调用clSetKernelArg，这会降低主机端的性能。

【讨论】：

所以没有办法为每个内核函数设置 60 个参数？我意识到我可以保留这些参数，但是仍然，这意味着每个线程开始使用 240 字节的内存仅用于指针，它永远不会改变并且对于所有内核都是相同的。此外，传递参数还有潜在的性能开销——这些内核将被调用 60 次/秒。