fix(rocm): build on ROCm 6.x, and keep -O0 off the CLR hostcall path (#201, #132)

[localai-bot]

Closes #201. Closes #132.

Two independent ROCm build defects, both reported with the root cause already found. Neither reproduces here — no AMD GPU and no ROCm toolchain on any maintainer machine — so the reporters are the verification.

#201 — hipblasGemmEx overload mismatch (gfx1100, ROCm 6.4)

hipBLAS ships two generations of the *Ex entry points and chooses in the header: the legacy one takes hipblasDatatype_t (HIPBLAS_R_*), the current one takes the HIP-wide hipDataType (HIP_R_*). Distinct enums, no implicit conversion.

rocm_matmul_hipblaslt.hip is written entirely against the current generation — all 18 type constants are HIP_R_*, the compute type is already hipblasComputeType_t — but it never asked for it, so it compiled only where that generation is the default. That is ROCm 7.x, where the row was developed (gfx1201, #140). On ROCm 6.4 all six *Ex call sites fail at once.

Selecting the generation the file is written for, at the include, fixes all six without touching a call site. The alternative — a dual mapper plus six switched enums — is more code to say the same thing and would leave the file half-legacy. Nothing else in the TU is generation-sensitive: hipblasCreate/hipblasDestroy/hipblasSetStream are unchanged across both, and everything else is hipBLASLt.

#132 — -O0 device code deadlocks CLR teardown (gfx1100, ROCm 7.14)

@VikashLoomba's report is unusually complete and this follows its stated fix boundary.

The documented first ROCm build passes no -O at all, so hipcc compiles at -O0. At -O0 all eight RmsNorm specialisations come out with .uses_dynamic_stack: true plus hidden_hostcall_buffer / hidden_heap_v1 despite performing no hostcall. CLR then creates a hostcall buffer and starts a listener thread on first launch, and its startup/termination handshake is racy — the shared state is volatile, not atomic, and the upstream source still carries FIXME_lmoriche: fix termination handshake. Under host CPU saturation a delayed listener overwrites the kExit that terminate() wrote, and DSO finalisation spins forever on a thread that has already exited.

The symptom is test_backend_cross_device printing 11/11 cases, 39/39 assertions, Status: SUCCESS! — and never exiting.

Their controlled A/B relinked the same library with only rocm_rmsnorm.hip.o changed, on a 48-thread host:

RmsNorm object	metadata	result
default -O0	dynamic stack; hidden hostcall + heap	14/20 teardown timeouts
-O1	no dynamic stack, no hostcall/heap	20/20 clean exits

They also reproduced the same finaliser loop in a standalone raw-HIP program linking neither vllm.cpp nor doctest, which is what places the defect in CLR rather than in our harness.

So the definitive fix is upstream in CLR; this keeps our documented build off the path that triggers it, by flooring HIP device code at -O1 when nothing else sets a level. It yields to any explicit choice.

Verified here (condition logic only, no HIP):

CMAKE_BUILD_TYPE	CMAKE_HIP_FLAGS	result
(empty)	(empty)	floor -O1
Debug / debug	(empty)	floor -O1
Release / RelWithDebInfo	(empty)	unchanged
(empty)	-O0	unchanged
Debug	-O0	unchanged
(empty)	-Ofast	unchanged
(empty)	--rocm-path=/opt/rocm	floor -O1 — correctly not read as an -O

That last row matters: the ROCm block itself sets --rocm-path in CMAKE_HIP_FLAGS.

@dpblnt @VikashLoomba — a build report either way would be the gate here.

[VikashLoomba]

Verified on the #132 reporting hardware (4× RX 7900 XTX gfx1100, ROCm 7.14.0, HIP 7.14.60850, AMD Clang 23, kernel 7.0.0-28-generic).

Build config under test: the documented no-build-type command (cmake -S . -B build -DVLLM_CPP_HIP=ON -DROCM_PATH=/opt/rocm, no CMAKE_BUILD_TYPE) — the exact #132 trigger.

Check	base (7533eeb2)	PR branch (2b812abc)
-O on HIP compile lines	none (the defect precondition)	-O1 present (compile_commands.json, all .hip TUs)
Configure log	—	prints the floor message
ctest -R 'rocm|cross_device'	—	3/3 PASS
Saturated teardown, load shape A (48 pure spinners, 20 sequential runs)	20/20 clean	20/20 clean
Saturated teardown, load shape B (4 concurrent processes × 10 rounds, 48 dd churn workers, 40 exposures)	33/40 teardown hangs (82.5%)	40/40 clean

Notes, for the record:

1. Load shape matters. My original repro (14/20) was on an older tree; on current main the pure-spinner load no longer reproduces the race at all (20/20 on the pre-fix base). The load shape that re-establishes it reliably is concurrent process teardown + syscall churn — 4 concurrent instances (four listener/finalizer pairs racing) with 48 dd workers stirring the run queues. Under that load the pre-fix base fails 33/40, matching the original mechanism (delayed listener startup losing to the finalizer).
2. The fix side is clean under both shapes, including 40/40 under the load that fails the base 82.5% of the time. -O1 is confirmed present on the HIP compile lines, and the truth-table claim in the commit message holds on a real no-build-type configure (the floor fires; explicit -DCMAKE_BUILD_TYPE=Release builds from my earlier runs were already clean and stay unchanged).
3. rocm_matmul_hipblaslt.hip:384:13: error: no matching function for call to 'hipblasGemmEx' #201 side: this box is ROCm 7.14, so I can only confirm the include-level generation select is compile- and runtime-neutral on 7.x (full-tree build clean, gates green above) — the ROCm 6.4 fix itself needs the rocm_matmul_hipblaslt.hip:384:13: error: no matching function for call to 'hipblasGemmEx' #201 reporter's toolchain.

#132 verification: PASS as far as this hardware can take it — the documented first-build path no longer enters the hostcall-listener race, verified by controlled A/B on the same machine and load.