From b65a2dce12c14a489e19a059cb6ee59112f1b733 Mon Sep 17 00:00:00 2001 From: danielhanchen Date: Mon, 31 Aug 2026 09:23:42 +0000 Subject: [PATCH 1/2] ggml-cuda: make GGML_CUDA_ENABLE_UNIFIED_MEMORY=0 actually disable it ggml tests getenv(...) != nullptr, so =0 enables managed allocation. The name is spelled as a positive, so people reasonably write =0 to turn it off and get the opposite. GGML_CUDA_DISABLE_FUSION in this file already parses its value for the same reason. Empty and the usual falsy spellings are off; =1, =true, =on and =yes are unchanged. This is no longer a correctness bug. Before the cpy 2D fast path was fixed, =0 left the user corrupted while they believed they had opted out. It now costs throughput and capacity instead: on an integrated GPU managed allocation draws host RAM instead of the device carve-out rather than adding to it, measured at about 15% slower, and a model that loads without the variable can be OOM-killed with it. The HIP warning is rewritten accordingly. The earlier draft warned about corrupted output and linked ggml-org#26148; shipping that alongside the fast path fix would tell users to chase a defect their build does not have. It now states the cost, which is what remains true. Note this changes behaviour for every CUDA and HIP user who currently sets =0, not only AMD: they lose managed allocation, which is the point, but it is not a no-op. Rebased onto an upstream ancestor of both master and the base tag so it can be pinned; a branch cut from fork master drags master's divergence into the pin merge and conflicts on the deleted upstream workflow files. --- ggml/src/ggml-cuda/ggml-cuda.cu | 46 +++++++++++++++++++++++++++++++-- 1 file changed, 44 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 2456f7dcc621..220c0dc39dff 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -85,6 +85,7 @@ #include #include #include +#include #include #include #include @@ -95,6 +96,28 @@ static_assert(sizeof(half) == sizeof(ggml_fp16_t), "wrong fp16 size"); #define GGML_LOG_WARN_ONCE(str) \ { static std::once_flag warn_flag; std::call_once(warn_flag, []() { GGML_LOG_WARN(str); }); } +// Whether an environment variable is set to a value that means "on". +// +// Most ggml switches test presence only, and for a name like GGML_CUDA_NO_PINNED +// that is right: the only sensible use is to set it. GGML_CUDA_ENABLE_UNIFIED_MEMORY +// is spelled as a positive, so people reasonably write =0 to turn it off, and under a +// presence test that enabled the very path they were trying to leave. GGML_CUDA_DISABLE_FUSION +// in this file already parses its value for the same reason. +// +// Empty and the usual falsy spellings are off, anything else is on, so =1, =true, +// =on and =yes all keep working exactly as before. +static bool ggml_cuda_env_enabled(const char * name) { + const char * val = getenv(name); + if (val == nullptr || val[0] == '\0') { + return false; + } + std::string v(val); + for (char & c : v) { + c = (char) std::tolower((unsigned char) c); + } + return !(v == "0" || v == "false" || v == "no" || v == "off"); +} + [[noreturn]] void ggml_cuda_error(const char * stmt, const char * func, const char * file, int line, const char * msg) { int id = -1; // in case cudaGetDevice fails @@ -139,7 +162,26 @@ int ggml_cuda_get_device() { static cudaError_t ggml_cuda_device_malloc(void ** ptr, size_t size, int device) { ggml_cuda_set_device(device); cudaError_t err; - if (getenv("GGML_CUDA_ENABLE_UNIFIED_MEMORY") != nullptr) { + if (ggml_cuda_env_enabled("GGML_CUDA_ENABLE_UNIFIED_MEMORY")) { +#if defined(GGML_USE_HIP) + // Say so once. Every report of this path going wrong has come from someone + // who did not know a front-end had set the variable for them. + // + // This deliberately does NOT warn about corrupted output. An earlier draft + // did, pointing at ggml-org#26148, but that defect is the cpy 2D fast path + // using cudaMemcpyDeviceToDevice on managed pages, and it is fixed + // separately. Warning about it in a tree that carries the fix would send + // users chasing a bug their build does not have. + // + // What remains true is the cost. On an integrated GPU managed allocation + // draws host RAM instead of the device carve-out rather than adding to it, + // so it is usually slower and can lower the ceiling: a model that loads + // without the variable may be OOM-killed with it. + GGML_LOG_WARN_ONCE("GGML_CUDA_ENABLE_UNIFIED_MEMORY is set, allocating device memory as managed. " + "On integrated GPUs this draws host RAM instead of the device carve-out, which is " + "usually slower and can reduce the largest model that will load. Unset the variable " + "to disable it.\n"); +#endif // defined(GGML_USE_HIP) err = cudaMallocManaged(ptr, size); #if defined(GGML_USE_HIP) if (err == hipSuccess) { @@ -4789,7 +4831,7 @@ static void ggml_backend_cuda_device_get_memory(ggml_backend_dev_t dev, size_t * CUDA_CHECK(cudaGetDeviceProperties(&prop, ggml_cuda_get_physical_device(ctx->device))); // Check if UMA is explicitly enabled via environment variable - bool uma_env = getenv("GGML_CUDA_ENABLE_UNIFIED_MEMORY") != nullptr; + bool uma_env = ggml_cuda_env_enabled("GGML_CUDA_ENABLE_UNIFIED_MEMORY"); bool is_uma = prop.integrated > 0 || uma_env; if (is_uma) { From efa8b0e689ac8ba13f5840d9c26aa227ffd07153 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Wed, 14 Oct 2026 12:24:04 -0700 Subject: [PATCH 2/2] HIP: warn about managed memory only once it was allocated On hipErrorNotSupported the allocation falls back to hipMalloc, so the warning printed before it described memory that was never managed. Also name =0 as the way to turn it off. --- ggml/src/ggml-cuda/ggml-cuda.cu | 25 ++++++------------------- 1 file changed, 6 insertions(+), 19 deletions(-) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 7bc8b4d7160a..2f5abd5e055d 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -186,28 +186,15 @@ static cudaError_t ggml_cuda_device_malloc(void ** ptr, size_t size, int device) ggml_cuda_set_device(device); cudaError_t err; if (ggml_cuda_unified_memory_enabled()) { -#if defined(GGML_USE_HIP) - // Say so once. Every report of this path going wrong has come from someone - // who did not know a front-end had set the variable for them. - // - // This deliberately does NOT warn about corrupted output. An earlier draft - // did, pointing at ggml-org#26148, but that defect is the cpy 2D fast path - // using cudaMemcpyDeviceToDevice on managed pages, and it is fixed - // separately. Warning about it in a tree that carries the fix would send - // users chasing a bug their build does not have. - // - // What remains true is the cost. On an integrated GPU managed allocation - // draws host RAM instead of the device carve-out rather than adding to it, - // so it is usually slower and can lower the ceiling: a model that loads - // without the variable may be OOM-killed with it. - GGML_LOG_WARN_ONCE("GGML_CUDA_ENABLE_UNIFIED_MEMORY is set, allocating device memory as managed. " - "On integrated GPUs this draws host RAM instead of the device carve-out, which is " - "usually slower and can reduce the largest model that will load. Unset the variable " - "to disable it.\n"); -#endif // defined(GGML_USE_HIP) err = cudaMallocManaged(ptr, size); #if defined(GGML_USE_HIP) if (err == hipSuccess) { + // said once: reports of this path going wrong come from users who did not know a front-end set the variable + GGML_LOG_WARN_ONCE("GGML_CUDA_ENABLE_UNIFIED_MEMORY is set, allocating device memory as managed. " + "On integrated GPUs this draws host RAM instead of the device carve-out, which is " + "usually slower and can reduce the largest model that will load. Unset the variable " + "or set it to 0 to disable it.\n"); + // hipMemAdviseSetCoarseGrain is an optional performance hint; // ignore errors (e.g. hipErrorInvalidValue on some APU/iGPU configs). (void)cudaMemAdvise(*ptr, size, hipMemAdviseSetCoarseGrain, device);