Skip to content

Blackwell dGPU bringup - #5

Closed
richiejp wants to merge 15 commits into
mudler:mainfrom
richiejp:local-blackwell-main-transplant
Closed

Blackwell dGPU bringup#5
richiejp wants to merge 15 commits into
mudler:mainfrom
richiejp:local-blackwell-main-transplant

Conversation

@richiejp

Copy link
Copy Markdown
Contributor

Continuation of RTX 5070 bringup. Presently there is a 6% regression in performance.

  • spike(loader): rebase plain-BF16 direct loading
  • feat(loader): transplant plain-BF16 Qwen3.5 loading
  • docs(bench): activate local Qwen3.5-4B comparison
  • bench: restore exact 4B corpus evidence
  • bench: script exact Qwen3.5-4B comparison
  • docs(bench): bind Qwen3.5-4B series root
  • fix(bench): initialize 4B leg names safely
  • docs(bench): bind corrected 4B retry root
  • fix(cuda): restore discrete Blackwell memory classification
  • docs(bench): record corrected Qwen3.5 4B comparison
  • bench: void mismatched non-AOT comparison
  • bench: record corrected 4B Triton AOT comparison

richiejp added 15 commits July 21, 2026 08:42
Record the selected local-Blackwell transplant scope on current main and reject superseded branch experiments before implementation.

FOLLOWING_AGENTS_PROTOCOL

Assisted-by: Codex:GPT-5 [Codex]
Port the still-relevant local Blackwell loader, residency, environment, and diagnostic changes onto current main while preserving the superseding engine paths. Add real 4B retained-host/direct-device CUDA equivalence coverage and keep oracle, gate-model, sanitizer, trace, memory, and performance closure explicit as GATING.

FOLLOWING_AGENTS_PROTOCOL

Assisted-by: Codex:GPT-5 [Codex]
Record the user-directed local 4B benchmark as the active plain-BF16 direct-loading gate, preserve the exact historical comparison recipe, and leave the 27B/35B validation explicitly pending external hardware.

FOLLOWING_AGENTS_PROTOCOL

Assisted-by: Codex:GPT-5 [Codex]
Restore ShareGPT prompt loading and submission-order token-ID capture in the benchmark-only path so the transplanted loader can be compared against the preserved Qwen3.5-4B series with an exact correctness precondition.

FOLLOWING_AGENTS_PROTOCOL

Assisted-by: Codex:GPT-5 [Codex]
Add a production-vLLM closed-loop metric and token collector, one-lock cache/thermal/memory driver, and current-versus-historical summarizer for the user-directed 4B benchmark.

FOLLOWING_AGENTS_PROTOCOL

Assisted-by: Codex:GPT-5 [Codex]
Bind the active claim and public checkpoint surfaces to the immutable b0a520f driver output directory before the lock-held GPU series.

FOLLOWING_AGENTS_PROTOCOL

Assisted-by: Codex:GPT-5 [Codex]
Split Bash local declarations so set -u cannot expand composed benchmark leg names before their phase arguments are assigned; record the first root as void before model launch.

FOLLOWING_AGENTS_PROTOCOL

Assisted-by: Codex:GPT-5 [Codex]
Bind the active benchmark claim and public surfaces to corrected driver 98dc954 and its immutable retry directory.

FOLLOWING_AGENTS_PROTOCOL

Assisted-by: Codex:GPT-5 [Codex]
Record the failed 4B classifier diagnostic and restore the integrated-device guard required for direct loading.

FOLLOWING_AGENTS_PROTOCOL

Assisted-by: Codex:GPT-5 [Codex]
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Codex:GPT-5 [Codex]
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Codex:GPT-5 [Codex]
FOLLOWING_AGENTS_PROTOCOL
Assisted-by: OpenAI Codex:GPT-5 [Codex]
Attribute the local current-versus-previous slowdown with matched node-mode nsys captures and record the H=32 Triton AOT repair target. No engine behavior changes.

FOLLOWING_AGENTS_PROTOCOL

Assisted-by: OpenAI Codex:GPT-5 [Codex]
Mirror the current vLLM FLA raw-packed FP32-state ABI for the dense Qwen3.5-4B geometry, dispatch the exact Hv=32 cubin, and extend CPU/AOT/rollback coverage plus generated artifact records. Performance gating follows from this immutable checkpoint.

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: OpenAI Codex:GPT-5 [Codex]
Record the immutable same-binary A/B, matched node-mode profile, vLLM 0.25 comparison, deterministic eager-oracle result, safety gates, and remaining parity blockers for the dense Qwen3.5-4B H32 repair.

FOLLOWING_AGENTS_PROTOCOL

Assisted-by: OpenAI Codex:GPT-5 [Codex]
@richiejp

Copy link
Copy Markdown
Contributor Author
  • Packed recurrence: 2.585 s → 1.467 s (−43.25%)
  • Fresh vLLM 0.25 comparison: 6418 vs 6744 tok/s (0.9517×), so overall parity remains open.

mudler added a commit that referenced this pull request Jul 21, 2026
…full dgx gate + 6.84% adjudication

Record-only change on top of richiejp's 12 rebased commits. No code touched.

REBASE. `local-blackwell-main-transplant` moved from merge-base a1611c7 onto
current origin/main a63c497 with `git rebase --onto`; authorship preserved on
all 12 commits, none squashed. Only .agents/parity-ledger.md and
.agents/state.md conflicted (append-only tails, commit d8c4164), resolved as a
union keeping main's Qwen3-Coder W5/W6 entries verbatim and appending the PR's
Blackwell entries. The other 11 commits applied clean. Verified surviving
intact: main's entire new bf16 grouped-MoE GEMM in cuda_matmul_nvfp4.cu
(MoeGroupedGemmBf16 / ...Wmma / ...WmmaPipe / ...NaiveSplitK, EnsureMoeScratch,
EnsureMoePartials, cp.async multistage) and main's MoeBlockBf16Cuda plus the
per-layer ReleaseHost() memory fix in qwen3_5.cpp. ENGINE_ROWS reconciled
109 -> 110 against the merged engine matrix.

GATES (dgx GB10, clean full rebuild, CUDA -Werror 0/0, goldens md5-verified,
FINAL binary). 27B test_qwen27_paged_engine 235/235; 35B
test_qwen36_paged_engine 315/315; Qwen3-Coder test_qwen3coder_paged_engine 6/6
(max teacher-forced gap 0.0000 nats); test_ops_moe_grouped_bf16 7/7; NVFP4
test_ops_moe_grouped 9/9; test_qwen3_moe_forward 3/3; PR tests
test_model_registry 18/18, test_cuda_backend 5/5, test_bench 4/4. Record
checkers green. test_qwen35_plain_weights passes 3/3 but its real-checkpoint
leg SKIPS — Qwen/Qwen3.5-4B is absent from dgx — so the PR's headline
direct-device gate is NOT exercised on this hardware.

REGRESSION ADJUDICATION. The declared -6.84% is NOT a code regression against
main: a same-arch (121a) interleaved same-box A/B of origin/main vs this head
on the 27B NVFP4 (1024/128, c8/32 prompts, 4 pairs, one flock, cold pair
dropped) is neutral on every axis — output throughput 0.9995x, median TPOT
+0.03%, median TTFT -0.14%, all inside +/-0.15% against a ~0.3% per-arm spread.
The CMake arch suffix is independently confirmed a first-order lever:
CMakeLists.txt gates fp4-MMA, CUTLASS NVFP4, Marlin MoE and FlashAttention-2 on
MATCHES "12[01]a", and cmake/TritonAOT.cmake derives the vendored AOT subdir as
sm_<arch>; measured same-binary on the 27B, FA2 alone is +3.9% throughput and
-36.4% median TTFT. RESIDUAL recorded, not papered over: the specific -6.84% is
not reproducible here (no 4B checkpoint, dgx is sm_121 not sm_120), the PR's
own stated argument proves only that the box was unchanged, and the stated
direction is counter-intuitive because 120a compiles IN more fast kernels than
120. The same-source 120-vs-120a A/B that would close it is still owed, but it
is scoped entirely to the PR's own new plain-BF16 4B path.

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
@localai-bot

Copy link
Copy Markdown
Collaborator

Merged to main as ee5e749..e7fd53b (rebased from merge-base a1611c7 onto a63c497 with git rebase --onto; no squash, authorship preserved on all 12 commits).

Conflicts: only .agents/parity-ledger.md and .agents/state.md, resolved as a union (main's Qwen3-Coder W5/W6 entries kept verbatim, Blackwell entries appended). The three files main rewrote the same day auto-merged with zero textual overlap, verified by inspection: main's bf16 grouped-MoE GEMM in cuda_matmul_nvfp4.cu and the per-layer ReleaseHost() fix in qwen3_5.cpp both survived intact. ENGINE_ROWS reconciled to 110.

Gates run on dgx GB10 (clean full rebuild, 121a, goldens md5-verified, final binary), which this PR had listed as PENDING:

Gate Result
27B test_qwen27_paged_engine 235/235
35B test_qwen36_paged_engine 315/315
Qwen3-Coder test_qwen3coder_paged_engine 6/6 (max gap 0.0000 nats)
test_model_registry / test_cuda_backend / test_bench / test_qwen35_plain_weights 18/18, 5/5, 4/4, 3/3
compute-sanitizer memcheck (27B engine path) 0 errors
CUDA -Werror 0 warnings

test_cuda_backend logs pageable=1 integrated=1 UnifiedMemory=true on GB10, so the discrete-Blackwell classification fix is confirmed behavior-neutral on integrated hardware — it correctly stops a discrete Blackwell (which reports PageableMemoryAccess via HMM/UVM) from being handed the zero-copy DeviceScratch/KV/GDN contract.

On the 6.84%: a same-arch interleaved A/B of origin/main vs this branch on GB10 (27B NVFP4, 1024/128, c8, cold pair dropped) came out neutral to within ±0.15% on every axis against a ~0.3% spread — so there is no code regression reaching existing paths. That is also structural: DirectDeviceLoadEligible requires !is_unified_memory(), and the 35B/Coder are is_dense_model=false, so none of them enter the new branch.

Two caveats recorded rather than glossed:

  1. The 120-vs-120a attribution is asserted, not measured. The stated argument ("local vLLM is 1.002x the previous reference, so the residual is project/build-side") doesn't carry — an unchanged vLLM arm shows the box is unchanged, not that our delta is build- rather than code-side. The direction is also counter-intuitive, since 120a compiles in more fast kernels. The mechanism is real and large enough (CMakeLists.txt:71,88,97,137,612 gate four kernel families on 12[01]a; FA2 alone measured +3.9% throughput / -36.4% TTFT on the 27B), but the owed experiment is a same-source 120-vs-120a A/B on the 5070 — dgx is sm_121 and has no 4B checkpoint, so it can't substitute.
  2. test_qwen35_plain_weights passes, but its real-checkpoint leg skips here (Qwen/Qwen3.5-4B isn't on dgx), so the headline direct-device path is not exercised by our gates.

Merged on the basis that the gates are green and the 6.84% is scoped entirely to a newly-added path that has no main baseline it could regress. Thanks for the unusually rigorous benchmark record — voiding your own mismatched comparison rather than reporting it is the right instinct.

mudler added a commit that referenced this pull request Jul 22, 2026
…rget, first arch through the additive seams

USER PRIORITY 1 had the framework built and zero architectures added. This adds
the first one: BACKEND-CUDA-SM120 moves PARTIAL -> ACTIVE, meaning CLAIMED and
BUILD-SUPPORTED. It is deliberately NOT a runtime-support claim.

The headline is what sm_120a did NOT need: zero kernel, model, runner, sampler or
FEATURE-TABLE edits. The arch-additivity seams already covered it — the table
cells were written 12.0a,12.1a at W1 (mirroring vLLM FP4_SM120_ARCHS), the fp4
tactic predicate is caps.sm_major == 12, attention priority is
LookupAttnPriority(major), and residency is pageable && integrated, which is
exactly the discrete-sm_120 case PR #5 fixed. The audit prediction (breadth-sweep
§A.3) is confirmed by measurement.

CRUX RESOLVED: the recorded "heterogeneous fat build cannot compile" obstacle is
FAMILY-SCOPED. "120a;121a" configures, compiles and links -Werror clean (0 warn),
and cuobjdump shows 38 TUs with sm_120a cubins alongside 38 sm_121a, the fp4-MMA
TU among them; single-arch "120a" likewise builds clean (38 sm_120a). The
sm12x-only PTX, its __CUDA_ARCH_SPECIFIC__ guard and the CUTLASS Sm120 arch tag
are satisfied by both targets. "90a;121a" is unchanged and still cannot compile;
per-source gencode narrowing stays W7.

Landed:
  * cmake/CudaArchFeaturesTest.cmake (NEW) — 35 hard expectations on the per-arch
    feature resolution, cmake -P, no CUDA toolkit and no GPU, mutant-checked;
    wired into CI as the cuda-arch-features job.
  * cmake/TritonAOT.cmake — detect the multi-arch case instead of deriving a
    nonexistent sm_120a_121a directory. AOT cubins are single-arch by nature;
    the build degrades gracefully via VLLM_CPP_TRITON=OFF. No cubins were faked:
    triton_aot_vendored/sm_120a/ must be regenerated ON the card.
  * CMakeLists.txt — VLLM_CPP_CUDA_ARCHITECTURES documents supported values.
    The DEFAULT stays 121a; a fat default would double every build for no gain.

GB10 REGRESSION PROOF (the real risk — a fat build changes codegen for the arch
we DO run). Full set re-run ON the fat binary at production flags, differing from
production in the arch list alone, each gate STANDALONE, GPU idle, under
flock $HOME/gpu.lock: 27B 235/235, 35B 315/315, Qwen3-Coder 138/138 (6/6),
Qwen3-dense 664/664 (16/16 on 0.6B and 4B), OPT 36/36 (6/6), DeepSeek-V2-Lite
223/223 (8/8) — ALL UNCHANGED. Clean CUDA -Werror 0 warnings. Golden corpus (475
files) md5 identical before and after: 2965ef5772b556d3f3f86fedf4221b2f.

NOT PROVEN, stated plainly: no gate model, kernel suite or benchmark has ever
executed on sm_120 hardware, because none exists here. A green fatbinary link is
not execution evidence. The only sm_120 runtime evidence in the tree remains an
external RTX 5070 Ti loader/backend-leaf result (PR #5), cited not re-derived.
docs/BENCHMARKS.md records NOT APPLICABLE. The validation steps for a 50-series
owner are written down in the spec.

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
mudler added a commit that referenced this pull request Jul 29, 2026
… store+read

HIGH-priority feature-gap #5 (KV-FP8): the standard memory/throughput lever that
halves the KV footprint by storing K/V as fp8 with a per-tensor dequant scale.
CPU-only W1 brick, mirroring vLLM 1:1.

W0: .agents/specs/fp8-kv-cache.md spikes the whole vLLM fp8-KV path (the
cache_dtype config, BaseKVCacheMethod scale handling, the fp8 store in
reshape_and_cache_flash, the read dequant, the halved-block memory accounting)
with the exact scale convention FP8=Quantize(HP/scale); Dequant(FP8)*scale=HP.

W1 (all additive, CPU-only):
- include/vt/fp8_kv.h (NEW): Fp8KVCacheDataType interpretation enum (mirror
  csrc/attention/dtype_fp8.cuh:9-13) + the fp8-e4m3 codec + StoreKvFp8E4M3/
  LoadKvFp8E4M3, bit-identical to the landed vllm::F32ToF8E4M3/F8E4M3ToF32 (vt
  does not depend on vllm).
- vt::ReshapeAndCacheFp8 (new OpId kReshapeAndCacheFp8 + wrapper + CPU kernel):
  the fp8 store, mirroring reshape_and_cache_flash_kernel's fp8 branch
  (cache_kernels.cu:241-252,314-401).
- additive default-inert PagedAttentionArgs.{kv_cache_dtype,k_scale,v_scale} +
  the CPU paged-attention read dequant (Dequant(fp8)*scale).
- vllm::v1::ParseCacheDType/IsQuantizedKvCache: the cache_dtype config parse
  (mirror CacheDType, config/cache.py:19-36,76 + is_quantized_kv_cache).

fp8 storage rides DType::kI8 (1 byte + the interpretation enum, mirroring vLLM's
cache_t=uint8_t + KV_DTYPE template param), NOT a new DType enumerator. A new op
(not a widened ReshapeAndCacheFn) keeps the CUDA/Metal registrations untouched.
CUDA/non-CPU + fp8_e5m2 are refused at the wrapper (named later bricks).

Gate: tests/vt/test_ops_fp8_kv_cache.cpp 8 cases / 511 assertions GREEN (ported
from tests/kernels/attention/test_cache.py::test_reshape_and_cache fp8 branch):
store->read within the e4m3 band, fp8-vs-bf16 NMSE<1%, paged-attention e2e within
5%, ParseCacheDType mirror. RED-first proven: a wrong store direction (hp*scale)
fails 3 cases/480 assertions; a wrong read v_scale diverges >0.05. No sibling
regression (reshape 12/12, paged 14/14). CPU -Wall -Wextra -Werror clean.

Residuals (named W2-W5 in the spec): the CUDA fp8 store + fp8 paged-attention
read (the GPU memory-halving path, DGX-blocked); the runner/spec integration
(half-sized KV blocks + checkpoint k/v_scale threading + --kv-cache-dtype/
--calculate-kv-scales); fp8_e5m2 CPU compute + per-attention-head scales.

Records: KV-FP8 engine row INVENTORIED->ACTIVE (rollup ACTIVE 42->43, INVENTORIED
35->34), QUANT-KV-FP8 INVENTORIED->PARTIAL, CLAIM-KV-FP8, feature-matrix, roadmap
parity-ledger, state. No counted-matrix row added.

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
mudler added a commit that referenced this pull request Jul 31, 2026
…ble/token — 3.9x decode

The W7 speed profile measured Laguna decode at 32.7% GPU-active, 67% host-orchestration.
LagunaEmbed dequantized the ENTIRE [Vsz,H] embed table to f32 (~311M element-converts,
~1.23 GB) on every token just to gather T rows — the DOMINANT decode cost, not the minor
"#5 cleanup" the profile filed it as. Now gathers only the T needed rows directly from
the (f32/bf16) table bytes.

BIT-IDENTICAL by construction: same per-element f32/bf16->f32 conversion, same rows
selected. GATED on the real 3-shard UD-Q4_K_XL GGUF (GB10 sm_121a, --gpu, W6 cached
path, drop_caches cold, prompt "The capital of France is", 24 tokens):
- TOKEN-IDENTICAL to the W5/W6 golden (byte-equal ids `22345 83 350 785 ...`, coherent
  " Paris." continuation).
- decode 0.66 -> 0.17 s/tok = 3.9x (Laguna 1.5 -> 5.9 tok/s; 18x -> 4.7x vs llama.cpp
  27.8). Prefill 0.98s, peak 69.96 GiB.

Syntax-verified with production build flags; CUDA laguna-gen build clean on GB10.

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]
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.

2 participants