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

CMakeLists.txt

@@ -333,6 +333,45 @@ elseif(VLLM_CPP_HIP)
   # in the RmsNorm reduction is exactly the kind of drift that budget is meant to
   # catch rather than absorb.
   add_compile_options($<$<COMPILE_LANGUAGE:HIP>:-ffp-contract=off>)
+  # Floor HIP DEVICE code at -O1 when nothing else sets an optimisation level.
+  #
+  # This is not a performance preference, it is a correctness workaround for the
+  # ROCm runtime (issue #132, VikashLoomba, 4x gfx1100 / ROCm 7.14). The
+  # documented first build — `cmake -S . -B build-hip -DVLLM_CPP_HIP=ON`, no
+  # CMAKE_BUILD_TYPE — passes no -O at all, so hipcc compiles at -O0. At -O0 all
+  # eight RmsNorm specialisations come out with AMDHSA metadata
+  # `.uses_dynamic_stack: true` plus `hidden_hostcall_buffer` / `hidden_heap_v1`,
+  # even though the kernel performs no hostcall. That makes CLR create a hostcall
+  # buffer and start a listener thread on the first launch.
+  #
+  # CLR's listener startup/termination handshake is racy (the shared state is
+  # `volatile`, not atomic, and the upstream source still carries its own
+  # `FIXME_lmoriche: fix termination handshake`): under host CPU saturation a
+  # delayed listener can overwrite the `kExit` written by terminate(), and DSO
+  # finalisation then spins forever waiting for a thread that has already exited.
+  # The symptom is `test_backend_cross_device` hanging AFTER printing
+  # `Status: SUCCESS!` — 11/11 cases, 39/39 assertions, then no exit.
+  #
+  # The report's controlled A/B, relinking the same library with ONLY
+  # rocm_rmsnorm.hip.o changed, on a 48-thread host:
+  #   -O0 object (dynamic stack, hidden hostcall + heap): 14/20 teardown timeouts
+  #   -O1 object (no dynamic stack, no hostcall/heap):    20/20 clean exits
+  # The same finaliser loop also reproduces in a standalone raw-HIP program that
+  # links neither vllm.cpp nor doctest, so the defect is CLR's, not the harness'.
+  # The definitive fix belongs upstream in CLR; this keeps our documented build
+  # off the path that triggers it.
+  #
+  # Deliberately yields to any explicit choice: if the developer set a build type
+  # or put an -O in CMAKE_HIP_FLAGS, that wins and nothing is added here.
+  string(TOUPPER "${CMAKE_BUILD_TYPE}" _vllm_hip_build_type)
+  if((_vllm_hip_build_type STREQUAL "" OR _vllm_hip_build_type STREQUAL "DEBUG")
+     AND NOT "${CMAKE_HIP_FLAGS}" MATCHES "(^| )-O")
+    add_compile_options($<$<COMPILE_LANGUAGE:HIP>:-O1>)
+    message(STATUS
+      "HIP device code floored at -O1: -O0 marks kernels as dynamic-stack users "
+      "and trips the ROCm CLR hostcall-listener teardown deadlock (issue #132)")
+  endif()
+  unset(_vllm_hip_build_type)
   set(VLLM_CPP_HIP ON)
 else()
   set(VLLM_CPP_HIP OFF)

docs/BUILD.md

@@ -147,9 +147,16 @@ real install, the configure now derives the compiler hints from it
 (`CMAKE_HIP_COMPILER_ROCM_ROOT`, `--rocm-path` in `CMAKE_HIP_FLAGS`, the
 `ROCM_PATH` environment variable — each only if you have not set it), which is
 what makes Arch and TheRock dist-tarball layouts configure without the manual
-flags issue #41's gfx1151 report needed. An explicit `-DCMAKE_BUILD_TYPE` (any
-optimizing one) matters on ROCm: a `-O0` device build trips a CLR teardown race
-([#132](https://github.com/mudler/vllm.cpp/issues/132)).
+flags issue #41's gfx1151 report needed.
+
+A build with no `CMAKE_BUILD_TYPE` now floors **HIP device code** at `-O1`
+automatically, and says so at configure time. At `-O0` hipcc marks the kernels
+as dynamic-stack users, which makes the ROCm runtime start a hostcall listener
+the kernels never use, and its teardown handshake can deadlock at process exit —
+the tests all pass and then the process never returns
+([#132](https://github.com/mudler/vllm.cpp/issues/132)). Setting a build type,
+or putting your own `-O` in `CMAKE_HIP_FLAGS`, overrides this and is respected
+as-is.
 
 `-DVLLM_CPP_HIP=ON` **fails the configure** when no HIP compiler is found rather
 than quietly producing a CPU-only build, for the same reason the CUTLASS note

docs/USAGE.md

@@ -22,6 +22,16 @@ example targets are named after the directories they are built from, so an
 in-source build makes the linker write each executable over its own source
 directory (issue #85).
 
+### One ROCm-specific behaviour
+
+Worth knowing before you read a hang as a bug in the tests: a build that sets no
+`CMAKE_BUILD_TYPE` floors **HIP device code** at `-O1` and prints a configure
+line saying so. At `-O0` the ROCm runtime starts a hostcall listener the kernels
+never use, and its teardown can deadlock at process exit — every test passes,
+`Status: SUCCESS!` prints, and the process never returns
+([#132](https://github.com/mudler/vllm.cpp/issues/132)). Setting a build type,
+or putting your own `-O` in `CMAKE_HIP_FLAGS`, overrides it.
+
 ## Confirming which CUDA architecture a build targets
 
 `CMakeCache.txt` is now a reliable answer. Configuring with

src/vt/rocm/rocm_matmul_hipblaslt.hip

@@ -6,6 +6,26 @@
 // Matmul:   out[M,N] = a[M,K] @ b[K,N]
 #include <hip/hip_bf16.h>
 #include <hip/hip_runtime.h>
+// hipBLAS ships TWO generations of the *Ex entry points and picks between them
+// in the header. The legacy generation takes `hipblasDatatype_t` (HIPBLAS_R_*);
+// the current one takes the HIP-wide `hipDataType` (HIP_R_*). This file is
+// written entirely against the current one — every data type below is HIP_R_*,
+// and the compute type is already the current `hipblasComputeType_t` — but it
+// never ASKED for it, so it only compiled on a ROCm new enough to make that
+// generation the default (7.x, where the row was developed on gfx1201).
+//
+// On ROCm 6.4 the legacy generation is the default, the two enums are distinct
+// types with no implicit conversion, and all six *Ex call sites fail with
+// "no matching function for call to 'hipblasGemmEx' ... no known conversion
+// from 'const hipDataType' to 'hipblasDatatype_t' for 9th argument" (issue #201,
+// gfx1100). Selecting the generation the file is written for fixes all six at
+// the include, rather than teaching six call sites to switch enums.
+//
+// Must precede <hipblas/hipblas.h>: it is what the header dispatches on. Newer
+// hipBLAS, where this generation is already the default, ignores it.
+#ifndef HIPBLAS_V2
+#define HIPBLAS_V2
+#endif
 #include <hipblas/hipblas.h>
 #include <hipblaslt/hipblaslt.h>