Results 1 to 5 of 5

Thread: __constant array initialization

  1. #1
    Junior Member
    Join Date
    Aug 2011
    Posts
    3

    __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

    Code :
    __kernel void test()
     
    {
         __constant int ow = 1; 
    }

    doesn't compile with error

    Code :
    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?

  2. #2
    Member
    Join Date
    Jul 2011
    Location
    Moscow, Russia
    Posts
    41

    Re: __constant array initialization

    > 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.
    Blog (in russian)

  3. #3
    Junior Member
    Join Date
    Aug 2011
    Posts
    3

    Re: __constant array initialization

    Oh, thanks. Then I even more don't want to use it 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:
    Code :
       __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.

  4. #4
    Member
    Join Date
    Jul 2011
    Location
    Moscow, Russia
    Posts
    41

    Re: __constant array initialization

    Quote Originally Posted by Michail
    Oh, thanks. Then I even more don't want to use it
    4xxx were not designed with OpenCL in mind. 5xxx (and later) were.
    Blog (in russian)

  5. #5
    Junior Member
    Join Date
    Aug 2011
    Posts
    3

    Re: __constant array initialization

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

Similar Threads

  1. Array initialization failed
    By pelangi15 in forum OpenCL
    Replies: 3
    Last Post: 06-13-2011, 07:14 PM
  2. Fast array initialization
    By nachovall in forum OpenCL
    Replies: 4
    Last Post: 03-16-2011, 12:15 PM

Posting Permissions

  • You may not post new threads
  • You may not post replies
  • You may not post attachments
  • You may not edit your posts
  •