feat(intrinsics): admit bf16x2 to FP8x2 pack conversions

.github/workflows/examples-compile.yml

@@ -76,30 +76,31 @@ jobs:
       # The device pipeline shells out to LLVM's llc/opt for PTX
       # generation. This pin is the FLOOR GATE: if someone lands IR that
       # needs a newer llc, this job is where the floor bump becomes
-      # visible instead of on contributors' machines. LLVM 22 is the true
-      # floor: the ldmatrix/stmatrix family (llvm.nvvm.stmatrix.*, used by
-      # gemm_sol and gemm_sol_final) postdates llc-21, which "lowers" an
-      # unknown intrinsic by leaking its dotted LLVM name into the PTX as
-      # an `.extern .func` (exit 0, no diagnostic): unassemblable output
-      # that only surfaced once the ptxas verdict gate started assembling
-      # what this job produces. The pipeline now rejects such leaks with a
-      # named error, so an under-floor llc fails loudly here.
-      # If a future ubuntu-latest codename has no llvm-toolchain-*-22
+      # visible instead of on contributors' machines. LLVM 23 is the true
+      # floor: the bf16x2 to packed-FP8 conversions (PTX 9.1 forms, used
+      # by generated_intrinsics_blackwell) need `+ptx91`, which llc-22
+      # does not know. It ignored the unknown feature and emitted
+      # `.version 8.8` PTX that failed only at JIT; the pipeline now
+      # refuses a PTX 9.1 request below LLVM 23 with a named error, so an
+      # under-floor llc fails loudly here. The previous floor, LLVM 22 for
+      # the ldmatrix/stmatrix family (llc-21 leaked unknown intrinsics
+      # into the PTX as `.extern .func`), is covered by the same gate.
+      # If a future ubuntu-latest codename has no llvm-toolchain-*-23
       # suite, apt-get update fails loudly; recover by pinning runs-on to
       # the previous image or bumping the LLVM major here.
-      - name: Install LLVM 22 (llc/opt for the NVPTX backend)
+      - name: Install LLVM 23 (llc/opt for the NVPTX backend)
         run: |
           set -euo pipefail
           source /etc/os-release
           sudo install -d -m 0755 /etc/apt/keyrings
           wget -qO- https://apt.llvm.org/llvm-snapshot.gpg.key \
             | sudo tee /etc/apt/keyrings/apt-llvm-org.asc >/dev/null
-          echo "deb [signed-by=/etc/apt/keyrings/apt-llvm-org.asc] https://apt.llvm.org/${VERSION_CODENAME}/ llvm-toolchain-${VERSION_CODENAME}-22 main" \
-            | sudo tee /etc/apt/sources.list.d/llvm-22.list
+          echo "deb [signed-by=/etc/apt/keyrings/apt-llvm-org.asc] https://apt.llvm.org/${VERSION_CODENAME}/ llvm-toolchain-${VERSION_CODENAME}-23 main" \
+            | sudo tee /etc/apt/sources.list.d/llvm-23.list
           sudo apt-get update
-          sudo apt-get install -y --no-install-recommends llvm-22
-          llc-22 --version | head -2
-          opt-22 --version | head -2
+          sudo apt-get install -y --no-install-recommends llvm-23
+          llc-23 --version | head -2
+          opt-23 --version | head -2
 
       - name: Compile every example (smoketest --compile-only)
         # Step-level timeout, below the job's 75: a job-level timeout
@@ -108,7 +109,7 @@ jobs:
         timeout-minutes: 50
         env:
           CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }}
-          CUDA_OXIDE_LLC: /usr/bin/llc-22
+          CUDA_OXIDE_LLC: /usr/bin/llc-23
           # Each example is its own cargo workspace; without a shared
           # target dir the dep tree rebuilds ~85 times and blows the
           # runner's ~14 GB disk (same trick as clippy.yml). smoketest.sh
@@ -136,7 +137,7 @@ jobs:
         working-directory: cuda-oxide
         env:
           CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }}
-          CUDA_OXIDE_LLC: /usr/bin/llc-22
+          CUDA_OXIDE_LLC: /usr/bin/llc-23
           CARGO_TARGET_DIR: ${{ runner.temp }}/examples-shared-target
           CARGO_INCREMENTAL: "0"
         run: |

cuda-oxide/crates/cuda-device/src/generated/async_copy.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::async_copy` to keep existing paths stable.

cuda-oxide/crates/cuda-device/src/generated/atomic.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::atomic` to keep existing paths stable.

cuda-oxide/crates/cuda-device/src/generated/bf16.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::bf16` to keep existing paths stable.

cuda-oxide/crates/cuda-device/src/generated/bf16x2.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::bf16x2` to keep existing paths stable.

cuda-oxide/crates/cuda-device/src/generated/clc.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::clc` to keep its public API stable.

cuda-oxide/crates/cuda-device/src/generated/cluster_barrier.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::cluster` to keep public paths stable.

cuda-oxide/crates/cuda-device/src/generated/cluster_memory.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::cluster` to keep its public API stable.

cuda-oxide/crates/cuda-device/src/generated/cluster_sreg.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // This file is included inside `cuda_device::cluster`.

cuda-oxide/crates/cuda-device/src/generated/convert.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::convert`.
@@ -81,6 +81,15 @@ pub fn cvt_rn_relu_f16x2_f32(lo: f32, hi: f32) -> u32 {
     )
 }
 
+/// Converts packed bf16x2 to packed e4m3x2 with nearest-even finite saturation, preserving half order.
+#[inline(never)]
+pub fn cvt_rn_satfinite_e4m3x2_bf16x2(packed: u32) -> u16 {
+    let _ = packed;
+    unreachable!(
+        "generated CUDA intrinsic `cuda_device::convert::cvt_rn_satfinite_e4m3x2_bf16x2` executed outside device compilation"
+    )
+}
+
 /// Converts packed f16x2 to packed e4m3x2 with nearest-even finite saturation, preserving half order.
 #[inline(never)]
 pub fn cvt_rn_satfinite_e4m3x2_f16x2(packed: u32) -> u16 {
@@ -99,6 +108,15 @@ pub fn cvt_rn_satfinite_e4m3x2_f32(lo: f32, hi: f32) -> u16 {
     )
 }
 
+/// Converts packed bf16x2 to packed e5m2x2 with nearest-even finite saturation, preserving half order.
+#[inline(never)]
+pub fn cvt_rn_satfinite_e5m2x2_bf16x2(packed: u32) -> u16 {
+    let _ = packed;
+    unreachable!(
+        "generated CUDA intrinsic `cuda_device::convert::cvt_rn_satfinite_e5m2x2_bf16x2` executed outside device compilation"
+    )
+}
+
 /// Converts packed f16x2 to packed e5m2x2 with nearest-even finite saturation, preserving half order.
 #[inline(never)]
 pub fn cvt_rn_satfinite_e5m2x2_f16x2(packed: u32) -> u16 {
@@ -117,6 +135,15 @@ pub fn cvt_rn_satfinite_e5m2x2_f32(lo: f32, hi: f32) -> u16 {
     )
 }
 
+/// Converts packed bf16x2 to packed e4m3x2 with nearest-even finite saturation and ReLU, preserving half order.
+#[inline(never)]
+pub fn cvt_rn_satfinite_relu_e4m3x2_bf16x2(packed: u32) -> u16 {
+    let _ = packed;
+    unreachable!(
+        "generated CUDA intrinsic `cuda_device::convert::cvt_rn_satfinite_relu_e4m3x2_bf16x2` executed outside device compilation"
+    )
+}
+
 /// Converts packed f16x2 to packed e4m3x2 with nearest-even finite saturation and ReLU, preserving half order.
 #[inline(never)]
 pub fn cvt_rn_satfinite_relu_e4m3x2_f16x2(packed: u32) -> u16 {
@@ -135,6 +162,15 @@ pub fn cvt_rn_satfinite_relu_e4m3x2_f32(lo: f32, hi: f32) -> u16 {
     )
 }
 
+/// Converts packed bf16x2 to packed e5m2x2 with nearest-even finite saturation and ReLU, preserving half order.
+#[inline(never)]
+pub fn cvt_rn_satfinite_relu_e5m2x2_bf16x2(packed: u32) -> u16 {
+    let _ = packed;
+    unreachable!(
+        "generated CUDA intrinsic `cuda_device::convert::cvt_rn_satfinite_relu_e5m2x2_bf16x2` executed outside device compilation"
+    )
+}
+
 /// Converts packed f16x2 to packed e5m2x2 with nearest-even finite saturation and ReLU, preserving half order.
 #[inline(never)]
 pub fn cvt_rn_satfinite_relu_e5m2x2_f16x2(packed: u32) -> u16 {

cuda-oxide/crates/cuda-device/src/generated/counted_barrier.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::barrier` to keep its public API stable.

cuda-oxide/crates/cuda-device/src/generated/debug_control.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::debug` to keep its public API stable.

cuda-oxide/crates/cuda-device/src/generated/debug_sreg.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::debug` to keep existing paths stable.

cuda-oxide/crates/cuda-device/src/generated/dotprod.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::dotprod` to keep existing paths stable.

cuda-oxide/crates/cuda-device/src/generated/f16.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::f16` to keep existing paths stable.

cuda-oxide/crates/cuda-device/src/generated/f16x2.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::f16x2` to keep existing paths stable.

cuda-oxide/crates/cuda-device/src/generated/f32x2.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::f32x2` to keep existing paths stable.

cuda-oxide/crates/cuda-device/src/generated/fence.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // This file is included inside `cuda_device::fence` so existing paths stay stable.

cuda-oxide/crates/cuda-device/src/generated/float.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::float`.

cuda-oxide/crates/cuda-device/src/generated/grid_dependency.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::grid` to keep its public API stable.

cuda-oxide/crates/cuda-device/src/generated/grid_sreg.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::grid` to keep existing paths stable.

cuda-oxide/crates/cuda-device/src/generated/i16x2.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::i16x2` to keep existing paths stable.

cuda-oxide/crates/cuda-device/src/generated/int.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::int` to keep existing paths stable.

cuda-oxide/crates/cuda-device/src/generated/ldmatrix.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::wmma` to keep public paths stable.

cuda-oxide/crates/cuda-device/src/generated/mbarrier_basic.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::barrier` to keep existing paths stable.

cuda-oxide/crates/cuda-device/src/generated/mbarrier_extended.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::barrier` to keep existing paths stable.

cuda-oxide/crates/cuda-device/src/generated/movmatrix.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::wmma` to keep its public API stable.

cuda-oxide/crates/cuda-device/src/generated/prmt.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::prmt` to keep the public paths stable.

cuda-oxide/crates/cuda-device/src/generated/register_control.rs

@@ -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: c243e568fc16c9f907c9fd87fbacbe9f823043a9badd8633e43e9faa87adbd10.
+// Catalog schema/version: 46/0.1.0; intrinsic ABI: v1; catalog SHA-256: 31b5238818203ac4c636984d416ea7791c5861feb0ad30d1f5fc693ebefb4399.
 // LLVM source: aab11ee0ea5ea9a1d462ad0bc1fae9d9bb643019.
 
 // Included inside `cuda_device::thread` to keep its public API stable.