recipes/docker-images/glm53-kindling/image/patches

Vendored SM121 kernel patcher

glm53-flash_SM121.py is taken verbatim from MiaAI-Lab/GLM-5.3-Flash-NVFP4-Dual-DGX-Spark, MIT licensed — the licence is kept beside it as LICENSE.MiaAI-Lab.

It is what makes GLM-5.3-Flash run on GB10 at all. The stock glm53-flash-arm64-cu130 image offers only FLASHINFER_MLA_SPARSE_SM120 on compute capability 12, and that backend mandates the packed fp8_ds_mla layout with pe_dim == 64. This checkpoint is NoPE (qk_rope_head_dim: 0), so the kernel refuses and the engine dies in concat_and_cache_mla — after a full ~10 minute weight load.

The seven patches gate paths that already exist onto GB10 rather than writing new kernels. Two are the crux:

• bake_sm90_on_capability_12 lists FLASHINFER_MLA_SPARSE_SM90 for capability 12, ahead of the SM120 entry.
• bake_sm90_wrapper_for_gb10 relaxes that backend's capability.major == 9 to in (9, 12) and swaps FA3 for FA2 off Hopper.

That second half is the part worth internalising: the SM90 sparse-MLA backend is not inherently Hopper-only. Its FA3 kernel is, but FA2 is portable to sm_12x. An earlier reading of this tree concluded the SM90 path was unreachable on GB10 because of wgmma; that was wrong, and it cost a working configuration.

The rest handle FlashInfer's FA2 fp8 gate and CTA tile for GB10's ~101 KB smem, turn PDL off (KDA race on this hardware), and fix indexer/kpool details.

Two of its targets have moved since it was written: the indexer allocation now lives in models/glm5next/nvidia/sparse_indexer.py, and FlashInfer 0.7.0 moved the FA2 fp8 gate into mla/_batch_mla/_backends/_fa_common.py. mia_retarget.py rewrites those paths and anchors in a copy before it runs, so this file stays verbatim.

The script asserts every replacement matches exactly once and aborts otherwise, so a base-image change fails the build rather than silently producing a half-patched tree. Drop it once these land upstream.