From 3e28c4374cbad6ee00931ee1486d38ea3465f588 Mon Sep 17 00:00:00 2001 From: moonshadow-25 Date: Sat, 14 Mar 2026 12:47:58 +0800 Subject: [PATCH 1/2] ggml/hip: fix APU compatibility - soft error handling for hipMemAdviseSetCoarseGrain On AMD APU/iGPU devices (unified memory architecture), hipMemAdviseSetCoarseGrain returns hipErrorInvalidValue because the hint is not applicable to UMA systems. The previous CUDA_CHECK() call treated this as a fatal error, causing crashes on APU systems such as AMD Strix Halo (gfx1151). Fix: treat hipMemAdviseSetCoarseGrain as an optional performance hint - call it without error checking and clear any resulting error with hipGetLastError(). Also add pre-allocation debug logging (GGML_LOG_DEBUG) to help diagnose memory issues on APU systems, and store totalGlobalMem in device info. Context: AMD APUs on Windows are affected by a ROCm runtime bug that limits hipMallocManaged to ~64GB regardless of available system RAM. A fix has been submitted upstream: https://github.com/ROCm/rocm-systems/pull/4077 Co-Authored-By: Claude Sonnet 4.6 --- ggml/src/ggml-cuda/ggml-cuda.cu | 12 +++++++++++- 1 file changed, 11 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 9d2aacf4b2c2..aa6e63d749af 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -119,12 +119,21 @@ int ggml_cuda_get_device() { static cudaError_t ggml_cuda_device_malloc(void ** ptr, size_t size, int device) { ggml_cuda_set_device(device); + { + size_t mem_free = 0, mem_total = 0; + cudaMemGetInfo(&mem_free, &mem_total); + GGML_LOG_DEBUG("%s: before alloc %.2f MiB, free %.2f MiB / total %.2f MiB\n", + __func__, size/1024.0/1024.0, mem_free/1024.0/1024.0, mem_total/1024.0/1024.0); + } cudaError_t err; if (getenv("GGML_CUDA_ENABLE_UNIFIED_MEMORY") != nullptr) { err = cudaMallocManaged(ptr, size); #if defined(GGML_USE_HIP) if (err == hipSuccess) { - CUDA_CHECK(cudaMemAdvise(*ptr, size, hipMemAdviseSetCoarseGrain, device)); + // hipMemAdviseSetCoarseGrain is an optional performance hint; + // ignore errors (e.g. hipErrorInvalidValue on some APU/iGPU configs). + cudaMemAdvise(*ptr, size, hipMemAdviseSetCoarseGrain, device); + (void)hipGetLastError(); // clear any error } // fall back to cudaMalloc if not supported (e.g. on Windows) @@ -238,6 +247,7 @@ static ggml_cuda_device_info ggml_cuda_init() { info.default_tensor_split[id] = total_vram; total_vram += prop.totalGlobalMem; + info.devices[id].total_vram = prop.totalGlobalMem; info.devices[id].integrated = false; // Temporarily disabled due to issues with corrupted output (e.g. #15034) info.devices[id].nsm = prop.multiProcessorCount; info.devices[id].smpb = prop.sharedMemPerBlock; From 4c72b964f5ff11fdcd69fb2b8741894a8f37e935 Mon Sep 17 00:00:00 2001 From: moonshadow-25 Date: Sun, 15 Mar 2026 09:11:27 +0800 Subject: [PATCH 2/2] ggml/hip: remove unrelated changes, keep only hipMemAdviseSetCoarseGrain fix --- ggml/src/ggml-cuda/ggml-cuda.cu | 7 ------- 1 file changed, 7 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index aa6e63d749af..57e2ee597ef0 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -119,12 +119,6 @@ int ggml_cuda_get_device() { static cudaError_t ggml_cuda_device_malloc(void ** ptr, size_t size, int device) { ggml_cuda_set_device(device); - { - size_t mem_free = 0, mem_total = 0; - cudaMemGetInfo(&mem_free, &mem_total); - GGML_LOG_DEBUG("%s: before alloc %.2f MiB, free %.2f MiB / total %.2f MiB\n", - __func__, size/1024.0/1024.0, mem_free/1024.0/1024.0, mem_total/1024.0/1024.0); - } cudaError_t err; if (getenv("GGML_CUDA_ENABLE_UNIFIED_MEMORY") != nullptr) { err = cudaMallocManaged(ptr, size); @@ -247,7 +241,6 @@ static ggml_cuda_device_info ggml_cuda_init() { info.default_tensor_split[id] = total_vram; total_vram += prop.totalGlobalMem; - info.devices[id].total_vram = prop.totalGlobalMem; info.devices[id].integrated = false; // Temporarily disabled due to issues with corrupted output (e.g. #15034) info.devices[id].nsm = prop.multiProcessorCount; info.devices[id].smpb = prop.sharedMemPerBlock;