Reconstructing Fine-grained Memory Reuse Structure in GPU Kernels

arXiv:2611.14741v1 [cs.DC] 10 Nov 2026

Reconstructing Fine-grained Memory Reuse Structure in GPU Kernels Yanbo Zhao

Yueming Hao

Zecheng Li

North Carolina State University Raleigh, NC, USA yzhao62@ncsu.edu

Meta Menlo Park, CA, USA yhao@meta.com

North Carolina State University Raleigh, NC, USA zli94@ncsu.edu

Mao Lin

Shuyin Jiao

Keren Zhou

University of California, Merced Merced, CA, USA mlin59@ucmerced.edu

North Carolina State University Raleigh, NC, USA sjiao2@ncsu.edu

George Mason University Fairfax, VA, USA kzhou6@gmu.edu

Xu Liu

Jiajia Li

North Carolina State University Raleigh, NC, USA xliu88@ncsu.edu

North Carolina State University Raleigh, NC, USA jiajia.li@ncsu.edu shared and reused during execution. This requires understanding which instructions access the same data, how those accesses are distributed across threads, and across which execution scopes reuse occurs. These facts, however, are what today’s software stack fails to surface. GPU hardware provides mechanisms suited to reuse at different execution scopes: registers within a thread, shuffles within a warp, shared memory within a block, and the shared L2 cache across blocks. Consider a buffer placed in shared memory: it may carry communication between threads, or each thread may merely write its own elements and read them back. Both uses occupy the same memory space, but the first involves inter-thread communication while the second exhibits only thread-private reuse. This difference changes which storage mechanisms are appropriate, and telling the two apart requires relating the accessing instructions to the threads and scopes that reuse their data. Programming frameworks [8, 26, 29] do not represent these reuse relations in an explicit, queryable form. Judging whether a placement matches actual use therefore requires exposing these instruction-level sharing relations.

Abstract Optimizing GPU data movement requires knowing which instructions access the same data and across which execution scopes reuse occurs. Existing profilers can attribute memory behavior to individual instructions but do not expose the reuse relations among them, while aliasing, dynamic indexing, and runtime parameters prevent static analysis from establishing these relations in general. Cycle-level simulation can reconstruct such relations, but at costs prohibitive for production workloads. We introduce two abstractions for representing this structure: Topological Distance, which labels each dynamic memory access with the smallest execution scope (thread, warp, block, or grid) containing its most recent reuse, and PCDepGraph, which organizes these labels into a PC-level graph of the kernel’s memory reuse structure. Defined over logical execution scopes rather than transient hardware state, the representation remains stable across runs and GPU targets while directly indicating the sharing scope that candidate optimization mechanisms must cover. Our tool, cuVein, reconstructs PCDepGraph directly from optimized binaries, without sampling, at generally under 10× overhead. Across HPC and AI workloads on NVIDIA RTX 4090 and H100, optimizations guided by PCDepGraph eliminate up to 90.18% of memory instructions and yield up to 2.84× kernel speedup; reuse-scope distributions remain consistent across platforms, supporting that the representation captures program structure rather than transient hardware behavior.

1

1.1

Fundamental Mismatch in Existing Solutions

The example generalizes to a question that recurs across memory optimizations: what is the actual reuse scope of the data an instruction touches? Answering it requires linking each access to its reuse origin and the execution scope containing both. For optimization, this relational information must further be combined with the access’s own attributes. Static analysis knows the program’s thread organization, its explicit synchronization, and each access’s declared memory space. When aliasing, indexing, and control flow are statically resolvable, it can prove how data is shared. In the general case, however, which accesses actually touch the

INTRODUCTION

Over the past decade, GPU compute throughput has grown far faster than off-chip memory bandwidth, making data movement a dominant bottleneck in many modern workloads [2, 5, 14, 35, 36]. Moving data efficiently depends on matching hardware mechanisms to how data is actually 1

Zhao et al.

same address is determined only at runtime, so the actual reuse scope cannot be established at compile time. Dynamic observation is the natural source, and profilers [3, 9, 16, 23] are its standard form: memory-transaction volumes, cache behavior, and stall breakdowns help identify memory-bound kernels, and recent tools attribute metrics such as redundancy rates or value patterns to individual instructions [22, 38–40]. These metrics characterize individual instructions; the missing information is relational. Knowing how many accesses a shared-memory PC issues still does not say whether the data it writes is re-read by other threads or only by the thread that wrote it. Existing metrics do not link accessing PCs, reuse origins, and reuse scopes into something analyses can query together. Cycle-level simulators [20] can reconstruct such relations from full execution records, but at over six orders of magnitude slowdown and typically on scaled-down inputs [4]. Our starting point is that the reuse facts memory-behavior analyses need, which accesses reuse whose data and within which logical execution scope, can be collected directly, without modeling hardware timing. Organizing them, linking each access’s attributes to its reuse origin and scope, produces a representation whose attributes and relations different analyses can jointly query and combine. 1.2

11 ... // four more fields interpolated likewise 12 }

Listing 1. XSBench example: across different blocks, high and low derived from nuclide_grids are heavily reused. 1.3

Our Approach

We characterize memory reuse through the GPU’s logical execution hierarchy. Threads, warps, blocks, and grids form nested scopes defined by the program and its launch configuration, with each scope corresponding to hardware mechanisms suited to serving reuse within it. We classify each dynamic reuse by the smallest execution scope containing both accesses. This captures how data is shared across execution contexts while abstracting away transient physical residency. Aggregating these scope-labeled reuse relations by PC yields a compact representation that connects sharing behavior to the responsible instructions and potential optimization mechanisms, with a size independent of dynamic trace length. We realize this design as two constructs. Topological Distance labels each dynamic memory access with the smallest execution scope (thread, warp, block, or grid) containing its most recent reuse. The Program Counter Dependency Graph (PCDepGraph) organizes these observations into a graph whose nodes are memory-access PCs annotated with access attributes (type, memory space, and per-execution efficiency distributions), and whose edges link each access to its reuse origin, labeled with Topological Distance and weighted by traffic. Because both are defined over logical execution scopes rather than transient hardware state, the representation remains directly interpretable by programmers and compilers. PCDepGraph supports multiple analyses by composing node attributes with reuse relations. Node attributes alone support per-instruction access-efficiency analysis (coalescing, broadcast, bank behavior). Joining node attributes with scope-labeled edges supports reuse-pair analysis: which accesses are redundant and which execution scope contains their reuse. Groups of nodes and their relations expose sharing structure: closed subgraphs identify instruction groups whose reuse relations remain within the group, which in our softmax case study separates two logical buffers within a single allocation (Section 6.1). A taxonomy of memory-behavior patterns is then expressible as mechanical rules over these attributes and relations, relating observed structure to the hardware mechanisms suited to serve it and narrowing the candidate optimizations (Section 3.3). Our tool, cuVein, makes the representation obtainable in practice: lightweight binary instrumentation and a parallel offline analyzer reconstruct PCDepGraph, without sampling, directly from fully optimized binaries, without source modification or specialized hardware support, with end-to-end overhead generally under 10×.

Motivation Example

Listing 1, the key computational kernel of the Monte Carlo neutron transport proxy XSBench [30], illustrates why reuse structure matters for optimization. Hardware counters report heavy global-memory traffic and a low L2 hit rate, a symptom with two structurally different explanations: the kernel may stream a large working set with little reuse, or substantial reuse may exist that the cache fails to retain. Our representation distinguishes them directly. Over 90% of the bytes the kernel accesses re-read data whose most recent access came from a different thread block, and the reuse volume is large relative to the accessed footprint: each interpolated field of nuclide_grids touches tens of megabytes of distinct data, yet generates repeated accesses totaling hundreds of times that footprint from across the grid. Reuse at this scope can be captured only by the L2 cache shared among all SMs, yet the touched set exceeds what L2 retains under default eviction. Marking the region for preferential L2 residency through a host-side access-policy window raises the hit rate and reduces the measured global-memory traffic by 29.13% (79.26 to 56.17 GB), without modifying the kernel. 1 __device__ void calculate_micro_xs (...) { 2 NuclideGridPoint * low , * high ; 3 // searching for target points from global nuclide_grids 4 ... 5 idx = grid_search_nuclid (...) 6 low = & nuclide_grids [... idx ...] ; 7 ... 8 high = low + 1; 9 ... 10 xs_vector [0] = high - > total_xs - f * ( high - > total_xs -

low -> total_xs ) ;

2

Reconstructing Fine-grained Memory Reuse Structure in GPU Kernels

memory with a direct-mapped, lock-free layout (Sec. 4), we reduce the tens-fold slowdown typical of shadow-memory GPU analysis to generally under 10×. Locality Metrics and Reuse Distance: Reuse distance (LRU stack distance) [11] is the canonical machine-independent locality metric on CPUs: the number of distinct elements between two accesses to the same datum directly predicts LRU cache behavior, and CUDAAdvisor [28] brings its measurement to GPUs via LLVM instrumentation. The metric presumes a single sequential access stream, in which the intervening count reflects the program’s own locality. On a GPU, tens of thousands of concurrent threads interleave into that stream, so the count is dominated by whatever unrelated accesses the parallelism happens to interpose, and, since the interleaving is decided by warp scheduling, it varies from run to run for the same kernel. Collection cost compounds the problem: a search-structure update on every access already pushes CPU tools toward sampling. These observations motivate Topological Distance, which characterizes reuse by execution scope rather than temporal distance (Section 3.1). Control/Data/Value Flow Analysis: GPU SIMT semantics and hierarchical memory spaces pose additional challenges to traditional data flow analysis. NCU [23] provides a static register dependency tree reflecting the compiler’s instruction selection and scheduling. GPA [41] uses data flow analysis over GPU PC samples to attribute instruction stalls to root causes. HPCToolkit [1] has been extended to support GPU profiling through PC sampling and binary analysis of CUDA binaries, enabling metric attribution across heterogeneous CPU–GPU contexts. GVProf [39] and ValueExpert [40] perform value-flow analysis that tracks how value patterns propagate to memory operands to identify value-related inefficiencies; foundational CPU-side predecessors include DeadSpy [6] and RedSpy [34]. Compiler transformations that exploit uniform or redundant SIMT structure [27, 32, 37] are natural downstream consumers of the reuse structure PCDepGraph exposes.

Contribution Summary. • We introduce Topological Distance and PCDepGraph, a runtime representation of GPU memory reuse structure defined over execution scopes. Its node attributes and scope-labeled reuse relations jointly express the dynamic facts that static analysis must conservatively forgo and aggregate counters cannot express. (Sections 3.1 and 3.2) • We define a taxonomy of memory-behavior patterns as mechanical rules over the representation, spanning the instruction, reuse-pair, and data-object levels and relating observed structure to candidate hardware mechanisms. (Section 3.3) • We implement cuVein, which reconstructs PCDepGraph from optimized binaries, without sampling, at generally under 10× overhead. We validate the representation for fidelity (against hardware counters and analytical microbenchmarks), stability (across runs and GPU platforms), and actionability: guided fixes eliminate up to 90.18% of memory instructions and reach up to 2.84× kernel speedup. (Sections 4 and 5)

2

RELATED WORK

Prior work on GPU performance diagnosis has mainly advanced along two complementary directions. One focuses on memory behavior analysis, characterizing inefficient accesses, redundancy, and locality loss from the perspective of memory transactions and data objects. The other emphasizes control/data/value flow analysis, explaining how dependencies, value propagation, and SIMT semantics shape performance behavior. Memory Behavior Analysis: Memory behavior recognition is a key approach for observing and diagnosing practical GPU inefficiencies, and instruction-level metrics are usually more actionable than kernel-level aggregates because developers can directly map abnormal transactions back to concrete instructions and data structures. DrGPUM [22] observes data-object level memory inefficiency patterns through a runtime map of objects. MDM [33] measures memory divergence to detect poor spatial locality under SIMT execution. DrGPU [15] adopts a top-down methodology that attributes GPU inefficiencies to hierarchical bottleneck categories. Accel-Sim [20] enables detailed modeling of GPU cache and memory behavior through simulation. Our work, by contrast, is memory-centric: it describes data flow across threads directly from observed memory reuse, without requiring hardware simulation. RedSan [38] pinpoints only redundant memory transactions. This is a narrow subset of the patterns our work detects (Sec. 3.3); more fundamentally, a redundancy verdict discards which instruction the reuse repeats and which execution-scope boundary it crosses—precisely the information that determines the appropriate fix. Our design is also faster: by replacing RedSan’s serial hashmap shadow

3

METHODOLOGY

In this section, we first define Topological Distance as an instruction-level abstraction of fine-grained memory reuse across GPU execution scopes. We then introduce PCDepGraph, which aggregates these dynamic reuse observations at the PC level into a structured representation that preserves both reuse origin and reuse scope. Building on PCDepGraph, we further define analysis attributes and representative memory-behavior patterns, showing how different access behaviors manifest as different reuse structures and efficiency implications on GPU hardware. 3.1

Topological Distance

GPU kernels execute over a hierarchical organization of threads, warps (32 threads sharing a SIMT schedule), blocks, 3

Zhao et al.

1 2 3 4 5 6 7

(e.g., thread-block clusters add a cluster level between IntraBlock and Intra-Grid); our current implementation targets the four classical scopes. Defining Topological Distance over logical scopes rather than physical residency is deliberate, for three reasons. First, it remains effective: a larger Topological Distance generally corresponds to a longer physical data-movement path, from thread-private registers, to SM-local storage, to inter-SM L2, and ultimately to off-chip global memory. Second, it improves interpretability and actionability, because thread, warp, block, and grid are stable execution scopes that programmers and compilers can reason about directly: a thread’s warp and block membership is fixed at launch, and each block is bound to a single Streaming Multiprocessor (SM) for its lifetime.1 Third, it is more reliable: physical residency of data and thread are difficult to characterize accurately, because cache-sharing effects depend on hardware management and are neither deterministic nor guaranteed by the programming model.

int gid = blockIdx . x * blockDim . x + tid ; float s = a[ gid ]; // Cold Miss s += a[ gid ]; // Intra - Thread s += a[ gid & ~31]; // Intra - Warp - Same - Inst s += a[ gid ^ 1]; // Intra - Warp - Diff - Inst s += a[ gid ^ 32]; // Intra - Block s += a[ gid ^ blockDim . x ]; // Intra - Grid

Listing 2. A minimal kernel fragment illustrating every Topological Distance level on a global array a (tid: thread index within a block). Each access is annotated with the label of (most of) its lanes, assuming a power-of-two block size of at least 64. and the kernel-wide grid, whose nested levels form natural scopes over which data reuse can arise; we adopt NVIDIA terminology throughout, other vendors exposing analogous scopes (e.g., AMD’s wavefronts and workgroups). Topological Distance is an abstraction of the fine-grained reuse cost of an individual memory instruction. For each dynamic memory instruction instance, we locate the most recent prior access to the same tracked memory address, and assign the current instruction a Topological Distance according to the smallest execution scope that contains both accesses. Intuitively, Topological Distance reflects how far data must travel through the GPU execution hierarchy to serve the current instruction from prior reuse. We define Topological Distance over logical execution scopes rather than dynamic physical residency, because logical scopes are stable, interpretable, and more actionable for optimization (discussed below). We define Topological Distance using the following six levels, which later serve as edge attributes in our PCDepGraph (Section 3.2); Listing 2 illustrates each level with a minimal annotated access sequence:

untracked memory 0x790 st.shared 4B

0x860 st.shared 4B

0x890 ld.shared 4B

0x880 ld.shared 4B

0x980 st.shared 4B

0xa30 st.shared 4B

0x720 ld.global 4B

0x6b0 ld.global 2B

0xa50 ld.global 4B

0x6f0 ld.global 4B

Group 3

• Cold Miss: the current instruction accesses a memory region for which no prior access exists in the current analysis context. • Intra-Thread: the most recent prior access to the tracked memory region comes from the same thread. • Intra-Warp-Same-Inst: the most recent prior access comes from another thread in the same warp during the same dynamic instruction issue. This case corresponds to warp-uniform accesses that can be naturally served by broadcast-style hardware behavior. • Intra-Warp-Diff-Inst: the most recent prior access comes from another thread in the same warp, but from a different dynamic instruction issue. • Intra-Block: the most recent prior access comes from a different warp in the same block. • Intra-Grid: the most recent prior access comes from a different block in the same grid.

0xa80 ld.shared 4B

0xa70 ld.shared 4B

Group 1

Group 2

0xba0 st.global 4B

intra-thread

intra-warp-same-inst

intra-grid

intra-block

intra-warp-diff-inst

cold miss

Figure 1. An example PCDepGraph from the soft_max_f32 kernel of llama.cpp/softmax. The blue circle node represents first accesses to untracked memory regions (Cold Miss); edge colors encode the Topological Distance of each reuse (see legend), and edge width is logarithmically scaled by edge weight to maintain readability over a wide dynamic range. 3.2

PCDepGraph

3.2.1 Definition. We aggregate instruction-level Topological Distance observations at the PC level into PCDepGraph, a graph structure. Intuitively, each node in PCDepGraph represents a memory-access PC, and each edge summarizes where the data used by that PC most recently came from. Formally, PCDepGraph is a directed multi-attribute graph 𝐺 = (𝑉 , 𝐸, Φ, 𝑤), where:

These levels follow the GPU execution hierarchy rather than any particular hardware generation, so scopes introduced by hardware evolution extend the taxonomy naturally

1 Rare block rebinding (e.g., under debugging or UVM page faults) does not

materially affect scope attribution.

4

Reconstructing Fine-grained Memory Reuse Structure in GPU Kernels

• 𝑉 : a set of nodes, where each 𝑣 ∈ 𝑉 represents a Program Counter (PC) corresponding to a memory instruction. Each node carries instruction attributes including access type (Read, Write, or Atomic), memory space (Shared, Global, or Local), and access size. We introduce a special node, “untracked memory addresses”, representing first accesses to memory regions and encoding the Cold Miss level of Topological Distance.2 • 𝐸: a multiset of directed edges 𝑒 = (𝑃𝐶 cur → 𝑃𝐶 prev ). An edge exists when the current access issued at 𝑃𝐶 cur touches a tracked memory region whose most recent prior access was issued at 𝑃𝐶 prev ; note that each edge points from an access back to its reuse origin, the reverse of the producer-to-consumer orientation of dataflow graphs. • Φ: a labeling function Φ : 𝐸 → 𝐿, where 𝐿 = {Thread, Warpsame_inst , Warpdiff _inst , Block, Grid, ColdMiss}. For each edge, Φ(𝑒) records one of the six levels of Topological Distance associated with the current access. • 𝑤: a weight function 𝑤 : 𝐸 → R>0 , where 𝑤 (𝑒) denotes the data volume accumulated on edge 𝑒.

graph is determined by the program’s memory access behavior and grid configuration, independent of the executing hardware up to a single schedule-dependent tie-break at scope boundaries, which we quantify in Section 5.1. PCDepGraph is intentionally designed for GPU programs rather than as a generic cross-platform memory-behavior graph: its interpretability rests on the stability of GPU execution scopes established in Section 3.1. By contrast, on CPUs, unless threads are explicitly pinned, the OS scheduler may migrate threads across physical cores, weakening fixed hardware-scope interpretations of access locality.

Figure 2. Example distributions of one 8-byte-access PC: a perfectly coalesced execution would touch 8 sectors; the collected distribution shows a scattered access pattern.

Figure 1 shows an example of PCDepGraph. The graph contains 13 memory-access PCs plus the special Cold Miss node, connected by 24 edges, 6 of them cold-miss edges. Edge colors indicate different Topological Distance labels, and edge width represents the corresponding edge weight. Self-loops are allowed, because repeated executions of the same PC may access the same tracked memory region. For a subgraph 𝐺 ′ = (𝑉 ′, 𝐸 ′ ) ⊆ 𝐺, we call 𝐺 ′ a closed subgraph if, for every node 𝑣 ∈ 𝑉 ′ , all non-Cold Miss edges of 𝑣 remain within 𝑉 ′ ; Figure 1 contains three, and reuse happens only within them. Because reuse edges follow addresses, the nodes of a closed subgraph are exactly the instructions that share a set of memory regions, in practice a data object or a tile of one. Its coldmiss edge weight is thus the object’s unique footprint, and its total edge weight the object’s actual access traffic, decomposed by reuse scope. This attribution operates on logical objects rather than allocations: the two shared-memory subgraphs in Figure 1 are two logical buffers residing in the same physical allocation (Section 6.1). Hardware counters offer no such object-level attribution, and unlike allocation-tracking profilers [22], it requires neither allocation interception nor source information. PCDepGraph turns dynamic, instruction-level reuse observations into a compact graph representation that preserves both reuse origin and reuse scope, revealing directly from observed reuse (without hardware simulation) which instructions carry reuse-intensive traffic, which are dominated by first-touch accesses, and which repeatedly obtain data from more distant execution scopes. As a result, the

3.2.2 Analysis Attributes. PCDepGraph captures reuse structure, but that alone is not sufficient for diagnosing memory efficiency: two PCs may exhibit similar reuse topology while differing substantially in coalescing quality or activelane participation. We therefore augment each PC node with dynamic attributes characterizing how efficiently the instruction executes, kept as full distributions rather than the averages hardware counters report. Dynamic access distributions. For each node 𝑣 ∈ 𝑉 , we maintain three dynamic distributions (Figure 2) collected over all executions of the corresponding PC: • Active-lane count distribution (𝐿𝑎𝑛𝑒 𝑣 ): the number of active lanes in each warp-level execution, reflecting control-flow divergence and boundary effects. • Distinct-address count distribution (𝐴𝑑𝑑𝑟 𝑣 ): the number of unique element addresses accessed per execution, capturing intra-sector duplication such as broadcast-like accesses or partial-lane contention. • Distinct-sector count distribution (𝑆𝑒𝑐𝑡𝑜𝑟 𝑣 ): the number of distinct memory sectors touched per execution, directly reflecting coalescing quality. For shared-memory PCs, these distributions are reinterpreted at bank-conflict granularity: shared memory has no sectors, so 𝐴𝑑𝑑𝑟 𝑣 becomes a distinct 32-bit-word count, and in place of 𝑆𝑒𝑐𝑡𝑜𝑟 𝑣 we record a per-instance wavefront count distribution (𝑊 𝑎𝑣𝑒 𝑣 ). For each warp-level execution, every active lane’s access footprint is mapped to its 32-bit words, duplicates are removed, each word is assigned to one of the 32 banks ((addr ≫ 2) & 31 under the 32-bank, 4byte-wide organization of NVIDIA shared memory), and the

2 Asynchronous copies (LDGSTS), which touch two memory spaces, are han-

dled in Section 4.

5

Zhao et al.

wavefront count is the maximum number of distinct words that fall on any single bank, so broadcasting one word to any number of lanes adds no cost; wide accesses that move more than 4 bytes per lane are processed by the hardware in fixed lane groups (e.g., a 128-bit store handles eight lanes per group), and their count sums the per-group maxima. This distribution underlies the Bank Conflict pattern (Section 3.3). Together, these distributions distinguish whether low memory efficiency arises from scattered addressing, intrasector duplication, or limited lane utilization. The lane, sector, and address attributes are collected on-the-fly by warp-level primitives in the Trace Collector, while the wordgranularity wavefront counts are derived by the Offline Analyzer from the recorded per-lane addresses (Section 4). Weighted Out-Degree. For a node 𝑣 ∈ 𝑉 , let 𝐸𝑜𝑢𝑡 (𝑣) and 𝐸𝑖𝑛 (𝑣) denote the sets of outgoing and incoming edges of 𝑣, respectively. For a Topological Distance label 𝑙 ∈ L, we define the Weighted Out-Degree of 𝑣 at scope 𝑙 and the Total Weighted Out-Degree of 𝑣 as Í W𝑜𝑢𝑡 (𝑣, 𝑙) = 𝑒 ∈𝐸𝑜𝑢𝑡 (𝑣), Φ(𝑒 )=𝑙 𝑤 (𝑒), Í W𝑡𝑜𝑡𝑎𝑙 (𝑣) = 𝑙 ∈ L W𝑜𝑢𝑡 (𝑣, 𝑙).

match its scope leaves performance on the table, and it is precisely such mismatches that the inefficient patterns below diagnose. Together, they provide a practical vocabulary for interpreting PCDepGraph and connecting observed reuse structure to likely optimization implications. I. Block/Grid level. Hierarchical Reuse. For a subgraph 𝐺 ′ ⊆ 𝐺, the reuse ratio at scope 𝑙 is Í 𝑣 ∈𝑉 (𝐺 ′ ) W𝑜𝑢𝑡 (𝑣, 𝑙) 𝑟 (𝑙) = Í , c 𝑣 ∈𝑉 (𝐺 ′ ) W (𝑣) c (𝑣) is W𝑡𝑜𝑡𝑎𝑙 (𝑣) except its 𝑊 𝑎𝑟𝑝𝑠𝑎𝑚𝑒_𝑖𝑛𝑠𝑡 and Coldwhere W Miss components, to exclude non-reused and broadcast volumes (Section 3.1). We say that 𝐺 ′ exhibits a dominant reuse pattern at scope 𝑙 if 𝑟 (𝑙) exceeds a threshold 𝛿 (e.g., 𝛿 = 0.8 in our experiments). When 𝑙 = 𝐺𝑟𝑖𝑑, the subgraph is dominated by intra-grid reuse, indicating reuse across block boundaries and potential pressure on the L2 cache. When 𝑙 ∈ {𝑇ℎ𝑟𝑒𝑎𝑑,𝑊 𝑎𝑟𝑝𝑑𝑖 𝑓 𝑓 _𝑖𝑛𝑠𝑡 , 𝐵𝑙𝑜𝑐𝑘 }, the subgraph is dominated by reuse within a block, suggesting potential opportunities for shared-memory promotion. II. Warp level. Poor Coalescing. For a node 𝑣, let 𝐸 (𝑆𝑒𝑐𝑡𝑜𝑟 𝑣 ) denote the weighted average distinct sector count derived from the distribution 𝑆𝑒𝑐𝑡𝑜𝑟 𝑣 defined in Section 3.2. The theoretical optimal sector count is   𝐿𝑎𝑛𝑒 𝑣 · AccSize 𝑆𝑜𝑝𝑡 (𝑣) = , SecSize where 𝐿𝑎𝑛𝑒 𝑣 is the active-lane count, AccSize is the accesssize attribute of the node in PCDepGraph, and SecSize is the hardware sector size (32 bytes on NVIDIA GPUs). We quantify coalescing quality as 𝐸 (𝑆𝑒𝑐𝑡𝑜𝑟 𝑣 ) 𝑅𝑐𝑜𝑎𝑙 (𝑣) = ; 𝑆𝑜𝑝𝑡 (𝑣)

Intuitively, W𝑜𝑢𝑡 (𝑣, 𝑙) measures how much data volume issued by PC 𝑣 obtains its most recent reuse from scope 𝑙, while W𝑡𝑜𝑡𝑎𝑙 (𝑣) measures the total outgoing data volume associated with 𝑣 across all Topological Distance labels. We focus on outgoing edges because, under the current-to-previous edge direction in PCDepGraph, they directly characterize the reuse sources of accesses issued by PC 𝑣. Reuse Intensity. Beyond per-node access-efficiency attributes, we define a subgraph-level metric, Reuse Intensity (𝑅𝐼 ), to summarize how strongly a set of PCs depends on reused data rather than first-touch accesses. For a closed subgraph 𝐺 ′ ⊆ 𝐺 with edge multiset 𝐸 (𝐺 ′ ), we define Í 𝑒 ∈𝐸 (𝐺 ′ ) 𝑤 (𝑒) ′ 𝑅𝐼 (𝐺 ) = Í . 𝑒 ∈𝐸 (𝐺 ′ ), Φ(𝑒 )=ColdMiss 𝑤 (𝑒)

𝑅𝑐𝑜𝑎𝑙 (𝑣) > 1 means the active lanes touch more sectors than theoretically necessary (sparse or unaligned accesses that leave fetched sector data unused); in the worst case, each lane touching a different sector, the ratio approaches SecSize/AccSize (8 for 4-byte accesses). Broadcasting/Multicasting/Atomic Contention. A node 𝑣 is classified as Broadcasting when it satisfies: (1) 𝑇𝑦𝑝𝑒 (𝑣) = Read, (2) 𝐿𝑎𝑛𝑒 𝑣 > 1, and (3) Mode(𝐴𝑑𝑑𝑟 𝑣 ) = 1, i.e., multiple active lanes read the same address in the same instruction instance, a warp-level broadcast that hardware mechanisms such as the constant-cache path serve efficiently. If instead Mode(𝐴𝑑𝑑𝑟 𝑣 ) ∈ (1, 𝐿𝑎𝑛𝑒 𝑣 ), 𝑣 is classified as Multicasting, or as Atomic Contention if 𝑇𝑦𝑝𝑒 (𝑣) = Atomic.

Intuitively, 𝑅𝐼 increases when a subgraph has substantial reuse traffic relative to cold-miss footprint. Under the datacentric interpretation of closed subgraphs above, 𝑅𝐼 (𝐺 ′ ) is the traffic amplification of the corresponding data object relative to its unique footprint, making it particularly useful for identifying hot memory regions. 3.3

Identifiable Memory Behavior Patterns

In this section, we present representative memory-behavior patterns identifiable in PCDepGraph and organize them by their most relevant execution scope: Thread, Warp, Block/ Grid, and Others. These patterns cover both common memory behaviors and potentially inefficient ones. Each scope has hardware mechanisms suited to serving its reuse (registers for intra-thread, warp shuffles for intra-warp, shared memory or the SM-local L1 for intra-block, and the SM-shared L2 for intra-grid); reuse served by a mechanism that does not 6

Reconstructing Fine-grained Memory Reuse Structure in GPU Kernels Worker pool

GPU binary

Bank Conflict. The hardware serializes a shared-memory instruction instance with wavefront count 𝑘 (the 𝑊 𝑎𝑣𝑒 𝑣 attribute of Section 3.2) into 𝑘 shared-memory cycles. We define the conflict level of a shared-memory node 𝑣 as 𝐸 (𝑊 𝑎𝑣𝑒 𝑣 ) 𝑅𝑐𝑜𝑛𝑓 (𝑣) = , 𝐸 (𝑊 𝑎𝑣𝑒 𝑖𝑑𝑒𝑎𝑙 ) 𝑣 where the ideal count of an instance, its bandwidth floor, is ⌈footprint words/32⌉ accumulated over the same lane groups. A conflict-free access pattern yields that 𝑅𝑐𝑜𝑛𝑓 (𝑣) = 1; the excess over 1 measures pure serialization waste, approaching the warp width in the worst case (all active lanes addressing distinct words in a single bank). The computed wavefronts are validated against the hardware’s shared-memory wavefront counters in Section 5.1. III. Thread level. Redundant Access. A Redundant Access is defined on two different nodes. For an edge 𝑒 = (𝑢 → 𝑣) in PCDepGraph, 𝑢 denotes the current PC and 𝑣 denotes the most recent prior PC. We identify the following three cases: •RAR: ∀𝑒 ∈ 𝐸𝑜𝑢𝑡 (𝑢), if Φ(𝑒) = 𝑇ℎ𝑟𝑒𝑎𝑑, 𝑇𝑦𝑝𝑒 (𝑢) = 𝑅𝑒𝑎𝑑, and 𝑇𝑦𝑝𝑒 (𝑣) = 𝑅𝑒𝑎𝑑, then 𝑢 is a redundant read: it re-reads values already read by the same thread. •RAW: ∀𝑒 ∈ 𝐸𝑜𝑢𝑡 (𝑢), if Φ(𝑒) = 𝑇ℎ𝑟𝑒𝑎𝑑, 𝑇𝑦𝑝𝑒 (𝑢) = 𝑅𝑒𝑎𝑑, and 𝑇𝑦𝑝𝑒 (𝑣) = 𝑊 𝑟𝑖𝑡𝑒, then 𝑢 is a redundant read: it re-reads values already produced by the same thread. •WAW: ∀𝑒 ∈ 𝐸𝑖𝑛 (𝑣), if Φ(𝑒) = 𝑇ℎ𝑟𝑒𝑎𝑑, 𝑇𝑦𝑝𝑒 (𝑢) = 𝑊 𝑟𝑖𝑡𝑒, and 𝑇𝑦𝑝𝑒 (𝑣) = 𝑊 𝑟𝑖𝑡𝑒, then 𝑣 is a redundant write: every value it produces is overwritten by the same thread with no reuse in between, making 𝑣 a dead store. A redundant PC indicates a strong candidate for REG promotion, but the REG budget should be considered to avoid decreasing occupancy. Abuse SMEM. A closed subgraph 𝐺 ′ is identified as exhibiting only thread-private SMEM use when all nodes in 𝐺 ′ are shared-memory instructions and all edges in 𝐺 ′ are intrathread reuse edges. This pattern indicates that the data placed in shared memory is reused only within individual threads, without inter-thread sharing in the subgraph. IV. Others. Streaming Access. A node 𝑣 is identified as exhibiting the Streaming Access pattern when its outgoing data flow in PCDepGraph is dominated by ColdMiss edges, i.e., when the ratio W𝑜𝑢𝑡 (𝑣, ColdMiss) / W𝑡𝑜𝑡𝑎𝑙 (𝑣) is sufficiently large (e.g., above 0.95 in our experiments): first-touch traffic with negligible temporal locality. Its sector-count distribution is typically close to 𝑆𝑜𝑝𝑡 (𝑣), indicating spatially dense, wellcoalesced accesses. Small, Read-Only Broadcast Region. A subgraph 𝐺 ′ ⊆ 𝐺 is identified as exhibiting the Broadcasting Constant Region pattern when it satisfies the following conditions in PCDepGraph: (1) every node 𝑣 ∈ 𝑉 (𝐺 ′ ) exhibits the Broadcasting pattern, and (2) the total cold-miss footprint of the subgraph fits within the capacity of the GPU read-only constant-cache

dispatch

patch signals

Trace Collector

trace buffers

Event Monitor

events trace batches

reduce

Offline Analyzer

JSON reports

GUI

Figure 3. The overview of cuVein workflow. path. Formally, 𝐺 ′ must satisfy the capacity constraint Í 𝑣 ∈𝑉 (𝐺 ′ ) W𝑜𝑢𝑡 (𝑣, ColdMiss) ≤ C𝑐𝑜𝑛𝑠𝑡 , where C𝑐𝑜𝑛𝑠𝑡 is the effective capacity of the constant-cache hierarchy; the unique footprint of 𝐺 ′ is exactly its coldmiss traffic. Such a subgraph is a strong candidate for __constant__ memory promotion, most beneficial when 𝑅𝐼 (𝐺 ′ ) is high.

4

THE cuVein FRAMEWORK

cuVein is a practical, general-purpose framework that reconstructs the Topological Distance and PCDepGraph representations of Section 3 from real-world GPU applications: it directly instruments fully optimized GPU binaries at runtime, collects fine-grained memory traces, and performs efficient offline analysis, requiring neither operating-system support nor source-level modifications. Figure 3 illustrates the end-to-end workflow of cuVein, which consists of four core components: Trace Collector, Event Monitor, Offline Analyzer, and GUI. Trace Collector patches fully optimized CUBINs at module load with SASS-level callbacks that collect runtime metrics and per-lane addresses on every memory instruction issue. We use the NVIDIA Compute Sanitizer API [24] because it is a maintained CUDA-toolkit component that reports memory space (global, shared, or local) and access type at the callback boundary. Any instrumentation framework exposing the same facts (PC, address, access size, memory space, access type, block/warp/lane identity, and active mask), such as NVBit [31], could feed the same analysis pipeline. The trace is organized at warp granularity, each buffer corresponding to one instruction instance of a single warp. Buffers are aggregated into batches and transferred to system memory for analysis through a producer–consumer queue between the GPU and CPU. The Trace Collector also supports asynchronous globalto-shared copies (cp.async, SASS LDGSTS), which modern kernels increasingly rely on for software pipelining. A single such instruction performs both a global read and a shared write, which would otherwise break the single-memoryspace property of PCDepGraph nodes. cuVein therefore splits each async-copy instance into two trace records attributed to synthetic PCs 𝑝𝑐 and 𝑝𝑐 + 1: since SASS instruction addresses are 16-byte aligned, the low bits are free to distinguish the global-read half from the shared-write half. Both records are tagged as async copies, so the global→shared staging dataflow of pipelined kernels appears explicitly in 7

Zhao et al.

PCDepGraph: subsequent shared-memory reads form reuse edges to the shared-write half. Event Monitor handles driver events: module loading triggers instrumentation of whitelisted kernels; GPU memory (de)allocation triggers corresponding shadow-memory (de)allocation; and buffer-full notifications trigger trace transfer to system memory and analyzer invocation. It also tracks kernel launches and completion, exporting each completed kernel’s profile as JSON keyed by a unique per-launch ID. Offline Analyzer maintains GMEM and SMEM shadowmemory state on the host and reconstructs each kernel’s PCDepGraph from the collected accesses. Constructing an instruction-level, full-lifecycle dependency graph incurs substantial time and memory overhead: the last-access state must be consulted and updated on every dynamic memory access, over footprints that scale with the application’s allocations. cuVein keeps this high-fidelity analysis practical through three system-level designs. (1) Direct-mapped, demand-paged shadow memory with an interleaved layout. A hashmap keyed by address suffers from indirection, collisions, and rehashing that dominate at large footprints [38]. Instead, cuVein reserves a byte-wise host shadow region for every GPU allocation via mmap, making each lookup a single O(1) index computation. Demand paging commits physical pages only when the corresponding shadow bytes are first touched, and resetting a region is a page-table-level madvise rather than an O(size) clear. Tracking each access’s metadata at its start address and every fourth byte thereafter suffices; interleaving the region into four byte-lane partitions (same-offset bytes packed contiguously) keeps these entries dense in host caches and commits physical pages in only one partition for aligned accesses. (2) Compact lock-free shadow entries. Each shadow entry packs the identity of the last access (flat block/warp/lane ID, PC offset, and a kernel-generation tag) into a single 64-bit word. A single atomic exchange then retrieves the previous access and installs the current one in one step, so concurrent workers never block on locks; the generation tag invalidates stale entries across kernel launches implicitly, so no perlaunch reset of the shadow region is needed. (3) Contention-free parallel workers. Trace batches are partitioned by CTA ID across a pool of persistent worker threads. Each worker holds private analysis state, including graph-edge statistics, per-PC attributes, and CTA-local shared-memory shadow objects, so per-access statistics updates never contend, and shared-memory tracking requires no atomics at all because each CTA is processed by exactly one worker. Per-worker results are merged in a short post-batch sequential reduction, keeping all synchronization off the hot path. Together, these designs keep the end-to-end overhead generally under 10× on production workloads (Section 5.3). GUI is a self-contained browser application that renders the JSON reports interactively at PC granularity; an optional

Kernel Base PC input converts PC offsets into absolute instruction addresses for direct cross-reference with NCU or nvdisasm output. Diagnostic Workflow. cuVein is designed to complement hardware profilers such as Nsight Compute (NCU). In the recommended workflow, developers first use NCU to identify bottleneck kernels and confirm memory-bound behavior through hardware counters (cache hit rates, stall breakdowns, throughput metrics). cuVein is then applied to the suspicious kernel to diagnose the root cause: its PCDepGraph reveals which PCs are responsible, what data reuse patterns exist, and at which execution scope the reuse occurs. After applying an optimization, developers return to NCU to validate whether the hardware-level metrics have improved. Our implementation is open-source, and its overhead characteristics have been independently reproduced by external users across multiple GPU platforms. It is available at https://github.com/AccelProf/AccelProf.

5

EVALUATION

We evaluate cuVein on the two GPU workstations listed in Table 1, with an RTX 4090 (sm_89) and an H100 (sm_90a). The evaluation covers applications and benchmarks from multiple domains: Rodinia [7], XSBench [30], LULESH [18], llama.cpp [13] (commit a8304b4d), and PASTA [21]. Kernels with different template parameters or launch configurations are reported separately, since these choices can lead to substantially different performance behavior. All reported inefficiencies were manually confirmed and optimized. 5.1

Representation Fidelity and Stability

Pattern-recognition accuracy via microbenchmarks. We first validate that cuVein recognizes every memory behavior of Section 3.3, using 20 microbenchmark kernels that each implant one behavior whose PCDepGraph is analytically derivable, covering the six Topological Distance levels, the three redundancy variants, and known conflict and coalescing degrees. Every reported graph matches the derived expectation on schedule-invariant aggregates (edge labels and endpoint types, per-scope traffic volumes, sector and wavefront counts). Because the pattern rules are deterministic functions of these aggregates, every implanted behavior is classified correctly. Cross-validation against hardware counters. By design, every executed warp-level memory instruction produces exactly one trace record with no sampling (two for asynchronous copies, one per memory space, matching how hardware counters attribute them; Section 4), so the counts implied by PCDepGraph are directly comparable against independent hardware counters. We run every kernel of the paper’s 19 applications (61 kernels, a strict superset of those optimized in Table 2) under both cuVein and NCU with identical inputs, excluding two spin-polling kernels whose load counts 8

Reconstructing Fine-grained Memory Reuse Structure in GPU Kernels

Table 1. Platform Configuration. CPU AMD 5955WX (16c) Intel(R) Xeon 8468 (48c)

GPU RTX 4090 H100 SXM

#SM GPU Memory 128 24 GB GDDR6X 132 80 GB HBM3

System Linux 5.15.0 Linux 5.15.0

are not reproducible even across NCU runs. On 58 of the remaining 59 kernels, every cuVein-derived traffic count (global load/store requests, 32-byte sectors, shared-memory instructions, and atomic requests) agrees with NCU to the last count, and the last kernel deviates by at most 0.03%; shared-memory wavefront predictions agree within 2.5% on all 34 shared-memory kernels, exactly on 19. Since replaybased counter collection has small error sources of its own, we read this agreement as mutual corroboration rather than either side serving as absolute ground truth. Counters corroborate only aggregate volumes; the reuse edges and scope labels PCDepGraph adds are covered by the microbenchmarks above, and since edge weights accumulate from the same per-access records whose totals match NCU , the scopedecomposed traffic inherits this fidelity. Cross-platform stability. Section 3.2 argues that PCDepGraph is determined by the program and its launch configuration, up to a single schedule-dependent tie-break at scope boundaries. We test this on both platforms of Table 1: for each kernel and launch configuration run with identical inputs, we compare the byte-weighted distribution of traffic across the six reuse scopes by total variation distance (TVD). Across the 48 configurations with identical executed code and inputs, the distributions are identical on 19 and within 0.1% TVD on 42; the mean is 0.08% and the maximum 1.75%. Every deviation above 0.1% is a compensating shift between adjacent scopes (PASTA’s mttkrp relabels 1.75% of its bytes from intra-grid to intra-block), exactly the predicted tie-break: warp scheduling can change which samedata candidate is the temporally nearest reuse, and thus the boundary label, but not the volume of reuse. The remaining differences track genuine platform differences: llama.cpp’s parallelize_cols launches one block per SM, so its grid grows from 128 to 132 blocks, yet its signature matches within 0.1%; a CUB-based cumsum kernel selects different code on sm_89 and sm_90a and yields a correspondingly different signature (6% TVD), which cuVein reports faithfully. 5.2

GPU Driver CUDA Version GCC Compiler Options 580.65.06 13.0 11.4.0 -lineinfo -O3 -arch sm_89 550.163.01 12.4 11.4.0 -lineinfo -O3 -arch sm_90a

primary validation criterion is structural as well: whether the guided fix removes exactly the waste the graph attributes, on the meter of the mechanism involved. We therefore measure each fix by its reduction in memory-instruction count for register-serviceable reuse, in shared-memory wavefronts for bank conflicts, and in global-memory (L2–DRAM) traffic for grid-scope contention; these quantities are deterministic, independent of warp scheduling, and verifiable against hardware counters at the exactness established in Section 5.1. Kernel speedup, measured as the ratio of kernel cycles (Table 2), is reported alongside as corroboration, but it is an incomplete indicator of a memory fix in isolation: the removed operations may lie off the critical path or be covered by latency hiding and concurrent instruction streams, so a structurally confirmed fix can leave runtime nearly unchanged on one platform while its eliminated traffic still matters for contention, energy, and other occupancy regimes. As summarized in Table 2, the guided fixes eliminate up to 90.18% of memory instructions across the table’s six optimization categories (28.7% on average over the memory-instruction entries of both platforms, excluding the wavefront- and traffic-validated rows), with kernel speedups reaching 2.84×. The two exceptions follow their own meters: the bank-conflict fixes restructure shared-memory addressing rather than remove instructions, and their starred cells report wavefront reductions of 41–70%; XSBench’s dagger-marked fix cuts global-memory traffic by 29.13% while leaving instruction counts essentially unchanged. 5.3

Overhead

We also report the end-to-end overhead of cuVein on the RTX 4090 platform. No sampling is used: every instrumented memory instruction is traced and analyzed, so the reported overhead reflects the full cost of exact dependency reconstruction. As shown in Table 3, cuVein remains practical for high-fidelity analysis on real workloads, with overhead generally under 10× (median 3.2×). For calibration, RedSan, which tracks the same last-access state through a serial hashmap, reports 34.36× on rodinia/backprop where cuVein incurs 4.80× (Section 2). The two major outliers are rodinia/srad_v1 (53.75×), which issues 502 launches of six small kernels so instrumentation dominates, and XSBench (25.28×), whose memory-instruction volume implies a pre-cache analysis footprint of about 415 GB. A kernel whitelist lets users bound analysis effort on such complex applications. We further evaluate the scalability of the parallel offline analyzer under different worker-thread counts using two representative workloads: XSBench (input -s large -m event -l 17000000), whose traffic is global-memory only, and a

Actionability of PCDepGraph

PCDepGraph provides a logical view of memory behavior, from which some observed patterns can be identified as inefficient. However, not all such inefficiencies are equally actionable: some may be largely absorbed by hardware mechanisms or compiler optimizations, while others remain performancecritical in the final execution. Across diverse benchmarks and real applications, we use cuVein to locate these actionable inefficiency patterns and validate targeted fixes with only minor code modifications, typically fewer than 10 lines. Because the graph diagnoses structural mismatch (reuse served by a mechanism that does not match its scope), the 9

Zhao et al.

Table 2. Kernel and memory performance comparison on RTX 4090 and H100. Speedup is original over optimized kernel cycles (frequency-independent), averaged over 10 Nsight Compute runs; memory instruction reduce counts load, store, and atomic instructions across local, shared, and global memory. XSBench’s grid-scope contention exceeds the H100 L2 (50 MB); its fix is instead evaluated on an RTX 5090 (96 MB L2), with speedup and traffic reduction similar to the RTX 4090 results reported here. ∗ Bank-conflict fixes do not remove memory instructions (counts change by at most 0.39%); starred cells report the reduction in serialized shared-memory wavefronts instead. † Fixes targeting grid-scope contention are validated by global-memory (L2–DRAM) traffic; XSBench’s drops by 29.13%. Application

rodinia/backprop rodinia/bfs rodinia/hotspot rodinia/huffman rodinia/lud rodinia/nw

rodinia/srad_v1

Inefficiency Pattern

bpnn_layerforward_CUDA bpnn_adjust_weights_cuda Kernel calculate_temp vlc_encode_kernel_sm64huff prescan<(bool)1, (bool)0> lud_diagonal lud_perimeter needle_cuda_shared_1 needle_cuda_shared_2 reduce srad srad2

redundant redundant L2 contention / redundant abuse SMEM bank conflict bank conflict bank conflict bank conflict bank conflict bank conflict redundant constant constant

1.02× 1.55× 1.08× 1.04× 1.04× 1.01× 1.04× 1.15× 1.07× 1.07× 2.62× 1.01× 0.97×

12.03% 33.33% 7.56% 31.65% 50.37%∗ 42.60%∗ 41.35%∗ 69.83%∗ 68.85%∗ 68.85%∗ 90.18% 14.28% 9.09%

0.98× 1.03× 1.07× 1.05× 1.01× 1.02× 1.03× 1.14× 1.04× 1.04× 2.08× 1.00× 1.00×

12.03% 33.33% 7.03% 31.65% 50.32%∗ 41.80%∗ 41.35%∗ 69.83%∗ 68.85%∗ 68.85%∗ 90.18% 14.29% 9.09%

redundant redundant redundant redundant redundant redundant redundant / abuse SMEM redundant / abuse SMEM redundant / abuse SMEM redundant / abuse SMEM abuse SMEM abuse SMEM constant redundant redundant

0.99× 1.08× 1.01× 1.01× 1.00× 1.07× 1.03× 1.06× 1.05× 1.05× 1.03× 0.97× 1.00× 1.01× 1.22×

21.14% 20.25% 15.13% 9.09% 40.00% 22.53% 33.33% 33.33% 34.29% 34.29% 16.67% 20.83% 29.58% 26.36% 62.07%

1.05× 1.06× 1.02× 1.00× 0.99× 1.09× 1.00× 1.02× 1.00× 1.02× 1.01× 0.97× 1.01× 1.13× 1.10×

21.14% 20.25% 15.13% 9.09% 40.00% 22.53% 33.33% 33.33% 34.29% 34.29% 16.67% 20.00% 29.58% 51.71% 62.07%

cumsum_kernel rms_norm_f32<(int)1024, (bool)0, (bool)0>≪(5,4,3), (1024,1,1)≫ llama.cpp/rms_norm rms_norm_f32<(int)256, (bool)0, (bool)0>≪(2,2,1),(256,1,1)≫ rms_norm_f32<(int)256, (bool)0, (bool)0>≪(5,4,3),(256,1,1)≫ rms_norm_back_f32<(int)32> llama.cpp/rms_norm_back rms_norm_back_f32<(int)1024> soft_max_f32<(bool)1, (int)1024, (int)1024, __half>≪(16,1,1),(1024,1,1)≫ soft_max_f32<(bool)1, (int)1024, (int)1024, __half>≪(1024,1,1),(1024,1,1)≫ soft_max_f32<(bool)1, (int)1024, (int)1024, float>≪(16,1,1),(1024,1,1)≫ llama.cpp/soft_max soft_max_f32<(bool)1, (int)1024, (int)1024, float>≪(1024,1,1),(1024,1,1)≫ soft_max_f32<(bool)1, (int)32, (int)32, float>≪(2,32,1),(32,1,1)≫ soft_max_f32<(bool)1, (int)32, (int)32, __half>≪(2,32,1),(32,1,1)≫ ApplyMaterialPropertiesAndUpdateVolume_kernel LULESH CalcTimeConstraintsForElems_kernel CalcMinDtOneBlock llama.cpp/cumsum

XSBench

xs_lookup_kernel_baseline

L2 contention

1.03×

0%†

–

–

PASTA

spt_MTTKRPKernelRankSplitNnz4DOneKernel spt_TTMRankRBNnzKernelSM

L1 contention / constant L1 contention / abuse SMEM

1.00× 2.84×

27.53% 32.04%

1.00× 1.67×

27.53% 32.04%

Table 3. Overhead comparison between the original and cuVein-enabled executions on the RTX 4090 platform in Table 1 with 32 worker threads.

cuBLAS 4096 × 4096 GEMM, where 82% of the accesses target shared memory. As shown in Figure 4, the analyzer runs 9.1× and 13.7× faster than the single-worker configuration at 32 threads, respectively. Notably, the host CPU has only 16 physical cores (Table 1), so the 16→32-thread step runs entirely on SMT yet still gains another 30% and 43%; the 13.7× speedup corresponds to 86% of the physical core count.

Normalized runtime

RTX 4090 H100 Kernel Mem Inst. Kernel Mem Inst. Speedup Reduce Speedup Reduce

Kernel Name

1 1/2 1/4 1/8 1/16 1/32

Ideal linear scaling XSBench cuBLAS GEMM

1

2

4 8 16 Worker threads

9.1× 13.7×

Application Name

origin (ms)

cuVein_enabled (ms)

Overhead

rodinia/backprop rodinia/bfs rodinia/hotspot rodinia/srad_v1 llama.cpp/cumsum llama.cpp/rms_norm llama.cpp/rms_norm_back llama.cpp/soft_max XSBench LULESH pasta/mttkrp pasta/ttm_hicoo

204 2912 314 206 460 511 433 4824 4432 395 2080 3492

979 9408 1013 11073 1445 1417 1212 10501 112042 3809 13592 5094

4.80× 3.23× 3.23× 53.75× 3.14× 2.77× 2.80× 2.18× 25.28× 9.64× 6.53× 1.46×

32

6

CASE STUDIES AND DISCUSSION

This section demonstrates PCDepGraph as an analysis substrate rather than a fixed report: a recurring redundancy diagnosed and fixed in real applications (Section 6.1), a path from

Figure 4. Analyzer runtime with different worker-thread counts, normalized to one worker, for XSBench (GMEMonly) and cuBLAS GEMM (SMEM-heavy). 10

Reconstructing Fine-grained Memory Reuse Structure in GPU Kernels

example). Because the compiler cannot determine how many elements each thread will process, it may conservatively generate redundant loads or spill values to local memory, even when register pressure is not the limiting factor. This scope misplacement is not specific to AI libraries. In the Rodinia suite, hotspot allocates the shared-memory objects power_on_device and temp_t that are each touched by only one thread, and backprop stages weights in a sharedmemory array that each thread writes and then immediately reads back itself.

1 template <... int ncols_template ... > 2 static __global__ void soft_max_f32 (...) { 3 const int ncols = ncols_template == 0 ? p . ncols : 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 }

ncols_template ; float * vals = use_shared ? buf_iw + WARP_SIZE : dst ; ... for ( int col0 = 0; col0 < ncols ; col0 += block_size ) { const int col = col0 + tid ; // col is thread unique const float val = x [ col ]* p . scale + ...; vals [ col ] = val ; max_val = max ( max_val , val ) ; } max_val = block_reduce < MAX ,... >( max_val , buf_iw ) ; for (...) { // same thread - unique col indexing const float val = expf ( vals [ col ] - max_val ) ; tmp += val ; vals [ col ] = val ; } tmp = block_reduce < SUM ,... >( tmp , buf_iw ) ; for (...) dst [ col ] = vals [ col ] * inv_sum ;

6.2

While cuVein currently informs manual restructuring, nothing about consuming PCDepGraph is inherently manual. The patterns in Section 3.3 are mechanical predicates over the graph, and several of the transformations they suggest (register promotion for redundant accesses, __constant__ promotion for broadcast regions, shared-memory demotion for thread-private SMEM) are mechanical. A natural next step is to close the loop: treat PCDepGraph as a memoryplacement profile fed back into compilation, in the same way edge and path profiles guide code layout and inlining in classical profile-guided optimization. Because the representation is stable across GPU targets, such profiles are reusable across hardware generations; and because it is extracted from optimized binaries, the feedback reflects the code the compiler emitted rather than a source-level approximation. The same information could also flow downward to hardware. Semantic-aware access annotations (e.g., threadunique, warp-broadcast, streaming) derived from PCDepGraph could be lowered by the compiler onto specialized hardware paths (constant caches, L2 eviction policies, or automated shared-memory mapping), replacing the conservative assumptions that today lead to cache thrashing and wasted bandwidth.

Listing 3. A softmax kernel listing showing max-reduction, exponentiation, and normalization. these developer-driven analyses to profile-guided optimization (Section 6.2), and a discussion of what PCDepGraph carries beyond performance, and its limitations (Section 6.3). 6.1

Toward Profile-Guided Memory Optimization

A Common Redundancy Across AI Libraries

During evaluation of cuVein, we identified a recurring redundancy pattern in multi-pass operators (softmax, layer_norm, and rms_norm, etc.) across multiple AI libraries (including llama.cpp [13], ONNXruntime [10], PyTorch [25], llm.c [19], and Triton [29] tutorial). These multi-pass operators typically require at least two scans of the same input: the first pass computes a global statistic (e.g., maximum) and broadcasts it to threads in the same block, while the second pass re-scans the input for final outputs using that statistic. In Listing 3, when the template parameter use_shared is true, the vals pointer is placed in shared memory but accessed in a thread-private order. Figure 1 in Section 3.2 is the PCDepGraph of this kernel, and the misuse is legible as graph structure: the four SHARED PCs accessing vals form the leftmost closed subgraph (Group 1), and every edge inside it is intra_thread: each element staged in shared memory is written and read back only by the thread that produced it. The same figure supplies the control case: the second closed subgraph (Group 2), formed by the broadcast buffer buf_iw, carries the intra_block and intra_warp edges that justify block-scope storage. Both subgraphs live in the same shared-memory allocation (vals is buf_iw + WARP_SIZE), yet only one exhibits cross-thread reuse; this distinction is invisible to bandwidth- or transaction-level metrics and is precisely the scope information PCDepGraph attaches to every reuse pair. Thread-private data held in block-scope storage forfeits the core purpose of shared memory (low-latency, software-managed sharing within a thread block), introduces redundant memory instructions, and reduces the available L1 capacity, since shared memory is carved out of the same physical array. Another important limitation leading to redundancy comes from the statically unknown loop bound (ncols in this

6.3

Discussion

Beyond performance. Although PCDepGraph is designed for a performance tool, its edges carry correctness information. A cross-thread reuse edge with a write on either endpoint is a pair of conflicting accesses, and its Topological Distance names the synchronization scope that must order the pair: any such edge not separated by a barrier of at least that scope is a potential race. Race candidates thus become an offline query over the same profile. The graph in Figure 1 contains one: the cross-thread write-after-read edges from 0xa30 to 0x880 in Group 2 record a block-level reduction overwriting the broadcast buffer a previous reduction is still reading, with no intervening __syncthreads(); the race is genuine and was later fixed upstream (commit 9bd4c09e). Because the store’s shadow comparison meets the earlier read under every interleaving, the pair is captured on every run rather than only under a racy schedule, already attributed to both PCs and their scope. The same rule read in reverse 11

Zhao et al.

degenerates to an inefficiency check: a barrier spanned by no cross-thread reuse edge orders nothing, which is how PCDepGraph exposes the removable __syncthreads() in rodinia/backprop (Table 2). We deliberately stop short of building a race checker; dedicated GPU race detectors address the general problem [12, 17]. The observation is narrower: the conflicting-pair structure that constitutes the expensive half of such analyses is already present in the profile cuVein extracts once, so extending the analysis surface means writing an offline query, not new instrumentation. Limitations. As a dynamic analysis, cuVein inherits the classical boundaries of profiling. First, the reconstructed PCDepGraph reflects the executed control flow and data distribution: reuse behavior that a given input does not exercise remains invisible, and strongly input-dependent kernels should be profiled under representative inputs. Second, the analysis is diagnostic rather than prescriptive: the patterns of Section 3.3 name the responsible instructions, scopes, and applicable hardware mechanisms, but applying a fix still involves developer judgment (e.g., weighing register pressure or occupancy), and closing this loop automatically is the profile-guided direction outlined above. Third, while the representation itself is vendor-neutral, our implementation targets NVIDIA GPUs through the Compute Sanitizer API; porting requires a substrate that reports the same per-lane facts (Section 4). Finally, tracing every access has a cost that grows with dynamic memory-instruction count: most workloads stay under 10×, but launch-heavy or extremely instruction-heavy applications reach 25–54× and rely on the kernel whitelist to bound analysis effort (Section 5.3).

7

[2] Amey Agrawal, Nitin Kedia, Ashish Panwar, Jayashree Mohan, Nipun Kwatra, Bhargav S. Gulavani, Alexey Tumanov, and Ramachandran Ramjee. 2024. Taming throughput-latency tradeoff in LLM inference with sarathi-serve. In Proceedings of the 18th USENIX Conference on Operating Systems Design and Implementation (Santa Clara, CA, USA) (OSDI’24). USENIX Association, USA, Article 7, 18 pages. [3] AMD.Corp. 2025. RocProfiler. https://rocm.docs.amd.com/projects/ rocprofiler/en/latest/ [4] Cesar Avalos Baddouh, Mahmoud Khairy, Roland N Green, Mathias Payer, and Timothy G Rogers. 2021. Principal kernel analysis: A tractable methodology to simulate scaled gpu workloads. In MICRO-54: 54th Annual IEEE/ACM International Symposium on Microarchitecture. 724–737. [5] Shiyi Cao, Shu Liu, Tyler Griggs, Peter Schafhalter, Xiaoxuan Liu, Ying Sheng, Joseph E. Gonzalez, Matei Zaharia, and Ion Stoica. 2025. MoE-Lightning: High-Throughput MoE Inference on Memoryconstrained GPUs. In Proceedings of the 30th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Volume 1 (Rotterdam, Netherlands) (ASPLOS ’25). Association for Computing Machinery, New York, NY, USA, 715–730. doi:10.1145/3669940.3707267 [6] Milind Chabbi and John Mellor-Crummey. 2012. Deadspy: a tool to pinpoint program inefficiencies. In Proceedings of the Tenth International Symposium on Code Generation and Optimization. 124–134. [7] Shuai Che, Michael Boyer, Jiayuan Meng, David Tarjan, Jeremy W Sheaffer, Sang-Ha Lee, and Kevin Skadron. 2009. Rodinia: A benchmark suite for heterogeneous computing. In 2009 IEEE international symposium on workload characterization (IISWC). IEEE, 44–54. [8] Tianqi Chen, Thierry Moreau, Ziheng Jiang, Lianmin Zheng, Eddie Yan, Meghan Cowan, Haichen Shen, Leyuan Wang, Yuwei Hu, Luis Ceze, Carlos Guestrin, and Arvind Krishnamurthy. 2018. TVM: an automated end-to-end optimizing compiler for deep learning. In Proceedings of the 13th USENIX Conference on Operating Systems Design and Implementation (Carlsbad, CA, USA) (OSDI’18). USENIX Association, USA, 579–594. [9] AMD Corporation. [n. d.]. Omniperf documentation. https://rocm. docs.amd.com/projects/omniperf/en/docs-6.2.0/. Accessed: April 2026. [10] ONNX Runtime developers. 2021. ONNX Runtime. https:// onnxruntime.ai/. [11] Chen Ding and Yutao Zhong. 2003. Predicting Whole-Program Locality through Reuse Distance Analysis. In Proceedings of the ACM SIGPLAN 2003 Conference on Programming Language Design and Implementation (PLDI ’03). ACM, New York, NY, USA, 245–257. doi:10.1145/781131.781159 [12] Ariel Eizenberg, Yuanfeng Peng, Toma Pigli, William Mansky, and Joseph Devietti. 2017. BARRACUDA: Binary-level Analysis of Runtime RAces in CUDA programs. In Proceedings of the 38th ACM SIGPLAN Conference on Programming Language Design and Implementation. [13] Georgi Gerganov. 2026. llama.cpp: Port of Facebook’s LLaMA model in C/C++. https://github.com/ggml-org/llama.cpp. [14] Amir Gholami, Zhewei Yao, Sehoon Kim, Coleman Hooper, Michael W. Mahoney, and Kurt Keutzer. 2024. AI and Memory Wall. IEEE Micro 44, 3 (2024), 33–39. doi:10.1109/MM.2024.3373763 [15] Yueming Hao, Nikhil Jain, Rob Van der Wijngaart, Nirmal Saxena, Yuanbo Fan, and Xu Liu. 2023. Drgpu: A top-down profiler for gpu applications. In Proceedings of the 2023 ACM/SPEC International Conference on Performance Engineering. 43–53. [16] Intel.Corp. 2025. Intel VTune Profiler. https://www.intel.com/content/ www/us/en/developer/tools/oneapi/vtune-profiler.html#gs.llu2ig [17] Aditya K Kamath and Arkaprava Basu. 2021. iGUARD: In-GPU Advanced Race Detection. In Proceedings of the ACM SIGOPS 28th Symposium on Operating Systems Principles.

CONCLUSION

We presented Topological Distance and PCDepGraph, a runtime representation of GPU memory reuse structure defined over execution scopes, together with cuVein, which reconstructs it from optimized binaries at profiler-class overhead. PCDepGraph reports how data is actually shared across instructions and execution scopes, the facts that optimization decisions consume. Across HPC and AI workloads on two GPU generations, the representation matched hardware counters, remained stable across platforms, and guided fixes that eliminated up to 90.18% of memory instructions. We see this as a step toward reuse structure as a first-class program profile: today’s manual-fix analyses are already mechanical rules over the graph, pointing toward profile-guided memory optimization.

References [1] Laksono Adhianto, Sinchan Banerjee, Mike Fagan, Mark Krentel, Gabriel Marin, John Mellor-Crummey, and Nathan R Tallent. 2010. HPCToolkit: Tools for performance analysis of optimized parallel programs. Concurrency and Computation: Practice and Experience 22, 6 (2010), 685–701.

12

Reconstructing Fine-grained Memory Reuse Structure in GPU Kernels

[18] Ian Karlin, Jeff Keasler, and Rob Neely. 2013. LULESH 2.0 Updates and Changes. Technical Report LLNL-TR-641973. 1–9 pages. [19] Andrej Karpathy. 2026. llm.c: LLM training in simple, raw C/CUDA. https://github.com/karpathy/llm.c. [20] Mahmoud Khairy, Zhesheng Shen, Tor M Aamodt, and Timothy G Rogers. 2020. Accel-sim: An extensible simulation framework for validated gpu modeling. In 2020 ACM/IEEE 47th Annual International Symposium on Computer Architecture (ISCA). IEEE, 473–486. [21] Jiajia Li, Yuchen Ma, Xiaolong Wu, Ang Li, and Kevin Barker. 2019. PASTA: a parallel sparse tensor algorithm benchmark suite. CCF Transactions on High Performance Computing 1 (2019), 111–130. [22] Mao Lin, Keren Zhou, and Pengfei Su. 2023. Drgpum: Guiding memory optimization for gpu-accelerated applications. In Proceedings of the 28th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Volume 3. 164–178. [23] NVIDIA.Corp. 2025. Nsight Compute. https://docs.nvidia.com/nsightcompute/NsightCompute/index.html [24] NVIDIA.Corp. 2025. Sanitizer API. https://docs.nvidia.com/computesanitizer/SanitizerApiGuide/index.html [25] Adam Paszke, Sam Gross, Francisco Massa, Adam Lerer, James Bradbury, Gregory Chanan, Trevor Killeen, Zeming Lin, Natalia Gimelshein, Luca Antiga, Alban Desmaison, Andreas Köpf, Edward Yang, Zach DeVito, Martin Raison, Alykhan Tejani, Sasank Chilamkurthy, Benoit Steiner, Lu Fang, Junjie Bai, and Soumith Chintala. 2019. PyTorch: an imperative style, high-performance deep learning library. Curran Associates Inc., Red Hook, NY, USA. [26] Jonathan Ragan-Kelley, Connelly Barnes, Andrew Adams, Sylvain Paris, Frédo Durand, and Saman Amarasinghe. 2013. Halide: a language and compiler for optimizing parallelism, locality, and recomputation in image processing pipelines. SIGPLAN Not. 48, 6 (June 2013), 519–530. doi:10.1145/2499370.2462176 [27] Charitha Saumya, Kirshanthan Sundararajah, and Milind Kulkarni. 2022. DARM: Control-Flow Melding for SIMT Thread Divergence Reduction. In 2022 IEEE/ACM International Symposium on Code Generation and Optimization (CGO). 1–13. doi:10.1109/CGO53902.2022. 9741285 [28] Du Shen, Shuaiwen Leon Song, Ang Li, and Xu Liu. 2018. Cudaadvisor: Llvm-based runtime profiling for modern gpus. In Proceedings of the 2018 International Symposium on Code Generation and Optimization. 214–227. [29] Philippe Tillet, H. T. Kung, and David Cox. 2019. Triton: an intermediate language and compiler for tiled neural network computations. In Proceedings of the 3rd ACM SIGPLAN International Workshop on Machine Learning and Programming Languages (Phoenix, AZ, USA) (MAPL 2019). Association for Computing Machinery, New York, NY, USA, 10–19. doi:10.1145/3315508.3329973 [30] John R Tramm, Andrew R Siegel, Tanzima Islam, and Martin Schulz. 2014. XSBench - The Development and Verification of a Performance Abstraction for Monte Carlo Reactor Analysis. In PHYSOR 2014 - The Role of Reactor Physics toward a Sustainable Future. Kyoto. https: //www.mcs.anl.gov/papers/P5064-0114.pdf [31] Oreste Villa, Mark Stephenson, David Nellans, and Stephen W Keckler. 2019. Nvbit: A dynamic binary instrumentation framework for nvidia gpus. In Proceedings of the 52nd Annual IEEE/ACM International Symposium on Microarchitecture. 372–383. [32] Kai Wang and Calvin Lin. 2017. Decoupled affine computation for SIMT GPUs. ACM SIGARCH Computer Architecture News 45, 2 (2017), 295–306. [33] Lu Wang, Magnus Jahre, Almutaz Adileho, and Lieven Eeckhout. 2020. MDM: The GPU memory divergence model. In 2020 53rd Annual IEEE/ACM International Symposium on Microarchitecture (MICRO). IEEE, 1009–1021. [34] Shasha Wen, Milind Chabbi, and Xu Liu. 2017. Redspy: Exploring value locality in software. In Proceedings of the Twenty-Second International

Conference on Architectural Support for Programming Languages and Operating Systems. 47–61. [35] Samuel Williams, Andrew Waterman, and David Patterson. 2009. Roofline: an insightful visual performance model for multicore architectures. Commun. ACM 52, 4 (April 2009), 65–76. doi:10.1145/ 1498765.1498785 [36] Wm A Wulf and Sally A McKee. 1995. Hitting the memory wall: Implications of the obvious. ACM SIGARCH computer architecture news 23, 1 (1995), 20–24. [37] Tsung Tai Yeh, Roland N Green, and Timothy G Rogers. 2020. Dimensionality-aware redundant simt instruction elimination. In Proceedings of the Twenty-Fifth International Conference on Architectural Support for Programming Languages and Operating Systems. 1327– 1340. [38] Yanbo Zhao, Yueming Hao, Zecheng Li, Shuyin Jiao, Xu Liu, and Jiajia Li. 2025. RedSan: A Redundant Memory Instruction Sanitizer for GPU Programs. In Proceedings of the International Conference for High Performance Computing, Networking, Storage and Analysis. 368–382. [39] Keren Zhou, Yueming Hao, John Mellor-Crummey, Xiaozhu Meng, and Xu Liu. 2020. GVProf: A value profiler for GPU-based clusters. In SC20: international conference for high performance computing, networking, storage and analysis. IEEE, 1–16. [40] Keren Zhou, Yueming Hao, John Mellor-Crummey, Xiaozhu Meng, and Xu Liu. 2022. ValueExpert: Exploring value patterns in GPUAccelerated applications. In Proceedings of the 27th ACM International Conference on Architectural Support for Programming Languages and Operating Systems. 171–185. [41] Keren Zhou, Xiaozhu Meng, Ryuichi Sai, and John Mellor-Crummey. 2021. GPA: a GPU performance advisor based on instruction sampling. In Proceedings of the 2021 IEEE/ACM International Symposium on Code Generation and Optimization (Virtual Event, Republic of Korea) (CGO ’21). IEEE Press, 115–125. doi:10.1109/CGO51591.2021.9370339

13

Related documents

• Reconstructing Fine-grained Memory Reuse Structure in GPU Kernels#1648771
• Over the Memory Wall, Into the Instruction Wall: The New Bottleneck in GPU Data Processing#458359
• Adaptive Reuse#473681
• Rotary GPU: Exploring Local Execution Paths for Large Mixture-of-Experts Models Under Limited GPU Memory#238571
• GPU-Accelerated First-Order Method with Randomized Sampling for Binary Integer Programs#1062806