Estimating Kernel Data Movement from Access Patterns

Hello! Last lesson distinguished a warp waiting for a load result from a memory system that has reached its throughput limit. To investigate either possibility, you need an estimate of the traffic the kernel’s accesses ask the memory hierarchy to handle. This lesson builds that estimate from source-level indices, one warp instruction at a time. The method is useful both for forming performance hypotheses and for checking whether a profiler’s memory counts are plausible.

Start with bytes the program asks to use

Consider a vector-add kernel that computes one output element per thread:

int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
    c[i] = a[i] + b[i];
}

Each valid thread loads two 4-byte floats and stores one. For $N$ valid elements, the logical traffic is

This is the right first count, but it does not tell you how many bytes the memory hierarchy must handle. A thread requesting four bytes may cause a larger unit of data to be accessed. Equally, two threads or warps may request data already present in a cache. Keep three quantities separate:

• Logical bytes: the operands and results requested by the program.
• Warp-level sector demand: the 32-byte sectors touched by each executed global-memory instruction.
• Traffic at a particular boundary: bytes observed between, for example, a cache and device memory.

These quantities can differ substantially. In particular, a sector-count estimate is not automatically a DRAM-byte estimate.

The memory-analysis walkthrough provides a useful visual model of those boundaries.

Turn warp addresses into sectors

For a single global load or store instruction, write down the addresses requested by its active lanes. Group the bytes touched by those addresses into aligned 32-byte sectors. Count each distinct sector once for that warp instruction. For naturally aligned 4-byte words, a quick calculation is to group addresses by

If an access crosses a sector boundary, count both sectors. Repeat this analysis separately for each load and store instruction, then across the warps that execute it. In compact form, if $S_{w,j}$ is the set of sectors touched by warp $w$ during memory instruction $j$, the estimated sector demand is

The sum counts sectors per instruction, not unique sectors across the entire kernel. The same sector can appear again when another instruction requests it; a cache may satisfy that later request without fetching it again from device memory.

Suppose all 32 lanes of a warp load consecutive floats starting at a 32-byte boundary. They request 128 distinct useful bytes in four sectors:

Now suppose neighboring lanes instead load floats 128 bytes apart. Each lane touches a different sector: the warp requests only 128 useful bytes but touches 32 sectors, or 1,024 sector bytes. That is eight times the sector demand for this load, not necessarily eight times the kernel’s runtime or DRAM traffic.

The official programming guide develops both cases, then applies the same reasoning to a transpose.

Two qualifications make this method more useful than the slogan “adjacent threads need adjacent addresses.” First, alignment matters: 32 consecutive floats beginning four bytes past a sector boundary touch five sectors rather than four. Second, addresses need not be ordered by lane. A permutation of 32 floats that still covers the same aligned 128-byte span touches the same four sectors. The question is which sectors the active lanes collectively touch.

Account for the whole kernel, not just its worst instruction

Return to vector add. Assume $N$ is a multiple of 32, the three arrays are suitably aligned, and all lanes participate. Each warp’s load from a touches four sectors; its load from b touches four more; its store to c touches four. The total is 12 sectors, or 384 sector bytes, for 32 results. This equals $32\times 12=384$ logical bytes.

A gather shows why counting every instruction matters:

int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
    out[i] = data[index[i]];
}

Assume index and out are contiguous and aligned. Also assume, for this estimate, that the 32 indices in each warp select 32 floats in different sectors. For one full warp, the index load touches four sectors, the gather from data touches 32, and the output store touches four. The total is 40 sectors, or 1,280 sector bytes, although the logical count remains 384 bytes.

Thus the gather load alone has eight times the sector demand of a contiguous load. The whole kernel’s estimated demand is $40/12$, about 3.3 times that of three coalesced accesses. If indices repeat within a warp, fewer than 32 data sectors may suffice; if data is reused across warps, caches may reduce traffic farther down the hierarchy. State those assumptions rather than presenting 1,280 bytes as an unconditional measurement.

For a concrete scale, with $N=1,048,576$ valid elements under these assumptions, logical traffic is 12 MiB. The estimated sector demand is 12 MiB for three coalesced accesses, versus 40 MiB for the gather pattern. Those are demands at the warp-access level, not predictions of bytes read and written to DRAM.

Read a two-dimensional access as a one-dimensional address

The transpose in the programming guide provides a common trap. For a row-major float matrix with leading dimension $m$, element $(r,c)$ has element index $rm+c$. In the guide’s kernel, neighboring threadIdx.x values increase myCol. Its input read has the form

a[myRow * m + myCol]

so neighboring lanes read neighboring floats. Its output write has the form

c[myCol * m + myRow]

so neighboring lanes write addresses $4m$ bytes apart. A visually orderly mapping of threads to a matrix is not enough; expand the address expression and inspect its change from one lane to the next.

Suppose $m=1024$, a warp has 32 active lanes, and its input row begins on a sector boundary. The read touches four sectors. Output addresses are 4,096 bytes apart, so its store touches 32 distinct sectors. At the per-warp-instruction level, that is 36 sectors for a transpose that logically reads and writes 256 bytes per warp. If both accesses touched four sectors, the corresponding count would be eight.

This 36-versus-eight comparison locates the inefficient instruction. It does not mean every output store moves 1,024 fresh bytes to DRAM: nearby warps may update different words in the same sectors, and lower levels may combine or cache those writes. Later optimization lessons will address changing such access patterns; for now, the objective is to quantify what the current pattern requests.

Choose the measurement boundary before comparing with a profile

A prediction is most convincing when it names the counter it could be compared against. The following technical-blog analysis reports both load sectors per warp-level request and device-memory measurements.

For the blog’s contiguous load, four sectors per warp-level load request matches the address calculation. For its scattered load, 32 sectors per request matches the assumption that each lane reaches a distinct sector. The ratio is a direct check of coalescing at that measured level. DRAM counters answer a different question: how much traffic reached device memory under that run’s cache behavior and other conditions.

The distinction also clarifies the previous lesson’s stall analysis. A warp waiting on a load result could suffer from latency even when its access uses only four sectors. A kernel issuing many 32-sector loads could put heavy pressure on memory throughput even if scheduling hides much of each load’s latency. Neither conclusion follows from the sector estimate alone, but the estimate tells you what evidence to seek.

A reusable estimation procedure

For a new kernel, make a short accounting sheet before looking at counters:

1. List executed memory operations. Include loads, stores, index-array reads, and any repeated accesses in loops. Note element widths and which threads are active at boundaries.
2. Compute logical bytes. Count requested operands and results, stating whether repeated reads count again. This gives a workload-level reference.
3. Expand addresses across a warp. Hold block and warp coordinates fixed; vary lane position. Check base alignment, strides, permutations, and whether lanes select the same address.
4. Count sectors per instruction and scale up. Give a typical full-warp estimate and, where necessary, separate estimates for partial boundary warps or irregular indices.
5. Name the memory boundary. Treat the sector calculation as request-side demand. To predict lower-level or DRAM bytes, add explicit assumptions about cache misses, reuse, and stores; otherwise report them as unknowns to measure.

This separates a fact visible in code—“these lanes touch 32 sectors”—from a hypothesis about the machine—“those sectors are fetched from DRAM every time.” That separation is especially important when evaluating a claimed optimization: reducing sector demand at one interface is promising, but the runtime benefit depends on whether that traffic was limiting performance.

Takeaways

Estimate data movement in layers. First count logical bytes, then count the distinct 32-byte sectors touched by each warp memory instruction, and only then reason about traffic at L1, L2, or DRAM using stated cache and reuse assumptions. Alignment, active lanes, and address expressions matter as much as the number of elements processed. A poor pattern in one instruction need not inflate the entire kernel by the same factor.

Next lesson, you will read public architecture documentation for performance-relevant changes. The habit developed here—specifying the level at which a number applies—will help you decide which architectural changes should alter a kernel’s predicted behavior.