Repository navigation
perf(qwen4exp): enable guarded GPU argmax on gfx1151 - #823
Draft
dusterbloom wants to merge 202 commits into
Draft
dusterbloom wants to merge 202 commits into
dusterbloom wants to merge 202 commits into
Conversation
- QWEN4EXP_DENSE_TABLE: per-shape measured dense-GEMM dispatch (gfx1151, T=16366), routing by measured winner (mostly MMQ over MMB/cuBLAS). - QWEN4EXP_HC_TILE16: 16-row HC gate/mix tile on RDNA3.5, byte-exact. - expert-major MoE pipeline implemented but neutral; rejected, kept default-off behind QWEN4EXP_MOE_PIPELINE. - harness BENCH_MAX_TOKENS / BENCH_THINKING / BENCH_EFFORT experiment knobs. - delivered @16366 tokens: 1063.1 t/s median / 1065.5 best (control 994.0 / 1002.2), min-of-8 all >1000; quality HE 10/10, GSM 10/10, Math 9/10, recall 2/2; GSQ + IQ4_NL reference profiles byte-exact. - evidence under docs/performance/qwen4exp-1100/.
…N by default These three were validated wins in 52a40e5 but shipped default-off, so the released binary ran the slow path. Now on by default (opt out with =0). - QWEN4EXP_DENSE_TABLE=1, QWEN4EXP_HC_TILE16=1, QWEN4EXP_LAST_TOKEN_FFN=1 - QWEN4EXP_UPSTREAM=1 still excludes the two dispatch candidates, so the reference profile is unchanged. Verified on IQ4_NL with no candidate flags: - quality HE 10/10, GSM 10/10, Math 9/10, recall 2/2 - prefill min-of-8 @16366 tokens: TTFT 15.27 s -> ~1072 t/s - reference differential (GSQ, QWEN4EXP_UPSTREAM=1): byte-exact, 0/160 expert-ID mismatches, no onset.
QWEN4EXP_DECODE_REUSE retains a T=1 decode metadata context and graph allocator in Qwen4ExpCache: each token resets the metadata arena, rebuilds the graph, remeasures allocator assignments, and reuses the allocator backing buffers. Default-on; opt out with =0; excluded under QWEN4EXP_UPSTREAM=1. Decode min-of-N (IQ4_NL, 1 warm-up + 8): - ~2k: 27.87 -> 28.55 t/s (+2.43%) - ~16k: 24.62 -> 25.34 t/s (+2.93%) Gates: exact constrained greedy output/token count at both depths; reference differential byte-exact (ratios 1.0000, 0/48 expert-ID mismatches); quality HE 10/10, GSM 10/10, Math 9/10, recall 2/2; prefill min-of-8 unchanged (15.30 s @16366).
The T=1 decode step now runs a stable, per-KV-bucket cgraph (fixed views, bucketed kv_len, single input buffer) that HIP captures once and replays: ~2,236 replays / 2,240 calls, one hipGraphLaunch per steady-state token carrying all ~2,800 kernels. Build/allocation per token drops from 0.95->0.011 ms (2k) and 1.21->0.010 ms (16k). Decode min-of-N (IQ4_NL, 1 warm-up + 8): - ~2k: 28.58 -> 29.89 t/s (+4.58%) - ~16k: 25.32 -> 28.06 t/s (+10.85%) Default-on: QWEN4EXP_DECODE_STABLEGRAPH (opt out =0) and the HIP graph replay build option (server/CMakeLists.txt now defaults GGML_HIP_GRAPHS=ON; -DGGML_HIP_GRAPHS=OFF to disable). Excluded under QWEN4EXP_UPSTREAM=1. Gates: exact greedy token stream; quality HE 10/10, GSM 10/10, Math 9/10, recall 2/2; reference differential byte-exact (no onset); prefill min-of-8 unchanged (~1068 t/s @16366).
Adds 22 measured GSQ tuples to the default-on QWEN4EXP_DENSE_TABLE, selecting MMQ (instead of the slow MMB path) for exact-16366-token, gfx1151 qwen4exp runs. Excluded under QWEN4EXP_UPSTREAM=1. GSQ prefill min-of-N @16,366: 579.63 [579.33-583.97] -> 757.99 [754.33-764.33] t/s (+30.77%). Gates: exact planted recall, greedy "Paris", GSQ quality HE10/GSM8/Math10/recall2, GSQ+IQ4 upstream differentials no onset 0/ID mismatches, IQ4 prefill 1073.69 t/s and decode thermally matched. A BF16 rocBLAS extension reached 857.93 t/s but corrupted long-context output and was rejected (rocBLAS/f16-fallback risk; see log E26).
T12 (all-arch regression) found that with HIP graphs on, qwen4exp missed one standard 13k recall fact (7/8 matched cases vs 8/8 with graphs off) and the first greedy response was not byte-stable on first capture; fresh repeats passed 6/6 but the signal blocks an unconditional ship. Restore GGML_HIP_GRAPHS=OFF (server/CMakeLists.txt) and make the stable T=1 graph opt-in (QWEN4EXP_DECODE_STABLEGRAPH=1). No throughput regressions were seen on qwen35/deepseek4/qwen4exp, but correctness gates win. The E23/E24 code stays for triage (T14).
T14 root-caused the E30 quality variance to pre-existing prefill numerics (full-vocab logit hashes already differ at step 0 even with HIP graphs OFF, QSA/MMB/tiled-GDN off, and the upstream profile) rather than the stable graph, and showed HIP capture/replay is throughput-neutral in a matched stable-path A/B. So re-enable only the stable T=1 graph (cached cgraph + allocator, which carries the decode win) while keeping GGML_HIP_GRAPHS=OFF. QWEN4EXP_DECODE_STABLEGRAPH is default-on (opt out =0); excluded under QWEN4EXP_UPSTREAM=1. Decode min-of-N (IQ4_NL, 1 warm-up + 8), same graphs-off build, paired: - ~2k: 27.85 -> 29.70 t/s - ~16k: 22.40 -> 26.60 t/s (E23 order-controlled: 25.32 -> 28.06, +10.85%) Gates: HE 10/10, GSM 10/10, Math 9/10, recall 2/2; prefill min-of-8 15.31 s @16366; reference differential no onset; T14 graph-on recall 16/16.
- fattn: move the sparse-mask assertion after the DeepSeek4 maskless dispatch (the qwen4exp guard was aborting DS4 maskless prefill). - backend: drop the process-global GGML_CUDA_MMB setenv; implement real park/unpark (release + reload weights/cache, propagate errors). - graph: require every full-attention layer to pass qsa_layer_ok before eliding the dense mask; parse QWEN4EXP_QSA by value (=0 now disables). - loader/cache: free buffer + close PLE reader on failure; clear the PLE n-gram window on reset/free; reject malformed GDN metadata. - hip_compat: add cudaErrorNoDevice; drop a self-referential __launch_bounds__. - remove rejected/diagnostic-only paths (QWEN4EXP_MOE_PIPELINE + check, QWEN4EXP_KQ_MASK_DEV); CUDA_CHECK HIP calls; exhaustive CLAMP rejection list. - tools/docs/tests: derive repo paths (no hardcoded /home/duster); full flag inventory in docs/qwen4exp.md + ENVIRONMENT.md; qwen4exp capability regression tests; fix two stale tests. Validated: ROCm build PASS; ctest 730/730 (13 expected skips); quality HE 10/10, GSM 10/10, Math 9/10, recall 2/2; upstream differential no onset; prefill N=8 1065.84 [1063.42-1066.88] t/s (+0.79%).
Replaces the previous approach (b4bdf73), which fixed a read-after-free by dropping the CONCAT-time interception. That forced full materialization of the conv-input concat (37 tensors ~[16369,10240], about 24.8 GB at T=16,366) and cost +2.4 s prefill (IQ4_NL 17.10 vs 14.49 s; GSQ 23.46 vs 20.95 s). This restores the fast interception and makes the dependency explicit to the allocator through unused input slots, so the fused kernel reads tensors whose lifetime is graph-encoded at the node where it runs: - GDN: ordinary SSM_CONV ignores src[3]; it is set to the transpose input. - PLE: CONT ignores src[1]; it is set to the normalized input for the first shift. Gates: 16,366 prefill at/above the pre-fix parent (IQ4_NL 14.340 s / 1141 t/s, GSQ 20.715 s / 790 t/s); C=1024 N=3 fresh-process snapshots byte-identical (1f34fa...926c, matching the correct output); quality HE10/GSM10/Math9/recall2; QWEN4EXP_UPSTREAM=1 differential no onset, 0 expert-ID mismatches; build OK.
Qwen4ExpPleReader::gather() partitioned the serial-fallback range by the configured thread-pool size, so a T=1 gather (fewer rows than workers) filled only the first worker's slice and left the remaining per-layer-embedding rows zero. Partition by the actual worker count (1 on the serial path). Found via the batched-decode layer-by-layer diff; after the fix, solo quality (HE 10/10, GSM 10/10, Math 9/10, recall 2/2) and the QWEN4EXP_UPSTREAM=1 differential (no onset, 0 expert-ID mismatches) are unchanged.
Adds qwen4exp_forward_batched(): one graph with T=active_slots independent one-token rows, running the shared embedding/projection/MoE/HC ops batched (weights read once) and the per-slot full-attention, GDN recurrence and PLE against each slot's own Qwen4ExpCache, emitting one logits row per slot. Gated by QWEN4EXP_BATCHED_DECODE=1 (default off; excluded under QWEN4EXP_UPSTREAM=1). The serving scheduler does not call it yet. Uses one shared caller-owned workspace; N=1 delegates to the single-sequence path. Validated: N=1 bit-identical to the existing path; 4 identical slots mutually bit-identical; untouched-slot isolation, row permutation and cancel/reset/reuse pass; batched-vs-solo residual logit deltas ~1e-3 at boundaries growing to ~2 through routing (margin policy). N=4 @32768 fits the 96 GiB GTT pool with ~12 GiB headroom.
Adds Qwen4ExpSeqEngine implementing the shared SeqEngine contract: one full-context Qwen4ExpCache per slot (F16, no paging), one shared batched-decode workspace, a single FIFO prefill owner with 512-token slices, per-slot reset, and bounded-progress plan handling. Decode batch runs via qwen4exp_forward_batched(). Fixed the engine to advance the per-slot Qwen4ExpCache::cur_pos after each successful prefill/decode (the single-sequence forward leaves it at zero). Gated behind LUCE_QWEN4EXP_SEQ_ENGINE=1 + QWEN4EXP_BATCHED_DECODE=1 (both default off; excluded under QWEN4EXP_UPSTREAM=1). When enabled, the feature gate permits --max-concurrency 2..4 at --max-ctx 32768 with one local device and no --kv-pool-tokens; otherwise the existing --paged-attention requirement is unchanged. Wired through ModelBackend::seq_engine(). Validated at N=2/N=4: real-engine SeqEngine contract PASS, randomized admit/retire soak PASS (peak GTT 75.4 GB -> 18.6 MB teardown), and 4 concurrent distinct requests all coherent (distinct-concurrent=PASS); only slot 0 diverged from solo at step 7 (epsilon 0.509 vs solo margin 0.021, within the margin allowance). The decisive quality-under-batch gate is P3.
…engine server_main.cpp auto-enabled --paged-attention whenever --max-concurrency > 1, which conflicts with the qwen4exp full-cache sequence engine that requires paging off. Add a narrow exception: skip the auto-enable only when LUCE_QWEN4EXP_SEQ_ENGINE=1 and QWEN4EXP_BATCHED_DECODE=1 (and not QWEN4EXP_UPSTREAM); every other path is unchanged. Validated end-to-end (N=4, --max-ctx 32768): the server starts in concurrent mode, 4 simultaneous distinct requests return correct independent answers (no cross-talk), and the sequential quality suite still passes (HE 10/10, GSM 10/10, Math 9/10, recall 2/2). Goodput N=1 vs N=4: per-request decode 27.9 -> 7.5 t/s, aggregate ~1.08x -- MoE expert streams do not amortize across rows.
The Qwen3.8-Flash-Next kernel suite used AMD-specific intrinsics that broke
non-gfx1151 targets:
- CUDA/NVIDIA: ggml-cuda.cu/fattn.cu referenced the HIP-only MMB/HC/QSA symbols
unconditionally, but the CUDA build excludes those TUs -> link failure. All 65
references are now behind #if defined(GGML_USE_HIP) with the generic
matmul/attention path restored for CUDA.
- gfx1201 (RDNA4 / ROCm 7.2.2): mmb.cu and qsa.cu failed to compile ('Cannot
select intrinsic llvm.amdgcn.wmma.f32.16x16x16.f16', plus the BF16 WMMA
builtin). The real MMB/QSA implementations are now compiled only for gfx1151,
other targets get link-compatible unsupported stubs so the supported_* checks
fall back to generic kernels, and MMB runtime enablement is restricted to
RDNA3.5.
Verified: gfx1151 build (luce_server + smoke), IQ4_NL forward smoke (T=16
prefill + decode, finite logits), HTTP greedy smoke (Paris), and compile-only
checks for mmb.cu/qsa.cu at gfx1201. CUDA and full gfx1201 link are validated by
CI.
The concurrent engine pinned a single FIFO prefill owner and capped the whole cohort at 512 prefill tokens, so prompts never shared a weight read with each other or with the live decode rows. Decode was already width-batched (all dense and routed matmuls ran at T = live rows). Packed ragged prefill gives each slot its own 512-token chunk, up to the graph's 2048-row budget minus live decode rows, so independent sequences share one traversal for dense projections, HC, router and MoE, while attention KV, GDN recurrence and PLE history stay per-slot. Whole-wave aggregate goodput (Halo, 3 fresh servers per width): N=1 11.45, N=2 14.22, N=3 15.86 (1.39x), N=4 15.42 tok/s, +22% at N=4 over 12.64. Quality, contract, soak, HTTP edge gates and QWEN4EXP_UPSTREAM parity all pass. Opt-in via LUCE_QWEN4EXP_SEQ_ENGINE=1 + QWEN4EXP_BATCHED_DECODE=1.
Issue Luce-Org#675 asked how to see what is responsible for growing host RAM: KV-cache size in tokens and bytes, device versus host allocations, whether memory is released after compaction, and DFlash host buffers. None of it was visible without reading the code. /props gains a `memory` section: - `cache`: the live cache the model decodes from, allocated once for --max-ctx: location (device or host), capacity and live tokens, and bytes split into attention K/V, recurrent state, DFlash target features and the rest of the allocation, plus live host-side decode state (DeepSeek4 logits and DSpark feature window). - `snapshots`: every saved snapshot with its slot, role (prefix, agent_turn, prefill_cache, disk_staging), tokens, bytes and location, with host and device totals. On discrete GPUs this is where RAM grows when an agent conversation keeps one snapshot per turn. - `process`: resident and peak RSS, read when /props is served. ModelBackend::memory_report() builds the numbers from the real buffers (Qwen and DeepSeek4; other backends report nothing). In DeepSeek4 concurrent serving the cache is the paged cache plus the per-slot image staging caches, and host vectors count by capacity, since trimming them keeps their allocation. The generation thread publishes the report at startup and after every request, in the classic worker and in the concurrent scheduler, and /props serves the stored copy, so it never touches backend state from the client thread. The OpenAPI spec and props-endpoint.md describe the new section; the change is additive, so props_schema stays 2. Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…uce-Org#779) * perf(qwen35): adaptive DFlash2 verify width in the concurrent engine The concurrent engine verified every chain at the full draft block. With several lanes the batched verify is compute-bound, so a round of 8 lanes x 16 rows mostly verifies drafts that are rejected early. Each batch bucket now keeps an AdaptiveSpecWidth and picks the verify width per round from {4, 8, block}: expected committed tokens from per-depth conditional acceptance (a depth counts only when offered with every shallower candidate accepted, so narrow rounds do not bias deeper estimates) divided by the measured round cost. Each width is offered three times per bucket first; the cold first offer is not a cost sample. The drafter still proposes the full block and a round verifies its prefix. On by default for batched serving; LUCE_ADAPTIVE_SPEC_WIDTH=0 restores the fixed block width. Single-request speculative decoding is unchanged. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * fix(qwen35): review fixes for the adaptive verify width - Key the width controllers by the exact speculative lane count instead of the graph bucket, which spans several lane counts. - Take cost samples only from rounds without AR peers and without a graph rebuild, so build time and mixed loads do not bias the costs. - Scope the default-on note to DFlash2 chain speculation. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * fix(qwen35): calibrate widths on clean samples; learn from verified prefixes - Calibration offers a width until it has two clean cost samples (capped at eight offers), so the cost model is active after calibration even when rounds often carry AR peers or graph rebuilds. - Acceptance statistics use the target-verified prefix before the room and min-token clamps, which are serving policy, not rejections. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * fix(qwen35): saturate width calibration counters; reuse per-round buffers Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * fix(qwen35): observe the round's majority lane last The controller's clean-draft flag comes from its last observe(); the lane whose outcome matches the round majority now goes last, carrying the round's cost sample. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * fix(qwen35): key adaptive width by the verify graph bucket The verify graph is built for tree_bucket lanes, with missing lanes as padding, so a round's cost is set by the bucket and width rather than the number of real lanes. Keying by bucket keeps each cost model consistent and calibrates each bucket once. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> --------- Co-authored-by: mrciffa <davide@lucebox.com> Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…ged serving (Luce-Org#781) * perf(deepseek4): stage long text prompts through sparse prefill in paged serving Paged DeepSeek4 serving prefilled every prompt row in the gathered graph, one row per lane per step on the R9700 + Strix Halo deployment: a 1843-token prompt took 85 s to its first token. With --ds4-prefill sparse (now accepted with --paged-attention), text prompts of at least 64 tokens take the path image prompts already use: prefill into the slot's staging cache on the single-request sparse layer-major path, copy into the paged slot, then prefill the last token and decode in the gathered graph. The monolithic deployment shares the layer-sliced staged pass. The heterogeneous one runs one chunk of the regular hybrid prefill per step (2048 rows idle, 1024 while others decode), since the shared pass needs the whole model on one GPU. Staged requests are identified by request id, not by image handle. Without vision the staging caches are optional: if they do not fit, text prompts keep the gathered prefill. --ds4-prefill exact is unchanged. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * fix(deepseek4): review fixes for staged text prefill - Keep the image handle in in-flight staged-pass members so a request retired mid-pass cannot free image rows the pass still reads; the request id stays the identity check. - Image blocks take the shared pass budget before text prompts. - Heterogeneous staged chunks use the same placement-aware long-context cap as single-request sparse prefill. - Free the slot-0 cache too when text staging does not fit. - The 64-token threshold counts the whole prompt. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * fix(deepseek4): fall back to gathered prefill when slot 0 staging does not fit Text-only sparse paged serving now treats a failed slot-0 staging cache like the other slots: text prompts keep the gathered prefill. Image serving still fails startup, since it needs the staging caches. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * fix(deepseek4): stage text on a hybrid placement only with the in-process expert path run_staged_text_chunk() drives the heterogeneous prefill through the in-process expert backend or expert runtime. A hybrid placement without either now keeps the gathered prefill instead of staging text prompts that would then fail. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * fix(deepseek4): take short staged text tails whole; document chunk failure - In the shared pass, a text tail below twice the layer-major minimum goes in whole when it fits the remaining budget instead of waiting for a later pass. - A heterogeneous staged chunk failure stays terminal for the request, now with the reason in the code. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> --------- Co-authored-by: mrciffa <davide@lucebox.com> Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…arget The portability guard used #if defined(__gfx1151__), but the HIP host pass defines no arch macro, so in every HIP build since then MMB's host dispatch compiled to the stub and QSA reported "unsupported". Guard the device code with defined(__gfx1151__) || !defined(__HIP_DEVICE_COMPILE__) and gate the host decision on the runtime cc (GGML_CUDA_CC_IS_RDNA3_5). Restores pre-regression behaviour on gfx1151 single-arch and fat builds: MMB live at T>=512 (routed MoE per first chunk 1474 -> 391 ms), QSA=1 runs, quality HE10/GSM10/Math9/recall2, QWEN4EXP_UPSTREAM clean with 0 expert-id mismatches.
Adds mmb-q8f16.cuh: a Q8_0 dense GEMM for MMB prefill (T>=512, gfx1151) that dequantizes weights while staging LDS into an F16 WMMA 256-row tile with whole-K F32 accumulation, ported from gufo's MIT design. Opt-in via LUCE_MMB_Q8F16=1; changes numerics vs the bf16 tile, so it stays default-off. UD-Q4_K_XL: dense MUL_MAT per d0 chunk 2299 -> 841 ms, 16,366 TTFT 25.69 -> 17.28 s (637 -> 947 tok/s), decode unchanged. Microbench 32.9-35.9 TFLOPS on qkv/gate/q/ ssm_out (gufo F16 route ~36-37). Quality HE10/GSM10/Math10/recall2 with it on.
…mmitted header f6b4dfe declared Qwen4ExpHeadOffload and an eight-argument qwen4exp_forward with a defaulted head parameter, but the implementation lives only in the unshipped layer-split work. The committed .cpp defines the seven-argument overload, so every seven-argument call was ambiguous and luce_common failed to compile on Linux, Windows and GB10.
mmb.cu includes mmb-small-m.cuh for the opt-in LUCE_MMB_SMALL_M bandwidth kernel, but the header was never committed, so both ROCm builds stopped at the missing include. This is the same file the gfx1151 measurements were built with.
… dense matmul The probe timed every dense route inline in ggml_cuda_mul_mat and forced routes through a global MMB tile override. It was an offline capture tool for building the frozen QWEN4EXP_DENSE_TABLE, which stays. With the flag unset the route and the tile override were always zero, so default dispatch is unchanged.
The trace wrote selected activations to /tmp from all three forward paths and duplicated QWEN4EXP_DUMP, which the upstream differential harness uses. Every trace call acted only when the flag was set, so the graphs are unchanged. The batched smoke test's trace-prefix and first-decode hooks go with it.
The MMB sources are built for HIP only, so a CUDA build has no ggml_cuda_mmb_release_all and test_server_unit failed to link on the DGX runner. Guard the call the way ggml_backend_cuda_free does. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Reverts 98b69d7. Upstream ggml numbers Q2_0 as type 42, and the ISTA-DASLab Qwen3.8-Flash-Next GSQ-RCO GGUFs store tensors with it (IQ3_S: 9 routed-expert tensors, Q2_0: 202). With TQ3_0 at 42 those tensors would load as the wrong type. TQ3_0 is a KV cache type that no GGUF stores, so it stays at 43 with the disk cache version at 3. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
# Conflicts: # server/CMakeLists.txt # server/docs/QWEN4EXP.md # server/src/common/backend_factory.cpp # server/src/common/model_capabilities.h # server/src/qwen4exp/qwen4exp_backend.cpp # server/src/qwen4exp/qwen4exp_backend.h # server/src/qwen4exp/qwen4exp_cache.cpp # server/src/qwen4exp/qwen4exp_cache.h # server/src/qwen4exp/qwen4exp_graph.cpp # server/src/qwen4exp/qwen4exp_graph.h # server/src/qwen4exp/qwen4exp_internal.h # server/src/qwen4exp/qwen4exp_loader.cpp # server/test/smoke/smoke_qwen4exp_forward.cpp # server/test/unit/test_feature_gate.cpp # server/test/unit/test_qwen4exp_indexer_score.cpp # server/test/unit/test_qwen4exp_qsa_ids.cpp # server/test/unit/test_server_unit.cpp
Merging main back in restored what Luce-Org#774 removed from the backend: the reference/upstream mode, the dump instrumentation, qsa_rebuild_reference, the float qsa_cell_ids fallback, docs/handoffs/ and scripts/qwen4exp_forward_diff.py. This keeps MTP (sidecar loading, the draft and verify graphs and workspaces, adaptive width, rollback, stable QSA decode reuse, fused QSA cell ids, server and feature-gate changes) on main's single path: - the MTP draft graph uses the default path; verify no longer checks a reference cache and requires at most one PLE layer (the rollback snapshot covers one PLE state; every published GGUF has one); - the loader is main's plus sidecar loading and the PLE image token id the draft vocabulary needs; - the smoke test's --stable compares the replayed stable QSA graph with a fresh graph each step; the MTP profile test restores to the generic mode. Units pass on the R9700 and the Strix Halo (test_qwen4exp_mtp: 4236 cases). UD-Q4_K_XL on the Strix Halo, MTP off -> on: code decode 22.3 -> 40.4 tok/s (accept 0.84), thinking 22.4 -> 34.8, 16K prose 20.1 -> 26.8; greedy outputs identical; 6/6 checks; prefill 5-13% lower (draft-layer K/V fill). DeepSeek V4.1 output unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
# Conflicts: # server/src/qwen4exp/qwen4exp_backend.cpp
dusterbloom
changed the title
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
4 participants
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.
Qwen4Exp decode on gfx1151 still has two avoidable serial costs in eligible
requests: reading the full vocabulary row to the CPU for greedy selection, and
running the routed and shared expert branches on one stream.
This PR adds two guarded paths:
graph. It is selected automatically for gfx1151. Verify rows, processed-logit
sampling, debug checks, invalid device results, and terminal-logit retention
keep the existing full-logit path.
joins it before the existing combine when the existing
LUCE_QWEN_SHARED_OVERLAP=1switch explicitly enables it. The closed48-layer gfx1151 stable-graph guard still applies. Ordinary one-token prefill
tails may use it; hidden-export, verify, and MTP-prefill forwards remain
serial.
The existing
LUCE_GPU_ARGMAX=0control disables automatic GPU argmax.LUCE_QWEN_SHARED_OVERLAP=1enables overlap; unset or0keeps the serialshared branch. The final default-policy change adds no further environment
names and keeps all rejected plans on the existing path.
The exact candidate is
9aec87f3a8e19df431bfbc8c135f9234eda5c349with tree
140d365ec9bb71b29e35fa4977ed4d0a322b21f9. Relative to the previousfresh-built
8b78284dcandidate, it changes only the shared-overlap selector inqwen4exp_graph.cpp; all HIP sources and the automatic per-instance gfx1151argmax selector are byte-identical. The implementation remains the reviewed
#810-relative 10-file delta.
Validation
Final exact-head qualification is in progress. Before this draft is marked
ready, this section will be replaced with the retained receipts for:
as an explicit opt-in;
No exact-head throughput, quality, or merge-readiness claim is made until those
receipts complete.