Eval bug:

[torhve]

Name and Version

v0.4.1 and also self compiled crashes on ROCm R9700: mmf.cuh:636: GGML_ASSERT(stride_col_y % 2 == 0) on first inference.

also saw mmvf.cu:423: GGML_ASSERT(stride_col_y % 2 == 0)

Operating systems

Linux

GGML backends

HIP

Hardware

1x Radeon R9700 (gfx1201)

Models

No response

Problem description & steps to reproduce

crash on first inference on qwen 3.7 27b q6 k xl

First Bad Commit

No response

Relevant log output

Logs

[Anbeeld]

Pushed v0.4.2 branch with a number of fixes for ROCm, please try it.

[torhve]

swapping out q8_0 to kvarn8 (or kvarn4)
ggml_backend_cuda_buffer_type_alloc_buffer: allocating 51032.79 MiB on device 0: cudaMalloc failed: out of memory

[Anbeeld]

Pushed v0.4.3 with what should be a fix, please try it.

[torhve]

KVarN KV Cache Comparison: 27B Q4_K_XL on 32GB R9700 (ROCm)

Setup:

• GPU: AMD Radeon R9700 (32GB VRAM)
• Model: Qwen3.6-27B-UD-Q4_K_XL.gguf (17GB)
• Binary: beellama.cpp (KVarN fork) on ROCm 10.0
• All tests: -ngl 99, --spec-type none, -fa on, --kv-unified, --fit-ctx 64000 --fit-target 0

Results

KV Cache Type	Max ctx (auto-fit)	VRAM Used	Notes
q4_0 / q4_0	262,144 (full)	22.32 GB	Full training context, ~9.5 GB headroom
kvarn4 / kvarn4	134,400	31.84 GB (~full)	Only half the context, VRAM fully consumed

Observations

1. q4_0 outperforms both KVarN variants for this model. It reaches the full 262K training context with 9.5 GB to spare, while kvarn4 fills the entire 32GB VRAM at only 134K context.

2. KVarN is more memory-hungry than q4_0 at equivalent context lengths.

3. KVarN is "target-context-only" — auxiliary contexts (MTP draft) fall back to default types ?

4. Flash attention (-fa on) works with all three types, no silent fallback observed.

5. With MTP enabled, the max usable context drops further (kvarn8 + MTP maxed at ~64K before MTP draft OOM'd).

[Anbeeld]

Please send logs from your runs. I don't have AMD GPU and can't examine the problems from vibe analysis.

[raufaser]

Maybe I can provide releavnt input. I started from the same error as OP in post 1. Ending up having a different error in preview v0.4.3

Name and Version

Started with: v0.4.1 getting -> mmvf.cu:423: GGML_ASSERT(stride_col_y % 2 == 0)
Now: v0.4.3 -> fattn-mma-f16.cuh:1846: ERROR: HIP kernel flash_attn_ext_f16 has no device code compatible with HIP arch 1300.

Operating systems

Windows 11

GGML backends

HIP

Hardware

1x AMD Radeon RX 7900 XTX

Models

Qwen3.6-27B-Fable-Fusion-711-Uncensored-Heretic-NM-DAU-NEO-MAX-MTP-GGUF-Q4_K_M

Problem description & steps to reproduce

crash when using kvarn quants.

Relevant log output

v4.3.0 KV q8_0 q8_0 - no problems

> .\Downloads\beellama-preview-v0.4.3-bin-win-hip-radeon-x64\llama-server.exe -m .lmstudio\models\DavidAU\Qwen3.6-27B-Fable-Fusion-711-Uncensored-Heretic-NM-DAU-NEO-MAX-MTP-GGUF\Qwen3.6-27B-Fable-Fus-711-UnHeretic-NM-DAU-NEO-MAX-NEO-MTP-Q4_K_M.gguf -ngl 99 -fa 1 -ctk q8_0 -ctv q8_0
HIP Library Path: C:\WINDOWS\SYSTEM32\amdhip64_7.dll
0.00.050.877 I cmn  common_param: common_params_print_info: verbosity = 3 (adjust with the `-lv N` CLI arg)
0.00.128.205 W srv  llama_server: -----------------
0.00.128.211 W srv  llama_server: CORS is set to allow all origins ('*') and no API key is set
0.00.128.211 W srv  llama_server: this can be a security risk (cross-origin attacks)
0.00.128.213 W srv  llama_server: more info: https://github.com/ggml-org/llama.cpp/pull/25655
0.00.128.213 W srv  llama_server: -----------------
0.00.142.051 I srv    load_model: loading model '.lmstudio\models\DavidAU\Qwen3.6-27B-Fable-Fusion-711-Uncensored-Heretic-NM-DAU-NEO-MAX-MTP-GGUF\Qwen3.6-27B-Fable-Fus-711-UnHeretic-NM-DAU-NEO-MAX-NEO-MTP-Q4_K_M.gguf'
0.08.056.250 I srv    load_model: initializing, n_slots = 4, n_ctx_slot = 134400, kv_unified = 'true'
0.08.070.959 I srv          init: chat template supports preserving reasoning, consider enabling it via --reasoning-preserve
0.08.070.975 I srv  llama_server: model loaded
0.08.070.977 I srv  llama_server: listening on http://127.0.0.1:8080
0.17.565.155 I slot get_availabl: id  3 | task -1 | selected slot by LRU, t_last = -1
0.17.565.205 I slot launch_slot_: id  3 | task 0 | processing task, is_child = 0
0.20.992.893 I slot print_timing: id  3 | task 0 | n_decoded =    100, tg =  32.50 t/s, tg_3s =  32.50 t/s
0.24.007.438 I slot print_timing: id  3 | task 0 | n_decoded =    197, tg =  32.34 t/s, tg_3s =  32.18 t/s
0.27.037.474 I slot print_timing: id  3 | task 0 | n_decoded =    295, tg =  32.34 t/s, tg_3s =  32.34 t/s
0.30.062.029 I slot print_timing: id  3 | task 0 | n_decoded =    392, tg =  32.27 t/s, tg_3s =  32.07 t/s
0.33.091.942 I slot print_timing: id  3 | task 0 | n_decoded =    490, tg =  32.29 t/s, tg_3s =  32.34 t/s
0.36.095.257 I slot print_timing: id  3 | task 0 | n_decoded =    587, tg =  32.29 t/s, tg_3s =  32.30 t/s
0.39.123.528 I slot print_timing: id  3 | task 0 | n_decoded =    684, tg =  32.25 t/s, tg_3s =  32.03 t/s
0.42.145.568 I slot print_timing: id  3 | task 0 | n_decoded =    781, tg =  32.23 t/s, tg_3s =  32.10 t/s
0.45.168.165 I slot print_timing: id  3 | task 0 | n_decoded =    878, tg =  32.22 t/s, tg_3s =  32.09 t/s
0.48.171.834 I slot print_timing: id  3 | task 0 | n_decoded =    974, tg =  32.19 t/s, tg_3s =  31.96 t/s
0.51.197.286 I slot print_timing: id  3 | task 0 | n_decoded =   1071, tg =  32.18 t/s, tg_3s =  32.06 t/s
0.54.227.690 I slot print_timing: id  3 | task 0 | n_decoded =   1168, tg =  32.17 t/s, tg_3s =  32.01 t/s
0.57.229.919 I slot print_timing: id  3 | task 0 | n_decoded =   1264, tg =  32.15 t/s, tg_3s =  31.98 t/s
1.00.231.410 I slot print_timing: id  3 | task 0 | n_decoded =   1360, tg =  32.14 t/s, tg_3s =  31.98 t/s
1.02.269.618 I slot print_timing: id  3 | task 0 | prompt eval time =     350.31 ms /   143 tokens (    2.45 ms per token,   408.21 tokens per second)
1.02.269.623 I slot print_timing: id  3 | task 0 |        eval time =   44354.08 ms /  1425 tokens (   31.13 ms per token,    32.13 tokens per second)
1.02.269.624 I slot print_timing: id  3 | task 0 |       total time =   44704.38 ms /  1568 tokens
1.02.269.625 I slot print_timing: id  3 | task 0 |    graphs reused =       1417
1.02.269.656 I slot      release: id  3 | task 0 | stop processing: n_tokens = 1567, truncated = 0
3.09.074.186 I srv    operator(): operator(): cleaning up before exit...

v4.3.0 KV kvarn6 kvarn6 - CRASH

> .\Downloads\beellama-preview-v0.4.3-bin-win-hip-radeon-x64\llama-server.exe -m .lmstudio\models\DavidAU\Qwen3.6-27B-Fable-Fusion-711-Uncensored-Heretic-NM-DAU-NEO-MAX-MTP-GGUF\Qwen3.6-27B-Fable-Fus-711-UnHeretic-NM-DAU-NEO-MAX-NEO-MTP-Q4_K_M.gguf -ngl 99 -fa 1 -ctk kvarn6 -ctv kvarn6
HIP Library Path: C:\WINDOWS\SYSTEM32\amdhip64_7.dll
0.00.048.335 I cmn  common_param: common_params_print_info: verbosity = 3 (adjust with the `-lv N` CLI arg)
0.00.126.606 W srv  llama_server: -----------------
0.00.126.611 W srv  llama_server: CORS is set to allow all origins ('*') and no API key is set
0.00.126.611 W srv  llama_server: this can be a security risk (cross-origin attacks)
0.00.126.613 W srv  llama_server: more info: https://github.com/ggml-org/llama.cpp/pull/25655
0.00.126.613 W srv  llama_server: -----------------
0.00.132.266 I srv    load_model: loading model '.lmstudio\models\DavidAU\Qwen3.6-27B-Fable-Fusion-711-Uncensored-Heretic-NM-DAU-NEO-MAX-MTP-GGUF\Qwen3.6-27B-Fable-Fus-711-UnHeretic-NM-DAU-NEO-MAX-NEO-MTP-Q4_K_M.gguf'
D:/a/beellama.cpp/beellama.cpp/ggml/src/ggml-cuda/template-instances\..\fattn-mma-f16.cuh:1846: ERROR: HIP kernel flash_attn_ext_f16 has no device code compatible with HIP arch 1300.
D:/a/beellama.cpp/beellama.cpp/ggml/src/ggml-cuda/template-instances\..\fattn-mma-f16.cuh:1846: ERROR: HIP kernel flash_attn_ext_f16 has no device code compatible with HIP arch 1300.
D:/a/beellama.cpp/beellama.cpp/ggml/src/ggml-cuda/template-instances\..\fattn-mma-f16.cuh:1846: ERROR: HIP kernel flash_attn_ext_f16 has no device code compatible with HIP arch 1300.
D:/a/beellama.cpp/beellama.cpp/ggml/src/ggml-cuda/template-instances\..\fattn-mma-f16.cuh:1846: ERROR: HIP kernel flash_attn_ext_f16 has no device code compatible with HIP arch 1300.
D:/a/beellama.cpp/beellama.cpp/ggml/src/ggml-cuda/template-instances\..\fattn-mma-f16.cuh:1846: ERROR: HIP kernel flash_attn_ext_f16 has no device code compatible with HIP arch 1300.
D:/a/beellama.cpp/beellama.cpp/ggml/src/ggml-cuda/template-instances\..\fattn-mma-f16.cuh:1846: ERROR: HIP kernel flash_attn_ext_f16 has no device code compatible with HIP arch 1300.
D:/a/beellama.cpp/beellama.cpp/ggml/src/ggml-cuda/template-instances\..\fattn-mma-f16.cuh:1846: ERROR: HIP kernel flash_attn_ext_f16 has no device code compatible with HIP arch 1300.

longlog.txt

[Anbeeld]

@raufaser based on your logs, you were not losing the latest v0.4.3, and some related fixes landed after your version.

Please try the latest v0.4.3 and report back.

[raufaser]

From my logs: build 10999 (9d59db2)
From releases: Rolling preview build from [v0.4.3] [9d59db2]

Isn't that the same?
Where can i get a newer one?

[Anbeeld]

From my logs: build 10999 (9d59db2) From releases: Rolling preview build from [v0.4.3] [9d59db2]

Isn't that the same? Where can i get a newer one?

It got stuck in CI/CD. v0.4.3 should be updated now, please check it out.

[raufaser]

OK. Now the model is loaded successfully, but inference fails.

0.29.785.429 I srv  llama_server: listening on http://127.0.0.1:8080
0.29.785.448 I srv  update_slots: all slots are idle
0.40.522.912 I srv  server_strea: conv_id=3036328d-c627-4bd9-8208-c0ea3d8d237a (empty=0)
0.40.526.336 I srv    operator(): chat format: peg-native
0.40.526.700 I slot get_availabl: id  3 | task -1 | selected slot by LRU, t_last = -1
0.40.526.700 I srv  get_availabl: updating prompt cache
0.40.526.716 I srv          load:  - looking for better prompt, base f_keep = -1.000, sim = 0.000
0.40.526.721 I srv        update:  - cache state: 0 prompts, 0.000 MiB (limits: 8192.000 MiB, 194560 tokens, 8589934592 est)
0.40.526.722 I srv  get_availabl: prompt cache update took 0.02 ms
0.40.526.782 I cmn  common_reaso: activated, budget=2147483647 tokens
0.40.526.808 I slot launch_slot_: id  3 | task -1 | sampler chain: logits -> ?penalties -> ?dry -> ?top-n-sigma -> top-k -> ?typical -> top-p -> min-p -> ?xtc -> ?temp-ext -> dist
0.40.526.824 I slot launch_slot_: id  3 | task -1 | sampler params:
        repeat_last_n = 64, repeat_penalty = 1.000, frequency_penalty = 0.000, presence_penalty = 0.000
        dry_multiplier = 0.000, dry_base = 1.750, dry_allowed_length = 2, dry_penalty_last_n = 194560
        top_k = 20, top_p = 0.950, min_p = 0.050, xtc_probability = 0.000, xtc_threshold = 0.100, typical_p = 1.000, top_n_sigma = -1.000, temp = 1.000
        mirostat = 0, mirostat_lr = 0.100, mirostat_ent = 5.000, adaptive_target = -1.000, adaptive_decay = 0.900
0.40.526.827 I slot launch_slot_: id  3 | task 0 | processing task, is_child = 0
0.40.526.831 I slot process_sing: id  0 | task -1 | saving idle slot to prompt cache
0.40.526.832 I slot process_sing: id  1 | task -1 | saving idle slot to prompt cache
0.40.526.832 I slot process_sing: id  2 | task -1 | saving idle slot to prompt cache
0.40.526.854 I slot   operator(): id  3 | task 0 | new prompt, n_ctx_slot = 194560, n_keep = 0, task.n_tokens = 16
0.40.526.866 I slot   operator(): id  3 | task 0 | cached n_tokens = 0, memory_seq_rm [0, end)
0.40.658.143 I slot   operator(): id  3 | task 0 | cached n_tokens = 12, memory_seq_rm [12, end)
0.40.658.570 I slot init_sampler: id  3 | task 0 | init sampler, took 0.00 ms, tokens: text = 16, total = 16
0.41.048.183 I slot create_check: id  3 | task 0 | created context checkpoint 1 of 32 (pos_min = 11, pos_max = 11, n_tokens = 12, size = 174.378 MiB)
D:/a/beellama.cpp/beellama.cpp/ggml/src/ggml-cuda/template-instances\..\mma.cuh:796: ERROR: HIP kernel ggml_cuda_mma::load_ldmatrix has no device code compatible with HIP arch 1300.
D:/a/beellama.cpp/beellama.cpp/ggml/src/ggml-cuda/template-instances\..\mma.cuh:796: ERROR: HIP kernel ggml_cuda_mma::load_ldmatrix has no device code compatible with HIP arch 1300.

longlog1.txt
longlog2.txt

[Anbeeld]

Implemented fixes. Please try latest v0.4.3 preview.

[raufaser]

Works!

Only thing i noticed: When MTP enabled, ctx size is much smaller than on my Nvidia 4090. Both have 24GB VRAM. But that is a different topic. Thanks so far.