* CUDA: Fuse MMVQ for NVFP4 and BS 1
TODO:
1. Add tests to test-backend-ops (did verify correctness manually for
one model)
2. Reorder bias/scale once PRs for NVFP4 are merged/landed
* Add dense MMVQ fusion as well
Perf numbers on B4500. Note qwen35 is FP8->Q8
+ ./scripts/compare-llama-bench.py -b master -c osimons/nvfp4_fuse_mmvq --tool llama-bench -i llama-bench.sqlite
| Model | Test | t/s master | t/s osimons/nvfp4_fuse_mmvq | Speedup |
|:-------------------------|:-------------|-------------:|------------------------------:|----------:|
| qwen35moe 35B.A3B NVFP4 | tg128@d32768 | 150.15 | 156.29 | 1.04 |
| qwen35moe 35B.A3B Q4_K_M | tg128@d32768 | 157.91 | 157.64 | 1.00 |
Perf numbers on DGX Spark
+ ./scripts/compare-llama-bench.py -b master -c osimons/nvfp4_fuse_mmvq --tool llama-bench -i llama-bench.sqlite
| Model | Test | t/s master | t/s osimons/nvfp4_fuse_mmvq | Speedup |
|:-------------------------|:-------------|-------------:|------------------------------:|----------:|
| qwen35moe 35B.A3B NVFP4 | tg128@d32768 | 58.31 | 59.69 | 1.02 |
| qwen35moe 35B.A3B Q4_K_M | tg128@d32768 | 54.94 | 54.79 | 1.00 |
* Add tests for the added fusion ops
* Cleanup test-backend-ops
* Cleanup ggml-cuda/mmvq
1. Unrestrict post-scale fusion
2. Rename names accordingly
3. Remove env variable to disable fusion
* Merge old mul_mat patterns into the lane-based approach
* Enable fusion for MoE in shared MMVQ
* Restrict scale_view_nodes, enroll MM + ADD into lane-matcher
* Refactor mmvq loads, still does not help non-nvfp4 kernels
* Restrict scale-fusion to NVFP4
This is necessary, as the prolog is quite heavy in GEMV for some
quants/model configs, leading to net perf regression.
We should really be looking to refactor this such that ratio of
prologue/hot-loop/epilogue is better on the hot-loop
front:
+ ./scripts/compare-llama-bench.py -b master -c
c1b9381d327e063cc846b46b59708444b66dc4d8 --tool llama-bench -i llama-bench.sqlite
| CPU | Model | Test | t/s master | t/s
c1b9381d3 | Speedup |
|:----------------------------|:-------------------------|:-------------|-------------:|----------------:|----------:|
| INTEL(R) XEON(R) GOLD 6542Y | gemma4 26B.A4B NVFP4 | tg128@d32768 | 151.70 | 154.32 | 1.02 |
| INTEL(R) XEON(R) GOLD 6542Y | gemma4 26B.A4B Q4_K_M | tg128@d32768 | 187.95 | 185.73 | 0.99 |
| INTEL(R) XEON(R) GOLD 6542Y | gpt-oss 20B MXFP4 MoE | tg128@d32768 | 304.62 | 300.69 | 0.99 |
| INTEL(R) XEON(R) GOLD 6542Y | qwen35moe 35B.A3B NVFP4 | tg128@d32768 | 193.72 | 211.99 | 1.09 |
| INTEL(R) XEON(R) GOLD 6542Y | qwen35moe 35B.A3B Q4_K_M | tg128@d32768 | 217.76 | 218.15 | 1.00
* Reorder scale & bias-add to adhere to #24331
* Restrict lane scale to NVFP4
Don't need to test unfused combinations
* Cleanup
* Merge single-lane mm-fusion helpers
* Refactor and clean-up host-side fusion logic
* Move gate_bias and scale into the same active-thread guard
Latest perf numbers:
B6000
build:
5b7d9f272 (9578)
+ ./scripts/compare-llama-bench.py -b master -c osimons/nvfp4_fuse_mmvq --tool llama-bench -i llama-bench.sqlite
| CPU | Model | Test | t/s master | t/s osimons/nvfp4_fuse_mmvq | Speedup |
|:----------------------------|:-------------------------|:-------------|-------------:|------------------------------:|----------:|
| INTEL(R) XEON(R) GOLD 6542Y | gemma4 26B.A4B NVFP4 | tg128@d32768 | 151.79 | 154.10 | 1.02 |
| INTEL(R) XEON(R) GOLD 6542Y | gemma4 26B.A4B Q4_K_M | tg128@d32768 | 187.90 | 187.27 | 1.00 |
| INTEL(R) XEON(R) GOLD 6542Y | gpt-oss 20B MXFP4 MoE | tg128@d32768 | 303.77 | 306.56 | 1.01 |
| INTEL(R) XEON(R) GOLD 6542Y | qwen35moe 35B.A3B NVFP4 | tg128@d32768 | 193.41 | 207.99 | 1.08 |
| INTEL(R) XEON(R) GOLD 6542Y | qwen35moe 35B.A3B Q4_K_M | tg128@d32768 | 217.60 | 218.58 | 1.00 |
DGX Spark
build:
5b7d9f272 (9578)
+ ./scripts/compare-llama-bench.py -b master -c osimons/nvfp4_fuse_mmvq --tool llama-bench -i llama-bench.sqlite
| CPU | Model | Test | t/s master | t/s osimons/nvfp4_fuse_mmvq | Speedup |
|:------|:-------------------------|:-------------|-------------:|------------------------------:|----------:|
| CPU | gemma4 26B.A4B NVFP4 | tg128@d32768 | 34.61 | 34.84 | 1.01 |
| CPU | gemma4 26B.A4B Q4_K_M | tg128@d32768 | 46.95 | 46.90 | 1.00 |
| CPU | gpt-oss 20B MXFP4 MoE | tg128@d32768 | 64.84 | 64.62 | 1.00 |
| CPU | qwen35moe 35B.A3B NVFP4 | tg128@d32768 | 59.63 | 60.72 | 1.02 |
| CPU | qwen35moe 35B.A3B Q4_K_M | tg128@d32768 | 56.53 | 56.55 | 1.00 |
PPL values for 5 chunks:
this PR
model mode ppl uncertainty log
/mnt/share/gguf/unsloth/Qwen3.6-35B-A3B-GGUF/Qwen3.6-35B-A3B-UD-Q4_K_M.gguf fusion_enabled 5.2892 0.35389 ppl-value-checks/Qwen3.6-35B-A3B-UD-Q4_K_M.fusion_enabled.log
/mnt/share/gguf/unsloth/Qwen3.6-35B-A3B-GGUF/Qwen3.6-35B-A3B-UD-Q4_K_M.gguf fusion_disabled 5.2742 0.35215 ppl-value-checks/Qwen3.6-35B-A3B-UD-Q4_K_M.fusion_disabled.log
/mnt/share/gguf/nvidia/Qwen3.6-35B-A3B-2.06GB-per-token-CT/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.gguf fusion_enabled 5.4487 0.36866 ppl-value-checks/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.fusion_enabled.log
/mnt/share/gguf/nvidia/Qwen3.6-35B-A3B-2.06GB-per-token-CT/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.gguf fusion_disabled 5.4403 0.36782 ppl-value-checks/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.fusion_disabled.log
/mnt/share/gguf/nvidia/Gemma-4-26B-A4B-NVFP4/Gemma-4-26B-A4B-NVFP4_fp8_q8.gguf fusion_enabled 17342.4348 3703.13932 ppl-value-checks/Gemma-4-26B-A4B-NVFP4_fp8_q8.fusion_enabled.log
/mnt/share/gguf/nvidia/Gemma-4-26B-A4B-NVFP4/Gemma-4-26B-A4B-NVFP4_fp8_q8.gguf fusion_disabled 18627.0624 3998.42475 ppl-value-checks/Gemma-4-26B-A4B-NVFP4_fp8_q8.fusion_disabled.log
/mnt/share/gguf/ggml-org/gpt-oss-20b-GGUF/gpt-oss-20b-mxfp4.gguf fusion_enabled 363.8913 33.14007 ppl-value-checks/gpt-oss-20b-mxfp4.fusion_enabled.log
/mnt/share/gguf/ggml-org/gpt-oss-20b-GGUF/gpt-oss-20b-mxfp4.gguf fusion_disabled 363.8913 33.14007 ppl-value-checks/gpt-oss-20b-mxfp4.fusion_disabled.log
/mnt/share/gguf/unsloth/gemma-4-26B-A4B-it-GGUF/gemma-4-26B-A4B-it-UD-Q4_K_XL.gguf fusion_enabled 17330.3926 3716.70472 ppl-value-checks/gemma-4-26B-A4B-it-UD-Q4_K_XL.fusion_enabled.log
/mnt/share/gguf/unsloth/gemma-4-26B-A4B-it-GGUF/gemma-4-26B-A4B-it-UD-Q4_K_XL.gguf fusion_disabled 17933.9524 3883.17066 ppl-value-checks/gemma-4-26B-A4B-it-UD-Q4_K_XL.fusion_disabled.log
master:
summary: ppl-value-checks/summary.tsv
model mode ppl uncertainty log
/mnt/share/gguf/unsloth/Qwen3.6-35B-A3B-GGUF/Qwen3.6-35B-A3B-UD-Q4_K_M.gguf fusion_enabled 5.2892 0.35389 ppl-value-checks/Qwen3.6-35B-A3B-UD-Q4_K_M.fusion_enabled.log
/mnt/share/gguf/unsloth/Qwen3.6-35B-A3B-GGUF/Qwen3.6-35B-A3B-UD-Q4_K_M.gguf fusion_disabled 5.2742 0.35215 ppl-value-checks/Qwen3.6-35B-A3B-UD-Q4_K_M.fusion_disabled.log
/mnt/share/gguf/nvidia/Qwen3.6-35B-A3B-2.06GB-per-token-CT/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.gguf fusion_enabled 5.4487 0.36866 ppl-value-checks/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.fusion_enabled.log
/mnt/share/gguf/nvidia/Qwen3.6-35B-A3B-2.06GB-per-token-CT/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.gguf fusion_disabled 5.4403 0.36782 ppl-value-checks/Qwen3.6-35B-A3B-2.06GB-per-token-CT_fp8_q8.fusion_disabled.log
/mnt/share/gguf/nvidia/Gemma-4-26B-A4B-NVFP4/Gemma-4-26B-A4B-NVFP4_fp8_q8.gguf fusion_enabled 17342.4348 3703.13932 ppl-value-checks/Gemma-4-26B-A4B-NVFP4_fp8_q8.fusion_enabled.log
/mnt/share/gguf/nvidia/Gemma-4-26B-A4B-NVFP4/Gemma-4-26B-A4B-NVFP4_fp8_q8.gguf fusion_disabled 18627.0624 3998.42475 ppl-value-checks/Gemma-4-26B-A4B-NVFP4_fp8_q8.fusion_disabled.log
/mnt/share/gguf/ggml-org/gpt-oss-20b-GGUF/gpt-oss-20b-mxfp4.gguf fusion_enabled 363.8913 33.14007 ppl-value-checks/gpt-oss-20b-mxfp4.fusion_enabled.log
/mnt/share/gguf/ggml-org/gpt-oss-20b-GGUF/gpt-oss-20b-mxfp4.gguf fusion_disabled 363.8913 33.14007 ppl-value-checks/gpt-oss-20b-mxfp4.fusion_disabled.log
/mnt/share/gguf/unsloth/gemma-4-26B-A4B-it-GGUF/gemma-4-26B-A4B-it-UD-Q4_K_XL.gguf fusion_enabled 17330.3926 3716.70472 ppl-value-checks/gemma-4-26B-A4B-it-UD-Q4_K_XL.fusion_enabled.log
/mnt/share/gguf/unsloth/gemma-4-26B-A4B-it-GGUF/gemma-4-26B-A4B-it-UD-Q4_K_XL.gguf fusion_disabled 17933.9524 3883.17066 ppl-value-checks/gemma-4-26B-A4B-it-UD-Q4_K_XL.fusion_disabled.log
* Allow views to weights in ggml_can_fuse_subgraph
* Remove gate_first from test_mul_mat_vec_fusion
* Ditch lane-parsing approach in favor of hard-coded patterns
* Apply suggestions from code review
Co-authored-by: Georgi Gerganov <redacted>
* Rename ggml_is_constant_view_src to ggml_is_constant
* Finish renaming of
0905129e9d12e2bc6f16d6d3cc4e6b40606fc893
* Readd descriptive prints for fusion debugging
* Add weight-buffer pre-allocation to `test_case`
This is required so we correctly test fusion of NVFP4.
* Update ggml/src/ggml.c
Co-authored-by: Johannes Gäßler <redacted>
* Add 2nd context for weights as suggested by @JohannesGaessler
This reflects more natural use of ggml compared to artifically
pre-allocating weights into the same context
* Exclude fused tests from gradient mode
I'm unsure of the current state, but naively every fusion pattern
should require its own backpropagation implementation. I don't see these
implemented for the CUDA backend, so we can disable tests to avoid
triggering GGML_ASSERT for
ggml_tensor * build_graph(ggml_context * ctx) override {
GGML_ASSERT(!use_weight_context());
return build_graph(ctx, nullptr);
}
* Apply suggestions from code review
Co-authored-by: Johannes Gäßler <redacted>
---------
Co-authored-by: Georgi Gerganov <redacted>
Co-authored-by: Johannes Gäßler <redacted>
const ggml_tensor * x_bias = nullptr;
const ggml_tensor * gate = nullptr;
const ggml_tensor * gate_bias = nullptr;
+ const ggml_tensor * x_scale = nullptr;
+ const ggml_tensor * gate_scale = nullptr;
ggml_glu_op glu_op;
};
struct ggml_cuda_mm_fusion_args_device {
const void * x_bias = nullptr;
const void * gate = nullptr;
const void * gate_bias = nullptr;
+ const void * x_scale = nullptr;
+ const void * gate_scale = nullptr;
ggml_glu_op glu_op;
};
const ggml_tensor * ffn_gate,
const ggml_tensor * glu,
const ggml_tensor * ffn_up_bias = nullptr,
- const ggml_tensor * ffn_gate_bias = nullptr) {
+ const ggml_tensor * ffn_gate_bias = nullptr,
+ const ggml_tensor * ffn_up_scale = nullptr,
+ const ggml_tensor * ffn_gate_scale = nullptr) {
const bool has_bias = ffn_up_bias != nullptr || ffn_gate_bias != nullptr;
+ const bool has_scale = ffn_up_scale != nullptr || ffn_gate_scale != nullptr;
if (has_bias && (!ffn_up_bias || !ffn_gate_bias)) {
return false;
}
+ if (has_scale && (!ffn_up_scale || !ffn_gate_scale)) {
+ return false;
+ }
const bool is_mul_mat = ffn_up->op == GGML_OP_MUL_MAT && ffn_gate->op == GGML_OP_MUL_MAT && glu->op == GGML_OP_GLU;
const bool is_mul_mat_id = ffn_up->op == GGML_OP_MUL_MAT_ID && ffn_gate->op == GGML_OP_MUL_MAT_ID && glu->op == GGML_OP_GLU;
}
const ggml_op expected_bias_op = is_mul_mat ? GGML_OP_ADD : GGML_OP_ADD_ID;
+ const ggml_tensor * ffn_up_bias_src = has_scale ? ffn_up_scale : ffn_up;
+ const ggml_tensor * ffn_gate_bias_src = has_scale ? ffn_gate_scale : ffn_gate;
+ const ggml_tensor * ffn_up_out = has_bias ? ffn_up_bias : ffn_up_bias_src;
+ const ggml_tensor * ffn_gate_out = has_bias ? ffn_gate_bias : ffn_gate_bias_src;
- if (has_bias) {
- if (ffn_up_bias->op != expected_bias_op || ffn_gate_bias->op != expected_bias_op) {
+ if (glu->src[0] != ffn_gate_out || glu->src[1] != ffn_up_out) {
+ return false;
+ }
+
+ if (has_scale) {
+ if (ffn_up_scale->op != GGML_OP_MUL || ffn_gate_scale->op != GGML_OP_MUL) {
+ return false;
+ }
+ const bool up_has_mm = ffn_up_scale->src[0] == ffn_up || ffn_up_scale->src[1] == ffn_up;
+ const bool gate_has_mm = ffn_gate_scale->src[0] == ffn_gate || ffn_gate_scale->src[1] == ffn_gate;
+ if (!up_has_mm || !gate_has_mm) {
return false;
}
+ }
- if (glu->src[0] != ffn_gate_bias || glu->src[1] != ffn_up_bias) {
+ if (has_bias) {
+ if (ffn_up_bias->op != expected_bias_op || ffn_gate_bias->op != expected_bias_op) {
return false;
}
if (expected_bias_op == GGML_OP_ADD) {
- const bool up_has_mul = ffn_up_bias->src[0] == ffn_up || ffn_up_bias->src[1] == ffn_up;
- const bool gate_has_mul = ffn_gate_bias->src[0] == ffn_gate || ffn_gate_bias->src[1] == ffn_gate;
+ const bool up_has_mul = ffn_up_bias->src[0] == ffn_up_bias_src || ffn_up_bias->src[1] == ffn_up_bias_src;
+ const bool gate_has_mul = ffn_gate_bias->src[0] == ffn_gate_bias_src || ffn_gate_bias->src[1] == ffn_gate_bias_src;
if (!up_has_mul || !gate_has_mul) {
return false;
}
} else { // GGML_OP_ADD_ID
- if (ffn_up_bias->src[0] != ffn_up || ffn_gate_bias->src[0] != ffn_gate) {
+ if (ffn_up_bias->src[0] != ffn_up_bias_src || ffn_gate_bias->src[0] != ffn_gate_bias_src) {
return false;
}
if (ffn_up_bias->src[2] != ffn_up->src[2] || ffn_gate_bias->src[2] != ffn_gate->src[2]) {
return false;
}
}
- } else {
- if (glu->src[0] != ffn_gate && glu->src[1] != ffn_up) {
- return false;
- }
}
if (ffn_up->src[0]->type != ffn_gate->src[0]->type || !ggml_are_same_shape(ffn_up->src[0], ffn_gate->src[0]) ||
return false;
}
- if (ffn_up->src[2] && (ffn_up->src[2] != ffn_gate->src[2])) {
+ if (is_mul_mat_id && ffn_up->src[2] != ffn_gate->src[2]) {
return false;
}
bool fused_mul_mat_vec = false;
int fused_node_count = 0;
- // gate + glu + up
+ auto get_mul_mat_scale = [](const ggml_tensor * scale_node, const ggml_tensor * mm_node) -> const ggml_tensor * {
+ const bool scale_lhs_mm = scale_node->src[0] == mm_node;
+ const bool scale_rhs_mm = scale_node->src[1] == mm_node;
+ if (!scale_lhs_mm && !scale_rhs_mm) {
+ return nullptr;
+ }
+
+ const ggml_tensor * scale = scale_lhs_mm ? scale_node->src[1] : scale_node->src[0];
+ if (mm_node->src[0]->type != GGML_TYPE_NVFP4 || scale_node->type != GGML_TYPE_F32 ||
+ scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != 1 ||
+ !ggml_are_same_shape(scale_node, mm_node)) {
+ return nullptr;
+ }
+
+ return scale;
+ };
+
+ auto get_mul_mat_id_scale = [](const ggml_tensor * reshape, const ggml_tensor * repeat, const ggml_tensor * getrows,
+ const ggml_tensor * scale_node, const ggml_tensor * mm_node) -> const ggml_tensor * {
+ if (repeat->src[0] != reshape || getrows->src[0] != repeat || getrows->src[1] != mm_node->src[2]) {
+ return nullptr;
+ }
+ if (!((scale_node->src[0] == mm_node && scale_node->src[1] == getrows) ||
+ (scale_node->src[0] == getrows && scale_node->src[1] == mm_node))) {
+ return nullptr;
+ }
+
+ const ggml_tensor * scale = reshape->src[0];
+ if (mm_node->src[0]->type != GGML_TYPE_NVFP4 || scale_node->type != GGML_TYPE_F32 ||
+ scale->type != GGML_TYPE_F32 || !ggml_is_contiguous(scale) || ggml_nelements(scale) != mm_node->src[0]->ne[2] ||
+ !ggml_are_same_shape(scale_node, mm_node)) {
+ return nullptr;
+ }
+
+ return scale;
+ };
+
+ auto get_bias_tensor = [](const ggml_tensor * bias_node, const ggml_tensor * mul_node, ggml_op op_bias) -> const ggml_tensor * {
+ if (op_bias == GGML_OP_ADD) {
+ if (bias_node->src[0] == mul_node) {
+ return bias_node->src[1];
+ }
+ if (bias_node->src[1] == mul_node) {
+ return bias_node->src[0];
+ }
+ return nullptr;
+ }
+ GGML_ASSERT(op_bias == GGML_OP_ADD_ID);
+ GGML_ASSERT(bias_node->src[0] == mul_node);
+ return bias_node->src[1];
+ };
+
+ // gate + glu + up, with optional scale/bias on both lanes.
for (ggml_op op : { GGML_OP_MUL_MAT, GGML_OP_MUL_MAT_ID }) {
const ggml_op bias_op = op == GGML_OP_MUL_MAT ? GGML_OP_ADD : GGML_OP_ADD_ID;
+ if (op == GGML_OP_MUL_MAT) {
+ for (const bool with_bias : { false, true }) {
+ const int gate_idx = i;
+ const int gate_scale_idx = i + 1;
+ const int gate_bias_idx = with_bias ? i + 2 : -1;
+ const int up_idx = with_bias ? i + 3 : i + 2;
+ const int up_scale_idx = up_idx + 1;
+ const int up_bias_idx = with_bias ? up_idx + 2 : -1;
+ const int glu_idx = with_bias ? up_idx + 3 : up_idx + 2;
+
+ const int out_nodes[] = { glu_idx };
+ ggml_op ops[7];
+ if (with_bias) {
+ ops[0] = op;
+ ops[1] = GGML_OP_MUL;
+ ops[2] = bias_op;
+ ops[3] = op;
+ ops[4] = GGML_OP_MUL;
+ ops[5] = bias_op;
+ ops[6] = GGML_OP_GLU;
+ } else {
+ ops[0] = op;
+ ops[1] = GGML_OP_MUL;
+ ops[2] = op;
+ ops[3] = GGML_OP_MUL;
+ ops[4] = GGML_OP_GLU;
+ }
+ const int n_ops = with_bias ? 7 : 5;
+
+ if (!ggml_can_fuse_subgraph(cgraph, i, n_ops, ops, out_nodes, 1) ||
+ !ggml_cuda_check_fusion_memory_ranges(cgraph, i, n_ops, out_nodes, 1)) {
+ continue;
+ }
+
+ ggml_tensor * gate_n = cgraph->nodes[gate_idx];
+ ggml_tensor * gate_scale_n = cgraph->nodes[gate_scale_idx];
+ ggml_tensor * gate_out_n = with_bias ? cgraph->nodes[gate_bias_idx] : gate_scale_n;
+ ggml_tensor * up_n = cgraph->nodes[up_idx];
+ ggml_tensor * up_scale_n = cgraph->nodes[up_scale_idx];
+ ggml_tensor * up_out_n = with_bias ? cgraph->nodes[up_bias_idx] : up_scale_n;
+ const ggml_tensor * glu = cgraph->nodes[glu_idx];
+
+ if (!ggml_cuda_should_fuse_mul_mat(up_n, gate_n, glu,
+ with_bias ? up_out_n : nullptr, with_bias ? gate_out_n : nullptr, up_scale_n, gate_scale_n)) {
+ continue;
+ }
+
+ const ggml_tensor * gate_scale = get_mul_mat_scale(gate_scale_n, gate_n);
+ const ggml_tensor * up_scale = get_mul_mat_scale(up_scale_n, up_n);
+ if (!gate_scale || !up_scale) {
+ continue;
+ }
+
+ const ggml_tensor * up_bias = with_bias ? get_bias_tensor(up_out_n, up_scale_n, bias_op) : nullptr;
+ const ggml_tensor * gate_bias = with_bias ? get_bias_tensor(gate_out_n, gate_scale_n, bias_op) : nullptr;
+ if (with_bias && (!ggml_are_same_shape(gate_out_n->src[0], gate_out_n->src[1]) ||
+ !ggml_are_same_shape(up_out_n->src[0], up_out_n->src[1]))) {
+ continue;
+ }
+
+ const ggml_tensor * src0 = up_n->src[0];
+ const ggml_tensor * src1 = up_n->src[1];
+ const ggml_tensor * ids = up_n->src[2];
+
+ ggml_cuda_mm_fusion_args_host fusion_data{};
+ fusion_data.gate = gate_n->src[0];
+ fusion_data.x_bias = up_bias;
+ fusion_data.gate_bias = gate_bias;
+ fusion_data.x_scale = up_scale;
+ fusion_data.gate_scale = gate_scale;
+ fusion_data.glu_op = ggml_get_glu_op(glu);
+
+ if (ggml_cuda_should_fuse_mul_mat_vec_q(up_n)) {
+ ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, cgraph->nodes[glu_idx], &fusion_data);
+ fused_mul_mat_vec = true;
+ fused_node_count = n_ops;
+ break;
+ }
+ }
+
+ if (fused_mul_mat_vec) {
+ break;
+ }
+ } else {
+ for (const bool with_bias : { false, true }) {
+ const int gate_idx = i;
+ const int gate_scale_idx = i + 4;
+ const int gate_bias_idx = with_bias ? i + 5 : -1;
+ const int up_idx = with_bias ? i + 6 : i + 5;
+ const int up_scale_idx = up_idx + 4;
+ const int up_bias_idx = with_bias ? up_idx + 5 : -1;
+ const int glu_idx = with_bias ? up_idx + 6 : up_idx + 5;
+
+ const int out_nodes[] = { glu_idx };
+ ggml_op ops[13];
+ if (with_bias) {
+ ops[0] = op;
+ ops[1] = GGML_OP_RESHAPE;
+ ops[2] = GGML_OP_REPEAT;
+ ops[3] = GGML_OP_GET_ROWS;
+ ops[4] = GGML_OP_MUL;
+ ops[5] = bias_op;
+ ops[6] = op;
+ ops[7] = GGML_OP_RESHAPE;
+ ops[8] = GGML_OP_REPEAT;
+ ops[9] = GGML_OP_GET_ROWS;
+ ops[10] = GGML_OP_MUL;
+ ops[11] = bias_op;
+ ops[12] = GGML_OP_GLU;
+ } else {
+ ops[0] = op;
+ ops[1] = GGML_OP_RESHAPE;
+ ops[2] = GGML_OP_REPEAT;
+ ops[3] = GGML_OP_GET_ROWS;
+ ops[4] = GGML_OP_MUL;
+ ops[5] = op;
+ ops[6] = GGML_OP_RESHAPE;
+ ops[7] = GGML_OP_REPEAT;
+ ops[8] = GGML_OP_GET_ROWS;
+ ops[9] = GGML_OP_MUL;
+ ops[10] = GGML_OP_GLU;
+ }
+ const int n_ops = with_bias ? 13 : 11;
+
+ if (!ggml_can_fuse_subgraph(cgraph, i, n_ops, ops, out_nodes, 1) ||
+ !ggml_cuda_check_fusion_memory_ranges(cgraph, i, n_ops, out_nodes, 1)) {
+ continue;
+ }
+
+ ggml_tensor * gate_n = cgraph->nodes[gate_idx];
+ ggml_tensor * gate_scale_n = cgraph->nodes[gate_scale_idx];
+ ggml_tensor * gate_out_n = with_bias ? cgraph->nodes[gate_bias_idx] : gate_scale_n;
+ ggml_tensor * up_n = cgraph->nodes[up_idx];
+ ggml_tensor * up_scale_n = cgraph->nodes[up_scale_idx];
+ ggml_tensor * up_out_n = with_bias ? cgraph->nodes[up_bias_idx] : up_scale_n;
+ const ggml_tensor * glu = cgraph->nodes[glu_idx];
+
+ if (!ggml_cuda_should_fuse_mul_mat(up_n, gate_n, glu,
+ with_bias ? up_out_n : nullptr, with_bias ? gate_out_n : nullptr, up_scale_n, gate_scale_n)) {
+ continue;
+ }
+
+ const ggml_tensor * gate_scale = get_mul_mat_id_scale(cgraph->nodes[gate_idx + 1], cgraph->nodes[gate_idx + 2],
+ cgraph->nodes[gate_idx + 3], gate_scale_n, gate_n);
+ const ggml_tensor * up_scale = get_mul_mat_id_scale(cgraph->nodes[up_idx + 1], cgraph->nodes[up_idx + 2],
+ cgraph->nodes[up_idx + 3], up_scale_n, up_n);
+ if (!gate_scale || !up_scale) {
+ continue;
+ }
+
+ const ggml_tensor * up_bias = with_bias ? get_bias_tensor(up_out_n, up_scale_n, bias_op) : nullptr;
+ const ggml_tensor * gate_bias = with_bias ? get_bias_tensor(gate_out_n, gate_scale_n, bias_op) : nullptr;
+
+ const ggml_tensor * src0 = up_n->src[0];
+ const ggml_tensor * src1 = up_n->src[1];
+ const ggml_tensor * ids = up_n->src[2];
+
+ ggml_cuda_mm_fusion_args_host fusion_data{};
+ fusion_data.gate = gate_n->src[0];
+ fusion_data.x_bias = up_bias;
+ fusion_data.gate_bias = gate_bias;
+ fusion_data.x_scale = up_scale;
+ fusion_data.gate_scale = gate_scale;
+ fusion_data.glu_op = ggml_get_glu_op(glu);
+
+ if (ggml_cuda_should_fuse_mul_mat_vec_q(up_n)) {
+ ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, cgraph->nodes[glu_idx], &fusion_data);
+ fused_mul_mat_vec = true;
+ fused_node_count = n_ops;
+ break;
+ }
+ }
+
+ if (fused_mul_mat_vec) {
+ break;
+ }
+ }
+
if (ggml_cuda_can_fuse(cgraph, i, { op, bias_op, op, bias_op, GGML_OP_GLU }, {})) {
ggml_tensor * glu = cgraph->nodes[i + 4];
ggml_tensor * gate_bias_n = glu->src[0];
continue;
}
- auto get_bias_tensor = [](const ggml_tensor * bias_node, const ggml_tensor * mul_node, ggml_op op_bias) {
- if (op_bias == GGML_OP_ADD) {
- if (bias_node->src[0] == mul_node) {
- return bias_node->src[1];
- }
- if (bias_node->src[1] == mul_node) {
- return bias_node->src[0];
- }
- return (ggml_tensor *) nullptr;
- }
- GGML_ASSERT(op_bias == GGML_OP_ADD_ID);
- GGML_ASSERT(bias_node->src[0] == mul_node);
- return bias_node->src[1];
- };
-
- ggml_tensor * up_bias_tensor = get_bias_tensor(up_bias_n, up_n, bias_op);
- ggml_tensor * gate_bias_tensor = get_bias_tensor(gate_bias_n, gate_n, bias_op);
+ const ggml_tensor * up_bias_tensor = get_bias_tensor(up_bias_n, up_n, bias_op);
+ const ggml_tensor * gate_bias_tensor = get_bias_tensor(gate_bias_n, gate_n, bias_op);
if (!up_bias_tensor || !gate_bias_tensor) {
continue;
fused_mul_mat_vec = false;
fused_node_count = 0;
- // gate + add + glu + up + add
+ // mul_mat + scale + optional bias
+ for (ggml_op op : { GGML_OP_MUL_MAT, GGML_OP_MUL_MAT_ID }) {
+ const ggml_op bias_op = op == GGML_OP_MUL_MAT ? GGML_OP_ADD : GGML_OP_ADD_ID;
+
+ for (const bool with_bias : { false, true }) {
+ const int n_ops = op == GGML_OP_MUL_MAT ? (with_bias ? 3 : 2) : (with_bias ? 6 : 5);
+ const int out_nodes[] = { i + n_ops - 1 };
+ ggml_op ops[6];
+ if (op == GGML_OP_MUL_MAT) {
+ if (with_bias) {
+ ops[0] = op;
+ ops[1] = GGML_OP_MUL;
+ ops[2] = bias_op;
+ } else {
+ ops[0] = op;
+ ops[1] = GGML_OP_MUL;
+ }
+ } else {
+ if (with_bias) {
+ ops[0] = op;
+ ops[1] = GGML_OP_RESHAPE;
+ ops[2] = GGML_OP_REPEAT;
+ ops[3] = GGML_OP_GET_ROWS;
+ ops[4] = GGML_OP_MUL;
+ ops[5] = bias_op;
+ } else {
+ ops[0] = op;
+ ops[1] = GGML_OP_RESHAPE;
+ ops[2] = GGML_OP_REPEAT;
+ ops[3] = GGML_OP_GET_ROWS;
+ ops[4] = GGML_OP_MUL;
+ }
+ }
+
+ if (!ggml_can_fuse_subgraph(cgraph, i, n_ops, ops, out_nodes, 1) ||
+ !ggml_cuda_check_fusion_memory_ranges(cgraph, i, n_ops, out_nodes, 1)) {
+ continue;
+ }
+
+ ggml_tensor * mm_node = cgraph->nodes[i];
+ ggml_tensor * scale_node = op == GGML_OP_MUL_MAT ? cgraph->nodes[i + 1] : cgraph->nodes[i + 4];
+ ggml_tensor * out_node = with_bias ? cgraph->nodes[i + n_ops - 1] : scale_node;
+
+ const ggml_tensor * scale = nullptr;
+ if (op == GGML_OP_MUL_MAT) {
+ scale = get_mul_mat_scale(scale_node, mm_node);
+ } else {
+ scale = get_mul_mat_id_scale(cgraph->nodes[i + 1], cgraph->nodes[i + 2], cgraph->nodes[i + 3], scale_node, mm_node);
+ }
+ if (!scale) {
+ continue;
+ }
+
+ const ggml_tensor * bias = with_bias ? get_bias_tensor(out_node, scale_node, bias_op) : nullptr;
+ if (with_bias && !bias) {
+ continue;
+ }
+ if (with_bias && bias_op == GGML_OP_ADD && !ggml_are_same_shape(out_node->src[0], out_node->src[1])) {
+ continue;
+ }
+ if (with_bias && bias_op == GGML_OP_ADD_ID && out_node->src[2] != mm_node->src[2]) {
+ continue;
+ }
+
+ const ggml_tensor * src0 = mm_node->src[0];
+ const ggml_tensor * src1 = mm_node->src[1];
+ const ggml_tensor * ids = mm_node->src[2];
+
+ ggml_cuda_mm_fusion_args_host fusion_data{};
+ fusion_data.x_bias = bias;
+ fusion_data.x_scale = scale;
+
+ if (ggml_cuda_should_fuse_mul_mat_vec_q(mm_node)) {
+ ggml_cuda_mul_mat_vec_q(*cuda_ctx, src0, src1, ids, out_node, &fusion_data);
+ fused_mul_mat_vec = true;
+ fused_node_count = n_ops;
+ break;
+ }
+ }
+ if (fused_mul_mat_vec) {
+ break;
+ }
+ }
+
+ if (fused_mul_mat_vec) {
+ return fused_node_count - 1;
+ }
+
+ // mul_mat + add
for (ggml_op op : { GGML_OP_MUL_MAT, GGML_OP_MUL_MAT_ID }) {
const ggml_op bias_op = op == GGML_OP_MUL_MAT ? GGML_OP_ADD : GGML_OP_ADD_ID;
}
}
-#ifdef GGML_CUDA_DEBUG
- const int nodes_fused = i - prev_i - 1;
- if (nodes_fused > 0) {
- GGML_LOG_INFO("nodes_fused: %d\n", nodes_fused);
- }
-#endif
prev_i = i;
if (ggml_cuda_is_view_or_noop(node)) {
int nodes_to_skip = ggml_cuda_try_fuse(cuda_ctx, cgraph, i);
if (nodes_to_skip != 0) {
+#ifdef GGML_CUDA_DEBUG
+ const int last_fused = i + nodes_to_skip;
+ GGML_LOG_INFO("nodes_fused: %d, first: %s (%s), last: %s (%s)\n",
+ nodes_to_skip + 1, ggml_op_name(node->op), node->name,
+ ggml_op_name(cgraph->nodes[last_fused]->op), cgraph->nodes[last_fused]->name);
+#endif
i += nodes_to_skip;
continue;
}
bool use_gate = false;
bool use_bias = false;
bool use_gate_bias = false;
+ bool use_scale = false;
+ bool use_gate_scale = false;
[[maybe_unused]] const void * vgate = nullptr;
const float * x_bias = nullptr;
const float * gate_bias = nullptr;
+ const float * x_scale = nullptr;
+ const float * gate_scale = nullptr;
ggml_glu_op active_glu;
if constexpr (has_fusion) {
x_bias = (const float *) fusion.x_bias;
gate_bias = (const float *) fusion.gate_bias;
active_glu = fusion.glu_op;
+ if constexpr (type == GGML_TYPE_NVFP4) {
+ use_scale = fusion.x_scale != nullptr;
+ use_gate_scale = fusion.gate_scale != nullptr && use_gate;
+ x_scale = (const float *) fusion.x_scale;
+ gate_scale = (const float *) fusion.gate_scale;
+ }
}
[[maybe_unused]] float x_biases[ncols_dst] = { 0.0f };
[[maybe_unused]] float gate_biases[ncols_dst] = { 0.0f };
+ [[maybe_unused]] float x_scales;
+ [[maybe_unused]] float gate_scales;
if constexpr (has_fusion) {
+ // 1. Hide latency by prefetching bias, gates and scales here
+ // 2. load only on threads that won't die after partial sum calculation
const uint32_t channel_bias = ids ? channel_x : channel_dst;
- if (use_bias) {
- x_bias = x_bias + sample_dst*stride_sample_dst + channel_bias*stride_channel_dst + row0;
- // 1. Hide latency by prefetching bias and gate here
- // 2. load only on threads that won't die after partial sum calculation
- if (threadIdx.x < rows_per_cuda_block && threadIdx.y == 0 &&
- (rows_per_cuda_block == 1 || uint32_t(row0 + threadIdx.x) < stride_col_dst)) {
+ if (threadIdx.x < rows_per_cuda_block && threadIdx.y == 0 &&
+ (rows_per_cuda_block == 1 || uint32_t(row0 + threadIdx.x) < stride_col_dst)) {
+ if (use_bias) {
+ x_bias = x_bias + sample_dst * stride_sample_dst + channel_bias * stride_channel_dst + row0;
#pragma unroll
for (int j = 0; j < ncols_dst; ++j) {
x_biases[j] = x_bias[j * stride_col_dst + threadIdx.x];
}
}
- }
- if (use_gate_bias) {
- gate_bias = gate_bias + sample_dst*stride_sample_dst + channel_bias*stride_channel_dst + row0;
- if (threadIdx.x < rows_per_cuda_block && threadIdx.y == 0 &&
- (rows_per_cuda_block == 1 || uint32_t(row0 + threadIdx.x) < stride_col_dst)) {
+ if (use_gate_bias) {
+ gate_bias = gate_bias + sample_dst * stride_sample_dst + channel_bias * stride_channel_dst + row0;
#pragma unroll
for (int j = 0; j < ncols_dst; ++j) {
gate_biases[j] = gate_bias[j * stride_col_dst + threadIdx.x];
}
}
+ if constexpr (type == GGML_TYPE_NVFP4) {
+ if (use_scale) {
+ x_scales = x_scale[ids ? channel_x : 0];
+ }
+ if (use_gate_scale) {
+ gate_scales = gate_scale[ids ? channel_x : 0];
+ }
+ }
}
}
if (threadIdx.x < rows_per_cuda_block && (rows_per_cuda_block == 1 || uint32_t(row0 + threadIdx.x) < stride_col_dst)) {
float result = tmp[j][threadIdx.x];
if constexpr (has_fusion) {
+ if constexpr (type == GGML_TYPE_NVFP4) {
+ if (use_scale) {
+ result *= x_scales;
+ }
+ }
if (use_bias) {
result += x_biases[j];
}
if (use_gate) {
float gate_value = tmp_gate[j][threadIdx.x];
+ if constexpr (type == GGML_TYPE_NVFP4) {
+ if (use_gate_scale) {
+ gate_value *= gate_scales;
+ }
+ }
if (use_gate_bias) {
gate_value += gate_biases[j];
}
}
if constexpr (!has_fusion) {
- GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, active_glu, gate_bias, x_bias, tmp_gate);
+ GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, use_scale, use_gate_scale, active_glu, gate_bias, x_bias, x_scale, gate_scale, tmp_gate);
+ }
+ if constexpr (type != GGML_TYPE_NVFP4) {
+ GGML_UNUSED_VARS(use_scale, use_gate_scale, x_scale, gate_scale, x_scales, gate_scales);
}
}
const dim3 & block_nums, const dim3 & block_dims, const int nbytes_shared,
const uint32_t ids_stride, cudaStream_t stream) {
- const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr;
+ const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr ||
+ fusion.x_scale != nullptr || fusion.gate_scale != nullptr;
if constexpr (c_ncols_dst == 1) {
if (has_fusion) {
const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(block_nums, block_dims, nbytes_shared, stream);
const int warp_size = ggml_cuda_info().devices[device].warp_size;
const mmvq_parameter_table_id table_id = get_device_table_id(cc);
- const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr;
const bool has_ids = ids != nullptr;
const auto should_use_small_k = [&](int c_ncols_dst) {
GGML_ABORT("fatal error");
break;
}
-
- GGML_UNUSED(has_fusion);
}
static void mul_mat_vec_q_switch_type(
const void * vx, const ggml_type type_x, const void * vy, const int32_t * ids, const ggml_cuda_mm_fusion_args_device fusion, float * dst,
if (fusion) {
GGML_ASSERT( !ids || dst->ne[2] == 1);
GGML_ASSERT( ids || dst->ne[1] == 1);
+ // Scale fusion is only allowed for NVFP4 currently as the cost of checking this at run-time in the prologue is
+ // non-negligible for some models such as gpt-oss-20b
+ GGML_ASSERT((fusion->x_scale == nullptr && fusion->gate_scale == nullptr) || src0->type == GGML_TYPE_NVFP4);
if (fusion->x_bias) {
GGML_ASSERT(fusion->x_bias->type == GGML_TYPE_F32);
GGML_ASSERT(!ids || fusion->gate_bias->ne[1] == src0->ne[2]);
fusion_local.gate_bias = fusion->gate_bias->data;
}
+ if (fusion->x_scale) {
+ GGML_ASSERT(fusion->x_scale->type == GGML_TYPE_F32);
+ GGML_ASSERT(ggml_is_contiguous(fusion->x_scale));
+ GGML_ASSERT(ggml_nelements(fusion->x_scale) == (ids ? src0->ne[2] : 1));
+ fusion_local.x_scale = fusion->x_scale->data;
+ }
+ if (fusion->gate_scale) {
+ GGML_ASSERT(fusion->gate_scale->type == GGML_TYPE_F32);
+ GGML_ASSERT(ggml_is_contiguous(fusion->gate_scale));
+ GGML_ASSERT(ggml_nelements(fusion->gate_scale) == (ids ? src0->ne[2] : 1));
+ fusion_local.gate_scale = fusion->gate_scale->data;
+ }
fusion_local.glu_op = fusion->glu_op;
}
return -1;
}
+static bool ggml_is_constant(const struct ggml_tensor * tensor) {
+ return tensor->buffer != NULL && ggml_backend_buffer_get_usage(tensor->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS && (tensor->flags & GGML_TENSOR_FLAG_PARAM) == 0;
+}
+
bool ggml_can_fuse_subgraph_ext(const struct ggml_cgraph * cgraph,
const int * node_idxs,
int count,
return false;
}
- // if node is a view, check if the view_src and all it's parent view_srcs are within the subgraph
+ // if node is a view, check if the view_src and all its parent view_srcs are within the subgraph.
+ // external view sources are allowed only for weight tensors, which are constant for this graph execution.
struct ggml_tensor * view_src = node->view_src;
while (view_src) {
- if (ggml_node_list_find_tensor(cgraph, node_idxs, count, view_src) == -1) {
+ if (ggml_node_list_find_tensor(cgraph, node_idxs, count, view_src) == -1 && !ggml_is_constant(view_src)) {
return false;
}
view_src = view_src->view_src;
}
virtual ggml_tensor * build_graph(ggml_context * ctx) = 0;
+ virtual ggml_tensor * build_graph(ggml_context * ctx, ggml_context * ctx_weights) {
+ GGML_UNUSED(ctx_weights);
+ return build_graph(ctx);
+ }
virtual double max_nmse_err() {
return 1e-7;
virtual bool run_whole_graph() { return false; }
virtual std::vector<ggml_tensor *> fusion_test_nodes() { return {}; }
+ virtual bool use_weight_context() { return false; }
ggml_cgraph * gf = nullptr;
ggml_cgraph * gb = nullptr;
/* .mem_base = */ NULL,
/* .no_alloc = */ true,
};
+ const bool use_weights = use_weight_context();
+
ggml_context * ctx = ggml_init(params);
GGML_ASSERT(ctx);
+ ggml_context * ctx_weights = use_weights ? ggml_init(params) : nullptr;
+ GGML_ASSERT(!use_weights || ctx_weights);
gf = ggml_new_graph(ctx);
// pre-graph sentinel
add_sentinel(ctx);
+ if (ctx_weights) {
+ add_sentinel(ctx_weights);
+ }
- ggml_tensor * out = build_graph(ctx);
+ ggml_tensor * out = build_graph(ctx, ctx_weights);
current_op_name = op_desc(out);
check_for_f16_tensor(ctx);
if (!matches_filter(out, op_names_filter)) {
//printf(" %s: skipping\n", op_desc(out).c_str());
+ ggml_free(ctx_weights);
ggml_free(ctx);
return test_status_t::SKIPPED;
}
print_test_result_locked(output_printer, result);
+ ggml_free(ctx_weights);
ggml_free(ctx);
return test_status_t::NOT_SUPPORTED;
}
// post-graph sentinel
add_sentinel(ctx);
+ if (ctx_weights) {
+ add_sentinel(ctx_weights);
+ }
+
+ ggml_backend_buffer_t buf_weights = nullptr;
+ if (ctx_weights) {
+ buf_weights = ggml_backend_alloc_ctx_tensors(ctx_weights, backend1);
+ if (buf_weights == NULL) {
+ printf("failed to allocate weight tensors [%s] ", ggml_backend_name(backend1));
+ ggml_free(ctx_weights);
+ ggml_free(ctx);
+ return test_status_t::FAIL;
+ }
+ ggml_backend_buffer_set_usage(buf_weights, GGML_BACKEND_BUFFER_USAGE_WEIGHTS);
+ }
// allocate
ggml_backend_buffer_t buf = ggml_backend_alloc_ctx_tensors(ctx, backend1);
if (buf == NULL) {
printf("failed to allocate tensors [%s] ", ggml_backend_name(backend1));
+ ggml_backend_buffer_free(buf_weights);
+ ggml_free(ctx_weights);
ggml_free(ctx);
return test_status_t::FAIL;
}
// randomize tensors
initialize_tensors(ctx);
+ if (ctx_weights) {
+ initialize_tensors(ctx_weights);
+ }
// compare
struct callback_userdata {
fused_nodes_to_verify.size());
ggml_backend_buffer_free(buf);
-
+ ggml_backend_buffer_free(buf_weights);
+ ggml_free(ctx_weights);
ggml_free(ctx);
// Create test result
/* .mem_base = */ NULL,
/* .no_alloc = */ true,
};
+ const bool use_weights = use_weight_context();
+
ggml_context_ptr ctx(ggml_init(params)); // smart ptr
GGML_ASSERT(ctx);
+ ggml_context_ptr ctx_weights(use_weights ? ggml_init(params) : nullptr);
+ GGML_ASSERT(!use_weights || ctx_weights);
- ggml_tensor * out = build_graph(ctx.get());
+ ggml_tensor * out = build_graph(ctx.get(), ctx_weights.get());
current_op_name = op_desc(out);
if (!matches_filter(out, op_names_filter)) {
//printf(" %s: skipping\n", op_desc(out).c_str());
return true;
}
+ ggml_backend_buffer_ptr buf_weights(nullptr);
+ if (ctx_weights) {
+ buf_weights.reset(ggml_backend_alloc_ctx_tensors(ctx_weights.get(), backend));
+ if (buf_weights == NULL) {
+ printf("failed to allocate weight tensors\n");
+ return false;
+ }
+ ggml_backend_buffer_set_usage(buf_weights.get(), GGML_BACKEND_BUFFER_USAGE_WEIGHTS);
+ }
+
// allocate
ggml_backend_buffer_ptr buf(ggml_backend_alloc_ctx_tensors(ctx.get(), backend)); // smart ptr
// randomize tensors
initialize_tensors(ctx.get());
+ if (ctx_weights) {
+ initialize_tensors(ctx_weights.get());
+ }
// build graph
ggml_cgraph * gf = ggml_new_graph_custom(ctx.get(), graph_nodes, false);
const bool b; // broadcast b matrix (only for use_id)
const bool with_bias;
const bool with_gate;
+ const bool with_lane_scale;
std::array<int64_t, 2> batch_dims;
test_mul_mat_vec_fusion(ggml_type type, ggml_glu_op op, int64_t m, int64_t n, int64_t k,
bool use_id = false, int n_mats = 1, int n_used = 1, bool b = false, bool with_bias = false, bool with_gate = true,
- std::array<int64_t, 2> batch_dims = {4, 2})
- : type(type), glu_op(op), m(m), n(n), k(k), use_id(use_id), n_mats(n_mats), n_used(n_used), b(b), with_bias(with_bias), with_gate(with_gate), batch_dims(batch_dims) {
+ bool with_lane_scale = false, std::array<int64_t, 2> batch_dims = {4, 2})
+ : type(type), glu_op(op), m(m), n(n), k(k), use_id(use_id), n_mats(n_mats), n_used(n_used), b(b), with_bias(with_bias),
+ with_gate(with_gate), with_lane_scale(with_lane_scale), batch_dims(batch_dims) {
if (use_id) {
GGML_ASSERT(n_used <= n_mats);
}
}
std::string vars() override {
- return VARS_TO_STR12(type, glu_op, m, n, k, use_id, n_mats, n_used, b, with_bias, with_gate, batch_dims);
+ return VARS_TO_STR13(type, glu_op, m, n, k, use_id, n_mats, n_used, b, with_bias, with_gate, with_lane_scale, batch_dims);
}
std::string op_desc(ggml_tensor * t) override {
}
bool run_whole_graph() override { return true; }
+ bool use_weight_context() override { return use_id && with_lane_scale; }
ggml_tensor * build_gate(ggml_context * ctx, ggml_tensor * ffn_gate, ggml_tensor * ffn_up) {
ggml_tensor * out = nullptr;
return out;
}
+ ggml_tensor * build_lane_scale_dense(ggml_context * ctx, ggml_tensor * out) {
+ ggml_tensor * scale = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 1);
+ return ggml_mul(ctx, out, scale);
+ }
+
+ ggml_tensor * build_lane_scale_id(ggml_context * ctx, ggml_context * ctx_weights, ggml_tensor * out, ggml_tensor * ids) {
+ GGML_ASSERT(ctx_weights);
+ ggml_tensor * scale = ggml_new_tensor_1d(ctx_weights, GGML_TYPE_F32, n_mats);
+ ggml_tensor * s = ggml_reshape_3d(ctx, scale, 1, n_mats, 1);
+ s = ggml_repeat_4d(ctx, s, 1, n_mats, m, 1);
+ s = ggml_get_rows(ctx, s, ids);
+ return ggml_mul(ctx, out, s);
+ }
+
ggml_tensor * build_graph(ggml_context * ctx) override {
+ GGML_ASSERT(!use_weight_context());
+ return build_graph(ctx, nullptr);
+ }
+
+ ggml_tensor * build_graph(ggml_context * ctx, ggml_context * ctx_weights) override {
if (!use_id) {
const int channels = batch_dims[0];
const int samples = batch_dims[1];
ggml_tensor * gate = with_gate ? ggml_new_tensor(ctx, type, 4, ne0.data()) : nullptr;
ggml_tensor * up = ggml_new_tensor(ctx, type, 4, ne0.data());
- ggml_tensor * ffn_up = ggml_mul_mat(ctx, up, cur);
- if (with_bias) {
- std::array<int64_t, 4> bias_ne = { ffn_up->ne[0], 1, channels, samples };
- ggml_tensor * up_bias = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, bias_ne.data());
- ffn_up = ggml_add(ctx, ffn_up, up_bias);
- }
+ auto build_lane_up = [&]() {
+ ggml_tensor * ffn_up = ggml_mul_mat(ctx, up, cur);
+ if (with_lane_scale) {
+ ffn_up = build_lane_scale_dense(ctx, ffn_up);
+ }
+ if (with_bias) {
+ std::array<int64_t, 4> bias_ne = { ffn_up->ne[0], 1, channels, samples };
+ ggml_tensor * up_bias = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, bias_ne.data());
+ ffn_up = ggml_add(ctx, ffn_up, up_bias);
+ }
+ return ffn_up;
+ };
- ggml_tensor * ffn_gate = with_gate ? ggml_mul_mat(ctx, gate, cur) : nullptr;
- if (with_bias && with_gate) {
- std::array<int64_t, 4> bias_ne = { ffn_gate->ne[0], 1, channels, samples };
- ggml_tensor * gate_bias = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, bias_ne.data());
- ffn_gate = ggml_add(ctx, ffn_gate, gate_bias);
- }
+ auto build_lane_gate = [&]() {
+ ggml_tensor * ffn_gate = ggml_mul_mat(ctx, gate, cur);
+ if (with_lane_scale) {
+ ffn_gate = build_lane_scale_dense(ctx, ffn_gate);
+ }
+ if (with_bias) {
+ std::array<int64_t, 4> bias_ne = { ffn_gate->ne[0], 1, channels, samples };
+ ggml_tensor * gate_bias = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, bias_ne.data());
+ ffn_gate = ggml_add(ctx, ffn_gate, gate_bias);
+ }
+ return ffn_gate;
+ };
+
+ ggml_tensor * ffn_up = build_lane_up();
+ ggml_tensor * ffn_gate = with_gate ? build_lane_gate() : nullptr;
ggml_tensor * out = with_gate ? build_gate(ctx, ffn_gate, ffn_up) : ffn_up;
ggml_tensor * cur = ggml_new_tensor_3d(ctx, GGML_TYPE_F32, k, this->b ? 1 : n_used, m);
ggml_set_name(cur, "cur");
- ggml_tensor * ffn_up = ggml_mul_mat_id(ctx, ups, cur, ids);
- if (with_bias) {
- ggml_tensor * up_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_up->ne[0], n_mats);
- ffn_up = ggml_add_id(ctx, ffn_up, up_bias_param, ids);
- }
+ auto build_lane_up = [&]() {
+ ggml_tensor * ffn_up = ggml_mul_mat_id(ctx, ups, cur, ids);
+ if (with_lane_scale) {
+ ffn_up = build_lane_scale_id(ctx, ctx_weights, ffn_up, ids);
+ }
+ if (with_bias) {
+ ggml_tensor * up_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_up->ne[0], n_mats);
+ ffn_up = ggml_add_id(ctx, ffn_up, up_bias_param, ids);
+ }
+ return ffn_up;
+ };
- ggml_tensor * ffn_gate = with_gate? ggml_mul_mat_id(ctx, gates, cur, ids) : nullptr;
- if (with_bias && with_gate) {
- ggml_tensor * gate_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_gate->ne[0], n_mats);
- ffn_gate = ggml_add_id(ctx, ffn_gate, gate_bias_param, ids);
- }
+ auto build_lane_gate = [&]() {
+ ggml_tensor * ffn_gate = ggml_mul_mat_id(ctx, gates, cur, ids);
+ if (with_lane_scale) {
+ ffn_gate = build_lane_scale_id(ctx, ctx_weights, ffn_gate, ids);
+ }
+ if (with_bias) {
+ ggml_tensor * gate_bias_param = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, ffn_gate->ne[0], n_mats);
+ ffn_gate = ggml_add_id(ctx, ffn_gate, gate_bias_param, ids);
+ }
+ return ffn_gate;
+ };
+
+ ggml_tensor * ffn_up = build_lane_up();
+ ggml_tensor * ffn_gate = with_gate ? build_lane_gate() : nullptr;
ggml_tensor * out = with_gate ? build_gate(ctx, ffn_gate, ffn_up) : ffn_up;
if (!with_gate && glu_op != GGML_GLU_OP_SWIGLU) {
continue;
}
- test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256,
- use_id, 16, 8, b, with_bias, with_gate));
- test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256,
- use_id, 16, 8, b, with_bias, with_gate, {1, 1}));
+ for (bool with_lane_scale : {false, true}) {
+ if (with_lane_scale && type != GGML_TYPE_NVFP4) {
+ continue;
+ }
+ test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256,
+ use_id, 16, 8, b, with_bias, with_gate, with_lane_scale));
+ test_cases.emplace_back(new test_mul_mat_vec_fusion(type, glu_op, 1, 32, 256,
+ use_id, 16, 8, b, with_bias, with_gate, with_lane_scale, {1, 1}));
+ }
}
}
}
}
if (mode == MODE_GRAD) {
+ test_cases.erase(
+ std::remove_if(test_cases.begin(), test_cases.end(), [](const std::unique_ptr<test_case> & tc) {
+ return tc->run_whole_graph();
+ }),
+ test_cases.end()
+ );
+
size_t n_ok = 0;
for (auto & test : test_cases) {
if (test->eval_grad(backend, op_names_filter, output_printer)) {
overview / pkg / ggml / sources / llama.cpp / commitdiff