Eval bug: GLM-5.2 terrible performance on ROCm/HIP after Indexer PR #25407: prefill ~6x slower, load time ~40x longer

[forgit]

Name and Version

88bfee1

Operating systems

Linux

GGML backends

HIP

Hardware

8x AMD MI300X (gfx942), 2x AMD EPYC 9555 (256 threads total)

Models

unsloth/GLM-5.2-GGUF:Q2_K_XL (https://huggingface.co/unsloth/GLM-5.2-GGUF)

Problem description & steps to reproduce

Summary

Commit 88bfee142 / PR #25407 ("model: add GLM 5.2 Indexer support") causes a severe
regression for GLM-5.2 on ROCm (HIP). Prefill drops ~6x (using 2 of 8 GPUs instead of all 8)
and model load takes ~40x longer. The regression is still present on the latest commit
272700b36. The parent commit 95a923a64 works correctly.

Environment

• OS: Rocky Linux 9.8 (x86_64)
• Hardware: 8x AMD MI300X (gfx942)
• ROCm 7.2
• llama.cpp: regression in 88bfee142, reproduced on 272700b36
• Build: -DGGML_HIP=ON -DAMDGPU_TARGETS=gfx942 -DGGML_HIP_GRAPHS=ON -DGGML_HIP_RCCL=ON -DGGML_HIP_MMQ_MFMA=ON
• KV cache: kv-unified = false (separate KV per slot)

To Reproduce

1. Build llama.cpp as above and start llama-server.
2. Load a GLM-5.2 GGUF (e.g. unsloth/GLM-5.2-GGUF:Q2_K_XL or Q4_K_XL, parallel=5).
3. Send a 32k-token prompt and measure prompt processing and decode.

Expected behavior

As on the parent commit 95a923a64: model loads in ~30 s, and prefill uses all 8 GPUs at 100%.

Actual behavior

• On all builds, the text_model stage reaches value 1.0 in ~13 s. On 95a923a64 the model finishes loading in under 30 s. On 88bfee142 / 272700b36, after the n_ctx is not divisible by n_seq_max - rounding down warning the server takes 17-41 minutes before load_model: initializing / model loaded.
• Prefill uses only ~2 of 8 GPUs and is ~6x slower.
• Decode is ~1.5-2x slower (26.7 t/s on 95a923a64 vs 12-17 t/s on the regressed builds).

Measurements (same hardware, same model, 32k-token prompt)

On all three builds, text_model reaches value 1.0 in ~13 s; the difference is the time between the n_ctx is not divisible warning and the model finishing loading.

Build	Time to model loaded	Prefill @32k	Decode
95a923a64 (parent, working)	< 30 s	1197 t/s	~26.7 t/s
88bfee142 (indexer PR)	~17 min	202 t/s	12-17 t/s
272700b36 (latest)	~41 min	202 t/s	~17 t/s

Relevant logs

Working (95a923a64) — finishes loading quickly:

0.13.333.433 W llama_context: n_ctx is not divisible by n_seq_max - rounding down to 3600384
0.41.097.610 I srv    load_model: initializing, n_slots = 12, n_ctx_slot = 300032, kv_unified = 'false'
0.41.111.773 I srv  llama_server: model loaded

Regressed (272700b36) — hangs 41 minutes after the same warning:

0.13.302.538 W llama_context: n_ctx is not divisible by n_seq_max - rounding down to 3600384
41.49.430.952 I srv    load_model: initializing, n_slots = 12, n_ctx_slot = 300032, kv_unified = 'false'
41.49.447.986 I srv  llama_server: model loaded

First Bad Commit

88bfee1

Relevant log output

Logs

llama-server invocation (from the router):

/home/.../build/bin/llama-server
  --chat-template-kwargs {"preserve_thinking":true,"reasoning_effort":"high"}
  --host 127.0.0.1 --jinja
  --min-p 0.01
  --no-mmap
  --port 46629
  --reasoning-preserve
  --temperature 1.0 --top-k 0 --top-p 0.95
  --alias unsloth/GLM-5.2-GGUF:Q2_K_XL
  --batch-size 8192 --ctx-size 3600000 --cont-batching
  --cache-ram 150000
  --device ROCm0,ROCm1,ROCm2,ROCm3,ROCm4,ROCm5,ROCm6,ROCm7
  --flash-attn on
  --hf-repo unsloth/GLM-5.2-GGUF:Q2_K_XL
  --no-kv-unified --parallel 12
  --seed 42 --slot-prompt-similarity 0.5
  --ubatch-size 1024

Working (95a923a64) — finishes loading quickly and prefills on all 8 GPUs:

[46629] 0.13.333.433 W llama_context: n_ctx is not divisible by n_seq_max - rounding down to 3600384
[46629] 0.41.097.610 I srv    load_model: initializing, n_slots = 12, n_ctx_slot = 300032, kv_unified = 'false'
[46629] 0.41.111.773 I srv  llama_server: model loaded
[46629] 0.44.693.431 I slot print_timing: id 11 | task 0 | prompt processing, n_tokens =   8192,  progress = 0.24, t =   3.52 s / 2328.73 tokens per second
[46629] 0.50.070.299 I slot print_timing: id 11 | task 0 | prompt processing, n_tokens =  16384,  progress = 0.48, t =   8.89 s / 1842.00 tokens per second
[46629] 1.08.529.072 I slot print_timing: id 11 | task 0 | prompt processing, n_tokens =  32768,  progress = 0.97, t =  27.35 s / 1197.95 tokens per second
[46629] 1.20.267.783 I slot print_timing: id 11 | task 0 | n_decoded =    100, tg =  26.71 t/s, tg_3s =  26.71 t/s

Regressed (88bfee142 and 272700b36) — same warning, then hangs 41 minutes before loading, and prefill is ~6x slower:

0.13.302.538 W llama_context: n_ctx is not divisible by n_seq_max - rounding down to 3600384
41.49.430.952 I srv    load_model: initializing, n_slots = 12, n_ctx_slot = 300032, kv_unified = 'false'
41.49.447.986 I srv  llama_server: model loaded
[36635] 42.08.891.775 I slot print_timing: id 11 | task 0 | prompt processing, n_tokens =   8192,  progress = 0.24, t =  19.38 s / 422.66 tokens per second
[36635] 42.41.910.454 I slot print_timing: id 11 | task 0 | prompt processing, n_tokens =  16384,  progress = 0.48, t =  52.40 s / 312.67 tokens per second
[36635] 43.29.534.737 I slot print_timing: id 11 | task 0 | prompt processing, n_tokens =  24576,  progress = 0.73, t = 100.02 s / 245.70 tokens per second
[36635] 44.31.467.204 I slot print_timing: id 11 | task 0 | prompt processing, n_tokens =  32768,  progress = 0.97, t = 161.96 s / 202.32 tokens per second
[36635] 44.47.145.061 I slot print_timing: id 11 | task 0 | n_decoded =    100, tg =  17.35 t/s, tg_3s =  17.35 t/s

[fairydreaming]

That may be related to #26399
Another reason is that lightning indexer uses vector kernel on HIP that is likely much slower than WMMA kernel. I don't have access to any AMD GPUs, so I'm not able to work on this. Shouldn't be hard, needs something like:

#elif defined(GGML_USE_HIP)
#include <rocwmma/rocwmma.hpp>
namespace wmma = rocwmma;
#endif // !defined(GGML_USE_HIP)

but without hardware access I'm not working on this.

[forgit]

Thanks, happy to test any ROCm fix if someone writes one.

[fairydreaming]

Thanks, happy to test any ROCm fix if someone writes one.

@forgit You can try kubakomu@42b262a right now. This shall fix slow top-k.

[Geramy]

rocWMMA Has been deprecated as far as I know in exchange for MMA which produces better results and is more flexible.

[fairydreaming]

rocWMMA Has been deprecated as far as I know in exchange for MMA which produces better results and is more flexible.

@Geramy If you feel the calling I'm not stopping you. I don't feel strong enough for multi-platform MMA fights.

[forgit]

Thanks, happy to test any ROCm fix if someone writes one.

@forgit You can try kubakomu@42b262a right now. This shall fix slow top-k.

It helped with load time (it loads almost right away), but on prefill only two GPUs are used at a time instead of all 8. Decode improved to ~20t/s.

[fairydreaming]

Thanks, happy to test any ROCm fix if someone writes one.

@forgit You can try kubakomu@42b262a right now. This shall fix slow top-k.

It helped with load time (it loads almost right away), but on prefill only two GPUs are used at a time instead of all 8. Decode improved to ~20t/s.

@forgit There's now a PR that uses hipCUB to implement ggml_top_k(): #26592 so another thing to test.

[Geramy]

@Geramy If you feel the calling I'm not stopping you. I don't feel strong enough for multi-platform MMA fights.

@fairydreaming I didnt see this, that's a good one haha, I walked right into that lol.

[forgit]

@fairydreaming @Geramy Tested PR #26592 on 88bfee1: prefill unchanged (~200 t/s, still only 1–2 GPUs busy in a sequential wave), decode improved to ~24 t/s.

[fairydreaming]

@forgit how about the current master? (today ROCm/HIP TOP_K kernel was merged)