ci: route GPU verification by machine

.github/ci/hip_smoke.cpp

@@ -1,10 +1,11 @@
-// Self-contained HIP kernel correctness smoke for the lucebox3 Strix Halo iGPU
-// (Radeon 8060S, gfx1151). Compiled and run by the gpu-tests-amd CI job to
+// Self-contained HIP kernel correctness smoke for the AMD GPUs on lucebox3.
+// Compiled and run once for gfx1201 and once for gfx1151 by gpu-tests-amd to
 // prove the ROCm/HIP compute path actually executes on AMD hardware (GitHub
 // hosted runners have no GPU). No model weights, deterministic, fast.
 #include <hip/hip_runtime.h>
 #include <cstdio>
 #include <cstdlib>
+#include <cstring>
 
 __global__ void vadd(const float* a, const float* b, float* c, int n) {
     int i = blockIdx.x * blockDim.x + threadIdx.x;
@@ -15,10 +16,28 @@ __global__ void vadd(const float* a, const float* b, float* c, int n) {
     fprintf(stderr, "HIP error at line %d: %s\n", __LINE__, hipGetErrorString(e)); \
     return 2; } } while (0)
 
-int main() {
+static bool arch_matches(const char* actual, const char* expected) {
+    const size_t expected_len = std::strlen(expected);
+    return std::strncmp(actual, expected, expected_len) == 0 &&
+           (actual[expected_len] == '\0' || actual[expected_len] == ':');
+}
+
+int main(int argc, char** argv) {
+    if (argc != 2) {
+        std::fprintf(stderr, "usage: %s <expected-gfx-arch>\n", argv[0]);
+        return 2;
+    }
+
     hipDeviceProp_t p;
     CK(hipGetDeviceProperties(&p, 0));
     printf("HIP device 0: %s (%s, %d CUs)\n", p.name, p.gcnArchName, p.multiProcessorCount);
+    if (!arch_matches(p.gcnArchName, argv[1])) {
+        std::fprintf(stderr,
+                     "HIP device selection mismatch: expected %s after "
+                     "HIP_VISIBLE_DEVICES filtering, got %s (%s)\n",
+                     argv[1], p.gcnArchName, p.name);
+        return 3;
+    }
 
     const int n = 1 << 20;
     const size_t sz = n * sizeof(float);

.github/workflows/ci.yml

@@ -133,75 +133,100 @@ jobs:
           "
 
   gpu-tests:
-    name: GPU tests (self-hosted RTX 3090, sm_86)
-    # Runs only after the GitHub-hosted jobs pass. Those compile the dflash
-    # kernels for sm_86 but execute on GPU-less VMs, so the kernel never runs.
-    # This job lands on lucebox3 (RTX 3090) and actually executes it.
+    name: ${{ matrix.job_name }}
+    # Execute the shared CUDA verifier suite on both NVIDIA architectures.
+    # Machine-specific labels prevent stale labels on another box from accepting
+    # a job after a hardware change. The sparse-attention allocator smoke remains
+    # 3090-only; all portable verifier tests are defined once below.
     #
     # SECURITY: fork PRs DO run here, but only after a maintainer approves the
     # workflow run. The repo enforces approval_policy=all_external_contributors,
     # so every push from an outside collaborator parks the run as
     # action_required until a maintainer reviews the diff and approves it.
     # That human checkpoint is the line of defense for the self-hosted box;
     # do not weaken the approval policy while this job runs fork code.
-    # Gate only on the 1-minute workspace check, NOT the ~18-minute hosted
-    # CPU build: the GPU job compiles the same tree itself in ~2 minutes on
-    # real hardware, so serializing it behind the cloud build only delayed
-    # the strongest signal this CI produces.
+    # Gate only on the fast workspace check, not the hosted CPU build. Each GPU
+    # compiles and executes the focused suite independently on real hardware.
     needs: [uv-workspace]
-    runs-on: [self-hosted, gpu, sm86]
+    strategy:
+      fail-fast: false
+      matrix:
+        include:
+          - job_name: GPU tests (self-hosted RTX 3090, sm_86)
+            runner: lucebox-rtx3090
+            arch: "86"
+            compute_cap: "8.6"
+            concurrency_group: lucebox-rtx3090-gpu-runner
+            sparse_attention: "true"
+          - job_name: GPU tests (self-hosted DGX GB10, sm_121 verifier)
+            runner: dgx-gb10
+            arch: "121"
+            compute_cap: "12.1"
+            concurrency_group: dgx-gb10-a197-gpu-runner
+            sparse_attention: "false"
+    runs-on: [self-hosted, "${{ matrix.runner }}"]
     timeout-minutes: 30
-    # Serialize CUDA jobs across PRs (one RTX 3090). The ROCm job has its
-    # own group: different physical GPU, no contention.
+    env:
+      CUDA_HOME: /usr/local/cuda
+      CUDACXX: /usr/local/cuda/bin/nvcc
+      CUDA_ARCH: ${{ matrix.arch }}
+      EXPECTED_COMPUTE_CAP: ${{ matrix.compute_cap }}
+      RUN_SPARSE_ATTENTION_TEST: ${{ matrix.sparse_attention }}
     concurrency:
-      group: lucebox3-gpu-runner
+      group: ${{ matrix.concurrency_group }}
       cancel-in-progress: false
     steps:
       - uses: actions/checkout@df4cb1c069e1874edd31b4311f1884172cec0e10 # v6.0.3
         with:
           submodules: recursive
           token: ${{ secrets.SUBMODULE_PAT || secrets.GITHUB_TOKEN }}
 
-      - name: GPU smoke (nvidia-smi)
-        run: nvidia-smi --query-gpu=name,driver_version,memory.total,power.limit --format=csv
+      - name: GPU and CUDA smoke (${{ matrix.job_name }})
+        run: |
+          nvidia-smi --query-gpu=name,compute_cap,driver_version,memory.total,power.limit --format=csv
+          test "$(nvidia-smi --query-gpu=compute_cap --format=csv,noheader | tr -d '[:space:]')" = "$EXPECTED_COMPUTE_CAP"
+          "$CUDACXX" --version
 
-      - name: Build GPU test binary (sm_86, runner-local CUDA)
-        # Uses the runner's own nvcc (/usr/bin/nvcc); no Jimver download since
-        # lucebox3 already has the CUDA toolkit installed.
+      - name: Build focused NVIDIA verifier tests (sm_${{ matrix.arch }})
         run: |
           cd server
           cmake -B build \
-            -DCMAKE_CUDA_ARCHITECTURES="86" \
+            -DCMAKE_CUDA_COMPILER="$CUDACXX" \
+            -DCMAKE_CUDA_ARCHITECTURES="$CUDA_ARCH" \
+            -DDFLASH27B_USER_CUDA_ARCHITECTURES="$CUDA_ARCH" \
             -DDFLASH27B_ENABLE_BSA=OFF \
             -DDFLASH27B_FA_ALL_QUANTS=OFF \
             -DCMAKE_BUILD_TYPE=Release
-          cmake --build build \
-            --target test_flash_attn_sparse test_deepseek4_mmid_grouped_cuda \
-              test_deepseek4_unit test_rocmfp3_mix_registry \
-              test_rocmfp_mix_slice_matvec test_rocmfp_mix_gateup_glu \
-              test_ds4_mix_registry_teardown \
-            -j"$(nproc)"
-
-      - name: Run flash-attn sparse kernel test on the 3090
-        # Self-contained kernel correctness test (no model weights), so it is
-        # safe and fast in CI. This is the execution the hosted jobs cannot do.
-        run: ./server/build/test_flash_attn_sparse
-
-      - name: Run grouped MMID dispatch and parity test on the 3090
-        run: ./server/build/test_deepseek4_mmid_grouped_cuda
-
-      - name: Run DeepSeek4 graph-generation regression on the 3090
-        run: ./server/build/test_deepseek4_unit
+          if [[ "$CUDA_ARCH" == "121" ]]; then
+            grep -q '^DFLASH27B_EFFECTIVE_BLACKWELL_CONSUMER:INTERNAL=OFF$' build/CMakeCache.txt
+          fi
+          targets=(
+            test_deepseek4_mmid_grouped_cuda
+            test_deepseek4_unit
+            test_rocmfp3_mix_registry
+            test_rocmfp_mix_slice_matvec
+            test_rocmfp_mix_gateup_glu
+            test_ds4_mix_registry_teardown
+          )
+          if [[ "$RUN_SPARSE_ATTENTION_TEST" == "true" ]]; then
+            targets+=(test_flash_attn_sparse)
+          fi
+          cmake --build build --target "${targets[@]}" -j"$(nproc)"
 
-      - name: Run mixed-quantization registry safety tests on the 3090
+      - name: Run focused NVIDIA verifier tests
         run: |
+          if [[ "$RUN_SPARSE_ATTENTION_TEST" == "true" ]]; then
+            ./server/build/test_flash_attn_sparse
+          fi
+          ./server/build/test_deepseek4_mmid_grouped_cuda
+          ./server/build/test_deepseek4_unit
           ctest --test-dir server/build --output-on-failure \
             -R 'rocmfp3_mix_registry|rocmfp_mix_slice_matvec|rocmfp_mix_gateup_glu|ds4_mix_registry_teardown'
 
       # Optional model-backed end-to-end smoke (real spec-decode on the 3090),
       # disabled by default because it builds dflash_server and lazy-loads the
       # ~16 GB Qwen3.6-27B target + draft (~1-2 min). The weights are already
-      # staged at /opt/models on lucebox3 (override with repo var
+      # staged at /opt/models on lucebox-rtx3090 (override with repo var
       # LUCEBOX_MODELS_DIR). Verified working on the runner. To enable, uncomment;
       # continue-on-error keeps a heavy/slow run from ever blocking a PR.
       # - name: dflash end-to-end smoke (model-backed)
@@ -220,32 +245,36 @@ jobs:
       #       -d '{"prompt":"The capital of France is","max_tokens":8}' | grep -q '"text"'
 
   gpu-tests-amd:
-    name: GPU tests (self-hosted Radeon 8060S, gfx1151 / ROCm)
-    # Companion to gpu-tests: exercises the OTHER half of lucebox3 - the Strix
-    # Halo iGPU - via ROCm/HIP. GitHub-hosted runners cannot touch an AMD GPU at
-    # all. The runner is pinned to ROCm 6.4.4 on kernel 6.14 (ROCm 7.2 page-
-    # faults on gfx1151); hipcc is not on the runner's minimal PATH, so it is
-    # invoked by absolute path.
+    name: GPU tests (${{ matrix.device_name }}, ${{ matrix.arch }} / ROCm)
+    # Companion to the NVIDIA jobs: lucebox3 carries both an R9700 and the Strix
+    # Halo iGPU. Run the same HIP suite once per device with an explicit visibility
+    # mask. GitHub-hosted runners cannot execute these kernels.
+    # hipcc is not on the runner's minimal PATH, so invoke it by absolute path.
     #
     # SECURITY: same model as gpu-tests - fork PRs run only after a maintainer
     # approves the workflow run (approval_policy=all_external_contributors).
-    # Gate only on the 1-minute workspace check and the NVIDIA GPU job, NOT the
-    # ~18-minute hosted CPU build: the GPU job compiles the same tree itself in
-    # ~2 minutes on real hardware. The NVIDIA dependency keeps the two
-    # self-hosted GPU jobs from entering the same concurrency group
-    # simultaneously; GitHub cancels older pending jobs in a group even when
-    # cancel-in-progress is false.
-    needs: [uv-workspace, gpu-tests]
-    if: ${{ always() && needs['uv-workspace'].result == 'success' }}
-    runs-on: [self-hosted, rocm, gfx1151]
+    # Gate only on the fast workspace check, not the hosted CPU build. The two AMD
+    # matrix jobs share one physical runner and therefore execute one at a time.
+    needs: [uv-workspace]
+    strategy:
+      fail-fast: false
+      matrix:
+        include:
+          - device_name: Radeon AI PRO R9700
+            arch: gfx1201
+            device_index: "0"
+            concurrency_key: r9700
+          - device_name: Strix Halo Radeon 8060S
+            arch: gfx1151
+            device_index: "1"
+            concurrency_key: strix-halo
+    runs-on: [self-hosted, lucebox3]
     timeout-minutes: 20
-    # Serialize across PRs per GPU. NOT the same group as the CUDA job:
-    # the combo box has two distinct GPUs (RTX 3090 + Strix iGPU), and a
-    # shared group only holds one waiting job, so the Radeon leg was
-    # chronically displaced ("higher priority waiting request") by every
-    # new CUDA job entering the queue.
+    env:
+      HIP_VISIBLE_DEVICES: ${{ matrix.device_index }}
+      EXPECTED_HIP_ARCH: ${{ matrix.arch }}
     concurrency:
-      group: lucebox3-rocm-runner
+      group: lucebox3-${{ matrix.concurrency_key }}-runner
       cancel-in-progress: false
     steps:
       - uses: actions/checkout@df4cb1c069e1874edd31b4311f1884172cec0e10 # v6.0.3
@@ -282,22 +311,23 @@ jobs:
           fi
           wait $PROBE && echo "KFD healthy" || { echo "::error::rocminfo exited non-zero"; cat /tmp/rocminfo.out | tail -5; exit 1; }
 
-      - name: ROCm smoke (rocminfo sees gfx1151)
-        run: cat /tmp/rocminfo.out | grep -E "Name:|Marketing Name:" | grep -iE "gfx1151|Radeon 8060S"
+      - name: ROCm inventory includes ${{ matrix.arch }}
+        run: cat /tmp/rocminfo.out | grep -E "Name:|Marketing Name:" | grep -F "${{ matrix.arch }}"
 
-      - name: Build + run HIP vector-add on the Radeon 8060S
+      - name: Build + run HIP vector-add on ${{ matrix.device_name }}
         # Self-contained HIP kernel correctness test (no model weights). This is
-        # the execution the GitHub-hosted jobs cannot do.
+        # the execution the GitHub-hosted jobs cannot do. The binary also checks
+        # that HIP_VISIBLE_DEVICES selected the matrix's expected architecture.
         run: |
-          /opt/rocm/bin/hipcc --offload-arch=gfx1151 -O2 \
+          /opt/rocm/bin/hipcc --offload-arch="$EXPECTED_HIP_ARCH" -O2 \
             -o "$RUNNER_TEMP/hip_smoke" .github/ci/hip_smoke.cpp
-          "$RUNNER_TEMP/hip_smoke"
+          "$RUNNER_TEMP/hip_smoke" "$EXPECTED_HIP_ARCH"
 
       - name: Build + test ROCm formats and inference core
         run: |
           cmake -S server -B "$RUNNER_TEMP/rocmfp-build" \
             -DDFLASH27B_GPU_BACKEND=hip \
-            -DDFLASH27B_HIP_ARCHITECTURES=gfx1151 \
+            -DDFLASH27B_HIP_ARCHITECTURES="$EXPECTED_HIP_ARCH" \
             -DDFLASH27B_SERVER=OFF \
             -DDFLASH27B_TESTS=ON \
             -DGGML_HIP_GRAPHS=ON \

.github/workflows/docker.yml

@@ -191,13 +191,13 @@ jobs:
           # dispatch builds keep the full consumer-GPU list so the published
           # image runs on every supported card.
           DFLASH_CUDA_ARCHES: ${{ github.event_name == 'pull_request' && '86' || '75;80;86;89;90;120' }}
-          # Same split for HIP: PR builds compile gfx1151 only (the verified
-          # reference, Strix Halo); main / release builds widen to consumer
+          # Same split for HIP: PR builds compile gfx1201 only (the verified
+          # R9700); main / release builds widen to consumer
           # RDNA so the published :rocm runs on RX 7900 (gfx1100) and both RDNA4
           # code objects — gfx1200 (RX 9060) and gfx1201 (RX 9070 / Radeon AI
           # PRO R9700), which are NOT binary-compatible with each other. CDNA
           # datacenter parts (gfx90a/gfx942) stay out until someone can test them.
-          DFLASH_HIP_ARCHES: ${{ github.event_name == 'pull_request' && 'gfx1151' || 'gfx1151;gfx1100;gfx1200;gfx1201' }}
+          DFLASH_HIP_ARCHES: ${{ github.event_name == 'pull_request' && 'gfx1201' || 'gfx1151;gfx1100;gfx1200;gfx1201' }}
         with:
           # bake-action v6+ defaults `source` to the remote git context, which
           # makes the local metadata-action bake-file (written to the runner

.github/workflows/speed-profile.yml

@@ -1,7 +1,7 @@
 name: Speed Profile
 
 # Report-only speed profile for the inference engine. Runs on the self-hosted
-# RTX 3090 (lucebox3) on PRs that touch the engine or the optimizations, and on
+# RTX 3090 (lucebox-rtx3090) on PRs that touch the engine or the optimizations, and on
 # manual dispatch. It NEVER blocks a PR (continue-on-error: true) — it publishes a
 # report to the run summary + uploads the JSON / markdown / nsys trace as artifacts.
 #
@@ -22,13 +22,13 @@ on:
 # fight over the GPU. cancel-in-progress=false: let a queued profile finish rather
 # than killing it mid-measurement.
 concurrency:
-  group: lucebox3-gpu-runner
+  group: lucebox-rtx3090-gpu-runner
   cancel-in-progress: false
 
 jobs:
   speed-profile:
     name: Speed profile (self-hosted RTX 3090, sm_86)
-    runs-on: [self-hosted, gpu, sm86]
+    runs-on: [self-hosted, lucebox-rtx3090]
     timeout-minutes: 30
     continue-on-error: true   # report-only: a slow/failed profile must not block the PR
 
@@ -38,6 +38,8 @@ jobs:
     # draft path; keep the older LUCEBOX_SPEED_PROFILE_* variable names as aliases
     # so existing repo settings continue to work.
     env:
+      CUDA_HOME: /usr/local/cuda
+      CUDACXX: /usr/local/cuda/bin/nvcc
       MODELS: ${{ vars.LUCEBOX_MODELS_DIR || '/opt/models' }}
       TARGET_MODEL: ${{ vars.LUCEBOX_TARGET_MODEL || vars.LUCEBOX_SPEED_PROFILE_TARGET || 'Qwen3.6-27B-Q4_K_M.gguf' }}
       DRAFT_MODEL: ${{ vars.LUCEBOX_DRAFT_MODEL || vars.LUCEBOX_SPEED_PROFILE_DRAFT || 'draft/dflash-draft-3.6-q4_k_m.gguf' }}
@@ -51,7 +53,8 @@ jobs:
 
       - name: GPU info (and pin clocks to cut variance, if permitted)
         run: |
-          nvidia-smi --query-gpu=name,driver_version,memory.total,power.limit --format=csv
+          nvidia-smi --query-gpu=name,compute_cap,driver_version,memory.total,power.limit --format=csv
+          test "$(nvidia-smi --query-gpu=compute_cap --format=csv,noheader | tr -d '[:space:]')" = "8.6"
           # Locking clocks makes the numbers comparable run-to-run. Safe to skip if the
           # runner user can't run nvidia-smi -lgc; the profiler still records the power cap.
           sudo nvidia-smi -lgc 1395 2>/dev/null || echo "clock lock not permitted; continuing"
@@ -99,7 +102,9 @@ jobs:
         run: |
           cd server
           cmake -B build \
+            -DCMAKE_CUDA_COMPILER="$CUDACXX" \
             -DCMAKE_CUDA_ARCHITECTURES="86" \
+            -DDFLASH27B_USER_CUDA_ARCHITECTURES="86" \
             -DDFLASH27B_ENABLE_BSA=OFF \
             -DDFLASH27B_FA_ALL_QUANTS=OFF \
             -DCMAKE_BUILD_TYPE=Release