Feature request: support NVIDIA GB10 (sm121) with FFA_FA4 backend #321

Summary

Please add NVIDIA GB10 / compute capability 12.1 (sm121) support to the
FFA_FA4 backend.

MagiAttention builds and imports successfully on GB10 with CUDA 13.0, but
neither native FFA nor the current FFA_FA4 backend can execute. The current
SM100 path uses tcgen05 instructions that are not available on GB10.

I opened an inference-first draft implementation in the current FA4 submodule:

• Draft: add SM120/SM121 CpAsync forward path demonatic/flash-attention#14

The draft adds an explicit SM120/SM121 CpAsync + warp-MMA forward path and
preserves arbitrary-mask semantics using the existing SM80-style kernel. It is
intentionally limited and still requires real GB10 compile/correctness
validation before it can be consumed here.

NVIDIA CUTLASS PR #3030 provides the relevant SM120/SM121 implementation
reference:

• [CuTeDSL] Flash Attention v2 for SM120 (Blackwell GeForce) NVIDIA/cutlass#3030

Environment

• GPU: NVIDIA GB10
• Compute capability: (12, 1) / sm121
• Architecture: Linux aarch64
• Driver: 580.95.05
• CUDA toolkit: 13.0 (nvcc 13.0.88)
• PyTorch: 2.11.0+cu130
• MagiAttention: 2c64135
• Flash-Attention submodule: ee1d151
• flash_attn_cute: 0.0.2
• Downstream workload: NVIDIA LocateAnything-3B MTP inference

Current failure

(12, 1)
AssertionError: Unsupported compute capability. Supported: 9.x, 10.x

Forcing the SM100 dispatch confirms that an allow-list-only change is not
sufficient:

cutlass.cute.nvgpu.common.OpError:
expects arch to be one of [Arch.sm_100a, Arch.sm_100f, Arch.sm_101a,
Arch.sm_101f, Arch.sm_103a, Arch.sm_103f, Arch.sm_110a, Arch.sm_110f],
but got Arch.sm_121a

Proposed initial scope

• inference-only BF16/FP16 forward
• head dimensions 64 and 128
• MagiAttention arbitrary-mask/HSTU correctness
• explicit GB10 capability detection and dispatch
• end-to-end LocateAnything-3B fast and hybrid validation

Backward, varlen, paged KV, SplitKV, and performance tuning can remain out of
scope for the first integration.