diff --git a/.github/actions/setup-cuda-container/action.yml b/.github/actions/setup-cuda-container/action.yml
new file mode 100644
index 0000000000..90a95c5a19
--- /dev/null
+++ b/.github/actions/setup-cuda-container/action.yml
@@ -0,0 +1,36 @@
+# 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. The caller checks the
+  repository out first: a local action is resolved from the workspace, so
+  this action cannot be the job's first step. 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
+        apt-get install -y -qq --no-install-recommends \
+          build-essential ca-certificates curl git pkg-config libclang-dev 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
+      run: |
+        curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs \
+          | sh -s -- -y --no-modify-path --profile minimal --default-toolchain none
+        echo "${HOME}/.cargo/bin" >> "${GITHUB_PATH}"
+        # Root rust-toolchain.toml is the stable pin. cuda-oxide/ selects the
+        # nightly when a later step sets that working directory.
+        "${HOME}/.cargo/bin/rustup" show
+        (cd cuda-oxide && "${HOME}/.cargo/bin/rustup" show)
diff --git a/.github/workflows/book.yml b/.github/workflows/book.yml
index 732ff8b667..ec0746e717 100644
--- a/.github/workflows/book.yml
+++ b/.github/workflows/book.yml
@@ -1,21 +1,14 @@
 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]
+    branches:
+      - main
+      - "pull-request/[0-9]+"
     paths:
       - 'cuda-oxide/cuda-oxide-book/**'
       - 'cutile-rs/cutile-book/**'
@@ -39,11 +32,11 @@ 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:
   build:
-    runs-on: ubuntu-latest
+    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 +69,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 +79,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 +89,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..13c1df6c1e 100644
--- a/.github/workflows/ci.yml
+++ b/.github/workflows/ci.yml
@@ -4,20 +4,15 @@ 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]
+    branches:
+      - main
+      - "pull-request/[0-9]+"
     paths:
       - 'cuda-oxide/**'
       - 'cuda-bindings/**'
diff --git a/.github/workflows/clippy.yml b/.github/workflows/clippy.yml
index 018779dbb9..b40aeaa069 100644
--- a/.github/workflows/clippy.yml
+++ b/.github/workflows/clippy.yml
@@ -1,39 +1,50 @@
 ---
 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.1.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.1.1-devel-ubuntu24.04
+      env:
+        DEBIAN_FRONTEND: noninteractive
+    env:
+      CUDA_TOOLKIT_PATH: /usr/local/cuda
     timeout-minutes: 30
     steps:
+      # A local action is loaded from the workspace, so checkout comes first.
+      # The CUDA image has no git; without it checkout stores a tarball.
+      - name: Install git
+        run: |
+          apt-get update -qq
+          apt-get install -y -qq --no-install-recommends ca-certificates git
       - name: Clone repository
         uses: actions/checkout@v4
-
-      - name: Install CUDA toolkit
-        uses: Jimver/cuda-toolkit@v0.2.35
-        id: cuda-toolkit
-        with:
-          cuda: "13.1.1"
-          method: network
-          sub-packages: '["nvcc", "cudart-dev"]'
-          non-cuda-sub-packages: '["libcurand-dev"]'
+      - uses: ./.github/actions/setup-cuda-container
 
       # `--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 +66,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 (rustc-codegen-cuda examples)
         env:
-          CUDA_TOOLKIT_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }}
           CARGO_TARGET_DIR: ${{ runner.temp }}/clippy-shared-target
         run: |
           set -e
diff --git a/.github/workflows/codeql.yml b/.github/workflows/codeql.yml
index 60aaba1f3d..2c84a2e2c3 100644
--- a/.github/workflows/codeql.yml
+++ b/.github/workflows/codeql.yml
@@ -1,21 +1,17 @@
 ---
 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]
+    branches:
+      - main
+      - "pull-request/[0-9]+"
     paths:
       - '**/*.rs'
       - '**/Cargo.toml'
@@ -28,6 +24,11 @@ on:
 
 # 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' }}
@@ -35,7 +36,14 @@ concurrency:
 jobs:
   analyze:
     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,6 +62,16 @@ 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:
diff --git a/.github/workflows/cutile-rs.yml b/.github/workflows/cutile-rs.yml
index 665ed14296..9afa02a557 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 c924c78d7f..b611f9df43 100644
--- a/.github/workflows/docs.yml
+++ b/.github/workflows/docs.yml
@@ -3,19 +3,12 @@ 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]
+    branches:
+      - main
+      - "pull-request/[0-9]+"
     paths:
       - 'cuda-oxide/**'
       - 'cuda-bindings/**'
@@ -35,33 +28,42 @@ 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:
   rustdoc-and-doctests:
     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.1.1, the same pin the toolkit
+    # install used, and it has no rustup.
+    container:
+      image: nvidia/cuda:13.1.1-devel-ubuntu24.04
+      env:
+        DEBIAN_FRONTEND: noninteractive
+    env:
+      CUDA_TOOLKIT_PATH: /usr/local/cuda
     timeout-minutes: 30
     steps:
+      # A local action is loaded from the workspace, so checkout comes first.
+      # The CUDA image has no git; without it checkout stores a tarball.
+      - name: Install git
+        run: |
+          apt-get update -qq
+          apt-get install -y -qq --no-install-recommends ca-certificates git
       - 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@v0.2.35
-        id: cuda-toolkit
-        with:
-          cuda: "13.1.1"
-          method: network
-          sub-packages: '["nvcc", "cudart-dev"]'
-          non-cuda-sub-packages: '["libcurand-dev"]'
+      - uses: ./.github/actions/setup-cuda-container
 
       - 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 0f8544cb63..cca315e663 100644
--- a/.github/workflows/examples-compile.yml
+++ b/.github/workflows/examples-compile.yml
@@ -3,19 +3,12 @@ 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]
+    branches:
+      - main
+      - "pull-request/[0-9]+"
     paths:
       - 'cuda-oxide/**'
       - 'cuda-bindings/**'
@@ -35,10 +28,25 @@ 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:
   examples-compile:
     name: compile all examples (device pipeline, GPU-less)
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
+    # CUDA 13.1.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.1.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,33 +54,27 @@ jobs:
     # focused 10-minute materialization gate, plus toolkit/LLVM setup.
     timeout-minutes: 75
     steps:
+      # A local action is loaded from the workspace, so checkout comes first.
+      # The CUDA image has no git; without it checkout stores a tarball.
+      # sync-example-locks.sh --check is a python3 stdlib script. ubuntu-latest
+      # had python3; this image does not.
+      - name: Install git and Python
+        run: |
+          apt-get update -qq
+          apt-get install -y -qq --no-install-recommends ca-certificates git python3
       - name: Clone repository
         uses: actions/checkout@v4
+      - uses: ./.github/actions/setup-cuda-container
 
       # Every example workspace must pin the same dependency versions. The
       # shared CARGO_TARGET_DIR below keys build caches on the resolved
       # 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@v0.2.35
-        id: cuda-toolkit
-        with:
-          cuda: "13.1.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
@@ -84,20 +86,20 @@ jobs:
       # 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
-      # 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-22, this update fails
+      # loudly; recover by bumping the LLVM major here.
       - name: Install LLVM 22 (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}-22 main" \
-            | sudo tee /etc/apt/sources.list.d/llvm-22.list
-          sudo apt-get update
-          sudo apt-get install -y --no-install-recommends llvm-22
+            | tee /etc/apt/sources.list.d/llvm-22.list
+          apt-get update
+          apt-get install -y --no-install-recommends llvm-22
           llc-22 --version | head -2
           opt-22 --version | head -2
 
@@ -107,7 +109,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-22
           # Each example is its own cargo workspace; without a shared
           # target dir the dep tree rebuilds ~85 times and blows the
@@ -135,7 +136,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-22
           CARGO_TARGET_DIR: ${{ runner.temp }}/examples-shared-target
           CARGO_INCREMENTAL: "0"
diff --git a/.github/workflows/fmt.yml b/.github/workflows/fmt.yml
index 4d8035a248..c3c2d96ec1 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 nightly as the rest of CI.
diff --git a/.github/workflows/naming-guard.yml b/.github/workflows/naming-guard.yml
index 00eb30f37e..09c1f11081 100644
--- a/.github/workflows/naming-guard.yml
+++ b/.github/workflows/naming-guard.yml
@@ -9,14 +9,12 @@ 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]
+    branches:
+      - main
+      - "pull-request/[0-9]+"
     paths:
       - 'cuda-oxide/crates/**'
       - 'cuda-oxide/scripts/check-reserved-prefixes.sh'
@@ -34,7 +32,7 @@ permissions:
 jobs:
   reserved-symbols-guard:
     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 1bb635e1df..7e423bfcee 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
@@ -69,7 +76,7 @@ jobs:
   # book quotes behind.
   toolchain-pin-parity:
     name: toolchain pin parity
-    runs-on: ubuntu-latest
+    runs-on: linux-amd64-cpu4
     timeout-minutes: 5
     steps:
       - name: Clone repository
@@ -85,7 +92,7 @@ jobs:
   # 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
@@ -103,7 +110,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
@@ -124,12 +131,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
       # nightly's. No cargo, no CUDA toolkit, no network beyond the toolchain.
@@ -144,7 +156,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
@@ -163,7 +175,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
@@ -181,7 +193,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
@@ -201,12 +213,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
@@ -220,7 +237,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
@@ -238,7 +255,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
@@ -254,7 +271,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 cf9a33179b..9ef31fba2e 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.
@@ -55,9 +62,6 @@ jobs:
           - package: dialect-mir
           - package: dialect-nvvm
           - package: mir-importer
-          # Includes a real llc -> ptxas spine-kernel validation.
-          - package: cuda-oxide-codegen
-            needs_cuda: true
           - package: mir-lower
           - package: mir-transforms
           - package: nvvm-transforms
@@ -69,10 +73,6 @@ jobs:
           - package: oxide-artifacts
             test_args: --all-targets --features object
           - 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;
@@ -83,9 +83,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.
@@ -95,34 +92,10 @@ jobs:
       - name: Clone repository
         uses: actions/checkout@v4
 
-      - name: Install CUDA toolkit
-        if: matrix.needs_cuda
-        uses: Jimver/cuda-toolkit@v0.2.35
-        id: cuda-toolkit
-        with:
-          cuda: "13.1.1"
-          method: network
-          sub-packages: '["nvcc", "cudart-dev"]'
-          non-cuda-sub-packages: '["libcurand-dev"]'
-
-      - name: Install LLVM 21 (llc/opt for the NVPTX backend)
-        if: matrix.package == 'cuda-oxide-codegen'
-        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}-21 main" \
-            | sudo tee /etc/apt/sources.list.d/llvm-21.list
-          sudo apt-get update
-          sudo apt-get install -y --no-install-recommends llvm-21
-          llc-21 --version | head -2
-          opt-21 --version | head -2
+      - 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)"
@@ -148,14 +121,77 @@ 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.1.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
+          # 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:
+      # A local action is loaded from the workspace, so checkout comes first.
+      # The CUDA image has no git; without it checkout stores a tarball.
+      - name: Install git
+        run: |
+          apt-get update -qq
+          apt-get install -y -qq --no-install-recommends ca-certificates git
+      - name: Clone repository
+        uses: actions/checkout@v4
+      - uses: ./.github/actions/setup-cuda-container
+
+      - 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 }} --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
@@ -170,7 +206,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
@@ -178,6 +216,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.