title: How do I make a strided copy from global to local memory?
body: <p>I want to copy some data from a buffer in the global device memory to the local memory of a processing core - but, with a twist.</p>
  
  <p>I know about <a href="http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/async_work_group_copy.html" rel="nofollow">async_work_group_copy</a>, and it's nice (or rather, it's klunky and annoying, but working). However, my data is not contiguous - it is strided, i.e. there might be X bytes between every two consecutive Y bytes I want to copy. </p>
  
  <p>Obviously I'm not going to copy all the useless data - and it might not even fit in my local memory. What can I do instead? I want to avoid writing actual kernel code to do the copying, e.g.</p>
  
  <pre><code>threadId = get_local_id(0);
  if (threadId &lt; length) {
      unsigned offset = threadId * stride;
      localData[threadId] = globalData[offset];
  }
  </code></pre>