<p>You already have the fastest solution.</p>
<p>To understand why, we need to dig a bit deeper.</p>
<p>The memory bandwidh of a GPU is different for coalesced/misaligned read, and this is different for read/write operations. <a href="https://doi.org/10.48550/arXiv.2112.08926" rel="nofollow noreferrer">On most GPUs</a>, while misaligned writes are almost as fast as coalesced reads, for misaligned writes there is a large performance penalty. So misaligned reads are ok, but misaligned writes should be avoided.<a href="https://i.stack.imgur.com/wgzGE.png" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/wgzGE.png" alt="enter image description here" /></a></p>
<p>In your example, you have coalesced writes to <code>A</code> and misaligned reads from <code>B</code>, so you already get peak memory bandwidth.</p>
<pre class="lang-c prettyprint-override"><code>kernel void reverse_kernel(global float* A, global float* B) { // equivalent to &quot;for(uint i=0u; i&lt;N; i++) {&quot;, but executed in parallel
    const uint i = get_global_id(0);
    A[i] = B[N-1-i];
}
</code></pre>
<p>For coalesced memory access, generally contiguous threads must access contiguous memory addresses. But there is some special cases where coalesced access is still allowed in differend access patterns: strided access and broadcasting. I am not sure if reverse access also falls into that class, but for the reads it doesn't matter anyways.</p>
<p>To your initial question with shared memory: This is how you would do it. But there probably won't be any significant speedup here. Speedup with shared memory is only large if the shared memory is accessed multiple times; here you write and read it only once.</p>
<pre class="lang-c prettyprint-override"><code>kernel void reverse_kernel(global float* A, global float* B) { // equivalent to &quot;for(uint i=0u; i&lt;N; i++) {&quot;, but executed in parallel
    const uint i = get_global_id(0);
    const uint lid = get_local_id(0);
    const uint gid = get_group_id(0);
    local float cache[64]; // workgroup size on C++ size must be set to 64 as well
    cache[lid] = B[64*gid+lid]; // coalesced read from B
    barrier(CLK_LOCAL_MEM_FENCE);
    A[i] = cache[64-1-lid]; // coalesced write to A
}
</code></pre>
Reverse of an array in OpenCL avoiding Uncoalasced Memory Access
<p>I have this function written in C, that reverse an array:</p>
<pre><code>for (int i=0; i&lt;N;i++){
  A[i] = B[N-i-1];
}
</code></pre>
<p>I have to write a kernel function that not suffer of Uncoalasced Memory Access (given at B[N-i-1]) using tiling and the local memory. So the idea is: Doing reverse in a local memory and write result back in the array A. How i can do it? Im a newebie.
Assumption: input size match with global size.</p>