[kernel] Address review: guard sm_121a on CUDA>=13, downgrade main arch, wheel+README

- Gemini(high): the main extension (fastvideo_kernel_ops) used the raw
  TORCH_CUDA_ARCH_LIST, so a CUDA<13 build with 12.1a in the list would fail to
  compile. Downgrade 121a->120a for the main archs too on CUDA<13 (sm_120a PTX
  forward-JITs to sm_121); mirrors the fp4* guard.
- Copilot: gate sm_121a on CUDA>=13.0 (not 12.9) to match the doc/comment text;
  fix the "[per-arch]" status message to not always claim sm_121a.
- Copilot: make the "sm_120a-only" claim true, not just reworded — update
  fastvideo-kernel/README.md support matrix, and add 12.1a to the aarch64 cu130
  wheel arch list so the published Spark wheel actually includes sm_121a FP4.

Author: Mister-Raggs

.github/workflows/publish-kernel.yml

@@ -176,7 +176,7 @@ jobs:
           # the main extension for the full arch list. CMAKE_BUILD_PARALLEL_LEVEL caps
           # Ninja so heavy CUTLASS/TK template TUs don't OOM the 16 GB runner (exit 143).
           if [ "${{ matrix.platform.arch }}" = "aarch64" ]; then
-            export TORCH_CUDA_ARCH_LIST="10.0a;12.0a"
+            export TORCH_CUDA_ARCH_LIST="10.0a;12.0a;12.1a"
             export CMAKE_ARGS="${CMAKE_ARGS:-} -DFASTVIDEO_KERNEL_BUILD_TK=OFF -DFASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER=ON"
             export CMAKE_BUILD_PARALLEL_LEVEL=1
           elif [ "${{ matrix.torch-cuda.torch-cuda-short }}" = "cu130" ]; then

fastvideo-kernel/CMakeLists.txt

@@ -333,7 +333,15 @@ if(ENABLE_ATTN_QAT_INFER AND TORCH_CUDA_ARCH_LIST)
     set(FASTVIDEO_MAIN_CUDA_ARCHS "${TORCH_CUDA_ARCH_LIST}")
     string(REPLACE "sm_" "" FASTVIDEO_MAIN_CUDA_ARCHS "${FASTVIDEO_MAIN_CUDA_ARCHS}")
     string(REPLACE "." "" FASTVIDEO_MAIN_CUDA_ARCHS "${FASTVIDEO_MAIN_CUDA_ARCHS}")
-    message(STATUS "[per-arch] main extension archs=${FASTVIDEO_MAIN_CUDA_ARCHS}, fp4* archs=sm_120a/sm_121a")
+    # sm_121a needs CUDA 13's ptxas. On older toolchains, downgrade any requested
+    # 121a to 120a for the *main* extension too (not just the fp4* targets below),
+    # so a CUDA-12.x build with 12.1a in the arch list doesn't fail to compile —
+    # sm_120a PTX forward-JITs to sm_121 at runtime.
+    if(CUDAToolkit_VERSION VERSION_LESS 13.0)
+        string(REPLACE "121a" "120a" FASTVIDEO_MAIN_CUDA_ARCHS "${FASTVIDEO_MAIN_CUDA_ARCHS}")
+        list(REMOVE_DUPLICATES FASTVIDEO_MAIN_CUDA_ARCHS)
+    endif()
+    message(STATUS "[per-arch] main extension archs=${FASTVIDEO_MAIN_CUDA_ARCHS}, fp4* archs set per-target below")
 endif()
 
 # Compiler flags
@@ -520,11 +528,11 @@ if(ENABLE_ATTN_QAT_INFER)
     )
     # FP4 kernels target consumer/workstation Blackwell: sm_120a (RTX 5090 /
     # PRO 6000) and sm_121a (DGX Spark GB10). sm_121a needs CUDA 13's ptxas
-    # (12.8 lacks it), so it is added only on CUDA >= 12.9 — pre-13 toolchains
-    # keep the original sm_120a-only build unchanged. sm_121a GB10-verified
-    # (cos ~0.98 vs bf16).
+    # (CUDA 12.x lacks it), so it is added only on CUDA >= 13.0 — older
+    # toolchains keep the original sm_120a-only build unchanged. sm_121a
+    # GB10-verified (cos ~0.98 vs bf16).
     set(FASTVIDEO_FP4_CUDA_ARCHS "120a")
-    if(NOT CUDAToolkit_VERSION VERSION_LESS 12.9)
+    if(NOT CUDAToolkit_VERSION VERSION_LESS 13.0)
         list(APPEND FASTVIDEO_FP4_CUDA_ARCHS "121a")
     endif()
     message(STATUS "fp4* CUDA architectures: ${FASTVIDEO_FP4_CUDA_ARCHS}")

fastvideo-kernel/README.md

@@ -12,7 +12,7 @@ Compiled CUDA extensions (CMake, see the build summary printed at the end of eve
 | same extension, optional part | ThunderKittens sliding-tile attention (`sta_fwd`) and VSA block-sparse (`block_sparse_fwd/bwd`) | `csrc/attention/*_h100.cu` | Hopper `sm_90a` only | `FASTVIDEO_KERNEL_BUILD_TK` (AUTO = ON iff `9.0a` is in the arch list; always OFF on aarch64 hosts — TK headers don't compile there) |
 | same extension, optional part | MiniMax-H3 block-sparse VSA forward (64- and 128-token blocks) | `csrc/attention/block_sparse*_sm100a.cu` | Blackwell `sm_100a` only | ON iff `10.0a` is in `TORCH_CUDA_ARCH_LIST` |
 | same extension, optional part | fused NVLink Ulysses all-to-all | `csrc/comm/ulysses_all_to_all.cu` | CUDA | `FASTVIDEO_KERNEL_BUILD_ULYSSES_A2A` (AUTO = ON with NCCL 2.29+ device headers and library; always OFF on ROCm) |
-| `fp4attn_cuda`, `fp4quant_cuda` | FP4 attention + quantization ("attn_qat_infer", modified SageAttention3) | `attn_qat_infer/` | consumer Blackwell `sm_120a` only, CUDA ≥ 12.8 | `FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER` (AUTO = ON iff `12.0a` is in the arch list) |
+| `fp4attn_cuda`, `fp4quant_cuda` | FP4 attention + quantization ("attn_qat_infer", modified SageAttention3) | `attn_qat_infer/` | consumer/workstation Blackwell `sm_120a` (CUDA ≥ 12.8) + `sm_121a` (DGX Spark GB10; CUDA ≥ 13) | `FASTVIDEO_KERNEL_BUILD_ATTN_QAT_INFER` (AUTO = ON iff `12.0a` or `12.1a` is in the arch list) |
 
 Runtime-JIT kernels (no build step, ship in every wheel/image):
 
@@ -28,7 +28,7 @@ Runtime-JIT kernels (no build step, ship in every wheel/image):
 |---|---|---|---|---|---|---|
 | PyPI wheels (`.github/workflows/publish-kernel.yml`) | version bump in `fastvideo-kernel/pyproject.toml` on main, or manual dispatch | x86_64 cu126 | `9.0a` | ON | AUTO | — (CUDA < 12.8) |
 | | | x86_64 cu130 | `9.0a;10.0a;12.0a` | ON | AUTO | ON |
-| | | aarch64 cu130 | `10.0a;12.0a` | — | AUTO | ON |
+| | | aarch64 cu130 | `10.0a;12.0a;12.1a` | — | AUTO | ON |
 | Docker images `ghcr.io/hao-ai-lab/fastvideo/fastvideo-dev` (`.github/workflows/infra-build-image.yml`) | `docker/Dockerfile` changes on main, or manual dispatch | amd64 cuda12.6.3 + cuda13.0.0 | `9.0a` | ON | AUTO | — |
 | | | arm64 cuda12.6.3 (GH200) | `9.0a` | — (aarch64) | AUTO | — |
 | | | arm64 cuda13.0.0 (GB10 / DGX Spark) | `12.1` | — | AUTO | — |