Launch configuration in practice

01 · Awkward sizes are the normal case

Nobody's problem is a power of two

Day 13's example launched exactly 1,048,576 threads because it was built to. Real work arrives as 1,000,000 elements, or 999,983, or a 1023×767 image. The formula survives: blocks = ceil(n / threads) — in CUDA idiom, (n + threads - 1) / threads — and day 13's if (i < n) guard absorbs the overshoot. Overshooting by at most one block's worth of threads, and no-oping the extras, is the idiom.

The same trick extends to multiple dimensions, which is where grids stop being arrays and start being pictures. A grid is 3D — gridDim.x/y/z — and for images or matrices you let x and y carry the two axes:

// launch: dim3 grid(W/16, H/16), dim3 block(16, 16)
__global__ void blur(const float* img, float* out,
                        int w, int h){
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x >= w || y >= h) return;

    float s = 0.0f;
    // sum the 3×3 neighborhood — row-major means y*w + x
    for (int dy = -1; dy <= 1; ++dy)
        for (int dx = -1; dx <= 1; ++dx)
            s += img[(y + dy) * w + (x + dx)];
    out[y * w + x] = s / 9.0f;
}

One note worth its weight: threads within a warp still get their x from the fastest-varying index, so a 16×16 block maps 32 consecutive threads to two consecutive image rows — and coalescing (day 07) works out naturally. If you'd mapped y to threadIdx.x instead, every lane would jump a full row and the loads would scatter. Dimension order is a performance decision.

02 · Grid-stride loops

The loop every production kernel opens with

The day-13 mapping hardwires one thread = one element, which means grid size scales with problem size. Production kernels decouple the two with a grid-stride loop:

__global__ void vecadd(const float* x, const float* y,
                              float* out, int n){
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = gridDim.x * blockDim.x;
    for (; i < n; i += stride)
        out[i] = x[i] + y[i];
}

The grid now covers however many threads you chose to launch — say 46 SMs' worth — and each thread loops, hopping forward by the total thread count until the array is done. Three properties make this the house style: it works for any n with no ceil-division bookkeeping in the launch; the same compiled kernel runs efficiently on a 2 GB laptop GPU and a 80 GB data-center card; and each thread processes multiple elements, amortizing its setup cost and — as day 28 will show — opening the door to processing 2–4 elements per iteration for more memory-level parallelism.

03 · Choosing block size

256, and stop worrying — usually

Block size is the knob beginners obsess over and experts mostly set to a multiple of 32 near 128–512. The constraints stack up quickly: it must be a multiple of 32 or you waste lanes (day 03/04); it changes the warp/block trade-off that day 06 described — more threads per block means fewer blocks, and blocks are the unit of scheduling flexibility; and it interacts with day 09's register budget, since a bigger block needs more registers in total before the SM accepts it. The 4070 caps a block at 1,024 threads and an SM at 48 resident warps.

Rules of thumb that survive contact with production: 256 is a fine default; use 128 when a kernel has heavy shared-memory needs per block (more blocks per SM helps scheduling); use 512–1,024 when threads cooperate a lot within the block. Then measure — day 22's autotuner exists precisely because the analytic answer loses to a 30-second search.

04 · A worked mapping

From problem size to launch, once, slowly

Say the job is 4096×4096 matmul with 16×16 blocks. Blocks needed: (4096/16) × (4096/16) = 256 × 256 = 65,536 blocks, each with 256 threads — 16.7 million threads total. Sounds absurd; it's routine. The GPU's scheduler cycles waves of blocks onto the 46 SMs — 16,896 threads in flight per SM (day 06's occupancy math) — and keeps issuing waves until the grid drains. You don't manage this. You state the grid; the hardware finds the parallelism. That's the mental shift from CPU threading, where you'd never dare spawn a million threads.

Next: day 15 — the payoff. Tiling with shared memory: fixing day 12's naive matmul with one idea, 16× the ceiling for free.