I would like to understand how to correctly use the async_work_group_copy() call in OpenCL. Let's have a look on a simplified example:

__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);
}

The reference http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/async_work_group_copy.html says "Perform an async copy of num_elements gentype elements from src to dst. The async copy is performed by all work-items in a work-group and this built-in function must therefore be encountered by all work-items in a workgroup executing the kernel with the same argument values; otherwise the results are undefined."

But that doesn't clarify my questions...

I would like to know, if the following assumptions are correct:

  1. The call to async_work_group_copy() must be executed by all work-items in the group.
  2. The call should be in a way, that the source address is identical for all work-items and points to the first element of the memory area to be copied.
  3. As my source address is relative based on the global work-item id of the first work-item in the work-group. So I have to subtract the local id to have the address identical for all work-items...
  4. Is the third parameter really the number of elements (not the size in bytes)?

Bonus questions:

a. Can I just use barrier(CLK_LOCAL_MEM_FENCE) instead of wait_group_events() and ignore the return value? If so, would that be probably faster?

b. Does a local copy also make sense for processing on CPUs or is that overhead as they share a cache anyway?

Regards, Stefan

opencl
Edit
Report