chore: move bloomery to NVIDIA/cuda-rust 9df9e7ce and speed up the CuTe verifier

.github/workflows/book.yml

@@ -35,6 +35,8 @@ permissions:
 # Scope this to the ref, and keep the single-deploy guarantee on the deploy job
 # itself. Superseded PR runs are cancelled; pushes to main are not, so a
 # publish in flight is never interrupted.
+# The group deliberately has no commit SHA. Parallel builds on main would let
+# a slower older commit deploy after a newer one and replace the published site.
 concurrency:
   group: ${{ github.workflow }}-${{ github.ref }}
   cancel-in-progress: ${{ github.event_name == 'pull_request' }}

.github/workflows/ci-status.yml

@@ -0,0 +1,71 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+---
+name: CI status
+
+# One status for the whole repository, so the README needs one badge.
+#
+# Badges are per workflow file, and the two pipelines have to stay separate
+# for now: ci.yml (SIMT and the shared host crates) runs on GitHub-hosted
+# runners on push and pull_request; cutile-rs.yml (Tile) runs on NVIDIA
+# self-hosted runners, which cannot run on pull_request events, so it
+# triggers on push to main and to the copy-pr-bot mirror branches.
+#
+# This workflow runs whenever either of them completes on main. For each
+# pipeline it finds the newest completed push run on main that actually
+# executed jobs and fails if that run failed. Two kinds of run say nothing
+# about the pipeline and are skipped over: a run whose jobs were all skipped
+# (cutile-rs.yml's `changes` gate does that on a non-Tile push) and a
+# cancelled run. Taking either at face value would turn the badge green while
+# the last real run was red.
+#
+# The badge in README.md points here. To cover another pipeline, add its
+# workflow name to the trigger list and its file name to the loop below.
+on:
+  workflow_run:
+    workflows: [CI, cutile-rs-pr]
+    types: [completed]
+    branches: [main]
+
+permissions:
+  actions: read
+  contents: read
+
+jobs:
+  status:
+    name: combined CI status on main
+    runs-on: ubuntu-latest
+    timeout-minutes: 5
+    steps:
+      - name: Report the newest executed main run of each pipeline
+        env:
+          GH_TOKEN: ${{ github.token }}
+        run: |
+          set -euo pipefail
+          failed=0
+          for workflow in ci.yml cutile-rs.yml; do
+            verdict=none
+            url=-
+            while read -r id conclusion run_url; do
+              [ -n "${id}" ] || continue
+              case "${conclusion}" in
+                success|failure|timed_out) ;;
+                *) continue ;;
+              esac
+              executed="$(gh api "repos/${GITHUB_REPOSITORY}/actions/runs/${id}/jobs?per_page=100" \
+                --jq '[.jobs[] | select(.name != "changes" and .conclusion != "skipped")] | length')"
+              if [ "${executed}" -gt 0 ]; then
+                verdict="${conclusion}"
+                url="${run_url}"
+                break
+              fi
+            done < <(gh api \
+              "repos/${GITHUB_REPOSITORY}/actions/workflows/${workflow}/runs?branch=main&event=push&status=completed&per_page=10" \
+              --jq '.workflow_runs[] | "\(.id) \(.conclusion) \(.html_url)"')
+            echo "${workflow}: ${verdict} (${url})"
+            case "${verdict}" in
+              success|none) ;;
+              *) failed=1 ;;
+            esac
+          done
+          exit "${failed}"

.github/workflows/ci.yml

@@ -28,9 +28,11 @@ on:
       - 'Cargo.lock'
       - '.github/workflows/**'
 
+# Each commit on main gets its own group, so it starts at once. Any other ref
+# still shares one group and cancels the older run.
 concurrency:
-  group: ci-gate-${{ github.ref_name }}
-  cancel-in-progress: true
+  group: ci-gate-${{ github.ref_name }}${{ github.ref_name == 'main' && format('-{0}', github.sha) || '' }}
+  cancel-in-progress: ${{ github.ref_name != 'main' }}
 
 # Every other workflow in the repo pins this rather than inheriting whatever the
 # repository default happens to be. This one is a pure caller today, so the four

.github/workflows/clippy.yml

@@ -19,10 +19,10 @@ jobs:
         uses: actions/checkout@v4
 
       - name: Install CUDA toolkit
-        uses: Jimver/cuda-toolkit@v0.2.35
+        uses: Jimver/cuda-toolkit@0ca15bb027fc96f7f61ddbb9c5221d699da06bdb  # v0.2.36
         id: cuda-toolkit
         with:
-          cuda: "13.1.1"
+          cuda: "13.3.1"
           method: network
           sub-packages: '["nvcc", "cudart-dev"]'
           non-cuda-sub-packages: '["libcurand-dev"]'
@@ -37,7 +37,7 @@ jobs:
         working-directory: cuda-oxide
         run: cargo clippy --workspace --all-targets -- -D warnings
 
-      # `rustc-codegen-cuda` and each example under examples/ are their own
+      # `rustc-codegen-cuda`, the standard and CuTe examples, and CuTe models have
       # standalone workspaces (the codegen backend is swapped in via
       # RUSTFLAGS, so they can't live in the root [workspace]). That's 60+
       # cargo invocations from different workspace roots. By default each
@@ -49,7 +49,7 @@ jobs:
       # Point all of those invocations at one shared CARGO_TARGET_DIR under
       # ${RUNNER_TEMP}. Cargo fingerprints each unit by package + features
       # + workspace_root + toolchain, so co-located workspaces using the
-      # same pinned nightly deduplicate on identical deps and compile each
+      # same pinned toolchain deduplicate on identical deps and compile each
       # one once. This is intentionally scoped *not* to cover the
       # `cargo clippy --workspace` step above, which already has a
       # well-defined `target/` under cuda-oxide.
@@ -60,18 +60,26 @@ jobs:
         working-directory: cuda-oxide/crates/rustc-codegen-cuda
         run: cargo clippy --all-targets -- -D warnings
 
-      - name: Run clippy (rustc-codegen-cuda examples)
+      - name: Run clippy (standard and CuTe examples and models)
         env:
           CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }}
           CARGO_TARGET_DIR: ${{ runner.temp }}/clippy-shared-target
+          RUSTC_BOOTSTRAP: "1"
         run: |
           set -e
-          for ex in cuda-oxide/crates/rustc-codegen-cuda/examples/*/; do
-            if [ ! -f "$ex/Cargo.toml" ]; then continue; fi
-            name=$(basename "$ex")
-            echo "::group::clippy: $name"
-            (cd "$ex" && cargo clippy --all-targets -- -D warnings)
-            echo "::endgroup::"
+          example_roots=(
+            cuda-oxide/crates/rustc-codegen-cuda/examples
+            cuda-oxide/cute-rs-preview/examples
+            cuda-oxide/cute-rs-preview/models
+          )
+          for example_root in "${example_roots[@]}"; do
+            for ex in "$example_root"/*/; do
+              if [ ! -f "$ex/Cargo.toml" ]; then continue; fi
+              name=$(basename "$ex")
+              echo "::group::clippy: $name"
+              (cd "$ex" && cargo clippy --all-targets -- -D warnings)
+              echo "::endgroup::"
+            done
           done
 
           # The loop above is one level deep. A few examples nest a second
@@ -97,11 +105,11 @@ jobs:
                grep -q "\"${sub}\""; then
               continue
             fi
-            echo "::group::clippy: ${dir#cuda-oxide/crates/rustc-codegen-cuda/examples/}"
+            echo "::group::clippy: ${dir#cuda-oxide/}"
             (cd "${dir}" && cargo clippy --locked --all-targets -- -D warnings)
             echo "::endgroup::"
           done < <(
-            find cuda-oxide/crates/rustc-codegen-cuda/examples -mindepth 3 -name Cargo.toml \
+            find "${example_roots[@]}" -mindepth 3 -name Cargo.toml \
               -not -path '*/target/*' | sort
           )
 

.github/workflows/codeql.yml

@@ -1,17 +1,36 @@
 ---
 name: codeql
 
+# Push and pull_request run only when Rust, workflow, or script inputs change.
+# A Markdown-only commit skips both analyzers. The weekly schedule has no path
+# filter, so that scan still covers the whole tree.
 on:
   push:
     branches: [main]
+    paths:
+      - '**/*.rs'
+      - '**/Cargo.toml'
+      - '**/Cargo.lock'
+      - '.github/**'
+      - '**/*.sh'
+      - '**/*.py'
   pull_request:
     branches: [main]
+    paths:
+      - '**/*.rs'
+      - '**/Cargo.toml'
+      - '**/Cargo.lock'
+      - '.github/**'
+      - '**/*.sh'
+      - '**/*.py'
   schedule:
     - cron: "44 3 * * 6"
 
+# Each commit on main gets its own group, so it starts at once. Any other ref
+# still shares one group and cancels the older run.
 concurrency:
-  group: ${{ github.workflow }}-${{ github.ref_name }}
-  cancel-in-progress: true
+  group: ${{ github.workflow }}-${{ github.ref_name }}${{ github.ref_name == 'main' && format('-{0}', github.sha) || '' }}
+  cancel-in-progress: ${{ github.ref_name != 'main' }}
 
 jobs:
   analyze:

.github/workflows/cutile-rs.yml

@@ -25,9 +25,11 @@ on:
       - "pull-request/[0-9]+"
   workflow_dispatch:
 
+# Each commit on main gets its own group, so it starts at once. Any other ref
+# still shares one group and cancels the older run.
 concurrency:
-  group: ${{ github.workflow }}-${{ github.ref_name }}
-  cancel-in-progress: true
+  group: ${{ github.workflow }}-${{ github.ref_name }}${{ github.ref_name == 'main' && format('-{0}', github.sha) || '' }}
+  cancel-in-progress: ${{ github.ref_name != 'main' }}
 
 permissions:
   contents: read
@@ -39,8 +41,11 @@ defaults:
 
 jobs:
   # Tile lane gate: cutile-rs tree plus shared host crates. Skips when only
-  # cuda-oxide changes. When copy-pr-bot yields no detectable file diff, run
-  # anyway so the mirror push is never silently skipped.
+  # cuda-oxide (or any other non-tile path) changes. `any_changed` is a real
+  # filter: dorny/paths-filter sets an output only for filters declared here,
+  # and an undeclared name stays empty, which would make `!= 'true'` always
+  # succeed. When copy-pr-bot yields no detectable file diff, both outputs are
+  # false and the jobs run anyway so the mirror push is never silently skipped.
   changes:
     runs-on: ubuntu-latest
     outputs:
@@ -50,7 +55,7 @@ jobs:
       - uses: actions/checkout@v7
         with:
           fetch-depth: 0
-      - uses: dorny/paths-filter@v3
+      - uses: dorny/paths-filter@0e4a8c6effa4802afeda77dc8d303f8176d7dfad  # v3.0.4
         id: changes
         with:
           filters: |
@@ -65,6 +70,10 @@ jobs:
               - 'rust-toolchain.toml'
               - '.github/workflows/cutile-rs.yml'
               - '.github/copy-pr-bot.yaml'
+            # `*` covers a root file such as AGENTS.md. `**` covers every nested path.
+            any_changed:
+              - '*'
+              - '**'
 
   # Exercise the oldest Tile IR assembler, including its lack of
   # --list-versions. The separate 13.0 lane only checks shared host crates.
@@ -199,7 +208,7 @@ jobs:
 
       # Prebuilt binary instead of `cargo install` (which compiled the tool
       # from source on every run).
-      - uses: taiki-e/install-action@9114bf4d891761788c546334fd37538eae1bf8b3 # v2.87.16
+      - uses: taiki-e/install-action@9114bf4d891761788c546334fd37538eae1bf8b3  # v2.87.16
         with:
           tool: cargo-semver-checks
 

.github/workflows/docs.yml

@@ -26,9 +26,11 @@ on:
       - 'Cargo.lock'
       - '.github/workflows/**'
 
+# Each commit on main gets its own group, so it starts at once. Any other ref
+# still shares one group and cancels the older run.
 concurrency:
-  group: ${{ github.workflow }}-${{ github.ref_name }}
-  cancel-in-progress: true
+  group: ${{ github.workflow }}-${{ github.ref_name }}${{ github.ref_name == 'main' && format('-{0}', github.sha) || '' }}
+  cancel-in-progress: ${{ github.ref_name != 'main' }}
 
 permissions:
   contents: read
@@ -48,10 +50,10 @@ jobs:
       - name: Install CUDA toolkit
         # `cargo doc --workspace` compiles cuda-bindings' build-time bindgen,
         # which needs `cuda.h`. Same toolkit pin as the unit-tests job.
-        uses: Jimver/cuda-toolkit@v0.2.35
+        uses: Jimver/cuda-toolkit@0ca15bb027fc96f7f61ddbb9c5221d699da06bdb  # v0.2.36
         id: cuda-toolkit
         with:
-          cuda: "13.1.1"
+          cuda: "13.3.1"
           method: network
           sub-packages: '["nvcc", "cudart-dev"]'
           non-cuda-sub-packages: '["libcurand-dev"]'

.github/workflows/examples-compile.yml

@@ -26,9 +26,11 @@ on:
       - 'Cargo.lock'
       - '.github/workflows/**'
 
+# Each commit on main gets its own group, so it starts at once. Any other ref
+# still shares one group and cancels the older run.
 concurrency:
-  group: ${{ github.workflow }}-${{ github.ref_name }}
-  cancel-in-progress: true
+  group: ${{ github.workflow }}-${{ github.ref_name }}${{ github.ref_name == 'main' && format('-{0}', github.sha) || '' }}
+  cancel-in-progress: ${{ github.ref_name != 'main' }}
 
 permissions:
   contents: read
@@ -59,10 +61,10 @@ jobs:
       # Needed for cuda.h and curand.h: bindgen runs in the shared
       # cuda-bindings build script. Same setup as unit-tests.yml / clippy.yml.
       - name: Install CUDA toolkit
-        uses: Jimver/cuda-toolkit@v0.2.35
+        uses: Jimver/cuda-toolkit@0ca15bb027fc96f7f61ddbb9c5221d699da06bdb  # v0.2.36
         id: cuda-toolkit
         with:
-          cuda: "13.1.1"
+          cuda: "13.3.1"
           method: network
           sub-packages: '["nvcc", "cudart-dev"]'
           # nvJitLink is packaged separately from cuda-nvcc on Linux. Keep it
@@ -74,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
@@ -106,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
@@ -153,7 +156,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: |