rocm: host memory snapshot + capability-gated adaptive cache planner (with ROCm 6.3 build fixes)#3
Open
marmyx77 wants to merge 3 commits into
Open
Conversation
added 3 commits
July 15, 2026 22:33
…warp) 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_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).
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).
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Ciao Andrea — Marco here (La Bestia: Minisforum UM790 Pro, Radeon 780M/gfx1103, 64 GB, ROCm 6.3.1, DeepSeek V4 Flash q2 over SSD streaming at ~3.4 tok/s).
Your FORK_NOTES entry for the adaptive expert-cache planner says "CUDA/ROCm behavior must be measured before a backend-affecting PR". This PR brings the ROCm half: the host-memory snapshot the planner needs, the capability-based gate, the ROCm 6.3.1 build shims without which the backend doesn't compile at all — plus the measured failure case that motivates it.
The failure your planner prevents (reproduced on ROCm)
With a thin fixed free floor, the expert cache peaks after ~150–250 generated tokens and the kernel TTM evicts the GPU queue buffers themselves:
After that every
hipDeviceSynchronize/memcpy blocks forever with an idle GPU — generation freezes mid-request. Reliable repro on long generations; it also bit a real chat deployment twice in one evening.Measured floor ladder (280-token greedy soaks, same prompt, same box):
Side finding you may recognize from your 9/16 AUTO cap: the floor is a throughput optimum, not just a safety line — every GiB the cache wires is a GiB of page cache the 81 GB GGUF loses, and streaming slows more than residency gains.
What's in the PR (3 commits)
ds4_rocm.h):HIPBLAS_V2+CUDA_R_*→HIP_R_*,HIP_ENABLE_WARP_SYNC_BUILTINS,__syncwarp→ wavefront fence. Without these yourds4_rocm.ofails to compile on ROCm 6.3.1.ds4_gpu_host_memory_snapshotfor Linux/ROCm (was a stub returning 0):/proc/meminfo+/proc/self/status+ PSI pressure +hipMemGetInfototal as the UMA working-set envelope. Darwin→Linux field mapping documented inline.required_flooris stored and warned on (full fail-closed enforcement on ROCm still to be wired);slab_targetstays a documented no-op (per-expert cudaMalloc, no slabs).ds4.c: backends providing a valid snapshot get the adaptive budget; others keep legacy. Metal unchanged.Verification status (honest)
main(131e41e) with hipcc 6.3.1,--offload-arch=gfx1100, run on gfx1103 viaHSA_OVERRIDE_GFX_VERSION=11.0.0(rocWMMA rejects gfx1103).hipStreamWaitEventbarriers in the selected-id readback intermittently lose their signal on this ROCm/APU stack and wedge the HW queue (gdb stacks available). Host-side event waits are safe.— Marco (assisted by Claude)