fix(docker): build the ROCm image on ROCm 7.2.4

Dockerfile.rocm

@@ -3,18 +3,18 @@
 # ─── Stage 1: builder ───────────────────────────────────────────────────────
 # ROCm/HIP sibling of the CUDA Dockerfile. ROCM_VERSION / UBUNTU_VERSION /
 # LUCE_HIP_ARCHES are build args so the same Dockerfile can be repinned.
-#   • lucebox-hub:rocm  — ROCm 6.4.x, gfx1151 (+ optional fat gfx list)
+#   • lucebox-hub:rocm  — ROCm 7.2.4, gfx1151 (+ optional fat gfx list)
 # See docker-bake.hcl for the canonical invocation.
 #
-# NOTE: gfx1151 (Strix Halo / Ryzen AI MAX) needs ROCm >= 6.4.1. The default
-# stays on 6.4.1 because the 7.2.x stack has shown intermittent problems on
-# Strix Halo. The flip side: against a ROCm 7.x HOST driver the 6.4.x
-# userspace can segfault at model load (seen on gfx1151 + host ROCm 7.2.2:
-# SIGSEGV in backend creation, bogus 1.28 TB VRAM report) — in that case
-# rebuild with ROCM_VERSION=7.2.2 to match the host. Rule of thumb: keep the
-# base's major version aligned with the host driver.
-ARG ROCM_VERSION=6.4.1
-ARG UBUNTU_VERSION=22.04
+# NOTE: ROCm 7.2.4 is the stack the engine is built and tested on (the
+# Lucebox boxes and the AMD CI runner); the engine no longer builds on 6.x.
+# Keep the base's major version aligned with the host driver: a 6.4.x
+# userspace segfaults at model load against a ROCm 7.x host driver (seen on
+# gfx1151 + host ROCm 7.2.2: SIGSEGV in backend creation, bogus 1.28 TB VRAM
+# report). Ubuntu 24.04: ROCm 7.2's CMake package files need CMake >= 3.25,
+# and 22.04 ships 3.22.
+ARG ROCM_VERSION=7.2.4
+ARG UBUNTU_VERSION=24.04
 FROM rocm/dev-ubuntu-${UBUNTU_VERSION}:${ROCM_VERSION} AS builder
 
 ARG DEBIAN_FRONTEND=noninteractive
@@ -164,8 +164,11 @@ RUN apt-get update && apt-get install -y --no-install-recommends \
 # version is fixed and no remote installer script runs at build time.
 COPY --from=ghcr.io/astral-sh/uv:0.11.2 /uv /uvx /usr/local/bin/
 
+# Ubuntu 24.04 ships a Python 3.12 of its own; take uv's managed build anyway,
+# as the CUDA image does, so the interpreter lives under /opt/uv.
 ENV UV_PYTHON_INSTALL_DIR=/opt/uv/python \
-    UV_TOOL_DIR=/opt/uv/tools
+    UV_TOOL_DIR=/opt/uv/tools \
+    UV_PYTHON_PREFERENCE=only-managed
 
 WORKDIR /opt/lucebox-hub
 

docker-bake.hcl

@@ -54,12 +54,10 @@ variable "LUCE_CUDA_ARCHES" { default = "75;80;86;89;90;120" }
 # gfx1201 explicitly.
 variable "LUCE_HIP_ARCHES" { default = "gfx1151" }
 
-# ROCm base-image tag for the rocm variant. gfx1151 needs >= 6.4.1. Default
-# stays 6.4.1 (7.2.x has shown intermittent problems on Strix Halo), but on a
-# ROCm 7.x HOST driver the 6.4.x userspace can segfault at model load — set
-# ROCM_VERSION=7.2.2 there. Keep the base aligned with the host driver (see
-# Dockerfile.rocm).
-variable "ROCM_VERSION" { default = "6.4.1" }
+# ROCm base-image tag for the rocm variant: 7.2.4, the stack the engine is
+# built and tested on. Keep it aligned with the host driver's major version
+# (see Dockerfile.rocm).
+variable "ROCM_VERSION" { default = "7.2.4" }
 
 # Image identity stamped into /opt/lucebox-hub/IMAGE_INFO at build time and
 # surfaced under /props.build at runtime (git_sha, image_tag, build_time).
@@ -141,7 +139,7 @@ target "_rocm-base" {
     dockerfile = "Dockerfile.rocm"
     args = {
         ROCM_VERSION      = ROCM_VERSION
-        UBUNTU_VERSION    = "22.04"
+        UBUNTU_VERSION    = "24.04"
         LUCE_HIP_ARCHES = LUCE_HIP_ARCHES
         GIT_SHA           = GIT_SHA
         IMAGE_TAG         = IMAGE_TAG

server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-indexer.cu

@@ -979,8 +979,10 @@ void ggml_cuda_op_ds4_indexer_score(
         ? ds4_env_flag_enabled(packed_small_name)
         : GGML_CUDA_CC_IS_RDNA3_5(device_info.cc) ||
           GGML_CUDA_CC_IS_RDNA4(device_info.cc);
+#if DS4_INDEXER_WMMA_AVAILABLE
     // Tokenwise (ggml_ds4_indexer_score_tokenwise): each token scored by the
-    // one-token kernel, whatever the multi-token kernels would pick.
+    // one-token kernel, whatever the multi-token kernels would pick. Without
+    // rocWMMA 2.x every score takes the scalar kernel below, one token each.
     if (ggml_get_op_params_i32(dst, 2) != 0 && wmma_capable && q->type == GGML_TYPE_F32) {
         const dim3 grid((unsigned) ((n_comp + 127) / 128), (unsigned) n_tokens, 1);
         ds4_indexer_score_decode_wmma_kernel<<<grid, 256, 0, stream>>>(
@@ -994,7 +996,6 @@ void ggml_cuda_op_ds4_indexer_score(
         CUDA_CHECK(cudaGetLastError());
         return;
     }
-#if DS4_INDEXER_WMMA_AVAILABLE
 #if defined(GGML_USE_HIP)
     // The M32 switches are read per call on purpose: the unit tests toggle
     // them between graph computes inside one process. M32 defaults on for

server/hip_compat/cuda_runtime.h

@@ -99,8 +99,8 @@ using cudaDeviceProp        = hipDeviceProp_t;
 // Pointer attributes — used by geometric_draft_topk_cuda.cu to confirm the
 // logits pointer is device memory and find which device it lives on. The
 // hipPointerAttribute_t `.type` member (an hipMemoryType) is spelled `type` on
-// ROCm >= 5.0; the AMD targets in scope (ROCm 6.4.x on the CI runner, 7.x per
-// the maintainer's discrete-RDNA3 guidance) all use `type`.
+// ROCm >= 5.0; the AMD targets in scope (ROCm 7.2 on the CI runner and in the
+// published image) all use `type`.
 using cudaPointerAttributes = hipPointerAttribute_t;
 #define cudaPointerGetAttributes            hipPointerGetAttributes
 #define cudaMemoryTypeDevice                hipMemoryTypeDevice