Hello CUDA Example

The traditional first kernel — every GPU thread prints its coordinates

This is the smallest CUDA program that proves your toolchain works end-to-end: the driver loads the kernel, the runtime launches threads, the threads execute device code, and cudaDeviceSynchronize flushes their output back to the CPU stdout.

Read the source carefully. Every line corresponds to one of the three CUDA extensions you saw in lesson 1:

__global__ void say_hello() — function qualifier marking this as a kernel.
say_hello<<<1, 4>>>() — triple-angle launch: 1 block, 4 threads.
blockIdx.x, threadIdx.x — SIMT built-ins giving each thread its identity.

The launch syntax <<<grid, block>>> is one of the few CUDA constructs that looks like template syntax but isn't. NVCC parses it as a special launch expression. After NVCC's pre-pass, what reaches your normal C++ compiler is a regular function call into the runtime API with grid/block dimensions packed as arguments.

Code

hello.cu — the smallest CUDA program·cuda

#include <cstdio>
#include <cuda_runtime.h>

__global__ void say_hello() {
    printf("Hello from block %d, thread %d\n",
           blockIdx.x, threadIdx.x);
}

int main() {
    say_hello<<<1, 4>>>();           // 1 block, 4 threads
    cudaError_t err = cudaDeviceSynchronize();
    if (err != cudaSuccess) {
        fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err));
        return 1;
    }
    return 0;
}

Build and run·bash

# Linux / WSL — replace sm_89 with your compute capability
nvcc -arch=sm_89 hello.cu -o hello
./hello
# Hello from block 0, thread 0
# Hello from block 0, thread 1
# Hello from block 0, thread 2
# Hello from block 0, thread 3

# Windows (Developer Command Prompt for VS 2022)
nvcc -arch=sm_89 hello.cu -o hello.exe
hello.exe

Exercise

Save hello.cu, compile, and run. Then change the launch from <<<1, 4>>> to <<<2, 8>>> (2 blocks, 8 threads each = 16 prints). Recompile, run, and notice that the lines may not be in the order (block 0 thread 0, block 0 thread 1, ..., block 1 thread 7) you'd expect. Read the blockIdx.x + threadIdx.x values to confirm the GPU scheduler interleaved the blocks. This is your first taste of why GPU debugging requires fences and prints with thread IDs.

The traditional first kernel — every GPU thread prints its coordinates

Code

External links

Exercise

Progress

Comments 0