Repository navigation
[Bugfix] qwen4_exp: keep the PLE table BF16 under online quantization - #4
Closed
0xSandberg wants to merge 164 commits into
Closed
0xSandberg wants to merge 164 commits into
0xSandberg wants to merge 164 commits into
Conversation
Net change of the 318 commits from 31e2443 to 5fc3e2c310.
Upstream now indexes the align state copies by request slot: the pre-copy and post-copy kernels read row idx_mapping[i] of the per-request-slot block tables instead of batch row i. The control-path test passed its one-row batch table to precopy_mamba_align_fused_kernel with the request at slot 2, so the kernel read past the table and the Triton-interpreted child crashed. The test now places the request row at its state slot of a table with one row per slot, the other slots holding the null block.
Upstream moved the align state copies to the per-request-slot block tables, set up once in initialize_kv_cache, because under PP the batch-order tables are re-gathered by later steps before a step post-copies. The VLLM_STATE_INDEX_CHECK hooks still read the batch-order tables by batch row and named mamba_group_ids in preprocess_state, where it is no longer defined. initialize_kv_cache now captures the per-request-slot tables, the number of blocks and the mamba block size when the check is on, and check_align_copies reads row idx_mapping[i] of them for batch row i. _full_table, which recovered the persistent batch-order table, is removed.
Upstream (2274e86, DFlash2 with GLM-5.3-Flash) made Glm5NextForCausalLM and Glm5NextForConditionalGeneration SupportsEagle3 and Glm5NextModel an EagleModelMixin with its own aux hidden state capture (_aux_hidden_state, backed by a new _aux_post_op MHCPostOp). 2a54f6b and 9f3d70e added ROCm-only fused shared experts to Glm5NextModel.__init__. The fork classes no longer list these interfaces again; they inherit them. The fork backbone keeps its own forward: the upstream capture takes the same positions (state entering layer i, plus end_layer) but has no pipeline boundary, ignores aux_hidden_state_tensor and contracts by config.mhc_num_residual_streams, which the GLM-5.3-Flash text config does not carry; the fork contracts by hc_mult. adopt still receives an upstream-built Glm5NextModel; the extra _aux_post_op child has no parameters and is unused by the fork forward. The fused shared experts are only enabled on ROCm gfx950. The registration test asserted that the upstream classes do not support EAGLE3. It now checks what that guarded: the registered classes are EAGLE3 targets whose backbone carries aux states across pipeline stages. Seams glm-model-init, glm-model-forward, glm-upstream-causal-lm-class and glm-upstream-conditional-class are refreshed with their notes updated.
Each pinned object changed upstream; the fork code relying on it is unchanged: - attn-register-backend: a4b14e5 lets register_backend take a string name and creates a dynamic enum member for it. The fork still passes AttentionBackendEnum.CUSTOM with a class path, which takes the old path. - comm-cuda-communicator-init: a6f43e5 passes register_graph_buffers to the ROCm AITER all-reduce constructors. The attribute names the fork communicator reads, the CustomAllreduce it swaps and the final _log_all_reduce_backend_selection call are unchanged. - config-vllm-config-post-init: 194bd72 removed the sequence parallelism / async TP handling and a6f43e5 touched the cuMem pool check. The plugin guard is still called first. - glm-gdn-metadata-class, glm-gdn-builder-init, glm-gdn-builder-build: 5485973 adds the ROCm aiter_flydsl GDN prefill backend: a defaulted aiter_prefill_metadata field, a _check_chunk_metadata_override call in __init__, and a build branch that skips _build_chunk_metadata. The metadata is still a plain dataclass built fresh per build (the GLM subclass inherits the new defaulted field), use_spec_decode and supports_update_block_table are still set in __init__, and the fork builder does not override _build_chunk_metadata, so the check passes. - glm-mamba-state-prepare-attn: 8a26869 creates the Mamba align context in initialize_kv_cache (_ensure_align_ctx is gone); prepare_attn now asserts it exists. The RecoverSSM step is still recorded through self.recoverssm. - mla-update-kv-cache: 4ed4c44 drops the HiSparse prefill mirror write. The early return and the do_kv_cache_update of the gathered inputs that the GLM override defers are unchanged; HiSparse never defers. - moe-factory: 9f3d70e adds shared_expert_prefix. MoERunner and RoutedExperts are still the defaults when no class is passed.
Upstream's Mamba prefill split now places the prefill checkpoint on the scheduler block size (Scheduler.block_size) when fine-grained hits are off. The stub that stands in for the scheduler in the DFlash boundary tests carried only the cache config's block size, so every split raised AttributeError. It now carries the scheduler block size too.
Upstream reworked the hybrid Mamba prefix-cache boundaries. Without the EAGLE block drop a prefill checkpoint now leaves the last token to recompute, so a block-aligned prompt keeps its checkpoint one block below its end and the scheduler stops the prefill there. EAGLE lookups read the whole prompt before dropping a block. A block-aligned prompt checkpoint is also registered again through the partial tail path, which marks its Mamba block as written in that step. The DFlash boundary tests asserted the old rule. At an exact prompt edge both modes now reuse the block below the edge. A cold prompt one token short of the next block serves that resend only in the drop mode, whose sparse checkpoint sits one block lower. A block-aligned prompt yields its first token after two chunks. The warm request in the cache-hit test is admitted in a later step, as the scheduler admits it. Past the edge the boundary mode still reuses the prompt's last full block, one more than the drop mode.
Upstream now seeds an admitted request's Mamba state column with cache_config.mamba_block_size rather than cache_config.block_size, which a drafter group with a smaller block pulls below the Mamba block. The test asserted the old source text. It now builds the model state, raises the Mamba block size after construction as the late alignment does while block_size stays smaller, and checks the column seeded for a hit at 11520 tokens.
The fork now sits on upstream efd1410 instead of 31e2443, so the fork map compares with that base. Upstream renamed the autoregressive speculator to target_dependent_ar, which moves that patch in the list; the scheduler and worker patches shrank where upstream now covers what they did (prefill checkpoint tail shift, sequence-parallel gather); the GLM model, registration, state index and DFlash boundary test lines follow the adaptations to the new base.
Under RecoverSSM (TP4) every step replays the accepted tokens over each request's KDA state after sampling: one more read and write of every state per layer. With VLLM_AMPERE_KDA_FOLD=1 that replay moves into the next verify, which loads the same state anyway: - the commit plan splits the rows; a row whose state nobody reads before its next verify (align mode: no block boundary crossed or reached, the commit stays in the block the verify read, the next step cannot spill into the next block; "none" mode: every row) commits nothing and records the length it owes and its state block per KDA cache group and request slot; the other rows commit exactly as before; - records are kept per request slot; the verify (kda_decode_v2 FOLD) replays the owed records with the commit's fma(c, k, h*e) sequence before its gate arrival, stores the state back and reads the conv window at the owed position, so every state, window and checkpoint a consumer reads is the same bytes; - the fold's buffers are static per layer, on the layer's records (set by the GDN builder's RecoverSSM state before capture); - the GLM model state (Glm5NextModelState) holds the fold's hooks (None with the fold off) and calls them after the hybrid state's own work: per step after the metadata is built (verify rows to request slots and owed lengths), on add_request (a slot taken owes nothing) and on remove_request. A streaming-input session re-added in the step it was removed in keeps its state, so its owed commit is written before its next chunk reads it, with the commit's own kernels and plan values; a preempted request is never written, its blocks may already belong to another request. Off by default: the switch off adds no call and the off variant of the verify compiles to the same PTX. No effect where RecoverSSM is off, under data parallelism or under PP. Tests: tests/kernels/test_ampere_kda_fold.py (CPU: the plan against the upstream plan, the gather, multi-step byte identity of what consumers read through the real commit and pre-copy path incl. slot reuse with a spec-row first step, streaming-input re-adds and a preempted request whose blocks are reused, deliberately broken folds fail, host cost, the off path, CUDA-graph buffers, the real GLM-5.3-Flash TP4 config's metadata; GPU part prepared).
Scheduler.reset_prefix_cache(reset_running_requests=True) preempts and resumes running requests in the same step: the worker removes such a request in finish_requests, removes it again in add_requests and re-adds it on blocks that were just freed. The fold's remove hook recorded the first removal as if the request kept its state, so a re-add with computed tokens kept (a KV connector hit) would have written the old owed commit into a freed block. A second removal of the same id in one step now drops the record; only a request removed once and re-added in the same step (a streaming-input session) has its owed commit written. Tests: the streaming scenario also preempts and resumes a request in one step, and catches the old hook (8 differences); it is in the script-mode CPU test list now.
The CUDA-graph test of the KDA commit fold ran one eager step, kept its results, then restored the state, conv window and owed lengths before each replay, but not the record pool. The FOLD verify replays a row's owed tokens from its records and writes the step's records back to the pool, so every replay started from records the previous run had written and could not match the eager run. The test now restores the pool too.
The CUDA-graph test of the KDA commit fold started from an all-zero record pool, so the owed-token replay under capture multiplied the state by zero decay gates and never did real arithmetic. The pool now holds finite records (small corrections, unit keys, decay gates in (0.5, 1)). The verify test covers nseq 7 too: 7 and 8 are the sequence counts whose block config differs from the rest.
thin_gemm_group runs two or three projections that read the same input in one launch. Each member keeps the schedule its solo thin_gemm call would use (BLOCK_N, BLOCK_K, SPLIT_K, member-local rotation start), its own split-K partials and arrival counters, and its own output pointer and strides, so every output is bitwise the separate calls. Split-K members are scheduled first so their reduction tail overlaps the unsplit member. thin_gemm_group writes members into columns of one [M, sum N] output; thin_gemm_group_split gives each member its own contiguous tensor. warmup() takes the groups to compile and allocate. The solo kernel, thin_gemm and the workspace helpers are unchanged.
VLLM_THIN_GEMM_CONCAT (default off) runs q_b_proj + indexer.wq_b and fused_qkv_a_proj + the indexer kpool gate as one thin GEMM launch each. Warmup registers the groups by member weight and compiles them. Member 0 launches the group and stashes member 1 output keyed on its input; member 1 takes the stash, or runs alone and logs once. Outputs are bitwise the separate calls. The gate needs VLLM_THIN_GEMM, sm_80 and compilation mode NONE and logs why when it stays closed. With the flag unset the pairing module is not imported and the package exports its own ampere_thin_gemm and thin_linear.
…M logger The served group checks now cover the TP4 capture sizes 5, 6, 7 and 10 as well, and one test drives the concat banner and a gate-closed line through the info_once and warning_once of the vLLM logger.
With VLLM_AMPERE_KDA_SMALL_STACK=1, every RecoverSSM verify at nseq 7-8 (16 heads), the commit fold's included, runs as one 8-row warp per CTA (four times the CTAs) instead of one 32-row warp. The 32-row warp spills 2,184 B of stack per thread, 2,624 B in the fold's variant; the 8-row warp needs at most 72 B on the largest plans (sm_80 compile). It keeps the 32-row warp's layouts, so every reduction runs in the same order and the results are bitwise the same. The driver reserves the largest stack launched (at least its default 1,024 B) for every thread an SM can hold, and the first verify launch falls in the CUDA-graph memory profile, so that reservation comes out of the KV cache. Measured on one card, today's verify raises it by 168 MiB and the fold's by a further 62 MiB. On the same card the 8-row config is bitwise on every case and faster at nseq 7-8: 15-37% less kernel time for today's verify, 21-47% for the fold's (the least at T 1). Off by default, with a banner where a verify config changes (max_num_seqs 7 or more at 16 heads) and "set but off" lines; read at the first launch; no effect where RecoverSSM is off or the fused KDA decode is not used. Tests: tests/kernels/test_ampere_kda_small_stack.py (CPU: the switch, the configs it changes and keeps for today's verify and the fold's, banner and set-but-off lines, the stack compiled for sm_80; GPU: bitwise against the switch off for today's verify and the fold's, nseq 1..8, T 1..8).
The concat groups bound impl() at import, so a re-resolved or patched impl was bypassed with VLLM_THIN_GEMM_CONCAT set. Call it, and use_ampere_thin_gemm, through the package as the base dispatch does.
Upstream now indexes the align state copies by request slot: the pre-copy kernel reads row idx_mapping[i] of the per-request-slot block tables instead of batch row i. The fold harness passed its batch-order tables, so the kernel read rows past the table for any request whose slot exceeded the batch size and the Triton-interpreted child crashed. The harness now builds, per group, a table with one row per request slot (every live request's blocks, free slots NULL) for the pre-copy, as the model runner keeps them; the attention metadata keeps the batch-order tables.
Upstream 3d6853b seeds the align-mode state column with cache_config.mamba_block_size instead of cache_config.block_size, which a smaller drafter group can pull below the Mamba block. add_request still only takes the request slot (acceptance count, state column) and calls nothing that reads the KDA commit fold, so the Glm5NextModelState hook after it stays correct.
Adds the KDA commit fold, the thin GEMM concat groups and the small-stack RecoverSSM verify config, cherry-picked onto upstream efd1410.
With VLLM_AMPERE_KDA_SMALL_STACK=1 the RecoverSSM verify at nseq 7-8 (16 heads) takes BV 16 over one warp instead of BV 8. The warp keeps the 32-row layouts, so outputs, records and states are bitwise those of BV 8. Compiled for sm_80 the stack stays at most 872 B per thread (the commit fold's verify), within the 1,024 B default. Verify at 8 requests, T 4: 53.5 -> 39.9 us per call on a CMP 170HX. Tests pin _select for nseq 1..8 with and without small_stack, check the 16-row config's stack at every token tile (skipped without the sm_80 toolchain) and, on a GPU, compare 16 rows with 8 rows bitwise at nseq 7-8, T 1..8, today's verify and the fold's.
The nightly advance replayed the stack onto upstream 81c9936; the fork map now names that base. No other line changes.
The test builds the model state from a bare config; with the fold switch set in the environment, the real fold check reached that config's parallel settings and failed. The fold is not under test there.
…DA_RECOVER_OFF) With VLLM_AMPERE_KDA_RECOVER_OFF=1, kda_recover_enabled() returns False where its gate holds, so the KDA layers, metadata builders, model state, warmup and commit fold take the path a closed gate takes: one KDA state per draft position (draft-position state pages, the per-position verify), no records, commit or fold. It gives a reference without state recovery from the same tree. Never for serving. Off by default; unset or 0 runs exactly what ran before. A warning banner names the switch when it turns RecoverSSM off, and a line says when it is set where RecoverSSM is off anyway.
The fold tests expect RecoverSSM on; with VLLM_AMPERE_KDA_RECOVER_OFF=1 exported in the shell they failed. kda_fold_world.env and the real-config metadata test now set the switch to 0.
A second revision of the compiled MoE decode GEMM for the tensor-parallel N=512 shard: one bf16x2 FMA per two weights for the exact (q-8)*s dequant, the first work item per warp without an atomic, 4-row chunks for 2 to 7 rows and zero activation fragments for padding tokens at 32 rows. Item split, partial planes and per-output order are those of decode_gemm_orig, which is unchanged. The fp32 sums are carried at 2^-16 scale, so the output equals decode_gemm_orig bit for bit for weight scales 0 or in [2^-110, 2^10) as long as no nonzero weight-activation product falls below 2^-110.
With VLLM_AMPERE_MARLIN_DECODE_CUDA=1, the switch replaces decode_gemm_orig and its per-row-count configs by decode_gemm_v2 on the N=512 shard. The op is optional; without it, at other widths, or with a weight scale outside the exact range (checked once per layer outside capture), decode_gemm_orig runs and one line says why, as it does when the compiled decode itself is off. Warmup launches every variant before CUDA-graph capture. Default off.
…ODE_V2 With VLLM_AMPERE_SPARSE_MLA_DECODE_V2=1, triton_mla_sparse_fwd sends bf16 decode calls in the 512-wide NoPE layout with 1..64 query rows, 16 or 64 heads, 2048 or 2176 index slots and contiguous last dimensions on sm_80 to triton_mla_sparse_decode_v2.py. Each CTA reads its index row once and walks only the tiles between the first and last valid entry; with four or more splits a first pass finds each sparse row's exact logit max and a second pass spreads its valid tiles over every split; partials are combined by a parallel merge. At 16-head blocks of 64-key tiles the walk runs as a Gluon kernel with a double-buffered cp.async gather. Every other call, and every call with the flag unset, keeps the shipped kernel unchanged. Off by default.
…x helper The refusal test sliced one element off a buffer of exactly the trellis size, so the view could not hold the matrix. The split-K helper already builds a 2-byte aligned trellis with one spare element.
The Gluon thin GEMM's kernels are compiled in the pre-capture warmup for contiguous operands. Triton specializes each launch on whether a pointer is 16-byte aligned and whether an integer argument is 1 or a multiple of 16. A table call whose x has another layout therefore launched a new variant at serve time. For a row stride that is 8 mod 16 elements, that variant does not compile at all: the cp.async operand loads become 2-byte transfers, which cp.async does not support. The dispatch now sends a table call to the Gluon module only when x has a unit K stride, a row stride that is a multiple of 16 elements and a 16-byte aligned base, and the weight a row stride that is a multiple of 16 and an aligned base. Every other table call takes F.linear (cuBLAS), with one warning naming the strides. Every call site of the seven shapes passes a freshly allocated or row-sliced contiguous activation, so served calls keep the Gluon path. The kernel module is unchanged. Tests: the layout check on admitted and refused views of x and of the weight, on the CPU. On the GPU: admitted views give the contiguous call's result bit for bit, and refused views give F.linear's through the dispatch. These replace the awkward-input test, which fed the module directly with views the dispatch no longer passes.
…the extension loader ops.ensure_extension calls dense_prefill_tc.resolve before it marks the library loaded, and with VLLM_EXL3_DENSE_PREFILL_TC set and the gate open resolve warms the operator up. The warm-up sized its workspace with ops.linear_workspace_size, which calls ensure_extension again; the library was not marked loaded yet, so the first EXL3 use of such a process recursed until RecursionError. The warm-up now reads the size from torch.ops._exl3_C.linear_workspace_size. ensure_extension has imported the library and checked that it has linear_workspace_size before resolve runs. Marking the library loaded before resolve instead would leave a process whose resolve raised looking loaded, with the gate half set up; a re-entry guard would add module state for one call. A CPU test runs ensure_extension through the real resolve and warm-up with the gate open (tensor parallel size 2) and a stand-in library; it fails with RecursionError on the old warm-up. A GPU test does the same with the built library on the card.
… the Gluon thin GEMM
The Gluon thin GEMM's graph test read the allocator's allocation counter before torch.cuda.graph() and after it. Since torch 2.13, capture_begin allocates the default CUDA generator's RNG seed and offset tensors for every new capture, so the count rose by two on a capture that holds only the GEMM calls. The test now reads the counter inside the capture, around the calls. The 74-SM thin GEMM's graph test read the allocated and reserved bytes after torch.equal had compared the replayed outputs. torch.equal allocates its eq and all() temporaries, and whether those take a new 2 MiB segment depends on what earlier tests left in the pool. The test now reads the bytes right after the replays. Both tests still require zero allocations from the GEMM calls in capture and zero growth across replays.
… compile nothing new thin_gemm sent x and w rows with a stride of 8 mod 16 elements to the Gluon kernels, which compile only for strides Triton specializes as multiples of 16; such a direct call failed to compile instead of taking the Triton kernel its comment names. The host check now asks for strides that are multiples of 16 for both operands. A GPU test calls thin_gemm with a row stride of K + 24 and checks that it runs the Triton kernel: no Gluon launch, the same bits as an unaligned view, within the FP64 gate. The admitted-layout GPU test now also checks the property the layout gate exists for: after the contiguous call, an admitted view adds no compiled variant to any of the module's JIT kernels. A view with a 16-byte column offset, the tightest base the gate admits, joins the admitted layouts.
…summary The module summary now names the 16-byte base offset among the admitted views, the check that admitted views compile no variant beyond the contiguous call's, and the direct call with a row stride of K + 24 that takes the Triton kernel.
…ls for captured calls Launched outside a CUDA graph, a Gluon thin GEMM call pays its Triton launch on the host every time: on the 74-SM part the drafter's eager decode passes took 1.7-3.2x cuBLAS's time per call at 1-8 rows. With the gate open, the dispatch now asks the current stream's capture status first (the binding behind torch.cuda.is_current_stream_capturing, called directly) and sends every call made outside capture to F.linear, the call it makes with both flags off. Calls captured into a graph keep the previous routing. The status holds for breakable PIECEWISE graphs, for FULL graphs, whose forward runs with runtime mode NONE inside torch.cuda.graph, and for side streams joined to a capture. CPU tests: with the gate open, eager calls take F.linear before any other test, for any shape, row count, bias or layout; a closed gate never asks for the capture status; the capture test is torch's binding; the routing and layout tests run as captured calls. GPU tests: the end-to-end dispatch test captures its calls into a graph (the Gluon kernels' bits) and checks that eager calls give F.linear's bits without a launch; the real-input test captures every case and checks the eager call against F.linear.
…be the eager route No eager forward runs the Gluon kernels before a capture any more (the warm-up forward takes F.linear), so a call captured at a row count the warmup was not given would compile a variant and allocate its split-K workspace inside the capture. The warmup now readies every M from 1 to 32 for the seven shapes, whatever capture sizes it is given. The flag's description now says that once it is active every call outside capture, of any shape, takes F.linear, the Triton calls of VLLM_THIN_GEMM included, and the banner names the captured-only route. Tests: a GPU test warms with the graph sizes 1-32 of the served list and captures at 3, 5, 12, 20 and 28 rows: no new compiled variant, no allocation but the output inside the capture, replay bitwise the direct call. The eager CPU test also makes the row-count and table tests raise, so the capture test has to come before them. The real-input test counts kernel launches (none for an eager call, one per capture) and records, instead of raising, a call that leaves the table path inside a capture.
…interpreter The test for warming every row count up to 32 ran after the FP64 test, which had already compiled every variant and allocated every workspace buffer at M 1..32 in the same process, so a warmup that skipped the unlisted row counts would still have passed it. Its body now runs in a fresh interpreter, as the binding tests do: there, only the warmup can have readied the kernels and buffers that the captures at 3, 5, 12, 20 and 28 rows need. A warmup restricted to the given sizes makes 5 of those 35 calls compile a new variant and 15 allocate a split-K buffer inside the capture.
ensure_extension registered the five base fake implementations on every call that found the library not yet loaded. When dense_prefill_tc.resolve raised (for example in its warm-up), the library stayed unloaded, as intended, but the next call registered the same fakes again; torch refuses a second fake implementation, so the retry failed with that error and never reached resolve. The base fakes are now registered once per process, before the first resolve. A CPU test makes resolve raise once with a stand-in library and checks that the next ensure_extension reaches resolve and loads.
Three tests stand in for torch.library.register_fake and call the real ensure_extension. They restored the loaded flag but not the flag that records the base fake registration, so a later real load in the same process skipped the fakes it never registered. Each test now patches that flag too.
VLLM_EXL3_PLE_NATIVE_READ (default off) replaces the Python row loop on the eight-thread pool with one host operator, _exl3_C::ple_read_rows, called through torch.ops so the GIL is released for the whole read. Rows already in the page cache are read on the calling thread with preadv2(RWF_NOWAIT); the others are read on up to eight threads so their storage latencies overlap. Ownership, zero-filled unowned and padding rows, retries with the same backoff, the failed rows and their errors are those of Exl3PLERowReader.read_rows, so the staging buffer, the poisoned rows and every kernel after them are unchanged. The pool path stays as it is and is used whenever the switch is off; the switch can change between steps. A startup line names the native reader when it is on and the placement is ssd.
Reserve the per-thread failure lists and the thread list before the first thread starts, so nothing can throw while a started thread is still joinable. State that the non-blocking probe also starts the read of a row that is not cached, one read attempt more than the Python reader makes for that row.
…ny reason The native PLE reader caught only std::system_error from the std::thread constructor; a std::bad_alloc while earlier workers were joinable ended the process. Catch any std::exception and read the chunk on the calling thread. The texts on the native path also said its failures equal the pool path's. Its non-blocking probe is one extra read attempt for rows not in the page cache, so its failures are a subset: say so.
…r retry and the native PLE reader
…t a stream query Asking the current stream whether it is capturing costs ~0.3 us of C++ and driver work per call on the 74-SM part, and every eager call of the dispatch paid it before taking F.linear; on back-to-back eager calls that left the dispatch 3-7 % behind the plain F.linear call. When the gate opens, the dispatch now wraps torch.cuda.CUDAGraph.capture_begin and capture_end once and counts the captures under way; the captured-call test is a read of that count. Every capture in this tree goes through those two methods: torch.cuda.graph for FULL graphs and the segments of the breakable capture for PIECEWISE graphs. A side stream joined to a capture counts as captured, and the eager segments of a breakable graph, which run between two segments and are replayed outside any capture, read zero. Tests: the count on a stand-in graph class (nested captures, arguments passed through, a single wrap, a begin that raises counts nothing, an end that raises still counts down, an end without a counted begin does not), the default class being torch.cuda.CUDAGraph, the warmup keeping the count from before the gate opens and not when it stays closed; on the GPU, the count against the stream's own capture status through a torch.cuda.graph capture, a joined side stream, a breakable capture with an eager break at capture and at replay, and after a capture whose body raised.
…stream PleHostReader reads the rows of a step with the native reader core inside a host function queued on the current stream, poisons the rows that fail and reports results through take(). The native reader core now takes preallocated scratch, so the host function allocates nothing except the threads it starts for rows not in the page cache. Not yet called from Python.
VLLM_EXL3_PLE_HOSTFN_READ=1 queues the row read of a step as a host function (PleHostReader) on the PLE side stream, between the copy of the ids to the host and the copy of the rows to the device. The prefetch no longer waits on the host for the ids or for the previous output copy: the slot queues that wait on the side stream instead. Rows that fail are poisoned in the staging buffer, so their tokens fail in the same step through the existing mask; the error line is logged when the next step takes the finished reads. Up to four reads may be queued or not yet taken; a prefetch that finds all four taken waits for the side stream once. Close drains the stream, logs the last reads and drops the reader. Off by default; with it off the step is unchanged.
NOTICE records that this is a fork of vLLM under the Apache License 2.0 and that the sm80 support was developed by Morrowmake; CITATION.cff gives the citation metadata.
Describe the vLLM 0.32.0 dev base, both supported models and their recipes, the current switch naming and sm80 plugin, the per-release fork tags, and the NOTICE and citation metadata.
Qwen4ExpPLEEmbeddingMethod.from_quant_config raised NotImplementedError for OnlineQuantizationConfig, so any --quantization online shorthand (int8_per_channel_weight_only, mxfp4, fp8_per_block, ...) failed to load a BF16 Qwen3.8-Flash-Next / qwen4_exp checkpoint. Online quantization converts only Linear and MoE weights at load time, so the PLE table stays BF16, the same as the existing QuarkConfig branch. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011MGSr6DMigthpK1b8MCxc3 Signed-off-by: 0xSandberg <HomeAI7950@proton.me>
6 tasks done
Owner
|
Thanks for the clear report and fix. Upstream vLLM has since rewritten this function (vllm-project#61454): any quantization config it doesn't handle, |
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.
Overview
Online quantization (
--quantization int8_per_channel_weight_only,mxfp4,fp8_per_block, ...) can't load a BF16qwen4_expcheckpoint (Qwen3.8-Flash-Next and finetunes such as Blackfrost-AI/CYBER-FROST-3.8-BF16):Qwen4ExpPLEEmbeddingMethod.from_quant_configraisesNotImplementedError: Qwen4Exp PLE embedding does not support quantization config OnlineQuantizationConfig. This adds anOnlineQuantizationConfigbranch that keeps the PLE table BF16, the same way the existingQuarkConfigbranch does.Claims
--quantization int8_per_channel_weight_onlynow loads and serves CYBER-FROST-3.8-BF16 (qwen4_exp) at TP8 on 8x CMP 170HX; before this change every online-quant launch failed at model construction.Validation
Rig: 8x CMP 170HX 64 GB (sm_80, 74 SMs), PCIe Gen2 x4 links, 180 W cap, driver 610.57.04, image
ghcr.io/morrowmake/vllm-cmp170hx@sha256:5b45c4d6…(1830594). Theamperecopy ofngram_embedding.pyat f02c7a7 is byte-identical to the image's, so the patched file was bind-mounted over it.NotImplementedError ... OnlineQuantizationConfigfromngram_embedding.py, engine core init fails.Using TRITON Int8 MoE backend), correct answers on a smoke prompt, and single-user decode (512 output tokens, median of 5) of 65.1 / 58.1 tok/s (reasoning / code prompts) with MTP 2, vs 57.8 / 53.5 for the BF16 baseline on the same rig. An earlier run with a functionally identical local patch measured 61.8 / 54.7.vllm.model_executor.layers.quantization.online.baseintroduces no import cycle;from_quant_config(OnlineQuantizationConfig(int8_per_channel_weight_only), prefix)returnsQwen4ExpPLEUnquantizedEmbeddingMethod.Details
from_quant_confighandles ModelOpt, CompressedTensors, INC and Quark, then requiresFp8Config; the online-quantization config path added later was never routed here.vllm/models/qwen4_exp/common/ngram_embedding.py. No behaviour change for checkpoint-quantized models.mainhas the same code in this file, so the same change applies there.mxfp4also tries to quantize the vision tower and, at TP8, the shared expert (widths not divisible by 32); worked around with--language-model-onlyand--quantization-config '{"ignore":["*shared_expert*"]}';ModelConfigcopies--quantizationbut not--quantization-config, so the ignore list is lost for the draft (vllm/config/speculative.py).Pull Request Checklist
I used vLLM's
/pr-checklistskill. (Mandatory for agents, optional for humans).AI assistance was used during the creation of this PR.
Design Fit: Minimizes impact on core components, reuses existing functionality, and justifies added complexity.
Testing and Validation: Validates the change and ensures any added tests are meaningful and reliable, with CI coverage or documented CI resource constraints and validation performed outside CI.
Code Quality and Style: Keeps code and comments clear and concise, and updates relevant documentation and examples.
Pull Request Contents: Includes a brief summary and relevant links, supports claims with evidence, explains root causes and implementation trade-offs, and follows the contributing guide.
🤖 Generated with Claude Code
https://claude.ai/code/session_011MGSr6DMigthpK1b8MCxc3