Skip to content

feat(tma): add CTA-local G2S tile intrinsics #1336

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 20 commits into from
Oct 19, 2026
Merged

feat(tma): add CTA-local G2S tile intrinsics #1336

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
20 commits
Select commit Hold shift + click to select a range
26043ae
feat(cuda-device): add SharedPtr, a CTA-shared pointer typed addrspac…
lucifer1004 Sep 26, 2026
8287263
feat(tma): add CTA-local G2S tile intrinsics
lucifer1004 Sep 25, 2026
87b6a5a
Merge origin/main into the CTA-local TMA branch
nihalpasham Oct 19, 2026
4fd958c
fix(cuda-intrinsics-gen): gate the CTA-local TMA admission behind sha…
nihalpasham Oct 19, 2026
e786af0
docs(tma): state CTA locality in the raw-layer safety text and refres…
nihalpasham Oct 19, 2026
0f42431
chore(intrinsics): re-record the CTA TMA evidence under LLVM 23.1.3 a…
nihalpasham Oct 19, 2026
74609ee
test(examples): cover two-field and returned SharedPtr shapes in shar…
nihalpasham Oct 19, 2026
d7a3e3a
fix(intrinsics): move the CTA-local TMA entries to ABI ids i1034 to i…
nihalpasham Oct 19, 2026
b8a29a8
Merge origin/main into the CTA-local TMA branch
nihalpasham Oct 19, 2026
e20462a
chore(intrinsics): regenerate the catalog outputs after the second merge
nihalpasham Oct 19, 2026
a5f14c7
docs(book): refresh the dialect-nvvm catalog stamp after the second m…
nihalpasham Oct 19, 2026
c28c841
fix(cuda-device): make SharedArray::shared_ptr an unsafe constructor
nihalpasham Oct 19, 2026
a09a080
fix(llvm-export): refuse kernel parameters that hold shared or local …
nihalpasham Oct 19, 2026
7d79f82
test(examples): make the shared_memory_helper code check discriminating
nihalpasham Oct 19, 2026
e7d7d47
test(examples): add shared_ptr_api covering SharedPtr across helper b…
nihalpasham Oct 19, 2026
4aa4923
docs(book): state the SharedPtr limits and make the snippet compilable
nihalpasham Oct 19, 2026
5c3f168
docs(mir-importer): say why ThreadIndex needs no projection strip
nihalpasham Oct 19, 2026
a9768ec
Merge remote-tracking branch 'origin/main' into HEAD
nihalpasham Oct 19, 2026
fc7e160
docs(book): show the CTA-local TMA form next to the cluster form
nihalpasham Oct 19, 2026
fbe27a6
docs(book): wrap the CTA-form TMA snippet in unsafe
nihalpasham Oct 19, 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
  •  
  •  
  •  
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::async_copy` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::atomic` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::bf16` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::bf16x2` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::clc` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::cluster` to keep public paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::cluster` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// This file is included inside `cuda_device::cluster`.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::convert`.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::barrier` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::debug` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::debug` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::dotprod` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::f16` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::f16x2` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::f32x2` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// This file is included inside `cuda_device::fence` so existing paths stay stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::float`.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::grid` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::grid` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::i16x2` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::int` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::wmma` to keep public paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::barrier` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::barrier` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::wmma` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::prmt` to keep the public paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::thread` to keep its public API stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::wmma` to keep public paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::shared` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::wmma` to keep public paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// This file is included lexically inside `cuda_device::thread` so existing
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::tcgen05` to keep public paths stable.
Expand Down
2 changes: 1 addition & 1 deletion cuda-oxide/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: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::tcgen05` to keep its public API stable.
Expand Down
Loading
Loading