HRX CI: run on self-hosted ROCm GPU runners for gfx1151/gfx1201 (#22)

* HRX CI: run on self-hosted ROCm GPU runners for gfx1151/gfx1201

Replace the ubuntu-latest CPU job with a matrix on
linux-gfx1151-gpu-rocm and linux-gfx120X-gpu-rocm, building for the
runner's gfx target and running the sample MUL_MAT correctness/perf
configs on the HRX0 backend. Fetch per-arch rocwmma dev artifacts.

BENCH_REF temporarily points at the p066-multi-arch-ci branch of the
bench tooling; restore to main when this merges.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Options should already be set up in the container

* debug

* HRX CI: run container as root, write build outputs off the bind mount

Run the job container with --user 0:0 so checkout can write the
runner's _work tree on runners whose agent owns it as root (the plain
shark65 runner); the image's default tester user (uid 1001) only works
on the prolense fleet where _work is 1001-owned. Add --device kfd/dri
for GPU access on runners that don't inject devices.

Point HRX_WORK_DIR and LLAMA_BUILD_DIR at a container-local /work path so
ROCm downloads and build trees are reclaimed on container removal instead
of accumulating in the persistent runner's bind-mounted workspace.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* cleanup comments

* HRX CI: source runtime-env.sh in op-test/perf steps

The bench tooling moved runtime env setup (LD_LIBRARY_PATH incl. the HRX
install lib, and HSA_FORCE_FINE_GRAIN_PCIE=1 for gfx12) out of
fetch-rocm-assets.sh into runtime-env.sh. Build/validate scripts source
it internally; the op-test and op-perf steps invoke the python tools
directly, so source runtime-env.sh there too. Without it those steps
would now run with neither the HRX libs on LD_LIBRARY_PATH nor the gfx12
fine-grained-memory flag, which is what produced the MUL_MAT NaNs.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* ggml-hrx: port to renamed HRX export-info field (constant_byte_length)

Upstream HRX renamed hrx_executable_export_info_t.constant_count ->
constant_byte_length (now the byte length of the constants buffer, not a uint32
count). Update the catalog-kernel ABI check accordingly:
  constant_count * sizeof(uint32_t) == constants_size
becomes
  constant_byte_length == constants_size
and fix the corresponding diagnostic log. Needed to build ggml-hrx against the
bumped HRX (ROCm/hrx upstream main); verified building against e8275fbb5.

* cleanup

* ci: on hrx-integration, run only build-hrx.yml (push + pull_request)

The fork's hrx-integration branch should run only the HRX CI, not the upstream
llama.cpp matrix.
- build-hrx.yml: add a push trigger for hrx-integration (was pull_request-only).
- Every other workflow whose push/pull_request/pull_request_target could match
  hrx-integration now carries branches-ignore: ['hrx-integration'] (build.yml,
  build-vulkan/cann/apple/android/riscv/self-hosted, server*, check-vendor,
  hip-quality-check, python-*, labeler, copilot-setup-steps, etc.).
- gguf-publish is tag-only (can't fire on a branch push) and the rest were
  already scoped to master.

Verified: only build-hrx.yml fires on hrx-integration for push and PR events.
Caveat: these touch upstream workflow files, so upstream merges into
hrx-integration may conflict on them (build-hrx.yml itself does not).

* ignore

* ci(hrx): runner-identity debug block + GPU health-check step

- Debug (runner identity): first step, if: always(), logs RUNNER_NAME, hostname,
  uname, /dev/kfd|dri, lspci, RUNNER_*/GITHUB_* env -- so the specific
  self-hosted runner behind a failure is identifiable from the log alone.
- GPU health check: after 'Fetch ROCm assets' (prefix binaries present), runs
  the bench runner-health.sh (rocminfo/amd-smi/rocm-smi diagnostics + fail-fast
  on /dev/kfd missing, dead rocminfo, or expected gfx arch not enumerated).

Caveat: neither detects the gfx1151 coherence NaN flake (GPU reports healthy);
the debug block enables post-hoc identification. runner-health.sh must be on the
bench at BENCH_REF (AaronStGeorge/llamacpp_ci) for the health step to resolve.

* ci(hrx): add MUL_MAT failure-rate loop (per-runner diagnostic)

Run the failing mul_mat_f16 config N times (MULMAT_LOOP_N, default 50) in one
job and print the fail/pass count + RUNNER_NAME, to answer: on a bad machine
does it fail every iteration or only some, and what's the rate per machine?
Advisory (continue-on-error) and placed BEFORE the gating correctness step so it
still runs on a failing machine. Correlate the count with RUNNER_NAME from the
Debug step.

* ci(hrx): move runner-identity dump into bench script

Replace the inline "Debug (runner identity)" block with a call to
runner-identity.sh (in the bench repo, alongside runner-health.sh). The step
now runs right after the bench checkout (earliest the script is available)
instead of first; matrix values are passed via MATRIX_* env since a script
can't read ${{ matrix.* }}. if: always() retained.

Requires runner-identity.sh on the bench BENCH_REF branch (committed on
AaronStGeorge/llamacpp_ci runner-health).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* cleanup

---------

Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Author: AaronStGeorge

.github/workflows/build-android.yml

@@ -16,6 +16,7 @@ on:
     ]
 
   pull_request:
+    branches-ignore: ['hrx-integration']
     types: [opened, synchronize, reopened]
     paths: [
       '.github/workflows/build-android.yml',

.github/workflows/build-apple.yml

@@ -19,6 +19,7 @@ on:
     ]
 
   pull_request:
+    branches-ignore: ['hrx-integration']
     types: [opened, synchronize, reopened]
     paths: [
       '.github/workflows/build-apple.yml',

.github/workflows/build-cann.yml

@@ -16,6 +16,7 @@ on:
     ]
 
   pull_request:
+    branches-ignore: ['hrx-integration']
     types: [opened, synchronize, reopened]
     paths: [
       '.github/workflows/build-cann.yml',

.github/workflows/build-hrx.yml

@@ -1,6 +1,9 @@
 name: HRX CI
 
 on:
+  push:
+    branches:
+      - hrx-integration
   pull_request:
     types: [opened, synchronize, reopened]
     branches:
@@ -16,40 +19,95 @@ concurrency:
 
 jobs:
   hrx-build:
-    name: ubuntu-latest
-    runs-on: ubuntu-latest
+    name: ${{ matrix.name }} (${{ matrix.gpu_target }})
+    strategy:
+      # Per-arch failures are independent; keep both signals.
+      fail-fast: false
+      matrix:
+        include:
+          - name: gfx1151_strix-halo
+            gpu_target: gfx1151
+            runs_on: linux-gfx1151-gpu-rocm
+            extra_rocm_artifacts: rocwmma_dev_gfx1151
+          - name: gfx1201_9070
+            gpu_target: gfx1201
+            runs_on: linux-gfx120X-gpu-rocm
+            extra_rocm_artifacts: rocwmma_dev_gfx120X-all
+    runs-on: ${{ matrix.runs_on }}
+    timeout-minutes: 10
+    defaults:
+      run:
+        # --noprofile skips loading ~/.bash_profile
+        # --norc      skips loading ~/.bashrc
+        # -exo pipefail:
+        #   -e exit immediately if any command fails
+        #   -x print each command before executing
+        #   -o pipefail ensure pipeline fails if any command in it fails
+        # {0} github actions placeholder for the command to run
+        shell: bash --noprofile --norc -exo pipefail {0}
+    container:
+      image: ghcr.io/rocm/no_rocm_image_ubuntu24_04@sha256:fba5f55a122dbb15925e98c51fe65bffe88c36e11ebb25b73daf2bea04202dc3
+      # --user 0:0: actions running inside the container need to write to some
+      #             files set up outside the container by the runner agent.  In
+      #             June 2026 some runner agents set those files up with
+      #             permissions for a "tester" user with UID/GID 1001, which
+      #             matches the user in no_rocm_image_ubuntu24_04, and some are
+      #             set up as root; accessing a file owned by root with user
+      #             "tester" gives an EACCES. Running as root is the common
+      #             denominator.
+      # --device kfd/dri: GPU access.
+      options: >-
+        --user 0:0
+        --device /dev/kfd
+        --device /dev/dri
     env:
-      HRX_WORK_DIR: ${{ github.workspace }}
+      # Write out of tree build/install to a non-bind-mounted directory. In June
+      # 2026 a debug test showed some runners had assets dated from 6 months ago
+      # in the bind-mounted github workspace, it seems like nothing cleans it up
+      # currently. Writing to an non-bind-mounted directory ensures that
+      # rocm-artifacts/builds are cleaned up when docker container is removed.
+      HRX_WORK_DIR: /work
       # Public location for bench tooling (rsuderman/llamacpp_ci and fork
       # AaronStGeorge/llamacpp_ci)
       #TODO: switch to ROCm/llamacpp-hrx-bench once it is open sourced.
       BENCH_REPOSITORY: AaronStGeorge/llamacpp_ci
       BENCH_REF: 'main'
+      # actions/checkout@v6 does auto-cleanup, an in-tree build would be
+      # auto-cleaned as well.
       BENCH_DIR: ${{ github.workspace }}/bench
       LLAMA_SRC_DIR: ${{ github.workspace }}/llama-src
-      HRX_ARTIFACT_SET: 'core-with-upstream-hip'
+      LLAMA_BUILD_DIR: /work/llama-build
+      HRX_EXTRA_ROCM_ARTIFACTS: ${{ matrix.extra_rocm_artifacts }}
+      GGML_HRX_AMDGPU_TARGETS: ${{ matrix.gpu_target }}
       CCACHE_COMPILERCHECK: content
 
     steps:
-      - name: Checkout llama.cpp (under test)
-        uses: actions/checkout@v6
-        with:
-          path: llama-src
-
       - name: Checkout bench tooling
         uses: actions/checkout@v6
         with:
           repository: ${{ env.BENCH_REPOSITORY }}
           ref: ${{ env.BENCH_REF }}
           path: bench
 
+      - name: Runner info
+        env:
+          MATRIX_NAME: ${{ matrix.name }}
+          MATRIX_GPU_TARGET: ${{ matrix.gpu_target }}
+          MATRIX_RUNS_ON: ${{ matrix.runs_on }}
+        run: "${BENCH_DIR}/scripts/hrx/runner-info.sh"
+
+      - name: Checkout llama.cpp (under test)
+        uses: actions/checkout@v6
+        with:
+          path: llama-src
+
       - name: Install ROCm build dependencies
         run: "${BENCH_DIR}/scripts/hrx/install-rocm-deps.sh"
 
       - name: ccache
         uses: ggml-org/ccache-action@v1.2.21
         with:
-          key: ubuntu-latest-${{ env.HRX_ARTIFACT_SET }}
+          key: ${{ matrix.gpu_target }}
           evict-old-files: 1d
           save: ${{ github.event_name == 'pull_request' }}
 
@@ -68,28 +126,30 @@ jobs:
       - name: Build llama.cpp with HRX
         run: "${BENCH_DIR}/scripts/hrx/build-llama-hrx.sh"
 
-      - name: Run sample MUL_MAT correctness config on CPU
+      - name: Run sample MUL_MAT correctness config on HRX
         run: |
           . "${BENCH_DIR}/scripts/hrx/env.sh"
+          . "${BENCH_DIR}/scripts/hrx/runtime-env.sh"
           "${BENCH_DIR}/tools/run-op-test.py" \
             --test-backend-ops "${LLAMA_BUILD_DIR}/bin/test-backend-ops" \
             --test-file "${BENCH_DIR}/benchmark-configs/test/mul_mat_f16.txt" \
             --op MUL_MAT \
-            --backend CPU \
-            --output benchmark-results/sample-mul-mat-f16-cpu-test.jsonl
+            --backend HRX0 \
+            --output benchmark-results/sample-mul-mat-f16-hrx-test.jsonl
 
-      - name: Run sample MUL_MAT benchmark config on CPU
+      - name: Run sample MUL_MAT benchmark config on HRX
         run: |
           . "${BENCH_DIR}/scripts/hrx/env.sh"
+          . "${BENCH_DIR}/scripts/hrx/runtime-env.sh"
           "${BENCH_DIR}/tools/run-op-perf.py" \
             --test-backend-ops "${LLAMA_BUILD_DIR}/bin/test-backend-ops" \
             --test-file "${BENCH_DIR}/benchmark-configs/test/mul_mat_f16.txt" \
             --op MUL_MAT \
-            --backend CPU \
-            --output benchmark-results/sample-mul-mat-f16-cpu-perf.jsonl
+            --backend HRX0 \
+            --output benchmark-results/sample-mul-mat-f16-hrx-perf.jsonl
 
-      - name: Upload CPU benchmark results
+      - name: Upload benchmark results
         uses: actions/upload-artifact@v5
         with:
-          name: cpu-benchmark-results
+          name: benchmark-results-${{ matrix.name }}
           path: benchmark-results/

.github/workflows/build-riscv.yml

@@ -16,6 +16,7 @@ on:
     ]
 
   pull_request:
+    branches-ignore: ['hrx-integration']
     types: [opened, synchronize, reopened]
     paths: [
       '.github/workflows/build-riscv.yml',

.github/workflows/build-self-hosted.yml

@@ -24,6 +24,7 @@ on:
     ]
 
   pull_request:
+    branches-ignore: ['hrx-integration']
     types: [opened, synchronize, reopened]
     paths: [
       '.github/workflows/build-self-hosted.yml',

.github/workflows/build-vulkan.yml

@@ -18,6 +18,7 @@ on:
     ]
 
   pull_request:
+    branches-ignore: ['hrx-integration']
     types: [opened, synchronize, reopened]
     paths: [
       '.github/workflows/build-vulkan.yml',

.github/workflows/build.yml

@@ -25,6 +25,7 @@ on:
     ]
 
   pull_request:
+    branches-ignore: ['hrx-integration']
     types: [opened, synchronize, reopened]
     paths: [
       '.github/workflows/build.yml',

.github/workflows/check-vendor.yml

@@ -11,6 +11,7 @@ on:
     ]
 
   pull_request:
+    branches-ignore: ['hrx-integration']
     types: [opened, synchronize, reopened]
     paths: [
       'vendor/**',

.github/workflows/copilot-setup-steps.yml

@@ -5,9 +5,11 @@ name: "Copilot Setup Steps"
 on:
   workflow_dispatch:
   push:
+    branches-ignore: ['hrx-integration']
     paths:
       - .github/workflows/copilot-setup-steps.yml
   pull_request:
+    branches-ignore: ['hrx-integration']
     paths:
       - .github/workflows/copilot-setup-steps.yml