Results 1 to 8 of 8

Thread: What happens if one of kernel argument is const int ?

  1. #1
    Junior Member
    Join Date
    Apr 2010
    Posts
    27

    What happens if one of kernel argument is const int ?

    Hello. I have a question about kernel arguments.

    Assuming my kernel is defined as something like this:

    Code :
    __kernel void 
    my_kernel_func(__global const int w)

    I pass the argument:

    Code :
    clSetKernelArg(kernel_obj, 0, sizeof(cl_int), &w) ;

    Ok, then, how exactly the compiler treats the const int w ?
    I mean, is the argument 'w' treated as constant integer and optimized out in compile time ?
    Or, is it just in a global memory and is accessing the variable become slow ?

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

    Re: What happens if one of kernel argument is const int ?

    Are you sure you didn't mean "__global const int *w"?

    This code:

    Code :
    __kernel void 
    my_kernel_func(__global const int w)

    doesn't appear to be valid OpenCL C. See section 6.5 of the spec:

    All function arguments shall be in the __private address space.
    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
    Junior Member
    Join Date
    Apr 2010
    Posts
    27

    Re: What happens if one of kernel argument is const int ?

    Quote Originally Posted by david.garcia
    Are you sure you didn't mean "__global const int *w"?
    Well, my kernel is as it is and it work without problems until now.

    I couldn't exactly understand what the spec says...

    should I use pointer to pass a constant to a kernel ?

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

    Re: What happens if one of kernel argument is const int ?

    Well, my kernel is as it is and it work without problems until now.
    Ehrm... okay. There's a bug in that compiler

    should I use pointer to pass a constant to a kernel ?
    You can pass a 'const int' to a kernel. Don't declare it as '__global' and everything should be fine.

    Going back to your original question:

    Ok, then, how exactly the compiler treats the const int w ?
    That will depend on the OpenCL implementation. I can say that it will not be treated as a compile-time constant because its value is not known at compilation time. I wouldn't worry about its performance.
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

  5. #5
    Junior Member
    Join Date
    Apr 2010
    Posts
    27

    Re: What happens if one of kernel argument is const int ?

    Thanks !
    You helped me understand OpenCL better.

  6. #6

    Re: What happens if one of kernel argument is const int ?

    Quote Originally Posted by david.garcia
    Well, my kernel is as it is and it work without problems until now.
    Ehrm... okay. There's a bug in that compiler
    Could you elaborate? I have plenty of "const unsigned int X" as parameters that work on several OCL compilers, do you mean that any __global parameters must be pointers?

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

    Re: What happens if one of kernel argument is const int ?

    Could you elaborate? I have plenty of "const unsigned int X" as parameters that work on several OCL compilers
    Yeah, but that example you've given is syntactically correct. The issue only occurs with declarations like "__global int X" or "__global const int X".

    do you mean that any __global parameters must be pointers?
    This is the thing: when you declare a variable as "__global int* foo", it doesn't mean that variable "foo" is in global memory. It means that it points to global memory. In other words, the value of "foo" is the address of some data in the __global address space.

    "foo" itself will be in whatever is the default address space for variables in that scope since the declaration of "foo" didn't specify any address space. In the case of kernel arguments, the default (and only valid) address space for "foo" is __private memory.

    I hope it's a bit less confusing now.
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

  8. #8
    Junior Member
    Join Date
    Apr 2010
    Posts
    27

    Re: What happens if one of kernel argument is const int ?

    Ehrm... okay. There's a bug in that compiler
    By the way, I tested the example above on a iMac with ATI Radeon HD 5750 and OS X 10.6.5.

Similar Threads

  1. Compiling Kernel With Struct Pointer in Argument
    By ConstantinS in forum OpenCL
    Replies: 1
    Last Post: 03-29-2013, 02:58 PM
  2. Kernel argument size
    By perlish in forum OpenCL
    Replies: 5
    Last Post: 09-03-2012, 05:55 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
  •