epic: AMD GPU (ROCm) backend on Linux via mlxcelverse #1801

Summary

Program of work to run mlxcel model inference with GPU acceleration on Linux hosts with AMD GPUs (ROCm/HIP). The first target is RDNA 3.5 (gfx1151, Strix Halo), where a feasibility spike already produced correct results and usable 4-bit decode speed.

The approach is to vendor an existing ROCm backend for MLX into mlxcel as a set of source overlays and apply it on top of the same pinned ml-explore/mlx commit that the Metal and CUDA builds use. It becomes the ROCm part of mlxcelverse, the name for mlxcel's MLX-side layer (see Design decisions). The ROCm overlay is copied into the MLX tree only when building with a new rocm cargo feature, so Apple Silicon and CUDA builds are untouched.

The ROCm backend comes from the rocm-support branch of NripeshN/mlx (MIT), which is the head of the upstream draft ml-explore/mlx#2300 and the only actively developed ROCm line for MLX today (tracking request: ml-explore/mlx#2556). Downstream projects such as lemonade-sdk/lemon-mlx-engine already build against it.

Why an in-repo overlay and not a separate MLX fork repository

The fork was measured against its upstream merge base (39886de4 to 75915908):

mlxcel's C++ bridge is written against the API of its pinned MLX commit, so any ROCm source has to follow that pin no matter where it lives. A separate repository would add a second pin, an ancestry rule between the two, and changes to both pin parsers (they require ml-explore/mlx in the repository URL), while the per-bump work stays the same. Vendoring keeps one pin and lets a pin bump fix the ROCm side in the same PR, using the overlay discipline mlxcel already applies to its Metal and CUDA patches.

This is different from porting MLX's CUDA backend to HIP by hand inside the overlays, which would mean replacing CUTLASS, CCCL and cudnn-frontend (10k+ lines). That option is not pursued.

Design decisions

0. mlxcelverse. mlxcelverse names everything mlxcel builds on top of upstream MLX, in two kinds: (a) per-backend source overlays that replace or add MLX files (today src/lib/mlx-cpp/patches/ for Metal and CUDA and src/lib/mlx-cpp/patches-cuda/; this epic adds patches-rocm/), maintained by a 3-way merge and a line-by-line review on every pin bump; and (b) mlxcel's own kernels, fusions and extension functions built on MLX's public custom-kernel APIs (fast::metal_kernel, fast::cuda_kernel, and on ROCm fast::hip_kernel), today in src/lib/mlx-cpp/turbo/ and the bridge, which only need public-API compatibility across bumps. This epic stays scoped to the ROCm members: the overlay (build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802) and ROCm kernel ports (perf(rocm): allocator footprint and ROCm ports of mlxcel fused kernels #1814). Reorganizing the existing tree under the mlxcelverse name, with no change to any build output, is tracked separately in refactor(mlx-cpp): organize mlxcel's MLX-side layer as mlxcelverse (backend overlays and mlxcel kernels) #1816.
1. Layout. src/lib/mlx-cpp/patches-rocm/ mirrors the MLX tree: the whole mlx/backend/rocm/ directory plus 15 core files, all as whole-file overlays (no diff patches). Whole-file configure_file COPYONLY is idempotent across the reconfigures that build.rs triggers, which diff patches are not.
2. ROCm-only copy. The overlay is copied only when MLX_BUILD_ROCM is on, following the patches-cuda/ precedent, so Metal and CUDA builds never compile a ROCm-modified core file. MLX_BUILD_ROCM together with MLX_BUILD_CUDA is rejected at configure time.
3. Enable by cargo feature. --features rocm (root) forwards to mlxcel-core/rocm. Feature combinations get separate build-script output directories, so a ROCm-patched _deps/mlx-src is never reused by another build.
4. Single MLX pin. The ROCm build fetches the same ml-explore/mlx commit as every other build. The pin parsers do not change.
5. Provenance. patches-rocm/UPSTREAM records the source repository, branch, commit and license; NOTICE gets an entry; vendored files keep their original headers and never receive a Lablup header.
6. Quantization modes ROCm cannot run are converted where possible. Load-time policy converts unsupported modes to affine (for example NVFP4 to affine 4-bit, following the existing dense repack path) and rejects with a clear message only when conversion is impossible. The pre-Ampere CUDA load policy in src/models/sanitize.rs is the precedent.

Feasibility spike (measured)

Host: AMD Ryzen AI MAX+ 395 with Radeon 8060S (gfx1151, RDNA 3.5, 40 CUs), 96 GiB VRAM carve-out, Debian with kernel 6.18, ROCm 10.0.0 packages (HIP 7.15, AMD clang 23). GPU otherwise idle.

Step 1: the fork as-is (NripeshN/mlx@75915908, Python bindings, -DMLX_BUILD_ROCM=ON -DCMAKE_HIP_ARCHITECTURES=gfx1151). Builds in about 3 minutes.

• Op correctness: 42/42 checks pass against an f32 CPU reference (matmul f32/f16/bf16, softmax, sum, logsumexp, RMS/layer norm, RoPE, argmax, sort, quantized_matmul affine 4/8-bit group 32/64 GEMV and GEMM, SDPA causal prefill and decode with GQA).
• Decode-shaped GEMV (8192x8192, including per-call sync): fp16 885 us (~152 GB/s), q4 254 us (~148 GB/s), q8 469 us (~152 GB/s).
• mlx-lm benchmark -p 512 -g 128:

Step 2: the ROCm overlay on mlxcel's pin (ml-explore/mlx@81ba1c6a plus the fork's backend directory plus the 17 core files). 14 core files applied cleanly; 3 needed a merge (mlx/backend/common/compiled.cpp, mlx/fast_primitives.h, mlx/io/safetensors.cpp). Six API-drift fixes were needed in the ROCm sources because upstream moved on after the fork's June merge base:

1. Upstream a124ac09 (Propagate NaN in arg reductions ml-explore/mlx#4291) added a host-only mlx::core::isnan template that hides the device isnan overloads inside the ROCm namespace. 23 call sites now use ::isnan.
2. compiled_collapse_contiguous_dims returns a 4-tuple with negative_strides; negative strides force the large-index kernel, as on CUDA.
3. fast::CustomKernel keeps both upstream compile_options and the fork's output_input_aliases; aliases stay out of state() because export serializes it.
4. SDPA use_fallback gained force_fused (CUDA semantics: throw if forced and no fused kernel applies).
5. New upstream primitives without ROCm kernels get NO_GPU stubs: GatherQQMM, SearchSorted, fast::CrossEntropy (+VJP, with fallback).
6. Event::error() storage ( Propagate CPU errors to events ml-explore/mlx#3742). ROCm does not populate it yet (see Phase 1).

Result: 42/42 op checks, same GEMV bandwidth, and the same mlx-lm numbers (Qwen3-0.6B tg 224, Llama-3.1-8B tg 32.3, Qwen3-30B-A3B tg 58.7). The assembled overlay is 121 files (106 backend + 15 core); copying it onto a fresh 81ba1c6a checkout reproduces the trial tree exactly. mlx/backend/{metal,cuda}/custom_kernel.cpp from the fork are dropped because they are not compiled in a ROCm build.

Quantization mode coverage on the ROCm GPU:

Known mlxcel-side gaps (from code review of main, re-verified on 2026-09-29)

• Resolved (refactor(core): route custom kernels by GPU backend kind #1877, fix(rocm): refuse instead of aborting at the sampler entry points #2010, fix(rocm): guard the four remaining custom-kernel launchers #2018, refactor(core): choose custom-kernel ports through one helper #2026, refactor: finish the kernel port standardization #2029): Ten call sites choose kernels with use_cuda = !metal::is_available() (src/lib/mlxcel-core/cpp/mlx_cxx_kernels.cpp:166,1483,1991 and src/lib/mlx-cpp/turbo/{fused_rope_append,sampling,fused_norm,paged_attention,paged_attention_v2,paged_attention_v2_merge,sampling_rejection}.cpp). On ROCm they call fast::cuda_kernel, which throws "No CUDA back-end". Most have a graph fallback; fused MoE and BitNet bitlinear_matmul do not. Every launcher now resolves its kernel through select_kernel_port (src/lib/mlx-cpp/turbo/kernel_port.h), which refuses with a catchable error naming the fallback when a backend has no port, instead of calling fast::cuda_kernel on ROCm.
• Resolved (feat(rocm): add a GPU vendor concept and AMD device reporting #1883): Pre-load memory estimation on Linux reads host MemAvailable (src/execution/memory_estimate.rs), which misjudges a UMA carve-out (the spike host sees about 30 GiB of host RAM next to 96 GiB of VRAM). The estimate now takes the ROCm allocator's nonzero memory_limit() (comment at src/execution/memory_estimate.rs:756).
• Resolved (feat(rocm): add a GPU vendor concept and AMD device reporting #1883): Hardware detection (src/lib/mlxcel-core/src/hardware.rs) only knows Apple sysctl and CUDA. It now has GpuVendor, GpuBackendKind, gpu_backend_kind() and the device_architecture and device_memory_bytes fields, and src/lib/mlxcel-core/src/rocm_arch.rs enforces compiled-versus-device gfx target coverage.
• Resolved (perf(bench): ROCm support in the benchmark harness and a published gfx1151 baseline #1810, PR perf(bench): ROCm support in the benchmark harness and a gfx1151 baseline #2056): The bench harness tagged a non-NVIDIA Linux host as metal (detect_backend() at scripts/bench_decode.sh:297-305, which had no rocm case), and scripts/compare_bench_csv.py:107-110 hardcoded host and runtime names (m5max/m1ultra and pylm/metal), so a ROCm host's scan returned {}. PR perf(bench): ROCm support in the benchmark harness and a gfx1151 baseline #2056 added ROCm support to the harness and published the gfx1151 baseline at docs/benchmark_results/rocm-baseline-gfx1151-2026-09-30.md.
• Resolved (feat(rocm): vendor the ROCm backend into mlxcelverse, add rocm feature #1818): build.rs only watches ../mlx-cpp/patches and ../mlx-cpp/patches-cuda (src/lib/mlxcel-core/build.rs:241-242). It now also watches ../mlx-cpp/patches-rocm (src/lib/mlxcel-core/build.rs:280-282).
• Partly resolved (ci(rocm): build, link and generate on a self-hosted gfx1151 runner #1989, test(rocm): add a local verify-rocm gate with a shared build-and-generate smoke #2008): No AMD runner exists in any workflow. ci(rocm): build, link and generate on a self-hosted gfx1151 runner #1989 added a rocm-build job to .github/workflows/ci.yml, but it is parked behind vars.ROCM_CI_ENABLED because no self-hosted AMD runner is registered, and a rocm-ci-status job reports that state (ci(rocm): self-hosted gfx1151 runner that builds, links and smoke-tests ROCm changes #1811 stays open for the runner); test(rocm): add a local verify-rocm gate with a shared build-and-generate smoke #2008 added make verify-rocm as the local substitute, mirroring make verify plus scripts/ci/rocm_smoke.sh.

Kernel-port standard (#2026, #2029): every custom-kernel launcher resolves its kernel from a KernelPorts table through select_kernel_port, and scripts/ci/check_kernel_port_dispatch.py enforces four rules (no backend comparison, no hand-rolled refusal, no direct kernel-holder access, no Rust gate spelled metal_is_available() || cuda_is_available()) in make verify, make verify-rocm and an unconditional hosted CI job. Of the 19 port tables on main, only bitlinear_ports() (src/lib/mlxcel-core/cpp/mlx_cxx_kernels.cpp:240, #1862) has a ROCm port; the other 18 have .rocm = nullptr and are the remaining surface for #1814, where filling a slot also needs that kernel's support predicate to agree with the table. The reasoning is in TECHNICAL_REPORTS/2026-kernel-port-dispatch-standard-20260929.en.md and TECHNICAL_REPORTS/2029-finish-kernel-port-standardization-20260929.en.md.

Non-goals

• Windows on ROCm (the fork's CMake assumes /opt/rocm and GCC libstdc++).
• Multi-GPU and distributed inference.
• CDNA (wave64, MI300) tuning. The build may target it, but no tuning or validation is in scope.
• Native NVFP4 kernels on ROCm. NVFP4 checkpoints are served through conversion.
• Any change to Metal or CUDA behavior. Every item must be a no-op there, verified by the existing gates.

Sub-issues

Phases map to execution waves. depends on edges override the phase default; items without an edge can run in parallel.

The ROCm gate is green: make verify-rocm on a fresh build of f9aefa39 passed 11,785 tests with 0 failed (378 ignored), plus the smoke generation. The four open items are external or tracking: #1811 waits for a spare AMD runner, #1873 for a GB10 node, #1813 for the maintainer's manual upstream submission, and #1814 tracks the kernel ports #2063 to #2069.

Phase 0: foundation

• build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802 build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature, landed in feat(rocm): vendor the ROCm backend into mlxcelverse, add rocm feature #1818

Phase 1: runtime correctness (parallel after Phase 0)

• refactor(core): route custom kernels by GPU backend kind instead of treating every non-Metal GPU as CUDA #1803 Route custom kernels by GPU backend kind instead of treating every non-Metal GPU as CUDA (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802), landed in refactor(core): route custom kernels by GPU backend kind #1877
• fix(rocm): surface ROCm GPU failures through Event::error instead of NaN or hangs #1804 Surface ROCm GPU failures through Event::error (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802), landed in fix(rocm): surface GPU failures through Event::error instead of NaN or hangs #2033
• feat(rocm): memory estimation, device info and hardware detection on AMD UMA hosts #1805 ROCm platform integration: memory estimation, device info, hardware detection on UMA hosts (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802), landed in feat(rocm): add a GPU vendor concept and AMD device reporting #1883
• feat(quant): backend quantization capability table and load-time convert-or-reject policy #1806 Backend quantization capability table and load-time convert-or-reject policy (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802), landed in feat(quant): add backend quant capability table and NVFP4 load policy #2030
• feat(rocm): implement FFT, Hadamard and SearchSorted on ROCm (audio, TurboQuant KV cache, stochastic speculative decoding abort today) #1825 feat(rocm): implement FFT, Hadamard and SearchSorted on ROCm (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802), landed in feat(rocm): implement Hadamard and SearchSorted, fix narrow index dtypes #1856, feat(rocm): implement FFT through hipFFT #1861
• fix(rocm): Gather reads narrow index dtypes as int64 and faults the GPU queue #1853 Gather reads narrow index dtypes as int64 and faults the GPU queue (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802), landed in feat(rocm): implement Hadamard and SearchSorted, fix narrow index dtypes #1856
• fix(rocm): the Gumbel-max and rejection sampler entry points have no ROCm arm and abort on metal_kernel #1885 Gumbel-max and rejection sampler entry points have no ROCm arm and abort on metal_kernel (depends on refactor(core): route custom kernels by GPU backend kind instead of treating every non-Metal GPU as CUDA #1803), landed in fix(rocm): refuse instead of aborting at the sampler entry points #2010
• fix(rocm): the WebUI catalog reports every architecture unsupported on ROCm #1886 WebUI catalog reports every architecture unsupported on ROCm (depends on feat(rocm): memory estimation, device info and hardware detection on AMD UMA hosts #1805), landed in fix(rocm): select the catalog backend column from the resolved backend #1921
• fix(rocm): CUDA and ROCm hosts share one MTP policy cache key #1887 CUDA and ROCm hosts share one MTP policy cache key (depends on feat(rocm): memory estimation, device info and hardware detection on AMD UMA hosts #1805), landed in fix(rocm): distinguish CUDA and ROCm in the MTP policy cache key #1933
• fix(rocm): get_launch_args clamps the grid without a grid-stride contract #1874 get_launch_args clamps the grid without a grid-stride contract (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802), landed in fix(rocm): remove get_launch_args, which capped the grid silently #2046
• fix(rocm): hipFFT blocks when too many plans are alive, root cause unknown #1876 hipFFT blocks when too many plans are alive, root cause unknown (depends on feat(rocm): implement FFT, Hadamard and SearchSorted on ROCm (audio, TurboQuant KV cache, stochastic speculative decoding abort today) #1825), landed in fix(rocm): set blocking-sync before first HIP queue, FFT cache to 128 #2049
• fix(rocm): SliceUpdate reduce path clamps the grid at 65535 blocks while slice_update_op_kernel is not grid-stride #2050 SliceUpdate reduce path clamps the grid at 65535 blocks while slice_update_op_kernel is not grid-stride (depends on fix(rocm): get_launch_args clamps the grid without a grid-stride contract #1874), landed in fix(rocm): make slice_update_op_kernel grid-stride #2053
• fix(rocm): validate MLX_ROCM_FFT_CACHE_SIZE (0 is UB, junk throws) #2051 Validate MLX_ROCM_FFT_CACHE_SIZE (0 is UB, junk throws) (depends on fix(rocm): hipFFT blocks when too many plans are alive, root cause unknown #1876), landed in fix(rocm): validate MLX_ROCM_FFT_CACHE_SIZE (0 is UB, junk throws) #2073
• fix(rocm): SliceUpdate and DynamicSliceUpdate donate a source array that is still referenced #2052 SliceUpdate and DynamicSliceUpdate donate a source array that is still referenced (depends on fix(rocm): SliceUpdate reduce path clamps the grid at 65535 blocks while slice_update_op_kernel is not grid-stride #2050), landed in fix(rocm): stop SliceUpdate donating a still-referenced source #2070
• fix(build/rocm): HIP objects do not rebuild when an included overlay header changes #2075 HIP objects do not rebuild when an included overlay header changes (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802), landed in fix(build/rocm): rebuild HIP objects when an included header changes #2076
• fix(rocm): bf16 prefill dense GEMM is not byte-identical to quantized_matmul on gfx1151 #2081 bf16 prefill dense GEMM is not byte-identical to quantized_matmul on gfx1151 (depends on feat(quant): backend quantization capability table and load-time convert-or-reject policy #1806), landed in fix(rocm): route large bf16 qmm to hipBLASLt and gate dense prefill #2085

Phase 2: quantization coverage

• feat(rocm/quant): mxfp8 end-to-end on ROCm, including FP8 block checkpoints and the MoE gather path #1807 mxfp8 end-to-end on ROCm: FP8 block checkpoints and the MoE gather path (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802, feat(quant): backend quantization capability table and load-time convert-or-reject policy #1806), landed in test(rocm): prove mxfp8 end to end, incl. FP8 block checkpoints #2071
• fix(rocm/quant): mxfp4 qmm hang and GPU quantize launch failure, with an affine 4-bit fallback #1808 mxfp4 on ROCm: fix the qmm hang and GPU quantize launch, with an affine fallback (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802, fix(rocm): surface ROCm GPU failures through Event::error instead of NaN or hangs #1804, feat(quant): backend quantization capability table and load-time convert-or-reject policy #1806), landed in fix(rocm/quant): commit mxfp4 regression tests and log the native route #2057
• fix(rocm/quant): mxfp4 quantized_matmul at 2880x2880 M=1 intermittently exceeds tolerance #2072 mxfp4 quantized_matmul at 2880x2880 M=1 intermittently exceeds tolerance (depends on fix(rocm/quant): mxfp4 qmm hang and GPU quantize launch failure, with an affine 4-bit fallback #1808), landed in fix(rocm): run CPU-stream BLAS single-threaded over fine-grained memory #2079

Phase 3: validation, tooling, operations

• test(rocm): correctness matrix against a Metal baseline and a verify-test-rocm gate #1809 Correctness matrix against a Metal baseline and a verify-test-rocm gate (depends on refactor(core): route custom kernels by GPU backend kind instead of treating every non-Metal GPU as CUDA #1803, feat(rocm): memory estimation, device info and hardware detection on AMD UMA hosts #1805, feat(quant): backend quantization capability table and load-time convert-or-reject policy #1806), landed in update(rocm): #1809 matrix rows, paged-attention skips, scan fix #2059 and test(rocm): complete the #1809 matrix rows with Metal M5 traces #2083, closed after the final gate
• perf(bench): ROCm support in the benchmark harness and a published gfx1151 baseline #1810 Benchmark harness support for ROCm and a published gfx1151 baseline (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802, feat(rocm): memory estimation, device info and hardware detection on AMD UMA hosts #1805), landed in perf(bench): ROCm support in the benchmark harness and a gfx1151 baseline #2056
• ci(rocm): self-hosted gfx1151 runner that builds, links and smoke-tests ROCm changes #1811 Self-hosted gfx1151 CI runner: build, link and smoke on ROCm-relevant changes (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802). Parked: no spare AMD runner.
• docs(rocm): installation guide and platform matrix for Linux with AMD GPUs #1812 Installation guide and platform matrix for Linux + ROCm (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802, feat(rocm): memory estimation, device info and hardware detection on AMD UMA hosts #1805), landed in docs(rocm): document the experimental AMD ROCm build and mlxcelverse #1819
• chore(rocm): mlxcelverse ROCm fork sync script, MLX pin-bump procedure, and upstreaming local fixes #1813 mlxcelverse ROCm maintenance: fork sync script, MLX pin-bump procedure, upstreaming local fixes (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802), landed in update(rocm): add mlxcelverse sync, drift checks and fork PR packages #2078 and closed under the 2026-10-06 mlxcelverse maintenance policy.
• test(cuda): run the deferred CUDA arm for the merged ROCm routing work on a GB10 node #1873 Run the deferred CUDA arm for the merged ROCm routing work on a GB10 node (depends on refactor(core): route custom kernels by GPU backend kind instead of treating every non-Metal GPU as CUDA #1803). Needs a GB10 node.
• chore(ci): kernel dtype-key checker cannot see launches that move into headers #1875 Kernel dtype-key checker cannot see launches that move into headers (depends on build(rocm): vendor the ROCm backend into mlxcelverse and add the rocm cargo feature #1802), landed in update(ci): make the kernel dtype-key checker's scope fail-closed #2055
• ci: stop waking the shared GB10 runner for changes that cannot affect what it builds #1992 Stop waking the shared GB10 runner for changes that cannot affect what it builds, landed in ci: stop waking the shared GB10 runner for changes it cannot build #1993

Phase 4: performance

• feat(rocm): port the BitNet bitlinear_matmul kernel to mlxcelverse #1862 Port the BitNet bitlinear_matmul kernel to mlxcelverse (depends on refactor(core): route custom kernels by GPU backend kind instead of treating every non-Metal GPU as CUDA #1803), landed in feat(rocm): port the BitNet bitlinear_matmul kernel #1870
• perf(rocm): allocator footprint and ROCm ports of mlxcel fused kernels #1814 ROCm performance: allocator footprint and ROCm ports of mlxcel fused kernels (depends on refactor(core): route custom kernels by GPU backend kind instead of treating every non-Metal GPU as CUDA #1803, test(rocm): correctness matrix against a Metal baseline and a verify-test-rocm gate #1809). All native sub-issues (perf(rocm): profile decode on gfx1151 and settle MLX ROCm gather_mm #2061 to perf(rocm): port Metal-only kernels (xielu, add3 norm, mamba1, relu2) #2069) closed (perf(rocm): profile decode on gfx1151 and settle MLX ROCm gather_mm #2061 in perf(rocm): profile gfx1151 decode per kernel and rank the #1814 ports #2086, perf(rocm): explain and bound allocator peak memory on UMA hosts #2062 in fix(rocm): bound in-flight batch memory and enforce the cache limit #2084, perf(rocm): port fused_add_rms_norm and fused_rope_qk_append to HIP #2063 in fix(tts): cast bf16 tables explicitly where the CUDA overlay demotes #2093, perf(rocm): port gumbel_max_sample and rejection_sample to HIP #2064 in fix(server): pool-backed KVCache::trim is a no-op, so tile-padded prefill leaves pad rows in the cache and default M5 serving returns wrong greedy output #2096, perf(rocm): port the fused MoE decode kernels (moe_gateup, moe_down) #2065 in update(rocm): port the Gumbel-max and rejection samplers to HIP #2100, perf(rocm): fix expert-batched gather_qmm for bf16 MoE prefill #2066 in update(rocm): port fused_add_rms_norm and fused_rope_qk_append to HIP #2107, perf(rocm): port ssm_update_kernel to HIP and fix ssm_kernel_available #2067 in perf(speculative): prompt lookup slows prose decoding 3-21% on GB10 #2091, perf(rocm): port paged attention (v1 decode, v2 partial, merge) to HIP #2068 in update(rocm): port paged attention (v1, v2 partial, merge) to HIP #2103, perf(rocm): port Metal-only kernels (xielu, add3 norm, mamba1, relu2) #2069 in fix: CPU-only Linux build fails to link kv_inplace_write #2108).

Acceptance criteria

• cargo build --release --features rocm produces mlxcel and mlxcel-server on a Linux AMD host, and the binaries run without LD_LIBRARY_PATH. Evidenced by scripts/ci/rocm_smoke.sh (test(rocm): add a local verify-rocm gate with a shared build-and-generate smoke #2008) for mlxcel generate and the 2026-10-15 gfx1151 verification (epic: AMD GPU (ROCm) backend on Linux via mlxcelverse #1801 comment 6193786952, PR fix(rocm): update Fence::wait and Fence::update signatures for MLX 243bdf09 #2317 at fe98588f), where both target/release/mlxcel and target/release/mlxcel-server carry DT_RPATH = [/opt/rocm/lib], resolve all ROCm shared libraries under env -u LD_LIBRARY_PATH ldd, and mlxcel-server serves /health and /v1/chat/completions on gfx1151 with LD_LIBRARY_PATH unset.
• Affine 4/8-bit checkpoints from each model family in the correctness matrix generate on the AMD GPU and meet the decided-position mismatch threshold against the Metal baseline. Complete under test(rocm): correctness matrix against a Metal baseline and a verify-test-rocm gate #1809: the twelve-pair dense and MoE matrix landed in test(rocm): add the gfx1151 correctness matrix and its gate #1826 (docs/benchmark_results/rocm-correctness-gfx1151-2026-09-12.md), and PRs update(rocm): #1809 matrix rows, paged-attention skips, scan fix #2059 and test(rocm): complete the #1809 matrix rows with Metal M5 traces #2083 added the remaining rows with Metal M5 traces (docs/benchmark_results/rocm-correctness-gfx1151-2026-09-30.md).
• mxfp8, mxfp4 and NVFP4 checkpoints either run (natively or through conversion) or are rejected at load with an actionable message. None produce NaN or hang. mxfp8 and mxfp4 run natively on gfx1151 (feat(rocm/quant): mxfp8 end-to-end on ROCm, including FP8 block checkpoints and the MoE gather path #1807, PR test(rocm): prove mxfp8 end to end, incl. FP8 block checkpoints #2071; fix(rocm/quant): mxfp4 qmm hang and GPU quantize launch failure, with an affine 4-bit fallback #1808, PR fix(rocm/quant): commit mxfp4 regression tests and log the native route #2057, with the intermittent 2880x2880 tolerance failure fixed in fix(rocm/quant): mxfp4 quantized_matmul at 2880x2880 M=1 intermittently exceeds tolerance #2072, PR fix(rocm): run CPU-stream BLAS single-threaded over fine-grained memory #2079), and NVFP4 converts to affine 4-bit at load through the capability table (feat(quant): backend quantization capability table and load-time convert-or-reject policy #1806, PR feat(quant): add backend quant capability table and NVFP4 load policy #2030).
• mlxcel-server serves /v1/chat/completions on the AMD GPU. Evidenced by f9ece5d9 (test(rocm): verify mlxcel-server chat completions on gfx1151 #1831): Meta-Llama-3.1-8B-Instruct-4bit and Qwen3-30B-A3B-4bit answered coherently on the Radeon 8060S, streaming and non-streaming agreed, and neither server log carried an error; the MoE run used MLXCEL_FUSED_MOE=0, which fix(rocm): guard the four remaining custom-kernel launchers #2018 and refactor(core): choose custom-kernel ports through one helper #2026 made unnecessary by turning the fused MoE abort into a catchable refusal that takes the graph path.
• A ROCm CI job builds, links and runs a smoke generation on every change to patches-rocm/, the MLX pin, build.rs or the mlx-cpp CMake. The rocm-build job exists (ci(rocm): build, link and generate on a self-hosted gfx1151 runner #1989) but is parked behind vars.ROCM_CI_ENABLED until a runner is registered (ci(rocm): self-hosted gfx1151 runner that builds, links and smoke-tests ROCm changes #1811).
• Installation docs, the platform matrix and a gfx1151 benchmark page are published. Docs and platform matrix landed in docs(rocm): document the experimental AMD ROCm build and mlxcelverse #1819 (docs(rocm): installation guide and platform matrix for Linux with AMD GPUs #1812); the benchmark page docs/benchmark_results/rocm-baseline-gfx1151-2026-09-30.md landed in perf(bench): ROCm support in the benchmark harness and a gfx1151 baseline #2056 (perf(bench): ROCm support in the benchmark harness and a published gfx1151 baseline #1810).
• Metal and CUDA builds and gates are unchanged. Not verifiable from the ROCm host; the CUDA arm is test(cuda): run the deferred CUDA arm for the merged ROCm routing work on a GB10 node #1873.

References

• ROCm backend source: https://github.com/NripeshN/mlx/tree/rocm-support (spike used 75915908dfe5028335d318b10340313744fd3a8d)
• Upstream draft PR: [Experiment] ROCm backend ml-explore/mlx#2300
• Upstream tracking issue: Add ROCm Support for AMD GPUs ml-explore/mlx#2556
• Reference engine on the same fork: https://github.com/lemonade-sdk/lemon-mlx-engine
• Overlay mechanism: src/lib/mlx-cpp/CMakeLists.txt (mlx_apply_source_overlays, lines 18-82; pin at line 122)
• Build glue to mirror: src/lib/mlxcel-core/build.rs (build_mlx 343, detect_cuda_arch 444, link_cuda 492)

Refresh log

2026-09-29

• Added fix(rocm): the Gumbel-max and rejection sampler entry points have no ROCm arm and abort on metal_kernel #1885, fix(rocm): the WebUI catalog reports every architecture unsupported on ROCm #1886, fix(rocm): CUDA and ROCm hosts share one MTP policy cache key #1887 and ci: stop waking the shared GB10 runner for changes that cannot affect what it builds #1992 as native sub-issues, and made ## Sub-issues list exactly the 24 native sub-issues (newly listed: fix(rocm): Gather reads narrow index dtypes as int64 and faults the GPU queue #1853, feat(rocm): port the BitNet bitlinear_matmul kernel to mlxcelverse #1862, test(cuda): run the deferred CUDA arm for the merged ROCm routing work on a GB10 node #1873, fix(rocm): get_launch_args clamps the grid without a grid-stride contract #1874, chore(ci): kernel dtype-key checker cannot see launches that move into headers #1875, fix(rocm): hipFFT blocks when too many plans are alive, root cause unknown #1876 and those four).
• Checkboxes now match issue state (11 closed, 13 open), each checked item names its landing PR, and feat(quant): backend quantization capability table and load-time convert-or-reject policy #1806 is marked as the only red item in make verify-rocm.
• Re-verified the known gaps against main at d8d34e2b: four resolved, the bench harness gap still open under perf(bench): ROCm support in the benchmark harness and a published gfx1151 baseline #1810 with corrected line references, the AMD runner partly resolved, and a paragraph on the kernel-port standard added.
• Acceptance criteria: checked the chat-completions criterion on f9ece5d9 (test(rocm): verify mlxcel-server chat completions on gfx1151 #1831) and stated the status of every unchecked criterion.
• refactor(mlx-cpp): organize mlxcel's MLX-side layer as mlxcelverse (backend overlays and mlxcel kernels) #1816, fix(bench): bench_mlxlm.py hardcodes a 128 GB memory budget #1858 and perf(moe): attribute the MoE decode gap at the M = 1 expert gather #1859 are deliberately not sub-issues: refactor(mlx-cpp): organize mlxcel's MLX-side layer as mlxcelverse (backend overlays and mlxcel kernels) #1816 touches the Metal and CUDA paths, and fix(bench): bench_mlxlm.py hardcodes a 128 GB memory budget #1858 and perf(moe): attribute the MoE decode gap at the M = 1 expert gather #1859 are not ROCm work.

2026-10-01

• Recorded the /epic-impl 1801 run of 2026-09-29 to 2026-10-01: checked fix(rocm): surface ROCm GPU failures through Event::error instead of NaN or hangs #1804, feat(quant): backend quantization capability table and load-time convert-or-reject policy #1806, feat(rocm/quant): mxfp8 end-to-end on ROCm, including FP8 block checkpoints and the MoE gather path #1807, fix(rocm/quant): mxfp4 qmm hang and GPU quantize launch failure, with an affine 4-bit fallback #1808, test(rocm): correctness matrix against a Metal baseline and a verify-test-rocm gate #1809, perf(bench): ROCm support in the benchmark harness and a published gfx1151 baseline #1810, chore(ci): kernel dtype-key checker cannot see launches that move into headers #1875 and fix(rocm): hipFFT blocks when too many plans are alive, root cause unknown #1876 with their landing PRs, named fix(rocm): remove get_launch_args, which capped the grid silently #2046 on the already-checked fix(rocm): get_launch_args clamps the grid without a grid-stride contract #1874, and left ci(rocm): self-hosted gfx1151 runner that builds, links and smoke-tests ROCm changes #1811, chore(rocm): mlxcelverse ROCm fork sync script, MLX pin-bump procedure, and upstreaming local fixes #1813, perf(rocm): allocator footprint and ROCm ports of mlxcel fused kernels #1814 and test(cuda): run the deferred CUDA arm for the merged ROCm routing work on a GB10 node #1873 open with the reason each waits.
• Listed the six sub-issues filed and closed during the run: fix(rocm): SliceUpdate reduce path clamps the grid at 65535 blocks while slice_update_op_kernel is not grid-stride #2050, fix(rocm): validate MLX_ROCM_FFT_CACHE_SIZE (0 is UB, junk throws) #2051, fix(rocm): SliceUpdate and DynamicSliceUpdate donate a source array that is still referenced #2052, fix(build/rocm): HIP objects do not rebuild when an included overlay header changes #2075 and fix(rocm): bf16 prefill dense GEMM is not byte-identical to quantized_matmul on gfx1151 #2081 in Phase 1 and fix(rocm/quant): mxfp4 quantized_matmul at 2880x2880 M=1 intermittently exceeds tolerance #2072 in Phase 2, with depends on edges to the issue each was found under. Noted the perf(rocm): allocator footprint and ROCm ports of mlxcel fused kernels #1814 split into perf(rocm): profile decode on gfx1151 and settle MLX ROCm gather_mm #2061 to perf(rocm): port Metal-only kernels (xielu, add3 norm, mamba1, relu2) #2069 under perf(rocm): allocator footprint and ROCm ports of mlxcel fused kernels #1814's item: perf(rocm): profile decode on gfx1151 and settle MLX ROCm gather_mm #2061 (PR perf(rocm): profile gfx1151 decode per kernel and rank the #1814 ports #2086) and perf(rocm): explain and bound allocator peak memory on UMA hosts #2062 (PR fix(rocm): bound in-flight batch memory and enforce the cache limit #2084) are closed, and the open ports are listed in the order perf(rocm): profile gfx1151 decode per kernel and rank the #1814 ports #2086 measured.
• Replaced the "feat(quant): backend quantization capability table and load-time convert-or-reject policy #1806 is the only red item" paragraph: make verify-rocm on a fresh build of f9aefa39 passed 11,785 tests with 0 failed. PR fix(test): repair gate failures left by #2034 and #2037 #2080 repaired the main gate failures left by feat(nemotron_parse): port Nemotron-Parse on the seq2seq worker #2034 and feat(nemotron_voicechat): load VoiceChat and run offline duplex inference #2037.
• Marked the bench harness gap resolved by PR perf(bench): ROCm support in the benchmark harness and a gfx1151 baseline #2056, and noted that PR update(rocm): add mlxcelverse sync, drift checks and fork PR packages #2078 delivered chore(rocm): mlxcelverse ROCm fork sync script, MLX pin-bump procedure, and upstreaming local fixes #1813's in-repo part while chore(rocm): mlxcelverse ROCm fork sync script, MLX pin-bump procedure, and upstreaming local fixes #1813 stays open (status:blocked) for manual upstream submission.
• Acceptance criteria: checked the correctness matrix (test(rocm): correctness matrix against a Metal baseline and a verify-test-rocm gate #1809), quantization coverage (feat(quant): backend quantization capability table and load-time convert-or-reject policy #1806, feat(rocm/quant): mxfp8 end-to-end on ROCm, including FP8 block checkpoints and the MoE gather path #1807, fix(rocm/quant): mxfp4 qmm hang and GPU quantize launch failure, with an affine 4-bit fallback #1808) and benchmark page (perf(bench): ROCm support in the benchmark harness and a gfx1151 baseline #2056) criteria. The CI-job criterion stays open on ci(rocm): self-hosted gfx1151 runner that builds, links and smoke-tests ROCm changes #1811 and the Metal/CUDA-unchanged criterion on test(cuda): run the deferred CUDA arm for the merged ROCm routing work on a GB10 node #1873.

2026-10-15

• Checked off chore(rocm): mlxcelverse ROCm fork sync script, MLX pin-bump procedure, and upstreaming local fixes #1813 and perf(rocm): allocator footprint and ROCm ports of mlxcel fused kernels #1814 (all native sub-issues perf(rocm): profile decode on gfx1151 and settle MLX ROCm gather_mm #2061 to perf(rocm): port Metal-only kernels (xielu, add3 norm, mamba1, relu2) #2069 closed).
• Merged PR fix(rocm): update Fence::wait and Fence::update signatures for MLX 243bdf09 #2317 (fe98588f, updating mlx/backend/rocm/fence.cpp for MLX 243bdf09), ran make verify-rocm-held on gfx1151 (12,384 passed, 0 failed, 411 ignored, rocm_smoke.sh 241.69 tok/s), verified target/release/mlxcel-server serves /health and /v1/chat/completions on gfx1151 with LD_LIBRARY_PATH unset (DT_RPATH = [/opt/rocm/lib]), and checked off the LD_LIBRARY_PATH-free release binary acceptance criterion.