Your first CUDA kernel

A CUDA kernel is a function that runs on the GPU. In this lesson, you will write one, launch eight threads, and explain why the program needs cudaDeviceSynchronize().

You need basic C++ syntax and a working CUDA compiler. If either is new, start with Day 0.

The one idea to learn

The CPU does not wait when it launches a kernel. It places the work in a queue, then runs the next host statement.

This matters because device printf does not write straight to your terminal. It writes to a buffer that CUDA flushes at a synchronization point, but not at process exit.

Read the launch

This kernel runs once per GPU thread:

__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));
}

__global__ marks printThreadIds as a kernel. The CPU can launch it, but the GPU runs its body.

The launch below starts two blocks with four threads in each block:

printThreadIds<<<2, 4>>>();

CUDA calls <<<2, 4>>> the execution configuration. It creates eight threads, and every thread runs printThreadIds once.

Check your reading

For printThreadIds<<<3, 2>>>(), how many lines should the kernel produce? What global thread id does block 2, thread 1 get?

Follow the host code

The lesson program checks the launch, prints a host line, waits for the GPU, then prints a second host line.

int main() {
    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");
    return EXIT_SUCCESS;
}

Read it in order. The launch queues work, cudaGetLastError() checks whether CUDA accepted the launch, and the first host line prints without waiting for the kernel.

cudaDeviceSynchronize() stops the host until earlier device work finishes. It also reports errors that occurred while the kernel ran.

The second host line can print only after all eight threads finish. This gives you a clear boundary between queued work and completed work.

Build and run it

The full source is in code/day01-first-kernel/hello.cu.

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

Do not compare the order of the eight device lines. CUDA does not promise that blocks will run or print in block-number order.

Read one verified run

This transcript comes from the evidence file listed in the page metadata:

host: launch returned, kernel may not have run yet
device: block 0, thread 0, global thread 0
device: block 0, thread 1, global thread 1
device: block 0, thread 2, global thread 2
device: block 0, thread 3, global thread 3
device: block 1, thread 0, global thread 4
device: block 1, thread 1, global thread 5
device: block 1, thread 2, global thread 6
device: block 1, thread 3, global thread 7
host: sync returned, every device line is above this

The first host line appears before the device output because the launch returned at once. The last host line appears after the device output because the synchronization call waited.

The run used a Tesla T4, driver 580.173.02, and CUDA 13.0 on 2026-09-02. Those details record the test, but the launch and synchronization rules apply to CUDA programs on other supported GPUs.

Remove the wait

Delete this line, rebuild, and run the program again:

CUDA_CHECK(cudaDeviceSynchronize());

The program can still return exit code 0 with no device lines. CUDA did not flush the device printf buffer before the process ended.

Exit code 0 proves only that the host process returned success. It does not prove that the kernel finished or that you checked its run-time errors.

Practice

Work in three steps. Run the program after each change.

1. Change the launch to three blocks of two threads. Predict every global id before you run it.
2. Print only even global ids. Keep all six threads in the launch and put the test inside the kernel.
3. Remove cudaDeviceSynchronize(). Write two sentences that explain the missing device output.

The checker treats device output as a set of lines because print order can change. It expects ids 0 through 5 in step 1, ids 0, 2, and 4 in step 2, and no device lines in step 3.

if (t % 2 == 0) {
    printf("device: block %u, thread %u, global thread %llu\n", blockIdx.x,
           threadIdx.x, static_cast<unsigned long long>(t));
}

Common mistakes

If the kernel prints nothing, check for a synchronization call after the launch. Reinstalling CUDA will not fix a program that exits before its buffered output is flushed.

If %zu prints the wrong value in device code, cast size_t to unsigned long long and use %llu. Device printf supports a smaller set of size fields than the host version.

If cudaGetLastError() reports invalid configuration argument, check the block size and grid dimensions. CUDA rejects a block with more threads than the device allows.

Sources

• CUDA C++ basics: synchronizing the CPU and GPU
• CUDA C++ language support: device printf
• CUDA Samples: simplePrintf

Next

Day 2 explains how the GPU groups these threads into warps and assigns blocks to streaming multiprocessors.