如何在 OpenCL 中使用 async_work_group_copy？

Question

我想了解如何在 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() 并忽略返回值吗？如果是这样，那可能会更快吗？

b.本地副本是否也适用于 CPU 上的处理，还是因为它们共享缓存而产生的开销？

问候， 斯蒂芬

Answer 1

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

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

对于您的具体问题：

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

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

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

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

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

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