Shared Memory Occupancy: Size the Reservation to What the Kernel Uses

This kernel stages each block's inputs in two regions of dynamic shared memory before adding them. The launch reserves 49,152 bytes per block, which caps how many blocks fit on an SM. Trace the shared-memory indices, derive the smallest safe reservation for the configured block, and remove the unused reservation without making either region out of bounds.

Languages: CUDA C++, Python, Rust

Lesson

Dynamic shared memory is reserved per block by the launch's sharedMemBytes value; the extern declaration does not allocate a fixed amount. Read the pointer offsets and the largest indexed element to determine the required byte range. Reserving more than the kernel can reach wastes a scarce per-SM resource, while reserving less makes a valid-looking index run beyond the allocation.

Any illustrated 50-block packing is shared-memory-only capacity, not simultaneous residency. Actual resident blocks = min(shared-memory capacity, register capacity, thread/warp capacity, hardware block limit). The excessive reservation is a teaching warm-up.

What you will learn

Read shared-memory occupancy as a theoretical launch-resource metric.
Work out how many bytes of dynamic shared memory a kernel actually accesses from its indexing.
See that under-reserving corrupts results (out-of-bounds) while over-reserving lowers occupancy - the reservation must match usage exactly.
Relate the dynamic shared-memory launch argument to the kernel's extern __shared__ usage.

Starter code

    // Stages this block's A and B values into two back-to-back regions of
    // dynamic shared memory. Derive the launch reservation from these offsets.
    extern __shared__ float s[];
    float* sA = s;
    float* sB = s + blockDim.x;

    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        sA[threadIdx.x] = A[i];
        sB[threadIdx.x] = B[i];
        __syncthreads();
        C[i] = sA[threadIdx.x] + sB[threadIdx.x];
    }