15

我想了解如何在 OpenCL 中正确使用 async_work_group_copy() 调用。让我们看一个简化的例子：

__kernel void test(__global float *x) {
  __local xcopy[GROUP_SIZE];

  int globalid = get_global_id(0);
  int localid = get_local_id(0);
  event_t e = async_work_group_copy(xcopy, x+globalid-localid, GROUP_SIZE, 0);
  wait_group_events(1, &e);
}

参考http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/async_work_group_copy.html说“从 src 到 dst 执行 num_elements gentype 元素的异步复制。异步复制由因此，工作组中的所有工作项和此内置函数必须由执行具有相同参数值的内核的工作组中的所有工作项遇到；否则结果未定义。”

但这并不能澄清我的问题......

我想知道，如果以下假设是正确的：

  1. 对 async_work_group_copy() 的调用必须由组中的所有工作项执行。
  2. 调用的方式应该是，所有工作项的源地址都相同，并指向要复制的内存区域的第一个元素。
  3. 因为我的源地址是相对的，基于工作组中第一个工作项的全局工作项 ID。所以我必须减去本地 id 以使所有工作项的地址都相同......
  4. 第三个参数真的是元素的数量（不是字节大小）吗？

奖励问题：

一个。我可以只使用 barrier(CLK_LOCAL_MEM_FENCE) 而不是 wait_group_events() 并忽略返回值吗？如果是这样，那可能会更快吗？

湾。本地副本对于在 CPU 上进行处理是否也有意义，还是因为它们共享缓存而产生的开销？

问候，斯特凡

4

1 回答 1

12

存在此功能的主要原因之一是允许驱动程序/内核编译器有效地复制内存，而无需开发人员对硬件做出假设。

您将需要复制的内存描述为单线程副本，async_work_group_copy 使用并行硬件为您完成。

对于您的具体问题：

  1. 我从未见过 async_work_group_copy 仅被组中的某些工作项使用。我一直认为这是因为它需要。我认为 wait_group_events 的阻塞性质迫使所有工作项成为副本的一部分。

  2. 是的。所有工作项的源（和目标）地址必须相同。

  3. 您可以减去您的本地 ID 以获得正确的地址，但我发现基于 groupId 的地址也可以解决此问题。(get_group_id)

  4. 是的。最后一个参数是元素的数量，而不是字节大小。

一个。不会。基于事件的你会发现你的障碍几乎立即被工作项击中，数据不一定会被复制。这是有道理的，因为一些 opencl 硬件甚至可能根本不使用计算单元来执行实际的复制操作。

湾。我认为当您使用本地内存时，cpu opencl 实现可能会保证 L1 缓存的使用。确定这是否表现更好的唯一方法是使用各种设置对您的应用程序进行基准测试。

于 2013-03-21T13:10:49.100 回答