Tail Handling and Bounds Checks in Partial Tiles

Almost no real problem divides evenly by the block size. A million-element buffer with 256 threads per block uses 3907 blocks, which covers 1,000,192 threads — the last 192 of which have no element and must not touch memory. The usual response is an if at the top of the kernel that returns when the thread’s index exceeds the problem size. That response is correct on its own but interacts badly with every other construct a kernel might use: barriers, shared-memory tiling, warp-level primitives, and index arithmetic on neighbours. The tail is where a kernel stops being a uniform bulk computation and becomes a collection of special cases.

The grid size is a ceiling, not a count

Rounding up is unavoidable, and the interesting choice is only how to compute the rounded value without overflow. (n + t - 1) / t overflows when $n$ is within $t-1$ of the maximum of its type; n / t + (n % t != 0) never does. With 64-bit sizes the overflow is unreachable in practice, but a 32-bit element count and a 1024-thread block put the danger zone within 1023 elements of $2^{32}$ — a 4 GiB buffer of single-byte elements lands there exactly.

Having rounded up, the launched thread count satisfies $\text{blocks} \times t \ge n$ and the excess is at most $t-1$. Those excess threads are all in the final block. Every other block is full, so the bounds check on all of them is a comparison that always succeeds — but the compiler cannot know that, and the check is still emitted for every thread in the grid. This is the hidden cost of the uniform check: for a large problem it is a compare and a branch on every thread, against a benefit that applies to at most one block. The cost per thread is a few instructions; the benefit is that the kernel does not need to specialise the last block.

The alternative is to split the launch: one kernel over the full blocks, with no check at all, and a second small launch whose block size is exactly the remainder. This removes the check from the bulk path but doubles the number of launches and forces the bulk kernel to be compiled without the check. The trade is instruction count against launch overhead and code duplication, and it only pays when the per-element body is a handful of instructions — for a heavy body the check is in the noise.

Early return versus predication

An early return inside a warp becomes a predicate on the remaining lanes. The hardware does not remove those lanes from the warp; it masks them off, and the warp continues to occupy the same execution slots with fewer active lanes. A warp that is entirely beyond the bound still exists, still fetches and executes the branch, and then exits, so the block still holds that warp’s resources until the block completes.

The performance consequence is concentrated in the last block. If the block size is 256 and the problem has 1,000,000 elements, the last block’s first 64 threads are in range and the remaining 192 are not: two warps do work and six do nothing, so that block runs at a quarter of its nominal width. Spread over 3907 blocks the loss is irrelevant; the point is that it is confined to one block rather than paid by every block in the grid. Choosing a block size that is not a multiple of 32 is far worse, because every block then has a partially masked warp on every instruction, not just the last one.

There is a second difference, in code that has to run at all. A lane that returned early does not execute the rest of the kernel, so any operation that requires all lanes of a warp to participate — a warp shuffle, a collective reduction issued as a warp intrinsic, or an instruction that borrows stack space under a reconvergence assumption — cannot be placed after the return unless its arguments are arranged to tolerate inactive lanes. Predicating the memory access rather than the thread keeps the lanes alive:

const size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
const bool live = i < n;

// Out-of-range lanes still execute, contributing the identity value,
// so the warp's shuffle below has all 32 lanes in its mask.
float v = live ? in[i] : 0.0f;
for (int off = 16; off > 0; off >>= 1)
    v += __shfl_down_sync(0xffffffffu, v, off);

// After the tree, lane 0 of each warp holds that warp's total.
// The warp count is exact because the block size is a multiple of 32.
if ((threadIdx.x & 31) == 0)
    partial[blockIdx.x * (blockDim.x / 32) + (threadIdx.x >> 5)] = v;

The trade here is that every lane now runs the arithmetic of the body, including lanes with no data, against the early return which lets inactive lanes skip it. For a short body the predicated form is simpler and avoids the divergence penalty of a branch; for a long body with expensive per-element work the early return wins. What neither form can do is let a thread skip a __syncthreads().

Barriers cannot be predicated away

A __syncthreads() executed by some threads and not others is undefined behaviour, and in practice it hangs. The reason is mechanical: the barrier is a count of arrivals. If a block of 256 threads is arranged so that the 64 in-range threads reach the barrier and the 192 out-of-range threads returned early, the arrival count never reaches the expected number and the block waits forever.

This makes the early-return idiom illegal in any kernel that tiles through shared memory, because the tiled loop is bracketed by barriers. The pattern has to change shape: the loading of the tile becomes predicated, with out-of-range positions filled with a neutral value, and the barrier is reached by every thread in the block.

__shared__ float tile[256];
const int t = threadIdx.x;
const size_t base = (size_t)blockIdx.x * blockDim.x;
const size_t gi = base + t;

tile[t] = (gi < n) ? in[gi] : 0.0f;   // every thread stores something
__syncthreads();                       // every thread arrives

const size_t oi = base + (blockDim.x - 1 - t);
if (oi < n) out[oi] = tile[t];

Filling with zero is not neutral for every operator — for a maximum reduction the identity is the smallest representable value, and for a product it is one. Choosing the wrong identity is a silent correctness bug that appears only when the problem size is not a multiple of the tile size, which is exactly the case that a test with a power-of-two buffer never exercises. Designing the identity into the load is therefore part of the tail handling, not a detail of the reduction.

Neighbour access and the edge halo

Stencils widen the problem. A one-dimensional three-point stencil reads i-1, i and i+1, so the thread at index 0 reads -1 and the thread at the last index reads n. Both are outside the buffer. Wrapping the index modulo $n$ gives periodic boundaries, clamping it gives replicating boundaries, and substituting a constant gives fixed boundaries; the choice is mathematics, not implementation, and it must be made explicit before the code is written.

The subtle failure is the offset that is in range but belongs to a different row. For a two-dimensional stencil, i-1 at column 0 of row $y$ is the last column of row $y-1$; the address is valid, the load succeeds, and the result is wrong. Bounds checks in two dimensions must test the column and the row separately, which is why the neighbour pattern is normally expressed in tile coordinates with explicit halo loads rather than as an offset on the linear index. The horizontal halo is described in tiled stencils and halo exchange in shared memory.

Degenerate sizes

A launch with a zero-length grid is invalid: the number of blocks per dimension must be at least one. A problem of size zero therefore has to be handled on the host by not launching at all, or by launching one block whose single predicate is false. The second option costs a launch and is easy to get wrong when the kernel later assumes at least one active thread per block.

Sizes near the limits of the index type deserve the same treatment as zero. A problem larger than $2^{32}$ elements has a linear index of $2^{32}$ for its first out-of-range element, and any 32-bit arithmetic anywhere in the index computation wraps that value to 0. Because the wrap lands on the first element rather than outside the buffer, the access is legal and the bug is data corruption rather than a fault.

Rectangular problems add a case that a linear launch hides. If a 2-D grid is used with a check per axis, then a thread may pass the x check and fail the y check, or the reverse, and a kernel that computes its linear index before the checks will compute an in-range index for an out-of-range coordinate whenever the width is not a multiple of the block’s x dimension. The sheared output of thread hierarchy and index arithmetic is the typical symptom. The fix is to check both coordinates before forming any offset.

Testing the tail deliberately

The tail is exercised by problem sizes that are one less than a multiple of the block size, one more than a multiple, and smaller than a single block. Sizes of 1, $t-1$, $t$, $t+1$ and $\text{prime} \times t + 1$ cover the transitions. A buffer filled with a sentinel value and checked afterwards catches both the missing elements and the overwritten ones; a size chosen as a power of two, which is the natural choice when testing a bulk kernel, exercises none of these cases.

The general lesson is that the bounds check is the least interesting part of tail handling and the part most likely to be correct. The failure modes concentrate where the check meets a construct that assumes uniformity — a barrier, a collective, a fill value, or an index formed from two axes — and those are the places worth designing for first.