Eval bug: GGML_OP_TOP_K falls back to CPU on HIP/ROCm above ~3–4K context — 6.4× token-generation loss on DeepSeek-V4-Flash

[stegorg]

Name and Version

Tested on two prebuilt gfx906 images:

mixa3607/llama.cpp-gfx906:b10187-rocm-7.2.4-mxxm-9fa3fef
mixa3607/llama.cpp-gfx906:b10207-rocm-7.2.4-9fa3fef

ROCm 7.2.4. Note that llama-server --version on these images reports version: 1 (81a8712) — the fork stamps its own build number and commit, so the upstream base is identifiable only from the image tag (b10187 / b10207).

Both builds reproduce identically. b10207 is ~20 upstream builds newer than b10187, so this is not a regression introduced in that window and not specific to the gfx906 fork.

Operating systems

Linux

GGML backends

HIP

Hardware

6 × AMD Instinct MI50, gfx906:sramecc+:xnack- (0x66a1), 32 GiB each
Uniform PCIe, 2 hops, all NUMA node 0, no XGMI
Host: 8-core AMD CPU, 503 GiB RAM

Models

bartowski DeepSeek-V4-Flash-0731-MXFP4 — 4 shards, 145.63 GiB, 284.33 B params, arch deepseek4.

Fully GPU-resident across all 6 cards. No CPU expert offload. deepseek4 is layer-split only on this backend (--split-mode tensor errors, row fails on gfx906).

Problem description & steps to reproduce

Command line

llama-bench \
  -m DeepSeek-V4-Flash-0731-MXFP4-00001-of-00004.gguf \
  -ngl 999 -sm layer -mg 0 -ts 0.95/1.00/1.00/0.95/1.00/1.10 \
  -fa 1 -b 2048 -ub 2048 -t 8 -ctk f16 -ctv f16 \
  -p 512 -n 64 -d 2048,2560,3072,4096 -r 2

Environment: HSA_OVERRIDE_GFX_VERSION=9.0.6, HIP_VISIBLE_DEVICES=0,1,2,3,4,5, ROCBLAS_USE_HIPBLASLT=0 (hipBLASLt ships no gfx906 Tensile kernels).

Problem description

Token generation is flat through depth 3072 and then collapses by 6.4× at depth 4096. Prompt processing degrades smoothly across the same range, with no corresponding discontinuity.

depth | pp512 (t/s) | tg64 (t/s) -- | -- | -- 2048 | 150.28 ± 14.07 | 15.61 ± 0.12 2560 | 150.93 ± 14.04 | 15.76 ± 0.16 3072 | 148.09 ± 12.95 | 15.78 ± 0.05 4096 | 123.67 ± 8.86 | 2.47 ± 0.03

Wider sweep (-p 2048 -n 128 -r 3) shows the post-cliff value is a flat floor rather than continuing decay: tg128 = 16.14 (d0), 16.11 (d1024), 15.68 (d2048), 2.38 (d4096), 2.32 (d8192). pp2048 over the same range: 231.86 → 215.43 → 210.03 → 170.97 → 153.94 → 129.58 (d16384), smooth throughout.

Confirmed on llama-server as well, not just llama-bench:

• 6-token prompt → tg 16.09 t/s
• 7001-token prompt → tg 2.37 t/s (pp 196.03 t/s, unaffected)

Root cause

GGML_SCHED_DEBUG=2 with llama-server -v. Op × backend assignment across the whole graph:

    189 CPU TOP_K
     10 CPU GET_ROWS

This is the sched_reserve dump — the worst-case graph at full declared context, not what runs per request. Those are the only CPU-assigned nodes there. Everything else — MUL_MAT, MUL_MAT_ID, ADD, CONCAT, MUL, SCALE, CONT, RMS_NORM, GET_ROWS, SET_ROWS, ROPE, FLASH_ATTN, LIGHTNING_INDEXER, DSV4_HC_* — is on ROCm0–5.

sched_reserve: graph splits = 49 for a 6-device layer split (≈6 expected). 22 of the 49 splits are CPU splits, each with exactly one input, and it is always the lightning-indexer score tensor:

## SPLIT #2:  CPU # 1 inputs : [lid_score_masked-2  (32K)]
## SPLIT #4:  CPU # 1 inputs : [lid_score_masked-4  (32K)]
...
## SPLIT #47: CPU # 1 inputs : [lid_score_masked-42 (32K)]

So the graph leaves the GPUs ~22 extra times per eval to run TOP_K on the host and copy indices back. lid_score_masked scales with token count (32K → 512K between the two reserve passes, tracking embd 16K → 256K), so at decode the transferred tensor scales with n_kv.

This accounts for every observed symptom:

• pp survives — one round-trip amortised across a 2048-token batch
• tg collapses — 22 round-trips per single token
• hard threshold, then a flat floor — the op is either accepted or rejected, with no middle ground
• immune to -b / -ub — -b 512 -ub 512 reproduces identically
• GPU utilisation sits near 0% with sporadic 11–16% blips during the collapse
• ROCm_Host compute buffer size = 5736.68 MiB at load, on an otherwise fully-resident model

Under AMD_LOG_LEVEL=1 the collapsed state shows a tight retry loop (~4 ms apart) of 32 MiB host allocations:

Can't allocate memory size - 0x02000000 bytes!
rocdevice.cpp:1878: Failed creating memory
memory.cpp:357: Video memory allocation failed!

with dmesg concurrently flooding amdgpu: init_user_pages: Failed to get user pages: -1. This is the host staging traffic for the CPU excursions, not an independent fault — there is 5–8.5 GiB free VRAM on every card and 380 GiB available host RAM at the time.

Confirmation: the op switches backend between the two depths

Per-request scheduler dumps from a single llama-server -v process (GGML_SCHED_DEBUG=2, --ctx-size 16384), log sliced by request:

prompt_n = 2501 → tg 15.15 t/s. All TOP_K nodes on GPU:

     12 ROCm0 TOP_K      16 ROCm1 TOP_K      12 ROCm2 TOP_K
     16 ROCm3 TOP_K      12 ROCm4 TOP_K      16 ROCm5 TOP_K

prompt_n = 6001 → tg 2.35 t/s. Most TOP_K nodes moved to CPU:

    105 CPU TOP_K
      6 ROCm0 TOP_K       8 ROCm1 TOP_K       6 ROCm2 TOP_K
      8 ROCm3 TOP_K       6 ROCm4 TOP_K       8 ROCm5 TOP_K

Same process, same weights, same flags — only the prompt length differs. The fallback is partial: 42 nodes stay on GPU while 105 move, consistent with a per-node ne[0] threshold rather than a whole-op rejection.

Two further checks rule out competing explanations:

• The AMD_LOG_LEVEL=1 32 MiB allocation-failure count is identical at depth 3072 and 4096 (25368 each), so the host-allocation retry loop is a fixed load-time artifact and not the cause of the collapse.
• Cutting --ctx-size from 204800 to 32768 does not move the threshold — it tracks actual n_kv, not declared context.

Root cause in upstream source: CUB gate excludes HIP

ggml_cuda_op_top_k is present as a defined symbol in the HIP backend:

$ nm -D --defined-only /app/libggml-hip.so | grep -i top_k
00000000032a2590 T _Z18ggml_cuda_op_top_kR25ggml_backend_cuda_contextP11ggml_tensor

The kernel is compiled in. The fallback comes from supports_op in ggml/src/ggml-cuda/ggml-cuda.cu:

c

case GGML_OP_TOP_K:
case GGML_OP_ARGSORT:
#ifndef GGML_CUDA_USE_CUB
    return op->src[0]->ne[0] <= 1024;
#else
    return true;
#endif

and GGML_CUDA_USE_CUB is defined in ggml/src/ggml-cuda/common.cuh as:

c

#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11070
#    define GGML_CUDA_USE_CUB
#endif

So CUDA builds with CUDART >= 11.7 return true unconditionally and never fall back. HIP and MUSA are explicitly excluded from the CUB path and are therefore subject to the size gate. CUB is NVIDIA-only; there is no in-tree rocPRIM / hipCUB equivalent wired up for the HIP backend.

This is not gfx906-specific and not fork-specific — it affects every AMD GPU on the HIP backend, and MUSA. Any ROCm user running DeepSeek-V4-Flash (or any other architecture whose graph contains a large TOP_K) past a few thousand tokens of context will hit it.

Note one unexplained detail: master's gate is ne[0] <= 1024, but TOP_K ran entirely on GPU here at n_kv = 2501. So either ne[0] on the indexer's TOP_K input is not n_kv directly (chunking or stride in the deepseek4 graph), or the tested builds carry a different threshold than current master. The CUB-vs-no-CUB split is the operative cause either way — the op demonstrably switches backends as n_kv grows — but the exact threshold value in these builds is unconfirmed.

Related HIP-backend strings, confirming the non-CUB bitonic path is what gets used:

_Z28argsort_f32_i32_cuda_bitonicPKfPiii15ggml_sort_orderP12ihipStream_t
_Z20ggml_cuda_op_argsortR25ggml_backend_cuda_contextP11ggml_tensor
_ZL17k_argsort_f32_i32IL15ggml_sort_order0EEvPKfPiii
/build/llamacpp/ggml/src/ggml-cuda/argsort.cu
hipMemcpy2DAsync(dst_d, k * sizeof(int), tmp_dst, ncols * sizeof(int), k * sizeof(int), nrows, hipMemcpyDeviceToDevice, stream)

Suggested fix

Route the HIP backend through hipCUB / rocPRIM for GGML_OP_TOP_K and GGML_OP_ARGSORT, or extend the non-CUB bitonic path to handle larger ncols via a multi-pass / global-memory sort. Either removes the CPU excursion entirely.

Ruled out

• Compute-buffer / ubatch pressure — -b 512 -ub 512 reproduces identically
• Device VRAM exhaustion — 5–8.5 GiB free per card during the collapse (~46 GiB total)
• Host RAM exhaustion — 503 GiB total, 380 GiB available, swap unused
• RLIMIT_MEMLOCK — --ulimit memlock=-1:-1 --cap-add IPC_LOCK changes nothing; --privileged already grants CAP_IPC_LOCK
• GGML_CUDA_NO_VMM=1 — no effect
• Mid-run graph buffer reallocation — no realloc/gallocr output after either a short or a long request
• Missing lightning-indexer kernel — GGML_OP_LIGHTNING_INDEXER runs on ROCm here (189 nodes, none on CPU), so cuda : CUDA GGML_OP_LIGHTNING_INDEXER implementation (generic vector kernel + wmma kernel) #25545 is present and working
• Declared context size — cutting --ctx-size from 204800 to 32768 does not move the threshold; it tracks actual depth
• Build version — identical on b10187 and b10207

GGML_CUDA_NO_PINNED=1 is not a workaround; the server still collapses with it set.

Impact

DeepSeek-V4-Flash is unusable past ~3K context on ROCm despite full GPU residency, on hardware that otherwise sustains ~16 t/s. There is no runtime or configuration mitigation — thread count is capped by an 8-core host, batch size is irrelevant, and no newer prebuilt image changes the behaviour.

Because the gate is #ifndef GGML_CUDA_USE_CUB rather than anything architecture-specific, the impact is backend-wide: all AMD GPUs on HIP (CDNA and RDNA alike) and all MUSA devices. CUDA is unaffected.

Possibly related

• cuda : CUDA GGML_OP_LIGHTNING_INDEXER implementation (generic vector kernel + wmma kernel) #25545 — CUDA GGML_OP_LIGHTNING_INDEXER implementation (present and working here)
• Eval bug: DeepSeek-V4-Flash lightning-indexer materialises a ctx-proportional (4 GiB/layer) score buffer #25468 — DeepSeek-V4-Flash lightning-indexer ctx-proportional score buffer
• Misc. bug: NVIDIA CUB implementations of GGML argsort and top-k cause CUDA graph capture failure for some tensor shapes #21682 — CUB argsort / top-k shape issues on CUDA

First Bad Commit

No response

Relevant log output

Logs

[pwilkin]

@Geramy want to take a stab?

[kubakomu]

I noticed this issue when during inference of DS4 on my R9700 cards, the CPU had a very high utilization.
I did some profiling and basically asked AI to fill in the gaps, related to what was found here.
kubakomu@42b262a
Above commit improved my TG from about ~11 to more than ~24, so about a 100% which I could hardly believe.

Since its AI generated and I am not super aware with this codebase I dont think its worth trying to make a PR out of it, but maybe it could be interesting for someone as a reference.

[Geramy]

@pwilkin yeah I’ll take a stab at it.

[shruggr]

Very similar issue happening on 4 x R9700 GPUs with unsloth UD-IQ3_XSS. TG is falling back to CPU with 2475% CPU utilization (24 core/48 thread Threadripper).

[shruggr]

I applied @kubakomu fix, and it fixed the CPU utilization, but unfortunately didn't increase my TG performance any.

[Geramy]

I'm downloading the model now and I'll start looking into this, I was a little bit busy this weekend.

[Geramy]

@pwilkin Can confirm this is happening. The op is coming back as unsupported, so it falls off the GPU — I'm seeing 189 TOP_K nodes end up on CPU, and it's because ne[0]=16384, well over the 1024 limit. Still digging into it.

UPDATE:
One GPU doesn't seem to see any slowdowns at all from top-k being on the CPU. But it could be the configuration of the machine, I am going to test tensor split now.

[Geramy]

I've split the load between my gfx1151 and gfx1201, the 1201 is connected over a pcie 4.0 x4 lane but the card runs at pcie 5.0 x16 which doesn't really matter or mean anything here since the real bandwidth is x4. I am not seeing a steep tps / pp's fall off when doing tensor split myself. I am running unsloth/DeepSeek-V4-Flash-GGUF:UD-IQ1_S I will download the exact same model, I think maybe I have enough usable VRAM / UFM RAM. I am seeing 200% CPU usage but it isn't effecting it on my 2 GPU tensor split system.

The next step is probably introducing a new kernel on rocm for top-k so that it can go beyond 1024, I will start working on this.

[Geramy]

@stegorg please try this PR #26493 I don't believe this is the most correct fix and its going to take a lot of performance testing of main against this PR to guarantee there are no regressions on other models, but its a start. The better way would be to implement a hip equivalent of the CUB argsort / top-k

[stegorg]

@stegorg please try this PR #26493 I don't believe this is the most correct fix and its going to take a lot of performance testing of main against this PR to guarantee there are no regressions on other models, but its a start. The better way would be to implement a hip equivalent of the CUB argsort / top-k

Worked like a charm! TY

[stegorg]

Performance stays healthy all the way through large context now.