perf: Qwen3.5 levers3; GLM-5.3 decode and MTP-verify levers (routed-expert ring, top-k at 2-4 rows, L2 issue stop, hyper-connection front) - #84
Open
marcospaulo wants to merge 48 commits into
Open
marcospaulo wants to merge 48 commits into
marcospaulo wants to merge 48 commits into
Conversation
…tate update, and the chain's weights and states requested before the PDL waits The kernels between the qkv group and the Gated DeltaNet recurrence are a chain of short launches, each one's DRAM misses paid on the critical path (rig's roofline research, 2026-09-27). Each lever below has its own switch; set to 1, the switch restores the code before this commit. The alpha/beta fold (GGML_CUDA_SSM_CONV_AB_LEGACY=1): - Qwen3.5's two bf16 matvecs on the layer's normed input run inside the conv-state update (ggml_cuda_try_ssm_conv_ab, ssm_conv_ab_rows): one launch and one boundary fewer on the chain, 48 times a token. The blocks compute the rows before their dependency wait, mul_mat_vec_f's 256-thread arithmetic, so the values are the pair launch's bit for bit, and write them after it. - The normed input is read before the wait, so it arrives by a release/acquire handoff, not the PDL chain. griddepcontrol.wait makes a prerequisite grid's writes visible to the waiting grid only. rms_norm_fwht_cuda releases to a handoff slot (ggml_cuda_ssm_conv_ab_slots) once its blocks have stored the input, and every conv block acquires it. The pair folds in only when the writer released this evaluation; otherwise it keeps its own launch. - The PQ2_0 group writing the conv's inputs takes at most 152 registers and triggers its dependents after its own wait. At 152, a sub-partition keeps room for the conv's 56-register blocks beside it; at 168 it did not. Every other group keeps 168 and its early trigger; GGML_CUDA_PQ2_MMA_GROUP_REGS_LEGACY=1 puts that group at 168 too. Requested before the waits: - GGML_CUDA_PQ2_PREFETCH_BETWEEN_LEGACY=1: a PQ2_0 launch's L2 prefetch covers, before the next launch's heads, the weights of the kernels between the two (up to the budget). GET_ROWS tables, expert stacks and Hadamard rotation tables are skipped, since no kernel reads them whole. - GGML_CUDA_SSM_CONV_PREWAIT_LEGACY=1: the conv's weights are loaded before its dependency wait. - GGML_CUDA_SSM_CONV_STATE_PREFETCH_LEGACY=1: the conv's cache row goes into L2 before its wait. - GGML_CUDA_GDN_STATE_PREFETCH_LEGACY=1: with the fused gather, the recurrence requests its head's state into L2 before its wait. - GGML_CUDA_GDN_TRIGGER_LEGACY=1 and GGML_CUDA_MMVF_TRIGGER_LEGACY=1: the recurrence and mul_mat_vec_f let the next launch start at their start (mul_mat_vec_f_vec already did). Also: - The conv-state update is capped at 56 registers and has a 4-token instance (decode, an MTP verify of up to 3 drafts). - llama-bench puts an NVTX range "prompt" around the measured prompt pass. An adversarial review of the earlier version (29 agents, three votes a finding) confirmed three defects, fixed here: - the fold never registered; - every group took the 152 cap; - the between-prefetch spent about 1.9 MB a segment on the rotation table. Measured on the GLM-5.3 proxy, which has no bf16 alpha/beta pair and so no fold: tg64 59.72 / 59.72 / 59.45 on the RTX 5070 Ti, same session. On Ternary Bonsai 2 27B, the Qwen3.5 model the fold is for, it is still to be measured, and the bit-identity check is still to be run.
…ph before the recurrence, g_a beside f_a The output gate g_b(g_a(x)) reads the layer input, not the recurrence's output, but the graph built it after the recurrence: two small mat-vec launches on the chain between the recurrence and the output projection. It is now expanded into the graph right after the convolution: g_a beside f_a, two mat-vecs of one input, which ggml-cuda's MUL_MAT pair runs as one launch (5d58cc2), and g_b before the recurrence. The 44-layer GLM-5.3 proxy on an RTX 5070 Ti alone, -cts f16, tg16 node trace: one small mat-vec launch fewer a KDA layer, the token's chain -0.36 % (about 60 us of 16.6 ms); PPL at -ub 1 386825.5220 before and after, to the last digit. It is also the order the next step needs: the recurrence's kernel applying the gated norm itself reads the gate.
…rnel, two launches fewer a KDA layer After a KDA layer's GATED_DELTA_NET, the output gate normalizes the attention rows and multiplies them by sigmoid(gate): RMS_NORM -> MUL(ssm_o_norm), then SIGMOID(gate) -> MUL. Before this, ggml-cuda ran them as two launches, a fused rms_norm + weight and a fused sigmoid + mul. The first waited on the recurrence's just-written rows, the second on the first, both on the layer's critical path. ggml_cuda_try_gdn_gated_norm now registers the four nodes for the GDN node, which skips them. The recurrence's kernel then writes the gated norm itself: - Each (sequence, head) has a ticket. The last of its blocks to finish reads the head's rows back and normalizes them, a warp a row, and sets the ticket back to 0 (ggml_cuda_gdn_norm_tickets, zeroed once a context). - The arithmetic is the unfused kernels'. rms_norm_f32 sums squares by a butterfly over each 32 columns and combines the 32-column sums as its block_reduce does. Then its scale, the weight, and unary_gated_op_kernel's sigmoid product. So the output is the four nodes' output bit for bit. - It fuses only up to 16 tokens a sequence (decode and the MTP verify), for 32- to 128-wide heads that the unfused norm reduces in one pass. Not on the chunked prefill path, not while the graph forks streams, and not when the output shares memory with what the GDN reads or writes, except exactly its attention rows or the gate. glm5next builds the chain in place over the attention rows (ggml_rms_norm_inplace, ggml_mul_inplace), which nothing reads after the norm. Built as separate tensors, the allocator put the output where it pleased, often on the GDN's own inputs, which die at the GDN while other heads' blocks still read them. The matcher declined those (4 of 33 layers at decode, all 33 in the 3-token verify graph). In place, the output is the attention rows in every graph, and all 33 fuse. In place, the norm's RMS_NORM -> MUL no longer fuses on the unfused path (GGML_CUDA_GDN_GATED_NORM_LEGACY=1, or another backend): one launch more there, the same values. On the 44-layer GLM-5.3 proxy, RTX 5070 Ti, the norm on against GGML_CUDA_GDN_GATED_NORM_LEGACY=1: - test-backend-ops: GATED_DELTA_NET_GATED_NORM 13/13 on and legacy (in place and as separate tensors: decode and verify over 1-2 seqs, f16/q8_0/f32 caches, the gathered state, 64- and 32-wide scalar-gate heads, and 20 tokens, which decline). GATED_DELTA_NET_CACHE_FUSION 77/77 and GATED_DELTA_NET 71/71. A kernel without the sigmoid fails the 8 fused cases of the out-of-place build and passes the declining one. - PPL at -ub 1 (386825.5220) and at -c 128 -ub 8 over two sequences (389464.8783): every chunk the same. - llama-server at one slot, greedy, 128 tokens with top-5 log-probabilities: bit-identical, plain and with the MTP draft, and the same over repeated runs of each arm. A kernel that flips the last bit of every fused output differs from position 0 in both. - The node trace of tg16: rms_norm 23 a token against 56, unary_gated 0 against 33, all 33 KDA layers; llama-server's trace, every one of its 594 GDN launches followed by the wo projection's quantize. - test-llama-archs: 0 FAIL (glm5next OK on the GPU, the CPU and both Meta configs). GGML_CUDA_GDN_GATED_NORM_LEGACY=1 keeps the four nodes.
…536ed7), which measured neutral 7536ed7 had the recurrence's kernel write glm5next's KDA output gate: two launches fewer a KDA layer, bit-identical. Measured against e74ed87 (the same graph built out of place, the norm on its own two launches), the 44-layer GLM-5.3 proxy, llama-bench rotated over 4 rounds: - RTX 5070 Ti, tg128: +0.24 % (95 % CI -0.95 to +1.42); pp3 at ubatch 3, the MTP verify's shape: -1.34 % (-3.33 to +0.65). - RTX 5080 + 5070 Ti under -sm tensor, tg128: -0.71 % (-1.75 to +0.33) for a restructured epilogue (the ticket before the state stores, loads hoisted), which on the 5070 Ti measured +0.83 % (-0.67 to +2.33) tg128 and -0.43 % pp3. The node trace says why. The norm's first launch took about 4 us a layer, but that was the wait for the recurrence's state stores to drain, which stays on the chain wherever it lands. Under PDL and CUDA graphs the two small launches cost about 1 us each, and the epilogue's tail cost as much: the ticket's round trip, then the rows reloaded from L2, then the stores. A change measured neutral doesn't earn its place on the default path. It also built the chain in place, which cost the unfused path (other backends, the legacy switch) its RMS_NORM -> MUL fusion. The tree is e74ed87's again. The restructured epilogue is kept as a patch in rig's research notes (glm-matvec-l2-2026-09-29/gatednorm-v3-restructure.diff), with the served-bits check and the A/B that measured it.
…s order (e74ed87) and the gated norm taken out (7536ed7, b850d70) The levers' row records what the served-bits check found beside the bits: the alpha/beta fold does not engage at 331d727 (the node trace's launch count is the same with its switch and without), and the register caps leave 16-32 B stacks. The gated norm's row records why it measured neutral: the rms_norm's time after the recurrence is the wait for its state stores to drain, which stays on the chain wherever the norm runs.
…its first tiles before its PDL wait The routed-expert ring (mmvq-moe.cu) triggered the next launch at its start and read ids only past its dependency wait, so a down projection's first rows were asked for only after its gate/up had ended and the q8_1 of the GLU's output had run. The 44-layer GLM-5.3 proxy under -sm tensor, graph node trace: the down charged 18.94 us a layer on the RTX 5070 Ti, against 14.5 at its DRAM peak. - The ring triggers the next launch only past its own wait, so whatever the kernels before it wrote (the layer's top-k) is whole before any later kernel on the stream starts. - A ring launch whose ids an earlier ring launch on the same stream read in this evaluation (the down after its gate/up) lists its experts and issues its block's first own tiles before its wait: a ring's worth at most, and never a ticket from the stream's counter, which the launch before may still be taking. ids are read past L1. GGML_CUDA_MMVQ_MOE_IDS_EARLY_LEGACY=1 keeps the old order. Measured on the proxy, -sm tensor, -cts f16, RTX 5080 + 5070 Ti; graph node traces, 2 rounds alternated, the median of every MoE layer: - The down: 18.94 -> 17.25 us on the 5070 Ti, 17.76 -> 16.19 on the 5080. - The gate/up is level (34.24 against 34.30); the quantize between is +0.12. - A token's charged kernel time: 5080 7,779.5 -> 7,724.9 us (-0.70 %), 5070 Ti 8,885.7 -> 8,830.1 (-0.63 %). - With the switch set, both cards are level with the base, so moving the trigger alone costs nothing. - llama-bench tg128, rotated over 4 rounds: +0.75 % (95 % CI -1.43 to +2.94 over the rounds), too noisy to resolve a change this size. Bit-identical: PPL at -ub 1 and -ub 3, and under -sm tensor, the same to the last digit in all three arms. Tests: - MOE_FFN_CHAIN, new: GLM's routed FFN, gate/up with and without the SwiGLU limit, then down on one ids; 1 and 3 tokens; a 2,048 and a 1,024 FFN. 8/8. - A build whose tiles before the wait read the next expert's rows fails 8 of 8. It passes under the switch, and passes MUL_MAT_ID 1066/1066, whose single launches never take the path. - MUL_MAT_ID 1066/1066, MUL_MAT_ID_FUSION 13/13, MUL_MAT_VEC_FUSION 1032/1032. - The IQ3_XXS instances stay at 96 registers, with no stack.
…he activation fold that lost The row records the mechanism measure (median charged time per MoE layer from node traces), each card's kernel time a token, the round-paired throughput interval (with the finding that ab-arms.sh's old interval treated a run's samples as independent, about 3x too narrow), bit-identical PPL, and the new MOE_FFN_CHAIN test with its mutants. It also records the down quantizing its own activations, measured after it and not landed: +3.1 and +2.4 us a MoE layer, because every block quantizes every pair's vector after its wait while its landed tiles wait in the ring.
…it the MUL and the ADDs it replaces build_moe_ffn writes the weighted sum as MUL(experts, weights), a view of each slot, and the views added in slot order. The MUL ran as one k_bin_bcast and the ADDs as a second, fused. ggml_cuda_op_moe_weighted_sum does both in one kernel. Each thread takes one element of a token: e0*w0 + e1*w1 + ..., in slot order, with each product and each sum rounded on its own (__fmul_rn, __fadd_rn), as the nodes round them. The matcher checks the MUL's shape, each view's offset and row stride, the ADD chain's order, the subgraph's use counts, and that the output lies over neither input. GGML_CUDA_MOE_WSUM_LEGACY=1 runs the two launches. test-backend-ops MOE_WEIGHTED_SUM: n_embd 4096 and 2880, 8/4/2 slots, 1/3/8 tokens, compared bit for bit (the error is the count of elements whose bits differ, the tolerance 0). 18/18 pass fused and 18/18 under the switch. An nsys trace of the run shows k_moe_weighted_sum 18 times and no k_bin_bcast. Summing the slots in reverse order fails the 12 cases with 4 or 8 slots (1,304 to 20,226 elements differ); the 2 slot cases pass, as swapping two addends is the same add. 44-layer GLM-5.3 proxy under -sm tensor on a 5080 + 5070 Ti, -cts f16, tg32 node traces over two rotated rounds: - the weighted sum, 1.83 and 1.76 us a MoE layer in two launches, is 1.25 and 1.22 in one; - the chain from topk through the shared expert is -0.56 and -0.58 us a layer; - each card's kernel time a token, without the all-reduce, is -0.33 % (7,730.9 -> 7,705.1 us) and -0.34 % (8,309.1 -> 8,280.7). PPL is the same to the last digit: 386825.5220 at -ub 1 and 385507.0946 at -ub 3 on one card, 387892.2627 under -sm tensor.
…a68863) and the three levers that lost The row records the bit-for-bit MOE_WEIGHTED_SUM test with its reverse-order mutant, the sum's charged time a MoE layer and each card's kernel time a token under -sm tensor (-0.33 % and -0.34 %), and bit-identical PPL. It also records three levers measured the same way and not landed: mul_mat_vec_q triggering the next launch past its PDL wait (it only moves charged time between launches), the shared expert's gate/up moved ahead of the routed experts with topk's ids marked whole (+2.3 to +3.7 % a token, through the L2 issuer's plan losing the next layer's q, k and v), and the issuer counting a MUL_MAT_ID as heavy (+4.82 / +3.95 % on the landed order).
…opy, in place of a quantize launch at its reader The front's second kernel (dsv4_hc_pre_gram_f32) makes the sublayer's normed mix a warp to a q8_1 block, so it quantizes each block as it writes it, with quantize_q8_1's arithmetic on the values it stores: the same bits. The evaluation makes the copy for every such mix with a quantized reader on mul_mat_vec_q before any node runs (one copy a key, which each layer's front writes in turn, as ggml-alloc gives every layer's mix the same bytes), and a reader of one or more reads it. A copy remembers the index of the node that last wrote it, so the writes of the nodes in the front's fused group, or before it, leave it whole. Where the front does not run fused, or n_embd is not a multiple of MATRIX_ROW_PADDING, the first reader quantizes it as before. GGML_CUDA_MMVQ_Q8_1_PRODUCER_LEGACY=1 turns it off. test-backend-ops DSV4_HC_PRE_Q8_1, new: the front's mix read by one or two Q8_0 or Q4_K MUL_MATs, 16/16. A doubled block scale in the front's copy fails the 12 cases whose readers read it; the 4 that pass are the two where the reader quantizes it (n_embd 256, 12 tokens) and Q4_K at 8 tokens, which Blackwell runs on MMQ. The 44-layer GLM-5.3 proxy under -sm tensor on a 5080 + 5070 Ti, -cts f16, tg32 node traces: two quantize launches a layer fewer (2,640 a run of 30 evaluations); each card's kernel time a token past the all-reduce -0.76 to -0.81 % and -0.30 to -0.34 % over three runs, -0.68 / -0.56 % with the L2 issuer off. PPL bit for bit: 386825.5220 at -ub 1, 385507.0946 at -ub 3, 387892.2627 under -sm tensor.
) The row records the new DSV4_HC_PRE_Q8_1 test and its doubled-scale mutant, two quantize launches fewer a layer, each card's kernel time a token over three runs with the L2 issuer and one without, why the 5070 Ti (no all-reduce slack, its fronts beside an L2 issue) keeps about half of the gain, the trigger variant that did not help, bit-identical PPL, and the review's open item: no test fails if the copy silently stops being read.
…the chain to the sublayer's projections The front's second kernel ran the comb's Sinkhorn (20 row and column normalizations in one warp) before it could end, and the sublayer's first projections wait for it to end, though only the sublayer's DSV4_HC_POST reads the comb. It now leaves the comb's inputs in their slots of the weights, and a warp a token on a stream of its own (forked after it) makes the comb from them with the same function on the same values, so the same bits, and writes it over them. The evaluation's stream waits for it before any node that reads those weights, in every DSV4_HC_POST, and at the evaluation's end; not while the graph runs concurrent streams. GGML_CUDA_HC_COMB_SIDE_LEGACY=1 keeps it on the chain. test-backend-ops DSV4_HC_PRE_POST, new: the front, then a DSV4_HC_POST reading its weights at once, 5/5. Without the waits the post races the comb: the 1,000-iteration case fails 5 runs of 5, the 20-iteration ones by chance. Without the side kernel, DSV4_HC_PRE_FUSED fails 31 of 36 (the 5 that pass run unfused), DSV4_HC_PRE_POST 5 of 5 and DSV4_HC_PRE_Q8_1 16 of 16. The 44-layer GLM-5.3 proxy under -sm tensor on a 5080 + 5070 Ti, tg32 node traces over two rotated rounds: the front's second kernel ends 1.00 and 0.96 us sooner past its first, and the next launch still starts before it ends (PDL kept across the fork); each card's kernel time a token past the all-reduce -1.03 % and -0.69 % against the switch set. PPL bit for bit: 386825.5220 at -ub 1, 385507.0946 at -ub 3, 387892.2627 under -sm tensor.
… so work it leaves running shows ggml_backend_compare_graph_backend computed the backend under test, then the reference, then read both. The reference's evaluation sat between the backend's return and the read of its outputs, so a kernel the backend left running past the return (a side stream's, not waited for at the evaluation's end) had finished by the read. The reference now runs first. The results of a backend that finishes its work by the return are the same. test-backend-ops DSV4_HC_PRE_FUSED, one case more: a comb of 1,000 iterations with no DSV4_HC_POST in the graph, so only the evaluation's end waits for the comb CUDA makes beside the stream. The front's weights are marked an output, as the tests read them back. With the evaluation's end not waiting (ggml-cuda.cu's ggml_cuda_dsv4_hc_comb_join after the node loop removed), the new case fails 5 runs of 5 and 7 to 17 of the 20-iteration cases fail each run; with the old order all 37 passed 5 runs of 5.
…nto weights read past the front The side kernel writes the comb into the front's weights after the front's launch, and only a reader of the weights or the evaluation's end waits for it. Weights read by the front's own DSV4_HC_PRE alone are free to the allocator once that node is placed, so a later node's output may lie over them and the comb would land in it. The front now makes the comb beside the stream only where the weights outlive it: a view of them besides pre's is in the graph (the post and comb views, read by a DSV4_HC_POST that waits, or past the evaluation's end, which waits), or they are an output; otherwise its second kernel makes the comb as before. The use count is the whole graph's, which the meta backend's and the scheduler's subgraphs keep, so under -sm tensor, where the post is past the all-reduce in the next subgraph, the comb stays beside the stream. No model builds a front without its post, so nothing changes where one runs: the 44-layer GLM-5.3 proxy under -sm tensor launches 100 dsv4_hc_comb_side a token as before, and PPL is bit for bit (386825.5220 at -ub 1, 385507.0946 at -ub 3, 387892.2627 under -sm tensor). test-backend-ops cannot fail for the case this closes: it gives every tensor its own memory, so nothing lies over the weights.
… to lie over its logits ggml_cuda_check_fusion_memory_ranges let the top-k fusion's outputs overlap its logits at one row only, and ggml-alloc places a layer's weights and ids over the logits whenever it can, so at an MTP verify's 3 tokens every MoE layer of the GLM-5.3 proxy ran the unfused chain: 8 kernels (sigmoid, bias add, argsort, get_rows, sum, clamp, div, scale) where one runs at a decode. topk_moe_cuda now takes 4 rows a block, a warp each, and its warps meet at a barrier once every row's logits are read, before a weight or an id is written; ggml_cuda_topk_moe_reads_before_writes says so for the rows that fit one block, and the check lets those overlap. GGML_CUDA_TOPK_MOE_ALIAS_LEGACY=1: at one row only. The proxy at -p 3 -ub 3 (nsys, CUDA graphs): the routing chain from the hyper-connection front to the routed gate/up 8.32 -> 3.90 us a MoE layer on an RTX 5070 Ti, 8.32 -> 3.39 on an RTX 5080 and 8.19 -> 3.58 on the 5070 Ti under -sm tensor; an evaluation's kernel time under -sm tensor -1.73 % (5080) and -1.55 % (5070 Ti), and its wall time -1.93 +- 0.49 % over 4 rotated rounds of 1,500 (the switch on: +0.01 +- 0.52). On the 5070 Ti alone the wall time moved -0.24 +- 0.39 %: the paced L2 issuer's request after the attention output runs ~25 us into the routed gate/up, which then ends where the issue lets it, whenever it starts (with GGML_CUDA_L2_ISSUE_LEGACY=1 the evaluation's kernel time is -0.96 %). A decode is unchanged (+0.10 % and -0.10 % a token). PPL is bit for bit at -ub 1 (386825.5220) and under -sm tensor (387892.2627). At -ub 3 it is 385634.0270 against 385507.0946: the fused kernel normalizes the weights with its own rounding, as a decode's always did. KLD against the -ub 1 logits: 0.006763 +- 0.000061 against the unfused chain's 0.006756 +- 0.000061. test-backend-ops TOPK_MOE (320), MOE_WEIGHTED_SUM, MOE_FFN_CHAIN, MUL_MAT_ID_FUSION, MUL_MAT_ID and ARGSORT pass on the 5080, test-llama-archs too. test-backend-ops gives every tensor its own memory, so no case overlaps the outputs and the logits; the proxy's layers at 3 rows put the weights over the logits at the same rows, and without the barrier PPL was the same.
The plan sizes an issue by the graph nodes of the chain beside it (GGML_CUDA_L2_ISSUE_NODE_US each), but where the evaluation fuses (a hyper-connection front, a top-k) many nodes are one launch, so the issue after the attention output ran on into the routed gate/up. The ring then shared DRAM with it and ended where the issue let it: on the 44-layer GLM-5.3 proxy at -p 3 -ub 3 on an RTX 5070 Ti, 97 % of the gate/up launches started with an issue running, which ended 27.2 us into them, and the gate/up took 125.2 us a KDA layer against 107.9 with no issuer, where the shared expert it had brought into L2 gained back 9.2 us. A ring launch (mmvq_moe) now bumps a word of the context's (ggml_cuda_l2_issue_stop, made before any capture as the tile counters are) as its reads start, and the issuer reads the word when it starts and every 4th piece after, and stops requesting once it has changed. GGML_CUDA_L2_ISSUE_STOP_LEGACY=1: the issuer ignores it. The proxy at -p 3 -ub 3, nsys over 4 rotated rounds, the switch set against not: on the 5070 Ti alone the issue ends 9.3 us into the gate/up (p90 10.6) against 27.2, the gate/up 125.15 -> 108.26 us a KDA layer (107.94 with no issuer), the shared expert's gate +3.2 and up +1.95, the layer -11.5 us, and an evaluation's kernel time -2.30 % (-0.58 % with no issuer at all). Under -sm tensor no issue is running when a ring starts, with or without the switch, and every issue lasts as long: the stop never fires there. PPL is bit for bit: 386825.5220 at -ub 1, 385634.0270 at -ub 3, 387892.2627 under -sm tensor. test-backend-ops MUL_MAT_ID, MUL_MAT_ID_FUSION, MOE_FFN_CHAIN, TOPK_MOE, MOE_WEIGHTED_SUM and ARGSORT pass on the RTX 5080 and the 5070 Ti, test-llama-archs too.
… fragment once for its expert's pairs At 3 tokens (an MTP verify's shape) the ring was bound by its math, not by DRAM: each pair routed to an expert decoded every IQ3_XXS fragment of the expert's rows again (8 grid gathers from shared memory, 4 sign lookups and the sign ops) before its 8 dot products. On the 44-layer GLM-5.3 proxy, where every token routes to each of its 8 experts, the gate/up took 113-116 us at 51-53 % of DRAM with its memory pipes 72-74 % busy (ncu, RTX 5070 Ti). A launch of several tokens now runs its own instance (mmvq_moe<type, nmat, rpw, pairs_once>, IQ3_XXS at rpw*nmat<=2): the expert's pairs 4 at a time (MMVQ_MOE_PB), each fragment decoded once into signed bytes (iq3_xxs_frag_decode) and met with each pair's vector in turn (vec_dot_iq3_xxs_frag_q8), each pair's sums added in the same order, so the same bits. Its own instance, as both paths in one kernel spilled (96 registers a thread at 17 warps an SM): the one-pair path's gate/up went 108 -> 139 us. The one-token instances are unchanged: all 8,302 kernels of b3fc826's libggml-cuda have the same SASS encodings (control words included), and the 3 kernels added are the pairs_once instances. iq3_xxs_frag_pair holds the sign logic both use; vec_dot_iq3_xxs_frag meets each pair of ints as it is decoded, as before, where decoding the whole fragment first put the one-pair down's 18 predicated q8_1 loads behind branches and took it 61.4 -> 66.9 us a layer. GGML_CUDA_MMVQ_MOE_PAIRS_LEGACY=1: a pair at a time. The proxy at -p 3 -ub 3, -cts f16, nsys over 4 rotated rounds: on the 5070 Ti alone the gate/up 107.81 -> 75.65 us a KDA layer, the down 61.41 -> 49.66, an evaluation's kernel time -12.12 %; under -sm tensor -7.01 % on the RTX 5080 and -6.52 % on the 5070 Ti. llama-bench pp3 at -ub 3, 1,500 repetitions over 4 rotated rounds: 140.41 -> 154.00 t/s on one card (+9.68 %, 95 % CI +8.58 to +10.79), 229.27 -> 244.48 under -sm tensor (+6.63 %, +4.33 to +8.93). tg128 reads +1.30 % (-1.39 to +3.99) on one card and -0.31 % (-2.43 to +1.82) under -sm tensor, measured with a draft whose one-token instances were slower; these are b3fc826's. The down stays 21 us over its DRAM floor (28.7 us): its tokens' vectors, 55 KB at 3 tokens of 8 slots, do not fit beside the ring and are read from global memory. PPL is bit for bit, with the switch and without: 386825.5220 at -ub 1, 385634.0270 at -ub 3, 387892.2627 under -sm tensor. test-backend-ops gains MUL_MAT_ID at IQ3_XXS with 8 experts all used at 2, 3, 5 and 8 tokens (past 4 pairs a pass at 5 and 8), gate/up and down shapes, and MOE_FFN_CHAIN on those experts at 3 and 5 tokens; a mutant that stores pair q's sums from pair q+1's fails all 16 IQ3_XXS cases of several tokens on the ring (10 MUL_MAT_ID, 6 MOE_FFN_CHAIN) and passes the one-token ones. MUL_MAT_ID, MUL_MAT_ID_FUSION, MOE_FFN_CHAIN, TOPK_MOE and MUL_MAT pass on the RTX 5080 and the 5070 Ti, test-llama-archs too.
…wait With ids whole before the launch, a ring launch bumps the L2 issue's stop word (b3fc826) before its PDL wait; the note on what the ring touches before that wait now says so. Only the issuer on its own stream reads the word.
marcospaulo
changed the title
… tables and its tile and slot by fastdiv ncu of the ring at 3 tokens (the 44-layer GLM-5.3 proxy, RTX 5070 Ti) put the down at 48 % of DRAM, its loads hitting L1 99.5 % of the time and no warp eligible 55 % of the cycles, and ~15 % of its instructions in integer divisions: each pair of each tile found its vector and its dst row from p / n_used, p % n_used and slot % nchannels_y, and each tile its expert, row tile, slot and phase from tile / ntr and i / nslots, 11 division sequences a warp a tile. The host now makes pair p's vector offset and dst offset (mmvq_moe_dev_args::pair_y, pair_dst: 64 pairs at most, a constant load at an index the warp shares), and the kernel no longer takes the strides they came from; tile / ntr and i / nslots are fast_div_modulo on values made on the host. Every address is the same, so the same bits. The down's instance for several tokens goes 3,296 -> 2,792 instructions, the gate/up's 3,920 -> 3,408, the one-token down's 2,072 -> 1,872; 96 registers or fewer, no stack. The proxy at -p 3 -ub 3, -cts f16, nsys over 4 rotated rounds, against 3dbefa8: on the 5070 Ti alone the down 49.92 -> 40.51 us a KDA layer, the gate/up 75.01 -> 74.66, an evaluation's kernel time -3.42 %; with GGML_CUDA_MMVQ_MOE_PAIRS_LEGACY=1 the gate/up 101.86 (108 before). Under -sm tensor the gate/up 36.35 -> 34.82 us on the RTX 5080 and 39.2 -> 37.3 on the 5070 Ti, the down 30.18 -> 29.5 and 33.7 -> 33.15, the same in each of the 4 rounds; the 5080's kernel time -0.68 %, the 5070 Ti's +2.21 % pooled from rounds that jump in both arms. pp3 at -ub 3 on one card +2.57 % (95 % CI +0.36 to +4.78, host load 12-22). A decode's ring launches read the same or less: the down -2.2 % and -2.8 % under -sm tensor, the gate/up -0.5 % on one card. On the 4-layer proxy with GLM-5.3-Flash's 288 experts (the 5080, random routing: nearly every expert meets one pair) the rings move -0.2 % (gate/up) and -1.2 % (down), and the several-token instance stays 1.2 % and 2.1 % over the one-pair one there, as it was. PPL is bit for bit, with the switch and without: 386825.5220 at -ub 1, 385634.0270 at -ub 3, 387892.2627 under -sm tensor. test-backend-ops MUL_MAT_ID, MUL_MAT_ID_FUSION, MOE_FFN_CHAIN, TOPK_MOE and MUL_MAT pass on the RTX 5080 and the 5070 Ti, test-llama-archs too; adjacent pairs' dst offsets swapped fail 37 MUL_MAT_ID cases and MOE_FFN_CHAIN 10 of 10.
…Ms takes the vector kernel at 2-8 columns At an MTP verify's 3 tokens glm5next's KDA gate projections ran on mul_mat_f, which launches a block for each 32 rows of each dst channel and reads a row's whole K in it: ssm_f_a and ssm_g_a (4096 -> 128 bf16) 8.6 us each on 4 blocks and ssm_beta (4096 -> 64) 6.4 us on 2, RTX 5070 Ti, where at one token mul_mat_vec_f ran f_a and g_a as one launch. On NVIDIA from Ampere on, at 2 to 8 columns and rows x cols <= 1024, where mul_mat_f would launch fewer blocks than the device has SMs, mul_mat_vec_f now takes the product, and ggml_cuda_mul_mat_runs_mmvf says so, so a pair of them on one src1 (f_a and g_a) runs as one launch (ggml_cuda_mul_mat_vec_f_pair). GGML_CUDA_MMVF_UNDERFILLED_LEGACY=1: mul_mat_f. The limit is measured: test-backend-ops perf over f16/bf16 weights of 64-2048 rows, K 1024-8192 and 2-8 columns (RTX 5070 Ti) had the vector kernel at 0.21-0.67 of mul_mat_f's time at rows x cols <= 1024 but 1.1-4.4x slower past 2048 (mul_mat_f's time there follows K, the vector kernel's rows x cols). With it, all 74 such shapes take 0.18-0.68 of mul_mat_f's time, 3 rounds (128 x 4096 bf16 at 3 columns 2.06 us against 7.62); the 118 other shapes 0.94-1.03. On the RTX 5080, one round: 0.23-0.63. The 44-layer GLM-5.3 proxy at -p 3 -ub 3, -cts f16, nsys over 4 rotated rounds on the 5070 Ti alone: a KDA layer's span from its conv to its recurrence 33.63 -> 18.18 us (its float mat-muls 32.67 -> 16.90), an evaluation's kernel time -2.76 % against 2606ae5 and -3.03 % against the switch; pp3 +4.26 % (95 % CI +1.51 to +7.02). Under -sm tensor on the 5080 and the 5070 Ti, -5.71 % and -4.46 % against 2606ae5, the switch within 0.32 % of it. A decode launches the same 1,247 kernels as before. Not the same bits at several tokens: the vector kernel keeps the activations in f32 where mul_mat_f rounds them to bf16, and sums in another order. PPL at -ub 3 385634.0270 -> 388534.0622, and bit for bit at -ub 1 (386825.5220) and with the switch; against the switch's logits the mean KL divergence is 0.006795, where -ub 1 against -ub 3 (summation order alone, on this random-weight proxy) is 0.006763. test-backend-ops MUL_MAT (new cases on both sides of each limit, batched and broadcast), MUL_MAT_PAIR (glm5next's f_a/g_a pair), MUL_MAT_VEC_FUSION, SSM_CONV, SSM_CONV_STATE_UPDATE, GATED_DELTA_NET_CACHE_FUSION and test-llama-archs pass on the RTX 5080 and the 5070 Ti; dst zeroed after the new launch fails 31 MUL_MAT cases, all shapes the rule takes, and MUL_MAT_PAIR none (its pairs fuse).
…ctor kernel at 2-8 columns (18fa139)
… only where they are as many as the experts 3dbefa8 gave the ring an instance for launches of several tokens that decodes each IQ3_XXS fragment once and meets it with every pair of its expert. That pays for an expert that meets several pairs and costs the instance's bookkeeping on one that meets one. Where 3 tokens route to 8 of GLM-5.3-Flash's 288 experts, nearly every expert a launch reads meets one pair: on a 4-layer proxy with its 288 experts (random routing, RTX 5080) that instance ran the gate/up 1.2 % and the down 2.1 % slower than the one-pair one, and GLM-5.3-Flash on 2 RTX PRO 6000 under -sm tensor was 2.4 % and 1.4 % slower end to end at 3 and 4 tokens (GGML_CUDA_MMVQ_MOE_PAIRS_LEGACY=1 against the default, llama-bench -p 3 and -p 4, two rounds each). The host now takes that instance only where the launch's pairs (ntokens * n_used) are at least as many as the experts (ggml_cuda_mmvq_moe_args::n_experts, src0's channels): not for GLM-5.3-Flash at up to 8 tokens, still for the 44-layer proxy, whose 3 tokens each route to all 8 of its 8 experts. GGML_CUDA_MMVQ_MOE_PAIRS_LEGACY=1 keeps the one-pair instance whatever the counts. At -p 3 -ub 3, nsys over 2 rotated rounds on the RTX 5080: on the 288-expert proxy the rings now launch the one-pair instance (their demangled names in the traces) and a layer's gate/up takes 171.67 and 175.54 us where it took 174.23 and 176.95, its down 83.50 and 84.83 where 84.00 and 85.71 (the two layer kinds; the switch's arm, the same instance, within 0.9 % of it). On a 4-layer proxy with 16 experts, where an expert that any pair meets meets ~1.7 of them at random routing, the rule keeps the several-pair instance, which beats the one-pair one there: gate/up 107.17 and 109.71 against 108.58 and 113.03 us, down 51.36 and 51.71 against 55.79 and 57.22. The 44-layer proxy launches the same kernels and instances as before. The two instances make the same bits: PPL at -ub 3 388534.0622 on the 44-layer proxy and 337649.7140 on the 288-expert one, before and after. test-backend-ops MUL_MAT_ID (a new case at 4 tokens on 32 experts, the rule's edge), MUL_MAT_ID_FUSION, MOE_FFN_CHAIN and test-llama-archs pass on the RTX 5080 and the 5070 Ti.
…y where they are as many as the experts (2c5fb99)
…n (GGML_CUDA_HC_COMB_SIDE=1) 4bce940 made each fused hyper-connection front's comb on a stream of its own, off the chain to the sublayer's projections, measured on node traces only: each card's kernel time a token 1.03 % and 0.69 % lower under -sm tensor (the 44-layer GLM-5.3 proxy, RTX 5080 + 5070 Ti). On GLM-5.3-Flash on 2 RTX PRO 6000 under -sm tensor, tg64 was 3.7 % faster with GGML_CUDA_HC_COMB_SIDE_LEGACY=1 in 4 of 4 interleaved pairs (rig-glm, 2e699cc, llama-bench), and on the proxy locally the stream is only +0.70 % over it (95 % CI -0.15 to +1.55, tg64, 4 rotated rounds). So by default the front's second kernel makes the comb on the chain, as under the old switch, which is gone; GGML_CUDA_HC_COMB_SIDE=1 makes it beside the stream as before. The same function on the same values either way: PPL bit for bit (387892.2627 under -sm tensor by default and with the switch, 386825.5220 at -ub 1, 388534.0622 at -ub 3). A tg node trace under -sm tensor launches no dsv4_hc_comb_side by default and 88 an evaluation with the switch, one for each of the proxy's fronts. test-backend-ops DSV4_HC_POST, DSV4_HC_PRE_FUSED, DSV4_HC_PRE_POST and DSV4_HC_PRE_Q8_1 pass on the RTX 5080 and the 5070 Ti each way, test-llama-archs too.
98a6656 made the comb beside the stream opt-in (GGML_CUDA_HC_COMB_SIDE=1) after GLM-5.3-Flash on 2 RTX PRO 6000 under -sm tensor ran tg64 3.7 % faster with the comb on the front's chain (4 of 4 interleaved pairs, rig-glm, 2e699cc), while the 44-layer proxy put the stream at +0.70 % with a 95 % CI through zero (-0.15 to +1.55). No default runs it, and it kept a stream, two events, a list of pending weights that every node was checked against, a join in every DSV4_HC_POST and at each evaluation's end, and c6b4ecf's check that the weights outlive the front, none of which anything else needs. They are gone with the side kernel and the switch: dsv4-hc.cu and dsv4-hc.cuh are as they were before 4bce940, and the front's second kernel makes the comb, as it has by default since 98a6656. Kept: 6588e1d's order in ggml_backend_compare_graph_backend (the reference first, so work a backend leaves running past its return reads as a wrong result, which the paced L2 issuer's stream still could), the front's weights marked an output in its tests, and DSV4_HC_PRE_POST with its 1,000-iteration cases, whose comments now say what they catch without the side stream: a read of the comb before the front has written it. The CUDA library holds no dsv4_hc_comb_side (strings: 0 lines, 8 in 5b9b8e5's). test-backend-ops DSV4_HC_POST (4/4), DSV4_HC_PRE_FUSED (37/37), DSV4_HC_PRE_POST (5/5), DSV4_HC_PRE_Q8_1 (16/16) and test-llama-archs pass on the RTX 5070 Ti.
… not every cell each token llama_kv_cache_set_input_kpool made about eight passes over the cache's cells each token (the pool range, the positions, the pool of each cell, its completeness, the two mask rows) to rebuild the pool maps and the two masks. At 64K cached tokens that was the largest host stack of the server's main thread (evidence.md, decode at depth). llama_kv_cells now logs every cell a mutator changes (llama_kv_cells_log: one consumer takes the list; a reset, a resize, an assignment, or a list past 1/16 of the cells tells it to rebuild) and counts the cells of each sequence. llama_kpool_views, kept by the cache, holds for each sequence the position of each cell, a position -> cell table, the cells of each pool and the pools that are complete, and follows the log: a token's step reads the cells it changed. Removals are applied before insertions, so a cell that takes the position another leaves in the same list is not a collision. The maps and masks are then written from the tables, byte for byte as before. A view does not serve, and the maps are built from the cells as before, when two cells share a position, when the positions span more pools than a map holds, when the cells used reach past n_kv, when a ubatch holds more than one sequence of a stream, past 8 sequences, and when the view's count of cells differs from the cells' own (a missed change is rebuilt, not served). LLAMA_KPOOL_INPUT_LEGACY=1 builds the maps from the cells every call; LLAMA_KPOOL_INPUT_CHECK=1 builds them both ways on every call and aborts on the first differing byte. llama_kpool_set_input(cells_of, views, ...) is the old function over a cells accessor (views == nullptr is the old behaviour); llama_kv_cache_set_input_kpool calls it with the cache's views. llama_kpool_set_input at 1 token, 1 sequence, kpool 4, per call, views against the cells: 2048 cached tokens 1.1 against 12.9 us, 8192 3.8 against 49.1, 32768 16.5 against 223.7, 65536 37.7 against 458.4; at 3 tokens (an MTP verify) and 65536 77.2 against 449.1. The 44-layer proxy on the RTX 5080 + 5070 Ti under -sm tensor at depth 32768, 3 interleaved pairs, tg64: 85.37, 84.22, 84.93 tok/s with the views against 83.69, 83.12, 84.79 with the switch (+2.0, +1.3, +0.2 %; the microbenchmark's 207 us of an 11.8 ms token is 1.7 %). pp3 at -ub 3 in the same runs: 111.2, 108.4, 122.2 against 103.1, 116.1, 98.8, too spread to read (another job held both cards at 98 % load). LLAMA_KPOOL_INPUT_CHECK=1 on the 44-layer proxy, llama-server -c 4096 -sm tensor: 6 requests that decode, reuse a cached prompt by LCP, trim its tail, reset the slot and, once, shift the context, 1536 checked calls, all from views, 23 views rebuilt, no difference; the same with --spec-type draft-mtp --spec-draft-n-max 3 (verify ubatches of 4 tokens and rollbacks): 8192 checked calls, no difference. tests/test-kpool-input.cpp: 420 random rounds (tokens appended, drafts rolled back, a range out of the middle, seq_cp, prepare's save and restore, shifts and divisions, two cells at one position, positions sparse past a map, clears, the cells copied over; one unified pool or one stream each; f16 and f32 masks; with and without the key cache) build the inputs from one llama_kpool_views that lives through the round and from the cells, and compare the 7 tensors byte for byte; 200 rounds check the log against a snapshot of the cells; a decode loop with rollbacks, and one that also moves the newest token into a lower hole (a defrag), must not rebuild the view after the first build. Each of these mutants fails it: a mutator not logging (seq_rm, pos_set, set, seq_keep, pos_add), reset not telling the consumer to rebuild, a removal that leaves the pool's mark, an insertion that does not see a taken position, the count guard removed, a view served with no list, the padding row, the run's base, the rep of a pool, and removals and insertions interleaved. test-kv-mask and test-kpool-can-reuse pass.
…elected (upstream ggml-org#27970) GLM-5.3-Flash's DSA layers mask the whole cache down to the cells their indexer picked (2048 of them and a tail of at most 3 on the proxy), and the MMA kernel read every cell under that mask: at 32K cached tokens the flash attention node took 595 / 667 us a decode token on the RTX 5080 / 5070 Ti under -sm tensor. Ported by hand from ggml-org/llama.cpp 8e93a97 (ggml-org#27970, Aman Gupta): FLASH_ATTN_EXT takes a bound n_kv_max on each mask row's finite cells (ggml_flash_attn_ext_set_n_kv_max, op_params[5], since this fork's [4] is mask_prefix), a kernel compacts each row's finite cells into an index list, and the MMA kernel for <512,512,1,8> and <576,512,1,16> loads K, V and the mask by index over n_kv_max cells, one query a tile, instead of the whole cache. Adapted to this fork: launch_fattn takes the sparse flag last, so the tile and vector kernels are unchanged; the sparse path needs an f16 mask (the packed I16 mask never reaches it: build_attn_sparse's mask is the cache's width), turns off kv_range, kv_live and the KV_max scans, and runs one stage. build_attn_sparse passes GGML_PAD(top_k + kpool - 1, 32) = 2080 on the proxy: the top-k pools' cells and the tail, positions [(q + 1)/r*r, q], at most r - 1 cells unless a sequence holds two cells at one position (rig-glm), which the 32-cell pad keeps at no cost. The inherited DSA path passes top_k as upstream does. deepseek4's compressed attention is not wired here (no DSV4 check past 4K on this box). When: an NVIDIA MMA device, n_kv_max set, the f16 mask of one head, no ALiBi or softcap, and K >= max(4096, n_gather) with n_gather = min(n_tokens, 64/ncols2) * n_kv_max, the cells the gather reads for the queries one dense pass covers. Upstream asks 2 n_gather; on these cards (test-backend-ops perf, GLM's shape with 32 heads on the latent, kv 8448, 16640, 33280 x batch 1 to 512) the gather ran at 0.25-0.99x the dense time wherever K >= n_gather and 1.13-1.27x wherever it is under, and the 2x margin gave up a 3-token verify at 8K cells (0.73x) and prefills at 16K (0.63-0.72x). The 44-layer proxy, both cards, -sm tensor, against GGML_CUDA_FATTN_SPARSE_LEGACY=1 (bands declared before each run): - the FLASH_ATTN_EXT node a decode token (nsys, NVTX, graphs off), card 0 / card 1: -d 2048 197/223 us against 195/224 (under the gate, the same kernels); -d 8192 260/259 against 407/446 (0.64/0.58x); -d 32768 337/335 against 595/667 (0.57/0.50x). The attention kernel itself is 17 us a layer at both depths (dense: 31-35 at 8K, 48-55 at 32K); the rest is the compaction, one block of 256 threads a row, 3.6 us at 8K and 11.2 at 32K (next). - tg32 @ d32768, graphs on: a process's reps 2-6 at 89.3 against 88.2 t/s (median; +1.2 %, where the node's 258-332 us of an 11.3 ms token predicted 2.3-2.9 %); its first rep is slower each way and more so here (73.9 against 78.2), so 3 interleaved pairs of -r 3 means came out 79.3/75.6/82.2 against 81.1/82.1/80.3 (the band, above in each pair, missed). - pp512 (3 interleaved pairs): -d 32768 1235.6 against 975.0 t/s (+26.7 %); -d 8192 1623.9 against 1603.3, dense each way under the gate (before it, the sparse kernel there cost 4.0 %). - PPL at -c 8192 with the sparse kernel taking every batch past 4096 cells (before the gate; one-token batches are the same either way): 348974.1045 against 349039.3242 (0.019 %); KLD mean 0.001611, max 0.002327, same top token 88.8 % (the proxy's order noise: ub 1 against ub 3 is 0.31 % and 0.0068). - test-backend-ops FLASH_ATTN_EXT: every case (3229/3229) and the sparse cases pass on both cards, each way. Upstream's cases, GLM's (512/512, 64 heads, 8192 cells, 2051-cell bound, 1 and 3 queries, 2 streams) and a dense fallback under its gather. Mutants that fail them: the gather reading row k_VKQ_0 + i instead of its index (7 of 7 sparse cases), the compaction dropping each row's last index (upstream's 4; the GLM cases' 1 in 2051 is under the error bound). - the crossover grid is in test-backend-ops perf (GLM's shape at 3 cache lengths and 8 batches). Co-authored-by: Aman Gupta <amangupta052@gmail.com>
…r slots and the pooled-key write's unused slots ggml_compute_forward_set_rows splits a scatter's slots over threads, so two slots naming one row are a data race on the CPU backend (the same value, but a race). - The sparse mask (build_attn_sparse): a pool the top-k took with no finite score (fewer than select_k pools are live) carried cell 0 x r for a padded pool, or cells that overlap the tail. build_indexer now returns the slots' liveness (exp of the pool bias at the selected pools) and build_attn_sparse sends a filler slot to its own dump column past n_kv, in a mask widened by those columns that is viewed back to n_kv before the cand and kq masks. - The pooled-key write (llama_kpool_set_input): an unused slot of the fixed-size write repeated a complete pool's row, which a new slot names too. It now names a spare cell, one each: empty, or not the last of its block (the pooled key is read at the last cell of a block only) and not a cell a real slot writes. With no spare left (kpool 1) it repeats a pool as before. The reference path also emits a rep once, for sequences that share cells. Red on 0b7316a: test-llama-archs -a glm5next under a CPU ThreadSanitizer build reports 3 races in ggml_compute_forward_set_rows, all in the indexer key cache; test-kpool-input gains a check that the rows of the write differ and that a slot naming the last cell of a block computes that cell's pool, and it fails 20 rounds. Green here: 0 reports, every round.
The meta backend (-sm tensor) derives a node's split from its sources' splits. handle_generic returned UNKNOWN for a node with none, and allocating the graph aborted: ggml-backend-meta.cpp GGML_ASSERT(ret.axis != GGML_BACKEND_SPLIT_AXIS_UNKNOWN). ARANGE is routed there, and the sparse mask's dump columns (4a88ac9, build_attn_sparse) are an arange, so every GLM-5.3 decode under -sm tensor aborted in ggml_gallocr_alloc_graph. Every device computes the same values for such a node: it is mirrored. Upstream master has the same UNKNOWN. test-backend-meta-sourceless computes y = x + arange(n) through the meta backend over one GPU twice. Red on the parent's libggml-base: the assert, exit 134. Green here: OK, exit 0.
…barrier (upstream b74f590) process_tile's combine for np > 1 had one __syncthreads() inside threadIdx.y % np == 0 and another in the else branch: a barrier the warps reached at different instructions. Every warp now reaches one barrier, and the combine and the write-back run in the np == 0 warps around it. Ported by hand from ggml-org/llama.cpp b74f590 (ggml-org#27870). No arithmetic changes. Measured against c986f0c's build: - compute-sanitizer synccheck on the RTX 5070 Ti: upstream's repro (hsk=192 hsv=128 nh=4 [8,1] kv=512 nb=3) 3680 errors before, 0 after; GLM-5.3's cases (512/512 [64,1] 8192, sparse, nb 1 and 3) 512 before, 0 after. - test-backend-ops FLASH_ATTN_EXT 3229/3229 on both cards. - PPL bit for bit: 386825.5220 (-c 256 -ub 1), 388534.0622 (-ub 3), 387892.2627 (-sm tensor), 349039.3239 (-c 8192 -sm tensor). - Kernel time, ncu at base clocks with L2 flushed (gpu__time_duration of flash_attn_ext_f16 over GLM's grid, 3 cache lengths x 8 batches, sparse and dense, 48 cells): new/base within 1.79 %; the instrument's base/base spread 1.44 %. Co-authored-by: Siavash Norouzi <35790025+siavashnorouzi@users.noreply.github.com>
…reads, four 16-byte loads each flash_attn_mask_to_sparse_indices tested 2048 mask columns a round with 256 threads issuing scalar loads, which cost 11.2 us a layer at 32K cached tokens on a GLM-5.3-Flash proxy's decode. It now tests 32 consecutive columns a thread from four 16-byte loads issued together, 32768 columns a round, and places each thread's cells with a block-wide scan of the counts. Upstream's scan stays as flash_attn_mask_to_sparse_indices_legacy, selected by GGML_CUDA_FATTN_SPARSE_SCAN_LEGACY=1 and automatically for a mask row that is not 16-byte aligned. The index lists the two produce are identical, so attention's result is bit for bit unchanged. Measured on an RTX 5080 + RTX 5070 Ti pair, -sm tensor -ts 1/1 -fa 1, on a 45-block proxy (33 KDA layers, 11 DSA, 1 dense lead): - test-backend-ops FLASH_ATTN_EXT 3229/3229 on both cards, by default and with the scan switch. Two new cases at 40960 mask cells cover a second scan round. - Bit for bit: PPL at -c 8192 -b 8192 -ub 3 --chunks 1 is identical by default, with the scan switch, and from the parent build. - Three mutants each fail the sparse cases and only those: a thread's offset taken inclusive of its own count; a round's count not carried into the next (fails only the 40960-cell cases); and the earlier last-index drop. - Scan kernel, median a launch over the last 31 decode tokens, NVTX with graphs off: 2.56 us (5080) and 3.26 us (5070 Ti) at -d 8192, 4.54 and 4.42 us at -d 32768, against 3.6 and 11.2 us before. The FLASH_ATTN_EXT node at -d 8192 is 0.62x and 0.58x of the dense path's, all three kernels counted. - Wall, tg32 at -d 32768, -r 6, reps 2-6 median, two interleaved runs: 89.92 and 88.36 tok/s against 84.22 and 84.92 with the scan switch, +6.8 % and +4.1 %. The declared bands asked for <= 3 us at d8192 and <= 4 us at d32768; three of those four points miss by 8-13 %, recorded beside the lever's bands. The d32768 node ratio is not quoted here: its dense-path arm produced no capture and is being re-run.
…res a key on a quad of lanes Two changes to GLM-5.3-Flash's sparse indexer. They land together because the second rewrites the kernel the first teaches to read rows, and splitting them would leave a commit that builds while computing the wrong thing. ggml_lightning_indexer_rows: GGML_OP_LIGHTNING_INDEXER takes an optional src[4], I32 rows, so key i of stream s is k's row rows[i, s]. glm5next's fused path hands the f16 index cache's pooled head and pool_reps straight to the indexer instead of materialising get_rows's f32 copy of every pool's key -- at 32K cached tokens that copy was an ~20 us gather a layer a token, 8258 blocks and 4.2 MB written, 11 times a token. The CPU backend reads the rows too; Metal and SYCL refuse the variant; the meta backend mirrors it, since every device reads the same keys through the same rows. LLAMA_INDEXER_GATHER_LEGACY=1 restores the gather. lightning_indexer_kernel_quad: the vector kernel's cases now score a key on a quad of lanes. Lane s of a quad holds the dims lanes 4j + s held before and adds its 8 products in registers in the order warp_reduce_sum's xor 16, 8 and 4 added them, then reduces across the quad with xor 2 and 1: two shuffles a key and head where the vector kernel took five and served one key with them. A block holds every head's q in shared memory at once; same grid, 64 keys a block. The scores are the vector kernel's bit for bit. GGML_CUDA_LIGHTNING_INDEXER_VEC_LEGACY=1 restores the vector kernel, and GGML_CUDA_LIGHTNING_INDEXER_CHECK=1 runs both and traps on the first score whose bits differ. Measured on an RTX 5080 + RTX 5070 Ti pair, -sm tensor -ts 1/1 -fa 1, on a 45-block proxy (33 KDA layers, 11 DSA, 1 dense lead): - test-backend-ops LIGHTNING_INDEXER passes on both cards for every case, by default, with either switch, and under the check with no trap. 26 new rows cases: f16 and f32 keys, 32 and 64 heads, batch 1, 3 and 64, one and two streams, and a 77-key count deliberately off the kernel's block size. - R3, each mutant failing only what it should. The kernel ignoring the rows (row i for key i) fails the new cases. Adding the in-register products in another order (xor 4's pairs first) traps under the check and passes test-backend-ops without it, since a comparison against the CPU cannot see the summation order. - Bit for bit: PPL at -c 8192 -b 8192, ub 512 and ub 3 --chunks 1, two runs each, is identical to the last digit by default and with either switch (349190.7774 and 348976.2995). - Quad against vector kernel, ncu gpu__time_duration.sum at base clocks with all caches flushed, 20 launches a side on the 5070 Ti: 0.516x at 8258 pooled keys batch 1, 0.424x at batch 3, 0.536x at 65536 f32 keys, with 0.0-0.9 % drift on a repeated arm. The same shapes at boost clocks with a warm cache read 0.449x, 0.445x and 0.350x. Both conditions are real: the warm one is how it serves, the flushed one is the conservative floor. - Wall, tg32 at -d 32768, -r 6, reps 2-6 median, two interleaved runs. Reading pooled keys in place: 89.28 and 88.86 tok/s against 86.80 and 87.77 with the gather, +2.86 % and +1.24 %. The quad kernel: 90.97 and 84.10 tok/s against 91.37 and 87.89 with the vector kernel, so no gain in either run. That is not a readable result rather than a regression: the two runs of the same arm differ by 8 % (90.97 and 84.10, other work sharing the cards during the second), while this lever's ceiling is near 0.5 %. Two band notes, recorded rather than smoothed over. The rows lever was declared at >= 1.0 % of wall in both runs and earned it. The quad kernel was declared at <= 0.4x the vector kernel's time and >= 0.8 % of wall; it misses both, and the second was unreachable by construction: the indexer_pool_score node is 196.8 us of a 21.7 ms token on the 5070 Ti (11 launches, measured by NVTX attribution), so 0.91 % of the wall, and halving it caps the wall gain near 0.5 %. A band set from a shuffle count for a kernel whose roofline is memory is the defect; perf bands belong at a margin over a measured prior.
…vectors, bit for bit Measured first, on an nsys capture at -d 32768 on the RTX 5070 Ti's main stream: a quantize_q8_1 launch ends 3.14 us (median) after the gate/up ring it follows, 43 of them a token, 167 us of serial chain; the ring's down projection then copies that q8_1 into shared memory anyway. Where a routed launch's vectors are each its own (ne11 == n_used, which is a down projection), no other MUL_MAT reads their q8_1, and the plan keeps them in shared memory, the host now passes the f32 vectors and the ring's 512 consumer threads quantize them into that shared copy past the dependency wait: a thread a float4 (eight of them a round trip to L2), 8 lanes a q8_1 block, the max and sum taken in warp_reduce_max/sum<QK8_1>'s own tree (xor 16, 8, 4 across the 8 lanes as xor 4, 2, 1, then xor 2 and 1 within the lane's float4), then d, the rounding and ds exactly as quantize_q8_1 computes them. So the results are the launch's bit for bit. GGML_CUDA_MMVQ_MOE_QUANTIZE_LEGACY=1 restores the launch, and GGML_CUDA_MMVQ_MOE_QUANTIZE_CHECK=1 runs the ring from both and traps on the first result whose bits differ. Measured on an RTX 5080 + RTX 5070 Ti pair, -sm tensor -ts 1/1 -fa 1, on a 45-block proxy (33 KDA layers, 11 DSA, 1 dense lead): - test-backend-ops MUL_MAT_ID 1076/1076 and MOE_FFN_CHAIN 10/10 on both cards, by default and with the legacy switch, and under the check with no trap. - R3: a mutant adding the float4's elements in another order ((0 + 1) + (2 + 3)) traps under the check on the down cases and passes test-backend-ops without it. - Bit for bit: PPL under the check with the legacy switch's value, at -c 1024 -ub 1 --chunks 1 and -c 2048 -ub 3 --chunks 1 (the ring taking 1 to 8 tokens), and identical to the last digit at -c 8192 -b 8192 ub 512, ub 3 and ub 1 (349190.7774, 348976.2995, 321534.0329). - Wall, tg32 at -d 32768, -r 6, reps 2-6 median, two interleaved runs: 90.19 and 91.67 tok/s against 92.75 and 90.30 with the launch. The two runs disagree in sign, -2.76 % and +1.52 %. The wall band asked for >= 0.6 % in both runs and is a miss, but the number to act on is the disagreement, not the mean: 167 us of removed serial chain is about 1 % of a token here, and two runs of one arm span more than that on this host while other work shares its CPU. tg32 at -r 6 does not resolve a 1 % lever. The node attribution this lever was declared against -- no quantize_q8_1 after an mmvq_moe, and the ffn_down_exps node at least 100 us lower on card 1 -- is the instrument that can see it, and it has not run yet; it is recorded as owed beside the declaration rather than papered over with the one positive run.
The front (dsv4_hc_front's Gram path) was two launches a sublayer, 88 a decode token on the 44-layer proxy: dsv4_hc_mix_gram (64 blocks, the dot products and Gram partials over a 64-column slice) and dsv4_hc_pre_gram_f32 (16 blocks, each summing every partial, then the pre weights, the RMS, block 0's Sinkhorn, and its slice of the normed mix and its q8_1). An earlier capture on the RTX 5070 Ti at -d 32768 with graphs off put them at 2.39 and 3.90 us a front, 210 and 343 us a token, while ncu read 0.23 and 0.05 waves with 93 % and 91 % of cycles holding no eligible warp. Latency, not work. dsv4_hc_front_one runs mix_gram's blocks; each block's writers fence, thread 0 takes the token's ticket (ggml_cuda_hc_front_tickets, one unsigned int a stream and token, zeroed before a graph evaluation the way the PQ2_0 tile counters are), and the block holding the last ticket sets it back to 0 and does pre_gram's work for the whole token through the same device functions dsv4_hc_pre_gram_f32 used -- the partial sums, the pre weights and RMS, the weights, each element and its q8_1 -- so the bits are the same. It reads the other blocks' partials streaming through L2 (__ldcg). One launch a front. GGML_CUDA_HC_FRONT_ONE_LEGACY=1 launches the two kernels; GGML_CUDA_HC_FRONT_CHECK=1 runs the two into scratch before the one launch and traps on any differing bit of the normed mix, the weights or the q8_1 copy. Measured on an RTX 5080 + RTX 5070 Ti pair, -sm tensor -ts 1/1 -fa 1, on a 45-block proxy (33 KDA layers, 11 DSA, 1 dense lead): - test-backend-ops DSV4_HC_PRE_FUSED 37/37, DSV4_HC_PRE_POST 5/5, DSV4_HC_PRE_Q8_1 16/16 and DSV4_HC_POST 4/4 on both cards, by default and with the legacy switch, and under the check with 0 check lines. - R3, two mutants, each failing band 1 in its own way: the last ticket taken as n_slices - 2 produces 320 check lines and 5 test failures; the ticket never set back to 0 leaves the next launch on the stream with no last block, 544 check lines and 56 failures, collapsing the suites to 6/37, 1/5 and 1/16. - Bit for bit: PPL identical to the last digit by default and with the legacy switch at -c 8192 -b 8192 ub 512, ub 3 --chunks 1 and ub 1 --chunks 1 (349190.7774, 348976.2995, 321534.0329). - Wall, tg32 at -d 32768, -r 6, reps 2-6 median, two interleaved runs: 86.17 and 86.72 tok/s against 87.23 and 85.69 with the two launches. The runs disagree in sign, -1.22 % and +1.20 %. The wall band asked for >= 0.8 % in both and is a miss. It was also the wrong gate for a lever this size: the two kernels are 553 us of a token on card 1, so cutting them to 0.75x saves about 138 us, near 0.6 % of the wall, while two runs of one arm on this host span 0.5 % at best and several percent when another build shares the CPU. The band this lever should be judged on is its node attribution -- the one launch's median against the two medians, and 88 fewer launches a token on each card -- which has not run yet and is recorded as owed beside the declaration.
…r, the wide mask scan, the indexer reading pooled keys in place with its quad kernel, the ring quantizing its down vectors, and the one-launch hyper-connection front
… measured slower than the two kernels 6fbde0b landed the front as one launch with its wall result unresolved and its node attribution owed. That attribution has now run, and it is a regression. NVTX with graphs off, the last 31 decode tokens at -d 32768, 87.8 fronts a token in both arms: one launch (median, a token) two kernels (medians, a token) ratio RTX 5080 8.48 us, 754.5 us 2.11 + 4.22 us, 563.2 us 1.34x RTX 5070 Ti 8.38 us, 759.4 us 2.15 + 3.01 us, 535.4 us 1.62x against a declared <= 0.75x. The launches a token do fall as designed, 1561.1 to 1473.3, exactly 88 fewer. A sum of two kernel times omits the gap between them, which is the very thing fusing removes, so the NVTX issue span was measured too -- it does include that gap. The DSV4_HC_POST ranges total 1475.4 us a token with one launch against 1395.0 us with two, on card 1. Both denominators agree, so the conclusion holds under the correction. The cause is the ticket's shape, not the fusion idea. dsv4_hc_pre_gram_f32 summed every partial on 16 blocks in parallel; in the fused kernel the block holding the last ticket does all of that work alone for the whole token, after fencing on all 64 blocks. Replacing 16-way parallelism with one block costs more than the launch it removes. The two kernels' 93 % and 91 % of cycles with no eligible warp, which motivated the change, were a symptom of a small grid rather than of the launch boundary. So the two kernels become the default and the one launch is opt-in behind GGML_CUDA_HC_FRONT_ONE=1 -- the same move this fork made for the comb beside the stream at 98a6656 when it measured worse on the cards. GGML_CUDA_HC_FRONT_CHECK=1 now implies the one launch, so the check always has something to compare. Everything 6fbde0b verified is kept: DSV4_HC_PRE_FUSED 37/37, DSV4_HC_PRE_POST 5/5, DSV4_HC_PRE_Q8_1 16/16 and DSV4_HC_POST 4/4 on the 5070 Ti by default, with the one launch, and under the check with 0 check lines, re-run against this change. A next attempt should keep the 16-way reduction: the last-ticket block spreading pre_gram's work across its own warps, or a reduction over more than one block, rather than one block doing all of it.
…anges, the ub-512 logits part on a fusion the layout flips, and its decode cost
…he top-k, not a liveness chain in every layer 4a88ac9 made the rows of the mask's scatter unique with a chain per DSA layer: the exp of the pool bias at the selected pools, repeated to cells, an arange of dump columns, the index arithmetic in f32, casts, and a concat of zero columns onto a copy of sel_mask. It cost ~1.1% of decode kernel time and 128 launches, and its per-layer view of pool_bias, an input, was a split input of its own (24 + 11 > 30), so the meta backend cut a second split and captured two CUDA graphs a token. The top-k now does it. llm_graph_input_kpool carries select_k dump pools after the real ones: pool_dump, an f32 input of -FLT_MAX concatenated onto the pool scores, above a dead pool's -inf and under every live score, so a row takes a dump pool exactly where fewer than select_k pools are live, never a dead one; pool_cells names their cells n_kv + d*kpool + t, and sel_mask carries those columns, -inf in every row, so the scatter writes them, no two slots of a row name one cell, and the view back to n_kv drops them. build_attn_sparse scatters top_k as it is into a dup of sel_mask; `live` is gone from build_indexer and build_attn_sparse. pool_cells' 3-D view is built once in build_inp_kpool, not in each indexer layer: a view of an input made per layer was a split input, and a host copy each token, of its own. test-kpool-input builds the maps with 0, a few or n_pools dump pools and checks, both paths, that the dump cells are n_kv + c and the real ones under n_kv, and that every dump column is -inf in every row; it fails with a dump cell repeated (n_kv + c/r) and with the dump columns left unwritten. test-kpool-can-reuse refuses a change of the dump count.
…rom the top-k, the parent's logits at ub 3 and 1, 95.7 fewer launches a token and the split inputs halved
requires_imatrix was asked of the target type alone, and raised before the cur_type != new_type check that copies an already-correct tensor verbatim. So requantizing a model that holds IQ-family tensors was impossible without an imatrix even when those tensors were untouched: forcing the experts of an IQ3_XXS model to IQ3_XXS, to hold them while the non-expert tensors move, failed on blk.1.ffn_down_exps.weight with "this quantization requires an importance matrix" for a tensor that would have been copied byte for byte. --dry-run only records the requirement in will_require_imatrix instead of raising it, so the dry run passed and the real run failed on identical arguments -- the divergence that makes a dry run worthless as a check. Gate the requirement on the type actually changing. The two later guards (will_require_imatrix at the quantize decision, and the imatrix lookup) read the same flag and would have fired spuriously for the same tensors.
The 8.9 GB of non-expert weight is 70 % of what a glm5next token reads, so the byte lever lives there, and a format that cuts bytes 47 % while losing 30 % of achieved bandwidth is not a 47 % win. Adds the four real shapes -- two KDA projections, the DSA output and the 4096 x 154880 head -- across Q8_0, Q4_K, IQ4_XS, MXFP4 and NVFP4, at one row and at an MTP verify's width. bs=1 takes the vector kernel, where the block-scaled MMA that NVFP4 and MXFP4 exist for cannot apply; bs=3 is where it can, so the ranking may differ between them and both are measured. Caveat for whoever reads the output: perf reports us/run and TFLOPS, never bandwidth, and at these shapes every arm but the head fits in an RTX 5080's 67 MB L2 -- Q8_0 there appears to reach 2816 GB/s against a 960 GB/s DRAM peak. Only the head exceeds L2 in every format, and there all five formats land within 0.7 % of each other at 905-912 GB/s with time tracking bytes to 0.4 pp. The smaller shapes need ncu --cache-control all to say anything about DRAM.
KL divergence was scored only over the second half of each chunk, and the per-token values were sorted before reporting, so position was discarded. For a recurrent architecture that hides the thing worth measuring: a KDA layer carries its state forward, so a quantized tensor written INTO the state (attn_k, attn_v) has its error compound with distance from the start of the sequence, while a tensor feeding only the readout (attn_q) stays flat. The two are indistinguishable in a corpus mean and are not remotely the same risk. At n_ctx 32768 the default also means positions 0-16384 are never evaluated, so no amount of context length produces an early-position bin. - LLAMA_KLD_FIRST sets the first scored position. One helper serves both the base-logits writer and the reader, which must agree: it sets how many values a chunk occupies in the base file. - The three buffers were sized from the hardcoded n_ctx/2 before first was defined; a lower first overran all of them. They now size from first. - The header carries n_ctx, n_vocab and n_chunk but not first, so a base file written at the default and read at 0 would misalign every chunk and report a plausible, meaningless dKLD. The payload length is now checked against what first implies and the run refuses with both counts. - dKLD is reported per position bin (0-512, 512-4k, 4k+) before the sort: count, mean, median, p99 and mean |dp|. An empty bin prints why rather than vanishing, because an empty late bin means the run cannot see accumulation at all. Verified on a 4-layer proxy, three claims: an identical model gives dKLD 0.000000 with p99 0.000002 in both populated bins, so the binning indexes correctly; the bins fill with 2048 and 2044 of 4092 values over 4 chunks and 4k+ reports itself empty; and reading a first=0 base file at the default refuses, "holds 1267570656 payload bytes but this run expects 633165792" -- the 2.0020 ratio of 1023 to 511 values a chunk.
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.
Commits on top of #83 (82a23c9). Every kernel change is behind a
*_LEGACYswitch (the graph-order and harness changes have none), and every change has its row in TORAD.md with its numbers, its tests and what stays open.Qwen3.5 and glm5next KDA
331d727: Qwen3.5 levers between the qkv group and the recurrence, each behind its own switch.
mul_mat_vec_ftrigger the next launch at their start.Served bits (Ternary Bonsai 2 27B, one slot and four, 128 greedy tokens with top-5 log-probabilities) are identical to all switches set. Open, in its row: the fold does not engage at this commit; the caps leave 16–32 B stacks; the Qwen speed sweep that sets the defaults is still to run.
e74ed87: glm5next's KDA output gate goes into the graph before the recurrence.
g_aruns besidef_aas one mat-vec launch,g_bbefore the recurrence: one small mat-vec launch fewer a KDA layer, the chain −0.36 % (RTX 5070 Ti).7536ed7, then b850d70: the KDA gated norm written by the recurrence's kernel, then taken out. Bit-identical, and neutral against e74ed87 over 4 rotated rounds (tg128 +0.24 %, CI −0.95 to +1.42); the norm's time is the wait for the state stores to drain, which stays on the chain wherever the norm runs. b850d70's tree is e74ed87's.
GLM-5.3 decode (44-layer proxy, RTX 5080 + 5070 Ti)
-sm tensor. NewMOE_FFN_CHAINtest.MOE_WEIGHTED_SUMtest, compared bit for bit.DSV4_HC_PRE_Q8_1test.DSV4_HC_PRE_POSTtest. Opt-in since 98a6656 (below).test-backend-opsruns the reference first, so a kernel a backend leaves running past its return shows as a failure.GLM-5.3 MTP verify (3 tokens)
-sm tensor. PPL at-ub 3moves 385507.0946 → 385634.0270: the fused kernel's own normalization rounding, as a decode's always had (KLD vs-ub 10.006763 against 0.006756).-sm tensor: kernel time −7.01 % (5080) and −6.52 % (5070 Ti).-ub 3: +9.68 % (CI +8.58 to +10.79) on one card, +6.63 % (CI +4.33 to +8.93) under-sm tensor.-sm tensor: both rings faster in every round on both cards.mul_mat_fwould run on fewer blocks than the GPU has SMs takesmul_mat_vec_f, and pairs on one input fuse. That covers glm5next's KDA f_a/g_a (4096 → 128 bf16, 8.6 us each on 4 blocks) and beta (6.4 us on 2 blocks).mul_mat_f's time on the 5070 Ti and 0.23–0.63 on the 5080.-sm tensor. pp3 +4.26 % (CI +1.51 to +7.02).-d 2048, where the DSA indexer scores: −3.1 % kernel time.mul_mat_frounds them to bf16, and sums in another order. PPL at-ub 3385634.0270 → 388534.0622. The KL divergence against the switch, 0.006795, matches-ub 1against-ub 3(order alone), 0.006763. Bit for bit at-ub 1and under-sm tensor.-p 3/-p 4on 2 RTX PRO 6000. Measured locally:GGML_CUDA_HC_COMB_SIDE=1). GLM-5.3-Flash's tg64 was 3.7 % faster without it on 2 RTX PRO 6000 (4 of 4 pairs). Locally it is only +0.70 % (CI −0.15 to +1.55). Same bits.PPL is bit for bit wherever a row says so: 386825.5220 at
-ub 1, 387892.2627 under-sm tensor, and at-ub 3385507.0946 before 5c03ed9, 385634.0270 after it, and 388534.0622 from 18fa139.Gates
test-backend-opsMUL_MAT 1475/1475, MUL_MAT_ID 1076/1076, MUL_MAT_VEC_FUSION 1032/1032, GATED_DELTA_NET 71/71, GATED_DELTA_NET_CACHE_FUSION 77/77 (and with the L2 persistence, PDL on and off), DSV4_HC_POST 4/4, DSV4_HC_PRE_FUSED 37/37, LORA_RANK1 3/3, SSM_CONV 45/45, SSM_CONV_STATE_UPDATE 48/48, MOE_FFN_CHAIN 10/10, MUL_MAT_ID_FUSION 13/13, MOE_WEIGHTED_SUM 18/18, DSV4_HC_PRE_Q8_1 16/16, DSV4_HC_PRE_POST 5/5, TOPK_MOE 320/320, ARGSORT 98/98, MUL_MAT_PAIR 120/120; DSV4_HC_PRE_FUSED, DSV4_HC_PRE_POST and DSV4_HC_PRE_Q8_1 again withGGML_CUDA_HC_COMB_SIDE=1; test-llama-archs; test-backend-meta split, views and capture on both cards.test-backend-opsMUL_MAT 1433/1433, MUL_MAT_ID 1074/1074, MUL_MAT_VEC_FUSION 1032/1032, GATED_DELTA_NET 71/71, GATED_DELTA_NET_CACHE_FUSION 77/77 (and with the L2 persistence, PDL on and off), DSV4_HC_POST 4/4, DSV4_HC_PRE_FUSED 37/37, LORA_RANK1 3/3, SSM_CONV 45/45, SSM_CONV_STATE_UPDATE 48/48, MOE_FFN_CHAIN 10/10, MUL_MAT_ID_FUSION 13/13, MOE_WEIGHTED_SUM 18/18, DSV4_HC_PRE_Q8_1 16/16, DSV4_HC_PRE_POST 5/5, TOPK_MOE 320/320, ARGSORT 98/98; test-llama-archs; test-backend-meta split, views and capture on both cards.test-backend-opsgroup they touch, test-llama-archs, and the meta split/views/capture tests on both cards.GLM-5.3-Flash: the sparse indexer, the mask scan and the hyper-connection front (Sep 30)
Eight commits on a 45-block proxy (33 KDA layers, 11 DSA, 1 dense lead), RTX 5080 + RTX 5070 Ti,
-sm tensor -ts 1/1 -fa 1. Every lever carries a*_LEGACYswitch, and the three that claim bitequality carry a
*_CHECKswitch that runs both paths and traps on the first differing bit.write's unused slots. Its own evidence is still being settled by its author and its row says so; it is
not claimed bit-identical here.
handle_genericreturnedUNKNOWNfor such a nodeand allocation aborted on
GGML_ASSERT(ret.axis != GGML_BACKEND_SPLIT_AXIS_UNKNOWN); ARANGE routesthere and the sparse mask's dump columns are an arange, so every GLM-5.3 decode under
-sm tensoraborted in
ggml_gallocr_alloc_graph. Newtest-backend-meta-sourceless: red on the parent (exit134), green here. Upstream master carries the same
UNKNOWN.ggml-org/llama.cpp
b74f590ea(ggml-cuda: fix divergent barrier in f16 flash attention ggml-org/llama.cpp#27870). compute-sanitizer synccheck: upstream's repro 3680 errors →0, GLM-5.3's cases 512 → 0. PPL bit for bit;
flash_attn_ext_f16within 1.79 % of base by ncu at baseclocks with L2 flushed, against the instrument's own 1.44 % spread.
each, a block-wide scan placing each thread's cells) where upstream's read 2048 with scalar loads.
Index lists identical, so attention is bit for bit. Scan median a launch: 2.56 / 3.26 us at
-d 8192and 4.54 / 4.42 at
-d 32768, against 3.6 and 11.2 before. TheFLASH_ATTN_EXTnode at-d 8192is0.62x / 0.58x of the dense path's, all three kernels counted. Wall tg32 @ d32768: +6.8 % and
+4.1 % over two interleaved runs. Two new
FLASH_ATTN_EXTcases at 40960 mask cells cover a secondscan round; three mutants fail the sparse cases and only those.
teaches to read rows.
ggml_lightning_indexer_rows:GGML_OP_LIGHTNING_INDEXERgains an optionalsrc[4], I32 rows, sokey i of stream s is k's row
rows[i, s]. glm5next passes the f16 index cache's pooled head andpool_repsinstead of materialising aget_rowsf32 copy of every pool's key — at 32K cachedtokens an ~20 us gather a layer, 8258 blocks and 4.2 MB written, 11 times a token. CPU reads the
rows; Metal and SYCL refuse the variant; the meta backend mirrors it. Wall +2.86 % and +1.24 %.
lightning_indexer_kernel_quad: a key scored on a quad of lanes, two shuffles a key and head wherethe vector kernel took five and served one key with them; the vector kernel's scores bit for bit.
By ncu at base clocks with all caches flushed, 20 launches a side: 0.516x / 0.424x / 0.536x at
8258 pooled keys batch 1 and 3 and at 65536 f32 keys, 0.0–0.9 % drift on a repeated arm; the same
shapes at boost clocks with a warm cache read 0.449x / 0.445x / 0.350x.
LIGHTNING_INDEXERrows cases; PPL identical to the last digit at ub 512 and ub 3, two runseach; two mutants, one of which (products summed xor-4-first) is invisible to a CPU comparison and
is caught only by the check kernel.
quantize_q8_1launch ended 3.14 us after the ring it follows, 43 a token, 167 us of serial chain, and the ring then
copied that q8_1 into shared memory anyway; its 512 consumer threads now produce it there past the
dependency wait, through the same reduction tree and rounding, bit for bit.
Each block's writers fence, thread 0 takes the token's ticket, and the block with the last ticket does
the second kernel's work for the whole token through the same device functions, so the bits are the
same. The two kernels were 2.39 + 3.90 us a front with 93 % / 91 % of cycles holding no eligible warp
— latency, not work. Two mutants fail band 1 in distinct ways (320 and 544 check lines).
Three wall bands missed, and why the misses are the instrument's
c40b9d5ce(declared ≥ 0.6 %),6fbde0b8e(≥ 0.8 %) and the quad kernel in252f29c0a(≥ 0.8 %) allmiss, and two of the three produced runs disagreeing in sign (−2.76 % / +1.52 %, and −1.22 % / +1.20 %).
That is the instrument, not the code:
indexer_pool_score196.8 us, halvedTwo runs of the same arm span 0.5 % at best on this host and 8.0 % when another build shares its CPU
(90.97 and 84.10 tok/s on one arm).
tg32 -r 6therefore cannot resolve a sub-2 % lever, and the quadkernel's ≥ 0.8 % band was unreachable before a single run — the whole node is 0.91 % of the wall.
Below ~2 %, node attribution is the deciding instrument and the wall is a sanity check. The bands were
declared from each change's shape rather than from the node's measured share, which was already in hand;
each miss is recorded beside its declaration rather than dropped.
Where the two cards actually wait
From an NVTX capture at
-d 32768with graphs off (so host time is an upper bound), over the last 31decode tokens, every >100 us idle gap charged to the nodes on either side of it:
(55.1 % idle).
is 4.3 ms (card 0) and 8.7 ms (card 1).
DSV4_HC_POSTnode, 6.58 ms in total; card 0 waits afterits
dsa_out/kda_out/ffn_shexpmatmuls, 3.48 ms. The same per-layer tensor-parallel dependencyseen from both ends.
3.5 % of a token, so not where the time goes.
perfect overlap is 21.7 → 13.8 ms a token, 1.57x. Beating it needs a
-tsother than 1/1, measured.This is why single-kernel levers of the sizes above cannot move the wall much however good the kernel is,
and it names the next lever: fewer, later sync points a layer, and a weighted split.
Gates (continued)
test-backend-opsMUL_MAT 1475/1475, MUL_MAT_ID 1076/1076, MUL_MAT_VEC_FUSION 1032/1032, GATED_DELTA_NET 71/71,
GATED_DELTA_NET_CACHE_FUSION 77/77, DSV4_HC_POST 4/4, DSV4_HC_PRE_FUSED 37/37, LORA_RANK1 3/3,
DSV4_HC_PRE_POST 5/5, DSV4_HC_PRE_Q8_1 16/16, MUL_MAT_PAIR 120/120, TOPK_MOE 320/320,
FLASH_ATTN_EXT 3229/3229; the cache-fusion persist leg at PDL on and off (77/77 each);
test-llama-archs; test-backend-meta split, views, capture and sourceless; and the proxy under
-sm tensorwith CUDA graphs on and off, 4 of 4 rows and 0 asserts.*_LEGACYswitch, andunder its
*_CHECKswitch where it has one.Final node attributions, and one lever turned off by them
The node attribution the three missed wall bands pointed at has now run (NVTX, graphs off, the last 31
decode tokens at
-d 32768, both arms from the same bin). It confirms two levers, bounds a third, andreverses the fourth:
f8352b6c2mask scanFLASH_ATTN_EXT268.4 → 666.4 us, 0.403x (card 0 0.452x)252f29c0akeys in placeGET_ROWSa token are gone252f29c0aquad kernelindexer_pool_score164.7 → 230.0 us, 0.716x (card 0 0.598x)c40b9d5cering quantizeMUL_MAT_ID ffn_moe_down862.5 → 1085.9 us, −223.4 us; the 43quantize_q8_1-after-mmvq_moea token are gone6fbde0b8eHC fronte1c1ca7d8Two notes on method, both of which changed a conclusion:
ffn_moe_gatenode is the ring lever's control: it is unchanged (1545.1 vs 1548.4 us), so the nodethat should not have moved did not.
removes — the wrong denominator for a latency lever. The NVTX issue span, which includes that gap, was
measured too: 1475.4 us a token with one launch against 1395.0 with two. Both agree, so the regression
stands. Cause:
dsv4_hc_pre_gram_f32summed every partial on 16 blocks in parallel, and the fusedkernel's last-ticket block does all of it alone after fencing on 64 blocks. The 93 % / 91 % of cycles
with no eligible warp that motivated the change were a small grid, not the launch boundary.
e1c1ca7d8therefore makes the two kernels the default and the one launch opt-in behindGGML_CUDA_HC_FRONT_ONE=1— the same move98a6656aemade for the comb beside the stream. Every property6fbde0b8everified is kept and was re-run against the flip: DSV4_HC_PRE_FUSED 37/37, DSV4_HC_PRE_POST5/5, DSV4_HC_PRE_Q8_1 16/16, DSV4_HC_POST 4/4 by default, with the one launch, and under the check with 0
check lines.
GGML_CUDA_HC_FRONT_CHECK=1now implies the one launch so the check still has something tocompare.
Also earned here:
252f29c0a's KLD band, with its own control — default against the switch reads mean KLD0.000000, max 0.000003, same-top 99.992 %, and the switch against itself reads exactly the same, so the
3e-6 maximum is the run-to-run floor rather than the lever.