Results 1 to 5 of 5

Thread: passing constant value to kernel

  1. #1
    Junior Member
    Join Date
    Sep 2009
    Posts
    7

    passing constant value to kernel

    The OpenCL 1.0 does not allow to pass a number, e.g. __constant int val, to kernel from the constant memory. Only pointers to constant memory space may be passed. As a result, one has to pass the number as private (__private const int val). This clearly increases the register pressure. I wonder why this is the case.

    In NVIDIA's implementation, constant memory is cached and is as fast as local registers if there is no cache miss.

  2. #2
    Senior Member
    Join Date
    Jul 2009
    Location
    Northern Europe
    Posts
    311

    Re: passing constant value to kernel

    I don't quite follow. Are you talking about passing values in to kernels via clSetKernelArgs or calling functions from within a kernel? If it's the former I'm a bit confused because you can't pass in private values because the private space is per-work-item, not per kernel launch. Constant values can be accessed by all work-items in the kernel launch, but private values are unique to one work-item and can not be shared. If it's the latter case, then you should be able to pass a private variable between function calls within a single work-item.

  3. #3
    Junior Member
    Join Date
    Sep 2009
    Posts
    7

    Re: passing constant value to kernel

    Suppose we have a kernel like this
    Code :
    __kernel void foo(const int val, ...) 
    { ... }
    In this example, the same value of "val" is passed from the host program to each work item. But it is located in the private space. In other words, each work item has a register holding the same "val". This is a waste of private space. It would be nice to change "const int val" to "__constant int val" which is not allowed by OpenCL 1.0.

  4. #4
    Member
    Join Date
    Sep 2009
    Posts
    35

    Re: passing constant value to kernel

    __constant variables are allocated in global memory, and the constant memory space is cached. Accessing private memory is faster and should remain as only option where to store function arguments (as they are top priority of things you want to access from inside functions).

  5. #5
    Junior Member
    Join Date
    Sep 2009
    Posts
    7

    Re: passing constant value to kernel

    Quote Originally Posted by jbasic
    __constant variables are allocated in global memory, and the constant memory space is cached. Accessing private memory is faster and should remain as only option where to store function arguments (as they are top priority of things you want to access from inside functions).
    I understand your argument. However, at least in NVIDIA's implementation, constant memory is as fast as private memory as long as it is cached. So I wonder why the standard did not allow it since it does not hurt.

Similar Threads

  1. Passing struct to kernel
    By helloOpenCL in forum OpenCL
    Replies: 6
    Last Post: 08-29-2011, 02:56 PM
  2. Passing NULL pointer to the kernel ?
    By viewon01 in forum OpenCL
    Replies: 2
    Last Post: 07-26-2010, 06:21 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
  •