Skip to content

fix: Metal backend performance on AMD discrete GPUs - #19527

Open
ai-janitor wants to merge 1 commit into
ggml-org:masterfrom
ai-janitor:fix/metal-amd-simd-reduction
Open

fix: Metal backend performance on AMD discrete GPUs#19527
ai-janitor wants to merge 1 commit into
ggml-org:masterfrom
ai-janitor:fix/metal-amd-simd-reduction

Conversation

@ai-janitor

@ai-janitor ai-janitor commented Feb 11, 2026

Copy link
Copy Markdown

Summary

Vibe coded a fix for Metal backend on AMD/Intel discrete GPUs. Performance went from 5.3 t/s → 60.4 t/s (11x) on AMD Radeon Pro 5300M. Output went from garbage → correct.

What Was Broken

Bug 1 (PRIMARY): Broken MTLDispatchTypeConcurrent cache coherency on AMD

  • memoryBarrierWithScope:MTLBarrierScopeBuffers doesn't flush AMD Radeon's L2 cache between concurrent compute dispatches
  • Layer N+1 reads stale output from layer N's buffer
  • Corruption scales with -ngl: more GPU layers = more garbage
  • Proof: GGML_METAL_CONCURRENCY_DISABLE=1 alone fixed all garbage output

Bug 2: Uncached PCIe reads on discrete GPUs

  • MTLResourceStorageModeShared uses uncached PCIe (3 MB/s)
  • Vulkan uses eHostCached → MoltenVK translates to Managed mode (3 GB/s, cached)
  • 1000x bandwidth gap

How I Vibe Coded This

Binary timestamp tracing

  • mach_absolute_time() at 4 instrumentation points
  • Identified 99.9% "overhead" was actually GPU execution time, not CPU sync

GPU timestamp profiling

  • Metal's GPUStartTime/GPUEndTime on command buffers
  • Found 12-20x shader performance gap between Metal and Vulkan on same hardware

Isolation testing

  • 9 separate Swift Metal compute tests to rule out simd hardware bugs
  • simd_sum, simd_max, simd_shuffle_down all work correctly in isolation
  • Corruption is concurrent dispatch, not broken arithmetic

Cross-backend validation

  • CPU: correct output, 22.1 t/s
  • Vulkan: correct output, 68.2 t/s
  • Metal (before): garbage output, 4.7 t/s
  • Tested at -ngl 0, 1, 2, 5+, 99 to identify corruption threshold

MoltenVK source analysis

  • Cloned entire MoltenVK repo for shader translation comparison
  • MVKCommandBuffer.mm:1063: MoltenVK uses MTLDispatchTypeSerial for all compute
  • MVKDevice.mm:790-800: Disables placementHeaps explicitly on AMD
  • Dumped translated MSL via MVK_CONFIG_SHADER_DUMP_DIR to compare memory access patterns

The Fix

Auto-detect non-Apple GPUs and force serial dispatch + managed buffers:

File Change
ggml-metal-context.m:133 Disable concurrency for !supports_gpu_family_apple7
ggml-metal-device.h:229 Add use_managed_buffers property
ggml-metal-device.m:790-800 Initialize based on !has_unified_memory
ggml-metal-device.m:1429,1512,1543 Conditional storage mode (3 allocation sites)
ggml-metal-device.m:1614,1649 Add didModifyRange sync after CPU writes

No configuration needed — auto-detects discrete GPUs and applies fix.

Environment variable overrides for testing:

  • GGML_METAL_MANAGED_BUFFERS_DISABLE / ENABLE
  • GGML_METAL_CONCURRENCY_DISABLE (now redundant — auto-disabled)

Benchmark Results

Qwen2.5-1.5B-Instruct Q4_K_M, -ngl 99 (all 28 layers on GPU)

AMD Radeon Pro 5300M (4GB, RDNA 1) on Intel Mac:

Backend Prompt t/s Gen t/s Output Notes
Vulkan 30.2 68.2 Correct Via MoltenVK, best option
Metal (fixed) 49.5 60.4 Correct 89% of Vulkan parity
Metal (before) 7.0 4.7 Garbage Concurrent dispatch broken
CPU (6 threads) 29.1 22.1 Correct Memory bandwidth bound

11x generation speedup, correct output.

Testing

  • Hardware: Intel Mac, AMD Radeon Pro 5300M 4GB (RDNA 1), Intel UHD 630
  • Correctness: Math, Knowledge, Reasoning, Long-form generation
  • Corruption pattern verified at -ngl 0, 1, 2, 5+, 99
  • F16 model also tested (rules out quantization as cause)

Why MoltenVK Works (Vulkan on Same Hardware)

  1. Serial dispatch by default — never opted into MTLDispatchTypeConcurrent
  2. Managed buffer mode — Vulkan's eHostCached translated to MTLResourceStorageModeManaged
  3. Explicit AMD workarounds — isAMDRDNAGPU() detection, disabled placementHeaps

See ggml-vulkan.cpp:4769 — Vulkan backend already has AMD + macOS workarounds.

Related Issues

  • #15228 — Metal3 + AMD Radeon (closed stale)
  • #4004 — Metal garbage on Intel Macs (closed unresolved)
  • #15846 — Vulkan subgroup arithmetic on AMD (fixed)

100% vibe coded. Works on my machine — AMD Radeon Pro 5300M.

🤖 Generated with Claude Code

hung @claude
fix: Metal backend performance on AMD discrete GPUs
e459796 
Fixes two critical issues with Metal backend on AMD/Intel discrete GPUs:

1. Auto-disable concurrent dispatch on non-Apple GPUs
   - MTLDispatchTypeConcurrent has broken memory barriers on AMD
   - Causes L2 cache coherency issues leading to corrupted output
   - Now auto-detects via supports_gpu_family_apple7 and uses serial dispatch

2. Use Managed buffer mode on discrete GPUs for cached PCIe reads
   - Shared mode uses uncached PCIe reads (~3 MB/s) on discrete GPUs
   - Managed mode enables cached reads (~3 GB/s) with explicit sync
   - Adds didModifyRange after CPU writes, synchronizeResource for reads
   - Improves generation speed from ~11 t/s to ~60 t/s on AMD Radeon Pro 5300M

Performance on AMD Radeon Pro 5300M:
- Before: 11.5 t/s generation (corrupted output with concurrent dispatch)
- After: 60+ t/s generation (correct output, 88% of Vulkan performance)

Environment variable overrides for testing:
- GGML_METAL_MANAGED_BUFFERS_DISABLE/ENABLE
- GGML_METAL_CONCURRENCY_DISABLE (still works)

Also adds bench_metal.sh with benchmark and conversation modes.

Co-Authored-By: Claude Sonnet 4.5 <noreply@anthropic.com>
@github-actions github-actions Bot added ggml changes relating to the ggml tensor library for machine learning Apple Metal https://en.wikipedia.org/wiki/Metal_(API) labels Feb 11, 2026
@taronaeo taronaeo added the vibe-coded Created with heavy use of LLM assistants, requires human verification label Feb 11, 2026
lvalen91 added a commit to lvalen91/llama.cpp that referenced this pull request Mar 30, 2026
@lvalen91 @claude
metal: AMD discrete GPU support (non-Apple-Silicon)
83395c7 
Cherry-picked fixes from ai-janitor's Metal AMD work:

1. Disable concurrent dispatch on non-UMA GPUs
   - MTLDispatchTypeConcurrent + barriers don't sync AMD L2 cache
   - Causes garbage output scaling with -ngl

2. Use MTLResourceStorageModeManaged for discrete GPUs
   - StorageModeShared uses uncached PCIe reads (3 MB/s)
   - Managed mode uses cached reads (3 GB/s) — 1000x bandwidth

3. Add scalar Flash Attention for non-simdgroup_mm GPUs
   - Enables FA on AMD/Intel without simd_matrix intrinsics
   - Fix simd_max(half) bug with -INF on RDNA

4. Use mul_mv_ext for large n on non-simdgroup_mm GPUs
   - Removes ne11<=8 upper bound for AMD, reducing threadgroup
     count from ~1.1M to ~280 for pp512

Tested on AMD Radeon PRO W6900X (32GB, RDNA2):
- Qwen3.5-27B Q4_K_M: 21 tok/s generation, correct output
- Qwen3-14B Q8_0: 29 tok/s gen, 138 tok/s prompt

Credit: github.com/ai-janitor (PRs ggml-org#19527, ggml-org#19600, ggml-org#19601)

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
@lvalen91 lvalen91 mentioned this pull request Mar 30, 2026
Closed
@kuroko3 kuroko3 mentioned this pull request Aug 12, 2026
Closed
lvalen91 added a commit to lvalen91/llama.cpp that referenced this pull request Aug 14, 2026
@lvalen91 @claude
metal: AMD discrete GPU support (non-Apple-Silicon)
bb3abba 
Cherry-picked fixes from ai-janitor's Metal AMD work:

1. Disable concurrent dispatch on non-UMA GPUs
   - MTLDispatchTypeConcurrent + barriers don't sync AMD L2 cache
   - Causes garbage output scaling with -ngl

2. Use MTLResourceStorageModeManaged for discrete GPUs
   - StorageModeShared uses uncached PCIe reads (3 MB/s)
   - Managed mode uses cached reads (3 GB/s) — 1000x bandwidth

3. Add scalar Flash Attention for non-simdgroup_mm GPUs
   - Enables FA on AMD/Intel without simd_matrix intrinsics
   - Fix simd_max(half) bug with -INF on RDNA

4. Use mul_mv_ext for large n on non-simdgroup_mm GPUs
   - Removes ne11<=8 upper bound for AMD, reducing threadgroup
     count from ~1.1M to ~280 for pp512

Tested on AMD Radeon PRO W6900X (32GB, RDNA2):
- Qwen3.5-27B Q4_K_M: 21 tok/s generation, correct output
- Qwen3-14B Q8_0: 29 tok/s gen, 138 tok/s prompt

Credit: github.com/ai-janitor (PRs ggml-org#19527, ggml-org#19600, ggml-org#19601)

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
@AndrewGDX

AndrewGDX commented Aug 15, 2026

Copy link
Copy Markdown

Can confirm that MacBook Pro 15 Vega 20 is also fixed thanks to this!

Results for Qwen3-4B-Instruct-2507-Q4_0.gguf:
Vulkan - now generates correct ~24 t/s
Metal - before generated broken ~4 t/s
CPU - before generated correct ~9 t/s

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@ggerganov ggerganov Awaiting requested review from ggerganov

At least 2 approving reviews are required to merge this pull request.

Assignees

No one assigned

Labels

Apple Metal https://en.wikipedia.org/wiki/Metal_(API) ggml changes relating to the ggml tensor library for machine learning vibe-coded Created with heavy use of LLM assistants, requires human verification

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

3 participants

@ai-janitor @AndrewGDX @taronaeo