通知OpenCL内核多个内存对象的正确方法是什么？

Question

在我的OpenCL程序中，我需要使用60多个全局内存缓冲区，并且每个内核都需要能够访问这些缓冲区。有没有推荐的方法来让每个内核知道这些缓冲区的位置？

这些缓冲区在应用程序的整个生命周期中都是稳定的，也就是说，我们将在应用程序启动时分配这些缓冲区，在调用多个内核后，仅在应用程序结束时释放这些缓冲区。但是，它们的内容可能会随着内核对它们进行读/写而发生更改。

在CUDA中，我是通过在CUDA代码中创建60多个程序范围的全局变量来实现这一点的。然后，在主机上，我将分配的设备缓冲区的地址写入这些全局变量中。然后内核只需使用这些全局变量即可找到其需要处理的缓冲区。

那么，在OpenCL中该如何做呢？似乎CL的全局变量与CUDA的有所不同，但我找不到明确的答案，无法确定我的CUDA方法是否可行，如果可行，如何将缓冲指针传输到全局变量中。如果这种方法行不通，还有其他最佳方法吗？

- int3h

Answer 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和offsets传递给您的内核，就可以访问每个数组。例如：

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，它可能更快。它肯定会减少参数列表的长度。

- Ryan Marcus

---

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

Answer 2

您需要做两件事。首先，每个使用全局内存缓冲区的内核应该为每个 声明一个参数，类似于这样：

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

这样每个内核使用的缓冲区都会被列出。然后，在主机端，您需要创建每个缓冲区，并使用clSetKernelArg将给定的内存缓冲区附加到给定的内核参数上，然后调用clEnqueueNDRangeKernel来启动程序。

请注意，如果内核将在每个内核执行中继续使用相同的缓冲区，则只需要设置内核参数一次。我经常看到的一个常见错误是，在完全不必要的情况下反复调用clSetKernelArg，这可能会影响主机性能。

- James

---

那么每个内核函数都必须有60个参数，这是无法避免的吗？我知道我可以保留这些参数，但这意味着每个线程一开始就使用240字节的内存来存储指针，而这些指针从不改变且对于所有内核函数都相同。此外，传递参数可能会带来潜在的性能开销——这些内核函数将被调用60次/秒。 - int3h