diff --git a/CMakeLists.txt b/CMakeLists.txt index 452f5bd07..e6fc56d6f 100644 --- a/CMakeLists.txt +++ b/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($<$:-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($<$:-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) diff --git a/docs/BUILD.md b/docs/BUILD.md index f92e41870..4bb6ab041 100644 --- a/docs/BUILD.md +++ b/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 diff --git a/docs/USAGE.md b/docs/USAGE.md index 2a1fb6ba3..f65c93b90 100644 --- a/docs/USAGE.md +++ b/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 diff --git a/src/vt/rocm/rocm_matmul_hipblaslt.hip b/src/vt/rocm/rocm_matmul_hipblaslt.hip index 70b21f239..2e6334d39 100644 --- a/src/vt/rocm/rocm_matmul_hipblaslt.hip +++ b/src/vt/rocm/rocm_matmul_hipblaslt.hip @@ -6,6 +6,26 @@ // Matmul: out[M,N] = a[M,K] @ b[K,N] #include #include +// 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 : 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 #include