c++ - 如何避免 OpenCL 中的常量内存复制

Question

我编写了模拟简单热流的 C++ 应用程序。它使用 OpenCL 进行计算。 OpenCL 内核采用二维 (n x n) 温度值数组及其大小 (n)。它在每个循环后返回包含温度的新数组:

伪代码:

int t_id = get_global_id(0);
if(t_id < n * n)
{
    m_new[t_id / n][t_id % n] = average of its and its neighbors (top, bottom, left, right) temperatures
}

如您所见，每个线程都在计算矩阵中的单个单元格。当主机应用程序需要执行 X 个计算周期时，它看起来像这样

• 对于 1 ... X
  1. 将内存复制到 OpenCL 设备
  2. 调用内核
  3. 复制内存

我想重写内核代码以执行所有 X 周期，而无需向/从 OpenCL 设备复制常量内存。

1. 将内存复制到 OpenCL 设备
2. 调用内核 X 次或调用内核一次并使其计算 X 周期。
3. 复制内存

我知道内核中的每个线程都应该在所有其他线程都在完成它们的工作时锁定，之后 - m[][] 和 m_new[][] 应该被交换。我不知道如何实现这两个功能中的任何一个。

或者也许有另一种最佳方式来做到这一点？

Answer 1

Copy memory to OpenCL device
Call kernel X times
Copy memory back

这行得通。确保 call kernel 没有阻塞(因此每个周期节省 1-2 毫秒)并且没有任何主机可访问的缓冲区属性，例如 USE_HOST_PTR 或 ALLOC_HOST_PTR。

如果调用内核 X 次没有获得令人满意的性能，您可以尝试使用单个工作组(例如只有 256 个线程)循环 X 次，每个循环在末尾都有一个 barrier() 以便所有 256 个线程在开始之前同步下一个周期。通过这种方式，您可以同时计算 M 个不同的热流问题，其中 M 是计算单元(或工作组)的数量，如果它是一个服务器，它可以处理那么多的计算。

全局同步是不可能的，因为当最新的线程启动时，第一个线程已经消失了。它与(计算单元数)(每个工作组的线程数)(每个工作组的波前数)线程同时工作。例如，具有 5 个计算单元且 local-range=256 的 R7-240 gpu，它一次可以运行 5*256*20=25k 个线程。

然后，为了进一步提高性能，您可以应用本地内存优化。