chore(train): engine-10 continued: the SwiGLU limit fused, routed experts through a ring of bulk copies - #76
Merged
Merged
Conversation
…ec's GLU
GLM-5.3 clamps before the GLU (a SwiGLU limit of 10, in its routed experts, shared experts and dense FFN):
llama-graph puts a CLAMP of the gate to [-inf, 10] and of the up projection to [-10, 10] between the mat-vecs and the
SWIGLU, and those two nodes kept the gate/up mat-vec + GLU fusion from matching, so a decode ran the two mat-vecs
apart, then two clamps and the GLU: five launches where the fusion takes one. ggml_cuda_can_fuse now takes {MUL_MAT,
CLAMP, MUL_MAT, CLAMP, GLU} and its MUL_MAT_ID twin (either mat-vec first: the GLU's first source is the gate's clamp)
when the clamps are exactly a limit L (gate [-inf, L], up [-L, L], L finite and above 0), nothing sits between a
mat-vec and its clamp, and the GLU is a SwiGLU; mul_mat_vec_q applies the limit in its epilogue as one float,
glu_limit. The PQ2_0 tensor-core path takes no limit and stays unfused. GGML_CUDA_GLU_LIMIT_FUSE_LEGACY=1 leaves the
clamps as nodes.
In place (moe-graph: GLM-5.3-Flash's routed FFN, 8 layers of 32 experts, CUDA graphs and PDL; RTX 5080, legs
A-B-B-A): 807/808 us at 1 token against 865/857 with the clamps as nodes. 3 tokens are unchanged (1,956/1,955 against
1,953/1,974): the fused mat-vec takes one column, and a verify's experts run apart as before.
test-backend-ops: MUL_MAT_VEC_FUSION builds the limit as llama-graph does, at 0.5 so that most products clamp and a
kernel that drops it fails, for Q8_0, IQ3_XXS and Q4_0, with and without ids, at 1 and 3 columns and with a batch.
…ared expert read once A decode's routed experts are a few thousand rows each (GLM-5.3-Flash: 8 of 288, gate and up 2,048 rows of 4,096 weights at IQ3_XXS, down 4,096 of 2,048), and mul_mat_vec_q reads them a row or two a block: a block's loads are its only bytes in flight, and at an MTP verify every token reads its experts again where the tokens share them. mmvq-moe.cu runs one block an SM: a producer warp lists the distinct experts once, each with every token/slot pair that routes to it, and streams tiles of their rows through a ring of shared-memory slots with cp.async.bulk, while two teams of eight warps take the dot products for all of a tile's pairs. An IQ3_XXS warp takes its lanes' fragments of a tile into registers and frees the slot before the math (vecdotq.cuh: iq3_xxs_frag), so the slots stay in flight, and keeps the grid in shared memory a copy a lane. The last tiles go by tickets on the stream's tile counter (ggml_cuda_pq2_tile_counters), so the blocks end within a tile of each other. It takes what it measured faster at in place (moe-graph: GLM-5.3-Flash's routed FFN on 32 experts, CUDA graphs and PDL; RTX 5080, legs A-B-B-A, ring against mul_mat_vec_q): - IQ3_XXS, 8 layers: 815/814 us against 818/879 at 1 token, 1,823/1,833 against 1,953/1,992 at 3. - Q8_0 (an MTP layer's experts), 4 layers: past 1 token only, 2,270/2,260 against 2,384/2,381 at 3; at 1 token it stays on mul_mat_vec_q (988/991 against 979/977), its dot products holding the slot. - IQ4_XS not at all: 527/525 against 506/505 at 1 token, 1,243/1,240 against 1,199/1,200 at 3. GGML_CUDA_MMVQ_MOE_LEGACY=1 turns it off. Per-tile %globaltimer stamps put its steady state at DRAM's ceiling (a 12.5 KB tile every 1.15 us an SM, ~913 GB/s); what it loses is each launch's start: a team's first q8_1 vector is loaded from global memory behind the stream's queues, 3-6 us, as is a down projection's at each new expert. The IQ3_XXS dot product takes its signs as byte masks from ksigns64, 3 integer ops for 4 weights where __vcmpne4 and __vsub4 are emulated (GGML_CUDA_IQ3_XXS_SIGNS_LEGACY builds the old ones), in mul_mat_vec_q too, and get_vec_dot_q_cuda and get_vdr_mmvq move to vecdotq.cuh for both kernels. test-backend-ops: GLM-5.3-Flash's shapes through the ring, MUL_MAT_ID at 1 and 3 tokens (gate/up and down) and MUL_MAT_VEC_FUSION gated at a SwiGLU limit; MUL_MAT_ID 1066/1066, MUL_MAT_VEC_FUSION 1030/1030.
… routed experts' ring (b8d44d4)
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
The shared branch's commits after #75:
GGML_CUDA_GLU_LIMIT_FUSE_LEGACY=1restores the separate clamps.mmvq-moe.cu. It runs one block per SM. A producer warp lists the distinct experts once, so a shared expert is read once, and streams their row tiles through a shared-memory ring ofcp.async.bulkcopies to two teams of consumer warps. The ring takes IQ3_XXS at any token count and Q8_0 past 1 token, where it measured faster. In place: IQ3_XXS 815 vs 818 us at 1 token and 1,823 vs 1,953 at 3; Q8_0 2,265 vs 2,382 at 3. IQ4_XS measured slower and stays onmul_mat_vec_q.GGML_CUDA_MMVQ_MOE_LEGACY=1turns the ring off.Checks: