From e2728c8547f5762f3e6b732ef609f6c7e3e5dd5d Mon Sep 17 00:00:00 2001 From: Marvin Date: Sat, 5 Sep 2026 20:14:29 -0300 Subject: [PATCH] qwen4exp: PP/TG opts for Ampere+Zen3 (merged experts, fused PLE taps, Zen3 tiling) --- ggml/src/ggml.c | 7 +++++-- ggml/src/iqk/iqk_gemm_kquants.cpp | 6 ++++++ ggml/src/iqk/iqk_mul_mat.cpp | 7 +++++-- src/graphs/build_qwen4exp.cpp | 7 ++++--- src/llama-load-tensors.cpp | 9 +++++++-- 5 files changed, 27 insertions(+), 9 deletions(-) diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index 74e9d866..fb0a8bc0 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -18324,7 +18324,9 @@ static void ggml_compute_forward_mul_mat_id( const void * wdata_mm = (src1->type == vec_dot_type) ? src1->data : params->wdata; const size_t row_size_mm = ggml_row_size(vec_dot_type, ne10); - const int chunks_per_expert = MAX(1, MIN(nth, (int)(ne01 / 32))); + // 64 for Zen3 5900XT (2-CCX): halves atomics/barriers vs /32 on narrow + // qwen4exp experts; keeps MIN(nth,...) cap so TG load balance is safe. + const int chunks_per_expert = MAX(1, MIN(nth, (int)(ne01 / 64))); int total_chunks = 0; for (int a = 0; a < n_as; a++) { @@ -18655,7 +18657,8 @@ static void ggml_compute_forward_mul_mat_id_up_gate( const size_t row_size_ug = ggml_row_size(vec_dot_type, ne10); const int64_t nr0_base = src0_2 ? ne01 : ne01/2; - const int chunks_per_expert_ug = MAX(1, MIN(nth, (int)(nr0_base / 32))); + // See above: /64 for Zen3 to cut scheduling overhead on PP-heavy MoE. + const int chunks_per_expert_ug = MAX(1, MIN(nth, (int)(nr0_base / 64))); int total_chunks_ug = 0; for (int a = 0; a < n_as; a++) { diff --git a/ggml/src/iqk/iqk_gemm_kquants.cpp b/ggml/src/iqk/iqk_gemm_kquants.cpp index 8b481b10..cee7ba74 100644 --- a/ggml/src/iqk/iqk_gemm_kquants.cpp +++ b/ggml/src/iqk/iqk_gemm_kquants.cpp @@ -646,6 +646,12 @@ static void mul_mat_qX_K_q8_K_T(int n, const void * vx, size_t bx, const DataInf for (int i = 0; i < nb; ++i) { + // DDR4 latency hiding on Zen3 TG: weight stream is ~270B/block, + // prefetch 4 blocks ahead (~1KB) stays in L1 streamer window. + if (i + 4 < nb) { + __builtin_prefetch(deq.x + i + 4, 0, 3); + } + auto all_scales = deq.new_block(i, q8, accd); __m256i sumi[nrc_y]; diff --git a/ggml/src/iqk/iqk_mul_mat.cpp b/ggml/src/iqk/iqk_mul_mat.cpp index 680f2639..21e472f0 100644 --- a/ggml/src/iqk/iqk_mul_mat.cpp +++ b/ggml/src/iqk/iqk_mul_mat.cpp @@ -61,7 +61,9 @@ struct MulMat { #ifdef __aarch64__ constexpr int k_x_step = 64; //8192; // Tiling does not seem to help on my M2 Max (but difference to tiling is small) #else - constexpr int k_x_step = 64; // This works best on my Ryzen-7950X (but differences to other tile size are small) + // 32 for Zen3 (2x32MB L3, DDR4): 64x8x256xf32 reuse set thrashes CCX on 5900XT. + // Glove-fit for znver3-only builds; Zen4/Intel still fine at 32 (small diff). + constexpr int k_x_step = 32; #endif if (func16 && nrc_y >= 16) { int n_step = (nrc_y - info.cur_y)/16; @@ -139,7 +141,8 @@ struct MulMat { #ifdef __aarch64__ constexpr int k_x_step = 64; //8192; // Tiling does not seem to help on my M2 Max (but difference to tiling is small) #else - constexpr int k_x_step = 64; // This works best on my Ryzen-7950X (but differences to other tile size are small) + // See mul_mat_NxM above: 32 for Zen3 CCX locality. + constexpr int k_x_step = 32; #endif auto op = ggml_unary_op(unary_op); float tmp[k_x_step*16]; diff --git a/src/graphs/build_qwen4exp.cpp b/src/graphs/build_qwen4exp.cpp index 30232065..1cecce61 100644 --- a/src/graphs/build_qwen4exp.cpp +++ b/src/graphs/build_qwen4exp.cpp @@ -135,13 +135,14 @@ static ggml_tensor * qwen4exp_ple_conv( ggml_view_2d(ctx0, model.layers[il].ple_conv1d, 1, hc_dim, model.layers[il].ple_conv1d->nb[1], k * model.layers[il].ple_conv1d->nb[0])); - wk = ggml_reshape_1d(ctx0, wk, hc_dim); + // [1, hc_dim] broadcasts over rows: same math as transpose-mul-transpose + // but saves 2 transposes + 2 conts per tap (2*K nodes per PLE layer). + wk = ggml_reshape_2d(ctx0, wk, 1, hc_dim); if (wk->type != GGML_TYPE_F32) { wk = ggml_cast(ctx0, wk, GGML_TYPE_F32); } - ggml_tensor * term = ggml_mul(ctx0, ggml_cont(ctx0, ggml_transpose(ctx0, shifted)), wk); - term = ggml_cont(ctx0, ggml_transpose(ctx0, term)); // [n_tokens, hc_dim] + ggml_tensor * term = ggml_mul(ctx0, shifted, wk); // [n_tokens, hc_dim] conv_out = conv_out ? ggml_add(ctx0, conv_out, term) : term; } diff --git a/src/llama-load-tensors.cpp b/src/llama-load-tensors.cpp index 447411bf..dbe58f9c 100644 --- a/src/llama-load-tensors.cpp +++ b/src/llama-load-tensors.cpp @@ -4820,7 +4820,11 @@ bool create_tensors_helper::create_std_ffn_exps(int64_t n_embd, const LLM_TN & t if (ug_meta) { layer.ffn_up_gate_exps = create_tensor(ffn_ctx, ug_name, { n_embd, 2*n_ff_exp, n_expert}, flags); } else { - merged = flags == 0 && ml.merge_up_gate_exps && merge_up_gate_exps(tn, i, 0); + // qwen4exp (Qwen3.8-Flash-Next): merge up/gate by default for fused MoE + // (halves expert GEMM launches on both CUDA mmq_id and CPU IQK paths). + // Falls back gracefully if types/shapes differ inside merge_up_gate_exps(). + const bool want_merge = ml.merge_up_gate_exps || model.arch == LLM_ARCH_QWEN4EXP; + merged = flags == 0 && want_merge && merge_up_gate_exps(tn, i, 0); if (!merged) { layer.ffn_up_exps = create_tensor(ffn_ctx, tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), { n_embd, n_ff_exp, n_expert}, flags); layer.ffn_gate_exps = create_tensor(ffn_ctx, tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), { n_embd, n_ff_exp, n_expert}, flags); @@ -4843,7 +4847,8 @@ bool create_tensors_helper::create_std_ffn_exps_from_meta(const LLM_TN & tn, int if (ug_meta) { layer.ffn_up_gate_exps = create_tensor(ffn_ctx, ug_name, { ug_meta->ne[0], ug_meta->ne[1], ug_meta->ne[2]}, flags); } else { - merged = flags == 0 && ml.merge_up_gate_exps && merge_up_gate_exps(tn, i, 0); + const bool want_merge = ml.merge_up_gate_exps || model.arch == LLM_ARCH_QWEN4EXP; + merged = flags == 0 && want_merge && merge_up_gate_exps(tn, i, 0); if (!merged) { auto u_name = tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i); auto g_name = tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i);