Name and Version
./llama-cli --version
version: (2d8d612)
built with ROCm 7.14 for x86_64-linux-gnu
Operating systems
Linux
GGML backends
HIP
Hardware
AMD Ryzen AI Max+ 395 (Strix Halo), Radeon 8060S iGPU, gfx1151 / RDNA3.5, 128 GB unified memory.
Two nodes, the second attached via rpc-server.
Models
A 35B-A3B MoE at Q8_0. Any MoE model should do — the fault is in the MUL_MAT_ID path.
Problem description & steps to reproduce
Since #27621 (41ef91f7c), MoE models emit nothing but repeated punctuation (///////…) on
gfx1151. The model loads normally, token generation runs at full speed (41 t/s), and every
sampled token is garbage. So this is a numerical fault, not a crash and not a scheduling
failure.
I bisected it to a single condition in ggml_cuda_should_fuse_mul_mat_vec_q()
(ggml/src/ggml-cuda/ggml-cuda.cu):
- if (tensor->op == GGML_OP_MUL_MAT_ID && dst->ne[2] != 1) {
+ if (tensor->op == GGML_OP_MUL_MAT_ID && dst->ne[2] > get_mmvq_mmid_max_batch(src0->type, cc)) {
return false;
}Restoring dst->ne[2] != 1 on HIP makes the output correct again. Toggling only that
condition flips the behaviour reliably in both directions, so the multi-token
mul_mat_vec_q_moe path is where it goes wrong. The single-token path is fine.
Steps:
- Build
master for gfx1151: cmake -B build -DGGML_HIP=ON -DGPU_TARGETS=gfx1151
- Run any MoE model and send a prompt of a few hundred tokens (the fusion only engages
when several tokens share an expert matmul, so a one-word prompt shows nothing)
- Output is
///////…
- Change the condition above back to
dst->ne[2] != 1, rebuild, output is correct
It reproduces the same way with the RPC backend, where the affected matmuls execute on the
remote worker.
What I already ruled out:
- Host and device batch limits agree for RDNA3.5.
GGML_CUDA_CC_IS_RDNA3(cc) and
defined(RDNA3) both select get_mmvq_mmid_max_batch_rdna3(), so
block_dims(warp_size, ncols_dst) stays inside the kernel's __launch_bounds__.
- The MMVQ parameter tables are not involved.
mul_mat_vec_q_moe and its launch wrapper
never call calc_nwarps() or get_device_table_id(), and use a fixed
rows_per_block = 2.
I have no theory beyond that. Can anyone with RDNA3.0 hardware (gfx1100/gfx1101) check
whether it happens there too? The device path takes the same
get_mmvq_mmid_max_batch_rdna3() branch, so if it affects both, this is an RDNA3 issue
rather than an RDNA3.5 one.
First Bad Commit
41ef91f — CUDA: extend MOE fusion to specdec (#27621)
Relevant log output
can't provide as the model behaves normally in the logs, just output gets '/////////////////'
Name and Version
./llama-cli --version
version: (2d8d612)
built with ROCm 7.14 for x86_64-linux-gnu
Operating systems
Linux
GGML backends
HIP
Hardware
AMD Ryzen AI Max+ 395 (Strix Halo), Radeon 8060S iGPU, gfx1151 / RDNA3.5, 128 GB unified memory.
Two nodes, the second attached via
rpc-server.Models
A 35B-A3B MoE at Q8_0. Any MoE model should do — the fault is in the
MUL_MAT_IDpath.Problem description & steps to reproduce
Since #27621 (
41ef91f7c), MoE models emit nothing but repeated punctuation (///////…) ongfx1151. The model loads normally, token generation runs at full speed (41 t/s), and every
sampled token is garbage. So this is a numerical fault, not a crash and not a scheduling
failure.
I bisected it to a single condition in
ggml_cuda_should_fuse_mul_mat_vec_q()(
ggml/src/ggml-cuda/ggml-cuda.cu):Restoring
dst->ne[2] != 1on HIP makes the output correct again. Toggling only thatcondition flips the behaviour reliably in both directions, so the multi-token
mul_mat_vec_q_moepath is where it goes wrong. The single-token path is fine.Steps:
masterfor gfx1151:cmake -B build -DGGML_HIP=ON -DGPU_TARGETS=gfx1151when several tokens share an expert matmul, so a one-word prompt shows nothing)
///////…dst->ne[2] != 1, rebuild, output is correctIt reproduces the same way with the RPC backend, where the affected matmuls execute on the
remote worker.
What I already ruled out:
GGML_CUDA_CC_IS_RDNA3(cc)anddefined(RDNA3)both selectget_mmvq_mmid_max_batch_rdna3(), soblock_dims(warp_size, ncols_dst)stays inside the kernel's__launch_bounds__.mul_mat_vec_q_moeand its launch wrappernever call
calc_nwarps()orget_device_table_id(), and use a fixedrows_per_block = 2.I have no theory beyond that. Can anyone with RDNA3.0 hardware (gfx1100/gfx1101) check
whether it happens there too? The device path takes the same
get_mmvq_mmid_max_batch_rdna3()branch, so if it affects both, this is an RDNA3 issuerather than an RDNA3.5 one.
First Bad Commit
41ef91f — CUDA: extend MOE fusion to specdec (#27621)
Relevant log output
can't provide as the model behaves normally in the logs, just output gets '/////////////////'