Skip to content

draft(mma): investigate counted register-MMA preservation #1237

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

Closed
lucifer1004 wants to merge 2 commits into from
Closed

draft(mma): investigate counted register-MMA preservation #1237

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
64e23e1
feat(wmma): add native shared-address ldmatrix forms
lucifer1004 Sep 5, 2026
4284c72
refactor(wmma): name ldmatrix u32 carriers explicitly
lucifer1004 Sep 5, 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 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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/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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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/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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// 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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::warp` to keep existing paths stable.
Expand Down
2 changes: 1 addition & 1 deletion crates/cuda-device/src/generated/wgmma_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: 20bc41c2ab93828b7f97b6ee90efccfc118317a590803216a3a045b1c2879889.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3df52944985ca17502f6d46dac02c640ef6a6abb7784a1f72fb0dced7c24e3ad.
// LLVM source: 16696adcd119e6ba9cc175207d984d7021211acb.

// Included inside `cuda_device::wgmma` to keep its public API stable.
Expand Down
20 changes: 20 additions & 0 deletions crates/cuda-device/src/shared.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -526,4 +526,24 @@ pub unsafe fn cvta_generic_to_shared_offset(ptr: *const u8) -> u64 {
unreachable!("cvta_generic_to_shared_offset called outside CUDA kernel context")
}

/// Convert a generic address into its 32-bit `.shared::cta` state-space address.
///
/// Unlike [`cvta_generic_to_shared_offset`], which returns the `u64` carrier
/// used to construct WGMMA and tcgen05 descriptors, this form matches the
/// `u32` address operand consumed by CTA-shared instructions such as `ldmatrix`.
/// Convert a shared base once, then keep byte-offset and swizzle arithmetic in
/// `u32` when the complete addressed allocation fits the CTA shared window.
///
/// # Safety
///
/// `ptr` must be a generic pointer to memory in the current CTA's shared-memory
/// window. Converting any other pointer produces an unspecified address.
//
// Keep this uninlined: mir-importer intercepts the exact rendered def-path.
#[inline(never)]
pub unsafe fn cvta_generic_to_shared_u32(ptr: *const u8) -> u32 {
let _ = ptr;
unreachable!("cvta_generic_to_shared_u32 called outside CUDA kernel context")
}

include!("generated/shared_sreg.rs");
Loading