Optimizing Global Memory Access¶
Upstream NVIDIA / H200 tuning reference
This NVIDIA-platform tutorial is retained from tile-ai/TileOPs.github.io. All measurements, figures, hardware parameters and rules involving SMs, warps and shared memory belong to the original H200 / CUDA context. They are not Ascend 910B1 results or hardware specifications. These cases illustrate tuning methods; hardware rules and performance must be validated again when porting to Ascend. See Timing for this fork's Ascend measurements.
When a thread reads several elements from a row, the access can be written four ways. This page measures all four on two workloads and explains how to choose among them.
Checking whether DRAM bandwidth is the current limit¶
Elementwise and Reduction are the typical memory-bound kernels. Each recommendation below 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, another factor may set the limit, and some of the conclusions here can reverse.
The table below uses those two conditions to divide the space into three regimes. It gives the test for each regime and how the conclusions of these two pages apply there. The named limit is the dominant factor; more complex kernels usually have more factors 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 every change |
- 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. Changing the access pattern has no benefit in this regime, because another factor is setting the pace.
- With few blocks, load width matters more than the coalescing rules predict. There are too few warps to hide memory latency behind concurrent requests, so the remaining lever is to make each request wider and keep more bytes in flight per thread. Any change that trades load width for something else can reverse here.
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. The quality 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 requires three things at once:
- Contiguous addresses — the 32 threads of one instruction address a run with no holes in it;
- 32-byte alignment — the start address is a multiple of 32, so a run does not spill into one more sector;
- 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 requirement, and sector
utilization is \(1/V\):
striped — adjacent threads take adjacent elements. The addresses are now contiguous, but each thread fetches only one element per instruction, which breaks the third requirement. Reading \(V\) elements takes \(V\) instructions.
blocked + vectorized — still one contiguous run per thread, 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 consumption reads 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.
T.vectorized, T.Parallel, and T.copy hand that decision to TileLang's
layout inference. They differ in how much the programmer still specifies:
T.vectorized specifies the access width per thread and infers the thread
mapping; T.Parallel specifies neither, so it decides both how loop dimensions
are split across threads and how wide each read is; T.copy specifies only a
source region and a destination region, then generates the whole copy
(coalesced_width and loop_layout are available when the inferred result
needs to be overridden). Layout inference handles vectorization, address
alignment, and avoiding bank conflicts on the shared-memory side. Those are the
hardware-friendly details that are easy to get wrong by hand.
Writing indices by hand with T.serial means guaranteeing those three
requirements directly; handing the copy to layout inference means specifying
only the copy extent. The measurements below show which one to choose and what
bandwidth each reaches.
Measurements¶
Both workloads were measured upstream on an H200. The comparison is the memory bandwidth reached by each of the four access patterns: bytes moved divided by kernel time, in TB/s. The two workloads impose 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¶
-
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. -
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.
-
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. -
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, computations that require a contiguous run per thread, such as a serial prefix, cannot use it.
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 at the top of this page, element-by-element
blocked is the wrong form and should not be copied. striped works, but it is
not the fastest option (see point 4 above).
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 above):
@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]