The Pivot Point: Examining Source Code to Determine sm_120 Kernel Portability

In the high-stakes world of custom CUDA kernel development for large language model inference, few moments are as consequential as the one captured in message 12209 of this opencode session. After an extensive multi-message investigation that definitively established that all optimized MLA (Multi-head Latent Attention) kernels in the SGLang ecosystem were compiled exclusively for architectures other than the RTX PRO 6000's sm_120 — and after receiving the user's explicit greenlight to compile for that architecture — the assistant takes the critical step of moving from remote binary inspection to local source code examination. This message represents the pivot point between diagnosis and action, between understanding that a problem exists and understanding how to fix it.

The Investigation That Led Here

To appreciate the significance of message 12209, one must understand the trail of evidence that preceded it. Beginning with message 12202, the assistant had been systematically investigating why the SGLang service on the RTX PRO 6000 Blackwell GPUs was defaulting to the Triton attention backend rather than using any of the optimized MLA kernels. The initial hypothesis was that the DecisionTree's custom masking requirements forced Triton usage, but the investigation quickly revealed a more fundamental hardware compatibility problem.

Through a series of remote SSH commands, the assistant discovered that the flashmla_ops.abi3.so binary contained cubins only for sm_90a, sm_100a, and sm_103a — architectures corresponding to Hopper and Blackwell datacenter GPUs. The flash_ops.abi3.so file, a 1GB behemoth containing most other kernels, supported sm_80, sm_86, and sm_90a but also lacked sm_120. Only the small spatial_ops.abi3.so included sm_120a support. The "a" suffix on these architecture names is crucial: it denotes architecture-specific accelerated features that cannot forward-port to newer architectures through PTX JIT compilation. Code compiled for sm_90a uses Hopper-specific instructions like wgmma and TMA that simply do not exist on sm_120's instruction set architecture (ISA), which is closer to Ada than to Blackwell datacenter.

This explained why the service was stuck on Triton: no packaged optimized MLA kernel could physically execute on the RTX PRO 6000. The user's response in message 12207 — "Definitely compile for sm120 if there is a need. Note SM120 has smaller shared memory vs 100/103" — was both an authorization and a warning. The smaller 100KB shared memory per SM (versus 228KB on datacenter Blackwell) meant that simply recompiling existing kernels might not work; tile sizes designed for the larger pool would need to shrink, potentially requiring architectural changes to the kernels themselves.

Message 12209: The Source Code Examination

The subject message itself is concise but dense with intent. The assistant's reasoning block reveals a clear three-part plan:

I need to dig into the sgl-kernel source to understand how attention mechanisms are implemented, particularly looking at the CUTLASS MLA and merge_state implementations, then trace through the FlashMLA submodule to see how it's integrated, and finally check the build configuration to figure out how to add sm_120a support.

This is the moment where the assistant transitions from being a diagnostician to being an engineer. The question it seeks to answer is binary: can the existing MLA kernels be made to work on sm_120 by simply adding sm_120a to the build architecture list, or do they use sm_100-specific tensor core instructions (like tcgen05) that would require a fundamental rewrite?

The three checks encoded in the bash command are carefully chosen:

Check 1: List the attention CUDA source files. The output reveals five entries: cascade.cu, cutlass_mla_kernel.cu, a directory cutlass_sm100_mla, merge_attn_states.cu, and vertical_slash_index.cu. The presence of cutlass_sm100_mla as a directory (not a file) is immediately telling — it suggests that CUTLASS MLA has a dedicated subdirectory for sm_100-specific code, which likely contains kernels using tcgen05 instructions that are unique to Blackwell datacenter GPUs. This is a strong signal that CUTLASS MLA cannot simply be recompiled for sm_120.

Check 2: Look for FlashMLA as a third-party submodule. The command ls /home/theuser/sglang-pr/sgl-kernel/3rdparty/ | grep -iE 'mla|flash|cutlass' returns nothing. The output section labeled "flashmla 3rdparty" is empty. This is a significant finding: FlashMLA is not a git submodule or extracted third-party library in the local source tree. Combined with the earlier observation that flashmla_ops.abi3.so exists as a compiled binary but no matching source directory appears in the attention csrc listing, this suggests that FlashMLA's source code may be embedded differently — perhaps within the flashmla_extension.cc file seen in the top-level csrc listing, or pulled from a separate repository at build time. The absence of a visible 3rdparty source means the assistant cannot easily inspect FlashMLA's kernel internals without locating the actual source files.

Check 3: Examine the CMakeLists.txt build architecture configuration. The grep output shows only partial results, but the key lines are visible:

Assumptions and Potential Pitfalls

Several assumptions underpin the assistant's approach in this message:

Assumption 1: The local source tree at /home/theuser/sglang-pr/ is a complete, buildable copy of the sgl-kernel repository. The assistant assumes that all source files needed to understand the MLA kernels are present locally. The empty 3rdparty listing for FlashMLA challenges this assumption — if FlashMLA's source is in a submodule that hasn't been fetched, the assistant may need to clone additional repositories to complete the analysis.

Assumption 2: The build system (CMake) is the sole determinant of which architectures are compiled. The assistant assumes that adding sm_120a to the CMake gencode list is sufficient to produce working binaries. However, the kernel source files themselves may contain #ifdef or #if __CUDA_ARCH__ guards that explicitly exclude sm_120, or may use intrinsics that don't exist on that architecture. The CMake configuration is necessary but not sufficient.

Assumption 3: The "a" suffix architecture exclusivity is absolute. The assistant correctly understands that sm_90a code cannot run on sm_120, but there's a subtlety: some CUDA kernels embed both PTX (portable) and SASS (architecture-specific) code paths. The cuobjdump output showed only SASS cubins, but PTX could be embedded separately. If FlashMLA's binary contains PTX for sm_90 that could be JIT-compiled to sm_120, the situation might be more nuanced than the assistant assumes.

Assumption 4: The shared memory constraint is the primary architectural limitation. The user's warning about 100KB shared memory is correct and important, but there may be other sm_120 limitations that affect kernel portability — different warp scheduling behavior, different L1 cache configuration, different tensor core generations, or different register file sizes. The assistant focuses on shared memory because it's the most visible and well-documented difference, but the full picture may be more complex.

Input Knowledge Required

To fully understand this message, a reader needs knowledge spanning several domains:

CUDA architecture naming conventions: Understanding that sm_90 = Hopper (H100), sm_100 = Blackwell datacenter (B200/B300), sm_103 = a Blackwell variant, and sm_120 = Blackwell consumer/pro (RTX PRO 6000). The "a" suffix denotes architecture-specific accelerated features that don't forward-port.

SGLang's attention backend architecture: The service uses a registry pattern where attention backends (FlashMLA, CUTLASS MLA, FlashInfer MLA, Triton) are registered and auto-selected based on GPU capability. The assistant has been tracing through this selection logic for several messages.

MLA (Multi-head Latent Attention): The attention mechanism used by DeepSeek models, which decomposes the attention computation into query, key, and value components with a latent representation. MLA kernels are specialized for this computation pattern and are distinct from standard FlashAttention kernels.

CUDA build tooling: Understanding gencode flags, the difference between PTX and SASS, and how cuobjdump --list-elf reveals compiled architecture targets. The assistant has been using these tools extensively in the preceding messages.

The project context: The assistant is building a custom DDTree (Decision Tree) speculative decoding drafter for the Kimi K2.6 model, running on RTX PRO 6000 Blackwell GPUs with SGLang. The verify attention kernel is the performance bottleneck, and the assistant is exploring whether existing optimized MLA kernels can be used or whether a custom kernel must be written.

Output Knowledge Created

This message creates several forms of knowledge:

1. A confirmed local source tree structure: The assistant now knows that the sgl-kernel source is available at /home/theuser/sglang-pr/sgl-kernel/csrc/attention/ with five entries: cascade.cu, cutlass_mla_kernel.cu, cutlass_sm100_mla/, merge_attn_states.cu, and vertical_slash_index.cu. This directory listing is the foundation for all subsequent source-level analysis.

2. Evidence of architecture-specific code organization: The cutlass_sm100_mla directory name strongly suggests that CUTLASS MLA has sm_100-specific implementations, implying non-portability to sm_120. This is a key data point for the porting decision.

3. A confirmed build system structure: The CMakeLists.txt uses sm_90 as its baseline with an ENABLE_BELOW_SM90 option. The absence of sm_100/sm_120 configuration in the visible output indicates that these architectures may need to be added to the build system, not just enabled.

4. A gap in the FlashMLA source trace: The empty 3rdparty listing means the assistant cannot immediately inspect FlashMLA's source code. This creates a follow-up task: locate the actual FlashMLA source files, either elsewhere in the repository or in a separate repository that needs to be cloned.

5. A clear next step: The assistant now knows it needs to examine the actual kernel source files (particularly cutlass_mla_kernel.cu and the cutlass_sm100_mla directory) to determine whether they use sm_100-specific instructions. The CMakeLists.txt also needs a full reading to understand how to add sm_120a support.

The Thinking Process: A Study in Engineering Discipline

What makes this message remarkable is not just the technical content but the thinking process it reveals. The assistant is operating under significant uncertainty — it has confirmed that the compiled binaries don't support sm_120, but it doesn't yet know why or how hard it would be to fix that. Rather than guessing or jumping to conclusions, it systematically designs three checks that will provide the information needed to make an informed decision.

The reasoning shows a clear hierarchy of concerns:

  1. First, understand the source code structure (what files exist, how they're organized)
  2. Second, trace the dependencies (where FlashMLA's source actually lives)
  3. Third, understand the build system (how to add new architecture targets) This is classic engineering methodology: decompose the problem into independent sub-questions, gather data for each, then synthesize the results into a decision. The assistant is not rushing to implement a solution; it's building the understanding necessary to choose the right solution. The message also demonstrates the assistant's ability to learn from the conversation history. It has absorbed the user's warning about shared memory constraints and is now factoring that into its analysis. It has integrated the findings from the binary inspection (cuobjdump results) and is using them to guide the source code investigation. This cumulative reasoning — where each message builds on the conclusions of previous ones — is the hallmark of effective technical problem-solving.

Conclusion

Message 12209 is a pivot point in a complex engineering effort. It represents the transition from remote diagnosis to local investigation, from understanding what is broken to understanding how to fix it. The assistant's systematic approach — decomposing the portability question into source structure, dependency tracing, and build configuration — provides a model for how to approach similar hardware compatibility problems. The output of this message will directly inform whether the project takes the "cheap recompile" path or the "full kernel rewrite" path, a decision with major implications for the project timeline and technical direction.

In the broader narrative of this opencode session, this message is where the assistant stops being a detective and starts being an engineer. The investigation is complete; the construction is about to begin.