[elibol]

What this does

This gets cuda-oxide and cutile-rs onto the same three host-crate package identities before the monorepo move, without making high-level API unification part of that move.

The shape

cutile-rs#233 temporarily carries both API families:

shared cuda-core
├── cuda_core::*          Cutile / Tile API
└── cuda_core::simt::*    cuda-oxide SIMT API

shared cuda-async
├── cuda_async::*         Cutile / Tile API
└── cuda_async::simt::*   cuda-oxide SIMT API

This PR points cuda-oxide at those shared crates:

Before                         After

cuda-oxide                     cuda-oxide
  └─ local host crates   ──►      └─ shared host crates ◄── cutile-rs

This is a namespaced compatibility bridge. It does not rewrite cuda-host or cuda-macros around the Tile APIs, and it does not decide the eventual unified API.

Most Rust edits are path changes. For example:

// before
cuda_async::device_operation::DeviceOperation

// after
cuda_async::simt::device_operation::DeviceOperation

One mechanical edit needed semantic correction: cutile_inter_kernel composes a Tile DeviceOp, so its five Tile-side imports must remain under root cuda_async. That is fixed at the current head, and the example now compiles and runs on an RTX 5090.

What is actually still missing upstream

• Peer access: carry can_access_peer, enable_peer_access, and disable_peer_access, plus cuda-oxide's real two-GPU VMM/P2P test. The VMM implementation itself is already in #233; peer access and the cross-GPU test are the missing pieces.

• DeviceCopy derive: preserve the existing one-import API, or agree on a compatible replacement, and carry its trybuild tests:

  use cuda_core::DeviceCopy; // imports the trait and derive today
  
  #[derive(DeviceCopy)]
  struct Pair {
      x: f32,
      y: f32,
  }

One compatibility question, not a merge gate

Primitive Rust f16 is nightly-only, so making its DeviceCopy impl an opt-in feature in a stable shared crate is reasonable. Can we confirm that this is the intended policy?

• half::f16 remains available normally.
• atomic_f16 and f16_stress explicitly enable the new f16 feature.
• gemm_sol_final, the tcgen05 examples, and the matmul examples do not depend on primitive f16 or this feature.

Compatibility and release gates

• CUDA/toolkit support: preserve cuda-oxide's CUDA 12.x/13.1 and toolkit-layout support, or explicitly agree that the shared crates now require CUDA 13.2+ and Cutile's toolkit layout.

• Use one published oxide-artifacts crate: oxide-artifacts 0.2.1 is now published. Make both cuda-oxide and the shared cuda-core use that release. The shared cuda-core should re-export the dependency instead of copying its source into cuda_core::simt::artifacts. The copied reader understands the writer's v2 bytes today, but it is a second implementation and its public Rust types have a different identity:

  let original: oxide_artifacts::OwnedArtifactBundle = todo!();
  let _: cuda_core::simt::artifacts::OwnedArtifactBundle = original;
  // E0308 today: same fields and byte format, but different Rust types

  Keeping both projects on the same crates.io package gives the compiler and shared runtime one Rust type and one format implementation. No copied implementation, Git dependency, or transitional Cargo patch is needed.

• Preserve the bindings C function type: ordinary calls through the shared dynamic loader work, but the generated shim currently changes the function item's Rust type from extern "C" to the Rust ABI:

  unsafe { cuda_bindings::cuInit(0) }; // works today
  
  type CuInit = unsafe extern "C" fn(u32) -> cuda_bindings::CUresult;
  let _: CuInit = cuda_bindings::cuInit;
  // E0308 today: expected extern "C" fn, found Rust fn

  Preserve bindgen's ABI when generating the public loader shims and keep this assignment as a compile-time regression test. Dynamic loading remains unchanged.

• Scaffolds and docs: update cargo oxide new, repository/book docs, and any examples that still point at the retired local crates.

• Release/reproducibility: publish oxide-artifacts 0.2.1 and cuda-bindings 0.3.1 first, followed by cuda-core 0.3.1 → cuda-async 0.3.1. Verify each package, pin the released versions, regenerate affected locks, and keep DCO clean.

• Final runtime gate: run the peer/VMM P2P path on both GPUs and keep the representative SIMT, async, cooperative, cluster, and interop examples green at the final release head.

What is green now

• Hosted CI is green, including the full 216/216 compile-only example matrix.
• Shared cuda-core and cuda-async: 196 passed, 1 ignored on CUDA 13.3 / RTX 5090.
• Cutile GPU script: 187 tests passed on the two RTX 5090s against #233.
• cuda-oxide runtime checks passed: vecadd, async_vecadd, cluster, coop_groups_demo, atomic_f16, and cutile_inter_kernel.
• nvls_all_reduce still lives in cuda-oxide, builds against the shared VMM surface, and correctly skips on these 5090s because they do not have NVSwitch multicast support.
• cutile-rs#202 now separately provides Cutile's root VMM API and runs its vmm / vmm_multicast targets from the GPU script. That is useful shared coverage, but it is not a substitute for the missing peer/P2P test.

Merge picture

Already green
  216/216 compile cases + hosted CI + representative RTX 5090 runs
                              │
                              ▼
Close the two API gaps
  peer/P2P + DeviceCopy derive compatibility
                              │
                              ▼
Fix the two shared boundaries and update scaffolds/docs
  publish/reuse oxide-artifacts + preserve the bindings C ABI
                              │
                              ▼
Publish and pin the verified releases; DCO clean
                              │
                              ▼
                            merge

Once those items are closed, this remains a packaging move rather than a capability trade. Safety issues already present in either implementation remain separately tracked work; this migration does not claim to solve them.

[nihalpasham]

I updated the PR description to be clearer and more precise about what is already working and what is still needed.

I also fixed the one failing example. cutile_inter_kernel was accidentally pointed at the SIMT async surface even though it composes a Tile DeviceOp; those five imports are back on root cuda_async. The current head is green in hosted CI, including all 216 compile cases, and the example compiles and runs on an RTX 5090.

I published oxide-artifacts 0.2.1 from cuda-oxide. Could you update #233 so the shared cuda-core depends on and re-exports that crate instead of carrying a copied simt::artifacts implementation?

The remaining upstream items I see are:

• add the peer-access wrappers and the existing two-GPU vmm_p2p test;
• preserve the DeviceCopy derive re-export through a publishable solution and carry its trybuild tests;
• preserve the bindings extern "C" function type and add the function-pointer assignment regression;
• explicitly agree whether the shared releases preserve cuda-oxide's CUDA 12.x/13.1 and toolkit-layout support, or intentionally move to CUDA 13.2+.

The VMM implementation itself is already present in #233. I am also treating primitive f16 being opt-in as policy for nightly additions.

Once those upstream pieces are in, I will replace the temporary branch dependencies here with the published 0.3.1 crates, regenerate every affected lockfile, update cargo oxide new and the remaining docs/examples, and rerun the final single-/two-GPU matrix. We should also clean up the missing DCO trailers before release.

If we close those items without dropping existing cuda-oxide capability, I think this is in good shape to merge.

[elibol]

Thanks for the fixes and for publishing oxide-artifacts. The interop example is looking good.

Follow-ups now on #233:

• cuda-core depends on oxide-artifacts 0.2.1 and re-exports it as simt::artifacts.
• Peer wrappers are in, along with the two-GPU vmm_p2p and vmm_multicast tests. They self-skip on single-device machines.
• #[derive(DeviceCopy)] moved to a new cuda-core-derive crate so we can publish it. It's re-exported next to the trait, so use cuda_core::DeviceCopy; brings both into scope. The trybuild suite came with it.
• The bindings' loaded driver table is now public. Its fields are the raw unsafe extern "C" fn pointers, so the function-pointer type is preserved without changing the flat wrapper API. Added a regression test for the assignment.

Re: the CUDA floor: the shared host-side crates keep cuda-oxide's 12.x support using your probe approach. The bindings floor drops to 12.0, the elapsed-time helper goes back to the v1/v2 dispatch under the probed cfg, multicast goes back to a header probe (your cuda-core build script, verbatim), and CI here goes back to 13.1.1. cutile enforces its own 13.2 floor in cutile-compiler at tool discovery, so SIMT users below that never see it. We added a 12.x compile lane for the shared crates so we don't silently break 12.x.

One thing we should settle now rather than later: the cuTe DSL will have similar versioning constraints to cuTile (it has a tileiras-like binary for compiling cute dsl). I think the pattern should be per-stack floors enforced at each stack's tool boundary, with the shared crates at the lowest common floor. Does that match what you expect?

I'll clean up the DCO trailers before release.

[elibol]

Correction on the floor: 12.0 is lower than what the current wrappers support. The 12.x lane found what the shim calls directly: cuCtxSetFlags (12.1), cuMemAdvise_v2 and cuMemPrefetchAsync_v2 (12.2), cuEventElapsedTime_v2 (12.8), and cuLaunchHostFunc_v2 plus the CUhostTaskSyncMode enum (13.x only). So the floor is 12.8 for now. The 13.x-only host-function launch goes through a probed helper in cuda-bindings. On older toolkits it falls back to the blocking default instead of spin-wait.

Going below 12.8 will require some work... We'd need to route the elapsed-time call through the existing helper and add fallbacks for the three 12.1/12.2 symbols. The commits list them.

Will cuda-oxide require CUDA versions < 12.8?

[nihalpasham]

Thanks, I went through the current #233 head again.

CUDA 12.8+ is fine for cuda-oxide; we do not need support below 12.8. The oxide-artifacts, peer/P2P, and DeviceCopy derive changes look good.

The only remaining changes we'd need on #233 are:

• CUDA toolkit discovery: Please copy cuda-oxide’s existing lookup logic and tests into the shared cuda-bindings. We have users with custom, Jetson/Grace, cross-build, and targets/<arch> toolkit layouts. Those work in cuda-oxide today, but Add cuda-oxide host surface to host-side crates NVlabs/cutile-rs#233 only accepts <toolkit>/include/cuda.h, so these users would fail during the build.

• Shared-crate tests: The current cargo test --no-run step compiles the copied tests but does not execute their assertions or trybuild cases. Can we please wire the CPU-safe cuda-core/cuda-async tests into the CPU test path and the GPU-dependent tests into the GPU path, including simt_vmm_p2p and simt_vmm_multicast. This is ordinary Cargo testing; it does not involve cargo oxide. I will still run the final integrated single-/two-GPU matrix from Use the shared cutile-rs host crates in cuda-oxide #1102.

Could you also refresh the #233 description? It still describes peer/VMM as omitted, oxide-artifacts as vendored, and the derive as remaining in cuda-oxide.

Once those items and the planned DCO cleanup are done, I think #233 is ready to be the upstream release piece. I’ll handle the dependency switch, locks, local artifact identity, scaffolds/docs, and final integration here after 0.3.1 is published.

[elibol]

@nihalpasham any thoughts on how to address the CodeQL alerts?

[nihalpasham]

@nihalpasham any thoughts on how to address the CodeQL alerts?

@elibol - The 7 new alerts are pre-existing main-branch alerts (all 13 are open on main, created by the Aug 28 scan) that CodeQL misattributed to this PR because the diff was large.

They're all rust/access-invalid-pointer on #[cuda_module] kernel code, i.e. device kernels dereferencing GPU buffer pointers by design. We'll dismiss them as false positives / exclude example kernels from the CodeQL config on our side; nothing to change in this PR, and CodeQL isn't a required check.

[nihalpasham]

@elibol - where we are, after a full review and a runtime pass:

• The switch is correct: every path resolves on v0.3.1, 222/222 examples pass, and the simt GPU suites pass here too (two-GPU P2P included).
• I rebased the branch onto current main (your 17 commits, same authorship) and pushed four commits on top:
  • cargo oxide new templates and the docs sweep now use the simt paths and the crates.io release
  • the retired local cuda-core / cuda-async / cuda-bindings are deleted, with the CI plumbing and license guards that referenced them
  • oxide-artifacts is consumed from crates.io everywhere (one crate for writer and reader), with a parity guard
  • two more guards: host-API paths in docs/templates, and shared-crate pin parity across the example manifests
• CodeQL is fixed at the root on main (cuda-macros: keep generic kernel specializations alive with black_box, not a null-pointer volatile pair #1226), so the check is clean here now.

What we need for 0.3.1:

• publish the whole set to crates.io, cutile crates included
• leave v0.3.1 alone once published; any fix is a v0.3.2
• a git rebase --signoff on this branch after my last push, since 14 commits have no Signed-off-by (not a blocker)

Then I repoint the dependency lines to the release, relock, rerun the smoketest, and we merge.