AMD GPU support via HIP (single-source backend) + CUDA_EXTEND VRAM-only expert tier

[noobdev-ph]

Summary

Two additions, both opt-in, CPU default untouched:

1. AMD GPU support via HIP/ROCm (make HIP=1, make hip-test) — the existing CUDA expert backend compiles unchanged for AMD GPUs through a 29-line backend_gpu_compat.h mapping header, following the same one-shim-header pattern compat.h uses for the Windows port. One source, two vendors: nvcc path is byte-identical to today; hipcc maps the exact CUDA runtime surface backend_cuda.cu uses onto HIP 1:1 (the kernels use no vendor-specific intrinsics). First AMD datapoint below.

2. CUDA_EXTEND=1: the VRAM tier holds experts beyond the RAM pin — today the VRAM tier promotes experts already resident in the RAM pin (a compute mirror; ideal for matmul-bound machines). On disk-bound machines it adds no cache coverage. With CUDA_EXTEND=1, the VRAM budget pins the next experts in the frequency ranking: loaded once at startup, uploaded, host slab freed — extra pinned capacity at zero RAM cost (RSS measured flat). Legacy mirror stays the default.

Safety for VRAM-only slots: cached device tensors are callable without live host pointers (upload check-order fix, pinned by a new test assertion); if the GPU ever refuses a VRAM-only expert mid-run, the matmul flags it and emits zeros, the expert loop repairs in place with one disk reload + CPU recompute (not even that token is wrong), and the pin lookup skips the dead slot from then on; REPIN won't swap into slab-less slots.

Validation

• make -C c check (CPU build + all C/Python tests pass; default path untouched)
• make -C c hip-test HIP_ARCH=gfx1201 — q8/q4/q2/f32 kernel correctness on AMD, including the new cached-tensor-without-host-pointers assertion
• Performance: RX 9070 XT 16 GB (gfx1201, ROCm 7.2.4), Ryzen 7 5700G (PCIe Gen3: GPU x16 = 13.4 GB/s measured, NVMe x4 = 3.5 GB/s), 60 GB RAM, GLM-5.2 int4 + int8 MTP, --ram 40 --topp 0.7 --temp 0 --ngen 32, greedy fixed prompt, 10 runs per config (full per-run logs in the linked issue thread; same methodology as [Performance]: Ryzen 7 5700G + Gen5 NVMe capped at Gen3 — int8 MTP confirmed at 40%, learning-cache curve measured over 20 controlled runs #93):

config	median tok/s	expert hit	MTP acceptance
CPU only	0.22	31–38%	40% steady
HIP mirror (CUDA_EXPERT_GB=12)	0.22	33%	33%
HIP CUDA_EXTEND=1	0.24 (best 0.30)	44–49%	23–46%, mean 31%

Honest reading: the extension structurally lifts hit-rate ~10 points (620 → 1,254 pinned experts) but GPU float numerics shift the greedy trajectory, dropping MTP acceptance from the CPU's steady 40% to ~31% mean — the extra forwards absorb most of the caching win. When acceptance happened to recover (run 9: 46%), throughput hit 0.30 tok/s, +36% over CPU median. Follow-up that would make ~0.30 the steady state: a numerics-matched GPU kernel (int8 activation quant + integer dot, mirroring the CPU IDOT path). Happy to iterate on that in a second PR.

The mirror-mode neutrality on this box also confirms the README's design note: streaming stays CPU-side, and the mirror tier targets matmul-bound machines — this machine is disk-bound (51%), which is exactly what CUDA_EXTEND=1 addresses.

Compatibility

• The default CPU build remains dependency-free (HIP=1 is opt-in, mutually exclusive with CUDA=1, Linux-only, errors early)
• No model files, generated binaries, or benchmark artifacts are included

Context: hardware and baseline measurements are the same machine as #93.

[rajpratham1]

Thanks for the contribution—this is an ambitious feature that adds HIP support while also introducing VRAM-only expert caching and build system changes.

Before this can be merged, I think it would benefit from additional validation:

• Please add CI or documented test results for both CUDA and HIP builds to demonstrate that the shared backend works correctly across both toolchains.
• The new VRAM-only expert path introduces significant changes to memory ownership and fallback behavior. It would be helpful to include tests covering upload failures, CPU fallback, and repeated execution after host memory has been released.
• Since this changes the build system substantially (CUDA/HIP selection, new compatibility layer, Makefile updates), adding documentation describing the supported environments and expected build commands would make the feature easier to adopt and maintain.

The direction looks promising, but I'd like to see these validations before approving such a large cross-platform feature.

[cb88]

I built the code directly from noobdev-ph's repo, and it doesn't detect my MI50s

hip-test seems to pass but with some warnings.

Edit also tried pulling the PR just incase I messed up same result

model  144 shards · 378.8 GB
  disk   backing store · 673.5 GB free
  RAM    115.7 GB budget · 10.7 GB dense · 6.1 GB runtime · cap 68/layer
  VRAM   no NVIDIA device detected · CPU path

[cb88@MZ31AR0 c]$ make hip-test
make cuda-test HIP=1
make[1]: Entering directory '/home/cb88/colibri2/colibri/c'
/opt/rocm/bin/hipcc -O3 -std=c++17 -x hip --offload-arch=native -Wall -Wextra backend_cuda.cu tests/test_backend_cuda.cu -o backend_cuda_test
backend_cuda.cu:63:71: warning: unused parameter 'S' [-Wunused-parameter]
   63 |                                     const float *scales, int fmt, int S, int I, int O,
      |                                                                       ^
backend_cuda.cu:77:25: warning: comparison of integers of different signs: 'unsigned int' and 'int' [-Wsign-compare]
   77 |         if (threadIdx.x < n) partial[threadIdx.x] += partial[threadIdx.x + n];
      |             ~~~~~~~~~~~ ^ ~
backend_cuda.cu:86:15: warning: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Wunused-value]
   86 |     if (*ptr) cudaFree(*ptr);
      |               ^~~~~~~~~~~~~~
./backend_gpu_compat.h:21:34: note: expanded from macro 'cudaFree'
   21 | #define cudaFree                 hipFree
      |                                  ^
backend_cuda.cu:128:21: warning: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Wunused-value]
  128 |         if (ctx->x) cudaFree(ctx->x);
      |                     ^~~~~~~~~~~~~~~~
./backend_gpu_compat.h:21:34: note: expanded from macro 'cudaFree'
   21 | #define cudaFree                 hipFree
      |                                  ^
backend_cuda.cu:129:21: warning: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Wunused-value]
  129 |         if (ctx->y) cudaFree(ctx->y);
      |                     ^~~~~~~~~~~~~~~~
./backend_gpu_compat.h:21:34: note: expanded from macro 'cudaFree'
   21 | #define cudaFree                 hipFree
      |                                  ^
backend_cuda.cu:222:26: warning: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Wunused-value]
  222 |     if (tensor->weights) cudaFree(tensor->weights);
      |                          ^~~~~~~~~~~~~~~~~~~~~~~~~
./backend_gpu_compat.h:21:34: note: expanded from macro 'cudaFree'
   21 | #define cudaFree                 hipFree
      |                                  ^
backend_cuda.cu:223:25: warning: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Wunused-value]
  223 |     if (tensor->scales) cudaFree(tensor->scales);
      |                         ^~~~~~~~~~~~~~~~~~~~~~~~
./backend_gpu_compat.h:21:34: note: expanded from macro 'cudaFree'
   21 | #define cudaFree                 hipFree
      |                                  ^
7 warnings generated when compiling for gfx906.
backend_cuda.cu:63:71: warning: unused parameter 'S' [-Wunused-parameter]
   63 |                                     const float *scales, int fmt, int S, int I, int O,
      |                                                                       ^
backend_cuda.cu:77:25: warning: comparison of integers of different signs: 'unsigned int' and 'int' [-Wsign-compare]
   77 |         if (threadIdx.x < n) partial[threadIdx.x] += partial[threadIdx.x + n];
      |             ~~~~~~~~~~~ ^ ~
backend_cuda.cu:86:15: warning: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Wunused-value]
   86 |     if (*ptr) cudaFree(*ptr);
      |               ^~~~~~~~~~~~~~
./backend_gpu_compat.h:21:34: note: expanded from macro 'cudaFree'
   21 | #define cudaFree                 hipFree
      |                                  ^
backend_cuda.cu:128:21: warning: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Wunused-value]
  128 |         if (ctx->x) cudaFree(ctx->x);
      |                     ^~~~~~~~~~~~~~~~
./backend_gpu_compat.h:21:34: note: expanded from macro 'cudaFree'
   21 | #define cudaFree                 hipFree
      |                                  ^
backend_cuda.cu:129:21: warning: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Wunused-value]
  129 |         if (ctx->y) cudaFree(ctx->y);
      |                     ^~~~~~~~~~~~~~~~
./backend_gpu_compat.h:21:34: note: expanded from macro 'cudaFree'
   21 | #define cudaFree                 hipFree
      |                                  ^
backend_cuda.cu:222:26: warning: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Wunused-value]
  222 |     if (tensor->weights) cudaFree(tensor->weights);
      |                          ^~~~~~~~~~~~~~~~~~~~~~~~~
./backend_gpu_compat.h:21:34: note: expanded from macro 'cudaFree'
   21 | #define cudaFree                 hipFree
      |                                  ^
backend_cuda.cu:223:25: warning: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Wunused-value]
  223 |     if (tensor->scales) cudaFree(tensor->scales);
      |                         ^~~~~~~~~~~~~~~~~~~~~~~~
./backend_gpu_compat.h:21:34: note: expanded from macro 'cudaFree'
   21 | #define cudaFree                 hipFree
      |                                  ^
7 warnings generated when compiling for host.
./backend_cuda_test
[CUDA] device 0: AMD Radeon Graphics, 34.3 GB VRAM, sm_90
cuda backend: q8/q4/q2/f32 correctness ok on 1 device(s)
make[1]: Leaving directory '/home/cb88/colibri2/colibri/c'
[cb88@MZ31AR0 c]$ [CUDA] device 0: AMD Radeon Graphics, 34.3 GB VRAM, sm_90
cuda backend: q8/q4/q2/f32 correctness ok on 1 device(s)
make[1]: Leaving directory '/home/cb88/colibri2/colibri/c'

[cb88]

So I pointed qwen 3.6 27B at this problem and it added a rocm smi parser function to resource_plan.py ... so not sure how it works at all on your system without that! my local copy detects both GPUS and functions after that.

Also for some reason usage of the GPUS does not show up in rocm-smi or nvtop (I forgot to pass CUDA_EXPERT_GB)

colibri doctor · /home/cb88/colibri/glm52_i4
[  ok] model.path         model directory is readable
[  ok] model.config       config.json is valid
[  ok] model.tokenizer    tokenizer.json found
[  ok] storage.persistence model directory can store usage and KV state
[  ok] engine.binary      engine executable is ready
[warn] accelerator.cuda   NVIDIA GPU detected but the engine is CPU-only
[  ok] model.shards       safetensors headers are valid
[  ok] storage.disk       model backing store is available
[  ok] memory.ram         RAM budget is viable
[  ok] placement.plan     tier placement has no warnings

model  144 shards · 378.8 GB
disk   backing store · 673.2 GB free
RAM    115.7 GB budget · 10.7 GB dense · 6.1 GB runtime · cap 68/layer
VRAM   64.7 GB hot tier · ~3418 experts · 0:AMD Instinct MI60 / MI50, 1:AMD Radeon Graphics

def _discover_amd_gpus():
    """Discover AMD GPUs via rocm-smi (ROCm).

    rocm-smi outputs {"card0": {...}, "card1": {...}} with keys like "GPU Device Name" and
    "VRAM Total Memory (B)".  We need two invocations because --showid and --showmeminfo
    vram don't combine in a single JSON call."""
    # -- showid: device name
    try:
        result = subprocess.run(
            ["rocm-smi", "--showid", "--json"],
            text=True, capture_output=True, check=True, timeout=5)
        ids = json.loads(result.stdout)
    except (OSError, subprocess.SubprocessError, json.JSONDecodeError, ValueError):
        return []

    # -- showmeminfo vram: total / used in bytes
    try:
        result = subprocess.run(
            ["rocm-smi", "--showmeminfo", "vram", "--json"],
            text=True, capture_output=True, check=True, timeout=5)
        mem = json.loads(result.stdout)
    except (OSError, subprocess.SubprocessError, json.JSONDecodeError, ValueError):
        return []

    devices = []
    # Both dicts share the same keys: "card0", "card1", ...
    for key in ids:
        card_id = ids[key]
        card_mem = mem.get(key, {})
        name = card_id.get("Device Name", "AMD GPU").strip()
        total = int(card_mem.get("VRAM Total Memory (B)", 0))
        used = int(card_mem.get("VRAM Total Used Memory (B)", 0))
        devices.append({"index": int(key.replace("card", "")),
                        "name": name,
                        "total_bytes": total,
                        "free_bytes": total - used})
    return devices


def discover_gpus():
    """Discover GPUs from NVIDIA (nvidia-smi) and AMD (rocm-smi) backends."""
    devices = _discover_nvidia_gpus()
    devices += _discover_amd_gpus()
    return devices

[noobdev-ph]

Thanks for the thorough review — all three points were fair, and addressing them made the branch better. Everything below is pushed (36a8c3c).

1. CI + documented results for both toolchains

Added .github/workflows/gpu-build.yml: on every push/PR it compiles the shared backend, the test binary, and the full engine link under nvcc (nvidia/cuda container, sm_80) and hipcc (rocm/dev container, gfx1100), plus the standard CPU make check. A new gpu-compile Makefile target supports toolchain-only runners. First run on my fork: all green — https://github.com/noobdev-ph/colibri/actions/runs/29471442499

Honest scoping: hosted runners have no GPUs, so CI is compile-level. Runtime execution is covered by make hip-test on my RX 9070 XT (documented matrix in GPU_BACKENDS.md, 5/5 stable repeat runs, including concurrently with a full 12 GB expert tier resident). I don't own NVIDIA hardware — since the nvcc path is a pass-through include (byte-identical preprocessed source to the pre-HIP tree), a single make cuda-test run on your hardware would close that gap; the test source is vendor-neutral.

2. Tests for the VRAM-only path

tests/test_backend_cuda.cu now covers: upload failures (invalid device, invalid format, missing scales, missing host data, ~16 TB allocation) with stats-integrity assertions; 64× sustained matmul reuse after host pointers are gone; upload-from-a-freed-temporary (the exact VRAM-only slot lifecycle: load slab → upload → scribble → free → reuse); and a COLI_GPU_FAIL_AFTER=N fault-injection hook (env-gated, zero effect unset).

You were right to ask — the new failure tests immediately caught a real bug: a failed allocation left a sticky runtime error that poisoned the next healthy launch's cudaGetLastError() check, so one failed upload could false-fail subsequent valid GPU calls. Fixed in cuda_ok() (error consumed on the failure path) and now regression-covered.

For the engine-level repair/fallback behavior, the fault hook enables an end-to-end validation with a strong invariant: with CUDA_EXTEND=1 and COLI_GPU_FAIL_AFTER=0, every GPU matmul fails, so every VRAM-only expert must take the repair path (reload from disk, recompute on CPU, skip the slot thereafter) — and the greedy output must be byte-identical to a pure CPU run. Measured on the real 744B model: 1,689 tensors degraded, output cmp-identical to the CPU reference. Mid-run injection (N=2000, 1,423 tensors degraded mid-generation) completes cleanly with per-tensor notices. Procedure documented so anyone can reproduce it.

3. Documentation

Added GPU_BACKENDS.md: supported environments and build matrix for both vendors, runtime knobs (CUDA_EXTEND mirror-vs-extension semantics included), the validation procedures above, the hardware test matrix, and known-behavior notes (GPU/CPU numerics divergence and its MTP-acceptance effect, serial extension startup, coli plan not yet modeling the extension tier).

Happy to iterate further — and if you'd like the workflow to live elsewhere or trigger differently in your repo, easy to adjust.

[noobdev-ph]

Synced the branch with 54cfe56 (9be216e + follow-ups; CI green: https://github.com/noobdev-ph/colibri/actions/runs/29523184492). The reconciliation simplified this PR considerably, and the validation run on the merged tree produced findings you'll want to see.

CUDA_EXTEND is retired — your design won. PIN=auto + PIN_FILL + CUDA_RELEASE_HOST achieve the same VRAM capacity extension with much deeper integration, and expert_host_ensure supersedes the in-loop repair this branch carried. c/glm.c in this PR is now byte-identical to upstream main — the diff is down to the backend, compat header, Makefile, tests, CI, and docs.

One functional addition worth your attention: WMMA gating for HIP. Your new tensor-core kernels guard their bodies with __CUDA_ARCH__ >= 700, but the host dispatch checks compute_major >= 7 at runtime — and gfx GPUs report compute_major = 12, so under HIP the dispatch would select empty kernel bodies and produce garbage silently. The compat header now exports COLI_GPU_HAS_WMMA (1 under nvcc, 0 under hipcc) and the two dispatch sites honor it. rocWMMA matrix-core support would be the follow-up that lifts this.

Validation on the merged tree (RX 9070 XT, ROCm 7.2.4, GLM-5.2 int4; the fault hook now gates all 19 GPU compute entry points):

• make check and hip-test green, including your new tensor_update test (ordered after the reuse tests in the merged test file, since it mutates the shared tensor).
• Total-GPU-failure (COLI_GPU_FAIL_AFTER=0) with CUDA_RELEASE_HOST=1: no crash, coherent output, every host-released expert rematerialized via expert_host_ensure — your repair design holds under a stress mode it presumably never saw.
• Determinism: verified byte-identical across repeat runs given a frozen .coli_usage (across unfrozen runs, pin membership evolves by design and, in GPU mode, feeds grouped-path selection — worth knowing when comparing runs).
• Numerics identity limit, documented: a GPU-enabled run under total failure is not byte-identical to a pure-CPU run — the fallbacks select different kernel shapes/paths, consistent with your Greedy decoding is not reproducible: MTP (3/5 prompts) and the CUDA expert tier (2/5) both change the output #100. Stable and coherent on both sides; GPU_BACKENDS.md states the calibrated claims.
• Observation: under total GPU failure, MTP speculation disengages entirely (0 proposals) — graceful and lossless, but a failing GPU also silently costs the speculation speedup. Might deserve a log line.

Two CI/housekeeping findings that apply upstream too: (1) .gitignore misses tests/test_kv_alloc (and a few other new test binaries) — an accidentally committed binary shadows recompilation via mtime and fails on other machines' glibc; (2) rocm/dev:latest no longer fits hosted-runner disks — the workflow pins 6.2 on the 24.04 base and validates HIP via gpu-compile (backend + test binary through the hipcc driver); the gcc-side engine link doesn't work with containerized ROCm libs on either Ubuntu base and is documented as hardware-validated instead.

Performance on the merged tree: your CUDA_RELEASE_HOST through the HIP port measures 0.32 tok/s at 61% expert hit (vs 0.22–0.27 CPU on the same machine) — better than this branch's original extension mode, and the best result this machine has produced. Happy to run any additional configuration you'd like numbers for.

[JustVugg]

Could you rewrite this against current dev? I'm keeping it open rather than closing it, but it can't land in its present shape.

It's 24 commits behind dev. On glm.c, which changes several times a day, that isn't a rebase — it's a rewrite, and doing it for you would mean rewriting your design decisions without understanding why you made them. So I'd rather ask.

Yours is the closest to landing — only 24 commits behind. An AMD/HIP backend has the same problem as #84 though: nobody here can run it, so the CI can only prove it compiles. Worth thinking about what would let it be verified rather than trusted.

What changed under you, which makes a rewrite cheaper than it sounds:

• CI now exists (ci: make check on ubuntu / windows (MSYS2 UCRT64) / macos — and it already caught a Windows test bug (#140) #143 + ci: GitHub Actions — engine, web, Python, CUDA syntax (free for public repos) #144, merged today): make check runs on ubuntu/windows/macos and the engine/web/python/CUDA-syntax jobs on Linux, for every PR. You'll get a verdict in ~2 minutes instead of waiting on a maintainer who can't test your platform. It caught a bug in my own test on its first Windows run.
• All PRs now target dev, not main. main is locked to admins.
• The backlog is being cleared (31 → 17 today). Your PR sat long enough to rot, and that's this repo's fault, not yours — two contributors independently rebuilt the same Windows CUDA loader this week because nobody triaged in time.

What would help it land fast: the smallest version that does one thing. A 300–700 line PR touching glm.c needs a maintainer who can run it, and for most platforms here that maintainer doesn't exist. A focused change with a test the CI can execute gets reviewed in an hour.

If you'd rather not, say so and I'll close it with thanks — no hard feelings either way. And if you think I've misjudged and it should go in as-is, push back: I've been wrong twice today already and both times a contributor caught it.

[noobdev-ph]

Done as requested — rewritten against dev and split into the smallest pieces that each do one thing:

• GPU backend: failure-path hardening + tests (sticky-error fix, cached-tensor contract, fault-injection hook) #338 — vendor-neutral backend hardening (sticky-error bugfix + tests + fault-injection hook). Everything in it runs on your NVIDIA hardware: fully verifiable, nothing trusted.
• AMD GPU support: HIP/ROCm via single-source backend_gpu_compat.h (+ WMMA dispatch gate) #339 — the HIP/ROCm port itself (5 files, +218; glm.c untouched — the sync retired this branch's engine changes in favor of your PIN=auto/CUDA_RELEASE_HOST, which measure better than what they replaced: 0.33 tok/s vs 0.26 on the same card). The verifiability question is addressed head-on in its description, including what I commit to maintaining.

Closing this one in favor of those two. Thanks for the direct feedback — the 24-commit rot was real, and the split genuinely improved both halves. The full history (benchmarks, the WMMA silent-garbage finding, validation methodology) lives in this thread and both new PRs link back.