Occupancy Warm-Up: Vector Add
This vector-add kernel is already correct, but its launch starts with one thread per block. In the launch configuration below the kernel, increase the block's x dimension to give each block more useful warps and use GPUFlight metrics to reach the occupancy target.
Sign in to run this lab. GPU availability and credits are shown in the app.
Languages: CUDA C++, Python, Rust
Lesson
A CUDA launch can be functionally correct and still use the GPU poorly. This exercise teaches the difference between the static occupancy limit reported from launch resources and achieved occupancy measured while the kernel actually runs.
The one-thread block is deliberately inefficient teaching scaffolding. Occupancy diagnoses resource use; elapsed time is the optimization objective.
What you will learn
- Read theoretical occupancy and achieved occupancy as different signals.
- See why 1-thread blocks waste most lanes in every scheduled warp.
- Use the launch configuration to choose a multi-warp block and a grid expression that adapts to different N values.
Starter code
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) {
C[i] = A[i] + B[i];
}