Results 1 to 3 of 3

Thread: constant memory

  1. #1
    Junior Member
    Join Date
    Sep 2010
    Posts
    11

    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

  2. #2
    Senior Member
    Join Date
    May 2010
    Location
    Toronto, Canada
    Posts
    845

    Re: constant memory

    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.
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

  3. #3
    Member
    Join Date
    Oct 2010
    Location
    Vancouver, Canada
    Posts
    66

    Re: constant memory

    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).

Similar Threads

  1. Constant memory
    By Adam Simpson in forum OpenCL
    Replies: 1
    Last Post: 12-07-2009, 12:32 PM
  2. Constant Memory latency
    By PaulS in forum OpenCL
    Replies: 3
    Last Post: 10-21-2009, 08:06 AM

Posting Permissions

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