1D Stencil: Reuse Neighbor Values With A Halo Tile

The naive three-point stencil is correct, but every output reloads values its neighboring threads also requested. Stage one block of input plus a one-element halo on each side in shared memory, then reuse those values for the stencil. GPUFlight judges the reduction in issued global-memory sectors rather than cache-dependent wall time.

Languages: CUDA C++, Python, Rust

Lesson

A three-point stencil reads input[i-1], input[i], and input[i+1] for every interior output. Neighboring threads therefore ask for overlapping values. Hardware caches may satisfy much of that reuse and make a short benchmark look deceptively healthy, but the kernel still issues three global loads per output. A shared-memory halo tile makes the reuse explicit: every thread loads its center element once, the edge threads load the two neighboring halo elements, the block synchronizes, and the stencil reads its three values from shared memory. The profile launch uses 256 threads, so the tile needs 258 floats, or 1,032 bytes, of dynamic shared memory. GPUFlight's SASS global-sector total measures the issued traffic before cache capacity can hide the redundant requests.

Metric contract: algorithmic reads/writes count source-level sweeps. GLOBAL_MEMORY_SECTORS counts SASS-derived issued/requested sectors in the selected profiling scope, including repeated launches. It is cache-blind: it is neither L2 traffic nor device-memory bytes. Trace elapsed time is a separate measurement; a global load may hit cache.

Neighbor references are global requests, not necessarily separate device-memory fetches. The profile repeats launches; inspect the runner repetition count before interpreting aggregate sectors.

What you will learn

Identify overlapping neighbor loads in a stencil access pattern.
Stage a block-sized tile plus left and right halo elements in dynamic shared memory.
Place synchronization between cooperative loading and tile consumption.
Use issued global-memory sectors to measure explicit data reuse.

Starter code

    // Correct but redundant: adjacent threads repeatedly load the same neighbors.
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        if (i == 0 || i == N - 1) {
            output[i] = input[i];
        } else {
            output[i] = 0.25f * input[i - 1] + 0.5f * input[i] + 0.25f * input[i + 1];
        }
    }