Parent: #204 — Blackwell Phase 2 kernel port
Spec: blackwell-kernel-port-v1.md § Sub-task 2
Branch: tairan/blackwell-02-02-recompile-generic
Depends on: sub-task 1
What
Populate _sm100_extensions with 9 kernels that use only generic CUDA or SM80 mma.sync (safe to recompile for -arch=sm_100a without source changes):
marlin_grouped_gemm, fp8_blockwise_ops, marlin_transform, routing generics (4×, excluding fused_gate.cu which uses WGMMA), dispatch_scatter_3d, fused_kv_norm_rope, fused_q_absorb, fused_q_split.
Acceptance
- All 9 compile with ptxas reporting
sm_100a code
- Import smoke tests pass
Parent: #204 — Blackwell Phase 2 kernel port
Spec: blackwell-kernel-port-v1.md § Sub-task 2
Branch:
tairan/blackwell-02-02-recompile-genericDepends on: sub-task 1
What
Populate
_sm100_extensionswith 9 kernels that use only generic CUDA or SM80 mma.sync (safe to recompile for-arch=sm_100awithout source changes):marlin_grouped_gemm,fp8_blockwise_ops,marlin_transform, routing generics (4×, excluding fused_gate.cu which uses WGMMA),dispatch_scatter_3d,fused_kv_norm_rope,fused_q_absorb,fused_q_split.Acceptance
sm_100acode