Skip to content

Optimizing Global Memory Access

When a thread reads several elements from a row, the access can be written in four ways. This page compares the four on two workloads and gives the basis for choosing one.

Checking whether DRAM bandwidth is the current limit

Elementwise and Reduction are the typical memory-bound kernels. Each recommendation on this page states when it applies, why it applies, and what the wrong and right code look like.

Every measurement on this page uses the same conditions: an input larger than the 60 MiB L2, and enough blocks to fill the whole card (an H200 has 132 SMs). Under those conditions, DRAM bandwidth is the main bottleneck, and differences in access pattern show up directly in performance.

Where this applies

Outside these conditions, the main limit may be another factor, and some of the conclusions on this page can reverse.

The table below uses those two conditions to divide the space into three regimes. Each row gives the test for the regime and how the conclusions of the global memory and shared memory pages apply there. The table lists the dominant limit; the more complex the kernel, the more factors are active at the same time.

Regime Test Main limit How to use the conclusions
Bandwidth saturated Input > 60 MiB, blocks at more than twice the SM count DRAM bandwidth, that is sector utilization Apply directly
Small data Input fits in L2, one call takes tens of microseconds or less Fixed launch overhead, cache state Avoid the wrong forms; changing access pattern buys nothing
Few blocks Blocks fewer than twice the SM count The width of each load instruction, bytes in flight Keep load width first, and measure each change separately
  • With small data, launch overhead and cache state dominate. On 65536 × 4096 (512 MB), the four access patterns of the same row-reduction kernel (fp16, 256 threads, clocks unlocked) measure 4.20 to 4.43 TB/s, within 6% of each other. On 2048 × 4096 (16 MB, which fits in L2), one call takes a dozen or so microseconds, and two measurements of the same access pattern can differ by threefold. In this regime another factor limits performance, and changing the access pattern has no benefit.
  • With few blocks, load width matters more than the coalescing rules predict. With too few warps, memory latency cannot be hidden behind concurrent requests. The only option is to make each request wider and keep more bytes in flight per thread. In this regime, any change that gives up load width for another gain can produce the opposite result.

Coalescing global memory accesses

A global memory read moves data in units of a 32-byte sector.

Name Size What it is
cache line 128 bytes The L1 and L2 line, and the unit a cache lookup uses
sector 32 bytes A cache line is 4 sectors; L1 and L2 transfer whole sectors

Lookups operate on cache lines, while transfers operate on sectors. On a sector miss, L1 requests only that sector from L2 instead of pulling the whole line. Therefore fetching 1 byte costs the same as fetching all 32, and the efficiency of a memory instruction is measured by its sector utilization: bytes actually used / (sectors touched × 32).

The hardware coalesces a warp's 32 accesses into as few 32-byte transactions as it can. Reaching the minimum number of transactions requires three conditions at once:

  1. Contiguous addresses — the 32 threads of one instruction address a run with no holes in it;
  2. 32-byte alignment — the start address is a multiple of 32, so a run does not spill into one more sector;
  3. 16 bytes fetched per thread per instruction — one instruction then covers \(32 \times 16 = 512\) contiguous bytes, which is 16 fully used sectors.

An instruction that satisfies all three is as friendly to the hardware as it gets.

A thread reading \(V\) elements (\(V\) = elements per row / threads) has four access patterns available.

blocked — each thread takes one contiguous run. For a fixed c, adjacent threads are \(V\) elements apart. This breaks the first condition, and sector utilization is \(1/V\):

for c in T.serial(V):
    acc[0] = acc[0] * X[row, tx * V + c]

striped — adjacent threads take adjacent elements. The addresses are contiguous, but each thread fetches only one element per instruction, which breaks the third condition. Reading \(V\) elements takes \(V\) instructions.

for c in T.serial(V):
    acc[0] = acc[0] * X[row, c * threads + tx]

blocked + vectorized — each thread still takes one contiguous run, but T.vectorized reads a full 16 bytes at a time, satisfying all three:

buf = T.alloc_local((V,), dtype)

for c in T.vectorized(V):
    buf[c] = X[row, tx * V + c]
for c in T.serial(V):
    acc[0] = acc[0] * buf[c]

staged — T.Parallel performs the copy, and the consumption phase reads from shared memory, which also satisfies all three:

sh = T.alloc_shared((threads, V + pad), dtype)

for t, c in T.Parallel(threads, V):
    sh[t, c] = X[row, t * V + c]
T.sync_threads()
for c in T.serial(V):
    acc[0] = acc[0] * sh[tx, c]

The four patterns differ in who decides which thread reads which elements, and how wide each read is.

With T.serial, the loop body runs sequentially on a single thread. The index expression is translated directly into memory instructions, with no coalescing or vectorization. The pattern in the source is the pattern the hardware sees.

For T.vectorized, T.Parallel, and T.copy, TileLang's layout inference decides the mapping. The three differ in what the programmer writes:

  • T.vectorized states the access width per thread, and layout inference derives the thread mapping;
  • T.Parallel does not state the width; layout inference decides both how the loop dimensions are split across threads and how wide each read is;
  • T.copy states only a source region and a destination region, and generates the whole copy; coalesced_width and loop_layout are available when the inferred result needs to be specified by hand.

Layout inference handles the rest: vectorization, address alignment, and avoiding bank conflicts on the shared memory side. Those are the parts that are hardware-friendly but easy to get wrong by hand.

When indices are written by hand with T.serial, the programmer guarantees the three conditions one by one; when the copy is left to layout inference, only the copy extent is stated. The measurements below show which form to choose and the bandwidth each reaches.

blocked tx * V + c 0 1 2 3 4 5 6 7 4 sectors fetched, 2 of 8 elements used in each blocked + vectorized T.vectorized(V) 0 1 2 3 4 5 6 7 one 16-byte vector read: 4 sectors fetched, all fully used striped c * threads + tx 0 1 2 3 4 5 6 7 1 sector fetched, all 8 of 8 elements used staged T.Parallel row copy 4 sectors fetched, all fully used staged[tx * V + c] consumed in shared memory, where there are no sectors Diagram: fp32, 8 threads, V = 4, 8 elements to a sector. The measurements below use bf16, 256 threads, V = 16 — the same shape.

One read instruction. A violet dot marks an element read by a thread in this instruction. The teal background marks a sector fetched by the hardware as a result. Thread numbers appear above the cells, and alternating shades mark thread boundaries. The dashed violet row at the bottom is shared memory, which is not organized in sectors.

Measurements

This section measures two workloads on an H200 and compares the memory bandwidth reached by each of the four access patterns: bytes moved divided by kernel time, in TB/s. The two workloads place different requirements on the order in which elements are processed, and that requirement determines which access patterns are available.

The SM clock is locked at 1830 MHz. The input is bf16 \(65536 \times 4096\) (512 MB, which must exceed the 60 MiB L2; otherwise the measured number is L2 bandwidth). Each configuration runs three times, with the three results agreeing to within ±0.5%. staged uses two columns: one without padding, where the stride is exactly \(V\) words and produces bank conflicts (see Optimizing Shared Memory Access), and one using the best value among several measured padding choices.

Workload 1, product along a row: the elements of each row are multiplied together. Order does not affect the result, and each row writes one value. All four access patterns are available.

Threads \(V\) blocked
TB/s		 striped
TB/s		 blocked + vectorized
TB/s		 staged, no pad
TB/s		 staged, padded
TB/s
512 8 3.02 3.04 3.22 3.05 3.02
256 16 1.83 3.31 3.81 3.79 3.79
128 32 0.95 3.36 3.99 3.81 3.98
64 64 0.48 3.49 3.32 2.80 4.02
32 128 0.47 3.43 3.11 2.83 3.88

Workload 2, serial prefix product along a row: each position depends on every element to its left. The order cannot change, and the whole row is written back. striped is unavailable here because a thread cannot hold a contiguous run.

Threads \(V\) blocked
TB/s		 blocked + vectorized
TB/s		 staged, no pad
TB/s		 staged, padded
TB/s
512 8 0.92 4.20 2.47 3.15
256 16 0.42 3.85 1.63 3.29
128 32 0.34 2.76 0.89 3.35
64 64 0.26 1.80 0.46 3.69
32 128 0.27 1.60 0.46 3.24

Choosing an access pattern

  1. Element-by-element blocked is the worst access pattern whenever \(V > 1\). For a fixed c, adjacent threads are \(V\) elements apart and sector utilization is \(1/V\), so performance drops as \(V\) grows. In workload 1, it falls from 3.02 at \(V = 8\) to 0.48 at \(V = 64\). That relation follows from the coalescing rules and does not change with shape.

  2. Use vectorized blocked at small \(V\); once \(V\) grows enough that register pressure cuts occupancy, switch to padded staged. Vectorized blocked keeps the run in registers (\(V/2\) registers per thread for bf16). staged puts the run in shared memory and pays one synchronization to recover those registers. The crossover depends on the register budget left by the rest of the kernel, not on a fixed \(V\): at the same row width, the two workloads above cross at \(V = 64\) and \(V = 32\). Measure that crossover on the target kernel.

  3. The shared buffer of staged has to avoid bank conflicts. Declared as (threads, V) the stride is exactly \(V\) words, which conflicts whenever \(V\) is a power of two; the padding rule and its candidates are in Optimizing Shared Memory Access. In workload 2 the same configuration measures 0.46 unpadded and 3.69 padded.

  4. striped coalesces fully, but spends one instruction per element. That puts it above element-by-element blocked and below vectorized blocked: 3.31 versus 1.83 and 3.81 at \(V = 16\). It fits cases where minimizing the code change matters more than extracting the last bit of bandwidth. Because each thread holds non-contiguous elements, it cannot be used by computations that require each thread to hold a contiguous run, such as a serial prefix product.

The two listings below are complete templates for the recommended access patterns. M, N, V, threads, pad, and dtype are all compile-time constants. Among the four listings in Coalescing global memory accesses, element-by-element blocked is the wrong form. striped works, but it is not the fastest option (see point 4 in Choosing an access pattern).

Recommended at small \(V\) — vectorized blocked:

@T.prim_func
def main(X: T.Tensor((M, N), dtype), Out: T.Tensor((M, threads), "float32")):
    with T.Kernel(M, threads=threads) as row:
        tx = T.get_thread_binding()
        buf = T.alloc_local((V,), dtype)
        acc = T.alloc_local((1,), "float32")
        acc[0] = T.cast(1.0, "float32")

        for c in T.vectorized(V):                  # one 16-byte vector read
            buf[c] = X[row, tx * V + c]

        for c in T.serial(V):                      # consume, in whatever order the computation needs
            acc[0] = acc[0] * T.cast(buf[c], "float32")

        Out[row, tx] = acc[0]

Recommended at large \(V\) — padded staged (for the crossover, see point 2 in Choosing an access pattern):

@T.prim_func
def main(X: T.Tensor((M, N), dtype), Out: T.Tensor((M, threads), "float32")):
    with T.Kernel(M, threads=threads) as row:
        tx = T.get_thread_binding()
        sh = T.alloc_shared((threads, V + pad), dtype)   # for pad, see the shared memory page
        acc = T.alloc_local((1,), "float32")
        acc[0] = T.cast(1.0, "float32")

        for t, c in T.Parallel(threads, V):        # copy; the mapping comes from layout inference
            sh[t, c] = X[row, t * V + c]
        T.sync_threads()

        for c in T.serial(V):                      # consume
            acc[0] = acc[0] * T.cast(sh[tx, c], "float32")

        Out[row, tx] = acc[0]