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:

__kernel void 
my_kernel_func(__global const int w) 

I pass the argument:

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 ?

Are you sure you didn’t mean “__global const int *w”?

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

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 ?

Well, my kernel is as it is and it work without problems until now.

Ehrm… okay. There’s a bug in that compiler :slight_smile:

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.

Thanks !
You helped me understand OpenCL better.

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?

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.

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.