Repository navigation
Living Weights on Qwen3.8 Flash Next (CUDA, GB10): the pocket in every pass, the trainer, and --slide - #896
ZackSample wants to merge 57 commits into
Conversation
A resident prompt state died (its slot's rows overwritten, or its pool blocks given up) left the Store serving the dead entry: every later lookup found it, its restore failed, and the pass prefilled the whole prompt from the start — the repeated full-refill cascade seen on live servers with long conversations. The backend now tells the prompt cache the moment a state goes stale (Slots.expire), so the Store drops that entry immediately and the next turn keeps a fresh state instead of paying for the same failure again.
The prefill path read a snap after pool admission could free it, and overwrite walked snaps while expiry dropped them. A resident state that dies now expires its cache entry at the host boundary; each lookup borrows its entry and snap until resumed or release ends that ownership; reclamation runs at host boundaries and skips active borrowers.
# Conflicts: # zig/src/core/prompt_cache.zig
…unchanged The learner, its whole-lesson and routed paths, the pocket and its fits and helpers move from families/nemotron to a `living` module that both builds import, as 19 renames with no change to their contents. The family-neutral half of learn_shape.zig (block sizes, Adam settings, direction lengths) moves with them; Nemotron's learn_shape.zig keeps where its blocks sit and names the shared constants from the module. Nemotron's files import the moved ones through the module. The learner's tests now also run on Linux, as `zig build test-living` and under test-cuda-host. No behaviour change.
… layout, neuron kind and shard format A family now describes itself once, as data (living/families.zig): where its shared experts sit and their tensor names, the neuron's activation (relu^2, or SwiGLU with an optional per-row gate), the quantization, the config keys and the object they live in, and its MTP head's expert. Its host gives the learner a layout and the pocket hooks (living/family.zig). Nemotron is the one family listed, and its lessons, routed blocks, edits and probes are unchanged. - The pocket reads its layers and width from the family's layout, not from Nemotron's config, and the learner offers routed blocks, edits and probes only to a trainer with `sites` (Nemotron's). - living/shards.zig widens any family's experts at its group and width; Nemotron's writer is now a wrapper over it. shard_edit sets and reads nested config keys (text_config). - A SwiGLU neuron fitted as u is written as gate row 16u and up row u/16: relu(u x)^2 on its fact's rows, at most 8e-5 of it at a 0.5 margin; a gated family's outputs fit gate * act(x). No listed family uses it yet; the tests use a stand-in family shaped as Qwen3.8 Flash Next's shared expert. - The server's view reads layers, pockets and shards through the family its config's model_type names, and native/nemotron_slide.zig is native/slide_adapter.zig. - pocket.zig's quiet refit and its loading move to pocket_quiet.zig and pocket_load.zig, for the 600-line rule. - Tests: the generic writer equals affine4_host at 4 bits in groups of 64, Nemotron's save test over the wrapper, a SwiGLU neuron fitted and stored at 4 bits in groups of 32, and a whole pocket lesson over a stand-in host for each activation (test-living, also under test-cuda-host).
…roup-32 token, so the CUDA replay starts from the row the Python engine reads.
… so the linear-attention input is that mixed row.
…f the mixed row on one token, so the DeltaNet input is that projection.
…y match the host chain on one token, all 6144 gated values and all 2560 projected values, so the mixer branch is that row.
…, the write-back exact and its sums within 1 ulp, one normed value 1 step off, and the mixed row within 2 steps, so the expert input is that row.
…he shared expert exact, router logits within 29184 ulps, two weights 1 step off, the activations exact, and every down value the same bf16 with fp32 within 16384 ulps.
…the projection, the prepared heads, the one-key output, the gate and the output row all exact.
…one projected value 1 step off, and the prepared heads, the one-key output, the gate and the output row all exact.
… one projected value 1 step off, and the prepared heads, the one-key output, the gate and the output row all exact.
… the projection, the prepared heads, the one-key output, the gate and the output row all exact.
… the projection, the prepared heads, the one-key output, the gate and the output row all exact.
… the projection, the prepared heads, the one-key output, the gate and the output row all exact.
… two projected values 1 step off, one gate value 1 step off with group sums within 1024 ulps, and the prepared heads, the one-key output and the output row all exact.
… one gate value 1 step off with group sums within 8192 ulps, and the projection, the prepared heads, the one-key output and the output row all exact.
… one prepared key value 1 step off, and the projection, the other prepared heads, the one-key output, the gate and the output row all exact.
… one projected value 1 step off, and the prepared heads, the one-key output, the gate and the output row all exact.
… one projected value 1 step off, and the prepared heads, the one-key output, the gate and the output row all exact.
… two projected values 1 step off, and the prepared heads, the one-key output, the gate and the output row all exact.
…all 2560 values and the group sums exact.
…token, all 2560 mixed values exact, and three of 248320 scores 1 step off.
…ing every expert onto the GPU once.
…m; prompts past 512 keys; LR7 capture. The ids and position uploads were synchronous copies, which do not wait for the non-blocking compute stream, so a later row overwrote a queued row's token id or position: wrong tokens on most prompts. Both now go on the stream. attnStep launches one 512-key chunk per started chunk of the row's keys (as attention.attention does) instead of always one, so prompts run to the captured NCH * 512 keys; the lanes engine defaults to that window and refuses a larger one. The prompt check prints each run's top five scores and its timing. zig/tests/cuda/flashnext: capture.py and capture.sh (Python engine 0.6.1, drafts off, greedy, 1536 cache slots) and the Unbound-LR7 prompts with their references; flashnext-prompt matches the teacher paths on all three.
…ules and device weight bytes from what it declares; Flash Next registers The registry gains flashnext.native. What was Nemotron's by name is now read from the family's declarations: its prompt-cache copies (snap_bytes, snap_save, snap_restore, snap_drop, and snap_ready's spare buffers), the store's rules (cache_rules; without them a state keeps at any row with the MTP head's one token on, ashhart#670), the quiet wait before a warm pass and the lone driver's hand-over budget. Nemotron declares the copies, the wait and the budget, so its paths are unchanged; a family without copies gets no prompt cache, and the startup line says so. A family that keeps part of its checkpoint off the device says what it uploads (deviceWeightBytes) and admission counts that, not every shard: Flash Next leaves its n-gram tables mapped and loads no vision tower, so the MLX 4-bit checkpoint uploads 73.3 of its 105.4 GiB, where counting every shard made the default reserve on a GB10 refuse it by 0.1 GiB. --slide is refused before anything loads for a family that declares no learner. The Flash Next family names what the host reads: one segment, its prompt rows, each stream's bytes, its refusals in words, and the native host's lone-driver signature.
…the decode kernel's sums: experts_rows.cu experts_rows.cu is experts.cu's decode expert kernel with two 16-row tiles a warp, so a warp reads each group of an expert's weights once for up to 32 pair rows. Every row keeps the decode kernel's arithmetic: the same group sums, the same two m16n8k16 steps from zero a group, the same fma epilogue order and the same SwiGLU, built with experts.cu's flags. A row's bits therefore do not depend on how many rows share its window or chunk, which is what the lane core's exactness contract needs of a lane kernel. It is instantiated at groups of 32 (gate and up stacked with SwiGLU, and the down projection to fp32), the MLX 4-bit group-32 experts Flash Next reads. cuda.experts.RowsExperts resolves the two instances and launches them on the shared plan (cuda.grouped.Router at its 32-pair tile). `tf-cuda-test experts-rows` packs random group-32 experts with experts_pack.cu, routes 1, 2, 5, 16, 33, 100 and 300 rows (three routed experts and a shared one a row) and checks every pair row of the rows kernel against experts.cu's one-row launches of the same row, byte for byte, for the SwiGLU rows and the fp32 down rows.
…te_rows.cu) tf_route_logits gives OUT[m, e] = x[m] . w[e] in fp32 for bf16 rows and router rows on mma tiles: one chain over K in 16-wide steps from zero, the inputs staged with cp.async and read back with ldmatrix. A row's sums never depend on the row count or the tile, so the few-row form (up to 16 rows, three 256-wide stages in flight) and the prompt form (two 128-wide stages) give the same bits. tf_route_topk takes a row's top k of its first E logits (the lower id among equal ones) with weights bf16(e^(l_k - l_0) / their sum, summed in pick order), and optionally a shared expert as slot k at bf16(sigmoid(bf16(logit E))), the softmax routing Qwen-style MoE families use. core/route_ref.zig is the float64 reference, and tf-cuda-test route-rows checks both kernels against it at 1, 3, 16, 17 and 40 rows (logits within the fp32 dot's bound, picks equal to the reference rule's on the kernel's own logits, a row of ties too, weights and the shared gate within half a bf16 ulp plus 2^-18), and every row count's bits equal.
…ith their group sums embed_rows.cu's affine 4-bit rows take the scale group (32 or 64) and a copy count (EmbedRows.affine4Groups; affine4 keeps groups of 64 and one copy, so Nemotron's launches are unchanged). norms.cu gains tf_rms_rows_sums (Norms.rmsSums): bf16 rows read x_ld apart, y = bf16(x * rsqrt(mean(x^2) + eps) * w) with fp32 weights and IEEE division and square root, then each group's sum of the stored y, for a group-quantized matmul's input. tf-cuda-test norm-rows checks both against float64 at every row count (Flash Next's 2,560-wide embedding in 4 copies, its MTP head's 2,560- and 10,240-wide norms) with the bits of every row count equal.
…attn.cu, against float64 Our kernels for every entry point the batched pass took from a captured Triton set, with the Python engine's formulas and bf16 rounding points, IEEE division and square root and the core's within-one-ulp exp (glue_math.cuh), built without contraction: - flashnext_hc.cu: the hyper-connection write-back (stream sums alone, the MoE slots, K slices), the down projection with its normed rows (the tiled 4-bit dots on mma tiles, groups in order, per K slice), the slices' sum with the bf16 SiLU and inject gates, the up projection with the stream mix, the n-gram gate and dilated conv, and the MTP head's stream adds; - flashnext_attn.cu: the q, k and indexer heads' norm and RoPE with the cache writes, the output gate, the indexer's pooled blocks and block scores, and attention over 512-key chunks (q . k and p . v on mma tiles, an online softmax with fp32 sums, dense or through a sparse row's list) with the chunks merged in order. Every kernel keeps a row's own sums in a fixed order, so a row's bits are the same at any row count; attention has a form that reads the next tile ahead for a decode round's few rows, with the same sums. cuda_glue.zig launches them (Glue), with the core's token rows, RMS rows and route rows for the embedding, the MTP head's norms and the router. tf-cuda-test flashnext-hc and flashnext-attn check each kernel at the model's shapes against a float64 chain that rounds where the kernel rounds: each bf16 output within one bf16 ulp (an intermediate rounding can tip either way), fp32 sums within their sum-of-|terms| bound, attention within 2^-8 of sum p |v| / l plus half an ulp (p enters p . v as bf16), copies and the stream adds exact, and the bits of 1, 3 and 17 rows (attention: 1, 3 and 20, both forms) equal.
…y entry point is used for comparison The engine no longer needs TENSORFOLD_CUDA_KERNELS: every launch Tri made from a captured cubin now runs our kernels (cuda_glue.zig) unless a set is loaded. A set given by TENSORFOLD_CUDA_KERNELS, or found beside the binary, replaces them only when it holds all 15 entry points the pass launches (17 with the MTP head); one without them, Nemotron's say, is passed over with a log line, and the startup line says "own glue kernels" (native/cuda.zig asks the family's Loaded.own_glue; aot.Set.has names a set's functions). flashnext-prompt takes native in place of a set's folder. flashnext-compare loads the engine once and runs a prompt on our glue and on a captured set: every prompt row's scores (top-1 agreement, the largest score difference, KL, each path's perplexity on the prompt's own next tokens), then each path's greedy reply and the token where they part.
…ue beside a captured set With a captured set loaded, TF_FLASHNEXT_OWNED (all, or a list of Triton names such as _chunks,_merge) makes those entry points run our kernels and the rest the set's, so a difference can be put down to one kernel. Glue reads it once at load into a mask; without a set every entry is ours, as before. flashnext-compare takes it on its own side, so it compares the captured set against the captured set with the named entries swapped.
…for comparison The recipe, the RUNBOOK line, the CUDA and Flash Next pages and the CHANGELOG say the engine needs no capture: every kernel is built from the repository, TENSORFOLD_CUDA_KERNELS names a captured set only to compare against the Python engine's teacher path, and TF_FLASHNEXT_OWNED swaps single entry points. The recipe lists the new tf-cuda-test cases and flashnext-compare, the semantic settings the config check refuses, measurements on our kernels beside the captured set's, and where our tokens part from the teacher path. The family recipe's lane-kernel row names route_rows.cu.
|
The same five lessons with the learner fixes applied The lesson receipt above runs the learner as it is on
So on this sequence, the fixes bring back what #870 took (lesson 2), make the take-out exact and halve the failing |
…rnels run The check compared each label only with linear_attention, so a full-attention layer labelled sliding_attention (or any other word) was accepted and served as full attention. Every label now has to be linear_attention or full_attention, as the layer's place in the stack says. The new test relabels full and linear layers in turn and expects layer_types back; on the old check the full layer relabelled sliding_attention came back accepted.
…he staging is reused or freed A pass queues its uploads from the scratch's pinned staging, then records an event behind them, and the next pass waits on that event before it writes the staging again. If a later upload or the record failed, an earlier upload of the same pass stayed queued while the event still marked the previous pass, so the next pass, or the scratch's free, could change or release bytes a copy had yet to read. The staging now notes an upload queued after the last record and, until a record covers it, waits by draining the stream instead of on the event. Scratch.free waits the same way before it frees the staging. The new test runs a clean pass and then one whose second upload, or whose record, fails, against a stand-in queue, and expects nothing queued at the next wait; waiting on the event alone, as before, leaves one copy queued.
…t within float64's tolerance (tf-cuda-test gdn-back)
…erts with a shared one, and softmax top-k routing train_q4.cu with cuda.train_q4: a linear's input gradient on tensor cores (split K for a head), the experts' products over MLX stacks or experts_pack.cu's packed table, each pair's SwiGLU, and the top_k softmax and sigmoid shared gate back through the router. tf_train_sort_pairs takes up to 1024 experts. tf-cuda-test q4-back and moe-back check them against float64 at 64 and 256 rows, each reference against central differences; packed experts give the stacks' bits.
…d's q then its gate, RMS-normed q and k, partial RoPE) train_attn.cu with cuda.train_attn: the prepared heads again, the output gate, causal attention to 2048 keys with train.cu's key and value part, the heads' RoPE and norms undone, and the q, k and v projections' input gradients through cuda.train_q4. tf-cuda-test attn-back checks every gradient against float64 at 64 and 256 rows (24 query heads, 2 kv heads, head dim 256), the reference against central differences.
…nts, and fn_back over the core's layers flashnext_train.cu: the write-back's and read-out's backward (mix, activations, down and inject, the streams' RMS norms), and the pocket's SwiGLU neurons beside the shared expert behind its gate: the forward in two kernels (tf_fn_pocket_in, then _out, tiled for many rows; a layer with no neuron launches nothing), its input gradient, and its output vectors' gradient. flashnext.fn_back composes them with cuda.train_q4 and cuda.train_attn on one stream. tf-cuda-test fn-hc-back and fn-moe-back check against float64 at 64 and 256 rows; fn-pocket checks the two kernels against tf_fn_pocket_rows byte for byte over 444 cases; fn-back runs every input-gradient case the trainer uses.
…w-c-fnback (ashhart#869 and the backward kernels) Flash Next's trainer needs both open PRs: the CUDA family (ashhart#869) and the learner in zig/src/living (ashhart#878). The two touched zig/build/cuda.zig and native/cuda.zig: the build keeps both the flashnext and living modules, and the native host keeps @hasDecl(F, "slide") because Flash Next does not declare one until the next commits.
…hed pass's keeps for a backward A folder Living Weights taught holds wider shared experts: the engine packs their own 640 rows as before and keeps the pocket's neurons in fp32 (cuda_pocket.zig), and every pass (prompts, drafted windows, rounds) adds them behind the shared gate on the layers that have one, through tf_fn_pocket_in and _out. Such a layer's MoE output is summed in fp32 (fn_moe_sum, MODE 2's order), the pocket added, and written back rounded once (MODE 4 of one slice); every other layer writes back as before, so an untaught folder's bits are unchanged. The config check takes the shared expert's width with the pocket's neurons (tensorfold_learned_neurons). A pass may keep what a backward reads: each read-out's streams, normed rows, v, up rows (fn_hc_upmix_up, the same mixed bits) and inject gates, the mixer and MoE branches, the MoE's input and slots, attention's projections and gated rows, and the final mixer's read-out.
…lashnext-learn against the serving forward cuda_train.zig runs the learner's passes (keys, loss, target) with sequences side by side on the engine's own batched pass (cuda_batch_train.zig: the row-wise launches over every row, the n-gram block, DeltaNet and attention one sequence at a time from zeros), the open fact's pocket neurons scaled per sequence and each nudge added to its layer's MoE output; then the loss's gradient from the 4-bit head down to the lowest nudged layer with flashnext.fn_back and cuda.gdn_train, reading the engine's packed experts where they sit and MLX copies of the small linears (cuda_train_buf.zig, 10.2 GiB beside the engine). A captured set's glue keeps no up rows, so the trainer runs on our own. zig build flashnext-learn: a row's scores never depend on later rows; the trainer's MoE inputs, shared gates and scores equal the serving forward's byte for byte (untaught and with a pocket neuron); packed sequences give each one's loss as alone; the push against finite differences over four windows of the prompt (the direct path's median must agree at small steps; the cross-row paths are reported); a pocket neuron moves the forward as gate act v.
…: qwen4_exp in the core's families living/families.zig lists qwen4_exp: a SwiGLU shared expert at every layer behind its own gate, 4-bit groups of 32, keys under text_config, the MTP layer's expert in the index; the stand-in test family goes, its tests now read it. flashnext/living.zig gives its layout and writes its pocket through the core's shard writer (pocket_shards_test.zig: the first save widens every expert and the MTP head's with zeros, a later one lands in place). cuda_slide.zig is the learner's host over the trainer, and the lane backend opens with --slide: the trainer's pack of rows in the pass's scratch, the layout from config.json, its bytes kept out of the streams' room. A captured comparison set refuses --slide; Macs refuse it for this model as before.
47748f7 to
bad3a2a
Compare
…'s and judge's reads see the model as shipped A lane stream's untaught flag rides its Seq. A prompt or window pass of an untaught stream runs without the pockets; a decode round of untaught streams alone leaves them out, and a mixed one scales the pocket by 0 on the untaught rows, as Nemotron's pocket skips untaught segments. fn-pocket checks that a 0-scaled row keeps every output bit.
|
514e4ee makes Flash Next honour |
|
The same five lessons with the two lesson-writer fixes added My last comment found that the name fact's loss at lesson 3 was the corrections rule (#678) matching two different
I added both commits to the local comparison build from my last comment (this branch with #876, #877 and #891 ported).
So on this sequence, the headline goes from 4 of 12 to 9 of 12 once the learner fixes and the two writer fixes are both |
|
Thanks @ZackSample for separating the learner and writer effects. The latest held-out result is 9/12 with both sets of fixes, not the earlier 4/12; lesson five still fails, and unrelated replies are not all byte-identical. We are carrying the learner repairs onto the selected Flash Next CUDA engine while retaining the landed writer safeguards, conversation checks, exact take-out and untaught streams. The zero-scaled-row control is useful, but the taught-folder served untaught check is still owed. Retain that case, restart parity, retention and save/reload in the integrated gate. The improved recall receipt does not qualify the remaining fit failure. Posted by Ash's AI Assistant Synthena |
…son gate (#895, #896) Flash Next on CUDA takes --slide. A lesson's neurons go into its MoE layers' pockets through the shared learner and the gate Nemotron already uses: the conversation reader, the thresholded quiet refit, up-row neuron counting, family validation, context-aware rechecks, and delete rounding and put-back. A rejected repair restores every neuron, positive row and retired lesson, and the final fit is rounded by the shard encoder and checked for stability before it is judged. A save that fails after an accepted lesson keeps the live values and leaves an admission marker, so no engine opens the folder until a checked retry or a rollback; close drains the trainer and its sources before it frees them. The gradient, train_q4, train_attn and Flash Next training kernels are new images whose launches stay unqualified until their records are measured. Existing kernels are unchanged; Flash Next's two serving images keep their old bytes and gain the training entry points after them. Co-authored-by: ZackSample <134018356+ZackSample@users.noreply.github.com> Co-authored-by: akol1 <15747200+akol1@users.noreply.github.com> Co-authored-by: BobClawblaw <288991054+BobClawblaw@users.noreply.github.com>
…896, #911), which move to 1.1.1 This reverts a291b68de and 4e0d7927c with the follow-ups that only served them: ca7a0882d (Nemotron's learning drain), 493cdf5d4 (the save marker), 650fac5d7 (the MoE backward pack), the learning records in cda734f7d, ec141ed97 and d9566f585, and the trainer tests that 2708455af moved into cuda_close_tests.zig. The training images leave the kernel list and the CUDA contract's image list, gradient_pending, qwen_learning_pending and registry_learning go, and the 13.2 set and the live audit keep the serving, prompt, GLM and paging records. The lean moves that later commits build on stay: the kernel list remains in cuda_kernels.zig and the close-test file goes only because every test in it was a trainer's.
4 participants
Stack
This sits on two open PRs, so it opens as a draft until they land:
lw-c-fnback, 4 commits on Flash Next on CUDA: the qwen4_exp family on the native CUDA host, every stream through the lane rounds #869's head 7746bb9);lw-b-core75978a1).The branch is
lw-c-fnbackplus one merge commit oflw-b-core(46627a8), then the 4 commits below (the fourth is docs). That merge isthe only one: the trainer needs the CUDA family from #869 and the learner in
zig/src/livingfrom #878, and neithercontains the other. It touched
zig/build/cuda.zig(both modules kept) andnative/cuda.zig(the host keeps@hasDecl(F, "slide")until Flash Next declares one). Once #869 and #878 land, I will rebase this ontomainand themerge goes away; the diff to review now is
lw-c-fnback...lw-d-fnlearnafter the merge, or the 4 commits.What this changes
Weights neurons. Now the engine packs each shared expert's own 640 rows as before and keeps the pocket's neurons in
fp32 (
cuda_pocket.zig); every pass (prompt chunks, drafted windows, decode rounds) adds them behind the shared gateon the layers that have any, through nemotron cuda: skip empty pocket layers and spread the pocket over many blocks #787's two-kernel form ported to this family (
tf_fn_pocket_in, then_out;a layer with no neuron launches nothing). On such a layer the MoE output is summed in fp32 in MODE 2's order
(
fn_moe_sum), the pocket added, and written back rounded once (the write-back's one-slice MODE 4). Every otherlayer writes back exactly as before, so an untaught folder's bits do not change (receipt below). The pocket's
kernels are row by row, so drafted equals
"draft": falseand concurrent equals solo with a pocket too.The batched pass can also keep what a backward reads: each read-out's input streams, normed rows, v, up rows and
inject gates, the mixer's and the MoE's branches, the MoE's input and slots, attention's projections and gated rows,
the final mixer's read-out. The up rows come from
fn_hc_upmix_up, our read-out kernel writing the rows its mixreads (the same mixed bits as
fn_hc_upmix). A captured comparison set's_qmm_upmixkeeps none, so--sliderefuses a captured set.
Flash Next on CUDA: the qwen4_exp family on the native CUDA host, every stream through the lane rounds #869's config check refused a taught folder's wider shared expert, which the first restart in the receipt below
found:
cuda_config.zignow takes the width with the pocket's neurons (tensorfold_learned_neurons, a whole numberof groups, 128 at most), with a test.
cuda_train.zigruns the learner's passes (keys, loss, target) with sequences side byside on the engine's own batched pass:
cuda_batch_train.zigruns the row-wise launches over every row, and then-gram block, DeltaNet and attention one sequence at a time from position 0 and zeros. The open fact's neurons are
scaled per sequence, and each nudge joins its layer's MoE output. The loss's gradient then runs from the 4-bit head
down to the lowest nudged layer with
flashnext.fn_backandcuda.gdn_train, reading the engine's packed expertswhere they sit and MLX copies of the small linears (
cuda_train_buf.zig). Both files are under 600 lines (the old676-line trainer is split in two).
zig build flashnext-learnchecks it on the GPU (receipt below).--slidefor qwen4_exp on CUDA (0986c5a).living/families.ziglistsqwen4_exp: a SwiGLU shared expert atevery layer behind its own gate, 4-bit groups of 32, keys under
text_config, the MTP layer's expert in the index.This replaces Living Weights as a core module: the learner, pocket and fits in zig/src/living, a family described as data (Nemotron only) #878's stand-in test family, whose tests now read the real entry.
flashnext/living.ziggives thelayout and writes the pocket through the core's shard writer (
pocket_shards_test.zig: the first save widens everyexpert, the MTP head's with zeros, and tells
text_config; a later save lands in place; every other byte kept).cuda_slide.zighosts the core learner over the trainer, and the lane backend opens with--slide: the trainer's1,024 rows in the pass's scratch, the layout from config.json, and the trainer's bytes kept out of the streams' room.
The native host's existing rules still apply (
--slideon a GB10 only, no--segments, a folder of links refused).Macs keep refusing
--slidefor this model.this receipt as measured), and a CHANGELOG line.
Receipt
Environment:
lw-d-fnlearn47748f7 (this PR's 4 commits on the merge 46627a8 oflw-b-core75978a1 intolw-c-fnbacka5d183d, itself on Flash Next on CUDA: the qwen4_exp family on the native CUDA host, every stream through the lane rounds #869's 7746bb9), Zig 0.17.0; NVIDIA GB10 (DGX Spark, sm_121), driver 580.178.04,CUDA 13.0.88, Ubuntu 24.04.5; checkpoint
leimroth-lab/Unbound-LR7revision 2ae3a2ff (Qwen3.8 Flash Next, MLX affine4-bit group 32, MTP head). "Probes" below are 23 greedy chat requests of up to 120 tokens: 12 questions about the
user (two or three wordings for each fact taught below) and 11 unrelated prompts (capitals, arithmetic, a haiku, a
Python one-liner and so on).
An untaught folder serves as before. Flash Next on CUDA: the qwen4_exp family on the native CUDA host, every stream through the lane rounds #869's head and this branch's head on the same snapshot: the 23 probes give
the same replies, drafted (23/23) and with
"draft": false(23/23). Decode speed (tools/bench_openai.py, 256tokens, median of 3): drafted 83.9 and 62.3 tok/s on code and chat against Flash Next on CUDA: the qwen4_exp family on the native CUDA host, every stream through the lane rounds #869's 84.1 and 62.6;
--no-drafts35.1and 35.1 against 35.0 and 34.9. A cold 3,001- and 12,001-token prompt (
tools/bench_cache.py --require-cold, time to the first token): 2.66 and 10.27 sagainst Flash Next on CUDA: the qwen4_exp family on the native CUDA host, every stream through the lane rounds #869's 2.71 and 10.28 s. The model takes 75.82 GiB against Flash Next on CUDA: the qwen4_exp family on the native CUDA host, every stream through the lane rounds #869's 75.81.
The trainer's forward is the serving forward.
zig build flashnext-learnon the snapshot (one prompt of thePython engine's reference set, its first 256 tokens, the last 32 the answer):
and 254 rows);
Engine.forward, every row'sscores): every layer's MoE input and shared gate and every row's scores byte for byte, untaught and with a pocket
neuron at layer 31;
The gradient against finite differences. In bf16 a finite difference is noisy: below about a bf16 ulp of the
residual the rounding decides, above it the loss's curvature does. The check nudges each layer's MoE output at the
last row (its nudge reaches only its own loss: the direct path) and at a row before the answer (it reaches the loss
only through later rows' attention and DeltaNet: the cross-row paths), at layers 12, 23, 31 and 42, along the push and
half along a random direction, in four windows of the prompt (offsets 0 to 3), with least-squares slopes over
steps of plus and minus one to six times 0.4% and 3% of the residual's size (32 checks a path and a step). Over all
four windows the direct path's median relative error is 10.6% at the small steps and 111% at the large ones; the
cross-row paths' is 43% at the small steps and 16% at the large ones (worst 204% and 69%). The first window's last
row disagrees at every layer (2% to 73% at the small steps, about 2x at the large ones); in the other three windows
the direct path agrees with a median of 6.2% at the small steps (worst 25%, 71% in one check of 8, and 17%), and
the last row of the same prompt's first 255 tokens agreed within 0.3% to 11% at layers 12, 23, 42 and 46. So the
direct path holds where a finite difference can see it, and the cross-row paths stay noisy, as they were with the
captured kernels before Flash Next on CUDA: the qwen4_exp family on the native CUDA host, every stream through the lane rounds #869's own. The kernels themselves are each checked against
float64 in the previous PR.
Lessons. A fresh copy of the snapshot (the files a lesson rewrites copied, the rest hard-linked), served with
--slide --context 8192 --no-thinking --temperature 0, taught these five facts one learn request each ("source": "chat"), with the probes after each lesson:None of the 12 probe wordings was among the 85 questions the lessons trained on, so every probe is held out: 4 of 12
after the five lessons. The first fact came back on all three wordings after its lesson and on one after the
second lesson, none after the third. After the second lesson that is Living Weights: a later lesson's quiet refit can silence an earlier fact on its own questions, and nothing asks it again #870 (a refit silences the earlier fact). After
the third it is the corrections rule: the server took "My GitHub name is ZackSample." as replacing "My name is Zack."
and retired the name fact's neurons (the graph says
replaced by fact 3), because both answer a "what is my ... name"question. This PR uses Living Weights as a core module: the learner, pocket and fits in zig/src/living, a family described as data (Nemotron only) #878's learner as it is; with Living Weights: a quiet refit holds an earlier fact's neuron on its own questions as read now #876/Living Weights: every earlier fact asked again after a lesson, and a take-out that puts the pocket back bit for bit #877 applied the name fact survives the second lesson (see
the comment below), and the false replacement is a separate lesson-writer fix. The 11
unrelated prompts kept the untaught model's replies byte for byte. A kept lesson took 8.7 to 10.5 minutes: 2.5 to
4.1 writing the lesson, 2.9 to 4.1 the trainer's first round (one step), then 0.6 to 2.1 for the second round and
about 25 s of checks a round.
A taught folder serves. The copy after the five lessons (27 neurons in its pockets) restarted without
--slideand with it: the 23 probes give the same replies as the teaching server did after the last lesson (23/23 each).
Drafted equals
"draft": false(23/23), and four streams at once (--parallel 4) equal each alone (23/23). A1,670-token prompt asked again resumes 1,663 tokens from the prompt cache and gives the fresh reply, drafted and
with
"draft": false. Decode:drafted 79.8 and 59.1 tok/s on code and chat (5% below untaught),
--no-drafts34.7 and 34.5 (1% below). Coldprompts: 2.67 and 10.30 s for 3,001 and 12,001 tokens.
Memory. The server reports the model at 75.82 GiB and 10.59 GiB kept for the trainer (its buffers take
10.23 GiB: MLX copies of the small linears it reads, the keep lists for 1,024 rows, the backward's scratch);
flashnext-learnholds 84.24 GiB of device memory with the engine and the trainer, 84.59 GiB at its peak. During alesson the GB10's 121 GiB showed 92 GiB in use.
Tests:
zig build test-cuda-hostfrom a fresh cache: 33/33 steps, 546/546 tests pass;zig build test-living: 39/39(Living Weights as a core module: the learner, pocket and fits in zig/src/living, a family described as data (Nemotron only) #878: 38, the one more is Flash Next's shard writer test). Every commit after the merge builds (
zig build,native,test-cuda-host,test-living) and passeszig fmt --check build.zig zig;python3 tools/zig/lean_check.pygives527 problems against 527 on
main(0 new); all 46 fatbins build for sm_86, sm_120 and sm_121 andcheck_fatbinspasses. On the head build,
tf-cuda-testflashnext-hc(its read-out kernel now also has the form that keeps the uprows),
flashnext-attn,route-rows,norm-rows,experts-rows,fn-backandfn-pocketall pass.Not run: Metal (Macs refuse
--slidefor this model as before; whether a Mac serves a folder taught here is notchecked), sm_86 and sm_120 cards.
Checklist
"draft": false, a resumed prompt equals a fresh one, and concurrent equals solo.no bits change. Taught: a layer with a pocket rounds its MoE output once after the pocket is added in fp32.)
Not in this PR, said plainly
facts. This PR runs Living Weights as a core module: the learner, pocket and fits in zig/src/living, a family described as data (Nemotron only) #878's learner as it is, and the receipt shows what that gives.
without the learned facts; verify windows run the main model's pockets, so drafted replies still equal undrafted ones.
flashnext-prompt's row-by-row reference path (Zig CUDA: one drafts-off Flash Next prompt matches #476's check tool) does not add a pocket; only the batched pass serves.Questions
Trainer.scores(a loss pass that stops at the scores) exists forflashnext-learn's byte check against theserving forward. Is a check-only method on the trainer acceptable, or should the check live behind a build option?
qwen4_explive inliving/families.zig(as here), or in the family folder (flashnext/living.zigalreadyholds its layout and writer), as asked on Living Weights as a core module: the learner, pocket and fits in zig/src/living, a family described as data (Nemotron only) #878?