Repository navigation
HIP: say once that GGML_CUDA_ENABLE_UNIFIED_MEMORY is on - #149
danielhanchen wants to merge 3 commits into
Conversation
Codex Review SummaryThis comment shows the latest Codex review activity on this pull request.
ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings. |
|
One note on how this reaches users, since I got it wrong earlier today on #145. The nightly tree is the upstream tag plus Happy to open the pin as a follow-up, or to take it upstream first if that is preferred, in which case the pin points at the ggml-org PR and retires itself. |
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.
a78bbb1 to
b65a2dc
Compare
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: b65a2dce12
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| 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"); |
There was a problem hiding this comment.
On HIP configurations where cudaMallocManaged returns hipErrorNotSupported, such as the Windows case handled by the fallback below, this warning claims that memory is being allocated as managed and describes managed-memory costs before the allocation is attempted; the code then uses ordinary cudaMalloc instead. Emit this warning only from the err == hipSuccess branch so users do not receive an inapplicable warning immediately followed by the fallback warning.
Useful? React with 👍 / 👎.
| // 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. |
There was a problem hiding this comment.
Replace this implementation-history narrative and the similar long block before the managed allocation with concise comments that retain only non-obvious invariants. These blocks hard-wrap sentences and document prior drafts, issue history, and obvious environment-variable behavior across many lines, contrary to this repository's explicit requirement that comments usually remain 1-2 lines, avoid excessive commentary, and not be hard-wrapped.
AGENTS.md reference: AGENTS.md:L74-L78
Useful? React with 👍 / 👎.
…les it (#162) ggml tests getenv(...) != nullptr, so =0 enables managed allocation. With #157 pinned this is no longer a correctness bug, but it still costs about 15% throughput and can lower the ceiling: managed draws host RAM instead of the device carve-out rather than adding to it, so a model that loads without the variable can be OOM-killed with it. Verified locally to co-merge onto b10705 with #157 and #158, net +79/-5 across 4 files. Co-authored-by: danielhanchen <elliegouldingstuff@gmail.com>
|
| build | env | outcome |
|---|---|---|
ggml-org b10766 win-rocm-7.14 |
GGML_CUDA_ENABLE_UNIFIED_MEMORY=0 |
host starved to 5.21 GiB free, killed |
| unsloth b10798-mix-659e406 | GGML_CUDA_ENABLE_UNIFIED_MEMORY=0 |
completed normally, pp128 158.71 t/s |
| unsloth b10798-mix-659e406 | GGML_CUDA_ENABLE_UNIFIED_MEMORY=1 |
free RAM to 0.01 GiB, killed |
| unsloth b10798-mix-659e406 | unset | completed normally |
So on mainline, =0 behaves like =1 — the presence test this PR describes. On this fork's current prebuilt, =0 behaves like unset, which is what a user typing =0 is asking for.
The =1 row is worth keeping in mind for the warning half of this PR: at a large carve-out the variable is not a performance trade, it is a hard denial of service on the host. A user who set it once for a big model on a small carve-out, then raised their carve-out, would hit this with no indication of why. That is precisely the population the GGML_LOG_WARN_ONCE is aimed at, and it argues for keeping the warning rather than gating it further.
Caveats
- One host, one driver, one model. Windows only.
- The mainline package needed its ROCm 7.14.0 runtime staged beside it to use the GPU at all (
rocblas,hipblas,libhipblaslt,amdhip64_7and the kernel directories); without it,--list-devicesreports(none)and it silently falls back to CPU. Backend was verified asROCm : NO_VMM = 1insystem_infofor every GPU cell. - This measures only the switch behaviour. It says nothing about the underlying corruption in ggml-org#26148, which this PR is explicit about not fixing.
…_enabled() b11609 parses GGML_CUDA_ENABLE_UNIFIED_MEMORY by value (ggml-org#30327), so =0 already disables it upstream. Drop ggml_cuda_env_enabled() and keep only the once-per-process HIP warning.
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.
danielhanchen
changed the title
This branch has not been deployed
1 participant
Overview
One small change to
ggml-cuda.cu: on HIP, say once thatGGML_CUDA_ENABLE_UNIFIED_MEMORYis on. Users on AMD APUs often run with unified memory without knowing it, because a front-end set the variable for them.The other half of the original change, making
GGML_CUDA_ENABLE_UNIFIED_MEMORY=0turn it off, is upstream since ggml-org#30327 (ggml_cuda_env_flag, which parses the value withatoi, so1enables it and0, empty or a word such astruedo not). This branch now carries only the warning on top of it.The warning
After a successful
hipMallocManaged, once per process:It is printed only when the managed allocation succeeds. When
hipMallocManagedreturnshipErrorNotSupported(for example on Windows), the existing fallback tohipMallocand its own warning are unchanged, so the two messages no longer contradict each other. HIP only, since that is where the reports come from. It does not point at the old corruption report (ggml-org#26148): that defect is in the 2D copy fast path on managed pages and is fixed separately.Verification
HIP is not built on this host. The change moves an existing
GGML_LOG_WARN_ONCEcall into the existingerr == hipSuccessbranch; CUDA and non-HIP builds are unaffected (the code is inside#if defined(GGML_USE_HIP)).