chore(train): engine-10 continued: the SwiGLU limit fused, routed experts through a ring of bulk copies

TORAD.md

@@ -104,6 +104,8 @@ which pins a commit of this branch as a submodule.
 | `881c823c5` | `ggml_cuda_op_top_k` with a k of 64 or more past 1,024 columns selects by radix (`topk_radix`): four 8-bit histogram passes find the k-th largest key, then one ordered pass writes every column above it and the lowest columns equal to it, the tiled path's set and tie-break. Up to 16,384 columns a block keeps its row in registers; a wider row is cut into 8,192-column tiles selected in parallel, then one block a row selects among their candidates; one wide row stays with CUB's top-k. The GLM-5.3 DSA indexer (512 pools a layer) paid the tiled path's k serial block reductions: RTX 5080, k 512, 3,520 columns 194.5 -> 6.1 us, 109,020 x 3 66.3 -> 21.8, a 512-row ubatch 13.8-18x. TOP_K 523/523; two mutants fail it (the last equal key unwritten; stage 2 writing a candidate's position). | `GGML_CUDA_TOPK_RADIX_LEGACY` |
 | `33eb70bc0` | `build_attn_mha` takes `mask_is_prefix` (true by default) and the two sparse-attention callers (the DSA layers' top-k mask, and the kpool path's) pass false: every one-sequence causal mask was tagged as a prefix of the cells, the hint under which the CUDA flash attention skips its range scan and live tiles and applies the mask over the whole range, so a GLM-5.3 DSA decode read every cell's K and V to keep the indexer's 2,048. The dense layers keep the hint. Same attention by another split of the cells, bits can move at float rounding; the gain sits at long caches (436k) and is measured on the served model. | `LLAMA_ATTN_SPARSE_MASK_PREFIX_LEGACY` |
 | `78e0fadf7` | `topk_radix` also takes a row of up to 1,024 columns at any k: with CUB's top-k available such a row had no path but CUB one row after another, four launches a row (a DSA indexer's ubatch under 4,096 cells: 2,048 launches a layer). RTX 5080: 3.3-3.9 us against 8.2-12.7 (1 row) and 129-171 (16 rows). | `GGML_CUDA_TOPK_RADIX_LEGACY` |
+| `640a12c41` | A SwiGLU limit (GLM-5.3's routed experts, shared experts and dense FFN: the gate clamped to `[-inf, 10]`, the up projection to `[-10, 10]`, then the GLU) fuses into the quantized mat-vec's GLU epilogue: `ggml_cuda_can_fuse` takes `{MUL_MAT(_ID), CLAMP, MUL_MAT(_ID), CLAMP, GLU}` when the clamps are exactly a limit and the GLU a SwiGLU, and `mul_mat_vec_q` applies it as one float, where the two nodes had kept the gate/up fusion from matching (five launches for one). In place (GLM-5.3-Flash's routed FFN, 8 layers of 32 experts, RTX 5080): 807 against 861 us at 1 token; a verify's 3 tokens are unchanged, their experts unfused. | `GGML_CUDA_GLU_LIMIT_FUSE_LEGACY` |
+| `b8d44d4b8` | Routed experts at 1-8 tokens stream through `mmvq-moe.cu`: one block an SM, a producer warp listing the distinct experts once (every token/slot pair that routes to one reads it once) and streaming tiles of their rows through a ring of shared-memory slots by `cp.async.bulk`, two teams of eight warps on the dot products; an IQ3_XXS warp frees its slot before the math, and the last tiles go by tickets on the stream's tile counter. It takes IQ3_XXS, and Q8_0 past 1 token: in place (GLM-5.3-Flash's routed FFN on 32 experts, CUDA graphs and PDL, RTX 5080) IQ3_XXS 815 against 818 us for 8 layers at 1 token and 1,823 against 1,953 at 3, Q8_0 2,265 against 2,382 for 4 layers at 3; IQ4_XS measured slower and keeps `mul_mat_vec_q`. Its steady state runs at DRAM's ceiling (~913 GB/s); a launch's first q8_1 vector, loaded from global memory behind the stream, costs 3-6 us. | `GGML_CUDA_MMVQ_MOE_LEGACY` |
 
 Every switch in the last column is read once per process and parses as an integer: a `*_LEGACY` switch set to `0`
 is the same as unset (the change stays on), and `=0` turns off `GGML_CUDA_LORA_RANK1_FUSE` and

ggml/src/ggml-cuda/common.cuh

@@ -1512,10 +1512,11 @@ struct ggml_cuda_pq2_prefetch {
     int64_t      bytes = 0;
 };
 
-// One tile counter a stream for the PQ2_0 tensor-core launches (mmvq-pq2-mma.cu): past its own first tiles a block takes
-// the next by an atomic add on it, and the block that takes the launch's last ticket sets it back to 0. A launch takes
-// tickets only after its dependency wait, when every launch before it on the stream has ended, so one counter serves
-// every launch of a stream, CUDA graph replays included. Made zeroed before a graph evaluation, never inside a capture.
+// One tile counter a stream for the ring launches (the PQ2_0 tensor-core ones, mmvq-pq2-mma.cu, and the routed experts',
+// mmvq-moe.cu): past its own first tiles a block takes the next by an atomic add on it, and the block that takes the
+// launch's last ticket sets it back to 0. A launch takes tickets only after its dependency wait, when every launch before
+// it on the stream has ended, so one counter serves every launch of a stream, CUDA graph replays included. Made zeroed
+// before a graph evaluation, never inside a capture.
 struct ggml_cuda_pq2_tile_counters {
     int * ptr = nullptr; // GGML_CUDA_MAX_STREAMS ints
 
@@ -1881,6 +1882,7 @@ struct ggml_cuda_mm_fusion_args_host {
     const ggml_tensor * x_scale = nullptr;
     const ggml_tensor * gate_scale = nullptr;
     ggml_glu_op glu_op;
+    float glu_limit = 0.0f; // > 0: before the GLU, the gate clamped to at most glu_limit and x to +-glu_limit (a SwiGLU limit)
 };
 struct ggml_cuda_mm_fusion_args_device {
     const void * x_bias = nullptr;
@@ -1889,6 +1891,7 @@ struct ggml_cuda_mm_fusion_args_device {
     const void * x_scale = nullptr;
     const void * gate_scale = nullptr;
     ggml_glu_op glu_op;
+    float glu_limit = 0.0f;
 };
 
 struct ggml_cuda_kernel_launch_params {

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -1701,13 +1701,38 @@ static void ggml_cuda_mul_mat_cublas(ggml_backend_cuda_context & ctx, const ggml
     }
 }
 
+// A SwiGLU limit (GLM-5.3, DeepSeek-V4): the gate clamped to [-inf, L] and the up projection to [-L, L] before the GLU,
+// two CLAMP nodes the mat-vec epilogue takes as one float. 0 when the pair is not exactly that.
+static float ggml_cuda_glu_limit(const ggml_tensor * gate_clamp, const ggml_tensor * up_clamp) {
+    if (gate_clamp->op != GGML_OP_CLAMP || up_clamp->op != GGML_OP_CLAMP) {
+        return 0.0f;
+    }
+    const float limit = ggml_get_op_params_f32(up_clamp, 1);
+    if (!(limit > 0.0f) || !std::isfinite(limit) || ggml_get_op_params_f32(up_clamp, 0) != -limit ||
+        ggml_get_op_params_f32(gate_clamp, 0) != -INFINITY || ggml_get_op_params_f32(gate_clamp, 1) != limit) {
+        return 0.0f;
+    }
+    return limit;
+}
+
 static bool ggml_cuda_should_fuse_mul_mat(const ggml_tensor * ffn_up,
                                           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_up_scale = nullptr,
-                                          const ggml_tensor * ffn_gate_scale = nullptr) {
+                                          const ggml_tensor * ffn_gate_scale = nullptr,
+                                          const ggml_tensor * ffn_up_clamp = nullptr,
+                                          const ggml_tensor * ffn_gate_clamp = nullptr) {
+    const bool has_clamp = ffn_up_clamp != nullptr || ffn_gate_clamp != nullptr;
+    if (has_clamp) {
+        // a limit fuses alone: no bias or scale between the mat-vec and its clamp, and only into a SwiGLU
+        if (!ffn_up_clamp || !ffn_gate_clamp || ffn_up_bias || ffn_gate_bias || ffn_up_scale || ffn_gate_scale ||
+            ffn_up_clamp->src[0] != ffn_up || ffn_gate_clamp->src[0] != ffn_gate ||
+            ggml_get_glu_op(glu) != GGML_GLU_OP_SWIGLU || ggml_cuda_glu_limit(ffn_gate_clamp, ffn_up_clamp) == 0.0f) {
+            return false;
+        }
+    }
     const bool has_bias = ffn_up_bias != nullptr || ffn_gate_bias != nullptr;
     const bool has_scale = ffn_up_scale != nullptr || ffn_gate_scale != nullptr;
 
@@ -1730,8 +1755,8 @@ static bool ggml_cuda_should_fuse_mul_mat(const ggml_tensor * ffn_up,
     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;
+    const ggml_tensor * ffn_up_out        = has_clamp ? ffn_up_clamp   : has_bias ? ffn_up_bias   : ffn_up_bias_src;
+    const ggml_tensor * ffn_gate_out      = has_clamp ? ffn_gate_clamp : has_bias ? ffn_gate_bias : ffn_gate_bias_src;
 
     if (glu->src[0] != ffn_gate_out || glu->src[1] != ffn_up_out) {
         return false;
@@ -3738,6 +3763,26 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph *                cgraph,
         }
     }
 
+    std::initializer_list<enum ggml_op> mul_mat_id_clamp_glu_ops = { GGML_OP_MUL_MAT_ID, GGML_OP_CLAMP, GGML_OP_MUL_MAT_ID, GGML_OP_CLAMP, GGML_OP_GLU };
+    std::initializer_list<enum ggml_op> mul_mat_clamp_glu_ops    = { GGML_OP_MUL_MAT,    GGML_OP_CLAMP, GGML_OP_MUL_MAT,    GGML_OP_CLAMP, GGML_OP_GLU };
+
+    if ((is_equal(mul_mat_id_clamp_glu_ops, ops) || is_equal(mul_mat_clamp_glu_ops, ops)) &&
+        ggml_can_fuse_subgraph(cgraph, node_idx, ops, { node_idx + 4 })) {
+        // each mat-vec is followed by its clamp; the GLU's first source names which pair is the gate
+        const ggml_tensor * glu        = cgraph->nodes[node_idx + 4];
+        const bool          gate_first = glu->src[0] == cgraph->nodes[node_idx + 1];
+        const ggml_tensor * ffn_gate   = cgraph->nodes[node_idx + (gate_first ? 0 : 2)];
+        const ggml_tensor * gate_clamp = cgraph->nodes[node_idx + (gate_first ? 1 : 3)];
+        const ggml_tensor * ffn_up     = cgraph->nodes[node_idx + (gate_first ? 2 : 0)];
+        const ggml_tensor * up_clamp   = cgraph->nodes[node_idx + (gate_first ? 3 : 1)];
+
+        if (ggml_cuda_should_fuse_mul_mat(ffn_up, ffn_gate, glu, nullptr, nullptr, nullptr, nullptr, up_clamp, gate_clamp)) {
+            int out_nodes[] = { node_idx + 4 };
+            return ggml_cuda_check_fusion_memory_ranges(cgraph, node_idx, (int)ops.size(), out_nodes, 1, false,
+                                                        ggml_cuda_mmvq_staged_src1(ffn_up));
+        }
+    }
+
     if ((is_equal(mul_mat_id_glu_ops, ops) || is_equal(mul_mat_glu_ops, ops)) &&
         ggml_can_fuse_subgraph(cgraph, node_idx, ops, { node_idx + 2 })) {
         const ggml_tensor * ffn_gate = cgraph->nodes[node_idx];
@@ -4426,7 +4471,9 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
         return bias == nullptr || ids != nullptr || ggml_cuda_mmvq_fusion_operand_ok(bias, out);
     };
 
-    // gate + glu + up, with optional scale/bias on both lanes.
+    // gate + glu + up, with optional scale/bias on both lanes, or a SwiGLU limit's two clamps.
+    // GGML_CUDA_GLU_LIMIT_FUSE_LEGACY=1: the clamps stay nodes of their own and the mat-vecs run apart.
+    static const bool glu_limit_legacy = ggml_env_switch("GGML_CUDA_GLU_LIMIT_FUSE_LEGACY");
     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;
 
@@ -4712,6 +4759,25 @@ static int ggml_cuda_try_fuse(ggml_backend_cuda_context * cuda_ctx, ggml_cgraph
                 fused_node_count  = 3;
                 break;
             }
+        } else if (!glu_limit_legacy && ggml_cuda_can_fuse(cgraph, i, { op, GGML_OP_CLAMP, op, GGML_OP_CLAMP, GGML_OP_GLU }, {})) {
+            // a SwiGLU limit (GLM-5.3's experts, shared experts and dense FFN): the two clamps ride the quantized
+            // mat-vec's epilogue with the GLU, one launch for five nodes
+            ggml_tensor * glu        = cgraph->nodes[i + 4];
+            const bool    gate_first = glu->src[0] == cgraph->nodes[i + 1];
+            ggml_tensor * gate       = cgraph->nodes[i + (gate_first ? 0 : 2)];
+            ggml_tensor * up         = cgraph->nodes[i + (gate_first ? 2 : 0)];
+
+            if (ggml_cuda_should_fuse_mul_mat_vec_q(up, /*with_gate =*/ true)) {
+                ggml_cuda_mm_fusion_args_host fusion_data{};
+                fusion_data.gate      = gate->src[0];
+                fusion_data.glu_op    = ggml_get_glu_op(glu);
+                fusion_data.glu_limit = ggml_cuda_glu_limit(glu->src[0], glu->src[1]);
+
+                ggml_cuda_mul_mat_vec_q(*cuda_ctx, up->src[0], up->src[1], up->src[2], glu, &fusion_data);
+                fused_mul_mat_vec = true;
+                fused_node_count  = 5;
+                break;
+            }
         }
     }