Repository navigation
Conversation
…lashML-org#638) * fix(kernel): correct jit matcher options and pdl arch guards - with_dtype<Ts...>(sym) and with_device<Codes...>(sym) applied no options, so their dtype and device checks accepted anything - index and store list kDLROCM, which the ignored check let through - emit griddepcontrol only for !FREETOKEN_USE_ROCM && __CUDA_ARCH__ >= 900; it failed ptxas below sm_90 and does not exist on HIP - declare the launch attribute cache only on CUDA - use kWarpThreads for the index and store launch geometry Co-authored-by: zihaomu <zihaomu@outlook.com> * fix(kernel): treat zero as a matched size in tensor matcher - a zero dimension counted as unset, so empty inputs failed with "Size value is not set" and a zero-row tensor did not catch a mismatched one * fix(kernel): return early on empty index and store batches - an empty batch launched a zero-sized grid, which CUDA and HIP reject; the leftover error then failed the next GPU call Co-authored-by: zihaomu <zihaomu@outlook.com> --------- Co-authored-by: zihaomu <zihaomu@outlook.com>
…lashML-org#502) * feat(deepseek_v4): serve image input on the V4-Flash-Vision release The block layout depends on where the image lands: lead pads align its first row token to a 4-token compression stride, so the item hash carries the layout and two alignments never share a radix key. The block's tokens all carry one content pad id (the shared pad scheme); the sparse attention reads its spans from batch.mm_spans, which the scheduler recovers from the flat block ends its existing plan already reports. * fix(mm): keep the shared mmprocessor signature and own the dsv4 offset in its apply * fix(models): build dsv4 image spans in the model and start them at image_start * fix(models): mount the dsv4 vision tower under visual and keep its weights unquantized * fix(models): keep dsv4 image blocks whole across chunked prefill
* fix(rocm): gate cuda-only optional kernel integrations * test(rocm): narrow optional-backend regression coverage
…#652) * test(kernel): cover triton activation and norm fallbacks * test(kernel): allow fp32 activation approximation error
…e and read mxfp4 gemv weights as column pairs (FlashML-org#653)
* feat(server): require an API key on every route but /health * feat(shell): send the server's API key * feat(ctl): send the server's API key * feat(launch): hand the server's API key to the agent * feat(daemon): pass the serve's API key as apiKey The daemon reaches the serve's /v1/stats and prepare-stop itself, so the key travels in the start request, reaches the serve through its environment and is kept in the state file for re-adoption. The daemon's own FREETOKEN_API_KEY is no longer passed down to the serve. * fix(launch): keep the API key out of the variable dsh's web search reads dsh's web search reads DEEPSEEK_API_KEY and sends it to api.deepseek.com, so dsh now reads the key from FREETOKEN_API_KEY through the apiKeyEnv of its llm-deepseek section, and DEEPSEEK_API_KEY stays the dummy. * fix(server): check the API key on OPTIONS requests * fix(launch): write agent configs that hold the API key owner-only * chore: drop comments that restate the code
* fix(sampling): apply presence and frequency penalties * fix(server): reject non-finite sampling penalties * fix(server): bound presence and frequency penalties
…FlashML-org#662) * feat(cpu-moe): block-FP8 expert GEMV (W8A16) for the cpu/hybrid backends Block-fp8 checkpoints (Qwen3.5/3.6-FP8, DeepSeek-V3-style, GLM) had no CPU expert path at all: NotImplementedError: --moe-backend cpu/hybrid ... supports ['bf16', 'ds_fp4', 'mxfp4_triton', 'nvfp4', 'q4_0'] formats, but this checkpoint's experts are 'fp8_block' so every expert miss had to stream over PCIe at ~25 GB/s while the CPU sat idle. The bank is fp8-e4m3 row-major [rows, K] with one bf16 scale per 128x128 block -- K contiguous per output row, unlike mxfp4's transposed bank, so it maps straight onto a bf16 dot product. e4m3 -> bf16 is *exact* (3 mantissa bits into 7, and bf16's exponent range covers e4m3's whole 2^-9..448 span), so the AVX-512 path widens weights in-register and feeds `_mm512_dpbf16_ps` against the bf16 activations: identical products to the reference's fp32 math, differing only in summation order -- the same latitude the bf16 dot already takes. The normal range rebases the exponent with a shift and a constant +120; exp == 0 is subnormal (m * 2^-9) and does not obey that, so those lanes blend in from a small table. AVX2 does the same in fp32; scalar goes through the existing 256-entry e4m3 LUT. `ft bench bw` gains an fp8_block CPU measurement (it could already build the banks). Measured on main with this commit alone, `--dtype fp8_block --reps 5`, 2x Xeon Gold 6526Y: CPU MoE PCIe ratio backend picked main not measurable 25.1 -- offload main + this 53.7 GB/s 25.1 2.14x hybrid Before this commit the format reports no CPU number, so the backend picker has nothing to weigh and selects offload. The ratio is what decides the backend, so adding the kernel is what lets a block-fp8 deployment use the CPU at all. End to end on Qwen/Qwen3.6-35B-A3B-FP8, median of 3 generations at matched GPU expert residency: residency offload hybrid delta 10% 43.94 50.94 +15.9% 25% 82.65 55.12 -33% The win is real but conditional: hybrid pays off once the miss rate is high enough that PCIe is the bottleneck, and costs throughput when enough experts are resident that the CPU becomes the limiter. See the PR discussion -- the ratio makes `--moe-backend auto` select hybrid for every block-fp8 checkpoint, which is right for the oversubscribed case this project targets and wrong for a model that nearly fits, and the bench models bandwidth but not miss rate. Correctness is checked against the Triton block-fp8 decode kernel on identical banks and routing (bs 1/2/5), plus an exactness check on the e4m3 -> bf16 widening over all 256 codes. A weight-format ABI probe (`max_weight_format_id`) makes a stale prebuilt `.so` fail loudly instead of falling through to the wrong dequant branch and returning silently wrong numbers. Assisted-by: Claude * fix(cpu-moe): honor padded fp8 scale strides The scale banks' trailing dimension is padded so per-expert rows land on a 16-byte boundary for the fused copy (kernel/aot_models.py: fp8_block_scale_pad), so the row stride is not ceil(K/128). Recomputing it from the unpadded block count reads scales at the wrong offset. Qwen3.8 is the case in the tree: models/qwen3_5_moe/weight.py records down_scale as 20x5 bf16 = 200 B, and 200 % 16 != 0 pads cols 5 -> 6. The benchbw workload (H=2048, I=512) and DSV4-Flash (H=7168, I=2048) both happen to give rows*cols divisible by 8, so neither trips it. Read the stride off the tensor and pass it through to the kernel, and let the shape check assert the unpadded extent is present rather than an exact shape. The fixture now allocates padded banks, so the existing GPU cross-check covers this too; the GPU kernel was already stride-aware. * test(cpu-moe): add float64 parity reference and an avx512f fp8 tier The existing GPU-vs-CPU check compares two fp32 accumulations at a 2e-2 tolerance, which cannot see accumulator precision loss. Add a float64 reference over mixed-sign, full-range e4m3 weights and zero-mean activations, and force each ISA tier explicitly so the sweep is real. Also add an avx512f tier between avx512bf16 and avx2, so CPUs with AVX-512 but no BF16 extension (Skylake-SP through Ice Lake-SP) stop falling back to 8-wide fp32. Measured on Xeon Platinum 8368 (Ice Lake-SP, no avx512_bf16), three repetitions each: +19% single-thread cache-resident (3.31 -> 3.94 GB/s, disjoint intervals), but only +4% under `ft bench bw --dtype fp8` at 75 threads (60.4 -> 62.7 GB/s, overlapping intervals). At that thread count the kernel is memory-bound, so the project's own benchmark cannot resolve the tier; the single-thread gain is the honest figure. * test(cpu-moe): cover fp8 dispatch and cpu-only regressions The new tier shipped without test coverage: deleting the avx512f branch from select_fp8dot left every test green, because isa_name() reports select_dot() (the bf16 chain) rather than the fp8 dispatch. Injecting a ragged-tail bug into dot_fp8_block_avx2 was invisible for the same reason -- no test ran numerics under a forced tier. Force each available fp8 tier through the float64 parity GEMV, expose the selected fp8 tier so the assertion can see it, and make the padded bank test run a GEMV instead of only constructing an executor. Also move the module-level CUDA guard onto the tests that actually need a GPU. A CPU MoE kernel's most relevant hosts are the ones without a usable GPU, and all ten tests were skipping there. * review: force each ISA tier in the fp8_block test, and stop the docstring overclaiming The comparison test only varied batch size, so all three cases ran whatever tier this machine auto-selects and the scalar / AVX2 / AVX-512-BF16 fp8_block kernels were never compared against the GPU at all. Parametrise FREETOKEN_CPU_MOE_ISA over the four tiers. Setting the variable alone would not be enough: pick_isa() caps *downward*, so asking for a tier the CPU or build lacks silently returns a lower one and the test would pass four times over the same kernel. The executor latches its choice into isa_name() at construction, so read it back and skip when it does not match what was asked for. On a machine with all four, that is 4 x 3 cases against a 8-expert, 256-hidden problem. test_subnormal_and_sign_decode_exactly said it validated that "every e4m3 byte must widen to the value the reference LUT gives", but it never calls the CPU kernel -- it checks that bf16 round-trips e4m3 losslessly and that subnormals are finite and small. That is the premise the widening rests on, which is worth pinning, but it is not the widening. Renamed to test_bf16_carries_every_e4m3_byte and the docstring now says what the body asserts. Not run here: this file is skipif(not torch.cuda.is_available()), and this machine has no GPU and no torch. Compiles clean and ruff reports the same two pre-existing E741 findings before and after. Assisted-by: Claude * fix(moe): let fp8_block experts decode on the cpu executor * fix(cpu-moe): build the bf16 fp8 tier only when the compiler has avx512bf16 * fix(bench): version the bw profile and ignore older ones * perf(cpu-moe): read fp8 weights as fp16 bits in the cpu dot kernels * refactor(cpu-moe): write the fp8_block bank checks like the other formats * refactor(cpu-moe): name the fp8 cpu symbols fp8_block * docs(server): list fp8_block in the --moe-cpu-layers help --------- Co-authored-by: Gabriel A. Devenyi <gdevenyi@gmail.com> Co-authored-by: ChenyuHe <hechenyu@zju.edu.cn>
# Conflicts: # docs/models.md # python/freetoken/attention/triton.py # python/freetoken/kernel/triton/attention.py # python/freetoken/kvcache/__init__.py # python/freetoken/kvcache/qsa_pool.py # Conflicts: # docs/models.md # python/freetoken/attention/triton.py # python/freetoken/kernel/triton/attention.py # python/freetoken/kvcache/__init__.py # python/freetoken/kvcache/qsa_pool.py # Conflicts: # docs/models.md # python/freetoken/attention/triton.py # python/freetoken/kernel/triton/attention.py # python/freetoken/kvcache/__init__.py # python/freetoken/kvcache/qsa_pool.py # Conflicts: # docs/models.md # python/freetoken/attention/triton.py # python/freetoken/kernel/triton/attention.py # python/freetoken/kvcache/__init__.py # python/freetoken/kvcache/qsa_pool.py # tests/kernels/test_e4m3_compat.py # Conflicts: # docs/cli.md # docs/models.md # python/freetoken/attention/triton.py # python/freetoken/kernel/triton/attention.py # python/freetoken/kvcache/__init__.py # python/freetoken/kvcache/base.py # python/freetoken/kvcache/qsa_pool.py # tests/kernels/test_e4m3_compat.py # Conflicts: # docs/cli.md # docs/models.md # python/freetoken/attention/triton.py # python/freetoken/kernel/triton/attention.py # python/freetoken/kvcache/__init__.py # python/freetoken/kvcache/base.py # python/freetoken/kvcache/qsa_pool.py # tests/kernels/test_e4m3_compat.py # Conflicts: # docs/cli.md # docs/models.md # python/freetoken/attention/triton.py # python/freetoken/kernel/triton/attention.py # python/freetoken/kvcache/__init__.py # python/freetoken/kvcache/base.py # python/freetoken/kvcache/qsa_pool.py # tests/kernels/test_e4m3_compat.py
# Conflicts: # tests/tokenizer/test_tokenize.py # Conflicts: # tests/tokenizer/test_tokenize.py # Conflicts: # tests/tokenizer/test_tokenize.py
…d-request # Conflicts: # python/freetoken/engine/config.py # Conflicts: # python/freetoken/engine/config.py # Conflicts: # python/freetoken/engine/config.py # Conflicts: # python/freetoken/engine/config.py # Conflicts: # python/freetoken/engine/config.py # Conflicts: # python/freetoken/engine/config.py
… is measurable # Conflicts: # python/freetoken/engine/engine.py # Conflicts: # python/freetoken/engine/engine.py # Conflicts: # python/freetoken/engine/engine.py # Conflicts: # python/freetoken/engine/engine.py # Conflicts: # python/freetoken/engine/engine.py # Conflicts: # python/freetoken/engine/engine.py
…acked on FlashML-org#385 TP) # Conflicts: # python/freetoken/models/qwen4_exp/attention.py # python/freetoken/models/qwen4_exp/config.py # python/freetoken/models/qwen4_exp/weight.py # Conflicts: # python/freetoken/models/qwen4_exp/attention.py # python/freetoken/models/qwen4_exp/config.py # python/freetoken/models/qwen4_exp/weight.py # Conflicts: # python/freetoken/models/qwen4_exp/attention.py # python/freetoken/models/qwen4_exp/config.py # python/freetoken/models/qwen4_exp/weight.py # Conflicts: # python/freetoken/layers/quantization/moe/nvfp4.py # python/freetoken/models/qwen4_exp/attention.py # python/freetoken/models/qwen4_exp/config.py # python/freetoken/models/qwen4_exp/weight.py # Conflicts: # python/freetoken/layers/quantization/moe/nvfp4.py # python/freetoken/models/qwen4_exp/attention.py # python/freetoken/models/qwen4_exp/config.py # python/freetoken/models/qwen4_exp/weight.py # Conflicts: # python/freetoken/layers/quantization/moe/nvfp4.py # python/freetoken/models/qwen4_exp/attention.py # python/freetoken/models/qwen4_exp/config.py # python/freetoken/models/qwen4_exp/weight.py
… bound, not an oracle
…lit differs from the fp8 env flags
…when its slot floor does not fit
Upstream FlashML-org#517 tests without its docs and pyproject dev-extra: the SDK tests importorskip anthropic/openai/httpx so the server suite still collects without them. (cherry picked from commit b044276, tests only) Assisted-by: Claude Opus 5.5 (1M context) (cherry picked from commit 50f99d1)
…lp, which it does not serve Assisted-by: Claude Opus 5.5 (1M context) (cherry picked from commit c6b7ce7)
…rging system messages render_messages hoisted system ahead of a leading developer message, which the tokenizer then merged as "S\n\nD"; it now leaves developer requests to the tokenizer's per-template mapping. Both merge sites joined str() of an image part list into the prompt, data URL included; they now raise the template's 400 instead. Assisted-by: Claude Opus 5.5 (1M context) (cherry picked from commit e60b2c3)
…is served here Assisted-by: Claude Opus 5.5
… in every layer An unpinned layer's copy_missing is the whole-layer pageable copy, which cannot serve the LRU slot remap, so it keeps streaming. The per-layer fallback when a layer routes to more than half the cache is gone: the engine's num_experts slot floor already keeps every chunk resident, and a uniform path cannot leave a later layer in an overlap buffer that was never begun. Assisted-by: Claude Opus 5.5 (1M context) (cherry picked from commit 37c2479)
… kwargs Assisted-by: Claude Opus 5.5
`torch.equal` fails for all three cut points on torch 2.11.0+cu130 / flashinfer 0.6.18 / triton 3.6.0 (sm_89): cuBLAS picks a different kernel for M=5000 and M=5000-cut, so the indexer's q and k round differently, a few rows flip one block at the top-k margin, and the layer output differs by up to 2.6e-3 (a few bf16 ulps). Both paths are deterministic run to run. Use the tolerance of the neighbouring oracle tests. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0173pf9k9fSVtwbm3f898HDt (cherry picked from commit 3849e3c) Assisted-by: Claude Opus 5.5 (1M context) (cherry picked from commit e700e8f)
…o the one-shot run At bf16 output tolerance a broken dual-source compress only swaps a block in a few rows and passes; the slab keys are unit-scale, so a wrong straddling-group key stands out. Assisted-by: Claude Opus 5.5 (1M context) (cherry picked from commit cedf179)
…ights host) The embedding is a row lookup, so a single-GPU engine can leave the table in page-locked host RAM and gather each step's rows over UVA with the PLE row kernel (5 KiB per decode token, graph-safe): 1.2 GiB of VRAM back for Qwen3.8-Flash-Next. Assisted-by: Claude Opus 5.5 (1M context) (cherry picked from commit 9f2b7a2)
…lacement Assisted-by: Claude Opus 5.5 (1M context) (cherry picked from commit ccfb2b8)
No regression test: overflowing it needs an output of 2^31+ elements (4 GiB+ in bf16). Assisted-by: Claude Opus 5.5 (1M context) (cherry picked from commit cc5af7a)
Assisted-by: Claude Opus 5.5 (1M context) (cherry picked from commit a8b2f0a)
… 1 at config time Assisted-by: Claude Opus 5.5 (1M context) (cherry picked from commit 579e03a)
…lashML-org#573 passes Assisted-by: Claude Opus 5.5
…am so a post-launch fence fails as a hang Upstream FlashML-org#588's _graph_consumed fence is left out of the merge: our graph fill runs on the filler thread before its own launch passes its WAIT, so that fence deadlocks the second step. Assisted-by: Claude Opus 5.5
…t placement and load activations contiguously # Conflicts: # tests/kernels/test_e4m3_compat.py # Conflicts: # python/freetoken/kernel/triton/e4m3_compat.py # tests/kernels/test_e4m3_compat.py # Conflicts: # python/freetoken/kernel/triton/e4m3_compat.py # python/freetoken/kernel/triton/nvfp4_fused_moe.py # tests/kernels/test_e4m3_compat.py # Conflicts: # python/freetoken/kernel/triton/e4m3_compat.py # python/freetoken/kernel/triton/nvfp4_fused_moe.py # tests/kernels/test_e4m3_compat.py
…es the one FlashML-org#603 added to main Assisted-by: Claude Opus 5.5
…m policy instead of hoisting them FlashML-org#487's hoist ran before FlashML-org#608's fold on /v1/messages and moved a late system reminder into the system head, so every reminder missed the prefix cache; the OpenAI and Responses paths keep the hoist the qwen templates need. Assisted-by: Claude Opus 5.5
…nd the count fake the inline-system kwarg Assisted-by: Claude Opus 5.5
…stops the server
…ctory that already holds an FTW
…d experts The expert spec now takes its tensor names and global-scale direction from the checkpoint's quant config, as qwen3_5_moe does, so RedHatAI/Qwen3.8-Flash-Next-NVFP4 loads and converts; the nvidia modelopt checkpoint keeps the same spec. Assisted-by: Claude Opus 5.5
…g#610's blocks_per_bank through Assisted-by: Claude Opus 5.5
…ore growing ones in live rebuild (Closes FlashML-org#643)
… NVFP4 host banks
… ties in the ds_fp4 transcription check
…ader and gdn group the fit check reads The carried host kv tier (FlashML-org#499) rejects a pages rebuild through engine.kv_offloader, and the fork's gdn workspace reserve asks the model config for its linear attention group; the SimpleNamespace shell had neither. 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
Main (FlashML-org#527) bounds presence/frequency penalties to [-2, 2] on the request models, so they never reach the handler's 400 path. 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 wide prefill profile (BLOCK_N 64, 2 warps) that the nvfp4 path inherits from the bf16 tuning has load_nvfp4 decode a 64 x 256 fp32 tile for K and V on 64 threads. Triton 3.8 compiles it to 40 registers with ~1.4 KB of spills: 439 ms instead of 45 ms at 4096 rows on sm_120, and a ~556 MiB driver local-memory reservation outside the torch allocator that OOMs this deployment on its first prefill. BLOCK_N 16 at 4 warps does not spill: 12.3 ms on Triton 3.8 (33 ms on 3.6), same output to bf16 rounding. Upstream: FlashML-org#675, root cause posted on FlashML-org#408. 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
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.
Stacked on #84. The base
5080/base-1639600bis the previous 5080deploy tip. This PR rebuilds 5080deploy on upstream main9c635a15, which adds 11 commits since07813248, including FlashML-org#630 (torch 2.14.1, Triton 3.8, sglang-kernel 0.4.9), FlashML-org#654, FlashML-org#527, FlashML-org#639 and FlashML-org#669.Superseded and dropped:
0067996b, the Triton sampling-softmax NaN fix. fix(kernel): avoid nan in triton sampling softmax when a chunk starts with -inf FlashML-org/FreeToken#654 fixes it upstream, and the dropped fix's own tests pass against fix(kernel): avoid nan in triton sampling softmax when a chunk starts with -inf FlashML-org/FreeToken#654 (11/11).Added:
test_gemv_path_matches_the_transcriptionfailure).695bfc95fix(kernel): narrow NVFP4 QSA prefill tiles at head_dim 256. Under Triton 3.8, the carried feat(kvcache): add nvfp4 kv quantization FlashML-org/FreeToken#408 NVFP4 QSA prefill profile (BLOCK_N=64, 2 warps) compiles to 40 registers with 1.4 KB of spills per thread. The kernel runs ~10x slower (439 vs 45 ms at 4096 rows), and the driver reserves 556 MiB of local memory outside the torch allocator, which OOMed this deployment on its first prefill. WithBLOCK_N=16and 4 warps it does not spill: 12.3 ms on 3.8 and 33 ms on 3.6. Its output matches the old profile to bf16 rounding, and it has a regression test. Reported as NVFP4 QSA sparse attention (#408) spills under Triton 3.8: ~10x slower prefill and ~550 MiB local memory FlashML-org/FreeToken#675, with the root cause commented on feat(kvcache): add nvfp4 kv quantization FlashML-org/FreeToken#408.Replay: 101 first-parent steps, with rerere. Hand resolutions:
--api-key(feat(server): add --api-key authentication FlashML-org/FreeToken#639) and--dist-port(fix(server): loopback-bind the TP rendezvous store, add --dist-port FlashML-org/FreeToken#461): a union inargs.pyanddocs/cli.md.Environment: a new venv
.venv-t214with torch 2.14.1+cu130, Triton 3.8.0, sglang-kernel 0.4.9, flashinfer 0.6.18.post1 and transformers 5.16.1. The old torch 2.11.venvis kept for rollback.A/B (interleaved, 2 rounds, RTX 5080, production config):
1639600b(torch 2.11)Numerics: greedy outputs change (49/60 cold nonce prompts). New cold fingerprints: code
4ac964c44c6bc68b, ctx08996814db15744c, think8c72d19d0ae12d01, identical across both cold starts. Paired NLL on the frozen corpus (207k tokens):Tests:
Not merged:
🤖 Generated with Claude Code
https://claude.ai/code/session_01Bu6LgoxLR4wETqb7RPR2vt