fix(multi-gpu): disable torch's post-conv-algorithm-search global emptyCache - #196
Draft
lstein wants to merge 1 commit into
Draft
fix(multi-gpu): disable torch's post-conv-algorithm-search global emptyCache#196lstein wants to merge 1 commit into
lstein wants to merge 1 commit into
Conversation
…tyCache torch's conv algorithm search (MIOpen on every algo-cache miss; cuDNN in benchmark mode) ends each find with a process-global emptyCache() to release benchmarking workspace. On a multi-GPU box that call takes every device's allocator mutex and frees their cached blocks, freezing a peer GPU's worker for the remainder of its in-flight step (observed via py-spy on a dual-W7900 ROCm rig: chooseAlgorithm -> emptyCache -> hipFree waiting out a 40-100 s denoise step). It is the one convoy layer the Python-level peer-aware empty_cache wrapper cannot intercept. The call is gated on _cudnn_get_conv_benchmark_empty_cache(), and torch exposes the setter. When more than one CUDA/HIP generation device is registered, clear the flag at session-processor startup: benchmarking workspace blocks stay cached in the allocator for reuse instead of being returned to the driver - the same trade the peer-aware skips already make. Single-GPU installs are unchanged. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_014xFbHnmFTLsHE9e1PAdvme
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
The last convoy layer
After #159's peer-aware skips, one
empty_cacheconvoy remained, C++-side: torch's conv algorithm search (chooseAlgorithm→findAlgorithm) ends every find with a process-globalCUDACachingAllocator::emptyCache()to release benchmarking workspace. On MIOpen this runs on every algo-cache miss (each new conv shape per process), benchmark flag or not. py-spy caught it on the rig: one worker insidechooseAlgorithm -> emptyCache -> hipFree, waiting out the other worker's 40-100 s denoise step with every device's allocator mutex held.We had written this off as unfixable from Python — it isn't. The call is gated on
_cudnn_get_conv_benchmark_empty_cache(), and torch exposes the setter (torch._C._cudnn_set_conv_benchmark_empty_cache; verified honored byConv_miopen.cppin v2.13.0).Change
When more than one CUDA/HIP generation device is registered at session-processor startup, clear the flag. Benchmarking workspace blocks stay cached in the allocator for reuse instead of being returned to the driver — the same trade every peer-aware skip already makes, and #192's reclaimable-reserve crediting keeps the VRAM budget seeing them as available. Single-GPU installs are untouched. No-op guard for torch builds without the flag.
Testing
tests/backend/util/test_devices.py::test_disable_conv_benchmark_empty_cache_flips_torch_flag(flip + idempotence + CPU-build no-op).Rig test
With this + MIOPEN_FIND_MODE=FAST, a first-ever SDXL (or any new conv shape) run during an H3 denoise should no longer freeze at conv-heavy stages (VAE decode) until the peer's step boundary. This was the residual stall in the #159 testing (spy3/spy4/spy6 dumps).
🤖 Generated with Claude Code
https://claude.ai/code/session_014xFbHnmFTLsHE9e1PAdvme