如何在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()并忽略返回值吗？如果是这样,那会更快吗？

湾 本地副本是否对CPU上的处理有意义,还是因为它们共享缓存而产生的开销？

此致,斯特凡

Answer 1

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

您描述了需要复制的内存,就像它是单线程副本一样,async_work_group_copy使用并行硬件为您完成.

针对您的具体问题:

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

2. 是.所有工作项的源(和目标)地址必须相同.

3. 您可以减去本地ID以获取正确的地址,但我发现基于groupId的地址也解决了这个问题.(get_group_id)

4. 是.最后一个参数是元素的数量,而不是以字节为单位的大小.

一个.不会.基于事件,您会发现工作项几乎立即触及您的屏障,并且不一定会复制数据.这是有道理的,因为一些opencl硬件甚至可能根本不使用计算单元来进行实际的复制操作.

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