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