OpenCL: __constant vs. __local?

0
votes

Let's say I have a large array of values (still smaller than 64 kB), which is read very often in the kernel, but not written to. It can however change from outside. The array has two sets of values, lets call them left and right. So the question is, is it faster to get the large array as a __global and write it into __local left and __local right arrays; or get it as a constant __constant large and handle the accesing in the kernel? For example:

__kernel void f(__global large, __local left, __local right, __global x, __global y) {
    for(int i; i < size; i++) {
        left[i] = large[i];
        right[i] = large[i + offset];
    }
    ...
    x = foo * left[idx];
    y = bar * right[idx];
}

vs:

__kernel void f(__constant large, __global x, __global y) {
    ...
    x = foo * large[idx];
    y = bar * large[idx * offset];
}

(The indexing is a bit more complicated, but can be made with macros, for instance) I read that constant memory lives in the global space, so should it be slower? It will run in a Nvidia card.

1
openclgpgpu

1 Answers

1
votes

First of all in the second case you should have someway of making the result available for your CPU. I am assuming you copy back to a global space after computation.

I think it depends on what you do in the kernel. For example if you kernel computation is heavy (a lot of computations per thread) then the first option might pay of. Why?

  • You spend some time copying data from global large space to local spaces left and right - Acceptable
  • You do a lot of computation on the data on local space - OK
  • You spend some time copying back from local left and right to global large. - Acceptable.

However if you kernel is relatively light i.e. each thread will do some small computations, then

  • You do a few computations with data on constant space. Which most probably means you don't need to access it a lot.
  • You store intermediate results in local space.
  • You spend some time copying back from local space to global space. - Acceptable.

To sum it up for large kernels the first option is better. For small kernels the second.

P.S. One more thing to note is that if you have multiple kernels that wwork on large one after the other, then definitely go with the first option. Because then you can keep the data on global memory space and you don't have to do copy every time you launch a kernel.

EDIT: since you have said it is accessed very often then I think you should probably go with the first option.