The Architecture Decision That Unlocked Blackwell Consumer GPUs

In the high-stakes world of deploying large language models on novel hardware, a single sentence from a domain expert can redirect days of engineering effort. At message [msg 12207] in this opencode session, the user delivered exactly such a message — brief, technically dense, and carrying the weight of an architectural decision. The message reads:

Definitely compile for sm120 if there is a need. Note SM120 has smaller shared memory vs 100/103

To an outsider, this looks like a casual nod. To the assistant that received it, it was the green light to abandon a dead-end investigation and commit to a fundamentally different engineering strategy — one that would ultimately deliver a 3–6× decode speedup over the existing Triton-based implementation.

The Context: A Dead End at the Binary Level

The message arrived at a critical juncture. The assistant had been deep in a forensic investigation of why the DDTree speculative decoding verify attention was running at a mere 14 GB/s effective bandwidth — 130× below the 1.8 TB/s peak of the RTX PRO 6000 Blackwell GPUs. The root cause had been traced to SGLang's DDTree verify path being locked to the Triton MLA attention kernel, which was the only backend that accepted the custom tree visibility mask required for speculative decoding.

The natural question was: could the team simply use one of the existing optimized MLA kernels — FlashMLA, cutlass-MLA, or flashinfer-MLA — which were known to deliver dramatically better performance on other architectures? The assistant had spent several messages ([msg 12201] through [msg 12206]) methodically investigating this possibility, running cuobjdump on the compiled binaries, inspecting architecture guards in the Python source, and tracing the SGLang attention backend registry.

The finding was definitive and discouraging. The cuobjdump --list-elf output on flashmla_ops.abi3.so showed SASS cubins only for sm_90a, sm_100a, and sm_103a — architectures corresponding to Hopper (H100) and Blackwell datacenter (B200/B300) GPUs. There was no sm_120a target anywhere. The flash_ops.abi3.so (a 1 GB binary containing most other kernels) similarly lacked sm_120 support. Critically, the assistant also confirmed that these .so files contained no embedded PTX — meaning there was no fallback path for the CUDA driver to JIT-compile portable code to the new architecture. The kernels used non-portable instructions like wgmma (Hopper) and tcgen05 (Blackwell datacenter) that simply do not exist on sm_120's Ada-like ISA, which relies on mma.sync and cp.async.

The assistant's reasoning in [msg 12206] laid out the architectural reality: "The 'a' suffix means architecture-specific accelerated features that can't port forward, so sm_90a/100a/103a code simply won't run on sm_120." The assistant correctly identified that sm_120 (used by RTX PRO 6000 and consumer Blackwell GPUs like the 5070 Ti) is architecturally closer to Ada than to Hopper or Blackwell datacenter — it lacks the tensor core instructions that the optimized MLA kernels depend on.

The User's Decision: Compile for sm_120

The user's response at [msg 12207] cuts through the uncertainty with two clear directives. First, "Definitely compile for sm120 if there is a need" — this is authorization to abandon the search for off-the-shelf optimized kernels and instead build custom sm_120 code. The assistant had been weighing two options: a time-boxed "R1" spike to try recompiling FlashMLA for sm_120 (likely to fail due to ISA incompatibility), versus a full "R2" effort to write a custom sm_120 flash-MLA verify kernel from scratch. The user's response implicitly endorses the R2 path, since the assistant had already established that a simple recompile wouldn't work.

Second, and perhaps more importantly, the user provides a critical architectural constraint: "Note SM120 has smaller shared memory vs 100/103." This is domain knowledge that directly shapes kernel design. Blackwell datacenter GPUs (sm_100/sm_103) offer approximately 228 KB of shared memory per SM, while consumer Blackwell (sm_120) is limited to roughly 100 KB. This difference is not a minor detail — it determines tile sizes, register pressure, and whether certain algorithmic approaches are even feasible. A flash attention kernel designed for 228 KB of shared memory would simply not fit in 100 KB without retuning, and the user's warning ensures the assistant accounts for this from the start.

Assumptions and Knowledge

The message assumes several things that are worth examining. First, it assumes the assistant has already done the investigative work to determine that sm_120 compilation is necessary — which is correct, as the preceding messages show exhaustive binary analysis. Second, it assumes the assistant understands the implications of the shared memory constraint for kernel design — a reasonable assumption given the assistant's demonstrated expertise with CUDA and flash attention algorithms. Third, it assumes that building for sm_120 is technically feasible, which is true: while the optimized MLA kernels use non-portable instructions, a custom kernel using mma.sync and cp.async can target sm_120 directly.

There is one potential incorrect assumption embedded in the user's phrasing. The phrase "compile for sm120" could be read as suggesting that existing kernels merely need a recompile with sm_120 added to the architecture list. The assistant's earlier investigation had already disproven this: the issue is not just missing binaries but ISA incompatibility. The assistant correctly interprets the user's intent as authorization to build sm_120-native code, not as a naive instruction to flip a compiler flag.

The Impact: A Pivot That Unlocked 3–6× Speedup

This message was the turning point. Before it, the assistant was still investigating whether existing optimized kernels could be adapted. After it, the direction was clear: build a custom sm_120 flash-MLA verify kernel, implement it as a custom SGLang attention backend, and optimize for the 100 KB shared memory constraint.

The assistant immediately acted on this directive. In the following messages ([msg 12208] onward), it began inspecting the local sgl-kernel source tree, examining the CMake build configuration, and designing the custom kernel architecture. The resulting plan — documented in plans/0002-sm120-verify-kernel-defrag.md — specified a streaming flash-decode MLA verify kernel with online softmax, cp.async double-buffering, and tile sizes tuned for ≤100 KB shared memory. The kernel would consume SGLang's native static buffers for CUDA graph capture safety, use a fixed NSPLIT parameter, and validate token-exact correctness against a naive oracle.

The results, as documented in the segment summary for chunk 1, were dramatic: the custom sm_120 verify kernel delivered a 3–6× end-to-end decode speedup over Triton across all context lengths (4k–65k tokens), with CUDA graph capture succeeding in ~1.5 seconds and replay producing generations matching the Triton baseline. The shared memory constraint that the user flagged became a central design parameter — the kernel's tile sizing, reduction strategy, and occupancy tuning all revolved around fitting within the 100 KB limit.

The Thinking Process Revealed

The user's message reveals a thinking process that is both decisive and informed. The user had been following the assistant's investigation — the cuobjdump results, the architecture analysis, the ISA incompatibility findings — and had reached the same conclusion: off-the-shelf kernels won't work on sm_120. Rather than asking for more analysis or hedging with conditional approval, the user gives a definitive "yes" and immediately adds the most critical design constraint.

This is characteristic of a domain expert who understands both the hardware and the software stack. The user knows that sm_120 is not just "another GPU architecture" but one with specific resource limitations that will dominate kernel design decisions. By flagging the shared memory difference upfront, the user prevents the assistant from spending time on a kernel design that assumes datacenter-class resources.

The message also reveals trust in the assistant's technical judgment. The user doesn't specify how to compile for sm_120, what kernel design to use, or how to handle the shared memory constraint — they trust the assistant to translate these high-level directives into concrete implementation. This trust was well-placed: the assistant's subsequent work produced a kernel that not only worked but delivered the promised speedup.

Conclusion

Message [msg 12207] is a masterclass in efficient technical communication. In 18 words, the user accomplishes three things: authorizes a new technical direction, provides a critical architectural constraint, and implicitly confirms the assistant's analysis. The message's brevity belies its importance — it was the pivot point that transformed a stalled investigation into a successful engineering effort, ultimately delivering a 3–6× decode speedup on hardware that the existing software stack could not effectively utilize.