Source-linked AI summary
Towards Standardized Evaluation of GPU Memory Safety with GMSBench
Saurabh Singh, Jaewon Lee, Seonjin Na, Hyesoon Kim
TL;DR
GPU memory-safety evaluation lacks a comprehensive, standardized benchmark, and reliable assessment must distinguish whether a violation manifests from whether a tool detects it. GMSBench provides 149 self-contained tests across error classes and GPU memory spaces, and its evaluation across GPU platforms reveals systematic gaps in Compute Sanitizer coverage.
Problem
Existing evaluations use inconsistent, often technique-specific tests, while GPU violations can manifest without obvious crashes, making coverage comparison and ground-truth assessment difficult.
Source: §I · p. 1, B. Challenges in Evaluating GPU Memory… · p. 2Method
GMSBench provides 149 self-contained, mechanism-agnostic CUDA tests spanning spatial and temporal errors and data races, with baseline execution used to establish observable manifestations where possible.
Source: §III · p. 2, A. Design Principles · p. 2Results
Across three GPU platforms, GMSBench behaved consistently and revealed systematic Compute Sanitizer coverage gaps; Memcheck detected only 18/144 spatial and temporal error tests.
Source: §IV · pp. 4–5, §VI · p. 5Takeaways & Limitations
GMSBench provides a common foundation for reproducible evaluation and comparison of GPU memory-safety mechanisms while exposing gaps in detection coverage.
Source: §I · p. 1, §VI · p. 5
Abstract
from arXiv · showhide
As GPUs become increasingly integral to high-performance computing and machine learning, ensuring memory safety in GPU programs has become crucial for reliable and secure execution. However, evaluating GPU memory safety techniques remains challenging due to the lack of comprehensive and standardized benchmarks. In this paper, we present GMSBench, a GPU memory safety benchmark designed to evaluate a broad range of memory safety violations across different GPU memory spaces and execution scenarios. GMSBench comprises 149 self-contained CUDA tests spanning spatial, temporal, and concurrency errors. The suite provides a standardized foundation for the evaluation and comparative analysis of GPU memory safety mechanisms and helps expose gaps in their detection coverage. We demonstrate the utility of GMSBench by evaluating Compute Sanitizer, a widely used GPU memory error detection tool across multiple GPU architectures.
I. INTRODUCTION
GPU memory errors can cause erroneous results, crashes, and security vulnerabilities, yet inconsistent test collections make detection coverage difficult to compare. The paper introduces GMSBench to standardize evaluation across violation classes and GPU architectures.
I. INTRODUCTION
- Existing test collections vary in error types, memory spaces, and access patterns, making detection coverage difficult to compare and leaving classes of violations untested.GPU memory-safety errors can also be exploited to degrade performance, hijack control flow, or corrupt CPU memory –.Source: §I · p. 1
- Prior benchmarks differ in availability, organization, documentation, and coverage; cuCatch uses a closed-source suite, while CuSafe offers a relatively small suite.These differences make it difficult to determine which vulnerability classes, memory spaces, and access patterns are exercised.Source: §I · p. 1
- GMSBench is an open-source benchmark that systematically covers vulnerability classes, memory spaces, allocation mechanisms, and access behaviors for reproducible evaluation.Its broad coverage is intended to expose weaknesses and blind spots while enabling comparison across mechanisms.Source: §I · p. 1
- The paper classifies spatial, temporal, and concurrency-related errors, then evaluates a GPU memory-error detection tool across architectures to characterize its coverage.The benchmark is open-source and designed for evaluating hardware- and software-based mechanisms.Source: §I · p. 1
A. GPU Memory Safety
GPU memory safety spans distinct memory spaces and execution relationships, so violations can manifest differently from CPU errors. Evaluation must account for that diversity and distinguish actual manifestations from detector reports.
A. GPU Memory Safety
- GPU memory includes host- or device-allocated global memory, per-thread local memory, per-block shared memory, constant memory, and host-pinned memory, with distinct allocation and sharing properties.Violations can cross between valid allocations, corrupt neighboring shared-memory regions, depend on local stack organization, or affect host RAM through pinned memory.Source: A. GPU Memory Safety · pp. 1–2
B. Challenges in Evaluating GPU Memory Safety
- Detection behavior varies with violation type, memory space, allocation mechanism, object layout, and execution scope, so tests should isolate violations and remain mechanism-independent.For example, global overflows can differ by victim adjacency, while race detection may depend on whether threads share a warp, block, or neither.Source: B. Challenges in Evaluating GPU Memory… · p. 2
- A crash or non-zero return code does not establish that the intended violation was detected; evaluation must separately establish whether it manifested and whether a mechanism detected it.Some invalid accesses silently corrupt valid objects, while other semantically present violations show no observable native execution effect.Source: B. Challenges in Evaluating GPU Memory… · p. 2
III. GPU MEMORY SAFETY BENCHMARK
GMSBench comprises 149 self-contained CUDA tests covering spatial and temporal errors and data races across GPU memory spaces and object relationships. Its design aims for meaningful coverage, isolated violations, independent ground truth, and mechanism-agnostic tests.
III. GPU MEMORY SAFETY BENCHMARK
- 149 self-contained CUDA tests cover spatial errors, temporal errors, and data races across GPU memory spaces and object relationships.The benchmark builds on challenges in evaluating GPU memory safety and is summarized by a classification figure.Source: §III · p. 2
A. Design Principles
- GMSBench selects variants for meaningful semantic distinctions or differences in detector behavior, rather than exhaustively enumerating every dimension combination.Each test isolates one violation so results can be attributed to a specific detection capability.Source: A. Design Principles · p. 2
- Baseline GPU execution checks for observable manifestations such as sentinel corruption, planted-value disclosure, or incorrect results; tests without reliable manifestations are marked detection-only.Tests remain mechanism-independent, using shared harness helpers for allocation, layout, and ground truth.Source: A. Design Principles · p. 2
1) Taxonomy:
The taxonomy organizes GMSBench tests across four dimensions: error class, memory window, object context, and specific violation type.
1) Taxonomy:
- Error class distinguishes spatial violations, temporal violations, and data races.These are the three broad classes in the benchmark's classification.Source: 1) Taxonomy · p. 2, Figure 1 · p. 3
- Memory window identifies the memory domains involved, while object context captures allocation mechanisms and relationships among participating objects.Contexts include global allocation, shared-memory placement, local stack-frame placement, and intra-object placement.Source: 1) Taxonomy · p. 2
- Violation type names specific errors such as buffer over-reads or writes, use-after-free, invalid free, double free, uninitialized read, and lost update.Variants such as victim adjacency, access direction, and reuse timing are included when they represent meaningful behavioral or detection differences.Source: 1) Taxonomy · p. 2
2) Spatial Memory Safety:
GMSBench tests spatial violations across GPU memory spaces, varying access types, object relationships, and allocation contexts. Its coverage also includes intra-object and cross-space accesses, with specific scope boundaries.
2) Spatial Memory Safety:
- Spatial tests cover global, local, shared, and constant memory, including cross-space accesses and linear or targeted out-of-bounds reads and writes.They also vary object context where it may affect detector behavior.Source: 2) Spatial Memory Safety · p. 3
- Global-memory tests use cudaMalloc(), in-kernel malloc(), and pinned host-memory allocations accessible to GPU kernels.Source: 2) Spatial Memory Safety · p. 3
- Tests vary source-victim relationships and stack or shared-memory placement; constant-memory cases cover linear and targeted out-of-bounds reads.Adjacent and non-adjacent object relationships are included only for static-to-static cases where a live intervening object can be constructed reliably.Source: 2) Spatial Memory Safety · p. 3
- Intra-object tests cross field boundaries, while cross-space tests probe pointer provenance and bounded local or shared accesses can extend beyond their memory windows.Managed memory contributes no spatial tests because its tested behavior was equivalent to cudaMalloc().Source: 2) Spatial Memory Safety · p. 3
3) Temporal Memory Safety:
GMSBench covers temporal errors across global and local memory, with tests varying allocation mechanism, object lifetime, and whether expired addresses are reused.
3) Temporal Memory Safety:
- Global-memory tests cover use-after-free, invalid free, double free, and uninitialized reads; local-memory tests cover use-after-return and use-after-scope.Shared memory omits use-after-free and free errors because it has no explicit object deallocation during kernel execution; shared-memory uninitialized reads lack reliable ground truth.Source: 3) Temporal Memory Safety · p. 3
- Global-memory temporal tests span cudaMalloc(), in-kernel malloc(), managed-memory, and pinned host-memory allocations.Use-after-free cases distinguish immediate from delayed accesses and include pointer aliases; managed memory has only delayed cases because deallocation unmaps the region.Source: 3) Temporal Memory Safety · p. 3
- Invalid-free tests use interior or mismatched-allocator pointers, while double-free tests deallocate an object twice and uninitialized-read tests read before initialization.Local use-after-scope and use-after-return cases vary whether expired stack objects are accessed before or after subsequent reuse.Source: 3) Temporal Memory Safety · p. 3
4) Data Races:
GMSBench exercises races through unsynchronized non-atomic updates and varies the participating thread relationship and memory scope.
4) Data Races:
- The benchmark excludes linear overflows that reach the source memory-window boundary before reaching another memory window.Source: 4) Data Races · p. 3
- Data-race tests use unsynchronized, non-atomic read-modify-write operations on a common counter, producing deterministic lost updates.Source: 4) Data Races · pp. 3–4
- Shared-memory races cover intra-warp and intra-block relationships, while global-memory races also include cross-block cases.Source: 4) Data Races · p. 4
C. Test Design
Each GMSBench test isolates one violation in a self-contained CUDA file, with a shared harness and observable checks supporting consistent evaluation across targets.
C. Test Design
- Each test defines one gms_test() function, while a common harness supplies the entry point and allocation and layout helpers.A typical test constructs a memory layout, triggers one violation, and exposes an observable effect for baseline checking.Source: C. Test Design · p. 4
- Layout helpers construct controlled placements by measuring addresses, rather than relying on allocator behavior or declaration order.The helpers support adjacent, non-adjacent, pinned-memory, and reclaimed-allocation layouts.Source: C. Test Design · p. 4
- Ground truth checks sentinel corruption, disclosed values, deterministic race errors, runtime rejection of some illegal frees, or reliable stale-value observations.Tests without reliable native manifestations are marked detection-only and evaluated using the mechanism's diagnostic output.Source: C. Test Design · p. 4
- A config-driven runner compiles and executes tests across targets, matching diagnostic patterns while separating detector reports from benchmark ground truth.If a required layout or allocation condition cannot be established, it reports SETUP_FAIL rather than a detector miss.Source: C. Test Design · p. 4
IV. EVALUATION
The evaluation measures violation manifestation under native CUDA execution and detection coverage across three NVIDIA GPU platforms. Results show consistent behavior across platforms, alongside substantial coverage gaps in Compute Sanitizer tools.
Evaluation setup
- The evaluation measures native violation manifestation and detection coverage with memcheck, initcheck, racecheck, and apicheck on three NVIDIA GPU platforms.The systems are an RTX A5000 and A100 (Ampere) and a GH200 Grace Hopper system.Source: §IV · p. 4
Cross-platform invariance
- Native manifestations and tool verdicts were identical across all tested GPU architectures.The results table covers GMSBench results across GPU platforms.Source: §IV · p. 4, Table I · p. 4
Compute Sanitizer Coverage
- 18/144 spatial and temporal error tests were detected by memcheck.Memcheck detects immediate use-after-free accesses but misses delayed and aliased cases; it also misses several live-allocation, targeted out-of-bounds, intra-object, and cross-space violations.Source: §IV · p. 5
- Apicheck detected 10/11 invalid- and double-free tests, initcheck 4/4 uninitialized reads, and racecheck 2/5 race tests.Apicheck missed a device-heap double free performed through in-kernel free(); racecheck detected both shared-memory races but none of the three global-memory races.Source: §IV · p. 5
V. RELATED WORK
Prior GPU memory-safety benchmarks differ in availability, organization, documentation, and coverage, limiting direct comparison across tools. The related work contrasts those suites and surveys detection mechanisms that motivate a common benchmark.
Existing benchmarks
- cuCatch uses 56 tests but is not publicly available, while CuSafe offers 33 open-source tests with more limited size and taxonomy.Neither suite systematically covers several dimensions, including constant and host-pinned memory, cross-window provenance violations, uninitialized reads, aliased use-after-free, and data races under one taxonomy.Source: §V · p. 5
GPU memory-safety tools
- GPU memory-safety tools use redzones or canaries, object-bounds tracking, pointer metadata, compiler instrumentation, runtime checking, and hardware-software co-design,, –.These approaches differ in coverage across memory spaces, allocation mechanisms, and error classes, motivating a common basis for comparison.Source: §V · p. 5
Existing benchmarks
- Table II compares existing GPU memory-safety benchmarks; its legend indicates taxonomy coverage with circle fill and source-code availability with “a”.The discussion notes substantial differences among suites in availability and coverage.Source: §V · p. 5, Table II · p. 5
VI. CONCLUSION
The conclusion presents GMSBench as an open-source benchmark grounded in a systematic GPU memory-safety taxonomy. Evaluation across three GPU platforms shows consistent behavior and exposes systematic gaps in Compute Sanitizer coverage.
Conclusion
- The paper presents a systematic taxonomy and an open-source benchmark for evaluating GPU memory-safety mechanisms.GMSBench covers diverse memory spaces, allocation mechanisms, access patterns, and execution scopes.Source: §VI · p. 5
- Evaluation across three GPU platforms shows consistent benchmark behavior and reveals systematic gaps in Compute Sanitizer detection coverage.Source: §VI · p. 5
- GMSBench provides a common foundation for reproducible evaluation and comparison of GPU memory-safety mechanisms.Source: §VI · p. 5