Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
104 commits
Select commit Hold shift + click to select a range
31b4a9d
feat(cuteir): add shared semantic frontend
nihalpasham Aug 17, 2026
4341206
feat(cuteir): add CUTLASS translation backend
nihalpasham Aug 17, 2026
f33f241
docs(cuteir): record CUTLASS performance comparison
nihalpasham Aug 18, 2026
68f2d1b
test(cute-rs): refresh trybuild snapshot for the pinned nightly's dia…
nihalpasham Sep 2, 2026
60ff1ed
chore(vscode): track the project rust-analyzer settings
nihalpasham Sep 2, 2026
bb18255
feat(cuteir): add SM100 CuTe GEMM and compiler support
nihalpasham Sep 16, 2026
f766ecc
chore(vscode): include SM100 CuTe GEMM in rust-analyzer projects
nihalpasham Sep 16, 2026
e1644ef
fix(cuteir): correct SM100 GEMM license metadata
nihalpasham Sep 16, 2026
4368493
feat(cuteir): pass TMA descriptors as grid constants
nihalpasham Sep 18, 2026
ffffcbe
feat(bench): autotune cuBLASLt baselines and record B200 results
nihalpasham Oct 1, 2026
14e6f4d
merge: migrate CuTe translation to the CUDA Rust layout
nihalpasham Oct 2, 2026
2574f6c
feat: add CUTLASS primitives export profile
nihalpasham Oct 11, 2026
257713c
feat: select primitives profile through CUTLASS codegen
nihalpasham Oct 11, 2026
5e49461
feat: add typed primitive memory and tensor maps
nihalpasham Oct 11, 2026
7453e9a
feat: add explicit primitive resource lifecycles
nihalpasham Oct 11, 2026
d475be9
feat: expose typed TCGen05 primitive instructions
nihalpasham Oct 11, 2026
fd7e801
feat: expose typed TMA and synchronization primitives
nihalpasham Oct 11, 2026
9c684c1
test: verify primitives API through the CUTLASS compiler
nihalpasham Oct 11, 2026
a101bcc
feat(primitives): port persistent GEMMs and add SM120 MXFP4 support
nihalpasham Oct 11, 2026
787b13f
docs(primitives): record RTX 5090 GEMM validation and timing
nihalpasham Oct 11, 2026
f6e5519
perf(primitives): avoid checked shared-region result carriers
nihalpasham Oct 11, 2026
0809425
fix(translation): adapt examples and tests to rebased main
nihalpasham Oct 11, 2026
77fcde6
feat(examples): add C++ blockscale primitives comparison
nihalpasham Oct 11, 2026
c8d3337
refactor(examples): colocate blockscale GEMM references
nihalpasham Oct 11, 2026
96e7aa1
Extract typed CUDA primitives into cuteir/ctm-rs
nihalpasham Oct 11, 2026
96bf070
Rename the blockscale Rust example to blockscale_gemm_ctm
nihalpasham Oct 11, 2026
1755d95
Move the FP16 CTM example into cuteir/examples
nihalpasham Oct 11, 2026
5626ed6
Add a matching CUTLASS CTM blockscale reference
nihalpasham Oct 11, 2026
d108a74
Move the CUTLASS CTM reference beside the Rust CTM example
nihalpasham Oct 11, 2026
3a00962
Improve FP16 CTM codegen and migrate CuTe primitives
nihalpasham Oct 11, 2026
53cc755
Name FP16 benchmark implementations explicitly
nihalpasham Oct 11, 2026
01364c5
Add SM120a attention examples and refine CTM primitives
nihalpasham Oct 12, 2026
fc51afd
Compare FMHA examples against CuTeDSL by default
nihalpasham Oct 12, 2026
238a2ef
Add validated SM100 CuTe-RS attention examples
nihalpasham Oct 12, 2026
d642dcb
Modularize and document the CuTe implementation
nihalpasham Oct 14, 2026
ab462fd
Route SM100 CuTe attention primitives through CTM
nihalpasham Oct 14, 2026
8e8d3f8
Add typed CuTe fusion and BF16 RMSNorm foundations
nihalpasham Oct 15, 2026
bd038b3
Add checked BF16 CuTe feed-forward composition
nihalpasham Oct 15, 2026
c8db207
Route CuTe kernels through CTM and add typed RoPE cache
nihalpasham Oct 15, 2026
c3cbc49
Compose typed BF16 Qwen decoders and grouped-query attention
nihalpasham Oct 15, 2026
2d03cff
Implement typed Qwen speech generation and establish performance base…
nihalpasham Oct 15, 2026
5895bf7
Optimize Qwen GPU execution while preserving exact outputs
nihalpasham Oct 15, 2026
f1aa992
Express Qwen GPU kernels through typed CuTe abstractions
nihalpasham Oct 16, 2026
05a4257
Modularize CuTe IR crates and preserve defining identities
nihalpasham Oct 16, 2026
c128512
Merge upstream SM100 CTM work into CuTe refactor
nihalpasham Oct 16, 2026
1f87f3c
Organize CuTe models and document Qwen workloads
nihalpasham Oct 16, 2026
29466ea
Reduce Qwen projection and Talker execution overhead
nihalpasham Oct 16, 2026
f5f0e95
Clarify CuTe translation docs and clean local artifacts
nihalpasham Oct 16, 2026
6b1d0dc
Match the reference reduction order in Qwen decode projections
nihalpasham Oct 17, 2026
20063af
Document the Qwen workloads and the Vivian greedy investigation
nihalpasham Oct 17, 2026
3bdf00d
Record the post-fix Qwen results in the model docs
nihalpasham Oct 17, 2026
a0dfe62
cuda-oxide: migrate CuTeIR printing to upstream Pliron MLIR interop (…
nihalpasham Oct 18, 2026
904638f
cuda-oxide: add public FP16 GEMM comparison reference
nihalpasham Oct 18, 2026
029fda8
cuda-oxide: prepare FP16 comparison environment automatically
nihalpasham Oct 18, 2026
a0d5776
cuda-oxide: document SM100 attention kernel timings from Nsight
nihalpasham Oct 18, 2026
236cbc4
docs(cuda-oxide): clarify CuTeIR examples and record timing comparisons
nihalpasham Oct 18, 2026
eb0d90d
cuda-oxide: simplify GEMM guides and compare public CTM reference
nihalpasham Oct 18, 2026
a1576bd
Merge latest cuda-oxide CuTeIR translation updates
nihalpasham Oct 18, 2026
746f19b
cuda-oxide: use published Pliron MLIR interop
nihalpasham Oct 18, 2026
38f9940
docs(cuda-oxide): add visual CuTeIR README and CuTe in Rust book section
nihalpasham Oct 19, 2026
916662a
Merge remote-tracking branch 'origin/exp/cuteir-translation' into exp…
nihalpasham Oct 19, 2026
2d62c43
feat(cuteir): add true NVFP4 (E2M1+E4M3) support and qwen3_8b_cute model
nihalpasham Oct 20, 2026
8c729ff
Rename the CuTe stack folder from cuteir to cute-rs-preview
nihalpasham Oct 21, 2026
cba06af
Add the cute-rs preview support matrix to the book
nihalpasham Oct 21, 2026
3d52120
Replace the CuTe stack docs folder with two book pages and fold the c…
nihalpasham Oct 21, 2026
56384d4
Merge origin/main into the cute-rs preview branch
nihalpasham Oct 21, 2026
c6d2bd5
Draw the targets and feature-additions sections of the support matrix…
nihalpasham Oct 21, 2026
cd81346
Make every cute example and model pass the per-example clippy lane
nihalpasham Oct 21, 2026
faa49d6
Remove the recorded measurement evidence from the cute-rs preview
nihalpasham Oct 21, 2026
5f73b0f
Make the cute-rs agent guide state the CuTe-only, CuTeDSL-named, pari…
nihalpasham Oct 22, 2026
9437480
Recognise cute-rs compiler hooks from the importer registry in the co…
nihalpasham Oct 22, 2026
0640e19
Format the collector recognition tests
nihalpasham Oct 22, 2026
44dc384
Rename the cute-rs preview layout vocabulary to the CuTeDSL names
nihalpasham Oct 22, 2026
e2809c2
Rename the cute-rs preview tensor partition and shared-tile reads to …
nihalpasham Oct 22, 2026
7ae3473
Rename the cute-rs preview MMA fragments and atom call to the CuTeDSL…
nihalpasham Oct 22, 2026
143c9c9
Give the cute-rs preview pipelines and arch helpers their CuTeDSL names
nihalpasham Oct 22, 2026
de97b3b
Give the cute-rs preview TMA descriptor and ctm-rs primitives their C…
nihalpasham Oct 22, 2026
664daa7
Give the cute-rs preview pipeline internals one expect verb and one r…
nihalpasham Oct 22, 2026
bac04a9
Move the ctm-rs named-barrier and fence forwarders out of mbarrier
nihalpasham Oct 22, 2026
0782c91
Decode Rust tuple shapes and strides in the cute-rs layout marker dec…
nihalpasham Oct 22, 2026
836b386
Give the cute-rs preview layout markers and coordinate maps their CuT…
nihalpasham Oct 22, 2026
7afaaef
Give the cute-rs preview numeric and block-scaled types their CuTeDSL…
nihalpasham Oct 22, 2026
80650cb
Give the cute-rs preview TMA load pipeline its CuTeDSL name
nihalpasham Oct 22, 2026
99f8da0
Give the cute-rs preview TMA store pipelines their CuTeDSL name
nihalpasham Oct 22, 2026
5d36c96
Give the cute-rs preview pipeline cursor its CuTeDSL index accessor
nihalpasham Oct 22, 2026
a2c93e8
Give the cute-rs preview tile scheduler its CuTeDSL verbs and tile type
nihalpasham Oct 22, 2026
b2b626d
Give the cute-rs preview tile scalar store the store_at verb
nihalpasham Oct 22, 2026
d8a8cef
Give the cute-rs preview divisibility promise its CuTeDSL name
nihalpasham Oct 22, 2026
92e434a
Give the cute-rs preview arch module the CuTeDSL cluster coordinate t…
nihalpasham Oct 22, 2026
0459408
Name the current types in two cute-rs rustdoc examples and one dialec…
nihalpasham Oct 22, 2026
4c591a8
Add 1D local_tile to the cute-rs preview tensor views
nihalpasham Oct 22, 2026
5fe18f4
Add ThrCopy to the cute-rs preview tiled copy
nihalpasham Oct 22, 2026
b3c89b5
Add fence_sync_restrict to the cute-rs preview arch module
nihalpasham Oct 22, 2026
70e449b
Add Sub, Mul, Div and Neg to the cute-rs preview RegisterTile
nihalpasham Oct 22, 2026
55be3ca
Add the swizzle portion accessors to the cute-rs preview ComposedLayout
nihalpasham Oct 22, 2026
f9de624
Add the cluster grid coordinate triples to the cute-rs preview arch m…
nihalpasham Oct 22, 2026
effade4
Name the cute-rs preview block-scale coordinate maps crd2idx
nihalpasham Oct 22, 2026
5eaa6cb
Cover the local_tile route in an example and state three CuTeDSL diff…
nihalpasham Oct 22, 2026
9756630
Bring the cute book figures to the renamed cute-rs names
nihalpasham Oct 22, 2026
3fb9516
Retire the cute-rs root module aliases and move callers to the defini…
nihalpasham Oct 22, 2026
6ccbb52
Record the cute-rs naming deviations on the comparison page and refre…
nihalpasham Oct 22, 2026
7e91ad2
Align the cute docs after the documentation review
nihalpasham Oct 22, 2026
a29b22d
Retire the hidden epilogue compiler-hook re-export
nihalpasham Oct 22, 2026
6f2d0cc
cuda-oxide: move SIMT toolchain pin to stable 1.99.1
nihalpasham Oct 22, 2026
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
The table of contents is too big for display.
Diff view
Diff view
  •  
  •  
  •  
28 changes: 18 additions & 10 deletions .github/workflows/clippy.yml
Original file line number Diff line number Diff line change
Expand Up @@ -37,7 +37,7 @@ jobs:
working-directory: cuda-oxide
run: cargo clippy --workspace --all-targets -- -D warnings

# `rustc-codegen-cuda` and each example under examples/ are their own
# `rustc-codegen-cuda`, the standard and CuTe examples, and CuTe models have
# standalone workspaces (the codegen backend is swapped in via
# RUSTFLAGS, so they can't live in the root [workspace]). That's 60+
# cargo invocations from different workspace roots. By default each
Expand All @@ -60,18 +60,26 @@ jobs:
working-directory: cuda-oxide/crates/rustc-codegen-cuda
run: cargo clippy --all-targets -- -D warnings

- name: Run clippy (rustc-codegen-cuda examples)
- name: Run clippy (standard and CuTe examples and models)
env:
CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }}
CARGO_TARGET_DIR: ${{ runner.temp }}/clippy-shared-target
RUSTC_BOOTSTRAP: "1"
run: |
set -e
for ex in cuda-oxide/crates/rustc-codegen-cuda/examples/*/; do
if [ ! -f "$ex/Cargo.toml" ]; then continue; fi
name=$(basename "$ex")
echo "::group::clippy: $name"
(cd "$ex" && cargo clippy --all-targets -- -D warnings)
echo "::endgroup::"
example_roots=(
cuda-oxide/crates/rustc-codegen-cuda/examples
cuda-oxide/cute-rs-preview/examples
cuda-oxide/cute-rs-preview/models
)
for example_root in "${example_roots[@]}"; do
for ex in "$example_root"/*/; do
if [ ! -f "$ex/Cargo.toml" ]; then continue; fi
name=$(basename "$ex")
echo "::group::clippy: $name"
(cd "$ex" && cargo clippy --all-targets -- -D warnings)
echo "::endgroup::"
done
done

# The loop above is one level deep. A few examples nest a second
Expand All @@ -97,11 +105,11 @@ jobs:
grep -q "\"${sub}\""; then
continue
fi
echo "::group::clippy: ${dir#cuda-oxide/crates/rustc-codegen-cuda/examples/}"
echo "::group::clippy: ${dir#cuda-oxide/}"
(cd "${dir}" && cargo clippy --locked --all-targets -- -D warnings)
echo "::endgroup::"
done < <(
find cuda-oxide/crates/rustc-codegen-cuda/examples -mindepth 3 -name Cargo.toml \
find "${example_roots[@]}" -mindepth 3 -name Cargo.toml \
-not -path '*/target/*' | sort
)

Expand Down
10 changes: 8 additions & 2 deletions .github/workflows/fmt.yml
Original file line number Diff line number Diff line change
Expand Up @@ -43,13 +43,19 @@ jobs:
run: cargo fmt --all -- --check
working-directory: cuda-oxide/crates/cuda-macros/tests/device-only

- name: Check example crate formatting
- name: Check example and model crate formatting
env:
RUSTC_BOOTSTRAP: "1"
run: |
set -euo pipefail
# `cargo fmt --all` also formats local path dependencies. Running it
# per example repeatedly reformats the large shared workspaces.
shopt -s globstar nullglob
manifests=(cuda-oxide/crates/rustc-codegen-cuda/examples/**/Cargo.toml)
manifests=(
cuda-oxide/crates/rustc-codegen-cuda/examples/**/Cargo.toml
cuda-oxide/cute-rs-preview/examples/**/Cargo.toml
cuda-oxide/cute-rs-preview/models/**/Cargo.toml
)
((${#manifests[@]} > 0))
for manifest in "${manifests[@]}"; do
example_dir="$(dirname "$manifest")"
Expand Down
12 changes: 12 additions & 0 deletions .github/workflows/unit-tests.yml
Original file line number Diff line number Diff line change
Expand Up @@ -55,6 +55,18 @@ jobs:
- package: dialect-mir
- package: dialect-nvvm
- package: mir-importer
# CuTe types, semantic IR, and MLIR export need no CUDA toolkit.
# Facade macro tests also compile host metadata via cuda-host.
- package: cute-layout
- package: cute-rs
needs_cuda: true
- package: ctm-rs
needs_cuda: true
- package: ctm-rs
test_args: --all-targets --features host
needs_cuda: true
- package: dialect-cute
- package: cuda-oxide-mlir-export
# Includes a real llc -> ptxas spine-kernel validation.
- package: cuda-oxide-codegen
needs_cuda: true
Expand Down
3 changes: 3 additions & 0 deletions cuda-oxide/.cargo/config.toml
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,9 @@ oxide = "run --package cargo-oxide --"

# Environment variables for all builds
[env]
# Unlock `-Z` flags and `#![feature(rustc_private)]` on the pinned stable
# release (`rust-toolchain.toml`).
RUSTC_BOOTSTRAP = "1"
# libmathdx installation path (for libmathdx-sys crate)
# Required only when using --features libmathdx with mathdx_ffi_test.
# IMPORTANT: Replace the placeholder below with your libmathdx install dir,
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/.devcontainer/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,7 @@ This dev container provides the toolchain expected by cuda-oxide:
- CUDA Toolkit 13.0
- LLVM 21 with NVPTX support
- Clang 21 resource headers for `bindgen`
- Rust `nightly-2026-10-11` with `rust-src`, `rustc-dev`, `rust-analyzer`,
- Rust `1.99.1` with `rust-src`, `rustc-dev`, `rust-analyzer`,
`rustfmt` and `clippy`. `llvm-tools` is not preinstalled in the image;
rustup adds it on your first cargo command under `cuda-oxide/`, because
`cuda-oxide/rust-toolchain.toml` names it. The image ships LLVM 21 and
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/.devcontainer/devcontainer.json
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@
"upgradePackages": false
},
"ghcr.io/devcontainers/features/rust:1": {
"version": "nightly-2026-10-11",
"version": "1.99.1",
"profile": "minimal",
"components": "rust-src,rustc-dev,rust-analyzer,rustfmt,clippy"
}
Expand Down
2 changes: 2 additions & 0 deletions cuda-oxide/.gitattributes
Original file line number Diff line number Diff line change
@@ -1 +1,3 @@
intrinsics/probes/*.ll whitespace=-blank-at-eof
# Retain upstream fixture whitespace and unified patch context verbatim.
third_party/pliron-mlir-interop/** -whitespace
2 changes: 2 additions & 0 deletions cuda-oxide/.gitignore
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,7 @@
.smoketest-logs/
.regress-gemm-sol/
/.layer-target/
/lean/.lake/
*.tar.gz

# GPU compilation artifacts.
Expand All @@ -15,6 +16,7 @@
.ltoir_arch
*.ll
*.bc
*.mlir
*.cubin
*.identity
*.gpu
Expand Down
80 changes: 80 additions & 0 deletions cuda-oxide/Cargo.lock

Some generated files are not rendered by default. Learn more about how customized files appear on GitHub.

16 changes: 15 additions & 1 deletion cuda-oxide/Cargo.toml
Original file line number Diff line number Diff line change
@@ -1,6 +1,11 @@
[workspace]
resolver = "3"
members = [
# CuTe frontend and semantic IR shared by every CuTe backend.
"cute-rs-preview/cute-layout",
"cute-rs-preview/cute-rs",
"cute-rs-preview/ctm-rs",
"cute-rs-preview/dialect-cute",
# Core crates
"crates/cuda-target-spec",
"crates/cuda-intrinsics",
Expand All @@ -13,6 +18,7 @@ members = [
"crates/cuda-host",
"crates/cuda-macros",
"crates/llvm-export",
"crates/cuda-oxide-mlir-export",
"crates/dialect-mir",
"crates/dialect-iket",
"crates/dialect-nvvm",
Expand Down Expand Up @@ -44,7 +50,8 @@ repository = "https://github.com/NVIDIA/cuda-rust"
# Note: crates/rustc-codegen-cuda is NOT a workspace member because it requires
# special rustc nightly features and a different build process.
# Build and run examples with: cargo oxide run <example>
# Examples are in: crates/rustc-codegen-cuda/examples/
# Examples are in: crates/rustc-codegen-cuda/examples/ and cute-rs-preview/examples/
# Complete model workloads are in: cute-rs-preview/models/

[workspace.dependencies]
# Pliron MLIR framework, pinned to the published 0.20.0 release, retaining
Expand All @@ -63,6 +70,12 @@ pliron-llvm = { version = "=0.20.0", default-features = false, features = [
# Internal crates
cuda-target-spec = { path = "crates/cuda-target-spec" }
llvm-export = { path = "crates/llvm-export" }
# Published printer; version-specific compatibility stays in the export adapter.
pliron-mlir-interop = { version = "=0.1.0" }
cuda-oxide-mlir-export = { path = "crates/cuda-oxide-mlir-export" }
cute-layout = { path = "cute-rs-preview/cute-layout" }
ctm-rs = { path = "cute-rs-preview/ctm-rs" }
dialect-cute = { path = "cute-rs-preview/dialect-cute" }
dialect-mir = { path = "crates/dialect-mir" }
dialect-iket = { path = "crates/dialect-iket" }
dialect-ptx = { path = "crates/dialect-ptx" }
Expand Down Expand Up @@ -119,6 +132,7 @@ combine = "4.6.8"
log = "0.4"
half = "2.7"
libloading = "0.8"
object = { version = "0.36", default-features = false, features = ["read", "std", "write"] }
sha2 = "0.10"
serde = { version = "1", features = ["derive"] }
serde_json = "1"
Expand Down
10 changes: 8 additions & 2 deletions cuda-oxide/Justfile
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,10 @@ set windows-shell := ["powershell.exe", "-NoLogo", "-Command"]
# Use Bash on Unix
set shell := ["bash", "-c"]

# Unlock unstable Cargo (`cargo-features = ["profile-rustflags"]`) and rustc
# (`-Z`, `#![feature(rustc_private)]`) gates on the pinned stable toolchain.
export RUSTC_BOOTSTRAP := "1"

# Format all Rust code (SIMT workspace, codegen, examples)
fmt:
cargo oxide fmt
Expand Down Expand Up @@ -64,7 +68,8 @@ test:
-p oxide-tile-export -p oxide-tilebc -p cuda-device -p libnvvm-sys -p nvjitlink-sys \
-p cuda-artifact-finalizer -p cargo-oxide \
-p dialect-iket -p iket-lower -p ptx-parse -p dialect-ptx \
-p ptx-schedule
-p ptx-schedule -p cute-layout -p dialect-cute \
-p cuda-oxide-mlir-export
# `default = []`, but every consumer turns the object features on, and the
# default set alone skips the eight ELF emit/extract tests.
cargo test -p oxide-artifacts --all-targets --features object
Expand All @@ -87,7 +92,8 @@ test:
# Run the CUDA-toolkit-dependent packages (no driver required)
test-cuda:
cargo test --all-targets \
-p cuda-oxide-codegen -p cuda-macros -p cuda-host
-p cuda-oxide-codegen -p cuda-macros -p cuda-host -p cute-rs -p ctm-rs
cargo test -p ctm-rs --all-targets --features host

# Mirror unit-tests.yml's third job, `generated-intrinsics`: the three gates a
# change under crates/cuda-intrinsics-gen has to pass. Nothing else here ran
Expand Down
Loading
Loading