Kernels are highlighted by the __global__ keyword in .cu
files. When using the CUDA runtime, they can be invoked in-line with the
triple-angle-bracket <<< >>> syntax. Chapter 7 gives a
detailed description of how kernels can be invoked and how they execute
on the GPU.
The GPU executable code of the module comes in the form of kernels
that are invoked with the language integration features of the CUDA
runtime (<<< >>> syntax), or the cuLaunchKernel()
function in the driver API. Historically, loading a module copied
all of its kernels into device memory at once, whether or not
they were ever launched – a cost that grew painful as libraries such as
cuBLAS and cuDNN came to ship modules holding thousands of specialized
kernels. Lazy loading, the default since CUDA 12.2, defers each
kernel and its data until first use: a kernel is copied into device
memory when it is first launched rather than when its module loads,
which cuts both the startup time and the resident-memory footprint of a
program that touches only a fraction of a large module. The
CUDA_MODULE_LOADING environment variable selects the policy
– LAZY or EAGER – and the runtime defaults to
lazy wherever the driver supports it.
Once a module is loaded, kernels may be queried with
cuModuleGetFunction(); the kernel’s attributes can be queried with
cuFuncGetAttribute(); and the kernel may be launched with
cuLaunchKernel().
cuLaunchKernel() rendered a whole slew of API entry points obsolete:
functions such as cuFuncSetBlockShape() specified the block size to use
the next time a given kernel was launched; functions such as
cuParamSetv() specified the parameters to pass the next time a given
kernel was launched; and cuLaunch(), cuLaunchGrid(), and
cuLaunchGridAsync() launched a kernel using the previously-set state.
These APIs were inefficient because it took so many calls to set up a
kernel launch, and because parameters such as block size are best
specified atomically with the request to launch the kernel.
The cuFuncGetAttribute() function may be used to query specific
attributes of a function, such as:
the maximum number of threads per block,
the amount of statically-allocated shared memory,
the size of user-allocated constant memory,
the amount of local memory used by each function,
the number of registers used by each thread of the function,
the virtual (PTX) and binary architecture versions for which the function was compiled.
When using the driver API, it is usually a good idea to use extern
"C" to inhibit the default name-mangling behavior of C++. Otherwise, you
have to specify the mangled name to cuModuleGetFunction().
As executables that were built with the CUDA runtime are loaded, they create global data structures in host memory that describe the CUDA resources to be allocated when a CUDA device is created. Once a CUDA device is initialized, these globals are used to create the CUDA resources all at once. Because these globals are shared process-wide by the CUDA runtime, it is not possible to incrementally load and unload CUDA modules using the CUDA runtime.
Because of the way the CUDA runtime is integrated with the C++
language, kernels and symbols should be specified by name (i.e., not
with a string literal) to API functions such as cudaFuncGetAttributes()
and cudaMemcpyToSymbol().
In Fermi-class architectures, the Streaming Multiprocessors have L1
caches that can be split as 16K shared / 48K L1 cache, or 48K shared /
16K L1 cache8. Initially, CUDA allowed the cache
configuration to be specified on a per-kernel basis, using
cudaFuncSetCacheConfig() in the CUDA runtime or cuFuncSetCacheConfig()
in the driver API. Later, this state was moved to be more global:
cuCtxSetCacheConfig()/cudaDeviceSetCacheConfig() specifies the default
cache configuration.