Skip to content

Implement generated mbarrier_arrive_no_complete intrinsic #1101

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
nihalpasham merged 3 commits into from
Aug 30, 2026
Merged

Implement generated mbarrier_arrive_no_complete intrinsic #1101

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
9231490
feat(intrinsics): route mbarrier_arrive_no_complete through the gener…
uurl Aug 30, 2026
2026dcb
chore(intrinsics): regenerate catalog for mbarrier_arrive_no_complete…
nihalpasham Aug 30, 2026
594b082
docs(book): re-stamp the dialect page for the i1017 catalog
nihalpasham Aug 30, 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
  •  
  •  
  •  
26 changes: 0 additions & 26 deletions crates/cuda-device/src/barrier.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -106,32 +106,6 @@ include!("generated/mbarrier_basic.rs");
include!("generated/mbarrier_extended.rs");
include!("generated/counted_barrier.rs");

// =============================================================================
// Barrier Arrive Operations
// =============================================================================

/// Arrive at barrier without returning a token (fire-and-forget).
///
/// Use this when the thread will not wait on this barrier.
/// For producer threads that signal completion but don't need to wait.
///
/// # Safety
///
/// - `bar` must be initialized
/// - Only use when this thread will NOT call `mbarrier_wait`
///
/// # PTX
///
/// ```ptx
/// mbarrier.arrive.noComplete.shared.b64 _, [addr];
/// ```
#[inline(never)]
pub unsafe fn mbarrier_arrive_no_complete(bar: *const Barrier) {
let _ = bar;
// Lowered to: call void @llvm.nvvm.mbarrier.arrive.noComplete.shared(ptr %bar)
unreachable!("mbarrier_arrive_no_complete called outside CUDA kernel context")
}

// =============================================================================
// Barrier Wait Operations
// =============================================================================
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/async_copy.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::async_copy` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/atomic.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::atomic` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/bf16.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::bf16` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/bf16x2.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::bf16x2` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/clc.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::clc` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/cluster_barrier.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::cluster` to keep public paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/cluster_memory.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::cluster` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/cluster_sreg.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// This file is included inside `cuda_device::cluster`.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/convert.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::convert`.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/counted_barrier.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::barrier` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/debug_control.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::debug` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/debug_sreg.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::debug` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/dotprod.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::dotprod` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/f16.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::f16` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/f16x2.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::f16x2` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/f32x2.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::f32x2` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/fence.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// This file is included inside `cuda_device::fence` so existing paths stay stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/float.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::float`.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/grid_dependency.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::grid` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/grid_sreg.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::grid` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/i16x2.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::i16x2` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/int.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::int` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/ldmatrix.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::wmma` to keep public paths stable.
Expand Down
15 changes: 14 additions & 1 deletion crates/cuda-device/src/generated/mbarrier_basic.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::barrier` to keep existing paths stable.
Expand All @@ -22,6 +22,19 @@ pub unsafe fn mbarrier_arrive(bar: *const Barrier) -> u64 {
unreachable!("mbarrier_arrive called outside CUDA kernel context")
}

/// Arrives at a CTA shared-memory barrier by a dynamic count without completing the current phase and returns its prior opaque state.
/// Lowers to `mbarrier.arrive.noComplete.shared.b64 <register>, [<address>], <register>;`.
///
/// # Safety
/// `bar` must point to a live, eight-byte-aligned `Barrier` in shared memory.
/// The barrier must be initialized. `count` must be a valid PTX arrival count and this operation must not complete the current phase.
/// Use the returned opaque state only with this barrier and phase.
#[inline(never)]
pub unsafe fn mbarrier_arrive_no_complete(bar: *const Barrier, count: u32) -> u64 {
let _ = (bar, count);
unreachable!("mbarrier_arrive_no_complete called outside CUDA kernel context")
}

/// Initializes a CTA shared-memory barrier with the expected arrival count.
/// Lowers to `mbarrier.init.shared.b64 [<address>], <register-or-immediate>;`.
///
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/mbarrier_extended.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::barrier` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/movmatrix.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::wmma` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/prmt.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::prmt` to keep the public paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/register_control.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::thread` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/register_mma.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::wmma` to keep public paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/shared_sreg.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::shared` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/sparse_mma.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::wmma` to keep public paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/sreg.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// This file is included lexically inside `cuda_device::thread` so existing
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/stmatrix.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::tcgen05` to keep public paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/tcgen05.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::tcgen05` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/tma.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::tma` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/warp_sreg.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
*/

// @generated by cuda-intrinsics-gen 0.1.0. DO NOT EDIT.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 49bcd626b91f3261e80f77fb2817a54d96bddd6364f2564d113eee7ee62b8266.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 02a6a6b3415ab7468c40db4cf5becc52a6c53da5ef7d944933e9cfa5858b4397.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::warp` to keep existing paths stable.
Expand Down
Loading
Loading