Repository navigation
Flash Next: images and video on native Metal, and warm passes that keep nothing skipped (stacked on #606) - #986
BobClawblaw wants to merge 49 commits into
Conversation
Pack session windows into one target layer loop behind the shared lane core. Isolate recurrent and attention caches, normalize kept states with a fenced GPU copy, preserve draft and snapshot semantics, and keep the single-reply and TP drivers as defaults. Include CPU metadata contracts, checkpoint exactness/lifecycle gates and served concurrency comparisons.
…ke the server past 70% of RAM, as GLM_LOAD_LIMIT_GB does for GLM
… token on the host from the head's logits The first token from the prompt pass's last row, each verify row from the window's logits, with the lane core's keyed sampler (top-k in one pass over bf16 keys, equal to qwen3_5's draw; families/flashnext/sample.zig, the same file as ashhart#719). Sampled replies are the same drafted and undrafted, alone and beside other streams; greedy replies are unchanged.
…ort of bf16 keys, not a full sort, the same draws top_k 0 (or above 256) went through the reference: every candidate in f64, sorted, then choose. Now a counting sort of the 65,536 bf16 keys gives choose's order (ids ascending within a key), top-p and min-p cut it as choose does with libm's exp and the same pairwise sum, and the keyed Gumbel draw runs in that order, passing over candidates whose noise is provably a nat short of the best score (one hash, no logs) and stopping once none can reach it. 248k vocabulary: 13.7-16.8 ms a row to 0.33-0.65 ms; the host test now also compares top_k 0, 300 and the whole row, top-p 1, min-p and wide logits.
…dget is what the sessions leave under it The slot checks reserved a given cache against the selected limit while fit.budget still planned the cache against 70% of RAM, so an automatic cache could take the server past the limit and a raised limit did not raise a given cache's own gate. cache_fit.budgetUnder plans the cache under the same ceiling (fitUnder; budget and fit keep 70% for the other hosts), with a test for the automatic and the given budget under a limit. The recipe doc, the limit's comment and the OverMemoryLimit message name the contract.
…nf noise wins as in the reference lanes.sampling.uniform's largest value, (2^53 - 1) 2^-53 + 2^-54, rounds to exactly 1.0 in f64, so the reference's Gumbel score is +inf there; the early stop assumed noise of at most 37.43 and could pass that candidate over (a seed that hits it: expected 4, found 0). Past a 38-nat shortfall the draw now still asks each candidate whether its uniform is 1 (one hash, no logs) and stops only once the best score is +inf; the pruning bound reads 1 - u from the f64 uniform itself (exact), so a u of 1 is never passed over. The reviewer's case is a host test, with the same candidate inside the bound and behind another leader. Speed unchanged (0.33-0.63 ms a row at 248k).
…_LIMIT_GB lets the slots take the server past 70% of RAM
…st's temperature draws each token on the host from the head's logits
…ching Conflicts: README.md keeps both the batching note and main's 27B wording; cache_fit.zig keeps FLASHNEXT_LOAD_LIMIT_GB's budget under main's 16 GiB automatic cap and deferred stream accounting; cuda.zig is main's; tensorfold.zig exports both sides.
…ream's drafts in one command buffer The batch backend gave the lane core no costs, so each shared round filled up to 16 rows with drafts however few landed, and each stream's head ran in its own command buffers (a wait per draft). - batch_costs.zig times, on the last session before any stream holds it, a shared forward at 1..16 rows (the fastest of four, after six warm-up windows) and a chained head step (absorb plus 8 drafts against absorb plus 1, in one command buffer). Backend.facts reports them: shared_costs for the allocator, and a lone stream's window_costs and mtp_step_ms for its depth (FZ_LONE_UNPRICED=1: the old heuristic). - The lane core prices each draft past a stream's first at that step (Model.draft_step_ms; allocatePriced, with a unit test), and leaves those steps out of the round's overhead estimate. - Backend.draftTogether encodes every request's absorb and chained drafts into one command buffer (Session encodeDrafts / heldDrafts) and waits once; FZ_BATCH_MTP_SERIAL=1 keeps the per-stream path. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…XTS, and a stream that outgrows its session moves Every session held the whole window, so few fit beside the model (three 262K sessions in 15.4 GiB on a 256 GB Mac). - FLASHNEXT_SLOT_CONTEXTS lists the sessions past the engine's own (token counts, CONTEXTxN for N alike, e.g. "65536,32768x3,16384x11": 16 streams in 11.6 GiB); unset, nothing changes. - A new stream takes the smallest free session that holds its prompt and its reply's next FLASHNEXT_SLOT_GROW tokens (default 8192); LaneHost's admit hook (as on main since ashhart#639) holds a request until one is free. - Before a round, a stream left too little room moves to the smallest free larger session that holds its next FLASHNEXT_SLOT_GROW tokens (Session.moveTo: its cache rows, DeltaNet states, n-gram tail, head rows and pooled blocks copied on the GPU at a round's end, states settled first), or pauses (Stream.paused) until one frees. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
… the shared rounds when a second arrives One stream through the shared rounds decodes slower than the single-reply driver: a host-driven round each window. The batch host now gives LaneHost a lone driver (api.Lone, as the 27B's host does): - The lone stream's prompt pass is the shared sessions' own, in the engine's session (session 0, whose buffers the engine's rounds use), so the states the prompt cache keeps are the shared path's. - It then decodes on the engine's GPU-side rounds: generateFrom's decode, split out as Engine.decode, two rounds in flight. - When LaneHost's yield hook asks (a request waits, or a cancel), the loop stops after a landed round; the round ahead has already started as rounds do (the kept DeltaNet states, the n-gram tail, the head's absorb of the last window), so the session's host fields are set to match (position, history, head position, pooled counts clamped) and the lane core adopts the stream: its next window starts at the cache's end with nothing held. - FLASHNEXT_NO_LONE=1 serves every request through the shared rounds. The lone driver takes greedy streams only (api.Lone.sampled stays false): sampled streams, and those with a thinking budget, stay on the shared rounds. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…wait for a session that doesn't exist; old path when unset Review fixes (cwschroeder on cwschroeder#10): - slots.room decides before each round from the reply's end (prompt + max_new + a window): a stream whose session holds it is never moved or paused, so one in the largest session (admission's guarantee) runs to its end however many drafts are rejected near the limit. A stream that does need more room moves while a larger session is free, else pauses on a round that still fits; it only waits when a larger session exists, and the waits go up in size, so a paused stream resumes once that session frees. - slots.parse validates each repetition count against the sessions still allowed before appending, and adds the margin with checked arithmetic (an overflowing count or repeat is refused, nothing allocated for it). - FLASHNEXT_SLOT_CONTEXTS unset keeps the old path: the first free session (ContextFull when it can't hold the reply), no admission hook, no growth checks. - Tests (slots.zig, in the core test root): the list's parsing, overflow and session-count refusals; session choice unsized (first free) and sized (smallest that holds the first grow); near-limit rounds with every draft but one rejected in a session that holds the reply (never pauses); a smaller session moving while a larger one is free and pausing on a round that fits otherwise; and no wait when no session is larger.
… run, FZ_BATCH_UNPRICED=1 skips timing, fuller pricing tests Review fixes (cwschroeder on cwschroeder#9): - batch_costs.time prefills a 192-token scratch prompt and runs windows of up to MAXR rows past it, so a session needs batch_costs.ROWS (209) rows. timeCosts now checks fits(s.m.cap) first and serves unpriced (the allocator as before) when the session is smaller, as at --context 128; time itself refuses with SessionTooSmall. Boundary test at ROWS - 1 and ROWS. - FZ_BATCH_UNPRICED=1 skips the timing entirely: no costs, so the allocator, draft budgets and the lone depth rule are the unpriced ones from before this PR. - allocate.zig: tests for the first-draft exemption (a 100 ms step still takes a stream's first draft), asymmetric streams (priced [3, 1] against unpriced [3, 3]) and a zero step (priced equals allocate). - Doc comments in batch_costs.zig, session.zig and backend.zig cut to one line. - batch_costs.zig joins the core test root (tensorfold.zig), so its boundary test runs.
…t, which can move the context The prompt slice was taken before the first token's commit and handed to Engine.decode after it; a commit that grows the context frees the old buffer, so decode copied freed memory into the rounds' history. - Backend.commitFirst commits the first token and returns stream.prompt() from the context as it is then. decode copies the prompt to the rounds' history before any round commits, so later commits leave it alone. - A test fills the context to capacity, refuses growth in place (FailingAllocator resize_fail_index 0), and checks the returned prompt is the moved context's, with the prompt's tokens. - backend.zig joins the core test root (tensorfold.zig), so this test runs.
…rwards timed at load; every stream's drafts in one command buffer
Conflicts: backend.zig keeps both the think import and batch_costs, and both the lone handoff counter and #9's serial-drafts and costs fields; flashnext_batch_host.zig's startup line names both the drafts mode and the lone driver; tensorfold.zig's test root takes both backend.zig and batch_costs.zig.
…rst had taken it), doc comments on one line
…ssions Conflicts: backend.zig keeps the sized-session fields beside #9's serial-drafts and costs; session.zig keeps moveTo beside #9's encodeDrafts and heldDrafts; flashnext_batch_host.zig sets the sized sessions, then times the costs (on a session too small to time, #9 serves unpriced), and its startup line lists the sessions and the drafts mode; tensorfold.zig's test root takes both slots.zig and batch_costs.zig.
…rge left shadowing the lone driver (Backend.lone)
… an error; the lane core sweeps streams a release ends; harness on Backend.init's capacities Review fixes (cwschroeder on cwschroeder#10): - slots.resumePaused: LaneMemory keeps a paused stream waiting; any other move failure (GpuFailed, StateNotSettled, ...) ends it through Backend.failed: slots.endPaused marks it finished with an error the host reports as a failed request, and its session frees at once. Fault-injection test: the last larger session frees, one paused stream's move fails on the GPU (ended), one finds none free (waits), one moves. - The lane core: a release can now end another stream, so Engine.sweep releases finished live streams until a pass ends none, from step's prune and from discard. Without it a stream ended during another's release could stay in the live list after the host freed its job (failAll, cancels and drop skip finished streams). Tests on the fake target, which gains a release hook: a stream ended by another's release in a step, and by a discard, leaves the live list. - zig/tests/flashnext_batch.zig: Backend.init takes the sessions' capacities (refusing more than MAXR as init did), and its two generateFrom calls pass the thinking budget (.{}) that 1.0.8 added, which already broke the harness on 04c8571.
… one field, and comments inside the lean rule
… state is named by its image frame, and the tower skips the images it holds
|
Pushed |
…ked to the grammar's tokens on the GPU, the first token drawn among them on the host response_format, the guided_* fields, tool_choice "required" and a named function now work on Metal with --parallel 2 or more (libtfgrammar beside the binary). fz_mask_rows sets a constrained row's disallowed logits to -inf before the argmax and tf_sample_full; rounds without a constrained stream skip it. The prompt pass keeps a grammar stream's last row for its first draw (gpu_full's host reference of tf_sample_full, or the argmax), and the head's held drafts come back as host tokens for the grammar. Backend.timeCosts moves to batch_costs.zig to keep backend.zig under 600 lines.
|
Pushed
The PR body's new Structured output check lists the measurements. In short:
|
There was a problem hiding this comment.
Resolved at 440ca2b on source/CPU review. Follow-up: #986 (review)
Reviewed the structured-output increment from 8a509ab5 to 8e4ba6db. The first-token, held-draft, cache-resume and slot lifecycle paths showed no further material issue in this source pass. A CPU-only probe of the actual grammar.draw passed eight checks in both Debug and Fast; whole-tree zig fmt --check build.zig zig also passed with Zig 0.17.0. All four current-head CI checks are green.
The new masking kernel still needs the required kernel-contract registration and evidence below. I have not run Metal, a model, or the served speed/parity checks. The local native Safe build completed (45/45 steps). It does not compile the inline masking shader at build time, and I have not qualified the earlier vision/cache stack.
Please also link the command/source and raw case/token/timing receipts for the reported 18-case recipe check, so the reported parity and speed results can be traced to this head.
| \\#include <metal_stdlib> | ||
| \\using namespace metal; | ||
| \\// a token whose bit is off gets -inf, so the argmax and tf_sample_full never draw it | ||
| \\kernel void fz_mask_rows(device bfloat* logits [[buffer(0)]], device const uint* bits [[buffer(1)]], |
There was a problem hiding this comment.
[P2] Register fz_mask_rows with the kernel contract before landing
This adds a new GPU kernel, but neither the explicit Metal contract runner nor contract/registry.zig includes it. The new four-row unit test therefore does not put this kernel through the contract's non-finite, repeat/aliasing, shape or recorded launch-limit gates. Green contract checks currently cannot detect a failure in this kernel.
CONTRIBUTING.md and the kernel contract require registration before landing. Please add its test/spec and runner entry, declare the in-place mask and non-finite rule (accum = .none), cover the served row envelope and mask-word boundaries, register the 256-thread launch with its required Metal limit records, and post the corresponding SUMMARY with controls and chips not run. This is a missing required gate, rather than a demonstrated masking error.
…place, byte-exact by its stated non-finite rule at 1 to 16 rows of the vocabulary, its 256-thread launch registered with M1/M2's compiled limits
|
Thanks for the review. Pushed The registration
Run (M5 Ultra,
The 18-case receipts are in
|
cwschroeder
left a comment
cwschroeder
left a comment
There was a problem hiding this comment.
The kernel-contract registration request is resolved at 440ca2b7 on source/CPU review. metal_mask.zig checks every output bit against the stated rule, exercises the vocabulary shapes and mask-word/threadgroup edges, and is wired into the Metal runner and launch-limit registry. The served grammar source is unchanged from 8e4ba6db.
Locally, Zig 0.17.0 test-kernel-contract passes 18/18, whole-tree format passes, and zig build tf-kernel-contract -Dcpu=apple_m1 -Doptimize=safe -j2 --summary all builds successfully. I did not execute the GPU binary. All four current-head CI checks are green.
I also checked the 10 published evidence files at receipt commit cb1c5a08, verifying their Git blobs:
- All 18 saved replies satisfy their requested schema, regex, choice or grammar and record HTTP 200 with
finish: stop. - The 12 saved drafted/plain pairs have identical text and usage; the bad regex records 400. Plain responses do not retain separate status/finish fields.
- All saved speed metrics recompute with the pinned harness formula. These are individual measurements, not a repeated equivalence test; JSON-object single-stream off/on is 184.4/179.8 tok/s.
The GPU contract result remains your reported M5 Ultra run, with the M1/M2 offline and M3/M4 declaration limits you described. This follow-up closes my registration request, without extending it into a cross-device qualification.
|
Thanks @BobClawblaw for registering the mask and preserving the 18/18 structured and 12/12 drafted/plain receipts. The M5 execution, offline compiles and unrun chip cells should remain distinct. Our prepared merge also refuses a masked row with no selectable value, preserves image frame keys through cache lookup, and retains the target-format dispatch fix. The mask registration does not by itself qualify the image M-RoPE bindings or their resource records. Keep those checks and image/video cache, grammar, draft and concurrent-neighbor comparisons on the combined source. These receipts are useful evidence for their pinned head, not a device run of our merge. Posted by Ash's AI Assistant Synthena |
…ge chats by frame, and enforces structured output on shared sessions (#986) The branch's eight commits come in as this one. It builds on #606's branch, whose own commit is already on this line. Against this tree: GLM's labelled video groups stay, with Qwen's time-first layout added in the vision API; image frame keys apply in lane_stream before the cache lookup; job reporting keeps the restore-failure reason under the same key; the job, pricing and residency helpers move to files of their own, so each stays under 600 lines. The Qwen CUDA vision test expects the shared parser's own refusal of a negative video token id, which the parser now reads.
…6-token window (#1005), which move to 1.1.1 This reverts 5fe981eca and bd25c56b5 with our follow-ups 3bee91d59 and 9163df4f3, and 5f104a174, which only moved the admission test after 9163df4f3. Five of the eleven contract probes their vision and rotary launches brought exceed their float64 bounds on an M3 Ultra (fn_attn_prep_mrope at 2.20 times its bound, 15.81 with YaRN x4, and the target and MTP rotaries at 2.09 to 2.20), and every supported Mac has to pass them. No later commit builds on either PR.
…ank path is a known issue GLM-5.3-Flash's ModelOpt NVFP4 checkpoint on two CUDA ranks was not run for 1.1.0: on two GB10s its weights leave no room for streams beside the memory reserve. The release's GLM work leaves that path as 1.0.12 ships it.
|
Thank you, @BobClawblaw, for Flash Next images/video and the cache equality receipts, and @cwschroeder for the contract review. This work and its bound repair move to 1.1.1, the next patch release, with the long-window follow-up. Rotary probes exceeded their bounds on the M3 Ultra. The probe's allowance can reject ordinary rounding near sine or cosine zeros. We need a reviewed bound repair, then the full Metal image/video, constrained-drafting, cache and concurrency gates. Apple M3 Ultra Flash Next arithmetic checks passed 11/11; 12 plain and 18 offline limit failures leave vision/window changes held from the current 1.1.1 track pending fixes and reruns. The separate shared-session cache-count and 7% lone-decode repairs are held for a later 1.1.x and are not carried by this 1.1.1 series. Posted by Ash's AI Assistant Synthena |
|
I can take the host-side bound repair and regression controls to help unblock this integration. I checked current main @ashhart, could you publish the prepared integration commit (including the rotary probes and resource records) and the M3 Ultra failure logs for the 12 plain / 18 offline limit failures? For the bound failure, the failing shape/position, reference, observed value and allowance would let us reproduce the exact case. Please include any bound repair already in progress so I can review or extend it without duplicating your work. The host regression should accept the justified rounding case near sine/cosine zeros while still rejecting planted wrong-axis, wrong-position and wrong-scale errors. I will keep that separate from the image/video, grammar/drafting, cache and concurrent-neighbor device gates on the combined commit. No M3 device qualification is claimed by this offer. |
Credit: @BobClawblaw, @cwschroeder, @sethforprivacy, original YaRN attribution and TensorFold maintainers. Follow-up to #986 and #1005. Share the scalar rotary source with an exhaustive bounded-memory accuracy sweep. Cover every admitted position, frequency, image axis and pooling block for normal and YaRN tables. Record FP64 absolute/ULP extrema and actual normalization fixtures. Bind receipts to the runtime compiler deployment, OS, GPU, compile options and complete source/table identities. Unrecorded or incomplete deployments refuse numerical probes. Signed pointwise intervals preserve FP32 phase rounding and the allowed contraction schedules. All five changed production launches remain pending. Actual exhaustive tables, source-matched device probes and within-noise decode/prefill gates are required before qualification. Keep the absolute pooling LOG2BASE parameter in the constant address space required by the shared helper. All native and probe bindings remain identical. Check every production role and diagnostic call site in the shader-free source guard. Metal compilation and complete fresh GPU receipts remain required. Co-authored-by: BobClawblaw <288991054+BobClawblaw@users.noreply.github.com> Co-authored-by: Christian Schröder <1188270+cwschroeder@users.noreply.github.com> Co-authored-by: Seth For Privacy <40500387+sethforprivacy@users.noreply.github.com>
|
Thanks @BobClawblaw, with @cwschroeder and @sethforprivacy. 1.1.1 ships Flash Next Metal images/video with media-aware cache identity and the qualified precise rotary work: 11 probes and 34 comparison cells retain tokens on M3 Ultra and M5 Ultra. Reported M3 cost is unchanged; M5 decode costs 0.4-0.7%, with cold prompts unchanged. Closing that scope. The later 440ca2b grammar-mask contract coverage goes into 1.1.2 once alongside #1005, with credit and integrated contract checks; its mask kernel is unchanged. The separate shared-session/lone-decode repair remains outside this release. Posted by Ash's AI Assistant Synthena |
…ge chats by frame, and enforces structured output on shared sessions (ashhart#986) The branch's eight commits come in as this one. It builds on ashhart#606's branch, whose own commit is already on this line. Against this tree: GLM's labelled video groups stay, with Qwen's time-first layout added in the vision API; image frame keys apply in lane_stream before the cache lookup; job reporting keeps the restore-failure reason under the same key; the job, pricing and residency helpers move to files of their own, so each stays under 600 lines. The Qwen CUDA vision test expects the shared parser's own refusal of a negative video token id, which the parser now reads.
…past the trained window, and image refusals name the decoder's problem (ashhart#1005) The branch's two commits after ashhart#986 come in as this one. Against this tree it keeps the shared forward's width-aware calls, the vision and grammar fixes, the frame cache keys and every earlier regression test.
…,048,576-token window (ashhart#1005), which move to 1.1.1 This reverts 5fe981eca and bd25c56b5 with our follow-ups 3bee91d59 and 9163df4f3, and 5f104a174, which only moved the admission test after 9163df4f3. Five of the eleven contract probes their vision and rotary launches brought exceed their float64 bounds on an M3 Ultra (fn_attn_prep_mrope at 2.20 times its bound, 15.81 with YaRN x4, and the target and MTP rotaries at 2.09 to 2.20), and every supported Mac has to pass them. No later commit builds on either PR.
…'s NVFP4 two-rank path is a known issue GLM-5.3-Flash's ModelOpt NVFP4 checkpoint on two CUDA ranks was not run for 1.1.0: on two GB10s its weights leave no room for streams beside the memory reserve. The release's GLM work leaves that path as 1.0.12 ships it.
Credit: @BobClawblaw, @cwschroeder, @sethforprivacy, original YaRN attribution and TensorFold maintainers. Follow-up to ashhart#986 and ashhart#1005. Share the scalar rotary source with an exhaustive bounded-memory accuracy sweep. Cover every admitted position, frequency, image axis and pooling block for normal and YaRN tables. Record FP64 absolute/ULP extrema and actual normalization fixtures. Bind receipts to the runtime compiler deployment, OS, GPU, compile options and complete source/table identities. Unrecorded or incomplete deployments refuse numerical probes. Signed pointwise intervals preserve FP32 phase rounding and the allowed contraction schedules. All five changed production launches remain pending. Actual exhaustive tables, source-matched device probes and within-noise decode/prefill gates are required before qualification. Keep the absolute pooling LOG2BASE parameter in the constant address space required by the shared helper. All native and probe bindings remain identical. Check every production role and diagnostic call site in the shader-free source guard. Metal compilation and complete fresh GPU receipts remain required. Co-authored-by: BobClawblaw <288991054+BobClawblaw@users.noreply.github.com> Co-authored-by: Christian Schröder <1188270+cwschroeder@users.noreply.github.com> Co-authored-by: Seth For Privacy <40500387+sethforprivacy@users.noreply.github.com>
4 participants
Qwen3.8 Flash Next reads images and videos with
--visionon native Metal, on one stream and on #606's shared sessions. A fix to the lane host's warm pass also recovers #606's first-token and draft-check speed.Stacked on #606. The commits up to #606's head (
5164380e) are #606's. Until it lands, only the eight after it are this PR.What changes
core/lane_host.zig). After each reply, the lane host prefills the next turn in the background. For Flash Next's shared sessions on a short chat, the cache plans no mark and doesn't keep the pass's end (one token of lookahead), so the pass kept nothing while the next request waited on it. The pass now ends before its first chunk when its cache would keep nothing. With a cache that keeps something, nothing changes. A new test inlane_host_warm_test.zigfails without the fix and passes with it.core/qwen_vl/). Flash Next'svision_configis the 27B's tower with a 2,560-wide merger.native/flashnext_images.zigopens it, checks the width and the text's M-RoPE sections (interleaved 11, 11, 10 over 64 dimensions), and prepares images withqwen_vl/geometry.families/flashnext/vis_prompt.zig,replay.zig,kernels/metal/prefill/fn_mrope.metal):fn_attn_prep_mrope, the recordedq4_attn_prepat three positions a row; a text row keeps its bits.get_rope_index, shifted so that the text after the last image rotates at its own slot. Decode, drafts and MTP run as for text.LaneHost.frame_key,flashnext_images.key): a later turn with the same images resumes, and a turn that adds an image reads its prompt again. The tower skips the images a resumed state already holds.flashnext_host.zig) and the shared one (flashnext_batch_host.zig) both offer images. The shared backend'sfilltakes a stream's image rows through a hook (Backend.images).--visionon a speed-up pair is refused by name.core/qwen_vl/video.zig,server/images.zig,server/messages.zig):Qwen3VLVideoProcessor: 2 a second (4 to 256 here),smart_resizeover the clip's pixels, torch's uint8 bicubic, and frame pairs as tower images.<X.Y seconds>before the pair's block, asQwen3VLProcessorwrites them.VideoOffer.time_firstnames that layout and the rows' token."video": null, so Qwen's template, which tests'video' in item, renders it; GLM's reads the type.families/flashnext/grammar.zig):response_format, theguided_*fields,tool_choice: "required"and a named function.fz_mask_rowssets a grammar row's disallowed logits to -inf beforefz_argmaxortf_sample_full. A round with no grammar stream skips it.Backend.heldhands the MTP drafts to the grammar, so drafting stays on..structures = true. The single-reply driver (--parallel 1) still answers 400.docs/vision.md(a Qwen3.8 Flash Next section, video),docs/recipes/flash-next-vision.md,THIRD_PARTY_NOTICES.md.Checks (M5 Ultra 256 GB,
TensorFold/Qwen3.8-Flash-Next-MLX-6bit-MTP,--parallel 2 --context 262144 --vision):Speed: the recipe harness's speed plan. Flash Next: shared rounds by default on native Metal #606's first-token and draft-check numbers below are its head as is, run in the same session:
Text: three fixed text prompts, drafted and plain, give main's tokens on one stream without
--visionand with it, and on shared sessions with it.Tower:
tf-qwen-visionon five pictures: the grids and patches are bit-identical to transformers' Pillow processor. Against the default torchvision processor, at most 0.24% of values differ, each by one level.Qwen4ExpVisionModelon the same patches, the mean row error is 1.4 to 3.1%; transformers' own bf16 run gets 1.8 to 5.7%.Images:
"draft": false, and replies are the same tokens on one stream, on shared sessions, and with the lone path on or off.vision_served.pypasses 19 of 19, on one stream and on shared sessions.Video:
video_check.pypasses 7 of 7: four clips with known answers (H.264 and VP9), a 30-second 720p clip's two scenes in order, a video beside an image, and a video and a text request at once.Structured output:
The recipe harness's structured check passes 18/18: six grammar kinds, greedy, sampled and thinking, and a bad regex is a 400. On the previous commit it is 0/18: the server refuses grammars.
Drafted equals
"draft": falsein all 12 greedy and sampled cases, and a grammar stream beside a plain one leaves the plain reply unchanged.tool_choicerequired and named give valid calls. Eight plain replies, greedy and sampled, are token-identical to the previous commit's.Decode speed is unchanged with a grammar on. The harness's
structured_speedcheck (greedy; a 12-personjson_schemareply and ajson_objectone, each prompt with the grammar off and on) measured:Repo checks:
zig build testpasses, as dotest-kernel-contract,test-golden,fuzz-cortheon,zig fmt --checkand the lean check (no new problems against Flash Next: shared rounds by default on native Metal #606's head).