Repository navigation
fix(linear): repair B12X kernel registry after the b12x_block rebase - #50
Merged
randomvariable merged 3 commits intoAug 19, 2026
Merged
Conversation
The nvidia DeepseekV4DecoderLayer exposes hc_pre and hc_post_pre but no hc_post wrapper: the model calls mhc_post_tilelang directly from DeepseekV4Model.forward. The warmup sweep therefore compiled the pre and fused-post-pre kernels but left mhc_post to JIT on the first request that reached the end of the layer stack. Call mhc_post_tilelang alongside hc_post_pre in the sweep. Its token dimension is T.dynamic and it takes no split-k factor, so any shape warms every shape; it rides the existing loop rather than a separate pass. Add token-size selection tests pinning the invariant that the sweep covers both split-k families: mhc_pre/mhc_fused_post_pre derive split-k from hc_mult * hidden_size while mhc_pre_broadcast uses hidden_size alone, and the smaller K only merges plateaus, so the broadcast transitions are a subset of the swept ones. Adapted from vllm-project#52941. The PR's primary change - detecting nvidia layers that lack hc_pre/hc_post and calling tilelang directly - does not apply: this fork's nvidia layer does define those wrappers, so _find_first_mhc_layer already matches and the sweep already runs. Co-authored-by: OMP Agent <noreply@omp.local> Signed-off-by: Naadir Jeewa <naadir@randomvariable.co.uk>
Upstream 4ab5e50 (vllm-project#52368) consolidated scaled_mm/b12x_block.py and scaled_mm/b12x_tensor.py into scaled_mm/b12x.py. The rebase took the new module and the deletion of b12x_block.py, but left the old import paths and several duplicated registrations behind. scaled_mm/__init__.py still imported from the deleted b12x_block module, so importing vllm.model_executor.kernels.linear.scaled_mm raised ModuleNotFoundError. That aborted collection of every test touching the package, including tests/v1/attention/test_flashinfer_sparse_mla_sm120_api.py. Point it at b12x, matching linear/__init__.py, which was already correct. Three priority lists carried the same kernel twice, because the fork registered it before upstream did and the rebase kept both entries: _POSSIBLE_FP8_KERNELS MarlinFP8ScaledMMLinearKernel _POSSIBLE_NVFP4_KERNELS B12xNvFp4LinearKernel _POSSIBLE_MXFP4_KERNELS B12xMxFp4LinearKernel The duplicates sat at the head of each list, so they silently took priority: Marlin ahead of FlashInfer and Cutlass for FP8, and B12X ahead of every FlashInfer path for FP4. Drop them and keep upstream's positions, which the registration-priority tests already pin. Selection by explicit --linear-backend resolves through _LINEAR_BACKEND_KERNEL_MAP rather than list order, so the qualified B12X serving path is unaffected. _POSSIBLE_MXFP6_KERNELS was defined twice, identically; mypy rejects the redefinition as soon as the file is touched. mxfp8/b12x.py referenced b12x_warmup_token_counts, Iterable and Any without importing any of them, so warmup_b12x_mxfp8_linear raised NameError when it found a layer to warm. tests/model_executor/kernels/test_b12x_linear.py imported from b12x_block and from b12x_tensor; the latter's two symbols were unused, and one of them (_b12x_tensor_fp8_linear) no longer exists anywhere. Import from the package facade, as upstream's copy of this file does. tests/model_executor/kernels/test_b12x_linear.py: 33 passed (was: no collection). tests/v1/attention/test_flashinfer_sparse_mla_sm120_api.py: 5 passed (was: no collection). Co-authored-by: OMP Agent <noreply@omp.local> Signed-off-by: Naadir Jeewa <naadir@randomvariable.co.uk>
Upstream 4ab5e50 (vllm-project#52368) replaced the standalone warmup_b12x_{mxfp4,nvfp4}_linear functions with a get_b12x_warmup_unit provider that b12x_warmup discovers through layer.b12x_warmup_provider, and dropped the b12x_{mxfp4,nvfp4}_linear marker attributes. The rebase took the new kernels but left these two test files importing the deleted functions, so neither collected. Assert the contract where it now lives: two layers with the same weight signature produce equal unit keys and a differing one does not, and compiling a unit runs the kernel once per token count. The per-call weight, alpha and bias assertions carry over unchanged, so the same behaviour is still pinned - only the entry point moved. The marker assertions become provider-identity assertions. tests/model_executor/kernels/: 49 passed (was: 2 files uncollectable). Co-authored-by: OMP Agent <noreply@omp.local> Signed-off-by: Naadir Jeewa <naadir@randomvariable.co.uk>
This was referenced Aug 19, 2026
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.
Purpose
import vllm.model_executor.kernels.linear.scaled_mmraised:which aborted collection of every test touching the package — including
tests/v1/attention/test_flashinfer_sparse_mla_sm120_api.py, whose failure had been misread as an environment limitation.Root cause: upstream
4ab5e5012a(#52368) consolidatedscaled_mm/b12x_block.pyandscaled_mm/b12x_tensor.pyintoscaled_mm/b12x.py, and replaced the standalonewarmup_b12x_*_linearfunctions with aget_b12x_warmup_unitprovider protocol. The rebase took the new module and the file deletion but left the old import paths and several duplicated registrations behind.Pulling that thread surfaced five further defects, all from the same half-applied refactor.
Defects fixed
1. Dead import path (the reported P0).
scaled_mm/__init__.pyimported from the deletedb12x_block. Pointed atb12x, matchinglinear/__init__.py, which was already correct.2. Duplicate registrations silently reordering kernel priority. Three lists carried the same kernel twice, because the fork registered it before upstream did and the rebase kept both entries — with the duplicate at the head of the list, so it took priority:
_POSSIBLE_FP8_KERNELSMarlinFP8ScaledMMLinearKernel_POSSIBLE_NVFP4_KERNELSB12xNvFp4LinearKernel_POSSIBLE_MXFP4_KERNELSB12xMxFp4LinearKernelBlame confirms the provenance: the head entries are fork commits
e080b6ed0e(2026-07-21) and the pre-vllm-project#52182 Marlin position; the lower entries are upstream3c79b1a8bf(vllm-project#52016) andbca7bea240(vllm-project#52182). Upstream's positions are the maintained ones and the registration-priority tests already pin them, so the head duplicates are dropped.Selection via explicit
--linear-backendresolves through_LINEAR_BACKEND_KERNEL_MAP, not list order — verified by the test that monkeypatches_get_linear_backend— so the qualified B12X serving path is unaffected.3.
_POSSIBLE_MXFP6_KERNELSdefined twice, identically. mypy rejects the redefinition as soon as the file is touched.4.
mxfp8/b12x.pymissing three imports. It referencedb12x_warmup_token_counts,IterableandAnywithout importing any of them, sowarmup_b12x_mxfp8_linearraisedNameErroras soon as it found a layer to warm.5. Three test files importing deleted symbols.
test_b12x_linear.pyimported from bothb12x_blockandb12x_tensor; the latter's two symbols were unused and one (_b12x_tensor_fp8_linear) no longer exists anywhere.test_b12x_{mxfp4,nvfp4}_linear.pyimported the deleted standalone warmup functions.Test migration
The FP4 warmup tests moved onto the provider protocol rather than being deleted: the dedupe contract now lives in
unit.key, so they assert that two layers with the same weight signature produce equal keys and a differing one does not, then compile the distinct units. The per-call weight/alpha/bias assertions carry over unchanged, so the same behaviour is pinned — only the entry point moved. Marker assertions (layer.b12x_mxfp4_linear) became provider-identity assertions, matching whatprocess_weights_after_loadingnow sets.Test Plan
Duplicate check:
Test Result
tests/model_executor/kernels/: 49 passed — previously 3 of 4 files could not be collectedtests/v1/attention/test_flashinfer_sparse_mla_sm120_api.py: 5 passed — previously uncollectablepre-commit: exit 0 (ruff, ruff-format, mypy 3.10)Known remaining drift
Not fixed here, deliberately, because resolving it needs on-device evidence rather than a guess:
scaled_mm/b12x_tensor.pysurvives as a duplicate ofb12x.py'sB12xTensorFP8ScaledMMLinearKernel, andwarmup/kernel_warmup.pystill calls its standalonewarmup_b12x_tensor_fp8_linearwhile also callingb12x_warmup. Both select layers by theb12x_tensor_fp8_packed_weightattribute thatb12x.pysets, so the tensor-FP8 path is plausibly warmed twice. Confirming that — and removing the stale file — wants a GB10 warmup log, so it is left for the image build.AI assistance was used for this change.