From 4f9eac23a2d7e925fccdaa6a7ff7db18cc191993 Mon Sep 17 00:00:00 2001 From: Marco Armellino Date: Wed, 15 Jul 2026 22:33:42 +0200 Subject: [PATCH 1/3] rocm: ROCm 6.3.1 build shims (hipBLAS 2.x, warp-sync builtins, __syncwarp) Without these the ROCm backend does not compile on ROCm 6.3.1 at all: - HIPBLAS_V2 + CUDA_R_* -> HIP_R_*: hipBLAS 2.x only declares the hipDataType-based GemmEx/GemmStridedBatchedEx overloads behind HIPBLAS_V2; the deprecated hipblasDatatype_t ones no longer match. - HIP_ENABLE_WARP_SYNC_BUILTINS: ROCm 6.3 gates __shfl_*_sync/__ballot_sync behind this macro. - __syncwarp(...) -> wavefront-scoped fence: no HIP equivalent exists; on RDNA3 wave32 lanes execute in lockstep so reconvergence is implicit and the fence supplies the compiler ordering the call relies on. Verified: ds4_rocm.o builds clean with hipcc 6.3.1, --offload-arch=gfx1100 (run on a Radeon 780M / gfx1103 via HSA_OVERRIDE_GFX_VERSION=11.0.0, since rocWMMA 1.6.0 rejects gfx1103). --- ds4_rocm.h | 23 +++++++++++++++++++++-- 1 file changed, 21 insertions(+), 2 deletions(-) diff --git a/ds4_rocm.h b/ds4_rocm.h index 50481b8d5..c05b11259 100644 --- a/ds4_rocm.h +++ b/ds4_rocm.h @@ -1,6 +1,15 @@ #pragma once +/* Opt in to the __shfl_*_sync / __ballot_sync warp builtins, which ROCm 6.3 + * gates behind this macro (hip/amd_detail/amd_warp_sync_functions.h). The + * backend calls __shfl_sync/__shfl_down_sync with an 8-byte MASK_T mask. */ +#define HIP_ENABLE_WARP_SYNC_BUILTINS 1 + #include +/* hipBLAS 2.x hides the hipDataType-based GemmEx/GemmStridedBatchedEx behind + * HIPBLAS_V2; without it only the deprecated hipblasDatatype_t overloads are + * declared and the CUDA_R_* (now hipDataType) arguments don't match. */ +#define HIPBLAS_V2 #include #include #include @@ -79,8 +88,11 @@ #define CUBLAS_DEFAULT_MATH HIPBLAS_DEFAULT_MATH #define CUBLAS_COMPUTE_32F HIPBLAS_COMPUTE_32F #define CUBLAS_TF32_TENSOR_OP_MATH HIPBLAS_TF32_TENSOR_OP_MATH -#define CUDA_R_16F HIPBLAS_R_16F -#define CUDA_R_32F HIPBLAS_R_32F +/* hipBLAS 2.x (ROCm 6.3) hipblasGemmEx/GemmStridedBatchedEx take hipDataType + * (HIP_R_*), not the deprecated hipblasDatatype_t (HIPBLAS_R_*). The compute + * type stays hipblasComputeType_t (HIPBLAS_COMPUTE_32F). */ +#define CUDA_R_16F HIP_R_16F +#define CUDA_R_32F HIP_R_32F #define cublasCreate hipblasCreate #define cublasDestroy hipblasDestroy @@ -92,6 +104,13 @@ namespace cub = hipcub; +/* CUDA's __syncwarp has no HIP equivalent (not provided even with + * HIP_ENABLE_WARP_SYNC_BUILTINS on ROCm 6.3). On RDNA3 wave32 the lanes of a + * wave execute in lockstep, so warp reconvergence is implicit; a wavefront- + * scoped fence supplies the compiler ordering the call is relied on for. + * Variadic so both __syncwarp() and __syncwarp(mask) map through. */ +#define __syncwarp(...) __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "wavefront") + static __device__ __forceinline__ int32_t __vcmpne4(uint32_t a, uint32_t b) { // For each byte: 0xFF if a != b, 0x00 if a == b uint32_t diff = a ^ b; From 2722019d056f5564356a1a43a0f5831e2d2133d5 Mon Sep 17 00:00:00 2001 From: Marco Armellino Date: Wed, 15 Jul 2026 22:33:42 +0200 Subject: [PATCH 2/3] rocm: implement the host memory snapshot for the adaptive cache planner MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ds4_gpu_host_memory_snapshot was a stub returning 0 on ROCm, which kept the adaptive expert-cache planner permanently unavailable there. Implement it for Linux: /proc/meminfo (MemTotal, MemFree, Inactive(file), Cached), /proc/self/status VmRSS, PSI /proc/pressure/memory for the pressure flags, and hipMemGetInfo total as the working-set envelope (on UMA APUs the GTT pool is the closest analogue of recommendedMaxWorkingSetSize). The Darwin field mapping is documented inline. Also store the planner's required floor and warn when the configured budget is below it (full fail-closed enforcement on ROCm still to be wired); the slab target stays a documented no-op (the ROCm cache is per-expert cudaMalloc, no slab allocator). Motivation, measured on a Radeon 780M / 64 GB / 48 GiB GTT running DeepSeek V4 Flash q2 with SSD streaming: with a fixed thin free floor the cache peaks after ~150-250 generated tokens and the kernel TTM evicts the GPU queue buffers themselves — [drm:amdgpu_cs_ioctl] *ERROR* Not enough memory for command submission! amdgpu: Freeing queue vital buffer ..., queue evicted — after which every HIP sync blocks forever with an idle GPU. A host-aware floor prevents it: 5.8-6.6 GiB floors froze or ran 2.97 tok/s; a 12 GiB floor ran 280-token soaks clean twice at 3.47 tok/s (best measured on that box; the wired cache competes with the page cache of the 81 GB GGUF, so the floor is a throughput optimum too, not just a safety line). --- rocm/ds4_rocm_current_api_compat.cuh | 102 ++++++++++++++++++++++++++- 1 file changed, 100 insertions(+), 2 deletions(-) diff --git a/rocm/ds4_rocm_current_api_compat.cuh b/rocm/ds4_rocm_current_api_compat.cuh index 574f6e61b..ba9162850 100644 --- a/rocm/ds4_rocm_current_api_compat.cuh +++ b/rocm/ds4_rocm_current_api_compat.cuh @@ -174,9 +174,23 @@ extern "C" void ds4_gpu_set_streaming_expert_cache_budget(uint32_t experts) { g_stream_expert_cache_budget = experts; } +/* Fail-closed correctness floor from the adaptive planner (ds4_ssd.c). The + * ROCm cache does not yet shrink under live pressure the way Metal does, so + * the floor currently guards only the startup budget; storing it keeps the + * model contract visible to shared engine code and future enforcement. */ +static uint32_t g_stream_expert_cache_required_floor; + extern "C" void ds4_gpu_set_streaming_expert_cache_required_floor( uint32_t experts) { - (void)experts; + g_stream_expert_cache_required_floor = experts; + if (experts != 0 && g_stream_expert_cache_budget != 0 && + g_stream_expert_cache_budget < experts) { + fprintf(stderr, + DS4_GPU_LOG_PREFIX "streaming expert cache budget %u is below " + "the model's correctness floor %u\n", + g_stream_expert_cache_budget, + experts); + } } extern "C" void ds4_gpu_set_streaming_expert_cache_expert_bytes(uint64_t bytes) { @@ -199,9 +213,93 @@ extern "C" uint64_t ds4_gpu_recommended_working_set_size(void) { return (uint64_t)total_b; } +/* + * Linux/ROCm host memory snapshot for the adaptive SSD expert-cache planner. + * + * Field mapping from the Darwin collector (ds4_metal.m) to Linux sources: + * physical_bytes <- /proc/meminfo MemTotal + * recommended_bytes <- hipMemGetInfo() total: on UMA APUs the GTT pool is + * the platform's working-set envelope, the closest + * analogue of recommendedMaxWorkingSetSize + * task_footprint <- /proc/self/status VmRSS + * free_bytes <- MemFree (MemAvailable would double-count the + * file-backed credit the planner adds separately) + * purgeable_bytes <- 0 (no Linux analogue) + * inactive_bytes <- Inactive(file): reclaimable page-cache, the + * conservative counterpart of Mach inactive_count + * file_backed_bytes <- Cached + * pressure_normal <- PSI /proc/pressure/memory, "some avg10" < 5.0 + * + * Motivating failure (Radeon 780M, 64 GB, GTT 48 GiB): with the legacy fixed + * pool/8 floor the expert cache peaked after ~150-250 generated tokens and + * the kernel TTM evicted the GPU queue buffers themselves — dmesg shows + * "Not enough memory for command submission!" and "Freeing queue vital + * buffer, queue evicted" — freezing generation with an idle GPU. A + * host-aware budget prevents it; this snapshot is what the planner needs to + * compute one. + */ extern "C" int ds4_gpu_host_memory_snapshot(ds4_ssd_host_memory *out) { - if (out) memset(out, 0, sizeof(*out)); + if (!out) return 0; + memset(out, 0, sizeof(*out)); +#if !defined(__linux__) return 0; +#else + size_t gtt_free = 0; + size_t gtt_total = 0; + if (cudaMemGetInfo(>t_free, >t_total) != cudaSuccess || gtt_total == 0) { + (void)cudaGetLastError(); + return 0; + } + + uint64_t mem_total = 0, mem_free = 0, inactive_file = 0, cached = 0; + FILE *mi = fopen("/proc/meminfo", "r"); + if (!mi) return 0; + char line[160]; + while (fgets(line, sizeof(line), mi)) { + unsigned long long kb = 0; + if (sscanf(line, "MemTotal: %llu kB", &kb) == 1) mem_total = (uint64_t)kb << 10; + else if (sscanf(line, "MemFree: %llu kB", &kb) == 1) mem_free = (uint64_t)kb << 10; + else if (sscanf(line, "Inactive(file): %llu kB", &kb) == 1) inactive_file = (uint64_t)kb << 10; + else if (sscanf(line, "Cached: %llu kB", &kb) == 1) cached = (uint64_t)kb << 10; + } + fclose(mi); + if (mem_total == 0) return 0; + + uint64_t vm_rss = 0; + FILE *st = fopen("/proc/self/status", "r"); + if (st) { + while (fgets(line, sizeof(line), st)) { + unsigned long long kb = 0; + if (sscanf(line, "VmRSS: %llu kB", &kb) == 1) { + vm_rss = (uint64_t)kb << 10; + break; + } + } + fclose(st); + } + + out->physical_bytes = mem_total; + out->recommended_bytes = (uint64_t)gtt_total; + out->task_footprint_bytes = vm_rss; + out->free_bytes = mem_free; + out->purgeable_bytes = 0; + out->inactive_bytes = inactive_file; + out->file_backed_bytes = cached; + + FILE *psi = fopen("/proc/pressure/memory", "r"); + if (psi) { + while (fgets(line, sizeof(line), psi)) { + float avg10 = 0.0f; + if (sscanf(line, "some avg10=%f", &avg10) == 1) { + out->pressure_status_available = true; + out->pressure_normal = avg10 < 5.0f; + break; + } + } + fclose(psi); + } + return 1; +#endif } extern "C" uint32_t ds4_gpu_stream_expert_cache_configured_count(void) { From b60fbec8ddd73e50a8842b99e71bd06761a229ca Mon Sep 17 00:00:00 2001 From: Marco Armellino Date: Wed, 15 Jul 2026 22:33:42 +0200 Subject: [PATCH 3/3] ssd: gate the adaptive planner on snapshot capability, not backend Any backend that can produce a live host-memory snapshot now gets the adaptive budget; backends without one keep the legacy fixed-fraction budget. Metal behavior is unchanged. This is the cross-backend step the planner entry in FORK_NOTES was waiting on: ROCm UMA APUs need the host-aware floor as much as Metal does (see the previous commit for the queue-eviction failure it prevents). --- ds4.c | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/ds4.c b/ds4.c index e388b3686..325ed2ece 100644 --- a/ds4.c +++ b/ds4.c @@ -31037,11 +31037,18 @@ static bool ds4_engine_configure_streaming_auto_cache(ds4_engine *e) { return true; } - if (e->backend != DS4_BACKEND_METAL) { + /* The adaptive planner is capability-gated, not backend-gated: any + * backend that can produce a live host-memory snapshot gets the adaptive + * budget (ROCm UMA APUs need it as much as Metal does — without a + * host-aware floor the kernel TTM can evict the GPU queue buffers under + * peak cache pressure and freeze generation). Backends without a + * snapshot keep the legacy fixed-fraction budget. */ + ds4_ssd_host_memory memory; + if (e->backend != DS4_BACKEND_METAL && + !ds4_gpu_host_memory_snapshot(&memory)) { return ds4_engine_configure_streaming_legacy_auto_cache(e); } - ds4_ssd_host_memory memory; if (!ds4_gpu_host_memory_snapshot(&memory)) { fprintf(stderr, "ds4: SSD streaming auto cache: host memory snapshot unavailable; "