Repository navigation
Conversation
DeviceBuffer::cast_elem, DeviceBuffer::cast_chunks, and Tensor::reinterpret reinterpreted device bytes as any element type of the right size and alignment, which let a u8 buffer become a DeviceBuffer<bool> or Tensor<bool> holding values other than 0 and 1 (NVIDIA#1506). The casts now require bytemuck::AnyBitPattern of their target, so bool is rejected at compile time, and bool implements DeviceCopy: a typed buffer holds valid values by construction, so bool buffers can be allocated, filled, launched, and read back like any other element type. cuda-core re-exports bytemuck, implements Pod for the DType newtypes, and enables half's bytemuck feature for f16 and bf16. The atomic wrappers in cuda-device are not Copy, so they cannot be AnyBitPattern. DeviceBuffer gains wrap_elem and peel_elem, bounded on bytemuck::TransparentWrapper, to view a scalar buffer as its atomic wrapper and back. cuda-device implements TransparentWrapper for its atomic types and Pod for its vector types, and the examples that cast into atomics use the new methods. Every example lock gains the bytemuck dependency of cuda-device. Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
…nd SIMT Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
…uncher Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
…eBuffers Signed-off-by: nihalpasham <nihalp@nvidia.com>
…viceBox::from_raw_parts and update docs Signed-off-by: nihalpasham <nihalp@nvidia.com>
…NVIDIA#1602) Preserve source locations and inline caller chains for const-generic arrays, generated reduction bodies, and checks relocated to the host. Keep deduplicated check origins private to the compiler and render them through the existing diagnostic API. Part of NVIDIA#1475. The remaining frontend items and debugger/backend fixes remain tracked separately; this does not close the full source-location audit. Validated with the CPU/assembler suites, source-location regressions, live CUDA-GDB tests, GPU host-check enforcement, host/Tile semver checks, and all 25 active CI checks on 825de3e. Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
…NVIDIA#1602) Preserve source locations and inline caller chains for const-generic arrays, generated reduction bodies, and checks relocated to the host. Keep deduplicated check origins private to the compiler and render them through the existing diagnostic API. Part of NVIDIA#1475. The remaining frontend items and debugger/backend fixes remain tracked separately; this does not close the full source-location audit. Validated with the CPU/assembler suites, source-location regressions, live CUDA-GDB tests, GPU host-check enforcement, host/Tile semver checks, and all 25 active CI checks on 825de3e. Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com> (cherry picked from commit 9718b43) Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
elibol
requested review from
binarybana,
cryos,
kkraus14,
nihalpasham,
roivanov and
rparolin
as code owners
October 22, 2026 22:49
This was referenced Oct 22, 2026
…ybitpattern-casts Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com> # Conflicts: # cutile-rs/CHANGELOG.md
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
…/shared-cuda-core Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com> # Conflicts: # cutile-rs/cutile/src/tensor.rs
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com> # Conflicts: # cutile-rs/cutile/src/tensor.rs
elibol
marked this pull request as draft
October 22, 2026 23:16
|
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
…bug-integration Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Advance the shared host and Tile crates and their path dependency constraints to v0.5.0. Update local package versions in the affected lockfiles without changing registry dependencies. Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
…NVIDIA#1598) ## What this adds A preview of CuTe kernels in Rust, under `cuda-oxide/cute-rs-preview`. Layouts, swizzles, operand roles, TMA pipelines and tensor-core MMA are Rust types; the compiler decodes them, verifies the whole module, and hands CUTLASS MLIR to the pinned CUTLASS 4.7 compiler. Two targets, `sm_100a` and `sm_120a`; everything else is rejected before compilation. ```text before CuTe kernels need Python: @cute.jit traces live MLIR wrappers at run time after cargo oxide run blockscale_gemm_cute --arch sm_120a Rust types -> dialect-cute -> whole-module verification -> CUTLASS 4.7 -> cubin in the binary ``` ```rust type ASmem = Composed<Swizzle<3, 3, 3>, 0, RowMajor<128, 64>>; #[kernel] pub fn gemm(#[grid_constant] a: &TmaDesc<f16, ASmem>, m: u64, n: u64, k: u64) { // the same ASmem type encodes the host descriptor and checks the device copy } ```  - Four crates: `cute-rs` (tensors, tiles, copies, pipelines, tiled MMA), `ctm-rs` (typed instruction primitives: TMA, mbarrier, tcgen05, TMEM, mma.sync), `cute-layout` (layout algebra and compile-time proofs) and `dialect-cute` (the Pliron dialect and verifier). - Compiler side: `mir-importer` recognises `cute-rs` calls by definition path and emits `dialect-cute`; `cuda-oxide-mlir-export` maps MIR, NVVM and CuTe operations to CUTLASS MLIR through the published `pliron-mlir-interop` printer; `cuda-oxide-codegen` drives the pinned `libCutlassCompiler.so` and embeds the validated cubin. `cargo oxide toolchain install cutlass` installs the compiler and checks its digests. - 14 example crates (10 on `sm_120a`, 4 on `sm_100a`) and two end-to-end models on `sm_120a`: Qwen3-TTS 1.7B and Qwen3-8B in BF16 and NVFP4. - Four new intrinsics: the `m16n8k64` MXF4 register MMA and three relaxed cluster barrier operations. - Book chapter "CuTe in Rust (preview)" with a support matrix against CuTeDSL 4.7, build and contracts pages. The preview covers a validated subset of CuTeDSL 4.7 on the two targets, not feature parity: one of eleven tcgen05 MMA kinds, three of seven warp-level kinds, 32x32b TMEM only, fixed epilogues, a static scheduler. The support matrix page lists every row with what the preview has. ## What we fixed - The CuTe stack folder is named `cute-rs-preview` and the word preview appears where readers look: the folder README, the book chapter and index card, the four crate descriptions. - Two book snippets quoted a producer method that does not exist; they now show the acquire, elected expect and barrier sequence the kernel uses. - The elementwise README selected no backend and no target; the block-scaled GEMM README claimed `sm_121`. - The stack's docs folder (compiler history, a superseded Pliron 0.19 experiment, duplicated figures) is gone; two pages moved into the book and one paragraph into the mlir-export README. <details><summary>Implementation</summary> - `cuda-oxide-mlir-export/src/profile.rs` holds the target gate: `CutlassFullCuteMlir22::new` accepts `sm_100a` or `sm_120a`; `tcgen05` operations are refused on `sm_120a` and MXF4 register MMA off `sm_120a`. The primitives-only profile `cutlass-primitives-mlir22` has the same gate. - `dialect-cute/src/verify/` runs before any MLIR is built: tensor provenance, pipeline lifecycle, cursor provenance, scheduler grids, MMA placement, epilogue phases, and a ghost-escape audit so compiler-only types never reach runtime storage. - `mir-importer/src/translator/cute/` is a closed world: 88 recognised entry points over 60 operation kinds, 15 audited source-abstraction modules; any other `cute_rs::` call is an error. - The merge with `main` renumbered the four new ABI ledger entries to `i1039` to `i1042` behind main's `i1030` to `i1038`, kept main's lockfiles and let cargo add the new crates, and regenerated every catalog-derived file; `check`, `probe --all` and `check-abi-history --base-ref main` pass. - The folder rename is one scripted commit; package names did not change, so `cute_rs::` definition paths are unchanged. </details> ## Verification Checked at `56384d465`: - ✅ `just fmt-check`, `just clippy`, `just test`, `just test-cuda`, `just check-guards`, `just doc-check` and `just check-intrinsics origin/main` pass on the merged tree. - ✅ `cargo test --workspace --all-targets --no-fail-fast` over 115 test binaries, `cargo test -p oxide-artifacts --features object`, and the nested `rustc-codegen-cuda` workspace pass. One pre-existing race, `ptx::tests::debug_builds_fail_when_opt_drops_the_debug_graph`, failed once under parallel load (a just-written temp file reported busy) and passed three times alone. - ✅ The generator gates: outputs current, 2,427 probe configurations across 1,042 routes, ledger append-only against `main` with 1,042 entries. - ✅ `scripts/sync-example-locks.sh --check` after aligning the 13 cute example lockfiles with the standard examples. - ✅ `scripts/check-primitives-api-compile.py` for the `api`, `fp16` and `blockscale` fixtures, and the 13 CUTLASS consumer tests with the pinned 4.7.0 library. - ✅ The book builds with warnings as errors; 436 relative links resolve; the source-reference, host-API-path, SPDX and crate-inventory guards pass. - ✅ On an RTX 5090 through CUTLASS 4.7: `elementwise_cute`, `rmsnorm_cute`, `swiglu_cute`, `ffn_cute`, `rope_cache_cute`, `blockscale_gemm_cute`, `blockscale_gemm_ctm`, `nvfp4_gemv_cute`, `fmha_prefill_sm120a_cute` and `fmha_decode_sm120a_cute` run and pass their reference comparisons; the FMHA pair also times against CuTeDSL 4.7.0 (prefill 18.0 us against 23.1 us, decode 4.4 us against 12.0 us). The four `sm_100a` examples compile. `vecadd`, `mma_mxf8f6f4` and `cluster` run on the native path. - ⏳ Not run locally: `scripts/smoketest.sh --compile-only`, which publishes to the shared backend cache; the examples-compile lane covers it here. ## Caveats - Two targets only. `sm_100f`, `sm_103a`, `sm_110a`, `sm_120f`, `sm_121a`, Hopper and Ampere are rejected at the profile gate. - CuTe kernels build only with `CUDA_OXIDE_DEVICE_BACKEND=cutlass-mlir`; the native backend has no CuTe lowering and currently stops with a verification error rather than a dedicated message. - NVFP4 in the Qwen3-8B model runs on CUDA cores through scaled loads and `cute.dot`; the NVFP4 tensor-core atom exists in the dialect and exporter but has no producer yet. - No CI job compiles the cute examples; the unit-test, clippy and fmt lanes cover the crates and the example manifests, and GPU runs are manual. - APIs can change between releases. The compatibility aliases in `cute-rs` and `dialect-cute` stay for now. --------- Signed-off-by: nihalpasham <nihalp@nvidia.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
Signed-off-by: Melih Elibol <1884921+elibol@users.noreply.github.com>
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Draft: release integration remains on hold while launch-overhead follow-up is assessed.
Integrate the shared host runtime and accepted source-location fixes into
release/v0.5.xfor v0.5.0. The source tree matches #1604; the only additional changes enable the existing CI workflows for PRs targetingrelease/v*. Push and deployment triggers are unchanged.Included changes:
93dd50704): independent DeviceCopy/DType contracts, valid reinterpretation targets, host argument contracts, and correct tensor alias accounting.47042a8fa): shared contexts, streams and allocation ownership; removecuda_core::simtand fix allocation uniqueness accounting.b811b7067): shared async operations and launches; removecuda_async::simt. Includes Nihal's buffer/DeviceBox read-permission fixes.9718b4386: source origins for const-generic arrays, generated reductions and host-relocated/deduplicated checks.0143af5cc: the shared host and Tile v0.5.0 version baseline.The three host-runtime PRs retain their separate review flow and are stacked in the order #1516 → #1603 → #1604. Remaining debugging candidates are tracked in #1606; CUDA-GDB #29 and the Tile IR assertion-trap/coredump issue #1379 remain external dependencies.
Validation before the latest CuTe integration:
The tensor/allocation uniqueness races are fixed. The original cutile-rs caller-managed
unsafe async_on()behavior is restored: no automatic retention after return, while safe futures keep nonblocking completion-based cleanup. The existing individual-launch benchmark improves from 4.95 to 1.88 µs (earlier main baseline: 1.68 µs). No batching or new launch API is needed for this fix.Related: #1460, #1461, #1475, #1506, #1606.
Current-main integration: includes the CuTe preview import, reconciled dependency lockfiles, and its host consumers migrated to the shared core API. The combined cuda-host/cuda-macros/ctm-rs/cute-rs suites pass (497 tests, 43 intentionally ignored), and all 16 CuTe example/model workspace Clippy checks pass. The latest integration at
038c82d2c3includes the async_on fix fromb811b7067; fresh CI is running.Validation of the async_on restoration: 133 cuda-async tests pass (14 intentionally ignored), all 11 Tile submission-lifetime tests and four SIMT lifetime/concurrency GPU tests pass, and Clippy/formatting/rustdoc checks pass.