What is a CUDA kernel?

A function you mark __global__ and launch across a grid of threads, which runs on the GPU while the CPU keeps going.

Two pieces of syntax make one. __global__ in front of the function is one of the execution space specifiers, and it says the host launches this function while the device runs it. The execution configuration in triple chevrons after the name says how many copies to start. printThreadIds<<<2, 4>>>() starts eight copies of one function body, two blocks of four threads, and each copy reads its own blockIdx and threadIdx to work out which part of the job it owns. The grid is all of those blocks together.

A kernel is not a function you call. A __global__ function must return void, cannot be a class member and cannot recurse (CUDA Programming Guide 5.4.1.1). The triple chevron is an enqueue, not a call: the same guide says these calls "return to the host thread before the device completes execution". Control reaches your next line while the GPU may not have started, so there is nowhere for a return value to go and no error to hand back at the launch line. That is why every launch on this site is followed by cudaGetLastError() and then a synchronize, which is the shape error checking takes here.

The thing people get wrong is treating the launch as the run. A hello world whose kernel calls printf exits with code 0 and an empty terminal, and the usual conclusion is a broken install. Nothing is broken. Device printf writes into a buffer that the runtime copies to stdout at a listed set of moments, and process exit is not one of them: "the buffer is not automatically flushed when the program exits" (guide 5.3.7.2). Add cudaDeviceSynchronize() and the lines arrive. Exit code 0 is not evidence that your kernel ran.

Measured

On a Tesla T4 (driver 595.84, CUDA 12.6, built with nvcc -O3 -arch=sm_75), day 1 launched printThreadIds<<<2, 4>>>() and the host's launch returned line came out before any device line, with the sync returned line after all eight of them. That ordering is the launch being asynchronous, printed.

Order among the device lines is not stable. Across 20 runs on one card block 0 printed first 16 times; 20 runs of the same binary on a second T4 of the same model gave 18. Two cards, two ratios, which is a better argument for "not guaranteed" than either number alone. Blocks are scheduled independently and what you read is buffer order, so never write a check that depends on the sequence. Captured 2026-08-30; the transcript and the tally are in code/day01-first-kernel/evidence/run-2026-08-30.txt.

Day 1 puts no number on launch-to-return time. Timing a launch takes events and a warm-up run, which day 9 sets up.

Code

Two slices of code/day01-first-kernel/hello.cu: the kernel, and the three lines around its launch.

__global__ void printThreadIds() {
    const size_t t = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    printf("device: block %u, thread %u, global thread %llu\n", blockIdx.x,
           threadIdx.x, static_cast<unsigned long long>(t));
}
// ...
    printThreadIds<<<kBlocks, kThreadsPerBlock>>>();
    CUDA_CHECK(cudaGetLastError());
    std::printf("host: launch returned, kernel may not have run yet\n");
    CUDA_CHECK(cudaDeviceSynchronize());
    std::printf("host: sync returned, every device line is above this\n");

printf inside the kernel is the device printf, whose size fields are h, l and ll, so a size_t is cast to unsigned long long first. %zu prints something strange.

Related terms

Where you meet this

Day 1, your first CUDA kernel, which owns the measurement above
Day 5, vector addition, the first kernel that owns a piece of an array
Day 9, why your GPU looks slower than your CPU, where launches get timed
invalid configuration argument, the error a launch refuses with

Sources

CUDA Programming Guide, C++ language extensions, section 5.4.1.1 on __global__ and the asynchrony of launches: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html (checked 2026-08-29)
CUDA Programming Guide, C++ language support, section 5.3.7.2 on the printf buffer and its flush points: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-support.html (checked 2026-08-29)
CUDA Programming Guide 2.1, "Intro to CUDA C++", on synchronizing CPU and GPU: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html (checked 2026-08-29)

Byline

Written by: unassigned. Reviewed by: unassigned. This entry is a draft and cannot publish until both are named people, and two different ones. Written on: not set. Last checked: not set. Numbers captured 2026-08-30 on the project's verification node.