Skip to content

feat(cute-rs): accept sm_103a (B300) as a validated target - #1632

Open
0z5a wants to merge 1 commit into
NVIDIA:mainfrom
0z5a:feat/cute-rs-sm103a
Open

0z5a wants to merge 1 commit into
NVIDIA:mainfrom
0z5a:feat/cute-rs-sm103a

Conversation

@0z5a

@0z5a 0z5a commented Oct 24, 2026

Copy link
Copy Markdown

Summary

Add sm_103a to the cute-rs profile gates and let the sm_100a-class examples run on cc 10.3.

Refs #1631

Changes

  • profile.rs / primitives.rs: accept sm_100a | sm_103a | sm_120a; tests cover acceptance of sm_103a and rejection of sm_103.
  • Examples (fp16_gemm_256x352_cute, fp16_gemm_256x352_ctm, sm100_fmha_prefill_cute, sm100_mla_decode_fp16_cute): Rust and Python capability checks accept (10, 0) and (10, 3). The GEMM reference compiles for sm_103a when an SM103 device is attached; offline compilation keeps sm_100a. The FMHA/MLA reports record the real target instead of a fixed sm_100a. The MLA test stub gains torch.cuda.get_device_capability.
  • Docs: support matrix, preview README and build guide list three targets. tcgen05 operations need sm_100a or sm_103a. The two overview SVGs still show two targets, and their captions say so. I can redraw them if you prefer.
  • ctm-rs tcgen05 / TMEM module comments.

Testing

Environment: B300 (cc 10.3), driver 610.57.04, CUDA 13.3, CUTLASS 4.7.0 compiler (SHA-pinned toolchain), CuTeDSL 4.7.0 + torch 2.14.1 for the comparisons, Linux x86_64, CUDA_OXIDE_DEVICE_BACKEND=cutlass-mlir, --arch sm_103a.

elementwise_cute:        ✓ SUCCESS: f32 and f16 elementwise add are bit-exact!
fp16_gemm_256x352_cute:  verification both-pass
fp16_gemm_256x352_ctm:   CPU verification: PASS (90112 values)
sm100_fmha_prefill_cute --shape 1,257,385,4,2 --causal --check-only:  CuTeDSL: PASS / CuTe-RS: PASS
sm100_fmha_prefill_cute --shape 1,1024,1024,8,8:           CuTeDSL 72.2 us / CuTe-RS 15.7 us
sm100_fmha_prefill_cute --shape 4,4096,4096,8,2 --causal:  CuTeDSL 145.5 us / CuTe-RS 136.6 us
sm100_mla_decode_fp16_cute --check-only:  CuTeDSL: PASS / CuTe-RS: PASS
sm100_mla_decode_fp16_cute (batch 1, kv 1024):  CuTeDSL 28.8 us / CuTe-RS 15.5 us

Timings are "complete invocation" medians as the scripts define them. For the small shapes they include Python host submission on the CuTeDSL side, so they are not a kernel-only comparison.

  • cargo test -p cuda-oxide-mlir-export: the full suite passes with the new tests (lib 86/86 and all integration binaries); after a rustfmt-only change, profile_ lib tests 4/4 and primitives_profile 7/7 again; cargo fmt --all --check clean.
  • Example Python tests (fp16_gemm public_reference.test_* and test_reference, FMHA test_compare, MLA test_compare / test_contract): all OK.
  • just -f cuda-oxide/Justfile check passes (GH200 aarch64, with three separate test-flake fixes applied and RUST_TEST_THREADS=1)
  • cargo oxide run <example> passes on B300

Checklist

  • All commits signed off (git commit -s)
  • SPDX headers on new source files -- no new files

The CUTLASS 4.7.0 compiler emits valid cubins for sm_103a, and every
tcgen05/TMEM form cute-rs uses is listed for sm_103a in the intrinsic
catalog. Add sm_103a to the full-CuTe and primitives profile gates, let
the four sm_100a examples and their CuTeDSL comparison scripts accept
compute capability 10.3, compile the GEMM reference for the attached
device's arch, and record the actual target in the FMHA/MLA reports.

Docs: support matrix, preview README and build guide now list three
targets. The two overview SVGs still show two targets; their captions
say so.

Tested on B300 (cc 10.3, CUDA 13.3) with --arch sm_103a: elementwise
bit-exact; fp16_gemm_256x352_cute both-pass vs the CuTeDSL oracle;
fp16_gemm_256x352_ctm CPU verification pass; FMHA prefill and MLA decode
(opt level 2) check-only and timed runs pass against CuTeDSL.

Cover sm_103a in the full-CuTe and primitives profile tests (and keep
plain sm_103 rejected). The MLA comparison now records the device's
target, so its test stub provides torch.cuda.get_device_capability.

Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>
@0z5a
0z5a requested a review from nihalpasham as a code owner October 24, 2026 18:37
@copy-pr-bot

copy-pr-bot Bot commented Oct 24, 2026

Copy link
Copy Markdown

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant