cuThermo: GPU Memory Profiling

• cuThermo is a runtime profiler that detects fine-grained GPU memory inefficiencies by analyzing distinct warp counts at the word-sector level.
• It uses dynamic binary instrumentation on compiled CUDA binaries without requiring hardware, OS, or source code modifications.
• The tool provides granular heat-map visualizations that help diagnose memory issues like false sharing, misalignment, and strided access patterns.

Searching arXiv for the primary paper and closely related profiling/tooling context. cuThermo is a lightweight runtime profiler for GPU memory analysis that identifies fine-grained memory inefficiencies in CUDA applications via heat-map visualization of distinct visited warp counts at the word-sector level. It operates directly on compiled CUDA binaries, requires no changes to hardware, operating system, drivers, or application source code, and is implemented on NVIDIA’s NVBit dynamic binary instrumentation framework (Zhao et al., 24 Jul 2025). Its central objective is to expose in-SM memory behavior that is obscured by aggregate counters, particularly sharing, coalescing, reuse, and contention phenomena at the granularity of words within 32 B sectors.

1. Definition and scope

cuThermo was introduced to address the lack of comprehensive runtime and fine-grained memory profiling support on GPU architectures (Zhao et al., 24 Jul 2025). The tool targets CUDA binaries rather than source programs and performs dynamic binary instrumentation at the SASS level. This design makes it applicable to practitioners who require actionable visibility into GPU memory behavior without recompilation or compiler-assisted transformation.

The distinguishing methodological choice is its profiling granularity. Rather than reporting only per-address totals, object-level usage, or aggregate utilization rates, cuThermo profiles words inside sectors, aligning the analysis with NVIDIA transaction semantics. In the formulation used by the tool, a cache line is 128 B and is subdivided into four 32 B sectors; each sector contains eight 4 B words. This word-sector representation is intended to reveal where warps interact within the memory system and whether those interactions correspond to efficient coalescing, reuse, or pathological sharing patterns (Zhao et al., 24 Jul 2025).

The profiler is explicitly designed for sector-based NVIDIA architectures and was evaluated on Ada Lovelace RTX 4090 and Ampere RTX A4500 under CUDA Toolkit 12.1, NVBit 1.7.3, Linux 5.15, and NVIDIA Driver 535.183.01 (Zhao et al., 24 Jul 2025). A plausible implication is that its abstractions are tightly coupled to current NVIDIA transaction rules, although the paper states that only the Trace Collector and Analyzer mappings would require minor updates if sector size or cache behavior changed.

2. Binary instrumentation and runtime pipeline

cuThermo intercepts CUDA execution through NVBit, extracts cubins, and injects collection code into memory operations at the SASS level. The instrumented operations include load, store, and atomic instructions. For each issued memory instruction, the tool records, per warp, the program counter pc, address[32] for the 32 lanes, size in bytes, active_mask, access_flags indicating type and space, warp_id within the block, and block_id (Zhao et al., 24 Jul 2025).

To control overhead and avoid trace ambiguity, cuThermo uses block-level and kernel-level sampling. By default, only one user-selected block, with block_id = 0 unless overridden, is traced per kernel launch. This avoids cross-block pollution arising from reused warp identifiers and reduces trace volume. A kernel whitelist mode restricts profiling to selected kernels. The traces are buffered in a GPU-side queue; when buffers fill, the Analyzer copies them to the CPU and processes them into word- and sector-level warp-visit histories, eventually writing CSV and configuration files for the GUI (Zhao et al., 24 Jul 2025).

The runtime model therefore combines dynamic instrumentation with selective sampling rather than full execution tracing. The paper argues that this trade-off preserves representative in-SM patterns while keeping profiling practical. This suggests that cuThermo is optimized for diagnosing recurring structural access motifs rather than for exhaustive whole-application memory accounting.

3. Distinct visited warp counts and heat-map construction

The core metric in cuThermo is the distinct visited warp count, defined at both the sector and word levels. For a sector $s$,

$$V(s) = |\{\, w \in W : \exists \text{ thread } t \text{ in warp } w \text{ that accessed any word in } s \,\}|.$$

For a word $(s,i)$ with $i \in \{0,\ldots,7\}$,

$$V(s, i) = |\{\, w \in W : \exists \text{ thread } t \text{ in warp } w \text{ that accessed word } i \text{ of sector } s \,\}|.$$

cuThermo implements these quantities using bitmasks of warp identifiers and counts set bits after kernel completion (Zhao et al., 24 Jul 2025).

For each lane address, the Analyzer computes

$$\text{sector\_tag} = \lfloor \text{address} / 32 \rfloor, \qquad \text{offset} = \text{address} \bmod 32.$$

It maintains sector_history_map[tag], a 9-element array of bitmasks: eight entries for the words in the sector and one aggregate entry for the sector as a whole. Updates use bitwise OR with 1 << warp_id. After the kernel completes, the bit counts of these masks yield the sector-level and word-level temperatures shown in the heat map (Zhao et al., 24 Jul 2025).

The visualization encodes sector tags on the horizontal axis and word IDs 0–7 on the vertical axis, with an additional top row representing the aggregate sector-level count. Consecutive sectors with identical temperatures are compressed and annotated. Color intensity denotes the number of distinct warps touching the word or sector (Zhao et al., 24 Jul 2025).

Interpretation depends on the relation between $V(s)$ and $\max_i V(s,i)$. When sector-level and word-level counts are similarly high, the pattern indicates concentrated sharing or hot spots. When $V(s)$ is much larger than the per-word counts, the same sector is being fragmented across multiple warps, a signature associated with inter-warp contention or false sharing. This sector-versus-word comparison is the key analytical step that aggregate profilers do not provide (Zhao et al., 24 Jul 2025).

4. Memory access patterns identified by cuThermo

Across six applications, cuThermo identified five memory access patterns reported as portable across different GPU architectures (Zhao et al., 24 Jul 2025). These patterns arise from sector-based transaction rules, warp-level coalescing, and the co-location of shared memory and L1 cache.

Pattern	Heat-map signature	Representative examples
Hot spots	Every word in one or more sectors shows high $V(s,i)$, with $$V(s) = |\{\, w \in W : \exists \text{ thread } t \text{ in warp } w \text{ that accessed any word in } s \,\}|.$$0 near per-word maxima	GEMM; vector x in SpMV as “random hot”
Abuse of shared memory	SMEM sectors appear as straight lines; every word temperature is 1, and sector $$V(s) = |\{\, w \in W : \exists \text{ thread } t \text{ in warp } w \text{ that accessed any word in } s \,\}|.$$1 is 1 or small	PASTA spt_TTMRankRBNnzKernelSM; cuSZp compress/decompress
Memory false sharing	Each word has $$V(s) = |\{\, w \in W : \exists \text{ thread } t \text{ in warp } w \text{ that accessed any word in } s \,\}|.$$2 from different warps, while $$V(s) = |\{\, w \in W : \exists \text{ thread } t \text{ in warp } w \text{ that accessed any word in } s \,\}|.$$3	GEMM gemm_v00
Memory misalignment	Boundary sectors of a contiguous region show fewer words accessed and extra sectors loaded	SpMV rowOffsets[r] / rowOffsets[r+1]
Strided memory access	One word per sector is accessed, with high $$V(s) = |\{\, w \in W : \exists \text{ thread } t \text{ in warp } w \text{ that accessed any word in } s \,\}|.$$4 for a single word position across sectors	GRAMSCHM kernel3; GPUMD cell_count variants

Hot spots include a regular form and a “random hot” form. In both, temperatures remain high, but in the random form they vary non-uniformly across words and sectors. The implications reported include cache pressure, potential thrashing, and overuse of L1 or SMEM when reuse is not handled effectively (Zhao et al., 24 Jul 2025).

Abuse of shared memory is characterized by storing thread-private or warp-private temporaries in SMEM, often guarded by __syncthreads(). The paper associates this with reduced effective L1 capacity, unnecessary synchronization, and extra memory operations. False sharing in global memory is diagnosed when multiple warps each touch different words inside the same 32 B sector, forcing multiple sector transactions that would not occur under warp-contiguous accesses. Misalignment is identified through underutilized boundary sectors and extra sector fetches. Strided access is marked by systematic underutilization, with seven out of eight words per sector untouched in the cited example (Zhao et al., 24 Jul 2025).

The portability claim rests on the persistence of sector granularity, warp-level issuance, and SMEM/L1 co-location across NVIDIA generations. The empirical evidence reported is that the same optimizations derived from these signatures produced comparable gains on Ampere and Ada (Zhao et al., 24 Jul 2025).

5. Optimization workflow and case studies

cuThermo is presented as part of an iterative optimization workflow. The sequence given in the paper is: run the CUDA application under cuThermo with default block sampling and optional kernel whitelist; inspect the heat-map GUI for high-temperature sectors, false-sharing signatures, boundary misalignment, and striding; map sectors to data objects using the configuration file; apply targeted optimizations such as data-layout changes, index remapping, alignment, tiling, register or SMEM staging, and warp intrinsics; validate using Nsight Compute metrics; and re-profile to confirm that the heat-map signature has improved (Zhao et al., 24 Jul 2025).

The tool’s guidance is concretized through case studies. In GEMM, false sharing in gemm_v00 was addressed by swapping row and column indices to improve coalescing, producing gemm_v01. The reported result was a 6.83×–7.21× speedup despite a slightly lower L1 hit rate, from 99.22% to 94.93% (Zhao et al., 24 Jul 2025). This case is significant because it shows that higher aggregate cache-hit statistics do not necessarily correspond to better memory behavior when access fragmentation dominates.

In PASTA, a shared-memory thread-local accumulator was replaced with a register-based implementation, removing barriers and producing approximately 1.6× speedup. In cuSZp, warp-local SMEM communication was replaced with __shfl_sync, eliminating shared-memory usage and reducing stall_short_scoreboard. In GRAMSCHM, a strided access pattern in q[i*NJ + k] was corrected by transposing the indexing, yielding 1.20×–1.23× speedup, about 20% fewer instructions, and two fewer registers. In SpMV, misalignment involving rowOffsets was corrected by duplicating indices to enable ldg.s32.v2 vectorized loads, producing roughly 2% speedup (Zhao et al., 24 Jul 2025).

Performance was evaluated in cycles via Nsight Compute on RTX A4500 and RTX 4090. Reported improvements include +721.79% and +682.82% for GEMM gemm_v00, +26.07% and +20.27% for GEMM gemm_v01, +1.85% and +1.97% for SpMV spmv_csr, +163.56% and +159.62% for PASTA spt_TTMRankRBNnzKernelSM, +4.51% and +9.19% for GRAMSCHM kernel2, and +23.18% and +19.81% for GRAMSCHM kernel3 (Zhao et al., 24 Jul 2025). The paper defines

$$V(s) = |\{\, w \in W : \exists \text{ thread } t \text{ in warp } w \text{ that accessed any word in } s \,\}|.$$5

6. Overhead, accuracy, and comparison with prior tools

The overhead study compared original runtime, cuThermo, and Nsight Compute with all memory metrics enabled, averaged over 10 runs with kernel sampling enabled for both profilers. Reported overheads were 9.86× for GEMM, 1.14× for SpMV, 1.07× for PASTA, 57.53× for GRAMSCHM, 1.97× for cuSZp, and 14.49× for GPUMD. The corresponding Nsight Compute overheads were 6.20×, 1.75×, 1.50×, 755.19×, 2.08×, and 57.90×, respectively (Zhao et al., 24 Jul 2025).

These results support two claims made in the paper. First, overhead scales with memory-instruction volume and the number of kernel launches, which is expected for trace-based instrumentation. Second, block and kernel sampling keep cuThermo overhead below Nsight Compute in most cases, although GEMM remains costly because of large trace volume (Zhao et al., 24 Jul 2025).

Accuracy is discussed in terms of pattern fidelity rather than exact event reconstruction. Sampling a single representative block removes cross-block warp_id aliasing and lowers noise, improving the fidelity of in-SM sharing and coalescing analysis (Zhao et al., 24 Jul 2025). This suggests that the tool prioritizes structural interpretability over exhaustive coverage.

The comparison to prior tools situates cuThermo against Nsight Compute and Nsight Systems, CUPTI-based profilers, prior NVBit-based instrumentation, and DrGPUM. Nsight tools provide utilization, coalescing, hit-rate, and stall counters at kernel or device scope but do not reveal which specific words and sectors are shared by which warps. CUPTI-based tools similarly emphasize events and counters without word-sector spatial mapping. Earlier NVBit-based works focus on redundancy or instruction sampling, whereas cuThermo’s novelty is the heat map of warp-level sector and word sharing using distinct warp counts and block sampling. DrGPUM is described as object-centric memory wastage profiling with API correlation but without explicit in-SM cache or shared-memory sharing visualization (Zhao et al., 24 Jul 2025).

7. Limitations, assumptions, and prospective extensions

cuThermo profiles global and shared memory and recognizes unified memory through cudaMallocManaged callbacks, but its scope is explicitly limited to in-SM sharing and coalescing phenomena. Inter-block sharing is not modeled because block placement is non-deterministic. The current design also does not directly analyze tensor core usage, asynchronous cp.async transfers, or shared-memory bank conflicts, although the heat map may hint at some shared-memory misuse (Zhao et al., 24 Jul 2025).

The representativeness of the analysis depends on the sampled block. The paper states that users may need to choose blocks heuristically or use external profilers to identify hot blocks. Pattern recognition is manual: the GUI supports human-in-the-loop diagnosis, but there is no automated classifier for hot spots, false sharing, misalignment, striding, or SMEM abuse (Zhao et al., 24 Jul 2025).

The platform dependency is another constraint. Because the implementation relies on NVBit, support is tied to NVIDIA tooling; portability to other vendors would require re-implementing the Trace Collector on their instrumentation frameworks. The paper nevertheless argues that the modular structure eases portability once such instrumentation exists (Zhao et al., 24 Jul 2025).

Proposed extensions include incorporating SM identifiers and grid-level monitoring to relate patterns to bank locations and cross-SM behavior, exploring improved data placement and hardware designs for special workloads, and releasing both an open-source version and a lightweight variant, cuThermo_light, that would rely only on NVBit to reduce dependencies and simplify maintenance (Zhao et al., 24 Jul 2025).

In this form, cuThermo occupies a specific niche in GPU performance engineering: it is not a replacement for device-wide counter profilers, but a binary-level diagnostic instrument for word-sector spatial structure within a sampled block. Its contribution lies in making memory inefficiencies visually and operationally identifiable at the granularity where GPU transactions are actually formed, enabling targeted transformations whose measured gains range from marginal improvements to multi-fold speedups across NVIDIA architectures (Zhao et al., 24 Jul 2025).