rocm: host memory snapshot + capability-gated adaptive cache planner (with ROCm 6.3 build fixes)

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; "

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 <hip/hip_runtime.h>
+/* 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 <hipblas/hipblas.h>
 #include <hip/hip_fp16.h>
 #include <hipcub/hipcub.hpp>
@@ -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;

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(&gtt_free, &gtt_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) {