Skip to content

fix(intrinsics): admit FP16 m16n8k8 MMA on sm_75 #1484

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

fix(intrinsics): admit FP16 m16n8k8 MMA on sm_75 #1484

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
6b0a90d
fix(intrinsics): admit FP16 m16n8k8 MMA on sm_75
praveenperera Oct 10, 2026
9546a65
Merge origin/main into the FP16 m16n8k8 sm_75 floor branch
nihalpasham Oct 19, 2026
88ec9e5
fix(intrinsics): re-record the Turing FP16 MMA evidence under the cur…
nihalpasham Oct 19, 2026
ff44cfd
chore(intrinsics): regenerate the catalog outputs and refresh the boo…
nihalpasham Oct 19, 2026
72b8ead
Merge origin/main into the FP16 m16n8k8 sm_75 floor branch again
nihalpasham Oct 19, 2026
f5f35b4
chore(intrinsics): regenerate the catalog outputs after the second me…
nihalpasham Oct 19, 2026
a746e50
chore(examples): re-lock inline_intent to the aligned versions
nihalpasham Oct 19, 2026
ed558e1
Merge remote-tracking branch 'origin/main' into HEAD
nihalpasham Oct 19, 2026
a40d7f2
Merge origin/main into the FP16 m16n8k8 sm_75 floor branch a third time
nihalpasham Oct 19, 2026
9fe89ab
chore(intrinsics): regenerate the catalog outputs after the third mer…
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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::thread` to keep its public API stable.
Expand Down
16 changes: 8 additions & 8 deletions 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

// Included inside `cuda_device::wmma` to keep public paths stable.
Expand Down Expand Up @@ -43,7 +43,7 @@ pub unsafe fn mma_m16n8k128_s32_b1_xor_popc(c: [i32; 4], a: [u32; 2], b: u32) ->
unreachable!("mma_m16n8k128_s32_b1_xor_popc called outside CUDA kernel context")
}

/// Executes the mma_m16n8k16_f16_f16 Ampere floating-point MMA form.
/// Executes the mma_m16n8k16_f16_f16 dense floating-point MMA form.
/// Lowers to `mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16`. C, A, and B are this lane's fragments in PTX register order.
/// Requires PTX 7.0 and `sm_80+`.
///
Expand Down Expand Up @@ -2515,7 +2515,7 @@ pub unsafe fn mma_m16n8k32_s32_u8_satfinite(c: [i32; 4], a: [u32; 4], b: [u32; 2
unreachable!("mma_m16n8k32_s32_u8_satfinite called outside CUDA kernel context")
}

/// Executes the mma_m16n8k4_f32_tf32 Ampere floating-point MMA form.
/// Executes the mma_m16n8k4_f32_tf32 dense floating-point MMA form.
/// Lowers to `mma.sync.aligned.m16n8k4.row.col.f32.tf32.tf32.f32`. C, A, and B are this lane's fragments in PTX register order.
/// Requires PTX 7.0 and `sm_80+`.
///
Expand Down Expand Up @@ -2667,9 +2667,9 @@ pub unsafe fn mma_m16n8k64_s32_u4_satfinite(c: [i32; 4], a: [u32; 4], b: [u32; 2
unreachable!("mma_m16n8k64_s32_u4_satfinite called outside CUDA kernel context")
}

/// Executes the mma_m16n8k8_f16_f16 Ampere floating-point MMA form.
/// Executes the mma_m16n8k8_f16_f16 dense floating-point MMA form.
/// Lowers to `mma.sync.aligned.m16n8k8.row.col.f16.f16.f16.f16`. C, A, and B are this lane's fragments in PTX register order.
/// Requires PTX 7.0 and `sm_80+`.
/// Requires PTX 6.5 and `sm_75+`.
///
/// # Safety
/// All 32 lanes must execute the same instruction with the same qualifiers, and no lane may have exited.
Expand All @@ -2683,7 +2683,7 @@ pub unsafe fn mma_m16n8k8_f16_f16(c: [u32; 2], a: [u32; 2], b: u32) -> [u32; 2]
unreachable!("mma_m16n8k8_f16_f16 called outside CUDA kernel context")
}

/// Executes the mma_m16n8k8_f32_bf16 Ampere floating-point MMA form.
/// Executes the mma_m16n8k8_f32_bf16 dense floating-point MMA form.
/// Lowers to `mma.sync.aligned.m16n8k8.row.col.f32.bf16.bf16.f32`. C, A, and B are this lane's fragments in PTX register order.
/// Requires PTX 7.0 and `sm_80+`.
///
Expand All @@ -2699,9 +2699,9 @@ pub unsafe fn mma_m16n8k8_f32_bf16(c: [f32; 4], a: [u32; 2], b: u32) -> [f32; 4]
unreachable!("mma_m16n8k8_f32_bf16 called outside CUDA kernel context")
}

/// Executes the mma_m16n8k8_f32_f16 Ampere floating-point MMA form.
/// Executes the mma_m16n8k8_f32_f16 dense floating-point MMA form.
/// Lowers to `mma.sync.aligned.m16n8k8.row.col.f32.f16.f16.f32`. C, A, and B are this lane's fragments in PTX register order.
/// Requires PTX 7.0 and `sm_80+`.
/// Requires PTX 6.5 and `sm_75+`.
///
/// # Safety
/// All 32 lanes must execute the same instruction with the same qualifiers, and no lane may have exited.
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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// 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: 7e6b4b4b3b5958a0a36fb0c40486b8263e0799fbe1f23cb9be64b6687f27e3a0.
// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 3f56e2a1acd3d2b767a2a5c91c66a7f2e7276516404e8385da9d2dece37cf0d7.
// LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.

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

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