bug: Metal backend produces garbage output on AMD discrete GPUs (Radeon Pro 5300M)

Bug Description

The Metal backend produces incorrect (garbage) output on AMD discrete GPUs while benchmarks show plausible performance numbers. The same model works correctly with the Vulkan backend (via MoltenVK) on the same GPU, and with the Metal backend when using CPU-only (-ngl 0).

Environment

GPU: AMD Radeon Pro 5300M (4 GB VRAM, discrete)
Secondary GPU: Intel UHD Graphics 630 (integrated)
macOS: 26.2 (Build 25C56)
llama.cpp: commit 4d688f9 (current master)
Metal family: MTLGPUFamilyMetal3 (5001), MTLGPUFamilyCommon3 (3003)
Metal props: simdgroup reduction = true, simdgroup matrix mul. = false, has unified memory = false, use shared buffers = false

Reproduction

# Build with Metal
cmake -B build -DGGML_METAL=ON -DGGML_METAL_EMBED_LIBRARY=ON
cmake --build build -j -t llama-completion

# Metal GPU — BROKEN (garbage output)
./build/bin/llama-completion \
  -m Qwen2.5-1.5B-Instruct-Q4_K_M.gguf \
  -ngl 99 -n 32 -s 42 -no-cnv \
  -p "Explain gravity in one sentence:"
# Output: "Explain gravity in one sentence:@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@@"

# CPU only — WORKS
./build/bin/llama-completion \
  -m Qwen2.5-1.5B-Instruct-Q4_K_M.gguf \
  -ngl 0 -n 32 -s 42 -no-cnv \
  -p "Explain gravity in one sentence:"
# Output: "Gravity is the force that attracts two masses towards each other..."

Observations

Affects both Q4_K_M and F16 model formats — not quant-type specific
llama-bench reports plausible performance numbers (pp128: ~100 t/s, tg32: ~63 t/s) but the computation is incorrect
Vulkan backend (via MoltenVK → Metal → same GPU) produces correct output at ~68 t/s tg32
The GPU is discrete with no unified memory (has_unified_memory = false, use_shared_buffers = false)
simdgroup matrix mul. = false (no Apple7 GPU family), so flash attention falls back to CPU and matrix multiply uses the legacy kernel_mul_mv path

Possible Causes

Buffer management: Discrete AMD GPUs use private buffers with explicit CPU↔GPU copies (unlike Apple Silicon's shared/unified memory). Data may not be correctly transferred.
KV cache handling: kv_unified = false in the logs suggests a different KV cache path for non-unified memory GPUs.
Kernel correctness: The legacy mul_mv kernels may have assumptions that hold on Apple Silicon but break on AMD (e.g., memory ordering, simdgroup behavior).

Additional Context

This was discovered while benchmarking Metal vs Vulkan performance on AMD. The benchmarks appeared reasonable, masking the fact that Metal was producing incorrect results. This affects any AMD discrete GPU on macOS using the Metal backend.

🤖 Generated with Claude Code

Confirmed on ggml-org/llama.cpp HEAD (f488429).

Additional bisection findings:

ngl=1: correct output (1 layer on GPU works)
ngl=2: garbage starts (bug appears at 2nd GPU layer)
ngl=99, -n 1: first generated token is plausible
ngl=99, -n 2+: all garbage tokens (e.g. @@@@@)
ngl=0 (CPU only): fully correct

This suggests the prompt processing (PP) path works, but the decode (TG) path corrupts starting from the 2nd token. The KV cache write during PP appears fine, but KV cache read-back during subsequent token generation returns wrong data on discrete GPUs with private buffers.

metal : fix incorrect output on AMD discrete GPUs

Root cause identified and fix submitted: #19567

Root cause: MTLDispatchTypeConcurrent compute encoders + memoryBarrierWithScope:MTLBarrierScopeBuffers don't properly synchronize on AMD discrete GPUs. The barriers work on Apple Silicon but AMD hardware doesn't honor them for concurrent dispatches.

Fix: Disable concurrent encoders on non-UMA GPUs by gating use_concurrency on has_unified_memory. Also fixes an NSMakeRange buffer overwrite in memset_tensor on the private buffer path.

Verified on AMD Radeon Pro 5300M — all previously garbage outputs now correct, no perf regression.

Summary

Bug confirmed on AMD Radeon Pro 5300M (discrete GPU, non-UMA architecture).

Symptom: Garbage output (@@@@@) when using -ngl > 1

Root cause: MTLDispatchTypeConcurrent memory barriers don't properly synchronize on discrete
GPUs. Works on Apple Silicon (UMA) but fails on AMD/Intel discrete.

Immediate Workaround

GGML_METAL_CONCURRENCY_DISABLE=1 ./llama-cli ...

Forces serial dispatch - produces correct output.

Proposed Fix

PR #19567 addresses this by gating use_concurrency on has_unified_memory (same pattern already used
for use_shared_buffers).

3-line change, tested on AMD discrete GPU, no regression on Apple Silicon.

---
For maintainers: If someone wants to vet the technical approach, the complete analysis and solution
are in #19567.

Happy to provide additional testing or details if needed.

Second bug found: Flash Attention vec kernel on AMD

After fixing the concurrent encoder issue (#19567), a second correctness bug surfaces when enabling the FA vec kernel on AMD discrete GPUs.

Background

kernel_flash_attn_ext_vec is currently gated behind has_simdgroup_mm, which AMD GPUs lack (no Apple7 family). The vec kernel only requires has_simdgroup_reduction — which AMD supports. Changing the gate to has_simdgroup_reduction enables GPU-side flash attention on AMD, but produces garbage output.

Root cause

simd_max(half) returns incorrect results on AMD RDNA when operands include half-precision -INF (0xFC00). The -INF block skip optimization at line 6078 of ggml-metal.metal:

if (simd_max(sm[tiisg]) <= -MAXHALF) {
    continue;
}

...incorrectly skips chunks containing valid attention positions when the mask has -INF entries (causal masking). This causes the softmax denominator S = 0, producing 1/S = inf → garbage output.

Fix

Cast half to float before simd_max:

if (simd_max(float(sm[tiisg])) <= -65504.0f) {
    continue;
}

Plus gate change in ggml-metal-device.m — vec kernels check has_simdgroup_reduction instead of has_simdgroup_mm.

Why tests didn't catch it

init_tensor_kq_mask() generates random uniform [-1, 1] mask values. For small tensors (kv=256), zero -INF blocks are generated — the skip path was never exercised. Added causal mask test coverage that reproduces the bug.

Results (AMD Radeon Pro 5300M, 10 reps)

Test	Metal FA=0	Metal FA=1 (with fix)	Vulkan FA=0
pp128	98.74 ± 0.61	92.24 ± 0.99	87.17 ± 0.05
tg32	62.01 ± 0.67	64.19 ± 0.35	65.74 ± 1.52

1284/1298 FA tests pass (14 failures are hsk=576 precision — pre-existing, unrelated).

Branch

Fix is on my fork if anyone wants to review or submit a PR:
https://github.com/ai-janitor/llama.cpp/tree/fix/metal-amd-fa-vec-simd-max

3 files changed, 57 insertions, 7 deletions.

fix(ggml-org#19563): correct buffer range calculation in ggml_metal_b…

Workaround: use Vulkan backend via MoltenVK (Intel Mac + AMD discrete GPU)

Confirmed working on: MacBook Pro, Intel x86_64, AMD Radeon Pro 5500M 8GB, macOS 15.5.

The Metal direct path fails with kIOAccelCommandBufferCallbackErrorTimeout on every inference call — even a 1-token prompt. This is the MTLDispatchTypeConcurrent concurrency bug described in this issue. GGML_METAL_CONCURRENCY_DISABLE=1 does not resolve it.

Working path: Vulkan backend via MoltenVK

brew install molten-vk vulkan-headers vulkan-loader glslang shaderc

Build llama.cpp from source with Vulkan enabled and Metal disabled:

cmake -B build \
  -DGGML_METAL=OFF \
  -DGGML_VULKAN=ON \
  -DCMAKE_BUILD_TYPE=Release \
  -DCMAKE_OSX_SYSROOT=/Library/Developer/CommandLineTools/SDKs/MacOSX15.5.sdk \
  "-DCMAKE_CXX_FLAGS=-isystem /Library/Developer/CommandLineTools/SDKs/MacOSX15.5.sdk/usr/include/c++/v1"

CPLUS_INCLUDE_PATH=/Library/Developer/CommandLineTools/SDKs/MacOSX15.5.sdk/usr/include/c++/v1 \
SDKROOT=/Library/Developer/CommandLineTools/SDKs/MacOSX15.5.sdk \
cmake --build build --target llama-server

The SDK path flags are required because the Command Line Tools C++ headers outside the SDK are incomplete (only 4 files in /usr/local/include/c++/v1/).

Run with Flash Attention disabled — there's a separate correctness bug on this path:

VK_ICD_FILENAMES=/usr/local/Cellar/molten-vk/1.4.1/etc/vulkan/icd.d/MoltenVK_icd.json \
./build/bin/llama-server -m model.gguf --n-gpu-layers 99 -c 4096 -fa 0

The AMD GPU is detected and used for inference. All 32 layers offloaded. Inference is fast and stable.

Note: the Homebrew llama.cpp bottle on Intel Mac has no GPU backend at all (warning: no usable GPU found) — a source build is required.

This issue was closed because it has been inactive for 14 days since being marked as stale.

feat: v0.3.1 — image gen + TTS slots live, homebrew tap repo, Apple I…

For folks arriving from search: this issue and the fix PR #19567 both went stale-closed, not resolved — but the root cause is real and fixable, in two layers on AMD discrete Metal:

ggml enables simdgroup_reduction + bfloat on any MTLGPUFamilyMetal3 device (incl. AMD discrete on Intel Mac); those kernels compile but compute wrong. Gating them to Apple7/Apple6 (with GGML_METAL_FORCE_* opt-outs) restores correct output.
MTLDispatchTypeConcurrent ordering is unreliable on non-UMA AMD → GGML_METAL_CONCURRENCY_DISABLE fixes cross-call nondeterminism; a bounded staging-buffer pool fixes a sustained-load GGML_ASSERT SIGABRT.

We run these as build-time patches in production on a Vega II (32 GB) — chat + embeddings + rerank, correct output, ~4 tok/s on a 7B. Patches + write-up: https://github.com/Cerid-AI/quenchforge/tree/main/patches (disclosure: I maintain it). If a maintainer is open to it, I'm glad to send the gating patch as a fresh, minimal PR behind opt-out env flags.

docs: public-facing wave64 GCN/Vega writeup (WAVE64.md)

fix(engine): the macOS x64 binary was never CPU-only

Swap darwin-x64 Metal build for Vulkan+MoltenVK to fix AMD Radeon output

Navigation Menu

Repository navigation

bug: Metal backend produces garbage output on AMD discrete GPUs (Radeon Pro 5300M) #19563

Bug Description

Environment

Reproduction

Observations

Possible Causes

Additional Context

Summary

Immediate Workaround

Second bug found: Flash Attention vec kernel on AMD

Background

Root cause

Fix

Why tests didn't catch it

Results (AMD Radeon Pro 5300M, 10 reps)

Branch

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

bug: Metal backend produces garbage output on AMD discrete GPUs (Radeon Pro 5300M) #19563

Description

Bug Description

Environment

Reproduction

Observations

Possible Causes

Additional Context

Activity

ai-janitor commented on Feb 12, 2026

ai-janitor commented on Feb 12, 2026

ai-janitor commented on Feb 12, 2026

Summary

Immediate Workaround

ai-janitor commented on Feb 12, 2026

Second bug found: Flash Attention vec kernel on AMD

Background

Root cause

Fix

Why tests didn't catch it

Results (AMD Radeon Pro 5300M, 10 reps)

Branch

odinkirk commented on Mar 22, 2026

github-actions commented on May 7, 2026

sunrunnerfire commented on May 31, 2026

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions