Passing an argument as 'const '

Hello, developers.

Today, I have a question about passing an arguments to a kernel.

When my kernel receives an integer as ‘const int’, what actually happens when the kernel is built ?

__global__ my_kernel(float *src, float *dst, const int a) 

{

....

  dst[xx] = src[xx]  + a ;

}

I mean… Is it optimized out in compilation if the value of argument does not change ?

Or, is it treated as a variable in the global memory ?

See the ptx code to be sure about how this is handled…!

Sorry to ask but I’m learning CUDA and couldn’t find a way to see how that kind of variables are treated…can you tell me?

Arguments to global functions are passed via shared memory for devices with compute capability 1.x, and are passed via constant memory for device with compute capability 2.x. See section D.2.4.1 of the CUDA C Programming Guide. This is independent of any attributes applied to kernel arguments. The qualifier “const” in CUDA has the same semantics it has in C/C++, namely “read only”.

Whether the compiler can make any use of the additional information supplied by the use of the “const” qualifier or the “restrict” attribute for optimization purposes depends on numerous factors, and is not something that is architecturally defined. My personal philosophy is to supply the compiler with as much information as possible; even if it cannot take advantage of that information today, it may do so in the future.

I see little practical use in applying the “const” qualifier to scalar variables, its primary utility is for pointer arguments where the intention inside the function is to only ever read through a particular pointer. Many examples of this can be found among the function prototypes of various standard C/C++ functions, for example strcpy().

As an extension to standard C/C++, CUDA offers the constant qualifier. Section B.2.2 of the CUDA C Programming Guide describes its sematics.

Marking the argument as [font=“Courier New”]const[/font] is a hint to the compiler that you do not intend to change the argument inside the kernel and thus any (accidental) manipulation of its value leads to an error from the compiler.

It should not make a difference for optimization as the compiler relies on its own analysis for that.

[font=“Courier New”]nvcc -ptx[/font] if your question is how to see the PTX code.