Skip to content

codegen: support grid-constant kernel parameters - #1225

Merged
nihalpasham merged 6 commits into
NVlabs:mainfrom
lucifer1004:feat/grid-constant-tma-1223
Sep 18, 2026
Merged

nihalpasham merged 6 commits into
NVlabs:mainfrom
lucifer1004:feat/grid-constant-tma-1223

Conversation

@lucifer1004

@lucifer1004 lucifer1004 commented Sep 3, 2026 •

Copy link
Copy Markdown
Contributor

What this adds

TMA descriptors now travel with the kernel launch, removing their separate GPU allocation/upload in tma_copy.

before: launch carries a pointer to a descriptor in GPU memory
after:  launch carries the descriptor's 128 bytes
#[kernel]
pub fn copy(#[grid_constant] descriptor: &TmaDescriptor) {
    // Use the descriptor or pass it to a device helper.
}

The host passes a Copy descriptor by value; device threads borrow one read-only copy shared by the grid for that launch.

What we fixed

  • Legacy NVVM now preserves the 128-byte descriptor and its alignment, instead of emitting a one-byte parameter.
  • Argument positions stay correct when slices expand or zero-sized arguments disappear. Generic helpers retain their ordinary pointer ABI.
  • Invalid layouts, mutable parameter storage and invalid outer lifetimes fail compilation.
  • The TMA example exposed an existing legacy address-conversion bug; the fix preserves cluster-shared addresses correctly.

Launch contract

  • Grid-constant host launches require unsafe, including prepared/async calls, numeric payloads and kernels declared safe.
  • Copying a descriptor does not own its GPU allocation. The caller keeps referenced memory GPU-accessible, alive and correctly synchronized through completion.
  • The blanket rule is a conservative API choice, not a CUDA requirement for numeric data.
Implementation
  • Macro/MIR carry the entry parameter's type, index and Rust alignment. Checks inspect specialized types and original reference lifetimes.
  • Lowering shares argument mapping with existing reference/alignment handling and rejects unsupported storage, including target-dependent shared-pointer fields.
  • LLVM emits byval and NVVM grid_constant metadata. Legacy declarations, symbol references and entry adapters retain the complete pointee type.
  • All generated grid host launchers share the unsafe policy, including both async forms.
  • The legacy TMA fallback uses cvta.to.shared::cluster, avoiding unsupported LLVM address-space handling while preserving cluster semantics.

Verification

Checked at f27fd082:

  • ✅ just check: 5,393 tests/doctests, formatting, strict Clippy, guards and docs. Book passes with warnings as errors.
  • ✅ All 53 hosted checks passed; PR book deployment was intentionally skipped.
  • ✅ tma_copy memcheck/synccheck on RTX 5090 through LLVM NVPTX, modern libNVVM and legacy libNVVM: 4,096 copied values and all 256 pipeline threads verified.
  • ✅ GPU/PTX checks for mixed arguments, generic helpers, descriptor layout and parameter positions.
  • ✅ Compile-only probes reject all 15 unguarded grid launch variants and accept their explicit-unsafe counterparts.

Caveats

  • Native compiler only; the experimental CUTLASS exporter needs separate integration. No speedup has been measured.
  • Use &T or &'_ T for the parameter; named or 'static outer lifetimes are rejected.
  • Legacy GPU runs used SM90 output and CUDA 13.0 for compatibility with the installed driver.
  • Ordinary KernelScalar: Copy launches retain an existing nested-reference safety gap; this PR does not redesign that API.
  • Multicast: modern libNVVM memcheck passes, but synccheck reports 32 existing divergent-barrier errors with both old and repaired lowering. Native LLVM passes; legacy remains blocked by unchanged cluster-register intrinsics.

Closes #1223.

@nihalpasham

Copy link
Copy Markdown
Collaborator

The generic-kernel path needs to be part of the design before this lands:

  • For a generic #[kernel], the macro splits your fn into a generated #[inline(always)] helper that keeps the body and a concrete entry wrapper per instantiation that calls it. The marker lands in both. When rustc's MIR inliner folds the generated helper into the entry, detect_grid_constant_params hard-errors on the duplicate.
  • Repro:
#[kernel] 
pub fn k<T: Copy>(#[grid_constant] d: &Desc, out: *mut u32, _tag: T)  { 
   unsafe { *out = d.values[3] } 
 } 

fails with kernel contains duplicate grid-constant marker for source parameter 0. Bigger bodies only compile because they exceed the inliner's cost threshold.

  • More importanly, the helper's marker leaks: a plain kernel
other(d: &Desc, out: *mut u32) { 
  k::<u8>(d, out, 0) 
} 

with no attribute comes out with .param .align 64 .b8 other_param_0[128] and a grid_constant annotation. Host passes an 8-byte pointer, device expects 128 bytes by value.

  • unchecked_indexing already handles this by stripping its marker from the re-emitted helper, see strip_unchecked_indexing_config_marker. grid_constant needs the same in both generic paths. A duplicate-tolerant detector alone fixes the error but not the leak.

And we need tests in the PR exercise a generic kernel with #[grid_constant].

To land this: the strip in both generic paths, a generic #[kernel] + #[grid_constant] test with a body small enough to be inlined, and a negative test that a non-opted kernel calling the helper still gets a plain pointer param.

@nihalpasham nihalpasham added enhancement New feature or request codegen Device code-generation pipeline (Rust MIR to IR to PTX) cuda-feature A CUDA hardware/toolkit feature (TMA, tcgen05, WGMMA) labels Sep 6, 2026

@nihalpasham nihalpasham left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Changes requested, details in #1225 (comment)

@lucifer1004

Copy link
Copy Markdown
Contributor Author

Addressed the generic-helper ABI concern in f2b6b15.

The macro now captures the grid-constant marker for the generated entry and strips it from the re-emitted callable helper in both generic expansion paths. The contract example now covers both sides:

  • an instantiated generic grid-constant entry retains the 128-byte by-value PTX parameter and launches successfully on sm_120;
  • an ordinary kernel calling the same generic helper retains a .param .u64 pointer ABI and does not acquire the 128-byte parameter.

Measured validation:

  • pixi run cargo test -p cuda-macros
  • strict clippy for cuda-macros and cuda_module_contract
  • pixi run cargo oxide run cuda_module_contract --arch sm_120 -- --verify-ptx
  • pixi run cargo oxide run cuda_module_contract --arch sm_120
  • the full just check code/test portion passed; its cargo-deny guard initially hit the local NFS advisory-lock limitation, then all deny/license guards passed with the advisory DB placed on local /tmp storage.

@nihalpasham nihalpasham added llvm-export Textual LLVM/NVVM IR exporter (llvm-export crate) miscompile Produces incorrect PTX/IR (wrong generated code) needs-changes Review found changes required before this PR can land labels Sep 15, 2026

@nihalpasham nihalpasham left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Reviewed f2b6b150: changes requested.

  • The generic marker repair is present. The legacy NVVM path still loses the descriptor's pointee type:
host payload: 128 bytes -> i8* byval -> PTX parameter: 1 byte
  • Preserve the aggregate type through the legacy signature and body references, and add a real legacy NVVM control. The modern path passing does not cover this.
  • Define the read-only/lifetime contract across entry routes. Give the generic GPU test one writer; it currently has 32 threads writing one element.
  • Sign the revised stack and restore its missing DCO trailers.

Independent review and libNVVM output confirm the ABI blocker. #1223 remains open.

lucifer1004 and others added 5 commits September 17, 2026 18:33
Let #[kernel] declarations mark immutable reference parameters as grid constants, derive the generated host launch ABI from the same declaration, and carry pointee layout through MIR and LLVM lowering.

Emit LLVM byval alignment and NVVM grid_constant metadata, with compiler and SM120 contract coverage.

Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>
Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>
Convert generic destinations explicitly with cvta.to.shared::cluster in inline PTX instead of constructing AS7, which legacy NVVM does not support. Retain the LLVM intrinsic path and test all G2S ranks and multicast forms.

Signed-off-by: nihalp <nihalp@nvidia.com>
Retain complete typed byval pointees in legacy NVVM signatures, symbol references and entry adapters. Share physical parameter mapping with reference validity and reject malformed metadata and target-dependent storage layouts.

Check monomorphized Freeze and launch-scoped reference lifetimes, keep ABI markers out of callable helpers, and require unsafe use of the hidden ABI marker. Exercise mixed parameter shapes, preserved helper calls, real TMA descriptors and rejection paths across the native backends.

Signed-off-by: nihalp <nihalp@nvidia.com>
Reject an ordinary device-extern declaration that would suppress a grid-constant kernel declaration after erased-pointer shape matching. Cover both modern and legacy public exporter paths.

Signed-off-by: nihalp <nihalp@nvidia.com>
@nihalpasham
nihalpasham force-pushed the feat/grid-constant-tma-1223 branch from f2b6b15 to 7853230 Compare September 17, 2026 13:37
@nihalpasham nihalpasham added maintainer-owned and removed needs-changes Review found changes required before this PR can land labels Sep 17, 2026
Copy and immutable parameter storage do not prove device validity, lifetime or synchronization of nested references. Keep that obligation explicit at every grid host launch boundary, including prepared sync, borrowed async and owned async methods.

Preserve device function signatures and scalar marshalling. Add compile-fail and positive controls, generated safety documentation, and explicit contracts at the example call sites.

Signed-off-by: nihalp <nihalp@nvidia.com>
@nihalpasham nihalpasham added the safety Memory safety, soundness, or undefined behavior label Sep 17, 2026
@nihalpasham
nihalpasham dismissed their stale review September 18, 2026 02:17

maintainer fixed

@nihalpasham

nihalpasham commented Sep 18, 2026 •

Copy link
Copy Markdown
Collaborator

I added the needed follow-ups to the descriptor-by-value support:

  • Fixed legacy NVVM loosing the descriptor's size: a 128-byte parameter was becoming one byte.
  • Kept argument positions and alignment consistent when slices expand or zero-sized arguments disappear; generic helpers retain their normal pointer ABI.
  • Added layout, read-only and lifetime checks so unsupported payloads fail compilation.
  • Made grid-constant host launches unsafe, including prepared/async calls. The caller must keep referenced memory GPU-accessible, alive and properly synchronized. This conservative rule also covers numeric payloads and kernels declared safe.
  • Switched tma_copy to pass descriptor bytes with the launch, removing the separate descriptor allocation/upload. Also fixed the existing legacy TMA address-conversion bug this exposed.
  • Added regressions for mixed arguments, generic helpers, invalid payloads and the unsafe launch requirement.

At f27fd082, full local checks and all 53 CI checks passed. tma_copy also passed memcheck/synccheck through all three tested compiler paths.

This addresses #1223 for the native backend. CUTLASS integration is separate; no speedup has been measured yet.

@nihalpasham
nihalpasham merged commit e01c241 into NVlabs:main Sep 18, 2026
54 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

codegen Device code-generation pipeline (Rust MIR to IR to PTX) cuda-feature A CUDA hardware/toolkit feature (TMA, tcgen05, WGMMA) enhancement New feature or request llvm-export Textual LLVM/NVVM IR exporter (llvm-export crate) maintainer-owned miscompile Produces incorrect PTX/IR (wrong generated code) safety Memory safety, soundness, or undefined behavior

Projects

None yet

Development

Successfully merging this pull request may close these issues.

Support grid-constant by-value kernel parameters for TMA descriptors

2 participants