__constant array initialization

Hello!

I’m new to OpenCL and now I’m trying to port one video filter to OCL (primarily for studying :)).

All works fine except not very fast speed (actually it’s a little faster than previous CPU code). Now I’m working on it’s optimization. I need to highly reuse 2 constant arrays of float of size less than 128. My GPU is 4850, so I can’t use Images. Local memory works, but limits max. Work Group Size to 64. Now I’m trying to use __constant memory. The fastest way I found is to initialize them in kernel as constants (generated in host code before compilation). But I didn’t found any way to do it in OpenCL C. Moreover, I can’t find example to init. one non-array variable. Code

__kernel void test()

{
     __constant int ow = 1; 
}

doesn’t compile with error

error: non-kernel
          function: variable with automatic storage duration cannot be stored
          in a named address space
  __constant int ow = 1; 
                 ^

in AMD APP KernelAnalyzer.

If i replace __constant with __private all works, but GRP usage is unacceptable.

Is there any way of initializing __constant variable inside kernel?

> My GPU is 4850, Local memory works, but limits max

Actually, no. Local memory is emulated in global memory for ATI 4xxx devices. You can figure out this by querying device for CL_DEVICE_LOCAL_MEM_TYPE.

Oh, thanks. Then I even more don’t want to use it :slight_smile: Espessially when barrier for it limits Work Group Size.
I’ll answer to myself here, answer was found in AMD forum. __constant vars needs to be declared and initialized globally.
That example compiles fine:

   __constant int ow[] = {1,2,3,4,5,6,7,8,9,10}; 

__kernel void test(__global int *out)
{
   out[get_local_size(0)] = ow[get_local_size(0)];
}

Tomorrow I’ll try it in real app.

4xxx were not designed with OpenCL in mind. 5xxx (and later) were.

I know this, but now I don’t want to buy new video for just some experiments with OCL :slight_smile: