tests: make the GPU test suite build under HIP on gfx1151

[huppiflupp]

Follow-up to the verification I ran on #1505. That PR gets make hip-test linking; these four items get it compiling. Each was hit and fixed on real gfx1151 hardware — Ryzen AI MAX+ 395 / Radeon 8060S, ROCm 7.1 / HIP 7.1.52802, rocwmma-devel installed, Nobara 44.

The four blockers

1. #include <cuda_runtime.h> before any compat mapping is in scope — six files:

tests/test_fp8_cuda.cu       tests/test_router_cuda.cu      tests/test_grouped_g4_cuda.cu
tests/test_mxfp4_cuda.cu     tests/bench_dsv4_deepgemm.cu   tests/bench_dsv4_mxfp8.cu

__HIP_PLATFORM_AMD__ is not defined yet at that point; __HIPCC__ is. Guarded on the latter.

2. cudaMallocManaged has no mapping in backend_gpu_compat.h. test_router_cuda and test_grouped_g4_cuda use it and both include backend_cuda.cu, so the header is in scope — the define was simply missing. hipMallocManaged exists in ROCm.

3. test_mxfp4_cuda.cu is a different case and the include guard alone is not enough for it. It links against a separately compiled backend_cuda.cu and includes only the public ABI header backend_cuda.h, which correctly does not pull the compat mapping — so that TU never sees it. It needs backend_gpu_compat.h itself under HIP. Its bare isinf() calls are additionally ambiguous against HIP's __DEVICE__ isinf(float)/isinf(double); qualified as std::isinf.

4. mxfp4_ref.o is built without -fPIC while GPUFLAGS builds with -fPIE. Under HIP the link then fails with R_X86_64_32S relocation cannot be used against local symbol. Added -fPIC to that one rule; harmless under CUDA.

Verified per test on gfx1151

test	result
test_fp8_cuda	builds, runs, OK — 0 mismatches on oracle and API
test_router_cuda	builds
test_grouped_g4_cuda	builds
test_mxfp4_cuda	builds, runs, test_mxfp4_cuda: ok — all five exponent cases, worst rel 8.63e-07

Worth noting for item 4: the Makefile comment calls MXFP4 "expected to fail" on some hosts. On gfx1151 it passes cleanly once it builds, so that expectation may be worth revisiting.

I did not touch the test link lines — #1505 owns those. Applied on top of each other, the suite gets through every test I could exercise here; I verified these four individually with manual link lines rather than claim a green make hip-test run that needs both PRs.

The machine stays available if you want anything re-run. Same box as #1502.

🤖 Generated with Claude Code

https://claude.ai/code/session_019WbKLB3C3vuXSZWDw5w83j

[huppiflupp]

Updated: rebased on dev with #1505 merged — make hip-test now exits 0 on gfx1151

Now that #1505 is in dev, I could do what the original description could not: run the whole
suite instead of verifying tests one at a time. Doing that surfaced two more blockers, so this
PR grew from four items to six. make hip-test now completes green on real gfx1151.

$ make hip-test ROCM_HOME=/usr HIPCC=/usr/bin/hipcc HIP_ARCH=gfx1151
...
./backend_cuda_test
./ragged_attention_test
./cuda_fmt_trap_test        cuda fmt trap tests: ok
./fp8_warp_test             OK
./absorb_determinism_test   batch_differing=0/145  ragged_differing=0/29
./fp8_cuda_test             OK
./weights_owned_test        8 injected H2D failures, free-mem delta 0 bytes
./mxfp4_cuda_test           test_mxfp4_cuda: ok
./alloc_footprint_test      test_alloc_footprint_cuda: ok
$ echo $?
0

Nine of nine, including the two the Makefile comments flag as schedule-sensitive or
"expected to fail" on some hosts.

The two new items

5. mxfp4_ref.o is built without -fPIC while GPUFLAGS carries -fPIE → the link fails
with relocation R_X86_64_32S cannot be used against local symbol. One flag on that rule.

6. GPUFLAGS carries -x hip, and it also applies to tests/mxfp4_ref.o on the link line,
so hipcc tries to compile the object file as source:

tests/mxfp4_ref.o:1:1: error: expected unqualified-id
tests/mxfp4_ref.o:1:41: error: source file is not valid UTF-8

Fixed with -x none before the object. Under CUDA nvcc does not do this, which is why it has
never bitten.

And item 3 turned out to cover one more file: test_alloc_footprint_cuda.cu is built the same
way as test_mxfp4_cuda.cu — separate TU against a separately compiled backend_cuda.cu,
including only the public ABI header — so it too needs backend_gpu_compat.h directly. It was
the last thing standing between the suite and exit 0.

Conflict note

Your #1505 and this branch both touch the mxfp4_cuda_test recipe. Resolved keeping both sides:
$(GPU_TEST_LIBS) from yours, -fPIC and -x none from here.

Diff is 9 files, +28/−9, no behaviour change outside the test targets. (I briefly force-pushed a
revision that had stray build artifacts in it — that is cleaned up, the branch is 9 source files.)

Same machine as #1502, still available if you want anything re-run.

[Kenneth-Javier]

Hey! Slightly off-topic, but I came across your Strix Halo work through this PR and ended up looking through your strix-halo-deepseek-v4-flash repo as well.

I've also been working on gfx1151 in Colibri, including the native Windows HIP support, and I'm currently digging into more model-independent HIP/XDNA2/heterogeneous execution work.

Seems like we're exploring a lot of the same hardware from different angles. If you're interested, I'd love to compare notes and maybe cross-test/tweak some Strix Halo stuff together sometime. 🙂

[huppiflupp]

Thanks for reaching out — glad the Strix Halo work was useful.

Happy to compare notes. Let me lead with the measurement that seems most relevant to your #1261, since it is the one datapoint I have that bears directly on whether an XDNA2 lane pays.

NPU vs iGPU on this box, 2026-08-16, Gemma4 (~4B active), identical prompts:

backend	decode
iGPU — Radeon 8060S via Vulkan	~53.6 tok/s
NPU — XDNA2 via FastFlowLM (flm:npu)	~12.0 tok/s

About 4.5x in the iGPU's favour for raw decode throughput. Two honest caveats: this was Gemma4, not GLM, and via Lemonade Server rather than Colibri — so it does not transfer directly to your shared-expert lane. And I did not measure power, which is presumably where the NPU actually wins. If your lane's case is efficiency rather than throughput, this number says nothing against it.

Where I think we are complementary: your #788 and #1261 are the Windows side, and I run this box on Linux — XRT 2.26.0 and the amdxdna DKMS driver built from source, because Nobara's ID=nobara breaks the exact-string flavour checks in AMD's build tooling even though ID_LIKE says fedora.

I have just put that up publicly so the numbers above are checkable rather than merely asserted: https://github.com/huppiflupp/strix-halo-npu-linux — the flavour patch, the memlock/limits config, a setup script, and the benchmark harness the NPU/iGPU figures came from. The README also records the negative result about ryzenai-llm below.

One thing you may already know, but it cost me a while to establish: AMD's ryzenai-llm — the stack that would actually split prefill and decode across NPU and iGPU — was Windows-only with no Linux roadmap when I last checked in August. If your heterogeneous work is heading that way, the Linux side currently has no vendor path, only what you build yourself.

What I can offer concretely: this machine (Ryzen AI MAX+ 395, 128 GB, gfx1151, NPU working under Linux) and a willingness to run controlled A/Bs on it. Tonight's numbers in #1502 and #1513 came off it. If you want the Gemma4 NPU comparison repeated inside Colibri rather than Lemonade, or your #1261 approach sanity-checked against the Linux XRT path, say the word.