如何在OpenCL中使用async_work_group_copy？

我想了解如何在OpenCL中正确使用async_work_group_copy（）调用。让我们在一个简单的例子来看看：如何在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 num_elements gentype元素的异步复制到DST异步副本由所有工作项目在一个工作组进行，这个内置函数因此必须被执行内核的工作组中的所有工作项所遇到，并且具有相同的参数值;否则结果是不确定的。“

但是，这并不清楚我的问题......

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

的调用async_work_group_copy（）必须由所有工作执行 - 组中的项目。
调用的方式应该是，所有工作项的源地址都相同，并指向要复制的存储区的第一个元素。
由于我的源地址是相对基于工作组中第一个工作项的全局工作项ID。所以我必须减去本地ID以使所有工作项的地址相同...
第三个参数是否是真正的元素数量（而不是字节大小）？

奖金的问题：

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

b。本地副本对于CPU上的处理还是有意义的，还是因为它们共享高速缓存而造成的开销？

问候，斯特凡

来源

2013-03-21 SDwarfs

其中一个主要原因此功能是存在的，允许驱动器/内核编译器有效，而无需进行有关硬件的假设开发商的内存复制。

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

为了您的具体问题：

我从来没有见过async_work_group_copy只用了一些组中的工作项目。我一直认为这是因为它需要。我认为wait_group_events的阻塞性质会强制所有工作项目成为副本的一部分。
是的。源（和目标）地址对于所有工作项目都必须相同。
你可以减去你的本地ID来得到正确的地址，但是我发现基于groupId的地址解决了这个问题。（get_group_id）
是的。最后一个参数是元素的数量，而不是字节的大小。

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

b。我认为，使用本地内存时，cpu opencl实现可能会保证L1缓存使用率。唯一可以肯定的是，如果这种方法性能更好，则可以使用各种设置对应用程序进行基准测试

来源

2013-03-21 13:10:49 mfa

谢谢！关于3.get_group_id（） - 很好的提示，但在这个特殊的需要乘以工作组的大小，即使对于已知的工作组大小为2^n这可能会加快左移。。但是，在其他情况下，这可能是非常有用的知道它也存在！感谢您分享您的体验！ --- Stefan – SDwarfs 2013-03-21 15:05:54

我对你对“a”的回答有一个疑问：据我所知，复制应至少由一个工作项目处理。由于工作组中的所有其他节点都需要准备好本地数据，因此我认为它们都等待数据可用。等待事件应该做到这一点。但是障碍应该保证是一样的，因为小组的所有其他工作项目都会等待复制过程中的工作项目。障碍可能更简单（每个工作组最多1个障碍，而每个工作项目最多等待1个事件），因此速度更快。 – SDwarfs 2013-03-21 15:21:14

re：'a'...未来可能存在一个处理异步拷贝而不使用任何工作项目的设备。如果至少有一个工作项目被复制操作阻止，则障碍将起作用。 – mfa 2013-03-21 18:31:32

如何在OpenCL中使用async_work_group_copy？

回答

相关问题