What is lazy module loading in CUDA?
The driver deferring each kernel's load until its first launch, which moves a cost most benchmarks then measure by accident.
Your binary carries device code for every kernel you compiled, and the driver used to push all of it onto the card when the context came up. Now it does not. "Lazy loading reduces program initialization time by waiting to load CUDA modules until they are needed", and the version table in that page reads "12.2 Lazy loading enabled by default for Linux" and "12.3 ... Now enabled by default for Windows". Call it 12.2 on Linux and 12.3 everywhere. For a program that links cuBLAS or cuDNN and launches four of the thousands of kernels inside them, that is a large saving in startup time and in device memory, which is why it became the default rather than a flag.
The bill arrives per kernel, and that is the part that breaks benchmarks. "Do a warm-up run first" is a rule people learned when a program paid once, at startup. Under lazy loading a warm-up warms the kernels it launches and nothing else, so time three kernels after warming one and two of your numbers still have a module load inside them. It shows up as a first iteration that looks several times slower than the rest, gets read as cache warming or clock ramping, and is neither.
Be precise about what a CUDA event pair around a cold launch actually catches, because the load itself is host-side work inside the launch call and events do not see host code. What they see is the hole it leaves in the device timeline: the start event completes, the GPU sits idle while the host loads the module, and only then does the kernel run. The elapsed time between the two events therefore contains a gap that is not compute, which is exactly why the number is wrong to quote and exactly why it is visible at all.
CUDA_MODULE_LOADING=EAGER puts the old behaviour back. It does not make anything free, it moves the cost out of the first launches and into context creation, so the initialization line grows while the cold launches fall towards their warm values. Do not confuse any of this with JIT compilation, which has the same symptom and a different cause: JIT is the driver compiling embedded PTX because your binary holds no SASS for the card in front of it. Both make launch one expensive. One is a load, one is a compile, and -gencode fixes only the second.
Measured
Two kernels, two cold launches, one process. Day 9 timed vectorAdd over 16,777,827 floats with events, then did the same for a second kernel it had never launched. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), nvcc -O3 -arch=sm_75, captured 2026-08-30. The toolkit is past 12.2 on Linux, so lazy loading was the default and no environment variable was set.
| Kernel | first launch | after warm-up |
|---|---|---|
vectorAdd |
1.029 ms | 0.786 ms |
vectorScale |
0.657 ms | 0.548 ms |
The second row is the whole point. vectorScale ran after vectorAdd had been launched, warmed and timed repeatedly, in a process with a fully built context, and it still paid to go first. The load is per module, not per program, and a warm-up loop that only touches your first kernel leaves the rest of them cold.
For scale, the initialization this feature was built to shrink is still the largest single number in that run: cudaSetDevice(0), the first CUDA call in the process, took 255.966 ms.
One gap, stated rather than filled. This lesson has not yet been run with CUDA_MODULE_LOADING=EAGER, so there is no measured EAGER column here. It is day 9's exercise, and the prediction to check is that the cold rows fall while the 255.966 ms rises.
Diagram
timeline-host-device, preset whole-process-vs-events. The device lane shows the start event, then a shaded idle gap labelled "module load", then the kernel bar. A second launch of the same kernel below it has no gap. A third bar, a different kernel launched later, has its own gap.
Alt text: "The first launch of each kernel leaves an idle gap between the start event and the kernel, and a second kernel in the same process gets its own gap."
Code
From code/day09-timing/timing.cu. The warm-up lives inside the timing helper on purpose, because a warm-up at the call site is the one that goes missing from the fourth kernel you add six months later.
// Warm up this kernel, not just the first kernel in the program. Lazy
// module loading has been the default since CUDA 12.2 on Linux, so the
// first launch of each kernel pays its own load.
for (int i = 0; i < kWarmupRuns; ++i) {
launch();
}
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaGetLastError());
Related terms
Where you meet this
- Day 9, why your GPU code looks slower than your CPU, the lesson that owns this term and the one that measures it.
- Day 1, your first CUDA kernel, the launch this cost hides behind.
- Day 45, occupancy, where every configuration you compare needs its own warm-up.
- /setup/check-cuda-version, because whether lazy is your default is a question about which toolkit you have.
CUDA driver version is insufficient, the other way a version mismatch turns up at run time rather than at compile time.
Sources
- CUDA Programming Guide, "Lazy Loading", for the wording and the 12.2 / 12.3 version table: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/lazy-loading.html (checked 2026-08-30)
- CUDA Programming Guide, environment variables, for
CUDA_MODULE_LOADING: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/environment-variables.html (checked 2026-08-30) - "first kernel launch takes 7x longer than subsequent launches", answered with "Kernel binaries are lazy loaded on their first invocation": https://www.reddit.com/r/CUDA/comments/1n6rk27/first_kernel_launch_takes_7x_longer_than/ (checked 2026-08-29)
Byline
Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. This entry stays a draft until a named author and a different named reviewer sign it.