diff --git a/.github/actions/setup-cuda-container/action.yml b/.github/actions/setup-cuda-container/action.yml
new file mode 100644
index 0000000000..e01e26ab80
--- /dev/null
+++ b/.github/actions/setup-cuda-container/action.yml
@@ -0,0 +1,54 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+---
+name: Set up the CUDA devel container
+description: >-
+  Installs the host packages the nvidia/cuda devel image lacks, marks the
+  workspace safe for git, and installs rustup. Callers use the $/ form so
+  this runs before checkout and git is already installed. The image already
+  contains the CUDA toolkit. Later steps pick up rust-toolchain.toml from
+  their working directory.
+
+runs:
+  using: composite
+  steps:
+    - name: Install host packages
+      shell: bash
+      run: |
+        apt-get update -qq
+        # python3 is for sync-example-locks.sh. The image has neither git nor
+        # python3; git has to be present before actions/checkout.
+        apt-get install -y -qq --no-install-recommends \
+          build-essential ca-certificates curl git pkg-config libclang-dev \
+          python3 wget
+
+    - name: Mark the workspace as a safe git directory
+      shell: bash
+      run: git config --global --add safe.directory "$GITHUB_WORKSPACE"
+
+    - name: Install Rust
+      shell: bash
+      env:
+        RUSTUP_VERSION: "1.29.2"
+        # SHA-256 of rustup-init for x86_64-unknown-linux-gnu, published next to
+        # the binary as .../rustup/archive/1.29.2/x86_64-unknown-linux-gnu/rustup-init.sha256.
+        # Bump both values together.
+        RUSTUP_INIT_SHA256: "5be191e8c70a5007074254cb74d78dbfdd25cd5a61c27de7b8e5410122ec24f4"
+      run: |
+        set -euo pipefail
+        [ "$(uname -m)" = "x86_64" ] || { echo "the rustup-init pin is for x86_64"; exit 1; }
+        init="${RUNNER_TEMP}/rustup-init"
+        curl --proto '=https' --tlsv1.2 -sSfo "${init}" \
+          "https://static.rust-lang.org/rustup/archive/${RUSTUP_VERSION}/x86_64-unknown-linux-gnu/rustup-init"
+        echo "${RUSTUP_INIT_SHA256}  ${init}" | sha256sum -c -
+        chmod +x "${init}"
+        "${init}" -y --no-modify-path --profile minimal --default-toolchain none
+        rm -f "${init}"
+        echo "${HOME}/.cargo/bin" >> "${GITHUB_PATH}"
+        # $/ runs this before the workspace checkout. The action download is
+        # the same commit, and rustup finds rust-toolchain.toml by walking up
+        # from the working directory. The root file is the stable pin;
+        # cuda-oxide/ selects the nightly.
+        repo="${GITHUB_ACTION_PATH}/../../.."
+        (cd "${repo}" && "${HOME}/.cargo/bin/rustup" show)
+        (cd "${repo}/cuda-oxide" && "${HOME}/.cargo/bin/rustup" show)
diff --git a/.github/workflows/book.yml b/.github/workflows/book.yml
index 732ff8b667..bf698fcf89 100644
--- a/.github/workflows/book.yml
+++ b/.github/workflows/book.yml
@@ -1,27 +1,15 @@
 name: Deploy CUDA Rust Docs
 
 # Path-filtered lane: only documentation sources can change what this builds.
+# Self-hosted runners reject pull_request. copy-pr-bot mirrors trusted PRs to
+# pull-request/<n>, and the build job triggers on that push. `deploy` stays on
+# main: a mirror push is a push event, so an event-name check would publish it.
 on:
   push:
-    branches: [main]
-    paths:
-      - 'cuda-oxide/cuda-oxide-book/**'
-      - 'cutile-rs/cutile-book/**'
-      - 'cutile-rs/scripts/build_versioned_book.sh'
-      - 'docs/**'
-      - '.github/workflows/book.yml'
-  # Validate the book on PRs too. The build job alone runs there; `deploy` is
-  # gated below, so a pull request never publishes. Without this the book was
-  # only ever built after merge, so a PR that broke it could not be caught by
-  # review or by CI.
-  pull_request:
-    branches: [main]
-    paths:
-      - 'cuda-oxide/cuda-oxide-book/**'
-      - 'cutile-rs/cutile-book/**'
-      - 'cutile-rs/scripts/build_versioned_book.sh'
-      - 'docs/**'
-      - '.github/workflows/book.yml'
+    # No workflow-level paths: filter; the changes job below gates the jobs.
+    branches:
+      - main
+      - "pull-request/[0-9]+"
   workflow_dispatch:
 
 # Only checkout rights by default. The Pages deployment permissions live on
@@ -39,11 +27,47 @@ permissions:
 # 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' }}
+  cancel-in-progress: ${{ github.ref != 'refs/heads/main' }}
 
 jobs:
+  # Path gate. copy-pr-bot creates pull-request/<n> at an existing SHA, and
+  # GitHub often treats that push as having no path changes, so a workflow-level
+  # `paths:` filter would skip the whole run. This job diffs against main
+  # instead, as cutile-rs.yml does, and every job below runs when no file diff
+  # is detected. Other events (schedule, workflow_dispatch) skip the gate.
+  changes:
+    if: github.event_name == 'push'
+    runs-on: linux-amd64-cpu4
+    timeout-minutes: 5
+    permissions:
+      contents: read
+    outputs:
+      relevant: ${{ steps.changes.outputs.relevant }}
+      any_changed: ${{ steps.changes.outputs.any_changed }}
+    steps:
+      - uses: actions/checkout@v4
+        with:
+          fetch-depth: 0
+      - uses: dorny/paths-filter@0e4a8c6effa4802afeda77dc8d303f8176d7dfad  # v3.0.4
+        id: changes
+        with:
+          filters: |
+            relevant:
+              - 'cuda-oxide/cuda-oxide-book/**'
+              - 'cutile-rs/cutile-book/**'
+              - 'cutile-rs/scripts/build_versioned_book.sh'
+              - 'docs/**'
+              - '.github/workflows/book.yml'
+            # `*` covers a root file, `**` every nested path. Both false means
+            # no file diff was detected.
+            any_changed:
+              - '*'
+              - '**'
+
   build:
-    runs-on: ubuntu-latest
+    needs: changes
+    if: ${{ !cancelled() && (github.event_name != 'push' || needs.changes.outputs.relevant == 'true' || needs.changes.outputs.any_changed != 'true') }}
+    runs-on: linux-amd64-cpu4
     # An umbrella page plus two Sphinx books. Bound the job so a
     # wedged install or a sphinx that never returns fails instead of burning to
     # the 6-hour default.
@@ -76,9 +100,9 @@ jobs:
         run: |
           # Unset CUTILE_DOCS_TAGS → build every tag matching the pattern.
           # Empty CUTILE_DOCS_TAGS → current tree only (fast PR check).
-          if [ "${{ github.event_name }}" = "pull_request" ]; then
-            export CUTILE_DOCS_TAGS=""
-          fi
+          case "${GITHUB_REF}" in
+            refs/heads/pull-request/*) export CUTILE_DOCS_TAGS="" ;;
+          esac
           bash docs/build_all_docs.sh
 
       # Only `deploy` consumes this, and `deploy` is skipped on pull requests,
@@ -86,7 +110,7 @@ jobs:
       # job is to prove both Sphinx builds succeed, which the steps above already
       # did; gate the upload on the same condition `deploy` uses so the two
       # cannot drift apart.
-      - if: github.event_name != 'pull_request'
+      - if: github.ref == 'refs/heads/main'
         uses: actions/upload-pages-artifact@v3
         with:
           path: docs/build/html
@@ -96,7 +120,7 @@ jobs:
     # Publishing is for main only. A pull request builds the book to prove it
     # is sound and stops there; it never touches the Pages environment or
     # holds the deployment permissions below.
-    if: github.event_name != 'pull_request'
+    if: github.ref == 'refs/heads/main'
     permissions:
       pages: write
       id-token: write
diff --git a/.github/workflows/cargo-deny.yml b/.github/workflows/cargo-deny.yml
index 73f426cc77..85803cb758 100644
--- a/.github/workflows/cargo-deny.yml
+++ b/.github/workflows/cargo-deny.yml
@@ -1,8 +1,8 @@
 ---
 name: cargo-deny
 
-# Runs as a job of the top-level CI workflow (.github/workflows/ci.yml) on
-# push/pull_request, plus a standalone weekly security scan via schedule.
+# Runs as a job of the top-level CI workflow (.github/workflows/ci.yml),
+# plus a standalone weekly security scan via schedule.
 # Concurrency for the CI path is governed by CI's own group.
 #
 # Four parts of the same concern. `check` enforces the license *policy* over
@@ -30,14 +30,24 @@ on:
 permissions:
   contents: read
 
+# A container job's default shell is dash; use bash, as on ubuntu-latest.
+defaults:
+  run:
+    shell: bash
+
 jobs:
   check:
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
+    container:
+      image: rust:1.98.0-bookworm
     timeout-minutes: 10
     steps:
       - name: Clone repository
         uses: actions/checkout@v4
 
+      - name: Mark the workspace as a safe git directory
+        run: git config --global --add safe.directory "$GITHUB_WORKSPACE"
+
       - name: Install cargo-deny
         run: cargo install cargo-deny --locked
 
@@ -94,12 +104,17 @@ jobs:
 
   license-manifest:
     name: dependency-licenses.csv covers every declared crate
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
+    container:
+      image: rust:1.98.0-bookworm
     timeout-minutes: 10
     steps:
       - name: Clone repository
         uses: actions/checkout@v4
 
+      - name: Mark the workspace as a safe git directory
+        run: git config --global --add safe.directory "$GITHUB_WORKSPACE"
+
       # Only reads manifests: `cargo metadata` resolves the graph without
       # running build scripts, so this needs no CUDA toolkit and no LLVM.
       - name: Verify the license inventory covers the workspace
@@ -107,12 +122,17 @@ jobs:
 
   example-policy:
     name: deny.toml governs the example workspaces
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
+    container:
+      image: rust:1.98.0-bookworm
     timeout-minutes: 15
     steps:
       - name: Clone repository
         uses: actions/checkout@v4
 
+      - name: Mark the workspace as a safe git directory
+        run: git config --global --add safe.directory "$GITHUB_WORKSPACE"
+
       - name: Install cargo-deny
         run: cargo install cargo-deny --locked
 
@@ -125,12 +145,17 @@ jobs:
 
   source-headers:
     name: every source file carries the SPDX header
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
+    container:
+      image: rust:1.98.0-bookworm
     timeout-minutes: 5
     steps:
       - name: Clone repository
         uses: actions/checkout@v4
 
+      - name: Mark the workspace as a safe git directory
+        run: git config --global --add safe.directory "$GITHUB_WORKSPACE"
+
       # Pure text check over tracked files: no toolchain, no network. The OSRB
       # review fixed 32 missing headers in #812 and the next two scripts to
       # land (#819, #835) drifted again; this makes the third time the last.
diff --git a/.github/workflows/ci-status.yml b/.github/workflows/ci-status.yml
index 2c9f44b127..1b7a72791d 100644
--- a/.github/workflows/ci-status.yml
+++ b/.github/workflows/ci-status.yml
@@ -5,11 +5,9 @@ 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.
+# Badges are per workflow file. ci.yml and cutile-rs.yml both run on NVIDIA
+# self-hosted runners, which reject pull_request, so both trigger 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
@@ -34,7 +32,7 @@ permissions:
 jobs:
   status:
     name: combined CI status on main
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 5
     steps:
       - name: Report the newest executed main run of each pipeline
diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml
index bd0c764954..8026229ee0 100644
--- a/.github/workflows/ci.yml
+++ b/.github/workflows/ci.yml
@@ -4,29 +4,16 @@ name: CI
 # Path-filtered SIMT lane: skip when a change cannot affect cuda-oxide or the
 # shared host crates. `cuda-oxide/**` is deliberately broad (prose as well as
 # code) because guard scripts cross-check the book against the crates.
+#
+# Self-hosted runners reject pull_request. copy-pr-bot mirrors trusted PRs to
+# pull-request/<n>, and this workflow triggers on that push
+# (config: .github/copy-pr-bot.yaml).
 on:
   push:
-    branches: [main]
-    paths:
-      - 'cuda-oxide/**'
-      - 'cuda-bindings/**'
-      - 'cuda-core/**'
-      - 'cuda-core-derive/**'
-      - 'cuda-async/**'
-      - 'Cargo.toml'
-      - 'Cargo.lock'
-      - '.github/workflows/**'
-  pull_request:
-    branches: [main]
-    paths:
-      - 'cuda-oxide/**'
-      - 'cuda-bindings/**'
-      - 'cuda-core/**'
-      - 'cuda-core-derive/**'
-      - 'cuda-async/**'
-      - 'Cargo.toml'
-      - 'Cargo.lock'
-      - '.github/workflows/**'
+    # No workflow-level paths: filter; the changes job below gates the jobs.
+    branches:
+      - main
+      - "pull-request/[0-9]+"
 
 # 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.
@@ -44,14 +31,60 @@ permissions:
   contents: read
 
 jobs:
+  # Path gate. copy-pr-bot creates pull-request/<n> at an existing SHA, and
+  # GitHub often treats that push as having no path changes, so a workflow-level
+  # `paths:` filter would skip the whole run. This job diffs against main
+  # instead, as cutile-rs.yml does, and every job below runs when no file diff
+  # is detected. Other events (schedule, workflow_dispatch) skip the gate.
+  changes:
+    if: github.event_name == 'push'
+    runs-on: linux-amd64-cpu4
+    timeout-minutes: 5
+    permissions:
+      contents: read
+    outputs:
+      relevant: ${{ steps.changes.outputs.relevant }}
+      any_changed: ${{ steps.changes.outputs.any_changed }}
+    steps:
+      - uses: actions/checkout@v4
+        with:
+          fetch-depth: 0
+      - uses: dorny/paths-filter@0e4a8c6effa4802afeda77dc8d303f8176d7dfad  # v3.0.4
+        id: changes
+        with:
+          filters: |
+            relevant:
+              - 'cuda-oxide/**'
+              - 'cuda-bindings/**'
+              - 'cuda-core/**'
+              - 'cuda-core-derive/**'
+              - 'cuda-async/**'
+              - 'Cargo.toml'
+              - 'Cargo.lock'
+              - '.github/workflows/**'
+              - '.github/actions/setup-cuda-container/**'
+            # `*` covers a root file, `**` every nested path. Both false means
+            # no file diff was detected.
+            any_changed:
+              - '*'
+              - '**'
+
   fmt:
+    needs: changes
+    if: ${{ !cancelled() && (github.event_name != 'push' || needs.changes.outputs.relevant == 'true' || needs.changes.outputs.any_changed != 'true') }}
     uses: ./.github/workflows/fmt.yml
 
   clippy:
+    needs: changes
+    if: ${{ !cancelled() && (github.event_name != 'push' || needs.changes.outputs.relevant == 'true' || needs.changes.outputs.any_changed != 'true') }}
     uses: ./.github/workflows/clippy.yml
 
   unit-tests:
+    needs: changes
+    if: ${{ !cancelled() && (github.event_name != 'push' || needs.changes.outputs.relevant == 'true' || needs.changes.outputs.any_changed != 'true') }}
     uses: ./.github/workflows/unit-tests.yml
 
   cargo-deny:
+    needs: changes
+    if: ${{ !cancelled() && (github.event_name != 'push' || needs.changes.outputs.relevant == 'true' || needs.changes.outputs.any_changed != 'true') }}
     uses: ./.github/workflows/cargo-deny.yml
diff --git a/.github/workflows/clippy.yml b/.github/workflows/clippy.yml
index b5035eb9fa..dfb154aa47 100644
--- a/.github/workflows/clippy.yml
+++ b/.github/workflows/clippy.yml
@@ -1,39 +1,47 @@
 ---
 name: clippy
 
-# Runs as a job of the top-level CI workflow (.github/workflows/ci.yml),
-# which is the single push/pull_request entry point. Concurrency is governed
-# by CI's own group, so this workflow declares none of its own.
+# Runs as a job of the top-level CI workflow (.github/workflows/ci.yml).
+# Concurrency is governed by CI's own group, so this workflow declares none
+# of its own.
 on:
   workflow_call:
 
 permissions:
   contents: read
 
+# A container job's default shell is dash. The example loop uses bash
+# process substitution.
+defaults:
+  run:
+    shell: bash
+
 jobs:
   clippy:
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
+    # CUDA 13.3.1, the toolkit pin these jobs already used. The devel image
+    # carries nvcc, cudart, and libcurand. rustup is installed in the
+    # container; the image itself has no Rust.
+    container:
+      image: nvidia/cuda:13.3.1-devel-ubuntu24.04
+      env:
+        DEBIAN_FRONTEND: noninteractive
+    env:
+      CUDA_TOOLKIT_PATH: /usr/local/cuda
     timeout-minutes: 30
     steps:
+      # $/ installs git before checkout. The CUDA image has no git; without
+      # it checkout stores a tarball.
+      - uses: $/.github/actions/setup-cuda-container
+
       - name: Clone repository
         uses: actions/checkout@v4
 
-      - name: Install CUDA toolkit
-        uses: Jimver/cuda-toolkit@0ca15bb027fc96f7f61ddbb9c5221d699da06bdb  # v0.2.36
-        id: cuda-toolkit
-        with:
-          cuda: "13.3.1"
-          method: network
-          sub-packages: '["nvcc", "cudart-dev"]'
-          non-cuda-sub-packages: '["libcurand-dev"]'
-
       # `--all-targets` covers tests, benches and examples, matching the two
       # steps below and the `just clippy` recipe. Without it a lint that fires
       # only in a test target passes here and fails for anyone running the
       # recipe locally.
       - name: Run clippy (workspace crates)
-        env:
-          CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }}
         working-directory: cuda-oxide
         run: cargo clippy --workspace --all-targets -- -D warnings
 
@@ -55,14 +63,12 @@ jobs:
       # well-defined `target/` under cuda-oxide.
       - name: Run clippy (rustc-codegen-cuda)
         env:
-          CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }}
           CARGO_TARGET_DIR: ${{ runner.temp }}/clippy-shared-target
         working-directory: cuda-oxide/crates/rustc-codegen-cuda
         run: cargo clippy --all-targets -- -D warnings
 
       - 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: |
diff --git a/.github/workflows/codeql.yml b/.github/workflows/codeql.yml
index 60aaba1f3d..30d51c0e2e 100644
--- a/.github/workflows/codeql.yml
+++ b/.github/workflows/codeql.yml
@@ -1,41 +1,80 @@
 ---
 name: codeql
 
-# Push and pull_request run only when Rust, workflow, or script inputs change.
+# Push runs only when Rust, workflow, or script inputs change.
+# Self-hosted runners reject pull_request. copy-pr-bot mirrors trusted PRs to
+# pull-request/<n>, and this workflow triggers on that push
+# (config: .github/copy-pr-bot.yaml).
 # 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'
+    # No workflow-level paths: filter; the changes job below gates the jobs.
+    branches:
+      - main
+      - "pull-request/[0-9]+"
   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.
+# A container job's default shell is dash; use bash, as on ubuntu-latest.
+defaults:
+  run:
+    shell: bash
+
 concurrency:
   group: ${{ github.workflow }}-${{ github.ref_name }}${{ github.ref_name == 'main' && format('-{0}', github.sha) || '' }}
   cancel-in-progress: ${{ github.ref_name != 'main' }}
 
 jobs:
+  # Path gate. copy-pr-bot creates pull-request/<n> at an existing SHA, and
+  # GitHub often treats that push as having no path changes, so a workflow-level
+  # `paths:` filter would skip the whole run. This job diffs against main
+  # instead, as cutile-rs.yml does, and every job below runs when no file diff
+  # is detected. Other events (schedule, workflow_dispatch) skip the gate.
+  changes:
+    if: github.event_name == 'push'
+    runs-on: linux-amd64-cpu4
+    timeout-minutes: 5
+    permissions:
+      contents: read
+    outputs:
+      relevant: ${{ steps.changes.outputs.relevant }}
+      any_changed: ${{ steps.changes.outputs.any_changed }}
+    steps:
+      - uses: actions/checkout@v4
+        with:
+          fetch-depth: 0
+      - uses: dorny/paths-filter@0e4a8c6effa4802afeda77dc8d303f8176d7dfad  # v3.0.4
+        id: changes
+        with:
+          filters: |
+            relevant:
+              - '**/*.rs'
+              - '**/Cargo.toml'
+              - '**/Cargo.lock'
+              - '.github/**'
+              - '**/*.sh'
+              - '**/*.py'
+            # `*` covers a root file, `**` every nested path. Both false means
+            # no file diff was detected.
+            any_changed:
+              - '*'
+              - '**'
+
   analyze:
+    needs: changes
+    if: ${{ !cancelled() && (github.event_name != 'push' || needs.changes.outputs.relevant == 'true' || needs.changes.outputs.any_changed != 'true') }}
     name: Analyze (${{ matrix.language }})
-    runs-on: ubuntu-latest
+    # nv-cpu-general: m7i.xlarge, 4 vCPU / 16 GiB. Same size as ubuntu-latest.
+    # The VM image has no rustup. rust:1.98.0-bookworm matches the root
+    # toolchain pin and is buildpack-deps, so git is there for checkout and
+    # rustup is on PATH. The Oxide nightly still comes from
+    # cuda-oxide/rust-toolchain.toml when the extractor enters that tree.
+    runs-on: linux-amd64-cpu4
+    container:
+      image: rust:1.98.0-bookworm
     timeout-minutes: 120
     permissions:
       security-events: write
@@ -54,13 +93,40 @@ jobs:
       - name: Checkout repository
         uses: actions/checkout@v4
 
+      # The workspace is mounted from the runner, so git refuses it until this
+      # directory is marked safe. CodeQL's result upload calls git.
+      - name: Mark the workspace as a safe git directory
+        run: git config --global --add safe.directory "$GITHUB_WORKSPACE"
+
+      # A missing rustup makes the Rust extractor exit 0 with an empty database.
+      - name: Confirm Rust
+        if: matrix.language == 'rust'
+        run: rustup show active-toolchain
+
       - name: Initialize CodeQL
         uses: github/codeql-action/init@v4
         with:
           languages: ${{ matrix.language }}
           build-mode: ${{ matrix.build-mode }}
 
+      # On a copy-pr-bot mirror push the results would upload under
+      # refs/heads/pull-request/<n> and never reach the pull request's code
+      # scanning alerts. Point the upload at the pull request instead. The
+      # mirror branch carries the pull request's head commit. Empty outputs
+      # (main, schedule) leave the action on its default ref.
+      - name: Resolve the code scanning upload ref
+        id: upload
+        run: |
+          case "${GITHUB_REF_NAME}" in
+            pull-request/[0-9]*)
+              echo "ref=refs/pull/${GITHUB_REF_NAME#pull-request/}/head" >> "${GITHUB_OUTPUT}"
+              echo "sha=${GITHUB_SHA}" >> "${GITHUB_OUTPUT}"
+              ;;
+          esac
+
       - name: Perform CodeQL analysis
         uses: github/codeql-action/analyze@v4
         with:
           category: "/language:${{ matrix.language }}"
+          ref: ${{ steps.upload.outputs.ref }}
+          sha: ${{ steps.upload.outputs.sha }}
diff --git a/.github/workflows/cutile-rs.yml b/.github/workflows/cutile-rs.yml
index 521a54ef34..10c3a4793b 100644
--- a/.github/workflows/cutile-rs.yml
+++ b/.github/workflows/cutile-rs.yml
@@ -47,7 +47,7 @@ jobs:
   # 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
+    runs-on: linux-amd64-cpu4
     outputs:
       tile: ${{ steps.changes.outputs.tile }}
       any_changed: ${{ steps.changes.outputs.any_changed }}
diff --git a/.github/workflows/docs.yml b/.github/workflows/docs.yml
index c12004ba20..e5a2631af3 100644
--- a/.github/workflows/docs.yml
+++ b/.github/workflows/docs.yml
@@ -3,28 +3,13 @@ name: docs
 
 # Path-filtered SIMT lane (same roots as ci.yml).
 on:
+  # Self-hosted runners reject pull_request. copy-pr-bot mirrors trusted PRs
+  # to pull-request/<n>, and this workflow triggers on that push.
   push:
-    branches: [main]
-    paths:
-      - 'cuda-oxide/**'
-      - 'cuda-bindings/**'
-      - 'cuda-core/**'
-      - 'cuda-core-derive/**'
-      - 'cuda-async/**'
-      - 'Cargo.toml'
-      - 'Cargo.lock'
-      - '.github/workflows/**'
-  pull_request:
-    branches: [main]
-    paths:
-      - 'cuda-oxide/**'
-      - 'cuda-bindings/**'
-      - 'cuda-core/**'
-      - 'cuda-core-derive/**'
-      - 'cuda-async/**'
-      - 'Cargo.toml'
-      - 'Cargo.lock'
-      - '.github/workflows/**'
+    # No workflow-level paths: filter; the changes job below gates the jobs.
+    branches:
+      - main
+      - "pull-request/[0-9]+"
 
 # 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.
@@ -35,33 +20,79 @@ concurrency:
 permissions:
   contents: read
 
+# A container job's default shell is dash; these steps use bash.
+defaults:
+  run:
+    shell: bash
+
 # Gate rustdoc and doctests, which the `unit-tests` job does not cover (it runs
 # `--lib --tests`, excluding doctests, and never builds docs). Keeps broken
 # intra-doc links and uncompilable doc examples from accumulating silently.
 jobs:
+  # Path gate. copy-pr-bot creates pull-request/<n> at an existing SHA, and
+  # GitHub often treats that push as having no path changes, so a workflow-level
+  # `paths:` filter would skip the whole run. This job diffs against main
+  # instead, as cutile-rs.yml does, and every job below runs when no file diff
+  # is detected. Other events (schedule, workflow_dispatch) skip the gate.
+  changes:
+    if: github.event_name == 'push'
+    runs-on: linux-amd64-cpu4
+    timeout-minutes: 5
+    permissions:
+      contents: read
+    outputs:
+      relevant: ${{ steps.changes.outputs.relevant }}
+      any_changed: ${{ steps.changes.outputs.any_changed }}
+    steps:
+      - uses: actions/checkout@v4
+        with:
+          fetch-depth: 0
+      - uses: dorny/paths-filter@0e4a8c6effa4802afeda77dc8d303f8176d7dfad  # v3.0.4
+        id: changes
+        with:
+          filters: |
+            relevant:
+              - 'cuda-oxide/**'
+              - 'cuda-bindings/**'
+              - 'cuda-core/**'
+              - 'cuda-core-derive/**'
+              - 'cuda-async/**'
+              - 'Cargo.toml'
+              - 'Cargo.lock'
+              - '.github/workflows/**'
+              - '.github/actions/setup-cuda-container/**'
+            # `*` covers a root file, `**` every nested path. Both false means
+            # no file diff was detected.
+            any_changed:
+              - '*'
+              - '**'
+
   rustdoc-and-doctests:
+    needs: changes
+    if: ${{ !cancelled() && (github.event_name != 'push' || needs.changes.outputs.relevant == 'true' || needs.changes.outputs.any_changed != 'true') }}
     name: cargo doc + doctests
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
+    # `cargo doc` compiles cuda-bindings' build-time bindgen, which needs
+    # cuda.h. The devel image is CUDA 13.3.1, the same pin the toolkit
+    # install used, and it has no rustup.
+    container:
+      image: nvidia/cuda:13.3.1-devel-ubuntu24.04
+      env:
+        DEBIAN_FRONTEND: noninteractive
+    env:
+      CUDA_TOOLKIT_PATH: /usr/local/cuda
     timeout-minutes: 30
     steps:
+      # $/ installs git before checkout. The CUDA image has no git; without
+      # it checkout stores a tarball.
+      - uses: $/.github/actions/setup-cuda-container
+
       - name: Clone repository
         uses: actions/checkout@v4
 
-      - 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@0ca15bb027fc96f7f61ddbb9c5221d699da06bdb  # v0.2.36
-        id: cuda-toolkit
-        with:
-          cuda: "13.3.1"
-          method: network
-          sub-packages: '["nvcc", "cudart-dev"]'
-          non-cuda-sub-packages: '["libcurand-dev"]'
-
       - name: Build docs (warning-free) and run doctests
         working-directory: cuda-oxide
         env:
-          CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }}
           RUSTDOCFLAGS: "-D warnings"
         run: |
           sysroot="$(rustc --print sysroot)"
diff --git a/.github/workflows/examples-compile.yml b/.github/workflows/examples-compile.yml
index 6cee90749e..1731c2df25 100644
--- a/.github/workflows/examples-compile.yml
+++ b/.github/workflows/examples-compile.yml
@@ -3,28 +3,13 @@ name: examples-compile
 
 # Path-filtered SIMT lane (same roots as ci.yml).
 on:
+  # Self-hosted runners reject pull_request. copy-pr-bot mirrors trusted PRs
+  # to pull-request/<n>, and this workflow triggers on that push.
   push:
-    branches: [main]
-    paths:
-      - 'cuda-oxide/**'
-      - 'cuda-bindings/**'
-      - 'cuda-core/**'
-      - 'cuda-core-derive/**'
-      - 'cuda-async/**'
-      - 'Cargo.toml'
-      - 'Cargo.lock'
-      - '.github/workflows/**'
-  pull_request:
-    branches: [main]
-    paths:
-      - 'cuda-oxide/**'
-      - 'cuda-bindings/**'
-      - 'cuda-core/**'
-      - 'cuda-core-derive/**'
-      - 'cuda-async/**'
-      - 'Cargo.toml'
-      - 'Cargo.lock'
-      - '.github/workflows/**'
+    # No workflow-level paths: filter; the changes job below gates the jobs.
+    branches:
+      - main
+      - "pull-request/[0-9]+"
 
 # 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.
@@ -35,10 +20,65 @@ concurrency:
 permissions:
   contents: read
 
+# A container job's default shell is dash. The LLVM install and the
+# smoketest steps use `set -o pipefail`.
+defaults:
+  run:
+    shell: bash
+
 jobs:
+  # Path gate. copy-pr-bot creates pull-request/<n> at an existing SHA, and
+  # GitHub often treats that push as having no path changes, so a workflow-level
+  # `paths:` filter would skip the whole run. This job diffs against main
+  # instead, as cutile-rs.yml does, and every job below runs when no file diff
+  # is detected. Other events (schedule, workflow_dispatch) skip the gate.
+  changes:
+    if: github.event_name == 'push'
+    runs-on: linux-amd64-cpu4
+    timeout-minutes: 5
+    permissions:
+      contents: read
+    outputs:
+      relevant: ${{ steps.changes.outputs.relevant }}
+      any_changed: ${{ steps.changes.outputs.any_changed }}
+    steps:
+      - uses: actions/checkout@v4
+        with:
+          fetch-depth: 0
+      - uses: dorny/paths-filter@0e4a8c6effa4802afeda77dc8d303f8176d7dfad  # v3.0.4
+        id: changes
+        with:
+          filters: |
+            relevant:
+              - 'cuda-oxide/**'
+              - 'cuda-bindings/**'
+              - 'cuda-core/**'
+              - 'cuda-core-derive/**'
+              - 'cuda-async/**'
+              - 'Cargo.toml'
+              - 'Cargo.lock'
+              - '.github/workflows/**'
+              - '.github/actions/setup-cuda-container/**'
+            # `*` covers a root file, `**` every nested path. Both false means
+            # no file diff was detected.
+            any_changed:
+              - '*'
+              - '**'
+
   examples-compile:
+    needs: changes
+    if: ${{ !cancelled() && (github.event_name != 'push' || needs.changes.outputs.relevant == 'true' || needs.changes.outputs.any_changed != 'true') }}
     name: compile all examples (device pipeline, GPU-less)
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
+    # CUDA 13.3.1 devel image: nvcc, cudart, libcurand, and libnvJitLink.
+    # The materialization gate below needs nvJitLink, which the devel image
+    # ships. rustup is installed in the container.
+    container:
+      image: nvidia/cuda:13.3.1-devel-ubuntu24.04
+      env:
+        DEBIAN_FRONTEND: noninteractive
+    env:
+      CUDA_TOOLKIT_PATH: /usr/local/cuda
     # Every example crate: backend build + shared dep tree + per-example
     # device codegen. Typically well under an hour with the shared
     # CARGO_TARGET_DIR; the timeout leaves headroom for cold runners.
@@ -46,6 +86,11 @@ jobs:
     # focused 10-minute materialization gate, plus toolkit/LLVM setup.
     timeout-minutes: 75
     steps:
+      # $/ installs git before checkout. The CUDA image has no git; without
+      # it checkout stores a tarball. The same step installs python3 for
+      # sync-example-locks.sh.
+      - uses: $/.github/actions/setup-cuda-container
+
       - name: Clone repository
         uses: actions/checkout@v4
 
@@ -54,25 +99,10 @@ jobs:
       # dependency subtree, so one drifted Cargo.lock silently costs a full
       # dependency-tree rebuild (bindgen over cuda.h included) for that one
       # example -- a 15-70s spike in an otherwise ~2s-per-example run.
-      # Fails fast, before any toolkit install.
+      # Fails fast, before LLVM and the compile.
       - name: Example lockfiles resolve identical dependency versions
         run: bash cuda-oxide/scripts/sync-example-locks.sh --check
 
-      # 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@0ca15bb027fc96f7f61ddbb9c5221d699da06bdb  # v0.2.36
-        id: cuda-toolkit
-        with:
-          cuda: "13.3.1"
-          method: network
-          sub-packages: '["nvcc", "cudart-dev"]'
-          # nvJitLink is packaged separately from cuda-nvcc on Linux. Keep it
-          # in this already-toolkit-backed job so the materialization gate
-          # below exercises the real public command without adding work to
-          # the latency-sensitive fmt job.
-          non-cuda-sub-packages: '["libcurand-dev", "libnvjitlink-dev"]'
-
       # 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
@@ -85,20 +115,20 @@ jobs:
       # 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.
+      # The container is Ubuntu 24.04, so VERSION_CODENAME is noble.
+      # If apt.llvm.org drops llvm-toolchain-noble-23, this update fails
+      # loudly; recover by bumping the LLVM major here.
       - 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
+          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
+            | 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}-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-23
+            | tee /etc/apt/sources.list.d/llvm-23.list
+          apt-get update
+          apt-get install -y --no-install-recommends llvm-23
           llc-23 --version | head -2
           opt-23 --version | head -2
 
@@ -108,7 +138,6 @@ jobs:
         # logs. A step timeout fails the step, so the upload still runs.
         timeout-minutes: 50
         env:
-          CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }}
           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
@@ -136,7 +165,6 @@ jobs:
         timeout-minutes: 10
         working-directory: cuda-oxide
         env:
-          CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }}
           CUDA_OXIDE_LLC: /usr/bin/llc-23
           CARGO_TARGET_DIR: ${{ runner.temp }}/examples-shared-target
           CARGO_INCREMENTAL: "0"
diff --git a/.github/workflows/fmt.yml b/.github/workflows/fmt.yml
index 33ba68059d..31f29cd3f6 100644
--- a/.github/workflows/fmt.yml
+++ b/.github/workflows/fmt.yml
@@ -1,23 +1,33 @@
 ---
 name: fmt
 
-# Runs as a job of the top-level CI workflow (.github/workflows/ci.yml),
-# which is the single push/pull_request entry point. Concurrency is governed
-# by CI's own group, so this workflow declares none of its own.
+# Runs as a job of the top-level CI workflow (.github/workflows/ci.yml).
+# Concurrency is governed by CI's own group, so this workflow declares none
+# of its own.
 on:
   workflow_call:
 
 permissions:
   contents: read
 
+# A container job's default shell is dash, which rejects `set -o pipefail`.
+defaults:
+  run:
+    shell: bash
+
 jobs:
   fmt:
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
+    container:
+      image: rust:1.98.0-bookworm
     timeout-minutes: 10
     steps:
       - name: Clone repository
         uses: actions/checkout@v4
 
+      - name: Mark the workspace as a safe git directory
+        run: git config --global --add safe.directory "$GITHUB_WORKSPACE"
+
       # rustfmt comes from cuda-oxide/rust-toolchain.toml's component list,
       # like clippy,
       # so `cargo fmt` resolves against the same toolchain as the rest of CI.
diff --git a/.github/workflows/naming-guard.yml b/.github/workflows/naming-guard.yml
index 00eb30f37e..a834c168f6 100644
--- a/.github/workflows/naming-guard.yml
+++ b/.github/workflows/naming-guard.yml
@@ -9,18 +9,13 @@ name: naming-guard
 # here would rerun the same search; covering them means adding them to
 # SEARCH_ROOTS in the script first.
 on:
+  # Self-hosted runners reject pull_request. copy-pr-bot mirrors trusted PRs
+  # to pull-request/<n>, and this workflow triggers on that push.
   push:
-    branches: [main]
-    paths:
-      - 'cuda-oxide/crates/**'
-      - 'cuda-oxide/scripts/check-reserved-prefixes.sh'
-      - '.github/workflows/naming-guard.yml'
-  pull_request:
-    branches: [main]
-    paths:
-      - 'cuda-oxide/crates/**'
-      - 'cuda-oxide/scripts/check-reserved-prefixes.sh'
-      - '.github/workflows/naming-guard.yml'
+    # No workflow-level paths: filter; the changes job below gates the jobs.
+    branches:
+      - main
+      - "pull-request/[0-9]+"
 
 # 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.
@@ -32,9 +27,43 @@ permissions:
   contents: read
 
 jobs:
+  # Path gate. copy-pr-bot creates pull-request/<n> at an existing SHA, and
+  # GitHub often treats that push as having no path changes, so a workflow-level
+  # `paths:` filter would skip the whole run. This job diffs against main
+  # instead, as cutile-rs.yml does, and every job below runs when no file diff
+  # is detected. Other events (schedule, workflow_dispatch) skip the gate.
+  changes:
+    if: github.event_name == 'push'
+    runs-on: linux-amd64-cpu4
+    timeout-minutes: 5
+    permissions:
+      contents: read
+    outputs:
+      relevant: ${{ steps.changes.outputs.relevant }}
+      any_changed: ${{ steps.changes.outputs.any_changed }}
+    steps:
+      - uses: actions/checkout@v4
+        with:
+          fetch-depth: 0
+      - uses: dorny/paths-filter@0e4a8c6effa4802afeda77dc8d303f8176d7dfad  # v3.0.4
+        id: changes
+        with:
+          filters: |
+            relevant:
+              - 'cuda-oxide/crates/**'
+              - 'cuda-oxide/scripts/check-reserved-prefixes.sh'
+              - '.github/workflows/naming-guard.yml'
+            # `*` covers a root file, `**` every nested path. Both false means
+            # no file diff was detected.
+            any_changed:
+              - '*'
+              - '**'
+
   reserved-symbols-guard:
+    needs: changes
+    if: ${{ !cancelled() && (github.event_name != 'push' || needs.changes.outputs.relevant == 'true' || needs.changes.outputs.any_changed != 'true') }}
     name: reserved-oxide-symbols literal guard
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 5
     steps:
       - name: Clone repository
diff --git a/.github/workflows/status-guard.yml b/.github/workflows/status-guard.yml
index 56deee2e06..0182f2d2e5 100644
--- a/.github/workflows/status-guard.yml
+++ b/.github/workflows/status-guard.yml
@@ -1,11 +1,13 @@
 ---
 name: status-guard
 
+# Self-hosted runners reject pull_request. copy-pr-bot mirrors trusted PRs to
+# pull-request/<n>, and this workflow triggers on that push.
 on:
   push:
-    branches: [main]
-  pull_request:
-    branches: [main]
+    branches:
+      - main
+      - "pull-request/[0-9]+"
 
 # 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.
@@ -16,10 +18,15 @@ concurrency:
 permissions:
   contents: read
 
+# A container job's default shell is dash. The VM jobs already run bash.
+defaults:
+  run:
+    shell: bash
+
 jobs:
   error-example-status-guard:
     name: error-example status guard
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 5
     steps:
       - name: Clone repository
@@ -35,7 +42,7 @@ jobs:
   # marker and needed a follow-up commit on main to stop reporting FAIL.
   example-smoketest-contract:
     name: smoketest example contract
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 5
     steps:
       - name: Clone repository
@@ -50,7 +57,7 @@ jobs:
   # were both absent, 1,700 lines with no row between them.
   crate-inventory:
     name: crate inventory
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 5
     steps:
       - name: Clone repository
@@ -68,7 +75,7 @@ jobs:
   # #1609 bumped the pin to 1.99.1 and left `flake.lock` behind.
   toolchain-pin-parity:
     name: toolchain pin parity
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 5
     steps:
       - name: Clone repository
@@ -78,13 +85,17 @@ jobs:
       - name: Verify every copy of the toolchain pin agrees
         run: bash cuda-oxide/scripts/check-toolchain-parity.sh
 
+      # The Rust jobs run in rust:<version> images, a copy of the root pin.
+      - name: Verify the workflow container images match the root pin
+        run: bash cuda-oxide/scripts/check-ci-image-pins.sh
+
   # The book's Command reference is the only complete list of the CLI surface in
   # the docs, and adding a subcommand to the clap enum is a one-line change that
   # leaves it silently short -- `test`, `fmt`, `update` and `emit-ltoir` were all
   # missing when that table was written (#794).
   cli-doc-coverage:
     name: CLI command reference
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 5
     steps:
       - name: Clone repository
@@ -102,7 +113,7 @@ jobs:
   # the crate and example READMEs.
   book-api-names:
     name: book device-API names
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 5
     steps:
       - name: Clone repository
@@ -123,12 +134,17 @@ jobs:
   # bodies, and partial snippets with no enclosing item, never trip it.
   book-import-order:
     name: book import order
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
+    container:
+      image: rust:1.98.0-bookworm
     timeout-minutes: 10
     steps:
       - name: Clone repository
         uses: actions/checkout@v4
 
+      - name: Mark the workspace as a safe git directory
+        run: git config --global --add safe.directory "$GITHUB_WORKSPACE"
+
       # rustfmt comes from cuda-oxide/rust-toolchain.toml's component list through the
       # rustup proxy, as in fmt.yml, so the order checked is the pinned
       # toolchain's. No cargo, no CUDA toolkit, no network beyond the toolchain.
@@ -143,7 +159,7 @@ jobs:
   # in markdown and in `///` / `//!` doc comments. Git and Python, no toolchain.
   source-references:
     name: source references
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 10
     steps:
       - name: Clone repository
@@ -162,7 +178,7 @@ jobs:
   # because the page derives staleness from it.
   book-catalog-stamp:
     name: book catalog stamp
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 5
     steps:
       - name: Clone repository
@@ -180,7 +196,7 @@ jobs:
   # reading as complete is what hid them.
   test-matrix-coverage:
     name: test-matrix coverage
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 5
     steps:
       - name: Clone repository
@@ -200,12 +216,17 @@ jobs:
   # here.
   device-only-build:
     name: device-only kernel crate
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
+    container:
+      image: rust:1.98.0-bookworm
     timeout-minutes: 10
     steps:
       - name: Clone repository
         uses: actions/checkout@v4
 
+      - name: Mark the workspace as a safe git directory
+        run: git config --global --add safe.directory "$GITHUB_WORKSPACE"
+
       # Deliberately no CUDA toolkit step: a toolkit on the runner would mask
       # exactly the dependency this job exists to keep out of the graph.
       - name: Verify a kernel-only crate builds with no host surface
@@ -219,7 +240,7 @@ jobs:
   # snippets, which is how some forty stale lines outlived the crate switch.
   host-api-paths:
     name: host API paths in docs and templates
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 5
     steps:
       - name: Clone repository
@@ -237,7 +258,7 @@ jobs:
   # compile.
   shared-crate-pin-parity:
     name: shared host-crate pin parity
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 5
     steps:
       - name: Clone repository
@@ -253,7 +274,7 @@ jobs:
   # consumer, so the in-tree version, both manifests, and every lock must agree.
   oxide-artifacts-parity:
     name: oxide-artifacts release parity
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 5
     steps:
       - name: Clone repository
diff --git a/.github/workflows/unit-tests.yml b/.github/workflows/unit-tests.yml
index 066597a4da..69534b69c7 100644
--- a/.github/workflows/unit-tests.yml
+++ b/.github/workflows/unit-tests.yml
@@ -1,28 +1,35 @@
 ---
 name: unit-tests
 
-# Runs as a job of the top-level CI workflow (.github/workflows/ci.yml),
-# which is the single push/pull_request entry point. Concurrency is governed
-# by CI's own group, so this workflow declares none of its own.
+# Runs as a job of the top-level CI workflow (.github/workflows/ci.yml).
+# Concurrency is governed by CI's own group, so this workflow declares none
+# of its own.
 on:
   workflow_call:
 
 permissions:
   contents: read
 
+# A container job's default shell is dash, which rejects `set -o pipefail`.
+defaults:
+  run:
+    shell: bash
+
 jobs:
   rust-unit-tests:
     name: cargo test -p ${{ matrix.package }}
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
+    container:
+      image: rust:1.98.0-bookworm
     timeout-minutes: 30
     strategy:
       fail-fast: false
       matrix:
-        # Each entry must be runnable on a stock `ubuntu-latest` runner — i.e.
-        # no live GPU. `needs_cuda: true` installs the CUDA toolkit for build-
-        # time bindgen (every crate that transitively depends on
-        # `cuda-bindings` needs `cuda.h` to compile). `test_args` overrides
-        # the default `--all-targets` where a crate needs a narrower set.
+        # Each entry must be runnable on a driverless runner — no live GPU,
+        # and this image has no CUDA toolkit. `cargo-oxide` doctor asserts
+        # that absence. Crates that bindgen against cuda.h run in
+        # `rust-unit-tests-cuda`. `test_args` overrides the default
+        # `--all-targets` where a crate needs a narrower set.
         #
         # Crates intentionally not in this matrix:
         #   * `fuzzer` — differential-testing infra with no unit tests.
@@ -56,20 +63,11 @@ jobs:
           - package: dialect-nvvm
           - package: mir-importer
           # CuTe types, semantic IR, and MLIR export need no CUDA toolkit.
-          # Facade macro tests also compile host metadata via cuda-host.
+          # cute-rs and ctm-rs compile host metadata through cuda-host, so
+          # they run in rust-unit-tests-cuda.
           - package: cute-layout
-          - package: cute-rs
-            needs_cuda: true
-          - package: ctm-rs
-            needs_cuda: true
-          - package: ctm-rs
-            test_args: --all-targets --features host
-            needs_cuda: true
           - package: dialect-cute
           - package: cuda-oxide-mlir-export
-          # Includes a real llc -> ptxas spine-kernel validation.
-          - package: cuda-oxide-codegen
-            needs_cuda: true
           - package: mir-lower
           - package: mir-transforms
           - package: nvvm-transforms
@@ -86,10 +84,6 @@ jobs:
           # The Tile IR bytecode writer for those descriptions. Pure Rust.
           - package: oxide-tilebc
           - package: reserved-oxide-symbols
-          # launch_requires_unsafe trybuild fixture dev-depends on
-          # cuda-host/cuda-core, which transitively pull cuda-bindings
-          - package: cuda-macros
-            needs_cuda: true
           # trybuild compile-fail suite for ThreadIndex / DisjointSlice safety
           - package: cuda-device
           # Runtime CUDA compiler/linker discovery and file-identity checks;
@@ -100,9 +94,6 @@ jobs:
           # pure and validates options, provenance, and cubin structure; the
           # live-toolkit smoke test remains explicitly ignored.
           - package: cuda-artifact-finalizer
-          # pure-Rust kernel-launch helpers + tcgen05 tiling tables
-          - package: cuda-host
-            needs_cuda: true
           # cargo subcommand wrapper; pure Rust, no CUDA toolkit required.
           # Running it on a runner with NO toolkit installed is the
           # regression gate for issue #87.
@@ -112,19 +103,10 @@ jobs:
       - name: Clone repository
         uses: actions/checkout@v4
 
-      - name: Install CUDA toolkit
-        if: matrix.needs_cuda
-        uses: Jimver/cuda-toolkit@0ca15bb027fc96f7f61ddbb9c5221d699da06bdb  # v0.2.36
-        id: cuda-toolkit
-        with:
-          cuda: "13.3.1"
-          method: network
-          sub-packages: '["nvcc", "cudart-dev"]'
-          non-cuda-sub-packages: '["libcurand-dev"]'
+      - name: Mark the workspace as a safe git directory
+        run: git config --global --add safe.directory "$GITHUB_WORKSPACE"
 
       - name: Run crate unit tests
-        env:
-          CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }}
         working-directory: cuda-oxide
         run: |
           sysroot="$(rustc --print sysroot)"
@@ -150,14 +132,79 @@ jobs:
           printf '%s\n' "${out}"
           printf '%s\n' "${out}" | grep -q "CUDA headers (cuda.h)"
 
+  # cuda.h crates. A container is chosen per job, so these cannot share the
+  # rust image above without also handing cargo-oxide a toolkit.
+  rust-unit-tests-cuda:
+    name: cargo test -p ${{ matrix.package }}
+    runs-on: linux-amd64-cpu4
+    container:
+      image: nvidia/cuda:13.3.1-devel-ubuntu24.04
+      env:
+        DEBIAN_FRONTEND: noninteractive
+    env:
+      CUDA_TOOLKIT_PATH: /usr/local/cuda
+    timeout-minutes: 30
+    strategy:
+      fail-fast: false
+      matrix:
+        include:
+          # Includes a real llc -> ptxas spine-kernel validation.
+          - package: cuda-oxide-codegen
+          # Facade macro tests compile host metadata through cuda-host.
+          - package: cute-rs
+          - package: ctm-rs
+          - package: ctm-rs
+            test_args: --all-targets --features host
+          # launch_requires_unsafe trybuild fixture dev-depends on
+          # cuda-host/cuda-core, which transitively pull cuda-bindings.
+          - package: cuda-macros
+          # pure-Rust kernel-launch helpers + tcgen05 tiling tables.
+          # bindgen still needs cuda.h.
+          - package: cuda-host
+    steps:
+      # $/ installs git before checkout. The CUDA image has no git; without
+      # it checkout stores a tarball.
+      - uses: $/.github/actions/setup-cuda-container
+
+      - name: Clone repository
+        uses: actions/checkout@v4
+
+      - name: Install LLVM 21 (llc/opt for the NVPTX backend)
+        if: matrix.package == 'cuda-oxide-codegen'
+        run: |
+          set -euo pipefail
+          source /etc/os-release
+          install -d -m 0755 /etc/apt/keyrings
+          wget -qO- https://apt.llvm.org/llvm-snapshot.gpg.key \
+            | 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}-21 main" \
+            | tee /etc/apt/sources.list.d/llvm-21.list
+          apt-get update
+          apt-get install -y --no-install-recommends llvm-21
+          llc-21 --version | head -2
+          opt-21 --version | head -2
+
+      - name: Run crate unit tests
+        working-directory: cuda-oxide
+        run: |
+          sysroot="$(rustc --print sysroot)"
+          export LIBRARY_PATH="${sysroot}/lib:${LIBRARY_PATH:-}"
+          export LD_LIBRARY_PATH="${sysroot}/lib:${LD_LIBRARY_PATH:-}"
+          cargo test -p ${{ matrix.package }} ${{ matrix.test_args || '--all-targets' }}
+
   rustc-codegen-cuda-unit-tests:
     name: cargo test --lib (rustc-codegen-cuda)
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
+    container:
+      image: rust:1.98.0-bookworm
     timeout-minutes: 30
     steps:
       - name: Clone repository
         uses: actions/checkout@v4
 
+      - name: Mark the workspace as a safe git directory
+        run: git config --global --add safe.directory "$GITHUB_WORKSPACE"
+
       # This standalone workspace uses rustc-private APIs, so run its tests
       # directly rather than pretending the root package matrix covers them.
       # The live CUDA-tool tests remain ignored; these regressions need no GPU
@@ -172,7 +219,9 @@ jobs:
 
   generated-intrinsics:
     name: generated intrinsics are current and ABI-compatible
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
+    container:
+      image: rust:1.98.0-bookworm
     timeout-minutes: 30
     steps:
       - name: Clone repository with ABI history
@@ -180,6 +229,9 @@ jobs:
         with:
           fetch-depth: 0
 
+      - name: Mark the workspace as a safe git directory
+        run: git config --global --add safe.directory "$GITHUB_WORKSPACE"
+
       # cuda-intrinsics-gen shells out to rustfmt when it writes generated
       # sources; it comes from rust-toolchain.toml's component list.
 
@@ -198,9 +250,22 @@ jobs:
       - name: Check append-only intrinsic ABI ledger
         working-directory: cuda-oxide
         run: |
+          base_ref=""
           if [ "${{ github.event_name }}" = "pull_request" ]; then
             base_ref="${{ github.event.pull_request.base.sha }}"
           else
+            case "${GITHUB_REF_NAME}" in
+              pull-request/[0-9]*)
+                # copy-pr-bot mirror branch. On the first push `before` is all
+                # zeros and HEAD^ would check only the last commit, so a
+                # breaking ledger change in an earlier commit of a multi-commit
+                # pull request would pass. Compare against where the pull
+                # request forked from main instead.
+                base_ref="$(git merge-base HEAD origin/main || true)"
+                ;;
+            esac
+          fi
+          if [ -z "${base_ref}" ]; then
             before="${{ github.event.before }}"
             zero="0000000000000000000000000000000000000000"
             if [ -n "${before}" ] && [ "${before}" != "${zero}" ] && git cat-file -e "${before}^{commit}"; then
@@ -211,4 +276,5 @@ jobs:
               base_ref="HEAD"
             fi
           fi
+          echo "ABI history base: ${base_ref}"
           cargo run -p cuda-intrinsics-gen -- check-abi-history --base-ref "${base_ref}"
diff --git a/AGENTS.md b/AGENTS.md
index 85a13cdb67..0573f2f169 100644
--- a/AGENTS.md
+++ b/AGENTS.md
@@ -90,6 +90,8 @@ compiler components (`rustc-dev`, `llvm-tools`), including for its nested backen
 workspace. Do not float a pin to work around a failure.
 `cuda-oxide/scripts/check-toolchain-parity.sh` checks the SIMT pin against its
 scaffold, development container, and documentation.
+`cuda-oxide/scripts/check-ci-image-pins.sh` checks that the `rust:<version>`
+container images in the workflows match the root pin.
 
 rustup normally resolves the toolchain from the working directory; an explicit
 override can change that. `--manifest-path` does not change the working directory,
diff --git a/cuda-oxide/Justfile b/cuda-oxide/Justfile
index 2690e9f273..5b81a6c47c 100644
--- a/cuda-oxide/Justfile
+++ b/cuda-oxide/Justfile
@@ -219,6 +219,7 @@ check-guards:
     bash scripts/check-example-smoketest-contract.sh
     bash scripts/check-crate-inventory.sh
     bash scripts/check-toolchain-parity.sh
+    bash scripts/check-ci-image-pins.sh
     bash scripts/check-cli-doc-coverage.sh
     bash scripts/check-book-api-names.sh
     bash scripts/check-book-catalog-stamp.sh
diff --git a/cuda-oxide/scripts/check-ci-image-pins.sh b/cuda-oxide/scripts/check-ci-image-pins.sh
new file mode 100755
index 0000000000..2c6332d7c4
--- /dev/null
+++ b/cuda-oxide/scripts/check-ci-image-pins.sh
@@ -0,0 +1,53 @@
+#!/usr/bin/env bash
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+# Verify the `rust:<version>` container images in the workflows still match the
+# stable pin in the git-root rust-toolchain.toml.
+#
+# The NVIDIA runner VM image has no rustup, so the Rust jobs run in
+# `rust:<version>-bookworm`. That tag is a copy of the root pin, repeated in
+# many jobs, and nothing else keeps the copies together. A stale tag does not
+# fail a job (rustup installs the pinned toolchain on first use) but it spends
+# a toolchain download in every job and leaves the CodeQL Rust extractor on a
+# different compiler than the one it analyzes.
+#
+# Run this after bumping the root pin.
+set -euo pipefail
+
+cd "$(dirname "$0")/../.."
+
+PIN_FILE=rust-toolchain.toml
+test -s "${PIN_FILE}"
+
+pin="$(sed -n 's/^channel *= *"\([^"]*\)".*/\1/p' "${PIN_FILE}" | head -n 1)"
+if [ -z "${pin}" ]; then
+    echo "parse self-test failed: no channel found in ${PIN_FILE}" >&2
+    exit 1
+fi
+
+# Every `rust:<major.minor.patch>` image tag in the workflows and local actions.
+refs="$(grep -rnE 'rust:[0-9]+\.[0-9]+\.[0-9]+' .github || true)"
+if [ -z "${refs}" ]; then
+    echo "parse self-test failed: no rust:<version> image found under .github;" >&2
+    echo "either the jobs stopped using it (delete this check) or the pattern is stale" >&2
+    exit 1
+fi
+
+failed=0
+count=0
+while IFS= read -r line; do
+    [ -n "${line}" ] || continue
+    count=$((count + 1))
+    tag="$(printf '%s\n' "${line}" | sed -n 's/.*rust:\([0-9]*\.[0-9]*\.[0-9]*\).*/\1/p')"
+    if [ "${tag}" != "${pin}" ]; then
+        echo "  ${line%%:*}:$(printf '%s\n' "${line}" | cut -d: -f2) uses rust:${tag}, ${PIN_FILE} pins ${pin}" >&2
+        failed=1
+    fi
+done <<< "${refs}"
+
+if [ "${failed}" -ne 0 ]; then
+    echo "error: a workflow container image does not match the root toolchain pin" >&2
+    exit 1
+fi
+
+echo "OK: all ${count} rust:<version> image reference(s) match ${PIN_FILE} (${pin})."