Day 9Module 1

in-technical-review

Why your GPU code looks slower than your CPU

Someone on r/CUDA, a few days into their first kernel:

"My CPU completed the task in 0.009632 ms, whereas my GPU took 200.466284 ms. I don't understand what I'm doing wrong."

https://www.reddit.com/r/CUDA/comments/1iribin/cpu_outperforming_gpu_consistently/ (checked 2026-08-29)

The measurement included more than the kernel. A clock around a CUDA program includes process start, context creation, three allocations, two PCIe transfers, module loading, and the arithmetic.

This page times one vector add five ways. Each measurement answers a different question. You will learn which number belongs in a bug report and when a CPU is faster for the full task.

What the clock around your program is holding

A kernel launch does not wait.

"Kernel launches are asynchronous with respect to the host thread. That is, the kernel will be setup for execution on the GPU, but the host code will not wait for the kernel to complete (or even start) executing on the GPU before proceeding."

https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html (checked 2026-08-30)

A host clock that starts before <<<>>> and stops on the next line measures the cost of queueing work. This launch overhead is separate from kernel run time.

Add cudaDeviceSynchronize() before stopping a host timer. The sync waits for the kernel but also removes host-device overlap from the measured region.

The first CUDA call also creates a context. As of CUDA 12.0, cudaSetDevice "initialize[s] the runtime and the primary context associated with the specified device".

The guide adds: "This is important when timing runtime function calls and when interpreting the error code from the first call into the runtime." (Same page, checked 2026-08-30.) A whole-process clock includes that work before the arithmetic starts.

The driver may load each kernel's module on its first launch instead of at program start. "Lazy loading reduces program initialization time by waiting to load CUDA modules until they are needed."

The version table lists lazy loading as enabled by default in CUDA 12.2 on Linux and 12.3 on Windows (https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/lazy-loading.html , checked 2026-08-30).

Lazy module loading means a warm-up only prepares the kernels it launches. If you warm one kernel and time three, the other two may still pay their first-load cost.

cudaMemcpy adds another cost. It moves inputs to the GPU and results back across PCIe, not through the GPU's local memory bus.

Note. The widget replays a real Nsight Systems trace of this page's program rather than simulating one, so it ships when the trace does. Until then it renders a labelled fixture.

The intuition that makes it worse

GPUs and CPUs have different costs. Moving a loop to a GPU does not ensure that the full task runs faster. Vector addition makes this clear because it does little arithmetic per byte moved.

Each element reads two floats, writes one float, and performs one add. That is twelve bytes per arithmetic operation, which is low arithmetic intensity.

The kernel is memory-bound, so memory bandwidth limits its speed. Block size and occupancy cannot remove that byte count.

The full task also uses cudaMemcpy to move data over PCIe. Even a zero-time kernel would still pay for those transfers.

For small inputs, transfer cost can exceed the kernel time, so the CPU wins. GPU programs avoid repeated transfers by keeping data on the device across several kernels.

You can test whether a timer measures the kernel by changing the launch. In the NVIDIA forum thread used here, changing add<<<numBlocks, blockSize>>>(N, x, y); to add<<<1,1>>>(N, x, y); did not change the reported time (https://forums.developer.nvidia.com/t/cuda-slower-than-cpu/263339 , checked 2026-08-30).

add<<<numBlocks, blockSize>>>(N, x, y);  // thousands of threads
add<<<1, 1>>>(N, x, y);                  // one thread, one lane

One thread takes far longer than thousands of threads for this kernel. If the measured time does not change, the timer does not cover the kernel.

<<<1, 1>>> is a bad launch configuration: it uses only one lane of one warp on one SM, leaving the rest of the GPU idle.

Five stopwatches over one kernel

The full program is in code/day09-timing/timing.cu. It times one vectorAdd over 16,777,827 floats in five ways, then compares four smaller sizes with a CPU reference.

Three rules hold it together.

Every row uses the same kernel and data. Only the timed region changes, so the difference between rows comes from the work one region adds.

The order is part of the measurement. Row 3 must be the first vectorAdd launch because a cold launch happens once. An earlier launch would warm the module and make row 3 match row 4.

The program uses one host clock on purpose. The course normally bans host clocks in .cu files because they require a sync to measure kernel run time. This lesson uses clock_gettime(CLOCK_MONOTONIC) to show both the launch-only and launch-plus-sync values.

The hostMs() helper produces row 2:

    const double tLaunch0 = hostMs();
    vectorAdd<<<blocks, kThreadsPerBlock>>>(d_a, d_b, d_out, kElems);
    const double tLaunch1 = hostMs();
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());
    const double launchOnlyMs = tLaunch1 - tLaunch0;
    const double launchPlusSyncMs = hostMs() - tLaunch0;

Row 3 is a CUDA event on each side of a single launch, with the warm-up deliberately missing:

    CUDA_CHECK(cudaEventRecord(start));
    vectorAdd<<<blocks, kThreadsPerBlock>>>(d_a, d_b, d_out, kElems);
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));
    CUDA_CHECK(cudaGetLastError());
    float coldMs = 0.0f;
    CUDA_CHECK(cudaEventElapsedTime(&coldMs, start, stop));

The events do not measure module loading on the host. They measure its effect on the device timeline: the start event completes, the GPU waits while the host loads the module, then the kernel runs.

Rows 4 and 5 use timeKernel, as day 5 used CUDA_CHECK. Copy it without changes. It creates events, warms the kernel, runs it ten times, and returns the mean:

template <typename LaunchFn>
static float timeKernel(LaunchFn launch) {
    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));

    // 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());

    CUDA_CHECK(cudaEventRecord(start));
    for (int i = 0; i < kTimedRuns; ++i) {
        launch();
    }
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));
    CUDA_CHECK(cudaGetLastError());

    float ms = 0.0f;
    CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));
    return ms / kTimedRuns;
}

The helper contains the warm-up so each timed kernel gets one.

Note. The copies in row 5 are pageable, not pinned, because that is the default a first program gets. Pinned copies are faster and day 53 measures the difference. Nothing here is the best the hardware can do, it is what the code in front of you produces.

Results

Re-verified on a Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88) on 2026-09-02. The original CUDA 12.6 timing remains in the page's evidence array. Absolute timings moved, while every qualitative conclusion and ordering held; the block below is the CUDA 13.0 run.

GPU: Tesla T4 (compute capability 7.5)
n = 16777827 floats, 64.0 MiB per buffer, 256 threads/block, 65539 blocks

cudaSetDevice(0), the first CUDA call in this process:   424.730 ms
No arithmetic of yours has run yet. That is the context.

Five ways to time one vector add
  1  wall clock, whole process                  1594.357 ms
  2  wall clock, around the launch                 0.005 ms
     the same clock, plus a sync                   0.777 ms
  3  events, kernel, first launch                  0.940 ms
  4  events, kernel, after warm-up                 0.768 ms
  5  events, copies + kernel + copy back          45.362 ms
     CPU, same n, one thread, no SIMD             18.909 ms

Rows 4 and 5 are means of 10 runs after 3 warm-ups.
Rows 1, 2 and 3 are single measurements, because each one is about
something that happens once.

A second kernel is cold too
     vectorScale, first launch                     0.618 ms
     vectorScale, after warm-up                    0.550 ms

Where the CPU wins
           n     CPU (ms)  kernel (ms) round trip (ms)
        1024        0.002        0.004          0.028
       65536        0.067        0.005          0.255
     1048576        1.158        0.052          2.930
     4194304        4.326        0.195         11.331

Row 4 is the kernel. Row 5 is what a user waits for. Row 1 is what
`time ./timing` prints, and it is never the answer to how fast a
kernel is.

The five measurements span a factor of ninety thousand. The first includes the whole process, while the last includes one round trip.

Row 1 is 2,076 times the kernel time. time ./timing reports 1594.357 ms, while the kernel takes 0.768 ms. Row 1 includes process start, the dynamic linker, libcudart loading, and 424.730 ms in the first CUDA call before the program's arithmetic starts.

Row 2 measures launch overhead. A host clock around the launch reads 0.005 ms because the call returns after it queues the work. Add a sync and the same clock reads 0.777 ms, within 1.2 percent of the event timing.

Rows 3 and 4 show the first-launch cost. The first launch takes 0.940 ms, while the warm mean is 0.768 ms. vectorScale also changes from 0.618 ms cold to 0.550 ms warm, so warming one kernel does not warm every kernel.

Row 5 measures the full GPU task. The two input copies, kernel, and output copy take 45.362 ms. One CPU thread without SIMD takes 18.909 ms for the same arithmetic.

In this run, the GPU kernel is 25 times faster than the CPU loop, but the full GPU task is 2.4 times slower. About 44.6 of the task's 45.4 milliseconds lies outside the kernel, mostly in PCIe transfers.

A faster kernel cannot remove that transfer time. Move less data or keep it on the device across several kernels.

At 1024 elements, the CPU wins. As n grows, the GPU kernel becomes faster than the CPU loop before the full GPU round trip does. Between those points, the kernel is faster but the full GPU task is slower.

Run it yourself

Run the program on a CUDA GPU. The build line is in the repo's README:

nvcc -std=c++17 -O3 -arch=sm_75 -o timing timing.cu

This page has no Compiler Explorer example. The program allocates three 64 MiB device buffers and runs more than a hundred timed launches, which may exceed the 20-second run limit.

If you have no GPU, read /setup/learn-cuda-without-a-gpu.

Exercise

Run the program, then write one sentence per row naming what that row contains and the row below it does not. Then run CUDA_MODULE_LOADING=EAGER ./timing and say where the cost went.

Time: 25 to 40 minutes. Submit: one sentence per row, one on where the EAGER cost went, and your row 1 divided by row 4.

Check: the program uses branches that return EXIT_FAILURE instead of an assert, because Release builds remove an assert under NDEBUG. It compares all 16,777,827 GPU results with the CPU reference.

It also checks that row 5 is at least row 4 and row 1 is at least row 5 because each region contains the next. A failure means the timing regions are broken. Expect row 1 divided by row 4 to be in the hundreds or thousands.

Hint 1

Each row is the row below it plus one named thing. Name the thing, not the difference. What does row 3 do that row 4 does not, and how often can it do it?

Hint 2

For the EAGER run: the total work has not changed, so if one number falls another must rise. Which call is the only candidate, and why does the second kernel's cold line move too?

Solution

Row 1 includes process start, context creation, allocations, host input setup, and every timed run. Row 5 covers one round trip, while row 4 covers only the warm kernel.

Row 3 adds the device wait caused by first-time module loading. Row 2 measures only host-side queueing, so adding a sync changes it by orders of magnitude.

The EAGER run makes nothing free. It moves the module load out of the first launch and into cudaSetDevice, so the context line grows while row 3 and vectorScale's cold line fall towards their warm values. The second kernel moves for the reason it exists: the load is per module, not per process.

A timing describes a timed region. State where the timer starts and stops.

Pitfalls

You quote time ./a.out in a bug report. It contains process start, context creation, allocation and both copies, and on a small problem the kernel is the smallest term in it. Time the kernel with events and report both, because the difference is the finding.

A host clock around the launch, with no sync. The launch returns as soon as the work is queued, so you timed the queue. If the number does not move when you change the launch configuration, this is what happened.

You warm up one kernel and time four. Lazy loading is per module, so the first launch of each kernel pays its own load. Warm up every kernel you time, which is what timeKernel does for you.

You compare the wrong pair. Kernel against CPU flatters the GPU by hiding the copies. Round trip against CPU is honest for a one-shot job and pessimistic for a pipeline that keeps its data resident. Say which you measured; "is the GPU worth it" often answers differently for the two.

You launch with <<<1, 1>>> and blame the hardware. One thread uses one lane of one warp on one SM. It is a real bug and it is invisible to a whole-process clock, which is why the two failures turn up in the same forum posts.

You conclude the GPU is the wrong tool from one small-n measurement. At that size it is. The crossover is real and it is not a code problem, but it belongs to this kernel and no other. Day 30 is where you learn to predict it for a kernel you have not written yet.

Go deeper

CUDA C++ Best Practices Guide, "Using CPU Timers" and "Using CUDA GPU Timers", for the sync rule and the half-microsecond event resolution: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html (checked 2026-08-30)
CUDA Programming Guide, "Lazy Loading", for the version table and CUDA_MODULE_LOADING: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/lazy-loading.html (checked 2026-08-30)
cuda-samples, cpp/0_Introduction/asyncAPI, NVIDIA's own event-timing example: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Introduction/asyncAPI (checked 2026-08-30)
Programming Massively Parallel Processors, 4th edition, chapter 6, on performance considerations: https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0 (checked 2026-08-30)

Day 10 sweeps block sizes, now that you can measure one without fooling yourself. Day 11 takes row 4 and finds out where it went, which for a memory-bound kernel is the only question left. Both use timeKernel unchanged, and Nsight Systems arrives on day 41 to draw the timeline this page describes in words.