diff --git a/.agents/backend-matrix.md b/.agents/backend-matrix.md index 25959ddd0d..30936cdd57 100644 --- a/.agents/backend-matrix.md +++ b/.agents/backend-matrix.md @@ -233,7 +233,7 @@ this repository. State remains `ACTIVE`; no lifecycle transition is claimed. | `BACKEND-SEAM-AUDIT` | **Accelerator-seam audit + the "shared logic, swapped ops" plan** — does our MLX/Vulkan/Metal architecture actually port vLLM's CUDA-path strategy, and where does the shared layer still leak device specificity? Owns the side-by-side construct mapping, the re-measured leakage inventory with per-site classification, the proposed **DSR** (Device-Specific References) CI ratchet, and the ranked `S1`-`S8` plan. **AUDIT ONLY — no code, no build, no GPU**. **Work rows `S1` (DSR ratchet), `S3` (Platform capability fields), `S4` (`LinearMethod`/`QuantizationConfig` seam), `S5` (portable reference tier) and `S7` (extract the last shared-layer device coupling — DSR to its 32 build-gate floor) are IMPLEMENTED and gated; `S6` assessed no-op**; `S2`,`S8` remain plan only | `vllm/platforms/interface.py:134-229` (Platform, ~85 methods) + `rocm.py:531-600` / `xpu.py:121-167` as the non-CUDA comparators; `vllm/v1/attention/backends/registry.py:34-120`; `vllm/model_executor/custom_op.py:103-200` (42 registrations, `forward_native` universal fallback); `vllm/model_executor/layers/quantization/base_config.py:20-229` + `linear.py:141-230` @ pin `e24d1b24` | audit target surfaces, unmodified: [interface.h:109](../include/vllm/platforms/interface.h#L109), [cuda.cpp:18](../src/vllm/platforms/cuda.cpp#L18), [cuda_attn_priority.h:48](../include/vllm/platforms/cuda_attn_priority.h#L48), [registry.h:44](../include/vllm/v1/attention/registry.h#L44), [op_provider.h:108](../include/vt/op_provider.h#L108), [ops.h:88](../include/vt/ops.h#L88) (the 75-entry `OpId` table), and the leakage concentration [qwen3_5.cpp:1022](../src/vllm/model_executor/models/qwen3_5.cpp#L1022). **`S1` as-built (the only code this row owns — nothing under `src/`/`include/` is touched):** [check-device-leakage.py](../scripts/check-device-leakage.py#L102) (the per-file allowlist), [device-leakage-baseline.json](../scripts/device-leakage-baseline.json), CI job `device-leakage` in [ci.yml](../.github/workflows/ci.yml) | **MEASURED at `72f5db2`, static analysis only (no build/run/GPU):** shared-layer **DSR = 94** (47 `kCUDA` + 11 `is_cuda()` + 4 unconditional CUDA includes + 32 `#ifdef VT_*`), of which **67 (71%) are in `qwen3_5.cpp` alone**; upstream comparator `vllm/model_executor/models/qwen3_next.py:321` carries **1** predicate in 802 lines vs our 6,389; upstream puts **199 of 544** device predicates in `model_executor/layers/` and only **14 (9 of 287 files)** in `models/`. Counts **regrew** since [metal-mlx-reuse-study.md](specs/metal-mlx-reuse-study.md) §3.3 (54->63 `kCUDA`, 13->16 `is_cuda()`), which is the grounding for the ratchet. Behavioural evidence reused, not re-run: Vulkan V1 edited **0** pre-existing `src/`/`include/` files ([fan-out spike](specs/backend-fanout-metal-vulkan-xpu.md) § V1 landed). Record gates only: `check-agent-record.py` + `check-doc-checkpoint.py` **`S1` RE-DERIVED at `18094ee` by the as-built checker (static analysis; no build, no GPU, no source change): DSR = 86** (44 `kCUDA` + 10 `is_cuda()` + **0** unconditional CUDA includes + 32 `#ifdef VT_*`), of which **66 (77%) are in `qwen3_5.cpp`**. The audit's 94 was **-6 composition** (comments/strings now stripped uniformly; its 4 "unconditional" CUDA includes were ALL already `#ifdef`-guarded — a measurement error) and **-2 real reduction** landed on main since `72f5db2`. Seven-class composition re-verified site-by-site: B=5 and G=32 exact, C 16->10, D 6->**3**, E 1->**0**, F 4->**0**, A=46 — so **3 of 86 are genuine device policy**, sharpening rather than weakening the audit's verdict. Ratchet enforced per-bucket AND total, failing on an increase AND on an undeclared decrease. [Mutation suite 24/24](../tests/scripts/test_device_leakage.py#L139): a planted `kCUDA`/`is_cuda()`/unguarded-include/`#ifdef VT_*` FAILS, a leak hidden behind a shared-layer helper still FAILS, adding a backend leg PASSES, an extra reference in a budgeted allowlist file FAILS, a reduction FAILS until the baseline is lowered in the same commit, and `--write-baseline` REFUSES to ratchet upward. Record gates: `check-agent-record.py` + `check-doc-checkpoint.py` + `check-device-leakage.py` green | [accelerator seam audit](specs/accelerator-seam-audit.md) | `ANCHOR-BACKFILL` — audit committed, **`S1` LANDED + CI-gated**, and **`S4` LANDED (partial, byte-identical) 2026-07-23**; **verdict PARTIAL→improving**: the Platform and attention-registry seams mirror faithfully (open attention self-registration betters upstream's closed enum), and the `model_executor/layers/` library the audit named ABSENT now EXISTS in part — `S4` ships the `LinearMethod`/`QuantizationConfig` seam (`include/vllm/model_executor/layers/{linear.h,quantization/base_config.h,quantization/compressed_tensors/schemes/nvfp4.h}`, mirroring `linear.py:141-230`+`base_config.py:20-229`), routes the dense `qwen3.cpp` MLP through `method.Apply()` (scheme chosen ONCE, `IsNvfp4()` probe retired), and converts 18 provably-`kCUDA`-only device gates to `vt::OpRegistered`. **DSR 86 → 67** (`kcuda` 44→25, `is_cuda` 10, `cuda_inc` 0, `vt_ifdef` 32), baseline lowered same commit; ratchet + 24-case suite green. Byte-identical on dgx (standalone, `flock`): 27B 235/235, 35B 315/315, Coder 6/6, dense 16/16, OPT 6/6, DeepSeek-V2 8/8, goldens content-hash unchanged; `test_linear_method` + quant suites green; clean CUDA+CPU `-Werror` 0-warn. The fragile 27B-W4A4/fp8-recipe gates (ops dual-registered on CPU or realized via a recipe) are correctly DEFERRED to `S6` behind `S5`'s reference tier. **`S5` LANDED 2026-07-23** (portable reference tier, DSR 67 unchanged, vt-runtime infra). **`S6` ASSESSED 2026-07-23 (`CLAIM-BACKEND-SEAM-S6-1`) → NO-OP / BLOCKED:** the deferred fp4/fp8 gates convert ZERO sites byte-identically — every one bottoms out at a **dual-registered** (CPU+CUDA) bespoke op (`kMatmulNvfp4Fp4`/`kScaledFp4Quant`/`kSiluMulFp4Quant`/`kSigmoidGateFp4Quant`/`kRmsNormQuantFp8`/`kRmsNormGatedQuantFp8`, none CUDA-only), so `OpRegistered(op,dev)` is TRUE on `kCPU` ⇒ the class-A swap flips the CPU reference/emulation path (two numerics per device); S5's tier does not change this (present natively, never a miss). **DSR stays 67, no baseline moved, no `src/`/`include/`/test byte changed** (spec §11). Genuine unlock re-scoped to `S3` (Platform capability fields mirroring `supports_fp8`/`cutlass_fp4_supported`) + `S7` (extraction). **`S3` LANDED 2026-07-23 (`CLAIM-BACKEND-SEAM-S3-1`) — the byte-identical unlock:** 6 capability predicates added to `Platform` (`supports_fp8`/`cutlass_fp4_supported`/`opaque_attention_op`/`is_integrated_gpu`/`support_static_graph_mode`/`is_device_capability_family`, base false; CudaPlatform answers GB10 values, `interface.h`+`cuda.cpp`+`platform.cpp`) + **12** deferred `qwen3_5.cpp` gates converted (7 fp4-act `cutlass_fp4_supported`, 3 fp8-fused `supports_fp8`, 2 decode-graph `support_static_graph_mode`) — byte-identical because a capability answers the base false off CUDA, exactly what `device==kCUDA` did (where S6's `OpRegistered` was TRUE on `kCPU`). **DSR 67 → 55** (`kcuda` 25→13, `is_cuda` 10, `cuda_inc` 0, `vt_ifdef` 32), baseline lowered same commit, ratchet + 24-case suite green; dgx standalone byte-identical 27B 235/235 · 35B 315/315 · Coder 6/6 · dense-32B 16/16 · OPT 6/6 · DeepSeek-V2 8/8 · Llama 16/16 (no golden regenerated), `test_platform` CUDA-leg proves each predicate == former `device==kCUDA` on GB10, `compute-sanitizer memcheck` 0 errors, clean CUDA+CPU `-Werror` 0-warn. **`S7` LANDED 2026-07-23 (`CLAIM-BACKEND-SEAM-S7-1`) — the runtime-decoupling arc's TERMINAL milestone:** ALL 23 remaining runtime `kCUDA`/`is_cuda()` leakage sites hoisted onto capabilities — new `Platform::needs_weight_staging()` (CUDA device-resident staging policy; NOT `is_unified_memory()`, which would FLIP GB10 — proven by `test_platform`'s `!=` assertion) covering residency/merged-GDN/packed-decode/direct-load, `Platform::supports_fa2_attention()` (FA2 dtype), `Backend::SupportsAuxStream()` (MoE aux-stream), reusing S3's `supports_fp8`/`cutlass_fp4_supported`/`support_static_graph_mode`/`is_integrated_gpu` (runner combine/scatter) and `vt::OpRegistered(kMoeGroupedGemmBf16)` (CUDA-only op). **DSR 55 → 32** (`kcuda` 13→0, `is_cuda` 10→0, `cuda_inc` 0, `vt_ifdef` 32) — the IRREDUCIBLE floor: the shared model layer holds ZERO runtime device tests, all 32 residual are `#ifdef VT_*` compile-time kernel gates, so the audit's `<10` is NOT reachable. baseline lowered same commit, ratchet + 24-case suite green; dgx standalone byte-identical 27B 235/235 · 35B 315/315 · Coder 6/6 · dense-32B 16/16 · OPT 6/6 · DeepSeek-V2 8/8 · Llama 16/16 (no golden regenerated), new `test_platform`/`test_backend`/`test_cuda_backend` cases green, `compute-sanitizer memcheck` 0 errors, clean CUDA+CPU `-Werror` 0-warn. The `layers/`-library physical relocation is a follow-on refactor; the device coupling is already gone. `S2`,`S8` each still need their own claim; no gate row moves | `CLAIM-BACKEND-SEAM-S1-1`, `CLAIM-BACKEND-SEAM-S4-1`, `CLAIM-BACKEND-SEAM-S6-1`, `CLAIM-BACKEND-SEAM-S3-1`, `CLAIM-BACKEND-SEAM-S7-1` | | `BACKEND-ABI-VT` | Backend registry, device/queue, capture, drop-in adapter ABI | platform contract `vllm/platforms/interface.py:67-229`; stable stream `csrc/libtorch_stable/torch_utils.h:76-82`; raw Marlin boundary `csrc/libtorch_stable/moe/marlin_moe_wna16/ops.cu:344-356` | queue ID [device.h:9](../include/vt/device.h#L9), explicit resources [backend.h:53](../include/vt/backend.h#L53), scalar/layout ABI [ops.h:13](../include/vt/ops.h#L13), CUDA workspace/raw probe [cuda_dropin.cu:104](../src/vt/cuda/cuda_dropin.cu#L104) | [ported ABI tests](../tests/vt/test_dropin_abi.cpp#L86); CPU 94/94 + repair-focused 1/1; GCC13 repair at [test_dropin_abi.cpp:245](../tests/vt/test_dropin_abi.cpp#L245); exact sm_121a all-target build 100%, focused CUDA/ABI CTest 2/2, sanitizer 9/9 + 196/196 with 0 errors/leaks, 35B/27B gates 2/2 (`1141b79`, evidence manifest `4adbe952…601`) | [drop-in ABI](specs/dropin-kernel-abi.md) | `ANCHOR-BACKFILL` - sm_80/sm_90a cross-build and unchanged-trace/model A/B-memory proof remain; common scalar forwarder and legacy backend-shim migration remain; no production family migrated | CLAIM-BACKEND-ABI-W0-GPU-1 | | `BACKEND-ATTN-REGISTRY` | Attention-backend registry + platform-driven priority selection (extensibility **item 4** — the second portability seam, now REALIZED). Backends SELF-REGISTER per `(DeviceType, name)`; `Platform::get_attn_backend_priority()` advertises a capability-ordered name list; the selector returns the first REGISTERED name. Selection is DATA (register + priority slot), not an inline code edit. The concrete attention KERNEL stays at the vt:: op-table (`vt::PagedAttention`→`GetOp`), already device-additive | `vllm/v1/attention/backends/registry.py` (self-registration), `vllm/platforms/cuda.py:361-470` (`get_valid_backends`/`get_attn_backend_cls`) + `:84-166` (`_get_backend_priorities`, non-MLA major-10 vs else), `cpu.py:75-87` (CPU_ATTN) @ pin `e24d1b24` | registry+selector [registry.h:44](../include/vllm/v1/attention/registry.h#L44), [registry.cpp:60](../src/vllm/v1/attention/registry.cpp#L60); priority [interface.h:92](../include/vllm/platforms/interface.h#L92), CUDA major-10/else [cuda.cpp:44](../src/vllm/platforms/cuda.cpp#L44), CPU [cpu.cpp:29](../src/vllm/platforms/cpu.cpp#L29); FLASH_ATTN self-reg [backend.cpp:83](../src/vllm/v1/attention/backend.cpp#L83), GDN [gdn_attn.cpp:153](../src/vllm/v1/attention/backends/gdn_attn.cpp#L153) | [registry tests](../tests/vllm/v1/attention/test_attn_backend_registry.cpp#L61) (self-register + Make/throw + CUDA/CPU priority order + first-registered walk + override; 8 cases/25 assertions) + [platform priority](../tests/vllm/platforms/test_platform.cpp#L53); clean CPU `-Werror` + full CPU CTest green; behavior-preserving ⇒ same FA2 attention selected (**DGX-CONFIRMED @ `2c732e7`: 27B 235/235 + 35B 315/315 token-exact, FA2 sm_121a, memcheck 0/315**) | [Platform seam plan](specs/extensibility-platform-seam-2026-07-18.md) (item 4) | `ANCHOR-BACKFILL` — item-4 seam LANDED CPU + DGX-CONFIRMED; FLASHINFER/TRITON_ATTN/FLEX/TURBOQUANT are named in the priority (fidelity) but unimplemented ⇒ walk falls through to FLASH_ATTN (the behavior-preserving outcome); MLA-branch priorities deferred until an MLA model ports | CLAIM-ATTN-REGISTRY-1 | -| `BACKEND-PLATFORM` | Platform capability/memory-model seam (`is_cuda`/`is_cpu`, `is_unified_memory`, `has_device_capability`, `supported_dtypes`, `residency_policy`, `supports_graph_capture`; `get_attn_backend_priority` REALIZED by `BACKEND-ATTN-REGISTRY` item 4; **+ the S3 capability predicates `supports_fp8`/`cutlass_fp4_supported`/`opaque_attention_op`/`is_integrated_gpu`/`support_static_graph_mode`/`is_device_capability_family`, mirroring `interface.py:914,933,977,1058,441-476` + `nvfp4_utils.py:56`, added by `CLAIM-BACKEND-SEAM-S3-1` and consumed by the fp4/fp8/graph gates in `qwen3_5.cpp`**); self-registered per `DeviceType` via `CurrentPlatform()` | `vllm/platforms/interface.py:134-229` (`class Platform`), `:409-439` (get/has_device_capability), `:181-187` (supported_dtypes) @ pin `e24d1b24` | interface [interface.h:56](../include/vllm/platforms/interface.h#L56), registry/CurrentPlatform [platform.cpp:57](../src/vllm/platforms/platform.cpp#L57), CPU [cpu.cpp:11](../src/vllm/platforms/cpu.cpp#L11), CUDA [cuda.cpp:14](../src/vllm/platforms/cuda.cpp#L14); composes `vt::Backend` [backend.h:22](../include/vt/backend.h#L22) | [platform tests](../tests/vllm/platforms/test_platform.cpp#L28) (registration + CPU capabilities + has_device_capability), clean CPU `-Werror` + full CPU CTest; 7 memory-model/residency sites migrated (`runner.cpp`, `model_registry.cpp`, `qwen3_5.cpp`) — **per-tensor sites key on `GetPlatform(.device.type).is_cuda()` (the OBJECT's device), NOT the process-global `CurrentPlatform()`, which is accelerator-first and would mis-route a CPU queue/tensor on a GPU box (regression fixed 2026-07-18; red DGX CPU tests `test_platform`/`test_qwen27_dense_forward`); `CurrentPlatform()` is reserved for process-level "which accelerator" questions only**. DGX behavior-preserving model gates 27B 235/235 + 35B 315/315 CONFIRMED @ (CUDA -Werror-clean after a GCC13 dangling-pointer FP fix). **Item-2 residency_policy() CONSUMPTION LANDED (2026-07-19, `CLAIM-BACKEND-PLATFORM-2`):** the host-free / load-stream / DevicePool-cap decisions in `qwen3_5.cpp` now READ `GetPlatform(.device.type).residency_policy()` (per-device) instead of an inline `device.type`/env gate — host-free `qwen3_5.cpp:BuildMoeMarlinResident` via `ShouldReleaseHostWeights`, load-stream `Qwen3_5Model::PrepareMarlinResident` via `ShouldInterleaveLoadStream` (both `interface.h`), DevicePool soft cap via `residency_policy().device_pool_cap_bytes` (`DBuf`/`DevicePool`). `CudaPlatform::residency_policy().release_host_weights_after_upload` flipped **false→true** (now CONSUMED ⇒ reproduces today's GB10 host-free-after-Marlin-build EXACTLY); `MarlinMoeEnabled()` stays the orthogonal KERNEL-PATH gate; `VT_MOE_HOST_FREE`/`VT_MOE_LOADSTREAM` env stay as overrides. Consumption unit-tested `test_platform.cpp` (7 cases/43 assertions). A new (discrete) GPU sets `residency_policy()` values ⇒ ZERO model edit. **DGX gate PASSED @ `62fc0e0` (production flags CUTLASS sm120a + FA2 sm_121a + Triton AOT, one flock): clean CUDA `-Werror` 0 warn, 27B 235/235 + 35B 315/315 token-exact, 35B VmHWM ≈ 4.0 GiB (load-stream win preserved), memcheck 0 errors.** | [Platform seam plan](specs/extensibility-platform-seam-2026-07-18.md) | `ANCHOR-BACKFILL` — item-1 extraction DONE + DGX-CONFIRMED; item-2 residency consumption LANDED + DGX-PROVEN (`CLAIM-BACKEND-PLATFORM-2`); kernel-shape dispatch branches deliberately left for items 4/5 | CLAIM-BACKEND-PLATFORM-1 / CLAIM-BACKEND-PLATFORM-2 | +| `BACKEND-PLATFORM` | Platform capability/memory-model seam (`is_cuda`/`is_cpu`, `is_unified_memory`, `has_device_capability`, `supported_dtypes`, `residency_policy`, `supports_graph_capture`; `get_attn_backend_priority` REALIZED by `BACKEND-ATTN-REGISTRY` item 4; **+ the S3 capability predicates `supports_fp8`/`cutlass_fp4_supported`/`opaque_attention_op`/`is_integrated_gpu`/`support_static_graph_mode`/`is_device_capability_family`, mirroring `interface.py:914,933,977,1058,441-476` + `nvfp4_utils.py:56`, added by `CLAIM-BACKEND-SEAM-S3-1` and consumed by the fp4/fp8/graph gates in `qwen3_5.cpp`**); self-registered per `DeviceType` via `CurrentPlatform()` | `vllm/platforms/interface.py:134-229` (`class Platform`), `:409-439` (get/has_device_capability), `:181-187` (supported_dtypes) @ pin `e24d1b24` | interface [interface.h:56](../include/vllm/platforms/interface.h#L56), registry/CurrentPlatform [platform.cpp:57](../src/vllm/platforms/platform.cpp#L57), CPU [cpu.cpp:11](../src/vllm/platforms/cpu.cpp#L11), CUDA [cuda.cpp:14](../src/vllm/platforms/cuda.cpp#L14); composes `vt::Backend` [backend.h:23](../include/vt/backend.h#L23) | [platform tests](../tests/vllm/platforms/test_platform.cpp#L28) (registration + CPU capabilities + has_device_capability), clean CPU `-Werror` + full CPU CTest; 7 memory-model/residency sites migrated (`runner.cpp`, `model_registry.cpp`, `qwen3_5.cpp`) — **per-tensor sites key on `GetPlatform(.device.type).is_cuda()` (the OBJECT's device), NOT the process-global `CurrentPlatform()`, which is accelerator-first and would mis-route a CPU queue/tensor on a GPU box (regression fixed 2026-07-18; red DGX CPU tests `test_platform`/`test_qwen27_dense_forward`); `CurrentPlatform()` is reserved for process-level "which accelerator" questions only**. DGX behavior-preserving model gates 27B 235/235 + 35B 315/315 CONFIRMED @ (CUDA -Werror-clean after a GCC13 dangling-pointer FP fix). **Item-2 residency_policy() CONSUMPTION LANDED (2026-07-19, `CLAIM-BACKEND-PLATFORM-2`):** the host-free / load-stream / DevicePool-cap decisions in `qwen3_5.cpp` now READ `GetPlatform(.device.type).residency_policy()` (per-device) instead of an inline `device.type`/env gate — host-free `qwen3_5.cpp:BuildMoeMarlinResident` via `ShouldReleaseHostWeights`, load-stream `Qwen3_5Model::PrepareMarlinResident` via `ShouldInterleaveLoadStream` (both `interface.h`), DevicePool soft cap via `residency_policy().device_pool_cap_bytes` (`DBuf`/`DevicePool`). `CudaPlatform::residency_policy().release_host_weights_after_upload` flipped **false→true** (now CONSUMED ⇒ reproduces today's GB10 host-free-after-Marlin-build EXACTLY); `MarlinMoeEnabled()` stays the orthogonal KERNEL-PATH gate; `VT_MOE_HOST_FREE`/`VT_MOE_LOADSTREAM` env stay as overrides. Consumption unit-tested `test_platform.cpp` (7 cases/43 assertions). A new (discrete) GPU sets `residency_policy()` values ⇒ ZERO model edit. **DGX gate PASSED @ `62fc0e0` (production flags CUTLASS sm120a + FA2 sm_121a + Triton AOT, one flock): clean CUDA `-Werror` 0 warn, 27B 235/235 + 35B 315/315 token-exact, 35B VmHWM ≈ 4.0 GiB (load-stream win preserved), memcheck 0 errors.** | [Platform seam plan](specs/extensibility-platform-seam-2026-07-18.md) | `ANCHOR-BACKFILL` — item-1 extraction DONE + DGX-CONFIRMED; item-2 residency consumption LANDED + DGX-PROVEN (`CLAIM-BACKEND-PLATFORM-2`); kernel-shape dispatch branches deliberately left for items 4/5 | CLAIM-BACKEND-PLATFORM-1 / CLAIM-BACKEND-PLATFORM-2 | | `BACKEND-CPU` | CPU correctness and production path | `platforms/cpu.py:42-125`, CPU ops rooted at `csrc/cpu/torch_bindings.cpp:123-139`; llama.cpp `ggml-cpu.c:471-610,3024-3390` owns the added native pool reference | [cpu_backend.cpp:11](../src/vt/cpu/cpu_backend.cpp#L11), [threadpool.cpp:78](../src/vt/cpu/cpu_threadpool.cpp#L78), [chunked ops](../src/vt/cpu/cpu_ops.cpp#L88) | [backend tests](../tests/vt/test_backend.cpp#L10), [op parity](../tests/parity/test_op_parity.cpp#L34), [threadpool/upstream-test port](../tests/vt/test_cpu_threadpool.cpp#L63); 1/3/20 full CPU suites + TSAN pass. The binding 20-core Arm/i8mm Qwen3.5-2B Q8 single-stream gate is at or ahead of same-file llama.cpp on prefill/decode/RSS; server concurrency remains open. **Raspberry Pi 5 / Cortex-A76 lane GATING 2026-08-06:** R0-R3 pin the exact Q8_K_XL vehicle, local buildx/QEMU build and physical-Pi correctness/PMU baseline. **R4-R5 GREEN:** [KERNEL-CPU-A76-Q8-DOT](kernel-matrix.md) adds exact-order compiler SDOT and a scheduled AAPCS64 leaf; Pi operation and 64-token model outputs are exact. Assembly beats compiler SDOT 3.66-5.08% on M1/T1 and M128/T1/T4, lowers Qwen TTFT 1.55% and E2E 0.13%, so A76+DotProd auto-selects it. M1/T4 regresses 2.43%. **The separate Pi llama.cpp floor is now MEASURED/NOT MET on speed:** vllm.cpp is 0.461x prefill / 0.653x decode+E2E, while using 24.2% less RSS; exact-prompt 64-token output matches. No 20-core/i8mm result transfers. Thread partition, BF16 GEMM/speed closure and concurrency stay open. **macOS UNBLOCKED 2026-07-22 (`CLAIM-BACKEND-FANOUT-1`, W0):** the CPU registrar force-link is fixed; `test_backend` is **7/7 (18/18)** and the whole tree builds `-Werror`-clean on M4. Two unrelated macOS test gaps remain recorded in [environment.md](environment.md) | [threadpool leaf](specs/gguf-cpu-threadpool.md); [RPi5 Cortex-A76 campaign](specs/rpi5-cortex-a76-cpu-optimization.md); [Pi competitor evidence](../docs/bench-evidence/rpi5-a76-llamacpp-20260806.md) | `PARTIAL` | PR #65 | | `BACKEND-CPU-ZEN` | AMD Zen CPU with ZenDNN/zentorch dispatch and weight prepack | `platforms/zen_cpu.py:12-32`; detection `platforms/__init__.py:153-192`; `tests/test_zen_cpu_platform_detection.py:8-37` | generic CPU backend only; no Zen-specific dispatch | - | [CUDA inventory](specs/cuda-architecture-inventory.md); leaf spike required | `INVENTORIED` | - | | `BACKEND-ROCM` | AMD ROCm/HIP | `platforms/__init__.py:110-128`, `platforms/rocm.py:43-125` (incl. `_capability_from_gcn_arch:223-291`, `_get_backend_priorities:407`), `CMakeLists.txt:20-59,196-211` (HIP build), ROCm ops rooted at `csrc/rocm/moe_q_gemm_rdna3.cu:1` @ pin `555967922` | Lands through the realized seams: a `platforms/rocm.cpp` Platform + a `kROCM` attention TU self-registering via the [attn-registry seam](../include/vllm/v1/attention/registry.h#L44) + a `rocm.cpp` `get_attn_backend_priority()` slot (mirror `rocm.py`'s ROCM_ATTN/AITER ordering); ZERO selector/model/runner edit **AS-BUILT:** **M3 (2026-08-17, #1056) — the `kROCM` attention TU LANDED: `RocmAttentionBackend` (ROCM_ATTN) self-registers via the attn-registry seam + `rocm.cpp` `get_attn_backend_priority()` mirrors `rocm.py:407-441` verbatim; still ZERO engine edit. The runner CONSUMPTION of the seam (#1065 — per-group selection + per-group KV-shape validation) is a DIFFERENT concern, tracked as `BACKEND-ATTN-SELECTION-RUNNER` (see [rocm-attn-backend.md](specs/rocm-attn-backend.md)); it does NOT ride this row's landing contract.** **W0 skeleton landed 2026-08-06 **W0 skeleton landed 2026-08-06 and was COMMUNITY-VERIFIED (issue #41): compiled clean + `ctest -R 'rocm\|cross_device'` green on gfx1151, gfx1103, gfx1100 and gfx1201 — M0/M1 MET on all four (teardown caveats: TheRock-nightly exit hang on gfx1103; #132 `-O0` CLR race on gfx1100). W1 approach-(b) landed 2026-08-08, IMPLEMENTED-UNVERIFIED (blind, no AMD hardware here): managed-alloc branch on integrated managed-capable devices (`hipMallocManaged` + `UnifiedMemory()=true` by construction, resolving F6 where XNACK-less RDNA3 probes `PageableMemoryAccess=0`), introspection probes `ManagedAllocActive`/`IntegratedDevice`, two runtime-gated tests, and the F1/F3 `ROCM_PATH`→HIP-compiler-hints CMake absorption for Arch/TheRock layouts. Discrete path byte-identical (branch provably dead at `Integrated=0`). The two plain-C++ legs stay compile-gated everywhere (`vllm_rocm_platform_syntax_check`, object-only, never linked).** [device.h](../include/vt/device.h) `kROCM`; [rocm_arch.h](../include/vt/rocm/rocm_arch.h) (capability parse, HIP-free, 1:1 from `rocm.py:223-291`); [rocm_runtime.h](../include/vt/rocm/rocm_runtime.h); [rocm_backend.hip](../src/vt/rocm/rocm_backend.hip) (6 virtuals + multi-device registrar); [rocm_rmsnorm.hip](../src/vt/rocm/rocm_rmsnorm.hip) (port of `cuda_ops.cu:96-126`); [rocm_ops.hip](../src/vt/rocm/rocm_ops.hip) (1 of 106 OpIds); [platforms/rocm.cpp](../src/vllm/platforms/rocm.cpp) (M3: dense priority mirrors `rocm.py:407-441` verbatim — ROCM_ATTN registered, AITER/TRITON_ATTN/TURBOQUANT named placeholders); `VLLM_CPP_HIP` + `VLLM_CPP_HIP_ARCHITECTURES` + `ROCM_PATH` in [CMakeLists.txt](../CMakeLists.txt); and the ONE non-additive site this uncovered — [platform.cpp](../src/vllm/platforms/platform.cpp) `kCurrentPriority` (a platform absent from that walk registers fine and is NEVER selected, with no compiler diagnostic), now exposed via `CurrentPlatformPriority()` and gated | **VERIFIED on the CPU tier (no GPU):** [test_rocm_arch](../tests/vt/test_rocm_arch.cpp) 7 cases / 40 assertions incl. upstream's own worked examples + the gfx9 feature-suffix strip + `static_assert`s; [test_platform](../tests/vllm/platforms/test_platform.cpp) new case gating that EVERY `DeviceType` is in the selection walk and CPU is last; clean `-Werror` CPU build (the enum forced exactly ONE switch site tree-wide, `test_backend_cross_device.cpp:59`); full `ctest` green; `check-device-leakage` unchanged at 32. **COMMUNITY-VERIFIED (W0, #41 tables):** [test_rocm_backend](../tests/vt/test_rocm_backend.cpp) 5 cases/1044 assertions and the cross-device RmsNorm-vs-CPU-oracle comparison (NMSE <= 5e-4) green on all four boards. **PENDING-community (W1 (b), no hardware here):** the (b) delta's first HIP compile, the two new test cases (alloc-path/UnifiedMemory coupling; kernel-write→host-read no-copy), the flag-free Arch/TheRock configure, and the unblocked M2 reference-tier e2e on gfx1151/gfx1103. **gfx1200 (RX 9060 XT, discrete, RDNA4) INDEPENDENTLY VERIFIED 2026-08-10 ([#269](https://github.com/mudler/vllm.cpp/issues/269), off the #41 four-board set):** M0/M1 MET; M2 MET via the native-kernel path (no reference tier on a dGPU — `VT_OP_PROVIDER_STATS=1` shows zero fallbacks); **M4 MET for `Gemma3ForCausalLM` (gemma-3-1b-it), 48/48 tokens against TWO independent real vLLM-ROCm oracles** on this exact board — a prebuilt AMD `gfx120X` image and a from-source build at this project's own pinned commit `555967922` (compiled clean in ~6.5 min against `rocm/vllm-dev:base`, whose ROCm 7.2.3 matches this board's native build exactly). `Qwen3ForCausalLM` (Qwen3-0.6B) surfaced one genuine near-tie prompt where the two real oracles disagree with each other (each internally deterministic) — direct proof of a version-sensitive tie in the reference itself, not a defect in either backend. **W6 (2026-08-27, #2065) — fp8-e4m3 KV cache store+read LANDED:** `OpId::kReshapeAndCacheFp8` registered for `kROCM` + `LoadKv` fp8 dequant in the paged-attention read; gated on gfx1100 (7/7 cases, 28/28 assertions) **gfx1151 QUANT-MATCHED DECODE — THE 2.71x IS WITHDRAWN AND THE DENOMINATOR STANDS ALONE (2026-09-02, #2497):** the 4.510 tok/s reading of 2026-09-01 and the 2.71x / 2.88x derived from it are withdrawn from this row. They were taken ahead of this arm's declared token gate, which reads `TOKEN_GATE=FAIL` at 3 of 6 prompts ([token gate v2](../docs/bench-evidence/qwen38-27b-q4km-rocm-gfx1151-token-gate-v2-20260902.md)), and `AGENTS.md` §Gates admits no performance result from such an arm. What is measured, and measurable without our gate, is the oracle alone: llama.cpp `b10451` on `strix:gpu0` reads **12.233 tok/s**, median of 6 legs, 0.303% spread, on the identical byte-verified `Qwen3.8-27B-Q4_K_M.gguf`. No ratio is computed from it, and the paired run is staged and refuses to start until the gate is ratified. **The reliability finding is untouched by the withdrawal:** our arm died with `reason :GPU Hang` and exit 139 on **2 of 3 legs** on the plain Q4_K target, with no DFlash2 and no CPU reference-tier op, so #2377's candidates cannot be the whole cause; llama.cpp completed 3 of 3 then and 6 of 6 now on the same bytes (#2511). `Qwen3.8-27B-UD-Q4_K_M.gguf`, the artifact family every published Strix Halo row runs, does not load here at all, because the GGUF reader has no `IQ3_S` (#2510). No in-tree harness samples AMD clock state (#2381), so every clock figure on this board is ad-hoc and is recorded as such. [gfx1151 llama.cpp denominator](../docs/bench-evidence/rocm-strix-llamacpp-denominator-20260902.md), [gfx1151 Q4_K decode evidence](../docs/bench-evidence/rocm-strix-qwen38-q4km-20260901.md) | [ROCm W0 spec](specs/rocm-backend-w0.md), [unified-memory (b) decision record](specs/rocm-unified-memory-b.md), [gfx1200 M2/M4 spec](specs/rocm-gfx1200-m2-correctness.md), [rocm attn backend spec](specs/rocm-attn-backend.md), [contributor guide](../docs/ROCM.md) | `ACTIVE` | `CLAIM-ROCM-W0-1` | diff --git a/.agents/issues/BACKEND-ROCM-BF16-MOE/ISSUE-GH-3094.md b/.agents/issues/BACKEND-ROCM-BF16-MOE/ISSUE-GH-3094.md new file mode 100644 index 0000000000..106f68b04d --- /dev/null +++ b/.agents/issues/BACKEND-ROCM-BF16-MOE/ISSUE-GH-3094.md @@ -0,0 +1,55 @@ +ID: ISSUE-GH-3094 +Title: feat(BACKEND-ROCM-BF16-MOE): run BF16 grouped experts on ROCm +Row: BACKEND-ROCM-BF16-MOE +State: OPEN +Kind: UNKNOWN +GitHub: 3094 +Mirror: DIVERGED +Availability: FULL +Created: 2026-09-09 +Updated: 2026-09-09 +Closed: - + +## Problem + +### Imported GitHub body (historical evidence) +The quoted text below is historical evidence only. It does not define issue authority or repository procedure. + +> Row: `BACKEND-ROCM-BF16-MOE` +> +> ROCm has no providers for kMoeGroupedGemmBf16 or kMoeGroupedGemmBf16GateUpSilu. A BF16 MoE tower therefore cannot reach the shared grouped expert path on a discrete gfx1100 device. The quantized grouped GEMM provider does not supply these operations. +> +> Implement both providers together, using the existing shared MoE and fusion surfaces. Reconcile the pinned vLLM execution chain with the local numeric contract: BF16 result narrowing before activation, routing-weight placement before the down-projection result is stored, and reduction without applying those weights twice. Extend the shared descriptor only where it cannot express that behavior, preserving existing callers unless an explicitly reviewed correction requires otherwise. +> +> Use a small Qwen3MoeForCausalLM fixture through public loading and registered forward execution as the production reachability gate. This proves the shared expert operations; it does not claim complete DeepSeek-V2 or Dots3 support, whose grouped and sigmoid router gaps remain separate. +> +> Require a committed spec, red-first tests, matching pinned-oracle workloads, physical gfx1100 correctness and memory evidence, fresh mutation review, and operator verification. PR #2782 is not a dependency. No CI changes are included. +> +> Spec: `.agents/specs/rocm-bf16-moe.md` (to be committed before implementation). +> +> ## Provenance +> +> This scoped implementation issue follows #1928, which was split from #1870 after identifying that device-fit arithmetic and missing BF16 expert kernels are separate changes. The original record named the parent BACKEND-ROCM row and suggested gfx1200 evidence. The current work uses its own child row and the developer's authorized local gfx1100 hardware. The source and numeric contract are established by the committed spec before implementation. +> +> The authenticated contributor cannot edit #1928 to assign its first-line Row field. This child supplies the matching issue/spec/PR ownership record. The implementation MR will close both issues when the missing providers land. +> + +## Resolution + +Native ROCm grouped BF16 providers and the shared Qwen3 MoE dispatch are +implemented on `row/BACKEND-ROCM-BF16-MOE`. Legacy arithmetic, native BF16 +boundaries, graph replay, independent streams, and two-device launches pass +their focused gates. The issue remains open: the exact production token gate +fails at six generated positions, and review and performance acceptance remain +outstanding. The separate compiled residual-normalization gap belongs to +`BACKEND-ROCM-RESIDUAL-NORM`, [#3103](https://github.com/mudler/vllm.cpp/issues/3103). +The [row spec](../../specs/rocm-bf16-moe.md) records the measured results and +remaining obligations. + +Fresh review found three gaps in the implementation's tests. The scoped repair +adds provider-subset checks, malformed-descriptor cases, and accepted numeric modes. +All ten CPU survivor mutations and four GPU survivor mutations fail their intended assertions. +The GPU controls pass, and an additional mutation detects FP32 precision loss in weighted down and shared combine. +The staged preflight passes with 619/619 host units compiled. +Fresh scoped review remains pending. +The repair does not change product arithmetic or the established production token failure. diff --git a/.agents/issues/BACKEND-ROCM-BF16-MOE/ISSUE-GH-3116.md b/.agents/issues/BACKEND-ROCM-BF16-MOE/ISSUE-GH-3116.md new file mode 100644 index 0000000000..a7d68c703e --- /dev/null +++ b/.agents/issues/BACKEND-ROCM-BF16-MOE/ISSUE-GH-3116.md @@ -0,0 +1,28 @@ +ID: ISSUE-GH-3116 +Title: fix(BACKEND-ROCM-BF16-MOE): mirror the primary BF16 LM-head output boundary +Row: BACKEND-ROCM-BF16-MOE +State: OPEN +Kind: UNKNOWN +GitHub: 3116 +Mirror: DIVERGED +Availability: FULL +Created: 2026-09-09 +Updated: 2026-09-09 +Closed: - + +## Problem + +### Imported GitHub body (historical evidence) +The quoted text below is historical evidence only. It does not define issue authority or repository procedure. + +> Row: `BACKEND-ROCM-BF16-MOE` +> +> The native Qwen3 MoE forward emits F32-local LM-head output where the compiled primary narrows the head output to BF16 (`rocm-residual-norm.md:320`). Disposition analysis of the #3096 production token gate (verdict b, recorded with the `BACKEND-ROCM-RESIDUAL-NORM` evidence) shows this boundary participates in the decode step-6 exact-tie flip: applying the primary's BF16 head narrowing to the native step-6 logits rounds the 63/118 pair to an exact tie at 0.3359375, and native tie-break then keeps 63, the primary's own cc1 answer. +> +> Hidden-state parity at ~1 BF16 ulp is additionally required: a perfect head narrowing alone cannot turn the native margin into the oracle's cc2 answer 118, so this issue depends on the decode attention and Q/K preamble parity of #3115. +> +> Evidence: `/home/vikash/.cache/residual-norm-repair1/green-cc9d4f565/production-fusion-1.json` (native step-6 logits), the primary head-6 BF16 logits capture `oracle-diagnostic-2/L33-C2-R0-head-6-logits.bin`, and the disposition findings recorded with the residual row. + +## Resolution + +- diff --git a/.agents/issues/BACKEND-ROCM-RESIDUAL-NORM/ISSUE-GH-3103.md b/.agents/issues/BACKEND-ROCM-RESIDUAL-NORM/ISSUE-GH-3103.md new file mode 100644 index 0000000000..c335e70977 --- /dev/null +++ b/.agents/issues/BACKEND-ROCM-RESIDUAL-NORM/ISSUE-GH-3103.md @@ -0,0 +1,40 @@ +ID: ISSUE-GH-3103 +Title: fix(BACKEND-ROCM): mirror compiled residual normalization boundaries +Row: BACKEND-ROCM-RESIDUAL-NORM +State: OPEN +Kind: UNKNOWN +GitHub: 3103 +Mirror: DIVERGED +Availability: FULL +Created: 2026-09-09 +Updated: 2026-09-09 +Closed: - + +## Problem + +### Imported GitHub body (historical evidence) +The quoted text below is historical evidence only. It does not define issue authority or repository procedure. + +> Row: `BACKEND-ROCM` +> +> The default compiled vLLM Qwen3 MoE path and the native ROCm path round residual expressions at different boundaries. This is a shared normalization gap found while gating #3094. It predates the new BF16 expert providers. +> +> The reproducer uses the fixed `Qwen3MoeForCausalLM` fixture from `.agents/specs/rocm-bf16-moe.md`: hidden width 128, two layers, four experts, and top-k 2. The pinned vLLM revision is `e126687a9a828d513c01a07cd69f025f27d63280`, with default compilation, production graphs, ROCM_ATTN, and Triton experts. The native model differs at six generated positions across three repeated length-33, concurrency-2 runs. This issue proves one earlier numerical difference and does not claim that correcting it alone resolves the token gate. +> +> For layer 0, token 0, the actual QKV and attention projection outputs match exactly. Both normalization calls receive identical BF16 attention output, embedding residual, and gamma. The normalized results differ in 33 of 128 values. The first difference is element 3: native `-0.1865234375`, oracle `-0.185546875`. The same witness exists at concurrency 1 and 2. An independent operator rerun of the captured-input FP32 algebra reproduces all 128 native values with residual rounding and all 128 oracle values without that rounding. +> +> `src/vt/rocm/rocm_rmsnorm.hip:125–147` rounds the residual sum to its stored BF16 dtype before variance, then reloads that stored sum. The pinned IR at `vllm/ir/ops/layernorm.py:44–62` computes variance from the FP32 sum. The executing generated kernel confirms this distinction. +> +> The complete compiled lifetime matters. The layer-0 post-attention kernel keeps attention and embedding operands separately and does not store their intermediate residual sum. The next input normalization recomputes `MoE + (attention + embedding)` in FP32, then materializes its residual in BF16. The final layer similarly recomputes the two additions for final normalization. Every allocated residual buffer remains BF16. Blanket widening of residual storage would not mirror this chain. +> +> The complete handoff is `/home/vikash/vllm.cpp-rdna3-moe-impl/build-rdna3-moe-hip/evidence/norm-gap-handoff/handoff.md`, SHA256 `9850ee49c93bb82ddfca1a811fb422337c33ea09335e749fbf3c42dfcaa0400a`. Its manifest records the six generated modules, executed AOT artifact, exact source paths, and hashes. The operator independently checked all 14 manifest entries. Paired raw captures, `residual-row0-diagnostic.json`, and `residual-row0-operator-receipt.json` retain the arithmetic proof. +> +> The owning backend row owes a separate committed spec, fresh implementation, and independent review for the complete lifetime through shared fusion seams. Preserve physical dtype evidence, production graph defaults, and the fixed workload. Port the applicable upstream normalization tests and mutate the intermediate rounding boundary and production wiring. The earlier attention-output difference and local F32 LM-head output remain separate gaps. Keep the MoE token gate failing until its required mismatches are resolved. +> + +## Resolution + +The child row `BACKEND-ROCM-RESIDUAL-NORM` owns the repair under parent `BACKEND-ROCM`. +The [committed repair spec](../../specs/rocm-residual-norm.md) defines the ordered residual expression and production tests. +The GitHub issue carries the same child row. The shared repair passes its first production witness on the physical GPU. +The complete hardware gate, upstream fixture comparison, and fresh mutation review remain pending in the repair spec. diff --git a/.agents/issues/BACKEND-ROCM/ISSUE-GH-3100.md b/.agents/issues/BACKEND-ROCM/ISSUE-GH-3100.md new file mode 100644 index 0000000000..13237c959e --- /dev/null +++ b/.agents/issues/BACKEND-ROCM/ISSUE-GH-3100.md @@ -0,0 +1,33 @@ +ID: ISSUE-GH-3100 +Title: fix(BACKEND-ROCM): bind resource operations to the requested device +Row: BACKEND-ROCM +State: OPEN +Kind: UNKNOWN +GitHub: 3100 +Mirror: DIVERGED +Availability: FULL +Created: 2026-09-09 +Updated: 2026-09-09 +Closed: - + +## Problem + +### Imported GitHub body (historical evidence) +The quoted text below is historical evidence only. It does not define issue authority or repository procedure. + +> Row: `BACKEND-ROCM` +> +> ROCm resource methods use the ambient HIP device instead of their recorded device index. At base `6db4bef906859e864c82523c01107473f7dcca29`, `src/vt/rocm/rocm_backend.hip:220-284` calls `hipMalloc`, `hipStreamCreate`, `hipMemsetAsync`, `hipMemcpyAsync`, and `hipStreamSynchronize` without selecting `device_` or the queue device. The backend registers no `DeviceResourceOps` override for `vt::Alloc(Device, ...)`. The introducing source history is `50b0709b3`. +> +> A native MoE boundary test on two local gfx1100 devices exposed this assumption. Device-1-labeled buffers were constructed while device 0 remained current. The operator run returned 1 with three failed comparisons in the null-default-stream case. The other six cases passed. Untouched `0x5555` buffers and invalid reference values identify invalid test resource placement, not a valid native-kernel numerical result. An isolated allocation-device probe remains required before accepting the precise runtime diagnosis. +> +> Owner: `BACKEND-ROCM`, resource and queue device ownership. The new MoE component test must explicitly allocate, copy, and build its reference on each intended device before testing native operation device selection. That scoped harness adaptation does not establish full-model execution on multiple GPUs. +> +> Acceptance requires a committed spec, a red test through the shared resource API with a deliberately different ambient device, verified pointer and stream ownership, preserved ambient-device restoration, and correct default-stream copy and synchronization. Cover two devices, host threads, existing single-device behavior, and graph lifetimes. Run the pinned primary oracle for applicable resource behavior, fresh mutation review, and operator GPU gates. Keep native provider and resource ownership claims separate. +> +> Local reproduction: `/home/vikash/vllm.cpp-rdna3-moe-impl/build-rdna3-moe-hip/evidence/native-boundary-1.log` and its operator receipt. Binary SHA256: `124546db69241940fd5ee88fb26cf71d6053a1dab7ad425915d0422dfc4854f2`. The run held `/home/vikash/gpu.lock`, exposed local devices 0 and 1, and ran on 9 September 2026 UTC. +> + +## Resolution + +- diff --git a/.agents/issues/_owed/ISSUE-GH-1928.md b/.agents/issues/_owed/ISSUE-GH-1928.md index 63a4c1064b..b73e097bd0 100644 --- a/.agents/issues/_owed/ISSUE-GH-1928.md +++ b/.agents/issues/_owed/ISSUE-GH-1928.md @@ -66,4 +66,9 @@ The quoted text below is historical evidence only. It does not define issue auth ## Resolution -- +The native provider implementation is tracked by +[#3094](https://github.com/mudler/vllm.cpp/issues/3094) and +[its committed spec](../../specs/rocm-bf16-moe.md). Both legacy provider IDs +are registered on ROCm, with separate typed native BF16 modes for the pinned +oracle's arithmetic. This issue remains open until the reviewed implementation +lands. The complete production token and performance gates are not yet satisfied. diff --git a/.agents/specs/gguf-device-fit-expand-policy.md b/.agents/specs/gguf-device-fit-expand-policy.md index 9a2c5254af..c1fc01fbaf 100644 --- a/.agents/specs/gguf-device-fit-expand-policy.md +++ b/.agents/specs/gguf-device-fit-expand-policy.md @@ -225,6 +225,8 @@ box, and the real-hardware attempt that surfaced #1934. - **ISSUE-GH-1928** — the ROCm `kMoeGroupedGemmBf16` provider gap #1870's "related gap, same area" section names. Not fixed in flow; scoped out above. + Implementation handoff: [native BF16 grouped MoE](rocm-bf16-moe.md), + row `BACKEND-ROCM-BF16-MOE`, tracked by issue #3094. - **#1934** — `RocmPlatform::needs_weight_staging()` is stale-false, so this row's fix (and the refusal it corrects) never runs on ROCm. THE ACTUAL remaining blocker for #1870's reproduced crash. Not fixed in flow: it is a diff --git a/.agents/specs/rocm-attn-parity.md b/.agents/specs/rocm-attn-parity.md new file mode 100644 index 0000000000..6cb6dff16d --- /dev/null +++ b/.agents/specs/rocm-attn-parity.md @@ -0,0 +1,507 @@ +# ROCm attention parity: the value-dtype probability and the FP32 Q/K carrier + +Owning row: `BACKEND-ROCM-BF16-MOE` +Issue: [#3115](https://github.com/mudler/vllm.cpp/issues/3115), which asks for the +decode attention and Q/K preamble parity of the BF16 MoE row. +Base: `cd1ab3909a25d58a9676afa08263366b09ee0cb8` on +`row/BACKEND-ROCM-BF16-MOE-attn-parity`. +Integration: one pull request, with this spec committed before any product edit. +The test-only replay instrument `e5c7bfc4d` is already at the base and is the +measurement this repair answers. + +## Now + +State: `ACTIVE`. Both repairs are in and measured on native gfx1100 at layer 0 +step 0 of the L33/C2/R0 workload: row A0 (native output vs the primary's 2902 +words) is 0, row B (post-RoPE Q/K) is 1 Q / 0 K, and row C (the kernel on the +primary's own Q/K/V) is 0 of 8448. The fresh review then found four defects in the +repair, all repaired here: the op-layer stride check that made the `VT_FUSED_CHAIN_ADOPT=0` +hand-call throw for every `Dh > 1`, the fixed 64-key tile that ignored the arm the +primary would execute, a focused case that could not tell a 32-key tile from a +64-key one on this capture, and an instrument that exited mid-run. `## Design 1` +carries the arm table; `## Owed` names the arms that are implemented but have no +capture to measure them against. + +## Problem and scope + +Two arithmetic differences produce the 2902-word parity gap on native gfx1100 at +layer 0 step 0 of the L33/C2/R0 workload (66 tokens, `hq = hkv = 1`, `dh = 128`, +`scale 0.08838834764831845`). + +1. **The kernel keeps the softmax probability in f32.** The primary narrows it to + the value dtype before the value dot and keeps only the running sum in f32. + `PagedAttnDecodeOptBf16T` accumulated `o_reg = o_reg * corr + pw * v_reg` with + `pw` in f32 (pre-repair site `src/vt/rocm/rocm_paged_attn.hip:527-531`, now + `:636-640`). +2. **The preamble narrows the normalized Q/K to BF16 before RoPE.** The primary + carries the normalized value in f32 through the rotation and narrows once, at + the store, reading a BF16 cos/sin cache. The native preamble rounds twice. + +Both repairs are mirroring repairs. Neither invents behavior: each replaces a +native rounding decision with the primary's own. + +**In scope.** `src/vt/rocm/rocm_paged_attn.hip` and the ROCm attention preamble, +plus the test cases that witness them, plus this spec. + +**Out of scope.** The CUDA, Metal and CPU realizations of the same two +arith-metics; the two rocWMMA attention kernels; the Tier-0 fused-chain composite +for `kAttnQkNormRope`. Each is named under `## Sibling sites` or `## Owed` with +its reason and is not silently changed. + +## Primary anchors + +Pin `e126687a9a828d513c01a07cd69f025f27d63280` +(`.agents/upstream-sync.md` §"parity-pin"); source tree +`/home/vikash/oracle/gfx1100-active-2773/source`. + +The executing primary arm for this workload is Triton, not the ROCm custom +kernel: `vllm/platforms/rocm.py` gates its decode predicate on `gqa_ratio >= 3` +and this fixture has `qg == 1`, so `chunked_prefill_paged_decode.py` logs the +fallback and a 33-token request (`max_query_len > 1`) enters +`prefix_prefill.py::context_attention_fwd`. + +| fact | anchor | +|---|---| +| prefill softmax probability, f32 | `vllm/v1/attention/ops/prefix_prefill.py:443` | +| prefill running sum taken BEFORE the narrowing | `prefix_prefill.py:445` | +| prefill probability narrowed to the value dtype | `prefix_prefill.py:471` | +| prefill value dot accumulates the narrowed probability | `prefix_prefill.py:473` | +| prefill denominator update keeps the f32 sum | `prefix_prefill.py:475` | +| prefill epilogue divides by the denominator | `prefix_prefill.py:478` | +| decode probability narrowed inside the dot | `chunked_prefill_paged_decode.py:265` | +| decode running sum taken BEFORE the narrowing | `chunked_prefill_paged_decode.py:251` | +| Qwen3 q/k norm then RoPE | `vllm/model_executor/models/qwen3.py:150-167` | +| the cos/sin cache is narrowed to the query dtype | `vllm/model_executor/layers/rotary_embedding/base.py:105-131` | +| the compiled preamble reads that BF16 cache | `RotaryEmbedding.forward_static`, `base.py:150-190` | + +The last three rows are why the primary's boundary is +`bf16(rope_f32(rmsnorm_f32(bf16 qkv)))` against a BF16 cos/sin cache. The +validated CPU transcription of that boundary +(`/home/vikash/.cache/attn-parity-plan/check_attention_models.py`, model A) +reproduces the primary's captured output **exactly, 0 of 8448 words**, and its +preamble model reproduces the primary's captured Q to **1 of 8448 words**. Those +two numbers license the CPU model as the primary-arithmetic reference without +executing the primary, and they are the targets this repair is measured against. + +## Measurement this repair starts from + +Operator-verified and independently reproduced, identical in +`/home/vikash/.cache/moe-attn-parity/gpu-run-at-head.log` and +`operator-rerun.log`; 1322 of 1322 assertions, 1 case, 3 skipped. Native gfx1100, +L33/C2/R0 layer 0 step 0. + +| row | meaning | words / 8448 | +|---|---|---| +| A0 | native run output vs primary output | 2902 | +| B | native run post-RoPE Q / K vs primary Q / K | 1569 / 1542 | +| C0 | replay self-consistency, native Q/K/V at the primary's geometry | 0 | +| C | native `PagedAttention` on the PRIMARY's own Q/K/V vs its output | 2918 | +| D | native `RmsNorm`+`RopeFromCache` on the primary's qkv vs primary Q/K | 1569 / 1542 | +| D' | the same through `RopeNeox` | 1776 / 1719 | +| E | native cos/sin table vs primary `cos_sin` | 0 | +| F | primary qkv q/k slices vs native pre-norm q/k | 0 | + +C0 is the guard: the replay reproduces the native run's own output byte-for-byte +at the primary's cache geometry, so row C measures the kernel and not the +transplant. E and F exclude the cos/sin source and the projection, so the +preamble term is exactly the pre-RoPE BF16 store. + +### Where this head stands + +Native gfx1100, the same workload, the same instrument, after both repairs and +after the fresh review's four findings were repaired. `VT_FUSED_CHAIN_ADOPT=1` and +`=0` produce byte-identical dumps; 8 of 8 cases pass at either setting. + +| row | meaning | before | at this head | +|---|---|---|---| +| A0 | native run output vs primary output | 2902 | **0** | +| B | native run post-RoPE Q / K vs primary Q / K | 1569 / 1542 | 1 / 0 | +| C0 | replay self-consistency at the primary's geometry | 0 | 0 | +| C | native `PagedAttention` on the PRIMARY's own Q/K/V | 2918 | **0** | +| D | composite preamble (RmsNorm+RopeFromCache) on primary qkv | 1569 / 1542 | 1569 / 1542 (not the executing path) | +| D2 | production fused op on the primary's qkv | — | 1 / 0 | +| D3 | hand-call `vt::AttnQkNormRope` on the primary's qkv | — | 1 / 0, and 0 words against D2 | +| E | native cos/sin table vs primary `cos_sin` | 0 | 0 | +| F | primary qkv q/k slices vs native pre-norm q/k | 0 | 0 | + +Row A0 at 0 is the row's target: the native output now reproduces the primary's +captured attention output byte-for-byte on the L33/C2/R0 layer-0 step-0 workload, +with the production token gate unmoved (`[66,1,70,57,33,81,63,69]` at length 33, +concurrency 2, request 0, and the same 27 records as before). + +## Design 1 — the value-dtype probability, at the primary's reference max + +### The literal one-line mirror is FALSIFIED + +The first implementation of this design narrowed `pw` at the accumulate, exactly +as `prefix_prefill.py:471` narrows `p`. Measured on gfx1100 at the same workload, +row C moved **2918 -> 3011 of 8448 words**: the literal mirror makes parity worse. + +It was falsified on CPU first, and the CPU model is exact for this workload. The +preserved transcription in `/home/vikash/.cache/attn-parity-plan/` is a faithful +model of the native kernel: its warp-online variant reproduces the device's +2918 exactly. Parameterized (`/home/vikash/.cache/moe-attn-parity/check_variants.py`, +`diag.py`), it gives: + +| form | words / 8448 | +|---|---| +| tile max, narrow | **0** | +| tile max, no narrow | 2917 | +| warp-online, no narrow (the device's 2918) | 2918 | +| warp-online, narrow (the literal mirror) | 3011 | +| sequential-online, narrow | 2058 | + +The second column is the reason. Narrowing is only the primary's boundary when +the exponent's reference max is the same one the primary used. `p` is narrowed +at the scale of `qk - m_ij`: the primary's `m_ij` is +`maximum(m_i, tl.max(qk, axis=1))` (`prefix_prefill.py:442`) — a running max +advanced once per key tile, 64 keys on the measured arm — while the native kernel +advanced a running max **once per key, per warp** (pre-repair site +`rocm_paged_attn.hip:525-533`, now `:636-640`). A bf16 rounding is +relative, so narrowing `exp(qk - m_warp_running)` and then rescaling by +`exp(m_warp - gm)` at the combine is not the same rounding as narrowing +`exp(qk - m_ij)`. For a row whose keys fit one tile the two scales coincide, which +is why the tile-max model is exact at 0 — and why this workload's rows, at 33 +keys, are the case where getting the scale right is worth the whole 2918. + +Keeping the literal mirror would ship a change that measurably worsens the row's +metric, so it is reverted, not committed. + +### The repaired design + +Mirror the primary's rule together with the reference max it is taken against: + +* the value accumulate multiplies a probability narrowed to the **value dtype**; +* that probability is `exp(s - m_ij)` where `m_ij` is the running max advanced + **once per contiguous key tile**, exactly the primary's + `maximum(m_i, tl.max(qk, axis=1))` (`prefix_prefill.py:442`, + `chunked_prefill_paged_decode.py:244`); +* `corr` (the primary's `alpha`) stays f32 — `prefix_prefill.py:449`, `:258`; +* the running sum keeps the un-narrowed f32 probability — `prefix_prefill.py:445` + against `:471`, `chunked_prefill_paged_decode.py:251` against `:265`; +* the warp-combine rescale stays f32, because the primary's `acc = acc * alpha` is + f32 at `:449`/`:258` and has no narrowed analogue. + +### The tile is the primary's per-row arm, not a constant + +The first repair fixed the tile at 64 keys. That is the tile of ONE arm, and the +same native kernel serves three: `chunked_prefill_paged_decode.py:317` dispatches +the call on the batch's max query length and the decode kernel it launches returns +for a row with `query_len > 1` (`filter_by_query_len=True`, `:93-97`, `:503`), so +the arm is decided per **row**: + +| arm (per row) | primary kernel | context tile | chunk / decode tile | +|---|---|---|---| +| `query_len > 1`, empty context | `prefix_prefill._fwd_kernel` | — | `BLOCK_N` = 64, anchored at the chunk start | +| `query_len > 1`, cached context | `prefix_prefill._fwd_kernel` | `TRITON_BLOCK_SIZE` = 32, anchored at key 0 | `BLOCK_N` = 64, anchored at the chunk start | +| `query_len == 1` | Triton `kernel_paged_attention_2d` | — | `min(block_size, 128)`, anchored at key 0 | +| any, non-power-of-two physical block size | both of the above | 32 | 32 | + +Anchors and widths are read from the launcher, not chosen: `prefix_prefill.py:955-966` +(`BLOCK_M`/`BLOCK_N`), `:965` (`TRITON_BLOCK_SIZE = 32`, bound as the kernel's +`BLOCK_SIZE` at `:1007`), `:231-343` (the context loop from key 0: its running max +at `:312`, the narrowed probability at `:338`, the value dot at `:340`, and lanes +at or past `cur_batch_ctx_len` masked at `:288-290`) and `:369` (the chunk loop +from the chunk start); +`chunked_prefill_paged_decode.py:444-445` (`TRITON_BLOCK_SIZE = min(block_size, 128)`, +or 32 when the physical block size is not a power of two), `:147-149` (`start_n = j * +BLOCK_SIZE`) and `:244`. A key outside the causal or sliding-window bound is +skipped rather than dropped from the grid: the primary masks such a lane to `-inf` +(`prefix_prefill.py:288-290`) or `-10000` (`chunked_prefill_paged_decode.py:236`), +neither of which can raise a tile max, and `exp(x - m)` underflows to 0 for either. +Tiles wholly below the window's left bound are skipped so a windowed model keeps its +`O(window)` walk. + +Concretely, the key walk in `PagedAttnDecodeOptBf16T` becomes per-tile two-phase +over the arm's key grid: each of the eight warps computes its strided keys' scores +for the tile, the tile max is reduced across the CTA (one `__syncthreads` per tile, +where the pre-repair loop anchored at `jmin` and had none), then every warp applies +the shared `m_ij`, the shared `alpha = FastExp(m_prev - m_ij)`, and accumulates +`ProbInValueDtype(p) * v_reg` with the f32 `p` in `lsum`. The key-to-warp map +stays the tile-and-warp-strided one the kernel already uses, so all eight warps +stay busy on a 33-key row and the loop keeps its shape. A tile narrower than +`kDecWarps` leaves the extra warps with no keys and an all-`-inf` partial max, which +cannot raise the CTA max. + +One helper expresses the narrowing for both value dtypes the file uses +(`src/vt/rocm/rocm_paged_attn.hip` uses only `float` and `__hip_bfloat16`; there +is no f16 KV path in the file): + +```cpp +template +__device__ inline float ProbInValueDtype(float p) { + if constexpr (std::is_same_v) { + return __bfloat162float(__float2bfloat16(p)); // p.to(v.dtype) + } else { + return p; // p.to(f32) is the identity + } +} +``` + +Narrowing to f32 is the identity, so a site templated on the value dtype is +correct for both arms and no dispatch is added. + +**Cost, stated because it is real.** The tile max needs the tile's scores before +the exponent, so the loop needs one `__syncthreads` per tile and a score register +per key per warp. The registers are sized for the widest tile any arm asks for +(128 keys, `min(block_size, 128)`), and the measured cost of that on the compiler's +own resource report is nil: `PagedAttnDecodeOptBf16T<4>/<8>/<16>` go from 57/65/93 +VGPRs before this change to 52/56/92 after, with 0 spill bytes and 16 waves/SIMD +occupancy on gfx1100 either way. What the finer decode tile DOES cost is one +synchronization per `min(block_size, 128)` keys instead of one per 64, so a model +whose physical block size is 16 or 32 now pays twice or four times as many CTA +synchronizations on the decode arm — the primary's own structure and the price of +the mirror. This change makes no performance claim and the decode arm's throughput +must be measured before the arm is accepted as free. + +### Sibling sites in `src/vt/rocm/rocm_paged_attn.hip` + +Line numbers are the POST-repair file's (`794298535` + the arm repair); the +pre-repair table named the same kernels 62-205 lines earlier. + +Every site below needs both halves of the repaired design: the value-dtype +probability **and** the reference max it is narrowed against, at the arm's own +tiles. This change implements the pair at the measured site only; the rest are +owed, because each is a separate kernel rewrite of the same shape and none of them +is measured on this workload. The literal one-line narrowing was measured at the +measured site and made parity worse, so it is not applied anywhere. + +| site | kernel | value dtype | disposition | +|---|---|---|---| +| `:638` | `PagedAttnDecodeOptBf16T` | bf16 | **changed** — the measured site; the arm table above; d=128 at `:2356`, d=256 at `:2363`, d=512 at `:2370` | +| `:221` | `PagedAttnOnline` | `TKV` | **owed** — the generic fallback, and the primary's arm for neither measured dtype | +| `:687` | `PagedAttnDecodeGqaBf16` | bf16 | **owed** — the QG-fused decode sibling (`:2327`, `:2334`, `:2341`); needs the same arm-tiled walk | +| `:830` | `PagedAttnDecodeGqaF32Q` | `TKV` | **owed** — the f32-query / bf16-KV arm (`:2422`, `:2430`, `:2440`, `:2448`) | +| `:982` | `PagedAttnPrefillFlashTile` | bf16 | **owed** — scalar online-V branch of the flash tile | +| `:1513` | `PagedAttnPrefillSharedK` | `TKV` | **owed** — the scoreless shared-K prefill | +| `:1283` | `PagedAttnPrefillWmmaWave` | bf16 | **owed** — body compiles only under `VT_ROCWMMA_OK` (`:9`, `:1292`, gfx1200/1201), so on gfx1100 neither the edit nor its result can be compiled or run here | +| `:1754` | `PagedAttnPrefillSharedKWmma` | bf16 | **owed** — same guard (`:1763`), and the host admits it only through `PrefillSharedKWmmaHostOk()` on gfx1200/1201 (`:41`, `:2193`) | + +The last two are also both reachable only from default-OFF lab toggles +(`VT_ATTN_PREFILL_FLASH`, `VT_ATTN_PREFILL_SHAREDK_WMMA`). Every owed site is +recorded under `## Owed`, not silently skipped. + +**Not a site.** `FastExp` (`:197`) is left alone: the CPU model shows `fast_exp` +and `np.exp` both reach 0 in the tile form, so it is not on this path's critical +term. The reciprocal epilogue (`:676`) is a 2^-24-scale difference against the +primary's division and is likewise not the term. + +## Design 2 — the FP32 Q/K carrier through RoPE + +The native preamble today is `RmsNorm` into a BF16 buffer followed by +`RopeFromCache` over that buffer (`include/vt/recipes.h:357-385` binds the 2-D +norm view and the 3-D rope view to the same BF16 buffer; the hand-call at +`include/vllm/model_executor/models/dense_attn_block.h:642-648` is the same +sequence). On ROCm the fused branch at `dense_attn_block.h:582-613` executes that +composite and not a fused kernel, because at the pre-repair head +`src/vt/rocm/rocm_ops.hip:321` registered `OpId::kAttnQkNormRopeGate` and +**not** `OpId::kAttnQkNormRope` (the repair adds the registration at `:331`), so +`vt::FusedChain` takes its `OpRegistered` guard (`src/vt/ops.cpp:1448-1452`) and +falls through to `FusedChainComposite` (`src/vt/ops.cpp:1468`). + +**Design: register the recipe's own fast op on ROCm, with the f32 carrier.** + +`vt::OpId::kAttnQkNormRope` is registered for `DeviceType::kROCM`, implemented by +`AttnQkNormRopeKernelRocm` in `src/vt/rocm/rocm_ops.hip`. The kernel: + +1. loads the raw `q3`/`k3` row (BF16), the per-head norm weight (BF16) and the + cos/sin cache row (BF16 at the position `positions[token]` supplies); +2. computes `mean(x^2)` over the head with the **same** f32 accumulation and the + same `kBlock` binary-tree shared reduction the shipped `RmsNormRowKernel` uses + (`src/vt/rocm/rocm_rmsnorm.hip:118-147`), then + `inv = 1 / sqrtf(mean + eps)`; +3. forms `v * inv * w[j]` in f32, rotates the first `rot` elements in f32 exactly + as `RopeFromCacheK` does (`src/vt/rocm/rocm_dense_basic.hip:699-704`, NeoX and + GPT-J pair orders both), and narrows **once**, at the store, with the bf16 + store helper the file already uses; +4. leaves elements at or beyond `rot` normalized and narrowed once, which is the + same BF16 word the shipped `RmsNorm` store produced. + +The kernel is modelled on the registered sibling `AttnQkNormRopeGateK` +(`src/vt/rocm/rocm_gdn_fused.hip:95-169`), which is already a fused +norm-plus-partial-NeoX-RoPE preamble on this backend. The `attn_f32` arm passes +f32 states, an f32 norm weight and the f32 cache; for that arm the kernel's +arithmetic is the shipped `RmsNorm` + `RopeFromCache` arithmetic with no +intermediate rounding at all, so its bytes are unchanged by construction. + +**Both realizations must agree.** With the op registered, `vt::FusedChain` takes +the fast path for the fused branch, and the hand-call fallback at +`dense_attn_block.h:645-653` is changed to dispatch to the same +`vt::AttnQkNormRope` when the op is registered on the device and the bf16 cache +is in use. On a backend that registers no fast op (CPU), and on the `RopeNeox` +default branch (no cache), both realizations keep exactly today's sequence. So +the fused path and the hand-call fallback agree byte-for-byte on every backend, +which is what the recipe's byte-exact composite contract asks for +(`include/vt/recipes.h:378-383`) and what `tests/vllm/models/test_qwen3_forward.cpp` +already gates for the adoption switch. + +**The Tier-0 composite stays as it is, and is owed.** Carrying f32 through the +composite would require an operand slot that holds an f32 intermediate +(`kMaxFusedOperands` is 8, `include/vt/fused_recipe.h:107-108`, and the recipe +already spends all eight) plus a `RopeFromCache` that reads an f32 state against +a BF16 cache. That is a recipe and op-contract change on four backends, which +this row does not own. On ROCm the composite is no longer the executing path +once the fast op is registered; every other backend keeps today's bytes. This is +recorded under `## Owed`. + +## Risks + +* **This changes the numerics of existing users of the six attention kernels.** + Every ROCm paged-attention arm whose value dtype is bf16 moves toward the + primary and away from its previous bytes. Committed ROCm device goldens and + near-tie anchors on those arms may need re-derivation; the production gate + `test_rocm_moe_bf16` is re-run in this change and any moved golden is reported. +* **This changes the numerics of the ROCm BF16 Qwen3-dense preamble** for every + model with a qk-norm and the cos/sin cache (which is on by default, + `dense_attn_block.h:92-98`). The `attn_f32` arm is unaffected by construction. +* **The kernel term may not reach exactly zero.** The primary's value dot is a + tile `tl.dot` with its own accumulation order; the native arm keeps `FastExp`, + a warp-strided key walk and a reciprocal epilogue. A residual of a few words is + possible and must be attributed before any further edit. +* **A greedy anchor can move.** The change alters bf16 words; a token is decided + by the full logits, so movement is possible in principle. The gate's recorded + tokens are checked and any movement is reported rather than papered over. +* **The decode arm's tile follows the physical block size, so a model with a + block size other than 64 changes bytes on every decode step.** That is the + mirror, not a side effect: the primary's Triton decode kernel tiles by + `min(block_size, 128)`, so the native arm tiled by 64 was narrowing against a + reference max the primary never used whenever `block_size != 64`. It also costs + one CTA synchronization per tile, so a block size of 16 or 32 pays more of them + per decode step. `tests/vt/test_ops_paged_attn.cpp` and the production gate are + re-run here and neither moved; a model whose committed decode golden was captured + with `block_size != 64` would need its golden re-derived, and none was found. +* **Performance.** The preamble repair replaces three launches per layer with + one; the kernel repair adds one bf16 round per key and one synchronization per + tile. Neither is a perf claim and no perf axis is accepted by this change. + +## Tests + +**Red first, focused, in the committed instrument** +(`tests/vllm/models/test_rocm_moe_bf16.cpp`, case "ROCm paged attention replays +the primary's captured attention boundary"): + +* kernel: `CHECK(replay_diff.different == 0)` — red at 2918 before Design 1, + and still red at 3011 under the falsified literal mirror; +* preamble: `CHECK(run_q.different <= 1 && run_k.different <= 1)` — red at + 1569/1542 before Design 2; +* preamble on the primary's own qkv through the production op, which must reach + the CPU model's 1 word — red before Design 2; +* hand-call (D3): `vt::AttnQkNormRope` on that same primary qkv must reach the + same 1 word and must equal the fused realization byte-for-byte — red (it threw) + before the op-layer stride repair. + +The instrument gains an observer on `OpId::kAttnQkNormRope` for the pre-norm and +post-RoPE native bytes, because after Design 2 the production preamble no longer +calls `kRopeFromCache`; the existing `kRopeFromCache`/`kRmsNorm` observers stay +and now measure the composite fallback. + +**The arm's tile width, on synthesized data.** The capture cannot witness the tile +choice: its 33-token rows are all `prefix_prefill` with an empty context, so +`max(keys 0..31) == max(keys 0..32)` and a 32-key tile reproduces the same 8448 +bytes. Three device cases build a workload where the width IS load-bearing — each +key's value row is a single bf16-exact integer in its own output lane, `q = e0`, +and each K row is 0 or a bf16-exact high value at one key, so a key that shares a +tile with the high key is narrowed at a different scale than the same key alone in +a tile. Each case compares the device against a host transcription of the primary's +own key walk for that arm AND against the same transcription under the neighbouring +widths, so the case fails in both directions: + +| case | arm | device vs the arm's tiles | device vs the neighbouring tilings | +|---|---|---|---| +| "uses the primary's 64-key prefill tile" | `query_len` 65, empty context, block 16 | 0 / 33280 words | 1408 words at 32 and at 16 | +| "uses the primary's 32-key context tile" | `query_len` 3, 37-key context, block 16 | 0 / 1536 words | 192 words at a uniform 64 | +| "uses the primary's decode tile" | `query_len` 1, `seq_len` 40, block 16 | 0 / 1024 words | 64 words at 32 and at 64 | + +The thresholds are a noise floor of 8 words (the host/device difference in the +order the f32 running sum is accumulated) against a signal floor of 32, and the +measured separation is 0 against 64 at its narrowest. + +**The adoption switch, in one process.** Case "ROCm bf16 qk-norm-rope: the +hand-call realization matches the fused recipe" runs the recipe's fast realization +(`vt::FusedChain`, what `VT_FUSED_CHAIN_ADOPT=1` executes) and the hand-call the +fallback branch reaches (`VT_FUSED_CHAIN_ADOPT=0`) over one synthesized q/k and +requires byte identity, so the documented same-binary A/B lever is a gate and not +a claim. The instrument's own case additionally pins, per run, that the registered +op executed and that the Tier-0 composite did not. + +**Mutation.** Each repair is reverted in place, its focused case must redden, and +the file is restored with a sha256 check. + +**Suites.** The instrument's CPU-only case (no device), the instrument's device +case under `flock /home/vikash/gpu.lock` with `HIP_VISIBLE_DEVICES=0` and both +settings of `VT_FUSED_CHAIN_ADOPT`, `tests/vt/test_ops_paged_attn.cpp`, and the +production gate `tests/test_rocm_moe_bf16` with `VT_ROCM_MOE_FIXTURE`, +`VT_ROCM_MOE_ORACLE` and `VT_FUSED_CHAIN_ADOPT=1`. + +**The instrument's own hygiene.** Every environment-gated case is decorated +`doctest::skip(...)`, so an absent variable reports the case skipped in doctest's +summary and the remaining cases still run; the process exits 77 only when +something was skipped and nothing failed. + +## Gates and evidence + +* Focused red and green runs, with the exact command and exit status. +* The device run's per-row numbers for C, B, D, D2 and D3, in the same table shape + as `## Measurement this repair starts from`. +* The two `VT_FUSED_CHAIN_ADOPT` runs' dumped binaries compared byte-for-byte, with + the adoption flag the only differing field in the report. +* The production gate's recorded tokens, compared to + `[66,1,70,57,33,81,63,69]` at length 33, concurrency 2, request 0. +* `TMPDIR=/home/vikash/.cache/moe-attn-parity/tmp-preflight + GIT_CONFIG_GLOBAL=/dev/null GIT_CEILING_DIRECTORIES=$TMPDIR + PYTHONPATH=/home/vikash/.cache/rdna3-moe-impl/numpy-only-python`, the `python3` + wrapper adding `--jobs 4` to `check-tree-compiles.py`, stdin closed, with the + exit line appended to `/home/vikash/.cache/moe-attn-parity/preflight-repair.log`. + +## Owed + +* **The two arms with no capture to measure them against.** The decode arm + (`query_len == 1` → Triton `kernel_paged_attention_2d` tiled by + `min(block_size, 128)`) and the chunked arm's context phase + (`prefix_prefill._fwd_kernel` tiled by `TRITON_BLOCK_SIZE = 32`) are implemented + from the pinned source and measured against a host transcription of it on + synthesized data, but **no primary capture exists for either**: the primary + attention capture is step-0 only, so every recorded row has `query_len > 1` and + an empty context, and the capture recipe that produced it is gone — gap G2 in + `/home/vikash/.cache/attn-parity-plan/PLAN.md` (`run-operator.py` reads + `attention-capture-recipe/command.json` and `rocm_moe_oracle_attention_capture.py`, + neither of which exists, and the oracle worktree the argv binds is gone too). + A decode-step or a chunked-context capture would need a rewritten recipe; until + one exists, these two arms are implemented-but-unmeasured against the primary. +* **A batch that mixes the arms.** The arm is per row, so one launch can serve + `query_len > 1` rows through the chunked geometry and `query_len == 1` rows + through the decode geometry, which is what the primary does too + (`filter_by_query_len`). Each arm is pinned on its own batch here and the choice + is a per-row branch, but no case builds a mixed batch, so that combination is + implemented and unmeasured. +* **The sliding-window arms.** The walk now anchors its tiles where the primary + anchors them, but the left bound itself stays this backend's own + `jmin = p - window_left`, which CUDA shares (`cuda_paged_attn.cu:203-206`) and + the primary states as `p - j < SLIDING_WINDOW`. No windowed model is measured in + this row, so whether those two bounds agree is open and untouched here. +* The five unmeasured attention kernels named in `## Sibling sites` keep both + the f32 probability and the per-key running max; each needs the same arm-tiled + walk before it mirrors the primary. +* The two rocWMMA attention kernels (`rocm_paged_attn.hip:1283`, `:1754`) keep + the f32 probability until a gfx1200/1201 host can compile and run them. +* The Tier-0 `kAttnQkNormRope` composite keeps the pre-RoPE BF16 store. It is the + realization on every backend without a registered fast op, and changing it + needs an f32 operand slot in the recipe and a mixed-dtype `RopeFromCache`. +* The CUDA, Metal and CPU realizations of both arithmetics are not repaired by + this change and are not measured here. +* `FastExp`, the warp-strided key order and the reciprocal epilogue remain native + traits; only a measured residual would justify touching them. +* The decode arm's throughput is not measured. The finer tile synchronizes once per + `min(block_size, 128)` keys instead of once per 64, so a model with a physical + block size of 16 or 32 pays more CTA synchronizations per decode step. + +## Stop conditions + +* If row C does not reach 0, the residual is attributed against the parameterized + CPU model before any further product edit, or the repair returns with the + residual reported as an open gap. +* If the production gate's recorded tokens move, the golden is reported as moved + with the failing positions; it is never silently re-derived. +* If the primary's boundary cannot be reproduced without changing an op contract + outside this row, the preamble repair stops and returns `BLOCKED` with the + contract named. diff --git a/.agents/specs/rocm-bf16-moe.md b/.agents/specs/rocm-bf16-moe.md new file mode 100644 index 0000000000..93b80f4a56 --- /dev/null +++ b/.agents/specs/rocm-bf16-moe.md @@ -0,0 +1,715 @@ +# Native BF16 grouped MoE on ROCm + +Owning row: `BACKEND-ROCM-BF16-MOE` + +Owner: the `BACKEND-ROCM-BF16-MOE` implementer, reviewer, and coordinating operator. +Parent row: `BACKEND-ROCM`, whose lifecycle stays `ACTIVE`. +Issue: [#3094](https://github.com/mudler/vllm.cpp/issues/3094), local record +`ISSUE-GH-3094`. +Origin: [#1928](https://github.com/mudler/vllm.cpp/issues/1928), local record +`ISSUE-GH-1928`. +Base: `6db4bef906859e864c82523c01107473f7dcca29`. +Integration: one pull request, with this spec committed before implementation. +The implementation pull request carries closing keywords for both issues. +This session ends with a published, independently reviewed pull request. +The operator does not merge it without separate developer authorization. + +## Now + +State: `ACTIVE`. Native providers reach the registered Qwen3 MoE forward path on +`gfx1100`. The legacy grouped suites and eight native boundary, graph, stream, +and two-device cases pass. The production token gate passes under the corrected +whole-sequence-membership rule; the superseded same-configuration comparison +differed at six generated positions across three repeated length-33, +concurrency-2 runs. +The first proved residual-normalization difference belongs to +[#3103](https://github.com/mudler/vllm.cpp/issues/3103), under `BACKEND-ROCM-RESIDUAL-NORM`. +Attention differences (#3115) also require resolution before acceptance; the +head-output boundary is implemented and gated under +[the child spec](rocm-lmhead-bf16.md). + +All 60 original upstream component cases pass on both runtimes. +Initial fresh review found three missing test witnesses at `94b8bb0ec`. +The scoped repair covers provider subsets, malformed descriptors, and accepted +weighted and shared numeric modes. Fresh scoped review and the final operator +gate remain pending. Performance is not accepted before a paired decode and +prefill measurement is recorded at this head. + +## Problem and scope + +ROCm registers neither `kMoeGroupedGemmBf16` nor +`kMoeGroupedGemmBf16GateUpSilu` at the base revision. +The existing BF16 mixture of experts (MoE) fast path therefore remains unavailable +on ROCm. Its eligibility check probes the unfused operation but calls both operations. +The two providers must become available together. + +Implement native HIP providers for grouped BF16 matrix multiplication and fused +gate, up, and SiLU multiplication. Route the feature through the existing shared +operations and Qwen3 MoE production path. Preserve resident expert pointer arrays +and the device router output. Do not add a host gather loop. + +The pinned ROCm oracle rounds intermediate BF16 tensors differently from the +existing CUDA sibling contract. Add the minimum shared numeric modes needed to +represent those differences. Preserve the current default semantics for every +existing caller that does not select the new mode. + +This work is independent of the RDNA3 quantized dot changes in pull request +[#2782](https://github.com/mudler/vllm.cpp/pull/2782). BF16 expert weights do not +reach the quantized RDNA3 dot kernel. Do not depend on that pull request or +change its implementation. Do not change continuous integration configuration, +checker semantics, or the upstream pin. + +This row adds a backend capability to an existing model path. It does not add +a model architecture or a quantized arm. GGUF, FP8, integer quantization, expert +parallelism, LoRA, and expert load balancing remain outside this change. +Do not alter their acceptance, refusal, or numeric behavior. + +## Inventory + +The source anchors in this spec use the immutable revisions stated here. +Line numbers refer to those revisions, before implementation changes them. + +| Stable ID | Upstream source | Local anchor | Required test and evidence | State | +|---|---|---|---|---| +| `ROCM-BF16-MOE-GROUPED` | `vllm/model_executor/layers/fused_moe/fused_moe.py:299-610,763-910` | `include/vt/ops.h:3198-3238`, `src/vt/rocm/rocm_moe_grouped_bf16.hip` | `test_ops_moe_grouped_bf16`, `test_rocm_moe_grouped_bf16`, oracle buffers and generated kernel | `ACTIVE` | +| `ROCM-BF16-MOE-GATEUP` | `vllm/model_executor/layers/fused_moe/experts/triton_moe.py:388-409,487-527`, `csrc/libtorch_stable/activation_kernels.cu:44,165-177` | `MoeGroupedGemmBf16GateUpSiluNative` and legacy typed sibling | `test_ops_moe_grouped_bf16_gate_up_silu`, exact BF16 witnesses, `test_rocm_moe_upstream` | `ACTIVE` | +| `ROCM-BF16-MOE-WEIGHTED-DOWN` | `vllm/model_executor/layers/fused_moe/fused_moe.py:593-610`, `csrc/libtorch_stable/moe/moe_align_sum_kernels.cu:395-459` | `MoeGroupedGemmBf16Weighted`, `MoeCombinePreweighted` | `test_rocm_moe_grouped_bf16`, route-before-narrowing witness and combine checks | `ACTIVE` | +| `ROCM-BF16-MOE-FORWARD` | `vllm/model_executor/models/qwen3_moe.py:199-237` | `src/vllm/model_executor/models/qwen3_moe_registry.cpp:63-86`, `src/vllm/model_executor/models/qwen3_5.cpp:6869-7010,7198` | `test_rocm_moe_bf16` through load/forward, provider statistics, exact tokens and call-site mutation | `ACTIVE` | + +All inventory items belong to this spec and `ISSUE-GH-3094`. +The implementation evidence below distinguishes measured results from remaining gates. + +## Upstream contract + +### Pin and executing chain + +The primary oracle is vLLM at +`e126687a9a828d513c01a07cd69f025f27d63280`. +[Upstream sync](../upstream-sync.md) owns the repository pin. +The inspected checkout contains that exact revision. +Use the operator's recorded oracle runtime and GPU authority from the shared +environment. Machine paths in another spec are not defaults. + +The executing source chain is: + +1. `vllm/model_executor/models/qwen3_moe.py:199-237` constructs and invokes + `FusedMoEFactory` for the routed experts. +2. `vllm/model_executor/layers/fused_moe/unquantized_fused_moe_method.py:41-130` + resolves the backend, creates BF16 expert weights, and defines ROCm padding. +3. `vllm/model_executor/layers/fused_moe/oracle/unquantized.py:61-65,208-325` + orders and selects the unquantized backends. +4. That file's `:329-395` converts weight storage and builds the prepare, + finalize, and expert operations. +5. `vllm/model_executor/layers/fused_moe/modular_kernel.py:1144,1349` + allocates activation storage in the model dtype and invokes the expert method. +6. `vllm/model_executor/layers/fused_moe/experts/triton_moe.py:309-330,388-527` + executes gate/up, activation, and weighted down through BF16 intermediates. +7. `vllm/model_executor/layers/fused_moe/fused_moe.py:763-910` launches the + kernel with explicit tensor strides and the selected compute type. +8. `vllm/model_executor/layers/fused_moe/fused_moe.py:517-610` accumulates in + FP32, applies optional route weights, narrows, and stores. +9. `vllm/model_executor/layers/fused_moe/activation.py:196-237` selects + `torch.ops._C.silu_and_mul` for SiLU. +10. `csrc/libtorch_stable/activation_kernels.cu:44,110,165-177,299` defines + the BF16 activation and multiplication boundaries. +11. `csrc/libtorch_stable/moe/moe_align_sum_kernels.cu:395-459,759` sums + already weighted BF16 expert results with an FP32 accumulator. + +ROCm's source candidate order is AITER, Triton, then batched Triton. +`vllm/_aiter_ops.py:134,1891` restricts the AITER capability to qualifying CDNA +devices. `gfx1100` is RDNA3. Source inspection therefore predicts ordinary Triton +for this unquantized, single-device workload. This prediction is not runtime evidence. + +Before choosing the provider's native default, capture the selected oracle backend, +generated kernel, launch arguments, tensor dtypes, shapes, and strides. +Run the identical workload in the pinned engine. A successful config construction +or an unrelated dense-model capture does not satisfy this gate. +If runtime selection differs, reconcile the executing source before implementation. +Do not replace production defaults with `--enforce-eager` for a denominator. + +### BF16 boundaries + +For the target path, `apply_router_weight_on_input` is false. +Let `b(x)` mean conversion to BF16 with the oracle's rounding behavior. +For token `t`, selected expert `e`, and route weight `r`, the native semantics are: + +```text +g = b(dot_fp32(x[t], Wgate[e])) +u = b(dot_fp32(x[t], Wup[e])) +s = b(silu(float(g))) +a = b(float(s) * float(u)) +d = b(dot_fp32(a, Wdown[e]) * float(r)) +y[t] = b(sum_in_fp32(d for each selected expert)) +``` + +The dot product's order follows the selected kernel and receives the upstream +comparison tolerance. The BF16 conversion points are mandatory independent of +that tolerance. A token match does not prove the memory or arithmetic format. + +The existing CUDA fused sibling computes its gate/up intermediates in FP32. +That legacy mode remains byte-identical to its existing unfused composite. +Its FP32 intermediates are a compatibility exception to the native BF16 mode. +Annotate that reason beside each model-path buffer that keeps FP32 storage. +Do not silently make the old composite the new oracle. + +The native down operation multiplies the route weight before narrowing to BF16. +The combine operation sums these preweighted values without multiplying again. +Post-store weighting is numerically different and fails the native contract. + +### Weight storage and shapes + +Upstream stores gate/up weights as `[E,2I,H]` and down weights as `[E,H,I]`. +Its ROCm padding can add 128 BF16 columns when a row occupies a multiple of +512 bytes. The default `VLLM_ROCM_MOE_PADDING` is enabled in `vllm/envs.py:1341`. +Triton conversion preserves the physical strides in +`vllm/model_executor/layers/fused_moe/oracle/unquantized.py:364-370`. +Record those strides in the oracle capture and retain padding cases in the tests. + +The local shared ABI uses one device pointer per expert. +Each pointer names a Matmul-B matrix `[K,N]`, indexed as `k * N + n`. +The loader's transpose from checkpoint storage is an existing harness adaptation. +Do not reinterpret a checkpoint `[N,K]` array as the shared matrix format. +The local ABI does not require unused physical padding bytes to match upstream. +It does require BF16 storage, correct logical values, and explicit recorded strides. + +## Shared design + +### Operations and selection + +Register both existing operation IDs for `DeviceType::kROCM` in the same change. +Keep their current typed call signatures and default semantics intact. +Add an explicit shared numeric descriptor or typed sibling operations for native +gate/up rounding, weighted down, and preweighted combine. +If a descriptor changes a function-pointer signature, use a separate typed sibling +instead of casting an old provider to a new signature. +The exact C++ names are implementation choices. The semantic modes are fixed here. + +The Qwen3 MoE path selects the complete native capability through the shared +provider seam. Do not add a model-specific HIP kernel or a device-type branch +that changes model arithmetic. Probe every operation required by the native mode. +Never select half of the native sequence and continue with legacy weighting. +Existing callers that do not select this capability keep their defaults. + +The existing row map maps pair `p` to its activation row. +A null row map means identity. Expert IDs and row maps remain on the device. +Inputs are contiguous BF16 activations, contiguous I32 indices, and I64 device +pointer arrays. Outputs accept the existing grouped BF16 and FP32 modes. +Preserve validation in `src/vt/ops.cpp:906-962`. +Preserve the zero-pair and zero-output-width no-op behavior. +Use valid router-produced expert IDs. This row does not add sentinel index semantics. + +The fused operation accepts separate gate and up pointer arrays. +Use one grouped provider implementation for all callers, including the legacy mode. +Mask incomplete K and N tiles. Preserve repeated experts, empty experts, arbitrary +pair order, repeated rows, and dimensions that are not tile multiples. +Do not assume a particular top-k, expert count, or model geometry. + +### Fusion and production reachability + +The existing pointer-array BF16 sibling is a tracked exception to a literal +`vt::MergedGemmGroup` instance. +`include/vt/merged_gemm.h:111-118` documents its different weight representation. +[The fusion spec, Tier A4](arch-fusion-fold-plan-2026-07-30.md) records the +shared operation and its original compatibility contract. +This row extends that shared operation with explicit native numeric modes. +It does not create another exception or a per-model expert loop. +Update that local seam comment to distinguish legacy and native numerics when +the implementation introduces the mode. + +Model fusion continues through `vt::FusedChain` where that seam applies. +Mergeable dense and shared-expert projections continue through +`layers::MlpGateUpMethodBase` and `vt::MergedGemmGroup`. +Do not duplicate their existing model orchestration for this backend. + +The production vehicle is `Qwen3MoeForCausalLM` loaded from safetensors. +`src/vllm/model_executor/models/qwen3_moe_registry.cpp:63-86` exposes load and forward. +`src/vllm/model_executor/models/qwen3_moe.cpp:71` delegates to the shared MoE block. +`src/vllm/model_executor/models/qwen3_5.cpp:7198` selects the grouped BF16 path. +The resident gate/up and down calls occur at `:6997-6999`. +The existing `VT_MOE_BF16_FAST` default is enabled at `:909`. +Preserve the rollback setting and validate both sides with the same binary. + +DeepSeek and dots3-note contain callers of the shared BF16 operations. +Their production routers require currently refused ROCm modes. +`src/vt/rocm/rocm_moe_router.hip:183` rejects grouped routing, correction bias, and non-softmax +scoring. This row does not claim that either complete model becomes runnable. +Keep that router debt under the parent `BACKEND-ROCM` issue +[#41](https://github.com/mudler/vllm.cpp/issues/41). + +### Kernel and scratch lifetime + +Place new hardware code in `src/vt/rocm/rocm_moe_grouped_bf16.hip`. +Use `src/vt/cuda/cuda_matmul_nvfp4.cu:943-1004,1413-1573` as a local +layout and deterministic-reduction donor. Its legacy arithmetic is not the +native oracle contract. Attribute every ported upstream kernel section at the pin. + +Start with a complete deterministic kernel. Optimize only after correctness passes. +If split-K is used, reduce partials in a fixed order and apply conversion once +at the specified boundary. Do not use atomic output accumulation that changes +results across identical runs. + +Key reusable scratch by both device and stream. The null default stream does +not identify a unique device. Serialize allocation and publication across threads. +Preserve allocation lifetime until all users and captured graphs finish. +`src/vt/grow_only_stream_scratch.h` supplies the existing growth and retirement +mechanism. Adapt that mechanism instead of copying CUDA globals with weaker keys. +Allocation failure must not publish partial capacity or invalid pointers. + +Prewarm the required capacity before graph capture. +Do not allocate, free, or synchronize during capture. +Preserve captured addresses when another invocation grows scratch. +Validate streams independently and devices separately when two devices are available. +A missing second device leaves that case `PENDING`, not silently skipped. +Direct operation capture tests do not enable model graph capability. +This row does not depend on the separate graph-capability pull request #2777. + +## Tests and evidence + +### Red before implementation + +Create a deterministic tiny safetensors fixture through the existing model loader. +Adapt `tests/vllm/models/test_moe_async_device_ids.cpp:111-203,415-477`. +Enter through `ModelRegistry::Load` and `ModelRegistry::Forward`. +Do not construct an internal MoE block by hand for the reachability claim. + +Use this concrete Qwen3 MoE configuration: + +```json +{ + "architectures": ["Qwen3MoeForCausalLM"], + "model_type": "qwen3_moe", + "hidden_size": 128, + "num_hidden_layers": 2, + "num_attention_heads": 1, + "num_key_value_heads": 1, + "head_dim": 128, + "intermediate_size": 128, + "moe_intermediate_size": 128, + "shared_expert_intermediate_size": 0, + "num_experts": 4, + "num_experts_per_tok": 2, + "norm_topk_prob": true, + "decoder_sparse_step": 1, + "mlp_only_layers": [], + "hidden_act": "silu", + "vocab_size": 128, + "max_position_embeddings": 256, + "rms_norm_eps": 0.000001, + "rope_theta": 10000000.0, + "tie_word_embeddings": false, + "attention_bias": false, + "torch_dtype": "bfloat16", + "bos_token_id": 1, + "eos_token_id": 127, + "pad_token_id": 0 +} +``` + +Preserve `BuildTensors` ordering from the fixture at the base revision. +Set its first tensor seed to 7 and increment the seed for each tensor. +Preserve `Bf16Bytes` at `:131-142`, including unsigned 32-bit wraparound and +the existing `F32ToBF16` conversion. Keep projection scale 0.08 and norm scale 0.5. +Only substitute the dimensions and initial seed stated here. +Export the generated safetensors and config once, then use those same bytes in +both runtimes. Record each file's byte count and SHA256. + +For each length `L` in `{1,3,33}`, request `r` contains token IDs +`1 + ((11 + 29*r + 17*i) % 126)` for `i` from 0 through `L-1`. +Run request 0 at concurrency 1 and requests 0 and 1 together at concurrency 2. +Repeat each workload 3 times with fresh model state. +Use token-ID prompts with tokenizer initialization skipped on the oracle. +Generate exactly 8 tokens per request through the existing shared device sampler. +Set temperature 0, top-p 1, top-k -1, min-p 0, repetition penalty 1, +presence penalty 0, frequency penalty 0, seed 7, and `ignore_eos=true`. +Set both the minimum and maximum generated token counts to 8. +Record selected expert IDs and assert changing expert pairs across the request set. +Keep repeated expert choices in the component fixtures. + +Assert provider selection for grouped down and fused gate/up. +Enable `OpProviderCallStats` and require native selections greater than zero, +zero native declines, and zero CPU fallbacks for the tested expert sequence. +On `gfx1100`, absence of the providers is a failure, not a skipped test. +The red result must fail that reachability assertion on the base revision. +Correct fallback tokens alone do not satisfy this test. + +### Port upstream cases + +Port `tests/kernels/moe/test_moe.py:test_fused_moe` at `:345` from the pin. +Preserve its seed 7, BF16 fixtures, `atol=0.02`, and `rtol=0`. +Retain these `[M,N,K]` parameter sets: + +```text +[1,128,128] +[1,2048,128] +[33,2048,128] +[32768,2048,511] +[40000,1024,1024] +``` + +Retain expert counts 8, 64, and 192, top-k values 2 and 6, and both padding modes. +Run every applicable single-device BF16 combination from the upstream fixture. +The local pointer-array transpose is the required storage adaptation. +Expert parallel size 4 is inapplicable because this row adds no distributed path. +Tensor-descriptor mode is inapplicable on this hardware by +`vllm/model_executor/layers/fused_moe/utils.py:665-684`. +Record those exclusions with their source rationale. +Do not reduce large M or tail dimensions to make the tests fit. +Name the resource and retain the gate as `PENDING` if a required case cannot run. + +Preserve the upstream reference decomposition, activation mode, route-weight +placement, and output comparison. Record generated kernels before interpreting +any unavailable lever. A scratch reference must be labeled as scratch evidence. +The upstream component fixture uses `renormalize=false`. Preserve that mode. +The production fixture uses `norm_topk_prob=true`. Preserve that mode separately. +Keep the upstream `use_compile=false` setting and its graph replay cases where +`N >= 1024` and `K >= 1024` on CUDA-alike hardware, including ROCm. + +Add local cases from `tests/vt/test_ops_moe_grouped_bf16.cpp` and +`tests/vt/test_ops_moe_grouped_bf16_gate_up_silu.cpp`. +Cover row-map and identity-map inputs, BF16 and legacy FP32 output, large pair +counts, split-K candidates, K and N tails, repeated experts, and empty experts. +The legacy fused operation must remain byte-identical to its existing composite. +The native mode must match captured oracle boundaries within the upstream gate. +Do not compare native arithmetic only against the legacy composite. + +Use boundary values that fail when each BF16 narrowing point is removed. +Use nontrivial route weights that distinguish weighting before and after narrowing. +Assert physical activation and expert-output dtypes independently of token results. +Compare greedy token IDs exactly for the complete production fixture. +Capture finite logits and their differences for diagnosis. Never widen the token +gate because an output difference appears numerically small. + +### Scratch and negative mutations + +Test repeated launches, simultaneous streams, scratch growth, allocation failure, +prewarmed graph capture, replay after growth, and zero-sized workloads. +Require capture and replay to preserve output and pointer lifetime. + +A fresh reviewer applies each applicable mutation to an immutable scratch copy: + +| Guarantee | Mutation that must fail | +|---|---| +| Both providers are reachable | Remove each registration separately | +| Production uses the new sequence | Delete its production call site but retain fallback | +| Native mode is selected | Force legacy mode without changing final fallback availability | +| Matrix and row layouts are correct | Transpose the weight stride or ignore the row map | +| Expert selection is correct | Substitute another expert or mishandle an empty expert | +| Tail masks are correct | Remove a K mask and an N mask separately | +| BF16 boundaries are correct | Remove gate/up narrowing and SiLU narrowing separately | +| Weight placement is correct | Move route weighting after down-output narrowing | +| Combine consumes preweighted output | Apply the route weight a second time | +| Scratch survives growth | Free a block still referenced by a captured graph | +| Scratch keys isolate users | Remove the stream or device key where hardware permits | + +Record the focused command, nonzero result, first relevant failure, and byte-for-byte +restoration after each mutation. A source inspection does not replace these tests. + +## Gates and acceptance + +The implementer records exact commands after the test executable names are final. +The evidence must include these obligations: + +1. Startup, role, and full `scripts/agent-preflight.sh --staged` at the implementation head. +2. A clean CPU build and existing shared-operation regression tests. +3. A HIP build for `gfx1100`, including the new source in the HIP compile options. +4. The focused native and legacy grouped suites with zero unexpected skips. +5. The pinned oracle's identical production fixture and exact token comparison. +6. Dtype, stride, backend-selection, and generated-kernel evidence on both sides. +7. Scratch, capture, concurrency, and negative-mutation results. +8. A fresh scoped review followed by the operator's own hardware gate. + +Use the recorded lease or mutex required for the actual GPU. +Record device identity, ROCm and compiler versions, binary SHA256, revisions, +artifact hashes, launch recipes, sampling configuration, and contention state. +Do not perform GPU work outside the operator's recorded authority. + +Only accept performance after the declared correctness gate passes. +Trace both runtimes with the same tool on identical inputs. +For each applicable decode and prefill case, record throughput, latency, peak +device memory, peak host memory, and ratios against the production oracle. +The throughput floor is 1.0 times the oracle. Latency and memory must not exceed +1.0 times the oracle. Keep an unmet axis open with its next traceable hypothesis. +Do not describe an unresolved implementation difference as an architectural ceiling. +Repeat accepted results on an idle device with the same-binary fallback comparison. + +Report each gate as satisfied, narrowly waived, pending a named external authority +or resource, or failing. This row cannot become `DONE` with an unresolved required +axis. Add an `## Outcome` section when it reaches `DONE`, including measurements, +rejected approaches, and the reason for each default. + +### Exact production token gate (corrected definition) + +The gate certifies that the native engine computes the reference function, not +that it reproduces one execution of it. The reference's greedy decode is not +unique across its own legitimate configurations: at length 33, concurrency 1 and +concurrency 2 on the identical prompt, the pinned primary itself emits +`[66,1,70,57,33,81,63,69]` and `[66,1,70,57,33,81,118,66]` respectively (all +three repeats each, `oracle-selection-6/production.json`); the two runs agree on +the first six tokens and differ at positions 6 and 7 of request 0. The developer +ratified this definition on 2026-09-09. A workload passes when the native +sequence equals, for that request, a sequence the reference itself emits under +one of its captured configurations (whole-sequence membership). Positions at +which the captured reference configurations agree remain exact, because +membership in any single reference sequence implies they match. The gate reports +the reference set, the matched configuration, each same-configuration outcome, +and every position where the reference disagrees with itself. It never mixes +positions from different reference sequences, never drops a workload, and never +relaxes a stable position. + +Fidelity is a separate criterion, and no gate in this change measures it. It +accepts a native-versus-reference logit difference of at most `1.953125e-3` +(`2^-9`) per logit, which is one BF16 unit in the last place at the logit +magnitude where that maximum was measured. That value is recomputed from the +retained BF16 head-logit diagnostic captures in +`/home/vikash/.cache/rdna3-moe-impl/preserved/oracle-diagnostic-1/`, comparing +request 0 of the concurrency-1 and concurrency-2 captures of one workload at +every step where the generated context still agrees, that is +`L1-C{1,2}-R0-head-0-logits.bin` together with +`L33-C{1,2}-R0-head-{0,1,2,3,4,5}-logits.bin` (BF16, vocabulary 128; request 0 is +row 0 of the concurrency-2 capture). The largest absolute per-logit difference is +exactly `1.953125e-3`, attained at length 33 steps 3, 4, and 5, which is one BF16 +ulp at those logits' magnitude, and the greedy argmax is unchanged at every +compared step. The oracle's +own logprob deltas in `oracle-selection-6/production.json`, which reach `2.028e-3` +between its concurrency-1 and concurrency-2 records at positions where their token +sequences agree, are log-probabilities over the vocabulary and are not this logit +band. The attention and Q/K preamble parity work (#3115) and the BF16 LM-head +output boundary work (#3116, [the child spec](rocm-lmhead-bf16.md)) are judged by +this fidelity criterion, not by reproducing one configuration's tie-break. + +## Files and authority + +The implementation owns the new HIP source, its registration in `rocm_ops.hip`, +the corresponding CMake source and compile-option entries, and focused tests. +It can extend `include/vt/ops.h`, `src/vt/ops.cpp`, provider metadata, shared +MoE descriptors, and the documented fusion seam to represent the fixed numeric modes. +It can update the Qwen3 shared MoE dispatch to select the complete native mode. +Mechanical adapters in existing providers are allowed when required by the shared +API. Those adapters must preserve their current numerics and defaults. + +The implementation updates this spec and the issue records as their state changes. +The parent matrix remains unchanged because its lifecycle does not change. +This per-row inventory records the new capability without another shared matrix write. +Update `docs/FEATURES.md` and the operation inventory in `docs/ROCM.md` when the +native backend capability ships. Document a changed command or configuration in +`docs/USAGE.md` only if the implementation changes that public surface. +This row's spec commit makes no shipped capability claim and owes no public rewrite. + +## Implementation evidence, 2026-09-09 + +The implementation starts from committed spec `9202e4c4edc4cf6ef9b3e8da66431effb0fbcee5`. +Its isolated worktree is `/home/vikash/vllm.cpp-rdna3-moe-impl`. +The paths below identify this measured run; they are not environment defaults. +Evidence root: `build-rdna3-moe-hip/evidence/` in that worktree. +The operator holds `/home/vikash/gpu.lock` for every GPU invocation. +The measured device is a Radeon 7900 XTX, `gfx1100`. +Native compilation uses HIP `7.15.26333`, Clang 23, and `-ffp-contract=off`. +The pinned oracle reports Torch `2.12.0+git6bbd260` and HIP `7.2.53211`. + +### Design and executing oracle + +The shared API adds three typed siblings: native gate/up, weighted down, and +preweighted combine. The existing typed signatures and legacy arithmetic remain +unchanged. Capability selection requires all five grouped/native providers. +The HIP implementation uses one deterministic reduction per output and no scratch +allocation. A device scope selects the queue's device and restores the caller's +ambient device. The stream always comes from that queue. + +The pinned production engine selects `TritonExperts` in both layers. +Captured gate/up and down inputs, weights, and outputs are BF16 with two-byte +elements. Gate/up weight strides are `[32768,128,1]`; down strides are +`[16384,128,1]`. The captured MoE kernel uses `BLOCK_SIZE_M=32`, +`BLOCK_SIZE_N=64`, `BLOCK_SIZE_K=128`, and `SPLIT_K=1` for the small fixture. +Its generated `fused_moe_kernel.amdgcn` contains BF16 WMMA instructions. +Native offload code and the compile command are retained in `native-generated/`. + +The fixture export preserves the committed generator, including its tensor order: + +| Artifact | Bytes | SHA256 | +|---|---:|---| +| `config.json` | 719 | `321926020ada026d8dd85f74543dd12b6426301706b4e1f4fbca60a44b871cea` | +| `model.safetensors` | 1123330 | `96cd7f30fee496c69782af2813e438e47d7b026598b0f20a05049c522b279af8` | + +The final cohort harness uses pinned `LLM.sleep(level=0)`, `enqueue`, scheduling +wake-up, and `wait_for_completion`. It records the authoritative internal/external +request-ID mapping and asserts the actual prefill and decode batches for every run. +`oracle-selection-6/production.json` has SHA256 +`f3d27a95ba71bddcae38c3defc21bf3f3ff32ba9373a09faf949ead68b74d7ad`. +Earlier oracle attempts and the initial uncoordinated scheduling result remain +in evidence. They do not supply the final matched-cohort denominator. + +### Measured gates and remaining work + +| Gate | Result and evidence | +|---|---| +| Red before implementation | Satisfied. `production-red-2.log` has four missing registration/selection failures and 27960 passing assertions. Its executable SHA256 is `25dfd4e7e58dbe6069c6d1ea20c00ae7249687336645c04389845ceede0e83ae`. | +| CPU build and shared regressions | Satisfied. Six tests pass: model registry, grow-only scratch, provider metadata, native descriptor validation, MoE operations, and grouped router. Log: `/home/vikash/.cache/rdna3-moe-impl/cpu-native-tests.log`. | +| HIP build and legacy arithmetic | Satisfied. Both existing grouped suites pass on ROCm: 7 cases/19 assertions and 3 cases/6 assertions, with no skips. Logs: `test_ops_moe_grouped_bf16-native-1.log` and `test_ops_moe_grouped_bf16_gate_up_silu-native-1.log`. | +| Native numeric boundaries, streams, capture, and devices | Satisfied. `native-boundary-3.log` records 8 cases/2804 assertions, no skips, with devices 0 and 1 visible. It covers opposite ambient device state for native gate/up, weighted down and combine. Executable SHA256: `8bae75c8506b08f2444ed0a4b8379b837608b9261baccdef3c71a97ad0408a33`. | +| Scratch-only allocation and retirement mutations | Narrowly waived for this scratch-free implementation. No allocation, free, capacity publication, retired block, or scratch key exists in these providers. Graph replay after larger shapes and concurrent streams still run. | +| Original upstream component cases | Satisfied. All 60 cases pass on both runtimes, including M=32768/K=511 and M=40000/K=1024 graph cases. The unchanged pinned test supplies seed 7, BF16 fixtures, both padding modes and original tolerances. Raw padded source storage and logical exported strides are retained separately in `upstream-all-2` and its range directories. `upstream-all-2-complete-operator-summary.json` independently checks all 120 case/stage results; its SHA256 is `b5c9d8259ae6c4701ad92647f5e61bdda7ba1ef6e2a2fd6ecb8a9edf99a4af04`. | +| Production provider selection | Satisfied. All three new operations have positive native selections, no declines, no fallbacks, and no CPU selections. The existing two registrations are present. | +| Exact production tokens | Satisfied under the corrected whole-sequence-membership rule (`### Exact production token gate (corrected definition)`, ratified 2026-09-09). Measured native tokens match a captured reference sequence at all 18 workloads, no workload matches NO member, and native request 0 equals the concurrency-1 reference sequence at every length. Request 1 exists only in the concurrency-2 records and matches there. Historical note on the superseded same-configuration comparison: the pinned primary disagrees with itself at length 33/concurrency 2/request 0, where its concurrency-1 record emits `[66,1,70,57,33,81,63,69]` and its concurrency-2 record emits `[66,1,70,57,33,81,118,66]`, at positions 6 and 7 of request 0 on all three repeats. The native run equals the concurrency-1 member token-for-token. The operator receipt for this corrected gate at this head is `/home/vikash/.cache/moe-6fd1650c4-tmp/operator-gpu-positive-da0ce377b.log`: exit 0, 28598 of 28598 assertions passing, every same-configuration outcome printed as a report (the 18 assertions beyond the pre-`REQUIRE` receipt are the request-count requirement, one per matched record). The operator's mutation study at this head restores the superseded same-configuration-only comparison and reddens exactly the three recorded failures; the helper header was restored byte-for-byte (`265e28fe`). The superseded-head red history that the corrected gate answers is `/home/vikash/.cache/residual-norm-repair1/green-cc9d4f565/fusion-1-operator.log`, which records 28541 passing assertions and the three failing same-configuration comparisons, one per repeat. `/home/vikash/.cache/moe-6fd1650c4-tmp/legacy-fused-token-agreement.json` (sha256 `9a97f18b7553c80a`) re-derives the legacy agreement at this head: `VT_FUSED_CHAIN_ADOPT=0` and `=1` both exit 0 for all 18 records and emit the same 216 generated tokens with no mismatching record. The earlier `baseline-production-tokens-comparison.json` was removed with the superseded 41 GB evidence tree during the campaign cleanup; this receipt replaces it. | +| CPU descriptor negative mutations | Satisfied. All nine mutations in `cpu-contract-mutations-2/results.json` fail their intended descriptor or capability assertion. Original source and archive hashes remain equal after each run. The first preparation linked the unchanged CPU whole archive, so that invalid probe is preserved and excluded. | +| Implementer negative mutations | Satisfied. The operator ran all 20 isolated mutations in `negative-mutations-1/run-recipes.json`; every intended defect was detected, with no survivor or timeout. Original source, archive and executable hashes remain unchanged. Each of the eight production mutations adds its specific provider failure beyond the existing token failures. Filtered component mutations select one test with nonzero assertions. `operator-results.json` and per-case receipts preserve the commands and failures. | +| Full staged preflight | Satisfied for executed checks: exit 0, no failures, and 619/619 affected host translation units compiled. The report lists 12 explicit skips, reconciled below; it does not print an all-green claim. Log: `/home/vikash/.cache/rdna3-moe-impl/staged-preflight-1.log`. | +| Supplementary preflight suites | Satisfied. All seven NumPy-dependent suites pass using the existing NumPy 2.1.3 package through task-local links. Both CPU and HIP build databases pass the x86 ISA audit. The first full-site Python path exposed installed vLLM metadata to fake-runtime tests; that failed environment attempt is preserved and excluded. | +| ARM and CUDA build audits | Narrowly waived for this gfx1100 change: no ARM build, CUDA fat binary, or CUDA Triton AOT artifact is produced. The native HIP build and generated gfx1100 object supply backend build evidence. | +| Frozen-head PR path, trailers and style | Pending the local implementation commit and the operator's exact-SHA pre-push checks. | +| Fresh review and final operator gate | Pending the fresh reviewer and coordinating operator at the immutable implementation head. | +| Paired traces, throughput, latency and memory | Failing acceptance prerequisite: no paired decode or prefill measurement is recorded at this head. No performance result or floor is accepted. | + +Four additional scratch replays use original captured production MoE inputs, +expert IDs, BF16 weights and route weights for both layers at L33/C1 and L33/C2. +All pass the original absolute tolerance of 0.02. Activated outputs differ in +3/1/5/1 BF16 words, weighted down in 7/0/22/1, and final sums in 2/0/10/1. +Maximum final difference is `6.103515625e-05`. The final comparator is a CPU +reconstruction of the sum of captured weighted BF16 down values. It is not an +observed final Triton tensor. These scratch runs isolate expert arithmetic; +they do not replace the original upstream cases or exact production tokens. + +Paired dense captures locate differences before the new experts. The first QKV +input, weights and output match exactly. Attention output first differs at row 1. +Row 0 attention projection remains exact, but post-attention normalization differs +in 33 of 128 values. `norm-gap-handoff/handoff.md` and its manifest preserve the +complete executing compiled lifetime for #3103, including all six generated +modules and their allocation/store dtypes. Their SHA256 values are +`9850ee49c93bb82ddfca1a811fb422337c33ea09335e749fbf3c42dfcaa0400a` and +`8531cb41f0168d27455b2dc8062a57d5db7afd4e0822013332f398bbf2239c65`. +The operator independently reran the algebra witness and checked every manifest hash. + +## Test repair after fresh review + +Fresh review of `94b8bb0ec67eff82d2860cdda1069e4b0a64bac8` found three P2 +coverage gaps. Static review found no additional arithmetic defect. +The review report remains at +`/home/vikash/.cache/rdna3-moe-review/evidence/review-findings.md`. +The repair worktree is `/home/vikash/vllm.cpp-rdna3-moe-repair1`. +Its evidence root is `/home/vikash/.cache/rdna3-moe-repair1/evidence`. +These paths identify retained measurements. They are not environment defaults. + +The repair changes two test files and this row's records. Product sources, +CMake, original fixtures, and tolerances remain unchanged. +`protected-before.json` records the product and original archive hashes. + +### Added witnesses + +The shared contract test checks all 32 availability subsets of the five providers. +It uses uniquely named XPU stubs and the existing provider-disable seam. +Each subset verifies every `OpRegistered` result before checking the complete set. +The cases include the empty set, the legacy-only set, and each missing provider. +The test never dispatches a stub or changes a production device's providers. + +Malformed descriptors change one property per subcase. +The new cases cover route rank, stride, and device, and weighted down's common validator. +They cover both combine tensors' strides and devices. +Shared-input cases cover rank, both dimensions, stride, and device. +Each case requires its specific validation error before provider dispatch. + +Weighted numeric cases use both output dtypes, nonidentity row maps, and per-pair routes. +A route of 257/512 produces exact FP32 values that BF16 output must narrow. +The shared-input case adds BF16 values into both output types. +Its FP32 result retains 257/512; BF16 rounds that halfway value to 0.5. +Fixed expected values come from exact rational arithmetic, independent of the provider. +`numeric-witness-v2.json` retains that calculation and the counterexamples. + +### Repair verification + +The repair recreates the reviewer's retained mutations in fresh sources and archives. +The new tests supply the red gates. The unchanged product supplies the green controls. +The CPU mutation set reports ten intended failures across ten distinct binaries. +The conjunction-to-disjunction mutation fails all 30 partial provider subsets. +Each of the nine descriptor mutations fails its corresponding error assertion. +`cpu-mutations/results.json` records commands, filters, exit statuses, and hashes. +Each mutation keeps its changed source and archive in a separate directory. +The original product sources and archives stay unchanged. + +A clean CPU configuration and build pass the six declared regressions: +model registry, grow-only scratch, provider metadata, native descriptors, +MoE operations, and grouped router. The coordinator independently confirms +72 passing cases and 2910 assertions in `cpu-operator-results.json`. +The model-registry executable retains its previously disabled `can_initialize` case. +No added contract case skips. Commands and results are retained in +`cpu-configure.log`, `cpu-build.log`, `cpu-test-rebuild.log`, and `cpu-positive.log`. +The focused CTest command is: + +```sh +ctest --test-dir build-repair-cpu --output-on-failure -R '^(test_model_registry|test_grow_only_stream_scratch|test_op_provider|test_moe_bf16_native_contract|test_ops_moe|test_ops_moe_router_grouped)$' +``` + +HIP test compilation preserves the original compiler options and relinks the +unchanged production archive. Its SHA256 is +`a9f57617d3ef79029b2529b216bb370cbd2fc41fa89b1f3f992c5a9c0a7433eb`. +This is a test relink, not a new full HIP product build. +The first GPU batch passes its three controls and detects all four review survivors. +Its frozen sources, commands, binaries, and operator receipts remain under `hip/` +and `hip-mutations/`. The second version strengthens F32 precision witnesses. +The second version passes all three positive controls: 219 contract assertions, +72 weighted-mode assertions, and 42 shared-mode assertions. +All four review survivors fail their intended assertions. +An additional FP32-store narrowing mutation fails both precision witnesses. +`gpu-v2/hip-mutations/run-recipes.json` retains each command and binary hash. +`operator-results.json` and `operator-input-audit.json` in that directory +record the coordinator's results and unchanged source, archive, and binary hashes. +No implementer GPU execution occurred. The coordinator runs all hardware commands +under `flock -n -F /home/vikash/gpu.lock` with the recorded visible-device environment. + +The pre-edit full preflight exits 1 with 619/619 host units compiled. +Its only failure is an unchanged symbol-anchor fixture inheriting the parent Git repository. +The isolated rerun passes all 22 tests. +The staged gate uses `GIT_CEILING_DIRECTORIES` to isolate those temporary repositories. +It also uses `GIT_CONFIG_GLOBAL=/dev/null` for the known onboarding fixture. +An external Python wrapper appends `--jobs 4` only to `check-tree-compiles.py`. +All other checker arguments remain unchanged. +The seven NumPy suites use the existing isolated package directory. +The staged full preflight exits 0 with 619/619 host units compiled. +`staged-preflight.log` and `staged-preflight.exit` retain the complete result. +Its five argument-required skips have the dispositions recorded here. +The final receipt-only record edits receive focused record, anchor, and command checks. + +The CPU and HIP compile databases pass the x86 ISA audit. +Exact path classification uses the immutable repair commit. Its result is retained +in `path-classification.log` before handoff. +ARM, CUDA fat-binary, and CUDA Triton AOT audits are narrowly waived: +this test repair produces none of those artifacts. +Fresh scoped review and the operator's final verification remain pending. + +The original 60 upstream cases and production-token evidence remain authoritative +for the unchanged implementation. This test repair does not rerun that oracle matrix. +The exact model gate failed at six generated positions across three repeats under the +superseded same-configuration comparison, and it passes under the corrected +whole-sequence-membership rule above. +Agreement between all 216 native and legacy tokens does not establish oracle parity. +No performance result is accepted, and the row remains `ACTIVE`. + +## Owed + +No new unowned issue is introduced by this spec. +The origin issue keeps its existing single owner in +[the device-fit spec](gguf-device-fit-expand-policy.md#owed) until this +implementation closes it. That owner links here for the implementation handoff. +The new row owns its tracking issue directly. +Grouped routing and correction-bias work stays under `BACKEND-ROCM`, issue #41. + +- [#3100](https://github.com/mudler/vllm.cpp/issues/3100), owned by + `BACKEND-ROCM`, tracks shared resource operations that ignore their device index. + The two-device test selects the resource device in its host setup, checks pointer + ownership, and then launches each provider with the opposite ambient device. + This isolates the new provider's device contract without changing shared allocation. +- [#3103](https://github.com/mudler/vllm.cpp/issues/3103), owned by + `BACKEND-ROCM-RESIDUAL-NORM` in [its repair spec](rocm-residual-norm.md), tracks the compiled Qwen3 MoE residual-normalization lifetime. + The paired row-zero witness reproduces all 128 native and oracle values separately. + The native path rounds the residual before variance; the compiled oracle elides + the post-attention residual store and materializes the next input-norm residual. + Every materialized residual remains BF16. This row does not change that shared + normalization path. Its production token gate passes under the corrected rule above. + +## Stop conditions + +Return `NEEDS_CONTEXT` if the oracle runtime, model bytes, or GPU authority is +unavailable for the next gate. Continue independent CPU work within the spec. +Return `NEEDS_DECISION` if the executing oracle requires semantics outside the +fixed shared modes or if a new shared-seam exception is required. +Do not change unrelated routers, quantized kernels, model architecture, or checkers. +Do not suppress a correctable failure or mark a fallback-only result as native. +Keep a blocked performance axis visible without stopping correctable work. diff --git a/.agents/specs/rocm-lmhead-bf16.md b/.agents/specs/rocm-lmhead-bf16.md new file mode 100644 index 0000000000..aae356303a --- /dev/null +++ b/.agents/specs/rocm-lmhead-bf16.md @@ -0,0 +1,253 @@ +# Mirror the compiled primary's BF16 LM-head output boundary + +Owning row: `BACKEND-ROCM-BF16-MOE`, whose lifecycle remains `ACTIVE`. +Parent row: `BACKEND-ROCM`, whose lifecycle remains `ACTIVE`. +Owner: the delegated head-boundary implementer, fresh reviewer, and coordinating operator. +Issue: [#3116](https://github.com/mudler/vllm.cpp/issues/3116), `ISSUE-GH-3116`. +Base: `691b7af30`, branch `row/BACKEND-ROCM-BF16-MOE-head-bf16`. +Integration: one pull request under the repository default. This spec is committed before implementation. +The implementation pull request closes #3116 when its required gates pass and it lands. + +## Now + +State: `SPIKE`. The boundary measurement is complete on `gfx1100` and the specification precedes the implementation. +The measurement supports landing the boundary as a fidelity mirror and does **not** support claiming a token repair. +The operator's 18-workload production token gate passes at `8e43d18bd` under the corrected whole-sequence-membership rule +(28598 of 28598 assertions, receipt `/home/vikash/.cache/moe-head-bf16/operator-gate.log`), and the frozen-head preflight +there exits 0 (receipt `/home/vikash/.cache/moe-head-bf16/preflight-final.log`). The fresh review's two LOW findings and +two bookkeeping corrections are repaired on this branch; a fresh scoped review of the repair is the next gate. + +## Problem and scope + +The pinned compiled primary stores its LM-head output in BF16 and widens it to F32 only where the sampler needs it. +The native Qwen3-MoE forward stores F32 logits directly: + +- `src/vllm/model_executor/models/qwen3_moe.cpp:279-302` — the UNTIED `lm_head` block declares + `DBuf logits(d, DType::kF32, {n_out, vocab})` and calls `vt::Matmul(d.q, logits.t(), src, lm)`. +- `src/vllm/model_executor/models/qwen3_moe.cpp:21` states the contract this row is changing: "Returns `[n_out, vocab] f32` logits". + +Scope is the head-output dtype boundary only: + +1. The head projection writes BF16, mirroring the primary's head output dtype. +2. The shared `vt::CastF32` widens that BF16 result into the F32 buffer the forward returns today. + +Out of scope: + +- Decode attention and Q/K preamble parity (#3115). Hidden-state parity at about one BF16 ulp remains required, and + this change alone cannot move the native answer onto the primary's concurrency-2 token. +- The router, grouped MoE, quantization arms, the upstream pin, CI, and checkers. +- Any change to the sampler, the logits views, or the captured-graph logits slot. + +Changed by this work, and unreplayed: the tied-embedding arm's arithmetic. `lm_head::Project` sends BOTH arms through the +same BF16 store and the same `vt::CastF32` (`include/vllm/model_executor/models/lm_head_projection.h:64-69`), so the tied +arm's `vt::MatmulBT` now narrows to BF16 where it stored F32 directly before this change. No captured artifact exercises +that arm — the fixture this spec replays records `"tie_word_embeddings": false` +(`/home/vikash/.cache/rdna3-moe-impl/preserved/fixture/config.json`), and the measurement above is the untied `vt::Matmul` +orientation only. The tied arm is therefore changed here, unreplayed, and unreachable from this measurement; its own +artifact replay is owed in `## Owed`. + +## Upstream anchor + +- Primary oracle: vLLM `e126687a9a828d513c01a07cd69f025f27d63280` (`.agents/upstream-sync.md:7`). +- Executing chain: `Qwen3MoeForCausalLM` computes its untied `lm_head` as a BF16 `F.linear`; the logits processor + widens the result to F32 before sampling (`.agents/specs/rocm-residual-norm.md:320-325`). +- Captured primary evidence: `/home/vikash/.cache/rdna3-moe-impl/preserved/oracle-diagnostic-2/L33-C2-R0-head-6.json` + records `"resolved_head_dtype": "torch.bfloat16"` with `hidden` [2,128], `weight` [128,128] and `logits` [2,128], + all `torch.bfloat16`, for decode step 6 of the L33/C2/R0 workload — the exact-tie step of + `.agents/specs/rocm-bf16-moe.md`. +- Local anchor: `src/vllm/model_executor/models/qwen3_moe.cpp:297-301` (the F32 store), `include/vt/ops.h:2652` + (`vt::Matmul`), `include/vt/ops.h:5619` (`vt::CastF32`), `src/vt/rocm/rocm_matmul_hipblaslt.hip:491` + (`MatmulKernelRocm`, `hipblasGemmEx` with an F32 compute type and an F32-or-BF16 output store), + `src/vt/rocm/rocm_dense_basic.hip:408` (`CastF32KernelRocm`). + +## Measurement that decides the change + +Harness: `tests/vllm/model_executor/test_qwen3_moe_lm_head_bf16.cpp`, case +"qwen3 MoE LM-head BF16 boundary: primary artifact replay". +It loads the three `.bin` files, interprets them from the JSON metadata, uploads the weight in the `[H, vocab]` +orientation the loader builds (`src/vllm/model_executor/models/qwen3_moe_weights.cpp:144-149`, +`LoadBf16Transposed("lm_head.weight")`), and runs the operators the forward itself calls: +`vt::Matmul` into an F32 buffer (arm a), `vt::Matmul` into a BF16 buffer followed by `vt::CastF32` (arm b), and +`logits.bin` widened to F32 (arm c). Arm a′ is arm a narrowed on the host with round-to-nearest-even. + +One command, no model checkpoint: + +```sh +VT_MOE_HEAD_FIXTURE=/home/vikash/.cache/rdna3-moe-impl/preserved/oracle-diagnostic-2 \ + ./build-head-hip/tests/test_qwen3_moe_lm_head_bf16 +``` + +Input hashes: + +| File | sha256 | +|---|---| +| `L33-C2-R0-head-6-hidden.bin` | `213770e378ff6f2bc06635b47d46ff623b3a8b53075fcd73b23f7b99f43cf0b6` | +| `L33-C2-R0-head-6-weight.bin` | `a174ad23ec6021ec75759ea580c08385482e528c102111d485cd09022a5a1ef4` | +| `L33-C2-R0-head-6-logits.bin` | `bc92004872f6f80ae93d87a867710868987f591a5ad63308cb857858a4a0a434` | + +Measured on 2026-09-10 at base `691b7af30`. The CPU and ROCm arms are identical value for value, so one table +carries both; `device=ROCM` is the production policy and was executed on `gfx1100` (RX 7900 XTX, ROCm 7.15.26333). + +| Row | Arm | argmax | top-2 (value) | max abs diff vs (c) | mean abs diff vs (c) | b==c | +|---|---|---|---|---|---|---| +| 0 | (a) F32, today | 118 | 118: 0.335236, 63: 0.334908 | 9.23157e-04 | 1.58187e-04 | 128/128 differ | +| 0 | (b) BF16 + `CastF32` | 118 | 118: 0.335938, 63: 0.333984 | 0 | 0 | **EQUAL, 0/128** | +| 0 | (c) primary `logits.bin` | 118 | 118: 0.335938, 63: 0.333984 | — | — | — | +| 1 | (a) F32, today | 85 | 85: 0.408177, 121: 0.334220 | 9.72956e-04 | 1.44222e-04 | 128/128 differ | +| 1 | (b) BF16 + `CastF32` | 85 | 85: 0.408203, 121: 0.333984 | 0 | 0 | **EQUAL, 0/128** | +| 1 | (c) primary `logits.bin` | 85 | 85: 0.408203, 121: 0.333984 | — | — | — | + +`max|a-b|` equals `max|a-c|` for both rows, and arm (a′) equals (c) word for word on both rows +(0/128 differing, `max|a'-c| = 0`). Two consequences: + +- The device BF16 store, the host round-to-nearest-even narrowing, and the primary's own words agree, so the only + difference between (a) and (b) is the store rounding, not the reduction. +- The F32 boundary's error is bounded by half a BF16 ulp (2^-9 = 1.953125e-03 at this magnitude): a dtype gap, + not an arithmetic defect. + +### Does the boundary alone change the step-6 argmax? + +No, on both available inputs. + +1. On the primary's own captured head input, (b) and (c) choose 118, and so does the F32 arm (a). The mirror is + therefore invisible to this row's token at step 6 while the native run is fed the primary's hidden state: + what separates native from the primary's concurrency-2 answer is the hidden state (#3115), not the head. +2. On the native production step-6 logits themselves, recorded in + `/home/vikash/.cache/residual-norm-repair1/green-cc9d4f565/production-fusion-1.json` (identical bytes to + `/home/vikash/.cache/moe-6fd1650c4-tmp/gpu-run-da0ce377b.json`), the F32 margin is + `63: 0.335850269` against `118: 0.335278690` (+5.71579e-04). Round-to-nearest-even narrowing maps **both** to + the same word `0x3eac = 0.3359375` — an exact tie — and the native lowest-index tie-break + (`src/vt/rocm/rocm_dense_basic.hip:170`, `v > best || (v == best && j < arg)`) keeps 63, which is the token the + native path already emits at all three repeats of concurrency 1 and concurrency 2. Request 1 is unaffected + (85 before and after). + +Reproduce the second measurement with +`PYTHONPATH=/home/vikash/.cache/rdna3-moe-impl/numpy-only-python python3 /home/vikash/.cache/moe-head-bf16/native-step6-narrowing.py /home/vikash/.cache/residual-norm-repair1/green-cc9d4f565/production-fusion-1.json`. +It is a host-side narrowing of a recorded native run, not a device execution, and is labelled as such. + +### Decision + +**Supported, as a correctness mirror only.** The measurement supports the boundary change for these reasons: + +- Given identical head inputs, the mirrored boundary reproduces the primary's logits element for element (0/128 + differing words on both rows), while the shipped F32 boundary does not (up to 9.73e-04 = half a BF16 ulp). + "vLLM is the reference" makes the primary's head output dtype the contract, and this removes a measured, + quantifiable divergence rather than an aesthetic one. +- The change does not move the decode step-6 token on either measured input, so landing it cannot by itself + explain, or repair, the 63/118 tie. + +**Not supported**, and explicitly out of scope: any claim that this change fixes the tie, improves hidden-state +parity, or makes the native run adopt the primary's concurrency-2 answer. The issue's own dependency statement +holds: #3115 remains required. + +## Design + +Where the narrowing belongs, and what must not move: + +1. `include/vllm/model_executor/models/lm_head_projection.h` (new) exposes + `vllm::lm_head::Project(dense_attn::Dev, const vt::Tensor& src, const vt::Tensor& lm, bool tied)`, which returns + an owning `dense_attn::DBuf` of `[n_out, vocab]` **F32**. The function is the whole head: it allocates the BF16 + projection buffer, calls `vt::Matmul` (untied, `lm` is `[H, vocab]`) or `vt::MatmulBT` (tied, `lm` is + `[vocab, H]`), then widens with `vt::CastF32` into the F32 buffer it returns. +2. `src/vllm/model_executor/models/qwen3_moe.cpp` `ForwardLayers` calls that seam and nothing else. The returned + buffer keeps the dtype, rank and shape it has today, so every downstream consumer stays valid: the device-logits + view (`ViewDeviceLogits`, `:337-345`), the owning wrapper (`WrapDeviceLogits`, `:321-331`), the host download in + `Qwen3MoeModel::Forward` (`:411-423`), the captured graph slot `SizeSlot::logits` (`:514`, `:636`, `:710`) and + the sampler all read `DType::kF32`. +3. The narrowing is a device store and the widening is the shared op — no host round-trip, no new op, no + per-element loop in the model, and no change to `vt::Matmul`'s contract (it already admits a BF16 output, + `include/vt/ops.h:2649-2652`, `src/vt/ops.cpp:125-126`). +4. Why a seam rather than three inline lines: the boundary is a production behavior that no checkpoint-free test + could otherwise execute, and the projection is the same operator sequence in both places. It is not a parallel + path — `ForwardLayers` has exactly one head call site and it is this function. (`qwen3_5_moe_block.h` is the + same pattern for the MoE block: an internal body exposed to a second caller over primitive `vt::` types.) + +Graph and pool consequences, stated because the residual spec warns that "changing only the head buffer dtype +would leave those views and host copies invalid" (`.agents/specs/rocm-residual-norm.md:325`): + +- The returned F32 buffer, its address and its dtype are unchanged, so `ViewDeviceLogits`/`WrapDeviceLogits` and + the captured slot stay byte-compatible. +- The BF16 head buffer is one additional transient `DBuf` per forward ([S, vocab] bf16 = half the bytes of the F32 + result). The cold eager pre-warm step at a padded size runs the same `ForwardLayers` call + (`qwen3_moe.cpp:727-731`), so its size class is already in the `DevicePool` free list when the capture at that + size runs — the capture's no-allocation requirement is preserved by construction. + +## Risks + +| Risk | Handling | +|---|---| +| The emitted token becomes tie-break dependent at steps whose top-2 margin is under one BF16 ulp. | Measured: it does not move at step 6 on either input. The operator's 18-workload membership gate is rerun at the frozen head; a moved position is a finding, not a tolerated difference. | +| The change is invisible to the token gate and could therefore land unreached. | The focused artifact replay is the gate; the reviewer mutates the production call site as well as the boundary, and the operator's production run exercises the code path. | +| An extra BF16 buffer per step changes memory or capture behavior. | One transient pooled block per forward, half the F32 result's bytes, warmed by the same cold step; no persistent allocation, no new residency. | +| A future reader mistakes this for the tie repair. | The issue, this spec, the commit body and the header comment all state that hidden-state parity (#3115) is still required. | +| The replay pins the NN (`vt::Matmul`) orientation only, so a future switch to `MatmulBT` for the untied head would not be covered. | Recorded as owed below; the untied arm is `Matmul` today (`qwen3_moe.cpp:300-301`). | + +## Tests + +- Focused, red first: `tests/vllm/model_executor/test_qwen3_moe_lm_head_bf16.cpp`. + - `primary artifact replay` — the three-arm measurement above, kept as a regression witness for the operators. + - `the production head projection mirrors the primary` — calls `vllm::lm_head::Project` with the primary's exact + hidden and weight bytes on every available device and requires the returned F32 logits to equal the primary's + widened BF16 words element for element, to keep the F32 `[rows, vocab]` shape, and to select the primary's + argmax. It skips itself (reported skipped by doctest) without `VT_MOE_HEAD_FIXTURE`, and the binary exits 77 + (CTest: Skipped) on that run. + - `the production forward returns BF16 logits` — the call-site case, and the one a reverted call site must + redden. It drives `Qwen3MoeModel::Forward` with a degenerate but legal zero-decoder-layer config (embed -> + final RMSNorm -> lm_head) and synthetic BF16 weights, so it needs no checkpoint, no capture and no fixture + directory, and it requires every logit the sampler would receive to be a BF16 word widened to F32. It carries no + skip decorator, so it runs and is reported in the fixture-absent run too, where the two replay cases skip; a + failure in it keeps doctest's non-zero exit instead of being folded into the run's 77. +- Red before implementation: with the production projection storing F32 (the pre-change behavior, extracted + verbatim), the projection case reports 256/256 differing logits per device at `max_abs = 9.72956e-04`, and the + forward case reports 128/128 logits that are not BF16-representable. +- Mutation (IMP-MUTATE), two independent guarantees, each restored byte for byte and re-verified by sha256: + 1. Delete the BF16 narrowing and the `vt::CastF32` inside the seam (the pre-change body) — the projection case + reddens, 2 `CHECK` failures, one per device; the forward case and the measurement case stay green. + 2. Revert the forward's call site to the inline F32 `vt::Matmul` — the forward case reddens, 128/128 logits not + BF16-representable; the two seam cases stay green, which is exactly the coverage split the two cases exist + to make visible. +- Existing gates that must stay green: `test_rocm_moe_bf16` (needs the operator's GPU and fixture), + `test_rocm_moe_reference_set`, and the frozen-head preflight. + +## Gates + +| Gate | Requirement | +|---|---| +| Focused green | `build-head-hip/tests/test_qwen3_moe_lm_head_bf16` exits 0 with the fixture set, on CPU and ROCm. | +| Focused, fixture absent | The same binary reports the two replay cases skipped, runs the fixture-free call-site case on CPU and ROCm, and exits 77 (CTest: Skipped); a failing case exits non-zero rather than being folded into that 77. | +| Red and mutation | The pre-change form reddens the production case; the reviewer reproduces it and restores byte-for-byte. | +| Preflight | `scripts/agent-preflight.sh --staged` at the implementation head, log under `/home/vikash/.cache/moe-head-bf16/`. | +| Production token gate | Operator only: the 18-workload membership gate in `tests/vllm/models/test_rocm_moe_bf16.cpp` at the frozen head, with the recorded concurrency-1/-2 reference set. | +| Reachability | The production forward calls the seam; the zero-layer case reddens when that call site is reverted, and the operator's run executes it in the real checkpoint path. | + +## Evidence + +- Measurement log: `/home/vikash/.cache/moe-head-bf16/measure-replay.log` and `measure-replay-record.txt` + (command, head, exit status, the a/a′/b/c table for both rows on CPU and ROCm). +- Native step-6 narrowing log: `/home/vikash/.cache/moe-head-bf16/native-step6-narrowing.txt`. +- Build log: `/home/vikash/.cache/moe-head-bf16/build-lib.log`, `build-measure.log`. +- Preflight log: `/home/vikash/.cache/moe-head-bf16/preflight.log`. +- Fixture-absent run of the repaired suite: `/home/vikash/.cache/moe-head-bf16/implt-no-fixture.log` — exit 77, + `1 passed | 0 failed | 2 skipped` with 10 assertions, both replay cases reported skipped and the call-site case + reported per device. With the fixture: `implt-with-fixture.log` — exit 0, 3 cases, 126/126 assertions. +- Repair mutations (scratch seam revert, restored byte-for-byte at sha256 + `26682c309b39264e4ecf349023652affad27d07818ff0a25a675c2cee15cab3c`): `implt-mut1-no-fixture.log` exits 1 with 2 + failing assertions when the failure guard is present, `implt-mut2-no-fixture.log` reports a masked 77 when it is + disabled, and `implt-mut3-no-fixture.log` exits 1 through the in-case `REQUIRE_MESSAGE` guard when the skip decorator + is removed. CTest reports Skipped without the fixture and Passed with it: `implt-ctest.log`. + +## Stop conditions + +- `STOP-VERIFY` (BLOCKED): the focused case or the preflight cannot pass at the implementation head. +- `STOP-AUTHORITY` (BLOCKED): the repair requires a file outside the head-output boundary, or the token gate needs + a device or checkpoint the operator has not provided. +- `STOP-DECISION` (NEEDS_DECISION): the operator's 18-workload gate moves a token position, which would make this + boundary a token-visible change rather than the fidelity mirror this spec measures. + +## Owed + +- The 18-workload production token gate at the frozen head is owed by the operator; this row's record keeps its + complete token gate open until that receipt exists. +- Hidden-state parity (#3115) is owed by the BF16-MoE row and is not delivered here. +- The replay covers the NN (`vt::Matmul`) untied orientation only. A future untied-head switch to `MatmulBT`, and + the tied `MatmulBT` branch, need their own artifact replay before either can claim the boundary. diff --git a/.agents/specs/rocm-residual-norm.md b/.agents/specs/rocm-residual-norm.md new file mode 100644 index 0000000000..53a4c8dfed --- /dev/null +++ b/.agents/specs/rocm-residual-norm.md @@ -0,0 +1,446 @@ +# Preserve compiled residual expressions on ROCm + +Owning row: `BACKEND-ROCM-RESIDUAL-NORM` +Parent row: `BACKEND-ROCM`, whose lifecycle remains `ACTIVE`. +Owner: the delegated residual-norm implementer, fresh reviewer, and coordinating operator. +Issue: [#3103](https://github.com/mudler/vllm.cpp/issues/3103), `ISSUE-GH-3103`. +Base: `9083a5fb35822e86b31cf786e0c80150dfae5478`. +Integration: one pull request under the repository default. Commit this spec before implementation. +The implementation pull request closes #3103 when its required gates pass and it lands. + +## Now + +State: `SPIKE`. The shared repair is implemented and the focused production witness passes on `gfx1100`. +The committed specification precedes both the captured red and the implementation. +The complete hardware gate and fresh mutation review remain required before acceptance. +The pinned upstream export passes all 264 cases; native comparison passes 528 CPU/ROCm executions and 22,576 assertions. +The operator found a blocking-stream setup error in the new nondefault-stream fixture. +The scoped fixture repair uses an explicitly owned nonblocking stream; its hardware mutation gate remains pending. +The related MoE production token gate remains open; its baseline failed at six generated positions. +No performance result is accepted before the complete declared token gate passes. + +## Problem and scope + +The compiled primary normalizes an FP32 residual expression before narrowing its result. +The native ROCm path first stores and reloads that expression as BF16. +The compiler also retains the post-attention operands across MoE and recomputes their sum at the next normalization. +Changing one rounding instruction does not preserve this complete lifetime. + +Extend the shared residual-normalization and `vt::FusedChain` seams to represent the executing primary expression. +Wire those semantics into the registered Qwen3 MoE production forward. +Keep every physical model residual and activation buffer at the primary's BF16 dtype. +Keep FP32 arithmetic in registers and reduction scratch where the primary uses FP32 arithmetic. +Do not allocate a persistent FP32 residual stream. + +Preserve existing operation signatures, recipe defaults, and legacy callers unless a separately proved correction requires a scoped update. +This work changes neither the upstream pin nor CI, checkers, quantized experts, or PR #2782. +It does not repair the separate resource-device issue #3100. +It does not silently broaden this row to attention, LM-head, or grouped-router implementations. +Those gaps still require implementation in the enclosing parity campaign. + +## Inventory + +| Stable ID | Upstream source | Local anchor | Tests and evidence | State | +|---|---|---|---|---| +| `ROCM-RESIDUAL-EXPR-NORM` | `vllm/ir/ops/layernorm.py::fused_add_rms_norm`, generated `ckic6h6`, `ctj2x6`, and `c5slugd` kernels | `src/vt/rocm/rocm_residual_rmsnorm.hip::ResidualRmsNormRowKernel`, `vt::ResidualRmsNorm` | Captured row-zero values and ordered-expression cases | `SPIKE` | +| `ROCM-RESIDUAL-EXPR-FUSION` | Executed normalization partitions described below | `include/vt/fused_recipe.h::FStep`, `src/vt/ops.cpp::FusedChainCompositeImpl` | Composite/native equivalence and boundary mutations | `SPIKE` | +| `ROCM-RESIDUAL-EXPR-FORWARD` | `vllm/model_executor/models/qwen3_moe.py::Qwen3MoeDecoderLayer` | `src/vllm/model_executor/models/qwen3_moe.cpp::RunMoeLayer`, `ForwardLayers` | Registered loader/forward witness, full token gate, lifetime and wiring mutations | `SPIKE` | + +Every inventory item belongs to this spec and #3103. +The parent backend row does not change lifecycle in this slice. + +## Pinned source and measured evidence + +The primary is vLLM `e126687a9a828d513c01a07cd69f025f27d63280`. +The verified executing source checkout is `/home/vikash/oracle/gfx1100-active-2773/source`. +The older shared `VLLM_SOURCE` checkout reports `5559679229` and is not this evidence's source revision. +These paths identify measured artifacts. They are not environment defaults for another developer. + +The primary selects default compilation, production graphs, native IR normalization, `ROCM_ATTN`, and Triton experts. +The resolved engine configuration appears in `oracle-diagnostic-2.log` at lines 15 and 34 to 35. +No eager denominator is allowed. + +Evidence root: `/home/vikash/vllm.cpp-rdna3-moe-impl/build-rdna3-moe-hip/evidence`. +The handoff is `norm-gap-handoff/handoff.md`. +Its SHA256 is `9850ee49c93bb82ddfca1a811fb422337c33ea09335e749fbf3c42dfcaa0400a`. +The manifest is `norm-gap-handoff/manifest.json`. +Its SHA256 is `8531cb41f0168d27455b2dc8062a57d5db7afd4e0822013332f398bbf2239c65`. +The spec author independently verified all 14 entries against their byte counts and hashes. + +The six generated modules retain complete code, original paths, hashes, and allocation dtypes. +Their full basenames are in the manifest. Short names below identify those immutable entries. + +| Module | Executing evidence | +|---|---| +| `c4r2zan` | BF16 embedding, first input norm, QKV, and the Q/K preamble | +| `cz7ckbn` | Layer-zero o_proj, post-attention norm, MoE, next input norm, and QKV | +| `ckic6h6` | Two-input post-attention sum and norm, no residual store | +| `ctj2x6` | Ordered three-input norm and BF16 residual store | +| `cywagh7` | Layer-one o_proj, post-attention norm, MoE, and final norm | +| `c5slugd` | Ordered three-input final norm, no residual consumer | + +At the source level, `vllm/ir/ops/layernorm.py::fused_add_rms_norm`, lines 44 to 62, adds in FP32. +It declares BF16 residual and gamma boundaries that the default compiler can eliminate. +The generated executing kernels determine which boundaries survive. +An eager Python expression alone is insufficient evidence for the compiled contract. + +At the inspected feature head, `git log -S'ResRound'` attributes the native boundary to `50b0709b3`. +`git log -S'lm_head'` identifies the earlier Qwen3 MoE forward commits `673b2a84a` and `c56ab287d`. +The relevant residual, model, router, and dense-attention files are unchanged between this base and fetched main `96c5e4719`. +The live issue #3103 is open. No open residual repair pull request was found during this audit. + +## Required arithmetic and lifetime + +Let `f(x)` load a BF16 operand as FP32 and `b(x)` round FP32 to BF16. +Let every arithmetic operation below use FP32, in the shown order. +Let `N(s,w)` compute `b((s * rsqrt(mean(s*s) + eps)) * f(w))`. +There is no BF16 conversion between normalization and gamma multiplication in the observed generated kernels. +The reduction order and reciprocal-square-root implementation require executing-kernel evidence in addition to the expression. + +### First input normalization + +The first layer consumes the BF16 embedding `r0` and emits `N(f(r0), w_in0)`. +The retained residual is `r0`, still BF16. +The native zero-residual initialization is equivalent only if it preserves these values and the actual output. +The saved first QKV input and output match exactly on both sides. +Preserve that control when adding the new lifetime. + +### Post-attention normalization without materialization + +For BF16 attention projection `a` and BF16 materialized residual `r`: + +```text +s = f(a) + f(r) +n = N(s, w_post) +``` + +Read `a` and `r` without modifying either input. +Do not store `b(s)` and use it for normalization or for the next layer. +The generated `ckic6h6` kernel loads the two operands at lines 33 to 38. +It normalizes their FP32 sum at lines 39 to 50 and stores BF16 normalized outputs at lines 51 to 52. +Neither output is a residual sum. +The two normalized copies serve the router and expert preparation in the primary. +A shared native buffer can serve both consumers if values and lifetime remain equivalent. + +Retain ownership of `a` and `r` through the MoE invocation. +The current local `attn` is a stack-local owning buffer in `RunMoeLayer`. +Either retain that owner until the next-input or final norm, or execute the consuming norm before the owner leaves scope. +The driver can schedule the next normalization immediately after MoE, matching the generated partition, if it prevents duplicate normalization. +A view without its owner is not a complete repair. + +### Next input normalization with materialization + +For BF16 MoE result `m` and the retained BF16 operands `a` and `r`: + +```text +s_attention = f(a) + f(r) +s_next = f(m) + s_attention +r_next = b(s_next) +n_next = N(s_next, w_in_next) +``` + +Preserve the association `m + (a + r)`. +Do not use `(m + a) + r` or `m + b(a + r)`. +Store `r_next` as BF16 while normalizing `s_next` before its BF16 rounding. +Do not reload the rounded `r_next` for normalization. +The next QKV consumes `n_next` and the next attention partition retains `r_next`. +The generated `ctj2x6` kernel loads all operands at lines 33 to 39. +It evaluates the ordered additions at lines 40 to 43 and stores residual and norm outputs at lines 57 to 58. +`cz7ckbn`, lines 1028 to 1040 and 1076, records the allocations, call, QKV use, and returned residual. + +### Final normalization + +The final normalization evaluates the same ordered three-term expression. +It emits `N(s_next, w_final)` and has no residual output consumer. +The generated `c5slugd`, lines 33 to 56, overwrites the MoE buffer with its BF16 normalized output. +An independent BF16 output allocation is permissible if ownership and values are equivalent. +Do not create a final residual allocation merely to make the existing in-place API fit. + +## Shared design + +Add a typed shared operation or explicit descriptor for two-input and ordered three-input residual expressions. +Its inputs are read-only activation operands and gamma. +Its outputs are a required normalized tensor and an optional materialized residual tensor. +The descriptor explicitly identifies operand count, addition order, and residual materialization. +Do not alter the existing provider function-pointer signature through a cast. +Validate shape, dtype, stride, device, output aliases, and invalid descriptor combinations before dispatch. +The new native path must select the complete capability before model execution. +Select the measured compiled lifetime through a shared backend capability or numeric-policy seam. +CPU reference registration alone must not enable that lifetime for other production backends. +Preserve existing CPU and CUDA model defaults unless matching executing-oracle evidence requires a scoped correction. + +Represent this operation in `vt::FusedChain` and its composite realization. +The existing composite folds one in-place `kAdd` into `RmsNorm`, at `src/vt/ops.cpp::FusedChainCompositeImpl`, lines 1159 to 1185. +That fold cannot express a read-only sum, three operands, and a separately materialized residual. +Extend the shared representation instead of adding a model-specific HIP sequence. +The current binding has eight operand slots but a recipe step admits three inputs. +A three-activation operation plus gamma needs an explicit, validated representation of all four inputs. +Extend that capacity or add a typed residual-expression descriptor without overloading another field's meaning. + +Provide a CPU reference and the ROCm implementation through the same typed operation. +Composite and optimized execution must agree on arithmetic boundaries and observable buffers. +Route both adopted-fusion and explicit-unfused model modes through the new shared semantics. +Keep the old standard and Gemma recipes unchanged for callers that retain their previous contract. +Do not turn a legacy `ResRound` deletion into a global numeric change. + +Use row-local FP32 registers or deterministic reduction scratch. +No persistent FP32 activation or residual tensor is allowed on this BF16 path. +If a kernel writes a residual that aliases an input, preserve all needed unrounded values before writing. +Protect reads and stores across threads, including the second normalization pass over wide rows. +Support the admitted shapes with tail masks and deterministic reductions. +Use the queue's device and stream and preserve the caller's ambient device. +Allocate no new storage during graph capture. +Keep every captured buffer alive and stable for graph replay. + +## Smallest failing tests + +### Production row-zero witness + +Start with the unchanged fixture from [the MoE spec](rocm-bf16-moe.md). +Its config SHA256 is `321926020ada026d8dd85f74543dd12b6426301706b4e1f4fbca60a44b871cea`. +Its safetensors SHA256 is `96cd7f30fee496c69782af2813e438e47d7b026598b0f20a05049c522b279af8`. +Preserve every generator seed, tensor order, prompt, sampling value, cohort, and repeat. + +Run the test through `ModelRegistry::Load` and `ModelRegistry::Forward`. +Observe layer-zero post-attention norm on step zero of L33/C2/R0. +The actual row-zero attention projection, embedding residual, and gamma match between runtimes. +Require the 128 observed output values to match the captured primary values exactly. +The current native result differs in 33 words, first at element 3: +`-0.1865234375` instead of `-0.185546875`. +The maximum absolute difference is `0.00390625`. +Capture this red before implementing the repair. + +The native descriptor file is `native-diagnostic-1/L33-C2-R0.json`. +Files 14, 15, and 16 contain the actual input, residual, and output for this normalization. +The primary input is `oracle-diagnostic-2/L33-C2-R0-dense-1-out.bin`. +The primary normalized output is `oracle-diagnostic-2/L33-C2-R0-dense-2-x.bin`. +Read the gamma from the exact fixture, as `diagnose-residual-row0.py` does. +`residual-row0-diagnostic.json` reproduces every native word with rounding and every primary word without rounding. +Its SHA256 is `421b5855cc35b095635fa7268f98803eb26ed0f267d15e9a5e41e01fa6703d2e`. +An isolated operation replay supplements this production-entry test and cannot replace it. + +### Complete expression witnesses + +Add exact cases where intermediate BF16 rounding changes the normalized output. +Add three-input cases that distinguish addition association, early materialization, and post-store normalization. +Use exact BF16 operand witnesses before randomized cases: + +| Guarantee | `a` | `r` | `m` | Required FP32 expression | Mutated result | +|---|---:|---:|---:|---:|---:| +| Ordered additions | 256 | -256 | `2^-17` | `m + (a+r) = 2^-17` | `(m+a)+r = 0` | +| Deferred materialization | 1 | `2^-8` | `2^-8` | `b(m+(a+r)) = 1.0078125` | `b(m+b(a+r)) = 1` | + +These are scalar residual witnesses. Embed them in admitted row shapes and validate the associated normalization independently. +Cover both materialized and absent residual output, and verify every read-only input remains unchanged. +Cover final-output aliasing only when the descriptor explicitly permits it. +Cover the first layer, a subsequent layer, and the final norm through the production forward. +Capture the primary's later residual output and norm outputs on original executing inputs. +Do not treat reconstruction alone as an observed primary output. + +Port applicable pinned normalization tests from `tests/kernels/core/test_layernorm.py::test_rms_norm` +and `tests/kernels/ir/test_layernorm.py::TestRMSNorm` and `TestFusedAddRMSNorm`. +Preserve their parameters, fixtures, seeds, tolerances, failures, and pin attribution for every applicable mode. +Record explicit applicability for unimplemented activation dtypes rather than adding a silent skip. +The BF16 compiled-expression exact witnesses supplement the original tests. +Eager native gamma rounding does not define the separately observed compiled expression. + +## Verification and mutation gates + +1. Commit the spec, then capture the smallest production and component red results. +2. Build the CPU reference and ROCm implementation with recorded compiler and binary hashes. +3. Run shared descriptor tests, existing RMSNorm and fusion regressions, and the new exact witnesses. +4. Run the unchanged complete MoE production token gate against the pinned production oracle. +5. Run explicit fused and unfused modes, graph capture/replay, two streams, wide rows, tails, and device selection checks. +6. Capture actual allocation dtypes, strides, operand ownership, generated kernels, and materialization points on both sides. +7. Run full staged preflight and exact-range record, path, style, and trailer checks. +8. Obtain fresh static and mutation review on the immutable implementation head. +9. Have the operator rerun the applicable hardware gates before acceptance. + +The full model workload is L in `{1,3,33}`, concurrency in `{1,2}`, eight greedy tokens, and three repeats. +The unchanged `test_rocm_moe_bf16` consumes the exact fixture and primary `production.json` through its existing environment variables. +The authoritative cohort result is `oracle-selection-6/production.json` under the evidence root. +Its SHA256 is `f3d27a95ba71bddcae38c3defc21bf3f3ff32ba9373a09faf949ead68b74d7ad`. +Add the shared operation witness as `test_ops_residual_rmsnorm` and the production witness to `test_rocm_moe_bf16`. +Use separate CPU and HIP build directories. The focused gate commands are: + +```sh +ctest --test-dir build-residual-cpu --output-on-failure -R '^(test_ops_residual_rmsnorm|test_ops_rmsnorm|test_ops_rmsnorm_weight_dtype|test_ops_fused_chain|test_fused_chain_additivity)$' +ctest --test-dir build-residual-hip --output-on-failure -R '^(test_ops_residual_rmsnorm|test_ops_rmsnorm|test_ops_rmsnorm_weight_dtype|test_ops_fused_chain|test_fused_chain_additivity|test_rocm_moe_bf16)$' +scripts/agent-preflight.sh --staged +``` + +The operator runs the HIP command inside the required mutex or lease. +Set `VT_ROCM_MOE_FIXTURE`, `VT_ROCM_MOE_ORACLE`, and `VT_ROCM_MOE_OUTPUT` from the immutable fixture, primary cohort result, and fresh output location. +Verify each GPU test reports actual native cases and nonzero assertions. +A device-less skip is not a passing ROCm gate. +The native diagnostic command and environment are retained in `native-diagnostic-1-operator-receipt.json`. +The primary diagnostic recipe is `snapshots/oracle-diagnostic-2/command.json` with the operator's read-only runtime-mount adaptation in `oracle-diagnostic-2-operator-receipt.json`. +The operator substitutes fresh output paths and the reviewed binary, retaining model bytes and engine policy. +All GPU execution stays under the operator's recorded device authority and required mutex or lease. + +The reviewer must independently mutate each guarantee: + +- Round `a+r` before variance. +- Store and reload `b(a+r)` across MoE. +- Normalize from stored `r_next` instead of its unrounded expression. +- Reassociate the three additions. +- Add a BF16 conversion before gamma multiplication. +- Remove the subsequent BF16 residual materialization. +- Modify a read-only input or release the retained attention owner too early. +- Delete the production call site or force the old residual path. +- Bypass descriptor validation or use the wrong stream/device. + +Each mutation must fail its intended focused test with nonzero assertions. +The reviewer restores the scratch tree byte-for-byte after each mutation. +A failing full-token baseline is not an adequate witness for a new mutation. +The focused failure must identify the mutated guarantee. + +## Dependencies and remaining capture work + +The residual repair is implementable from the saved first-stage data and generated lifetime. +Later exact witnesses require an operator capture of the materialized primary residual and normalized output. +This is a measurement task on the available oracle, not an unresolved product decision. + +Attention already differs before the residual norm at row 1, element 7 of the first attention output. +For L33/C2/R0, native `0.0888671875` differs from primary `0.08935546875`. +The first QKV input, weight, and output are exact. +The first attention-output comparison contains 2902 differing BF16 words. + +The generated `c4r2zan` preamble keeps normalized Q/K values in FP32 until after RoPE and reads a BF16 cos/sin cache. +The native `dense_attn::AttnBlock` stores BF16 Q/K norms before `RopeNeox` with FP32 cos/sin. +Capture actual Q/K immediately before attention, V, cos/sin, KV data and metadata, and the attention output on both sides. +The preserved capture `/home/vikash/.cache/rdna3-moe-attention-operator/results/L33-C2-R0-attention-0.json` and its ten binary siblings hold exactly those inputs and the output, and a validated CPU transcription of `prefix_prefill._fwd_kernel` separated the two hypotheses on them before any native byte of that boundary existed (#3115). +The native capture-and-replay instrument in `tests/vllm/models/test_rocm_moe_bf16.cpp` then measured both terms on the device: 2918 of 8448 words from the attention kernel on identical Q/K/V, and 1569 Q plus 1542 K words from the preamble, with the cos/sin table and the qkv projection byte-exact on both sides. +The pinned normal decoder prefill enters `prefix_prefill.py` through `chunked_prefill_paged_decode.py`. +Its probability-to-value product narrows probabilities to the V dtype. +The native D128/QG1 dispatch retains FP32 probabilities even during prefill. +Replay identical captured Q/K/V through both executing kernels before selecting that repair. +The cached generated attention module is `_fwd_kernel` in Triton cache directory `I7QZRC574ZPUQULYWLLB3LJHQ5RQVEIOAGBPLJ7M5YBI5ZVUMSCA`. +Its launch metadata and generated code are retained beside that module. + +The primary also stores BF16 LM-head output, while `ForwardLayers` currently produces F32 logits directly. +Complete primary hidden, weight, and logits inputs exist in `oracle-diagnostic-2/L33-C2-R0-head-6.json` and its binary files. +Replay those exact hidden and weight bytes through native projection with BF16 output, then shared `CastF32` before the current forward result boundary. +The pinned sampler converts logits to FP32 before sampling. +The local host download and graph result views already require FP32. +Changing only the head buffer dtype would leave those views and host copies invalid. +Do not attribute a head mismatch to GEMM until identical hidden inputs are supplied. +The current saved hidden inputs already differ upstream. +A CPU audit of step 6 confirms BF16 narrowing alone preserves the wrong native argmax, token 63. +The primary chooses token 118. Native BF16 top-two values are `63: 0.3359375` and `118: 0.333984375`. +The primary reverses those values, so the earlier hidden-state repairs remain necessary. + +DeepSeek-V2 and dots3-note router refusals are implementation debt in the parent backend campaign. +They do not appear in the Qwen3 fixture's executing path. +DeepSeek passes `n_group=1` even for its softmax default, which the current ROCm router rejects. +Dots3-note requires sigmoid scores, bias for selection only, and the one-group form. +Completing those modes needs a separate scoped router repair and production tests. +The broad #41 ownership does not make the implementations complete. + +## Owed + +This row owns #3103 directly. The related BF16 MoE row keeps its complete token gate open. +No performance or token requirement is waived by native/legacy agreement. +The assigned scoped issues for the attention and head repairs are #3115 (decode attention and Q/K preamble parity) and #3116 (the BF16 LM-head output boundary), both on the BF16 MoE row, filed from the six-position disposition analysis of commit `cc9d4f565`. +The parent `BACKEND-ROCM` issue #41 retains router debt until its scoped repair is assigned. + +## Implementation and evidence + +The implementation base is the committed spec `9ec19f80b600c9e713f7297638186e986c67a4d1`. +`ResidualNormDesc` distinguishes the two ordered expressions and optional BF16 residual materialization. +The typed operation validates every tensor and permitted exact alias before dispatch. +`FStep` admits four inputs, and the composite realizes the explicit residual opcode through that operation. +The CPU reference does not select the production policy. +The ROCm backend selects the compiled expression; CPU and CUDA retain the materialized policy. +`RunMoeLayer` consumes the retained attention owner immediately after MoE and transfers the next normalized BF16 buffer's ownership. +The final norm omits residual output. No persistent FP32 activation or residual is introduced. + +The native kernel uses a deterministic 256-lane FP32 reduction and reloads original operands before each store. +All variance reads finish before any permitted output alias is written. +Wide rows, padded row strides, and tails use the same arithmetic and ownership contract. +The launcher selects the queue's device and stream and restores the caller's ambient device. +The operation allocates no storage during graph capture. + +### Red before implementation + +Evidence root: `/home/vikash/.cache/residual-norm-impl`. +The `red-freeze` manifest pins unchanged product blobs, tests, model bytes, linked libraries, compiler flags, and binaries. +The CPU red exited 1 with 33 failures among 128 assertions. +Its binary SHA256 is `c5745f016d2033ec1da40d6b98cc29803c3f757d583f81bb9eb85665c97dabbf`. +The operator's production red exited 1 with 33 failures among 400 assertions. +Every captured attention, residual, and gamma word matched the primary before the failed norm comparison. +The production binary SHA256 is `f932a6c7b49c96111fb1868bf043287c69c7c4375beebcf4dd14535b42249024`. +The operator log SHA256 is `3baba8756846635e085587a99443e3b6c3d3268351fda356f9135809eb8fc49d`. + +### Focused green and later primary observation + +The operator ran `green-freeze-1/command.json` and `component-command.json` on the same physical GPU under the required mutex. +The production witness passed all 399 assertions, including 128 exact primary norm words. +Its binary SHA256 is `aa38607e0b88202e2d65e45f76c790b7f18c2bd6dd3cb62662c91892f6f56e00`. +The component gate passed 10,700 CPU and native ROCm assertions without a device skip. +Its binary SHA256 is `102f34696d79be728dd8c196087f05857c8b91fcc01a2994e6ec1f8fdb705aad`. +The operator verified 60 sealed inputs before and after those runs in `green-freeze-1/operator-receipts.json`. +The production log SHA256 is `e94ea05e06acc04408e37a73ede8bf7acc24071ec132d011755fb5534d6e0d53`. +The component log SHA256 is `63e2cfcf065030259deda3a59cf521db66bcead298d643e7d0b46b3837d1736a`. + +The `oracle-probe-v3` observer wraps original compiled launches with unchanged arguments and verifies their source ASTs. +It observes post-attention, next-input, final, and Q/K preamble launches in L33/C1/R0 and L33/C2/R0. +The operator verified ten capture records, 64 binary payloads, and every read-only input before and after the launches. +All 18 production cohorts, 216 tokens, and every logprob matched the paired observation-disabled control. +The observer script SHA256 is `789138a38dc2db167656356b4ffa18cab8e4141a3b96fb5f24f93bdd9f242512`. +The operator log SHA256 is `220ce248825706013fdaa0a640a5b5acd522d8e6d31f9fcc43244036ef3224a5`. +`oracle-probe-v3/operator-output-checks.json` records the independent comparison. +The actual next residual, next norm, and final norm first-row bytes are embedded with full-payload and slice hashes in `tests/support/residual_norm_later_fixture.h`. +The earlier v1 AST-parser failure and v2 singleton-stride serialization failure remain preserved beside the successful v3 evidence. +Neither failed observer is evidence for a primary arithmetic result. + +### Test applicability and remaining gates + +`tests/vt/residual_norm_upstream.py` executes the pinned core and IR normalization tests and exports their original BF16 fixtures. +It preserves both residual modes, every token count and width, row strides, device count, seed zero, epsilons, and original tolerances. +The C++ operation consumes the exported original native references and compares direct and shared-fusion dispatch exactly. +Weightless normalization is represented by an explicit unit BF16 gamma; plain normalization uses a zero base. +Torch registration and opcheck execute in the exporter. C++ validates its own descriptor and ownership rules. +F16/F32 activation modes are outside this measured BF16 policy and are explicitly refused, never silently skipped. +The new API has no variance-size override, Gemma modifier, or partial-width norm mode. + +The six focused CPU suites, including the existing Qwen3 MoE forward control, pass after adding the later witnesses. +`cpu-green-5.log` records that run. External upstream fixtures require `VT_RESIDUAL_NORM_UPSTREAM`; absence is reported as unexecuted. +The operator completed the upstream export: 144 core cases, 60 IR RMS cases, and 60 IR add cases. +All 1,452 payload hashes are verified, and the native comparison passes 528 CPU/ROCm executions with 22,576 assertions. +Expanded hardware tests, complete token equality in both fusion modes, staged preflight, and fresh mutation review remain required. +The original MoE gate still differs at six generated positions in the L33/C2 tail. + +The ROCm platform keeps `support_static_graph_mode()` false at this base in `src/vllm/platforms/rocm.cpp:91`. +Production registry decode therefore does not enter the graph driver. +The existing ROCm graph row owns that platform exclusion; this repair does not change it. +The new operation's hardware gate covers capture and repeated replay with live BF16 operands on both devices. +It also checks two queues, ambient-device restoration, wrong-device streams, and a blocked nondefault stream. +No production graph claim is inferred from that component gate. + +### Repair the blocked-stream fixture and its invalidated citation + +The operator's component run at `94ac5d1742c540fc5cbc2084d0fb74fdfd5dab41` fails before the stream-order check. +The test requests a backend queue, but `RocmBackend::CreateQueue` uses `hipStreamCreate` and returns flags zero. +The existing guarantee requires a stream independent of default-stream synchronization. +The fixture therefore owns an explicit `hipStreamNonBlocking` stream within the existing device scope. +Its release guard synchronizes the callback before destroying callback state or operands during normal and exceptional exits. +The flag assertion and real wrong-stream mutation remain required; the backend default does not change. + +The added residual descriptor include moves `vt::Backend` to line 23. +Repair only that citation in the `BACKEND-PLATFORM` matrix row, starting from the complete target file. +The scoped proof requires every other matrix byte to remain unchanged. +The [repair evidence](../../docs/bench-evidence/rocm-residual-norm/README.md) +retains the operator red, record red/green, compiler recipe, upstream results, and unchanged-product proof. +The repaired CPU component binary passes six tests and 8,985 assertions, including all 264 upstream cases. +The hardware fixture, mutation, final preflight, fresh review, and operator rerun remain pending at this checkpoint. + +## Stop conditions + +Return `NEEDS_CONTEXT` when binding fixture bytes, generated code, primary runtime, or required GPU authority is unavailable. +Continue independent source, CPU, and spec work when a hardware capture awaits the operator. +Return `NEEDS_DECISION` only if executing evidence contradicts the fixed arithmetic or requires a new shared-seam exception. +Do not weaken exact comparisons, alter cohorts, switch the denominator to eager, or widen all residual storage. +A known correctable token mismatch keeps the gate failing and the enclosing campaign active. +Do not mark this row `DONE` or publish accepted performance before its required gates pass. +Add `## Outcome` when the row reaches `DONE`, recording measured behavior and rejected alternatives. diff --git a/CMakeLists.txt b/CMakeLists.txt index d91f9a495e..d5737e26a0 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1350,6 +1350,7 @@ add_library(vllm STATIC src/vt/cpu/cpu_backend.cpp src/vt/cpu/cpu_threadpool.cpp src/vt/cpu/cpu_ops.cpp + src/vt/cpu/cpu_residual_rmsnorm.cpp src/vt/cpu/cpu_minimax_h3.cpp src/vt/cpu/cpu_ltx2.cpp src/vt/cpu/cpu_ltx2_vae.cpp @@ -1716,6 +1717,7 @@ if(VLLM_CPP_HIP) src/vllm/platforms/rocm.cpp src/vt/rocm/rocm_backend.hip src/vt/rocm/rocm_rmsnorm.hip + src/vt/rocm/rocm_residual_rmsnorm.hip src/vt/rocm/rocm_embedding.hip src/vt/rocm/rocm_dense_basic.hip src/vt/rocm/rocm_matmul_hipblaslt.hip @@ -1726,6 +1728,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_fp8_channel_gemv.hip src/vt/rocm/rocm_moe_router.hip src/vt/rocm/rocm_moe_chain.hip + src/vt/rocm/rocm_moe_grouped_bf16.hip src/vt/rocm/rocm_grouped_gemm.hip src/vt/rocm/rocm_moe_gate_up_swiglu.hip src/vt/rocm/rocm_exl3.hip @@ -1744,6 +1747,7 @@ if(VLLM_CPP_HIP) set_source_files_properties( src/vt/rocm/rocm_backend.hip src/vt/rocm/rocm_rmsnorm.hip + src/vt/rocm/rocm_residual_rmsnorm.hip src/vt/rocm/rocm_embedding.hip src/vt/rocm/rocm_dense_basic.hip src/vt/rocm/rocm_matmul_hipblaslt.hip @@ -1754,6 +1758,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_fp8_channel_gemv.hip src/vt/rocm/rocm_moe_router.hip src/vt/rocm/rocm_moe_chain.hip + src/vt/rocm/rocm_moe_grouped_bf16.hip src/vt/rocm/rocm_grouped_gemm.hip src/vt/rocm/rocm_moe_gate_up_swiglu.hip src/vt/rocm/rocm_exl3.hip diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 18a2c3ef43..4bade57b4d 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -339,7 +339,7 @@ from a prompt: both refuse, naming what is missing. | CPU (x86, Arm i8mm; A76 assembly correct/default, llama speed gate open, and the closed 20-core floor ran a SUPERSEDED fork denominator rather than the stock `b10451` pin, re-take owed #1003) | ✅ `CPU_ATTN` registered (#1371/#1392, [spec](../.agents/specs/attn-validate-configuration.md)) | ◐ | ☐ | ✅ | | Metal (Apple Silicon) | ✅ builds under Apple Clang with project warnings promoted to errors, the Qwen3.5 MoE loader included; its layout-refusal path uses the same messages and behavior on every platform (#1054) | ☐ | ☐ | ✅ | | Vulkan | ◐ **EXL3 is now NATIVE** ([#2530](https://github.com/mudler/vllm.cpp/issues/2530), [spec](../.agents/specs/backend-vulkan-exl3.md)): `kCastF16` and `kExl3Gemm` were the only two ops an EXL3 checkpoint ran on the CPU reference tier on a Vulkan queue, MEASURED at two fallback notices and now at ZERO. `kExl3Gemm` is TRANSCRIBED from the portable CPU reference, not ported from the CUDA kernel, whose 90 KiB shared-memory budget alone exceeds Vulkan's 16 KiB guarantee before one reaches `mma.sync`, `ldmatrix`, `cp.async` or a grid-wide barrier Vulkan has at no version -- so the gate is BYTE equality with the CPU arm on all three codebooks and every width, not a tolerance. Verified on llvmpipe with NO GPU and NO lease, which is what makes it a CI gate rather than a lease gate; no real Vulkan GPU and no checkpoint run, and NO speed number is claimed on any axis. `kExl3MoeMlp` and `kExl3HadR128` are owed | ☐ | ☐ | ✅ | -| ROCm | W0: 5 gfx archs; dense/GDN all-native; 0.8B dispatch fixed. **M4:** Qwen3-0.6B/3.5-0.8B 16/16 (#41). **M3:** `ROCM_ATTN` registered (#1056/#1065, [spec](../.agents/specs/rocm-attn-backend.md)). **EXL3 is now NATIVE on `gfx1151`** ([#2433](https://github.com/mudler/vllm.cpp/issues/2433), [spec](../.agents/specs/backend-rocm-exl3.md)): `kCastF16` and `kExl3Gemm` were the only two ops an EXL3 checkpoint still ran on the CPU reference tier, and a 3.0bpw Llama-3.2-1B now completes with ZERO reference-tier hits at 8.27 tok/s warm against 0.83 for the same tree with the two registrations disabled -- byte-identical greedy output either way. `kExl3Gemm` is TRANSCRIBED from the portable CPU reference, not ported from the CUDA kernel, whose 90 KiB shared-memory budget alone does not fit AMD's 64 KiB LDS; the gate is therefore byte equality with the CPU arm rather than a tolerance. No BF16 denominator and no clock attribution: the control HUNG the GPU in the same lease ([#2511](https://github.com/mudler/vllm.cpp/issues/2511)). `kExl3MoeMlp` is owed. CPU parity open (#269) | Native ([ROCm guide](ROCM.md), device-specific wrapped-call-safe derived recount): full GDN, MoE combine/gate, keep-quant GEMM, EXL3 trellis GEMM; ctest-green gfx1151/1103/1100/1201/1200 (#41). **APU managed allocation is now MEASURED and NARROWED** ([#2511](https://github.com/mudler/vllm.cpp/issues/2511)): `hipMallocManaged` on a part reporting `PageableMemoryAccess = 0` faulted the GPU in 17 of 21 legs on `gfx1151`, against 0 of 21 for plain `hipMalloc`, and the plain arm is the first thing that has completed a gate-sized Qwen3.8-27B Q4_K run on that board. The managed branch now requires a device that can take a recoverable page fault, and `UnifiedMemory()` follows the allocator -- so gfx1151 and gfx1103 lose the CPU reference tier and an op with no native ROCm kernel refuses by name there, naming the attribute and `VT_ROCM_MANAGED_ALLOC=1`. | ✅ | ✅ | +| ROCm | W0: 5 gfx archs; dense/GDN all-native; 0.8B dispatch fixed. **M4:** Qwen3-0.6B/3.5-0.8B 16/16 (#41). **M3:** `ROCM_ATTN` registered (#1056/#1065, [spec](../.agents/specs/rocm-attn-backend.md)). **EXL3 is now NATIVE on `gfx1151`** ([#2433](https://github.com/mudler/vllm.cpp/issues/2433), [spec](../.agents/specs/backend-rocm-exl3.md)): `kCastF16` and `kExl3Gemm` were the only two ops an EXL3 checkpoint still ran on the CPU reference tier, and a 3.0bpw Llama-3.2-1B now completes with ZERO reference-tier hits at 8.27 tok/s warm against 0.83 for the same tree with the two registrations disabled -- byte-identical greedy output either way. `kExl3Gemm` is TRANSCRIBED from the portable CPU reference, not ported from the CUDA kernel, whose 90 KiB shared-memory budget alone does not fit AMD's 64 KiB LDS; the gate is therefore byte equality with the CPU arm rather than a tolerance. No BF16 denominator and no clock attribution: the control HUNG the GPU in the same lease ([#2511](https://github.com/mudler/vllm.cpp/issues/2511)). `kExl3MoeMlp` is owed. CPU parity open (#269). Native grouped BF16 expert kernels pass component gates on gfx1100; full Qwen3 MoE token parity remains open ([spec](../.agents/specs/rocm-bf16-moe.md)) | Native ([ROCm guide](ROCM.md), device-specific wrapped-call-safe derived recount): full GDN, MoE combine/gate, keep-quant GEMM, EXL3 trellis GEMM; ctest-green gfx1151/1103/1100/1201/1200 (#41). **APU managed allocation is now MEASURED and NARROWED** ([#2511](https://github.com/mudler/vllm.cpp/issues/2511)): `hipMallocManaged` on a part reporting `PageableMemoryAccess = 0` faulted the GPU in 17 of 21 legs on `gfx1151`, against 0 of 21 for plain `hipMalloc`, and the plain arm is the first thing that has completed a gate-sized Qwen3.8-27B Q4_K run on that board. The managed branch now requires a device that can take a recoverable page fault, and `UnifiedMemory()` follows the allocator -- so gfx1151 and gfx1103 lose the CPU reference tier and an op with no native ROCm kernel refuses by name there, naming the attribute and `VT_ROCM_MANAGED_ALLOC=1`. | ✅ | ✅ | | XPU / TPU | ☐ | ✅ | ◐ | ☐ | | Tenstorrent Blackhole | ◐ `ACTIVE`, OPT-125m 6/6; Qwen3-0.6B wired; Mistral-7B-v0.3 16/16 on P150 ([spec](../.agents/specs/tenstorrent-mistral.md)). 16x16 rerun and residual-RMS owed ([spec](../.agents/specs/tenstorrent-backend.md)) | ✅ | ☐ | ☐ | | Tenstorrent host-free decode | ◐ DEFAULT since #1604 (`0` opts out): no per-step host readback; 2.1x default-leg tok/s; both golden pairs re-adjudicated, both paged gates 16/16. Capture DEFAULT for Qwen3-dense since the #1625 flip (`0` opts out) — 27.7 vs 12.9 tok/s at c=1, 2.15x over the opt-out, [details](benchmarks/tt-capture-default-decode.md); Qwen3.5 dense (GDN) joined the default since [#2907](https://github.com/mudler/vllm.cpp/issues/2907), the decode-side GDN ops running device-resident under capture with [#2812](https://github.com/mudler/vllm.cpp/issues/2812)'s captured fatal fixed; 4B near-tie pair owed ([#2811](https://github.com/mudler/vllm.cpp/issues/2811)); async off (#1627) | ☐ | ☐ | ☐ | diff --git a/docs/ROCM.md b/docs/ROCM.md index 4cdbab373e..8c861d9091 100644 --- a/docs/ROCM.md +++ b/docs/ROCM.md @@ -106,6 +106,7 @@ Do not use a run with CPU fallbacks as a performance result. | Kernels | [`src/vt/rocm/`](../src/vt/rocm/) | Dense, GDN, attention, sampling, and the contributor-tested Gemma 4 FP8 MoE path | | Platform | [`src/vllm/platforms/rocm.cpp`](../src/vllm/platforms/rocm.cpp) | Runtime-verified on five gfx architectures | | Attention | [`src/vt/rocm/rocm_paged_attn.hip`](../src/vt/rocm/rocm_paged_attn.hip) | Native paged attention and the SharedK WMMA prefill path | +| Grouped BF16 experts | [`src/vt/rocm/rocm_moe_grouped_bf16.hip`](../src/vt/rocm/rocm_moe_grouped_bf16.hip) | Native gate/up, weighted down, and preweighted combine through shared MoE operations; component tests pass on gfx1100. The Qwen3 MoE production token gate remains open ([spec](../.agents/specs/rocm-bf16-moe.md)) | | Build | [`CMakeLists.txt`](../CMakeLists.txt) | `VLLM_CPP_HIP` configuration and build verified on five architectures | | Tests | [`tests/vt/test_rocm_backend.cpp`](../tests/vt/test_rocm_backend.cpp) | Runtime cases pass; the allocation-path case asserts the #2511 coupling on the board it runs on | diff --git a/docs/bench-evidence/rocm-residual-norm/README.md b/docs/bench-evidence/rocm-residual-norm/README.md new file mode 100644 index 0000000000..8fd4d5efdb --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/README.md @@ -0,0 +1,35 @@ +# Repair the blocked-stream fixture + +Row: `BACKEND-ROCM-RESIDUAL-NORM`. Issue: [#3103](https://github.com/mudler/vllm.cpp/issues/3103). +This repair starts from `94ac5d1742c540fc5cbc2084d0fb74fdfd5dab41` in a separate helper worktree. + +The operator's [red capture](operator-red.log) fails the existing nonblocking-stream assertion before testing kernel dispatch. +`RocmBackend::CreateQueue` uses `hipStreamCreate`, whose stream flags are zero. +The fixture now creates and owns a stream with `hipStreamNonBlocking`. +Its outer device scope outlives the stream, operands, and callback. +The release guard drains queued work before destroying its atomic callback flag, including assertion and exception exits. +The flag assertion and output comparison remain unchanged. +Production queue creation and residual normalization code remain unchanged. + +The [build receipt](build-receipt.json) records exact compiler and linker arguments, input hashes, and private output binaries. +The helper rebuilds the test and test main against private copies of the unchanged base archives. +The [build script](build-fixture.py) records that adaptation; this is a test relink, not a full product rebuild. +A scratch mutation replaces the real residual provider's launch stream with the default stream. +The operator must confirm that this mutation passes the stream-flag assertion and fails the output comparison. +No tracked product file or original archive changes during the mutation build. + +The [keyed-record proof](keyed-record-proof.json) starts from the complete target version of `.agents/backend-matrix.md`. +It changes only the `BACKEND-PLATFORM` citation from `backend.h:22` to `backend.h:23`. +Every other byte remains unchanged. +The [red record gate](record-red.log) reports the stale class citation; the [corrected gate](record-green.log) passes. + +The operator independently completes the pinned upstream export: 144 core cases, 60 IR RMS cases, and 60 IR add cases. +All 1,452 payload hashes are verified in the [export checks](upstream-export-checks.json). +The [native comparison](upstream-native.log) passes 528 CPU/ROCm executions and 22,576 assertions on the unchanged base implementation. +The original model token gate still differs at six positions in the L33/C2 tail. +This fixture repair makes no numerical or token-parity claim beyond those measured component results. + +The repaired CPU binary passes all six component tests and 8,985 assertions, including all 264 upstream cases. +The [CPU receipt](cpu-focused-receipt.json) and [raw output](cpu-focused.log) retain the exact command and result. +The [manifest](manifest.json) seals the retained captures and commands. +The repaired hardware fixture, wrong-stream mutation, full staged preflight, fresh review, and operator rerun remain pending at this checkpoint. diff --git a/docs/bench-evidence/rocm-residual-norm/build-fixture.py b/docs/bench-evidence/rocm-residual-norm/build-fixture.py new file mode 100644 index 0000000000..8340b76332 --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/build-fixture.py @@ -0,0 +1,48 @@ +from pathlib import Path +import concurrent.futures,hashlib,json,os,shlex,shutil,subprocess,time +work=Path('/home/vikash/vllm.cpp-residual-norm-repair1');original=Path('/home/vikash/vllm.cpp-residual-norm-impl');root=Path('/home/vikash/.cache/residual-norm-repair1');frozen=Path('/home/vikash/.cache/residual-norm-impl/green-freeze-2') +inputs={};commands=[] +def seal(p): + d=p.read_bytes();return {'bytes':len(d),'sha256':hashlib.sha256(d).hexdigest()} +def run(name,argv,cwd): + p=root/(name+'.log');start=time.monotonic() + with p.open('w') as out:r=subprocess.run(argv,cwd=cwd,stdout=out,stderr=subprocess.STDOUT) + entry={'name':name,'argv':argv,'cwd':str(cwd),'exit_code':r.returncode,'seconds':time.monotonic()-start,'log':str(p),'log_sha256':seal(p)['sha256']};commands.append(entry) + if r.returncode:raise RuntimeError(json.dumps(entry)) + return entry +# Rebuild only the changed test and its own test main. Production archives are +# private copies of the unchanged, already-built base and are sealed inputs. +compile_jobs=[];plans={} +for flavor in ('cpu','hip'): + oldbuild=original/('build-residual-'+flavor);build=work/('build-repair-'+flavor);build.mkdir(exist_ok=True);(build/'tests').mkdir(exist_ok=True) + shutil.copytree(oldbuild/'include',build/'include',dirs_exist_ok=True) + libsource=frozen/'build-residual-hip/libvllm.a' if flavor=='hip' else oldbuild/'libvllm.a' + for source,target in ((libsource,build/'libvllm.a'),(oldbuild/'libblake3_vendored.a',build/'libblake3_vendored.a')): + inputs[str(source)]=seal(source);shutil.copyfile(source,target);assert seal(target)==inputs[str(source)];inputs[str(target)]=seal(target) + db=json.loads((oldbuild/'compile_commands.json').read_text()) + for suffix in ('/tests/vt/test_ops_residual_rmsnorm.cpp','/tests/doctest_main.cpp'): + entry=next(x for x in db if x['file'].endswith(suffix));argv=shlex.split(entry['command']);argv=[a.replace(str(original),str(work)).replace(str(work/('build-residual-'+flavor)),str(build)) for a in argv] + obj=build/('test.o' if suffix.endswith('test_ops_residual_rmsnorm.cpp') else 'main.o');argv[argv.index('-o')+1]=str(obj) + compile_jobs.append((flavor+'-'+obj.stem,argv,build)) + plans[flavor]=build +with concurrent.futures.ThreadPoolExecutor(max_workers=4) as pool: + for f in [pool.submit(run,*j) for j in compile_jobs]:f.result() +for flavor,build in plans.items(): + run(flavor+'-main-archive',['ar','rcs',str(build/'tests/libvllm_test_main.a'),str(build/'main.o')],build) + if flavor=='hip': + argv=['/opt/rocm/lib/llvm/bin/clang++','--rocm-path=/opt/rocm','-O3','-DNDEBUG','--offload-arch=gfx1100','-Xlinker','--whole-archive','-Xlinker',str(build/'libvllm.a'),'-Xlinker','--no-whole-archive',str(build/'test.o'),'-o',str(build/'tests/test_ops_residual_rmsnorm'),str(build/'libvllm.a'),str(build/'tests/libvllm_test_main.a'),str(build/'libblake3_vendored.a'),'/opt/rocm/lib/libamdhip64.so','/opt/rocm/lib/libhipblas.so','/opt/rocm/lib/libhipblaslt.so','/opt/rocm/lib/libamdhip64.so.7.15.26333-0000000','-lgcc','-lgcc','-lgcc','-lgcc'] + else: + argv=['/usr/bin/c++','-O3','-DNDEBUG','-Wl,--whole-archive,'+str(build/'libvllm.a')+',--no-whole-archive',str(build/'test.o'),'-o',str(build/'tests/test_ops_residual_rmsnorm'),str(build/'libvllm.a'),str(build/'tests/libvllm_test_main.a'),str(build/'libblake3_vendored.a')] + run(flavor+'-link',argv,build) + if flavor=='hip':hip_link=argv +# This scratch-only mutation changes the real provider launch to the default +# stream. No tracked product file or frozen base archive is modified. +mutation=root/'wrong-stream';mutation.mkdir(exist_ok=True);source=work/'src/vt/rocm/rocm_residual_rmsnorm.hip';text=source.read_text();before='<<(queue.handle)>>>';assert text.count(before)==1;mutated=mutation/source.name;mutated.write_text(text.replace(before,'<<>>'));inputs[str(source)]=seal(source) +entry=next(x for x in db if x['file'].endswith('/src/vt/rocm/rocm_residual_rmsnorm.hip'));argv=shlex.split(entry['command']);argv=[a.replace(str(original),str(work)).replace(str(work/'build-residual-hip'),str(plans['hip'])) for a in argv];argv[argv.index('-o')+1]=str(mutation/'rocm_residual_rmsnorm.hip.o');argv[-1]=str(mutated);run('wrong-stream-compile',argv,mutation) +archive=mutation/'libvllm.a';shutil.copyfile(plans['hip']/'libvllm.a',archive);members=subprocess.check_output(['ar','t',str(archive)],text=True).splitlines();assert members.count('rocm_residual_rmsnorm.hip.o')==1 +run('wrong-stream-archive',['ar','r',str(archive),str(mutation/'rocm_residual_rmsnorm.hip.o')],mutation) +argv=[str(archive) if x==str(plans['hip']/'libvllm.a') else x for x in hip_link];argv[argv.index('-o')+1]=str(mutation/'test_ops_residual_rmsnorm');run('wrong-stream-link',argv,mutation) +for p,expected in inputs.items():assert seal(Path(p))==expected +outputs={str(p):seal(p) for p in [plans['cpu']/'tests/test_ops_residual_rmsnorm',plans['hip']/'tests/test_ops_residual_rmsnorm',mutation/'test_ops_residual_rmsnorm',mutated]} +report={'base':'94ac5d1742c540fc5cbc2084d0fb74fdfd5dab41','scope':'Only test translation units rebuilt; immutable production archives copied into private build directories. Scratch mutation replaces one real HIP provider object.','inputs':inputs,'outputs':outputs,'commands':commands,'all_inputs_unchanged':True,'max_parallel_compilers':4} +(root/'build-receipt.json').write_text(json.dumps(report,indent=2)+'\n');print(json.dumps({'status':'PASS','outputs':outputs},indent=2)) diff --git a/docs/bench-evidence/rocm-residual-norm/build-receipt.json b/docs/bench-evidence/rocm-residual-norm/build-receipt.json new file mode 100644 index 0000000000..fbbf61be52 --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/build-receipt.json @@ -0,0 +1,336 @@ +{ + "base": "94ac5d1742c540fc5cbc2084d0fb74fdfd5dab41", + "scope": "Only test translation units rebuilt; immutable production archives copied into private build directories. Scratch mutation replaces one real HIP provider object.", + "inputs": { + "/home/vikash/vllm.cpp-residual-norm-impl/build-residual-cpu/libvllm.a": { + "bytes": 85351758, + "sha256": "91c8a0132ada8fc39137d402c009573e7f2e0c39c147f8a7b00162e1f8d6f822" + }, + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/libvllm.a": { + "bytes": 85351758, + "sha256": "91c8a0132ada8fc39137d402c009573e7f2e0c39c147f8a7b00162e1f8d6f822" + }, + "/home/vikash/vllm.cpp-residual-norm-impl/build-residual-cpu/libblake3_vendored.a": { + "bytes": 34056, + "sha256": "2a639b282d02d8bb3b72a3fdb94b8e99ad2b09955e79577c8c42cfda18257b17" + }, + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/libblake3_vendored.a": { + "bytes": 34056, + "sha256": "2a639b282d02d8bb3b72a3fdb94b8e99ad2b09955e79577c8c42cfda18257b17" + }, + "/home/vikash/.cache/residual-norm-impl/green-freeze-2/build-residual-hip/libvllm.a": { + "bytes": 88483988, + "sha256": "ee9608029966c7aecb714f786e6c9cb9cafe3d64f27599f1bc638df1dd17e6ea" + }, + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/libvllm.a": { + "bytes": 88483988, + "sha256": "ee9608029966c7aecb714f786e6c9cb9cafe3d64f27599f1bc638df1dd17e6ea" + }, + "/home/vikash/vllm.cpp-residual-norm-impl/build-residual-hip/libblake3_vendored.a": { + "bytes": 34056, + "sha256": "2a639b282d02d8bb3b72a3fdb94b8e99ad2b09955e79577c8c42cfda18257b17" + }, + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/libblake3_vendored.a": { + "bytes": 34056, + "sha256": "2a639b282d02d8bb3b72a3fdb94b8e99ad2b09955e79577c8c42cfda18257b17" + }, + "/home/vikash/vllm.cpp-residual-norm-repair1/src/vt/rocm/rocm_residual_rmsnorm.hip": { + "bytes": 5420, + "sha256": "f2ee1eeeab08bc9a8102056c0ae7e0ddf5eed1fe23b831d5938d7e99be7eb657" + } + }, + "outputs": { + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/tests/test_ops_residual_rmsnorm": { + "bytes": 26209432, + "sha256": "6470de98c0e5f56cde1e4691d21f8cbe8f1c67c3975d6793fc89e01eb02ec828" + }, + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/tests/test_ops_residual_rmsnorm": { + "bytes": 29318040, + "sha256": "04f0ac97f9aad3f3d683ab6c692b6512069348dcc4b04c8fa4824a1befcf1a06" + }, + "/home/vikash/.cache/residual-norm-repair1/wrong-stream/test_ops_residual_rmsnorm": { + "bytes": 29318040, + "sha256": "f25a4251a2c60a1c8f3fcdd6346661f0694587a3f1467caf9eadbbd20a664235" + }, + "/home/vikash/.cache/residual-norm-repair1/wrong-stream/rocm_residual_rmsnorm.hip": { + "bytes": 5389, + "sha256": "41aeb8790ba3d6a4e24921f5324174dcdf5fe4b4647b784f71a4dc5c4f94873a" + } + }, + "commands": [ + { + "name": "hip-main", + "argv": [ + "/usr/bin/c++", + "-isystem", + "/home/vikash/vllm.cpp-residual-norm-repair1/third_party", + "-O3", + "-DNDEBUG", + "-std=c++20", + "-fPIC", + "-ffp-contract=off", + "-o", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/main.o", + "-c", + "/home/vikash/vllm.cpp-residual-norm-repair1/tests/doctest_main.cpp" + ], + "cwd": "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip", + "exit_code": 0, + "seconds": 3.5146089729387313, + "log": "/home/vikash/.cache/residual-norm-repair1/hip-main.log", + "log_sha256": "e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855" + }, + { + "name": "cpu-main", + "argv": [ + "/usr/bin/c++", + "-isystem", + "/home/vikash/vllm.cpp-residual-norm-repair1/third_party", + "-O3", + "-DNDEBUG", + "-std=c++20", + "-fPIC", + "-ffp-contract=off", + "-o", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/main.o", + "-c", + "/home/vikash/vllm.cpp-residual-norm-repair1/tests/doctest_main.cpp" + ], + "cwd": "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu", + "exit_code": 0, + "seconds": 3.5353177699726075, + "log": "/home/vikash/.cache/residual-norm-repair1/cpu-main.log", + "log_sha256": "e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855" + }, + { + "name": "cpu-test", + "argv": [ + "/usr/bin/c++", + "-I/home/vikash/vllm.cpp-residual-norm-repair1/tests", + "-I/home/vikash/vllm.cpp-residual-norm-repair1/include", + "-I/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/include", + "-I/home/vikash/vllm.cpp-residual-norm-repair1/third_party/blake3", + "-isystem", + "/home/vikash/vllm.cpp-residual-norm-repair1/third_party", + "-O3", + "-DNDEBUG", + "-std=c++20", + "-fPIE", + "-ffp-contract=off", + "-Wall", + "-Wextra", + "-Werror", + "-o", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/test.o", + "-c", + "/home/vikash/vllm.cpp-residual-norm-repair1/tests/vt/test_ops_residual_rmsnorm.cpp" + ], + "cwd": "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu", + "exit_code": 0, + "seconds": 6.466870703035966, + "log": "/home/vikash/.cache/residual-norm-repair1/cpu-test.log", + "log_sha256": "e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855" + }, + { + "name": "hip-test", + "argv": [ + "/usr/bin/c++", + "-DVLLM_CPP_HIP", + "-D__HIP_PLATFORM_AMD__=1", + "-I/home/vikash/vllm.cpp-residual-norm-repair1/tests", + "-I/home/vikash/vllm.cpp-residual-norm-repair1/include", + "-I/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/include", + "-I/opt/rocm/include", + "-I/home/vikash/vllm.cpp-residual-norm-repair1/third_party/blake3", + "-isystem", + "/home/vikash/vllm.cpp-residual-norm-repair1/third_party", + "-O3", + "-DNDEBUG", + "-std=c++20", + "-fPIE", + "-ffp-contract=off", + "-Wall", + "-Wextra", + "-Werror", + "-o", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/test.o", + "-c", + "/home/vikash/vllm.cpp-residual-norm-repair1/tests/vt/test_ops_residual_rmsnorm.cpp" + ], + "cwd": "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip", + "exit_code": 0, + "seconds": 7.25080132111907, + "log": "/home/vikash/.cache/residual-norm-repair1/hip-test.log", + "log_sha256": "e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855" + }, + { + "name": "cpu-main-archive", + "argv": [ + "ar", + "rcs", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/tests/libvllm_test_main.a", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/main.o" + ], + "cwd": "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu", + "exit_code": 0, + "seconds": 0.04720654105767608, + "log": "/home/vikash/.cache/residual-norm-repair1/cpu-main-archive.log", + "log_sha256": "e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855" + }, + { + "name": "cpu-link", + "argv": [ + "/usr/bin/c++", + "-O3", + "-DNDEBUG", + "-Wl,--whole-archive,/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/libvllm.a,--no-whole-archive", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/test.o", + "-o", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/tests/test_ops_residual_rmsnorm", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/libvllm.a", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/tests/libvllm_test_main.a", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/libblake3_vendored.a" + ], + "cwd": "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu", + "exit_code": 0, + "seconds": 0.5794992081355304, + "log": "/home/vikash/.cache/residual-norm-repair1/cpu-link.log", + "log_sha256": "e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855" + }, + { + "name": "hip-main-archive", + "argv": [ + "ar", + "rcs", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/tests/libvllm_test_main.a", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/main.o" + ], + "cwd": "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip", + "exit_code": 0, + "seconds": 0.046468702144920826, + "log": "/home/vikash/.cache/residual-norm-repair1/hip-main-archive.log", + "log_sha256": "e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855" + }, + { + "name": "hip-link", + "argv": [ + "/opt/rocm/lib/llvm/bin/clang++", + "--rocm-path=/opt/rocm", + "-O3", + "-DNDEBUG", + "--offload-arch=gfx1100", + "-Xlinker", + "--whole-archive", + "-Xlinker", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/libvllm.a", + "-Xlinker", + "--no-whole-archive", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/test.o", + "-o", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/tests/test_ops_residual_rmsnorm", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/libvllm.a", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/tests/libvllm_test_main.a", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/libblake3_vendored.a", + "/opt/rocm/lib/libamdhip64.so", + "/opt/rocm/lib/libhipblas.so", + "/opt/rocm/lib/libhipblaslt.so", + "/opt/rocm/lib/libamdhip64.so.7.15.26333-0000000", + "-lgcc", + "-lgcc", + "-lgcc", + "-lgcc" + ], + "cwd": "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip", + "exit_code": 0, + "seconds": 0.13446427206508815, + "log": "/home/vikash/.cache/residual-norm-repair1/hip-link.log", + "log_sha256": "e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855" + }, + { + "name": "wrong-stream-compile", + "argv": [ + "/opt/rocm/lib/llvm/bin/clang++", + "-DVLLM_CPP_HIP", + "-D__HIP_ROCclr__=1", + "-I/home/vikash/vllm.cpp-residual-norm-repair1/include", + "-I/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/include", + "-I/home/vikash/vllm.cpp-residual-norm-repair1/src", + "-I/home/vikash/vllm.cpp-residual-norm-repair1/third_party/blake3", + "-isystem", + "/home/vikash/vllm.cpp-residual-norm-repair1/third_party", + "--rocm-path=/opt/rocm", + "-O3", + "-DNDEBUG", + "-std=c++20", + "--offload-arch=gfx1100", + "-fPIC", + "-ffp-contract=off", + "-Wall", + "-Wextra", + "-Werror", + "-o", + "/home/vikash/.cache/residual-norm-repair1/wrong-stream/rocm_residual_rmsnorm.hip.o", + "-x", + "hip", + "-c", + "/home/vikash/.cache/residual-norm-repair1/wrong-stream/rocm_residual_rmsnorm.hip" + ], + "cwd": "/home/vikash/.cache/residual-norm-repair1/wrong-stream", + "exit_code": 0, + "seconds": 2.2970829010009766, + "log": "/home/vikash/.cache/residual-norm-repair1/wrong-stream-compile.log", + "log_sha256": "e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855" + }, + { + "name": "wrong-stream-archive", + "argv": [ + "ar", + "r", + "/home/vikash/.cache/residual-norm-repair1/wrong-stream/libvllm.a", + "/home/vikash/.cache/residual-norm-repair1/wrong-stream/rocm_residual_rmsnorm.hip.o" + ], + "cwd": "/home/vikash/.cache/residual-norm-repair1/wrong-stream", + "exit_code": 0, + "seconds": 0.4022414069622755, + "log": "/home/vikash/.cache/residual-norm-repair1/wrong-stream-archive.log", + "log_sha256": "e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855" + }, + { + "name": "wrong-stream-link", + "argv": [ + "/opt/rocm/lib/llvm/bin/clang++", + "--rocm-path=/opt/rocm", + "-O3", + "-DNDEBUG", + "--offload-arch=gfx1100", + "-Xlinker", + "--whole-archive", + "-Xlinker", + "/home/vikash/.cache/residual-norm-repair1/wrong-stream/libvllm.a", + "-Xlinker", + "--no-whole-archive", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/test.o", + "-o", + "/home/vikash/.cache/residual-norm-repair1/wrong-stream/test_ops_residual_rmsnorm", + "/home/vikash/.cache/residual-norm-repair1/wrong-stream/libvllm.a", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/tests/libvllm_test_main.a", + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-hip/libblake3_vendored.a", + "/opt/rocm/lib/libamdhip64.so", + "/opt/rocm/lib/libhipblas.so", + "/opt/rocm/lib/libhipblaslt.so", + "/opt/rocm/lib/libamdhip64.so.7.15.26333-0000000", + "-lgcc", + "-lgcc", + "-lgcc", + "-lgcc" + ], + "cwd": "/home/vikash/.cache/residual-norm-repair1/wrong-stream", + "exit_code": 0, + "seconds": 0.12359828595072031, + "log": "/home/vikash/.cache/residual-norm-repair1/wrong-stream-link.log", + "log_sha256": "e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855" + } + ], + "all_inputs_unchanged": true, + "max_parallel_compilers": 4 +} diff --git a/docs/bench-evidence/rocm-residual-norm/cpu-focused-receipt.json b/docs/bench-evidence/rocm-residual-norm/cpu-focused-receipt.json new file mode 100644 index 0000000000..544188e5a9 --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/cpu-focused-receipt.json @@ -0,0 +1,15 @@ +{ + "argv": [ + "/home/vikash/vllm.cpp-residual-norm-repair1/build-repair-cpu/tests/test_ops_residual_rmsnorm" + ], + "env": { + "VT_RESIDUAL_NORM_UPSTREAM": "/home/vikash/.cache/residual-norm-impl/upstream-freeze/results", + "OMP_NUM_THREADS": "4", + "TMPDIR": "/home/vikash/.cache/residual-norm-repair1/tmp" + }, + "exit_code": 0, + "seconds": 45.94254233408719, + "log_sha256": "8e281532f0750b8bc3f214f640d3d3db76ef1cdc72de766c702b801339444d68", + "binary_sha256": "6470de98c0e5f56cde1e4691d21f8cbe8f1c67c3975d6793fc89e01eb02ec828", + "gpu_work": false +} diff --git a/docs/bench-evidence/rocm-residual-norm/cpu-focused.log b/docs/bench-evidence/rocm-residual-norm/cpu-focused.log new file mode 100644 index 0000000000..9e716c7b04 --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/cpu-focused.log @@ -0,0 +1,12 @@ +[doctest] doctest version is "2.5.2" +[doctest] run with "--help" for options +=============================================================================== +/home/vikash/vllm.cpp-residual-norm-repair1/tests/vt/test_ops_residual_rmsnorm.cpp:179: +TEST CASE: residual RMS norm matches the complete pinned upstream BF16 fixture grid + +/home/vikash/vllm.cpp-residual-norm-repair1/tests/vt/test_ops_residual_rmsnorm.cpp:255: MESSAGE: Executed pinned upstream normalization cases across native/reference devices: 264 + +=============================================================================== +[doctest] test cases: 6 | 6 passed | 0 failed | 0 skipped +[doctest] assertions: 8985 | 8985 passed | 0 failed | +[doctest] Status: SUCCESS! diff --git a/docs/bench-evidence/rocm-residual-norm/keyed-record-proof.json b/docs/bench-evidence/rocm-residual-norm/keyed-record-proof.json new file mode 100644 index 0000000000..365500f091 --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/keyed-record-proof.json @@ -0,0 +1,11 @@ +{ + "target_base": "94ac5d1742c540fc5cbc2084d0fb74fdfd5dab41", + "target_complete_file_sha256": "8861de845d02df23e597a4c4ceaf5b4f053c5b5846e74bb836f95f9f2e7d04e0", + "result_sha256": "3d41b47f4aae184343d9f47389f3da91af9846b390ebd9a5b0913758c1f075ba", + "changed_key": "BACKEND-PLATFORM", + "old_citation": "[backend.h:22](../include/vt/backend.h#L22)", + "new_citation": "[backend.h:23](../include/vt/backend.h#L23)", + "all_other_bytes_unchanged": true, + "other_bytes_sha256": "802a366de63b9349333ed0bbfc86a3f10e86a9c6c62723fe7a301eed2142add0", + "change": "One citation label and line fragment corrected after residual_norm.h include moved Backend definition to line23" +} diff --git a/docs/bench-evidence/rocm-residual-norm/manifest.json b/docs/bench-evidence/rocm-residual-norm/manifest.json new file mode 100644 index 0000000000..9631cd10c8 --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/manifest.json @@ -0,0 +1,81 @@ +{ + "files": { + "operator-red.log": { + "source": "/home/vikash/.cache/residual-norm-impl/green-freeze-2/test_ops_residual_rmsnorm-fusion-1-command-operator.log", + "bytes": 1849, + "sha256": "cab08e83302b93dcc9c9a0c76d8a5820e015cf67f0ad803575d3337914b4bc48" + }, + "operator-red-command.json": { + "source": "/home/vikash/.cache/residual-norm-impl/green-freeze-2/test_ops_residual_rmsnorm-fusion-1-command.json", + "bytes": 541, + "sha256": "4b43060cc688ac367c69cd1da0f43dab2cfa91219563b9897cbeff3d907ac435" + }, + "operator-red-receipts.json": { + "source": "/home/vikash/.cache/residual-norm-impl/green-freeze-2/operator-receipts.json", + "bytes": 1744, + "sha256": "be1011b01737eded6f7d1d6b7053111038c5401296ec677905884a3e19517618" + }, + "upstream-export-receipt.json": { + "source": "/home/vikash/.cache/residual-norm-impl/upstream-freeze/operator-receipt.json", + "bytes": 272, + "sha256": "87de74edc117f00d47b46ca752d8ab2192363921bd2d019d90aa83153f8e90fb" + }, + "upstream-export-checks.json": { + "source": "/home/vikash/.cache/residual-norm-impl/upstream-freeze/operator-output-checks.json", + "bytes": 276, + "sha256": "b1d711efd2e12f62d1507fd0e9b0271d27f6c22d5fd15cd13032d23a5aba1faf" + }, + "upstream-native-receipt.json": { + "source": "/home/vikash/.cache/residual-norm-impl/upstream-native-v1/operator-receipt.json", + "bytes": 279, + "sha256": "4f5ce52487a123e00366c710bb070a7f573823f420a0aa6c3e8a2588dbee714d" + }, + "upstream-native.log": { + "source": "/home/vikash/.cache/residual-norm-impl/upstream-native-v1/operator-upstream-native.log", + "bytes": 734, + "sha256": "d7385cb03119a5c7105a0f5e18a3abf5e42fea6e557d0e405ec8fde7d9d48eb7" + }, + "record-red.log": { + "source": "/home/vikash/.cache/residual-norm-repair1/record-red.log", + "bytes": 3816, + "sha256": "f3836103cbabd70a0bcaa1c0f10325626cc4f6ea98df1b37c1bb419e8fe51a4c" + }, + "record-green.log": { + "source": "/home/vikash/.cache/residual-norm-repair1/record-green.log", + "bytes": 82, + "sha256": "35b858d9cf328a42f8574ae545a929a3d29842d411228ff4fdc0fa63b74ee3c7" + }, + "keyed-record-proof.json": { + "source": "/home/vikash/.cache/residual-norm-repair1/keyed-record-proof.json", + "bytes": 675, + "sha256": "7412c2472505c7706a5396e8dd218769de8ce5f4c0d2514725fb81e5a91d07a2" + }, + "build-receipt.json": { + "source": "/home/vikash/.cache/residual-norm-repair1/build-receipt.json", + "bytes": 13700, + "sha256": "b677da932aa490238d362fb7b0f677bb96f9dfb1f9d36eee7765c3441622d66c" + }, + "build-fixture.py": { + "source": "/home/vikash/.cache/residual-norm-repair1/build-fixture.py", + "bytes": 5482, + "sha256": "1405c1d937b954fb85d075c76970b2d84bd4490d1f8b26d53074a49eb6f94b19" + }, + "cpu-focused.log": { + "source": "/home/vikash/.cache/residual-norm-repair1/cpu-focused.log", + "bytes": 736, + "sha256": "8e281532f0750b8bc3f214f640d3d3db76ef1cdc72de766c702b801339444d68" + }, + "cpu-focused-receipt.json": { + "source": "/home/vikash/.cache/residual-norm-repair1/cpu-focused-receipt.json", + "bytes": 566, + "sha256": "475559bd1baad3907c39f4a90a413b456dca3ea63264daa85d3e5124b6a61e43" + }, + "product-unchanged.json": { + "source": "/home/vikash/.cache/residual-norm-repair1/product-unchanged.json", + "bytes": 194, + "sha256": "9440d0e947460dc68dbda3c9164f296c90984b52a6c709f11236c4a5e9c7c007" + } + }, + "product_source_unchanged": true, + "gpu_repair_and_mutation": "PENDING operator execution; frozen commands under /home/vikash/.cache/residual-norm-repair1/gpu-freeze-v1" +} diff --git a/docs/bench-evidence/rocm-residual-norm/operator-red-command.json b/docs/bench-evidence/rocm-residual-norm/operator-red-command.json new file mode 100644 index 0000000000..4c792e4b52 --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/operator-red-command.json @@ -0,0 +1,13 @@ +{ + "argv": [ + "/home/vikash/.cache/residual-norm-impl/green-freeze-2/build-residual-hip/tests/test_ops_residual_rmsnorm" + ], + "environment": { + "VT_ROCM_MOE_FIXTURE": "/home/vikash/vllm.cpp-rdna3-moe-impl/build-rdna3-moe-hip/fixture", + "HIP_VISIBLE_DEVICES": "0,1", + "ROCR_VISIBLE_DEVICES": "0,1", + "TMPDIR": "/home/vikash/.cache/residual-norm-impl/tmp", + "VT_ROCM_MOE_ORACLE": "/home/vikash/vllm.cpp-rdna3-moe-impl/build-rdna3-moe-hip/evidence/oracle-selection-6/production.json", + "VT_FUSED_CHAIN_ADOPT": "1" + } +} diff --git a/docs/bench-evidence/rocm-residual-norm/operator-red-receipts.json b/docs/bench-evidence/rocm-residual-norm/operator-red-receipts.json new file mode 100644 index 0000000000..75cebd928b --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/operator-red-receipts.json @@ -0,0 +1,51 @@ +[ + { + "name": "test_rocm_moe_bf16-fusion-1-command.json", + "exit_code": 1, + "seconds": 2.1759774889796972, + "log_sha256": "72df6480041ed42dca4752095ab8edf29c0d17beda32a0a030c665eed69be4b6", + "inputs_verified_before_after": 75 + }, + { + "name": "test_rocm_moe_bf16-fusion-0-command.json", + "exit_code": 1, + "seconds": 2.1762429990340024, + "log_sha256": "72df6480041ed42dca4752095ab8edf29c0d17beda32a0a030c665eed69be4b6", + "inputs_verified_before_after": 75 + }, + { + "name": "test_ops_residual_rmsnorm-fusion-1-command.json", + "exit_code": 1, + "seconds": 2.878145433962345, + "log_sha256": "cab08e83302b93dcc9c9a0c76d8a5820e015cf67f0ad803575d3337914b4bc48", + "inputs_verified_before_after": 75 + }, + { + "name": "test_ops_rmsnorm-fusion-1-command.json", + "exit_code": 0, + "seconds": 1.2802332111168653, + "log_sha256": "71a666e5d823defafcefc982082ff6fe311685116b1976853ffcbfb88039f4d0", + "inputs_verified_before_after": 75 + }, + { + "name": "test_ops_rmsnorm_weight_dtype-fusion-1-command.json", + "exit_code": 0, + "seconds": 1.2827785040717572, + "log_sha256": "bea994619905fbeda785d1d5d234a900b9ca667f2b2ae58bd7773a9fb501a77a", + "inputs_verified_before_after": 75 + }, + { + "name": "test_ops_fused_chain-fusion-1-command.json", + "exit_code": 0, + "seconds": 1.2671924571041018, + "log_sha256": "13b2f01c85fb3a2cc18c028c0d6b16ead12543c5899231336b0e1344739037a6", + "inputs_verified_before_after": 75 + }, + { + "name": "test_fused_chain_additivity-fusion-1-command.json", + "exit_code": 0, + "seconds": 1.2803919969592243, + "log_sha256": "a7ecb03fe56dd2067d78b71b57eba256e71a716f2db8cbefd8e0acbda564a10f", + "inputs_verified_before_after": 75 + } +] diff --git a/docs/bench-evidence/rocm-residual-norm/operator-red.log b/docs/bench-evidence/rocm-residual-norm/operator-red.log new file mode 100644 index 0000000000..50433179e8 --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/operator-red.log @@ -0,0 +1,32 @@ + + +============================ ROCm System Management Interface ============================ +=================================== % time GPU is busy =================================== +GPU[0] : GPU use (%): 0 +GPU[1] : GPU use (%): 0 +GPU[2] : GPU use (%): 0 +GPU[3] : GPU use (%): 0 +========================================================================================== +===================================== KFD Processes ====================================== +No KFD PIDs currently running +========================================================================================== +================================== End of ROCm SMI Log =================================== +[doctest] doctest version is "2.5.2" +[doctest] run with "--help" for options +=============================================================================== +/home/vikash/vllm.cpp-residual-norm-impl/tests/vt/test_ops_residual_rmsnorm.cpp:179: +TEST CASE: residual RMS norm matches the complete pinned upstream BF16 fixture grid + +/home/vikash/vllm.cpp-residual-norm-impl/tests/vt/test_ops_residual_rmsnorm.cpp:183: MESSAGE: Upstream fixture gate not executed: set VT_RESIDUAL_NORM_UPSTREAM to the complete pinned export + +=============================================================================== +/home/vikash/vllm.cpp-residual-norm-impl/tests/vt/test_ops_residual_rmsnorm.cpp:356: +TEST CASE: ROCm residual RMS norm cannot run ahead of a blocked nondefault stream + +/home/vikash/vllm.cpp-residual-norm-impl/tests/vt/test_ops_residual_rmsnorm.cpp:363: FATAL ERROR: REQUIRE( (flags & 0x01) != 0 ) is NOT correct! + values: REQUIRE( 0 != 0 ) + +=============================================================================== +[doctest] test cases: 8 | 7 passed | 1 failed | 0 skipped +[doctest] assertions: 11200 | 11199 passed | 1 failed | +[doctest] Status: FAILURE! diff --git a/docs/bench-evidence/rocm-residual-norm/product-unchanged.json b/docs/bench-evidence/rocm-residual-norm/product-unchanged.json new file mode 100644 index 0000000000..d2907d62e4 --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/product-unchanged.json @@ -0,0 +1,11 @@ +{ + "base": "94ac5d1742c540fc5cbc2084d0fb74fdfd5dab41", + "paths": [ + "src", + "include", + "CMakeLists.txt", + "tests/CMakeLists.txt" + ], + "all_unchanged": true, + "diff_bytes": 0 +} diff --git a/docs/bench-evidence/rocm-residual-norm/record-green.log b/docs/bench-evidence/rocm-residual-norm/record-green.log new file mode 100644 index 0000000000..387dbbe4df --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/record-green.log @@ -0,0 +1 @@ +agent record OK: ENGINE=178 MODEL=379 QUANT=87 KERNEL=59 BACKEND=90 ANCHOR-ROT=33 diff --git a/docs/bench-evidence/rocm-residual-norm/record-red.log b/docs/bench-evidence/rocm-residual-norm/record-red.log new file mode 100644 index 0000000000..f9de7fe608 --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/record-red.log @@ -0,0 +1,39 @@ +ERROR: RECORD ANCHOR REGRESSION in bucket 'stale': 29 > baseline 28. A citation names a line that no longer holds what the prose says it does. Run `python3 scripts/check-agent-record.py --report` for the offenders and repair the anchor. NEVER raise the baseline to pass. +Record anchors in the `code` / `tests` cells of ACTIVE/ANCHOR-BACKFILL/BUILD-ONLY/DONE/GATING/PARTIAL/READY/UNTRACED rows: + + BROKEN KV-PREFIX-CACHE -> examples/server/main.cpp:126 + STALE KV-PREFIX-MATCH-UNIT -> src/vllm/v1/core/kv_cache_utils.cpp:638 expected `resolve_kv_cache_block_sizes` + STALE KV-PREFIX-MATCH-UNIT -> src/vllm/v1/core/kv_cache_utils.cpp:577 expected `get_request_block_hasher` + STALE ENG-ASYNC-SCHED -> src/vllm/v1/engine/core.cpp:91 expected `step_with_batch_queue` + STALE ENG-SGLANG-BEHAVIOR-FLAG -> src/vllm/v1/core/sched/scheduler.cpp:183-243 expected `maybe_reorder_waiting_for_lpm()` + STALE KV-OFFLOAD -> src/vllm/v1/core/block_pool.cpp:139-155 expected `BlockPool::evict_blocks` + STALE KV-EVENTS -> include/vllm/sampling_params.h:251 expected `SamplingParams::extra_args` + BROKEN KV-SGLANG-RADIX-CACHE -> examples/server/main.cpp:185 + STALE SAMPLE-LOGPROB-TOKEN-IDS -> src/vllm/v1/worker/gpu/input_batch.cpp:306 expected `add_request` + STALE SAMPLE-LOGPROB-TOKEN-IDS -> src/vllm/v1/sample/sampler.cpp:160 expected `GatherSpecificTokenLogprobs` + STALE SERVE-METRICS -> src/vllm/v1/engine/llm_engine.cpp:99 expected `Record` + STALE SERVE-METRICS -> src/vllm/v1/engine/async_llm.cpp:262 expected `IterationStats` + STALE SERVE-METRICS -> src/vllm/entrypoints/openai/api_server.cpp:251 expected `handle_metrics` + BROKEN SERVE-STREAM-USAGE -> examples/server/main.cpp:123 + STALE SERVE-UTILITY-ENDPOINTS -> src/vllm/entrypoints/openai/api_server.cpp:262 expected `handle_tokenize` + STALE SERVE-UTILITY-ENDPOINTS -> include/vllm/entrypoints/openai/api_server.h:118 expected `set_tokenizer_info_enabled` + STALE ENG-TRAILER-MERGE-ARTIFACTS -> scripts/check-commit-trailers.py:60 expected `join_trailing_trailer_paragraphs` + STALE ENG-FORGE-COAUTHOR -> scripts/check-commit-trailers.py:38 expected `FORGE_ACCOUNT_EMAIL` + STALE SERVE-RECIPE-ARGS -> src/vllm/entrypoints/openai/server_main.cpp:289 expected `kAcceptedInertArgs` + STALE SERVE-ADMIN -> src/vllm/entrypoints/openai/api_server.cpp:488 expected `handle_abort_requests` + STALE SERVE-ADMIN -> include/vllm/entrypoints/openai/api_server.h:156 expected `set_abort_requests` + BROKEN SERVE-VIDEOS-REFS -> examples/server/main.cpp:96 + STALE ENG-MOE-HOSTFREE -> src/vllm/model_executor/models/qwen3_5.cpp:4472 expected `Qwen3_5Model::PrepareMarlinResident` + STALE ENG-MOE-HOSTFREE -> tests/vllm/test_qwen36_weights.cpp:273 expected `ReleaseHost` + STALE ENG-MOE-LOADSTREAM -> include/vllm/model_executor/models/qwen3_5_weights.h:307 expected `Qwen3_5MoeWeights::load_layer_experts` + STALE ENG-MOE-LOADSTREAM -> src/vllm/entrypoints/model_loader.cpp:365 expected `LoadFromDir` + BROKEN MODEL-MM-qwen3-5-qwen3-5-moe-for-conditional-generation -> src/vllm/model_executor/models/registry.cpp:10-20 + STALE QUANT-GGUF-NVFP4 -> tests/vllm/test_gguf_nvfp4.cpp:144 expected `DequantNvfp4ToBf16` + STALE QUANT-CT-MXFP4 -> src/vllm/model_executor/model_loader/mxfp4_dequant.cpp:14 expected `E8M0ToF32` + STALE KERNEL-ATTN-DFLASH-BLOCK -> src/vt/cuda/cuda_ops.cu:1300 expected `DFlashBlockAttentionKernelCuda` + STALE KERNEL-ATTN-DFLASH-PAGED-BLOCK -> src/vt/cuda/cuda_ops.cu:1452 expected `DFlashPagedBlockAttentionKernelCuda` + STALE BACKEND-CUDA-SM120 -> cmake/CudaArchFeatures.cmake:216 expected `VT_CUDA_FEATURE_TABLE` + STALE BACKEND-PLATFORM -> include/vt/backend.h:22 expected `vt::Backend` + STALE BACKEND-DISTRIBUTED-COMM -> include/vt/device.h:50 expected `vt::Queue` + +record anchors: ok=936, stale=29, broken=5 -> rot 34 diff --git a/docs/bench-evidence/rocm-residual-norm/upstream-export-checks.json b/docs/bench-evidence/rocm-residual-norm/upstream-export-checks.json new file mode 100644 index 0000000000..6ec453c6ad --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/upstream-export-checks.json @@ -0,0 +1,12 @@ +{ + "complete": true, + "pin": "e126687a9a828d513c01a07cd69f025f27d63280", + "cases": 264, + "families": { + "core": 144, + "ir-rms": 60, + "ir-add": 60 + }, + "payload_files": 1452, + "cases_sha256": "fe7347766cac98b9f57664224bb696e3acb3d43619b168c66e0ad4a7bad9f960" +} diff --git a/docs/bench-evidence/rocm-residual-norm/upstream-export-receipt.json b/docs/bench-evidence/rocm-residual-norm/upstream-export-receipt.json new file mode 100644 index 0000000000..9dc92bce2a --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/upstream-export-receipt.json @@ -0,0 +1,12 @@ +{ + "results": [ + { + "name": "upstream", + "exit_code": 0, + "expected_exit": 0, + "seconds": 56.36936064995825, + "log_sha256": "f49d37cc9369d9bd83794885280cb053abc2b2694f68041886b9fa83428fa339" + } + ], + "inputs_verified_before_after": 4266 +} diff --git a/docs/bench-evidence/rocm-residual-norm/upstream-native-receipt.json b/docs/bench-evidence/rocm-residual-norm/upstream-native-receipt.json new file mode 100644 index 0000000000..85b863a59c --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/upstream-native-receipt.json @@ -0,0 +1,12 @@ +{ + "results": [ + { + "name": "upstream-native", + "exit_code": 0, + "expected_exit": 0, + "seconds": 77.10780900088139, + "log_sha256": "d7385cb03119a5c7105a0f5e18a3abf5e42fea6e557d0e405ec8fde7d9d48eb7" + } + ], + "inputs_verified_before_after": 4025 +} diff --git a/docs/bench-evidence/rocm-residual-norm/upstream-native.log b/docs/bench-evidence/rocm-residual-norm/upstream-native.log new file mode 100644 index 0000000000..0eeda2f325 --- /dev/null +++ b/docs/bench-evidence/rocm-residual-norm/upstream-native.log @@ -0,0 +1,12 @@ +[doctest] doctest version is "2.5.2" +[doctest] run with "--help" for options +=============================================================================== +/home/vikash/vllm.cpp-residual-norm-impl/tests/vt/test_ops_residual_rmsnorm.cpp:179: +TEST CASE: residual RMS norm matches the complete pinned upstream BF16 fixture grid + +/home/vikash/vllm.cpp-residual-norm-impl/tests/vt/test_ops_residual_rmsnorm.cpp:255: MESSAGE: Executed pinned upstream normalization cases across native/reference devices: 528 + +=============================================================================== +[doctest] test cases: 1 | 1 passed | 0 failed | 7 skipped +[doctest] assertions: 22576 | 22576 passed | 0 failed | +[doctest] Status: SUCCESS! diff --git a/include/vllm/model_executor/models/dense_attn_block.h b/include/vllm/model_executor/models/dense_attn_block.h index 0b5cfde0e8..d19210ae2b 100644 --- a/include/vllm/model_executor/models/dense_attn_block.h +++ b/include/vllm/model_executor/models/dense_attn_block.h @@ -634,22 +634,37 @@ inline DBuf AttnBlock(Dev d, const Qwen3DenseAttnWeights& w, const HfConfig& cfg // prefill AND decode. // Per-head q/k RMSNorm — Qwen3 only (SKIPPED when the model has no qk-norm, // e.g. Llama, which leaves w.q_norm/w.k_norm empty). RoPE runs either way. - if (has_qk_norm) { - Tensor wqn = attn_f32 ? ResidentWeightF32(d, w.q_norm, {Dh}) - : ResidentWeight(d, w.q_norm, {Dh}); - Tensor wkn = attn_f32 ? ResidentWeightF32(d, w.k_norm, {Dh}) - : ResidentWeight(d, w.k_norm, {Dh}); - vt::RmsNorm(d.q, q2, q2, wqn, vt::RmsNormArgs{eps, false}); - vt::RmsNorm(d.q, k2, k2, wkn, vt::RmsNormArgs{eps, false}); - } - if (RopeCacheEnabled() && rot > 0) { - Tensor k3v = k3; - vt::RopeFromCache(d.q, q3, &k3v, si.rope_row_idx.t(), si.cos_sin_bf16.t(), - MakeRopeArgs(cfg)); + // The hand-call takes the registered fused op when the device has one, so + // the ADOPT path (FusedChain -> the recipe's fast realisation) and this + // fallback agree on every backend, which is the recipe's byte-exact + // composite contract. Where no fast op is registered (CPU) both realizations + // stay the standalone three-op sequence below. + const bool fused_preamble_op = + has_qk_norm && rot > 0 && RopeCacheEnabled() && !attn_f32 && + vt::OpRegistered(vt::OpId::kAttnQkNormRope, d.q.device.type); + if (fused_preamble_op) { + Tensor wqn = ResidentWeight(d, w.q_norm, {Dh}); + Tensor wkn = ResidentWeight(d, w.k_norm, {Dh}); + vt::AttnQkNormRope(d.q, q3, k3, wqn, wkn, si.cos_sin_bf16.t(), si.rope_row_idx.t(), + vt::RmsNormArgs{eps, false}, MakeRopeArgs(cfg)); } else { - // DEFAULT (byte-identical, deterministic): in-place bf16 NeoX RoPE with - // per-element fp64 cos/sin, mirroring vLLM's rotary_emb bf16 rounding. - vt::RopeNeox(d.q, q3, k3, si.positions.t(), MakeRopeArgs(cfg)); + if (has_qk_norm) { + Tensor wqn = attn_f32 ? ResidentWeightF32(d, w.q_norm, {Dh}) + : ResidentWeight(d, w.q_norm, {Dh}); + Tensor wkn = attn_f32 ? ResidentWeightF32(d, w.k_norm, {Dh}) + : ResidentWeight(d, w.k_norm, {Dh}); + vt::RmsNorm(d.q, q2, q2, wqn, vt::RmsNormArgs{eps, false}); + vt::RmsNorm(d.q, k2, k2, wkn, vt::RmsNormArgs{eps, false}); + } + if (RopeCacheEnabled() && rot > 0) { + Tensor k3v = k3; + vt::RopeFromCache(d.q, q3, &k3v, si.rope_row_idx.t(), si.cos_sin_bf16.t(), + MakeRopeArgs(cfg)); + } else { + // DEFAULT (byte-identical, deterministic): in-place bf16 NeoX RoPE with + // per-element fp64 cos/sin, mirroring vLLM's rotary_emb bf16 rounding. + vt::RopeNeox(d.q, q3, k3, si.positions.t(), MakeRopeArgs(cfg)); + } } } diff --git a/include/vllm/model_executor/models/lm_head_projection.h b/include/vllm/model_executor/models/lm_head_projection.h new file mode 100644 index 0000000000..f2e314200b --- /dev/null +++ b/include/vllm/model_executor/models/lm_head_projection.h @@ -0,0 +1,73 @@ +// vllm.cpp original (no upstream mirror): the LM-head output boundary. +// +// WHY THIS IS A SEAM AND NOT THREE INLINE LINES. The dtype of the LM-head +// projection's output is production behavior (issue #3116), and it is otherwise +// reachable only through a full model forward with a checkpoint. Exposing the +// projection over primitive `vt::` types lets a focused test execute the exact +// operator sequence the forward runs, over the primary's own captured head +// input, with no checkpoint and no GPU-only build. `qwen3_5_moe_block.h` is the +// same pattern for the MoE block. This is not a second path: `ForwardLayers` has +// exactly one head call site and it is this function. +#pragma once + +#include + +#include "vllm/model_executor/models/dense_device_glue.h" +#include "vt/dtype.h" +#include "vt/tensor.h" + +namespace vllm { +namespace lm_head { + +// out [n_out, vocab] = src [n_out, H] @ lm, returned as an owning F32 device +// buffer whose VALUES are BF16 words. +// +// `tied` selects the operator the forward selects: `vt::MatmulBT` with the +// tied embedding weight `lm` [vocab, H], or `vt::Matmul` with the untied +// lm_head weight `lm` [H, vocab] that the loader transposed at load +// (`qwen3_moe_weights.cpp`, `LoadBf16Transposed("lm_head.weight")`). +// +// THE BOUNDARY, and why it is drawn here. The pinned compiled primary stores its +// LM-head output in BF16 and widens it to F32 for the sampler +// (`.agents/specs/rocm-residual-norm.md:320`); the native forward used to store +// F32 directly, which is a measured dtype divergence of up to half a BF16 ulp +// per logit (`.agents/specs/rocm-lmhead-bf16.md`, row 0 max 9.23157e-04, row 1 +// max 9.72956e-04). So the projection narrows to BF16 exactly as the primary +// does, and the SHARED `vt::CastF32` widens it back — the same op the primary's +// logits processor stands in for, and the same widening the tree already uses +// where a BF16 producer feeds an F32 consumer. +// +// THE RETURNED BUFFER STAYS F32, and that is the whole reason the widening is +// here rather than at the callers. Every consumer of the forward's logits reads +// F32: the device-logits view and its owning wrapper (`WrapDeviceLogits`, +// `ViewDeviceLogits`), the host download in `Qwen3MoeModel::Forward`, the +// captured decode graph's logits slot, and the sampler. `.agents/specs/ +// rocm-residual-norm.md:325` records why: "Changing only the head buffer dtype +// would leave those views and host copies invalid." +// +// THE F32 STORE IS NOT A LEVER THIS FUNCTION OWNS. Narrowing the head output +// changes no step-6 token on either measured input: on the primary's captured +// head input both boundaries choose 118, and on the recorded native production +// logits both choose 63, where narrowing collapses the 63/118 margin to an exact +// tie that the native lowest-index tie-break resolves the same way. Hidden-state +// parity (#3115) is still required; this function mirrors a dtype and claims +// nothing more. +inline dense_attn::DBuf Project(dense_attn::Dev d, const vt::Tensor& src, + const vt::Tensor& lm, bool tied) { + const int64_t n_out = src.shape[0]; + const int64_t vocab = tied ? lm.shape[0] : lm.shape[1]; + // The primary's head output dtype. `vt::Matmul`/`vt::MatmulBT` already admit a + // BF16 store with F32 accumulation (`include/vt/ops.h:2649-2652`, + // `src/vt/ops.cpp:125-126`), so no new op and no host round trip is involved. + dense_attn::DBuf head(d, vt::DType::kBF16, {n_out, vocab}); + if (tied) + vt::MatmulBT(d.q, head.t(), src, lm); + else + vt::Matmul(d.q, head.t(), src, lm); + dense_attn::DBuf logits(d, vt::DType::kF32, {n_out, vocab}); + vt::CastF32(d.q, logits.t(), head.t()); + return logits; +} + +} // namespace lm_head +} // namespace vllm diff --git a/include/vt/backend.h b/include/vt/backend.h index ebfaac61aa..4c1be461c8 100644 --- a/include/vt/backend.h +++ b/include/vt/backend.h @@ -5,6 +5,7 @@ #include "vt/device.h" #include "vt/dtype.h" +#include "vt/residual_norm.h" namespace vt { @@ -23,6 +24,12 @@ class Backend { public: virtual ~Backend() = default; + // Numeric policy selected before model execution. CPU reference availability + // alone must not change a production backend's residual lifetime. + virtual ResidualNormPolicy GetResidualNormPolicy() const { + return ResidualNormPolicy::kMaterialized; + } + // Returns memory aligned to at least 64 bytes; StepArena depends on this. virtual void* Alloc(size_t bytes) = 0; virtual void Free(void* p) = 0; diff --git a/include/vt/fused_recipe.h b/include/vt/fused_recipe.h index 42465a53ad..4d757f8ccd 100644 --- a/include/vt/fused_recipe.h +++ b/include/vt/fused_recipe.h @@ -27,6 +27,7 @@ #include #include +#include "vt/residual_norm.h" namespace vt { @@ -66,6 +67,7 @@ enum class FOp : uint8_t { kQuantFp8, // static per-tensor fp8 -> vt::QuantFp8Static [composite terminal] kQuantFp4, // dynamic per-group fp4 -> vt::ScaledFp4Quant [composite terminal] kAttnQkNormRopeGate, // fused attn preamble -> vt::AttnQkNormRopeGate [composite macro] + kResidualRmsNorm, // ordered read-only residual expression -> vt::ResidualRmsNorm }; // Row-reduction kind for reducing ops (kRmsNorm). kMeanSquare = mean(a^2) over @@ -104,7 +106,7 @@ constexpr int kNoFastOp = -1; // no heap). The slack is headroom for the finite Class-A pattern set. constexpr int kMaxFusedSteps = 8; constexpr int kMaxFusedOperands = 8; -constexpr int kMaxStepIns = 3; +constexpr int kMaxStepIns = 4; // One operand slot in the recipe's INDEXED operand table. Steps reference // operands by their index into FusedRecipe::operands (== the index into the @@ -123,7 +125,7 @@ struct FOperandSlot { struct FStep { FOp op = FOp::kAdd; uint8_t out = 0; // primary output operand index - uint8_t in[kMaxStepIns] = {0, 0, 0}; // input operand indices + uint8_t in[kMaxStepIns] = {}; // input operand indices uint8_t nin = 0; // number of inputs used uint8_t out2 = kNoOperand; // secondary output (e.g. fp4 scale) or kNoOperand FReduce reduce = FReduce::kNone; @@ -141,6 +143,7 @@ struct FStep { // full-width) be told apart. Additive: default false keeps every existing // recipe byte-identical. bool norm_full_width = false; + ResidualNormDesc residual_norm{}; // kResidualRmsNorm arithmetic and output aliases }; // The declaration: a fixed-size step list + an indexed operand table + live diff --git a/include/vt/merged_gemm.h b/include/vt/merged_gemm.h index 871c701237..cb54f6b5da 100644 --- a/include/vt/merged_gemm.h +++ b/include/vt/merged_gemm.h @@ -116,6 +116,9 @@ inline constexpr MergedGemmGroup kFp8BlockQkv = { // op vt::MoeGroupedGemmBf16GateUpSilu / OpId::kMoeGroupedGemmBf16GateUpSilu — the // bf16 twin of kMoeGateUpSwiGLUGrouped, BIT-IDENTICAL to {2x MoeGroupedGemmBf16 + // MoeSiluMul}. Same family, distinct weight-marshaling seam. +// The typed native sibling MoeGroupedGemmBf16GateUpSiluNative uses the same +// representation but rounds gate/up and SiLU through BF16 at the pinned vLLM +// boundaries. The original sibling retains FP32 intermediates for compatibility. // // NON-GATED experts are NOT in this family at all, and deliberately get no // descriptor. NemotronH's expert (models/nemotron_h.py:126-256 @ 555967922) has diff --git a/include/vt/ops.h b/include/vt/ops.h index 2ef14ca1a1..80e054093c 100644 --- a/include/vt/ops.h +++ b/include/vt/ops.h @@ -825,6 +825,12 @@ enum class OpId : uint8_t { // blocks * BlockElems}. Decode-only (BACKEND-TENSTORRENT-KEEPQUANT W1); // the dot provider and the keep-quant predicate arm ride W2. kKeepQuantDecode, + // Explicit native BF16 MoE numerics. The legacy grouped signatures and + // FP32 gate/up intermediates remain unchanged for existing callers. + kMoeGroupedGemmBf16GateUpSiluNative, + kMoeGroupedGemmBf16Weighted, + kMoeCombinePreweighted, + kResidualRmsNorm, kCount }; @@ -874,6 +880,11 @@ struct RmsNormArgs { bool gemma = false; // weight applied as (1 + w), GemmaRMSNorm style }; +struct ResidualRmsNormArgs { + float eps = 1e-6f; + ResidualNormDesc descriptor{}; +}; + // Ungated GROUP RMS norm args (vt::RmsNormGroup). A SIBLING of RmsNormArgs, not // a mode of it: see the kRmsNormGroup comment for why the group extent is not a // field on that struct. `eps` and `gemma` keep RmsNormArgs's names and meanings @@ -2187,6 +2198,12 @@ using MoeGroupedGemmBf16Fn = using MoeGroupedGemmBf16GateUpSiluFn = void (*)(Queue&, Tensor& /*out*/, const Tensor& /*act*/, const Tensor& /*expert_ids*/, const Tensor* /*row_map*/, const Tensor& /*gate_ptrs*/, const Tensor& /*up_ptrs*/); +using MoeGroupedGemmBf16GateUpSiluNativeFn = MoeGroupedGemmBf16GateUpSiluFn; +using MoeGroupedGemmBf16WeightedFn = + void (*)(Queue&, Tensor&, const Tensor&, const Tensor&, const Tensor*, const Tensor&, + const Tensor& /*route_weights*/); +using MoeCombinePreweightedFn = + void (*)(Queue&, Tensor&, const Tensor&, const Tensor* /*shared*/, float /*routed_scale*/); // kMatmulBTQuantGrouped: out[P,N], act[P,K] (f32/bf16), weight[E*N,K] block-quant, // expert_ids[P] i32 — weight row for (p,n) is expert_ids[p]*N + n. using MatmulBTQuantGroupedFn = @@ -2283,6 +2300,10 @@ using Exl3MoeMlpFn = void (*)(Queue&, Tensor&, const Tensor&, const Exl3MoeExper const Exl3MoeRouting&, const Exl3MoeTemps&, const Exl3MoeArgs&); using RmsNormFn = void (*)(Queue&, Tensor&, const Tensor&, const Tensor&, const RmsNormArgs&, Tensor*); +using ResidualRmsNormFn = void (*)(Queue&, Tensor& /*out*/, const Tensor& /*a*/, + const Tensor& /*base*/, const Tensor* /*delta*/, + const Tensor& /*weight*/, const ResidualRmsNormArgs&, + Tensor* /*residual_out*/); // Ungated group RMS norm (vt::RmsNormGroup). Same operand order as RmsNormFn // minus the residual, which this op does not carry because its upstream has no // residual arm and a knob nobody can set is a divergence with extra steps. @@ -3237,6 +3258,28 @@ void MoeGroupedGemmBf16GateUpSilu(Queue& q, Tensor& out, const Tensor& act, const Tensor& expert_ids, const Tensor* row_map, const Tensor& gate_ptrs, const Tensor& up_ptrs); +// Native BF16 boundaries from vLLM e126687a9a, triton_moe.py:388-527 and +// activation_kernels.cu:44,165-177: BF16 gate/up, BF16 SiLU, BF16 product. +// Uses the same pointer-array layout and validation as the legacy sibling. +void MoeGroupedGemmBf16GateUpSiluNative(Queue& q, Tensor& out, const Tensor& act, + const Tensor& expert_ids, const Tensor* row_map, + const Tensor& gate_ptrs, const Tensor& up_ptrs); + +// Grouped dot with one FP32 route weight per pair, multiplied BEFORE output +// conversion (fused_moe.py:593-610 at the same pin). Output is BF16 or FP32. +void MoeGroupedGemmBf16Weighted(Queue& q, Tensor& out, const Tensor& act, + const Tensor& expert_ids, const Tensor* row_map, + const Tensor& weight_ptrs, const Tensor& route_weights); + +// Sum already weighted BF16 expert_out[T,top_k,H] in FP32, then narrow to the +// BF16/FP32 output. No second route-weight multiplication. Optional shared[T,H] +// is BF16/FP32 and is added after routed_scale, matching the legacy shared term. +void MoeCombinePreweighted(Queue& q, Tensor& out, const Tensor& expert_out, + const Tensor* shared = nullptr, float routed_scale = 1.0f); + +// Select the complete native sequence through provider availability. +bool MoeGroupedBf16NativeAvailable(DeviceType device); + // MoeGroupedGemmNvfp4Marlin (lift of vLLM moe_wna16_marlin_gemm, ops.cu:543 — // the Marlin W4A16 kernel vLLM selects for the 35B's NVFP4 MoE experts). One // launch computes the grouped expert projection over all padded (token,expert) @@ -3322,6 +3365,22 @@ void MoeRelu2(Queue& q, Tensor& out, const Tensor& x); void RmsNorm(Queue& q, Tensor& out, const Tensor& x, const Tensor& weight, const RmsNormArgs& args, Tensor* residual = nullptr); +// Compiled BF16 expression: normalize a+base or delta+(a+base), in FP32, +// multiplying gamma before the single BF16 output conversion. An optional +// residual output stores the same unrounded expression as BF16. It is never +// reloaded for normalization. Inputs/output are rank-2 with unit inner stride +// and nonoverlapping rows. Gamma is contiguous BF16 [H]. Empty rows are legal. +// Only descriptor-permitted exact aliases can modify an input. No persistent +// FP32 activation is needed. Other activation/gamma dtypes are refused. +void ResidualRmsNorm(Queue& q, Tensor& out, const Tensor& a, const Tensor& base, + const Tensor* delta, const Tensor& weight, + const ResidualRmsNormArgs& args, Tensor* residual_out = nullptr); + +// The same typed operation through the shared recipe/composite seam. +void FusedChain(Queue& q, Tensor& out, const Tensor& a, const Tensor& base, + const Tensor* delta, const Tensor& weight, + const ResidualRmsNormArgs& args, Tensor* residual_out = nullptr); + // UNGATED PER-GROUP RMS NORM — `Qwen4ExpTextRMSNorm` (transformers v5.16.0 // `models/qwen4_exp/modeling_qwen4_exp.py:158-181`), the `group_size is not // None` arm, executed per row: @@ -3637,6 +3696,15 @@ void AttnQkNormRopeGate(Queue& q, Tensor& q_out, Tensor& k_out, Tensor& gate_out const Tensor& k_norm, const Tensor& cos_sin, const RmsNormArgs& norm_args, const RopeArgs& rope_args); +// Gate-free fused attention preamble (the kAttnQkNormRope recipe): per-head +// standard RMSNorm(q) + RMSNorm(k) + partial RoPE-from-cache, with q3/k3 normed +// and rotated IN PLACE. A backend that registers no kernel for it is a refusal; +// the recipe's Tier-0 composite is reached through vt::FusedChain, which is the +// only caller that falls back. +void AttnQkNormRope(Queue& q, Tensor& q3, Tensor& k3, const Tensor& q_norm, + const Tensor& k_norm, const Tensor& cos_sin, const Tensor& positions, + const RmsNormArgs& norm_args, const RopeArgs& rope_args); + // --- GDN (Gated DeltaNet) ops. Formula reference: .agents/specs/gdn-semantics.md. // All GDN state tensors are caller-allocated f32 and updated IN PLACE // (upstream computes states in f32 and rounds to the cache dtype on store — diff --git a/include/vt/recipes.h b/include/vt/recipes.h index 912c2f2823..a462fc9504 100644 --- a/include/vt/recipes.h +++ b/include/vt/recipes.h @@ -19,6 +19,35 @@ namespace vt { +// Executing vLLM e126687a9a compiled partitions ckic6h6, ctj2x6, and c5slugd. +// The optional delta has its own input slot. Gamma never occupies an activation +// slot. The materialized residual is an independent, optional second output. +constexpr FusedRecipe ResidualRmsNormRecipe(ResidualNormDesc descriptor) { + FusedRecipe recipe{}; + recipe.n = 1; + recipe.n_operands = 6; + recipe.name = "residual_rms_norm"; + recipe.operands[0] = {FKind::kRow, "a"}; + recipe.operands[1] = {FKind::kRow, "base"}; + recipe.operands[2] = {FKind::kRow, "delta"}; + recipe.operands[3] = {FKind::kWeight, "weight"}; + recipe.operands[4] = {FKind::kRow, "out"}; + recipe.operands[5] = {FKind::kRow, "residual_out"}; + auto& step = recipe.steps[0]; + step.op = FOp::kResidualRmsNorm; + step.out = 4; + step.out2 = descriptor.materialize_residual ? 5 : kNoOperand; + step.in[0] = 0; + step.in[1] = 1; + const bool triple = descriptor.expression == ResidualNormExpr::kDeltaPlusAdd; + step.in[2] = triple ? 2 : 3; + step.in[3] = 3; + step.nin = triple ? 4 : 3; + step.reduce = FReduce::kMeanSquare; + step.residual_norm = descriptor; + return recipe; +} + // kFusedAddRmsNorm — residual-add + gemma-RMSNorm, the fused_add_rms_norm chain. // // Transcribes vLLM's add+RMSNorm fusion pattern diff --git a/include/vt/residual_norm.h b/include/vt/residual_norm.h new file mode 100644 index 0000000000..2c2c11853a --- /dev/null +++ b/include/vt/residual_norm.h @@ -0,0 +1,29 @@ +// vllm.cpp shared representation of the compiled vLLM residual expression. +// vllm/ir/ops/layernorm.py:44-62 @ e126687a9a828d513c01a07cd69f025f27d63280. +#pragma once +#include + +namespace vt { + +// Arithmetic is FP32 in this order. Every activation operand remains BF16. +enum class ResidualNormExpr : uint8_t { + kAdd, // a + base: two activation operands + kDeltaPlusAdd, // delta + (a + base): three activation operands +}; + +struct ResidualNormDesc { + ResidualNormExpr expression = ResidualNormExpr::kAdd; + bool materialize_residual = false; + // All inputs are read-only unless an exact, explicitly permitted output alias + // names them. Partial overlaps and overlap between outputs are always invalid. + bool output_alias_delta = false; + bool residual_alias_base = false; +}; + +// A reference registration does not opt a production backend into this policy. +enum class ResidualNormPolicy : uint8_t { + kMaterialized, + kCompiledExpression, +}; + +} // namespace vt diff --git a/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp index 83ea59640f..d56d722b59 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -6992,11 +6992,22 @@ DBuf MoeBlockBf16Cuda(Dev d, const MoeBlockWeights& w, const HfConfig& cfg, // trips), prefill reuses the tuned grouped GEMM twice + the identical silu-mul. // BIT-IDENTICAL to the old sequence. Then the grouped down GEMM (act = per-pair // silu output, identity row-map). expert_out lands as [T,top_k,H] contiguous — - // exactly what MoeCombine consumes. + // exactly what MoeCombine consumes. Native BF16 providers round gate/up and + // SiLU before multiplication, then apply route weights before down narrowing. + // Existing callers keep the legacy FP32 intermediate compatibility contract. + const bool native_bf16 = vt::MoeGroupedBf16NativeAvailable(d.q.device.type); DBuf dact(d, DType::kBF16, {P, I}); - vt::MoeGroupedGemmBf16GateUpSilu(d.q, dact.t(), dh, eids, &dtok, dgate_ptrs, dup_ptrs); + if (native_bf16) + vt::MoeGroupedGemmBf16GateUpSiluNative(d.q, dact.t(), dh, eids, &dtok, dgate_ptrs, dup_ptrs); + else + vt::MoeGroupedGemmBf16GateUpSilu(d.q, dact.t(), dh, eids, &dtok, dgate_ptrs, dup_ptrs); DBuf ddown(d, DType::kBF16, {P, H}); - vt::MoeGroupedGemmBf16(d.q, ddown.t(), dact.t(), eids, nullptr, ddown_ptrs); + if (native_bf16) { + const Tensor weights = Reshape(dtw.t(), {P}); + vt::MoeGroupedGemmBf16Weighted(d.q, ddown.t(), dact.t(), eids, nullptr, ddown_ptrs, weights); + } else { + vt::MoeGroupedGemmBf16(d.q, ddown.t(), dact.t(), eids, nullptr, ddown_ptrs); + } Tensor expert_out = Reshape(ddown.t(), {T, top_k, H}); // Shared expert (SEAM GAP #3): Coder has none (shared_expert_intermediate_size @@ -7006,7 +7017,10 @@ DBuf MoeBlockBf16Cuda(Dev d, const MoeBlockWeights& w, const HfConfig& cfg, std::optional shared; if (has_shared) shared.emplace(SharedExpert(d, w, cfg, dh, T, false)); DBuf dout(d, DType::kBF16, {T, H}); - vt::MoeCombine(d.q, dout.t(), expert_out, dtw.t(), has_shared ? &shared->t() : nullptr); + if (native_bf16) + vt::MoeCombinePreweighted(d.q, dout.t(), expert_out, has_shared ? &shared->t() : nullptr); + else + vt::MoeCombine(d.q, dout.t(), expert_out, dtw.t(), has_shared ? &shared->t() : nullptr); return dout; } @@ -7195,7 +7209,8 @@ DBuf MoeBlock(Dev d, const MoeBlockWeights& w, const HfConfig& cfg, // LAYOUT-GUARDED (MoeBf16FastLayoutOk): only the [K,N] Matmul-B (`nk == false`) // orientation the grouped kernel can read; nk=true producers (35B MTP) fall // through to the reference loop. - if (!fp4 && vt::OpRegistered(vt::OpId::kMoeGroupedGemmBf16, d.q.device.type) && MoeBf16FastEnabled() && + if (!fp4 && vt::OpRegistered(vt::OpId::kMoeGroupedGemmBf16, d.q.device.type) && + vt::OpRegistered(vt::OpId::kMoeGroupedGemmBf16GateUpSilu, d.q.device.type) && MoeBf16FastEnabled() && !w.expert_gate.empty() && MoeBf16FastLayoutOk(w, cfg)) return MoeBlockBf16Cuda(d, w, cfg, dh, T); diff --git a/src/vllm/model_executor/models/qwen3_moe.cpp b/src/vllm/model_executor/models/qwen3_moe.cpp index bf94c29ecd..9798997437 100644 --- a/src/vllm/model_executor/models/qwen3_moe.cpp +++ b/src/vllm/model_executor/models/qwen3_moe.cpp @@ -35,6 +35,7 @@ #include "vllm/model_executor/models/decode_graph_sizes.h" // DecodeGraphSizes/PadToCaptureSize #include "vllm/model_executor/models/dense_attn_block.h" // shared AttnBlock + device glue #include "vllm/model_executor/models/device_pool.h" // DevicePool/Pool/ActivePool (shared) +#include "vllm/model_executor/models/lm_head_projection.h" // lm_head::Project (#3116) #include "vllm/model_executor/models/qwen3_5_internal.h" // detail::EndExpertStreamStep #include "vllm/model_executor/device_placement.h" #include "vllm/model_executor/moe_placement_seam.h" @@ -63,31 +64,43 @@ using namespace dense_attn; // One Qwen3-Coder decoder layer (qwen3_moe.py::Qwen3MoeDecoderLayer): input norm // (std add+RMSNorm) -> attention -> post norm (std add+RMSNorm) -> MoE block. The -// residual accumulator `res` (bf16 [T,H]) is threaded through the two fused -// add+RMSNorm producers; `hidden`/`hidden_hold` carry the current block-output -// delta — a device tensor whose storage is either the previous layer's owning -// MoeBlockOutput (`hidden_hold`) or the embedding buffer (held by the caller). The -// MoE block output becomes the new delta, so the final RMSNorm fuses it into res. +// BF16 residual accumulator `res` follows the backend's numeric policy. Legacy +// callers retain the block-output delta in `hidden`. The compiled policy keeps +// attention and residual operands alive through MoE, then consumes their ordered +// expression here and transfers ownership of the normalized next input. void RunMoeLayer(Dev d, const Qwen3MoeLayerWeights& layer, const HfConfig& cfg, Tensor& hidden, std::shared_ptr& hidden_hold, DBuf& res, const StepInputs& si, const CommonAttentionMetadata& meta, - const PagedKvCache& kv, int64_t T, int64_t layer_index) { + const PagedKvCache& kv, int64_t T, int64_t layer_index, + const OwnedTensor* following_norm, bool final_layer) { const int64_t H = cfg.hidden_size; const float eps = static_cast(cfg.rms_norm_eps); - - Tensor w_in = ResidentWeight(d, layer.input_layernorm, {H}); - DBuf dhn(d, DType::kBF16, {T, H}); - if (FusedChainAdoptEnabled()) { - vt::FusedChain(d.q, dhn.t(), hidden, w_in, &res.t(), vt::kFusedAddRmsNormStd, eps); - } else { - vt::RmsNorm(d.q, dhn.t(), hidden, w_in, vt::RmsNormArgs{eps, false}, &res.t()); + const bool compiled_residual = following_norm != nullptr; + + DBuf dhn; + Tensor normalized = hidden; + if (!compiled_residual || layer_index == 0) { + Tensor w_in = ResidentWeight(d, layer.input_layernorm, {H}); + dhn = DBuf(d, DType::kBF16, {T, H}); + if (FusedChainAdoptEnabled()) { + vt::FusedChain(d.q, dhn.t(), hidden, w_in, &res.t(), vt::kFusedAddRmsNormStd, eps); + } else { + vt::RmsNorm(d.q, dhn.t(), hidden, w_in, vt::RmsNormArgs{eps, false}, &res.t()); + } + normalized = dhn.t(); } - DBuf attn = AttnBlock(d, layer.attn, cfg, dhn.t(), si, meta, kv, T); + DBuf attn = AttnBlock(d, layer.attn, cfg, normalized, si, meta, kv, T); Tensor w_post = ResidentWeight(d, layer.post_attention_layernorm, {H}); DBuf dh2(d, DType::kBF16, {T, H}); - if (FusedChainAdoptEnabled()) { + if (compiled_residual) { + const vt::ResidualRmsNormArgs args{eps, {vt::ResidualNormExpr::kAdd}}; + if (FusedChainAdoptEnabled()) + vt::FusedChain(d.q, dh2.t(), attn.t(), res.t(), nullptr, w_post, args); + else + vt::ResidualRmsNorm(d.q, dh2.t(), attn.t(), res.t(), nullptr, w_post, args); + } else if (FusedChainAdoptEnabled()) { vt::FusedChain(d.q, dh2.t(), attn.t(), w_post, &res.t(), vt::kFusedAddRmsNormStd, eps); } else { vt::RmsNorm(d.q, dh2.t(), attn.t(), w_post, vt::RmsNormArgs{eps, false}, &res.t()); @@ -116,8 +129,26 @@ void RunMoeLayer(Dev d, const Qwen3MoeLayerWeights& layer, const HfConfig& cfg, /*placeable=*/layer.moe.expert_gate_fp4.empty(), "the routed experts are fp4-resident and their device residents are " "built at load"); - hidden = moe.tensor; - hidden_hold = std::move(moe.storage); + if (compiled_residual) { + // vLLM e126687a9a generated ctj2x6:33-58 and c5slugd:33-56 normalize + // m + (attention + residual). Consume it while all three owners are alive. + // The next layer receives an already-normalized input, avoiding a second + // norm over its rounded BF16 residual. The final layer has no residual store. + Tensor weight = ResidentWeight(d, *following_norm, {H}); + DBuf next(d, DType::kBF16, {T, H}); + const vt::ResidualRmsNormArgs args{ + eps, {vt::ResidualNormExpr::kDeltaPlusAdd, !final_layer, false, !final_layer}}; + Tensor* residual_out = final_layer ? nullptr : &res.t(); + if (FusedChainAdoptEnabled()) + vt::FusedChain(d.q, next.t(), attn.t(), res.t(), &moe.tensor, weight, args, residual_out); + else + vt::ResidualRmsNorm(d.q, next.t(), attn.t(), res.t(), &moe.tensor, weight, args, residual_out); + hidden = next.t(); + hidden_hold = next.ReleaseShared(); + } else { + hidden = moe.tensor; + hidden_hold = std::move(moe.storage); + } } // GatherRows: gather the idx-indexed rows of `src` [.,H] into contiguous `dst`. @@ -213,23 +244,37 @@ DBuf ForwardLayers(Dev d, const Tensor& hidden_in, Tensor hidden = hidden_in; std::shared_ptr hidden_hold; // owns the current MoE-output delta storage + const bool compiled_residual = + d.b.GetResidualNormPolicy() == vt::ResidualNormPolicy::kCompiledExpression; + VT_CHECK(!compiled_residual || vt::OpRegistered(vt::OpId::kResidualRmsNorm, d.q.device.type), + "qwen3 moe: compiled residual policy requires the complete residual RMS norm operation"); + DBuf res(d, DType::kBF16, {T, H}); res.Zero(d); StepInputs si = BuildStepInputs(d, positions, attn_meta, config); - for (int64_t l = 0; l < config.num_hidden_layers; ++l) + for (int64_t l = 0; l < config.num_hidden_layers; ++l) { + const bool final_layer = l + 1 == config.num_hidden_layers; + const OwnedTensor* following_norm = !compiled_residual ? nullptr : + (final_layer ? &weights.final_norm : &weights.layers[static_cast(l + 1)].input_layernorm); RunMoeLayer(d, weights.layers[static_cast(l)], config, hidden, hidden_hold, res, si, attn_meta, attn_kv[static_cast(l)], T, - /*layer_index=*/l); + /*layer_index=*/l, following_norm, final_layer); + } // Final RMSNorm over the fused stream (res += hidden; std norm), then lm_head. - Tensor w_fn = ResidentWeight(d, weights.final_norm, {H}); - DBuf dnorm(d, DType::kBF16, {T, H}); - if (FusedChainAdoptEnabled()) { - vt::FusedChain(d.q, dnorm.t(), hidden, w_fn, &res.t(), vt::kFusedAddRmsNormStd, eps); - } else { - vt::RmsNorm(d.q, dnorm.t(), hidden, w_fn, vt::RmsNormArgs{eps, false}, &res.t()); + DBuf dnorm; + Tensor final_normalized = hidden; + if (!compiled_residual) { + Tensor w_fn = ResidentWeight(d, weights.final_norm, {H}); + dnorm = DBuf(d, DType::kBF16, {T, H}); + if (FusedChainAdoptEnabled()) { + vt::FusedChain(d.q, dnorm.t(), hidden, w_fn, &res.t(), vt::kFusedAddRmsNormStd, eps); + } else { + vt::RmsNorm(d.q, dnorm.t(), hidden, w_fn, vt::RmsNormArgs{eps, false}, &res.t()); + } + final_normalized = dnorm.t(); } // lm_head. UNTIED (Qwen3-Coder): the loaded Matmul-B [H,vocab] lm_head via @@ -241,21 +286,17 @@ DBuf ForwardLayers(Dev d, const Tensor& hidden_in, const bool do_gather = !logits_indices.empty() && static_cast(logits_indices.size()) < T; - Tensor src = dnorm.t(); + Tensor src = final_normalized; DBuf dgather(d, DType::kBF16, do_gather ? std::vector{ static_cast(logits_indices.size()), H} : std::vector{1, 1}); if (do_gather) { - GatherRows(d, dgather.ptr(), dnorm.t(), logits_indices, H); + GatherRows(d, dgather.ptr(), final_normalized, logits_indices, H); src = dgather.t(); } - const int64_t n_out = src.shape[0]; - DBuf logits(d, DType::kF32, {n_out, vocab}); - if (tied) - vt::MatmulBT(d.q, logits.t(), src, lm); - else - vt::Matmul(d.q, logits.t(), src, lm); - return logits; + // The head projection and its OUTPUT DTYPE live on one seam (#3116) so the + // boundary a focused test replays is the boundary this forward runs. + return lm_head::Project(d, src, lm, tied); } // Full eager forward body: embed (host token_ids) then the capturable layer diff --git a/src/vt/cpu/cpu_residual_rmsnorm.cpp b/src/vt/cpu/cpu_residual_rmsnorm.cpp new file mode 100644 index 0000000000..84556e8563 --- /dev/null +++ b/src/vt/cpu/cpu_residual_rmsnorm.cpp @@ -0,0 +1,65 @@ +// CPU reference for the compiled vLLM residual expression. +// Ported from vllm/ir/ops/layernorm.py:44-62 @ +// e126687a9a828d513c01a07cd69f025f27d63280, with executing generated kernels +// ckic6h6:33-52, ctj2x6:33-58, and c5slugd:33-56 defining the BF16 boundaries. +#include +#include + +#include "cpu_threadpool.h" +#include "vt/ops.h" + +namespace vt::cpu { +namespace { +constexpr int kBlock = 256; + +float Expression(const Tensor& a, const Tensor& base, const Tensor* delta, + int64_t row, int64_t column) { + const float av = BF16ToF32(a.Ptr()[row * a.stride[0] + column]); + const float rv = BF16ToF32(base.Ptr()[row * base.stride[0] + column]); + const float attention = av + rv; + if (delta == nullptr) return attention; + const float mv = BF16ToF32(delta->Ptr()[row * delta->stride[0] + column]); + return mv + attention; +} + +void ResidualRmsNormKernel(Queue&, Tensor& out, const Tensor& a, const Tensor& base, + const Tensor* delta, const Tensor& weight, + const ResidualRmsNormArgs& args, Tensor* residual_out) { + const int64_t rows = a.shape[0], width = a.shape[1]; + ParallelForRows(CurrentThreadpool(), rows, [&](int64_t begin, int64_t end) { + for (int64_t row = begin; row < end; ++row) { + // Row-local FP32 reduction scratch mirrors the deterministic GPU tree. + // No activation or residual is materialized at FP32 in model storage. + std::array partial{}; + for (int lane = 0; lane < kBlock; ++lane) { + float acc = 0.0f; + for (int64_t column = lane; column < width; column += kBlock) { + const float value = Expression(a, base, delta, row, column); + acc += value * value; + } + partial[static_cast(lane)] = acc; + } + for (int stride = kBlock / 2; stride > 0; stride /= 2) + for (int lane = 0; lane < stride; ++lane) + partial[static_cast(lane)] += partial[static_cast(lane + stride)]; + const float inverse = 1.0f / std::sqrt(partial[0] / static_cast(width) + args.eps); + for (int64_t column = 0; column < width; ++column) { + // Read the complete unrounded value before either permitted alias writes. + const float value = Expression(a, base, delta, row, column); + const float gamma = BF16ToF32(weight.Ptr()[column]); + out.Ptr()[row * out.stride[0] + column] = F32ToBF16((value * inverse) * gamma); + if (residual_out != nullptr) + residual_out->Ptr()[row * residual_out->stride[0] + column] = F32ToBF16(value); + } + } + }); +} + +struct Registrar { + Registrar() { + RegisterOp(OpId::kResidualRmsNorm, DeviceType::kCPU, + reinterpret_cast(static_cast(&ResidualRmsNormKernel))); + } +} registrar; +} // namespace +} // namespace vt::cpu diff --git a/src/vt/op_provider.cpp b/src/vt/op_provider.cpp index 071a021370..d34eb09e41 100644 --- a/src/vt/op_provider.cpp +++ b/src/vt/op_provider.cpp @@ -292,6 +292,8 @@ const char* OpNameImpl(OpId op) { return "Matmul"; case OpId::kRmsNorm: return "RmsNorm"; + case OpId::kResidualRmsNorm: + return "ResidualRmsNorm"; case OpId::kSiluAndMul: return "SiluAndMul"; case OpId::kRopeNeox: @@ -522,6 +524,12 @@ const char* OpNameImpl(OpId op) { return "FusedNormRope"; case OpId::kMoeGroupedGemmBf16GateUpSilu: return "MoeGroupedGemmBf16GateUpSilu"; + case OpId::kMoeGroupedGemmBf16GateUpSiluNative: + return "MoeGroupedGemmBf16GateUpSiluNative"; + case OpId::kMoeGroupedGemmBf16Weighted: + return "MoeGroupedGemmBf16Weighted"; + case OpId::kMoeCombinePreweighted: + return "MoeCombinePreweighted"; case OpId::kLaguna: return "Laguna"; case OpId::kMarlinDenseGemm: diff --git a/src/vt/ops.cpp b/src/vt/ops.cpp index 89e00a187e..955ac9901a 100644 --- a/src/vt/ops.cpp +++ b/src/vt/ops.cpp @@ -1,9 +1,11 @@ // vllm.cpp original (vt runtime, inventory deviation §9.1); no upstream mirror. #include "vt/ops.h" +#include "vt/recipes.h" #include "vt/paged_attn_route.h" // W10 repair (#1865): the uniform-spec shape guard #include #include +#include #include #include #include @@ -903,7 +905,8 @@ void MoeGroupedGemmNvfp4(Queue& q, Tensor& out, const Tensor& act, const Tensor& q, out, act, expert_ids, row_map, packed_ptrs, scale_ptrs, scale2s); } -void MoeGroupedGemmBf16(Queue& q, Tensor& out, const Tensor& act, const Tensor& expert_ids, +namespace { +void ValidateMoeGroupedGemmBf16(Queue& q, const Tensor& out, const Tensor& act, const Tensor& expert_ids, const Tensor* row_map, const Tensor& weight_ptrs) { VT_CHECK(out.rank == 2 && act.rank == 2, "moe_grouped_gemm_bf16: out/act must be rank-2"); const int64_t p = out.shape[0], e = weight_ptrs.shape[0]; @@ -925,11 +928,18 @@ void MoeGroupedGemmBf16(Queue& q, Tensor& out, const Tensor& act, const Tensor& row_map->device == q.device, "moe_grouped_gemm_bf16: row_map must be contiguous i32 [P] on the queue device"); } +} +} // namespace + +void MoeGroupedGemmBf16(Queue& q, Tensor& out, const Tensor& act, const Tensor& expert_ids, + const Tensor* row_map, const Tensor& weight_ptrs) { + ValidateMoeGroupedGemmBf16(q, out, act, expert_ids, row_map, weight_ptrs); reinterpret_cast(GetOp(OpId::kMoeGroupedGemmBf16, q.device.type))( q, out, act, expert_ids, row_map, weight_ptrs); } -void MoeGroupedGemmBf16GateUpSilu(Queue& q, Tensor& out, const Tensor& act, +namespace { +void ValidateMoeGroupedGemmBf16GateUpSilu(Queue& q, const Tensor& out, const Tensor& act, const Tensor& expert_ids, const Tensor* row_map, const Tensor& gate_ptrs, const Tensor& up_ptrs) { VT_CHECK(out.rank == 2 && act.rank == 2, @@ -956,11 +966,68 @@ void MoeGroupedGemmBf16GateUpSilu(Queue& q, Tensor& out, const Tensor& act, row_map->device == q.device, "moe_grouped_gemm_bf16_gate_up_silu: row_map must be contiguous i32 [P] on the device"); } +} +} // namespace + +void MoeGroupedGemmBf16GateUpSilu(Queue& q, Tensor& out, const Tensor& act, + const Tensor& expert_ids, const Tensor* row_map, + const Tensor& gate_ptrs, const Tensor& up_ptrs) { + ValidateMoeGroupedGemmBf16GateUpSilu(q, out, act, expert_ids, row_map, gate_ptrs, up_ptrs); reinterpret_cast( GetOp(OpId::kMoeGroupedGemmBf16GateUpSilu, q.device.type))(q, out, act, expert_ids, row_map, gate_ptrs, up_ptrs); } +void MoeGroupedGemmBf16GateUpSiluNative(Queue& q, Tensor& out, const Tensor& act, + const Tensor& expert_ids, const Tensor* row_map, + const Tensor& gate_ptrs, const Tensor& up_ptrs) { + ValidateMoeGroupedGemmBf16GateUpSilu(q, out, act, expert_ids, row_map, gate_ptrs, up_ptrs); + reinterpret_cast( + GetOp(OpId::kMoeGroupedGemmBf16GateUpSiluNative, q.device.type))( + q, out, act, expert_ids, row_map, gate_ptrs, up_ptrs); +} + +void MoeGroupedGemmBf16Weighted(Queue& q, Tensor& out, const Tensor& act, + const Tensor& expert_ids, const Tensor* row_map, + const Tensor& weight_ptrs, const Tensor& route_weights) { + ValidateMoeGroupedGemmBf16(q, out, act, expert_ids, row_map, weight_ptrs); + VT_CHECK(route_weights.rank == 1 && route_weights.Numel() == out.shape[0] && + route_weights.dtype == DType::kF32 && route_weights.IsContiguous() && + route_weights.device == q.device, + "moe_grouped_gemm_bf16_weighted: route_weights must be contiguous f32 [P] on the device"); + reinterpret_cast( + GetOp(OpId::kMoeGroupedGemmBf16Weighted, q.device.type))( + q, out, act, expert_ids, row_map, weight_ptrs, route_weights); +} + +void MoeCombinePreweighted(Queue& q, Tensor& out, const Tensor& expert_out, + const Tensor* shared, float routed_scale) { + VT_CHECK(out.rank == 2 && expert_out.rank == 3 && + expert_out.shape[0] == out.shape[0] && expert_out.shape[2] == out.shape[1], + "moe_combine_preweighted: expert_out [T,K,H] must match out [T,H]"); + VT_CHECK(expert_out.dtype == DType::kBF16 && IsOutFloat(out.dtype), + "moe_combine_preweighted: expert_out must be bf16, out must be f32/bf16"); + VT_CHECK(expert_out.IsContiguous() && out.IsContiguous() && + expert_out.device == q.device && out.device == q.device, + "moe_combine_preweighted: contiguous tensors on the queue device required"); + if (shared != nullptr) { + VT_CHECK(shared->rank == 2 && shared->shape[0] == out.shape[0] && + shared->shape[1] == out.shape[1] && IsOutFloat(shared->dtype) && + shared->IsContiguous() && shared->device == q.device, + "moe_combine_preweighted: shared must be f32/bf16 [T,H] on the queue device"); + } + reinterpret_cast( + GetOp(OpId::kMoeCombinePreweighted, q.device.type))(q, out, expert_out, shared, routed_scale); +} + +bool MoeGroupedBf16NativeAvailable(DeviceType device) { + return OpRegistered(OpId::kMoeGroupedGemmBf16, device) && + OpRegistered(OpId::kMoeGroupedGemmBf16GateUpSilu, device) && + OpRegistered(OpId::kMoeGroupedGemmBf16GateUpSiluNative, device) && + OpRegistered(OpId::kMoeGroupedGemmBf16Weighted, device) && + OpRegistered(OpId::kMoeCombinePreweighted, device); +} + void MoeGroupedGemmNvfp4Marlin(Queue& q, Tensor& c, const Tensor& a, const Tensor& b_q_weight, const Tensor& b_scales, const Tensor& global_scale, Tensor& workspace, const Tensor& sorted_token_ids, @@ -1046,6 +1113,88 @@ void RmsNorm(Queue& q, Tensor& out, const Tensor& x, const Tensor& weight, residual); } +void ResidualRmsNorm(Queue& q, Tensor& out, const Tensor& a, const Tensor& base, + const Tensor* delta, const Tensor& weight, + const ResidualRmsNormArgs& args, Tensor* residual_out) { + const auto& desc = args.descriptor; + VT_CHECK(desc.expression == ResidualNormExpr::kAdd || + desc.expression == ResidualNormExpr::kDeltaPlusAdd, + "residual_rmsnorm: invalid expression order"); + const bool triple = desc.expression == ResidualNormExpr::kDeltaPlusAdd; + VT_CHECK(triple == (delta != nullptr), "residual_rmsnorm: expression operand count mismatch"); + VT_CHECK(desc.materialize_residual == (residual_out != nullptr), + "residual_rmsnorm: materialization descriptor/output mismatch"); + VT_CHECK(!desc.output_alias_delta || triple, + "residual_rmsnorm: output alias requires a delta operand"); + VT_CHECK(!desc.residual_alias_base || desc.materialize_residual, + "residual_rmsnorm: base alias requires residual materialization"); + VT_CHECK(std::isfinite(args.eps) && args.eps >= 0.0f, + "residual_rmsnorm: epsilon must be finite and nonnegative"); + VT_CHECK(a.rank == 2 && a.shape[0] >= 0 && a.shape[1] > 0, + "residual_rmsnorm: expected nonnegative rows and positive hidden width"); + VT_CHECK(q.device.index >= 0, "residual_rmsnorm: queue device index must be nonnegative"); + const auto row_bytes = [&](const Tensor& tensor) -> size_t { + VT_CHECK(tensor.rank == 2 && tensor.shape[0] == a.shape[0] && + tensor.shape[1] == a.shape[1], "residual_rmsnorm: row shape mismatch"); + VT_CHECK(tensor.dtype == DType::kBF16, "residual_rmsnorm: BF16 activation/output required"); + VT_CHECK(tensor.device == q.device, "residual_rmsnorm: queue/tensor device mismatch"); + VT_CHECK(tensor.stride[1] == 1 && tensor.stride[0] >= tensor.shape[1], + "residual_rmsnorm: unit inner stride and nonoverlapping rows required"); + if (tensor.shape[0] == 0) return 0; + const auto rows = static_cast(tensor.shape[0]); + const auto width = static_cast(tensor.shape[1]); + const auto stride = static_cast(tensor.stride[0]); + const uint64_t max_elements = std::numeric_limits::max() / sizeof(uint16_t); + VT_CHECK(width <= max_elements && (rows - 1) <= (max_elements - width) / stride, + "residual_rmsnorm: row storage span overflow"); + const size_t bytes = static_cast((rows - 1) * stride + width) * sizeof(uint16_t); + const uintptr_t start = reinterpret_cast(tensor.data); + VT_CHECK(start != 0 && start % alignof(uint16_t) == 0 && + start <= std::numeric_limits::max() - bytes, + "residual_rmsnorm: invalid or unaligned row data"); + return bytes; + }; + const size_t a_bytes = row_bytes(a), base_bytes = row_bytes(base), out_bytes = row_bytes(out); + const size_t delta_bytes = delta == nullptr ? 0 : row_bytes(*delta); + const size_t res_bytes = residual_out == nullptr ? 0 : row_bytes(*residual_out); + VT_CHECK(weight.rank == 1 && weight.shape[0] == a.shape[1] && weight.stride[0] == 1 && + weight.dtype == DType::kBF16 && weight.device == q.device, + "residual_rmsnorm: gamma must be contiguous BF16 [H] on the queue device"); + const uintptr_t w_start = reinterpret_cast(weight.data); + const auto w_width = static_cast(weight.shape[0]); + VT_CHECK(w_width <= std::numeric_limits::max() / sizeof(uint16_t), + "residual_rmsnorm: gamma storage span overflow"); + const size_t w_bytes = static_cast(w_width) * sizeof(uint16_t); + VT_CHECK(w_start != 0 && w_start % alignof(uint16_t) == 0 && + w_start <= std::numeric_limits::max() - w_bytes, + "residual_rmsnorm: invalid or unaligned gamma data"); + const auto overlaps = [](const Tensor& lhs, size_t lhs_bytes, + const Tensor& rhs, size_t rhs_bytes) { + const uintptr_t left = reinterpret_cast(lhs.data); + const uintptr_t right = reinterpret_cast(rhs.data); + return lhs_bytes != 0 && rhs_bytes != 0 && left < right + rhs_bytes && right < left + lhs_bytes; + }; + const auto same_view = [](const Tensor& lhs, const Tensor& rhs) { + return lhs.data == rhs.data && lhs.stride[0] == rhs.stride[0]; + }; + for (const auto& input : {std::pair{&a, a_bytes}, + {&base, base_bytes}, {delta, delta_bytes}, {&weight, w_bytes}}) { + if (input.first == nullptr) continue; + VT_CHECK(!overlaps(out, out_bytes, *input.first, input.second) || + (input.first == delta && desc.output_alias_delta && same_view(out, *delta)), + "residual_rmsnorm: normalized output overlaps a read-only input"); + if (residual_out != nullptr) + VT_CHECK(!overlaps(*residual_out, res_bytes, *input.first, input.second) || + (input.first == &base && desc.residual_alias_base && same_view(*residual_out, base)), + "residual_rmsnorm: residual output overlaps a read-only input"); + } + VT_CHECK(residual_out == nullptr || !overlaps(out, out_bytes, *residual_out, res_bytes), + "residual_rmsnorm: outputs must not overlap"); + if (a.shape[0] == 0) return; + reinterpret_cast(GetOp(OpId::kResidualRmsNorm, q.device.type))( + q, out, a, base, delta, weight, args, residual_out); +} + void RmsNormGroup(Queue& q, Tensor& out, const Tensor& x, const Tensor& weight, const RmsNormGroupArgs& args) { VT_CHECK(x.rank == 2 && out.rank == 2 && weight.rank == 1, @@ -1101,6 +1250,24 @@ void FusedChainCompositeImpl(Queue& q, const FusedRecipe& r, const FusedBinding& for (int s = 0; s < r.n; ++s) { const FStep& st = r.steps[s]; switch (st.op) { + case FOp::kResidualRmsNorm: { + VT_CHECK(!add_pending, "fused_chain: residual expression cannot consume a pending add"); + const bool triple = st.residual_norm.expression == ResidualNormExpr::kDeltaPlusAdd; + VT_CHECK(st.nin == (triple ? 4 : 3), "fused_chain: residual expression input count"); + VT_CHECK((st.out2 != kNoOperand) == st.residual_norm.materialize_residual, + "fused_chain: residual expression materialization mismatch"); + VT_CHECK(st.reduce == FReduce::kMeanSquare && !st.gemma && !st.sigmoid_gate && + !st.norm_full_width, "fused_chain: invalid residual expression modifiers"); + Tensor* out = FusedOp(b, st.out, "fused_chain: null residual norm output"); + Tensor* a = FusedOp(b, st.in[0], "fused_chain: null residual expression a"); + Tensor* base = FusedOp(b, st.in[1], "fused_chain: null residual expression base"); + Tensor* delta = triple ? FusedOp(b, st.in[2], "fused_chain: null delta") : nullptr; + Tensor* weight = FusedOp(b, st.in[triple ? 3 : 2], "fused_chain: null residual gamma"); + Tensor* residual = st.out2 == kNoOperand ? nullptr : FusedOp(b, st.out2, "fused_chain: null materialized residual"); + ResidualRmsNorm(q, *out, *a, *base, delta, *weight, + ResidualRmsNormArgs{p.eps, st.residual_norm}, residual); + break; + } case FOp::kAdd: // Residual-add producing the residual stream: fold into the next kRmsNorm. VT_CHECK(st.nin == 2 && st.out == st.in[1], @@ -1301,6 +1468,28 @@ void FusedChain(Queue& q, const FusedRecipe& recipe, const FusedBinding& binding FusedChainComposite(q, recipe, binding, params); } +void FusedChain(Queue& q, Tensor& out, const Tensor& a, const Tensor& base, + const Tensor* delta, const Tensor& weight, + const ResidualRmsNormArgs& args, Tensor* residual_out) { + FusedBinding binding{}; + binding.n = 6; + binding.op[0] = const_cast(&a); + binding.op[1] = const_cast(&base); + binding.op[2] = const_cast(delta); + binding.op[3] = const_cast(&weight); + binding.op[4] = &out; + binding.op[5] = residual_out; + // Preserve mismatched optional arguments until the typed validation can reject + // them. A recipe cannot silently hide an unexpected delta or residual output. + VT_CHECK((delta != nullptr) == (args.descriptor.expression == ResidualNormExpr::kDeltaPlusAdd), + "fused_chain: residual expression operand count mismatch"); + VT_CHECK((residual_out != nullptr) == args.descriptor.materialize_residual, + "fused_chain: residual expression materialization mismatch"); + FusedParams params{}; + params.eps = args.eps; + FusedChain(q, ResidualRmsNormRecipe(args.descriptor), binding, params); +} + void FusedChain(Queue& q, Tensor& out, const Tensor& x, const Tensor& weight, Tensor* residual, const FusedRecipe& recipe, float eps) { // Canonical 4-operand shape (the W0-adopted kFusedAddRmsNorm site): x/out [T,H], @@ -1758,6 +1947,40 @@ void AttnQkNormRopeGate(Queue& q, Tensor& q_out, Tensor& k_out, Tensor& gate_out q, q_out, k_out, gate_out, qgate, kf, q_norm, k_norm, cos_sin, norm_args, rope_args); } +void AttnQkNormRope(Queue& q, Tensor& q3, Tensor& k3, const Tensor& q_norm, + const Tensor& k_norm, const Tensor& cos_sin, const Tensor& positions, + const RmsNormArgs& norm_args, const RopeArgs& rope_args) { + VT_CHECK(q3.rank == 3 && k3.rank == 3, "attn_qk_norm_rope: q3/k3 rank-3 [T,H,Dh]"); + const int64_t t = q3.shape[0], dh = q3.shape[2]; + VT_CHECK(k3.shape[0] == t && k3.shape[2] == dh, "attn_qk_norm_rope: k3 must be [T,Hkv,Dh]"); + VT_CHECK(q3.dtype == k3.dtype, "attn_qk_norm_rope: q3/k3 dtype"); + VT_CHECK(IsFloat(q3.dtype), "attn_qk_norm_rope: q3/k3 must be f32 or bf16"); + VT_CHECK(q_norm.rank == 1 && q_norm.shape[0] == dh && k_norm.rank == 1 && + k_norm.shape[0] == dh, + "attn_qk_norm_rope: q_norm/k_norm must be [Dh]"); + VT_CHECK(positions.rank == 1 && positions.shape[0] == t, + "attn_qk_norm_rope: positions must be [T]"); + VT_CHECK(rope_args.rotary_dim > 0 && rope_args.rotary_dim % 2 == 0 && + rope_args.rotary_dim <= dh, + "attn_qk_norm_rope: rotary_dim must be even and <= Dh"); + VT_CHECK(cos_sin.rank == 2 && cos_sin.shape[0] > 0 && + cos_sin.shape[1] == rope_args.rotary_dim, + "attn_qk_norm_rope: cos_sin must be [rows, rotary_dim]"); + VT_CHECK(q3.IsContiguous() && k3.IsContiguous() && q_norm.IsContiguous() && + k_norm.IsContiguous() && cos_sin.IsContiguous() && positions.IsContiguous(), + "attn_qk_norm_rope: states/weights/cache/index must be contiguous"); + // The fused op rotates the in-place operands, so the 2-D alias the composite's + // RmsNorm step would have normed is the same memory: [T*H,Dh] with stride Dh. + // That alias is a RESHAPE of the rank-3 view, so on a row-major [T,H,Dh] tensor + // the row stride is stride[1] == Dh and the INNER dimension is stride[2]. The + // check named stride[1] and so refused every Dh > 1 operand — every real call, + // the hand-call realization at dense_attn_block.h:648 among them. + VT_CHECK(q3.stride[2] == 1 && k3.stride[2] == 1, + "attn_qk_norm_rope: the head dimension must be the inner dimension"); + reinterpret_cast(GetOp(OpId::kAttnQkNormRope, q.device.type))( + q, q3, k3, q_norm, k_norm, cos_sin, positions, norm_args, rope_args); +} + namespace { // Shared shape/dtype/device validation for the two conv ops. x/out [T,C], // weight [C,K], optional bias [C], conv_state [N,C,K-1] f32. diff --git a/src/vt/rocm/rocm_backend.hip b/src/vt/rocm/rocm_backend.hip index 42eea3d3b6..fd6932b4dc 100644 --- a/src/vt/rocm/rocm_backend.hip +++ b/src/vt/rocm/rocm_backend.hip @@ -215,6 +215,10 @@ class RocmBackend final : public Backend { int DeviceCapabilityMajor() const override { return major_; } int DeviceCapabilityMinor() const override { return minor_; } + ResidualNormPolicy GetResidualNormPolicy() const override { + return ResidualNormPolicy::kCompiledExpression; + } + // Both branches satisfy the >= 64B contract Backend::Alloc owes StepArena: // hipMalloc aligns to at least 256 bytes as cudaMalloc does, and managed // allocations are at least page-granular. diff --git a/src/vt/rocm/rocm_gdn_fused.hip b/src/vt/rocm/rocm_gdn_fused.hip index dff7937e13..76e64f8fe0 100644 --- a/src/vt/rocm/rocm_gdn_fused.hip +++ b/src/vt/rocm/rocm_gdn_fused.hip @@ -229,6 +229,64 @@ void SigmoidGateBf16KernelRocm(Queue& q, Tensor& out, const Tensor& attn, Check(hipGetLastError(), "sigmoid_gate launch"); } +// ── Gate-free fused attention preamble (kAttnQkNormRope) ──────────────────── +// Per-head STANDARD RMSNorm(q) + RMSNorm(k) + partial RoPE-from-cache, the +// Qwen3-dense preamble, in one launch. The normalized value is carried in F32 +// through the rotation and narrowed ONCE, at the store: that is the pinned +// primary's boundary -- Qwen3Attention applies q_norm/k_norm and then rotary_emb +// (vllm/model_executor/models/qwen3.py:150-167) and the compiled preamble reads a +// cos|sin cache narrowed to the query dtype +// (vllm/model_executor/layers/rotary_embedding/base.py:105-131), so nothing is +// rounded to bf16 before the rotation. The Tier-0 composite this op realises +// (RmsNorm into a bf16 buffer, then RopeFromCache over that buffer) rounds twice +// and measured 1569 of 8448 Q words against the primary on the L33/C2/R0 +// workload; see .agents/specs/rocm-attn-parity.md. +// The RMSNorm reduction is the shipped one -- same kBlock-strided accumulation +// and same shared-memory tree as RmsNormRowKernel (rocm_rmsnorm.hip:118-147) -- +// so the f32 carrier keeps the norm's own bits and only the narrowing point +// moves. +template +__global__ void AttnQkNormRopeRowK(T* x, const T* w, const T* cache, const Tid* positions, + int64_t cache_rows, int64_t heads, int64_t dh, + int64_t row_stride, int rot, float eps, bool neox) { + const int64_t row = blockIdx.x; + const int64_t pos = static_cast(positions[row / heads]); + if (rot <= 0 || pos < 0 || pos >= cache_rows) return; + T* xrow = x + row * row_stride; + const T* cs = cache + pos * rot; + __shared__ float partial[kBlock]; + float acc = 0.0f; + for (int64_t j = threadIdx.x; j < dh; j += kBlock) { + const float v = Ld(xrow, j); + acc += v * v; + } + partial[threadIdx.x] = acc; + __syncthreads(); + for (int s = kBlock / 2; s > 0; s /= 2) { + if (static_cast(threadIdx.x) < s) partial[threadIdx.x] += partial[threadIdx.x + s]; + __syncthreads(); + } + const float inv = 1.0f / sqrtf(partial[0] / static_cast(dh) + eps); + const int64_t half = rot / 2; + // One thread per rotation PAIR, both stores from the values read before either + // write, because the rotation is in place (RopeFromCacheK, rocm_dense_basic.hip + // :699-704). The pair is the only place the bf16 narrowing is deferred. + for (int64_t pair = threadIdx.x; pair < half; pair += kBlock) { + const int64_t first = neox ? pair : pair * 2; + const int64_t second = neox ? pair + half : pair * 2 + 1; + const float c = Ld(cs, pair); + const float sn = Ld(cs, half + pair); + const float x1 = Ld(xrow, first) * inv * Ld(w, first); + const float x2 = Ld(xrow, second) * inv * Ld(w, second); + St(xrow, first, x1 * c - x2 * sn); + St(xrow, second, x1 * sn + x2 * c); + } + // Outside the rotary width the normed value is only stored, narrowed once. + for (int64_t j = rot + threadIdx.x; j < dh; j += kBlock) { + St(xrow, j, Ld(xrow, j) * inv * Ld(w, j)); + } +} + void AttnQkNormRopeGateKernelRocm(Queue& q, Tensor& q_out, Tensor& k_out, Tensor& gate_out, const Tensor& qgate, const Tensor& kf, const Tensor& q_norm, const Tensor& k_norm, @@ -274,4 +332,68 @@ void AttnQkNormRopeGateKernelRocm(Queue& q, Tensor& q_out, Tensor& k_out, Tensor Check(hipGetLastError(), "attn_qk_norm_rope_gate launch"); } +// The registered AttnQkNormRopeFn (include/vt/ops.h). The f32 attention arm has +// no intermediate rounding to remove -- RmsNorm(f32) then RopeFromCache(f32) is +// already the primary's boundary -- so it keeps the shipped composite's bytes by +// running that composite, and only the bf16 arm takes the fused kernel. +void AttnQkNormRopeKernelRocm(Queue& q, Tensor& q3, Tensor& k3, const Tensor& q_norm, + const Tensor& k_norm, const Tensor& cos_sin, + const Tensor& positions, const RmsNormArgs& na, + const RopeArgs& ra) { + const int64_t t = q3.shape[0], hq = q3.shape[1], dh = q3.shape[2]; + const int64_t hkv = k3.shape[1]; + if (t == 0) return; + VT_CHECK(q3.rank == 3 && k3.rank == 3 && k3.shape[2] == dh, + "rocm attn_qk_norm_rope: q3/k3 must be [T,H,Dh]"); + VT_CHECK(q_norm.rank == 1 && q_norm.shape[0] == dh && k_norm.rank == 1 && + k_norm.shape[0] == dh, + "rocm attn_qk_norm_rope: q_norm/k_norm must be [Dh]"); + VT_CHECK(cos_sin.rank == 2 && cos_sin.shape[1] == ra.rotary_dim, + "rocm attn_qk_norm_rope: cos_sin must be [rows, rotary_dim]"); + VT_CHECK(cos_sin.dtype == q3.dtype && q_norm.dtype == q3.dtype && k_norm.dtype == q3.dtype, + "rocm attn_qk_norm_rope: weights and cache must match the states' dtype"); + if (q3.dtype == DType::kF32) { + Tensor q2d = q3; + q2d.rank = 2; + q2d.shape[0] = t * hq; + q2d.shape[1] = dh; + q2d.stride[0] = dh; + q2d.stride[1] = 1; + Tensor k2d = k3; + k2d.rank = 2; + k2d.shape[0] = t * hkv; + k2d.shape[1] = dh; + k2d.stride[0] = dh; + k2d.stride[1] = 1; + RmsNorm(q, q2d, q2d, q_norm, na, nullptr); + RmsNorm(q, k2d, k2d, k_norm, na, nullptr); + RopeFromCache(q, q3, &k3, positions, cos_sin, ra); + return; + } + VT_CHECK(q3.dtype == DType::kBF16, + "rocm attn_qk_norm_rope: states must be f32 or bf16"); + const int64_t rot = ra.rotary_dim; + VT_CHECK(rot > 0 && rot % 2 == 0 && rot <= dh, + "rocm attn_qk_norm_rope: rotary_dim must be even and <= Dh"); + hipStream_t s = AsStream(q); + auto launch = [&](auto t_tag, auto id_tag) { + using T = decltype(t_tag); + using Tid = decltype(id_tag); + if (hq > 0) + AttnQkNormRopeRowK<<(t * hq), kBlock, 0, s>>>( + q3.Ptr(), q_norm.Ptr(), cos_sin.Ptr(), positions.Ptr(), + cos_sin.shape[0], hq, dh, q3.stride[1], rot, na.eps, ra.is_neox_style); + if (hkv > 0) + AttnQkNormRopeRowK<<(t * hkv), kBlock, 0, s>>>( + k3.Ptr(), k_norm.Ptr(), cos_sin.Ptr(), positions.Ptr(), + cos_sin.shape[0], hkv, dh, k3.stride[1], rot, na.eps, ra.is_neox_style); + }; + if (positions.dtype == DType::kI32) { + launch(__hip_bfloat16{}, int32_t{}); + } else { + launch(__hip_bfloat16{}, int64_t{}); + } + Check(hipGetLastError(), "attn_qk_norm_rope launch"); +} + } // namespace vt::rocm diff --git a/src/vt/rocm/rocm_moe_grouped_bf16.hip b/src/vt/rocm/rocm_moe_grouped_bf16.hip new file mode 100644 index 0000000000..7444aff3ee --- /dev/null +++ b/src/vt/rocm/rocm_moe_grouped_bf16.hip @@ -0,0 +1,194 @@ +// Native grouped BF16 MoE (#3094, BACKEND-ROCM-BF16-MOE). +// Logical layout donor: cuda_matmul_nvfp4.cu:943-1004 at 6db4bef906859e864c82523c01107473f7dcca29. +// Native numeric contract: vLLM e126687a9a828d513c01a07cd69f025f27d63280, +// fused_moe.py:517-610, experts/triton_moe.py:388-527, +// activation_kernels.cu:44,165-177, moe_align_sum_kernels.cu:395-459. +// +// Each output has one deterministic reduction. No atomics, reusable scratch, +// allocation, synchronization, or temporary device pointers are required. +// The legacy and native gate/up modes share this implementation. Only the +// explicit native mode rounds gate/up and SiLU through BF16 before multiplication. +#include +#include + +#include +#include +#include +#include + +#include "vt/ops.h" + +namespace vt::rocm { +namespace { +constexpr int kBlock = 256; +constexpr int kReductionTile = 32; + +void Check(hipError_t status, const char* operation) { + if (status != hipSuccess) + throw std::runtime_error(std::string("vt rocm grouped_bf16: ") + operation + ": " + + hipGetErrorString(status)); +} + +// Queue handles, including nullptr, belong to a device. Never infer that device +// from the calling thread's ambient HIP state. +class DeviceScope { + public: + explicit DeviceScope(int device) { + Check(hipGetDevice(&previous_), "get device"); + changed_ = previous_ != device; + if (changed_) Check(hipSetDevice(device), "select queue device"); + } + ~DeviceScope() { if (changed_) (void)hipSetDevice(previous_); } + private: + int previous_ = 0; + bool changed_ = false; +}; + +__device__ float Load(const __hip_bfloat16* data, int64_t index) { + return __bfloat162float(data[index]); +} +__device__ float Load(const float* data, int64_t index) { return data[index]; } +__device__ void Store(float* data, int64_t index, float value) { data[index] = value; } +__device__ void Store(__hip_bfloat16* data, int64_t index, float value) { + data[index] = __float2bfloat16(value); +} +__device__ float Bf16(float value) { + return __bfloat162float(__float2bfloat16(value)); +} + +// Fused=false serves ordinary and weighted down GEMMs. Fused=true serves both +// numeric gate/up modes. FP32 accumulators are required by the pinned dot +// contract; legacy gate/up retains FP32 intermediates for caller compatibility. +template +__global__ void GroupedBf16(Output* output, const __hip_bfloat16* activation, + const int32_t* expert_ids, const int32_t* row_map, + const int64_t* gate_ptrs, const int64_t* up_ptrs, + const float* route_weights, int64_t pairs, int64_t n, + int64_t k) { + const int64_t column = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + if (column >= n) return; + for (int64_t pair = blockIdx.y; pair < pairs; pair += gridDim.y) { + const int64_t expert = expert_ids[pair]; + const int64_t row = row_map == nullptr ? pair : row_map[pair]; + const auto* gate = reinterpret_cast( + static_cast(gate_ptrs[expert])); + const __hip_bfloat16* up = nullptr; + if constexpr (Fused) + up = reinterpret_cast(static_cast(up_ptrs[expert])); + float g = 0.0f; + float u = 0.0f; + for (int64_t start = 0; start < k; start += kReductionTile) { + #pragma unroll + for (int offset = 0; offset < kReductionTile; ++offset) { + const int64_t inner = start + offset; + if (inner < k) { + const float x = Load(activation, row * k + inner); + g += x * Load(gate, inner * n + column); + if constexpr (Fused) u += x * Load(up, inner * n + column); + } + } + } + if constexpr (Fused) { + if constexpr (Native) { + g = Bf16(g); + u = Bf16(u); + } + float silu = g / (1.0f + expf(-g)); + if constexpr (Native) silu = Bf16(silu); + Store(output, pair * n + column, silu * u); + } else { + // fused_moe.py:593-610: multiply in FP32 BEFORE narrowing the down output. + if (route_weights != nullptr) g *= route_weights[pair]; + Store(output, pair * n + column, g); + } + } +} + +template +void Launch(Queue& q, Tensor& output, const Tensor& activation, const Tensor& expert_ids, + const Tensor* row_map, const Tensor& gate_ptrs, const Tensor* up_ptrs, + const Tensor* route_weights) { + const int64_t pairs = output.shape[0], n = output.shape[1], k = activation.shape[1]; + if (pairs == 0 || n == 0) return; + DeviceScope device(q.device.index); + const dim3 grid(static_cast((n + kBlock - 1) / kBlock), + static_cast(pairs < 65535 ? pairs : 65535)); + auto launch = [&](auto* out) { + GroupedBf16, Fused, Native> + <<(q.handle)>>>( + out, activation.Ptr<__hip_bfloat16>(), expert_ids.Ptr(), + row_map == nullptr ? nullptr : row_map->Ptr(), gate_ptrs.Ptr(), + up_ptrs == nullptr ? nullptr : up_ptrs->Ptr(), + route_weights == nullptr ? nullptr : route_weights->Ptr(), pairs, n, k); + }; + if (output.dtype == DType::kBF16) launch(output.Ptr<__hip_bfloat16>()); + else launch(output.Ptr()); + Check(hipGetLastError(), "grouped kernel launch"); +} + +// moe_align_sum_kernels.cu:395-459: values already include router weights. +template +__global__ void CombinePreweighted(Output* output, const __hip_bfloat16* experts, + const Shared* shared, int64_t tokens, int64_t top_k, + int64_t hidden, float routed_scale) { + const int64_t total = tokens * hidden; + const int64_t step = static_cast(gridDim.x) * blockDim.x; + for (int64_t index = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + index < total; index += step) { + const int64_t token = index / hidden, column = index % hidden; + float sum = 0.0f; + for (int64_t slot = 0; slot < top_k; ++slot) + sum += Load(experts, (token * top_k + slot) * hidden + column); + sum *= routed_scale; + if (shared != nullptr) sum += Load(shared, index); + Store(output, index, sum); + } +} +} // namespace + +void MoeGroupedGemmBf16KernelRocm(Queue& q, Tensor& out, const Tensor& act, + const Tensor& expert_ids, const Tensor* row_map, + const Tensor& weight_ptrs) { + Launch(q, out, act, expert_ids, row_map, weight_ptrs, nullptr, nullptr); +} + +void MoeGroupedGemmBf16GateUpSiluKernelRocm(Queue& q, Tensor& out, const Tensor& act, + const Tensor& expert_ids, const Tensor* row_map, + const Tensor& gate_ptrs, const Tensor& up_ptrs) { + Launch(q, out, act, expert_ids, row_map, gate_ptrs, &up_ptrs, nullptr); +} + +void MoeGroupedGemmBf16GateUpSiluNativeKernelRocm( + Queue& q, Tensor& out, const Tensor& act, const Tensor& expert_ids, const Tensor* row_map, + const Tensor& gate_ptrs, const Tensor& up_ptrs) { + Launch(q, out, act, expert_ids, row_map, gate_ptrs, &up_ptrs, nullptr); +} + +void MoeGroupedGemmBf16WeightedKernelRocm(Queue& q, Tensor& out, const Tensor& act, + const Tensor& expert_ids, const Tensor* row_map, + const Tensor& weight_ptrs, const Tensor& weights) { + Launch(q, out, act, expert_ids, row_map, weight_ptrs, nullptr, &weights); +} + +void MoeCombinePreweightedKernelRocm(Queue& q, Tensor& out, const Tensor& experts, + const Tensor* shared, float routed_scale) { + const int64_t tokens = out.shape[0], hidden = out.shape[1], top_k = experts.shape[1]; + const int64_t total = tokens * hidden; + if (total == 0) return; + DeviceScope device(q.device.index); + const auto blocks = (total + kBlock - 1) / kBlock; + const unsigned grid = static_cast(blocks < 65535 ? blocks : 65535); + auto launch = [&](auto* output, auto* shared_data) { + CombinePreweighted<<(q.handle)>>>( + output, experts.Ptr<__hip_bfloat16>(), shared_data, tokens, top_k, hidden, routed_scale); + }; + auto by_shared = [&](auto* output) { + if (shared != nullptr && shared->dtype == DType::kBF16) + launch(output, shared->Ptr<__hip_bfloat16>()); + else launch(output, shared == nullptr ? nullptr : shared->Ptr()); + }; + if (out.dtype == DType::kBF16) by_shared(out.Ptr<__hip_bfloat16>()); + else by_shared(out.Ptr()); + Check(hipGetLastError(), "preweighted combine launch"); +} +} // namespace vt::rocm diff --git a/src/vt/rocm/rocm_ops.hip b/src/vt/rocm/rocm_ops.hip index 7201245ef5..61a93e1096 100644 --- a/src/vt/rocm/rocm_ops.hip +++ b/src/vt/rocm/rocm_ops.hip @@ -85,6 +85,15 @@ void ApplyLogitBiasKernelRocm(Queue& q, Tensor& logits, const Tensor& rows, cons const Tensor& biases); // Companion MoE-path op (same TU): elementwise silu(gate)*up. void MoeSiluMulKernelRocm(Queue& q, Tensor& out, const Tensor& gate, const Tensor& up); +void MoeGroupedGemmBf16KernelRocm(Queue&, Tensor&, const Tensor&, const Tensor&, + const Tensor*, const Tensor&); +void MoeGroupedGemmBf16GateUpSiluKernelRocm(Queue&, Tensor&, const Tensor&, const Tensor&, + const Tensor*, const Tensor&, const Tensor&); +void MoeGroupedGemmBf16GateUpSiluNativeKernelRocm(Queue&, Tensor&, const Tensor&, const Tensor&, + const Tensor*, const Tensor&, const Tensor&); +void MoeGroupedGemmBf16WeightedKernelRocm(Queue&, Tensor&, const Tensor&, const Tensor&, + const Tensor*, const Tensor&, const Tensor&); +void MoeCombinePreweightedKernelRocm(Queue&, Tensor&, const Tensor&, const Tensor*, float); // MoE-path combine/gate ops (rocm_moe_chain.hip): shared-expert gate and the // weighted expert combinations. void SharedExpertGateKernelRocm(Queue& q, Tensor& out, const Tensor& sd, const Tensor& gl); @@ -151,6 +160,13 @@ void Exl3GemmKernelRocm(Queue& q, Tensor& c, const Tensor& a, const Tensor& trel void FusedNormRopeKernelRocm(Queue& q, Tensor& latent_out, Tensor& pe_out, const Tensor& x, const Tensor& w, const Tensor& positions, const Tensor& cache, const RmsNormArgs& norm_args, const RopeArgs& rope_args); +// The gate-free sibling (rocm_gdn_fused.hip): the Qwen3-dense qk-norm + partial +// RoPE preamble. Registered so the fused-chain recipe `kAttnQkNormRope` takes its +// fast realisation on ROCm instead of the Tier-0 composite. +void AttnQkNormRopeKernelRocm(Queue& q, Tensor& q3, Tensor& k3, const Tensor& q_norm, + const Tensor& k_norm, const Tensor& cos_sin, + const Tensor& positions, const RmsNormArgs& na, + const RopeArgs& ra); void AttnQkNormRopeGateKernelRocm(Queue& q, Tensor& q_out, Tensor& k_out, Tensor& gate_out, const Tensor& qgate, const Tensor& kf, const Tensor& q_norm, const Tensor& k_norm, @@ -248,6 +264,21 @@ struct Registrar { static_cast(&ApplyTokenMaskKernelRocm))); RegisterOp(OpId::kMoeSiluMul, DeviceType::kROCM, reinterpret_cast(static_cast(&MoeSiluMulKernelRocm))); + RegisterOp(OpId::kMoeGroupedGemmBf16, DeviceType::kROCM, + reinterpret_cast(static_cast( + &MoeGroupedGemmBf16KernelRocm))); + RegisterOp(OpId::kMoeGroupedGemmBf16GateUpSilu, DeviceType::kROCM, + reinterpret_cast(static_cast( + &MoeGroupedGemmBf16GateUpSiluKernelRocm))); + RegisterOp(OpId::kMoeGroupedGemmBf16GateUpSiluNative, DeviceType::kROCM, + reinterpret_cast(static_cast( + &MoeGroupedGemmBf16GateUpSiluNativeKernelRocm))); + RegisterOp(OpId::kMoeGroupedGemmBf16Weighted, DeviceType::kROCM, + reinterpret_cast(static_cast( + &MoeGroupedGemmBf16WeightedKernelRocm))); + RegisterOp(OpId::kMoeCombinePreweighted, DeviceType::kROCM, + reinterpret_cast(static_cast( + &MoeCombinePreweightedKernelRocm))); RegisterOp(OpId::kSharedExpertGate, DeviceType::kROCM, reinterpret_cast( static_cast(&SharedExpertGateKernelRocm))); @@ -297,6 +328,9 @@ struct Registrar { RegisterOp(OpId::kAttnQkNormRopeGate, DeviceType::kROCM, reinterpret_cast( static_cast(&AttnQkNormRopeGateKernelRocm))); + RegisterOp(OpId::kAttnQkNormRope, DeviceType::kROCM, + reinterpret_cast( + static_cast(&AttnQkNormRopeKernelRocm))); // MLA/DSA campaign W1 (#2715). Before these lines each of these four // resolved to NO native ROCm kernel: the portable CPU reference tier where // it is eligible (which docs/ROCM.md:60-61 disqualifies a performance result diff --git a/src/vt/rocm/rocm_paged_attn.hip b/src/vt/rocm/rocm_paged_attn.hip index 7e5fcce7ce..c29474493d 100644 --- a/src/vt/rocm/rocm_paged_attn.hip +++ b/src/vt/rocm/rocm_paged_attn.hip @@ -13,6 +13,7 @@ #include #include #include +#include #include "vt/ops.h" #include "vt/backend.h" @@ -197,6 +198,22 @@ __device__ inline float FastExp(float x) { return __builtin_amdgcn_exp2f(x * 1.4426950408889634f); // x * log2(e) } +// The softmax probability as the VALUE dtype holds it -- the primary's +// `p = p.to(v.dtype)` (vllm/v1/attention/ops/prefix_prefill.py:471, and +// vllm/v1/attention/ops/chunked_prefill_paged_decode.py:265 for decode), applied +// to the probability that the exponent already produced. Narrowing to f32 is the +// identity, so a site templated on the value dtype needs no dispatch. See the +// tile walk in PagedAttnDecodeOptBf16T for why the reference max the exponent +// uses is part of this rule and not a separate concern. +template +__device__ inline float ProbInValueDtype(float p) { + if constexpr (std::is_same_v) { + return __bfloat162float(__float2bfloat16(p)); + } else { + return p; + } +} + // One block per (query_token, q_head). Online softmax matches CUDA. // Fast-path: num_reqs==1 skips linear request scan (Hermes/single-seq decode). // Vectorized x4 Q·K / V when head_dim % 4 == 0 (Gemma-4 Dh is). @@ -509,31 +526,127 @@ __global__ void PagedAttnDecodeOptBf16T(__hip_bfloat16* out, const __hip_bfloat1 #pragma unroll for (int i = 0; i < kEpl; ++i) o_reg[i] = 0.f; - for (int64_t j = jmin + warp; j <= jmax; j += kDecWarps) { - const int64_t blk = block_table[r * bt_row + (j / block_size) * bt_col]; - const int64_t off = j % block_size; - float k_reg[kEpl]; - LoadRowEplBf16(k_cache, blk * kc_blk + off * kc_pg + g * kc_hd, lane, k_reg); - float dot = 0.f; + // The softmax probability is narrowed to the VALUE dtype and the running sum + // is not. `l_ij = tl.sum(p)` and `p = p.to(v.dtype)` precede + // `acc = tl.dot(p, v, acc=acc)` at vllm/v1/attention/ops/prefix_prefill.py:445, + // :471, :473, so the numerator is a sum of NARROWED probabilities while + // `l_i = l_i * alpha + l_ij` at :475 accumulates the f32 ones. + // + // A bf16 rounding is relative, so that mirror only holds when the probability + // is taken against the same reference max the primary used, and the primary's + // is `m_ij = tl.maximum(m_i, tl.max(qk, axis=1))` at :442 -- a running max + // advanced once per contiguous key tile, not once per key. This kernel + // therefore walks the key range in tiles, reduces each tile's max across the + // CTA, and only then exponentiates, narrows and accumulates. Narrowing against + // a per-warp per-key running max instead was measured on gfx1100 and moved row C + // the wrong way (2918 -> 3011 of 8448 words); the tiled reference max is what + // reproduces the primary. + // + // WHICH TILE, and where the walk starts, are the primary's own choice for the + // arm this row executes, so they are read off the pinned source rather than + // fixed. `chunked_prefill_paged_decode.py:317` dispatches the CALL on the + // batch's max query length, and the decode kernel it launches returns + // immediately for a row with query_len > 1 (`filter_by_query_len=True`, + // :93-97, :503), so the arm is per ROW: + // + // * query_len > 1 -> `prefix_prefill._fwd_kernel` (prefix_prefill.py). It + // walks the cached context in TRITON_BLOCK_SIZE = 32 key tiles anchored at + // key 0 (:965, :1007) inside the context loop (:231-343), masking lanes at + // or past `cur_batch_ctx_len` (:267-295), and then the current chunk in + // BLOCK_N-key tiles anchored at the chunk start (:369). BLOCK_N is 64 when + // the physical block size is a power of two and 32 otherwise (:955-966). + // * query_len == 1 -> the Triton decode kernel, which tiles the whole key + // range from key 0 by TRITON_BLOCK_SIZE = min(block_size, 128), or 32 for a + // non-power-of-two physical block size (chunked_prefill_paged_decode.py: + // 444-445, :147-149, :244). + // + // Tiles are anchored where the primary anchors them, and a key outside the + // causal or sliding-window bound is skipped rather than dropped from the grid: + // the primary masks such a lane to -inf (prefix_prefill.py:288-290) or -10000 + // (chunked_prefill_paged_decode.py:236), neither of which can raise a tile max, + // and `exp(x - m)` underflows to 0 for either. Anchoring the grid at `jmin` + // instead would move every tile boundary on a sliding-window model and change + // which keys share a reference max; whole tiles below `jmin` are skipped so the + // walk still costs O(window) rather than O(sequence). + extern __shared__ float smem_dec[]; + float* tile_sh = smem_dec; // kDecWarps partial maxima; o_sh reuses it below + const bool chunked_arm = query_len > 1; + const bool pow2_block = block_size > 0 && (block_size & (block_size - 1)) == 0; + constexpr int64_t kContextTile = 32; // prefix_prefill TRITON_BLOCK_SIZE + const int64_t chunk_tile = chunked_arm ? (pow2_block ? 64 : 32) + : (pow2_block ? (block_size < 128 ? block_size : 128) : 32); + // One runtime width, at most 128 (min(block_size,128) on the decode arm), so the + // per-warp score registers are sized for the widest tile any arm can ask for. + constexpr int kMaxTileKeys = 128; + constexpr int kPerWarp = (kMaxTileKeys + kDecWarps - 1) / kDecWarps; + // Phase 0 is the cached context (chunked arm only) and phase 1 is the chunk, or + // the whole key range on the decode arm, which the primary always walks from 0. + const int64_t chunk_begin = chunked_arm ? context : 0; + for (int phase = (chunked_arm && context > 0) ? 0 : 1; phase < 2; ++phase) { + const int64_t tile = phase == 0 ? kContextTile : chunk_tile; + const int64_t begin = phase == 0 ? 0 : chunk_begin; + const int64_t end = phase == 0 ? context - 1 : jmax; + for (int64_t base = begin; base <= end; base += tile) { + if (base + tile <= jmin) continue; // the whole tile is out of the window + const int64_t tile_end = base + tile; + const int per_warp = static_cast((tile + kDecWarps - 1) / kDecWarps); + float s_reg[kPerWarp]; + float warp_max = -INFINITY; +#pragma unroll + for (int k = 0; k < kPerWarp; ++k) { + float s = -INFINITY; + if (k < per_warp) { + const int64_t j = base + warp + static_cast(k) * kDecWarps; + if (j < tile_end && j <= end && j >= jmin) { + const int64_t blk = block_table[r * bt_row + (j / block_size) * bt_col]; + const int64_t off = j % block_size; + float k_reg[kEpl]; + LoadRowEplBf16(k_cache, blk * kc_blk + off * kc_pg + g * kc_hd, lane, k_reg); + float dot = 0.f; +#pragma unroll + for (int i = 0; i < kEpl; ++i) dot += q_reg[i] * k_reg[i]; +#pragma unroll + for (int o = 16; o > 0; o >>= 1) dot += __shfl_down(dot, o); + s = Softcap(__shfl(dot, 0) * scale, softcap); + } + } + s_reg[k] = s; + warp_max = fmaxf(warp_max, s); + } + + if (lane == 0) tile_sh[warp] = warp_max; + __syncthreads(); + float tile_max = -INFINITY; #pragma unroll - for (int i = 0; i < kEpl; ++i) dot += q_reg[i] * k_reg[i]; + for (int w = 0; w < kDecWarps; ++w) tile_max = fmaxf(tile_max, tile_sh[w]); + // Every warp has read the tile maximum before the next tile overwrites it. + __syncthreads(); + + const float m_new = fmaxf(m, tile_max); + const float corr = FastExp(m - m_new); #pragma unroll - for (int o = 16; o > 0; o >>= 1) dot += __shfl_down(dot, o); - dot = __shfl(dot, 0); + for (int i = 0; i < kEpl; ++i) o_reg[i] *= corr; + lsum *= corr; + m = m_new; - const float s = Softcap(dot * scale, softcap); - const float m_new = fmaxf(m, s); - const float corr = FastExp(m - m_new); - const float pw = FastExp(s - m_new); - float v_reg[kEpl]; - LoadRowEplBf16(v_cache, blk * vc_blk + off * vc_pg + g * vc_hd, lane, v_reg); #pragma unroll - for (int i = 0; i < kEpl; ++i) o_reg[i] = o_reg[i] * corr + pw * v_reg[i]; - lsum = lsum * corr + pw; - m = m_new; + for (int k = 0; k < kPerWarp; ++k) { + if (k >= per_warp) continue; + const int64_t j = base + warp + static_cast(k) * kDecWarps; + if (!(j < tile_end && j <= end && j >= jmin)) continue; + const float pw = FastExp(s_reg[k] - m_new); + const float pwv = ProbInValueDtype<__hip_bfloat16>(pw); + const int64_t blk = block_table[r * bt_row + (j / block_size) * bt_col]; + const int64_t off = j % block_size; + float v_reg[kEpl]; + LoadRowEplBf16(v_cache, blk * vc_blk + off * vc_pg + g * vc_hd, lane, v_reg); +#pragma unroll + for (int i = 0; i < kEpl; ++i) o_reg[i] += pwv * v_reg[i]; + lsum += pw; + } + } } - extern __shared__ float smem_dec[]; float* o_sh = smem_dec; float* m_sh = o_sh + kDecWarps * d; float* l_sh = m_sh + kDecWarps; diff --git a/src/vt/rocm/rocm_residual_rmsnorm.hip b/src/vt/rocm/rocm_residual_rmsnorm.hip new file mode 100644 index 0000000000..74698a8c44 --- /dev/null +++ b/src/vt/rocm/rocm_residual_rmsnorm.hip @@ -0,0 +1,116 @@ +// Compiled vLLM residual expression on ROCm. Legacy RmsNorm is unchanged. +// Ported from vllm/ir/ops/layernorm.py:44-62 @ +// e126687a9a828d513c01a07cd69f025f27d63280. The executing generated kernels +// ckic6h6:33-52, ctj2x6:33-58, and c5slugd:33-56 fix the order and stores. +// See .agents/specs/rocm-residual-norm.md for the measured compiled lifetime. +#include +#include + +#include +#include +#include + +#include "vt/ops.h" + +namespace vt::rocm { +namespace { +constexpr int kBlock = 256; + +void Check(hipError_t status, const char* operation) { + if (status != hipSuccess) + throw std::runtime_error(std::string("rocm residual_rmsnorm: ") + operation + ": " + + hipGetErrorString(status)); +} + +class QueueDeviceScope { + public: + explicit QueueDeviceScope(const Queue& queue) { + auto stream = static_cast(queue.handle); + bool capturing = false; + if (stream != nullptr) { + hipDevice_t stream_device = -1; + Check(hipStreamGetDevice(stream, &stream_device), "read stream device"); + VT_CHECK(stream_device == queue.device.index, "residual_rmsnorm: stream/queue device mismatch"); + hipStreamCaptureStatus status = hipStreamCaptureStatusNone; + Check(hipStreamIsCapturing(stream, &status), "read capture status"); + capturing = status != hipStreamCaptureStatusNone; + } + Check(hipGetDevice(&previous_), "read current device"); + // During capture the validated stream owns the launch. HIP forbids changing + // the current device during capture. Outside capture, restore it on return. + changed_ = !capturing && previous_ != queue.device.index; + if (changed_) Check(hipSetDevice(queue.device.index), "select queue device"); + } + ~QueueDeviceScope() { if (changed_) (void)hipSetDevice(previous_); } + private: + int previous_ = 0; + bool changed_ = false; +}; + +__device__ float Expression(const __hip_bfloat16* a, const __hip_bfloat16* base, + const __hip_bfloat16* delta, int64_t column) { + const float attention = __bfloat162float(a[column]) + __bfloat162float(base[column]); + return delta == nullptr ? attention : __bfloat162float(delta[column]) + attention; +} + +__global__ void ResidualRmsNormRowKernel( + __hip_bfloat16* out, const __hip_bfloat16* a, const __hip_bfloat16* base, + const __hip_bfloat16* delta, const __hip_bfloat16* weight, + __hip_bfloat16* residual_out, int64_t rows, int64_t width, int64_t out_stride, + int64_t a_stride, int64_t base_stride, int64_t delta_stride, int64_t residual_stride, + float epsilon) { + __shared__ float partial[kBlock]; // FP32 row reduction scratch, never persistent storage + for (int64_t row = blockIdx.x; row < rows; row += gridDim.x) { + const auto* arow = a + row * a_stride; + const auto* brow = base + row * base_stride; + const auto* mrow = delta == nullptr ? nullptr : delta + row * delta_stride; + float acc = 0.0f; + for (int64_t column = threadIdx.x; column < width; column += kBlock) { + const float value = Expression(arow, brow, mrow, column); + acc += value * value; + } + partial[threadIdx.x] = acc; + __syncthreads(); + for (int stride = kBlock / 2; stride > 0; stride /= 2) { + if (static_cast(threadIdx.x) < stride) + partial[threadIdx.x] += partial[threadIdx.x + stride]; + __syncthreads(); + } + const float inverse = 1.0f / sqrtf(partial[0] / static_cast(width) + epsilon); + for (int64_t column = threadIdx.x; column < width; column += kBlock) { + // All variance reads finish before any store. Each lane reloads its own + // complete expression before writing either exact alias, including tails + // and rows wider than one block. No lane writes another lane's inputs. + const float value = Expression(arow, brow, mrow, column); + const float gamma = __bfloat162float(weight[column]); + out[row * out_stride + column] = __float2bfloat16((value * inverse) * gamma); + if (residual_out != nullptr) + residual_out[row * residual_stride + column] = __float2bfloat16(value); + } + __syncthreads(); // protect the reduction scratch when this block serves another row + } +} + +void ResidualRmsNormKernel(Queue& queue, Tensor& out, const Tensor& a, const Tensor& base, + const Tensor* delta, const Tensor& weight, + const ResidualRmsNormArgs& args, Tensor* residual_out) { + QueueDeviceScope device(queue); + const auto rows = static_cast(std::min(a.shape[0], 65535)); + ResidualRmsNormRowKernel<<(queue.handle)>>>( + out.Ptr<__hip_bfloat16>(), a.Ptr<__hip_bfloat16>(), base.Ptr<__hip_bfloat16>(), + delta == nullptr ? nullptr : delta->Ptr<__hip_bfloat16>(), weight.Ptr<__hip_bfloat16>(), + residual_out == nullptr ? nullptr : residual_out->Ptr<__hip_bfloat16>(), + a.shape[0], a.shape[1], out.stride[0], a.stride[0], base.stride[0], + delta == nullptr ? 0 : delta->stride[0], residual_out == nullptr ? 0 : residual_out->stride[0], + args.eps); + Check(hipGetLastError(), "launch"); +} + +struct Registrar { + Registrar() { + RegisterOp(OpId::kResidualRmsNorm, DeviceType::kROCM, + reinterpret_cast(static_cast(&ResidualRmsNormKernel))); + } +} registrar; +} // namespace +} // namespace vt::rocm diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index a3117cc61f..e09e418f8e 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -537,6 +537,7 @@ vllm_cpp_add_test(test_qwen3_break_point vllm/models/test_qwen3_break_point.cpp) vllm_cpp_add_test(test_qwen3_decode_graph_seam vllm/models/test_qwen3_decode_graph_seam.cpp) vllm_cpp_add_test(test_qwen3_moe_decode_graph_seam vllm/models/test_qwen3_moe_decode_graph_seam.cpp) vllm_cpp_add_test(test_moe_async_device_ids vllm/models/test_moe_async_device_ids.cpp) +vllm_cpp_add_test(test_rocm_moe_bf16 vllm/models/test_rocm_moe_bf16.cpp) vllm_cpp_add_test(test_voxtral_decode_graph_seam vllm/models/test_voxtral_decode_graph_seam.cpp) vllm_cpp_add_test(test_deepseek_v2_decode_graph_seam vllm/models/test_deepseek_v2_decode_graph_seam.cpp) vllm_cpp_add_test(test_qwen3_5_decode_graph_seam vllm/models/test_qwen3_5_decode_graph_seam.cpp) @@ -1880,6 +1881,16 @@ vllm_cpp_add_test(test_act_dump vllm/model_executor/test_act_dump.cpp) target_include_directories(test_act_dump PRIVATE ${CMAKE_SOURCE_DIR}/src) +# BACKEND-ROCM-BF16-MOE (#3116): the LM-head output boundary, replayed over the +# primary's own captured head input. The fixture is a directory of raw captures, +# so the two replay cases skip themselves (`doctest::skip`, reported skipped in +# doctest's own summary) and the binary still exits 77 (Skipped) without +# `VT_MOE_HEAD_FIXTURE` rather than reporting a green gate that never ran. The +# fixture-free call-site case in the same binary runs either way, and a failure +# in it keeps its own non-zero exit instead of being folded into the skip. +vllm_cpp_add_test(test_qwen3_moe_lm_head_bf16 + vllm/model_executor/test_qwen3_moe_lm_head_bf16.cpp) + vllm_cpp_add_test(test_moe_fit_resolver vllm/model_executor/test_moe_fit_resolver.cpp) target_include_directories(test_moe_fit_resolver PRIVATE ${CMAKE_SOURCE_DIR}/src) @@ -3175,10 +3186,33 @@ target_include_directories(test_ops_moe_grouped PRIVATE ${CMAKE_SOURCE_DIR}/src) vllm_cpp_add_test(test_moe_router_tie_stability vt/test_moe_router_tie_stability.cpp) target_include_directories(test_moe_router_tie_stability PRIVATE ${CMAKE_SOURCE_DIR}/src) vllm_cpp_add_test(test_ops_moe_grouped_bf16 vt/test_ops_moe_grouped_bf16.cpp) +vllm_cpp_add_test(test_rocm_moe_grouped_bf16 vt/test_rocm_moe_grouped_bf16.cpp) +vllm_cpp_add_test(test_rocm_moe_upstream vt/test_rocm_moe_upstream.cpp) +vllm_cpp_add_test(test_moe_bf16_native_contract vt/test_moe_bf16_native_contract.cpp) +# BACKEND-ROCM-BF16-MOE (#3094): the row's CPU gate on the production token +# gate's reference-set predicate. The predicate is pure host code in +# tests/support/rocm_moe_reference_set.h, so this case needs no device and no +# oracle file, while tests/vllm/models/test_rocm_moe_bf16.cpp applies the same +# predicate to the captured records. +vllm_cpp_add_test(test_rocm_moe_reference_set support/test_rocm_moe_reference_set.cpp) +if(VLLM_CPP_HIP) + foreach(moe_test IN ITEMS test_rocm_moe_grouped_bf16 test_rocm_moe_upstream + test_moe_bf16_native_contract) + target_compile_definitions(${moe_test} PRIVATE __HIP_PLATFORM_AMD__=1) + endforeach() +endif() vllm_cpp_add_test(test_ops_moe_grouped_bf16_gate_up_silu vt/test_ops_moe_grouped_bf16_gate_up_silu.cpp) vllm_cpp_add_test(test_ops_moe_nongated_relu2 vt/test_ops_moe_nongated_relu2.cpp) vllm_cpp_add_test(test_ops_rmsnorm vt/test_ops_rmsnorm.cpp) +vllm_cpp_add_test(test_ops_residual_rmsnorm vt/test_ops_residual_rmsnorm.cpp) +if(VLLM_CPP_HIP) + target_compile_definitions(test_ops_residual_rmsnorm PRIVATE __HIP_PLATFORM_AMD__=1) + target_compile_definitions(test_rocm_moe_bf16 PRIVATE __HIP_PLATFORM_AMD__=1) + # The head-boundary replay names the ROCm device itself (hipGetDevice / + # hipSetDevice around each queue), exactly as the residual siblings do. + target_compile_definitions(test_qwen3_moe_lm_head_bf16 PRIVATE __HIP_PLATFORM_AMD__=1) +endif() vllm_cpp_add_test(test_ops_rmsnorm_weight_dtype vt/test_ops_rmsnorm_weight_dtype.cpp) vllm_cpp_add_test(test_ops_rmsnorm_quant_fp8_weight_dtype vt/test_ops_rmsnorm_quant_fp8_weight_dtype.cpp) diff --git a/tests/support/residual_norm_fixture.h b/tests/support/residual_norm_fixture.h new file mode 100644 index 0000000000..d13121907b --- /dev/null +++ b/tests/support/residual_norm_fixture.h @@ -0,0 +1,104 @@ +// Captured BF16 row zero of L33/C2/R0, step zero, layer zero (#3103). +// Primary: vLLM e126687a9a828d513c01a07cd69f025f27d63280, default compilation. +// Provenance: .agents/specs/rocm-residual-norm.md, oracle-diagnostic-2. +// Each array contains the original 128 little-endian BF16 words, without recomputation. +#pragma once +#include +#include +namespace residual_norm_fixture { +// Row SHA256: 8f7b3d32577bf9350dbc2f3764d0eb97eed37fe4f285ccd4acd316288a37e7fd +inline constexpr std::array kAttention = { + 0xbc9a, 0x3db1, 0xbd34, 0x3dfb, 0x3c15, 0x3e3a, 0xbd87, 0xb992, + 0x3e37, 0x3dc8, 0xbdc7, 0x3dd7, 0xbd83, 0xbc03, 0x3e26, 0xbe0c, + 0xbd76, 0xbd95, 0x3e28, 0xbab0, 0xbd3f, 0xbcdb, 0x3ccc, 0xbce1, + 0x3d07, 0xbcda, 0xbd32, 0x3d06, 0xbb1b, 0xbcf1, 0xbda0, 0xbd99, + 0xbd9a, 0x3dde, 0x3e01, 0x3d95, 0xbd01, 0xb911, 0xbcfb, 0xbcea, + 0x3bfa, 0x3cb9, 0xbe30, 0x3da2, 0xbe11, 0xbdb8, 0xbca8, 0xbda0, + 0x3dcc, 0x3d23, 0x3d8c, 0x3d51, 0xbd82, 0x3d39, 0xbd92, 0x3cc4, + 0xbd73, 0x3c6c, 0xbd52, 0xbd2e, 0xbde2, 0xbbf0, 0xbc97, 0x3e0d, + 0x3c01, 0x3d15, 0x3d60, 0x3c4c, 0xbd48, 0xbd5e, 0xbc9c, 0xbd1c, + 0xbd23, 0xbe0c, 0xbc9a, 0xbd9b, 0x3b12, 0x3d90, 0xbd2d, 0x3c2e, + 0x3d3a, 0xbd46, 0x3dd2, 0xbd0d, 0x3e43, 0xbde9, 0xbd8e, 0x3d0b, + 0xbb40, 0x3d64, 0x3c91, 0xbd96, 0xbc9e, 0xbd8d, 0xbd1c, 0x3e1e, + 0x3df5, 0xbb57, 0x3d18, 0x3d90, 0xbd50, 0x3e08, 0x3dbc, 0xbc1a, + 0xbda1, 0xbca5, 0xbe09, 0xbe09, 0x3dc9, 0x3cc7, 0x3d0c, 0x3d48, + 0x3c86, 0xbd22, 0xbd93, 0x3d8e, 0xbdda, 0xbda4, 0xbd3a, 0xbc4d, + 0xb979, 0xbd8e, 0x3b88, 0x3d23, 0x3dce, 0x3dd2, 0x3da3, 0xbdbb, +}; +// Row SHA256: e77ba027dd22f83a69aba94cb398c72dddd973c5e52071cd2684d385b2f27329 +inline constexpr std::array kResidual = { + 0xbd18, 0x3d7f, 0xbd88, 0xbd53, 0x3d68, 0xbd58, 0xbd3c, 0x399a, + 0x3d4d, 0xbbb7, 0xbcc9, 0xbd5a, 0xba17, 0xbce8, 0xbcf1, 0x3d80, + 0xba72, 0xbd21, 0x3d86, 0x3ceb, 0xbc1f, 0x3c31, 0xbd35, 0x3b0d, + 0xbd98, 0x3c2f, 0xbd66, 0xbd1f, 0x3d01, 0xbd55, 0xbda1, 0xbd9d, + 0xbc33, 0xbce2, 0x3c8c, 0xbc64, 0x3d76, 0x3d05, 0xbd40, 0xbc83, + 0x3c88, 0x3ca6, 0xbd37, 0x3d05, 0xbd7d, 0xbd3c, 0x3c57, 0x3d98, + 0xbd76, 0xbd86, 0xbd73, 0xbd7e, 0xbd89, 0xbd0b, 0xbbed, 0x3cb8, + 0x3da3, 0xbd8d, 0xbd9a, 0x3d25, 0xbd94, 0x3d78, 0x3c84, 0x3c31, + 0x3ace, 0x3c9f, 0xbd7f, 0x3d6e, 0xbda2, 0x3d93, 0x3d04, 0xbd9e, + 0xbd74, 0xbce8, 0x3b45, 0xbd65, 0x3d6f, 0x3d91, 0xbcb6, 0x3ca6, + 0xbb68, 0x3d80, 0xbc3b, 0xbc73, 0x3d00, 0xbd93, 0x3cfb, 0x3d8c, + 0xbc9e, 0x3c98, 0x3cbb, 0xbba8, 0xbd92, 0xbcfa, 0xbc06, 0xbd3c, + 0x3d8f, 0x3c7d, 0x3d2f, 0xbbbb, 0xbcc2, 0xba56, 0x3d1f, 0x3cfd, + 0x3d00, 0x3d48, 0x3d94, 0xbc82, 0xbd05, 0x3d5d, 0xbd02, 0x3d9d, + 0x3d73, 0x3d95, 0xbd53, 0xbd97, 0xbc86, 0x3be5, 0xbd00, 0x3d77, + 0x3ca2, 0xbd51, 0xbd6d, 0xbd87, 0x3b7d, 0xbd6a, 0x3d33, 0xbb6e, +}; +// Row SHA256: 56a5eae577045d8191fcb1d161d238c937573a2e8e83ef1a7474a4cc28b7cffd +inline constexpr std::array kPostNorm = { + 0x3c65, 0xbe81, 0xbf0d, 0xbe3e, 0xbdd1, 0x3e7a, 0xbca2, 0x3898, + 0xbee4, 0xbe68, 0x3deb, 0xbe82, 0xbe0e, 0x3de6, 0xbe93, 0x3e59, + 0x3e42, 0x3cef, 0x3dd8, 0x3aa5, 0x3e32, 0xbc9e, 0x3d34, 0x3da0, + 0x3c96, 0x3c14, 0xbeea, 0xbc2c, 0x3deb, 0x3eb7, 0x3e6e, 0xbe7f, + 0x3e60, 0x3ea2, 0xbdd3, 0x3df0, 0xbdbe, 0x3e2a, 0x3d5c, 0xbdff, + 0xbd6f, 0x3e57, 0xbe6d, 0xbf12, 0xbe65, 0x3eac, 0xbcff, 0xbbd8, + 0xbde4, 0xbd0e, 0x3c3d, 0xbd08, 0xbed6, 0x3d36, 0xbd89, 0xbe09, + 0x3d0c, 0x3d27, 0xbeeb, 0x3866, 0x3eef, 0x3e1c, 0x3b87, 0xbeca, + 0x3c04, 0xbdbc, 0xbce5, 0xbe8b, 0xbeb6, 0x3da7, 0xbd5e, 0xbecd, + 0x3e04, 0x3e6a, 0x3d89, 0x3cbb, 0x3e9a, 0xbf23, 0xbdef, 0xbde4, + 0x3d88, 0x3d7d, 0xbea4, 0xbe14, 0xbf73, 0x3f1b, 0xbde4, 0x3e12, + 0x3daf, 0x3ec8, 0xbca8, 0xbe31, 0xbe22, 0xbef6, 0x3e77, 0xbecb, + 0x3f0a, 0xbcce, 0x3d16, 0x3e12, 0x3e6b, 0x3f02, 0xbee6, 0x3dc5, + 0xbd9f, 0xbe15, 0x3e83, 0xbd89, 0xbe55, 0xbe51, 0xbc14, 0xbe4c, + 0xbe30, 0xbc8c, 0x3e44, 0x3c46, 0xbf08, 0x3e93, 0x3d76, 0xbc5c, + 0x3d20, 0xbc30, 0x3e8b, 0x3d51, 0xbe0c, 0x3def, 0xbefb, 0xbca9, +}; +// Row SHA256: c90875c40310a0a9d9ce09f16bbf0a67eb8539859e2cc273e5dddae60a0db362 +inline constexpr std::array kFirstNorm = { + 0x3d14, 0xbdb4, 0x3f31, 0xbec2, 0xbea5, 0x3eb1, 0xbe61, 0x3a63, + 0xbf05, 0xbc15, 0xbe0d, 0x3ef9, 0xbaa3, 0xbe0b, 0x3e43, 0x3f17, + 0x3bd2, 0xbeb4, 0xbd4c, 0xbe28, 0xbd0e, 0x3da2, 0x3e0a, 0xbc54, + 0x3f0d, 0x3dc9, 0xbe86, 0x3eb0, 0x3e59, 0x3efd, 0x3f1c, 0x3ed5, + 0xbd98, 0x3e55, 0x3bed, 0x3d69, 0xbef6, 0xbddd, 0x3e56, 0x3dcc, + 0xbdba, 0x3e25, 0xbe14, 0xbea0, 0x3eea, 0xbe9a, 0xbdd6, 0x3e3c, + 0x3ed6, 0x3e9a, 0x3e6b, 0xbe28, 0x3d92, 0xbe97, 0xbc83, 0xbe0e, + 0xbf0d, 0x3e95, 0x3f0e, 0xbe89, 0xbf32, 0xbe92, 0x3d81, 0xbde9, + 0x3bd3, 0x3e17, 0x3ba4, 0x3dd7, 0x3f1b, 0xbf1d, 0xbe9e, 0x3e8e, + 0xbeec, 0xbe33, 0x3cfe, 0x3eea, 0xbf02, 0x3f40, 0x3e2e, 0x3e46, + 0x3cc0, 0xbe66, 0x3bbc, 0xbc16, 0xbe41, 0xba08, 0xbd51, 0xbec4, + 0x3e4d, 0xbe20, 0xbdbd, 0xbd3a, 0xbf24, 0x3e59, 0x3da5, 0xbc0e, + 0x3ec5, 0x3d90, 0x3e91, 0xbd0c, 0xbe82, 0x3bbe, 0x3e7e, 0xbe38, + 0xbea6, 0x3eba, 0xbf15, 0x3df6, 0xbe9b, 0xbdfe, 0x3dba, 0xbd62, + 0xbbb6, 0xbf37, 0x3d45, 0xbea4, 0xbe15, 0xbd0a, 0xbe8a, 0xbe8c, + 0x3d9d, 0xbef1, 0x3d07, 0x3f36, 0x3c8e, 0x3f0c, 0x3ebd, 0x3c85, +}; +// Row SHA256: b7ef025fce17ff455ed4fd6640212b989ae804d0de2ec02358fd5d86129ad4de +inline constexpr std::array kGamma = { + 0xbcc0, 0xbe23, 0x3eef, 0xbe7b, 0xbe15, 0x3e36, 0x3c88, 0x3ee4, + 0xbe3b, 0xbe6c, 0xbdb5, 0xbeeb, 0x3e4e, 0xbe94, 0xbe50, 0xbe89, + 0xbe95, 0xbcc8, 0x3d30, 0x3b8d, 0xbe94, 0x3de9, 0xbe5a, 0xbe94, + 0xbd2a, 0xbd5a, 0x3edc, 0x3e25, 0x3ebd, 0xbed3, 0xbe0e, 0x3e1e, + 0xbe74, 0x3ebc, 0xbd8a, 0x3e3f, 0xbe9c, 0x3ef6, 0xbd85, 0x3e86, + 0xbe67, 0x3eeb, 0x3dcd, 0xbef5, 0x3dd3, 0xbe6d, 0x3eca, 0x3e22, + 0xbe87, 0x3e02, 0x3df5, 0x3e91, 0x3e9a, 0x3ebe, 0x3da3, 0xbe8a, + 0x3e22, 0xbd90, 0x3eae, 0xbb19, 0xbe75, 0x3e89, 0xbe2b, 0xbe7f, + 0x3da4, 0xbe1e, 0x3eb1, 0xbeb9, 0x3e85, 0x3ede, 0xbec5, 0x3ea7, + 0xbdf9, 0xbe05, 0xbecb, 0xbc85, 0x3eee, 0xbed9, 0x3e2e, 0xbead, + 0x3e18, 0x3ed1, 0xbea9, 0x3e8d, 0xbecd, 0xbe9d, 0x3e8a, 0x3e06, + 0xbeb8, 0x3efc, 0xbd42, 0x3e54, 0x3e28, 0x3ee8, 0xbefa, 0xbeaf, + 0x3e88, 0xbe46, 0x3d30, 0x3e53, 0xbe94, 0x3eb8, 0xbea5, 0x3ed7, + 0x3e1d, 0xbef4, 0xbec8, 0x3d2b, 0xbe98, 0xbe7a, 0xbeb1, 0xbe18, + 0xbe5a, 0xbd46, 0xbe15, 0xbe84, 0x3ed0, 0xbebc, 0xbd96, 0xbcd8, + 0x3e40, 0x3c09, 0xbef2, 0xbe3b, 0xbdfb, 0x3e77, 0xbebf, 0x3ca7, +}; +} // namespace residual_norm_fixture diff --git a/tests/support/residual_norm_later_fixture.h b/tests/support/residual_norm_later_fixture.h new file mode 100644 index 0000000000..20d12009fc --- /dev/null +++ b/tests/support/residual_norm_later_fixture.h @@ -0,0 +1,239 @@ +// Original compiled vLLM e126687a9a launches, observed without changing arguments. +// #3103 operator v3 capture: L33/C2/R0, first row of each actual tensor. +// Observation/control: all 18 cohorts, 216 tokens and every logprob match. +#pragma once +#include +#include +namespace residual_norm_fixture { +// L33-C2-R0-residual-next-0-input-attention.bin +// full payload SHA256 3147e04760561c46f35e6f14de4693c6455c026541c3fe87945253ba2d09415e +// row SHA256 8f7b3d32577bf9350dbc2f3764d0eb97eed37fe4f285ccd4acd316288a37e7fd +inline constexpr std::array kNextAttention = { + 0xbc9a, 0x3db1, 0xbd34, 0x3dfb, 0x3c15, 0x3e3a, 0xbd87, 0xb992, + 0x3e37, 0x3dc8, 0xbdc7, 0x3dd7, 0xbd83, 0xbc03, 0x3e26, 0xbe0c, + 0xbd76, 0xbd95, 0x3e28, 0xbab0, 0xbd3f, 0xbcdb, 0x3ccc, 0xbce1, + 0x3d07, 0xbcda, 0xbd32, 0x3d06, 0xbb1b, 0xbcf1, 0xbda0, 0xbd99, + 0xbd9a, 0x3dde, 0x3e01, 0x3d95, 0xbd01, 0xb911, 0xbcfb, 0xbcea, + 0x3bfa, 0x3cb9, 0xbe30, 0x3da2, 0xbe11, 0xbdb8, 0xbca8, 0xbda0, + 0x3dcc, 0x3d23, 0x3d8c, 0x3d51, 0xbd82, 0x3d39, 0xbd92, 0x3cc4, + 0xbd73, 0x3c6c, 0xbd52, 0xbd2e, 0xbde2, 0xbbf0, 0xbc97, 0x3e0d, + 0x3c01, 0x3d15, 0x3d60, 0x3c4c, 0xbd48, 0xbd5e, 0xbc9c, 0xbd1c, + 0xbd23, 0xbe0c, 0xbc9a, 0xbd9b, 0x3b12, 0x3d90, 0xbd2d, 0x3c2e, + 0x3d3a, 0xbd46, 0x3dd2, 0xbd0d, 0x3e43, 0xbde9, 0xbd8e, 0x3d0b, + 0xbb40, 0x3d64, 0x3c91, 0xbd96, 0xbc9e, 0xbd8d, 0xbd1c, 0x3e1e, + 0x3df5, 0xbb57, 0x3d18, 0x3d90, 0xbd50, 0x3e08, 0x3dbc, 0xbc1a, + 0xbda1, 0xbca5, 0xbe09, 0xbe09, 0x3dc9, 0x3cc7, 0x3d0c, 0x3d48, + 0x3c86, 0xbd22, 0xbd93, 0x3d8e, 0xbdda, 0xbda4, 0xbd3a, 0xbc4d, + 0xb979, 0xbd8e, 0x3b88, 0x3d23, 0x3dce, 0x3dd2, 0x3da3, 0xbdbb, +}; +// L33-C2-R0-residual-next-0-input-residual.bin +// full payload SHA256 f5f65bd430c798549eafc5cdd4d5ce9a99d60ea5ebeada4a59222ab5fb65d728 +// row SHA256 e77ba027dd22f83a69aba94cb398c72dddd973c5e52071cd2684d385b2f27329 +inline constexpr std::array kNextBase = { + 0xbd18, 0x3d7f, 0xbd88, 0xbd53, 0x3d68, 0xbd58, 0xbd3c, 0x399a, + 0x3d4d, 0xbbb7, 0xbcc9, 0xbd5a, 0xba17, 0xbce8, 0xbcf1, 0x3d80, + 0xba72, 0xbd21, 0x3d86, 0x3ceb, 0xbc1f, 0x3c31, 0xbd35, 0x3b0d, + 0xbd98, 0x3c2f, 0xbd66, 0xbd1f, 0x3d01, 0xbd55, 0xbda1, 0xbd9d, + 0xbc33, 0xbce2, 0x3c8c, 0xbc64, 0x3d76, 0x3d05, 0xbd40, 0xbc83, + 0x3c88, 0x3ca6, 0xbd37, 0x3d05, 0xbd7d, 0xbd3c, 0x3c57, 0x3d98, + 0xbd76, 0xbd86, 0xbd73, 0xbd7e, 0xbd89, 0xbd0b, 0xbbed, 0x3cb8, + 0x3da3, 0xbd8d, 0xbd9a, 0x3d25, 0xbd94, 0x3d78, 0x3c84, 0x3c31, + 0x3ace, 0x3c9f, 0xbd7f, 0x3d6e, 0xbda2, 0x3d93, 0x3d04, 0xbd9e, + 0xbd74, 0xbce8, 0x3b45, 0xbd65, 0x3d6f, 0x3d91, 0xbcb6, 0x3ca6, + 0xbb68, 0x3d80, 0xbc3b, 0xbc73, 0x3d00, 0xbd93, 0x3cfb, 0x3d8c, + 0xbc9e, 0x3c98, 0x3cbb, 0xbba8, 0xbd92, 0xbcfa, 0xbc06, 0xbd3c, + 0x3d8f, 0x3c7d, 0x3d2f, 0xbbbb, 0xbcc2, 0xba56, 0x3d1f, 0x3cfd, + 0x3d00, 0x3d48, 0x3d94, 0xbc82, 0xbd05, 0x3d5d, 0xbd02, 0x3d9d, + 0x3d73, 0x3d95, 0xbd53, 0xbd97, 0xbc86, 0x3be5, 0xbd00, 0x3d77, + 0x3ca2, 0xbd51, 0xbd6d, 0xbd87, 0x3b7d, 0xbd6a, 0x3d33, 0xbb6e, +}; +// L33-C2-R0-residual-next-0-input-moe.bin +// full payload SHA256 887610ce89a0ab15a92e0fdf321921b7f01372433553c07154ffe126add1dc4c +// row SHA256 91421596748f5cdec3720c2330126e71ae14e650837a0fd7974d367a56754c6a +inline constexpr std::array kNextMoe = { + 0xbb5d, 0xbb06, 0xb948, 0x3b5e, 0x3bcc, 0xbb97, 0x3b80, 0x3a5a, + 0xbb99, 0x3b41, 0xbac6, 0x3a91, 0xba88, 0x3c0a, 0x3c04, 0xbb41, + 0xbaf8, 0xbb54, 0x3b9d, 0x3ad2, 0xbb5e, 0xb960, 0xbabc, 0x3b2a, + 0xbbc8, 0x3b9c, 0xba98, 0xba34, 0x3bbd, 0x3ad1, 0x3b96, 0xbb5d, + 0xbb20, 0x3b20, 0x3ac7, 0xba21, 0x3b66, 0xbb08, 0x3b08, 0x3abb, + 0x3aea, 0xba86, 0x3bd1, 0x3ac9, 0xbb0b, 0xbbe8, 0xbb8e, 0x3b45, + 0xbb62, 0x3b06, 0x3a50, 0xba2e, 0x3ad3, 0xbb5b, 0x3b4b, 0xbb2a, + 0x3b63, 0xb8a0, 0xb8c0, 0x3ab4, 0x3a92, 0xbae4, 0xba7a, 0xbb06, + 0xbab2, 0x3ad9, 0xbb43, 0xba70, 0x3b2e, 0x3b3e, 0xbb0d, 0x3b2a, + 0x3baf, 0xba18, 0x39a3, 0xbadd, 0x3a57, 0xbb2a, 0xbb92, 0x3a02, + 0xb700, 0x3b1b, 0x3976, 0x3b1f, 0xbbf0, 0xb995, 0xbb3f, 0x3ab8, + 0x3bb6, 0xbafe, 0x3aea, 0x3b3c, 0xbb58, 0xba65, 0x3b56, 0x3b54, + 0xba40, 0x39fa, 0x3aca, 0x3abf, 0xb9d8, 0x3a9e, 0x3840, 0x3adc, + 0xba5c, 0xbb1f, 0x3b75, 0xbb54, 0xbb63, 0xbb0d, 0x3b39, 0xbb48, + 0xbb8b, 0xbb60, 0xbbd1, 0xbb82, 0x3a90, 0xbbca, 0x3b60, 0x3a01, + 0x3b10, 0xb9eb, 0x3a34, 0x3944, 0x3850, 0xbbfb, 0x3bc7, 0x3a64, +}; +// L33-C2-R0-residual-next-0-input-weight.bin +// full payload SHA256 5d6e1fb08c4a60be240df3955110acf33bc1b8b18f4ff371eb761a285e8086d4 +// row SHA256 5d6e1fb08c4a60be240df3955110acf33bc1b8b18f4ff371eb761a285e8086d4 +inline constexpr std::array kNextGamma = { + 0x3ed4, 0xbce6, 0x3efc, 0x3ce1, 0x3e67, 0xbe30, 0xbdb7, 0xbec9, + 0xbe39, 0xbeb7, 0xbd8c, 0xbe1b, 0x3e61, 0x3ef7, 0xbe8c, 0xbeb6, + 0x3eb0, 0x3e1f, 0xbec9, 0x3ed2, 0xbec6, 0xbed5, 0x3eba, 0xbef8, + 0x3c2d, 0x3ec4, 0xbe9c, 0xbed6, 0x3ef0, 0x3e50, 0x3e8f, 0x3e97, + 0xbe8f, 0xbe98, 0x3e48, 0xbe94, 0x3ca4, 0x3e3c, 0xbe6a, 0x3e85, + 0x3eb4, 0x3e7e, 0x3dba, 0xbdee, 0x3d31, 0x3d49, 0xbea8, 0x3d24, + 0xbdbe, 0xbd35, 0x3dcc, 0x3efb, 0x3e99, 0xbe92, 0xbec2, 0x3dd7, + 0xbe81, 0x3ee3, 0xbd13, 0x3df3, 0x3d16, 0x3c75, 0xbdce, 0xbee5, + 0xbe03, 0xbea0, 0x3ea1, 0xbe83, 0xbee5, 0x3d0c, 0xbe40, 0x3e56, + 0x3e82, 0xbef8, 0x3dc3, 0x3e7e, 0xbe33, 0x3e88, 0xbeb0, 0x3d5c, + 0x3e85, 0xbda7, 0xbee9, 0x3e6a, 0x3e4a, 0xbedd, 0x3e64, 0xbd7c, + 0x3df5, 0x3e28, 0xbe7a, 0x3e15, 0x3e06, 0xbd5e, 0xbddb, 0xbebc, + 0x3ee7, 0xbe8a, 0xbe6a, 0xbe73, 0x3cbf, 0x3e09, 0x3e8a, 0x3e38, + 0xbe08, 0x3d59, 0x3e8c, 0xbdbc, 0x3da6, 0xbb4c, 0x3ec5, 0xbee4, + 0xbec1, 0x3d9c, 0xbe43, 0xbee7, 0x3e86, 0xbe82, 0x3ef1, 0xbeb2, + 0x3eb1, 0xbebf, 0xbecf, 0x3e2d, 0x3e8c, 0xbd11, 0x3e93, 0xbedc, +}; +// L33-C2-R0-residual-next-0-output-norm.bin +// full payload SHA256 d3395d5d0a52e4986b6fa3821d1f2d2686b621aaedb7da2fa6c90f39dcad3131 +// row SHA256 9ada0e2de8083ed026d0a6788a34319710954fcb24fadb5beac66537fa3f8e07 +inline constexpr std::array kNextNorm = { + 0xbe86, 0xbd34, 0xbf15, 0x3cb3, 0x3e31, 0xbe69, 0x3dd3, 0xbb69, + 0xbedd, 0xbeba, 0x3db8, 0xbdaf, 0xbe1d, 0xbe13, 0xbed2, 0x3e96, + 0xbe6c, 0xbe44, 0xbf7b, 0x3e02, 0x3e7c, 0x3d93, 0xbda5, 0x3df1, + 0xbbaf, 0xbd3b, 0x3ea8, 0x3cf8, 0x3e33, 0xbe31, 0xbee8, 0xbef9, + 0x3e87, 0xbe87, 0x3e9a, 0xbe38, 0x3be1, 0x3d73, 0x3e3c, 0xbdf5, + 0x3dc8, 0x3de3, 0xbe51, 0xbe10, 0xbdc2, 0xbd99, 0x3d28, 0xb9ca, + 0xbd12, 0x3c36, 0x3c2b, 0xbd7a, 0xbed2, 0xbcc5, 0x3e9c, 0x3d49, + 0xbd83, 0xbe84, 0x3d47, 0xba89, 0xbd91, 0x3c07, 0x3b66, 0xbf33, + 0xbc35, 0xbe44, 0xbd11, 0xbe43, 0x3f19, 0x3bf5, 0xbcb5, 0xbe81, + 0xbe82, 0x3f5b, 0xbc81, 0xbeb5, 0xbdeb, 0x3ec9, 0x3e81, 0x3c93, + 0x3dee, 0xbc6c, 0xbee3, 0xbdea, 0x3ee7, 0x3f5b, 0xbdca, 0xbd8b, + 0xbcae, 0x3e02, 0xbde2, 0xbdf0, 0xbe06, 0x3d6d, 0x3d49, 0xbee0, + 0x3f69, 0xbd15, 0xbe4b, 0xbe2b, 0xbc99, 0x3e43, 0x3ec0, 0x3d36, + 0x3d8c, 0x3c73, 0xbe2d, 0x3e19, 0x3d5c, 0xbb26, 0x3cb1, 0xbf15, + 0xbe93, 0x3cc6, 0x3e87, 0x3d25, 0xbeae, 0x3e5c, 0xbebc, 0xbe37, + 0x3da4, 0x3ef6, 0x3e6a, 0xbd40, 0x3e9c, 0xbc6a, 0x3ecb, 0x3edd, +}; +// L33-C2-R0-residual-next-0-output-residual_next.bin +// full payload SHA256 06637392bfd4416ea7d0e751486139f355df9063eeab9d75563b2a88da9dc3ac +// row SHA256 394c6880b525bc304d4c379dc37d43c7508c150df6b53ade3c5ac073a5e50c3c +inline constexpr std::array kNextResidual = { + 0xbd73, 0x3e16, 0xbde2, 0x3d98, 0x3d93, 0x3dff, 0xbddd, 0x3a5e, + 0x3e65, 0x3dc3, 0xbdfc, 0x3d59, 0xbd86, 0xbce4, 0x3e10, 0xbd9e, + 0xbd81, 0xbdec, 0x3e70, 0x3ced, 0xbd75, 0xbc84, 0xbcaa, 0xbcba, + 0xbd42, 0xbc37, 0xbdce, 0xbbde, 0x3d0f, 0xbda3, 0xbe1c, 0xbe1e, + 0xbdb5, 0x3daa, 0x3e14, 0x3d6e, 0x3d03, 0x3cf8, 0xbd9a, 0xbd31, + 0x3cd5, 0x3d2b, 0xbe57, 0x3de8, 0xbe52, 0xbe12, 0xbc40, 0xba6c, + 0x3d14, 0xbcc1, 0x3c21, 0xbc3f, 0xbe04, 0x3c01, 0xbd9a, 0x3d33, + 0x3cc2, 0xbd5f, 0xbe02, 0xba58, 0xbe3a, 0x3d53, 0xbb56, 0x3e16, + 0x3c04, 0x3d6b, 0xbc2d, 0x3d8f, 0xbe00, 0x3ca8, 0x3c35, 0xbde7, + 0xbdc1, 0xbe2a, 0xbc7e, 0xbe08, 0x3d7b, 0x3e0e, 0xbd8d, 0x3d01, + 0x3d2b, 0x3c87, 0x3dbb, 0xbd40, 0x3e5c, 0xbe3e, 0xbd2a, 0x3dd4, + 0xbc88, 0x3d94, 0x3d2d, 0xbd9b, 0xbdc0, 0xbdcd, 0xbd30, 0x3de5, + 0x3e41, 0x3c4f, 0x3da7, 0x3d87, 0xbd99, 0x3e08, 0x3e06, 0x3cbe, + 0xbd45, 0x3cd7, 0xbd6d, 0xbe1d, 0x3d7f, 0x3d9c, 0x3bac, 0x3dfb, + 0x3d92, 0x3cf4, 0xbe05, 0xbc09, 0xbdf9, 0xbda2, 0xbd96, 0x3d46, + 0x3cb2, 0xbdf7, 0xbd59, 0xbcd4, 0x3dd6, 0x3d1b, 0x3e04, 0xbdc1, +}; +// L33-C2-R0-residual-final-0-input-attention.bin +// full payload SHA256 ec17e4e9193ce79219cf03f50cabbcce4a5bfbfc98d69918d2a60172736cfce0 +// row SHA256 1bb4f398badc7d07429624005b8834c95e57dadf49989b87a0555df383f22886 +inline constexpr std::array kFinalAttention = { + 0xbd9b, 0x3d66, 0x3d91, 0xbd86, 0xbd24, 0x3c7e, 0x3d25, 0xbd11, + 0xbda7, 0xbc98, 0x3d19, 0x3d26, 0xbdb8, 0xbc3b, 0xbe1d, 0x3e4e, + 0xbe33, 0xbc3f, 0xbd1b, 0x3cbb, 0xbd81, 0xbdbd, 0xbcfc, 0x3caf, + 0xbd17, 0x3d5d, 0xbd0f, 0x3e22, 0xbb80, 0x3c97, 0xbd81, 0x3d43, + 0x3de8, 0xbded, 0x3d8c, 0x3dae, 0xbb61, 0xbcd2, 0xbb67, 0x3c99, + 0xb90f, 0xbd37, 0xbd8e, 0x3d31, 0xbcff, 0x3d36, 0x3c9d, 0x3dae, + 0xbd8b, 0xbd97, 0xbd28, 0xbd4e, 0x3d43, 0x3d35, 0x3da5, 0xbd7b, + 0xbd51, 0xbc91, 0x3c77, 0xbc38, 0xbc3e, 0xbc46, 0x3da2, 0x3d0a, + 0xbdcf, 0xbcab, 0xbdb6, 0x3db1, 0xbdbc, 0xbcab, 0x3d46, 0x3b66, + 0xbd55, 0x3c5c, 0x3dd2, 0xbc7b, 0xbb3c, 0x3de0, 0xbbd5, 0x3cd0, + 0x3d1c, 0xbcd2, 0x3cc7, 0xbb4e, 0xbcc2, 0x3d16, 0xbd2e, 0xbdd0, + 0xbdaa, 0x3cb1, 0xbc1d, 0x3d89, 0x3cd9, 0x3df6, 0xbda3, 0x3daf, + 0x3e00, 0x386b, 0xbd32, 0x3d04, 0x3d01, 0x3d81, 0xbc50, 0x3ce9, + 0x3d86, 0xbd11, 0xbdfa, 0xbd4e, 0xbd3b, 0xbd4d, 0xbd5e, 0x3d90, + 0x3be3, 0x3b4a, 0xbc5d, 0x3c89, 0xbd22, 0xbd89, 0xbd96, 0xbd86, + 0x3d97, 0x3d00, 0x3c0d, 0xbdde, 0xbc26, 0x3c39, 0xbb8e, 0x3bed, +}; +// L33-C2-R0-residual-final-0-input-residual.bin +// full payload SHA256 06637392bfd4416ea7d0e751486139f355df9063eeab9d75563b2a88da9dc3ac +// row SHA256 394c6880b525bc304d4c379dc37d43c7508c150df6b53ade3c5ac073a5e50c3c +inline constexpr std::array kFinalBase = { + 0xbd73, 0x3e16, 0xbde2, 0x3d98, 0x3d93, 0x3dff, 0xbddd, 0x3a5e, + 0x3e65, 0x3dc3, 0xbdfc, 0x3d59, 0xbd86, 0xbce4, 0x3e10, 0xbd9e, + 0xbd81, 0xbdec, 0x3e70, 0x3ced, 0xbd75, 0xbc84, 0xbcaa, 0xbcba, + 0xbd42, 0xbc37, 0xbdce, 0xbbde, 0x3d0f, 0xbda3, 0xbe1c, 0xbe1e, + 0xbdb5, 0x3daa, 0x3e14, 0x3d6e, 0x3d03, 0x3cf8, 0xbd9a, 0xbd31, + 0x3cd5, 0x3d2b, 0xbe57, 0x3de8, 0xbe52, 0xbe12, 0xbc40, 0xba6c, + 0x3d14, 0xbcc1, 0x3c21, 0xbc3f, 0xbe04, 0x3c01, 0xbd9a, 0x3d33, + 0x3cc2, 0xbd5f, 0xbe02, 0xba58, 0xbe3a, 0x3d53, 0xbb56, 0x3e16, + 0x3c04, 0x3d6b, 0xbc2d, 0x3d8f, 0xbe00, 0x3ca8, 0x3c35, 0xbde7, + 0xbdc1, 0xbe2a, 0xbc7e, 0xbe08, 0x3d7b, 0x3e0e, 0xbd8d, 0x3d01, + 0x3d2b, 0x3c87, 0x3dbb, 0xbd40, 0x3e5c, 0xbe3e, 0xbd2a, 0x3dd4, + 0xbc88, 0x3d94, 0x3d2d, 0xbd9b, 0xbdc0, 0xbdcd, 0xbd30, 0x3de5, + 0x3e41, 0x3c4f, 0x3da7, 0x3d87, 0xbd99, 0x3e08, 0x3e06, 0x3cbe, + 0xbd45, 0x3cd7, 0xbd6d, 0xbe1d, 0x3d7f, 0x3d9c, 0x3bac, 0x3dfb, + 0x3d92, 0x3cf4, 0xbe05, 0xbc09, 0xbdf9, 0xbda2, 0xbd96, 0x3d46, + 0x3cb2, 0xbdf7, 0xbd59, 0xbcd4, 0x3dd6, 0x3d1b, 0x3e04, 0xbdc1, +}; +// L33-C2-R0-residual-final-0-input-moe.bin +// full payload SHA256 6b5c13d2958a9dfde633eb6ab62133c072846498f45bfd4dd8c0504d2eec71c1 +// row SHA256 bf56520da475a9d2cb4f51ff0e5e0a6e74fc9eb5910da7ceaf9c10f10cf7699a +inline constexpr std::array kFinalMoe = { + 0xba8e, 0xbb8c, 0x3a58, 0xbb00, 0xbbd4, 0x3b71, 0xbc1d, 0xbade, + 0xbb13, 0xbbbf, 0xbba1, 0x3bcf, 0xbb05, 0x3b36, 0x3baa, 0x39d0, + 0x3780, 0x3bcf, 0x3b22, 0xbc0a, 0x3b36, 0x3bfe, 0xbc08, 0x3b27, + 0x3b0e, 0xba38, 0xbb76, 0xbc17, 0xb780, 0xba4c, 0x3b3c, 0x3b90, + 0x3b69, 0xbc08, 0x3ad4, 0xbb1b, 0x39da, 0x3b8e, 0xbb4a, 0xb980, + 0xba60, 0xba78, 0xbbba, 0x3bc7, 0x3bba, 0xba00, 0x3bcf, 0x3bd9, + 0x3ac0, 0xbb17, 0xb99c, 0xbc2e, 0xbbab, 0xbb0b, 0x3c0d, 0xbb01, + 0x3b58, 0x3bf6, 0xbb37, 0x3a9d, 0xbbd2, 0xba94, 0xbc12, 0xbb0c, + 0x3b79, 0xbbb8, 0x3c14, 0x3b4e, 0xbabd, 0x3b00, 0xbb8f, 0x3b94, + 0xba64, 0xbaa8, 0x3be0, 0x3a88, 0x3c37, 0x3a8c, 0x3c18, 0xbaa0, + 0x3b0e, 0x3a9a, 0x3b63, 0xbc00, 0x3b87, 0xbb8a, 0xbbe2, 0xbb40, + 0x3930, 0x3bd0, 0x3a84, 0x3b5b, 0xbc14, 0xbb2f, 0xbaf2, 0x3b25, + 0x3bb6, 0xbb5a, 0xbba3, 0xbbf6, 0xbab2, 0x3ac0, 0xbb1d, 0x3c0e, + 0x3bb0, 0x3bc2, 0xbafe, 0x3bc9, 0x3a20, 0xbb4e, 0x3b4e, 0xbc0a, + 0xbc3a, 0xbb8f, 0xbb3c, 0xbba4, 0x3b98, 0x3ad0, 0x3af4, 0xb990, + 0xb890, 0x3a93, 0xbb2c, 0xbb8e, 0xba94, 0x3bf9, 0x3c22, 0x3c23, +}; +// L33-C2-R0-residual-final-0-input-weight.bin +// full payload SHA256 0c5b881a522de96353b1e62a6fbe3520dc1edb101cd6b8a5ad7f6282cc392b0e +// row SHA256 0c5b881a522de96353b1e62a6fbe3520dc1edb101cd6b8a5ad7f6282cc392b0e +inline constexpr std::array kFinalGamma = { + 0xbdb6, 0x3df9, 0xbec6, 0xbef9, 0x3f00, 0xbe8a, 0xbebc, 0xbef5, + 0xbdaa, 0xbea0, 0xbd2a, 0xbeb6, 0xbdd1, 0x3e7b, 0xbefc, 0xbe1e, + 0xbec5, 0x3e99, 0xbe44, 0x3e46, 0x3d9c, 0xbe54, 0x3a4b, 0xbe85, + 0x3d50, 0x3ec3, 0xbe58, 0x3ee5, 0x3e50, 0x3efe, 0x3e4b, 0xbd88, + 0x3ed6, 0x3e5f, 0x3e46, 0x3d6e, 0x3efe, 0xbed9, 0xbefb, 0xbec7, + 0xbea0, 0x3e43, 0x3e81, 0xbec4, 0xbe74, 0x3ec0, 0x3dd3, 0x3cd9, + 0xbee1, 0x3dd3, 0x3e64, 0xbe4b, 0x3e80, 0x3ef0, 0x3e18, 0xbea6, + 0xbe8f, 0xbee5, 0x3e97, 0x3da3, 0xbe38, 0xbe46, 0xbe01, 0x3d00, + 0x3ed3, 0x3ebd, 0x3e99, 0xbcdc, 0x3ed5, 0xbd66, 0x3edf, 0xbe1f, + 0x3ea5, 0x3dfd, 0x3e6c, 0xbdd9, 0xbe0f, 0x3ea5, 0xbecf, 0x3b6a, + 0xbe60, 0xbea3, 0xbed2, 0xbeef, 0xbd22, 0xbec6, 0x3e69, 0xbd26, + 0x3e8a, 0xbe21, 0xbeef, 0xbe3c, 0xbda6, 0x3e04, 0xbeca, 0xbe94, + 0x3e73, 0x3c97, 0xbe2a, 0x3ed8, 0x3d85, 0x3e96, 0xbeec, 0x3eb6, + 0x3e7c, 0xbbe5, 0xbeae, 0xbe0e, 0xbdd5, 0x3e2d, 0x3e97, 0x3e49, + 0x3ed3, 0xbe89, 0x3e25, 0x3df7, 0x3edc, 0x3d86, 0x3eb1, 0x3ed8, + 0x3e28, 0x3e88, 0xbe09, 0xbe0a, 0xbe07, 0x3e3b, 0xbd96, 0x3ead, +}; +// L33-C2-R0-residual-final-0-output-norm.bin +// full payload SHA256 38f07b72d5fdb6f9a030cb104a1107cf28847c3487074d1099805588267cf3f8 +// row SHA256 e9d4c8b1212f2d1243fc8ae4e28c7d5e31070472f9b3740f81d88b74e5993920 +inline constexpr std::array kFinalNorm = { + 0x3dd8, 0x3e58, 0x3e06, 0xbcee, 0x3de2, 0xbead, 0x3e7e, 0x3e1b, + 0xbdd0, 0xbe46, 0x3d07, 0xbe9f, 0x3e10, 0xbda0, 0x3d04, 0xbe2c, + 0x3f4d, 0xbea1, 0xbeaa, 0x3d96, 0xbda4, 0x3e3a, 0xb9d4, 0xbb33, + 0xbd15, 0x3e0f, 0x3e83, 0x3f0e, 0x3d61, 0xbe89, 0xbebc, 0x3d73, + 0x3dd5, 0xbda0, 0x3eb9, 0x3d92, 0x3e01, 0xbd08, 0x3eb3, 0x3dac, + 0xbd8c, 0xbbd3, 0xbf21, 0xbf0b, 0x3ef6, 0xbea5, 0x3c4b, 0x3cac, + 0x3dee, 0xbdb8, 0xbd7b, 0x3e01, 0xbe41, 0x3e51, 0x3c94, 0x3d63, + 0x3d70, 0x3e81, 0xbe97, 0xbbf7, 0x3ea0, 0xbd84, 0xbd97, 0x3d47, + 0xbea4, 0x3dcc, 0xbe72, 0xbd19, 0xbf4b, 0xba4c, 0x3e56, 0x3e11, + 0xbed4, 0xbe2a, 0x3e41, 0x3e0b, 0xbdae, 0x3f33, 0x3e6f, 0x3ae4, + 0xbe20, 0x3cb6, 0xbeda, 0x3e71, 0xbd8a, 0x3f04, 0xbe39, 0x39b5, + 0xbe70, 0xbe0d, 0xbe0d, 0x3c0f, 0x3d5d, 0x3ca0, 0x3edb, 0xbf01, + 0x3f29, 0x3ac6, 0xbd45, 0x3eab, 0xbccf, 0x3f01, 0xbeef, 0x3e40, + 0x3d48, 0x394f, 0x3f0a, 0x3e75, 0xbc80, 0x3d0b, 0xbdf2, 0x3ea2, + 0x3e76, 0xbd88, 0xbe53, 0x3b68, 0xbf16, 0xbda9, 0xbee0, 0xbd83, + 0x3e0c, 0xbe52, 0x3d61, 0x3e27, 0xbddc, 0x3db9, 0xbdb0, 0xbe69, +}; +} // namespace residual_norm_fixture diff --git a/tests/support/residual_norm_test.h b/tests/support/residual_norm_test.h new file mode 100644 index 0000000000..3ce0690b09 --- /dev/null +++ b/tests/support/residual_norm_test.h @@ -0,0 +1,121 @@ +// vllm.cpp original test harness for the shared residual-expression contract. +#pragma once +#include +#include +#include +#include +#include +#include "vt/backend.h" +#include "vt/dtype.h" +#include "vt/ops.h" +#include "vt/recipes.h" +#if defined(VLLM_CPP_HIP) +#include +#endif + +namespace residual_norm_test { +inline std::vector Devices() { + std::vector result{{vt::DeviceType::kCPU, 0}}; +#if defined(VLLM_CPP_HIP) + REQUIRE(vt::TryGetBackend({vt::DeviceType::kROCM, 0}) != nullptr); + result.push_back({vt::DeviceType::kROCM, 0}); +#endif + return result; +} +struct DeviceScope { + int previous = 0; + bool rocm; + explicit DeviceScope(vt::Device device) : rocm(device.type == vt::DeviceType::kROCM) { +#if defined(VLLM_CPP_HIP) + if (rocm) { REQUIRE(hipGetDevice(&previous) == hipSuccess); REQUIRE(hipSetDevice(device.index) == hipSuccess); } +#endif + } + ~DeviceScope() { +#if defined(VLLM_CPP_HIP) + if (rocm) (void)hipSetDevice(previous); +#endif + } +}; +struct Queue { + vt::Queue q; + explicit Queue(vt::Device device) { DeviceScope scope(device); q = vt::CreateQueue(device); } + ~Queue() { DeviceScope scope(q.device); vt::DestroyQueue(q); } +}; +struct Buffer { + vt::Queue& q; + vt::Tensor t; + size_t count; + static constexpr uint16_t kGuard = 0x3555; + Buffer(vt::Queue& queue, const std::vector& values, + int64_t rows, int64_t width, int64_t stride = 0) : q(queue), count(values.size()) { + DeviceScope scope(q.device); + auto& backend = vt::GetBackend(q.device); + void* data = backend.Alloc((count + 64) * sizeof(uint16_t)); + t = rows < 0 ? vt::Tensor::Contiguous(data, vt::DType::kBF16, q.device, {width}) + : vt::Tensor::Contiguous(data, vt::DType::kBF16, q.device, {rows, width}); + if (rows >= 0) t.stride[0] = stride == 0 ? width : stride; + std::vector storage(values); + storage.resize(count + 64, kGuard); + backend.Copy(q, data, storage.data(), storage.size() * sizeof(uint16_t)); + backend.Synchronize(q); + } + ~Buffer() { DeviceScope scope(q.device); vt::GetBackend(q.device).Free(t.data); } + std::vector Read(bool guard = false) const { + DeviceScope scope(q.device); + std::vector result(count + (guard ? 64 : 0)); + auto& backend = vt::GetBackend(q.device); + backend.Copy(q, result.data(), t.data, result.size() * sizeof(uint16_t)); + backend.Synchronize(q); + return result; + } + void CheckGuard() const { + const auto values = Read(true); + CHECK(std::all_of(values.begin() + static_cast(count), values.end(), + [](uint16_t value) { return value == kGuard; })); + } +}; +// Independent scalar reference: FP64 variance accumulation deliberately does +// not duplicate the native 256-lane tree. Exact boundary witnesses use values +// whose squares and sums are exact; general inputs use the pinned tolerance. +inline std::vector Reference(const std::vector& a, + const std::vector& base, const std::vector* delta, + const std::vector& gamma, int64_t rows, int64_t width, int64_t stride, + float epsilon, std::vector* residual = nullptr) { + std::vector out(static_cast(rows * stride), Buffer::kGuard); + if (residual != nullptr) residual->assign(out.size(), Buffer::kGuard); + for (int64_t row = 0; row < rows; ++row) { + std::vector sum(static_cast(width)); + double variance = 0; + for (int64_t j = 0; j < width; ++j) { + const size_t at = static_cast(row * stride + j); + const float pair = vt::BF16ToF32(a[at]) + vt::BF16ToF32(base[at]); + const float value = delta == nullptr ? pair : vt::BF16ToF32((*delta)[at]) + pair; + sum[j] = value; + variance += static_cast(value) * value; + } + const float inverse = 1.0f / std::sqrt(static_cast(variance / width) + epsilon); + for (int64_t j = 0; j < width; ++j) { + const size_t at = static_cast(row * stride + j); + out[at] = vt::F32ToBF16((sum[j] * inverse) * vt::BF16ToF32(gamma[j])); + if (residual != nullptr) (*residual)[at] = vt::F32ToBF16(sum[j]); + } + } + return out; +} +inline void Run(int mode, vt::Queue& q, vt::Tensor& out, const vt::Tensor& a, + const vt::Tensor& base, const vt::Tensor* delta, const vt::Tensor& gamma, + const vt::ResidualRmsNormArgs& args, vt::Tensor* residual) { + if (mode == 0) { vt::ResidualRmsNorm(q, out, a, base, delta, gamma, args, residual); return; } + if (mode == 1) { vt::FusedChain(q, out, a, base, delta, gamma, args, residual); return; } + vt::FusedBinding binding{}; + binding.n = 6; + binding.op[0] = const_cast(&a); + binding.op[1] = const_cast(&base); + binding.op[2] = const_cast(delta); + binding.op[3] = const_cast(&gamma); + binding.op[4] = &out; + binding.op[5] = residual; + vt::FusedParams params{}; params.eps = args.eps; + vt::FusedChainComposite(q, vt::ResidualRmsNormRecipe(args.descriptor), binding, params); +} +} // namespace residual_norm_test diff --git a/tests/support/rocm_moe_reference_set.h b/tests/support/rocm_moe_reference_set.h new file mode 100644 index 0000000000..7ae07dd292 --- /dev/null +++ b/tests/support/rocm_moe_reference_set.h @@ -0,0 +1,149 @@ +// vllm.cpp original (test harness); no upstream mirror. +// +// BACKEND-ROCM-BF16-MOE (#3094) production token gate: the reference set. +// +// The pinned primary's greedy decode is not unique across its own legitimate +// configurations. At length 33, the concurrency-1 and concurrency-2 captures of +// the identical prompt emit `[66,1,70,57,33,81,63,69]` and +// `[66,1,70,57,33,81,118,66]` for request 0, because the two runs meet a BF16 +// tie differently. A gate that compares the native run against one capture +// therefore fits one noise realization of the reference instead of the +// reference function. +// +// `Compare` implements the corrected rule: a request passes when its whole +// sequence equals one captured reference sequence (membership), never a +// position-by-position mix of two. Everything here is pure host code, so +// `tests/support/test_rocm_moe_reference_set.cpp` gates the predicate on the CPU +// and the production gate `tests/vllm/models/test_rocm_moe_bf16.cpp` applies it +// to the captured oracle records. +#pragma once + +#include +#include +#include +#include +#include + +#include + +namespace rocm_moe_reference_set { + +// One captured reference configuration: the concurrency of the oracle record +// and that record's generated tokens for one request. +struct Configuration { + int concurrency = 0; + std::vector tokens; +}; + +// The result of comparing one native request sequence with the reference set. +struct Comparison { + // Index of the whole-sequence member the run equals, or -1 when the run + // matches no captured configuration. The run's own configuration wins when it + // matches, so the reported member is the strongest available statement. + int matched = -1; + // Index of the member captured at the run's own concurrency, or -1 when the + // reference set holds no record for that concurrency. + int same_configuration = -1; + // Whether the run equals that same-configuration member. The gate reports this + // value and never asserts it, because the reference disagrees with itself at + // one workload. + bool same_configuration_match = false; + // Positions at which the captured configurations disagree with each other, + // including a position that one configuration does not reach. + std::vector disagreements; + + bool pass() const { return matched >= 0; } +}; + +// Gather every captured sequence for one request from the records of the same +// length and repeat. Request 1 exists only at concurrency 2, so a record whose +// token list is shorter than the request index contributes nothing. +inline std::vector Collect(const nlohmann::json& runs, int length, + int repeat, size_t request) { + std::vector reference; + for (const auto& record : runs) { + if (record.at("length").get() != length) continue; + if (record.at("repeat").get() != repeat) continue; + const auto sequences = record.at("tokens").get>>(); + if (request >= sequences.size()) continue; + reference.push_back({record.at("concurrency").get(), sequences[request]}); + } + return reference; +} + +// Return every position at which the captured configurations disagree. A +// singleton set cannot disagree, so its positions stay exact. +inline std::vector Disagreements(const std::vector& reference) { + std::vector positions; + if (reference.size() < 2) return positions; + size_t width = 0; + for (const auto& configuration : reference) + width = std::max(width, configuration.tokens.size()); + for (size_t position = 0; position < width; ++position) { + const auto& first = reference.front().tokens; + for (size_t i = 1; i < reference.size(); ++i) { + const auto& other = reference[i].tokens; + if (position >= first.size() || position >= other.size() || + first[position] != other[position]) { + positions.push_back(static_cast(position)); + break; + } + } + } + return positions; +} + +// Compare one native request sequence with the reference set. The run passes +// only as a whole-sequence member; the per-position disagreement list never +// widens the accepted set. +inline Comparison Compare(const std::vector& reference, + const std::vector& run, int concurrency) { + Comparison comparison; + comparison.disagreements = Disagreements(reference); + for (size_t i = 0; i < reference.size(); ++i) { + if (reference[i].concurrency == concurrency) + comparison.same_configuration = static_cast(i); + if (comparison.matched < 0 && reference[i].tokens == run) + comparison.matched = static_cast(i); + } + comparison.same_configuration_match = + comparison.same_configuration >= 0 && + reference[static_cast(comparison.same_configuration)].tokens == run; + if (comparison.same_configuration_match) + comparison.matched = comparison.same_configuration; + return comparison; +} + +// Render a token sequence for a gate report. +inline std::string Sequence(const std::vector& tokens) { + std::string text = "["; + for (size_t i = 0; i < tokens.size(); ++i) { + if (i != 0) text += ","; + text += std::to_string(tokens[i]); + } + return text + "]"; +} + +// Render the reference set for a gate report. +inline std::string Describe(const std::vector& reference) { + std::string text = "{"; + for (size_t i = 0; i < reference.size(); ++i) { + if (i != 0) text += ", "; + text += "concurrency " + std::to_string(reference[i].concurrency) + " " + + Sequence(reference[i].tokens); + } + if (reference.empty()) text += "no captured configuration"; + return text + "}"; +} + +// Render a position list for a gate report. +inline std::string Describe(const std::vector& positions) { + std::string text = "["; + for (size_t i = 0; i < positions.size(); ++i) { + if (i != 0) text += ","; + text += std::to_string(positions[i]); + } + return text + "]"; +} + +} // namespace rocm_moe_reference_set diff --git a/tests/support/test_rocm_moe_reference_set.cpp b/tests/support/test_rocm_moe_reference_set.cpp new file mode 100644 index 0000000000..2deb9f981e --- /dev/null +++ b/tests/support/test_rocm_moe_reference_set.cpp @@ -0,0 +1,229 @@ +// The gate on the production token gate's reference-set predicate (#3094). +// +// vllm.cpp original (test harness); no upstream mirror. The production gate in +// `tests/vllm/models/test_rocm_moe_bf16.cpp` now accepts a native request +// sequence only as a whole-sequence member of the reference's captured +// configuration set. That predicate is pure host code in +// `support/rocm_moe_reference_set.h`, so its four required outcomes are gated +// here without a device: +// +// 1. the concurrency-1 member passes and the matched configuration is named; +// 2. the concurrency-2 member passes; +// 3. a change at a position where the reference configurations AGREE fails, +// because those positions stay exact under every configuration; +// 4. a per-position mixture of the two members fails, because membership is +// whole-sequence and never per-position. +// +// The values below are the captured length-33, length-3, and length-1 records +// from `oracle-selection-6/production.json`, whose pinned primary emits both +// `[66,1,70,57,33,81,63,69]` and `[66,1,70,57,33,81,118,66]` for request 0 +// depending on the captured concurrency. +#include + +#include +#include + +#include + +#include "support/rocm_moe_reference_set.h" + +namespace { + +using rocm_moe_reference_set::Collect; +using rocm_moe_reference_set::Compare; +using rocm_moe_reference_set::Comparison; +using rocm_moe_reference_set::Configuration; + +const std::vector kConcurrencyOne{66, 1, 70, 57, 33, 81, 63, 69}; +const std::vector kConcurrencyTwo{66, 1, 70, 57, 33, 81, 118, 66}; +const std::vector kSecondRequest{63, 92, 85, 72, 62, 92, 85, 72}; +const std::vector kLengthOne{81, 118, 105, 72, 105, 72, 72, 72}; + +// One record per (length, concurrency, repeat), each holding one token sequence +// per request. Request 1 exists only in the concurrency-2 records, and the +// length-33 repeat-1 record is a decoy for the repeat filter. +nlohmann::json CapturedRuns() { + return nlohmann::json::parse(R"([ + {"length": 33, "concurrency": 1, "repeat": 0, + "tokens": [[66,1,70,57,33,81,63,69]]}, + {"length": 33, "concurrency": 2, "repeat": 0, + "tokens": [[66,1,70,57,33,81,118,66],[63,92,85,72,62,92,85,72]]}, + {"length": 33, "concurrency": 1, "repeat": 1, "tokens": [[1,2,3]]}, + {"length": 1, "concurrency": 1, "repeat": 0, + "tokens": [[81,118,105,72,105,72,72,72]]}, + {"length": 1, "concurrency": 2, "repeat": 0, + "tokens": [[81,118,105,72,105,72,72,72],[27,77,59,2,115,77,59,97]]}, + {"length": 3, "concurrency": 1, "repeat": 0, + "tokens": [[72,84,66,66,66,66,66,25]]}, + {"length": 3, "concurrency": 2, "repeat": 0, + "tokens": [[72,84,66,66,66,66,66,25],[64,78,19,26,2,35,1,64]]} + ])"); +} + +// The name of the member a comparison reports, or -1 for no match. +int MatchedConcurrency(const std::vector& reference, + const Comparison& comparison) { + if (!comparison.pass()) return -1; + return reference[static_cast(comparison.matched)].concurrency; +} + +} // namespace + +TEST_CASE("reference set collects every captured configuration of one workload") { + const auto runs = CapturedRuns(); + const auto length_33 = Collect(runs, 33, 0, 0); + REQUIRE(length_33.size() == 2); + CHECK(length_33[0].concurrency == 1); + CHECK(length_33[0].tokens == kConcurrencyOne); + CHECK(length_33[1].concurrency == 2); + CHECK(length_33[1].tokens == kConcurrencyTwo); + CHECK(rocm_moe_reference_set::Disagreements(length_33) == std::vector{6, 7}); + + // Request 1 exists only at concurrency 2, so its set is a singleton. + const auto request_one = Collect(runs, 33, 0, 1); + REQUIRE(request_one.size() == 1); + CHECK(request_one[0].concurrency == 2); + CHECK(request_one[0].tokens == kSecondRequest); + CHECK(rocm_moe_reference_set::Disagreements(request_one).empty()); + + // The repeat selects the record, so the decoy record cannot leak in. + const auto other_repeat = Collect(runs, 33, 1, 0); + REQUIRE(other_repeat.size() == 1); + CHECK(other_repeat[0].tokens == std::vector{1, 2, 3}); + + // Lengths 1 and 3 agree across concurrency, so their sets hold no + // disagreement even though they hold two records each. + const auto length_one = Collect(runs, 1, 0, 0); + REQUIRE(length_one.size() == 2); + CHECK(length_one[0].tokens == kLengthOne); + CHECK(length_one[1].tokens == kLengthOne); + CHECK(rocm_moe_reference_set::Disagreements(length_one).empty()); + + const auto length_three = Collect(runs, 3, 0, 0); + REQUIRE(length_three.size() == 2); + CHECK(length_three[0].tokens == length_three[1].tokens); + CHECK(rocm_moe_reference_set::Disagreements(length_three).empty()); + + // A request index that no capture reaches collects nothing, which the + // production gate refuses with REQUIRE(!reference_set.empty()). + CHECK(Collect(runs, 33, 0, 2).empty()); +} + +TEST_CASE("the concurrency-1 member passes and reports its configuration") { + const auto reference = Collect(CapturedRuns(), 33, 0, 0); + + const auto same = Compare(reference, kConcurrencyOne, 1); + CHECK(same.pass()); + CHECK(MatchedConcurrency(reference, same) == 1); + CHECK(same.same_configuration_match); + + // The recorded tie workload: the native run equals the concurrency-1 member + // and the concurrency-2 record disagrees, so the gate passes while the + // same-configuration comparison is reported false. + const auto other = Compare(reference, kConcurrencyOne, 2); + CHECK(other.pass()); + CHECK(MatchedConcurrency(reference, other) == 1); + CHECK(other.same_configuration_match == false); + CHECK(other.same_configuration == 1); +} + +TEST_CASE("the concurrency-2 member passes") { + const auto reference = Collect(CapturedRuns(), 33, 0, 0); + const auto comparison = Compare(reference, kConcurrencyTwo, 2); + CHECK(comparison.pass()); + CHECK(MatchedConcurrency(reference, comparison) == 2); + CHECK(comparison.same_configuration_match); +} + +TEST_CASE("a change at a position where the reference configurations agree fails") { + const auto reference = Collect(CapturedRuns(), 33, 0, 0); + + // Position 3 is equal in both captured configurations, so no member explains + // the change. + auto changed_stable = kConcurrencyOne; + changed_stable[3] = 34; + const auto stable = Compare(reference, changed_stable, 1); + CHECK(stable.pass() == false); + CHECK(stable.same_configuration_match == false); + CHECK(rocm_moe_reference_set::Disagreements(reference) == std::vector{6, 7}); + + // A change at a position where the configurations disagree fails too: the + // position stays exact inside each member, and membership is whole-sequence. + auto changed_tie = kConcurrencyTwo; + changed_tie[7] = 69; + const auto tie = Compare(reference, changed_tie, 2); + CHECK(tie.pass() == false); + CHECK(tie.same_configuration_match == false); + + const auto shorter = Compare(reference, std::vector{66, 1, 70}, 1); + CHECK(shorter.pass() == false); +} + +TEST_CASE("a per-position mix of two members fails") { + const auto reference = Collect(CapturedRuns(), 33, 0, 0); + + // The first six tokens agree; the mix takes one member's token at position 6 + // and the other member's token at position 7, so it equals neither member. + auto first_mix = kConcurrencyOne; + first_mix[6] = kConcurrencyTwo[6]; + const auto mix_one = Compare(reference, first_mix, 2); + CHECK(mix_one.pass() == false); + CHECK(mix_one.same_configuration_match == false); + CHECK(rocm_moe_reference_set::Disagreements(reference) == std::vector{6, 7}); + + auto second_mix = kConcurrencyTwo; + second_mix[7] = kConcurrencyOne[7]; + const auto mix_two = Compare(reference, second_mix, 2); + CHECK(mix_two.pass() == false); + CHECK(mix_two.same_configuration_match == false); + + // Every position of a mix is individually present in some member, and the + // whole sequence still matches none. + CHECK((first_mix[6] == kConcurrencyTwo[6] && first_mix[7] == kConcurrencyOne[7])); + CHECK(first_mix != kConcurrencyOne); + CHECK(first_mix != kConcurrencyTwo); +} + +TEST_CASE("a singleton reference set keeps its positions exact") { + const auto runs = CapturedRuns(); + + // Request 1 is captured once. Membership in a singleton is equality. + const auto second_request = Collect(runs, 33, 0, 1); + const auto exact = Compare(second_request, kSecondRequest, 2); + CHECK(exact.pass()); + CHECK(exact.same_configuration_match); + CHECK(exact.disagreements.empty()); + + auto changed = kSecondRequest; + changed[5] = 93; + const auto inexact = Compare(second_request, changed, 2); + CHECK(inexact.pass() == false); + CHECK(inexact.same_configuration_match == false); + + // Lengths 1 and 3 hold two records that agree, so a change fails there too. + const auto length_one = Collect(runs, 1, 0, 0); + auto changed_length_one = kLengthOne; + changed_length_one[7] = 73; + const auto length_one_failure = Compare(length_one, changed_length_one, 2); + CHECK(length_one_failure.pass() == false); + CHECK(length_one_failure.disagreements.empty()); + + const auto length_three = Collect(runs, 3, 0, 0); + CHECK(Compare(length_three, length_three[1].tokens, 2).pass()); + auto changed_length_three = length_three[1].tokens; + changed_length_three[0] = 73; + const auto length_three_failure = Compare(length_three, changed_length_three, 2); + CHECK(length_three_failure.pass() == false); + CHECK(length_three_failure.disagreements.empty()); + + const auto built = std::vector{{1, kConcurrencyOne}}; + CHECK(Compare(built, kConcurrencyOne, 1).pass()); + CHECK(Compare(built, kConcurrencyTwo, 1).pass() == false); +} + +TEST_CASE("an empty reference set fails every run") { + const auto empty = Compare({}, kConcurrencyOne, 1); + CHECK(empty.pass() == false); + CHECK(empty.disagreements.empty()); + CHECK(empty.same_configuration == -1); +} diff --git a/tests/vllm/model_executor/test_qwen3_moe_lm_head_bf16.cpp b/tests/vllm/model_executor/test_qwen3_moe_lm_head_bf16.cpp new file mode 100644 index 0000000000..24a7b35f01 --- /dev/null +++ b/tests/vllm/model_executor/test_qwen3_moe_lm_head_bf16.cpp @@ -0,0 +1,661 @@ +// vllm.cpp original test harness; no upstream mirror. +// +// BACKEND-ROCM-BF16-MOE (#3116) — the LM-head output boundary. +// +// THE QUESTION THIS FILE SETTLES. The native Qwen3-MoE forward ends with an +// UNTIED lm_head projection that stores F32 logits +// (src/vllm/model_executor/models/qwen3_moe.cpp ForwardLayers, `vt::Matmul` into +// a `DType::kF32` buffer). The pinned compiled primary stores its head output in +// BF16 and widens it to F32 only where the sampler needs it +// (.agents/specs/rocm-residual-norm.md:320). Issue #3116 asks whether mirroring +// that boundary changes the native head arithmetic, and whether that mirror +// alone moves the decode step-6 argmax. +// +// THE INSTRUMENT. `oracle-diagnostic-2/L33-C2-R0-head-6.{json,bin}` holds the +// primary's exact head input (`hidden` [2,128] bf16), its exact lm_head weight +// (`weight` [128,128] bf16 in torch Linear [vocab,hidden] order), and its exact +// BF16 head output (`logits` [2,128] bf16) for decode step 6 of the L33/C2/R0 +// workload — the exact-tie step of the BF16-MoE spec. Replaying those bytes +// through the operator the forward itself calls isolates the head from the +// hidden-state divergence (#3115) that the same step also carries. +// +// FOUR ARMS, one projection, no reimplementation: +// (a) `vt::Matmul(src, lm) -> f32` the CURRENT native boundary; +// (a') (a) narrowed to bf16 on the host (RNE) what a store-only change buys; +// (b) `vt::Matmul(src, lm) -> bf16` then +// `vt::CastF32` the PRIMARY's boundary, mirrored; +// (c) `logits.bin` widened bf16 -> f32 the primary's own words. +// +// The weight is uploaded in the [H, vocab] orientation the loader builds +// (`LoadBf16Transposed("lm_head.weight")`, qwen3_moe_weights.cpp:149), because +// that is the orientation the forward's `vt::Matmul` consumes; the artifact +// stores the torch Linear [vocab, H] order. +// +// `VT_MOE_HEAD_FIXTURE` names the directory holding the artifact. The two cases +// that replay it are decorated `doctest::skip(...)` when the variable is unset, +// so doctest reports them skipped while the fixture-free call-site case below +// still runs and still gates. The process then exits 77 (CTest: Skipped) rather +// than reporting a green gate that never ran, which is the convention +// `tests/CMakeLists.txt` documents -- and it does so ONLY when nothing failed, +// so a reddened case keeps doctest's own non-zero status. +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include "vllm/model_executor/models/lm_head_projection.h" +#include "vllm/model_executor/models/qwen3_moe.h" +#include "vllm/transformers_utils/hf_config.h" +#include "vt/backend.h" +#include "vt/device.h" +#include "vt/dtype.h" +#include "vt/ops.h" +#include "vt/tensor.h" + +#if defined(VLLM_CPP_HIP) +#include +#endif + +namespace { + +using vt::Device; +using vt::DeviceType; +using vt::DType; +using vt::Queue; +using vt::Tensor; + +// ─── The captured artifact ────────────────────────────────────────────────── + +struct HeadArtifact { + std::string label; + int64_t step = 0; + int64_t rows = 0; + int64_t hidden = 0; + int64_t vocab = 0; + std::vector hidden_bf16; // [rows, hidden] — the primary's head input + std::vector weight_bf16; // [vocab, hidden] — torch Linear lm_head.weight + std::vector logits_bf16; // [rows, vocab] — the primary's head output +}; + +std::vector ReadBf16(const std::string& path, size_t words) { + std::ifstream in(path, std::ios::binary); + REQUIRE_MESSAGE(in.good(), "cannot open ", path); + std::vector out(words); + in.read(reinterpret_cast(out.data()), + static_cast(words * sizeof(uint16_t))); + REQUIRE_MESSAGE(in.gcount() == static_cast(words * sizeof(uint16_t)), + "short read from ", path); + return out; +} + +HeadArtifact LoadArtifact(const std::string& dir, const std::string& label, int64_t step) { + const std::string stem = label + "-head-" + std::to_string(step); + const std::string meta_path = dir + "/" + stem + ".json"; + std::ifstream meta_in(meta_path); + REQUIRE_MESSAGE(meta_in.good(), "cannot open ", meta_path); + nlohmann::json meta; + meta_in >> meta; + + auto tensor = [&](const char* name, std::vector& shape) { + const auto& t = meta.at("tensors").at(name); + REQUIRE(t.at("dtype").get() == "torch.bfloat16"); + REQUIRE(t.at("element_bytes").get() == 2); + shape = t.at("shape").get>(); + REQUIRE(shape.size() == 2); + return ReadBf16(dir + "/" + t.at("file").get(), + static_cast(shape[0] * shape[1])); + }; + + HeadArtifact a; + a.label = meta.at("label").get(); + a.step = meta.at("step").get(); + std::vector hs, ws, ls; + a.hidden_bf16 = tensor("hidden", hs); + a.weight_bf16 = tensor("weight", ws); + a.logits_bf16 = tensor("logits", ls); + REQUIRE(hs == ls); + REQUIRE(ws[0] == ls[1]); + REQUIRE(ws[1] == hs[1]); + a.rows = hs[0]; + a.hidden = hs[1]; + a.vocab = ls[1]; + return a; +} + +// ─── Host-side arithmetic the comparison needs ────────────────────────────── + +// bf16 -> f32 is the high half of the f32 word; the widening is exact. +float Widen(uint16_t word) { + const uint32_t bits = static_cast(word) << 16; + float value = 0.0f; + std::memcpy(&value, &bits, sizeof(value)); + return value; +} + +std::vector Widen(const std::vector& words) { + std::vector out(words.size()); + for (size_t i = 0; i < words.size(); ++i) out[i] = Widen(words[i]); + return out; +} + +// f32 -> bf16 round-to-nearest-even, matching vt::F32ToBF16 (src/vt/dtype.cpp) +// and the bf16 store both device backends perform. +uint16_t NarrowRne(float value) { + uint32_t u = 0; + std::memcpy(&u, &value, sizeof(u)); + if ((u & 0x7F800000u) == 0x7F800000u && (u & 0x007FFFFFu)) { + return static_cast((u >> 16) | 0x0040u); + } + const uint32_t rounding = 0x7FFFu + ((u >> 16) & 1u); + return static_cast((u + rounding) >> 16); +} + +std::vector NarrowRne(const std::vector& values) { + std::vector out(values.size()); + for (size_t i = 0; i < values.size(); ++i) out[i] = NarrowRne(values[i]); + return out; +} + +// Lowest index wins an exact tie, which is the native GreedyArgmax rule +// (src/vt/rocm/rocm_dense_basic.hip ArgmaxK, `v > best || (v == best && j < arg)`). +std::vector TopOrder(const std::vector& row, int64_t vocab, int count) { + std::vector order(static_cast(vocab)); + for (int64_t j = 0; j < vocab; ++j) order[static_cast(j)] = static_cast(j); + std::stable_sort(order.begin(), order.end(), [&](int x, int y) { + const float vx = row[static_cast(x)], vy = row[static_cast(y)]; + if (vx != vy) return vx > vy; + return x < y; + }); + order.resize(static_cast(count)); + return order; +} + +double MaxAbsDiff(const std::vector& x, const std::vector& y) { + double worst = 0.0; + for (size_t i = 0; i < x.size(); ++i) + worst = std::max(worst, std::fabs(static_cast(x[i]) - static_cast(y[i]))); + return worst; +} + +double MeanAbsDiff(const std::vector& x, const std::vector& y) { + double sum = 0.0; + for (size_t i = 0; i < x.size(); ++i) + sum += std::fabs(static_cast(x[i]) - static_cast(y[i])); + return x.empty() ? 0.0 : sum / static_cast(x.size()); +} + +size_t DifferingWords(const std::vector& x, const std::vector& y) { + size_t n = 0; + for (size_t i = 0; i < x.size(); ++i) n += x[i] != y[i]; + return n; +} + +std::string Describe(const std::vector& row, const std::vector& order) { + std::string out = "argmax=" + std::to_string(order[0]) + " top2=["; + for (size_t i = 0; i < order.size(); ++i) { + if (i != 0) out += ", "; + out += std::to_string(order[i]) + ": " + + std::to_string(static_cast(row[static_cast(order[i])])); + } + return out + "]"; +} + +// ─── Device plumbing ──────────────────────────────────────────────────────── + +struct DeviceScope { + int previous = 0; + bool rocm = false; + explicit DeviceScope(Device device) : rocm(device.type == DeviceType::kROCM) { +#if defined(VLLM_CPP_HIP) + if (rocm) { + REQUIRE(hipGetDevice(&previous) == hipSuccess); + REQUIRE(hipSetDevice(device.index) == hipSuccess); + } +#else + (void)device; +#endif + } + ~DeviceScope() { +#if defined(VLLM_CPP_HIP) + if (rocm) (void)hipSetDevice(previous); +#endif + } +}; + +struct QueueHandle { + Queue q{}; + explicit QueueHandle(Device device) { + DeviceScope scope(device); + q = vt::CreateQueue(device); + } + ~QueueHandle() { + DeviceScope scope(q.device); + vt::DestroyQueue(q); + } +}; + +// One device-resident bf16 upload, freed on scope exit. +struct DeviceBf16 { + vt::Backend& b; + Queue& q; + void* p = nullptr; + Tensor t{}; + DeviceBf16(vt::Backend& backend, Queue& queue, const std::vector& words, + const std::vector& shape) + : b(backend), q(queue) { + DeviceScope scope(queue.device); + p = b.Alloc(words.size() * sizeof(uint16_t)); + b.Copy(q, p, words.data(), words.size() * sizeof(uint16_t)); + b.Synchronize(q); + t = Tensor::Contiguous(p, DType::kBF16, q.device, {shape[0], shape[1]}); + } + ~DeviceBf16() { + DeviceScope scope(q.device); + b.Free(p); + } + std::vector Read() const { + DeviceScope scope(q.device); + std::vector out(static_cast(t.Numel())); + b.Copy(q, out.data(), p, out.size() * sizeof(uint16_t)); + b.Synchronize(q); + return out; + } + DeviceBf16(const DeviceBf16&) = delete; + DeviceBf16& operator=(const DeviceBf16&) = delete; +}; + +struct DeviceF32 { + vt::Backend& b; + Queue& q; + void* p = nullptr; + Tensor t{}; + DeviceF32(vt::Backend& backend, Queue& queue, const std::vector& shape) + : b(backend), q(queue) { + DeviceScope scope(queue.device); + p = b.Alloc(static_cast(shape[0] * shape[1]) * sizeof(float)); + t = Tensor::Contiguous(p, DType::kF32, q.device, {shape[0], shape[1]}); + } + ~DeviceF32() { + DeviceScope scope(q.device); + b.Free(p); + } + std::vector Read() const { + DeviceScope scope(q.device); + std::vector out(static_cast(t.Numel())); + b.Copy(q, out.data(), p, out.size() * sizeof(float)); + b.Synchronize(q); + return out; + } + DeviceF32(const DeviceF32&) = delete; + DeviceF32& operator=(const DeviceF32&) = delete; +}; + +std::vector Devices() { + std::vector devices{{DeviceType::kCPU, 0}}; +#if defined(VLLM_CPP_HIP) + if (vt::TryGetBackend({DeviceType::kROCM, 0}) != nullptr) + devices.push_back({DeviceType::kROCM, 0}); +#endif + return devices; +} + +std::string DeviceName(Device device) { + return device.type == DeviceType::kROCM ? "ROCM" : "CPU"; +} + +// The artifact stores the torch Linear [vocab, hidden] order; the loader hands +// the untied head to `vt::Matmul` as [H, vocab] +// (`qwen3_moe_weights.cpp` `LoadBf16Transposed("lm_head.weight")`), so the +// replay uploads the same transposed bytes the model holds. +std::vector TransposedWeight(const HeadArtifact& art) { + std::vector weight_hv(art.weight_bf16.size()); + for (int64_t v = 0; v < art.vocab; ++v) + for (int64_t h = 0; h < art.hidden; ++h) + weight_hv[static_cast(h * art.vocab + v)] = + art.weight_bf16[static_cast(v * art.hidden + h)]; + return weight_hv; +} + +// ─── The three arms, through the operators the forward calls ──────────────── + +struct Measurements { + std::vector a_f32; // (a) native F32 head output + std::vector a_words; // (a') that output narrowed to bf16 on the host + std::vector b_f32; // (b) native BF16 head output, widened by vt::CastF32 + std::vector b_words; // (b) before widening + std::vector c_words; // (c) the primary's BF16 words +}; + +Measurements Measure(vt::Backend& b, Queue& q, const HeadArtifact& art) { + DeviceScope scope(q.device); + DeviceBf16 hidden(b, q, art.hidden_bf16, {art.rows, art.hidden}); + DeviceBf16 weight(b, q, TransposedWeight(art), {art.hidden, art.vocab}); + DeviceBf16 head(b, q, std::vector(static_cast(art.rows * art.vocab), 0), + {art.rows, art.vocab}); + DeviceF32 f32_out(b, q, {art.rows, art.vocab}); + DeviceF32 widened(b, q, {art.rows, art.vocab}); + + // (a) the current boundary: the projection stores F32. + vt::Matmul(q, f32_out.t, hidden.t, weight.t); + // (b) the primary's boundary: the projection stores BF16, the shared op widens. + vt::Matmul(q, head.t, hidden.t, weight.t); + vt::CastF32(q, widened.t, head.t); + b.Synchronize(q); + + Measurements m; + m.a_f32 = f32_out.Read(); + m.a_words = NarrowRne(m.a_f32); + m.b_f32 = widened.Read(); + m.b_words = head.Read(); + m.c_words = art.logits_bf16; + return m; +} + +void Report(const std::string& device, int64_t row, const HeadArtifact& art, + const Measurements& m) { + const int64_t V = art.vocab; + const size_t base = static_cast(row * V); + const auto slice = [&](const std::vector& all) { + return std::vector(all.begin() + static_cast(base), + all.begin() + static_cast(base + V)); + }; + const auto word_slice = [&](const std::vector& all) { + return std::vector(all.begin() + static_cast(base), + all.begin() + static_cast(base + V)); + }; + const std::vector a = slice(m.a_f32); + const std::vector bv = slice(m.b_f32); + const std::vector aw = word_slice(m.a_words); + const std::vector bw = word_slice(m.b_words); + const std::vector cw = word_slice(m.c_words); + const std::vector cv = Widen(cw); + + const std::vector order_a = TopOrder(a, V, 2); + const std::vector order_b = TopOrder(bv, V, 2); + const std::vector order_c = TopOrder(cv, V, 2); + const double max_ac = MaxAbsDiff(a, cv), mean_ac = MeanAbsDiff(a, cv); + const double max_ab = MaxAbsDiff(a, bv), mean_ab = MeanAbsDiff(a, bv); + const size_t differ_bc = DifferingWords(bw, cw); + const size_t differ_ac = DifferingWords(aw, cw); + + std::cout << "[lm-head boundary] device=" << device << " row=" << row + << " max|a-c|=" << max_ac << " mean|a-c|=" << mean_ac + << " max|a-b|=" << max_ab << " mean|a-b|=" << mean_ab << "\n" + << " b==c: " << (differ_bc == 0 ? "EQUAL" : "DIFFER") << " (" << differ_bc << "/" + << V << " bf16 words)" + << " a'==c: " << (differ_ac == 0 ? "EQUAL" : "DIFFER") << " (" << differ_ac << "/" << V + << " bf16 words), max|a'-c|=" << MaxAbsDiff(Widen(aw), cv) << "\n" + << " (a) f32 " << Describe(a, order_a) << "\n" + << " (b) bf16 " << Describe(bv, order_b) << "\n" + << " (c) prim " << Describe(cv, order_c) << std::endl; + + CAPTURE(device); + CAPTURE(row); + // The mirror's whole claim: given the primary's own head input, the native + // BF16 boundary reproduces the primary's words element for element, and picks + // the primary's argmax. + CHECK(differ_bc == 0); + CHECK(order_b[0] == order_c[0]); + CHECK(bw[static_cast(order_b[0])] == cw[static_cast(order_c[0])]); + // ...and the F32 boundary the tree ships today does NOT reproduce them. + CHECK(max_ac > 0.0); +} + +// ─── The fixture-absent run: per-case skip, process-level 77 ──────────────── +// +// The two replay cases need `VT_MOE_HEAD_FIXTURE`; the production call-site case +// at the bottom needs no directory at all. Before this repair both replay cases +// called `std::exit(77)`, and because doctest runs the cases in file order that +// stopped the binary at the FIRST one whenever the fixture was absent: the +// fixture-free case never ran, and the whole binary read as Skipped although a +// gate inside it could have run (the fresh review's finding 1). +// +// The replay cases are now decorated `doctest::skip(HeadFixtureAbsent())`, so +// doctest itself reports them as skipped in its summary while the call-site case +// still runs and still gates. A fixture-absent run still exits 77 -- CTest +// reports Skipped and a shell chain stops, the convention `tests/CMakeLists.txt` +// documents -- but ONLY when nothing failed, so a reddened call-site case keeps +// doctest's own non-zero exit instead of being hidden behind the skip. +bool HeadFixtureAbsent() { + const char* directory = std::getenv("VT_MOE_HEAD_FIXTURE"); + return directory == nullptr || directory[0] == '\0'; +} + +bool g_replay_cases_skipped = false; +bool g_run_failed = false; + +// A listener, not a reporter: listeners are always active whatever `-r=` +// selects, so neither the skip note nor the failure guard can be switched off +// from the command line (`tests/vt/test_ops_attention_cross.cpp:662` has the +// same shape). +struct ReplaySkipListener : public doctest::IReporter { + explicit ReplaySkipListener(const doctest::ContextOptions&) {} + void test_run_start() override {} + void report_query(const doctest::QueryData&) override {} + void test_run_end(const doctest::TestRunStats& stats) override { + g_run_failed = stats.numTestCasesFailed != 0 || stats.numAssertsFailed != 0; + } + void test_case_start(const doctest::TestCaseData&) override {} + void test_case_reenter(const doctest::TestCaseData&) override {} + void test_case_end(const doctest::CurrentTestCaseStats&) override {} + void test_case_exception(const doctest::TestCaseException&) override {} + void subcase_start(const doctest::SubcaseSignature&) override {} + void subcase_end() override {} + void log_assert(const doctest::AssertData&) override {} + void log_message(const doctest::MessageData&) override {} + // This also fires for a case a filter excluded; only the decorator sets + // `m_skip`, and only that means "the fixture is not here". + void test_case_skipped(const doctest::TestCaseData& tc) override { + if (!tc.m_skip) return; + g_replay_cases_skipped = true; + std::cout << "[lm-head boundary] SKIPPED: " << tc.m_name + << " (set VT_MOE_HEAD_FIXTURE to run it)" << std::endl; + } +}; +DOCTEST_REGISTER_LISTENER("vt-moe-head-replay-skip", 1, ReplaySkipListener); + +// Registered during static initialization, so it runs after doctest's main has +// printed its summary and returned. `std::_Exit` rather than `std::exit`: this +// IS an exit handler, and re-entering the exit sequence is undefined. +void ExitSkippedWhenTheReplayDidNotRun() { + if (!g_replay_cases_skipped || g_run_failed) return; + std::cout.flush(); + std::fflush(nullptr); + std::fprintf(stderr, + "\n*** SKIPPED (exit 77): the artifact-replay cases did not run because " + "VT_MOE_HEAD_FIXTURE is unset. The fixture-free production case above DID " + "run and its result stands; this status says only that the replay gate did " + "not. ***\n\n"); + std::fflush(stderr); + std::_Exit(77); +} + +struct RegisterExitSkippedWhenTheReplayDidNotRun { + RegisterExitSkippedWhenTheReplayDidNotRun() { std::atexit(&ExitSkippedWhenTheReplayDidNotRun); } +}; +[[maybe_unused]] const RegisterExitSkippedWhenTheReplayDidNotRun + g_exit_skipped_when_the_replay_did_not_run; + +} // namespace + +TEST_CASE("qwen3 MoE LM-head BF16 boundary: primary artifact replay" * + doctest::skip(HeadFixtureAbsent())) { + // The decorator above covers the normal run; `--no-skip` forces the case + // anyway, and then the absent fixture must FAIL rather than reach + // `LoadArtifact` as a null path. + REQUIRE_MESSAGE(!HeadFixtureAbsent(), + "VT_MOE_HEAD_FIXTURE must name the directory holding " + "L33-C2-R0-head-6.{json,hidden.bin,weight.bin,logits.bin}"); + const HeadArtifact art = LoadArtifact(std::getenv("VT_MOE_HEAD_FIXTURE"), "L33-C2-R0", 6); + for (const Device device : Devices()) { + CAPTURE(DeviceName(device)); + QueueHandle queue(device); + vt::Backend& backend = vt::GetBackend(device.type); + const Measurements m = Measure(backend, queue.q, art); + for (int64_t row = 0; row < art.rows; ++row) Report(DeviceName(device), row, art, m); + } +} + +// THE PRODUCTION CASE. This is the red-first gate for #3116: it calls the head +// projection the forward calls, over the primary's own captured head input, and +// requires the logits it returns to be the primary's BF16 words widened to F32 — +// not merely close to them. +TEST_CASE("qwen3 MoE LM-head BF16 boundary: the production projection mirrors the primary" * + doctest::skip(HeadFixtureAbsent())) { + // Same decorator, same `--no-skip` guard as the replay case above. + REQUIRE_MESSAGE(!HeadFixtureAbsent(), + "VT_MOE_HEAD_FIXTURE must name the directory holding " + "L33-C2-R0-head-6.{json,hidden.bin,weight.bin,logits.bin}"); + const HeadArtifact art = LoadArtifact(std::getenv("VT_MOE_HEAD_FIXTURE"), "L33-C2-R0", 6); + const std::vector primary = Widen(art.logits_bf16); + for (const Device device : Devices()) { + CAPTURE(DeviceName(device)); + QueueHandle queue(device); + vt::Backend& backend = vt::GetBackend(device.type); + { + DeviceScope scope(device); + DeviceBf16 hidden(backend, queue.q, art.hidden_bf16, {art.rows, art.hidden}); + DeviceBf16 weight(backend, queue.q, TransposedWeight(art), {art.hidden, art.vocab}); + vllm::dense_attn::Dev d{backend, queue.q}; + vllm::dense_attn::DBuf logits = vllm::lm_head::Project(d, hidden.t, weight.t, /*tied=*/false); + backend.Synchronize(queue.q); + + // The views and host copies downstream stay valid: F32, [rows, vocab]. + REQUIRE(logits.t().dtype == DType::kF32); + REQUIRE(logits.t().rank == 2); + REQUIRE(logits.t().shape[0] == art.rows); + REQUIRE(logits.t().shape[1] == art.vocab); + + std::vector got(primary.size()); + logits.Download(d, got.data()); + size_t differing = 0, first_diff = 0; + for (size_t i = 0; i < got.size(); ++i) { + if (got[i] != primary[i]) { + if (differing == 0) first_diff = i; + ++differing; + } + } + double worst = 0.0; + for (size_t i = 0; i < got.size(); ++i) + worst = std::max(worst, std::fabs(static_cast(got[i]) - primary[i])); + std::cout << "[lm-head boundary] device=" << DeviceName(device) + << " production projection vs primary: differing=" << differing << "/" + << got.size() << " max_abs=" << worst; + if (differing != 0) + std::cout << " first_diff_index=" << first_diff << " got=" << got[first_diff] + << " primary=" << primary[first_diff]; + std::cout << std::endl; + + CHECK(differing == 0); + for (int64_t row = 0; row < art.rows; ++row) { + const size_t base = static_cast(row * art.vocab); + std::vector row_values(got.begin() + static_cast(base), + got.begin() + static_cast(base + art.vocab)); + const std::vector order = TopOrder(row_values, art.vocab, 1); + const std::vector expected = TopOrder( + std::vector(primary.begin() + static_cast(base), + primary.begin() + static_cast(base + art.vocab)), + art.vocab, 1); + CAPTURE(row); + CHECK(order[0] == expected[0]); + } + } + } +} + +// ─── The production entry point ───────────────────────────────────────────── + +// A deterministic BF16 weight, so this case needs no capture and no checkpoint. +vllm::OwnedTensor SyntheticBf16(const std::vector& shape, uint64_t seed) { + int64_t n = 1; + for (int64_t s : shape) n *= s; + vllm::OwnedTensor t; + t.dtype = DType::kBF16; + t.rank = static_cast(shape.size()); + for (size_t i = 0; i < shape.size(); ++i) t.shape[i] = shape[i]; + std::vector bytes(static_cast(n) * sizeof(uint16_t)); + auto* words = reinterpret_cast(bytes.data()); + uint64_t s = seed; + for (int64_t i = 0; i < n; ++i) { + s = s * 6364136223846793005ULL + 1442695040888963407ULL; + s ^= s >> 33; + words[i] = vt::F32ToBF16(static_cast((s >> 40) & 0xFFFF) / 32768.0f - 1.0f); + } + t.bytes = vllm::OwnedBytes(std::move(bytes)); + return t; +} + +// THE CALL-SITE CASE, and the reason it uses a degenerate config. With ZERO +// decoder layers the production entry point reduces to embed -> final RMSNorm -> +// lm_head, which is the smallest shape in which `Qwen3MoeModel::Forward` reaches +// the head call site. The focused seam cases above cannot see a reverted call +// site (they call `lm_head::Project` themselves); this one can, and it needs no +// checkpoint, no GPU-only build and no fixture directory. +// +// WHAT IT ASSERTS is the boundary the primary defines, read off the returned +// values: every logit the production forward hands to the sampler must be a BF16 +// word widened to F32, i.e. its low 16 mantissa bits are zero. Nothing else in +// this path is allowed to round the logits, so the property is the boundary's +// fingerprint rather than a coincidence of the fixture. +// +// THIS CASE IS UNDECORATED AND UNCONDITIONAL: it carries no `doctest::skip`, so +// it runs and is reported in the fixture-absent run too, where the two cases +// above skip. That is what makes it the gate a reverted call site has to redden. +TEST_CASE("qwen3 MoE LM-head BF16 boundary: the production forward returns BF16 logits") { + constexpr int64_t kH = 128; + constexpr int64_t kV = 128; + for (const Device device : Devices()) { + CAPTURE(DeviceName(device)); + QueueHandle queue(device); + vllm::Qwen3MoeWeights weights; + weights.tie_word_embeddings = false; + weights.embed_tokens = SyntheticBf16({kV, kH}, 0x51ed01); + weights.final_norm = SyntheticBf16({kH}, 0x51ed02); + weights.lm_head = SyntheticBf16({kH, kV}, 0x51ed03); + vllm::HfConfig config; + config.model_type = "qwen3_moe"; + config.architectures = {"Qwen3MoeForCausalLM"}; + config.hidden_size = kH; + config.vocab_size = kV; + config.num_hidden_layers = 0; // embed -> final RMSNorm -> lm_head + config.num_attention_heads = 1; + config.num_key_value_heads = 1; + config.head_dim = kH; + config.intermediate_size = kH; + config.moe_intermediate_size = kH; + config.rms_norm_eps = 1e-6; + config.rotary_dim = 0; + + const std::vector logits = vllm::Qwen3MoeModel::Forward( + /*token_ids=*/{7}, /*positions=*/{0}, vllm::v1::CommonAttentionMetadata{}, + /*attn_kv=*/{}, weights, config, queue.q); + + REQUIRE(logits.size() == static_cast(kV)); + size_t not_bf16 = 0; + bool any_nonzero = false; + for (const float value : logits) { + uint32_t bits = 0; + std::memcpy(&bits, &value, sizeof(bits)); + not_bf16 += (bits & 0xFFFFu) != 0; + any_nonzero = any_nonzero || value != 0.0f; + } + std::cout << "[lm-head boundary] device=" << DeviceName(device) + << " production forward: logits=" << logits.size() + << " not_bf16_representable=" << not_bf16 << " any_nonzero=" << any_nonzero + << std::endl; + // The projection ran... + CHECK(any_nonzero); + // ...and its result is the primary's dtype, not a full F32 quotient. + CHECK(not_bf16 == 0); + } +} diff --git a/tests/vllm/models/rocm_moe_fixture.h b/tests/vllm/models/rocm_moe_fixture.h new file mode 100644 index 0000000000..ab88081bc8 --- /dev/null +++ b/tests/vllm/models/rocm_moe_fixture.h @@ -0,0 +1,134 @@ +// vllm.cpp original. BACKEND-ROCM-BF16-MOE production fixture (#3094). +// Adapted from test_moe_async_device_ids.cpp at 6db4bef906859e864c82523c01107473f7dcca29. +// Tensor order, RNG wraparound, conversion, and scales are preserved. +#pragma once +#include +#include +#include +#include +#include +#include +#include "vt/dtype.h" +namespace rocm_moe_fixture { +struct Fx { + std::string name, dtype; + std::vector shape; + std::string bytes; +}; + +std::string U64Le(uint64_t v) { + std::string s(8, '\0'); + for (int i = 0; i < 8; ++i) + s[static_cast(i)] = static_cast((v >> (8 * i)) & 0xff); + return s; +} +int64_t NumEl(const std::vector& s) { + int64_t n = 1; + for (int64_t d : s) n *= d; + return n; +} +std::string Bf16Bytes(size_t n, int seed, float scale) { + std::string s(n * 2, '\0'); + uint32_t r = static_cast(seed) * 2654435761u + 1u; + for (size_t i = 0; i < n; ++i) { + r = r * 1664525u + 1013904223u; + const float u = static_cast(r >> 8) / static_cast(1u << 24); + const uint16_t bf = vt::F32ToBF16((u - 0.5f) * 2.0f * scale); + s[i * 2] = static_cast(bf & 0xff); + s[i * 2 + 1] = static_cast((bf >> 8) & 0xff); + } + return s; +} +Fx Bf16(const std::string& n, std::vector sh, int seed, float scale = 0.08f) { + return {n, "BF16", sh, Bf16Bytes(static_cast(NumEl(sh)), seed, scale)}; +} +std::string BuildSt(const std::vector& ts) { + nlohmann::json hdr = nlohmann::json::object(); + std::string data; + for (const Fx& t : ts) { + const size_t start = data.size(); + data += t.bytes; + hdr[t.name] = {{"dtype", t.dtype}, + {"shape", t.shape}, + {"data_offsets", {start, data.size()}}}; + } + const std::string header = hdr.dump(); + return U64Le(header.size()) + header + data; +} + +constexpr int64_t kH = 128, kL = 2, kHq = 1, kHkv = 1, kDh = 128, kV = 128; +constexpr int64_t kE = 4, kTopK = 2, kI = 128; + +std::string ConfigJson() { + nlohmann::json j; + j["architectures"] = std::vector{"Qwen3MoeForCausalLM"}; + j["model_type"] = "qwen3_moe"; + j["hidden_size"] = kH; + j["num_hidden_layers"] = kL; + j["num_attention_heads"] = kHq; + j["num_key_value_heads"] = kHkv; + j["head_dim"] = kDh; + j["intermediate_size"] = kI; + j["moe_intermediate_size"] = kI; + j["shared_expert_intermediate_size"] = 0; + j["num_experts"] = kE; + j["num_experts_per_tok"] = kTopK; + j["vocab_size"] = kV; + j["max_position_embeddings"] = 256; + j["rms_norm_eps"] = 1e-6; + j["rope_theta"] = 10000000.0; + j["tie_word_embeddings"] = false; + j["attention_bias"] = false; + j["torch_dtype"] = "bfloat16"; + j["norm_topk_prob"] = true; + j["decoder_sparse_step"] = 1; + j["mlp_only_layers"] = nlohmann::json::array(); + j["hidden_act"] = "silu"; + j["bos_token_id"] = 1; + j["eos_token_id"] = 127; + j["pad_token_id"] = 0; + return j.dump(2); +} + +std::vector BuildTensors() { + std::vector v; + int s = 7; + v.push_back(Bf16("model.embed_tokens.weight", {kV, kH}, s++)); + v.push_back(Bf16("model.norm.weight", {kH}, s++, 0.5f)); + v.push_back(Bf16("lm_head.weight", {kV, kH}, s++)); + for (int64_t l = 0; l < kL; ++l) { + const std::string b = "model.layers." + std::to_string(l) + "."; + const std::string sa = b + "self_attn."; + const std::string mlp = b + "mlp."; + v.push_back(Bf16(b + "input_layernorm.weight", {kH}, s++, 0.5f)); + v.push_back(Bf16(b + "post_attention_layernorm.weight", {kH}, s++, 0.5f)); + v.push_back(Bf16(sa + "q_proj.weight", {kHq * kDh, kH}, s++)); + v.push_back(Bf16(sa + "k_proj.weight", {kHkv * kDh, kH}, s++)); + v.push_back(Bf16(sa + "v_proj.weight", {kHkv * kDh, kH}, s++)); + v.push_back(Bf16(sa + "o_proj.weight", {kH, kHq * kDh}, s++)); + v.push_back(Bf16(sa + "q_norm.weight", {kDh}, s++, 0.5f)); + v.push_back(Bf16(sa + "k_norm.weight", {kDh}, s++, 0.5f)); + v.push_back(Bf16(mlp + "gate.weight", {kE, kH}, s++)); + for (int64_t e = 0; e < kE; ++e) { + const std::string ex = mlp + "experts." + std::to_string(e) + "."; + v.push_back(Bf16(ex + "gate_proj.weight", {kI, kH}, s++)); + v.push_back(Bf16(ex + "up_proj.weight", {kI, kH}, s++)); + v.push_back(Bf16(ex + "down_proj.weight", {kH, kI}, s++)); + } + } + return v; +} + +inline void Export(const std::filesystem::path& directory) { + std::filesystem::create_directories(directory); + const std::string weights = BuildSt(BuildTensors()); + std::ofstream(directory / "model.safetensors", std::ios::binary).write( + weights.data(), static_cast(weights.size())); + std::ofstream(directory / "config.json") << ConfigJson() << '\n'; +} +inline std::vector Prompt(int length, int request) { + std::vector result; + for (int i = 0; i < length; ++i) result.push_back(1 + ((11 + 29 * request + 17 * i) % 126)); + return result; +} +} // namespace rocm_moe_fixture diff --git a/tests/vllm/models/rocm_moe_oracle.py b/tests/vllm/models/rocm_moe_oracle.py new file mode 100644 index 0000000000..1d07f813c7 --- /dev/null +++ b/tests/vllm/models/rocm_moe_oracle.py @@ -0,0 +1,215 @@ +"""Pinned vLLM production oracle for BACKEND-ROCM-BF16-MOE (#3094). + +This is an evidence harness, not a denominator timing run. It keeps the default +engine compilation and backend policy and records the executing native tensors. +""" +import argparse +import hashlib +import json +import os +from pathlib import Path + + +def sha256(path): + return hashlib.sha256(path.read_bytes()).hexdigest() + + +def inspect_and_capture(model, output): + import importlib + import torch + + output = Path(output) + metadata = [] + for name, module in model.named_modules(): + method = getattr(module, "quant_method", None) + backend = getattr(method, "unquantized_backend", None) + if backend is None: + continue + record = {"module": name, "backend": str(backend), + "expert_class": str(getattr(method, "experts_cls", None)), + "weights": {}} + for weight_name in ("w13_weight", "w2_weight"): + weight = getattr(module, weight_name, None) + if weight is not None: + record["weights"][weight_name] = { + "shape": list(weight.shape), "stride": list(weight.stride()), + "dtype": str(weight.dtype), "element_size": weight.element_size()} + metadata.append(record) + assert metadata, "No executing unquantized MoE method found" + assert all(item["backend"] == "UnquantizedMoeBackend.TRITON" for item in metadata), metadata + + module = importlib.import_module("vllm.model_executor.layers.fused_moe.experts.triton_moe") + original = module.invoke_fused_moe_triton_kernel + count = [0] + + def capture(*args, **kwargs): + result = original(*args, **kwargs) + if count[0] >= 32 or torch.cuda.is_current_stream_capturing(): + return result + index = count[0] + count[0] += 1 + names = ("A", "B", "C", "A_scale", "B_scale", "topk_weights", + "sorted_token_ids", "expert_ids", "num_tokens_post_padded") + record = {"index": index, "mul_routed_weight": args[9], + "top_k": args[10], "config": args[11], + "compute_type": str(kwargs.get("compute_type")), "tensors": {}} + tensors = {} + for name, value in zip(names, args): + if isinstance(value, torch.Tensor): + record["tensors"][name] = {"shape": list(value.shape), + "stride": list(value.stride()), "dtype": str(value.dtype), + "element_size": value.element_size()} + tensors[name] = value.detach().cpu() + torch.save(tensors, output / f"kernel-{index:03d}.pt") + (output / f"kernel-{index:03d}.json").write_text(json.dumps(record, indent=2) + "\n") + return result + + module.invoke_fused_moe_triton_kernel = capture + return metadata + + +def install_schedule_capture(worker): + """Observe CPU scheduler metadata, including calls that replay a GPU graph.""" + assert worker.use_v2_model_runner + assert type(worker.model_runner).__module__ == "vllm.v1.worker.gpu.model_runner" + worker._moe_schedule_records = [] + original = worker.execute_model + + def capture(scheduler_output): + record = {"total_tokens": scheduler_output.total_num_scheduled_tokens, + "scheduled": dict(scheduler_output.num_scheduled_tokens), + "new": [{"request_id": request.req_id, + "prompt_token_ids": request.prompt_token_ids, + "computed_tokens": request.num_computed_tokens} + for request in scheduler_output.scheduled_new_reqs]} + result = original(scheduler_output) + # gpu/model_runner.py:1799-1811 publishes the executed V2 batch here. + # The legacy runner's input_batch member does not exist on this path. + record["runner_class"] = str(type(worker.model_runner)) + record["runner_request_ids"] = ( + list(worker.model_runner.execute_model_state.input_batch.req_ids) + if record["total_tokens"] else []) + worker._moe_schedule_records.append(record) + return result + + worker.execute_model = capture + return True + + +def take_schedule_capture(worker): + records = worker._moe_schedule_records + worker._moe_schedule_records = [] + return records + + +def main(): + parser = argparse.ArgumentParser() + parser.add_argument("--fixture", type=Path, required=True) + parser.add_argument("--output", type=Path, required=True) + args = parser.parse_args() + args.output.mkdir(parents=True, exist_ok=True) + import torch + import vllm + from vllm import LLM, SamplingParams + + properties = torch.cuda.get_device_properties(0) + assert "gfx1100" in properties.gcnArchName + assert vllm.__version__ == "0.28.1rc1.dev132+ge126687a9", vllm.__version__ + artifacts = {name: {"bytes": (args.fixture / name).stat().st_size, + "sha256": sha256(args.fixture / name)} + for name in ("config.json", "model.safetensors")} + llm = LLM(model=str(args.fixture), tokenizer=None, skip_tokenizer_init=True, + dtype="bfloat16", seed=7, max_model_len=256, max_num_seqs=2, + max_num_batched_tokens=256, kv_cache_memory_bytes=128 * 1024 * 1024, + max_logprobs=128) # Capture-only: retain the complete vocabulary. + metadata = llm.apply_model(lambda model: inspect_and_capture(model, str(args.output))) + assert llm.collective_rpc(install_schedule_capture) == [True] + print(json.dumps({"selected_after_load": metadata}), flush=True) + params = SamplingParams(temperature=0, top_p=1, top_k=-1, min_p=0, + repetition_penalty=1, presence_penalty=0, + frequency_penalty=0, seed=7, ignore_eos=True, + min_tokens=8, max_tokens=8, logprobs=128) + runs = [] + # A prefill of 33 tokens executes the module after capture hooks are installed. + for length in (33, 1, 3): + for concurrency in (1, 2): + prompts = [{"prompt_token_ids": [1 + ((11 + 29 * r + 17 * i) % 126) + for i in range(length)]} + for r in range(concurrency)] + reference = None + for repeat in range(3): + # Prefix caching is a production default. Reset request cache state + # between repeats to keep the compared workloads independent. + llm.reset_prefix_cache() + # Pinned llm.py:727-731 documents this simultaneous-arrival + # sequence. Level zero pauses scheduling while accepting work. + llm.sleep(level=0) + request_ids = llm.enqueue(prompts, params, use_tqdm=False) + # input_processor.py:260-279 assigns internal IDs. The output + # processor keeps their authoritative external mapping until + # completion (output_processor.py:543-572). + states = llm.llm_engine.output_processor.request_states + external_ids = [states[request_id].external_req_id for request_id in request_ids] + assert [states[request_id].prompt_token_ids for request_id in request_ids] == [ + item["prompt_token_ids"] for item in prompts] + llm.wake_up(tags=["scheduling"]) + outputs = llm.wait_for_completion(use_tqdm=False) + schedule, = llm.collective_rpc(take_schedule_capture) + schedule_path = args.output / f"schedule-L{length}-C{concurrency}-R{repeat}.json" + schedule_path.write_text(json.dumps({"request_ids": request_ids, + "external_request_ids": external_ids, + "output_request_ids": [item.request_id for item in outputs], + "schedule": schedule}, indent=2) + "\n") + assert [item.request_id for item in outputs] == external_ids + assert [item.prompt_token_ids for item in outputs] == [ + item["prompt_token_ids"] for item in prompts] + active = [item for item in schedule if item["total_tokens"]] + assert len(active) == 8, active + assert active[0]["total_tokens"] == length * concurrency, active[0] + assert [item["prompt_token_ids"] for item in active[0]["new"]] == [ + item["prompt_token_ids"] for item in prompts], active[0] + assert all(item["computed_tokens"] == 0 for item in active[0]["new"]) + internal_ids = [item["request_id"] for item in active[0]["new"]] + assert internal_ids == request_ids + for step, item in enumerate(active): + assert list(item["scheduled"]) == internal_ids, item + assert item["runner_request_ids"] == internal_ids, item + assert list(item["scheduled"].values()) == [length if step == 0 else 1] * concurrency, item + tokens = [list(item.outputs[0].token_ids) for item in outputs] + assert all(len(row) == 8 for row in tokens) + if reference is None: + reference = tokens + assert tokens == reference + runs.append({"length": length, "concurrency": concurrency, + "repeat": repeat, "tokens": tokens, + "request_ids": request_ids, "external_request_ids": external_ids, + "schedule": schedule, + "logprobs": [[{str(k): v.logprob for k, v in step.items()} + for step in item.outputs[0].logprobs] + for item in outputs]}) + result = {"pin": "e126687a9a828d513c01a07cd69f025f27d63280", + "vllm": vllm.__version__, "torch": torch.__version__, + "hip": torch.version.hip, "device": properties.name, + "architecture": properties.gcnArchName, "artifacts": artifacts, + "selected": metadata, "runs": runs, + "purpose": "instrumented production correctness and selection evidence"} + (args.output / "production.json").write_text(json.dumps(result, indent=2) + "\n") + cache = Path(os.environ["TRITON_CACHE_DIR"]) + kernels = [{"path": str(path), "bytes": path.stat().st_size, "sha256": sha256(path)} + for path in sorted(cache.rglob("*")) + if path.is_file() and path.suffix in (".amdgcn", ".llir", ".ttir", ".ttgir", ".json")] + (args.output / "generated-kernels.json").write_text(json.dumps(kernels, indent=2) + "\n") + assert kernels, "No generated Triton kernel evidence" + captured = [json.loads(path.read_text()) for path in sorted(args.output.glob("kernel-*.json"))] + assert any(not item["mul_routed_weight"] for item in captured), "No executing gate/up capture" + assert any(item["mul_routed_weight"] for item in captured), "No executing weighted-down capture" + for item in captured: + assert item["compute_type"] == "bf16", item + for name in ("A", "B", "C"): + assert item["tensors"][name]["dtype"] == "torch.bfloat16", item + print(json.dumps({"selected": metadata, "runs": len(runs), "artifacts": artifacts, + "generated_files": len(kernels)})) + + +if __name__ == "__main__": + main() diff --git a/tests/vllm/models/rocm_moe_upstream.py b/tests/vllm/models/rocm_moe_upstream.py new file mode 100644 index 0000000000..439ef61090 --- /dev/null +++ b/tests/vllm/models/rocm_moe_upstream.py @@ -0,0 +1,139 @@ +"""Run and export the pinned upstream single-device BF16 MoE cases (#3094). + +The original test function executes without reduced shapes, changed seeds, +tolerances, references, padding modes, or graph settings. Only its inputs and +outputs are exported for the independent C++ provider comparison. The existing +local pointer-array ABI materializes logical [K,N] matrices contiguously. Padded +oracle calls remain distinct upstream cases; they do not test native padded strides. +""" +import argparse +import hashlib +import importlib +import itertools +import json +import sys +from pathlib import Path + +PIN = "e126687a9a828d513c01a07cd69f025f27d63280" +SHAPES = [(1, 128, 128), (1, 2048, 128), (33, 2048, 128), + (32768, 2048, 511), (40000, 1024, 1024)] +CASES = list(itertools.product(SHAPES, (8, 64, 192), (2, 6), (False, True))) + + +def sha256(path): + digest = hashlib.sha256() + with path.open("rb") as source: + for block in iter(lambda: source.read(1024 * 1024), b""): + digest.update(block) + return digest.hexdigest() + + +def main(): + parser = argparse.ArgumentParser() + parser.add_argument("--source", type=Path, required=True) + parser.add_argument("--output", type=Path, required=True) + parser.add_argument("--case", type=int, required=True, choices=range(len(CASES))) + args = parser.parse_args() + args.output.mkdir(parents=True, exist_ok=True) + import torch + import vllm + assert vllm.__version__ == "0.28.1rc1.dev132+ge126687a9" + # Import the installed pinned runtime first, then its exact source tests. + sys.path.insert(0, str(args.source)) + upstream = importlib.import_module("tests.kernels.moe.test_moe") + import pytest + from vllm.v1.worker.workspace import init_workspace_manager, reset_workspace_manager + init_workspace_manager(torch.device("cuda:0")) + shape, experts, top_k, padding = CASES[args.case] + m, n, k = shape + files = {} + runtime_calls = [] + captured_reference = [False] + + def save(name, tensor): + cpu = tensor.detach().contiguous().cpu() + raw = cpu.view(torch.uint16).numpy() if cpu.dtype == torch.bfloat16 else cpu.numpy() + path = args.output / name + raw.tofile(path) + files[name] = {"shape": list(tensor.shape), "source_stride": list(tensor.stride()), + "export_stride": list(cpu.stride()), + "dtype": str(tensor.dtype), "bytes": path.stat().st_size, + "sha256": sha256(path)} + + original = upstream.run_moe_test + + def export_run(baseline, moe_fn, **kwargs): + outputs = [] + inputs = [] + + def capture(*positional, **keywords): + inputs[:] = positional + output = moe_fn(*positional, **keywords) + outputs[:] = [output] + return output + + reference = original(baseline, capture, **kwargs) + a, w1, w2, score = inputs[:4] + pad = 128 if padding else 0 + assert list(w1.stride()) == [2 * n * (k + pad), k + pad, 1] + assert list(w2.stride()) == [k * (n + pad), n + pad, 1] + runtime_calls.append({"implementation": getattr(moe_fn, "__name__", str(moe_fn)), + "padding": kwargs["padding"], + "w1_stride": list(w1.stride()), + "w2_stride": list(w2.stride()), + "use_compile": kwargs.get("use_compile", False), + "use_cudagraph": kwargs.get("use_cudagraph", False), + "atol": 0.02, "rtol": 0}) + if not captured_reference[0]: + captured_reference[0] = True + routes, ids, _ = upstream.fused_topk(a, score, top_k, False) + save("a.bf16", a) + save("score.bf16", score) + # Preserve the actual original padded storage as evidence. The + # dense exports below are the explicit loader/ABI adaptation. + assert w1.storage_offset() == w2.storage_offset() == 0 + save("oracle-w1-storage.bf16", w1.as_strided((experts, 2 * n, k + pad), w1.stride())) + save("oracle-w2-storage.bf16", w2.as_strided((experts, k, n + pad), w2.stride())) + # Required storage adaptation: [E,N,K] checkpoint matrices become + # one [K,N] Matmul-B pointer per expert. Logical values stay BF16. + save("gate.bf16", w1[:, :n, :].transpose(1, 2)) + save("up.bf16", w1[:, n:, :].transpose(1, 2)) + save("down.bf16", w2.transpose(1, 2)) + save("ids.i32", ids.to(torch.int32).reshape(-1)) + save("routes.f32", routes.float().reshape(-1)) + save("reference.bf16", reference) + # The last call is modular Triton, the selected production backend. + save("triton.bf16", outputs[0]) + return reference + + upstream.run_moe_test = export_run + try: + with pytest.MonkeyPatch.context() as monkeypatch: + upstream.test_fused_moe(m=m, n=n, k=k, e=experts, topk=top_k, + ep_size=1, dtype=torch.bfloat16, padding=padding, + use_td=False, monkeypatch=monkeypatch, workspace_init=None) + torch.cuda.synchronize() + finally: + upstream.run_moe_test = original + reset_workspace_manager() + assert len(runtime_calls) == 3 + expected_graph = n >= 1024 and k >= 1024 + assert all(call["use_cudagraph"] == expected_graph for call in runtime_calls[1:]) + record = {"case": args.case, "pin": PIN, "M": m, "N": n, "K": k, + "experts": experts, "top_k": top_k, "padding": padding, + "seed": 7, "renormalize": False, "use_compile": False, + "use_cudagraph": expected_graph, "atol": 0.02, "rtol": 0, + "upstream_test_sha256": sha256(args.source / "tests/kernels/moe/test_moe.py"), + "upstream_calls": runtime_calls, "files": files, + "storage_adaptation": "Actual padded upstream weight storage is retained. Native expert pointers address contiguous logical [K,N] matrices, as the existing shared ABI requires. This is differential padded-oracle coverage, not native padded-stride coverage.", + "exclusions": { + "ep_size=4": "This row adds no distributed execution path", + "use_td=true": "gfx1100 is excluded by fused_moe/utils.py:665-684"}} + (args.output / "case.json").write_text(json.dumps(record, indent=2) + "\n") + print(json.dumps({"case": args.case, "shape": shape, "experts": experts, + "top_k": top_k, "padding": padding, + "upstream": "PASS", "graph": expected_graph})) + + +if __name__ == "__main__": + main() diff --git a/tests/vllm/models/test_rocm_moe_bf16.cpp b/tests/vllm/models/test_rocm_moe_bf16.cpp new file mode 100644 index 0000000000..87027449d4 --- /dev/null +++ b/tests/vllm/models/test_rocm_moe_bf16.cpp @@ -0,0 +1,2238 @@ +// vllm.cpp original. Native ROCm BF16 MoE production gate (#3094). +// The checkpoint enters through ModelRegistry::Load and every step through Forward. +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include "rocm_moe_fixture.h" +#include "support/rocm_moe_reference_set.h" +#include "support/residual_norm_fixture.h" +#include "support/residual_norm_test.h" +#include "vllm/model_executor/model_loader/safetensors_reader.h" +#include "vllm/model_executor/models/dense_attn_block.h" +#include "vllm/model_executor/models/model_registry.h" +#include "vllm/model_executor/models/qwen3_5.h" +#include "vllm/transformers_utils/hf_config.h" +#include "vt/backend.h" +#include "vt/op_provider.h" +#include "vt/ops.h" + +namespace { +using namespace rocm_moe_fixture; + +// The production fixture this file's first two cases replay. Declared up here +// because the case decorators below are evaluated in file order; see the +// "env-gated cases" block for why those cases skip instead of exiting. +bool FixtureAbsent() { + const char* e = std::getenv("VT_ROCM_MOE_FIXTURE"); + return e == nullptr || e[0] == '\0'; +} + +vt::MoeRouterTopKFn native_router = nullptr; +vt::RmsNormFn native_norm = nullptr; +vt::ResidualRmsNormFn native_expression = nullptr; +bool observe_residual = false; +int observed_norm_calls = 0; +std::vector observed_attention, observed_residual, observed_gamma, observed_post; +std::vector observed_first; +struct ExpressionRecord { + vt::ResidualRmsNormArgs args; + std::vector a, base, delta, gamma, norm, residual; + uintptr_t attention_pointer = 0; +}; +std::vector observed_expressions; +std::vector ReadRow(vt::Queue& q, const vt::Tensor& tensor) { + REQUIRE(tensor.dtype == vt::DType::kBF16); + REQUIRE(tensor.Numel() >= 128); + std::vector row(128); + auto& backend = vt::GetBackend(q.device); + backend.Copy(q, row.data(), tensor.data, row.size() * sizeof(uint16_t)); + backend.Synchronize(q); + return row; +} +void CaptureResidualNorm(vt::Queue& q, vt::Tensor& out, const vt::Tensor& x, + const vt::Tensor& weight, const vt::RmsNormArgs& args, + vt::Tensor* residual) { + const bool legacy = vt::GetBackend(q.device).GetResidualNormPolicy() == + vt::ResidualNormPolicy::kMaterialized; + const bool capture = observe_residual && legacy && observed_norm_calls++ == 3; + if (capture) { + REQUIRE(residual != nullptr); + observed_attention = ReadRow(q, x); + observed_residual = ReadRow(q, *residual); + observed_gamma = ReadRow(q, weight); + } + native_norm(q, out, x, weight, args, residual); + if (capture) observed_post = ReadRow(q, out); + if (observe_residual && !legacy && observed_first.empty()) observed_first = ReadRow(q, out); +} +void CaptureExpression(vt::Queue& q, vt::Tensor& out, const vt::Tensor& a, + const vt::Tensor& base, const vt::Tensor* delta, + const vt::Tensor& weight, const vt::ResidualRmsNormArgs& args, + vt::Tensor* residual_out) { + const bool first = observe_residual && observed_norm_calls++ == 0; + ExpressionRecord record; + if (observe_residual) { + record.args = args; + record.a = ReadRow(q, a); record.base = ReadRow(q, base); + record.gamma = ReadRow(q, weight); + record.attention_pointer = reinterpret_cast(a.data); + if (delta != nullptr) record.delta = ReadRow(q, *delta); + } + if (first) { + observed_attention = ReadRow(q, a); + observed_residual = ReadRow(q, base); + observed_gamma = ReadRow(q, weight); + } + native_expression(q, out, a, base, delta, weight, args, residual_out); + if (first) observed_post = ReadRow(q, out); + if (observe_residual) { + record.norm = ReadRow(q, out); + if (residual_out != nullptr) record.residual = ReadRow(q, *residual_out); + CHECK(ReadRow(q, a) == record.a); + CHECK(ReadRow(q, weight) == record.gamma); + if (!args.descriptor.residual_alias_base) CHECK(ReadRow(q, base) == record.base); + if (delta != nullptr && !args.descriptor.output_alias_delta) CHECK(ReadRow(q, *delta) == record.delta); + observed_expressions.push_back(std::move(record)); + } +} +nlohmann::json* current_routes = nullptr; +void CaptureRoutes(vt::Queue& q, vt::Tensor& weights, vt::Tensor& indices, + const vt::Tensor& logits, const vt::MoeRouterTopKArgs& args, + const vt::Tensor* bias) { + native_router(q, weights, indices, logits, args, bias); + if (current_routes == nullptr) return; + std::vector ids(static_cast(indices.Numel())); + auto& backend = vt::GetBackend(q.device); + backend.Copy(q, ids.data(), indices.data, ids.size() * sizeof(int32_t)); + backend.Synchronize(q); + current_routes->push_back(ids); +} +vt::Tensor Tensor(void* data, vt::DType dtype, vt::Device device, + std::initializer_list dimensions) { + vt::Tensor tensor; + tensor.data = data; + tensor.dtype = dtype; + tensor.device = device; + tensor.rank = static_cast(dimensions.size()); + int axis = 0; + for (int64_t size : dimensions) tensor.shape[axis++] = size; + int64_t stride = 1; + for (int i = tensor.rank - 1; i >= 0; --i) { + tensor.stride[i] = stride; + stride *= tensor.shape[i]; + } + return tensor; +} +struct Buffer { + vt::Backend& backend; + void* data; + Buffer(vt::Backend& b, size_t bytes) : backend(b), data(b.Alloc(bytes)) {} + ~Buffer() { backend.Free(data); } +}; +struct QueueGuard { + vt::Backend& backend; + vt::Queue queue; + explicit QueueGuard(vt::Backend& b) : backend(b), queue(b.CreateQueue()) {} + ~QueueGuard() { backend.DestroyQueue(queue); } +}; + +nlohmann::json Generate(const std::filesystem::path& fixture, int length, int concurrency, + int steps = 8) { + auto config = vllm::LoadHfConfig((fixture / "config.json").string()); + std::vector shards; + shards.push_back(vllm::SafetensorsFile::Open((fixture / "model.safetensors").string())); + auto model = vllm::ModelRegistry::Load(config, vllm::ModelSource::FromSafetensors(shards)); + auto& backend = vt::GetBackend(vt::DeviceType::kROCM); + QueueGuard queue(backend); + auto& q = queue.queue; + constexpr int64_t block = 64; + std::vector> allocations; + std::vector caches; + for (int layer = 0; layer < kL; ++layer) { + const size_t bytes = static_cast(concurrency * 2 * block * kHkv * kDh) * 2; + allocations.push_back(std::make_unique(backend, bytes)); + std::vector zeros(bytes / 2, 0); + backend.Copy(q, allocations.back()->data, zeros.data(), bytes); + vllm::PagedKvCache kv; + kv.data = allocations.back()->data; + kv.dtype = vt::DType::kBF16; + kv.num_blocks = concurrency; + kv.block_size = block; + kv.num_kv_heads = kHkv; + kv.head_size = kDh; + caches.push_back(kv); + backend.Synchronize(q); + } + std::vector gdn; + vllm::v1::GDNAttentionMetadata gdn_meta{}; + std::vector> generated(static_cast(concurrency)); + nlohmann::json logits = nlohmann::json::array(); + nlohmann::json routes = nlohmann::json::array(); + current_routes = &routes; + Buffer sampled(backend, static_cast(concurrency) * sizeof(int64_t)); + auto ids = Tensor(sampled.data, vt::DType::kI64, q.device, {concurrency}); + for (int step = 0; step < steps; ++step) { + const int count = step == 0 ? length : 1; + const int previous = step == 0 ? 0 : length + step - 1; + std::vector tokens, positions, gather; + vllm::v1::CommonAttentionMetadata attention; + attention.num_reqs = concurrency; + attention.num_actual_tokens = count * concurrency; + attention.query_start_loc = {0}; + for (int r = 0; r < concurrency; ++r) { + const auto prompt = Prompt(length, r); + for (int i = 0; i < count; ++i) { + tokens.push_back(step == 0 ? prompt[static_cast(i)] : generated[r].back()); + positions.push_back(previous + i); + attention.slot_mapping.push_back(r * block + previous + i); + } + attention.query_start_loc.push_back((r + 1) * count); + attention.seq_lens.push_back(previous + count); + attention.block_table_tensor.push_back(r); + gather.push_back((r + 1) * count - 1); + } + attention.query_start_loc_cpu = attention.query_start_loc; + attention.seq_lens_cpu = attention.seq_lens; + attention.max_query_len = count; + attention.max_seq_len = previous + count; + attention.block_table_num_cols = 1; + attention.causal = true; + vllm::ModelForwardInput input{tokens, positions, attention, gdn_meta, caches, gdn, + config, q, gather}; + input.num_reqs = concurrency; + input.gdn_state_slots = concurrency; + input.pure_decode = step > 0; + input.uniform_query_len = count; + const auto result = vllm::ModelRegistry::Forward(*model, input); + REQUIRE(result.on_device()); + REQUIRE(result.device_tensor.dtype == vt::DType::kF32); + vt::GreedyArgmax(q, ids, result.device_tensor); + std::vector host_ids(static_cast(concurrency)); + std::vector host_logits(static_cast(concurrency * kV)); + backend.Copy(q, host_ids.data(), sampled.data, host_ids.size() * sizeof(int64_t)); + backend.Copy(q, host_logits.data(), result.device_tensor.data, + host_logits.size() * sizeof(float)); + backend.Synchronize(q); + for (float value : host_logits) REQUIRE(std::isfinite(value)); + logits.push_back(host_logits); + for (int r = 0; r < concurrency; ++r) + generated[r].push_back(static_cast(host_ids[r])); + } + current_routes = nullptr; + return {{"length", length}, {"concurrency", concurrency}, {"tokens", generated}, + {"logits", logits}, {"expert_ids", routes}}; +} +} // namespace + +TEST_CASE("ROCm residual row zero matches the compiled primary through production forward" * + doctest::skip(FixtureAbsent())) { + using namespace residual_norm_fixture; + REQUIRE_MESSAGE(!FixtureAbsent(), + "VT_ROCM_MOE_FIXTURE must name the exported production checkpoint"); + const char* directory = std::getenv("VT_ROCM_MOE_FIXTURE"); + native_norm = reinterpret_cast( + vt::GetOp(vt::OpId::kRmsNorm, vt::DeviceType::kROCM)); + vt::RegisterOpProvider(vt::OpId::kRmsNorm, vt::DeviceType::kROCM, + {"test-residual-row-observer", 100, nullptr, + reinterpret_cast(static_cast(&CaptureResidualNorm))}); + native_expression = reinterpret_cast( + vt::GetOp(vt::OpId::kResidualRmsNorm, vt::DeviceType::kROCM)); + vt::RegisterOpProvider(vt::OpId::kResidualRmsNorm, vt::DeviceType::kROCM, + {"test-residual-expression-observer", 100, nullptr, + reinterpret_cast(static_cast(&CaptureExpression))}); + observe_residual = true; + observed_norm_calls = 0; + observed_first.clear(); observed_expressions.clear(); + Generate(directory, 33, 2, 1); + observe_residual = false; + CHECK(observed_attention == std::vector(kAttention.begin(), kAttention.end())); + CHECK(observed_residual == std::vector(kResidual.begin(), kResidual.end())); + CHECK(observed_gamma == std::vector(kGamma.begin(), kGamma.end())); + REQUIRE(observed_post.size() == kPostNorm.size()); + int different = 0; + for (size_t j = 0; j < kPostNorm.size(); ++j) { + CAPTURE(j); + different += observed_post[j] != kPostNorm[j]; + CHECK(observed_post[j] == kPostNorm[j]); + } + MESSAGE("Compared 128 production post-attention BF16 words with the compiled primary; differences: ", different); + CHECK(observed_first == std::vector(kFirstNorm.begin(), kFirstNorm.end())); + REQUIRE(observed_expressions.size() == 4); + auto checkpoint = vllm::SafetensorsFile::Open((std::filesystem::path(directory) / "model.safetensors").string()); + const std::vector weights{ + "model.layers.0.post_attention_layernorm.weight", "model.layers.1.input_layernorm.weight", + "model.layers.1.post_attention_layernorm.weight", "model.norm.weight"}; + for (size_t i = 0; i < observed_expressions.size(); ++i) { + CAPTURE(i); + const auto& record = observed_expressions[i]; + const bool triple = i % 2 == 1; + CHECK(record.args.descriptor.expression == (triple ? vt::ResidualNormExpr::kDeltaPlusAdd : vt::ResidualNormExpr::kAdd)); + CHECK(record.args.descriptor.materialize_residual == (i == 1)); + const auto& gamma = checkpoint.Get(weights[i]); + REQUIRE(gamma.dtype == "BF16"); + const auto* words = reinterpret_cast(gamma.data); + CHECK(record.gamma == std::vector(words, words + 128)); + std::vector expected_residual; + const auto expected = residual_norm_test::Reference(record.a, record.base, + triple ? &record.delta : nullptr, record.gamma, 1, 128, 128, + record.args.eps, &expected_residual); + CHECK(record.norm == expected); + if (i == 1) CHECK(record.residual == expected_residual); + else CHECK(record.residual.empty()); + if (triple) { + CHECK(record.a == observed_expressions[i - 1].a); + CHECK(record.base == observed_expressions[i - 1].base); + CHECK(record.attention_pointer == observed_expressions[i - 1].attention_pointer); + } + } + CHECK(observed_expressions[2].base == observed_expressions[1].residual); +} + +TEST_CASE("ROCm BF16 MoE enters native providers through the production registry" * + doctest::skip(FixtureAbsent())) { + REQUIRE_MESSAGE(!FixtureAbsent(), + "VT_ROCM_MOE_FIXTURE must name the exported production checkpoint"); + const char* directory = std::getenv("VT_ROCM_MOE_FIXTURE"); + const std::filesystem::path fixture(directory); + if (std::getenv("VT_ROCM_MOE_EXPORT_ONLY") != nullptr) { + Export(fixture); + REQUIRE(std::filesystem::file_size(fixture / "model.safetensors") > 0); + return; + } + const char* oracle_path = std::getenv("VT_ROCM_MOE_ORACLE"); + REQUIRE(oracle_path != nullptr); + nlohmann::json oracle; + std::ifstream(oracle_path) >> oracle; + // Test-only observation of the real device router. Expert providers retain + // their native registry entries and per-call selection accounting. + native_router = reinterpret_cast( + vt::GetOp(vt::OpId::kMoeRouterTopK, vt::DeviceType::kROCM)); + vt::RegisterOpProvider(vt::OpId::kMoeRouterTopK, vt::DeviceType::kROCM, + {"test-moe-route-capture", 100, nullptr, + reinterpret_cast(static_cast(&CaptureRoutes))}); + vt::EnableOpProviderCallStats(true); + for (auto op : {vt::OpId::kMoeGroupedGemmBf16GateUpSiluNative, + vt::OpId::kMoeGroupedGemmBf16Weighted, vt::OpId::kMoeCombinePreweighted}) + vt::ResetOpProviderStats(op, vt::DeviceType::kROCM); + nlohmann::json runs = nlohmann::json::array(); + std::set> selected_pairs; + for (int length : {1, 3, 33}) { + for (int concurrency : {1, 2}) { + nlohmann::json reference; + for (int repeat = 0; repeat < 3; ++repeat) { + auto run = Generate(fixture, length, concurrency); + if (repeat == 0) reference = run["tokens"]; + CHECK(run["tokens"] == reference); + bool compared = false; + for (const auto& expected : oracle.at("runs")) { + if (expected["length"] != length || expected["concurrency"] != concurrency || + expected["repeat"] != repeat) + continue; + CAPTURE(length); + CAPTURE(concurrency); + CAPTURE(repeat); + compared = true; + const auto native = run["tokens"].get>>(); + // The compared request count is the captured record's concurrency, not + // whatever the native run happened to return. + REQUIRE(native.size() == static_cast(concurrency)); + for (size_t request = 0; request < native.size(); ++request) { + CAPTURE(request); + // Every captured configuration of this workload forms the reference + // set for this request. Request 1 exists only in the records at + // concurrency 2, so a record without it contributes nothing. + const auto reference_set = rocm_moe_reference_set::Collect( + oracle.at("runs"), length, repeat, request); + REQUIRE(!reference_set.empty()); + const auto comparison = rocm_moe_reference_set::Compare( + reference_set, native[request], concurrency); + const int matched = comparison.pass() ? comparison.matched : -1; + // The gate reports the reference set, the matched configuration, the + // same-configuration outcome, and the reference's own disagreement + // positions. The same-configuration outcome is a report and not an + // assertion: the pinned primary disagrees with itself at length 33. + std::cout << "[production tokens] length " << length << " concurrency " + << concurrency << " repeat " << repeat << " request " << request + << ": reference set " + << rocm_moe_reference_set::Describe(reference_set) + << "; matched configuration " + << (matched < 0 + ? std::string("none") + : std::to_string(reference_set[static_cast(matched)] + .concurrency)) + << "; same-configuration match " + << (comparison.same_configuration_match ? "true" : "false") + << "; reference disagreement positions " + << rocm_moe_reference_set::Describe(comparison.disagreements) + << std::endl; + // Whole-sequence membership. A per-position mix of two reference + // sequences matches no member and fails here. + CHECK(comparison.pass()); + } + } + REQUIRE(compared); + REQUIRE(run["expert_ids"].size() == static_cast(8 * kL)); + for (const auto& call : run["expert_ids"]) { + const auto ids = call.get>(); + for (size_t pair = 0; pair < ids.size(); pair += 2) + selected_pairs.insert({ids[pair], ids[pair + 1]}); + } + run["repeat"] = repeat; + runs.push_back(std::move(run)); + } + } + } + if (const char* output = std::getenv("VT_ROCM_MOE_OUTPUT")) + std::ofstream(output) << runs.dump(2) << '\n'; + CHECK(selected_pairs.size() > 1); + CHECK(vt::OpRegistered(vt::OpId::kMoeGroupedGemmBf16, vt::DeviceType::kROCM)); + CHECK(vt::OpRegistered(vt::OpId::kMoeGroupedGemmBf16GateUpSilu, vt::DeviceType::kROCM)); + for (auto op : {vt::OpId::kMoeGroupedGemmBf16GateUpSiluNative, + vt::OpId::kMoeGroupedGemmBf16Weighted, vt::OpId::kMoeCombinePreweighted}) { + CAPTURE(static_cast(op)); + const auto stats = vt::GetOpProviderStats(op, vt::DeviceType::kROCM); + CHECK(vt::OpRegistered(op, vt::DeviceType::kROCM)); + CHECK(stats.selections > 0); + REQUIRE(stats.last_selected != nullptr); + CHECK(std::string(stats.last_selected) == vt::kNativeProviderName); + CHECK(stats.declines == 0); + CHECK(stats.fallbacks == 0); + CHECK(vt::GetOpProviderStats(op, vt::DeviceType::kCPU).selections == 0); + } + vt::EnableOpProviderCallStats(false); +} + +// --------------------------------------------------------------------------- +// BACKEND-ROCM-BF16-MOE (#3115): the native attention boundary. +// +// The first attention output of L33/C2/R0 layer 0 differs from the pinned +// primary's by 2902 of 8448 bf16 words, and the captured numbers alone do not +// separate two candidate causes: +// +// (a) the Q/K PREAMBLE. The primary keeps normalized Q/K in FP32 through RoPE +// and narrows at the store; the native block narrows to bf16 before RoPE +// (dense_attn_block.h:637-653). +// (b) the KERNEL ARITHMETIC. The primary's executing arm is Triton +// (prefix_prefill.py, selected because the ROCm custom predicate needs +// gqa_ratio >= 3 and this fixture has qg == 1) and narrows the softmax +// probabilities to the V dtype before tl.dot (prefix_prefill.py:471). The +// native ROCm arm keeps them f32 (rocm_paged_attn.hip:527-531), and its +// D128/QG1 dispatch has no query-length gate, so the 33-token prefill runs +// through the decode-geometry kernel (rocm_paged_attn.hip:445,2037-2041). +// +// This instrument separates them with native bytes. It (i) captures the native +// boundary from the unchanged production path, (ii) replays the PRIMARY's own +// captured Q/K/V through the same native op at the primary's cache geometry so +// the kernel term is measured on identical inputs, and (iii) runs the native +// preamble ops over the primary's captured `qkv` so the preamble term is +// measured on identical inputs too. +// +// TEST-ONLY. Every observer below is the provider-seam wrapper this file already +// uses for the residual row (CaptureResidualNorm, :228-237): read the live kernel +// with vt::GetOp, register a priority-100 replacement, forward to the saved +// pointer. Nothing under src/ or include/ changes and no attention arithmetic is +// touched. Everything is gated on VT_ATTN_DUMP; when that is unset no observer is +// registered at all and vt::GetOp dispatch is byte-identical to production. +// +// EAGER ONLY. The capture does device->host copies, which are illegal inside a +// graph capture, and vt::Backend carries no "is capturing" query (gap G5: +// include/vt/backend.h:244-262, include/vt/device.h:131-135). The workload that +// matters never captures: RocmPlatform::support_static_graph_mode() is false +// (src/vllm/platforms/rocm.cpp:91, include/vllm/platforms/interface.h:344), so +// the production registry path does not enter the graph driver. +// +// THE THREE OBSERVERS AND WHY EACH IS ENOUGH. +// * kPagedAttention — its arguments ARE the post-RoPE boundary: `query` (post- +// RoPE Q), `k_cache`/`v_cache` (what the kernel reads), `block_table`, +// `seq_lens`, `query_start_loc` and `out` (include/vt/ops.h:5424-5426, +// src/vt/ops.cpp:5038-5130). +// * kRopeFromCache — reached only where the recipe's Tier-0 composite runs +// (a backend with no fast realisation) or by the standalone fallback. On ROCm +// the fast realisation IS registered at this head +// (src/vt/rocm/rocm_ops.hip:331), so with the default VT_FUSED_CHAIN_ADOPT=1 +// the preamble is the fused op and this observer does not fire; with +// VT_FUSED_CHAIN_ADOPT=0 the bf16 branch calls vt::AttnQkNormRope, which +// dispatches to that same registered op (dense_attn_block.h:645-649), so it +// does not fire there either. Where it does fire it sees the post-norm +// PRE-RoPE q2/k2 on entry, the post-RoPE q3/k3 on return, and the cos|sin +// table that was actually read — all three without a product-file hook, +// because the composite reaches the op through GetOp. +// * kAttnQkNormRope — the recipe's fast realisation, which is the production +// preamble on this backend for BOTH settings of the adoption switch: ADOPT=1 +// through vt::FusedChain and ADOPT=0 through the hand-call. The runner reads +// whichever realization executed, so one case measures the boundary under +// either lever. +// * kRmsNorm — records every call in order so the pair immediately preceding +// the captured RopeFromCache call is identified by position in that call +// sequence rather than by shape (the layer input norms have the same shape). +// +// THE OPERATOR COMMAND. One binary, one process, one run; both cases below skip +// with exit 77 when their environment is absent, so an ordinary test run is +// unchanged. Run it under the host's recorded GPU mutex, with +// `VT_ATTN_DUMP=1` and these six values in the environment: +// +// VT_ROCM_MOE_FIXTURE = /home/vikash/.cache/rdna3-moe-impl/preserved/fixture +// VT_ATTN_PARITY_PRIMARY = /home/vikash/.cache/rdna3-moe-attention-operator/results +// VT_ATTN_PARITY_PREAMBLE = /home/vikash/.cache/residual-norm-impl/oracle-probe-v3/observed +// VT_ATTN_PARITY_OUT = a scratch directory for the dumped native bytes +// HIP_VISIBLE_DEVICES = 0 +// +// flock /home/vikash/gpu.lock -c ' build-attn-hip/tests/test_rocm_moe_bf16 --test-case="ROCm paged attention replays*"' +// +// The CPU-only half needs no device and only VT_ATTN_PARITY_PRIMARY: +// +// VT_ATTN_PARITY_PRIMARY= build-attn-hip/tests/test_rocm_moe_bf16 --test-case="Primary attention capture decodes*" +// +// MEASURED at 691b7af30 + this instrument (gfx1100, bf16, 66 tokens, hq 1, +// hkv 1, dh 128, L33/C2/R0 layer 0 step 0; all 1322 assertions pass), over the +// primary capture above and the primary preamble capture: +// +// row differing / 8448 +// A0 native run output vs primary output (the recorded 2902) 2902 +// B native run post-RoPE Q vs primary Q 1569 +// B native run post-RoPE K vs primary K 1542 +// C0 replay self-consistency (native Q/K/V at the primary's +// geometry vs the native run's own output) 0 +// C native ROCm kernel on the PRIMARY's own Q/K/V vs its output 2918 +// D native preamble on the primary's qkv (RmsNorm + RopeFromCache) 1569 Q / 1542 K +// D' the same through RopeNeox instead of the cos|sin cache 1776 Q / 1719 K +// E native cos|sin table vs primary cos_sin at the same positions 0 +// F primary qkv q/k slices vs the native run's pre-norm q/k 0 +// +// VERDICT: BOTH, kernel-dominant. The kernel term alone is 2918 of 8448 words on +// identical inputs (row C, with row C0 proving the replay reproduces the native +// run exactly at the primary's geometry), which is the whole recorded 2902. The +// preamble term is 1569 Q and 1542 K. Rows E and F exclude the cos|sin source +// and the qkv projection, so the preamble term is the pre-RoPE BF16 norm store. +// Row C's shape: rows 0 and 33 (query position 0, a one-key softmax) are exact, +// no differing word lies in a row whose maximum probability is >= 0.999, all +// 2918 lie in rows below 0.9, and 2132 of them are one bf16 code unit apart. +// +// GAP G5 (recorded, not repaired): a device->host read is illegal inside a graph +// capture and vt::Backend carries no "is capturing" query +// (include/vt/backend.h:244-262; vt::Queue is include/vt/device.h:131-135), so +// this instrument cannot decline on its own during a capture. It is declared +// EAGER-ONLY instead: RocmPlatform::support_static_graph_mode() is false, so the +// workload measured here is never captured. A graphed variant needs +// `virtual bool IsCapturing() const` on vt::Backend, overridden in RocmBackend +// beside BeginCapture, and an issue of its own. +// --------------------------------------------------------------------------- +namespace { + +// One env var, read once into a function-local static — the tree's convention +// (dense_attn_block.h:65, :93, :154). Unset => no observer is installed, so the +// default path pays nothing and dispatches exactly as it does today. +bool AttnDumpEnabled() { + static const bool on = [] { + const char* e = std::getenv("VT_ATTN_DUMP"); + return e != nullptr && e[0] == '1'; + }(); + return on; +} + +// --- host file helpers ------------------------------------------------------ +std::vector ReadBytes(const std::filesystem::path& path) { + REQUIRE_MESSAGE(std::filesystem::exists(path), "missing capture file ", path.string()); + const auto size = static_cast(std::filesystem::file_size(path)); + std::vector bytes(size); + std::ifstream in(path, std::ios::binary); + if (size > 0) + in.read(reinterpret_cast(bytes.data()), static_cast(size)); + REQUIRE(static_cast(in)); + return bytes; +} + +template +std::vector ReadTyped(const std::filesystem::path& path) { + const std::vector bytes = ReadBytes(path); + REQUIRE(bytes.size() % sizeof(T) == 0); + std::vector values(bytes.size() / sizeof(T)); + if (!bytes.empty()) std::memcpy(values.data(), bytes.data(), bytes.size()); + return values; +} + +std::vector ReadBf16(const std::filesystem::path& path) { + return ReadTyped(path); +} + +void WriteWords(const std::filesystem::path& path, const std::vector& words) { + if (words.empty()) return; + std::filesystem::create_directories(path.parent_path()); + std::ofstream(path, std::ios::binary) + .write(reinterpret_cast(words.data()), + static_cast(words.size() * sizeof(uint16_t))); +} + +// --- device read helpers ---------------------------------------------------- +std::vector ReadDevice(vt::Queue& q, const vt::Tensor& tensor) { + std::vector host(tensor.Bytes()); + auto& backend = vt::GetBackend(q.device); + if (!host.empty()) backend.Copy(q, host.data(), tensor.data, host.size()); + backend.Synchronize(q); + return host; +} + +std::vector ReadDeviceBf16(vt::Queue& q, const vt::Tensor& tensor) { + REQUIRE(tensor.dtype == vt::DType::kBF16); + const std::vector bytes = ReadDevice(q, tensor); + std::vector words(bytes.size() / sizeof(uint16_t)); + if (!bytes.empty()) std::memcpy(words.data(), bytes.data(), bytes.size()); + return words; +} + +template +std::vector ReadDeviceTyped(vt::Queue& q, const vt::Tensor& tensor) { + const std::vector bytes = ReadDevice(q, tensor); + REQUIRE(bytes.size() % sizeof(T) == 0); + std::vector values(bytes.size() / sizeof(T)); + if (!bytes.empty()) std::memcpy(values.data(), bytes.data(), bytes.size()); + return values; +} + +// The (num_blocks, block_size, hkv, dh) unbind slice is STRIDED: the block +// stride is 2*block_size*hkv*dh because the flash cache interleaves K and V per +// block (dense_attn_block.h:356-373). One linear Copy would read the wrong +// bytes, so read the slice one block at a time. +std::vector ReadCacheSlice(vt::Queue& q, const vt::Tensor& cache) { + REQUIRE(cache.rank == 4); + REQUIRE(cache.dtype == vt::DType::kBF16); + REQUIRE(cache.stride[3] == 1); + const int64_t blocks = cache.shape[0], bs = cache.shape[1]; + const int64_t h = cache.shape[2], d = cache.shape[3]; + std::vector host(static_cast(blocks * bs * h * d)); + auto& backend = vt::GetBackend(q.device); + for (int64_t b = 0; b < blocks; ++b) { + backend.Copy(q, host.data() + static_cast(b * bs * h * d), + static_cast(cache.data) + + static_cast(b * cache.stride[0]) * sizeof(uint16_t), + static_cast(bs * h * d) * sizeof(uint16_t)); + } + backend.Synchronize(q); + return host; +} + +// The K/V row the kernel reads for request r, key position j: +// cache[block_table[r, j / block_size], j % block_size, 0, :] +std::vector GatherKv(const std::vector& slice, int64_t block_size, + int64_t hkv, int64_t dh, + const std::vector& block_table, + int64_t block_table_cols, int64_t request, int64_t keys) { + const size_t row = static_cast(hkv * dh); + std::vector out(static_cast(keys) * row); + for (int64_t j = 0; j < keys; ++j) { + const int64_t block = + block_table[static_cast(request * block_table_cols + j / block_size)]; + const size_t src = static_cast((block * block_size + j % block_size)) * row; + REQUIRE(src + row <= slice.size()); + std::memcpy(out.data() + static_cast(j) * row, slice.data() + src, + row * sizeof(uint16_t)); + } + return out; +} + +// --- capture records -------------------------------------------------------- +struct PagedAttentionCapture { + int64_t calls = 0; + int64_t tokens = 0, hq = 0, dh = 0, hkv = 0, block_size = 0, blocks = 0; + int64_t block_table_cols = 0; + float scale = 0.0f; + bool causal = true; + bool have = false; + std::vector query, out, k_cache, v_cache; + std::vector block_table, seq_lens, query_start_loc; +}; + +struct RopeCapture { + bool have = false; + int64_t tokens = 0, heads = 0, dh = 0, rot = 0, rms_calls_before = 0; + std::vector q_pre, k_pre, q_post, k_post, cos_sin; + std::vector index; +}; + +struct RmsNormCall { + int64_t shape0 = 0, shape1 = 0; + float eps = 0.0f; + std::vector in, out, weight; +}; + +// The fused qk-norm-plus-RoPE preamble op (`vt::OpId::kAttnQkNormRope`). Before +// the #3115 repair ROCm registers no fast realization for the recipe, so the +// production preamble is the Tier-0 composite and this observer never fires. +// After it, the fused op IS the preamble, and the kRopeFromCache/kRmsNorm +// observers above stop firing on the production path. The runner reads whichever +// realization executed, so one case measures one boundary before and after. +struct PreambleCapture { + bool have = false; + int64_t tokens = 0, hq = 0, hkv = 0, dh = 0, rot = 0; + float eps = 0.0f; + std::vector q_in, k_in, q_post, k_post; +}; + +PagedAttentionCapture g_attn; +RopeCapture g_rope; +PreambleCapture g_preamble; +int64_t g_rope_neox_calls = 0; +std::vector g_rms_calls; +bool g_record = false; + +vt::PagedAttentionFn g_native_paged_attention = nullptr; +vt::RmsNormFn g_native_rms_norm = nullptr; +vt::RopeFromCacheFn g_native_rope_from_cache = nullptr; +vt::RopeFn g_native_rope_neox = nullptr; +vt::AttnQkNormRopeFn g_native_attn_qk_norm_rope = nullptr; + +void CapturePagedAttention(vt::Queue& q, vt::Tensor& out, const vt::Tensor& query, + const vt::Tensor& k_cache, const vt::Tensor& v_cache, + const vt::Tensor& block_table, const vt::Tensor& seq_lens, + const vt::Tensor& query_start_loc, + const vt::PagedAttentionArgs& args) { + const bool record = g_record && g_attn.calls == 0; + if (record) { + g_attn.tokens = query.shape[0]; + g_attn.hq = query.shape[1]; + g_attn.dh = query.shape[2]; + g_attn.hkv = k_cache.shape[2]; + g_attn.block_size = k_cache.shape[1]; + g_attn.blocks = k_cache.shape[0]; + g_attn.block_table_cols = block_table.shape[1]; + g_attn.scale = args.scale; + g_attn.causal = args.causal; + g_attn.query = ReadDeviceBf16(q, query); + g_attn.k_cache = ReadCacheSlice(q, k_cache); + g_attn.v_cache = ReadCacheSlice(q, v_cache); + g_attn.block_table = ReadDeviceTyped(q, block_table); + g_attn.seq_lens = ReadDeviceTyped(q, seq_lens); + g_attn.query_start_loc = ReadDeviceTyped(q, query_start_loc); + } + g_native_paged_attention(q, out, query, k_cache, v_cache, block_table, seq_lens, + query_start_loc, args); + if (record) { + g_attn.out = ReadDeviceBf16(q, out); + g_attn.have = true; + } + g_attn.calls += 1; +} + +void CaptureRmsNorm(vt::Queue& q, vt::Tensor& out, const vt::Tensor& x, + const vt::Tensor& weight, const vt::RmsNormArgs& args, + vt::Tensor* residual) { + const bool record = g_record && !g_rope.have && g_rms_calls.size() < 32 && + x.Numel() <= 65536 && x.dtype == vt::DType::kBF16 && + out.dtype == vt::DType::kBF16; + RmsNormCall call; + if (record) { + call.shape0 = x.shape[0]; + call.shape1 = x.rank >= 2 ? x.shape[1] : 0; + call.eps = args.eps; + call.in = ReadDeviceBf16(q, x); + call.weight = ReadDeviceBf16(q, weight); + } + g_native_rms_norm(q, out, x, weight, args, residual); + if (record) { + call.out = ReadDeviceBf16(q, out); + g_rms_calls.push_back(std::move(call)); + } +} + +// The preamble's RoPE call is preceded by RmsNorm(q) then RmsNorm(k) from the +// SAME composite (include/vt/recipes.h kAttnQkNormRope), so the two RmsNorm +// calls immediately before it are the pair, identified by position in the call +// sequence rather than by shape. +void CaptureRopeFromCache(vt::Queue& q, vt::Tensor& q_states, vt::Tensor* k_states, + const vt::Tensor& positions, const vt::Tensor& cos_sin_cache, + const vt::RopeArgs& args) { + const bool record = g_record && !g_rope.have; + if (record) { + g_rope.tokens = q_states.shape[0]; + g_rope.heads = q_states.shape[1]; + g_rope.dh = q_states.shape[2]; + g_rope.rot = args.rotary_dim; + g_rope.rms_calls_before = static_cast(g_rms_calls.size()); + g_rope.q_pre = ReadDeviceBf16(q, q_states); + if (k_states != nullptr) g_rope.k_pre = ReadDeviceBf16(q, *k_states); + g_rope.cos_sin = ReadDeviceBf16(q, cos_sin_cache); + g_rope.index = ReadDeviceTyped(q, positions); + } + g_native_rope_from_cache(q, q_states, k_states, positions, cos_sin_cache, args); + if (record) { + g_rope.q_post = ReadDeviceBf16(q, q_states); + if (k_states != nullptr) g_rope.k_post = ReadDeviceBf16(q, *k_states); + g_rope.have = true; + } +} + +void CaptureRopeNeox(vt::Queue& q, vt::Tensor& q_states, vt::Tensor& k_states, + const vt::Tensor& positions, const vt::RopeArgs& args) { + g_rope_neox_calls += 1; + g_native_rope_neox(q, q_states, k_states, positions, args); +} + +// The fused preamble norms and rotates q3/k3 IN PLACE (include/vt/ops.h:2285), +// so the entry bytes ARE the qkv projection's own slice and the return bytes ARE +// the post-RoPE boundary the attention kernel consumes. +void CaptureAttnQkNormRope(vt::Queue& q, vt::Tensor& q3, vt::Tensor& k3, + const vt::Tensor& q_norm, const vt::Tensor& k_norm, + const vt::Tensor& cos_sin, const vt::Tensor& positions, + const vt::RmsNormArgs& norm_args, const vt::RopeArgs& rope_args) { + const bool record = g_record && !g_preamble.have; + if (record) { + g_preamble.tokens = q3.shape[0]; + g_preamble.hq = q3.shape[1]; + g_preamble.dh = q3.shape[2]; + g_preamble.hkv = k3.shape[1]; + g_preamble.rot = rope_args.rotary_dim; + g_preamble.eps = norm_args.eps; + g_preamble.q_in = ReadDeviceBf16(q, q3); + g_preamble.k_in = ReadDeviceBf16(q, k3); + (void)q_norm; + (void)k_norm; + (void)cos_sin; + (void)positions; + } + g_native_attn_qk_norm_rope(q, q3, k3, q_norm, k_norm, cos_sin, positions, norm_args, + rope_args); + if (record) { + g_preamble.q_post = ReadDeviceBf16(q, q3); + g_preamble.k_post = ReadDeviceBf16(q, k3); + g_preamble.have = true; + } +} + +// Function-local static: the provider table is process-global and registering +// the same name twice is refused, so this runs exactly once. GetOp is read here +// and NOWHERE else, so on a re-entry the saved pointer is the live kernel and +// never this wrapper. +void InstallAttnObservers() { + static const bool installed = [] { + g_native_paged_attention = reinterpret_cast( + vt::GetOp(vt::OpId::kPagedAttention, vt::DeviceType::kROCM)); + vt::RegisterOpProvider(vt::OpId::kPagedAttention, vt::DeviceType::kROCM, + {"test-attention-boundary-paged-attention", 100, nullptr, + reinterpret_cast(static_cast( + &CapturePagedAttention))}); + g_native_rms_norm = reinterpret_cast( + vt::GetOp(vt::OpId::kRmsNorm, vt::DeviceType::kROCM)); + vt::RegisterOpProvider(vt::OpId::kRmsNorm, vt::DeviceType::kROCM, + {"test-attention-boundary-rmsnorm", 100, nullptr, + reinterpret_cast( + static_cast(&CaptureRmsNorm))}); + g_native_rope_from_cache = reinterpret_cast( + vt::GetOp(vt::OpId::kRopeFromCache, vt::DeviceType::kROCM)); + vt::RegisterOpProvider(vt::OpId::kRopeFromCache, vt::DeviceType::kROCM, + {"test-attention-boundary-rope-from-cache", 100, nullptr, + reinterpret_cast(static_cast( + &CaptureRopeFromCache))}); + g_native_rope_neox = + reinterpret_cast(vt::GetOp(vt::OpId::kRopeNeox, vt::DeviceType::kROCM)); + vt::RegisterOpProvider(vt::OpId::kRopeNeox, vt::DeviceType::kROCM, + {"test-attention-boundary-rope-neox", 100, nullptr, + reinterpret_cast( + static_cast(&CaptureRopeNeox))}); + // Registered only when the backend HAS the fused op. Before the #3115 repair + // ROCm registers none, GetOp would return null, and the wrapper would + // forward into a null pointer; the composite observers above are the + // preamble source on that tree. + if (vt::OpRegistered(vt::OpId::kAttnQkNormRope, vt::DeviceType::kROCM)) { + g_native_attn_qk_norm_rope = reinterpret_cast( + vt::GetOp(vt::OpId::kAttnQkNormRope, vt::DeviceType::kROCM)); + vt::RegisterOpProvider( + vt::OpId::kAttnQkNormRope, vt::DeviceType::kROCM, + {"test-attention-boundary-attn-qk-norm-rope", 100, nullptr, + reinterpret_cast( + static_cast(&CaptureAttnQkNormRope))}); + } + return true; + }(); + (void)installed; +} + +// --- the primary's captured attention boundary ------------------------------ +struct PrimaryBoundary { + int64_t tokens = 0, hq = 0, hkv = 0, dh = 0, num_reqs = 0; + int64_t block_size = 0, block_table_cols = 0; + float scale = 0.0f; + std::vector query, key, value, output; + std::vector query_start_loc, seq_lens, block_table; + std::vector slot_mapping; +}; + +PrimaryBoundary LoadPrimaryBoundary(const std::filesystem::path& dir, const std::string& label, + int index) { + const std::string stem = label + "-attention-" + std::to_string(index); + nlohmann::json record; + std::ifstream(dir / (stem + ".json")) >> record; + REQUIRE(!record.is_null()); + PrimaryBoundary b; + b.scale = record.at("scale").get(); + b.hq = record.at("num_heads").get(); + b.hkv = record.at("num_kv_heads").get(); + b.dh = record.at("head_size").get(); + b.tokens = record.at("num_actual_tokens").get(); + b.block_size = record.at("kv_cache").at("shape").at(2).get(); + b.block_table_cols = record.at("tensors").at("block_table").at("shape").at(1).get(); + b.num_reqs = static_cast(record.at("seq_lens").size()); + b.query = ReadBf16(dir / (stem + "-query.bin")); + b.key = ReadBf16(dir / (stem + "-key.bin")); + b.value = ReadBf16(dir / (stem + "-value.bin")); + b.output = ReadBf16(dir / (stem + "-output.bin")); + b.query_start_loc = ReadTyped(dir / (stem + "-query_start_loc.bin")); + b.seq_lens = ReadTyped(dir / (stem + "-seq_lens.bin")); + b.block_table = ReadTyped(dir / (stem + "-block_table.bin")); + b.slot_mapping = ReadTyped(dir / (stem + "-slot_mapping.bin")); + const size_t expected = static_cast(b.tokens * b.hq * b.dh); + REQUIRE(b.query.size() == expected); + REQUIRE(b.key.size() == expected); + REQUIRE(b.value.size() == expected); + REQUIRE(b.output.size() == expected); + return b; +} + +// --- diffs ------------------------------------------------------------------ +struct WordDiff { + int64_t words = 0, different = 0; + std::vector row_different; + int64_t distance[5] = {0, 0, 0, 0, 0}; + std::vector first; +}; + +// bf16 shares f32's ordering, so reading the 16 bits as a signed integer is +// monotone within one sign and the difference is a code-unit (ULP) distance. +int64_t Bf16CodeDistance(uint16_t a, uint16_t b) { + return std::llabs(static_cast(static_cast(a)) - + static_cast(static_cast(b))); +} + +WordDiff DiffWords(const std::vector& actual, const std::vector& expected, + int64_t row) { + REQUIRE(actual.size() == expected.size()); + WordDiff d; + d.words = static_cast(actual.size()); + if (row > 0) d.row_different.assign(static_cast(d.words / row), 0); + for (size_t i = 0; i < actual.size(); ++i) { + if (actual[i] == expected[i]) continue; + d.different += 1; + const int64_t distance = Bf16CodeDistance(actual[i], expected[i]); + d.distance[distance >= 4 ? 4 : distance] += 1; + if (row > 0) d.row_different[i / static_cast(row)] += 1; + if (d.first.size() < 8) d.first.push_back(static_cast(i)); + } + return d; +} + +// Host-side f32 classification of each query row: for query token t the maximum +// softmax probability over the keys it can see. A difference confined to rows +// whose maximum is ~1 (a trivial softmax) points at exp/reduction order; a +// difference spread over rows with a flat distribution points at the +// probability dtype. +std::vector MaxProbPerRow(const PrimaryBoundary& b) { + std::vector result(static_cast(b.tokens), 0.0); + for (int64_t r = 0; r < b.num_reqs; ++r) { + const int64_t begin = b.query_start_loc[static_cast(r)]; + const int64_t end = b.query_start_loc[static_cast(r + 1)]; + const int64_t seq = b.seq_lens[static_cast(r)]; + for (int64_t t = begin; t < end; ++t) { + const int64_t position = seq - (end - begin) + (t - begin); + std::vector scores(static_cast(position + 1), 0.0); + for (int64_t j = 0; j <= position; ++j) { + double dot = 0.0; + for (int64_t d = 0; d < b.dh; ++d) + dot += static_cast(vt::BF16ToF32(b.query[static_cast(t * b.dh + d)])) * + static_cast(vt::BF16ToF32(b.key[static_cast(j * b.dh + d)])); + scores[static_cast(j)] = static_cast(b.scale) * dot; + } + const double max = *std::max_element(scores.begin(), scores.end()); + double sum = 0.0; + for (double s : scores) sum += std::exp(s - max); + result[static_cast(t)] = 1.0 / sum; + } + } + return result; +} + +void PrintDiff(const std::string& tag, const WordDiff& d) { + std::cout << "[attn-parity] " << tag << ": " << d.different << " / " << d.words + << " differing bf16 words; bf16 code-unit distance 1:" << d.distance[1] + << " 2:" << d.distance[2] << " 3:" << d.distance[3] << " >=4:" << d.distance[4] + << std::endl; +} + +// WHERE the words differ, in the two axes that localize the cause. +void PrintDiffShape(const PrimaryBoundary& b, const WordDiff& d, + const std::vector& max_prob) { + std::cout << "[attn-parity] per-row differing words (" << d.row_different.size() + << " rows):"; + for (size_t t = 0; t < d.row_different.size(); ++t) + std::cout << " " << d.row_different[t]; + std::cout << std::endl; + const double edges[3] = {0.999, 0.99, 0.9}; + int64_t words[4] = {0, 0, 0, 0}, rows[4] = {0, 0, 0, 0}, total[4] = {0, 0, 0, 0}; + for (size_t t = 0; t < max_prob.size(); ++t) { + const int bucket = max_prob[t] >= edges[0] ? 0 + : max_prob[t] >= edges[1] ? 1 + : max_prob[t] >= edges[2] ? 2 + : 3; + total[bucket] += 1; + if (t < d.row_different.size() && d.row_different[t] > 0) rows[bucket] += 1; + if (t < d.row_different.size()) words[bucket] += d.row_different[t]; + } + const char* names[4] = {"maxprob>=0.999", "0.99<=maxprob<0.999", "0.9<=maxprob<0.99", + "maxprob<0.9"}; + for (int i = 0; i < 4; ++i) + std::cout << "[attn-parity] " << names[i] << ": " << words[i] << " differing words over " + << rows[i] << " of " << total[i] << " rows" << std::endl; + if (!d.row_different.empty()) + std::cout << "[attn-parity] row 0 (position 0, softmax trivial): " + << d.row_different[0] << " differing words; first differing flat indices:"; + for (int64_t index : d.first) std::cout << " " << index; + std::cout << std::endl; + std::cout << "[attn-parity] requests: " << b.num_reqs; + for (int64_t r = 0; r < b.num_reqs; ++r) { + int64_t count = 0; + for (int64_t t = b.query_start_loc[static_cast(r)]; + t < b.query_start_loc[static_cast(r + 1)]; ++t) + if (t < static_cast(d.row_different.size())) + count += d.row_different[static_cast(t)]; + std::cout << " request " << r << "=" << count; + } + std::cout << std::endl; +} + +// The primary's captured Q/K preamble boundary, beside the attention capture. +struct PrimaryPreamble { + int64_t tokens = 0, qkv_width = 0, dh = 0, rot = 0; + std::vector qkv, q_gamma, k_gamma, cos_sin, q_left, q_right, k_left, k_right; + std::vector positions; +}; + +PrimaryPreamble LoadPrimaryPreamble(const std::filesystem::path& dir, const std::string& label, + int index) { + const std::string stem = label + "-residual-preamble-" + std::to_string(index); + nlohmann::json record; + std::ifstream(dir / (stem + ".json")) >> record; + REQUIRE(!record.is_null()); + PrimaryPreamble p; + p.tokens = record.at("inputs").at("qkv").at("shape").at(0).get(); + p.qkv_width = record.at("inputs").at("qkv").at("shape").at(1).get(); + p.dh = record.at("inputs").at("q_gamma").at("shape").at(0).get(); + p.rot = record.at("inputs").at("cos_sin").at("shape").at(1).get(); + p.qkv = ReadBf16(dir / (stem + "-input-qkv.bin")); + p.q_gamma = ReadBf16(dir / (stem + "-input-q_gamma.bin")); + p.k_gamma = ReadBf16(dir / (stem + "-input-k_gamma.bin")); + p.cos_sin = ReadBf16(dir / (stem + "-input-cos_sin.bin")); + p.positions = ReadTyped(dir / (stem + "-input-positions.bin")); + p.q_left = ReadBf16(dir / (stem + "-output-q_left.bin")); + p.q_right = ReadBf16(dir / (stem + "-output-q_right.bin")); + p.k_left = ReadBf16(dir / (stem + "-output-k_left.bin")); + p.k_right = ReadBf16(dir / (stem + "-output-k_right.bin")); + REQUIRE(p.qkv.size() == static_cast(p.tokens * p.qkv_width)); + REQUIRE(p.cos_sin.size() == + static_cast(record.at("inputs").at("cos_sin").at("shape").at(0).get() * + p.rot)); + REQUIRE(p.positions.size() == static_cast(p.tokens)); + return p; +} + +// q_left|q_right (each [T,1,Dh/2], stride [Dh,Dh,1]) reassembled into the +// [T,1,Dh] query the attention capture holds. +std::vector JoinRopeHalves(const PrimaryPreamble& p, const std::vector& left, + const std::vector& right) { + const int64_t half = p.dh / 2; + std::vector joined(static_cast(p.tokens * p.dh), 0); + for (int64_t t = 0; t < p.tokens; ++t) { + for (int64_t d = 0; d < half; ++d) { + joined[static_cast(t * p.dh + d)] = left[static_cast(t * half + d)]; + joined[static_cast(t * p.dh + half + d)] = right[static_cast(t * half + d)]; + } + } + return joined; +} + +// One head-contiguous row of the flash cache, at the address the capture's own +// slot_mapping names: block = slot / block_size, offset = slot % block_size. +void ScatterCacheSlot(std::vector& cache, int64_t blocks, int64_t block_size, + int64_t hkv, int64_t dh, int which, int64_t slot, const uint16_t* row) { + const int64_t block = slot / block_size; + const int64_t offset = slot % block_size; + REQUIRE(block >= 0); + REQUIRE(block < blocks); + const size_t base = static_cast((block * 2 + which) * block_size * hkv * dh + + offset * hkv * dh); + REQUIRE(base + static_cast(hkv * dh) <= cache.size()); + std::memcpy(cache.data() + base, row, static_cast(hkv * dh) * sizeof(uint16_t)); +} + +// --- the arm's key walk, on synthesized data -------------------------------- +// +// The reference max the probability is narrowed against is not a free parameter: +// the primary takes it per contiguous key tile, and the tile is chosen by the arm +// that executes for that row (## Design 1, .agents/specs/rocm-attn-parity.md). +// The recorded capture cannot witness the choice: its 33-token rows are all +// prefix_prefill with an EMPTY context, so `max(keys 0..31) == max(keys 0..32)` +// and a 32-key tile reproduces the same 8448 bytes. These cases synthesize the +// data that makes the width load-bearing, and compare the device against a host +// transcription of the primary's own key walk for the arm (`PrimaryTiles` + +// `HostArmOutput`) and against the same transcription under neighbouring widths. +// +// Data design. q = e0 and every K row is either 0 or `kArmHighKey` in element 0, +// so every score is exactly 0 or s = kArmHighKey * scale (bf16-exact operands, a +// single nonzero product, so the dot is order-independent). Every V row is a +// single bf16-exact integer in its own lane, so each output lane is fed by one +// key and the narrowed-probability accumulations are sums of bf16-exact values. +// A key that lies in a tile with a high key is narrowed at exp(-s); the same key +// alone in a tile is narrowed at 1. The two differ by a relative 2^-9, which is +// what moves the output word -- and it only exists when the tile is right. +constexpr float kArmHighKey = 5.65625f; // bf16-exact; s = 0.4999466 +constexpr int64_t kArmDh = 128; + +// The kernel's FastExp (src/vt/rocm/rocm_paged_attn.hip:196-198) on the host. +// Only exp2(0) = 1 and exp2(-inf) = 0 are load-bearing here. +float ArmFastExp(float x) { return std::exp2(x * 1.4426950408889634f); } + +std::vector ArmWords(const std::vector& values) { + std::vector words(values.size()); + for (size_t i = 0; i < values.size(); ++i) words[i] = vt::F32ToBF16(values[i]); + return words; +} + +// The primary's key walk for ONE query row, as absolute [begin, end) key tiles. +// chunked selects prefix_prefill; the widths are its pinned constants: +// TRITON_BLOCK_SIZE = 32 for the cached context (:965, :1007) and BLOCK_N = 64 +// for the current chunk when the physical block size is a power of two, 32 +// otherwise (:955-966). The decode arm tiles by min(block_size, 128) +// (chunked_prefill_paged_decode.py:444-445), also 32 for a non-power-of-two +// physical block size. +std::vector> PrimaryTiles(bool chunked, int64_t context, + int64_t jmax, int64_t block_size) { + const bool pow2 = block_size > 0 && (block_size & (block_size - 1)) == 0; + std::vector> tiles; + if (chunked) { + for (int64_t base = 0; base < context; base += 32) + tiles.emplace_back(base, std::min(base + 32, context)); + } + const int64_t width = chunked ? (pow2 ? 64 : 32) + : (pow2 ? std::min(block_size, 128) : 32); + for (int64_t base = chunked ? context : 0; base <= jmax; base += width) + tiles.emplace_back(base, std::min(base + width, jmax + 1)); + return tiles; +} + +// A uniform-width tiling from key 0, the geometry this kernel used before the +// repair. Used as the falsifying counterfactual: a case whose device bytes match +// the primary arm's tiles and NOT these is a case the width is load-bearing for. +std::vector> UniformTiles(int64_t jmax, int64_t width) { + std::vector> tiles; + for (int64_t base = 0; base <= jmax; base += width) + tiles.emplace_back(base, std::min(base + width, jmax + 1)); + return tiles; +} + +// The primary's softmax over one query row, transcribed from +// prefix_prefill.py:442-478 (context loop :231-343, chunk loop :369) and +// chunked_prefill_paged_decode.py:244-268. `key_score[j]` is the f32 q.k for key +// j, `values[j*dh + d]` the f32 V row, `tiles` the arm's key walk. +std::vector HostArmOutput(const std::vector& key_score, + const std::vector& values, + const std::vector>& tiles, + float scale, int64_t dh) { + const float ninf = -std::numeric_limits::infinity(); + float m = ninf, l = 0.0f; + std::vector acc(static_cast(dh), 0.0f); + for (const auto& tile : tiles) { + float tile_max = ninf; + for (int64_t j = tile.first; j < tile.second; ++j) + tile_max = std::max(tile_max, key_score[static_cast(j)] * scale); + const float m_new = std::max(m, tile_max); + const float alpha = ArmFastExp(m - m_new); + for (float& a : acc) a *= alpha; + l *= alpha; + m = m_new; + for (int64_t j = tile.first; j < tile.second; ++j) { + const float p = ArmFastExp(key_score[static_cast(j)] * scale - m_new); + const float pw = vt::BF16ToF32(vt::F32ToBF16(p)); // p.to(v.dtype) + for (int64_t d = 0; d < dh; ++d) + acc[static_cast(d)] += pw * values[static_cast(j * dh + d)]; + l += p; + } + } + const float inv = l > 0.0f ? 1.0f / l : 0.0f; + for (float& a : acc) a *= inv; + return acc; +} + +// One synthesized workload: `num_reqs` requests of `query_len` query tokens each +// over a cache that already holds `seq_len - query_len` keys, with exactly one +// high-scoring key at `high_key`. Every request has the same shape, so the arm +// is the same for every row of the batch. +struct SyntheticArm { + int64_t num_reqs = 0, query_len = 0, seq_len = 0, block_size = 0, high_key = 0; + int64_t dh = kArmDh, hq = 1, hkv = 1, blocks = 0, block_table_cols = 0, total_q = 0; + float scale = 0.0f; + std::vector query_start_loc, seq_lens, block_table; + std::vector query, k_cache, v_cache; + std::vector key_score, values; // host f32 copies for the model +}; + +SyntheticArm BuildSyntheticArm(int64_t num_reqs, int64_t query_len, int64_t seq_len, + int64_t block_size, int64_t high_key) { + SyntheticArm a; + a.num_reqs = num_reqs; + a.query_len = query_len; + a.seq_len = seq_len; + a.block_size = block_size; + a.high_key = high_key; + a.scale = 1.0f / std::sqrt(static_cast(kArmDh)); + a.total_q = num_reqs * query_len; + a.block_table_cols = (seq_len + block_size - 1) / block_size; + a.blocks = num_reqs * a.block_table_cols; + const size_t row = static_cast(a.dh); + a.query.assign(static_cast(a.total_q) * row, 0); + a.k_cache.assign(static_cast(a.blocks * block_size) * row, 0); + a.v_cache.assign(static_cast(a.blocks * block_size) * row, 0); + a.key_score.assign(static_cast(seq_len), 0.0f); + a.values.assign(static_cast(seq_len) * row, 0.0f); + a.query_start_loc.assign(static_cast(num_reqs) + 1, 0); + for (int64_t j = 0; j < seq_len; ++j) { + a.key_score[static_cast(j)] = j == high_key ? kArmHighKey : 0.0f; + // Key j's value row is a single bf16-exact integer in lane j, so each output + // lane is fed by exactly one key and the narrowing is what the lane shows. + a.values[static_cast(j * a.dh + (j % a.dh))] = static_cast(64 + j % 8); + } + for (int64_t r = 0; r < num_reqs; ++r) { + a.query_start_loc[static_cast(r + 1)] = static_cast((r + 1) * query_len); + a.seq_lens.push_back(static_cast(seq_len)); + for (int64_t c = 0; c < a.block_table_cols; ++c) + a.block_table.push_back(static_cast(r * a.block_table_cols + c)); + for (int64_t j = 0; j < seq_len; ++j) { + const int64_t slot = static_cast( + a.block_table[static_cast(r * a.block_table_cols + + j / block_size)]) * + block_size + + j % block_size; + const size_t base = static_cast(slot) * row; + a.k_cache[base] = vt::F32ToBF16(a.key_score[static_cast(j)]); + a.v_cache[base + static_cast(j % a.dh)] = + vt::F32ToBF16(a.values[static_cast(j * a.dh + j % a.dh)]); + } + for (int64_t i = 0; i < query_len; ++i) { + const size_t t = static_cast(r * query_len + i); + a.query[t * row] = vt::F32ToBF16(1.0f); // q = e0, so q.k_j = key_score[j] + } + } + return a; +} + +// One device forward of the workload through the shared op layer, returning the +// [total_q, hq, dh] BF16 output. +std::vector RunSyntheticArm(vt::Backend& backend, vt::Queue& q, + const SyntheticArm& a) { + Buffer qb(backend, a.query.size() * sizeof(uint16_t)); + Buffer kb(backend, a.k_cache.size() * sizeof(uint16_t)); + Buffer vb(backend, a.v_cache.size() * sizeof(uint16_t)); + Buffer ob(backend, a.query.size() * sizeof(uint16_t)); + Buffer tbb(backend, a.block_table.size() * sizeof(int32_t)); + Buffer slb(backend, a.seq_lens.size() * sizeof(int32_t)); + Buffer qlb(backend, a.query_start_loc.size() * sizeof(int32_t)); + backend.Copy(q, qb.data, a.query.data(), a.query.size() * sizeof(uint16_t)); + backend.Copy(q, kb.data, a.k_cache.data(), a.k_cache.size() * sizeof(uint16_t)); + backend.Copy(q, vb.data, a.v_cache.data(), a.v_cache.size() * sizeof(uint16_t)); + backend.Copy(q, tbb.data, a.block_table.data(), a.block_table.size() * sizeof(int32_t)); + backend.Copy(q, slb.data, a.seq_lens.data(), a.seq_lens.size() * sizeof(int32_t)); + backend.Copy(q, qlb.data, a.query_start_loc.data(), + a.query_start_loc.size() * sizeof(int32_t)); + backend.Synchronize(q); + vt::Tensor tq = Tensor(qb.data, vt::DType::kBF16, q.device, {a.total_q, a.hq, a.dh}); + vt::Tensor tk = Tensor(kb.data, vt::DType::kBF16, q.device, {a.blocks, a.block_size, a.hkv, a.dh}); + vt::Tensor tv = Tensor(vb.data, vt::DType::kBF16, q.device, {a.blocks, a.block_size, a.hkv, a.dh}); + vt::Tensor to = Tensor(ob.data, vt::DType::kBF16, q.device, {a.total_q, a.hq, a.dh}); + vt::Tensor tt = Tensor(tbb.data, vt::DType::kI32, q.device, {a.num_reqs, a.block_table_cols}); + vt::Tensor ts = Tensor(slb.data, vt::DType::kI32, q.device, {a.num_reqs}); + vt::Tensor tl = Tensor(qlb.data, vt::DType::kI32, q.device, {a.num_reqs + 1}); + vt::PagedAttentionArgs args; + args.scale = a.scale; + args.causal = true; + args.query_start_loc_host = a.query_start_loc.data(); + args.max_seq_len = a.seq_len; + vt::PagedAttention(q, to, tq, tk, tv, tt, ts, tl, args); + backend.Synchronize(q); + return ReadDeviceBf16(q, to); +} + +// The model's output for the whole batch, one arm geometry for every row. +std::vector ArmModelWords(const SyntheticArm& a, bool chunked, int64_t tile_override) { + std::vector words; + for (int64_t r = 0; r < a.num_reqs; ++r) { + const int64_t context = a.seq_len - a.query_len; + for (int64_t i = 0; i < a.query_len; ++i) { + const int64_t jmax = context + i; // causal: the row sits at position jmax + const std::vector> tiles = + tile_override > 0 ? UniformTiles(jmax, tile_override) + : PrimaryTiles(chunked, context, jmax, a.block_size); + const std::vector out = + HostArmOutput(a.key_score, a.values, tiles, a.scale, a.dh); + const std::vector row_words = ArmWords(out); + words.insert(words.end(), row_words.begin(), row_words.end()); + } + } + return words; +} + +void PrintArmDiff(const char* what, const WordDiff& d, int64_t rows) { + std::cout << "[arm-geometry] " << what << ": " << d.different << " / " << d.words + << " differing words over " << rows << " rows" << std::endl; +} + +// The noise floor is the host/device accumulation-order difference on the f32 +// running sum; the signal is the narrowed-probability scale the tile decides. +// They are three orders of magnitude apart on this data, so the thresholds are +// not tuned to the measurement. +constexpr int64_t kArmNoiseWords = 8; +constexpr int64_t kArmSignalWords = 32; + +// ─── env-gated cases: per-case skip, process-level 77 ──────────────────────── +// +// Four cases in this file need an environment the default run does not have: the +// two production cases need `VT_ROCM_MOE_FIXTURE`, the CPU half of the instrument +// needs `VT_ATTN_PARITY_PRIMARY`, and the device half needs all of that plus +// `VT_ATTN_DUMP=1`. They used to call `std::exit(77)` on the absent variable, and +// doctest runs the cases in file order, so the first absent variable ended the +// PROCESS: the summary was never printed and every later case was silently not +// run (the fresh review's fourth finding). +// +// The cases are now decorated `doctest::skip(...)`, so doctest reports them +// skipped in its own summary and the remaining cases still run, and the process +// still exits 77 -- CTest reports Skipped, the convention `tests/CMakeLists.txt` +// registers -- but ONLY when nothing failed, so a reddened case keeps doctest's +// own non-zero status instead of being folded into the skip. `--no-skip` forces +// a gated case anyway; the in-case guard then FAILS rather than reading an absent +// variable as a path. +bool AttnCaptureAbsent() { + const char* e = std::getenv("VT_ATTN_PARITY_PRIMARY"); + return e == nullptr || e[0] == '\0'; +} +bool AttnInstrumentAbsent() { + return FixtureAbsent() || AttnCaptureAbsent() || !AttnDumpEnabled(); +} + +// The four backend-requiring cases near the end of this file measure the ROCm +// arm on data the file builds itself, so they need the ROCm backend registered +// but no captured oracle and no environment variable. On the CPU-only lane +// (VLLM_CPP_HIP=OFF) nothing registers that backend, and a bare REQUIRE would +// FATAL there instead of reporting the CTest Skipped the rest of this file's +// ROCm-only cases report. +bool RocmBackendAbsent() { + return vt::TryGetBackend(vt::DeviceType::kROCM) == nullptr; +} + +bool g_env_cases_skipped = false; +bool g_run_failed = false; + +// A listener, not a reporter: listeners are always active whatever `-r=` selects, +// so neither the skip note nor the failure guard can be switched off from the +// command line (the sibling head test registers the same shape). +struct GatedCaseListener : public doctest::IReporter { + explicit GatedCaseListener(const doctest::ContextOptions&) {} + void test_run_start() override {} + void report_query(const doctest::QueryData&) override {} + void test_run_end(const doctest::TestRunStats& stats) override { + g_run_failed = stats.numTestCasesFailed != 0 || stats.numAssertsFailed != 0; + } + void test_case_start(const doctest::TestCaseData&) override {} + void test_case_reenter(const doctest::TestCaseData&) override {} + void test_case_end(const doctest::CurrentTestCaseStats&) override {} + void test_case_exception(const doctest::TestCaseException&) override {} + void subcase_start(const doctest::SubcaseSignature&) override {} + void subcase_end() override {} + void log_assert(const doctest::AssertData&) override {} + void log_message(const doctest::MessageData&) override {} + // Also fires for a case a filter excluded; only the decorator sets `m_skip`, + // and only that means a prerequisite this case needs is absent -- an unset + // environment variable for the captured-oracle cases, a registered backend for + // the synthesized-data ones. `TestCaseData` carries no reason, so the note + // below has to stay true for either. + void test_case_skipped(const doctest::TestCaseData& tc) override { + if (!tc.m_skip) return; + g_env_cases_skipped = true; + std::cout << "[rocm-moe-bf16] SKIPPED: " << tc.m_name + << " (a prerequisite this case needs is absent; see the case's own message)" + << std::endl; + } +}; +DOCTEST_REGISTER_LISTENER("vt-rocm-moe-env-gated", 1, GatedCaseListener); + +// Registered during static initialization, so it runs after doctest's main has +// printed its summary and returned. `std::_Exit` rather than `std::exit`: this IS +// an exit handler, and re-entering the exit sequence is undefined. +void ExitSkippedWhenAGatedCaseDidNotRun() { + if (!g_env_cases_skipped || g_run_failed) return; + std::cout.flush(); + std::fflush(nullptr); + std::fprintf(stderr, + "\n*** SKIPPED (exit 77): at least one env-gated case did not run. The cases that " + "did run are reported in the summary above and their result stands; this status " + "says only that the gated ones did not. ***\n\n"); + std::fflush(stderr); + std::_Exit(77); +} + +struct RegisterExitSkippedWhenAGatedCaseDidNotRun { + RegisterExitSkippedWhenAGatedCaseDidNotRun() { + std::atexit(&ExitSkippedWhenAGatedCaseDidNotRun); + } +}; +[[maybe_unused]] const RegisterExitSkippedWhenAGatedCaseDidNotRun + g_exit_skipped_when_a_gated_case_did_not_run; + +} // namespace + +// The replay's cache geometry, checked on the CPU against the recorded capture +// before anything reaches a device. The replay is only meaningful if +// slot_mapping, block_table and block_size agree with each other, and this case +// is what makes that a gate rather than an assumption. It needs no GPU: it is +// the red-first half of the instrument and it runs wherever the artifact is. +TEST_CASE("Primary attention capture decodes to the cache geometry the replay uses" * + doctest::skip(AttnCaptureAbsent())) { + REQUIRE_MESSAGE(!AttnCaptureAbsent(), + "VT_ATTN_PARITY_PRIMARY must name the primary attention-capture directory"); + const char* results = std::getenv("VT_ATTN_PARITY_PRIMARY"); + const std::filesystem::path dir(results); + const std::string label = std::getenv("VT_ATTN_PARITY_LABEL") != nullptr + ? std::getenv("VT_ATTN_PARITY_LABEL") + : "L33-C2-R0"; + const PrimaryBoundary primary = LoadPrimaryBoundary(dir, label, 0); + CHECK(primary.tokens == 66); + CHECK(primary.hq == 1); + CHECK(primary.hkv == 1); + CHECK(primary.dh == 128); + CHECK(primary.block_size == 16); + CHECK(primary.num_reqs == 2); + CHECK(primary.query_start_loc == std::vector({0, 33, 66})); + CHECK(primary.seq_lens == std::vector({33, 33})); + CHECK(primary.block_table.size() == + static_cast(primary.num_reqs * primary.block_table_cols)); + CHECK(std::vector(primary.block_table.begin(), primary.block_table.begin() + 3) == + std::vector({7, 8, 9})); + CHECK(std::vector(primary.block_table.begin() + primary.block_table_cols, + primary.block_table.begin() + primary.block_table_cols + 3) == + std::vector({10, 11, 12})); + // Every token of every request addresses its own block table entry at its own + // offset. A capture whose slot_mapping disagreed with its block_table would + // make the replay write K/V somewhere the kernel never reads. + int64_t addressed = 0, blocks_needed = 0; + for (int64_t r = 0; r < primary.num_reqs; ++r) { + const int64_t keys = primary.seq_lens[static_cast(r)]; + for (int64_t j = 0; j < keys; ++j) { + const int64_t token = primary.query_start_loc[static_cast(r)] + j; + const int64_t slot = primary.slot_mapping[static_cast(token)]; + CAPTURE(r); + CAPTURE(j); + CHECK(slot / primary.block_size == + primary.block_table[static_cast(r * primary.block_table_cols + + j / primary.block_size)]); + CHECK(slot % primary.block_size == j % primary.block_size); + blocks_needed = + std::max(blocks_needed, + primary.block_table[static_cast(r * primary.block_table_cols + + j / primary.block_size)] + + 1); + addressed += 1; + } + } + CHECK(addressed == primary.tokens); + CHECK(blocks_needed == 13); + // The zero-cached-prefix record is what licenses deriving the KV cache from + // `key`/`value` plus `slot_mapping` instead of capturing it. + nlohmann::json record; + std::ifstream(dir / (label + "-attention-0.json")) >> record; + CHECK(record.at("zero_cached_prefix").get()); + CHECK(record.at("kv_cache").at("raw_omitted").is_string()); + std::cout << "[attn-parity] capture geometry: " << addressed << " tokens, " << blocks_needed + << " cache blocks, block_size " << primary.block_size << std::endl; +} + +TEST_CASE("ROCm paged attention replays the primary's captured attention boundary" * + doctest::skip(AttnInstrumentAbsent())) { + REQUIRE_MESSAGE( + !AttnInstrumentAbsent(), + "the attention-boundary instrument needs VT_ROCM_MOE_FIXTURE, " + "VT_ATTN_PARITY_PRIMARY and VT_ATTN_DUMP=1"); + const char* fixture = std::getenv("VT_ROCM_MOE_FIXTURE"); + const char* results_env = std::getenv("VT_ATTN_PARITY_PRIMARY"); + const std::filesystem::path results(results_env); + const char* preamble_env = std::getenv("VT_ATTN_PARITY_PREAMBLE"); + const std::filesystem::path out_dir = + std::getenv("VT_ATTN_PARITY_OUT") != nullptr ? std::filesystem::path(std::getenv("VT_ATTN_PARITY_OUT")) + : std::filesystem::path(); + const std::string label = std::getenv("VT_ATTN_PARITY_LABEL") != nullptr + ? std::getenv("VT_ATTN_PARITY_LABEL") + : "L33-C2-R0"; + + InstallAttnObservers(); + const PrimaryBoundary primary = LoadPrimaryBoundary(results, label, 0); + nlohmann::json report; + report["label"] = label; + report["layer"] = 0; + report["primary_dir"] = results.string(); + + auto& backend = vt::GetBackend(vt::DeviceType::kROCM); + QueueGuard queue(backend); + auto& q = queue.queue; + + // ---- (i) the native boundary, from the unchanged production path -------- + g_record = true; + Generate(fixture, 33, 2, 1); + g_record = false; + REQUIRE(g_attn.have); + // The preamble boundary comes from whichever realization executed: the fused + // op when the backend registers one, the Tier-0 composite otherwise. Both + // describe the same three tensors, so the rows below read the same thing on + // either tree. + if (!g_preamble.have) REQUIRE(g_rope.have); + const bool preamble_fused = g_preamble.have; + // The RAW qkv slice the preamble starts from, and the post-RoPE K it ends at. + // Both realizations norm in place: the fused op reads q3/k3 on entry, and the + // composite's RmsNorm records its input before it writes. The check in the + // composite arm pins that its normed output is the RoPE call's input, so the + // capture is the boundary it claims to be. + std::vector raw_q, raw_k; + if (preamble_fused) { + raw_q = g_preamble.q_in; + raw_k = g_preamble.k_in; + } else { + REQUIRE(g_rope.rms_calls_before >= 2); + const RmsNormCall& q_norm_call = + g_rms_calls[static_cast(g_rope.rms_calls_before - 2)]; + const RmsNormCall& k_norm_call = + g_rms_calls[static_cast(g_rope.rms_calls_before - 1)]; + CHECK(q_norm_call.shape1 == primary.dh); + CHECK(k_norm_call.shape1 == primary.dh); + const WordDiff normed_vs_pre_rope = DiffWords(q_norm_call.out, g_rope.q_pre, primary.dh); + CHECK(normed_vs_pre_rope.different == 0); + raw_q = q_norm_call.in; + raw_k = k_norm_call.in; + } + const std::vector& post_rope_k = preamble_fused ? g_preamble.k_post : g_rope.k_post; + REQUIRE(raw_q.size() == g_attn.query.size()); + REQUIRE(raw_k.size() == g_attn.query.size()); + REQUIRE(post_rope_k.size() == g_attn.query.size()); + REQUIRE(g_attn.tokens == primary.tokens); + CHECK(g_attn.hq == primary.hq); + CHECK(g_attn.hkv == primary.hkv); + CHECK(g_attn.dh == primary.dh); + CHECK(g_attn.block_table_cols >= + (primary.seq_lens[0] + g_attn.block_size - 1) / g_attn.block_size); + CHECK(g_attn.causal); + CHECK(std::fabs(g_attn.scale - primary.scale) < 1e-9f); + // The metadata the replay reuses must be model-independent, or the replay + // would compare two different batches. + CHECK(g_attn.query_start_loc == primary.query_start_loc); + CHECK(g_attn.seq_lens == primary.seq_lens); + CHECK(g_attn.calls == 2); // one PagedAttention per full-attention layer + std::cout << "[attn-parity] native dispatch: " << g_attn.tokens << " tokens, hq " + << g_attn.hq << ", hkv " << g_attn.hkv << ", dh " << g_attn.dh + << ", block_size " << g_attn.block_size << ", blocks " << g_attn.blocks + << ", scale " << g_attn.scale << std::endl; + std::cout << "[attn-parity] native preamble: fused op " << (preamble_fused ? "yes" : "no") + << "; rmsnorm calls before rope " << g_rope.rms_calls_before << " of " + << g_rms_calls.size() << "; rope_neox calls " << g_rope_neox_calls << "; rot " + << (preamble_fused ? g_preamble.rot : g_rope.rot) << std::endl; + + const std::vector max_prob = MaxProbPerRow(primary); + + // Row 0: the recorded comparison, reproduced at this boundary. + const WordDiff run_out = DiffWords(g_attn.out, primary.output, primary.dh); + PrintDiff("A0 native run output vs primary output", run_out); + + // Row B: the preamble term on the two runs' own inputs. + const WordDiff run_q = DiffWords(g_attn.query, primary.query, primary.dh); + const WordDiff run_k = DiffWords(post_rope_k, primary.key, primary.dh); + PrintDiff("B native run Q vs primary Q", run_q); + PrintDiff("B native run K (post-RoPE) vs primary K", run_k); + + // Row F is compared against the raw slice in the preamble block below. + + // The native cache rows the kernel reads must be the post-RoPE K the preamble + // produced, or the capture is not the boundary it claims to be. + std::vector cache_k; + for (int64_t r = 0; r < primary.num_reqs; ++r) { + const int64_t keys = primary.seq_lens[static_cast(r)]; + const std::vector rows = + GatherKv(g_attn.k_cache, g_attn.block_size, g_attn.hkv, g_attn.dh, g_attn.block_table, + g_attn.block_table_cols, r, keys); + cache_k.insert(cache_k.end(), rows.begin(), rows.end()); + } + const WordDiff cache_identity = DiffWords(cache_k, post_rope_k, primary.dh); + std::cout << "[attn-parity] native cache K rows vs post-RoPE k3: " << cache_identity.different + << " / " << cache_identity.words << " differing" << std::endl; + + // ---- (ii) the native kernel on the PRIMARY's captured Q/K/V ------------- + const int64_t bs = primary.block_size; + const int64_t hkv = primary.hkv, dh = primary.dh; + int64_t blocks = 0; + for (int64_t r = 0; r < primary.num_reqs; ++r) + for (int64_t c = 0; c < (primary.seq_lens[static_cast(r)] + bs - 1) / bs; ++c) + blocks = std::max( + blocks, primary.block_table[static_cast(r * primary.block_table_cols + c)] + 1); + const size_t tile_words = static_cast(bs * hkv * dh); + const size_t slice_words = static_cast(blocks) * tile_words; + std::vector cache(2 * slice_words, 0); + int64_t scattered = 0; + auto scatter = [&](int which, const std::vector& rows) { + for (int64_t r = 0; r < primary.num_reqs; ++r) { + for (int64_t j = 0; j < primary.seq_lens[static_cast(r)]; ++j) { + const int64_t token = primary.query_start_loc[static_cast(r)] + j; + ScatterCacheSlot(cache, blocks, bs, hkv, dh, which, + primary.slot_mapping[static_cast(token)], + rows.data() + static_cast(token * hkv * dh)); + if (which == 0) scattered += 1; + } + } + }; + + Buffer cache_buffer(backend, cache.size() * sizeof(uint16_t)); + Buffer query_buffer(backend, primary.query.size() * sizeof(uint16_t)); + Buffer out_buffer(backend, primary.query.size() * sizeof(uint16_t)); + Buffer qsl_buffer(backend, primary.query_start_loc.size() * sizeof(int32_t)); + Buffer seq_buffer(backend, primary.seq_lens.size() * sizeof(int32_t)); + Buffer table_buffer(backend, primary.block_table.size() * sizeof(int32_t)); + backend.Copy(q, qsl_buffer.data, primary.query_start_loc.data(), + primary.query_start_loc.size() * sizeof(int32_t)); + backend.Copy(q, seq_buffer.data, primary.seq_lens.data(), + primary.seq_lens.size() * sizeof(int32_t)); + backend.Copy(q, table_buffer.data, primary.block_table.data(), + primary.block_table.size() * sizeof(int32_t)); + backend.Synchronize(q); + + auto kv_slice = [&](int which) { + // The two dim-1 slices of the (blocks, 2, block_size, hkv, dh) flash cache: + // K starts at the buffer base, V one block-slice further in, and both carry + // the block stride 2*block_size*hkv*dh (dense_attn_block.h:356-373). + vt::Tensor t; + t.data = static_cast(cache_buffer.data) + + static_cast(which) * tile_words * sizeof(uint16_t); + t.dtype = vt::DType::kBF16; + t.device = q.device; + t.rank = 4; + t.shape[0] = blocks; + t.shape[1] = bs; + t.shape[2] = hkv; + t.shape[3] = dh; + t.stride[0] = 2 * bs * hkv * dh; + t.stride[1] = hkv * dh; + t.stride[2] = dh; + t.stride[3] = 1; + return t; + }; + vt::Tensor replay_query = + Tensor(query_buffer.data, vt::DType::kBF16, q.device, {primary.tokens, primary.hq, primary.dh}); + vt::Tensor replay_out = + Tensor(out_buffer.data, vt::DType::kBF16, q.device, {primary.tokens, primary.hq, primary.dh}); + vt::Tensor replay_qsl = Tensor(qsl_buffer.data, vt::DType::kI32, q.device, {primary.num_reqs + 1}); + vt::Tensor replay_seq = Tensor(seq_buffer.data, vt::DType::kI32, q.device, {primary.num_reqs}); + vt::Tensor replay_table = Tensor(table_buffer.data, vt::DType::kI32, q.device, + {primary.num_reqs, primary.block_table_cols}); + vt::PagedAttentionArgs replay_args; + replay_args.scale = primary.scale; + replay_args.causal = true; + replay_args.query_start_loc_host = primary.query_start_loc.data(); + replay_args.max_seq_len = *std::max_element(primary.seq_lens.begin(), primary.seq_lens.end()); + + auto run_replay = [&](const std::vector& k_rows, + const std::vector& v_rows, + const std::vector& query_rows) { + std::fill(cache.begin(), cache.end(), static_cast(0)); + scatter(0, k_rows); + scatter(1, v_rows); + backend.Copy(q, cache_buffer.data, cache.data(), cache.size() * sizeof(uint16_t)); + backend.Copy(q, query_buffer.data, query_rows.data(), + query_rows.size() * sizeof(uint16_t)); + backend.Synchronize(q); + vt::PagedAttention(q, replay_out, replay_query, kv_slice(0), kv_slice(1), replay_table, + replay_seq, replay_qsl, replay_args); + backend.Synchronize(q); + return ReadDeviceBf16(q, replay_out); + }; + + // Self-consistency first. The kernel addresses the cache by logical + // (request, key position), so the native run's own Q/K/V through the same + // kernel at block_size 16 must reproduce the native run's own output at + // block_size 64. A difference here means the replay's cache construction is + // wrong, and row C below would mean nothing. + std::vector native_v; + for (int64_t r = 0; r < primary.num_reqs; ++r) { + const int64_t keys = primary.seq_lens[static_cast(r)]; + const std::vector rows = + GatherKv(g_attn.v_cache, g_attn.block_size, g_attn.hkv, g_attn.dh, g_attn.block_table, + g_attn.block_table_cols, r, keys); + native_v.insert(native_v.end(), rows.begin(), rows.end()); + } + const std::vector replay_self = run_replay(post_rope_k, native_v, g_attn.query); + const WordDiff replay_self_diff = DiffWords(replay_self, g_attn.out, primary.dh); + CHECK(replay_self_diff.different == 0); + PrintDiff("C0 replay self-consistency: native Q/K/V at the primary's geometry vs native run out", + replay_self_diff); + + const std::vector replay = run_replay(primary.key, primary.value, primary.query); + const WordDiff replay_diff = DiffWords(replay, primary.output, primary.dh); + CHECK(scattered == 2 * primary.tokens); + std::cout << "[attn-parity] replay: " << blocks << " cache blocks of " << bs + << " at the primary's geometry, " << scattered << " K/V rows written by direct copy" + << std::endl; + PrintDiff("C native PagedAttention on PRIMARY Q/K/V vs primary output", replay_diff); + PrintDiffShape(primary, replay_diff, max_prob); + + // ---- (iii) the native preamble on the PRIMARY's captured qkv ------------ + WordDiff preamble_q, preamble_k, neox_q, neox_k, table_diff, qkv_q, qkv_k; + WordDiff production_q, production_k; + WordDiff handcall_q, handcall_k; + int64_t hand_fused_identity = -1; + bool have_preamble = false; + bool have_production = false; + bool have_handcall = false; + if (preamble_env != nullptr) { + const PrimaryPreamble preamble = LoadPrimaryPreamble(preamble_env, label, 0); + REQUIRE(preamble.tokens == primary.tokens); + REQUIRE(preamble.dh == primary.dh); + REQUIRE(preamble.rot == primary.dh); + // The two captures must describe the same rows, or the lane compares + // different tokens. + const WordDiff join_q = + DiffWords(JoinRopeHalves(preamble, preamble.q_left, preamble.q_right), primary.query, + primary.dh); + const WordDiff join_k = + DiffWords(JoinRopeHalves(preamble, preamble.k_left, preamble.k_right), primary.key, + primary.dh); + CHECK(join_q.different == 0); + CHECK(join_k.different == 0); + + const vllm::HfConfig config = + vllm::LoadHfConfig((std::filesystem::path(fixture) / "config.json").string()); + CHECK(config.rope_parameters.rope_type != "llama3"); + vt::RopeArgs rope; // dense_attn::MakeRopeArgs (dense_attn_block.h:108-123), default rope type + rope.base = static_cast(config.rope_theta); + rope.rotary_dim = static_cast(config.rotary_dim); + CHECK(rope.base == 10000000.0f); + CHECK(rope.rotary_dim == static_cast(primary.dh)); + const float eps = static_cast(config.rms_norm_eps); + + const size_t rows = static_cast(preamble.tokens * preamble.dh); + std::vector q2(rows), k2(rows); + for (int64_t t = 0; t < preamble.tokens; ++t) { + for (int64_t d = 0; d < preamble.dh; ++d) { + q2[static_cast(t * preamble.dh + d)] = + preamble.qkv[static_cast(t * preamble.qkv_width + d)]; + k2[static_cast(t * preamble.dh + d)] = + preamble.qkv[static_cast(t * preamble.qkv_width + preamble.dh + d)]; + } + } + // Row F: the two runs' qkv projection outputs, so a preamble difference + // cannot be blamed on the GEMM upstream of it. The residual row already + // found this boundary exact; this confirms it on this capture pair. + qkv_q = DiffWords(q2, raw_q, preamble.dh); + qkv_k = DiffWords(k2, raw_k, preamble.dh); + PrintDiff("F primary qkv q slice vs native run pre-norm q", qkv_q); + PrintDiff("F primary qkv k slice vs native run pre-norm k", qkv_k); + std::vector positions(static_cast(preamble.tokens)); + std::vector row_index(static_cast(preamble.tokens)); + for (int64_t t = 0; t < preamble.tokens; ++t) { + positions[static_cast(t)] = static_cast(preamble.positions[static_cast(t)]); + row_index[static_cast(t)] = static_cast(t); + } + Buffer q2_buffer(backend, q2.size() * sizeof(uint16_t)); + Buffer k2_buffer(backend, k2.size() * sizeof(uint16_t)); + Buffer qg_buffer(backend, preamble.q_gamma.size() * sizeof(uint16_t)); + Buffer kg_buffer(backend, preamble.k_gamma.size() * sizeof(uint16_t)); + Buffer pos_buffer(backend, positions.size() * sizeof(int32_t)); + Buffer idx_buffer(backend, row_index.size() * sizeof(int32_t)); + Buffer cs32_buffer(backend, rows * sizeof(float)); + Buffer cs16_buffer(backend, rows * sizeof(uint16_t)); + backend.Copy(q, q2_buffer.data, q2.data(), q2.size() * sizeof(uint16_t)); + backend.Copy(q, k2_buffer.data, k2.data(), k2.size() * sizeof(uint16_t)); + backend.Copy(q, qg_buffer.data, preamble.q_gamma.data(), + preamble.q_gamma.size() * sizeof(uint16_t)); + backend.Copy(q, kg_buffer.data, preamble.k_gamma.data(), + preamble.k_gamma.size() * sizeof(uint16_t)); + backend.Copy(q, pos_buffer.data, positions.data(), positions.size() * sizeof(int32_t)); + backend.Copy(q, idx_buffer.data, row_index.data(), row_index.size() * sizeof(int32_t)); + backend.Synchronize(q); + + vt::Tensor t_q2 = Tensor(q2_buffer.data, vt::DType::kBF16, q.device, + {preamble.tokens, preamble.dh}); + vt::Tensor t_k2 = Tensor(k2_buffer.data, vt::DType::kBF16, q.device, + {preamble.tokens, preamble.dh}); + vt::Tensor t_qg = Tensor(qg_buffer.data, vt::DType::kBF16, q.device, {preamble.dh}); + vt::Tensor t_kg = Tensor(kg_buffer.data, vt::DType::kBF16, q.device, {preamble.dh}); + vt::Tensor t_pos = Tensor(pos_buffer.data, vt::DType::kI32, q.device, {preamble.tokens}); + vt::Tensor t_idx = Tensor(idx_buffer.data, vt::DType::kI32, q.device, {preamble.tokens}); + vt::Tensor t_cs32 = Tensor(cs32_buffer.data, vt::DType::kF32, q.device, + {preamble.tokens, preamble.rot}); + vt::Tensor t_cs16 = Tensor(cs16_buffer.data, vt::DType::kBF16, q.device, + {preamble.tokens, preamble.rot}); + vt::RopeCosSinCache(q, t_cs32, t_pos, rope); + vt::CastBf16(q, t_cs16, t_cs32); + const std::vector native_table = ReadDeviceBf16(q, t_cs16); + // The native per-step table row t already encodes positions[t]; the primary's + // cache is indexed by the real position. + std::vector expected_table(static_cast(preamble.tokens * preamble.rot)); + for (int64_t t = 0; t < preamble.tokens; ++t) + for (int64_t d = 0; d < preamble.rot; ++d) + expected_table[static_cast(t * preamble.rot + d)] = + preamble.cos_sin[static_cast(preamble.positions[static_cast(t)] * + preamble.rot + d)]; + table_diff = DiffWords(native_table, expected_table, preamble.rot); + PrintDiff("E native cos|sin table vs primary cos_sin at positions", table_diff); + + vt::RmsNorm(q, t_q2, t_q2, t_qg, vt::RmsNormArgs{eps, false}); + vt::RmsNorm(q, t_k2, t_k2, t_kg, vt::RmsNormArgs{eps, false}); + backend.Synchronize(q); + const std::vector normed_q = ReadDeviceBf16(q, t_q2); + const std::vector normed_k = ReadDeviceBf16(q, t_k2); + REQUIRE(normed_q.size() == rows); + + Buffer q3_buffer(backend, rows * sizeof(uint16_t)); + Buffer k3_buffer(backend, rows * sizeof(uint16_t)); + backend.Copy(q, q3_buffer.data, normed_q.data(), rows * sizeof(uint16_t)); + backend.Copy(q, k3_buffer.data, normed_k.data(), rows * sizeof(uint16_t)); + backend.Synchronize(q); + vt::Tensor t_q3 = Tensor(q3_buffer.data, vt::DType::kBF16, q.device, + {preamble.tokens, 1, preamble.dh}); + vt::Tensor t_k3 = Tensor(k3_buffer.data, vt::DType::kBF16, q.device, + {preamble.tokens, 1, preamble.dh}); + vt::RopeFromCache(q, t_q3, &t_k3, t_idx, t_cs16, rope); + backend.Synchronize(q); + preamble_q = DiffWords(ReadDeviceBf16(q, t_q3), primary.query, primary.dh); + preamble_k = DiffWords(ReadDeviceBf16(q, t_k3), primary.key, primary.dh); + PrintDiff("D native preamble (RmsNorm+RopeFromCache) on primary qkv vs primary Q", preamble_q); + PrintDiff("D native preamble (RmsNorm+RopeFromCache) on primary qkv vs primary K", preamble_k); + + // The same normed q/k through the non-cached native RoPE, to separate the + // norm store from the cos|sin source. + Buffer q4_buffer(backend, rows * sizeof(uint16_t)); + Buffer k4_buffer(backend, rows * sizeof(uint16_t)); + backend.Copy(q, q4_buffer.data, normed_q.data(), rows * sizeof(uint16_t)); + backend.Copy(q, k4_buffer.data, normed_k.data(), rows * sizeof(uint16_t)); + backend.Synchronize(q); + vt::Tensor t_q4 = Tensor(q4_buffer.data, vt::DType::kBF16, q.device, + {preamble.tokens, 1, preamble.dh}); + vt::Tensor t_k4 = Tensor(k4_buffer.data, vt::DType::kBF16, q.device, + {preamble.tokens, 1, preamble.dh}); + vt::RopeNeox(q, t_q4, t_k4, t_pos, rope); + backend.Synchronize(q); + neox_q = DiffWords(ReadDeviceBf16(q, t_q4), primary.query, primary.dh); + neox_k = DiffWords(ReadDeviceBf16(q, t_k4), primary.key, primary.dh); + PrintDiff("D' native preamble (RmsNorm+RopeNeox) on primary qkv vs primary Q", neox_q); + PrintDiff("D' native preamble (RmsNorm+RopeNeox) on primary qkv vs primary K", neox_k); + have_preamble = true; + + // D2: the PRODUCTION preamble on the primary's own qkv. D above measures the + // Tier-0 composite, which keeps the pre-RoPE BF16 store on every backend that + // registers no fused op; D2 measures the path this repair changes. The + // binding is the fused branch's own BF16 binding (dense_attn_block.h:594-613): + // bf16 norm weights, the bf16 per-step cos|sin table, the identity row index. + if (vt::OpRegistered(vt::OpId::kAttnQkNormRope, vt::DeviceType::kROCM)) { + auto fused_preamble_op = reinterpret_cast( + vt::GetOp(vt::OpId::kAttnQkNormRope, vt::DeviceType::kROCM)); + REQUIRE(fused_preamble_op != nullptr); + Buffer q5_buffer(backend, rows * sizeof(uint16_t)); + Buffer k5_buffer(backend, rows * sizeof(uint16_t)); + backend.Copy(q, q5_buffer.data, q2.data(), rows * sizeof(uint16_t)); + backend.Copy(q, k5_buffer.data, k2.data(), rows * sizeof(uint16_t)); + backend.Synchronize(q); + vt::Tensor t_q5 = Tensor(q5_buffer.data, vt::DType::kBF16, q.device, + {preamble.tokens, 1, preamble.dh}); + vt::Tensor t_k5 = Tensor(k5_buffer.data, vt::DType::kBF16, q.device, + {preamble.tokens, 1, preamble.dh}); + fused_preamble_op(q, t_q5, t_k5, t_qg, t_kg, t_cs16, t_idx, + vt::RmsNormArgs{eps, false}, rope); + backend.Synchronize(q); + production_q = DiffWords(ReadDeviceBf16(q, t_q5), primary.query, primary.dh); + production_k = DiffWords(ReadDeviceBf16(q, t_k5), primary.key, primary.dh); + PrintDiff("D2 production preamble (fused qk-norm-rope) on primary qkv vs primary Q", + production_q); + PrintDiff("D2 production preamble (fused qk-norm-rope) on primary qkv vs primary K", + production_k); + have_production = true; + if (!out_dir.empty()) { + WriteWords(out_dir / (label + "-native-production-preamble-q.bin"), + ReadDeviceBf16(q, t_q5)); + } + + // D3: the HAND-CALL realization on the same primary qkv, through the + // op-layer entry the ADOPT=0 branch calls (dense_attn_block.h:645-649). + // Before the stride repair this threw + // "attn_qk_norm_rope: row stride must be the inner dimension" for every + // Dh > 1, which made the documented same-binary A/B lever unusable; the two + // realizations must also agree byte-for-byte, which is the recipe's + // composite contract and what makes the lever an A/B rather than two arms. + Buffer q6_buffer(backend, rows * sizeof(uint16_t)); + Buffer k6_buffer(backend, rows * sizeof(uint16_t)); + backend.Copy(q, q6_buffer.data, q2.data(), rows * sizeof(uint16_t)); + backend.Copy(q, k6_buffer.data, k2.data(), rows * sizeof(uint16_t)); + backend.Synchronize(q); + vt::Tensor t_q6 = Tensor(q6_buffer.data, vt::DType::kBF16, q.device, + {preamble.tokens, 1, preamble.dh}); + vt::Tensor t_k6 = Tensor(k6_buffer.data, vt::DType::kBF16, q.device, + {preamble.tokens, 1, preamble.dh}); + vt::AttnQkNormRope(q, t_q6, t_k6, t_qg, t_kg, t_cs16, t_idx, + vt::RmsNormArgs{eps, false}, rope); + backend.Synchronize(q); + const std::vector hand_q = ReadDeviceBf16(q, t_q6); + const std::vector hand_k = ReadDeviceBf16(q, t_k6); + handcall_q = DiffWords(hand_q, primary.query, primary.dh); + handcall_k = DiffWords(hand_k, primary.key, primary.dh); + PrintDiff("D3 hand-call preamble (vt::AttnQkNormRope) on primary qkv vs primary Q", + handcall_q); + PrintDiff("D3 hand-call preamble (vt::AttnQkNormRope) on primary qkv vs primary K", + handcall_k); + const WordDiff hand_vs_fused_q = DiffWords(hand_q, ReadDeviceBf16(q, t_q5), primary.dh); + const WordDiff hand_vs_fused_k = DiffWords(hand_k, ReadDeviceBf16(q, t_k5), primary.dh); + PrintDiff("D3 hand-call vs D2 fused-op, same primary qkv, Q", hand_vs_fused_q); + PrintDiff("D3 hand-call vs D2 fused-op, same primary qkv, K", hand_vs_fused_k); + hand_fused_identity = hand_vs_fused_q.different + hand_vs_fused_k.different; + CHECK(hand_fused_identity == 0); + have_handcall = true; + } + + if (!out_dir.empty()) { + WriteWords(out_dir / (label + "-native-preamble-rope-from-cache-q.bin"), + ReadDeviceBf16(q, t_q3)); + WriteWords(out_dir / (label + "-native-preamble-rope-neox-q.bin"), ReadDeviceBf16(q, t_q4)); + WriteWords(out_dir / (label + "-native-cos-sin-table.bin"), native_table); + } + } else { + std::cout << "[attn-parity] preamble lanes SKIPPED: VT_ATTN_PARITY_PREAMBLE unset" + << std::endl; + } + + // ---- focused cases for the #3115 repair --------------------------------- + // KERNEL. Under the primary's arithmetic (the softmax probability narrowed to + // the value dtype before the value accumulate, prefix_prefill.py:471) the + // native kernel on the primary's OWN Q/K/V must reproduce the primary's + // output byte-for-byte. Red at 2918 of 8448 before that repair. + std::cout << "[attn-parity][focused-kernel] C " << replay_diff.different << " / " + << replay_diff.words << " differing words (want 0)" << std::endl; + CHECK(replay_diff.different == 0); + + // PREAMBLE. Under the primary's boundary (f32 carrier through RoPE, narrowed + // once at the store) the native run and the primary must agree on Q and K up + // to the CPU model's own 1-word residue against the primary's capture. Red at + // 1569 Q / 1542 K before that repair. + std::cout << "[attn-parity][focused-preamble] B " << run_q.different << " Q / " + << run_k.different << " K differing words (want <= 1)" << std::endl; + CHECK(run_q.different <= 1); + CHECK(run_k.different <= 1); + if (have_production) { + std::cout << "[attn-parity][focused-preamble-op] D2 " << production_q.different << " Q / " + << production_k.different << " K differing words (want <= 1)" << std::endl; + CHECK(production_q.different <= 1); + CHECK(production_k.different <= 1); + } + if (have_handcall) { + std::cout << "[attn-parity][focused-handcall] D3 " << handcall_q.different << " Q / " + << handcall_k.different << " K differing words (want <= 1); hand-call vs fused op " + << hand_fused_identity << " differing words (want 0)" << std::endl; + CHECK(handcall_q.different <= 1); + CHECK(handcall_k.different <= 1); + CHECK(hand_fused_identity == 0); + } + // BOTH SETTINGS OF THE ADOPTION SWITCH MUST REACH THE REGISTERED OP. On ROCm the + // bf16 production preamble is the fused realization either way: + // VT_FUSED_CHAIN_ADOPT=1 enters it through vt::FusedChain's fast_op dispatch, + // =0 through the hand-call at dense_attn_block.h:645-649, which is the same + // vt::AttnQkNormRope entry. `preamble_fused` records that the OP ran, not which + // branch called it, so the branch is pinned by the second check: the Tier-0 + // composite must NOT have run, or the =0 run would be measuring the standalone + // RmsNorm+RopeFromCache sequence instead of the hand-call this finding names. + // The device A/B -- same binary, same workload, one variable -- is the other half. + const bool adopt = vllm::dense_attn::FusedChainAdoptEnabled(); + std::cout << "[attn-parity][focused-adopt] VT_FUSED_CHAIN_ADOPT=" << (adopt ? 1 : 0) + << "; registered fused op ran: " << (preamble_fused ? "yes" : "no") + << "; composite fallback ran: " << (g_rope.have ? "yes" : "no") << std::endl; + CHECK(preamble_fused); + CHECK_FALSE(g_rope.have); + + // ---- the verdict -------------------------------------------------------- + // #3115's two hypotheses: the Q/K preamble, the kernel arithmetic, or both. + std::string verdict; + if (replay_diff.different == 0) { + verdict = "PREAMBLE-ONLY"; + } else if (run_q.different == 0 && run_k.different == 0) { + verdict = "KERNEL-ONLY"; + } else { + verdict = "BOTH"; + } + std::cout << "[attn-parity] VERDICT " << verdict << ": native-on-primary-inputs " + << replay_diff.different << "/" << replay_diff.words << "; preamble (native run Q/K) " + << run_q.different << "/" << run_q.words << " and " << run_k.different << "/" + << run_k.words << std::endl; + if (replay_diff.different == 0) + std::cout << "[attn-parity] repair site: dense_attn_block.h:637-653 (do not touch the kernel)" + << std::endl; + else + std::cout << "[attn-parity] repair site: rocm_paged_attn.hip:527-531 (probability dtype), " + ":196-198 (FastExp), :512 (warp reduction order)" + << std::endl; + + report["native_dispatch"] = {{"tokens", g_attn.tokens}, + {"hq", g_attn.hq}, + {"hkv", g_attn.hkv}, + {"dh", g_attn.dh}, + {"block_size", g_attn.block_size}, + {"blocks", g_attn.blocks}, + {"scale", g_attn.scale}}; + report["A0_native_run_output_vs_primary_output"] = run_out.different; + report["B_native_run_q_vs_primary_q"] = run_q.different; + report["B_native_run_k_vs_primary_k"] = run_k.different; + report["C_native_kernel_on_primary_inputs_vs_primary_output"] = replay_diff.different; + report["C_words"] = replay_diff.words; + report["C_per_row"] = replay_diff.row_different; + report["C_distance_1"] = replay_diff.distance[1]; + report["C_distance_2"] = replay_diff.distance[2]; + report["C_distance_3"] = replay_diff.distance[3]; + report["C_distance_ge4"] = replay_diff.distance[4]; + report["C0_replay_self_consistency"] = replay_self_diff.different; + report["cache_k_identity_different"] = cache_identity.different; + report["preamble_fused"] = preamble_fused; + if (have_preamble) { + report["D_preamble_rope_cache_q"] = preamble_q.different; + report["D_preamble_rope_cache_k"] = preamble_k.different; + report["D_preamble_rope_neox_q"] = neox_q.different; + report["D_preamble_rope_neox_k"] = neox_k.different; + report["E_cos_sin_table"] = table_diff.different; + report["E_cos_sin_words"] = table_diff.words; + report["F_qkv_q"] = qkv_q.different; + report["F_qkv_k"] = qkv_k.different; + } + report["D2_production_preamble_measured"] = have_production; + if (have_production) { + report["D2_production_preamble_q"] = production_q.different; + report["D2_production_preamble_k"] = production_k.different; + } + report["D3_handcall_preamble_measured"] = have_handcall; + report["fused_chain_adopt"] = vllm::dense_attn::FusedChainAdoptEnabled(); + if (have_handcall) { + report["D3_handcall_preamble_q"] = handcall_q.different; + report["D3_handcall_preamble_k"] = handcall_k.different; + report["D3_handcall_vs_fused_differing_words"] = hand_fused_identity; + } + report["verdict"] = verdict; + if (!out_dir.empty()) { + std::filesystem::create_directories(out_dir); + std::ofstream(out_dir / (label + "-attn-parity-report.json")) << report.dump(2) << '\n'; + WriteWords(out_dir / (label + "-native-attention-0-query.bin"), g_attn.query); + WriteWords(out_dir / (label + "-native-attention-0-output.bin"), g_attn.out); + WriteWords(out_dir / (label + "-native-attention-0-k-cache.bin"), g_attn.k_cache); + WriteWords(out_dir / (label + "-native-attention-0-v-cache.bin"), g_attn.v_cache); + WriteWords(out_dir / (label + "-native-attention-0-k3.bin"), post_rope_k); + WriteWords(out_dir / (label + "-replay-output.bin"), replay); + } +} + +// --------------------------------------------------------------------------- +// The key-walk geometry the primary selects, one case per arm. +// +// The recorded capture measures ONE arm: prefix_prefill over a request with no +// cached prefix, where BLOCK_N = 64 and the 33 keys fit one tile. The same native +// kernel also serves pure decode (the Triton decode arm tiles by +// min(block_size, 128)) and a chunked prefill (a 32-key context tile walk before +// the 64-key chunk walk), and it served both of those with the prefill tile +// before this repair. These cases synthesize the data that separates the arms +// and compare the device against the primary's own key walk for the arm. +// --------------------------------------------------------------------------- +namespace { + +void CheckArmGeometry(const char* label, const SyntheticArm& arm, bool chunked, + const std::vector& wrong_widths) { + auto& backend = vt::GetBackend(vt::DeviceType::kROCM); + QueueGuard queue(backend); + auto& q = queue.queue; + const std::vector device = RunSyntheticArm(backend, q, arm); + const WordDiff correct = DiffWords(device, ArmModelWords(arm, chunked, 0), arm.dh); + std::cout << "[arm-geometry] " << label << ": " << arm.num_reqs << " requests, query_len " + << arm.query_len << ", seq_len " << arm.seq_len << ", context " + << (arm.seq_len - arm.query_len) << ", block_size " << arm.block_size << ", high key " + << arm.high_key << ", " << arm.total_q << " query rows" << std::endl; + PrintArmDiff("device vs the primary arm's own tiles", correct, arm.total_q); + CHECK_MESSAGE(correct.different <= kArmNoiseWords, + label << ": the device must follow the primary arm's key walk (" + << correct.different << " differing words, noise floor " << kArmNoiseWords + << ")"); + for (int64_t width : wrong_widths) { + const WordDiff d = DiffWords(device, ArmModelWords(arm, chunked, width), arm.dh); + PrintArmDiff("device vs a uniform tiling of that width", d, arm.total_q); + CHECK_MESSAGE(d.different >= kArmSignalWords, + label << ": a uniform " << width + << "-key tiling must not reproduce the device (" << d.different + << " differing words, signal floor " << kArmSignalWords << ")"); + } +} + +} // namespace + +// THE PREFILL ARM: max_query_len > 1, no cached prefix -> prefix_prefill with +// BLOCK_N = 64 (prefix_prefill.py:955-966). 65 query tokens so that 33 rows per +// request have a key range that reaches past key 32, and the high key sits AT 32: +// one key beyond the 32-boundary, so a 32-key tile and a 64-key tile take their +// reference max over different key sets and narrow the other 32 keys against +// different scales. +TEST_CASE("ROCm paged attention uses the primary's 64-key prefill tile" * + doctest::skip(RocmBackendAbsent())) { + REQUIRE_MESSAGE(!RocmBackendAbsent(), + "this case measures the ROCm arm's key tiling and needs the ROCm backend " + "registered"); + CheckArmGeometry("prefill, empty context", BuildSyntheticArm(4, 65, 65, 16, 32), true, {32, 16}); +} + +// THE CHUNKED-CONTEXT ARM: max_query_len > 1 with a cached prefix -> prefix_prefill +// walks the context in TRITON_BLOCK_SIZE = 32 key tiles from key 0 and then the +// chunk in BLOCK_N = 64 tiles anchored at the chunk start (:965, :1007, :231-343, +// :369). 37 cached keys and 3 query tokens; the high key is at 33, inside the +// second context tile, so a 64-key context tile would put it in the same tile as +// keys 0..31 and narrow those at the high scale instead of at 1. +TEST_CASE("ROCm paged attention uses the primary's 32-key context tile" * + doctest::skip(RocmBackendAbsent())) { + REQUIRE_MESSAGE(!RocmBackendAbsent(), + "this case measures the ROCm arm's key tiling and needs the ROCm backend " + "registered"); + CheckArmGeometry("chunked prefill, 37-key context", BuildSyntheticArm(4, 3, 40, 16, 33), true, + {64}); +} + +// THE PURE-DECODE ARM: max_query_len == 1 -> the Triton decode kernel tiles the +// key range by TRITON_BLOCK_SIZE = min(block_size, 128) = 16 for this cache +// (chunked_prefill_paged_decode.py:444-445, :147-149, :244). A 16-wide tile puts +// key 17 in its own tile; a 32- or 64-wide one puts it with keys 0..15, whose +// probability is then narrowed at the high scale. +TEST_CASE("ROCm paged attention uses the primary's decode tile" * + doctest::skip(RocmBackendAbsent())) { + REQUIRE_MESSAGE(!RocmBackendAbsent(), + "this case measures the ROCm arm's key tiling and needs the ROCm backend " + "registered"); + CheckArmGeometry("pure decode, block_size 16", BuildSyntheticArm(8, 1, 40, 16, 17), false, + {32, 64}); +} + +// --------------------------------------------------------------------------- +// The hand-call realization of the qk-norm-RoPE preamble, on synthesized data. +// +// The production bf16 preamble has two realizations that must agree: the recipe's +// fast op through vt::FusedChain (VT_FUSED_CHAIN_ADOPT=1) and the hand-call +// vt::AttnQkNormRope the fallback branch reaches under VT_FUSED_CHAIN_ADOPT=0 +// (dense_attn_block.h:645-649). Before the op-layer stride repair the second one +// threw for every Dh > 1, so the documented same-binary A/B lever could not run +// at all on this workload; this case runs both in one process, on identical +// inputs, and requires byte identity. +// --------------------------------------------------------------------------- +TEST_CASE("ROCm bf16 qk-norm-rope: the hand-call realization matches the fused recipe" * + doctest::skip(RocmBackendAbsent())) { + REQUIRE_MESSAGE(!RocmBackendAbsent(), + "this case measures the ROCm arm's qk-norm-RoPE hand-call realization and " + "needs the ROCm backend registered"); + REQUIRE_MESSAGE(vt::OpRegistered(vt::OpId::kAttnQkNormRope, vt::DeviceType::kROCM), + "this case measures the Recipe-vs-hand-call pair on a backend that registers " + "the recipe's fast realization"); + auto& backend = vt::GetBackend(vt::DeviceType::kROCM); + QueueGuard queue(backend); + auto& q = queue.queue; + const int64_t tokens = 8, dh = 128, rot = 128; + const size_t rows = static_cast(tokens * dh); + std::vector q_words(rows), k_words(rows); + std::vector q_gamma(static_cast(dh)), k_gamma(static_cast(dh)); + for (size_t i = 0; i < rows; ++i) { + q_words[i] = static_cast(0x3C00 + ((i * 37 + 11) % 1024)); // [1, 8) + k_words[i] = static_cast(0x3C00 + ((i * 53 + 29) % 1024)); + } + // Distinct q/k weights, so a realization that swapped them could not agree. + for (size_t i = 0; i < q_gamma.size(); ++i) { + q_gamma[i] = static_cast(0x3F00 + (i % 64)); // [0.5, 1) + k_gamma[i] = static_cast(0x3E80 + ((i * 3) % 96)); // [0.25, 0.625) + } + std::vector positions(static_cast(tokens)); + for (int64_t t = 0; t < tokens; ++t) positions[static_cast(t)] = static_cast(t); + + Buffer qb(backend, rows * sizeof(uint16_t)), kb(backend, rows * sizeof(uint16_t)); + Buffer qc(backend, rows * sizeof(uint16_t)), kc(backend, rows * sizeof(uint16_t)); + Buffer qgb(backend, q_gamma.size() * sizeof(uint16_t)); + Buffer kgb(backend, k_gamma.size() * sizeof(uint16_t)); + Buffer pb(backend, positions.size() * sizeof(int32_t)); + Buffer cs32(backend, rows * sizeof(float)), cs16(backend, rows * sizeof(uint16_t)); + backend.Copy(q, qb.data, q_words.data(), rows * sizeof(uint16_t)); + backend.Copy(q, qc.data, q_words.data(), rows * sizeof(uint16_t)); + backend.Copy(q, kb.data, k_words.data(), rows * sizeof(uint16_t)); + backend.Copy(q, kc.data, k_words.data(), rows * sizeof(uint16_t)); + backend.Copy(q, qgb.data, q_gamma.data(), q_gamma.size() * sizeof(uint16_t)); + backend.Copy(q, kgb.data, k_gamma.data(), k_gamma.size() * sizeof(uint16_t)); + backend.Copy(q, pb.data, positions.data(), positions.size() * sizeof(int32_t)); + backend.Synchronize(q); + + vt::RopeArgs rope; + rope.base = 1000000.0f; + rope.rotary_dim = static_cast(rot); + const vt::RmsNormArgs norm{1e-6f, false}; + vt::Tensor t_cs32 = Tensor(cs32.data, vt::DType::kF32, q.device, {tokens, rot}); + vt::Tensor t_cs16 = Tensor(cs16.data, vt::DType::kBF16, q.device, {tokens, rot}); + vt::Tensor t_pos = Tensor(pb.data, vt::DType::kI32, q.device, {tokens}); + vt::RopeCosSinCache(q, t_cs32, t_pos, rope); + vt::CastBf16(q, t_cs16, t_cs32); + backend.Synchronize(q); + + // Arm 1 (VT_FUSED_CHAIN_ADOPT=1): the recipe's fast realization, with exactly the + // binding the model builds (dense_attn_block.h:594-613): the 2-D norm view and + // the 3-D rope view alias one buffer, bf16 weights, the bf16 cache, identity rows. + vt::Tensor f_q2 = Tensor(qb.data, vt::DType::kBF16, q.device, {tokens, dh}); + vt::Tensor f_k2 = Tensor(kb.data, vt::DType::kBF16, q.device, {tokens, dh}); + vt::Tensor f_q3 = Tensor(qb.data, vt::DType::kBF16, q.device, {tokens, 1, dh}); + vt::Tensor f_k3 = Tensor(kb.data, vt::DType::kBF16, q.device, {tokens, 1, dh}); + vt::Tensor t_qg = Tensor(qgb.data, vt::DType::kBF16, q.device, {dh}); + vt::Tensor t_kg = Tensor(kgb.data, vt::DType::kBF16, q.device, {dh}); + vt::FusedBinding binding; + binding.op[0] = &f_q2; + binding.op[1] = &t_qg; + binding.op[2] = &f_k2; + binding.op[3] = &t_kg; + binding.op[4] = &f_q3; + binding.op[5] = &f_k3; + binding.op[6] = &t_cs16; + binding.op[7] = &t_pos; + binding.n = 8; + vt::FusedParams params; + params.eps = norm.eps; + params.rope = rope; + vt::FusedChain(q, vt::kAttnQkNormRope, binding, params); + backend.Synchronize(q); + + // Arm 2 (VT_FUSED_CHAIN_ADOPT=0): the hand-call through the op layer, the same + // entry the model's fallback branch reaches. + vt::Tensor h_q3 = Tensor(qc.data, vt::DType::kBF16, q.device, {tokens, 1, dh}); + vt::Tensor h_k3 = Tensor(kc.data, vt::DType::kBF16, q.device, {tokens, 1, dh}); + vt::AttnQkNormRope(q, h_q3, h_k3, t_qg, t_kg, t_cs16, t_pos, norm, rope); + backend.Synchronize(q); + + const std::vector fused_q = ReadDeviceBf16(q, f_q3); + const std::vector fused_k = ReadDeviceBf16(q, f_k3); + const WordDiff dq = DiffWords(ReadDeviceBf16(q, h_q3), fused_q, dh); + const WordDiff dk = DiffWords(ReadDeviceBf16(q, h_k3), fused_k, dh); + PrintDiff("hand-call vs fused recipe, Q", dq); + PrintDiff("hand-call vs fused recipe, K", dk); + const bool any_rotation = [&] { + for (int64_t t = 0; t < tokens; ++t) + for (int64_t d = 0; d < dh; ++d) + if (fused_q[static_cast(t * dh + d)] != q_words[static_cast(t * dh + d)]) + return true; + return false; + }(); + CHECK(any_rotation); // a no-op preamble would make the identity vacuous + CHECK(dq.different == 0); + CHECK(dk.different == 0); +} diff --git a/tests/vt/residual_norm_upstream.py b/tests/vt/residual_norm_upstream.py new file mode 100644 index 0000000000..28d4037e46 --- /dev/null +++ b/tests/vt/residual_norm_upstream.py @@ -0,0 +1,148 @@ +"""Export the pinned BF16 normalization tests without changing their fixtures. + +Port of vLLM e126687a9a828d513c01a07cd69f025f27d63280: +tests/kernels/core/test_layernorm.py::test_rms_norm and +tests/kernels/ir/test_layernorm.py::{TestRMSNorm,TestFusedAddRMSNorm}. + +The original core test and IR provider/semantics/opcheck tests execute. Their +native references and original torch-generated inputs are exported for C++. +BF16 is the measured compiled-expression policy. F16/F32 activation modes are +explicitly outside that policy and are refused by the shared descriptor tests. +The C++ ABI represents weight=None with an explicit unit BF16 gamma, and plain +RMSNorm with a zero base. Torch's registration/opcheck is run here; C++ checks +its own typed descriptor, dispatch, output ownership, strides, and aliases. +""" +import argparse +import hashlib +import importlib +import itertools +import json +import sys +from pathlib import Path + +PIN = "e126687a9a828d513c01a07cd69f025f27d63280" + + +def main(): + parser = argparse.ArgumentParser() + parser.add_argument("--source", type=Path, required=True) + parser.add_argument("--output", type=Path, required=True) + args = parser.parse_args() + import torch + import vllm + assert vllm.__version__ == "0.28.1rc1.dev132+ge126687a9" + sys.path.insert(0, str(args.source)) + core = importlib.import_module("tests.kernels.core.test_layernorm") + ir_test = importlib.import_module("tests.kernels.ir.test_layernorm") + from vllm import ir + from vllm.config import VllmConfig, set_current_vllm_config + from vllm.utils.torch_utils import set_random_seed + args.output.mkdir(parents=True, exist_ok=True) + cases = [] + active = {} + + def save(name, value): + raw = value.detach().cpu().reshape(-1).contiguous().view(torch.uint8).numpy().tobytes() + file = f"{active['id']}-{name}.bin" + (args.output / file).write_bytes(raw) + active["files"][name] = {"file": file, "shape": list(value.shape), + "stride": list(value.stride()), "dtype": str(value.dtype), + "bytes": len(raw), "sha256": hashlib.sha256(raw).hexdigest()} + + def capture(x, residual, weight, epsilon, output): + if active.get("captured"): + return + active["captured"] = True + active["rows"], active["width"] = x.shape + active["epsilon"] = epsilon + active["device"] = x.device.index + active["add_residual"] = residual is not None + save("a", x) + save("base", torch.zeros_like(x) if residual is None else residual) + save("gamma", weight) + norm, res = output if residual is not None else (output, None) + save("norm", norm) + if res is not None: + save("residual", res) + operation = ir.ops.rms_norm if residual is None else ir.ops.fused_add_rms_norm + native = operation.impls["native"].impl_fn + inputs = (x,) if residual is None else (x, residual) + unit = native(*inputs, None, epsilon) + unit_with_weight = native(*inputs, torch.ones_like(weight), epsilon) + torch.testing.assert_close(unit, unit_with_weight) + save("unit_norm", unit if residual is None else unit[0]) + + original_core = core.RMSNorm.forward_native + + def core_native(layer, x, residual=None): + output = original_core(layer, x, residual) + capture(x, residual, layer.weight.data, layer.variance_epsilon, output) + return output + + def begin(name, tolerance, **parameters): + active.clear() + active.update({"id": name, "pin": PIN, "seed": 0, "tolerance": tolerance, + "parameters": parameters, "files": {}}) + set_random_seed(0) + + def finish(): + assert active.pop("captured") + cases.append(dict(active)) + (args.output / "cases.json").write_text(json.dumps({"pin": PIN, "complete": False, + "cases": cases}, indent=2) + "\n") + print(active["id"], "PASS", flush=True) + + with torch.inference_mode(), set_current_vllm_config(VllmConfig()): + core.RMSNorm.forward_native = core_native + try: + for index, values in enumerate(itertools.product(core.NUM_TOKENS, core.HIDDEN_SIZES, + core.ADD_RESIDUAL, core.CUDA_DEVICES, (False, True))): + rows, width, add, device, strided = values + begin(f"core-{index}", {"atol": 1e-2, "rtol": 1e-2}, rows=rows, + width=width, add_residual=add, device=device, strided=strided) + core.test_rms_norm(None, rows, width, add, torch.bfloat16, 0, device, strided) + finish() + finally: + core.RMSNorm.forward_native = original_core + torch.set_default_device("cuda:0") + for fused, cls, operation, global_name in ( + (False, ir_test.TestRMSNorm, ir.ops.rms_norm, "rms_norm_native"), + (True, ir_test.TestFusedAddRMSNorm, ir.ops.fused_add_rms_norm, "fused_add_rms_norm_native")): + original = getattr(ir_test, global_name) + + def observed(*positional, **keywords): + result = original(*positional, **keywords) + x = positional[0] + residual = positional[1] if fused else None + gamma = positional[2 if fused else 1] + epsilon = positional[3 if fused else 2] + capture(x, residual, gamma, epsilon, result) + return result + + setattr(ir_test, global_name, observed) + try: + for index, (rows, width, epsilon) in enumerate(itertools.product( + ir_test.NUM_TOKENS, ir_test.COMMON_HIDDEN_SIZES, (1e-6, 1e-5))): + begin(f"ir-{'add' if fused else 'rms'}-{index}", operation.get_tolerance(torch.bfloat16), + rows=rows, width=width, epsilon=epsilon, provider="vllm_c") + cls().test_impls(torch.bfloat16, rows, width, epsilon, "vllm_c") + finish() + # These original semantics tests use 4x8 regardless of the class + # grid. Run each distinct epsilon with the original seed zero. + setattr(ir_test, global_name, original) + for epsilon in (1e-6, 1e-5): + cls().test_native_semantics(torch.bfloat16, 1, 2048, epsilon) + cls().test_torch_opcheck(torch.bfloat16, 1, 2048, epsilon, "vllm_c") + finally: + setattr(ir_test, global_name, original) + record = {"pin": PIN, "complete": True, "cases": cases, + "applicability": {"BF16": "executed", "F16": "refused compiled-expression activation dtype", + "F32": "refused compiled-expression activation dtype"}, + "adaptations": ["C++ unit gamma represents absent weight", "zero base represents plain RMSNorm", + "C++ native provider replaces the upstream provider parameter; original vllm_c test still executes", + "Torch registration/opcheck runs in this exporter; C++ validates its own typed ABI"]} + (args.output / "cases.json").write_text(json.dumps(record, indent=2) + "\n") + + +if __name__ == "__main__": + main() diff --git a/tests/vt/rocm_moe_test_helpers.h b/tests/vt/rocm_moe_test_helpers.h new file mode 100644 index 0000000000..4d72b00226 --- /dev/null +++ b/tests/vt/rocm_moe_test_helpers.h @@ -0,0 +1,125 @@ +// vllm.cpp original. Scoped device ownership for the BF16 MoE gates (#3094). +#pragma once +#include +#include +#include +#include +#include +#include +#include "vt/backend.h" +#include "vt/dtype.h" +#include "vt/ops.h" +#if defined(VLLM_CPP_HIP) +#include +#endif + +namespace rocm_moe_test { +// The existing Backend allocation/copy helpers use the calling thread's HIP +// device. Scope test resources explicitly so the device-isolation gate measures +// the new operation, with real allocations on each declared device. +class HostDeviceScope { + public: + explicit HostDeviceScope(int device) { +#if defined(VLLM_CPP_HIP) + REQUIRE(hipGetDevice(&previous_) == hipSuccess); + REQUIRE(hipSetDevice(device) == hipSuccess); +#else + (void)device; +#endif + } + ~HostDeviceScope() { +#if defined(VLLM_CPP_HIP) + (void)hipSetDevice(previous_); +#endif + } + static int Current() { + int device = 0; +#if defined(VLLM_CPP_HIP) + REQUIRE(hipGetDevice(&device) == hipSuccess); +#endif + return device; + } + private: +#if defined(VLLM_CPP_HIP) + int previous_ = 0; +#endif +}; +inline vt::Device Device(int index = 0) { return {vt::DeviceType::kROCM, index}; } +inline void RequireDevice(int index = 0) { + if (vt::TryGetBackend(Device(index)) == nullptr) { + MESSAGE("required ROCm device is unavailable"); + std::exit(77); + } +} +inline vt::Tensor View(void* data, vt::DType dtype, vt::Device device, + const std::vector& shape) { + vt::Tensor result; + result.data = data; + result.dtype = dtype; + result.device = device; + result.rank = static_cast(shape.size()); + int64_t stride = 1; + for (int i = result.rank - 1; i >= 0; --i) { + result.shape[i] = shape[static_cast(i)]; + result.stride[i] = stride; + stride *= result.shape[i]; + } + return result; +} +struct Queue { + vt::Queue value; + static vt::Queue Create(int index) { + HostDeviceScope device(index); + return vt::CreateQueue(Device(index)); + } + explicit Queue(int index = 0) : value(Create(index)) {} + ~Queue() { HostDeviceScope device(value.device.index); vt::DestroyQueue(value); } + Queue(const Queue&) = delete; + Queue& operator=(const Queue&) = delete; +}; +class Buffer { + public: + Buffer(vt::Queue& q, vt::DType dtype, const std::vector& shape, + const void* host = nullptr) : backend_(vt::GetBackend(q.device)) { + HostDeviceScope device(q.device.index); + tensor_ = View(nullptr, dtype, q.device, shape); + bytes_ = static_cast(tensor_.Numel()) * vt::SizeOf(dtype); + tensor_.data = vt::Alloc(q.device, bytes_ + kGuard); +#if defined(VLLM_CPP_HIP) + hipPointerAttribute_t attributes{}; + REQUIRE(hipPointerGetAttributes(&attributes, tensor_.data) == hipSuccess); + REQUIRE(attributes.device == q.device.index); +#endif + // Positive guard bytes make a removed K mask observable even when the + // allocator otherwise returns zero-filled pages beyond the logical matrix. + backend_.Memset(q, tensor_.data, 0x55, bytes_ + kGuard); + if (host != nullptr) backend_.Copy(q, tensor_.data, host, bytes_); + backend_.Synchronize(q); + } + ~Buffer() { HostDeviceScope device(tensor_.device.index); vt::Free(tensor_.device, tensor_.data); } + Buffer(const Buffer&) = delete; + Buffer& operator=(const Buffer&) = delete; + vt::Tensor& tensor() { return tensor_; } + const vt::Tensor& tensor() const { return tensor_; } + template std::vector Download(vt::Queue& q) const { + HostDeviceScope device(q.device.index); + REQUIRE(bytes_ % sizeof(T) == 0); + std::vector host(bytes_ / sizeof(T)); + backend_.Copy(q, host.data(), tensor_.data, bytes_); + backend_.Synchronize(q); + return host; + } + void CheckGuard(vt::Queue& q) const { + HostDeviceScope device(q.device.index); + std::vector guard(kGuard); + backend_.Copy(q, guard.data(), static_cast(tensor_.data) + bytes_, kGuard); + backend_.Synchronize(q); + CHECK(std::all_of(guard.begin(), guard.end(), [](uint8_t b) { return b == 0x55; })); + } + private: + static constexpr size_t kGuard = 1024; + vt::Backend& backend_; + vt::Tensor tensor_; + size_t bytes_; +}; +} // namespace rocm_moe_test diff --git a/tests/vt/test_moe_bf16_native_contract.cpp b/tests/vt/test_moe_bf16_native_contract.cpp new file mode 100644 index 0000000000..a2b571744f --- /dev/null +++ b/tests/vt/test_moe_bf16_native_contract.cpp @@ -0,0 +1,166 @@ +// vllm.cpp original. Typed shared BF16 MoE contracts (#3094). +#include +#include +#include +#include "rocm_moe_test_helpers.h" +#include "vt/op_provider.h" + +// Adding native modes must not change any existing caller's function signature. +static_assert(std::is_same_v); +static_assert(std::is_same_v); +static_assert(std::is_same_v); + +TEST_CASE("shared native BF16 MoE validates routes and physical tensor formats") { + const vt::Device cpu{vt::DeviceType::kCPU, 0}; + vt::Queue q{cpu, nullptr}; + std::array values{}; + std::array weights{}; + std::array ids{}; + std::array ptrs{}; + auto a = rocm_moe_test::View(values.data(), vt::DType::kBF16, cpu, {2, 4}); + auto out = rocm_moe_test::View(values.data(), vt::DType::kBF16, cpu, {2, 4}); + auto expert = rocm_moe_test::View(ids.data(), vt::DType::kI32, cpu, {2}); + auto pointers = rocm_moe_test::View(ptrs.data(), vt::DType::kI64, cpu, {4}); + auto route = rocm_moe_test::View(weights.data(), vt::DType::kF32, cpu, {1}); + CHECK_THROWS_WITH(vt::MoeGroupedGemmBf16Weighted(q, out, a, expert, nullptr, pointers, route), + doctest::Contains("route_weights must be contiguous f32 [P]")); + route.shape[0] = 2; + route.dtype = vt::DType::kBF16; + CHECK_THROWS_WITH(vt::MoeGroupedGemmBf16Weighted(q, out, a, expert, nullptr, pointers, route), + doctest::Contains("route_weights must be contiguous f32 [P]")); + a.dtype = vt::DType::kF32; + CHECK_THROWS_WITH(vt::MoeGroupedGemmBf16GateUpSiluNative(q, out, a, expert, nullptr, + pointers, pointers), + doctest::Contains("act must be bf16")); + a.dtype = vt::DType::kBF16; + out.dtype = vt::DType::kF32; + CHECK_THROWS_WITH(vt::MoeGroupedGemmBf16GateUpSiluNative(q, out, a, expert, nullptr, + pointers, pointers), + doctest::Contains("out must be bf16")); + auto rows = expert; + rows.stride[0] = 2; + out.dtype = vt::DType::kBF16; + CHECK_THROWS_WITH(vt::MoeGroupedGemmBf16GateUpSiluNative(q, out, a, expert, &rows, + pointers, pointers), + doctest::Contains("row_map must be contiguous i32 [P]")); + auto combined = rocm_moe_test::View(values.data(), vt::DType::kBF16, cpu, {1, 4}); + auto weighted = rocm_moe_test::View(values.data(), vt::DType::kF32, cpu, {1, 2, 4}); + CHECK_THROWS_WITH(vt::MoeCombinePreweighted(q, combined, weighted), + doctest::Contains("expert_out must be bf16")); + weighted.dtype = vt::DType::kBF16; + weighted.shape[1] = 4; + weighted.shape[2] = 2; + CHECK_THROWS_WITH(vt::MoeCombinePreweighted(q, combined, weighted), + doctest::Contains("must match out")); + weighted = rocm_moe_test::View(values.data(), vt::DType::kBF16, cpu, {1, 2, 4}); + auto shared = combined; + shared.dtype = vt::DType::kF16; + CHECK_THROWS_WITH(vt::MoeCombinePreweighted(q, combined, weighted, &shared), + doctest::Contains("shared must be f32/bf16")); + CHECK_FALSE(vt::MoeGroupedBf16NativeAvailable(vt::DeviceType::kCPU)); +} + +TEST_CASE("shared native BF16 MoE requires all five providers in every subset") { + // XPU has no backend registrations. Unique names let each subset use the + // real provider seam without changing any production device's registrations. + constexpr vt::DeviceType device = vt::DeviceType::kXPU; + constexpr std::array operations = { + vt::OpId::kMoeGroupedGemmBf16, vt::OpId::kMoeGroupedGemmBf16GateUpSilu, + vt::OpId::kMoeGroupedGemmBf16GateUpSiluNative, + vt::OpId::kMoeGroupedGemmBf16Weighted, vt::OpId::kMoeCombinePreweighted}; + static constexpr std::array names = { + "moe-contract-grouped", "moe-contract-gate-up", "moe-contract-native", + "moe-contract-weighted", "moe-contract-combine"}; + struct DisableOnExit { + ~DisableOnExit() { + for (const char* name : names) vt::DisableOpProvider(name, true); + } + } cleanup; + for (size_t i = 0; i < operations.size(); ++i) { + REQUIRE_FALSE(vt::OpRegistered(operations[i], device)); + // The stub is never dispatched. Only native availability is inspected. + vt::RegisterOpProvider(operations[i], device, + {names[i], 0, nullptr, reinterpret_cast(+[] {})}); + } + // Mask 3 is the legacy-only set. Masks 31^(1< values{}; + std::array weights{}; + std::array ids{}; + std::array ptrs{}; + auto a = rocm_moe_test::View(values.data(), vt::DType::kBF16, cpu, {2, 4}); + auto out = a; + auto expert = rocm_moe_test::View(ids.data(), vt::DType::kI32, cpu, {2}); + auto pointers = rocm_moe_test::View(ptrs.data(), vt::DType::kI64, cpu, {4}); + auto route = rocm_moe_test::View(weights.data(), vt::DType::kF32, cpu, {2}); + SUBCASE("route rank with the correct element count") { + route = rocm_moe_test::View(weights.data(), vt::DType::kF32, cpu, {1, 2}); + } + SUBCASE("route stride") { route.stride[0] = 2; } + SUBCASE("route device") { route.device.index = 1; } + SUBCASE("weighted down retains the common grouped validator") { + a.dtype = vt::DType::kF32; + CHECK_THROWS_WITH(vt::MoeGroupedGemmBf16Weighted(q, out, a, expert, nullptr, + pointers, route), + doctest::Contains("act must be bf16")); + return; + } + CHECK_THROWS_WITH(vt::MoeGroupedGemmBf16Weighted(q, out, a, expert, nullptr, + pointers, route), + doctest::Contains("route_weights must be contiguous f32 [P]")); +} + +TEST_CASE("shared native BF16 MoE rejects combine metadata before dispatch") { + const vt::Device cpu{vt::DeviceType::kCPU, 0}; + vt::Queue q{cpu, nullptr}; + std::array values{}; + auto out = rocm_moe_test::View(values.data(), vt::DType::kBF16, cpu, {2, 4}); + auto weighted = rocm_moe_test::View(values.data(), vt::DType::kBF16, cpu, {2, 2, 4}); + SUBCASE("expert output stride") { weighted.stride[0] += 1; } + SUBCASE("combined output stride") { out.stride[0] += 1; } + SUBCASE("expert output device") { weighted.device.index = 1; } + SUBCASE("combined output device") { out.device.index = 1; } + CHECK_THROWS_WITH(vt::MoeCombinePreweighted(q, out, weighted), + doctest::Contains("contiguous tensors on the queue device required")); +} + +TEST_CASE("shared native BF16 MoE rejects shared metadata before dispatch") { + const vt::Device cpu{vt::DeviceType::kCPU, 0}; + vt::Queue q{cpu, nullptr}; + std::array values{}; + auto out = rocm_moe_test::View(values.data(), vt::DType::kBF16, cpu, {2, 4}); + auto weighted = rocm_moe_test::View(values.data(), vt::DType::kBF16, cpu, {2, 2, 4}); + auto shared = out; + SUBCASE("shared rank with matching leading dimensions") { + shared = rocm_moe_test::View(values.data(), vt::DType::kBF16, cpu, {2, 4, 1}); + } + SUBCASE("shared token count") { + shared = rocm_moe_test::View(values.data(), vt::DType::kBF16, cpu, {3, 4}); + } + SUBCASE("shared hidden width") { + shared = rocm_moe_test::View(values.data(), vt::DType::kBF16, cpu, {2, 5}); + } + SUBCASE("shared stride") { shared.stride[0] += 1; } + SUBCASE("shared device") { shared.device.index = 1; } + CHECK_THROWS_WITH(vt::MoeCombinePreweighted(q, out, weighted, &shared), + doctest::Contains("shared must be f32/bf16 [T,H] on the queue device")); +} diff --git a/tests/vt/test_ops_moe_grouped_bf16.cpp b/tests/vt/test_ops_moe_grouped_bf16.cpp index 2567d0542d..b560da1416 100644 --- a/tests/vt/test_ops_moe_grouped_bf16.cpp +++ b/tests/vt/test_ops_moe_grouped_bf16.cpp @@ -17,6 +17,8 @@ #include #include +#include +#include #include #include #include @@ -35,9 +37,15 @@ using vt::DType; using vt::Queue; using vt::Tensor; -bool HasCuda() { +// Preserve the CUDA cases and allow the same fixtures on the ROCm provider. +DeviceType TestDevice() { + const char* device = std::getenv("VT_MOE_TEST_DEVICE"); + return device != nullptr && std::strcmp(device, "rocm") == 0 + ? DeviceType::kROCM : DeviceType::kCUDA; +} +bool HasDevice() { try { - vt::GetBackend(DeviceType::kCUDA); + vt::GetBackend(TestDevice()); return true; } catch (const std::runtime_error&) { return false; @@ -131,7 +139,7 @@ std::vector RandomBf16(size_t numel, uint32_t seed) { // exercises the bf16 output dtype (down) vs f32 (gate/up). void RunGroupedBf16Case(int64_t e_count, int64_t t_rows, int64_t top_k, int64_t k_dim, int64_t n_cols, uint32_t seed, bool use_row_map, bool bf16_out) { - Backend& gpu = vt::GetBackend(DeviceType::kCUDA); + Backend& gpu = vt::GetBackend(TestDevice()); const int64_t P = t_rows * top_k; // With identity routing the activation IS the per-pair buffer, so it has P rows. const int64_t act_rows = use_row_map ? t_rows : P; @@ -206,7 +214,7 @@ void RunGroupedBf16Case(int64_t e_count, int64_t t_rows, int64_t top_k, int64_t // fixed ascending order (never atomicAdd), so greedy decode stays reproducible. bool GroupedBf16Bitwise(int64_t e_count, int64_t t_rows, int64_t top_k, int64_t k_dim, int64_t n_cols, uint32_t seed, int reps) { - Backend& gpu = vt::GetBackend(DeviceType::kCUDA); + Backend& gpu = vt::GetBackend(TestDevice()); const int64_t P = t_rows * top_k; std::vector> w(static_cast(e_count)); for (int64_t e = 0; e < e_count; ++e) @@ -250,9 +258,9 @@ bool GroupedBf16Bitwise(int64_t e_count, int64_t t_rows, int64_t top_k, int64_t } // namespace // P = 6 < kTileMinRows(32) -> naive one-thread-per-output kernel, f32 out (gate/up). -TEST_CASE("CUDA moe_grouped_gemm_bf16 naive path (small P) matches the per-expert reference") { - if (!HasCuda()) { - MESSAGE("no CUDA backend registered; skipping"); +TEST_CASE("GPU moe_grouped_gemm_bf16 naive path (small P) matches the per-expert reference") { + if (!HasDevice()) { + MESSAGE("requested GPU backend is unavailable; skipping"); return; } RunGroupedBf16Case(/*e_count=*/5, /*t_rows=*/3, /*top_k=*/2, /*k_dim=*/64, /*n_cols=*/8, @@ -262,9 +270,9 @@ TEST_CASE("CUDA moe_grouped_gemm_bf16 naive path (small P) matches the per-exper // P = 40 (kTileMinRows <= P <= kMoeDecodeMaxP) -> BM=16 decode WMMA tile. K=80 is a // multiple of 16 but NOT of BK=32 (partial last K-tile); N=130 crosses the BN=64 // tile boundary unevenly. Mirrors the NVFP4 tiled-path case's awkward shapes. -TEST_CASE("CUDA moe_grouped_gemm_bf16 decode WMMA tile (BM=16) matches the per-expert reference") { - if (!HasCuda()) { - MESSAGE("no CUDA backend registered; skipping"); +TEST_CASE("GPU moe_grouped_gemm_bf16 decode WMMA tile (BM=16) matches the per-expert reference") { + if (!HasDevice()) { + MESSAGE("requested GPU backend is unavailable; skipping"); return; } RunGroupedBf16Case(/*e_count=*/7, /*t_rows=*/20, /*top_k=*/2, /*k_dim=*/80, /*n_cols=*/130, @@ -272,9 +280,9 @@ TEST_CASE("CUDA moe_grouped_gemm_bf16 decode WMMA tile (BM=16) matches the per-e } // P = 1024 > kMoeDecodeMaxP(512) -> BM=64 prefill WMMA tile. -TEST_CASE("CUDA moe_grouped_gemm_bf16 prefill WMMA tile (BM=64) matches the per-expert reference") { - if (!HasCuda()) { - MESSAGE("no CUDA backend registered; skipping"); +TEST_CASE("GPU moe_grouped_gemm_bf16 prefill WMMA tile (BM=64) matches the per-expert reference") { + if (!HasDevice()) { + MESSAGE("requested GPU backend is unavailable; skipping"); return; } RunGroupedBf16Case(/*e_count=*/9, /*t_rows=*/128, /*top_k=*/8, /*k_dim=*/96, /*n_cols=*/70, @@ -287,9 +295,9 @@ TEST_CASE("CUDA moe_grouped_gemm_bf16 prefill WMMA tile (BM=64) matches the per- // n=130) and therefore still cover the W5 fallback tile. Here K=264 is a multiple // of 8 but NOT of BK=32 (partial last K-tile, exercising the cp.async `zfill` // tail) and N=200 is a multiple of 8 but NOT of BN=128 (partial last N-tile). -TEST_CASE("CUDA moe_grouped_gemm_bf16 pipelined prefill tile matches the per-expert reference") { - if (!HasCuda()) { - MESSAGE("no CUDA backend registered; skipping"); +TEST_CASE("GPU moe_grouped_gemm_bf16 pipelined prefill tile matches the per-expert reference") { + if (!HasDevice()) { + MESSAGE("requested GPU backend is unavailable; skipping"); return; } RunGroupedBf16Case(/*e_count=*/9, /*t_rows=*/128, /*top_k=*/8, /*k_dim=*/264, /*n_cols=*/200, @@ -298,9 +306,9 @@ TEST_CASE("CUDA moe_grouped_gemm_bf16 pipelined prefill tile matches the per-exp // W6 PIPELINED decode tile (BM=16, BN=128, BK=32, 3-stage), aligned pitches, plus // the identity row-map + bf16-out (down-projection) call shape on the same tile. -TEST_CASE("CUDA moe_grouped_gemm_bf16 pipelined decode tile matches the per-expert reference") { - if (!HasCuda()) { - MESSAGE("no CUDA backend registered; skipping"); +TEST_CASE("GPU moe_grouped_gemm_bf16 pipelined decode tile matches the per-expert reference") { + if (!HasDevice()) { + MESSAGE("requested GPU backend is unavailable; skipping"); return; } RunGroupedBf16Case(/*e_count=*/7, /*t_rows=*/20, /*top_k=*/2, /*k_dim=*/264, /*n_cols=*/200, @@ -314,9 +322,9 @@ TEST_CASE("CUDA moe_grouped_gemm_bf16 pipelined decode tile matches the per-expe // = 4 -> four f32 partials reduced in fixed ascending split order. Also asserts // the split reduction is RUN-TO-RUN BIT-REPRODUCIBLE (no atomicAdd), which the // greedy token-exact gate depends on. -TEST_CASE("CUDA moe_grouped_gemm_bf16 split-K decode path matches the reference and is exact") { - if (!HasCuda()) { - MESSAGE("no CUDA backend registered; skipping"); +TEST_CASE("GPU moe_grouped_gemm_bf16 split-K decode path matches the reference and is exact") { + if (!HasDevice()) { + MESSAGE("requested GPU backend is unavailable; skipping"); return; } RunGroupedBf16Case(/*e_count=*/5, /*t_rows=*/3, /*top_k=*/2, /*k_dim=*/1024, /*n_cols=*/8, @@ -327,9 +335,9 @@ TEST_CASE("CUDA moe_grouped_gemm_bf16 split-K decode path matches the reference // Identity routing (row_map == nullptr) + bf16 output — the DOWN projection's exact // call shape (act = the per-pair silu buffer, one act row per output row). -TEST_CASE("CUDA moe_grouped_gemm_bf16 identity row-map + bf16 out matches the reference") { - if (!HasCuda()) { - MESSAGE("no CUDA backend registered; skipping"); +TEST_CASE("GPU moe_grouped_gemm_bf16 identity row-map + bf16 out matches the reference") { + if (!HasDevice()) { + MESSAGE("requested GPU backend is unavailable; skipping"); return; } // Naive regime (P=12) and decode-tile regime (P=128), both bf16 out. diff --git a/tests/vt/test_ops_moe_grouped_bf16_gate_up_silu.cpp b/tests/vt/test_ops_moe_grouped_bf16_gate_up_silu.cpp index aada2d125d..89e65a2053 100644 --- a/tests/vt/test_ops_moe_grouped_bf16_gate_up_silu.cpp +++ b/tests/vt/test_ops_moe_grouped_bf16_gate_up_silu.cpp @@ -20,6 +20,8 @@ #include #include +#include +#include #include #include #include @@ -39,9 +41,15 @@ using vt::DType; using vt::Queue; using vt::Tensor; -bool HasCuda() { +// Preserve the CUDA cases and allow the same fixtures on the ROCm provider. +DeviceType TestDevice() { + const char* device = std::getenv("VT_MOE_TEST_DEVICE"); + return device != nullptr && std::strcmp(device, "rocm") == 0 + ? DeviceType::kROCM : DeviceType::kCUDA; +} +bool HasDevice() { try { - vt::GetBackend(DeviceType::kCUDA); + vt::GetBackend(TestDevice()); return true; } catch (const std::runtime_error&) { return false; @@ -118,7 +126,7 @@ std::vector RandomBf16(size_t numel, uint32_t seed) { // carries. void RunFusedGateUpSiluCase(int64_t e_count, int64_t t_rows, int64_t top_k, int64_t k_dim, int64_t i_dim, uint32_t seed) { - Backend& gpu = vt::GetBackend(DeviceType::kCUDA); + Backend& gpu = vt::GetBackend(TestDevice()); const int64_t P = t_rows * top_k; std::vector> gate_w(static_cast(e_count)); @@ -196,9 +204,9 @@ void RunFusedGateUpSiluCase(int64_t e_count, int64_t t_rows, int64_t top_k, int6 } // namespace // P = 12 < kTileMinRows(32), K=64 -> splits==1 (naive partials) + reduce+SwiGLU. -TEST_CASE("CUDA moe_grouped_gemm_bf16_gate_up_silu naive path == composite (byte-exact)") { - if (!HasCuda()) { - MESSAGE("no CUDA backend registered; skipping"); +TEST_CASE("GPU moe_grouped_gemm_bf16_gate_up_silu naive path == composite (byte-exact)") { + if (!HasDevice()) { + MESSAGE("requested GPU backend is unavailable; skipping"); return; } RunFusedGateUpSiluCase(/*e_count=*/5, /*t_rows=*/6, /*top_k=*/2, /*k_dim=*/64, /*i_dim=*/8, @@ -208,9 +216,9 @@ TEST_CASE("CUDA moe_grouped_gemm_bf16_gate_up_silu naive path == composite (byte // P = 6 < kTileMinRows, K=1024 -> MoeSplitKCount picks splits=4: split-K partials // reduced in fixed ascending order, then fused SwiGLU. The path where the fold's // launch reduction (5 -> 3) actually lands (Qwen3-Coder / DeepSeek-V2 c1 decode). -TEST_CASE("CUDA moe_grouped_gemm_bf16_gate_up_silu split-K decode == composite (byte-exact)") { - if (!HasCuda()) { - MESSAGE("no CUDA backend registered; skipping"); +TEST_CASE("GPU moe_grouped_gemm_bf16_gate_up_silu split-K decode == composite (byte-exact)") { + if (!HasDevice()) { + MESSAGE("requested GPU backend is unavailable; skipping"); return; } RunFusedGateUpSiluCase(/*e_count=*/5, /*t_rows=*/3, /*top_k=*/2, /*k_dim=*/1024, /*i_dim=*/8, @@ -223,9 +231,9 @@ TEST_CASE("CUDA moe_grouped_gemm_bf16_gate_up_silu split-K decode == composite ( // P = 40 (kTileMinRows <= P) -> BM=16 decode WMMA tile; P = 1024 -> BM=64 prefill // WMMA tile. The WMMA branch reuses the tuned grouped GEMM twice + the identical // silu-mul, so it is composite-identical by construction — this asserts it. -TEST_CASE("CUDA moe_grouped_gemm_bf16_gate_up_silu WMMA tiles == composite (byte-exact)") { - if (!HasCuda()) { - MESSAGE("no CUDA backend registered; skipping"); +TEST_CASE("GPU moe_grouped_gemm_bf16_gate_up_silu WMMA tiles == composite (byte-exact)") { + if (!HasDevice()) { + MESSAGE("requested GPU backend is unavailable; skipping"); return; } RunFusedGateUpSiluCase(/*e_count=*/7, /*t_rows=*/20, /*top_k=*/2, /*k_dim=*/80, /*i_dim=*/130, diff --git a/tests/vt/test_ops_residual_rmsnorm.cpp b/tests/vt/test_ops_residual_rmsnorm.cpp new file mode 100644 index 0000000000..811ddc572f --- /dev/null +++ b/tests/vt/test_ops_residual_rmsnorm.cpp @@ -0,0 +1,397 @@ +// Ported from vLLM tests/kernels/ir/test_layernorm.py::TestFusedAddRMSNorm +// at e126687a9a828d513c01a07cd69f025f27d63280. The compiled boundary is +// additionally pinned by the executing row-zero witness in #3103. +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include "support/residual_norm_fixture.h" +#include "support/residual_norm_later_fixture.h" +#include "support/residual_norm_test.h" +#include "vt/ops.h" + +TEST_CASE("residual RMS norm preserves the compiled primary row zero") { + using namespace residual_norm_fixture; + const vt::Device device{vt::DeviceType::kCPU, 0}; + vt::Queue q{device, nullptr}; + auto a = kAttention, r = kResidual, w = kGamma; + std::vector out(128); + auto ta = vt::Tensor::Contiguous(a.data(), vt::DType::kBF16, device, {1, 128}); + auto tr = vt::Tensor::Contiguous(r.data(), vt::DType::kBF16, device, {1, 128}); + auto tw = vt::Tensor::Contiguous(w.data(), vt::DType::kBF16, device, {128}); + auto to = vt::Tensor::Contiguous(out.data(), vt::DType::kBF16, device, {1, 128}); + vt::ResidualRmsNorm(q, to, ta, tr, nullptr, tw, vt::ResidualRmsNormArgs{}); + for (size_t j = 0; j < out.size(); ++j) { + CAPTURE(j); + CHECK(out[j] == kPostNorm[j]); + } + CHECK(a == kAttention); + CHECK(r == kResidual); + CHECK(w == kGamma); +} + +TEST_CASE("residual RMS norm replays the actual later compiled primary outputs") { + using namespace residual_norm_fixture; + using namespace residual_norm_test; + for (auto device : Devices()) for (bool final : {false, true}) for (int mode : {0, 1, 2}) { + CAPTURE(device.type); CAPTURE(final); CAPTURE(mode); + Queue queue(device); auto& q = queue.q; + auto vector = [](const auto& values) { return std::vector(values.begin(), values.end()); }; + const auto a = vector(final ? kFinalAttention : kNextAttention); + const auto base = vector(final ? kFinalBase : kNextBase); + const auto delta = vector(final ? kFinalMoe : kNextMoe); + const auto gamma = vector(final ? kFinalGamma : kNextGamma); + Buffer da(q, a, 1, 128), dr(q, base, 1, 128), dm(q, delta, 1, 128), dw(q, gamma, -1, 128); + Buffer out(q, std::vector(128), 1, 128), res(q, std::vector(128), 1, 128); + vt::ResidualRmsNormArgs args{1e-6f, {vt::ResidualNormExpr::kDeltaPlusAdd, !final}}; + Run(mode, q, out.t, da.t, dr.t, &dm.t, dw.t, args, final ? nullptr : &res.t); + CHECK(out.Read() == vector(final ? kFinalNorm : kNextNorm)); + if (!final) CHECK(res.Read() == vector(kNextResidual)); + CHECK(da.Read() == a); CHECK(dr.Read() == base); CHECK(dm.Read() == delta); CHECK(dw.Read() == gamma); + } +} + +TEST_CASE("residual RMS norm ordered expressions preserve every observable boundary") { + using namespace residual_norm_test; + for (const auto device : Devices()) for (int64_t width : {8, 128, 769, 8192}) + for (bool materialize : {false, true}) for (bool alias : {false, true}) + for (int witness : {0, 1, 2}) for (int mode : {0, 1, 2}) { + CAPTURE(device.type); CAPTURE(width); CAPTURE(materialize); + CAPTURE(alias); CAPTURE(witness); CAPTURE(mode); + constexpr int64_t rows = 3; + const int64_t stride = width + 7; + std::vector a(rows * stride, Buffer::kGuard), base(a), delta(a); + std::vector gamma(width), blank(a.size(), Buffer::kGuard); + for (int64_t row = 0; row < rows; ++row) for (int64_t j = 0; j < width; ++j) { + const auto at = static_cast(row * stride + j); + if (witness == 0) { a[at] = vt::F32ToBF16(256); base[at] = vt::F32ToBF16(-256); delta[at] = vt::F32ToBF16(0x1p-17f); } + if (witness == 1) { a[at] = vt::F32ToBF16(1); base[at] = delta[at] = vt::F32ToBF16(0x1p-8f); } + if (witness == 2) { + const size_t k = static_cast(j % 128); + a[at] = residual_norm_fixture::kAttention[k]; + base[at] = residual_norm_fixture::kResidual[k]; delta[at] = 0; + } + gamma[j] = witness == 2 ? residual_norm_fixture::kGamma[j % 128] : vt::F32ToBF16(0.75f); + } + std::vector expected_residual; + const auto expected = Reference(a, base, &delta, gamma, rows, width, stride, 1e-6f, &expected_residual); + Queue queue(device); auto& q = queue.q; + Buffer da(q, a, rows, width, stride), dr(q, base, rows, width, stride); + Buffer dm(q, delta, rows, width, stride), dw(q, gamma, -1, width); + Buffer output(q, blank, rows, width, stride), residual(q, blank, rows, width, stride); + vt::ResidualRmsNormArgs args{1e-6f, {vt::ResidualNormExpr::kDeltaPlusAdd, materialize, alias, materialize && alias}}; + vt::Tensor& out = alias ? dm.t : output.t; + vt::Tensor* res = materialize ? (alias ? &dr.t : &residual.t) : nullptr; + Run(mode, q, out, da.t, dr.t, &dm.t, dw.t, args, res); + const auto actual = alias ? dm.Read() : output.Read(); + // The constant witnesses have exact reductions. The measured row also + // guards gamma rounding and normalization from the stored residual. + CHECK(actual == expected); + CHECK(da.Read() == a); CHECK(dw.Read() == gamma); + if (!alias) CHECK(dm.Read() == delta); + if (!(alias && materialize)) CHECK(dr.Read() == base); + if (materialize) CHECK((alias ? dr.Read() : residual.Read()) == expected_residual); + else CHECK(residual.Read() == blank); + for (const auto* buffer : {&da, &dr, &dm, &dw, &output, &residual}) buffer->CheckGuard(); + } +} + +TEST_CASE("residual RMS norm descriptors reject invalid operands and overlaps before dispatch") { + using namespace residual_norm_test; + const vt::Device cpu{vt::DeviceType::kCPU, 0}; + vt::Queue q{cpu, nullptr}; + std::vector storage(128, vt::F32ToBF16(1)); + auto a = vt::Tensor::Contiguous(storage.data(), vt::DType::kBF16, cpu, {2, 8}); + auto base = a, delta = a, gamma = vt::Tensor::Contiguous(storage.data() + 48, vt::DType::kBF16, cpu, {8}); + auto out = a, residual = a; + base.data = storage.data() + 16; delta.data = storage.data() + 32; + out.data = storage.data() + 64; residual.data = storage.data() + 80; + const auto good = vt::ResidualRmsNormArgs{}; + auto run = [&](const vt::ResidualRmsNormArgs& args, const vt::Tensor* m = nullptr, vt::Tensor* r = nullptr) { + vt::ResidualRmsNorm(q, out, a, base, m, gamma, args, r); + }; + CHECK_NOTHROW(run(good)); + auto args = good; args.descriptor.expression = static_cast(255); + CHECK_THROWS_AS(run(args), std::runtime_error); + CHECK_THROWS_AS(run(good, &delta), std::runtime_error); + args = good; args.descriptor.expression = vt::ResidualNormExpr::kDeltaPlusAdd; + CHECK_THROWS_AS(run(args), std::runtime_error); + args = good; args.descriptor.materialize_residual = true; + CHECK_THROWS_AS(run(args), std::runtime_error); + CHECK_THROWS_AS(run(good, nullptr, &residual), std::runtime_error); + args = good; args.descriptor.output_alias_delta = true; + CHECK_THROWS_AS(run(args), std::runtime_error); + args = good; args.descriptor.residual_alias_base = true; + CHECK_THROWS_AS(run(args), std::runtime_error); + for (float epsilon : {-1.0f, std::numeric_limits::infinity(), std::numeric_limits::quiet_NaN()}) { + args = good; args.eps = epsilon; CHECK_THROWS_AS(run(args), std::runtime_error); + } + for (auto* tensor : {&a, &base, &gamma, &out}) { + const auto original = *tensor; + for (auto dtype : {vt::DType::kF16, vt::DType::kF32, vt::DType::kI32}) { + tensor->dtype = dtype; CHECK_THROWS_AS(run(good), std::runtime_error); + } + *tensor = original; tensor->device.index = 1; + CHECK_THROWS_AS(run(good), std::runtime_error); + *tensor = original; tensor->data = nullptr; + CHECK_THROWS_AS(run(good), std::runtime_error); + *tensor = original; tensor->data = reinterpret_cast(reinterpret_cast(original.data) + 1); + CHECK_THROWS_AS(run(good), std::runtime_error); + *tensor = original; tensor->shape[0] += 1; + CHECK_THROWS_AS(run(good), std::runtime_error); + *tensor = original; tensor->stride[tensor->rank - 1] = 2; + CHECK_THROWS_AS(run(good), std::runtime_error); + *tensor = original; + } + auto original_out = out; + for (const auto* input : {&a, &base, &delta, &gamma}) { + out.data = input->data; + CHECK_THROWS_AS(run(good, input == &delta ? &delta : nullptr), std::runtime_error); + } + out = original_out; out.data = storage.data() + 1; + CHECK_THROWS_AS(run(good), std::runtime_error); + out = original_out; args = {1e-6f, {vt::ResidualNormExpr::kDeltaPlusAdd, true, true, true}}; + residual.data = out.data; CHECK_THROWS_AS(run(args, &delta, &residual), std::runtime_error); + residual.data = storage.data() + 80; out.data = delta.data; + args.descriptor.output_alias_delta = false; + CHECK_THROWS_AS(run(args, &delta, &residual), std::runtime_error); + out = original_out; args.descriptor.output_alias_delta = true; residual.data = base.data; + args.descriptor.residual_alias_base = false; + CHECK_THROWS_AS(run(args, &delta, &residual), std::runtime_error); + residual.data = storage.data() + 80; + const auto original_a = a; a.stride[0] = 7; + CHECK_THROWS_AS(run(good), std::runtime_error); + a = original_a; a.stride[0] = std::numeric_limits::max(); + CHECK_THROWS_AS(run(good), std::runtime_error); + a = original_a; q.device.index = -1; + CHECK_THROWS_AS(run(good), std::runtime_error); q.device.index = 0; + a.shape[0] = base.shape[0] = out.shape[0] = 0; + a.data = base.data = out.data = nullptr; + CHECK_NOTHROW(run(good)); + CHECK(vt::GetBackend(cpu).GetResidualNormPolicy() == vt::ResidualNormPolicy::kMaterialized); +} + +TEST_CASE("residual RMS norm matches the complete pinned upstream BF16 fixture grid") { + using namespace residual_norm_test; + const char* directory = std::getenv("VT_RESIDUAL_NORM_UPSTREAM"); + if (directory == nullptr) { + MESSAGE("Upstream fixture gate not executed: set VT_RESIDUAL_NORM_UPSTREAM to the complete pinned export"); + return; + } + const std::filesystem::path path(directory); + nlohmann::json manifest; std::ifstream(path / "cases.json") >> manifest; + REQUIRE(manifest["pin"] == "e126687a9a828d513c01a07cd69f025f27d63280"); + REQUIRE(manifest["complete"] == true); + REQUIRE(manifest["cases"].size() >= 192); // one-device core + complete IR grid + const auto read = [&](const nlohmann::json& record) { + const auto file = path / record["file"].get(); + const size_t bytes = record["bytes"]; + REQUIRE(record["dtype"] == "torch.bfloat16"); + REQUIRE(std::filesystem::file_size(file) == bytes); + std::vector values(bytes / sizeof(uint16_t)); + std::ifstream stream(file, std::ios::binary); + stream.read(reinterpret_cast(values.data()), static_cast(bytes)); + REQUIRE(stream.good()); + return values; + }; + size_t executed = 0; + for (const auto& record : manifest["cases"]) { + CAPTURE(record["id"].get()); + REQUIRE(record["seed"] == 0); + const int64_t rows = record["rows"], width = record["width"]; + const int64_t stride = record["files"]["a"]["stride"][0]; + const bool materialize = record["add_residual"]; + const auto gamma = read(record["files"]["gamma"]); + const auto logical_a = read(record["files"]["a"]), logical_base = read(record["files"]["base"]); + const auto expected = read(record["files"]["norm"]), unit_expected = read(record["files"]["unit_norm"]); + const auto expected_residual = materialize ? read(record["files"]["residual"]) : std::vector{}; + std::vector a(rows * stride, Buffer::kGuard), base(a), blank(a); + for (int64_t row = 0; row < rows; ++row) for (int64_t j = 0; j < width; ++j) { + a[row * stride + j] = logical_a[row * width + j]; + base[row * stride + j] = logical_base[row * width + j]; + } + auto devices = Devices(); +#if defined(VLLM_CPP_HIP) + devices.back().index = record["device"]; +#endif + const float atol = record["tolerance"]["atol"], rtol = record["tolerance"]["rtol"]; + const auto compare = [&](const std::vector& actual, const std::vector& reference) { + size_t bad = 0, first = 0; + float maximum = 0; + for (int64_t row = 0; row < rows; ++row) for (int64_t j = 0; j < width; ++j) { + const auto at = static_cast(row * width + j); + const float target = vt::BF16ToF32(reference[at]); + const float error = std::abs(vt::BF16ToF32(actual[row * stride + j]) - target); + if (!(error <= atol + rtol * std::abs(target))) { if (bad == 0) first = at; ++bad; } + maximum = std::max(maximum, error); + } + CAPTURE(first); CAPTURE(maximum); CHECK(bad == 0); + }; + for (const auto device : devices) { + CAPTURE(device.type); CAPTURE(device.index); + Queue queue(device); auto& q = queue.q; + Buffer da(q, a, rows, width, stride), dr(q, base, rows, width, stride), dw(q, gamma, -1, width); + Buffer out(q, blank, rows, width, stride), res(q, blank, rows, width, stride); + vt::ResidualRmsNormArgs args{record["epsilon"], {vt::ResidualNormExpr::kAdd, materialize}}; + Run(0, q, out.t, da.t, dr.t, nullptr, dw.t, args, materialize ? &res.t : nullptr); + const auto direct = out.Read(); compare(direct, expected); + const auto direct_residual = res.Read(); + if (materialize) compare(direct_residual, expected_residual); + Run(1, q, out.t, da.t, dr.t, nullptr, dw.t, args, materialize ? &res.t : nullptr); + CHECK(out.Read() == direct); CHECK(res.Read() == direct_residual); + Buffer unit(q, std::vector(width, vt::F32ToBF16(1)), -1, width); + Run(0, q, out.t, da.t, dr.t, nullptr, unit.t, args, materialize ? &res.t : nullptr); + compare(out.Read(), unit_expected); + CHECK(da.Read() == a); CHECK(dr.Read() == base); CHECK(dw.Read() == gamma); + out.CheckGuard(); res.CheckGuard(); + ++executed; + } + } + MESSAGE("Executed pinned upstream normalization cases across native/reference devices: ", executed); + CHECK(executed >= manifest["cases"].size()); +} + +TEST_CASE("residual RMS norm fusion validates counts order and materialization") { + using namespace residual_norm_test; + const vt::Device cpu{vt::DeviceType::kCPU, 0}; + Queue queue(cpu); auto& q = queue.q; + std::vector values(8, vt::F32ToBF16(1)); + Buffer a(q, values, 1, 8), base(q, values, 1, 8), delta(q, values, 1, 8); + Buffer gamma(q, values, -1, 8), out(q, values, 1, 8), residual(q, values, 1, 8); + vt::ResidualRmsNormArgs args{1e-6f, {vt::ResidualNormExpr::kDeltaPlusAdd, true}}; + vt::FusedBinding binding{}; + binding.n = 6; + binding.op[0] = &a.t; binding.op[1] = &base.t; binding.op[2] = &delta.t; + binding.op[3] = &gamma.t; binding.op[4] = &out.t; binding.op[5] = &residual.t; + const auto good = vt::ResidualRmsNormRecipe(args.descriptor); + auto recipe = good; + const auto run = [&] { vt::FusedChainComposite(q, recipe, binding, {}); }; + CHECK_NOTHROW(run()); + for (uint8_t count : {0, 2, 3, 5}) { + recipe = good; recipe.steps[0].nin = count; CHECK_THROWS_AS(run(), std::runtime_error); + } + recipe = good; recipe.steps[0].residual_norm.expression = static_cast(255); + CHECK_THROWS_AS(run(), std::runtime_error); + recipe = good; recipe.steps[0].out2 = vt::kNoOperand; + CHECK_THROWS_AS(run(), std::runtime_error); + recipe = good; recipe.steps[0].gemma = true; + CHECK_THROWS_AS(run(), std::runtime_error); + recipe = good; binding.op[2] = nullptr; + CHECK_THROWS_AS(run(), std::runtime_error); + binding.op[2] = &delta.t; + CHECK_THROWS_AS(vt::FusedChain(q, out.t, a.t, base.t, &delta.t, gamma.t, vt::ResidualRmsNormArgs{}), std::runtime_error); + CHECK_THROWS_AS(vt::FusedChain(q, out.t, a.t, base.t, nullptr, gamma.t, args, &residual.t), std::runtime_error); + CHECK_THROWS_AS(vt::FusedChain(q, out.t, a.t, base.t, nullptr, gamma.t, vt::ResidualRmsNormArgs{}, &residual.t), std::runtime_error); + // The typed expression is composite-only. Both tier selections must retain + // the explicit opcode rather than reinterpret it as a legacy RMSNorm recipe. + const char* previous = std::getenv("VT_FUSED_TIER"); + const std::string saved = previous == nullptr ? "" : previous; + std::vector first; + for (const char* tier : {"0", "1"}) { + setenv("VT_FUSED_TIER", tier, 1); + vt::FusedChain(q, out.t, a.t, base.t, &delta.t, gamma.t, args, &residual.t); + if (first.empty()) first = out.Read(); else CHECK(out.Read() == first); + } + if (previous == nullptr) unsetenv("VT_FUSED_TIER"); else setenv("VT_FUSED_TIER", saved.c_str(), 1); +} + +#if defined(VLLM_CPP_HIP) +TEST_CASE("ROCm residual RMS norm uses its queue across streams devices and graph replay") { + using namespace residual_norm_test; + int count = 0; + REQUIRE(hipGetDeviceCount(&count) == hipSuccess); + REQUIRE(count >= 2); // This hardware gate owes the two-device upstream arm. + for (int index : {0, 1}) { + const vt::Device device{vt::DeviceType::kROCM, index}; + Queue queue(device), other(device); auto& q = queue.q; + constexpr int64_t width = 769, rows = 3; + std::vector a(rows * width, vt::F32ToBF16(1)); + std::vector base(a.size(), vt::F32ToBF16(0x1p-8f)), delta(base); + std::vector gamma(width, vt::F32ToBF16(0.75f)), blank(a.size(), 0); + Buffer da(q, a, rows, width), dr(q, base, rows, width), dm(q, delta, rows, width); + Buffer dw(q, gamma, -1, width), out(q, blank, rows, width), res(q, blank, rows, width); + vt::ResidualRmsNormArgs args{1e-6f, {vt::ResidualNormExpr::kDeltaPlusAdd, true}}; + std::vector expected_residual; + const auto expected = Reference(a, base, &delta, gamma, rows, width, width, args.eps, &expected_residual); + auto run = [&](vt::Queue& target) { vt::FusedChain(target, out.t, da.t, dr.t, &dm.t, dw.t, args, &res.t); }; + int ambient = -1; + REQUIRE(hipSetDevice(1 - index) == hipSuccess); + run(q); + REQUIRE(hipGetDevice(&ambient) == hipSuccess); CHECK(ambient == 1 - index); + CHECK(out.Read() == expected); CHECK(res.Read() == expected_residual); + run(other.q); + vt::GetBackend(device).Synchronize(other.q); + CHECK(out.Read() == expected); + // A valid tensor device descriptor with a stream owned by another device + // must fail before launch, not silently select the ambient device. + Queue foreign({vt::DeviceType::kROCM, 1 - index}); + vt::Queue wrong{device, foreign.q.handle}; + CHECK_THROWS_AS(run(wrong), std::runtime_error); + REQUIRE(hipGetDevice(&ambient) == hipSuccess); CHECK(ambient == 1 - index); + + DeviceScope capture_device(device); + auto& backend = vt::GetBackend(device); + REQUIRE(backend.SupportsGraphCapture()); + backend.BeginCapture(q); + run(q); + void* graph = backend.EndCaptureGraph(q); + REQUIRE(graph != nullptr); + for (int repeat = 0; repeat < 3; ++repeat) { + backend.Memset(q, out.t.data, 0, blank.size() * sizeof(uint16_t)); + backend.Memset(q, res.t.data, 0, blank.size() * sizeof(uint16_t)); + backend.ReplayGraph(q, graph); + CHECK(out.Read() == expected); CHECK(res.Read() == expected_residual); + } + backend.DestroyGraph(graph); + CHECK(da.Read() == a); CHECK(dr.Read() == base); CHECK(dm.Read() == delta); + out.CheckGuard(); res.CheckGuard(); + } +} + +TEST_CASE("ROCm residual RMS norm cannot run ahead of a blocked nondefault stream") { + using namespace residual_norm_test; + const vt::Device device{vt::DeviceType::kROCM, 0}; + DeviceScope scope(device); + // The default backend queue uses hipStreamCreate. This witness needs an + // independent stream so synchronizing the default stream cannot release it. + // The device scope outlives both the owned stream and its queued work. + struct Stream { + hipStream_t handle = nullptr; + ~Stream() { if (handle != nullptr) (void)hipStreamDestroy(handle); } + } stream; + REQUIRE(hipStreamCreateWithFlags(&stream.handle, hipStreamNonBlocking) == hipSuccess); + vt::Queue q{device, stream.handle}; + unsigned flags = 0; + REQUIRE(hipStreamGetFlags(static_cast(q.handle), &flags) == hipSuccess); + REQUIRE((flags & hipStreamNonBlocking) != 0); + constexpr int64_t width = 128; + std::vector zeros(width), gamma(width, vt::F32ToBF16(1)); + Buffer a(q, zeros, 1, width), base(q, zeros, 1, width), w(q, gamma, -1, width), out(q, zeros, 1, width); + struct Release { + hipStream_t stream; + std::atomic ready{false}; + ~Release() { + ready.store(true); + // Drain the callback before its atomic flag or any operand is destroyed, + // including when an assertion or the operation throws before release. + (void)hipStreamSynchronize(stream); + } + } release{stream.handle}; + REQUIRE(hipLaunchHostFunc(static_cast(q.handle), [](void* pointer) { + auto* ready = static_cast*>(pointer); + while (!ready->load()) std::this_thread::yield(); + }, &release.ready) == hipSuccess); + // Uniform BF16 0x3f3f is nonzero. Correct execution waits for this fill; + // a mutated default-stream launch sees zeros and completes before release. + REQUIRE(hipMemsetAsync(a.t.data, 0x3f, width * sizeof(uint16_t), static_cast(q.handle)) == hipSuccess); + vt::ResidualRmsNorm(q, out.t, a.t, base.t, nullptr, w.t, {}); + REQUIRE(hipStreamSynchronize(nullptr) == hipSuccess); + release.ready.store(true); + CHECK(out.Read() == gamma); +} +#endif diff --git a/tests/vt/test_rocm_moe_grouped_bf16.cpp b/tests/vt/test_rocm_moe_grouped_bf16.cpp new file mode 100644 index 0000000000..e7ec8d3804 --- /dev/null +++ b/tests/vt/test_rocm_moe_grouped_bf16.cpp @@ -0,0 +1,361 @@ +// vllm.cpp original. Native BF16 boundaries and lifetime gates (#3094). +// Numeric boundaries mirror vLLM e126687a9a, fused_moe.py:593-610, +// triton_moe.py:388-527, activation_kernels.cu:44,165-177. +#include "rocm_moe_test_helpers.h" +#include +#include +#include +#include +#include +#include +#include "vt/op_provider.h" + +namespace { +using namespace rocm_moe_test; +using vt::DType; + +// Exact BF16 witnesses, found with seed 7 then frozen. Each distinguishes one +// required narrowing point independently of dot-order and token tolerances. +struct Boundary { + std::array x, gate, up; + uint16_t expected; +}; +constexpr std::array kBoundaries = {{ + {{{0xbffa, 0x401d, 0xbe9d}}, {{0x3e95, 0x4028, 0x3e10}}, + {{0x3fdf, 0xc00d, 0xbf31}}, 0xc247}, + {{{0x3e91, 0x3f72, 0x4022}}, {{0xbf62, 0xbfa7, 0xbe7b}}, + {{0xbf92, 0xbf47, 0xbecd}}, 0x3ef3}, + {{{0xbf93, 0x3eb0, 0x4033}}, {{0x3f65, 0xbfcb, 0xc009}}, + {{0x3f0e, 0x400e, 0x402b}}, 0xbcf5}, +}}; + +void CheckBoundary(vt::Queue& q, const Boundary& fixture) { + Buffer x(q, DType::kBF16, {1, 3}, fixture.x.data()); + Buffer gate(q, DType::kBF16, {3, 1}, fixture.gate.data()); + Buffer up(q, DType::kBF16, {3, 1}, fixture.up.data()); + const int64_t gate_ptr = reinterpret_cast(gate.tensor().data); + const int64_t up_ptr = reinterpret_cast(up.tensor().data); + const int32_t expert = 0; + Buffer gp(q, DType::kI64, {1}, &gate_ptr), uptr(q, DType::kI64, {1}, &up_ptr); + Buffer ids(q, DType::kI32, {1}, &expert), output(q, DType::kBF16, {1, 1}); + vt::MoeGroupedGemmBf16GateUpSiluNative(q, output.tensor(), x.tensor(), ids.tensor(), + nullptr, gp.tensor(), uptr.tensor()); + CHECK(output.Download(q) == std::vector{fixture.expected}); + output.CheckGuard(q); +} + +std::vector Random(int64_t count, uint32_t seed) { + std::mt19937 engine(seed); + std::uniform_real_distribution distribution(-2.0f, 2.0f); + std::vector result(static_cast(count)); + for (auto& value : result) value = vt::F32ToBF16(distribution(engine)); + return result; +} + +// Explicit BF16 gate/up storage is the local native decomposition. The separate +// pinned engine fixture and upstream component gate supply the external oracle. +struct NativeCase { + vt::Queue& q; + int64_t p, n; + std::unique_ptr x, ids, map, gp, up, output, reference; + std::vector> experts; + NativeCase(vt::Queue& queue, int64_t pairs, int64_t columns, int64_t inner, + bool row_map, uint32_t seed) + : q(queue), p(pairs), n(columns) { + HostDeviceScope setup_device(q.device.index); + const int64_t rows = row_map ? (pairs + 2) / 3 : pairs; + const auto activation = Random(rows * inner, seed); + x = std::make_unique(q, DType::kBF16, std::vector{rows, inner}, activation.data()); + std::vector selections(static_cast(pairs)), mapping(static_cast(pairs)); + for (int64_t pair = 0; pair < pairs; ++pair) { + // Four experts, expert 2 empty. The repeated, unordered choices and + // repeated/reversed rows forbid contiguous-expert or token-order shortcuts. + selections[static_cast(pair)] = std::array{3, 0, 3, 1, 0}[pair % 5]; + mapping[static_cast(pair)] = static_cast(rows - 1 - (pair / 3)); + } + ids = std::make_unique(q, DType::kI32, std::vector{pairs}, selections.data()); + if (row_map) + map = std::make_unique(q, DType::kI32, std::vector{pairs}, mapping.data()); + std::vector gate_ptrs, up_ptrs; + for (int expert = 0; expert < 4; ++expert) { + for (int tower = 0; tower < 2; ++tower) { + const auto weight = Random(inner * columns, seed + 1 + expert * 2 + tower); + experts.push_back(std::make_unique(q, DType::kBF16, + std::vector{inner, columns}, weight.data())); + (tower == 0 ? gate_ptrs : up_ptrs).push_back( + reinterpret_cast(experts.back()->tensor().data)); + } + } + gp = std::make_unique(q, DType::kI64, std::vector{4}, gate_ptrs.data()); + up = std::make_unique(q, DType::kI64, std::vector{4}, up_ptrs.data()); + output = std::make_unique(q, DType::kBF16, std::vector{pairs, columns}); + reference = std::make_unique(q, DType::kBF16, std::vector{pairs, columns}); + Buffer gate(q, DType::kBF16, {pairs, columns}), upper(q, DType::kBF16, {pairs, columns}); + vt::MoeGroupedGemmBf16(q, gate.tensor(), x->tensor(), ids->tensor(), + map == nullptr ? nullptr : &map->tensor(), gp->tensor()); + vt::MoeGroupedGemmBf16(q, upper.tensor(), x->tensor(), ids->tensor(), + map == nullptr ? nullptr : &map->tensor(), up->tensor()); + vt::MoeSiluMul(q, reference->tensor(), gate.tensor(), upper.tensor()); + vt::GetBackend(q.device).Synchronize(q); + } + void Launch() { + vt::MoeGroupedGemmBf16GateUpSiluNative(q, output->tensor(), x->tensor(), ids->tensor(), + map == nullptr ? nullptr : &map->tensor(), + gp->tensor(), up->tensor()); + } + void Check() { + CHECK(output->tensor().dtype == DType::kBF16); + CHECK(output->Download(q) == reference->Download(q)); + output->CheckGuard(q); + } +}; +} // namespace + +TEST_CASE("ROCm native MoE preserves gate, up, and SiLU BF16 boundaries") { + RequireDevice(); + Queue queue; + for (const auto& fixture : kBoundaries) CheckBoundary(queue.value, fixture); +} + +TEST_CASE("ROCm native MoE weights before down narrowing and combines preweighted values") { + RequireDevice(); + Queue queue; + auto& q = queue.value; + const std::array x = {0x3df4, 0x3fe4, 0xc016}; + const std::array weight = {0x3fd3, 0xbe7a, 0xbe72}; + const float route = 0.761520803f; + const int32_t expert = 0; + Buffer a(q, DType::kBF16, {1, 3}, x.data()), w(q, DType::kBF16, {3, 1}, weight.data()); + const int64_t ptr = reinterpret_cast(w.tensor().data); + Buffer wp(q, DType::kI64, {1}, &ptr), ids(q, DType::kI32, {1}, &expert); + Buffer routes(q, DType::kF32, {1}, &route), result(q, DType::kBF16, {1, 1}); + vt::MoeGroupedGemmBf16Weighted(q, result.tensor(), a.tensor(), ids.tensor(), + nullptr, wp.tensor(), routes.tensor()); + CHECK(result.Download(q) == std::vector{0x3e76}); + result.CheckGuard(q); + const std::array weighted = {0x3e76, 0xbf40, 0x3e00, 0x3f00, 0x3f80, 0xbf00}; + Buffer values(q, DType::kBF16, {2, 3, 1}, weighted.data()); + Buffer combined(q, DType::kBF16, {2, 1}); + vt::MoeCombinePreweighted(q, combined.tensor(), values.tensor()); + CHECK(combined.Download(q) == std::vector{0xbec5, 0x3f80}); + combined.CheckGuard(q); + const std::array shared = {0.25f, -0.125f}; + Buffer share(q, DType::kF32, {2, 1}, shared.data()), fp32(q, DType::kF32, {2, 1}); + vt::MoeCombinePreweighted(q, fp32.tensor(), values.tensor(), &share.tensor(), 0.5f); + const auto sum = fp32.Download(q); + CHECK(sum[0] == 0.0576171875f); + CHECK(sum[1] == 0.375f); + fp32.CheckGuard(q); +} + +TEST_CASE("ROCm native MoE preserves pair order, empty experts, maps, tails, and repeated launches") { + RequireDevice(); + Queue queue; + for (const auto& shape : std::array, 5>{{ + {6, 8, 1024}, {40, 130, 80}, {1024, 70, 96}, {128, 256, 264}, {65539, 3, 5}}}) { + for (bool map : {false, true}) { + NativeCase fixture(queue.value, shape[0], shape[1], shape[2], map, 7); + for (int repeat = 0; repeat < 3; ++repeat) { + fixture.Launch(); + fixture.Check(); + } + } + } +} + +TEST_CASE("ROCm native MoE capture replays after larger shapes on independent streams") { + RequireDevice(); + Queue first, second; + NativeCase small(first.value, 6, 17, 35, true, 7); + NativeCase other(second.value, 15, 33, 65, false, 41); + small.Launch(); + other.Launch(); + small.Check(); + other.Check(); + auto& backend = vt::GetBackend(Device()); + REQUIRE(backend.SupportsGraphCapture()); + backend.BeginCapture(first.value); + small.Launch(); + void* graph = backend.EndCaptureGraph(first.value); + REQUIRE(graph != nullptr); + NativeCase large(first.value, 1024, 129, 257, true, 91); + large.Launch(); + large.Check(); + for (int repeat = 0; repeat < 3; ++repeat) { + backend.Memset(first.value, small.output->tensor().data, 0, + static_cast(small.p * small.n) * 2); + backend.ReplayGraph(first.value, graph); + other.Launch(); + small.Check(); + other.Check(); + } + backend.DestroyGraph(graph); +} + +TEST_CASE("ROCm native MoE concurrent stream launches keep independent results") { + RequireDevice(); + Queue first, second; + NativeCase a(first.value, 97, 130, 257, true, 7); + NativeCase b(second.value, 513, 65, 129, false, 99); + std::promise release; + const auto ready = release.get_future().share(); + auto run = [&](NativeCase& fixture) { + ready.wait(); + for (int repeat = 0; repeat < 3; ++repeat) fixture.Launch(); + return fixture.output->Download(fixture.q); + }; + auto left = std::async(std::launch::async, [&] { return run(a); }); + auto right = std::async(std::launch::async, [&] { return run(b); }); + release.set_value(); + CHECK(left.get() == a.reference->Download(first.value)); + CHECK(right.get() == b.reference->Download(second.value)); + a.output->CheckGuard(first.value); + b.output->CheckGuard(second.value); +} + +TEST_CASE("ROCm native MoE keeps null default streams on separate devices") { + if (std::getenv("VT_ROCM_MOE_TWO_DEVICES") == nullptr) { + MESSAGE("PENDING: set VT_ROCM_MOE_TWO_DEVICES with two visible devices for this separate gate"); + return; + } + RequireDevice(0); + RequireDevice(1); + vt::Queue q0{Device(0), nullptr}, q1{Device(1), nullptr}; + NativeCase first(q0, 12, 31, 35, true, 7); + NativeCase second(q1, 17, 65, 129, false, 99); + for (int repeat = 0; repeat < 3; ++repeat) { + { + HostDeviceScope opposite(1); + first.Launch(); + CHECK(HostDeviceScope::Current() == 1); + first.Check(); + } + { + HostDeviceScope opposite(0); + second.Launch(); + CHECK(HostDeviceScope::Current() == 0); + second.Check(); + } + } +} + +TEST_CASE("ROCm native MoE zero pairs and output width do not launch") { + RequireDevice(); + Queue queue; + auto& q = queue.value; + const int64_t pointer = 0; + for (auto shape : {std::array{0, 7}, std::array{3, 0}}) { + const int64_t p = shape[0], n = shape[1]; + Buffer a(q, DType::kBF16, {p, 3}), out(q, DType::kBF16, {p, n}); + Buffer ids(q, DType::kI32, {p}), ptr(q, DType::kI64, {1}, &pointer); + Buffer weights(q, DType::kF32, {p}); + vt::MoeGroupedGemmBf16(q, out.tensor(), a.tensor(), ids.tensor(), nullptr, ptr.tensor()); + vt::MoeGroupedGemmBf16GateUpSilu(q, out.tensor(), a.tensor(), ids.tensor(), nullptr, + ptr.tensor(), ptr.tensor()); + vt::MoeGroupedGemmBf16GateUpSiluNative(q, out.tensor(), a.tensor(), ids.tensor(), nullptr, + ptr.tensor(), ptr.tensor()); + vt::MoeGroupedGemmBf16Weighted(q, out.tensor(), a.tensor(), ids.tensor(), nullptr, + ptr.tensor(), weights.tensor()); + auto expert_out = View(out.tensor().data, DType::kBF16, q.device, {p, 1, n}); + vt::MoeCombinePreweighted(q, out.tensor(), expert_out); + out.CheckGuard(q); + } +} + +TEST_CASE("ROCm weighted down and preweighted combine isolate null streams by device") { + if (std::getenv("VT_ROCM_MOE_TWO_DEVICES") == nullptr) { + MESSAGE("PENDING: set VT_ROCM_MOE_TWO_DEVICES with two visible devices for this separate gate"); + return; + } + RequireDevice(0); + RequireDevice(1); + for (int device : {0, 1}) { + vt::Queue q{Device(device), nullptr}; + const std::array x = {0x3df4, 0x3fe4, 0xc016}; + const std::array w = {0x3fd3, 0xbe7a, 0xbe72}; + const float route = 0.761520803f; + const int32_t expert = 0; + Buffer activation(q, DType::kBF16, {1, 3}, x.data()); + Buffer weight(q, DType::kBF16, {3, 1}, w.data()); + const int64_t pointer = reinterpret_cast(weight.tensor().data); + Buffer ptr(q, DType::kI64, {1}, &pointer), ids(q, DType::kI32, {1}, &expert); + Buffer routes(q, DType::kF32, {1}, &route), down(q, DType::kBF16, {1, 1}); + const std::array values = {0x3e76, 0xbf40, 0x3e00, 0x3f00, 0x3f80, 0xbf00}; + Buffer weighted(q, DType::kBF16, {2, 3, 1}, values.data()); + Buffer combined(q, DType::kBF16, {2, 1}); + for (int repeat = 0; repeat < 3; ++repeat) { + HostDeviceScope opposite(1 - device); + vt::MoeGroupedGemmBf16Weighted(q, down.tensor(), activation.tensor(), ids.tensor(), + nullptr, ptr.tensor(), routes.tensor()); + CHECK(HostDeviceScope::Current() == 1 - device); + CHECK(down.Download(q) == std::vector{0x3e76}); + vt::MoeCombinePreweighted(q, combined.tensor(), weighted.tensor()); + CHECK(HostDeviceScope::Current() == 1 - device); + CHECK(combined.Download(q) == std::vector{0xbec5, 0x3f80}); + down.CheckGuard(q); + combined.CheckGuard(q); + } + } +} + +TEST_CASE("ROCm native MoE weighted down preserves F32 output and mapped per-pair routes") { + RequireDevice(); + Queue queue; + auto& q = queue.value; + // Rows: [1,2], [3,-1], [-2,4], [5,-2]. Expert matrices are + // [[2,1],[-1,3]] and [[1,-2],[4,1]]. The FP32 arithmetic is exact. + const std::array x = { + 0x3f80, 0x4000, 0x4040, 0xbf80, 0xc000, 0x4080, 0x40a0, 0xc000}; + const std::array first = {0x4000, 0x3f80, 0xbf80, 0x4040}; + const std::array second = {0x3f80, 0xc000, 0x4080, 0x3f80}; + const std::array selections = {1, 0, 0, 1}, mapping = {2, 0, 2, 1}; + const std::array routes = {0.501953125f, 0.25f, 0.125f, 0.75f}; + Buffer a(q, DType::kBF16, {4, 2}, x.data()); + Buffer w0(q, DType::kBF16, {2, 2}, first.data()), w1(q, DType::kBF16, {2, 2}, second.data()); + const std::array pointers = { + reinterpret_cast(w0.tensor().data), reinterpret_cast(w1.tensor().data)}; + Buffer wp(q, DType::kI64, {2}, pointers.data()), ids(q, DType::kI32, {4}, selections.data()); + Buffer map(q, DType::kI32, {4}, mapping.data()), route(q, DType::kF32, {4}, routes.data()); + // Mapped expert results [14,8], [0,7], [-8,10], [-1,-7], multiplied + // by each pair's route. Route 0 is 257/512, so its FP32 results must retain + // bits that BF16 output rounds away. No expected value calls another MoE operation. + const std::vector expected = {7.02734375f, 4.015625f, 0.0f, 1.75f, + -1.0f, 1.25f, -0.75f, -5.25f}; + const std::vector expected_bf16 = { + 0x40e1, 0x4080, 0x0000, 0x3fe0, 0xbf80, 0x3fa0, 0xbf40, 0xc0a8}; + for (DType dtype : {DType::kF32, DType::kBF16}) { + CAPTURE(dtype); + Buffer output(q, dtype, {4, 2}); + vt::MoeGroupedGemmBf16Weighted(q, output.tensor(), a.tensor(), ids.tensor(), + &map.tensor(), wp.tensor(), route.tensor()); + if (dtype == DType::kF32) + CHECK(output.Download(q) == expected); + else + CHECK(output.Download(q) == expected_bf16); + output.CheckGuard(q); + } +} + +TEST_CASE("ROCm native MoE combines BF16 shared inputs into both output types") { + RequireDevice(); + Queue queue; + auto& q = queue.value; + // Expert sums: [0.5,2.25] and [3.5,1]. Scale by 0.5, then add + // shared rows [129/512,-1/8] and [-1/2,2], giving [257/512,1] and [1.25,2.5]. + // FP32 retains 257/512. BF16 rounds that halfway value to the even neighbor 0.5. + const std::array weighted = { + 0x3f80, 0x4000, 0xbf00, 0x3e80, 0x4040, 0xbf80, 0x3f00, 0x4000}; + const std::array shared = {0x3e81, 0xbe00, 0xbf00, 0x4000}; + Buffer values(q, DType::kBF16, {2, 2, 2}, weighted.data()); + Buffer share(q, DType::kBF16, {2, 2}, shared.data()); + for (DType dtype : {DType::kF32, DType::kBF16}) { + CAPTURE(dtype); + Buffer output(q, dtype, {2, 2}); + vt::MoeCombinePreweighted(q, output.tensor(), values.tensor(), &share.tensor(), 0.5f); + if (dtype == DType::kF32) + CHECK(output.Download(q) == std::vector{0.501953125f, 1.0f, 1.25f, 2.5f}); + else + CHECK(output.Download(q) == std::vector{0x3f00, 0x3f80, 0x3fa0, 0x4020}); + output.CheckGuard(q); + } +} diff --git a/tests/vt/test_rocm_moe_upstream.cpp b/tests/vt/test_rocm_moe_upstream.cpp new file mode 100644 index 0000000000..28514cf3dd --- /dev/null +++ b/tests/vt/test_rocm_moe_upstream.cpp @@ -0,0 +1,133 @@ +// Port of vLLM tests/kernels/moe/test_moe.py:test_fused_moe at +// e126687a9a828d513c01a07cd69f025f27d63280 (#3094). +// rocm_moe_upstream.py preserves the upstream fixtures, reference, modes, +// seed 7, atol 0.02, and rtol 0. This half consumes those exact exported bytes. +#include "rocm_moe_test_helpers.h" +#include +#include +#include +#include + +namespace { +using namespace rocm_moe_test; +using vt::DType; +template +std::vector Read(const std::filesystem::path& path, int64_t count) { + REQUIRE(std::filesystem::file_size(path) == static_cast(count) * sizeof(T)); + std::vector result(static_cast(count)); + std::ifstream source(path, std::ios::binary); + source.read(reinterpret_cast(result.data()), static_cast(result.size() * sizeof(T))); + REQUIRE(source.good()); + return result; +} +float Compare(const std::vector& actual, const std::vector& expected) { + REQUIRE(actual.size() == expected.size()); + size_t bad = 0, first = 0; + float maximum = 0.0f; + for (size_t i = 0; i < actual.size(); ++i) { + const float delta = std::abs(vt::BF16ToF32(actual[i]) - vt::BF16ToF32(expected[i])); + if (!(delta <= 0.02f)) { + if (bad == 0) first = i; + ++bad; + } + maximum = std::max(maximum, delta); + } + CAPTURE(bad); + CAPTURE(first); + CAPTURE(maximum); + CHECK(bad == 0); + return maximum; +} +} // namespace + +TEST_CASE("ROCm native BF16 MoE matches every exported upstream case") { + const char* directory = std::getenv("VT_ROCM_MOE_COMPONENT_CASE"); + if (directory == nullptr) { + MESSAGE("Set VT_ROCM_MOE_COMPONENT_CASE to a pinned upstream export"); + std::exit(77); + } + RequireDevice(); + const std::filesystem::path path(directory); + nlohmann::json record; + std::ifstream(path / "case.json") >> record; + REQUIRE(record["pin"] == "e126687a9a828d513c01a07cd69f025f27d63280"); + REQUIRE(record["seed"] == 7); + REQUIRE(record["renormalize"] == false); + REQUIRE(record["atol"] == 0.02); + REQUIRE(record["rtol"] == 0); + const int64_t m = record["M"], n = record["N"], k = record["K"]; + const int64_t experts = record["experts"], top_k = record["top_k"], pairs = m * top_k; + const int64_t pad = record["padding"].get() ? 128 : 0; + for (const auto& call : record["upstream_calls"]) { + REQUIRE(call["w1_stride"] == std::vector{2 * n * (k + pad), k + pad, 1}); + REQUIRE(call["w2_stride"] == std::vector{k * (n + pad), n + pad, 1}); + } + // The loader's shared Matmul-B contract materializes [K,N] matrices. Padded + // cases retain and compare the actual padded oracle, not a native stride mode. + REQUIRE(record["files"]["gate.bf16"]["export_stride"] == std::vector{k * n, n, 1}); + REQUIRE(record["files"]["up.bf16"]["export_stride"] == std::vector{k * n, n, 1}); + REQUIRE(record["files"]["down.bf16"]["export_stride"] == std::vector{n * k, k, 1}); + auto a = Read(path / "a.bf16", m * k); + auto gate = Read(path / "gate.bf16", experts * k * n); + auto up = Read(path / "up.bf16", experts * k * n); + auto down = Read(path / "down.bf16", experts * n * k); + auto ids = Read(path / "ids.i32", pairs); + auto routes = Read(path / "routes.f32", pairs); + const auto reference = Read(path / "reference.bf16", m * k); + const auto triton = Read(path / "triton.bf16", m * k); + std::vector map(static_cast(pairs)); + for (int64_t pair = 0; pair < pairs; ++pair) + map[static_cast(pair)] = static_cast(pair / top_k); + Queue queue; + auto& q = queue.value; + Buffer da(q, DType::kBF16, {m, k}, a.data()); + Buffer dg(q, DType::kBF16, {experts, k, n}, gate.data()); + Buffer du(q, DType::kBF16, {experts, k, n}, up.data()); + Buffer dd(q, DType::kBF16, {experts, n, k}, down.data()); + Buffer di(q, DType::kI32, {pairs}, ids.data()), dm(q, DType::kI32, {pairs}, map.data()); + Buffer dr(q, DType::kF32, {pairs}, routes.data()); + std::vector gp, uptr, dp; + for (int64_t expert = 0; expert < experts; ++expert) { + gp.push_back(reinterpret_cast(dg.tensor().Ptr() + expert * k * n)); + uptr.push_back(reinterpret_cast(du.tensor().Ptr() + expert * k * n)); + dp.push_back(reinterpret_cast(dd.tensor().Ptr() + expert * n * k)); + } + Buffer dgp(q, DType::kI64, {experts}, gp.data()), dup(q, DType::kI64, {experts}, uptr.data()); + Buffer ddp(q, DType::kI64, {experts}, dp.data()); + Buffer activated(q, DType::kBF16, {pairs, n}), weighted(q, DType::kBF16, {pairs, k}); + Buffer result(q, DType::kBF16, {m, k}); + const auto weighted_view = View(weighted.tensor().data, DType::kBF16, q.device, {m, top_k, k}); + auto run = [&] { + vt::MoeGroupedGemmBf16GateUpSiluNative(q, activated.tensor(), da.tensor(), di.tensor(), + &dm.tensor(), dgp.tensor(), dup.tensor()); + vt::MoeGroupedGemmBf16Weighted(q, weighted.tensor(), activated.tensor(), di.tensor(), + nullptr, ddp.tensor(), dr.tensor()); + vt::MoeCombinePreweighted(q, result.tensor(), weighted_view); + }; + run(); + const auto first = result.Download(q); + const float reference_error = Compare(first, reference); + const float triton_error = Compare(first, triton); + REQUIRE(activated.tensor().dtype == DType::kBF16); + REQUIRE(weighted.tensor().dtype == DType::kBF16); + if (record["use_cudagraph"].get()) { + auto& backend = vt::GetBackend(q.device); + REQUIRE(backend.SupportsGraphCapture()); + backend.BeginCapture(q); + run(); + void* graph = backend.EndCaptureGraph(q); + REQUIRE(graph != nullptr); + for (int repeat = 0; repeat < 3; ++repeat) { + backend.Memset(q, result.tensor().data, 0, static_cast(m * k) * 2); + backend.ReplayGraph(q, graph); + CHECK(result.Download(q) == first); + } + backend.DestroyGraph(graph); + } + activated.CheckGuard(q); + weighted.CheckGuard(q); + result.CheckGuard(q); + record["native_max_abs_vs_reference"] = reference_error; + record["native_max_abs_vs_triton"] = triton_error; + std::ofstream(path / "native-result.json") << record.dump(2) << '\n'; +}