constant memory

Hi there,

  1. What is the difference between “const” and “__constant”?
    Example:
    __kernel void function_cl (__constant float* var1, const uint var2)
    {

    }

  2. When I want to use constant variables in the function interface, do I have to define them differently, e.g. with CL_MEM_READ_ONLY in clCreateBuffer() ?

regards,
Fred

What is the difference between “const” and “__constant”?

Constant memory is physically separate from global memory on GPUs. On CPUs there’s no difference. Constant memory is typically faster to use than simply declaring a variable as “const”.

When I want to use constant variables in the function interface, do I have to define them differently, e.g. with CL_MEM_READ_ONLY in clCreateBuffer()

No, you can use any buffer. The OpenCL driver will then load the data from that buffer into constant memory transparently to the application.

Constant memory is physically separate from global memory on GPUs. On CPUs there’s no difference. Constant memory is typically faster to use than simply declaring a variable as “const”.

I would amend this slightly to say that constant memory may be physically separate from global memory on GPUs… and its possible that CPUs could implement this as well (by changing caching modes, memory banks, etc.), although I’m not aware of any OpenCL implementations that actually do this (so far).

Also note that “const” is a compiler level thing, and you can actually cast it away. In C/C++ const tends to be mostly informational for the programmer. Once the code is compiled into instructions, there is no evidence of what was marked ‘const’ and what wasn’t.

In OpenCL, however, __constant is a different memory space. You cannot cast it to a different memory space because it might be physically different memory, have different pointer sizes, use physical addresses that overlap with other spaces, etc. The runtime should behave differently in how it treats __constant memory (e.g. no writes allows, and while it sends buffers to __constant it should never need to bring them back from there).