Systematic Bug Hunting on Blackwell: How an AI Assistant Mapped the sm_120 Correctness Frontier

Introduction

When a state-of-the-art language model running on brand-new GPU hardware produces flawless output for a single user but reliably corrupts results at 60 concurrent requests, the debugging challenge is immense. The failure could lurk anywhere across a towering stack of dependencies: CUDA runtime libraries, kernel compilation toolchains, attention implementations, matrix multiplication templates, graph capture mechanisms, and sampling kernels. This article examines a pivotal subagent session within a larger coding conversation—a session that undertook precisely such a diagnosis [1]. Over six rounds of web research, an AI assistant systematically searched the entire LLM inference ecosystem for known correctness bugs on NVIDIA's Blackwell (sm_120) architecture, progressively narrowing from broad exploration to a focused, mechanistic hypothesis with a concrete triage protocol [6].

The context was a high-stakes deployment: a team serving DeepSeek-V4 (via the GLM-5-NVFP4-MTP variant) on eight RTX PRO 6000 Blackwell GPUs using a custom SGLang fork with bespoke sm_120 kernels for flashmla MMA (matrix multiply-accumulate attention), a Triton-based indexer for sparse attention routing, and NVFP4 MoE (Mixture-of-Experts) computation. The parent session had already resolved numerous infrastructure issues—driver installation, CUDA toolkit setup, flash-attn compilation, a PD deadlock, and HiCache corruption mitigations—but the core symptom remained: correct output at concurrency=1, reliably corrupted output at ~60 concurrent requests. The subagent was dispatched with a single mission: find every known correctness bug on Blackwell GPUs across the inference stack, especially anything batch-size or concurrency dependent.

What followed was a masterclass in systematic research methodology. The assistant's six-message arc—from broad parallel searches to targeted verification to a comprehensive ranked report—demonstrates how to approach architecture-dependent correctness failures in modern ML infrastructure. This article traces that arc, examining the key decisions, assumptions, and discoveries at each stage.

Phase 1: Broad Reconnaissance ([msg 1])

The assistant began with four parallel web searches, each targeting a different layer of the inference stack. This four-pronged approach revealed a sophisticated mental model of where correctness bugs could hide:

  1. SGLang issues (the serving framework itself): searching for "sglang github issue sm120 Blackwell RTX PRO 6000 wrong output corrupted results batch concurrency"
  2. sgl-kernel / CUTLASS (the low-level GEMM and MoE kernel implementations): searching for "sgl-kernel Blackwell sm120 NVFP4 MoE correctness garbage output batch size greater than 1"
  3. FlashInfer (attention kernels): searching for "flashinfer Blackwell sm120 attention correctness wrong results NaN fp8 fp4"
  4. Triton (compiler and code generation): searching for "triton Blackwell sm120 miscompilation wrong results incorrect output block size" The results were immediately productive. The assistant discovered four critical leads: - SGLang Issue #18954: NVFP4 models produce NaN outputs on RTX PRO 6000 Blackwell—directly relevant since the user was running NVFP4 on the exact same GPU. - CUTLASS Issue #3096: CUTLASS Grouped GEMM produces garbage output on Blackwell, with a fix involving FlashInfer SM120 patches and compute_120f (CUDA 13.0)—highly relevant to the MoE routing in DeepSeek-V4. - vLLM Issue #41651: FlashInfer attention + FP8 KV cache + CUDA graphs produces random output on Blackwell—relevant to the attention mechanism. - Triton PR #9734: Fix PTX codegen segfaults on consumer Blackwell caused by sm_120a misidentification—relevant to any Triton-compiled kernels in the stack. At this stage, the assistant had a set of promising leads but no clear diagnosis. The search had established that correctness bugs on Blackwell were a known phenomenon across multiple projects, but none of the results explicitly mentioned the concurrency-dependent pattern the user was experiencing.

Phase 2: Synthesis and Deepening ([msg 2])

In the second round, the assistant moved from discovery to analysis. The reasoning section reveals a deliberate cataloging process:

"I'm cataloging the most critical issues I've found, particularly a flashinfer_cutlass bug where FP4 GEMM operations produce NaN values in dense MLP layers when handling heterogeneous batches on SM120 architecture—it crashes consistently within a few rounds of concurrent requests, and the root cause appears to involve a missing component."

The assistant fetched full details from the most promising URLs: flashinfer issue #2708 (the PDL race condition), sglang PR #20047 (a workaround), and flashinfer PR #2716 (a fix). The fetched content revealed the core mechanism:

"The key finding is that missing GDC compile flags trigger a PDL race condition causing output tile corruption under concurrency, which directly matches the ~60 concurrent requests scenario I'm investigating."

This was the diagnostic breakthrough. The assistant had connected three things: (1) a specific compile-time flag (GDC—Grid Dependency Control), (2) a specific race condition mechanism (PDL—Programmatic Dependent Launch), and (3) a specific symptom (output tile corruption under concurrency).

The PDL race condition is a particularly insidious bug. In NVIDIA's CUTLASS library, GEMM operations can use persistent kernel designs where thread blocks persist across multiple tile iterations to amortize launch overhead. The PDL mechanism coordinates data movement between global memory and shared memory. When GDC compile flags are missing, the compiler does not emit the proper cache coherence instructions for the Blackwell memory hierarchy. Under low concurrency, the GPU's memory subsystem may never encounter a race condition because there is no contention for the same cache lines. But under high concurrency (60 requests), multiple thread blocks from different requests may attempt to load overlapping tiles from global memory simultaneously. Without proper GDC-enabled coherence, these concurrent loads can produce stale or partially-updated data in the cache, causing some output tiles to contain corrupted values while others remain correct.

The assistant also launched two additional searches to cover angles not yet explored: the DeepSeek NSA indexer workspace bounds issue (vLLM PR #42856) and CUDA graph batch padding on SM120 (vLLM Issue #35659). These revealed that the indexer and sparse MLA kernels were sized for H200's 140 GiB headroom but failed on consumer Blackwell's 95 GiB limit, and that CUDA graphs could produce cudaErrorIllegalAddress crashes under sustained parallel load.

Phase 3: Targeted Narrowing ([msg 3])

Message 3 represents a conscious shift from "what could be wrong?" to "what is most likely wrong given this specific stack?" The assistant's reasoning explicitly states:

"I'm narrowing in on some specific technical issues that directly match what the user is dealing with—particularly the DeepSeek-V4 indexer problems, the hardware-specific correctness bugs between H100 and other architectures with FP8 attention, and those vLLM issues around attention head scaling that are rooted in FlashInfer and Triton bugs."

The assistant formulated three targeted searches, each designed to probe a specific component of the user's custom kernel stack:

Search 1: The Triton Indexer and DeepSeek-V4 NSA Attention. The query targeted the triton indexer—a custom kernel handling sparse attention routing for DeepSeek-V4's compressed attention mechanism. The result was SGLang Issue #23657: "DeepSeek-V4 compressed attention backend: no SM120 fallback for Lightning Indexer." This revealed that SGLang's compressed attention backend calls get_paged_mqa_logits_metadata directly from DeepGEMM without a fallback path for SM120, meaning the indexer could be silently failing or producing incorrect results on Blackwell hardware.

Search 2: FP8 KV Cache Scaling in Triton Attention. The query probed whether FP8 KV cache scaling issues could contribute to the corruption. The result was vLLM PR #42580, which fixed a bug where k_scale and v_scale were not being applied in the triton attention backend. Even the recommended workaround (--attention-backend TRITON_ATTN) did not restore coherent output because the scaling factors were simply not being applied.

Search 3: The H100-to-SM120 Correctness Gap. The query targeted the most frustrating class of bugs: kernels that work perfectly on H100 (sm_90) or B200 (sm_100) but produce garbage on consumer Blackwell (sm_120). The result was CUTLASS Issue #3096, which documented that "all native FP4 MoE backends produce garbage output or crash on SM120 (compute_120) due to broken CUTLASS grouped GEMM templates." The fix required systematic patching of FlashInfer's SM120 capability checks and CuTe DSL architecture restrictions.

Phase 4: Final Verification ([msg 4])

Before compiling the final report, the assistant performed two final targeted searches to nail down remaining uncertainties. The reasoning shows a conscious weighing of evidence:

"Looking at the symptoms more carefully: the user has DeepSeek-V4 running on RTX PRO 6000 Blackwell with a custom SGLang fork that includes custom SM120 kernels for flashmla MMA, triton indexer, and NVFP4 MoE. The corruption appears at around 60 concurrent requests but works correctly at concurrency=1. The top match is the PDL race condition from flashinfer #2708 and the related PRs—the CUTLASS FP4 GEMM used in NVFP4 MoE produces NaN or garbage tiles under high concurrency (64+) but works fine at concurrency=1, and it's fixed by GDC compile flags or switching to flashinfer_cudnn. This matches the symptom exactly."

The two final searches confirmed critical details:

Phase 5: The Comprehensive Report ([msg 5])

The final message is a remarkable artifact: a ranked research report cataloging 17 distinct bug classes across seven tiers of relevance, spanning FlashInfer CUTLASS GEMM races, CUDA graph hazards, Triton miscompilation, DeepSeek-V4-specific indexer failures, NVFP4 numerical correctness issues, and sampling kernel bugs.

The seven tiers are organized by relevance to the specific symptom:

Tier 1 (direct match) covers the PDL race condition in both dense GEMM (flashinfer #2708) and fused MoE (flashinfer #2913) paths. The trigger conditions are documented with surgical precision: only under high concurrency (64+ simultaneous requests), never on standalone mm_fp4() or batch=1, batch-composition dependent, NaN/garbage appears in contiguous 128-aligned blocks matching CTA tile boundaries, crashes consistently within 1–4 rounds of 64 concurrent requests, and CUDA_LAUNCH_BLOCKING=1 makes it disappear.

Tier 2 covers CUDA-graph-related corruption, including cudaErrorIllegalAddress under sustained parallel load (vLLM #35659) and two independent attention-correctness bugs that are "right on sm_90/sm_100, wrong on sm_120" (vLLM #41651/#42650/#42580). The assistant provides a crucial architectural insight about why these are sm_120-specific: the sm_100 TRTLLM-Gen path reads heads at runtime (avoiding the plan-time mismatch), while sm_120 falls back to native FlashInfer prefill which has the plan/runtime split.

Tier 3 covers structural NVFP4/FP4 GEMM issues, including SMEM overflow in grouped GEMM (CUTLASS #3096), NVFP4 mm_fp4 returning all zeros (flashinfer #2577), and the critical finding that consumer Blackwell lacks the hardware E2M1 conversion PTX instruction that the FP4 path assumes (TensorRT-LLM #14154). The assistant provides a valuable triage heuristic: "deterministic garbage → scale/convert math bug; noisy/varying garbage → race or uninitialized scratch."

Tier 4 covers DeepSeek-V4-specific paths, including the lightning indexer producing NaN/all-zero on SM120 (ktransformers #2001), missing SM120 kernels for paged_mqa_logits (sglang #23657 + DeepGEMM #317), indexer top-k fallback correctness concerns (vLLM #41063), and a W4A16 + Marlin MoE token-stream corruption under concurrent decode (jasl/vllm #12).

Tier 5 covers Triton miscompilation, including the alarming finding that tl.dot() scratch-space allocation differs between CUDA graph capture and replay in Triton versions before 3.6, causing silent data corruption. Also covered are the sm_120 vs sm_120a target confusion causing invalid codegen and segfaults.

Tier 6 covers sampling correctness, including a native Triton top-k/top-p kernel that assumes contiguous logits and produces NaN with sliced/padded vocabularies, and top_k near vocab_size hanging on SM120.

Tier 7 catalogs additional CUDA-graph batch-padding hazards across sglang and vLLM.

The Bottom Line: A Diagnostic Framework

The report concludes with a concrete triage protocol that transforms the sprawling catalog into actionable steps:

  1. Run with CUDA_LAUNCH_BLOCKING=1. If corruption disappears, it's a PDL/launch-overlap race. Confirm CUTLASS_GDC_ENABLED expands to real griddepcontrol.wait/launch PTX, not no-ops, and that CUTLASS ≥ 4.4 is in use.
  2. A/B the FP4 GEMM backend. If routing dense/MoE FP4 through cudnn (or a non-PDL path) stops corruption, that pins it to the CUTLASS PDL path.
  3. Check determinism of the garbage. Byte-identical across runs → scale-inversion or E2M1-convert math bug; varying/noisy → race or uninitialized scratch.
  4. Verify Triton ≥ 3.6 for the custom Triton indexer if it uses tl.dot inside CUDA graphs, and numerically gate the indexer top-k against an fp32 reference at batch>1.
  5. Inspect plan-time vs runtime shape agreement (heads/topk/seqlen) in custom attention/indexer—the sm_120-specific "stops writing output → zeros → garbage" failure is concurrency-triggered and bs=1-clean.
  6. If a quantized MoE decode kernel is involved, beware a second concurrent-decode race independent of attention.

Methodological Lessons

This subagent session demonstrates several principles of effective bug research on new hardware architectures:

Parallel exploration, sequential narrowing. The assistant began with broad parallel searches across all relevant components, then progressively narrowed based on the strongest signal. This avoids premature commitment to a single hypothesis while still converging efficiently.

Symptom-to-mechanism mapping. The assistant didn't just collect bug reports; it connected symptoms (concurrency-dependent corruption) to mechanisms (missing GDC flags → PDL race → stale cache reads) to root causes (CUTLASS JIT omitting compile flags). This chain of reasoning is what transforms a catalog of issues into a diagnostic framework.

Cross-repository synthesis. The same fundamental bug (missing GDC compile flags) manifested across dense GEMM (flashinfer #2708), fused MoE (flashinfer #2913), and CUDA-graph replay (vLLM #35659). The assistant recognized this pattern and elevated it, showing that a single root cause can produce different symptoms in different code paths.

Architectural reasoning about specificity. The assistant explained why certain bugs are sm_120-specific: the sm_100 TRTLLM-Gen path reads heads at runtime (avoiding the plan-time mismatch), while sm_120 falls back to FlashInfer prefill which has the plan/runtime split. This kind of cross-architecture reasoning is essential for debugging on new hardware.

Practical triage design. The six-step triage protocol is arguably the most valuable output of the entire session. It transforms the research findings from a passive catalog into an active diagnostic tool, giving the team a clear path forward regardless of which specific bug is causing their corruption.

Conclusion

This subagent session is a testament to the power of systematic research in debugging. Over six rounds of web searching, fetching, and synthesizing, the assistant mapped the entire landscape of Blackwell correctness bugs—17 bug classes across 7 repositories—and organized them into a ranked, actionable framework. The primary suspect, the PDL race condition caused by missing GDC compile flags, provides a mechanistic explanation for the user's concurrency-dependent corruption that matches the symptom with remarkable precision.

For engineers deploying on Blackwell hardware, this session serves as both a diagnostic guide and a cautionary tale. The Blackwell software ecosystem is still maturing, and correctness bugs lurk at every layer of the stack—from CUTLASS GEMM templates to Triton compiler codegen to CUDA graph replay mechanics. The triage protocol developed here—start with CUDA_LAUNCH_BLOCKING=1, A/B the GEMM backend, check determinism, verify Triton version, inspect shape agreement—is a practical methodology that can be applied to any concurrency-dependent corruption on sm_120.

But the session also has limitations. It synthesizes second-hand information from GitHub issues, which may contain incorrect diagnoses or incomplete information. It assumes the user's custom kernels are similar enough to the open-source implementations for the analogies to hold. And it cannot account for bugs that have not been reported or that are specific to the user's unique software configuration. The true test will be whether the triage protocol, when applied to the actual deployment, isolates and resolves the corruption. That test lies beyond the scope of this research session—but the research has provided the map, the compass, and the trailhead.## References

[1] "The Blackwell Bug Hunt: A Deep Dive into a Web Research Task for sm_120 Correctness" — Analyzes the opening user message ([msg 0]) that launched the research mission, examining the strategic context, search architecture, and embedded assumptions.

[2] "The PDL Race Condition: How Missing GDC Compile Flags Explain Concurrency-Dependent Corruption on Blackwell GPUs" — Examines [msg 2] where the assistant identifies the PDL race condition as the primary suspect, connecting missing GDC flags to output tile corruption under concurrency.

[3] "Narrowing the Search: How an AI Assistant Refines Its Investigation into Blackwell GPU Correctness Bugs" — Examines [msg 3] where the assistant transitions from broad exploration to targeted verification of the DeepSeek-V4 indexer, FP8 KV cache scaling, and the H100-to-SM120 correctness gap.

[4] "The PDL Race Condition: How One Message Diagnosed a Concurrency-Dependent GPU Bug" — Examines [msg 4] where the assistant performs final verification searches and prepares to compile the ranked report.

[5] "Mapping the Blackwell Correctness Frontier: A Systematic Search for sm_120 Inference Bugs" — Examines [msg 1] where the assistant conducts initial broad parallel searches across SGLang, CUTLASS, FlashInfer, and Triton.

[6] "The PDL Race That Ate Your Inference: Dissecting a Blackwell Correctness Bug Hunt" — Examines [msg 5], the comprehensive ranked research report cataloging 17 bug classes across seven tiers of relevance.