Repository navigation
Conversation
…nsor TORCH_CHECK reports a failed CUDA call but does not read the status out of the runtime, so the error stays latched in the calling thread. The errors raised here are non-sticky and the context remains usable, but the next unrelated CUDA call picks up the stale status and fails as if it were its own. That turns a handled exception into a process-wide fault. host_banks.pin() catches the RuntimeError from host_register() and re-raises a friendlier message; any torch call made while handling that failure would report the stale cudaHostRegister error instead of its own. Route every CUDA call in the file through an FT_CUDA_CHECK macro that drains the status with cudaGetLastError() before raising. The two cudaDeviceGetAttribute calls in host_ptr_identity() were previously unchecked and are now checked as well. Check messages are unchanged. Add a test that drives the failure path directly and asserts the context survives it.
…on-fatal Two corrections to the previous commit. cudaFreeHost was the one call left uncovered, and it is the worst place to leak a status: it runs as a from_blob deleter during GC, with no exception to attribute the failure to. It drains unconditionally and never throws. host_ptr_identity's two cudaDeviceGetAttribute calls previously discarded their status with uva/reg pre-initialised to 0 -- the same idiom as driver_cuda_version's "stays 0 when no driver is installed", i.e. a deliberate fallback rather than an oversight. Raising there turned an unqueryable attribute into a fatal error on the offload path, and _host_ptr_identity is lru_cached, which does not cache exceptions, so it would have re-raised on every call. Drain without raising and keep the 0 fallback: "no identity" routes device_ptr through host_device_ptr, the real translation, which is correct everywhere. Also soften the macro comment: these calls report whatever is latched on the thread, so "every error below is non-sticky" was a stronger claim than holds. Draining leaves the context usable for the failures these calls originate; a sticky error latched elsewhere re-latches on the next call.
The expert-bank loader, the CPU bandwidth bench and the loader benchmark all sized host-RAM work from /proc/meminfo MemAvailable alone. Inside a container started with --memory=N that figure is the host's, not the cgroup's: the parallel expert reader was admitted against tens of GiB of "available" RAM while the cgroup could hold a fraction of it, and the process was OOM-killed mid-load. benchbw.py carried its own cgroup-v2-only parser; the loader had none. Rule: effective available memory = min(host MemAvailable, remaining finite cgroup budget). The cgroup budget is the tightest finite limit-minus-usage along the process's cgroup ancestry (v2 memory.max / memory.current with the `max` sentinel; v1 memory.limit_in_bytes / memory.usage_in_bytes with the unlimited sentinel as a fallback), resolved from /proc/self/cgroup with bounded reads and a bounded ancestor walk. Absent procfs or no visible memory controller yields None (unknown), so non-Linux hosts keep their previous best-effort behaviour and are never told they have 0 bytes; a present but malformed or unreadable control file fails closed to 0. The probe runs before any bank or buffer allocation: admission is the only decision point, because once the parallel reader's whole-shard anonymous buffers exist there is nothing reclaimable left to give back. freetoken/memory.py is stdlib-only (no torch import) so it can be loaded by path inside a slim container for diagnosis. Every production MemAvailable reader now goes through it: moe/expert_banks ._host_ram_fits_parallel, moe/benchbw._available_ram_bytes (the private _cgroup_mem_headroom parser is removed so there is one cgroup parser; the bench's clamping and synthetic-bank sizing are otherwise unchanged), and benchmarks/bench_load_weight_generic.MemSampler. Tests: tests/test_memory.py drives the resolver against fake /proc and cgroup trees under tmp_path -- unlimited, finite, nested ancestor-min, namespaced root, usage above limit, malformed, oversized, permission denied, v1 fallback and sentinel, absent files, and the hybrid `1:net_cls:/` + `0::/...` membership layout systemd and Docker both produce. tests/moe/test_expert_bank_memory_admission.py checks that the auto-admission path drops to the serial build when the fake cgroup budget (299 B) cannot hold the banks plus one shard transient (300 B) although MemAvailable is large, that an explicit --expert-load override bypasses the probe, and that the bench receives a finite budget unchanged while keeping its legacy 8 GiB fallback only for an unknown (None) probe. Validated on CPU-only Linux (Ubuntu 24.04, Docker 29.1, cgroup v2; no GPU is involved in this path): 47 tests pass under python and python -O; the full suite has the same two pre-existing failures as origin/main (tests/models/test_muse_glimmer.py) and no new ones; python:3.12-slim with --memory=2g / --memory=6g / no limit reports effective = memory.max - memory.current (2136084480 / 6432010240) and host MemAvailable respectively. Mount-namespace hardening of the cgroup lookup is deliberately out of scope for this change.
…shML-org#690) * build(kernel): vendor ggml-cuda from llama.cpp 8c2134a8 behind a tvm-ffi binding * build(deps): vendor the gguf-py reader * feat(quant): add the GGUF quant kind and an embedding quant method * feat(gguf): serve Flash-Next and Gemma4 GGUF on the GGUF quant kind and drop the old Q4_0 path * fix(kernel): build ggml-cuda for plain sm_120 by turning off the native fp4 MMA * perf(gguf): write ggml results in the activation dtype and keep MoE intermediates in bf16 * build(kernel): compress the ggml fatbin to keep the multi-arch kernel cache small * fix(gguf): read F32, F16 and BF16 PLE tables without ggml * fix(moe): pad the offload slot cache for ggml MMQ tail reads * fix(gguf): support hip builds and runtime devices * fix(gguf): copy same-type dequantize and check for missing ggml converters * fix(gguf): load quantized qwen4exp ple_key and ple_value * feat(gguf): run float parts of a fused GGUF linear as bf16 segments * fix(gguf): load quantized qwen4exp hyper-connection and gdn ba weights * fix(moe): zero the prefill staging area after the last prefill layer --------- Co-authored-by: zihaomu <zihaomu@outlook.com>
…L-org#587) The qsa_sparse_paged_attention tiles are tuned for the GB300's 228KB shared memory; RDNA gfx11/gfx12 have ~64KB/CU, so k+v staging overflows LDS. Query the device's actual shared-memory limit and shrink the KV tile only when it doesn't fit, ROCm-gated so NVIDIA tiles stay byte-identical. The gguf ROCm-compile fixes this branch also carried are now covered upstream by the ggml-cuda vendoring in FlashML-org#690 (separate DEFAULT_HIP_CFLAGS + vendors/hip.h), so this branch reduces to the QSA clamp only.
…d types (FlashML-org#702) * feat(gguf): run experts whose gate_up and down mix float and quantized types * fix(kernel): map an fp32 compute type to tl.float32 in fused_moe_kernel_triton * fix(gguf): write float expert GEMM output in the activation dtype
…ock the first launch (FlashML-org#697)
* fix(kernel): preserve fp4 scale rounding at subnormal ties * fix(kernel): preserve native e4m3 scale calculation
…ools (FlashML-org#700) * feat(moe): serve experts of several formats through per-format slot pools * refactor(cpu-moe): pass each layer's cpu format with its options * fix(moe): read only each format's layers when loading experts of several formats --------- Co-authored-by: andy-yang-1 <andy_yang@berkeley.edu>
) * perf(kernel): rewrite the triton nvfp4 moe prefill gemm * perf(moe): tune nvfp4 prefill tiles per gpu class GeForce runs fp32-accumulate tensor MMA at half rate, so RTX parts want smaller BLOCK_SIZE_M than H100 and their own BLOCK_N / BLOCK_KB / stages. Tables come from tile sweeps on RTX PRO 6000, RTX 5090, RTX 4090 and RTX 3080 Ti; H100-class GPUs and AMD keep the previous tuning.
…ugh ggml-cpu (FlashML-org#704) * feat(cpu-moe): serve GGUF experts on the cpu and hybrid backends through ggml-cpu * feat(cpu-moe): run float GGUF expert rows on the cpu through ggml-cpu's vec_dot * feat(cpu-moe): build ggml-cpu for each llama.cpp x86 variant and run the best one the CPU supports
…count (FlashML-org#729) * perf(kernel): split triton decode attention by context length and sm count * perf(kernel): merge triton decode splits a tile at a time - stage 2 loads SPLIT_TILE splits per step and merges them with one online-softmax update instead of one split per iteration - SPLIT_TILE is the launch's split cap rounded up to a power of two, at most 8192 fp32 elements and 64 splits; num_warps scales with it
…rg#733) * perf(moe): speed up the nvfp4 moe prefill gemm on rtx gpus * perf(moe): keep the byte e2m1 decode at every tile size on h100 --------- Co-authored-by: andy-yang-1 <andy_yang@berkeley.edu>
…rg#734) * feat(moe): bench hybrid fetch fractions per moe layer kind The ft bench bw profile (version 6) holds only layer_kinds rows, keyed by CPU format, options, expert shape, top_k and CPU worker threads; older profiles are ignored. ft bench bw --model also takes checkpoint paths and benches each of their layer kinds. Hybrid decode reads each layer's own row for its fetch fraction, else keeps a fixed cap of 1. --moe-strategy auto no longer upgrades to hybrid from the profile. * fix(bench): fill gguf mxfp4 scales, scope act_block to ds_fp4 and skip cpu-less rows
…sm_89 (FlashML-org#735) The pointer path of the Triton prefill GEMM loads the e4m3 scales one byte per thread, too small for cp.async, so every K step waits on that load and a CTA barrier. Read them as int32 words (4 scales each) on GeForce sm_86 / sm_89 so they join the pipeline. - SCALE_WORDS constexpr in `_prefill_nvfp4_moe_kernel` (off by default) - Triton prefill tuning as one record per (GPU class, cc); GeForce sm_86 / sm_89 turn scale words on, with 8 warps when K >= 2048, and sm_86 moves to BLOCK_SIZE_M 64 from 100 tokens per expert
…ashML-org#741) The NVFP4 release quantizes the attention gate with a different weight global scale than q/k/v, so one fused weight cannot carry a single global for all four parts. BREAKING CHANGE: an FTW checkpoint converted with the fused qkvg_proj no longer loads; reconvert it with `ft checkpoint`.
… the replay dropped The 5080deploy merge of FlashML-org#676 updated these tests inside the merge commit (the on-demand path runs the prefill GEMM over the slot cache, and the layer hands ensure_experts unique ids). Re-merging FlashML-org#676's head on the new base does not carry edits made inside a merge, so they are restored here. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Bu6LgoxLR4wETqb7RPR2vt
Upstream FlashML-org#689 replaced OffloadMoeCache.bank_views with per-pool and staging views. The FlashML-org#676 on-demand prefill GEMM runs over the whole slot cache (n=cache_size) and still calls bank_views, so the first short prefill raised AttributeError. Restore it for a cache with one pool and one bank layout, where a slot id is its bank row; the pool guard already refuses on-demand prefill with several pools. tests/moe/test_small_prefill_on_demand.py covers the call. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Bu6LgoxLR4wETqb7RPR2vt
…lashML-org#689 FlashML-org#650's engine fake and the decode-forward fake predate upstream FlashML-org#689: the rebuild now splits the MoE target into pool sizes (decode_target, _moe_pools, _moe_pool_bytes, validate_rebuild/rebuild take the sizes) and decode reads pool_bank_views. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Bu6LgoxLR4wETqb7RPR2vt
… gathers Upstream FlashML-org#689 keeps one fused-copy row-byte descriptor per bank layout (_copy_feat_bytes[layout]). The prefetch verify and slot-audit gathers still passed the whole list, so every enabled check failed at launch (TypeError: expected DLTensor*, got ffi.Array). Pass the layer's layout entry. The 83 failing test_skeleton/test_prefetch audit and verify tests pass again (385 passed). Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Bu6LgoxLR4wETqb7RPR2vt
Assisted-by: Claude Opus 5.5 (1M context) (cherry picked from commit 99c0757)
…t the API layer # Conflicts: # python/freetoken/server/api_server.py # python/freetoken/server/generation.py # tests/server/test_openai_api.py # Conflicts: # python/freetoken/server/api_server.py # python/freetoken/server/generation.py # tests/server/test_openai_api.py
…compatible servers do Upstream FlashML-org#223 made resolve_sampling reject top_k=0 with HTTP 400 ("top_k must be -1 (disabled) or >= 1"). Clients send top_k=0 to mean "off": Open WebUI with Ollama-style parameters and some SDKs do, and vLLM and the other OpenAI-compatible servers accept any top_k <= 0 as "no top-k filter". A checkpoint whose generation_config.json carries top_k=0 made every request that omits top_k fail the same way under --sampling-defaults=model. resolve_sampling is the one place that validates top_k (the chat, legacy completions, Responses and Anthropic converters and FlashML-org#393's n>1 fan-out all go through it; the pydantic request models put no bound on it), so it now maps any top_k <= 0, from the request or from the checkpoint defaults, to the engine's -1. The sampler already read top_k < 1 as "off". The NaN/inf/negative temperature and top_p (0, 1] checks stay. FlashML-org#223's test that expected ("top_k", 0) -> 400 now asserts that 0 and -5 are accepted and reach the engine as -1; tests/server/test_sampling_validation.py covers the request and checkpoint cases and that the other range checks still raise. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01RUamGLmkVPkWmPPBEUqQLP
…s and stream entries, host-built logprob row index, responses and truncated reasoning_tokens, logprobs null)
…d stream/non-stream parser consistency
With stop strings the detokenizer now holds a logprob entry until its text is
released, so one reply can carry several entries or none, and a per-reply
text_offset can no longer describe them. The detokenizer stamps each entry
with its own token's start instead, and UserReply.text_offset is gone;
completions reads the offset from the entry, keeps the pending carry-forward
for text-less replies, and drops stop_trimmed entries on both paths. The
non-stream chat choice keeps the always-present logprobs key, now with the
segmenter's {"content", "refusal"} object.
Assisted-by: Claude Opus 5.5
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Bu6LgoxLR4wETqb7RPR2vt
The stop-string hold block ran for every message, building and replacing empty lists even for requests that never asked for logprobs, about 0.25 us per message (6-14% of the detokenizer's per-message cost on the Qwen3.8 tokenizer). Move it into DecodeStatus.release_logprobs and call it only when the message carries an entry or entries are still held. The new test pins that held entries are still released by a message without its own entry. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Bu6LgoxLR4wETqb7RPR2vt
_generate_full_impl ran untrimmed() on every ack even for requests without logprobs, and its result then kept whatever raw, unsegmented entries the acks carried. Read the flag once and skip both, so the protocol-neutral result carries entries only for a request that asked for them, as the streaming path already does. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Bu6LgoxLR4wETqb7RPR2vt
Fork-only shim over FlashML-org#739: FTWReader warns once and the dense, host-bank and device-bank read paths upcast fp16 NVFP4 globals to the fp32 the kernels read, so an FTW converted before FlashML-org#739 serves with its old numerics until reconverted. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
…mpletions A set value of a result-changing vLLM extra (allowed_token_ids, return_tokens_as_token_ids, guided_*, ...) is a 400 instead of being dropped by extra="allow"; null/false/empty still pass. Empty prompts and more than 256 prompts x n are 400s, and an error reply's code (context_length_exceeded) now reaches the completions body. On c23d2d5 the new deny-list test fails for all 10 fields with 200 != 400. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
… logprobs Sampling is masked to the allowed ids at every position (a LogitsPlan mask like min_tokens'). A row that also asks for logprobs gets them renormalized over the set from the unprocessed logits, vLLM's processed_logprobs at temperature 0, so the combinations that mode would fold in are 400s. The -inf tail and padding are dropped from top entries, non-finite values are sent as -9999, and completions logprobs go up to 32. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
tokens and top_logprobs keys become "token_id:N", vLLM's id-keyed shape, which also keeps candidates that decode to the same text from collapsing. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
prompt as list[int] or list[list[int]] goes to the engine as is, one request per prompt x n in prompt-major order. Ids are range-checked at the API when a frontend tokenizer is up and always in tokenize(), so an id past the vocabulary never reaches the embedding. echo or suffix with one is a 400. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
The completions deny-list now covers /v1/chat/completions too, plus allowed_token_ids and return_tokens_as_token_ids, which only completions implements: a set value is a 400 instead of a 200 that ignored it. Before this, the new chat test fails for all 10 fields with 200 != 400. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
An id at or past len(tokenizer) (2**31 without a frontend tokenizer) in logit_bias or stop_token_ids is a 400 on chat and completions before anything is queued, and the tokenizer worker checks them again. Before, a huge id overflowed msgpack at send_one (a 500) and one past the vocabulary was silently ignored by the sampler. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
--logprob-read-snapshot-donate {all,boundary,none} (default all, unchanged)
applies to finished requests with max_tokens=1 and logprobs on hybrid GDN
models: boundary donates only the frozen x64 prefill snapshot, none donates
neither it nor the live state and frees the pages, since a node without a
snapshot can never be resumed. A burst of label reads then cannot evict a
long conversation's snapshots.
Assisted-by: Claude Opus 5.5
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
…izing expert banks
…re raising in pinned_tensor
…g#700's method loader FlashML-org#700 renamed _legacy_expert_banks to _method_expert_banks and replaced load_expert_banks' dtype argument with the quant method, so the tests patched a missing name and passed an unknown keyword. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
…rm offload capture test The fake graph has no CUDA graph exec for cuGraphUpload; the test fails the same way on upstream cc56b6b. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
… test host's gpu W4A4 asks the real device whether it is worth it, so on an sm_120 host auto picks w4a4 and the triton assertion fails (also on upstream cc56b6b). Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
…e_start Streamed /v1/messages opened with input_tokens 0 and no cache fields (FlashML-org#718). generate_events now yields PromptAdmitted on the scheduler's admission ack and message_start waits for it, so its input/cache split matches message_delta. A keepalive or error before admission still sends message_start first, with zeros. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
…starts The context check ran only in the scheduler, after the worker tokenized, so a streamed request whose prompt did not fit got HTTP 200 and an in-stream error (FlashML-org#749). Every endpoint now renders and tokenizes the prompt frontend-side before any response and rejects a prompt that leaves no room for one output token with the scheduler's "prompt is too long" text and code context_length_exceeded, against the /v1/models context. A max_tokens that does not fit is still clamped by the scheduler: agents always send a large one. The worker serves the counted ids instead of tokenizing again. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
A call cut by max_tokens inside </parameter> streamed {"city":"Lima and then
fell back to {} at close, while non-streaming returned {} (FlashML-org#710). The
qwen3_coder detector now finalizes an open call at end of stream, dropping the
partial closer and closing the JSON, and the one-shot parse replays that
finalize, so both modes return the same valid arguments with the value kept.
Assisted-by: Claude Opus 5.5
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
…e bank reader threads The engine allocates the host banks under torch.inference_mode(), which the reader pool's threads do not inherit, so the in-place fp32 copy raised and the server refused to start. The test now loads the banks under inference mode, as the engine does. Assisted-by: Claude Opus 5.5 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846
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
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.
Rebuild of
5080deploy(c23d2d5d) onto upstream maincc56b6bd(36 new upstream commits), plus this round's additions. Rollback tag:5080deploy-pre-r7(=c23d2d5d=5080/base-c23d2d5d).Rebuild
integ7/resolutions.md. The main ones:qwen4_exp/weight.py: our TP_shard,_fp8_denseandemit()re-ported on fix(qwen4_exp): read weights under the checkpoint's quant config FlashML-org/FreeToken#759's_DenseReader.layers/embedding.py: the host embedding gather and FP8_logitsre-ported on feat(gguf): serve GGUF checkpoints on vendored ggml-cuda kernels FlashML-org/FreeToken#690'squant_methodtable.moe/offload_cache.py: our guards plus feat(moe): serve experts of several formats through per-format slot pools FlashML-org/FreeToken#700's layouts, with the requiredlayer_id.checkpoint/ftw.py: the fix(checkpoint): validate FTW v1 indexes eagerly FlashML-org/FreeToken#459 eager validation plus the fix(quant)!: keep nvfp4 global scales in fp32 FlashML-org/FreeToken#739 checks.d705ebb5(perf(moe): dequant nvfp4 prefill experts by bit placement and load activations contiguously FlashML-org/FreeToken#578 nvfp4 prefill dequant), replaced by perf(kernel): rewrite the triton nvfp4 moe prefill gemm FlashML-org/FreeToken#723/perf(moe): speed up the nvfp4 moe prefill gemm on rtx gpus FlashML-org/FreeToken#733.52609c55(compressed-tensors experts), replaced by fix(qwen4_exp): read weights under the checkpoint's quant config FlashML-org/FreeToken#759.~/models/Qwen3.8-Flash-Next-NVFP4-FTW-fp8-r7, made withFREETOKEN_FP8_DENSE=1 FREETOKEN_FP8_LMHEAD=1 ft checkpoint --quant-backend moe.nvfp4=triton.018c4c99+e67bd429still loads an fp16-global FTW, upcasting it under inference mode. Upstream proposal: Proposal: load an FTW with fp16 NVFP4 global scales (upcast with a warning) instead of refusing it after #739 FlashML-org/FreeToken#769..venv-r7=.venv-t214+ transformers 5.18.0 (build(deps): upgrade transformers to 5.18 FlashML-org/FreeToken#698), withggufremoved..venv-t214is kept for rollback.Upstream PRs merged
d75874baadapts its tests to feat(moe): serve experts of several formats through per-format slot pools FlashML-org/FreeToken#700's loader.pinned_tensor.Test fixes
Both also fail on clean upstream cc56b6b on this GPU; filed as FlashML-org#765.
cbf345b7: stub fix(engine): upload cuda graphs at capture so a stream wait cannot block the first launch FlashML-org/FreeToken#697's graph upload in a fake-graph test.840aa6db: make the minimax W4A4 selection test independent of the host GPU.AnyJev / vLLM scoring
FlashML-org#768 documents the design.
24de7e68: completions return a 400 for result-changing vLLM fields that are not implemented (none, false and empty values stay unset). Also: empty prompts get a 400, prompts × n is capped at 256, andcontext_length_exceededis passed through.d4d2606c:allowed_token_idsmasks sampling. A row that sends bothallowed_token_idsandlogprobsgets logprobs renormalized over the allowed set. Restricted rows return a 400 for temperature > 0, penalties, min_tokens or logit_bias. The logprobs cap is now 32.c501d830:return_tokens_as_token_idsreturnstoken_id:Nkeys.7ade7500: token-id prompts, range-checked intokenize().2bfdae17: the same deny-list on chat completions.605b1ae0:logit_biasandstop_token_idsids are bounded.10ea0821:--logprob-read-snapshot-donate {all,boundary,none}, defaultall(no change). Not A/B-tested on GPU yet.integ7/anyjev_check.py:token_id:Nkeys, and token-id prompts are bit-identical to the same text prompts.Upstream issue fixes
76f6f1e7: streamed/v1/messagesmessage_startcarries the admitted prompt usage (Streaming /v1/messages: message_start reports input_tokens 0 and no cache fields FlashML-org/FreeToken#718; fork Streaming /v1/messages: message_start reports input_tokens 0 and no cache fields #93).27cd6f15: a prompt that does not fit gets a 400 before the stream starts, on all four endpoints. An explicitmax_tokensthat does not fit is still clamped, because agent clients always send a large one. The frontend now tokenizes once and sends the ids to the worker ([Bug] Context overflow is checked only after tokenization in the worker: a streamed request gets 200 plus an in-stream error, and an oversizedmax_tokensis silently clamped FlashML-org/FreeToken#749).a36322e5: a qwen3_coder call cut inside</parameterkeeps its value, as the same valid JSON in both modes (qwen3_coder: max_tokens truncation leaks raw tool-call markup into content and arguments FlashML-org/FreeToken#710).Results (production flags, quiet machine)
5080/r7-batchfix(not merged).Upstream issues filed this round
FlashML-org#765, FlashML-org#766, FlashML-org#767, FlashML-org#768, FlashML-org#769.
🤖 Generated with Claude Code
https://claude.ai/code/session_01Td2gieja9Da2rPTN17Z846