Decoding GPU Architecture Constraints: The sm_120 Compatibility Analysis for Custom MLA Verify Kernels

Subject Message (12208): The assistant responds to the user's greenlight to compile for sm_120 by analyzing GPU architecture constraints, shared memory budgets, and the portability of FlashMLA vs. CUTLASS MLA kernels. The message concludes with a source inspection command to determine whether the kernels can be recompiled or must be rewritten.

Introduction

In the high-stakes world of large language model inference optimization, few moments are as decisive as the discovery that the hardware you're targeting simply isn't supported by the software stack you're relying on. This article examines a single message from an opencode coding session—message 12208—where an AI assistant confronts exactly such a revelation and pivots from investigation to action. The message sits at a critical inflection point in a multi-session effort to deploy speculative decoding (DDTree/DFlash) on NVIDIA's RTX PRO 6000 Blackwell consumer GPUs (compute capability sm_120) for the Kimi K2.6 model.

The assistant had been systematically optimizing the DDTree verify attention kernel—the core computational bottleneck in speculative decoding—and had just discovered through binary inspection (cuobjdump) that every optimized MLA (Multi-head Latent Attention) kernel in the SGLang stack was compiled exclusively for architectures other than sm_120. FlashMLA targeted sm_90a (Hopper), sm_100a, and sm_103a (Blackwell datacenter). The larger flash_ops library covered sm_80, sm_86, and sm_90a. Neither included sm_120. The only .so file with sm_120 support was spatial_ops, which handles entirely different operations. This explained why the production service was stuck on the Triton JIT-compiled attention backend—it was literally the only path that worked on this hardware.

The user's response (message 12207) had been direct and decisive: "Definitely compile for sm120 if there is a need. Note SM120 has smaller shared memory vs 100/103." This single sentence greenlit the most consequential engineering decision of the session while planting a critical constraint that would shape every subsequent design choice. Message 12208 is the assistant's reasoning in response to that directive—a dense, multi-layered analysis that weighs architectural compatibility, shared memory budgets, tensor-core ISA differences, and the fundamental question of whether the required kernels can be ported or must be rewritten from scratch.

The Binary Revelation: What cuobjdump Told Us

The assistant opens message 12208 by recapping the cuobjdump results that had been gathered in the preceding messages (12205–12206). These results were the culmination of a multi-message investigation into why the SGLang service was defaulting to the Triton attention backend despite the availability of more optimized alternatives like FlashMLA and CUTLASS MLA.

The flashmla_ops.abi3.so binary contained compiled kernels for sm_90a, sm_100a, and sm_103a—but critically, not sm_120a. The flash_ops.abi3.so (a 1GB monolith bundling most attention kernels) supported sm_80, sm_86, and sm_90a, also missing sm_120. The spatial_ops.abi3.so was the only file with sm_120a support, but it handles spatial operations unrelated to MLA attention.

The "a" suffix on these architecture names is crucial. In NVIDIA's CUDA compilation model, the "a" suffix denotes architecture-specific accelerated code paths that use non-portable features—things like TMA (Tensor Memory Accelerator), wgmma (warp-group matrix multiply-accumulate), and tcgen05 tensor core instructions. These features cannot be forward-JIT-compiled to newer architectures. A kernel compiled with sm_90a features uses Hopper-specific instructions that simply don't exist on sm_120's ISA. By contrast, code compiled for sm_80 (without the "a" suffix) uses portable PTX that the driver can JIT-compile for newer GPUs—which explains why Marlin MoE kernels, compiled for sm_80, work fine on sm_120 even though no sm_120 binary is present.

This distinction is the root cause of the entire compatibility problem. The optimized MLA kernels were built using architecture-specific features that lock them to their target hardware generations. FlashMLA uses Hopper's wgmma and TMA instructions. CUTLASS MLA uses Blackwell datacenter's tcgen05 tensor cores. Neither set of instructions exists on sm_120, which uses an Ada-like ISA with traditional mma (matrix multiply-accumulate) instructions. The binaries physically cannot execute on the RTX PRO 6000, and no amount of driver-level JIT compilation can bridge the gap.

The User's Greenlight: A Constrained Mandate

The user's approval to "compile for sm120" was not a blank check. It came with a critical caveat: "Note SM120 has smaller shared memory vs 100/103." This single observation—that Blackwell consumer GPUs have only 100 KB of shared memory per SM compared to 228 KB on datacenter Blackwell (sm_100/sm_103)—fundamentally constrains the design space.

The assistant's reasoning in message 12208 immediately internalizes this constraint. The shared memory budget determines tile sizes: how many KV elements can be loaded per iteration, how many splits are needed, and ultimately, the arithmetic intensity of the kernel. A tile designed for 228 KB of shared memory on sm_100 would simply not fit in sm_120's 100 KB. The kernel would need resizing—fewer KV elements per iteration, more reduction steps, and some efficiency loss. But as the assistant notes, this is "manageable" compared to the alternative of writing a completely new kernel from scratch.

The user's greenlight also implicitly validates the assistant's investigative approach. Rather than continuing to speculate about whether FlashMLA or CUTLASS MLA might work on sm_120, the assistant had used cuobjdump to definitively answer the question with binary-level evidence. This evidence-driven methodology—replacing assumptions with measurements—is a consistent theme throughout the session and is explicitly rewarded by the user's directive to proceed with compilation.

The Architectural Chasm: tcgen05 vs. mma

One of the most technically nuanced sections of message 12208 is the assistant's analysis of the ISA (Instruction Set Architecture) differences between sm_100/sm_103 and sm_120. This is not merely a matter of recompiling with a different target flag—it's about fundamental architectural divergence.

NVIDIA's Blackwell architecture actually encompasses two distinct GPU families. The datacenter variants (B200, B300) use compute capability sm_100 and sm_103, and they include advanced tensor core features like tcgen05—a fifth-generation tensor core that can directly manipulate shared memory through TMA, enabling extremely efficient matrix multiply operations. The consumer/pro variants (RTX PRO 6000, RTX 5090) use sm_120, which is architecturally closer to Ada Lovelace (sm_90) than to datacenter Blackwell. Sm_120 lacks tcgen05 entirely and relies on the older mma (matrix multiply-accumulate) instruction set.

This has profound implications for kernel portability. CUTLASS MLA, which the assistant had identified as a potential candidate in earlier messages, is built on tcgen05. The entire kernel design—its tiling strategy, its data movement patterns, its synchronization primitives—assumes the presence of tcgen05 tensor cores. On sm_120, these instructions simply don't exist. The kernel cannot be recompiled for sm_120; it would need to be rewritten using mma-based approaches, which would fundamentally change its performance characteristics.

FlashMLA, by contrast, is built on Hopper-era techniques (wgmma and TMA) that are closer to sm_120's capabilities. While sm_120 doesn't have the exact same TMA implementation as Hopper, the wgmma instruction is available and the overall programming model is more portable. This makes FlashMLA the only viable candidate for recompilation—but even then, the shared memory constraint means the tile sizes and iteration strategy would need adjustment.

The assistant's reasoning here demonstrates a sophisticated understanding of GPU architecture. It's not enough to know that a kernel "supports" a given architecture; one must understand which instructions the kernel uses and whether those instructions exist on the target hardware. A kernel that relies on tcgen05 cannot be "ported" by changing a compiler flag—it requires structural changes to the algorithm itself.

The Shared Memory Constraint: 100 KB Changes Everything

The shared memory difference between sm_120 (100 KB) and sm_100/sm_103 (228 KB) is more than a quantitative reduction—it qualitatively changes what kinds of algorithms are feasible. The assistant's reasoning explores this in detail.

In the naive verify attention kernel, the entire score matrix is materialized in shared memory before being reduced. For a 64k-token prefix, this requires storing 64k × 64k partial scores—clearly impossible in 100 KB. The streaming flash attention approach, which processes KV elements in tiles and accumulates results incrementally, is the standard solution. But the tile size directly determines how much shared memory is needed: larger tiles mean fewer iterations but more shared memory per tile. With only 100 KB available, the tiles must be smaller, increasing the number of reduction steps and reducing arithmetic intensity.

The assistant calculates that this is manageable but not free. Fewer KV elements per iteration means more iterations, more global memory reads, and more synchronization overhead. The kernel will still be faster than the Triton baseline (which the assistant had already measured as severely bandwidth-limited at ~14 GB/s effective), but it won't achieve the same peak efficiency as a kernel designed for the larger shared memory pool.

This constraint also interacts with the CUDA graph capture requirement. CUDA graphs require that all kernel launches and memory operations be known at capture time—no dynamic allocations, no host-side synchronization. The assistant had already made the verify kernel capture-safe in the preceding chunk by using fixed buffer allocations and torch-managed workspaces. The shared memory constraint adds another dimension: the kernel's tile size must be fixed at graph capture time, and it must fit within the 100 KB budget while still achieving acceptable performance across all context lengths.

The Source Inspection: Determining Portability

The final section of message 12208 shifts from analysis to action. The assistant recognizes that the binary-level analysis has reached its limit—it can tell that sm_120 support is missing, but it cannot determine why without examining the source code. The critical question is whether FlashMLA's kernels use sm_100-specific instructions that won't port to sm_120, or whether they use portable primitives that just need a new compilation target.

To answer this, the assistant inspects the local sgl-kernel repository structure. The command at the end of the message lists the contents of /home/theuser/sglang-pr/sgl-kernel/csrc/, revealing the source tree: directories for allreduce, attention, cutlass_extensions, elementwise, flashmla_extension.cc, gemm, moe, quantization, spatial, speculative, and others. The presence of flashmla_extension.cc and the attention directory confirms that the FlashMLA source is available locally for inspection.

This inspection is the bridge between diagnosis and treatment. If the FlashMLA source uses only portable CUDA primitives (mma, shared memory, standard tensor cores), then adding sm_120a to the compilation target list is straightforward—a matter of modifying the build configuration and recompiling, with the shared memory caveat handled by adjusting tile sizes. If the source uses tcgen05 or other sm_100-specific features, then the kernel cannot be ported at all, and the assistant must either write a custom sm_120 verify kernel from scratch or find another approach.

The assistant's reasoning correctly identifies that merge_state_v2—the elementwise LSE (log-sum-exp) merging operation—is "trivially portable" since it's just elementwise arithmetic. Even if it's compiled into flash_ops (which lacks sm_120 support), a simple elementwise kernel can be written in minutes. The real question is the FlashMLA kernel itself.

Broader Implications: The MoE Bottleneck

While message 12208 focuses on the verify attention kernel, the assistant's reasoning is informed by a broader understanding of the system's performance bottlenecks. In the preceding chunk, the assistant had already identified that after the verify attention optimization, the bottleneck shifted to MoE (Mixture of Experts) expert imbalance at batch size 1. This is a structural limitation of tensor parallelism (TP) with small batch sizes: with 8 GPUs and only 1 request, the expert computation is spread across GPUs, but the load is inherently unbalanced because different experts handle different tokens.

This context explains why the assistant is willing to invest significant effort in the verify kernel despite the MoE bottleneck: fixing attention is necessary but not sufficient. The attention optimization removes one bottleneck, revealing the next. The MoE imbalance is a separate problem that will need its own solution (batching, expert parallelism, or dynamic scheduling), but it doesn't diminish the value of the attention work. Each layer of optimization reveals the next constraint.

The assistant's disciplined approach—measure, identify bottleneck, fix, measure again—is visible throughout the reasoning. The cuobjdump analysis is a measurement. The user's greenlight is a directive. The source inspection is the next measurement. Each step is grounded in evidence, not speculation.

Conclusion

Message 12208 represents a pivotal moment of synthesis in a complex optimization effort. The assistant takes the raw data from binary analysis (cuobjdump output), combines it with architectural knowledge (tcgen05 vs. mma, shared memory differences), contextualizes it with user input (the greenlight to compile for sm_120 with the shared memory caveat), and produces a clear path forward: inspect the FlashMLA source to determine whether this is a recompile or a rewrite, then proceed accordingly.

The message exemplifies several virtues of effective engineering reasoning: evidence-based diagnosis (using cuobjdump rather than guessing), architectural awareness (understanding ISA differences beyond simple capability numbers), constraint-driven design (internalizing the 100 KB shared memory limit), and systematic problem decomposition (separating the portability question from the shared memory question from the CUDA graph question).

In the broader narrative of the session, this message is the turning point where the team commits to building their own sm_120 kernel rather than waiting for upstream support. The assistant's careful analysis ensures that this commitment is made with full awareness of the constraints and risks, not out of naive optimism. The result—a custom sm_120 verify attention kernel that achieves 3–6× speedup over Triton—validates the reasoning in this message and demonstrates the power of deep architectural understanding in GPU kernel development.