Block Occupancy: Give Each Block More Warps

This vector-add kernel is correct, but the launch configuration starts with Block X = 32. That gives each block only one warp, so the hardware's resident-block-per-SM limit caps the number of active warps per SM even though warp capacity itself is nowhere near full. Increase the block's x dimension there and use GPUFlight's block occupancy metric to verify the change.

Languages: CUDA C++, Python, Rust

Lesson

Occupancy is ultimately about keeping enough warps resident on each SM, but a resident block is capped by SEVERAL independent hardware limits at once: registers, shared memory, threads, and a flat max-blocks-per-SM count. A one-warp block can hit that flat block-count ceiling long before the warp-count ceiling ever binds - and when that happens, the warp occupancy metric alone reads 100% and hides the problem, because it only models the warp-count ceiling. Block occupancy models the block-count ceiling directly.

The numeric SM limits in this example assume Ada/L4. Actual residency is bounded by the minimum of shared memory, registers, threads/warps and the hardware block limit.

What you will learn

Read block occupancy as a theoretical launch-resource metric distinct from warp occupancy.
See why one-warp blocks can under-fill an SM even when warp occupancy itself reads 100%.
Tune the block's x dimension in the launch configuration while keeping the grid's x dimension derived from N so every tested input size remains correct.

Starter code

    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    for (; i < N; i += stride) {
        C[i] = A[i] + B[i];
    }