Skip to content

build(vllmb12x): native MXFP8 MTP on B12X, CUDA dependency bump, Qwen3.8 step-5500 docs - #36

Merged
randomvariable merged 13 commits into
mainfrom
build/mxfp8-mtp-harbor
Sep 30, 2026
Merged

randomvariable merged 13 commits into
mainfrom
build/mxfp8-mtp-harbor

Conversation

@randomvariable

@randomvariable randomvariable commented Sep 30, 2026 •

Copy link
Copy Markdown
Owner

Change

Serve Qwen3.8-Flash-Next-NVFP4 QAD step 5500 (PLE 1000), whose MTP draft experts are MXFP8, entirely on B12X kernels. Carries the Karmic beta source rebase from #35.

  • vLLM: randomvariable/vllm dev/rv-mxfp8-mtp at 0a6739845a3249b07a30ad9b1e720e5e3fb6236f
    • da61df22f2 native MXFP8 MTP through the ModelOpt B12X backend
    • a749875731 test: disable established MXFP8 backends in the B12X fallback test
    • 0a6739845a prepare B12X MXFP8 TP=2 experts unpadded (replaces the interim 320->384 padding, b607be3600)
  • B12X: randomvariable/b12x feat/mxfp8-moe at 965e748f74fe9dac0a35c497a0947c5b108617c5
    • 5a6bcb27e6 native block-scaled MXFP8 W8A8 MoE (E4M3 weights, UE8M0 K32 scales)
    • 965e748f74 W8A8 MXFP8 experts at 32-aligned intermediate sizes via separate up/gate descriptors
  • Runtime dependencies (7272491d3d): CUDA pips to latest — nvidia-cutlass-dsl 4.8.0, quack-kernels 0.6.5, tilelang 0.1.15, cuda-python 13.4.1, cuda-tile 1.6.0, nvidia-cudnn-frontend 1.30.0, humming-kernels 0.1.16, tokenspeed-mla 0.2.15, nccl4py 0.6.0, nvtx 0.2.16, nvidia-ml-py 13.615.71 (apache-tvm-ffi 0.1.12, capped by tilelang). Added from upstream integration/karmic-kraken-beta: numba 0.65.0, PyNvVideoCodec 2.2.3, nixl 1.5.0, fastokens 0.3.2 (VLLM_USE_FASTOKENS=1); huggingface_hub floor raised to 1.31.0. Locks target manylinux_2_34.
  • CI (80328669e1, 0499544ea3, with GitOps 3dcd24307): both pipelines first release earlier finished runs' 256 GiB staging claims, which had starved the next clone on the build node.
  • Docs (b23e62b08d): Qwen3.8 Flash Next recipe and AI Gateway / llm-d routing docs match the live deployment; RecipeBuilder gains the routing target.

Verification

  • Script lane 55 pass; reproducibility lane 4 pass; included-changes check; recipes site 7/7 tests and full build.
  • SM120 (RTX PRO 6000, ARM64): B12X MXFP8 lane 29/29 (oracle parity at I=640/320/96, graph replay at 640 and 320), compute-sanitizer memcheck clean; vLLM MXFP8 lane 24/24 including the I=320 ModelOpt lifecycle. Same vLLM lane 24/24 with the bumped CUDA wheels overlaid.
  • MoE kernel time, SM120, I=320 unpadded vs I=384 padded (µs, median of 200, range over 3 processes): 1 tok 131.4-131.6 / 131.2-131.3; 4 tok 148.4-149.0 / 149.3-149.6; 16 tok 292.6-292.7 / 311.1-311.5; 64 tok 432.0-432.2 / 462.2-462.6.
  • Image pr-36-0a6739845a32-0499544ea3f9 (sha256:43aaf6d1b5e0…): both arch contracts pass; embedded source lock on arm64 and amd64 = vLLM 0a6739845a32, B12X 965e748f74fe.
  • Live GB10 (2 groups x TP=2, GitOps 818be2ad4): both groups boot; rank-0 log B12X NVFP4 target experts, B12X_MXFP8 MTP draft experts, nothing on Marlin. Adaptive shm spin enabled at 50 ms max grace; a 900 ms first boot deadlocked one group in B12X preparation (not isolated to the spin setting).

Benchmark

llm-inference-bench 0.7.5 against one idle group's rank-zero endpoint (0 running / 0 waiting before and after), 30 s cells, one sample per cell, temperature 1, max 8,192 output tokens. "Padded" = the previous image on the other group (different node pair).

Prefill:

context tok/s TTFT padded tok/s padded TTFT
8K 2,857 2.87 s 2,857 2.87 s
16K 2,882 5.61 s 2,898 5.65 s
64K 2,664 24.07 s 2,663 24.61 s
128K 2,416 52.98 s 1,873 69.98 s

Decode — aggregate tok/s / engine steps/s (MTP accept length):

context, concurrency this image padded
0, C1 52.5 / 23.4 (2.24) 45.1 / 23.4 (1.93)
0, C4 126.7 / 59.3 (2.14) 126.7 / 60.4 (2.10)
0, C8 181.8 / 86.8 (2.09) 178.6 / 84.7 (2.11)
16K, C1 51.0 / 23.6 (2.16) 52.5 / 23.3 (2.26)
16K, C4 129.4 / 61.1 (2.12) 116.6 / 58.9 (1.98)
16K, C8 175.2 / 84.4 (2.08) 170.3 / 81.7 (2.08)

Engine steps/s agree within 3% and prefill within 1% up to 64K: removing the padding saves nothing at the 1-4 tokens the draft processes per step. Aggregate differences track MTP accept length, which varies run to run at temperature 1. The 128K prefill cell ranged 1,873-2,416 tok/s across three single samples.

Limits

No SM121 kernel-level test beyond the live GB10 serving boot, no router/gateway-path benchmark, one sample per benchmark cell. The source forks' feature branches are not upstreamed.

@randomvariable

Copy link
Copy Markdown
Owner Author

/cancel-all

@randomvariable

Copy link
Copy Markdown
Owner Author

/retest

Both PR runs were evicted for node ephemeral-storage pressure during OCI materialization. Keep the checkout and Bazel output tree as separate directories on the 256Gi workspace instead of Tekton home emptyDir. Verify startup placement with a real Bazel output_base smoke run, the server dry-run, 55 script tests and four reproducibility tests.
@randomvariable

Copy link
Copy Markdown
Owner Author

The first candidate and its retest were evicted for node ephemeral-storage pressure during OCI materialization. Commit 6b8b64e moves publisher HOME, TMPDIR and Bazel output_user_root onto a 256Gi CSI workspace, with checkout and staging as sibling directories. Both publisher manifests pass server dry-run and shell/YAML validation. The real Bazel smoke confirms output_base under workspace staging. All 55 script tests and four reproducibility tests pass. PipelineRun vllmb12x-pull-request-tcmg6 is now building the unchanged source pair. No engine, release tag or latest tag is changed.

@randomvariable

Copy link
Copy Markdown
Owner Author

PipelineRun ci/vllmb12x-pull-request-tcmg6 succeeded. Multiarch assembly completed at 08:41:56 UTC; ARM64 and x86-64 image contracts passed at 09:19:34 and 09:48:04. Harbor publication completed at 09:54:06. Verified registry reference: harbor.services.home.internal.randomvariable.co.uk/custom/vllm-b12x-multi@sha256:628b6c74ec2b7244dc3de12f08beb7c25f29c73086389a5d23ff802d6deb7d63 (tag pr-36-da61df22f2b9-6b8b64e1673d). Both architecture manifests contain /opt/vllmb12x/sources.lock.json matching profiles/vllmb12x/profile.json, with hash-verified source-identity records for vLLM da61df22f2b99d5fce191c90fda808fbb3855c91 and B12X b600bf26c8e4775488faa079933ddb20303b8e66. This proves the CSI staging fix cleared the assembly/download step that evicted both earlier runs. No merge, release, latest-tag update, or deployment performed. Full step5500 serving remains untested pending isolated GPU capacity.

vLLM b607be3600 zero-pads the per-rank MXFP8 expert intermediate to
the B12X W8A8 K block, so the step5500 MTP experts load at TP=2
(640/2 = 320 -> 384).
CUDA packages now track their latest releases (pre-releases allowed):
nvidia-cutlass-dsl 4.8.0, quack-kernels 0.6.5, tilelang 0.1.15,
cuda-python 13.4.1, cuda-tile 1.6.0, nvidia-cudnn-frontend 1.30.0,
humming-kernels 0.1.16, tokenspeed-mla 0.2.15, nccl4py 0.6.0, nvtx 0.2.16,
nvidia-ml-py 13.615.71. tilelang 0.1.15 caps apache-tvm-ffi below 0.1.13,
so tvm-ffi resolves to 0.1.12.

Add what upstream integration/karmic-kraken-beta requires and the image
lacked: numba==0.65.0 (ngram proposer, JIT monitor, vision processors),
PyNvVideoCodec 2.2.3 (GPU video decoder), nixl 1.5.0 (object-store KV
tiering) and fastokens 0.3.2 (VLLM_USE_FASTOKENS=1). Raise the vendored
huggingface_hub floor to upstream's >=1.31.0 (arm64 had 1.30.0).

Locks now target manylinux_2_34, the tag pip.parse already downloads;
tilelang >=0.1.13 publishes aarch64 wheels only for 2_34.

Verified on the SM120 ARM64 host by overlaying the new arm64 wheels on
pr-36-b607be360037: every bumped module and its vLLM/B12X consumers
import, and the MXFP8 test lane passes 24/24 on CUTLASS DSL 4.8.0.
vLLM 0a6739845a drops the 320->384 zero padding; B12X 965e748f runs
W8A8 MXFP8 experts at 32-aligned intermediate sizes through separate
up/gate descriptors. SM120: B12X MXFP8 lane 29/29, vLLM lane 24/24,
memcheck clean; MoE kernel time equal at 1-4 tokens, 6-10% lower at
16-64 tokens.
Each run claims 256Gi on server21's thick-provisioned scratch pool and
kept it after finishing, so two finished runs blocked the next clone
pod for its whole 1h timeout. Both pipelines now run the
vllmb12x-release-staging Task (ci manifests) first, which deletes the
task pods and claims of earlier finished vllmb12x runs.
PAC inlines named taskRefs from the repository or hub and rejected the
run with 'cannot find referenced task'; a cluster-resolver reference is
left for Tekton to resolve from the ci namespace.
…deployment

Recipe: step-5500 revision 60215d26, 0.74 utilization, unpadded MXFP8 MTP
experts on B12X, adaptive shm spin at 50 ms, rank-zero KV events, the
image built from vLLM 0a6739845a32 + B12X 965e748f74fe, and the idle
benchmark on that image with the padded-draft run for comparison.

Routing docs: 3600 s route timeout, BackendTrafficPolicies, native
/v1/messages route, EPP -v=4, Envoy Gateway extension config and
GatewayClass controller name from the live cluster; RecipeBuilder
gains the routing target the guide already referenced.
@randomvariable randomvariable changed the title build(vllmb12x): native MXFP8 MTP source pair build(vllmb12x): native MXFP8 MTP on B12X, CUDA dependency bump, Qwen3.8 step-5500 docs Sep 30, 2026
@randomvariable
randomvariable merged commit b0cb13c into main Sep 30, 2026
3 checks passed
@randomvariable
randomvariable deleted the build/mxfp8-mtp-harbor branch September 30, 2026 22:40
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant