如何避免 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 个线程。

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