A Case For Host Code Guided GPU Data Race Detector

1. Problem Statement

Data races in GPU-accelerated programs pose a significant threat to software reliability. Detecting these races is challenging due to the massive parallelism, hierarchical programming models, and unique synchronization mechanisms of GPUs. Existing detection techniques suffer from critical limitations:

• Dynamic Analysis (Runtime): Tools like iGUARD trace memory accesses during execution. While they produce no false positives, they suffer from high false negatives. They can only detect races that manifest during the specific test inputs and execution order used. Since many critical races are rare or input-dependent, dynamic tools often miss them. Furthermore, they incur massive runtime overheads (e.g., 60× slowdown) and memory overheads, making them unsuitable for production.
• Static Analysis (Compile-time): Tools like GPUVerify and FaialAA analyze source code without execution. They can detect races that have not yet manifested but suffer from high false positives. Because they lack runtime information about concrete variable values, they must conservatively assume all possible values for parameters. This leads to reporting races that are theoretically possible but practically impossible given the program's logic. Additionally, many static tools fail to detect intra-warp races or handle acquire-release synchronization (locks) correctly.

2. Core Insight and Methodology

The authors propose a key observation: The host (CPU) code that launches GPU kernels contains crucial semantic information that restricts the values GPU kernel parameters can take. Existing static tools are "myopic," analyzing only the GPU kernel in isolation.

The authors introduce HGRD (Host code Guided GPU Race Detector), a state-of-the-art static analysis tool that performs a holistic analysis of both host (CPU) and kernel (GPU) code.

Key Methodological Steps:

1. Host Code Analysis & Constraint Generation:
   HGRD analyzes the host code to extract five specific classes of semantic information that limit kernel parameter values. These are encoded as constraints for a SAT solver:

   • Asserts: Constraints embedded in assert statements (e.g., ensuring a matrix is square) are extracted to rule out impossible parameter combinations.
   • Thread Grid Dimensions: The host code explicitly sets grid dimensions (e.g., gridDim.x). HGRD uses this to bound thread identifiers (e.g., blockIdx.x), preventing false alarms about races between non-existent threadblocks.
   • Parameter Relations: Host code often derives multiple kernel parameters from the same variable (e.g., image width determines both grid size and a kernel parameter). HGRD tracks these dependencies to ensure parameters remain consistent.
   • Loop Bounds: Parameters passed to kernels are often loop iterators. HGRD extracts loop bounds to constrain parameter ranges (e.g., an offset cannot be negative).
   • Allocation Size Relations: Memory allocation sizes (e.g., cudaMalloc) must be positive. If these sizes are passed as kernel parameters, HGRD infers that the parameters must also be positive.

2. Kernel Code Analysis:

   • Access Pairing: HGRD identifies pairs of memory accesses (loads/stores/atomics) that are not separated by synchronization barriers.
   • Intra-Warp Detection: Unlike prior tools, HGRD explicitly checks for intra-warp races by analyzing syncwarp barriers and thread divergence within a warp.
   • Acquire-Release Synchronization: HGRD handles fine-grain synchronization (locks). It traces atomicCAS (acquire) and atomicExch (release) patterns to determine if conflicting accesses are guarded. It uses the SAT solver to verify if the lock addresses match; if they mismatch, a race is possible.

3. SAT Solver Integration:
   HGRD generates constraints from both the kernel (address aliasing, thread IDs) and the host (parameter bounds, relations). It feeds these into a SAT solver (Google OR-Tools' CP-SAT) to check for satisfiability.

   • If constraints are satisfiable, a race is possible.
   • If constraints are unsatisfiable (due to host code limits), the potential race is discarded as a false positive.

3. Key Contributions

1. Host Code Guidance: First to demonstrate that analyzing host code is essential for accurate static race detection in GPUs, effectively bridging the gap between dynamic precision and static coverage.
2. Broadened Race Detection: Extends static analysis capabilities to detect intra-warp races and races protected by acquire-release synchronization, which previous static tools missed.
3. Zero False Positives/Negatives (in evaluation): HGRD achieves high accuracy by eliminating the conservative assumptions that plague other static tools and the execution dependency of dynamic tools.

4. Evaluation Results

The authors evaluated HGRD on 22 diverse GPU-accelerated programs (from Kaldi, HeCBench, NVIDIA samples, Rodinia, etc.) and compared it against iGUARD (dynamic), GKLEE (symbolic execution), GPUVerify, and FaialAA.

• Accuracy:
  • HGRD: Detected 15 true races with 0 false negatives and 0 false positives.
  • iGUARD: Missed 9 true races (false negatives) because they did not manifest during execution. It reported 0 false positives.
  • GPUVerify: Reported 35 false positives (~79% of its total reports) and missed no races. Most false positives were caused by ignoring host code constraints (e.g., asserting square matrices or specific grid dimensions).
  • FaialAA: Reported 6 false negatives (missing intra-warp races) and 16 false positives.
• Overhead:
  • HGRD: Adds no runtime overhead. Compilation time ranged from milliseconds to 5 minutes (mostly due to the SAT solver), which is a one-time cost during development.
  • iGUARD: Incurred runtime overheads of 1.5× to 60×, making it impractical for production.
  • GKLEE: Failed to terminate on several complex programs within a day.

5. Significance

This paper fundamentally shifts the paradigm of GPU race detection. By proving that host code is not just a launcher but a source of critical semantic constraints, HGRD solves the "false positive" problem inherent in static analysis and the "false negative" problem of dynamic analysis.

• Reliability: It enables developers to catch hard-to-find, infrequent data races before deployment without the performance penalty of runtime instrumentation.
• Scalability: It supports modern GPU features like independent thread scheduling and fine-grain locking, which are increasingly common in graph processing and AI workloads.
• Practicality: Being a compile-time tool with zero runtime overhead, HGRD is ready for integration into CI/CD pipelines and production debugging workflows, offering a robust solution for the reliability of GPU-accelerated software stacks.