2D Stencil: Pay the Halo Ring Once

Lesson

The computation first. For every interior point, output = 0.5 x centre + 0.125 x (left + right + up + down). Any point on the grid's border is copied through unchanged. The profile case is a 4096 x 4096 grid, and the harness verifies every element, including the borders and the odd sizes 33 and 257 whose tiles do not divide evenly.

Now count what the starter spends. A 32 x 8 block computes 256 outputs, and five reads each is 1,280 global-memory request requests. But those requests touch only 336 distinct points: the 256 the block owns plus a ring of 2 x (32 + 8) = 80 cells around them. Every interior point is fetched up to five times, by up to five different threads, when once would do.

The fix is the discipline from the 1D shared-halo lab, grown a dimension: stage the 336 points once in a static 10 x 34 shared tile, one barrier, then serve all five reads per output from shared memory.

Where each read is served from:
Starter, per output: centre global-memory request, left global-memory request, right global-memory request, up global-memory request, down global-memory request.
Solution, per output: all five from the shared tile.
Solution, per block: global-memory request is touched exactly once per point, 336 fetches for 256 outputs.

Two details are the actual test of understanding. First, a five-point rule never reads the four diagonal corners of the 34 x 10 patch, so the halo is a ring, not a full border; the solution still writes the corners once with a constant, because undefined shared memory is the kind of bug that only fires on someone else's machine. Second, cell count is not sector count: the halo rows above and below coalesce like any row, but each side column is a lone lane that drags in a whole 32-byte sector for one float. That is why staging only the x direction and leaving the up and down reads in global memory fails the gate: it removes the cheap half of the traffic and keeps two full sweeps of the expensive kind.

Contract notes: the block is pinned at 32 x 8 x 1 because the tile extents are literals and the arithmetic above assumes a warp is one tile row. The tile is a static __shared__ array, so sharedMemBytes stays 0. The harness derives the grid from N, so the Grid fields in the launch panel are inert. Measured on the profile runs, the starter issues about 1.36 billion sectors and the staged version about 578 million, roughly 2.4x apart; the gate at 740 million sits between them.

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.

Halo loading reduces repeated global requests, but neighboring blocks still duplicate halo values. Totals include the runner repetitions and are not physical DRAM bytes.

What you will learn

• Size a two-dimensional halo: see why a 32 x 8 output tile needs a 34 x 10 input patch.
• Weigh the halo tax as perimeter against area, and see why staging only one direction is not enough.
• Stage a static shared tile cooperatively, with one barrier between the load and the stencil.
• Use issued global-memory sectors to measure what the halo ring actually cost, in sectors rather than in cells.

Starter code

    // Naive five-point stencil: one output per thread and five separate global
    // loads. It is already well coalesced, and that is the point: the loads are
    // neatly shaped, there are simply five times more of them than needed.
    const int x = blockIdx.x * blockDim.x + threadIdx.x;
    const int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x >= N || y >= N) {
        return;
    }

    const int o = y * N + x;
    if (x == 0 || y == 0 || x == N - 1 || y == N - 1) {
        output[o] = input[o];
        return;
    }
    output[o] = 0.5f * input[o]
            + 0.125f * (input[o - 1] + input[o + 1] + input[o - N] + input[o + N]);