Kernel Args vs. Constants

It’s not uncommon to have a GPU kernel where a number of the kernel parameters are left as constants throughout the entire call chain and, if the target were a CPU, would have been implemented as traditional static const file scope variables.

For this reason, you might think that exploiting the __constant__ qualifier and its associated cudaXXXSymbol() routines is a good way of reducing the number of kernel arguments and generally cleaning up your code.

CUDA __constant__ variables can definitely simplify your kernel code however you’re not only shifting some complexity toward module symbol initialization but you’re potentially closing off the opportunity to launch back-to-back grids with different __constant__ variable values.

Let’s look at what is actually being generated by the compiler and see if there is any major difference between kernel args and __constant__ variables at the SASS level. The expectation is that there is not (see Section D.2.5.2 in the CUDA C Programming Guide).

These two kernels are functionally the same but one uses parameters and the other constants:

They produce identical SASS code:

Since kernel parameters are constants on sm_20+ I expected both kernels to produce similar SASS. The only difference is that it appears constants are being pulled from different constant banks.

If you look at pre-Fermi output you’ll see the SASS is not the same because kernel parameters are passed via shared memory.

So back to the original question. Are module-scoped __constant__ variables useful given that strictly using kernel parameters produces similar code? Sure, you can write simpler code as long as you’re sure you’ll never need to update the __constant__ variables in a context. This very issue and a workaround was recently discussed in the CUDA Developer Forums.

But what I really want is the ability to optionally declare that a const kernel parameter has been raised to kernel scope (file scope?) as if it had been declared as a CUDA __constant__ variable.

It would be convenient and powerful to be able to write something like:

__global__ foo(__constant__ int* bar, int* baz) { ... }

This would indicate that the parameter bar should be raised to kernel scope. The benefit being that the developer wouldn’t have to invoke cudaXXXSymbol() initialization routines after loading the kernel module and before kernel launch.

An earlier file scope declaration that matches the parameter name and type might also be appropriate.

This syntax won’t (ever) be appearing in CUDA C99/C++ but I think it’s conceptually interesting for anyone writing a performance-focused DSL for GPUs to understand where there are mismatches between C99/C++ and CUDA.

Update: A compromise might be to create a structure that contains all of the constants you wish to raise to kernel scope and pass it (by value) as a parameter to the kernel. This way only a single reference needs to be passed throughout the kernel to access a number of launch-time CTA-specific constants.

The idiom would look something like this:

typedef struct { int* bar; int* baz; } KernelEnv;

__global__ foo(const KernelEnv kenv) {
    myxl(kenv, ...);
    plyx(kenv, ...);
}

The code snippet and PTX/SASS dumps are here.