Your first kernel: the CPU tells the GPU to work
How host and device split the work, what __global__ means, and why a kernel launch is asynchronous
- Tell host (CPU) code from device (GPU) code, and know where each one runs
- Know the difference between the `__global__`, `__device__`, and `__host__` function qualifiers
- Understand that a kernel launch is asynchronous, and why you have to synchronize
- Compile and run your first CUDA program with nvcc
Before you write CUDA, get this straight: you write one source file, and it compiles into two binaries. One runs on the CPU (host): it allocates memory, moves data, and issues commands. The other runs on the GPU (device): it does the actual parallel work. The main job of nvcc is to split those two parts out of the same .cu file and hand them to the host compiler and the GPU backend.
Host (CPU) Device (GPU)
────────── ────────────
cudaMalloc ──── allocate device memory ────▶ [ global memory ]
cudaMemcpy H2D ──── copy inputs ───────────────▶ [ a, b ]
kernel<<<G,B>>> ──── submit work ───────────────▶ [ SM0 SM1 SM2 ... ]
│ │
(returns immediately, CPU keeps going) (kernel runs in parallel)
│ │
cudaDeviceSynchronize ◀──── wait until done ─────────┘
cudaMemcpy D2H ◀──── copy results back ──── [ c ]
cudaFree ──── free device memory ────▶Three function qualifiers
| Qualifier | Where it runs | Who can call it | Typical use |
|---|---|---|---|
__global__ | GPU | host (device can also call it via dynamic parallelism) | kernel entry point, must return void |
__device__ | GPU | GPU code only | helper functions inside a kernel |
__host__ | CPU | CPU | ordinary function, this is the default |
__host__ __device__ | compiled once for each side | callable from both | shared small helpers such as clamp |
1#include <cstdio>2 3// __global__ = runs on the GPU, launched from the CPU.4// A kernel must return void. Results come back only through device memory.5__global__ void hello() {6 int tid = blockIdx.x * blockDim.x + threadIdx.x;7 printf("block %d / thread %d -> global id %d\n",8 blockIdx.x, threadIdx.x, tid);9}10 11int main() {12 hello<<<2, 4>>>(); // 2 blocks, 4 threads each, 8 threads total13 14 // Kernel launch is asynchronous: the CPU has not waited for the GPU at this line.15 // Without this sync, main returns, the process exits, and you may see no output at all.16 cudaDeviceSynchronize();17 18 cudaError_t err = cudaGetLastError();19 if (err != cudaSuccess) {20 fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err));21 return 1;22 }23 return 0;24}What goes inside <<<>>>
<<<gridDim, blockDim>>> is CUDA-only syntax, the execution configuration. The first argument is how many blocks are in the grid. The second is how many threads are in each block. Both can be a 3D dim3. Written as an integer, they use only the .x dimension.
1kernel<<<2, 4>>>(); // grid=(2,1,1), block=(4,1,1)2 3dim3 grid(16, 16); // grid=(16,16,1)4dim3 block(32, 8); // block=(32,8,1) = 256 threads5kernel<<<grid, block>>>();6 7// The full form has two optional arguments:8// third: extra dynamic shared memory bytes per block9// fourth: which stream this launch goes into10kernel<<<grid, block, 48 * 1024, stream>>>();Compile and run
1# Check the driver and the GPU first2nvidia-smi3 4# -arch sets the target architecture. Match it to your GPU, or you get JIT compilation or a binary that will not run5# sm_70 Volta / V100 sm_80 Ampere / A1006# sm_75 Turing / T4 sm_86 Ampere / RTX 30xx7# sm_89 Ada / RTX 40xx sm_90 Hopper / H1008nvcc -O3 -arch=sm_80 -o hello 01_hello.cu9 10./helloThe print order can change from run to run. Threads in block 1 may print before block 0. That is not a bug. It is one of the most important CUDA rules: there is no execution-order guarantee between blocks. The GPU can schedule them in any order, on any SM, at any concurrency. Your kernel has to stay correct under that rule.
Self-check
Answer in your head first, then open the solution. If you cannot, reread that section.