From 91652903776c00aa51e774d6e75cfbd330ca5ddd Mon Sep 17 00:00:00 2001 From: Vikash Loomba Date: Mon, 10 Aug 2026 18:48:29 -0700 Subject: [PATCH 1/7] =?UTF-8?q?spec(rocm):=20BACKEND-ROCM-GDN-KERNELS=20?= =?UTF-8?q?=E2=80=94=20the=20gfx1100=20GDN=20kernel=20family=20spike=20(#4?= =?UTF-8?q?1)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The Qwen3.5-0.8B M2 blocker on discrete ROCm (throws at op 77; the reference tier cannot install on a dGPU) needs ten GDN ops with no ROCm kernel. This commit is the spike spec and the claim; implementation PRs follow one family per PR in throw-order (state-IO, conv, norm-gate, core, fused). Pre-claim validation on 4x gfx1100 / ROCm 7.14 (scratch, recorded in the spec): all ten kernels hand-translated from the CUDA donors and 68/68 standalone checks green (state paths and dtype conversions bit-exact vs an independent host reference); the drop-in TU compiles with the exact production flags; five red-first cross-device cases run 16/16 against the real library with unregistered ops skipping correctly. The key design finding: registering the full indexed set flips IndexedGdnOpsNative(kROCM) true, so the model takes the CUDA lane's device-resident path with zero model or runner edits. M2 disposition follows the #269/#273-ratified distributional gate for the near-tie regime. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: pi:kimi-k3 [pi] --- .agents/NOW.md | 2 +- .agents/coordination.md | 1 + .agents/specs/rocm-gdn-kernels.md | 120 ++++++++++++++++++++++++++++++ 3 files changed, 122 insertions(+), 1 deletion(-) create mode 100644 .agents/specs/rocm-gdn-kernels.md diff --git a/.agents/NOW.md b/.agents/NOW.md index d9c9e4c7b..bd796f22b 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -29,7 +29,7 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | `SERVE-METRICS` async (#277) | **`/metrics` was DEAD on the shipped server**: AsyncLLM folded nothing. Now live, ctest 366/366 | Config-gated families | | `ENG-LOAD-DIRECT-UPLOAD` (#150) | **default ON:** weights VIEW the mmap; 27B load **1.54x warm / 1.61x cold** | merged qkv/gate_up + lm_head | | Vulkan 27B | decode **MET 4.36 vs 4.35**. **LOADMEM: load held the model TWICE, 100.759 -> 53.413 GiB** | Load-phase host build is the new peak | -| `BACKEND-ROCM` | **(b) fix in; #140 gfx1201 hipBLAS + Gemma-4 MoE landed; W0 green** | compile + M2 ([spec](specs/rocm-unified-memory-b.md)) | +| `BACKEND-ROCM` | (b) in; #140 landed; W0 green; gfx1200 M0-M4 (#269) | APU compile+M2; GDN families claimed — `CLAIM-ROCM-GDN-KERNELS` ([spec](specs/rocm-gdn-kernels.md)) | | TP spike #287 (PR #143) | **TP-W1 LANDED**: rank-group table + TP handle (6/6); DSR leak FIXED (unblocks #127/#154/#155) | TP-W2 (linears + loader) | | Release | **ACTIVE; required W1-W11/W13 implemented in #196** | Finish hosted ten-SM proof; rebase/push; run full eight-tuple dry run | | `SAMPLE-PROMPT-LOGPROBS` (#223) | **LANDED** 21/21 | W2 `echo`; CUDA PENDING | diff --git a/.agents/coordination.md b/.agents/coordination.md index 86406ea52..0c29f309c 100644 --- a/.agents/coordination.md +++ b/.agents/coordination.md @@ -1913,6 +1913,7 @@ this claim will meet. The tiled row is speed-gatable on dgx. | `CLAIM-PARAKEET-MODEL-P4` | `MODEL-AUDIO-PARAKEET-ENCODER`, `MODEL-AUDIO-PARAKEET-TRANSDUCER` | Claude Opus 5 (1M context) | isolated worktree `/home/mudler/_git/vllm.cpp-parakeet` (CPU-only) | branch `row/MODEL-AUDIO-PARAKEET-ENCODER`, base `main` | The encoder, the CTC head with greedy collapse, the log-mel front end, an HF-safetensors loader and the RNN-T/TDT transducer. Owns `src/vllm/model_executor/models/parakeet_*`, `src/vllm/multimodal/parakeet_audio_processor.*`, `examples/parakeet_transcribe/`, the Parakeet tests and the two model-matrix rows. | `ACTIVE` | 2026-08-07 CPU-LANDED + gated against a HF `ParakeetForCTC` oracle with EXACT token ids; transcribes ctc-0.6b/1.1b, rnnt-0.6b, tdt-0.6b-v3. No CUDA, no aarch64, no speed number claimed. | | `CLAIM-BACKEND-TENSTORRENT-SPIKE` | `BACKEND-TENSTORRENT` (`INVENTORIED`→`ACTIVE`; W0 skeleton, one op) | Claude Code (claude-sonnet-5) | isolated worktree `/home/lu_zero/Sources/vllmcpp-tenstorrent`; real Blackhole (P150) hardware available all session | branch `row/BACKEND-TENSTORRENT`, base `origin/main` `7534da65`, draft PR [#197](https://github.com/mudler/vllm.cpp/pull/197) | New extension platform (no upstream vLLM equivalent, same class as Metal/Vulkan): `DeviceType::kTENSTORRENT`, `vt::tenstorrent::Backend` (host-staged, discrete PCIe), `TenstorrentPlatform`, ONE op (`kMatmul`, F32/rank-2) via `ttnn::operations::matmul::matmul`. Owns `include/vt/device.h`, `src/vllm/platforms/{platform,tenstorrent}.cpp`, `src/vt/tenstorrent/`, `tests/vt/test_tenstorrent_backend.cpp`, the `VLLM_CPP_TENSTORRENT` CMake wiring, the backend-matrix/porting-inventory/spec, and this claim. | `ACTIVE` | 2026-08-09 W0 LANDED: 3/3 test cases, 8/8 assertions PASS on real hardware (registration, Platform/Backend mirror, kMatmul vs host F32 reference within bf16 tolerance, max_abs_diff 0.03375). Two bugs found+fixed in the same pass (nlohmann-json ABI collision between vllm.cpp's vendored copy and tt-metal's, isolated via a separate OBJECT library; a process-exit segfault from static MeshDevice teardown ordering, fixed by a deliberate leak) — both documented in the spec's Risks/decisions. Not reviewed. No model runs; one op only; host round-trip per call, no perf claim. | | `CLAIM-ENG-LOAD-DIRECT-UPLOAD` | `ENG-LOAD-DIRECT-UPLOAD` (engine-matrix, NEW row -> `ACTIVE`) | Claude Code (opus-5) | isolated worktree `.claude/worktrees/agent-a122a5ddc3546bf31`; GB10 `dgx.casa` tree `~/work/load150/src` (Vulkan `build-vk` + CUDA `build-cuda`), every GPU stage under `flock $HOME/gpu.lock`, `local-ai-worker` parked and restored | branch `row/ENG-LOAD-DIRECT-UPLOAD`, base `origin/main` `375a471e` | Issue #150. Owns ONLY: the refcounted safetensors mapping + byte counters (`include/vllm/model_executor/model_loader/safetensors_reader.h`, `src/vllm/model_executor/model_loader/safetensors_reader.cpp`), `BorrowStTensorBytes` + the `AdoptDeviceBytesAsHost` mmap branch + `OwnedTensor::mmap_src` (`include/vllm/model_executor/models/qwen3_5_weights.h`, `src/vllm/model_executor/models/qwen3_5_weights.cpp`), the qualifying call sites in `include/vllm/model_executor/models/dense_weight_loaders.h` and `src/vllm/model_executor/models/qwen3_5_dense_weights.cpp`, the upload counter in `include/vllm/model_executor/models/dense_attn_block.h`, the `VT_LOAD_STATS` phase timing in `src/vllm/entrypoints/model_loader.cpp`, NEW `tests/vllm/test_load_direct_upload.cpp` + its `tests/CMakeLists.txt` line, NEW `.agents/specs/load-direct-upload.md`, the NEW engine-matrix row + its Loading/Total rollup, the `#150` roadmap intake cell, two `docs/ENVIRONMENT.md` rows, this claim, `.agents/NOW.md`, and the `docs/STATUS.md`/`docs/BENCHMARKS.md` one-liners. **NON-COLLISION:** no new kernel, no model registry change, no CMake target change beyond one test line; every non-verbatim loader helper is untouched and the lever is same-binary A/B-able (`VT_LOAD_DIRECT_UPLOAD=0`). MUST NOT touch README, Metal/SACRED/apex/darwin, or any model/kernel source | `ACTIVE` | 2026-08-09 — mechanism test 6/6 (77) and RED under two mutations; GB10 Vulkan gates on the changed tree green (`test_vulkan_backend` 35/35·2650, `test_backend_cross_device` 11/11·132, `test_opt_paged_engine` 6/6 prompts token-exact 96/96, 0 declines, device type 3); GB10 CUDA full `ctest` 383/393 with BOTH SACRED gates PASS and all 10 failures reproduced on a clean `origin/main` build; 27B bf16 load 1.54x warm / 1.61x cold, bytes moved 100.196 -> 81.260 GiB; `local-ai-worker` restored (`--restart=always`, Up). Fresh scoped review OWED | +| `CLAIM-ROCM-GDN-KERNELS` | `BACKEND-ROCM` (stays `ACTIVE` — M3 GDN kernel slice, issue #41) | pi (kimi-k3) | isolated worktree `/home/vikash/vllm.cpp-rocm-gdn-kernels`, base `origin/main` `5812b8b6`; 4x RX 7900 XTX gfx1100 discrete, ROCm 7.14, GPU jobs hold `~/gpu.lock` | `row/ROCM-GDN-KERNELS` | The M3 GDN kernel family unblocking Qwen3.5-0.8B M2 on discrete ROCm. Owns ONLY: NEW `.agents/specs/rocm-gdn-kernels.md`; NEW `src/vt/rocm/rocm_gdn.hip` (10 ops: kGdnStateGather/Scatter 77/78, kCausalConv1dFwd/Update 5/6, kRmsNormGated 9, kSigmoidGateBf16 58, kGdnPrefill/Decode 10/11, kGdnPostConv 65, kAttnQkNormRopeGate 67 — hand-translated from `src/vt/cuda/cuda_gdn.cu`/`cuda_ops.cu` donors); ADDITIVE registration entries in `src/vt/rocm/rocm_ops.hip`; ADDITIVE GDN cases + buffer helpers in `tests/vt/test_backend_cross_device.cpp`; this claim + `.agents/NOW.md`. **NON-COLLISION:** every new file is ROCm-only; the two shared-file edits are additive registration/case blocks that change no existing behavior (unregistered-op backends skip; CUDA/CPU/Vulkan paths byte-identical). Disjoint from #234 (sampling), #317 (gemma4 FP8/WMMA), #273 (near-tie docs). | `ACTIVE` | 2026-08-11 — **SPEC COMMITTED.** Pre-claim scratch validation on gfx1100 (recorded in the spec): all 10 kernels 68/68 standalone checks (state/conversions bit-exact); drop-in TU compiles with exact production flags; 5 red-first cross-device cases run-verified 16/16 against the real library. NEXT: family PRs in throw-order (GDN-STATE-IO first), each red-first cross-device case + focused gate + full gate + M2 rerun. | **CLOSED same-session claim — `CLAIM-TP-SPIKE-287` (task #287, 2026-08-08, records-only, helper `row/SPIKE-TENSOR-PARALLELISM` draft PR #143, base `b38f78a7`).** Not an active-claims row because every referenced row keeps its diff --git a/.agents/specs/rocm-gdn-kernels.md b/.agents/specs/rocm-gdn-kernels.md new file mode 100644 index 000000000..418042d24 --- /dev/null +++ b/.agents/specs/rocm-gdn-kernels.md @@ -0,0 +1,120 @@ +# ROCm GDN kernel family (gfx1100) — BACKEND-ROCM M3 slice + +**Issue:** [#41](https://github.com/mudler/vllm.cpp/issues/41) (ROCm backend umbrella; roadmap +issue-table row for `BACKEND-ROCM`). +**Claim:** `CLAIM-ROCM-GDN-KERNELS` (coordination.md). +**Base:** pinned `origin/main` `5812b8b6`. +**Hardware:** 4× RX 7900 XTX (`gfx1100`, discrete), ROCm 7.14, Linux Mint 22.3 (reporter's box). + +## The gap, verified + +Qwen3.5-0.8B (`Qwen3_5ForConditionalGeneration`; hybrid GDN — `layer_types` 3:1 +linear:full, `full_attention_interval=4`, `linear_conv_kernel_dim=4`) on discrete gfx1100 +throws `vt: no kernel for op 77 on device type 5` (`src/vt/op_provider.cpp`). The CPU +reference tier cannot install on a dGPU by design (`UnifiedMemory()=false`), so M2 for +this model needs native ROCm kernels for the ten ops the model calls that have none +(evidence: [#41 M2 attempt](https://github.com/mudler/vllm.cpp/issues/41#issuecomment-5230043080); +staged-state note: [#41 follow-up](https://github.com/mudler/vllm.cpp/issues/41#issuecomment-5237179696)). +Classic-dense M2 already runs all-native on this lane (Qwen3-0.6B, zero fallbacks — +that model needs none of these ops). + +## Op inventory (the ten) + +| OpId | Name | Semantics (gdn-semantics.md) | CPU oracle | CUDA donor | +|---|---|---|---|---| +| 77 | `kGdnStateGather` | indexed state rows → f32 working | cpu_ops.cpp:1666 | cuda_gdn.cu:248-335 | +| 78 | `kGdnStateScatter` | f32 working → indexed rows | cpu_ops.cpp:1708 | cuda_gdn.cu:278-335 | +| 5 | `kCausalConv1dFwd` | causal depthwise conv, prefill (§2) | cpu_ops.cpp:1021 | cuda_gdn.cu:479-520 (scalar) | +| 6 | `kCausalConv1dUpdate` | single-step decode conv (§3) | cpu_ops.cpp:1081 | cuda_gdn.cu:884-913 | +| 9 | `kRmsNormGated` | gated RMSNorm, norm_before_gate (§5) | cpu_ops.cpp:1210 | cuda_gdn.cu RmsNormGatedRowKernel | +| 58 | `kSigmoidGateBf16` | sigmoid·mul; **out bf16, attn f32/bf16, gate f32** (ops.cpp contract) | cpu_ops.cpp:2273 | no CUDA reg — CPU composite is the donor | +| 10 | `kGdnPrefill` | gated-delta recurrence (§7/§8) | cpu_ops.cpp:1331 | cuda_gdn.cu:1856 GdnScanKernel | +| 11 | `kGdnDecode` | single-step recurrence | cpu_ops.cpp:1368 | same scan, decode mode | +| 65 | `kGdnPostConv` | fused conv-split+l2norm+g/beta | cpu_ops.cpp:2337 | cuda_gdn.cu:1155 | +| 67 | `kAttnQkNormRopeGate` | fused full-attn preamble (gemma/plain) | cpu_ops.cpp:956 | cuda_ops.cu:1429 area | + +## Execution-path finding (decides the design) + +`qwen3_5.cpp` `IndexedGdnOpsNative()` gates the indexed state-I/O arm on +`kCausalConv1dUpdate` + `kGdnDecode` + `kGdnStateGather` + `kGdnStateScatter` all being +**natively** registered (`OpRegistered` excludes the reference tier). With the full set +registered, the model takes the CUDA lane's device-resident path with **zero model or +runner edits**. The row-copy arm is not a viable discrete fallback (host round trips per +GDN layer per token; the Vulkan record measured the cost). `needs_weight_staging()` for +discrete ROCm stays false in this row — the platform comment defers it to a measured +residency row after M2 runs. + +## Routing decision + +The portable `GdnScanKernel` serves both `kGdnPrefill` (qsl != null) and `kGdnDecode` +(qsl == null, optional state_idx + NULL-block zero-out). The CUDA lane's perf machinery — +WMMA chunked prefill (`#if __CUDA_ARCH__ >= 800`), fused/packed decode, register/tiled +conv variants, Triton-AOT cubins, spec-decode and fp8-quant variants — is NVIDIA-only or +not on this model's path and stays unported (docs/ROCM.md §6). Perf variants are M5 +levers, and upstream's ROCm answer (Triton/CK) is a separate later decision. + +## PR slicing (one family per PR, throw-order; each ends with an M2 rerun) + +1. **GDN-STATE-IO**: 77/78 (the current throw) — this PR's first slice. +2. **GDN-CONV**: 5/6. 3. **GDN-NORMGATE**: 9/58. 4. **GDN-CORE**: 10/11. +5. **GDN-FUSED**: 65/67. After 4 the model completes prefill+decode; 5 restores the fused +fast path (VT_GLUE_FUSE structure parity with the CUDA lane). + +## Pre-claim validation state (scratch, on gfx1100; landed through the lifecycle here) + +All ten kernels were hand-translated from the donors and validated BEFORE any tree +change, in standalone harnesses against independent host references (68/68 checks: +state paths and dtype conversions **bit-exact**; compute outputs at expf/FMA ulp level, +max rel ≤ 2e-5). A complete drop-in TU compiles clean with the build's exact production +flags (`-O3 -ffp-contract=off -std=c++20 --offload-arch=gfx1100`). Five red-first +cross-device cases were assembled into a scratch copy of `test_backend_cross_device.cpp` +and run against the real library: 16/16 green with GDN cases correctly skipping +unregistered ops. Writing them surfaced four `ops.cpp` contract requirements the kernels +satisfy (compact-arm per-token state rows; RmsNormGated rank-3 shape matching; rank-2 +preamble inputs; SigmoidGateBf16's f32 gate). Scratch root: `~/gdn-spike/` on the +reporter's box (harnesses + byte-exact logs). + +## Test plan (red-first per family) + +1. New `test_backend_cross_device` cases per family, CPU-oracle compared — byte paths + (gather/scatter, conv-state write-back/roll) **bit-exact**; arithmetic at NMSE ≤ 5e-4. + Written RED before each registration lands (skip-today, run-once-registered). +2. Model level: the §5.2 M2 rerun per family — Qwen3.5-0.8B vs `--device cpu`, greedy, + tokens + `VT_OP_PROVIDER_STATS` posted on #41; the final family carries the e2e claim + attempt. +3. Full `ctest` suite per family PR; gates reported exactly as observed. + +## M2 disposition (ratified context) + +The near-tie regime is on record: three boards (gfx1100, gfx1103, gfx1200) show +GPU-vs-CPU greedy near-tie flips on Qwen3-0.6B near-tied prompts, deterministic per +backend; joral's #269 analysis (uniform 0.01–0.6% per-layer drift; two real vLLM-ROCm +oracles disagreeing with *each other* at K=5) and the maintainer's #273 merge comment +ratified the **distributional gate** for this regime ("a token-exact bar cannot close +it"), and zero-fallbacks-under-`VT_OP_PROVIDER_STATS` as the discrete-M2 mechanism +evidence. This row's M2 claim for Qwen3.5-0.8B follows that disposition: e2e completion, +all-native op resolution, determinism, and the distributional comparison vs the CPU +backend — not strict token-exactness. + +## Risks + +- **Scan perf at M5** (portable scan vs chunked/fused): accepted — correctness first; + the perf lever is named, not hidden. +- **bf16 state arms at model level**: covered at op level by the bit-exact scratch + checks; the cross-device cases compare f32-state arms (matching the file's tier + structure) — model-level bf16 state behavior is observed in the M2 rerun. +- **Wave64 (gfx9/CDNA)**: none of these kernels use warp-width-sensitive primitives; + the scan's `__syncthreads` structure is width-agnostic. gfx9 boards remain untested + (unverbraucht's MI50 offer on #41) — a later row. + +## Stop conditions + +- A family fails its cross-device gate and the defect is not resolvable from the CPU + oracle + donor re-read → stop, post the failing evidence on #41. +- `IndexedGdnOpsNative(kROCM)` does not flip after family 4 (model path assumption + broken) → stop, NEEDS_DECISION on #41. +- Any maintainer redirection of scope/slicing → this spec is amended in the same change. + +## Outcome + +(pending — filled at DONE) From 152c9e0f94cd528f2557126234929ff44135b42d Mon Sep 17 00:00:00 2001 From: Vikash Loomba Date: Mon, 10 Aug 2026 19:16:04 -0700 Subject: [PATCH 2/7] =?UTF-8?q?feat(rocm):=20GDN=20indexed=20state=20I/O?= =?UTF-8?q?=20(kGdnStateGather/Scatter)=20=E2=80=94=20family=201=20of=20th?= =?UTF-8?q?e=20gfx1100=20GDN=20slice=20(#41)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The current Qwen3.5-0.8B M2 throw on discrete ROCm (op 77). Hand-translated from the cuda_gdn.cu:248-335 donors, gated by the new cross-device case: bit-exact vs the CPU oracle across f32/f16/bf16 cache arms, uniform and spec-widened layouts, and the has_initial_state i8/i32/absent forms (the bar is byte equality — indexed data movement plus the boundary RNE conversion both sides share). Evidence (4x gfx1100, ROCm 7.14, Release): - test_backend_cross_device GDN state case: 48/48 assertions (runs, not skips) - ctest -R 'rocm|cross_device': 3/3 - full ctest: same 7 failures as the unmodified base 5812b8b6 (pre-existing host/lane state — async-scheduling expectations on the ROCm lane and the GDN model tests that fail AT op 77 on base; on this branch the test_loaded_engine_dense throw advances 77 -> op 6) - M2 rerun (Qwen3.5-0.8B): op 77 now resolves vt-native; the throw advances to op 5 (kCausalConv1dFwd, family 2) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: pi:kimi-k3 [pi] --- CMakeLists.txt | 2 + docs/FEATURES.md | 2 +- docs/USAGE.md | 10 ++ src/vt/rocm/rocm_gdn_state.hip | 168 ++++++++++++++++++++++++ src/vt/rocm/rocm_ops.hip | 12 ++ tests/vt/test_backend_cross_device.cpp | 169 +++++++++++++++++++++++++ 6 files changed, 362 insertions(+), 1 deletion(-) create mode 100644 src/vt/rocm/rocm_gdn_state.hip diff --git a/CMakeLists.txt b/CMakeLists.txt index 3f0f009d5..288c056ff 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1312,6 +1312,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_gemma4_expert_geglu.hip src/vt/rocm/rocm_fp8_channel_gemv.hip src/vt/rocm/rocm_moe_router.hip + src/vt/rocm/rocm_gdn_state.hip src/vt/rocm/rocm_ops.hip) if(VLLM_CPP_HIP_ARCHITECTURES) set_source_files_properties( @@ -1326,6 +1327,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_gemma4_expert_geglu.hip src/vt/rocm/rocm_fp8_channel_gemv.hip src/vt/rocm/rocm_moe_router.hip + src/vt/rocm/rocm_gdn_state.hip src/vt/rocm/rocm_ops.hip PROPERTIES HIP_ARCHITECTURES "${VLLM_CPP_HIP_ARCHITECTURES}") endif() diff --git a/docs/FEATURES.md b/docs/FEATURES.md index dbd2d1455..6b56a7b52 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -297,7 +297,7 @@ CPU elementwise GEMM (f32/f16/bf16) runs AVX2 and AVX-512 tiers on x86 where the | LoRA end to end | CPU brick landed | Unwired standalone; not usable through the server | | Multimodal over HTTP | Image request path wired; forward + codec pending | `ROAD-V1-MM` W1-W3 landed (`server_main.cpp:826`). Open: no mm-forward consuming `Request.mm_features`; no image codec vendored (raw RGB only); video/audio/multi-image not started | | Reranking / classify models | Engine side only | Embeddings are LIVE (`LlamaModel`, `vllm_embed`, `/v1/embeddings`); the classify/score heads are landed ops with no registered arch | -| ROCm | W0 verified by community, model e2e pending | Backend + platform + 1 op, ctest-green on gfx1151/1103/1100/1201 ([#41](https://github.com/mudler/vllm.cpp/issues/41)). APU UnifiedMemory fix in (managed allocs, unverified); M2 unblocks with it. [ROCM.md](ROCM.md) | +| ROCm | W0-W1 community-verified on 5 gfx archs; classic-dense e2e runs all-native (near-tie regime, #269); GDN-hybrid blocked on remaining GDN kernels | Backend + platform + #140 op set + GDN state-I/O, ctest-green on gfx1151/1103/1100/1201/1200 ([#41](https://github.com/mudler/vllm.cpp/issues/41)). APU UnifiedMemory fix verified. [ROCM.md](ROCM.md) | | XPU, TPU | Not started | CUDA, CPU, Metal and Vulkan are the built backends | | Custom logits processors on CUDA | Open, not root-caused | Segfaults in a CUDA build, 232/232 green on CPU | | Memory budgeting (`ROAD-V1-MEM`, #83) | M1+M2 landed (absolute bytes) | `--kv-cache-memory` sizes the KV pool from an absolute byte budget (ABI v16, group-aware divisor); `--num-blocks` overrides; `--gpu-memory-utilization` needs the M3 profile run (dgx-gated). See `specs/kv-sizing.md` | diff --git a/docs/USAGE.md b/docs/USAGE.md index 89e172f76..5e26aba94 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -32,6 +32,16 @@ never use, and its teardown can deadlock at process exit — every test passes, ([#132](https://github.com/mudler/vllm.cpp/issues/132)). Setting a build type, or putting your own `-O` in `CMAKE_HIP_FLAGS`, overrides it. +### ROCm op coverage is incremental (and throws are by design) + +The ROCm backend registers native ops family by family +([#41](https://github.com/mudler/vllm.cpp/issues/41)); the GDN indexed state I/O +pair (`kGdnStateGather`/`kGdnStateScatter`) is the first GDN slice. On a +discrete card there is no CPU fallback tier, so a model whose layers call an op +that is not registered yet fails loudly with `vt: no kernel for op N on device +type 5` — that is the memory-safety design working, not a crash. Run with +`VT_OP_PROVIDER_STATS=1` to see which ops resolve native. + ### CUTLASS is fetched as headers only `-DVLLM_CPP_CUTLASS_FETCH=ON` downloads CUTLASS v4.5.0 and stops there: the diff --git a/src/vt/rocm/rocm_gdn_state.hip b/src/vt/rocm/rocm_gdn_state.hip new file mode 100644 index 000000000..23ebfcc34 --- /dev/null +++ b/src/vt/rocm/rocm_gdn_state.hip @@ -0,0 +1,168 @@ +// ROCm GDN indexed state I/O (BACKEND-ROCM-GDN-KERNELS family 1; issue #41). +// Hand-translation of src/vt/cuda/cuda_gdn.cu:248-335 (GdnStateGatherKernel / +// GdnStateScatterKernel + host dispatch), readable side by side against the +// donor. CPU oracle: src/vt/cpu/cpu_ops.cpp:1666/1708. kGdnStateGather/Scatter +// are the first two ops `IndexedGdnOpsNative()` (qwen3_5.cpp) requires for the +// device-resident GDN state path, and the current Qwen3.5-0.8B M2 throw on +// discrete boards (op 77). Semantics: .agents/specs/gdn-semantics.md; +// .agents/specs/rocm-gdn-kernels.md is this family's spec. +// +// Gate: the "GDN state gather/scatter" case in test_backend_cross_device.cpp — +// bit-exact vs the CPU oracle (pure indexed data movement + boundary RNE +// conversion both sides share). + +#include +#include +#include + +#include +#include +#include + +#include "vt/ops.h" + +namespace vt::rocm { +namespace { + +constexpr int kBlock = 256; + +inline void Check(hipError_t err, const char* what) { + if (err != hipSuccess) { + throw std::runtime_error(std::string("vt rocm gdn: ") + what + ": " + + hipGetErrorString(err)); + } +} +inline hipStream_t AsStream(const Queue& q) { + return static_cast(q.handle); +} +inline unsigned GridFor(int64_t n) { + if (n <= 0) return 1; + const int64_t g = (n + kBlock - 1) / kBlock; + return static_cast(g > 65535 ? 65535 : g); +} + +__device__ inline float Ld(const float* p, int64_t i) { return p[i]; } +__device__ inline float Ld(const __half* p, int64_t i) { return __half2float(p[i]); } +__device__ inline float Ld(const __hip_bfloat16* p, int64_t i) { + return __bfloat162float(p[i]); +} +__device__ inline void St(float* p, int64_t i, float v) { p[i] = v; } +__device__ inline void St(__half* p, int64_t i, float v) { p[i] = __float2half_rn(v); } +__device__ inline void St(__hip_bfloat16* p, int64_t i, float v) { + p[i] = __float2bfloat16(v); +} + +// ── State I/O (donor cuda_gdn.cu:248-298) ──────────────────────────────────── +template +__global__ void GdnStateGatherK(float* working, const TCache* cache, + const int32_t* state_idx, const int8_t* has_i8, + const int32_t* has_i32, int64_t row_elems, + int64_t n, int64_t work_inner, + int64_t cache_inner) { + const bool widened = cache_inner != work_inner; + const int64_t cache_row = (row_elems / work_inner) * cache_inner; + const int64_t step = static_cast(gridDim.x) * blockDim.x; + for (int64_t i = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + i < n; i += step) { + const int64_t row = i / row_elems; + const int64_t col = i - row * row_elems; + const bool keep = has_i8 != nullptr ? has_i8[row] != 0 + : has_i32 != nullptr ? has_i32[row] != 0 + : true; + int64_t cache_off; + if (!widened) { + cache_off = static_cast(state_idx[row]) * row_elems + col; + } else { + const int64_t m = col / work_inner, wcol = col - m * work_inner; + cache_off = static_cast(state_idx[row]) * cache_row + + m * cache_inner + wcol; + } + working[i] = keep ? Ld(cache, cache_off) : 0.0f; + } +} + +template +__global__ void GdnStateScatterK(TCache* cache, const float* working, + const int32_t* state_idx, int64_t row_elems, + int64_t n, int64_t work_inner, + int64_t cache_inner) { + const bool widened = cache_inner != work_inner; + const int64_t cache_row = (row_elems / work_inner) * cache_inner; + const int64_t step = static_cast(gridDim.x) * blockDim.x; + for (int64_t i = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + i < n; i += step) { + const int64_t row = i / row_elems; + const int64_t col = i - row * row_elems; + int64_t cache_off; + if (!widened) { + cache_off = static_cast(state_idx[row]) * row_elems + col; + } else { + const int64_t m = col / work_inner, wcol = col - m * work_inner; + cache_off = static_cast(state_idx[row]) * cache_row + + m * cache_inner + wcol; + } + St(cache, cache_off, working[i]); + } +} + +} // namespace + +void GdnStateGatherKernelRocm(Queue& q, Tensor& working, const Tensor& cache, + const Tensor& state_idx, const Tensor* has_initial_state) { + const int64_t rows = state_idx.shape[0]; + if (rows == 0) return; + const int64_t n = working.Numel(); + const int64_t row_elems = n / rows; + const int64_t work_inner = working.shape[working.rank - 1]; + const int64_t cache_inner = cache.shape[cache.rank - 1]; + const int8_t* has_i8 = + has_initial_state != nullptr && has_initial_state->dtype == DType::kI8 + ? has_initial_state->Ptr() + : nullptr; + const int32_t* has_i32 = + has_initial_state != nullptr && has_initial_state->dtype == DType::kI32 + ? has_initial_state->Ptr() + : nullptr; + hipStream_t s = AsStream(q); + if (cache.dtype == DType::kBF16) { + GdnStateGatherK<<>>( + working.Ptr(), cache.Ptr<__hip_bfloat16>(), state_idx.Ptr(), + has_i8, has_i32, row_elems, n, work_inner, cache_inner); + } else if (cache.dtype == DType::kF16) { + GdnStateGatherK<<>>( + working.Ptr(), cache.Ptr<__half>(), state_idx.Ptr(), + has_i8, has_i32, row_elems, n, work_inner, cache_inner); + } else { + GdnStateGatherK<<>>( + working.Ptr(), cache.Ptr(), state_idx.Ptr(), + has_i8, has_i32, row_elems, n, work_inner, cache_inner); + } + Check(hipGetLastError(), "gdn state gather launch"); +} + +void GdnStateScatterKernelRocm(Queue& q, Tensor& cache, const Tensor& working, + const Tensor& state_idx) { + const int64_t rows = state_idx.shape[0]; + if (rows == 0) return; + const int64_t n = working.Numel(); + const int64_t row_elems = n / rows; + const int64_t work_inner = working.shape[working.rank - 1]; + const int64_t cache_inner = cache.shape[cache.rank - 1]; + hipStream_t s = AsStream(q); + if (cache.dtype == DType::kBF16) { + GdnStateScatterK<<>>( + cache.Ptr<__hip_bfloat16>(), working.Ptr(), state_idx.Ptr(), + row_elems, n, work_inner, cache_inner); + } else if (cache.dtype == DType::kF16) { + GdnStateScatterK<<>>( + cache.Ptr<__half>(), working.Ptr(), state_idx.Ptr(), + row_elems, n, work_inner, cache_inner); + } else { + GdnStateScatterK<<>>( + cache.Ptr(), working.Ptr(), state_idx.Ptr(), + row_elems, n, work_inner, cache_inner); + } + Check(hipGetLastError(), "gdn state scatter launch"); +} + +} // namespace vt::rocm diff --git a/src/vt/rocm/rocm_ops.hip b/src/vt/rocm/rocm_ops.hip index bac34a990..c9b4dd173 100644 --- a/src/vt/rocm/rocm_ops.hip +++ b/src/vt/rocm/rocm_ops.hip @@ -38,6 +38,12 @@ void GeluTanhKernelRocm(Queue& q, Tensor& out, const Tensor& x); void GeluErfKernelRocm(Queue& q, Tensor& out, const Tensor& x); void MoeRouterTopKKernelRocm(Queue& q, Tensor& weights, Tensor& indices, const Tensor& logits, const MoeRouterTopKArgs& args, const Tensor* bias); +// BACKEND-ROCM-GDN-KERNELS family 1 (rocm_gdn_state.hip): the indexed state I/O +// pair `IndexedGdnOpsNative()` requires (issue #41, spec rocm-gdn-kernels.md). +void GdnStateGatherKernelRocm(Queue& q, Tensor& working, const Tensor& cache, + const Tensor& state_idx, const Tensor* has_initial_state); +void GdnStateScatterKernelRocm(Queue& q, Tensor& cache, const Tensor& working, + const Tensor& state_idx); namespace { @@ -95,6 +101,12 @@ struct Registrar { RegisterOp(OpId::kMoeRouterTopK, DeviceType::kROCM, reinterpret_cast( static_cast(&MoeRouterTopKKernelRocm))); + RegisterOp(OpId::kGdnStateGather, DeviceType::kROCM, + reinterpret_cast( + static_cast(&GdnStateGatherKernelRocm))); + RegisterOp(OpId::kGdnStateScatter, DeviceType::kROCM, + reinterpret_cast( + static_cast(&GdnStateScatterKernelRocm))); } } registrar; diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index 29e4b4abd..ae90d5abb 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -1064,6 +1064,175 @@ TEST_CASE("FusedChain matches the CPU oracle within NMSE <= 5e-4 (both tiers)") // registered unified devices (Metal M4, GB10 CUDA/Vulkan) the pointer is // host-accessible, so the fallback runs. On a plain CPU build there is no non-CPU // device and the case is inert. +// i32 device buffer (state_idx / query_start_loc / has_initial_state / +// conv_state_indices). Same staging discipline as DevBuf. +class DevBufI32 { + public: + DevBufI32(vt::Backend& b, Queue& q, size_t n) : b_(b), q_(q), n_(n) { + ptr_ = b_.Alloc(n * sizeof(int32_t)); + } + ~DevBufI32() { b_.Free(ptr_); } + DevBufI32(const DevBufI32&) = delete; + DevBufI32& operator=(const DevBufI32&) = delete; + void Upload(const std::vector& src) { + REQUIRE(src.size() == n_); + b_.Copy(q_, ptr_, src.data(), n_ * sizeof(int32_t)); + } + void* ptr() const { return ptr_; } + + private: + vt::Backend& b_; + Queue& q_; + size_t n_; + void* ptr_ = nullptr; +}; + +// Byte-addressed device buffer for i8 masks (has_initial_state) and u16 bf16 +// cache contents (sized in ELEMENTS of the templated width). +class DevBufBytes { + public: + DevBufBytes(vt::Backend& b, Queue& q, size_t nbytes) : b_(b), q_(q), n_(nbytes) { + ptr_ = b_.Alloc(nbytes); + } + ~DevBufBytes() { b_.Free(ptr_); } + DevBufBytes(const DevBufBytes&) = delete; + DevBufBytes& operator=(const DevBufBytes&) = delete; + void Upload(const void* src) { b_.Copy(q_, ptr_, src, n_); } + void Download(void* dst) { + b_.Synchronize(q_); + b_.Copy(q_, dst, ptr_, n_); + b_.Synchronize(q_); + } + void* ptr() const { return ptr_; } + + private: + vt::Backend& b_; + Queue& q_; + size_t n_; + void* ptr_ = nullptr; +}; + +// f32 -> bf16 bits through the CPU backend's own cast op, so the test never +// reimplements the codec it is comparing against. +std::vector Bf16Bits(const std::vector& src) { + vt::Backend& cpu = vt::GetBackend(DeviceType::kCPU); + Queue cq = cpu.CreateQueue(); + const Device cd{DeviceType::kCPU, 0}; + std::vector in = src; + std::vector out(src.size(), 0); + Tensor tin = T1(in.data(), cd, static_cast(in.size())); + Tensor tout = Tensor::Contiguous(out.data(), DType::kBF16, cd, + {static_cast(out.size())}); + vt::CastBf16(cq, tout, tin); + cpu.DestroyQueue(cq); + return out; +} + + +// --- GDN cases (BACKEND-ROCM-GDN-KERNELS) ------------------------------------- + +TEST_CASE("GDN state gather/scatter are BIT-EXACT against the CPU oracle") { + // Indexed data movement between an f32 working set and a persistent cache: + // no arithmetic anywhere, so the bar is byte equality — including the bf16 + // cache arm, where both sides apply the same RNE round at the boundary. + // Covers the uniform layout (cache_inner == work_inner), the spec-widened + // layout (leading work_inner cols per channel at the physical stride), the + // has_initial_state mask in i8/i32/absent forms, and scatter's untouched-row + // preservation. + const int64_t S = 8, R = 4, mid = 4, w_in = 6, c_in = 8; // c_in>w_in: widened + for (bool widened : {false, true}) { + const int64_t cache_inner = widened ? c_in : w_in; + CAPTURE(widened); + const size_t cache_n = static_cast(S * mid * cache_inner); + const size_t work_n = static_cast(R * mid * w_in); + const std::vector cache_f = RandomVec(cache_n, 910); + const std::vector cache_bf = Bf16Bits(cache_f); + const std::vector idx = {1, 0, 7, 6}; // unique slots + const std::vector has32 = {1, 0, 1, 0}; + const std::vector has8 = {1, 0, 1, 0}; + + for (int arm = 0; arm < 2; ++arm) { // 0 = f32 cache, 1 = bf16 cache + CAPTURE(arm); + for (int has = 0; has < 3; ++has) { // 0 = absent, 1 = i8, 2 = i32 + CAPTURE(has); + // ---- CPU oracle: gather, then scatter the gathered rows back. + std::vector ref_work(work_n, -1.0f); + std::vector ref_cache_f = cache_f; + std::vector ref_cache_bf = cache_bf; + std::vector ci = idx, ch32 = has32; + std::vector ch8 = has8; + { + vt::Backend& cpu = vt::GetBackend(DeviceType::kCPU); + Queue cq = cpu.CreateQueue(); + const Device cd{DeviceType::kCPU, 0}; + Tensor tidx = TI32(ci.data(), cd, R); + Tensor th8 = Tensor::Contiguous(ch8.data(), DType::kI8, cd, {R}); + Tensor th32 = TI32(ch32.data(), cd, R); + Tensor tw = Tensor::Contiguous(ref_work.data(), DType::kF32, cd, {R, mid, w_in}); + const Tensor* ph = has == 1 ? static_cast(&th8) + : has == 2 ? static_cast(&th32) + : nullptr; + if (arm == 0) { + Tensor tc = Tensor::Contiguous(ref_cache_f.data(), DType::kF32, cd, + {S, mid, cache_inner}); + vt::GdnStateGather(cq, tw, tc, tidx, ph); + vt::GdnStateScatter(cq, tc, tw, tidx); + } else { + Tensor tc = Tensor::Contiguous(ref_cache_bf.data(), DType::kBF16, cd, + {S, mid, cache_inner}); + vt::GdnStateGather(cq, tw, tc, tidx, ph); + vt::GdnStateScatter(cq, tc, tw, tidx); + } + cpu.DestroyQueue(cq); + } + + for (DeviceType dt : RegisteredDevices()) { + if (!OpAvailable(vt::OpId::kGdnStateGather, dt) || + !OpAvailable(vt::OpId::kGdnStateScatter, dt)) + continue; + CAPTURE(DeviceName(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + const Device d{dt, 0}; + DevBuf dwork(dev, q, work_n); + DevBufI32 didx(dev, q, R), dhas32(dev, q, R); + DevBufBytes dhas8(dev, q, R); + didx.Upload(idx); + dhas32.Upload(has32); + dhas8.Upload(has8.data()); + Tensor tidx = TI32(didx.ptr(), d, R); + Tensor th8 = Tensor::Contiguous(dhas8.ptr(), DType::kI8, d, {R}); + Tensor th32 = TI32(dhas32.ptr(), d, R); + Tensor tw = Tensor::Contiguous(dwork.ptr(), DType::kF32, d, {R, mid, w_in}); + const Tensor* ph = has == 1 ? &th8 : has == 2 ? &th32 : nullptr; + + const size_t cache_bytes = cache_n * (arm == 0 ? 4 : 2); + DevBufBytes dcache(dev, q, cache_bytes); + dcache.Upload(arm == 0 ? static_cast(cache_f.data()) + : static_cast(cache_bf.data())); + Tensor tc = Tensor::Contiguous(dcache.ptr(), + arm == 0 ? DType::kF32 : DType::kBF16, d, + {S, mid, cache_inner}); + vt::GdnStateGather(q, tw, tc, tidx, ph); + CHECK(dwork.Download() == ref_work); // gather: bit-exact + vt::GdnStateScatter(q, tc, tw, tidx); + if (arm == 0) { + std::vector got(cache_n); + dcache.Download(got.data()); + CHECK(got == ref_cache_f); // scatter round-trip: bit-exact + } else { + std::vector got(cache_n); + dcache.Download(got.data()); + CHECK(got == ref_cache_bf); + } + dev.DestroyQueue(q); + } + } + } + } +} + + TEST_CASE("reference tier: an op with no native kernel matches the CPU oracle (unified only)") { constexpr int64_t kRows = 7, kCols = 48; constexpr size_t kN = kRows * kCols; From 0f50c929e18ca9082c14d470d0580508e179afe3 Mon Sep 17 00:00:00 2001 From: Vikash Loomba Date: Mon, 10 Aug 2026 19:39:13 -0700 Subject: [PATCH 3/7] =?UTF-8?q?feat(rocm):=20causal=20conv1d=20fwd/update?= =?UTF-8?q?=20(kCausalConv1dFwd/Update)=20=E2=80=94=20family=202=20of=20th?= =?UTF-8?q?e=20gfx1100=20GDN=20slice=20(#41)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Hand-translated from the cuda_gdn.cu scalar donors (:479/:884). The exact-chunks descriptor form IS supported — the first draft's loud VT_CHECK against it fired in the M2 rerun, proving Qwen3.5 prefill passes VT_CONV_EXACT_CHUNKS descriptors on the live path; the kernel takes the donor Reg kernel's chunk addressing over the scalar body (same tap order; write-back owned by the block whose chunk reaches the sequence end). The register-resident/tiled CUDA perf variants stay unported (docs/ROCM.md §6). Evidence (4x gfx1100, ROCm 7.14, Release): - conv cross-device case: 33/33 assertions (legacy + descriptor arms; outputs NMSE, state write-back/roll bit-exact) - ctest -R 'rocm|cross_device': 3/3 - full ctest: same 7 pre-existing failures as base (zero delta) - M2 rerun (Qwen3.5-0.8B): ops 5/6 resolve vt-native; the throw advances 5 -> op 65 (kGdnPostConv, the fused post-conv glue — family 3, reordered to throw-order) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: pi:kimi-k3 [pi] --- CMakeLists.txt | 2 + docs/FEATURES.md | 2 +- docs/USAGE.md | 6 +- src/vt/rocm/rocm_gdn_conv.hip | 255 +++++++++++++++++++++++++ src/vt/rocm/rocm_ops.hip | 15 ++ tests/vt/test_backend_cross_device.cpp | 142 ++++++++++++++ 6 files changed, 419 insertions(+), 3 deletions(-) create mode 100644 src/vt/rocm/rocm_gdn_conv.hip diff --git a/CMakeLists.txt b/CMakeLists.txt index 288c056ff..dc4a580dd 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1313,6 +1313,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_fp8_channel_gemv.hip src/vt/rocm/rocm_moe_router.hip src/vt/rocm/rocm_gdn_state.hip + src/vt/rocm/rocm_gdn_conv.hip src/vt/rocm/rocm_ops.hip) if(VLLM_CPP_HIP_ARCHITECTURES) set_source_files_properties( @@ -1328,6 +1329,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_fp8_channel_gemv.hip src/vt/rocm/rocm_moe_router.hip src/vt/rocm/rocm_gdn_state.hip + src/vt/rocm/rocm_gdn_conv.hip src/vt/rocm/rocm_ops.hip PROPERTIES HIP_ARCHITECTURES "${VLLM_CPP_HIP_ARCHITECTURES}") endif() diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 6b56a7b52..f2c2260bc 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -297,7 +297,7 @@ CPU elementwise GEMM (f32/f16/bf16) runs AVX2 and AVX-512 tiers on x86 where the | LoRA end to end | CPU brick landed | Unwired standalone; not usable through the server | | Multimodal over HTTP | Image request path wired; forward + codec pending | `ROAD-V1-MM` W1-W3 landed (`server_main.cpp:826`). Open: no mm-forward consuming `Request.mm_features`; no image codec vendored (raw RGB only); video/audio/multi-image not started | | Reranking / classify models | Engine side only | Embeddings are LIVE (`LlamaModel`, `vllm_embed`, `/v1/embeddings`); the classify/score heads are landed ops with no registered arch | -| ROCm | W0-W1 community-verified on 5 gfx archs; classic-dense e2e runs all-native (near-tie regime, #269); GDN-hybrid blocked on remaining GDN kernels | Backend + platform + #140 op set + GDN state-I/O, ctest-green on gfx1151/1103/1100/1201/1200 ([#41](https://github.com/mudler/vllm.cpp/issues/41)). APU UnifiedMemory fix verified. [ROCM.md](ROCM.md) | +| ROCm | W0-W1 community-verified on 5 gfx archs; classic-dense e2e runs all-native (near-tie regime, #269); GDN-hybrid blocked on remaining GDN kernels | Backend + platform + #140 op set + GDN state-I/O and causal-conv1d ops, ctest-green on gfx1151/1103/1100/1201/1200 ([#41](https://github.com/mudler/vllm.cpp/issues/41)). APU UnifiedMemory fix verified. [ROCM.md](ROCM.md) | | XPU, TPU | Not started | CUDA, CPU, Metal and Vulkan are the built backends | | Custom logits processors on CUDA | Open, not root-caused | Segfaults in a CUDA build, 232/232 green on CPU | | Memory budgeting (`ROAD-V1-MEM`, #83) | M1+M2 landed (absolute bytes) | `--kv-cache-memory` sizes the KV pool from an absolute byte budget (ABI v16, group-aware divisor); `--num-blocks` overrides; `--gpu-memory-utilization` needs the M3 profile run (dgx-gated). See `specs/kv-sizing.md` | diff --git a/docs/USAGE.md b/docs/USAGE.md index 5e26aba94..cc1562978 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -35,8 +35,10 @@ or putting your own `-O` in `CMAKE_HIP_FLAGS`, overrides it. ### ROCm op coverage is incremental (and throws are by design) The ROCm backend registers native ops family by family -([#41](https://github.com/mudler/vllm.cpp/issues/41)); the GDN indexed state I/O -pair (`kGdnStateGather`/`kGdnStateScatter`) is the first GDN slice. On a +([#41](https://github.com/mudler/vllm.cpp/issues/41)); landed GDN slices so far: +the indexed state I/O pair (`kGdnStateGather`/`kGdnStateScatter`) and the causal +conv1d pair (`kCausalConv1dFwd`/`kCausalConv1dUpdate`, incl. the exact-chunks +descriptor form Qwen3.5 prefill passes). On a discrete card there is no CPU fallback tier, so a model whose layers call an op that is not registered yet fails loudly with `vt: no kernel for op N on device type 5` — that is the memory-safety design working, not a crash. Run with diff --git a/src/vt/rocm/rocm_gdn_conv.hip b/src/vt/rocm/rocm_gdn_conv.hip new file mode 100644 index 000000000..bec86c785 --- /dev/null +++ b/src/vt/rocm/rocm_gdn_conv.hip @@ -0,0 +1,255 @@ +// ROCm GDN causal conv1d (BACKEND-ROCM-GDN-KERNELS family 2; issue #41). +// Hand-translation of the cuda_gdn.cu scalar donors — CausalConv1dFwdKernel +// (:479, prefill, gdn-semantics.md §2) and CausalConv1dUpdateKernel (:884, +// decode seqlen==1, §3) — readable side by side against them. CPU oracles: +// src/vt/cpu/cpu_ops.cpp:1021/1081. The register-resident (VT_CONV_REG) and +// tiled (VT_CONV_TILED) CUDA variants are perf levers and stay unported +// (docs/ROCM.md §6). The exact-chunks descriptor form (VT_CONV_EXACT_CHUNKS) +// IS supported — same chunk addressing as the donor's Reg kernel over the +// scalar body — because Qwen3.5 prefill passes descriptors on the live path +// (the first draft's loud VT_CHECK guard against them fired in the M2 rerun, +// proving it). +// +// Gate: the "causal conv1d fwd/update" case in test_backend_cross_device.cpp — +// outputs NMSE <= 5e-4 vs the CPU oracle (silu epilogue arithmetic), conv_state +// write-back/roll BIT-EXACT (raw-x byte movement). + +#include +#include +#include + +#include +#include +#include + +#include "vt/ops.h" + +namespace vt::rocm { +namespace { + +constexpr int kBlock = 256; + +inline void Check(hipError_t err, const char* what) { + if (err != hipSuccess) { + throw std::runtime_error(std::string("vt rocm gdn: ") + what + ": " + + hipGetErrorString(err)); + } +} +inline hipStream_t AsStream(const Queue& q) { + return static_cast(q.handle); +} +inline unsigned GridFor(int64_t n) { + if (n <= 0) return 1; + const int64_t g = (n + kBlock - 1) / kBlock; + return static_cast(g > 65535 ? 65535 : g); +} + +__device__ inline float Ld(const float* p, int64_t i) { return p[i]; } +__device__ inline float Ld(const __half* p, int64_t i) { return __half2float(p[i]); } +__device__ inline float Ld(const __hip_bfloat16* p, int64_t i) { + return __bfloat162float(p[i]); +} +__device__ inline void St(float* p, int64_t i, float v) { p[i] = v; } +__device__ inline void St(__half* p, int64_t i, float v) { p[i] = __float2half_rn(v); } +__device__ inline void St(__hip_bfloat16* p, int64_t i, float v) { + p[i] = __float2bfloat16(v); +} +__device__ inline float Silu(float x) { return x / (1.0f + expf(-x)); } + +// ── Conv (donors cuda_gdn.cu:479 scalar fwd, :884 update) ─────────────────── +constexpr int kConvExactM = 8; // upstream compute_causal_conv1d_metadata BLOCK_M + +// exact mode (batch_ptr/token_chunk_offset_ptr non-null): blockIdx.y indexes +// PROGRAMS, program p owning sequence batch_ptr[p] and its +// token_chunk_offset_ptr[p]-th 8-token chunk — the upstream +// compute_causal_conv1d_metadata work split, numerically identical to the +// legacy whole-sequence mapping (same tap order; the state write-back belongs +// to the block whose chunk reaches the sequence end). +template +__global__ void CausalConv1dFwdK(Tout* out, const Tin* x, const Tin* w, + const Tin* bias, float* conv_state, + const int32_t* qsl, const THas* his, + int64_t c_dim, int64_t x_row_stride, int64_t k, + bool silu, const int32_t* batch_ptr, + const int32_t* token_chunk_offset_ptr) { + const bool exact = batch_ptr != nullptr; + const int64_t s = exact ? batch_ptr[blockIdx.y] : blockIdx.y; + const int64_t c = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + if (c >= c_dim) return; + const int64_t width = k - 1; + const int64_t begin = qsl[s]; + const int64_t t_len = qsl[s + 1] - begin; + const int64_t token_offset = + exact ? static_cast(token_chunk_offset_ptr[blockIdx.y]) * kConvExactM : 0; + // Chunks past the sequence end are skipped — except chunk 0, which also owns + // the t_len == 0 corner. + if (token_offset > 0 && token_offset >= t_len) return; + const int64_t token_end = + (exact && token_offset + kConvExactM < t_len) ? token_offset + kConvExactM : t_len; + const bool init = his[s] != 0; + float* srow = conv_state + (s * c_dim + c) * width; + const float b = bias != nullptr ? Ld(bias, c) : 0.0f; + for (int64_t t = token_offset; t < token_end; ++t) { + float acc = b; + for (int64_t j = 0; j < k; ++j) { + const int64_t ti = t - (k - 1 - j); + float v = 0.0f; + if (ti >= 0) { + v = Ld(x, (begin + ti) * x_row_stride + c); + } else if (init) { + v = srow[width + ti]; + } + acc += Ld(w, c * k + j) * v; + } + St(out, (begin + t) * c_dim + c, silu ? Silu(acc) : acc); + } + if (token_end == t_len) { + for (int64_t j = 0; j < width; ++j) { + const int64_t tj = t_len - width + j; + float v = 0.0f; + if (tj >= 0) { + v = Ld(x, (begin + tj) * x_row_stride + c); + } else if (init) { + v = srow[width + tj]; + } + srow[j] = v; + } + } +} + +template +__global__ void CausalConv1dUpdateK(Tout* out, const Tin* x, const Tin* w, + const Tin* bias, TState* conv_state, + const int32_t* cache_idx, int64_t n, + int64_t c_dim, int64_t x_row_stride, + int64_t k, int64_t state_len, bool silu) { + const int64_t width = k - 1; + const int64_t step = static_cast(gridDim.x) * blockDim.x; + for (int64_t idx = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + idx < n; idx += step) { + const int64_t bt = idx / c_dim; + const int64_t c = idx % c_dim; + int64_t srow_off = idx; + if (cache_idx != nullptr) { + const int32_t slot = cache_idx[bt]; + if (slot < 0) continue; + srow_off = static_cast(slot) * c_dim + c; + } + TState* srow = conv_state + srow_off * state_len; + const float xt = Ld(x, bt * x_row_stride + c); + float acc = bias != nullptr ? Ld(bias, c) : 0.0f; + for (int64_t j = 0; j < width; ++j) acc += Ld(w, c * k + j) * Ld(srow, j); + acc += Ld(w, c * k + width) * xt; + St(out, idx, silu ? Silu(acc) : acc); + for (int64_t j = 0; j + 1 < width; ++j) St(srow, j, Ld(srow, j + 1)); + if (width > 0) St(srow, width - 1, xt); + } +} + +template +void LaunchConvFwd(hipStream_t s, Tensor& out, const Tensor& x, const Tensor& w, + const Tensor* bias, Tensor& conv_state, const Tensor& qsl, + const Tensor& his, const CausalConv1dArgs& args) { + const int64_t c = x.shape[1], k = w.shape[1]; + const int64_t x_rs = x.stride[0]; + const int32_t* bp = args.batch_ptr != nullptr ? args.batch_ptr->Ptr() : nullptr; + const int32_t* tco = args.token_chunk_offset_ptr != nullptr + ? args.token_chunk_offset_ptr->Ptr() + : nullptr; + // exact mode: grid.y is the descriptor program count, not the sequence count. + const int64_t rows = bp != nullptr ? args.batch_ptr->shape[0] : conv_state.shape[0]; + const dim3 grid(static_cast((c + kBlock - 1) / kBlock), + static_cast(rows)); + if (his.dtype == DType::kI8) { + CausalConv1dFwdK<<>>( + out.Ptr(), x.Ptr(), w.Ptr(), + bias != nullptr ? bias->Ptr() : nullptr, conv_state.Ptr(), + qsl.Ptr(), his.Ptr(), c, x_rs, k, args.silu_activation, bp, tco); + } else { + CausalConv1dFwdK<<>>( + out.Ptr(), x.Ptr(), w.Ptr(), + bias != nullptr ? bias->Ptr() : nullptr, conv_state.Ptr(), + qsl.Ptr(), his.Ptr(), c, x_rs, k, args.silu_activation, bp, tco); + } + Check(hipGetLastError(), "causal_conv1d_fwd launch"); +} + +template +void LaunchConvUpdate(hipStream_t s, Tensor& out, const Tensor& x, const Tensor& w, + const Tensor* bias, Tensor& conv_state, const int32_t* cache_idx, + const CausalConv1dArgs& args) { + const int64_t n = x.shape[0] * x.shape[1], c = x.shape[1], k = w.shape[1]; + const int64_t x_rs = x.stride[0]; + const int64_t state_len = conv_state.shape[2]; + CausalConv1dUpdateK<<>>( + out.Ptr(), x.Ptr(), w.Ptr(), + bias != nullptr ? bias->Ptr() : nullptr, conv_state.Ptr(), + cache_idx, n, c, x_rs, k, state_len, args.silu_activation); + Check(hipGetLastError(), "causal_conv1d_update launch"); +} + +template +void LaunchConvUpdateS(hipStream_t s, Tensor& out, const Tensor& x, const Tensor& w, + const Tensor* bias, Tensor& conv_state, const int32_t* cache_idx, + const CausalConv1dArgs& args) { + if (conv_state.dtype == DType::kBF16) + LaunchConvUpdate(s, out, x, w, bias, conv_state, cache_idx, args); + else + LaunchConvUpdate(s, out, x, w, bias, conv_state, cache_idx, args); +} + +} // namespace + +void CausalConv1dFwdKernelRocm(Queue& q, Tensor& out, const Tensor& x, const Tensor& w, + const Tensor* bias, Tensor& conv_state, const Tensor& qsl, + const Tensor& his, const CausalConv1dArgs& args) { + VT_CHECK(x.dtype == DType::kF32 || x.dtype == DType::kBF16, + "rocm causal_conv1d_fwd: unsupported x dtype (f32/bf16 only)"); + VT_CHECK(w.dtype == x.dtype && (bias == nullptr || bias->dtype == x.dtype), + "rocm causal_conv1d_fwd: weight/bias dtype must match x"); + const int64_t n = conv_state.shape[0]; + if (n == 0 || x.shape[1] == 0) return; + hipStream_t s = AsStream(q); + if (x.dtype == DType::kF32) { + if (out.dtype == DType::kF32) { + LaunchConvFwd(s, out, x, w, bias, conv_state, qsl, his, args); + } else { + LaunchConvFwd(s, out, x, w, bias, conv_state, qsl, his, args); + } + } else { + if (out.dtype == DType::kF32) { + LaunchConvFwd<__hip_bfloat16, float>(s, out, x, w, bias, conv_state, qsl, his, args); + } else { + LaunchConvFwd<__hip_bfloat16, __hip_bfloat16>(s, out, x, w, bias, conv_state, qsl, his, args); + } + } +} + +void CausalConv1dUpdateKernelRocm(Queue& q, Tensor& out, const Tensor& x, const Tensor& w, + const Tensor* bias, Tensor& conv_state, + const Tensor* conv_state_indices, + const CausalConv1dArgs& args) { + VT_CHECK(x.dtype == DType::kF32 || x.dtype == DType::kBF16, + "rocm causal_conv1d_update: unsupported x dtype (f32/bf16 only)"); + VT_CHECK(w.dtype == x.dtype && (bias == nullptr || bias->dtype == x.dtype), + "rocm causal_conv1d_update: weight/bias dtype must match x"); + if (x.shape[0] * x.shape[1] == 0) return; + const int32_t* ci = + conv_state_indices != nullptr ? conv_state_indices->Ptr() : nullptr; + hipStream_t s = AsStream(q); + if (x.dtype == DType::kF32) { + if (out.dtype == DType::kF32) { + LaunchConvUpdateS(s, out, x, w, bias, conv_state, ci, args); + } else { + LaunchConvUpdateS(s, out, x, w, bias, conv_state, ci, args); + } + } else { + if (out.dtype == DType::kF32) { + LaunchConvUpdateS<__hip_bfloat16, float>(s, out, x, w, bias, conv_state, ci, args); + } else { + LaunchConvUpdateS<__hip_bfloat16, __hip_bfloat16>(s, out, x, w, bias, conv_state, ci, args); + } + } +} + +} // namespace vt::rocm diff --git a/src/vt/rocm/rocm_ops.hip b/src/vt/rocm/rocm_ops.hip index c9b4dd173..ffb649e5a 100644 --- a/src/vt/rocm/rocm_ops.hip +++ b/src/vt/rocm/rocm_ops.hip @@ -44,6 +44,15 @@ void GdnStateGatherKernelRocm(Queue& q, Tensor& working, const Tensor& cache, const Tensor& state_idx, const Tensor* has_initial_state); void GdnStateScatterKernelRocm(Queue& q, Tensor& cache, const Tensor& working, const Tensor& state_idx); +// BACKEND-ROCM-GDN-KERNELS family 2 (rocm_gdn_conv.hip): causal conv1d prefill +// fwd + decode update (gdn-semantics §2/§3). +void CausalConv1dFwdKernelRocm(Queue& q, Tensor& out, const Tensor& x, const Tensor& w, + const Tensor* bias, Tensor& conv_state, const Tensor& qsl, + const Tensor& his, const CausalConv1dArgs& args); +void CausalConv1dUpdateKernelRocm(Queue& q, Tensor& out, const Tensor& x, const Tensor& w, + const Tensor* bias, Tensor& conv_state, + const Tensor* conv_state_indices, + const CausalConv1dArgs& args); namespace { @@ -107,6 +116,12 @@ struct Registrar { RegisterOp(OpId::kGdnStateScatter, DeviceType::kROCM, reinterpret_cast( static_cast(&GdnStateScatterKernelRocm))); + RegisterOp(OpId::kCausalConv1dFwd, DeviceType::kROCM, + reinterpret_cast( + static_cast(&CausalConv1dFwdKernelRocm))); + RegisterOp(OpId::kCausalConv1dUpdate, DeviceType::kROCM, + reinterpret_cast( + static_cast(&CausalConv1dUpdateKernelRocm))); } } registrar; diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index ae90d5abb..88199413c 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -1233,6 +1233,148 @@ TEST_CASE("GDN state gather/scatter are BIT-EXACT against the CPU oracle") { } +TEST_CASE("causal conv1d fwd/update match the CPU oracle") { + // §2/§3 of gdn-semantics.md. Outputs are arithmetic (silu epilogue) -> NMSE; + // the conv_state write-back/roll moves RAW x values -> bit-exact. + const int64_t C = 24, K = 4, W = K - 1; + const std::vector qsl = {0, 5, 6, 15}; // 3 seqs: lens 5, 1, 9 + const std::vector has = {1, 0, 1}; + const int64_t T = 15, N = 3; + const size_t xn = static_cast(T * C), wn = static_cast(C * K); + const std::vector x = RandomVec(xn, 811); + const std::vector w = RandomVec(wn, 812, -0.5f, 0.5f); + const std::vector bias = RandomVec(static_cast(C), 813, -0.2f, 0.2f); + const std::vector st0 = RandomVec(static_cast(N * C * W), 814, -0.5f, 0.5f); + + // CPU oracle (fwd). + std::vector ref_out(xn, 0.0f), ref_state = st0; + { + vt::Backend& cpu = vt::GetBackend(DeviceType::kCPU); + Queue cq = cpu.CreateQueue(); + const Device cd{DeviceType::kCPU, 0}; + std::vector cx = x, cw = w, cb = bias; + std::vector cqsl = qsl, chas = has; + Tensor tx = T2(cx.data(), cd, T, C); + Tensor tw = T2(cw.data(), cd, C, K); + Tensor tb = T1(cb.data(), cd, C); + Tensor tst = Tensor::Contiguous(ref_state.data(), DType::kF32, cd, {N, C, W}); + Tensor tqsl = TI32(cqsl.data(), cd, N + 1); + Tensor this_ = TI32(chas.data(), cd, N); + Tensor tout = T2(ref_out.data(), cd, T, C); + vt::CausalConv1dFwd(cq, tout, tx, tw, &tb, tst, tqsl, this_, vt::CausalConv1dArgs{}); + cpu.DestroyQueue(cq); + } + + for (DeviceType dt : RegisteredDevices()) { + if (!OpAvailable(vt::OpId::kCausalConv1dFwd, dt)) continue; + CAPTURE(DeviceName(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + const Device d{dt, 0}; + DevBuf dx(dev, q, xn), dw(dev, q, wn), db(dev, q, static_cast(C)), + dout(dev, q, xn), dst(dev, q, static_cast(N * C * W)); + DevBufI32 dqsl(dev, q, N + 1), dhas(dev, q, N); + dx.Upload(x); + dw.Upload(w); + db.Upload(bias); + dst.Upload(st0); + dqsl.Upload(qsl); + dhas.Upload(has); + Tensor tx = T2(dx.ptr(), d, T, C); + Tensor tw = T2(dw.ptr(), d, C, K); + Tensor tb = T1(db.ptr(), d, C); + Tensor tst = Tensor::Contiguous(dst.ptr(), DType::kF32, d, {N, C, W}); + Tensor tqsl = TI32(dqsl.ptr(), d, N + 1); + Tensor this_ = TI32(dhas.ptr(), d, N); + Tensor tout = T2(dout.ptr(), d, T, C); + vt::CausalConv1dFwd(q, tout, tx, tw, &tb, tst, tqsl, this_, vt::CausalConv1dArgs{}); + CHECK(Nmse(ref_out, dout.Download()) <= kNmseTol); + CHECK(dst.Download() == ref_state); // raw-x write-back: bit-exact + + // Exact-chunks descriptor form (VT_CONV_EXACT_CHUNKS — the shape Qwen3.5 + // prefill actually passes on the live path): one program per (sequence, + // 8-token chunk). lens 5,1,9 -> programs (s0,c0), (s1,c0), (s2,c0), + // (s2,c1). Must equal the same CPU oracle (the CPU keeps the scalar + // mapping; the descriptors only re-slice the work). + const std::vector batch_ptr = {0, 1, 2, 2}; + const std::vector chunk_off = {0, 0, 0, 1}; + DevBufI32 dbp(dev, q, batch_ptr.size()), dtco(dev, q, chunk_off.size()); + dbp.Upload(batch_ptr); + dtco.Upload(chunk_off); + dst.Upload(st0); // reset state for the descriptor run + Tensor tbp = TI32(dbp.ptr(), d, static_cast(batch_ptr.size())); + Tensor ttco = TI32(dtco.ptr(), d, static_cast(chunk_off.size())); + vt::CausalConv1dArgs exact_args; + exact_args.batch_ptr = &tbp; + exact_args.token_chunk_offset_ptr = &ttco; + vt::CausalConv1dFwd(q, tout, tx, tw, &tb, tst, tqsl, this_, exact_args); + CHECK(Nmse(ref_out, dout.Download()) <= kNmseTol); + CHECK(dst.Download() == ref_state); // descriptor write-back: bit-exact + dev.DestroyQueue(q); + } + + // Update: B=4 tokens. Compact arm: conv_state [B,C,W] (one row per token, + // per the ops.cpp contract). Indexed arm: the full [SLOTS,C,W] cache with one + // NULL slot (-1 -> out row untouched). + const int64_t B = 4, SLOTS = 6; + const size_t un = static_cast(B * C); + const std::vector cidx = {3, -1, 0, 5}; + const std::vector ux = RandomVec(un, 821); + const std::vector ust0 = RandomVec(static_cast(SLOTS * C * W), 822, -0.5f, 0.5f); + for (bool indexed : {false, true}) { + CAPTURE(indexed); + const int64_t st_rows = indexed ? SLOTS : B; + const std::vector ust_arm(ust0.begin(), ust0.begin() + st_rows * C * W); + // CPU oracle. + std::vector ref_uout(un, -2.0f), ref_ust = ust_arm; + { + vt::Backend& cpu = vt::GetBackend(DeviceType::kCPU); + Queue cq = cpu.CreateQueue(); + const Device cd{DeviceType::kCPU, 0}; + std::vector cx = ux, cw = w, cb = bias; + std::vector cci = cidx; + Tensor tx = T2(cx.data(), cd, B, C); + Tensor tw = T2(cw.data(), cd, C, K); + Tensor tb = T1(cb.data(), cd, C); + Tensor tst = Tensor::Contiguous(ref_ust.data(), DType::kF32, cd, {st_rows, C, W}); + Tensor tci = TI32(cci.data(), cd, B); + Tensor tout = T2(ref_uout.data(), cd, B, C); + vt::CausalConv1dUpdate(cq, tout, tx, tw, &tb, tst, vt::CausalConv1dArgs{}, + indexed ? &tci : nullptr); + cpu.DestroyQueue(cq); + } + for (DeviceType dt : RegisteredDevices()) { + if (!OpAvailable(vt::OpId::kCausalConv1dUpdate, dt)) continue; + CAPTURE(DeviceName(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + const Device d{dt, 0}; + DevBuf dx(dev, q, un), dw(dev, q, wn), db(dev, q, static_cast(C)), + dout(dev, q, un), dst(dev, q, static_cast(st_rows * C * W)); + DevBufI32 dci(dev, q, B); + dx.Upload(ux); + dw.Upload(w); + db.Upload(bias); + dst.Upload(ust_arm); + dci.Upload(cidx); + // Untouched-row sentinel must match the oracle's initial -2 fill. + dout.Upload(std::vector(un, -2.0f)); + Tensor tx = T2(dx.ptr(), d, B, C); + Tensor tw = T2(dw.ptr(), d, C, K); + Tensor tb = T1(db.ptr(), d, C); + Tensor tst = Tensor::Contiguous(dst.ptr(), DType::kF32, d, {st_rows, C, W}); + Tensor tci = TI32(dci.ptr(), d, B); + Tensor tout = T2(dout.ptr(), d, B, C); + vt::CausalConv1dUpdate(q, tout, tx, tw, &tb, tst, vt::CausalConv1dArgs{}, + indexed ? &tci : nullptr); + CHECK(Nmse(ref_uout, dout.Download()) <= kNmseTol); + CHECK(dst.Download() == ref_ust); // roll: bit-exact + dev.DestroyQueue(q); + } + } +} + + TEST_CASE("reference tier: an op with no native kernel matches the CPU oracle (unified only)") { constexpr int64_t kRows = 7, kCols = 48; constexpr size_t kN = kRows * kCols; From 8ce0fcc8b4217cc91a99a980b812e7b229f7fab0 Mon Sep 17 00:00:00 2001 From: Vikash Loomba Date: Mon, 10 Aug 2026 19:53:45 -0700 Subject: [PATCH 4/7] =?UTF-8?q?feat(rocm):=20fused=20GDN=20post-conv=20glu?= =?UTF-8?q?e=20(kGdnPostConv)=20=E2=80=94=20family=203=20of=20the=20gfx110?= =?UTF-8?q?0=20GDN=20slice=20(#41)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The VT_GLUE_FUSE default path (conv-split + q/k l2norm + g/beta in one launch), hand-translated from cuda_gdn.cu:1155. Sequential per (token, head) threads matching the CPU oracle's reduction order; the donor's block-tree reduction and Fast/Split variants are M5 levers. Reordered to throw-order: this is the op the M2 rerun hit after the conv family landed. Evidence (4x gfx1100, ROCm 7.14, Release): - postconv cross-device case: 16/16 assertions, runs not skips (padded a/b row strides; q/k/v/g/beta all compared) - ctest -R 'rocm|cross_device': 3/3 - full ctest: same 7 pre-existing failures as base (zero delta) - M2 rerun (Qwen3.5-0.8B): op 65 resolves vt-native; the throw advances 65 -> op 10 (kGdnPrefill, the recurrence — family 4) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: pi:kimi-k3 [pi] --- CMakeLists.txt | 2 + docs/FEATURES.md | 2 +- docs/USAGE.md | 5 +- src/vt/rocm/rocm_gdn_postconv.hip | 154 +++++++++++++++++++++++++ src/vt/rocm/rocm_ops.hip | 9 ++ tests/vt/test_backend_cross_device.cpp | 86 ++++++++++++++ 6 files changed, 255 insertions(+), 3 deletions(-) create mode 100644 src/vt/rocm/rocm_gdn_postconv.hip diff --git a/CMakeLists.txt b/CMakeLists.txt index dc4a580dd..9481ee57e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1314,6 +1314,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_moe_router.hip src/vt/rocm/rocm_gdn_state.hip src/vt/rocm/rocm_gdn_conv.hip + src/vt/rocm/rocm_gdn_postconv.hip src/vt/rocm/rocm_ops.hip) if(VLLM_CPP_HIP_ARCHITECTURES) set_source_files_properties( @@ -1330,6 +1331,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_moe_router.hip src/vt/rocm/rocm_gdn_state.hip src/vt/rocm/rocm_gdn_conv.hip + src/vt/rocm/rocm_gdn_postconv.hip src/vt/rocm/rocm_ops.hip PROPERTIES HIP_ARCHITECTURES "${VLLM_CPP_HIP_ARCHITECTURES}") endif() diff --git a/docs/FEATURES.md b/docs/FEATURES.md index f2c2260bc..628e3706f 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -297,7 +297,7 @@ CPU elementwise GEMM (f32/f16/bf16) runs AVX2 and AVX-512 tiers on x86 where the | LoRA end to end | CPU brick landed | Unwired standalone; not usable through the server | | Multimodal over HTTP | Image request path wired; forward + codec pending | `ROAD-V1-MM` W1-W3 landed (`server_main.cpp:826`). Open: no mm-forward consuming `Request.mm_features`; no image codec vendored (raw RGB only); video/audio/multi-image not started | | Reranking / classify models | Engine side only | Embeddings are LIVE (`LlamaModel`, `vllm_embed`, `/v1/embeddings`); the classify/score heads are landed ops with no registered arch | -| ROCm | W0-W1 community-verified on 5 gfx archs; classic-dense e2e runs all-native (near-tie regime, #269); GDN-hybrid blocked on remaining GDN kernels | Backend + platform + #140 op set + GDN state-I/O and causal-conv1d ops, ctest-green on gfx1151/1103/1100/1201/1200 ([#41](https://github.com/mudler/vllm.cpp/issues/41)). APU UnifiedMemory fix verified. [ROCM.md](ROCM.md) | +| ROCm | W0-W1 community-verified on 5 gfx archs; classic-dense e2e runs all-native (near-tie regime, #269); GDN-hybrid blocked on remaining GDN kernels | Backend + platform + #140 op set + GDN state-I/O + conv1d + post-conv, ctest-green on gfx1151/1103/1100/1201/1200 ([#41](https://github.com/mudler/vllm.cpp/issues/41)). APU UnifiedMemory fix verified. [ROCM.md](ROCM.md) | | XPU, TPU | Not started | CUDA, CPU, Metal and Vulkan are the built backends | | Custom logits processors on CUDA | Open, not root-caused | Segfaults in a CUDA build, 232/232 green on CPU | | Memory budgeting (`ROAD-V1-MEM`, #83) | M1+M2 landed (absolute bytes) | `--kv-cache-memory` sizes the KV pool from an absolute byte budget (ABI v16, group-aware divisor); `--num-blocks` overrides; `--gpu-memory-utilization` needs the M3 profile run (dgx-gated). See `specs/kv-sizing.md` | diff --git a/docs/USAGE.md b/docs/USAGE.md index cc1562978..ae9577dab 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -36,9 +36,10 @@ or putting your own `-O` in `CMAKE_HIP_FLAGS`, overrides it. The ROCm backend registers native ops family by family ([#41](https://github.com/mudler/vllm.cpp/issues/41)); landed GDN slices so far: -the indexed state I/O pair (`kGdnStateGather`/`kGdnStateScatter`) and the causal +the indexed state I/O pair (`kGdnStateGather`/`kGdnStateScatter`), the causal conv1d pair (`kCausalConv1dFwd`/`kCausalConv1dUpdate`, incl. the exact-chunks -descriptor form Qwen3.5 prefill passes). On a +descriptor form Qwen3.5 prefill passes), and the fused post-conv glue +(`kGdnPostConv`). On a discrete card there is no CPU fallback tier, so a model whose layers call an op that is not registered yet fails loudly with `vt: no kernel for op N on device type 5` — that is the memory-safety design working, not a crash. Run with diff --git a/src/vt/rocm/rocm_gdn_postconv.hip b/src/vt/rocm/rocm_gdn_postconv.hip new file mode 100644 index 000000000..041eb4d77 --- /dev/null +++ b/src/vt/rocm/rocm_gdn_postconv.hip @@ -0,0 +1,154 @@ +// ROCm GDN post-conv glue (BACKEND-ROCM-GDN-KERNELS family 3; issue #41). +// kGdnPostConv: conv-split + q/k l2norm + g/beta in ONE launch (mirror of fla +// fused_gdn_prefill_post_conv; gdn-semantics.md §4/§6). Hand-translated from +// cuda_gdn.cu:1155, readable side by side against the donor; CPU oracle +// src/vt/cpu/cpu_ops.cpp:2337. This is the fused VT_GLUE_FUSE path the model +// calls by default — the M2 throw after conv lands. Sequential per +// (token, head) thread here, matching the CPU oracle's reduction order; the +// donor's block-tree reduction and the Fast/Split variants are M5 levers. +// +// Gate: the postconv arm of the fused-glue case in +// test_backend_cross_device.cpp — NMSE <= 5e-4 vs the CPU oracle, padded a/b +// row strides included. + +#include +#include +#include + +#include +#include +#include + +#include "vt/ops.h" + +namespace vt::rocm { +namespace { + +constexpr int kBlock = 256; + +inline void Check(hipError_t err, const char* what) { + if (err != hipSuccess) { + throw std::runtime_error(std::string("vt rocm gdn: ") + what + ": " + + hipGetErrorString(err)); + } +} +inline hipStream_t AsStream(const Queue& q) { + return static_cast(q.handle); +} +inline unsigned GridFor(int64_t n) { + if (n <= 0) return 1; + const int64_t g = (n + kBlock - 1) / kBlock; + return static_cast(g > 65535 ? 65535 : g); +} + +__device__ inline float Ld(const float* p, int64_t i) { return p[i]; } +__device__ inline float Ld(const __half* p, int64_t i) { return __half2float(p[i]); } +__device__ inline float Ld(const __hip_bfloat16* p, int64_t i) { + return __bfloat162float(p[i]); +} +__device__ inline void St(float* p, int64_t i, float v) { p[i] = v; } +__device__ inline void St(__half* p, int64_t i, float v) { p[i] = __float2half_rn(v); } +__device__ inline void St(__hip_bfloat16* p, int64_t i, float v) { + p[i] = __float2bfloat16(v); +} + +// ── Fused glue (donors cuda_gdn.cu:1155; cpu oracle :956 for the preamble) ── +template +__global__ void GdnPostConvK(Tqkv* q_out, Tqkv* k_out, Tqkv* v_out, float* g_out, + float* beta_out, const Tconv* conv, const Tgate* araw, + const Tgate* braw, const float* a_log, + const float* dt_bias, int64_t t, int64_t hk, int64_t dk, + int64_t hv, int64_t dv, int64_t a_row_stride, + int64_t b_row_stride, float eps) { + const int64_t key_dim = hk * dk, value_dim = hv * dv; + const int64_t conv_dim = 2 * key_dim + value_dim; + const int64_t items = t * (hk + 1); + const int64_t step = static_cast(gridDim.x) * blockDim.x; + for (int64_t item = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + item < items; item += step) { + const int64_t tok = item / (hk + 1); + const int64_t head = item % (hk + 1); + const int64_t row = tok * conv_dim; + if (head < hk) { + const Tconv* qin = conv + row + head * dk; + const Tconv* kin = conv + row + key_dim + head * dk; + float qss = 0.0f, kss = 0.0f; + for (int64_t j = 0; j < dk; ++j) { + const float qv = Ld(qin, j); + const float kv = Ld(kin, j); + qss += qv * qv; + kss += kv * kv; + } + const float qinv = 1.0f / sqrtf(qss + eps); + const float kinv = 1.0f / sqrtf(kss + eps); + Tqkv* qo = q_out + (tok * hk + head) * dk; + Tqkv* ko = k_out + (tok * hk + head) * dk; + for (int64_t j = 0; j < dk; ++j) { + St(qo, j, Ld(qin, j) * qinv); + St(ko, j, Ld(kin, j) * kinv); + } + } else { + const Tconv* vin = conv + row + 2 * key_dim; + Tqkv* vo = v_out + tok * value_dim; + for (int64_t j = 0; j < value_dim; ++j) St(vo, j, Ld(vin, j)); + for (int64_t h = 0; h < hv; ++h) { + const int64_t idx = tok * hv + h; + const float av = Ld(araw, tok * a_row_stride + h); + const float bv = Ld(braw, tok * b_row_stride + h); + const float x = av + dt_bias[h]; + const float sp = x > 20.0f ? x : log1pf(expf(x)); + g_out[idx] = -expf(a_log[h]) * sp; + beta_out[idx] = 1.0f / (1.0f + expf(-bv)); + } + } + } +} + +} // namespace + +void GdnPostConvKernelRocm(Queue& q, Tensor& q_out, Tensor& k_out, Tensor& v_out, + Tensor& g_out, Tensor& beta_out, const Tensor& conv, + const Tensor& araw, const Tensor& braw, const Tensor& a_log, + const Tensor& dt_bias, const L2NormArgs& args) { + const int64_t t = conv.shape[0]; + const int64_t hk = q_out.shape[1], dk = q_out.shape[2]; + const int64_t hv = v_out.shape[1], dv = v_out.shape[2]; + if (t == 0) return; + VT_CHECK(q_out.dtype == k_out.dtype && q_out.dtype == v_out.dtype, + "rocm gdn_post_conv: q/k/v out dtypes must match"); + VT_CHECK(q_out.dtype == DType::kF32 || q_out.dtype == DType::kBF16, + "rocm gdn_post_conv: q/k/v out must be f32 or bf16"); + VT_CHECK(conv.dtype == DType::kF32 || conv.dtype == DType::kBF16, + "rocm gdn_post_conv: conv must be f32 or bf16"); + const int64_t items = t * (hk + 1); + hipStream_t s = AsStream(q); + auto launch = [&](auto qkv_tag, auto conv_tag, auto gate_tag) { + using Tqkv = decltype(qkv_tag); + using Tconv = decltype(conv_tag); + using Tgate = decltype(gate_tag); + GdnPostConvK<<>>( + q_out.Ptr(), k_out.Ptr(), v_out.Ptr(), g_out.Ptr(), + beta_out.Ptr(), conv.Ptr(), araw.Ptr(), braw.Ptr(), + a_log.Ptr(), dt_bias.Ptr(), t, hk, dk, hv, dv, araw.stride[0], + braw.stride[0], args.eps); + }; + auto launch_gate = [&](auto qkv_tag, auto conv_tag) { + if (araw.dtype == DType::kBF16) + launch(qkv_tag, conv_tag, __hip_bfloat16{}); + else + launch(qkv_tag, conv_tag, float{}); + }; + auto launch_conv = [&](auto qkv_tag) { + if (conv.dtype == DType::kBF16) + launch_gate(qkv_tag, __hip_bfloat16{}); + else + launch_gate(qkv_tag, float{}); + }; + if (q_out.dtype == DType::kBF16) + launch_conv(__hip_bfloat16{}); + else + launch_conv(float{}); + Check(hipGetLastError(), "gdn_post_conv launch"); +} + +} // namespace vt::rocm diff --git a/src/vt/rocm/rocm_ops.hip b/src/vt/rocm/rocm_ops.hip index ffb649e5a..c79093a18 100644 --- a/src/vt/rocm/rocm_ops.hip +++ b/src/vt/rocm/rocm_ops.hip @@ -53,6 +53,12 @@ void CausalConv1dUpdateKernelRocm(Queue& q, Tensor& out, const Tensor& x, const const Tensor* bias, Tensor& conv_state, const Tensor* conv_state_indices, const CausalConv1dArgs& args); +// BACKEND-ROCM-GDN-KERNELS family 3 (rocm_gdn_postconv.hip): the fused +// post-conv glue (conv-split + q/k l2norm + g/beta), gdn-semantics §4/§6. +void GdnPostConvKernelRocm(Queue& q, Tensor& q_out, Tensor& k_out, Tensor& v_out, + Tensor& g_out, Tensor& beta_out, const Tensor& conv, + const Tensor& araw, const Tensor& braw, const Tensor& a_log, + const Tensor& dt_bias, const L2NormArgs& args); namespace { @@ -122,6 +128,9 @@ struct Registrar { RegisterOp(OpId::kCausalConv1dUpdate, DeviceType::kROCM, reinterpret_cast( static_cast(&CausalConv1dUpdateKernelRocm))); + RegisterOp(OpId::kGdnPostConv, DeviceType::kROCM, + reinterpret_cast( + static_cast(&GdnPostConvKernelRocm))); } } registrar; diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index 88199413c..cae719410 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -1375,6 +1375,92 @@ TEST_CASE("causal conv1d fwd/update match the CPU oracle") { } +TEST_CASE("GdnPostConv matches the CPU oracle within NMSE <= 5e-4") { + // The fused post-conv glue (the VT_GLUE_FUSE path the model calls by + // default): conv-split + q/k l2norm + g/beta in one launch, with padded a/b + // row strides. All arithmetic: NMSE. + const int64_t T = 4, HK = 2, DK = 16, HV = 4, DV = 24; + const int64_t key_dim = HK * DK, value_dim = HV * DV; + const int64_t conv_dim = 2 * key_dim + value_dim; + const int64_t a_outer = HV + 3, b_outer = HV + 5; + const std::vector conv = RandomVec(static_cast(T * conv_dim), 871, -0.5f, 0.5f); + const std::vector araw = RandomVec(static_cast(T * a_outer), 872, -0.4f, 0.4f); + const std::vector braw = RandomVec(static_cast(T * b_outer), 873, -0.4f, 0.4f); + const std::vector a_log = RandomVec(static_cast(HV), 874, -2.0f, -0.5f); + const std::vector dt_bias = RandomVec(static_cast(HV), 875, -0.1f, 0.1f); + vt::L2NormArgs l2a; + l2a.eps = 1e-6f; + + std::vector ref_q(static_cast(T * key_dim)), + ref_k(static_cast(T * key_dim)), ref_v(static_cast(T * value_dim)), + ref_g(static_cast(T * HV)), ref_b(static_cast(T * HV)); + { + vt::Backend& cpu = vt::GetBackend(DeviceType::kCPU); + Queue cq = cpu.CreateQueue(); + const Device cd{DeviceType::kCPU, 0}; + std::vector cc = conv, ca = araw, cb = braw, cal = a_log, cdt = dt_bias; + // araw/braw reach the op as padded [T, HV] rank-2 views (row stride honored). + auto Make = [](void* p, Device dd, std::initializer_list s) { + return Tensor::Contiguous(p, DType::kF32, dd, s); + }; + Tensor tq = Make(ref_q.data(), cd, {T, HK, DK}); + Tensor tk = Make(ref_k.data(), cd, {T, HK, DK}); + Tensor tv = Make(ref_v.data(), cd, {T, HV, DV}); + Tensor tg = Make(ref_g.data(), cd, {T, HV}); + Tensor tb = Make(ref_b.data(), cd, {T, HV}); + Tensor tc = Make(cc.data(), cd, {T, conv_dim}); + Tensor ta = Make(ca.data(), cd, {T, a_outer}); // logical HV cols, padded row + ta.shape[1] = HV; // view narrows the row; stride[0] stays a_outer + Tensor tb2 = Make(cb.data(), cd, {T, b_outer}); + tb2.shape[1] = HV; + Tensor tal = Make(cal.data(), cd, {HV}); + Tensor tdt = Make(cdt.data(), cd, {HV}); + vt::GdnPostConv(cq, tq, tk, tv, tg, tb, tc, ta, tb2, tal, tdt, l2a); + cpu.DestroyQueue(cq); + } + for (DeviceType dt : RegisteredDevices()) { + if (!OpAvailable(vt::OpId::kGdnPostConv, dt)) continue; + CAPTURE(DeviceName(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + const Device d{dt, 0}; + DevBuf dq(dev, q, ref_q.size()), dk(dev, q, ref_k.size()), dv(dev, q, ref_v.size()), + dg(dev, q, ref_g.size()), db(dev, q, ref_b.size()), dc(dev, q, conv.size()), + da(dev, q, araw.size()), db2(dev, q, braw.size()), dal(dev, q, HV), + ddt(dev, q, HV); + dq.Upload(std::vector(ref_q.size(), 0.0f)); + dc.Upload(conv); + da.Upload(araw); + db2.Upload(braw); + dal.Upload(a_log); + ddt.Upload(dt_bias); + auto Make = [&](void* p, std::initializer_list s) { + return Tensor::Contiguous(p, DType::kF32, d, s); + }; + Tensor tq = Make(dq.ptr(), {T, HK, DK}); + Tensor tk = Make(dk.ptr(), {T, HK, DK}); + Tensor tv = Make(dv.ptr(), {T, HV, DV}); + Tensor tg = Make(dg.ptr(), {T, HV}); + Tensor tb = Make(db.ptr(), {T, HV}); + Tensor tc = Make(dc.ptr(), {T, conv_dim}); + Tensor ta = Make(da.ptr(), {T, a_outer}); + ta.shape[1] = HV; + ta.stride[0] = a_outer; + Tensor tb2 = Make(db2.ptr(), {T, b_outer}); + tb2.shape[1] = HV; + tb2.stride[0] = b_outer; + Tensor tal = Make(dal.ptr(), {HV}); + Tensor tdt = Make(ddt.ptr(), {HV}); + vt::GdnPostConv(q, tq, tk, tv, tg, tb, tc, ta, tb2, tal, tdt, l2a); + CHECK(Nmse(ref_q, dq.Download()) <= kNmseTol); + CHECK(Nmse(ref_k, dk.Download()) <= kNmseTol); + CHECK(Nmse(ref_v, dv.Download()) <= kNmseTol); + CHECK(Nmse(ref_g, dg.Download()) <= kNmseTol); + CHECK(Nmse(ref_b, db.Download()) <= kNmseTol); + dev.DestroyQueue(q); + } +} + TEST_CASE("reference tier: an op with no native kernel matches the CPU oracle (unified only)") { constexpr int64_t kRows = 7, kCols = 48; constexpr size_t kN = kRows * kCols; From b6ab9382823bb6f18dd771e6cec57cd5f7954a78 Mon Sep 17 00:00:00 2001 From: Vikash Loomba Date: Mon, 10 Aug 2026 20:07:44 -0700 Subject: [PATCH 5/7] =?UTF-8?q?feat(rocm):=20GDN=20gated-delta=20recurrenc?= =?UTF-8?q?e=20(kGdnPrefill/kGdnDecode)=20=E2=80=94=20family=204=20of=20th?= =?UTF-8?q?e=20gfx1100=20GDN=20slice=20(#41)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The recurrence core over the portable scan (one kernel serves both ops: qsl != null => varlen prefill, qsl == null => single-step decode with optional state_idx + NULL-block zero-out), hand-translated from the cuda_gdn.cu:1856 donor. The WMMA chunked prefill and fused/packed/Triton-AOT decode machinery is NVIDIA-only and stays unported (docs/ROCM.md §6) — the scan is the correctness baseline; perf variants are M5 levers. Evidence (4x gfx1100, ROCm 7.14, Release): - recurrence cross-device case: 35/35 assertions, runs not skips (prefill two-sequence state carry; decode compact + indexed + NULL-slot zero-out; out AND in-place state compared) - ctest -R 'rocm|cross_device': 3/3 - full ctest: same 7 pre-existing failures as base (zero delta) - M2 rerun (Qwen3.5-0.8B): op 10 resolves vt-native; the throw advances 10 -> op 9 (kRmsNormGated — family 5, the final slice) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: pi:kimi-k3 [pi] --- CMakeLists.txt | 2 + docs/FEATURES.md | 2 +- docs/USAGE.md | 5 +- src/vt/rocm/rocm_gdn_scan.hip | 187 +++++++++++++++++++++++++ src/vt/rocm/rocm_ops.hip | 12 ++ tests/vt/test_backend_cross_device.cpp | 137 ++++++++++++++++++ 6 files changed, 342 insertions(+), 3 deletions(-) create mode 100644 src/vt/rocm/rocm_gdn_scan.hip diff --git a/CMakeLists.txt b/CMakeLists.txt index 9481ee57e..f6c5ad749 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1315,6 +1315,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_gdn_state.hip src/vt/rocm/rocm_gdn_conv.hip src/vt/rocm/rocm_gdn_postconv.hip + src/vt/rocm/rocm_gdn_scan.hip src/vt/rocm/rocm_ops.hip) if(VLLM_CPP_HIP_ARCHITECTURES) set_source_files_properties( @@ -1332,6 +1333,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_gdn_state.hip src/vt/rocm/rocm_gdn_conv.hip src/vt/rocm/rocm_gdn_postconv.hip + src/vt/rocm/rocm_gdn_scan.hip src/vt/rocm/rocm_ops.hip PROPERTIES HIP_ARCHITECTURES "${VLLM_CPP_HIP_ARCHITECTURES}") endif() diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 628e3706f..bcb8e121e 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -297,7 +297,7 @@ CPU elementwise GEMM (f32/f16/bf16) runs AVX2 and AVX-512 tiers on x86 where the | LoRA end to end | CPU brick landed | Unwired standalone; not usable through the server | | Multimodal over HTTP | Image request path wired; forward + codec pending | `ROAD-V1-MM` W1-W3 landed (`server_main.cpp:826`). Open: no mm-forward consuming `Request.mm_features`; no image codec vendored (raw RGB only); video/audio/multi-image not started | | Reranking / classify models | Engine side only | Embeddings are LIVE (`LlamaModel`, `vllm_embed`, `/v1/embeddings`); the classify/score heads are landed ops with no registered arch | -| ROCm | W0-W1 community-verified on 5 gfx archs; classic-dense e2e runs all-native (near-tie regime, #269); GDN-hybrid blocked on remaining GDN kernels | Backend + platform + #140 op set + GDN state-I/O + conv1d + post-conv, ctest-green on gfx1151/1103/1100/1201/1200 ([#41](https://github.com/mudler/vllm.cpp/issues/41)). APU UnifiedMemory fix verified. [ROCM.md](ROCM.md) | +| ROCm | W0-W1 community-verified on 5 gfx archs; classic-dense e2e runs all-native (near-tie regime, #269); GDN-hybrid blocked on remaining GDN kernels | Backend + platform + #140 ops + GDN state/conv/postconv/recurrence, ctest-green gfx1151/1103/1100/1201/1200 ([#41](https://github.com/mudler/vllm.cpp/issues/41)). APU UnifiedMemory fix verified. [ROCM.md](ROCM.md) | | XPU, TPU | Not started | CUDA, CPU, Metal and Vulkan are the built backends | | Custom logits processors on CUDA | Open, not root-caused | Segfaults in a CUDA build, 232/232 green on CPU | | Memory budgeting (`ROAD-V1-MEM`, #83) | M1+M2 landed (absolute bytes) | `--kv-cache-memory` sizes the KV pool from an absolute byte budget (ABI v16, group-aware divisor); `--num-blocks` overrides; `--gpu-memory-utilization` needs the M3 profile run (dgx-gated). See `specs/kv-sizing.md` | diff --git a/docs/USAGE.md b/docs/USAGE.md index ae9577dab..a7890dfa1 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -38,8 +38,9 @@ The ROCm backend registers native ops family by family ([#41](https://github.com/mudler/vllm.cpp/issues/41)); landed GDN slices so far: the indexed state I/O pair (`kGdnStateGather`/`kGdnStateScatter`), the causal conv1d pair (`kCausalConv1dFwd`/`kCausalConv1dUpdate`, incl. the exact-chunks -descriptor form Qwen3.5 prefill passes), and the fused post-conv glue -(`kGdnPostConv`). On a +descriptor form Qwen3.5 prefill passes), the fused post-conv glue +(`kGdnPostConv`), and the gated-delta recurrence (`kGdnPrefill`/`kGdnDecode`, +portable scan). On a discrete card there is no CPU fallback tier, so a model whose layers call an op that is not registered yet fails loudly with `vt: no kernel for op N on device type 5` — that is the memory-safety design working, not a crash. Run with diff --git a/src/vt/rocm/rocm_gdn_scan.hip b/src/vt/rocm/rocm_gdn_scan.hip new file mode 100644 index 000000000..6ffc49c75 --- /dev/null +++ b/src/vt/rocm/rocm_gdn_scan.hip @@ -0,0 +1,187 @@ +// ROCm GDN recurrence core (BACKEND-ROCM-GDN-KERNELS family 4; issue #41). +// kGdnPrefill + kGdnDecode over the PORTABLE scan — hand-translation of +// cuda_gdn.cu GdnScanKernel (:1856) + its launchers, readable side by side +// against the donor. CPU oracle: src/vt/cpu/cpu_ops.cpp:1331/1368 +// (GdnHeadTokenStep); formula reference gdn-semantics.md §7/§8. One kernel +// serves both ops (qsl != null => varlen prefill; qsl == null => single-step +// decode with optional state_idx + NULL-block zero-out). +// +// Deliberately NOT ported (docs/ROCM.md §6): the WMMA chunked prefill +// (`#if __CUDA_ARCH__ >= 800`), the fused/packed/Triton-AOT decode variants — +// NVIDIA-only machinery; the scan is the correctness baseline and the perf +// variants are M5 levers (upstream's ROCm answer, Triton/CK, is a later row). +// +// Gate: the recurrence case in test_backend_cross_device.cpp — NMSE <= 5e-4 +// vs the CPU oracle on out AND in-place state; decode NULL-slot zero-out +// EXACT. + +#include +#include +#include + +#include +#include +#include +#include + +#include "vt/ops.h" + +namespace vt::rocm { +namespace { + +constexpr int kBlock = 256; + +inline void Check(hipError_t err, const char* what) { + if (err != hipSuccess) { + throw std::runtime_error(std::string("vt rocm gdn: ") + what + ": " + + hipGetErrorString(err)); + } +} +inline hipStream_t AsStream(const Queue& q) { + return static_cast(q.handle); +} +inline unsigned GridFor(int64_t n) { + if (n <= 0) return 1; + const int64_t g = (n + kBlock - 1) / kBlock; + return static_cast(g > 65535 ? 65535 : g); +} + +__device__ inline float Ld(const float* p, int64_t i) { return p[i]; } +__device__ inline float Ld(const __half* p, int64_t i) { return __half2float(p[i]); } +__device__ inline float Ld(const __hip_bfloat16* p, int64_t i) { + return __bfloat162float(p[i]); +} +__device__ inline void St(float* p, int64_t i, float v) { p[i] = v; } +__device__ inline void St(__half* p, int64_t i, float v) { p[i] = __float2half_rn(v); } +__device__ inline void St(__hip_bfloat16* p, int64_t i, float v) { + p[i] = __float2bfloat16(v); +} + +// ── Recurrence scan (donor cuda_gdn.cu:1856; serves prefill AND decode) ───── +template +__global__ void GdnScanK(Tout* out, const Tin* q, const Tin* k, const Tin* v, + const float* g, const float* beta, TState* state, + const int32_t* qsl, const int32_t* state_idx, + int64_t state_slots, int64_t hk_n, int64_t dk, + int64_t hv_n, int64_t dv, float scale) { + const int64_t s = blockIdx.y; + const int64_t hv = blockIdx.x; + const int64_t hk = hv / (hv_n / hk_n); + const int64_t state_slot = state_idx != nullptr ? state_idx[s] : s; + if (state_slot < 0 || state_slot >= state_slots) { + const int64_t begin = qsl != nullptr ? qsl[s] : s; + const int64_t end = qsl != nullptr ? qsl[s + 1] : s + 1; + for (int64_t t = begin; t < end; ++t) + for (int64_t vi = threadIdx.x; vi < dv; vi += blockDim.x) + St(out, (t * hv_n + hv) * dv + vi, 0.0f); + return; + } + extern __shared__ float smem[]; // [dk] q' then [dk] k + float* q_sh = smem; + float* k_sh = smem + dk; + TState* s_head = state + (state_slot * hv_n + hv) * dv * dk; + const int64_t begin = qsl != nullptr ? qsl[s] : s; + const int64_t end = qsl != nullptr ? qsl[s + 1] : s + 1; + for (int64_t t = begin; t < end; ++t) { + for (int64_t i = threadIdx.x; i < dk; i += blockDim.x) { + q_sh[i] = Ld(q, (t * hk_n + hk) * dk + i) * scale; + k_sh[i] = Ld(k, (t * hk_n + hk) * dk + i); + } + __syncthreads(); + const float decay = expf(g[t * hv_n + hv]); + const float beta_t = beta[t * hv_n + hv]; + for (int64_t vi = threadIdx.x; vi < dv; vi += blockDim.x) { + TState* s_row = s_head + vi * dk; + float dot = 0.0f; + for (int64_t ki = 0; ki < dk; ++ki) { + const float decayed = Ld(s_row, ki) * decay; + dot += decayed * k_sh[ki]; + } + const float vp = (Ld(v, (t * hv_n + hv) * dv + vi) - dot) * beta_t; + float o = 0.0f; + for (int64_t ki = 0; ki < dk; ++ki) { + const float updated = Ld(s_row, ki) * decay + vp * k_sh[ki]; + St(s_row, ki, updated); + o += updated * q_sh[ki]; + } + St(out, (t * hv_n + hv) * dv + vi, o); + } + __syncthreads(); + } +} + +template +void LaunchGdnScanState(hipStream_t s, Tensor& out, const Tensor& q_in, + const Tensor& k, const Tensor& v, const Tensor& g, + const Tensor& beta, Tensor& state, const int32_t* qsl, + const int32_t* state_idx, int64_t n, + const GdnArgs& args) { + const int64_t hk_n = q_in.shape[1], dk = q_in.shape[2]; + const int64_t hv_n = v.shape[1], dv = v.shape[2]; + const dim3 grid(static_cast(hv_n), static_cast(n)); + const size_t shmem = 2 * static_cast(dk) * sizeof(float); + if (state.dtype == DType::kF16) { + GdnScanK<<>>( + out.Ptr(), q_in.Ptr(), k.Ptr(), v.Ptr(), + g.Ptr(), beta.Ptr(), state.Ptr<__half>(), qsl, state_idx, + state.shape[0], hk_n, dk, hv_n, dv, args.scale); + } else if (state.dtype == DType::kBF16) { + GdnScanK<<>>( + out.Ptr(), q_in.Ptr(), k.Ptr(), v.Ptr(), + g.Ptr(), beta.Ptr(), state.Ptr<__hip_bfloat16>(), qsl, + state_idx, state.shape[0], hk_n, dk, hv_n, dv, args.scale); + } else { + GdnScanK<<>>( + out.Ptr(), q_in.Ptr(), k.Ptr(), v.Ptr(), + g.Ptr(), beta.Ptr(), state.Ptr(), qsl, state_idx, + state.shape[0], hk_n, dk, hv_n, dv, args.scale); + } + Check(hipGetLastError(), "gdn scan launch"); +} + +template +void LaunchGdnScanIn(hipStream_t s, Tensor& out, const Tensor& q_in, + const Tensor& k, const Tensor& v, const Tensor& g, + const Tensor& beta, Tensor& state, const int32_t* qsl, + const int32_t* state_idx, int64_t n, const GdnArgs& args) { + if (out.dtype == DType::kBF16) + LaunchGdnScanState(s, out, q_in, k, v, g, beta, state, qsl, + state_idx, n, args); + else + LaunchGdnScanState(s, out, q_in, k, v, g, beta, state, qsl, + state_idx, n, args); +} + +} // namespace + +void GdnPrefillKernelRocm(Queue& q, Tensor& out, const Tensor& q_in, const Tensor& k, + const Tensor& v, const Tensor& g, const Tensor& beta, + Tensor& state, const Tensor& qsl, const GdnArgs& args) { + const int64_t n = state.shape[0]; + if (n == 0) return; + hipStream_t s = AsStream(q); + if (q_in.dtype == DType::kBF16) { + LaunchGdnScanIn<__hip_bfloat16>(s, out, q_in, k, v, g, beta, state, + qsl.Ptr(), nullptr, n, args); + } else { + LaunchGdnScanIn(s, out, q_in, k, v, g, beta, state, qsl.Ptr(), + nullptr, n, args); + } +} + +void GdnDecodeKernelRocm(Queue& q, Tensor& out, const Tensor& q_in, const Tensor& k, + const Tensor& v, const Tensor& g, const Tensor& beta, + Tensor& state, const Tensor* state_idx, const GdnArgs& args) { + const int64_t n = q_in.shape[0]; + if (n == 0) return; + const int32_t* si = state_idx != nullptr ? state_idx->Ptr() : nullptr; + hipStream_t s = AsStream(q); + if (q_in.dtype == DType::kBF16) { + LaunchGdnScanIn<__hip_bfloat16>(s, out, q_in, k, v, g, beta, state, nullptr, si, n, + args); + } else { + LaunchGdnScanIn(s, out, q_in, k, v, g, beta, state, nullptr, si, n, args); + } +} + +} // namespace vt::rocm diff --git a/src/vt/rocm/rocm_ops.hip b/src/vt/rocm/rocm_ops.hip index c79093a18..e2820c99d 100644 --- a/src/vt/rocm/rocm_ops.hip +++ b/src/vt/rocm/rocm_ops.hip @@ -59,6 +59,14 @@ void GdnPostConvKernelRocm(Queue& q, Tensor& q_out, Tensor& k_out, Tensor& v_out Tensor& g_out, Tensor& beta_out, const Tensor& conv, const Tensor& araw, const Tensor& braw, const Tensor& a_log, const Tensor& dt_bias, const L2NormArgs& args); +// BACKEND-ROCM-GDN-KERNELS family 4 (rocm_gdn_scan.hip): the gated-delta +// recurrence over the portable scan (gdn-semantics §7/§8). +void GdnPrefillKernelRocm(Queue& q, Tensor& out, const Tensor& q_in, const Tensor& k, + const Tensor& v, const Tensor& g, const Tensor& beta, + Tensor& state, const Tensor& qsl, const GdnArgs& args); +void GdnDecodeKernelRocm(Queue& q, Tensor& out, const Tensor& q_in, const Tensor& k, + const Tensor& v, const Tensor& g, const Tensor& beta, + Tensor& state, const Tensor* state_idx, const GdnArgs& args); namespace { @@ -131,6 +139,10 @@ struct Registrar { RegisterOp(OpId::kGdnPostConv, DeviceType::kROCM, reinterpret_cast( static_cast(&GdnPostConvKernelRocm))); + RegisterOp(OpId::kGdnPrefill, DeviceType::kROCM, + reinterpret_cast(static_cast(&GdnPrefillKernelRocm))); + RegisterOp(OpId::kGdnDecode, DeviceType::kROCM, + reinterpret_cast(static_cast(&GdnDecodeKernelRocm))); } } registrar; diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index cae719410..625d920f4 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -1461,6 +1461,143 @@ TEST_CASE("GdnPostConv matches the CPU oracle within NMSE <= 5e-4") { } } +TEST_CASE("GDN prefill/decode recurrence matches the CPU oracle within NMSE <= 5e-4") { + // §7/§8. All f32. NMSE on out AND on the in-place state (the recurrence is + // arithmetic end to end). Decode covers the compact arm, the indexed arm, + // and the NULL-slot zero-out. + const int64_t HK = 2, HV = 4, DK = 16, DV = 24; // HV = ratio*HK + const float scale = 0.25f; + vt::GdnArgs ga; + ga.scale = scale; + + // ---- prefill: two sequences, lens 4 and 1, fresh zero state. + const std::vector qsl = {0, 4, 5}; + const int64_t N = 2, T = 5; + const size_t qkn = static_cast(T * HK * DK), vn = static_cast(T * HV * DV); + const size_t gbn = static_cast(T * HV), stn = static_cast(N * HV * DV * DK); + const std::vector qin = RandomVec(qkn, 851, -0.5f, 0.5f); + const std::vector kin = RandomVec(qkn, 852, -0.5f, 0.5f); + const std::vector vin = RandomVec(vn, 853, -0.5f, 0.5f); + const std::vector gin = RandomVec(gbn, 854, -0.3f, -0.01f); // log-decay < 0 + const std::vector bin = RandomVec(gbn, 855, 0.0f, 0.5f); + + std::vector ref_out(vn, 0.0f), ref_st(stn, 0.0f); + { + vt::Backend& cpu = vt::GetBackend(DeviceType::kCPU); + Queue cq = cpu.CreateQueue(); + const Device cd{DeviceType::kCPU, 0}; + std::vector hq = qin, hk_ = kin, hv_ = vin, hg = gin, hb = bin; + std::vector cqsl = qsl; + Tensor tq = Tensor::Contiguous(hq.data(), DType::kF32, cd, {T, HK, DK}); + Tensor tk = Tensor::Contiguous(hk_.data(), DType::kF32, cd, {T, HK, DK}); + Tensor tv = Tensor::Contiguous(hv_.data(), DType::kF32, cd, {T, HV, DV}); + Tensor tg = T2(hg.data(), cd, T, HV); + Tensor tb = T2(hb.data(), cd, T, HV); + Tensor tst = Tensor::Contiguous(ref_st.data(), DType::kF32, cd, {N, HV, DV, DK}); + Tensor tqsl = TI32(cqsl.data(), cd, N + 1); + Tensor tout = Tensor::Contiguous(ref_out.data(), DType::kF32, cd, {T, HV, DV}); + vt::GdnPrefill(cq, tout, tq, tk, tv, tg, tb, tst, tqsl, ga); + cpu.DestroyQueue(cq); + } + for (DeviceType dt : RegisteredDevices()) { + if (!OpAvailable(vt::OpId::kGdnPrefill, dt)) continue; + CAPTURE(DeviceName(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + const Device d{dt, 0}; + DevBuf dq(dev, q, qkn), dk(dev, q, qkn), dv(dev, q, vn), dg(dev, q, gbn), + db(dev, q, gbn), dout(dev, q, vn), dst(dev, q, stn); + DevBufI32 dqsl(dev, q, N + 1); + dq.Upload(qin); + dk.Upload(kin); + dv.Upload(vin); + dg.Upload(gin); + db.Upload(bin); + dst.Upload(std::vector(stn, 0.0f)); + dqsl.Upload(qsl); + Tensor tq = Tensor::Contiguous(dq.ptr(), DType::kF32, d, {T, HK, DK}); + Tensor tk = Tensor::Contiguous(dk.ptr(), DType::kF32, d, {T, HK, DK}); + Tensor tv = Tensor::Contiguous(dv.ptr(), DType::kF32, d, {T, HV, DV}); + Tensor tg = T2(dg.ptr(), d, T, HV); + Tensor tb = T2(db.ptr(), d, T, HV); + Tensor tst = Tensor::Contiguous(dst.ptr(), DType::kF32, d, {N, HV, DV, DK}); + Tensor tqsl = TI32(dqsl.ptr(), d, N + 1); + Tensor tout = Tensor::Contiguous(dout.ptr(), DType::kF32, d, {T, HV, DV}); + vt::GdnPrefill(q, tout, tq, tk, tv, tg, tb, tst, tqsl, ga); + CHECK(Nmse(ref_out, dout.Download()) <= kNmseTol); + CHECK(Nmse(ref_st, dst.Download()) <= kNmseTol); + dev.DestroyQueue(q); + } + + // ---- decode: B=3 tokens over a 4-slot cache; slot -1 => zero out row, + // state untouched. Compact arm (no indices) alongside. + const int64_t B = 3, SLOTS = 4; + const size_t dqkn = static_cast(B * HK * DK), dvn = static_cast(B * HV * DV); + const size_t dgbn = static_cast(B * HV); + const std::vector sidx = {2, -1, 0}; + const std::vector dq_in = RandomVec(dqkn, 861, -0.5f, 0.5f); + const std::vector dk_in = RandomVec(dqkn, 862, -0.5f, 0.5f); + const std::vector dv_in = RandomVec(dvn, 863, -0.5f, 0.5f); + const std::vector dg_in = RandomVec(dgbn, 864, -0.3f, -0.01f); + const std::vector db_in = RandomVec(dgbn, 865, 0.0f, 0.5f); + const size_t dstn = static_cast(SLOTS * HV * DV * DK); + const std::vector dst0 = RandomVec(dstn, 866, -0.4f, 0.4f); + for (bool indexed : {false, true}) { + CAPTURE(indexed); + const int64_t st_rows = indexed ? SLOTS : B; + const std::vector dst_arm(dst0.begin(), dst0.begin() + st_rows * HV * DV * DK); + std::vector ref_dout(dvn, -7.0f), ref_dst = dst_arm; + { + vt::Backend& cpu = vt::GetBackend(DeviceType::kCPU); + Queue cq = cpu.CreateQueue(); + const Device cd{DeviceType::kCPU, 0}; + std::vector hq = dq_in, hk_ = dk_in, hv_ = dv_in, hg = dg_in, hb = db_in; + std::vector csi = sidx; + Tensor tq = Tensor::Contiguous(hq.data(), DType::kF32, cd, {B, HK, DK}); + Tensor tk = Tensor::Contiguous(hk_.data(), DType::kF32, cd, {B, HK, DK}); + Tensor tv = Tensor::Contiguous(hv_.data(), DType::kF32, cd, {B, HV, DV}); + Tensor tg = T2(hg.data(), cd, B, HV); + Tensor tb = T2(hb.data(), cd, B, HV); + Tensor tst = Tensor::Contiguous(ref_dst.data(), DType::kF32, cd, {st_rows, HV, DV, DK}); + Tensor tsi = TI32(csi.data(), cd, B); + Tensor tout = Tensor::Contiguous(ref_dout.data(), DType::kF32, cd, {B, HV, DV}); + vt::GdnDecode(cq, tout, tq, tk, tv, tg, tb, tst, ga, indexed ? &tsi : nullptr); + cpu.DestroyQueue(cq); + } + for (DeviceType dt : RegisteredDevices()) { + if (!OpAvailable(vt::OpId::kGdnDecode, dt)) continue; + CAPTURE(DeviceName(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + const Device d{dt, 0}; + DevBuf dq(dev, q, dqkn), dk(dev, q, dqkn), dv(dev, q, dvn), dg(dev, q, dgbn), + db(dev, q, dgbn), dout(dev, q, dvn), dst(dev, q, static_cast(st_rows * HV * DV * DK)); + DevBufI32 dsi(dev, q, B); + dq.Upload(dq_in); + dk.Upload(dk_in); + dv.Upload(dv_in); + dg.Upload(dg_in); + db.Upload(db_in); + dst.Upload(dst_arm); + dsi.Upload(sidx); + dout.Upload(std::vector(dvn, -7.0f)); // untouched-row sentinel + Tensor tq = Tensor::Contiguous(dq.ptr(), DType::kF32, d, {B, HK, DK}); + Tensor tk = Tensor::Contiguous(dk.ptr(), DType::kF32, d, {B, HK, DK}); + Tensor tv = Tensor::Contiguous(dv.ptr(), DType::kF32, d, {B, HV, DV}); + Tensor tg = T2(dg.ptr(), d, B, HV); + Tensor tb = T2(db.ptr(), d, B, HV); + Tensor tst = Tensor::Contiguous(dst.ptr(), DType::kF32, d, {st_rows, HV, DV, DK}); + Tensor tsi = TI32(dsi.ptr(), d, B); + Tensor tout = Tensor::Contiguous(dout.ptr(), DType::kF32, d, {B, HV, DV}); + vt::GdnDecode(q, tout, tq, tk, tv, tg, tb, tst, ga, indexed ? &tsi : nullptr); + CHECK(Nmse(ref_dout, dout.Download()) <= kNmseTol); + CHECK(Nmse(ref_dst, dst.Download()) <= kNmseTol); + dev.DestroyQueue(q); + } + } +} + + TEST_CASE("reference tier: an op with no native kernel matches the CPU oracle (unified only)") { constexpr int64_t kRows = 7, kCols = 48; constexpr size_t kN = kRows * kCols; From 1129849c7208d18df068f2115eb9d935896b0f0e Mon Sep 17 00:00:00 2001 From: Vikash Loomba Date: Mon, 10 Aug 2026 20:34:33 -0700 Subject: [PATCH 6/7] =?UTF-8?q?feat(rocm):=20GDN=20norm-gate=20+=20fused?= =?UTF-8?q?=20preamble=20+=20compressed-state=20capabilities=20=E2=80=94?= =?UTF-8?q?=20family=205,=20completing=20the=20gfx1100=20GDN=20slice=20(#4?= =?UTF-8?q?1)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The final three ops Qwen3.5-0.8B calls (kRmsNormGated, kSigmoidGateBf16, kAttnQkNormRopeGate), plus the capability seam the live model path forced: the model allocates the conv state bf16 (vLLM mamba_cache_dtype default) and the validators gated that on a hard-coded kCUDA. Adds Backend::SupportsCompressedGdnState (default false), converts CheckGdnCommon to the capability query (the pattern CheckConvCommon already used), keeps CUDA on exactly its prior branch, and advertises both capabilities on ROCm — the registered conv-update and scan kernels carry the compressed-state arms. Shared-file touches are behavior-preserving for every existing backend. Evidence (4x gfx1100, ROCm 7.14, Release): - the five GDN cross-device cases: 162/162 assertions, all run (none skip) - ctest -R 'rocm|cross_device': 3/3; full ctest: same 7 pre-existing failures as base (zero delta) - M2 E2E COMPLETE: Qwen3.5-0.8B runs end to end on discrete ROCm, every op vt-native, deterministic 3/3. 'The sky is' holds 8/8 CPU parity (11/16 at longer horizon). Open prompts diverge at token 1-3 into list-style tokens; VT_GDN_STATE_BF16=0 shifts but does not close it, so it is not a state- rounding artifact. Strict CPU token parity NOT claimed: consistent with the flat-top-logit near-tie regime joral measured on #269, but the GDN-path characterization (per-layer drift or a K-run oracle) is the named open follow-up, not a settled conclusion. - Named next blocker (MoE-bearing paths): the #140 MoeRouterTopK ROCm kernel admits f32 logits only; test_bench/test_loaded_engine_dense now fail there instead of at op 77. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: pi:kimi-k3 [pi] --- CMakeLists.txt | 2 + docs/FEATURES.md | 2 +- docs/USAGE.md | 10 +- include/vt/backend.h | 7 + src/vt/cuda/cuda_backend.cu | 4 + src/vt/ops.cpp | 13 +- src/vt/rocm/rocm_backend.hip | 8 + src/vt/rocm/rocm_gdn_fused.hip | 267 +++++++++++++++++++++++++ src/vt/rocm/rocm_ops.hip | 20 ++ tests/vt/test_backend_cross_device.cpp | 195 ++++++++++++++++++ 10 files changed, 521 insertions(+), 7 deletions(-) create mode 100644 src/vt/rocm/rocm_gdn_fused.hip diff --git a/CMakeLists.txt b/CMakeLists.txt index f6c5ad749..ed1068435 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1316,6 +1316,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_gdn_conv.hip src/vt/rocm/rocm_gdn_postconv.hip src/vt/rocm/rocm_gdn_scan.hip + src/vt/rocm/rocm_gdn_fused.hip src/vt/rocm/rocm_ops.hip) if(VLLM_CPP_HIP_ARCHITECTURES) set_source_files_properties( @@ -1334,6 +1335,7 @@ if(VLLM_CPP_HIP) src/vt/rocm/rocm_gdn_conv.hip src/vt/rocm/rocm_gdn_postconv.hip src/vt/rocm/rocm_gdn_scan.hip + src/vt/rocm/rocm_gdn_fused.hip src/vt/rocm/rocm_ops.hip PROPERTIES HIP_ARCHITECTURES "${VLLM_CPP_HIP_ARCHITECTURES}") endif() diff --git a/docs/FEATURES.md b/docs/FEATURES.md index bcb8e121e..8868892c1 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -216,7 +216,7 @@ HTTP are not started. | CPU (x86, Arm i8mm; A76 assembly correct/default, llama speed gate open) | ✅ | ◐ | ☐ | ✅ | | Metal (Apple Silicon) | ✅ | ☐ | ☐ | ✅ | | Vulkan | ◐ | ☐ | ☐ | ✅ | -| ROCm | ◐ (W0 community-verified on 4 gfx archs, #41; builds on ROCm 6.x with the CLR hostcall fix, #201/#132; APU unified-memory fix landed, unverified) | ✅ | ✅ | ✅ | +| ROCm | W0-W1 verified on 5 gfx archs; classic-dense AND GDN-hybrid e2e run all-native (strict CPU parity not met: near-tie regime, #269; GDN divergence characterization open) | Backend + platform + #140 ops + full GDN op set; ctest-green gfx1151/1103/1100/1201/1200 ([#41](https://github.com/mudler/vllm.cpp/issues/41)). APU UnifiedMemory fix verified. [ROCM.md](ROCM.md) | ✅ | ✅ | | XPU / TPU | ☐ | ✅ | ◐ | ☐ | | Tenstorrent Blackhole | ◐ `ACTIVE`, OPT-125m STRICT 6/6 e2e; Qwen3-0.6B gate wired with device goldens, full 16x16 rerun pending ([spec](../.agents/specs/tenstorrent-backend.md), `BACKEND-TENSTORRENT`) | ✅ | ☐ | ☐ | diff --git a/docs/USAGE.md b/docs/USAGE.md index a7890dfa1..9f6a226ca 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -39,8 +39,14 @@ The ROCm backend registers native ops family by family the indexed state I/O pair (`kGdnStateGather`/`kGdnStateScatter`), the causal conv1d pair (`kCausalConv1dFwd`/`kCausalConv1dUpdate`, incl. the exact-chunks descriptor form Qwen3.5 prefill passes), the fused post-conv glue -(`kGdnPostConv`), and the gated-delta recurrence (`kGdnPrefill`/`kGdnDecode`, -portable scan). On a +(`kGdnPostConv`), the gated-delta recurrence (`kGdnPrefill`/`kGdnDecode`, +portable scan), and the norm-gate/preamble ops (`kRmsNormGated`, +`kSigmoidGateBf16`, `kAttnQkNormRopeGate`) — the full set Qwen3.5-class +GDN-hybrid models call. Compressed conv/SSM state (bf16, the vLLM +`mamba_cache_dtype` default) is advertised via the +`SupportsCompressedConvState`/`SupportsCompressedGdnState` backend probes. +Known limit on a separate path: the ROCm `MoeRouterTopK` kernel takes f32 +logits only, so MoE-bearing models still throw there. On a discrete card there is no CPU fallback tier, so a model whose layers call an op that is not registered yet fails loudly with `vt: no kernel for op N on device type 5` — that is the memory-safety design working, not a crash. Run with diff --git a/include/vt/backend.h b/include/vt/backend.h index b90c0626b..bdbaafb6f 100644 --- a/include/vt/backend.h +++ b/include/vt/backend.h @@ -178,6 +178,13 @@ class Backend { // device, exactly as SupportsAsyncSampledTokenReadback did for the runner. virtual bool SupportsCompressedConvState() const { return false; } + // The GDN recurrent (SSM) state twin of the conv clause above: f16/bf16 + // [N,Hv,Dv,Dk] state addressed in place by the GdnPrefill/GdnDecode kernels, + // read/written in f32 registers (vLLM's mamba_cache_dtype default is bf16). + // CheckGdnCommon used to spell this as `device == kCUDA`; asking the backend + // keeps the shared op layer device-agnostic. + virtual bool SupportsCompressedGdnState() const { return false; } + // Optional graph/command capture (CUDA Graphs / Metal ICB / Vulkan CB). virtual bool SupportsGraphCapture() const { return false; } virtual void BeginCapture(Queue& q); diff --git a/src/vt/cuda/cuda_backend.cu b/src/vt/cuda/cuda_backend.cu index 8d953d027..75c5f63b1 100644 --- a/src/vt/cuda/cuda_backend.cu +++ b/src/vt/cuda/cuda_backend.cu @@ -116,6 +116,10 @@ class CudaBackend final : public Backend { // the branch it took before. bool SupportsCompressedConvState() const override { return true; } + // CheckGdnCommon's compressed-state clause used to name kCUDA directly; the + // capability query keeps CUDA on exactly the branch it took before. + bool SupportsCompressedGdnState() const override { return true; } + // --- Async-output primitives (ENG-ASYNC-SCHED W3, async_utils.py:12-70) ------ // Page-locked host memory the copy engine DMAs into without a staging bounce // (a pageable destination would force cudaMemcpyAsync to block), plus real diff --git a/src/vt/ops.cpp b/src/vt/ops.cpp index 171ed68e5..9d7a8af84 100644 --- a/src/vt/ops.cpp +++ b/src/vt/ops.cpp @@ -1642,12 +1642,17 @@ void CheckGdnCommon(const Queue& q, const Tensor& out, const Tensor& q_in, const VT_CHECK(g.dtype == DType::kF32 && beta.dtype == DType::kF32, std::string(name) + ": g/beta must be f32 (upstream keeps them f32)"); if (allow_compressed_state) { + // Asking the backend (Backend::SupportsCompressedGdnState) rather than + // naming a device — the same device-agnostic pattern CheckConvCommon + // already uses for the conv state. CUDA answers for its existing kernels; + // ROCm answers for the portable scan's f16/bf16 state arms. + const Backend* gdn_backend = TryGetBackend(q.device.type); VT_CHECK(state.dtype == DType::kF32 || - ((state.dtype == DType::kF16 || - state.dtype == DType::kBF16) && - q.device.type == DeviceType::kCUDA), + ((state.dtype == DType::kF16 || state.dtype == DType::kBF16) && + gdn_backend != nullptr && gdn_backend->SupportsCompressedGdnState()), std::string(name) + - ": state must be f32, or fp16/bf16 on CUDA (in/out, in place; " + ": state must be f32, or fp16/bf16 on a backend whose GDN kernels " + "support a compressed state in place (in/out, in place; " "read/written in f32 registers)"); } else { VT_CHECK(state.dtype == DType::kF32, diff --git a/src/vt/rocm/rocm_backend.hip b/src/vt/rocm/rocm_backend.hip index 10d09686e..ce3e00fbb 100644 --- a/src/vt/rocm/rocm_backend.hip +++ b/src/vt/rocm/rocm_backend.hip @@ -238,6 +238,14 @@ class RocmBackend final : public Backend { // host access is guaranteed by construction. This is the route // XNACK-less RDNA3 APUs take, where ground 1 reads false (issue #41 F6). // Probed, never inferred from the gfx name. + // The GDN conv-update and scan kernels registered for this backend carry + // bf16/f16 state arms (f32 register math, boundary conversion), so both + // compressed-state capabilities answer true. Validated on gfx1100 against + // the f32 oracle before advertisement (the M2 e2e run exercises the bf16 + // mamba_cache_dtype path end to end). + bool SupportsCompressedConvState() const override { return true; } + bool SupportsCompressedGdnState() const override { return true; } + bool UnifiedMemory() const override { return unified_memory_; } private: diff --git a/src/vt/rocm/rocm_gdn_fused.hip b/src/vt/rocm/rocm_gdn_fused.hip new file mode 100644 index 000000000..554912752 --- /dev/null +++ b/src/vt/rocm/rocm_gdn_fused.hip @@ -0,0 +1,267 @@ +// ROCm GDN norm/gate + fused attention preamble (BACKEND-ROCM-GDN-KERNELS +// family 5, the final slice; issue #41). kRmsNormGated (gdn-semantics §5), +// kSigmoidGateBf16 (ops.cpp contract: out bf16, attn f32/bf16, gate f32 — the +// sigmoid input stays unrounded), kAttnQkNormRopeGate (fused full-attention +// preamble: split q|gate + gemma/plain qk-RMSNorm + partial NeoX RoPE-from- +// cache + gate passthrough). Donors: cuda_gdn.cu RmsNormGatedRowKernel and +// cuda_ops.cu:1429 area; CPU oracles cpu_ops.cpp:1210/:2273/:956. +// kSigmoidGateBf16 has NO CUDA registration — the CPU composite semantics are +// the donor of record (the Vulkan lane's only native sibling). +// +// Gate: the norm-gate and preamble arms of test_backend_cross_device.cpp — +// NMSE <= 5e-4 for the norm/preamble arithmetic; SigmoidGateBf16 bit-exact +// (single multiply, same RNE store on both sides). + +#include +#include +#include + +#include +#include +#include + +#include "vt/ops.h" + +namespace vt::rocm { +namespace { + +constexpr int kBlock = 256; + +inline void Check(hipError_t err, const char* what) { + if (err != hipSuccess) { + throw std::runtime_error(std::string("vt rocm gdn: ") + what + ": " + + hipGetErrorString(err)); + } +} +inline hipStream_t AsStream(const Queue& q) { + return static_cast(q.handle); +} +inline unsigned GridFor(int64_t n) { + if (n <= 0) return 1; + const int64_t g = (n + kBlock - 1) / kBlock; + return static_cast(g > 65535 ? 65535 : g); +} + +__device__ inline float Ld(const float* p, int64_t i) { return p[i]; } +__device__ inline float Ld(const __half* p, int64_t i) { return __half2float(p[i]); } +__device__ inline float Ld(const __hip_bfloat16* p, int64_t i) { + return __bfloat162float(p[i]); +} +__device__ inline void St(float* p, int64_t i, float v) { p[i] = v; } +__device__ inline void St(__half* p, int64_t i, float v) { p[i] = __float2half_rn(v); } +__device__ inline void St(__hip_bfloat16* p, int64_t i, float v) { + p[i] = __float2bfloat16(v); +} +__device__ inline float Silu(float x) { return x / (1.0f + expf(-x)); } +__device__ inline float Sigmoid(float x) { return 1.0f / (1.0f + expf(-x)); } +__device__ inline float GemmaNormElem(float v, float inv, float w, bool gemma) { + float wj = w; + if (gemma) wj += 1.0f; + return v * inv * wj; +} + +// ── Norm/gate (donor cuda_gdn.cu RmsNormGatedRowKernel; CPU :1210/:2273) ───── +template +__global__ void RmsNormGatedK(Tout* out, const Tx* x, const Tx* gate, + const Tx* w, int64_t rows, int64_t d, + int64_t gate_group, int64_t gate_outer, float eps, + bool sigmoid_gate) { + const int64_t i = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + if (i >= rows) return; + float sumsq = 0.0f; + for (int64_t j = 0; j < d; ++j) { + const float v = Ld(x, i * d + j); + sumsq += v * v; + } + const float inv = 1.0f / sqrtf(sumsq / static_cast(d) + eps); + const int64_t gbase = (i / gate_group) * gate_outer + (i % gate_group) * d; + for (int64_t j = 0; j < d; ++j) { + const float z = Ld(gate, gbase + j); + const float act = sigmoid_gate ? Sigmoid(z) : Silu(z); + St(out, i * d + j, Ld(x, i * d + j) * inv * Ld(w, j) * act); + } +} + +// ops.cpp contract: out bf16, attn f32-or-bf16, gate f32 (unrounded sigmoid +// input). Tattn only. +template +__global__ void SigmoidGateK(__hip_bfloat16* out, const Tattn* attn, + const float* gate, int64_t n) { + const int64_t step = static_cast(gridDim.x) * blockDim.x; + for (int64_t i = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + i < n; i += step) { + St(out, i, Ld(attn, i) * Sigmoid(gate[i])); + } +} + +template +__global__ void AttnQkNormRopeGateK(Tqk* q_out, Tqk* k_out, Tgate* gate_out, + const Tsrc* qgate, const Tsrc* kf, + const float* q_norm, const float* k_norm, + const float* cos_sin, int64_t t, int64_t hq, + int64_t hkv, int64_t dh, int64_t qgate_stride, + int64_t kf_stride, int rot, float eps, + bool gemma) { + const int64_t half = rot / 2; + const int64_t items = t * (hq + hkv); + const int64_t step = static_cast(gridDim.x) * blockDim.x; + for (int64_t item = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + item < items; item += step) { + const int64_t tok = item / (hq + hkv); + const int64_t h = item % (hq + hkv); + const float* cs = cos_sin + tok * rot; + int64_t src_off, out_off; + if (h < hq) { + src_off = tok * qgate_stride + h * 2 * dh; + out_off = (tok * hq + h) * dh; + for (int64_t j = 0; j < dh; ++j) + St(gate_out, out_off + j, Ld(qgate, src_off + dh + j)); + float ss = 0.0f; + for (int64_t j = 0; j < dh; ++j) { + const float v = Ld(qgate, src_off + j); + ss += v * v; + } + const float inv = 1.0f / sqrtf(ss / static_cast(dh) + eps); + for (int64_t j = 0; j < dh; ++j) { + if (j < half) { + const float ni = GemmaNormElem(Ld(qgate, src_off + j), inv, q_norm[j], gemma); + const float nih = + GemmaNormElem(Ld(qgate, src_off + j + half), inv, q_norm[j + half], gemma); + St(q_out, out_off + j, ni * cs[j] - nih * cs[half + j]); + } else if (j < rot) { + const int64_t i = j - half; + const float ni = GemmaNormElem(Ld(qgate, src_off + i), inv, q_norm[i], gemma); + const float nih = + GemmaNormElem(Ld(qgate, src_off + i + half), inv, q_norm[i + half], gemma); + St(q_out, out_off + j, ni * cs[half + i] + nih * cs[i]); + } else { + St(q_out, out_off + j, + GemmaNormElem(Ld(qgate, src_off + j), inv, q_norm[j], gemma)); + } + } + } else { + const int64_t hk_i = h - hq; + src_off = tok * kf_stride + hk_i * dh; + out_off = (tok * hkv + hk_i) * dh; + float ss = 0.0f; + for (int64_t j = 0; j < dh; ++j) { + const float v = Ld(kf, src_off + j); + ss += v * v; + } + const float inv = 1.0f / sqrtf(ss / static_cast(dh) + eps); + for (int64_t j = 0; j < dh; ++j) { + if (j < half) { + const float ni = GemmaNormElem(Ld(kf, src_off + j), inv, k_norm[j], gemma); + const float nih = + GemmaNormElem(Ld(kf, src_off + j + half), inv, k_norm[j + half], gemma); + St(k_out, out_off + j, ni * cs[j] - nih * cs[half + j]); + } else if (j < rot) { + const int64_t i = j - half; + const float ni = GemmaNormElem(Ld(kf, src_off + i), inv, k_norm[i], gemma); + const float nih = + GemmaNormElem(Ld(kf, src_off + i + half), inv, k_norm[i + half], gemma); + St(k_out, out_off + j, ni * cs[half + i] + nih * cs[i]); + } else { + St(k_out, out_off + j, + GemmaNormElem(Ld(kf, src_off + j), inv, k_norm[j], gemma)); + } + } + } + } +} + +} // namespace + +void RmsNormGatedKernelRocm(Queue& q, Tensor& out, const Tensor& x, const Tensor& gate, + const Tensor& w, const RmsNormGatedArgs& args) { + VT_CHECK(x.dtype == DType::kF32 || x.dtype == DType::kBF16, + "rocm rmsnorm_gated: unsupported input dtype (f32/bf16 only)"); + VT_CHECK(gate.dtype == x.dtype && w.dtype == x.dtype, + "rocm rmsnorm_gated: gate/weight dtype must match x"); + const int64_t d = x.shape[x.rank - 1]; + const int64_t t = d == 0 ? 0 : x.Numel() / d; + if (t == 0 || d == 0) return; + const int64_t gate_group = gate.rank == 3 ? gate.shape[1] : 1; + const int64_t gate_outer = gate.stride[0]; + hipStream_t s = AsStream(q); + const unsigned grid = static_cast(t); + if (x.dtype == DType::kF32) { + if (out.dtype == DType::kF32) { + RmsNormGatedK<<>>( + out.Ptr(), x.Ptr(), gate.Ptr(), w.Ptr(), t, d, + gate_group, gate_outer, args.eps, args.sigmoid_gate); + } else { + RmsNormGatedK<<>>( + out.Ptr<__hip_bfloat16>(), x.Ptr(), gate.Ptr(), w.Ptr(), + t, d, gate_group, gate_outer, args.eps, args.sigmoid_gate); + } + } else { + if (out.dtype == DType::kF32) { + RmsNormGatedK<__hip_bfloat16, float><<>>( + out.Ptr(), x.Ptr<__hip_bfloat16>(), gate.Ptr<__hip_bfloat16>(), + w.Ptr<__hip_bfloat16>(), t, d, gate_group, gate_outer, args.eps, + args.sigmoid_gate); + } else { + RmsNormGatedK<__hip_bfloat16, __hip_bfloat16><<>>( + out.Ptr<__hip_bfloat16>(), x.Ptr<__hip_bfloat16>(), gate.Ptr<__hip_bfloat16>(), + w.Ptr<__hip_bfloat16>(), t, d, gate_group, gate_outer, args.eps, + args.sigmoid_gate); + } + } + Check(hipGetLastError(), "rmsnorm_gated launch"); +} + +void SigmoidGateBf16KernelRocm(Queue& q, Tensor& out, const Tensor& attn, + const Tensor& gate) { + // Contract (ops.cpp:3350): out bf16, attn f32/bf16, gate f32. + VT_CHECK(out.dtype == DType::kBF16, "rocm sigmoid_gate_bf16: out must be bf16"); + VT_CHECK(gate.dtype == DType::kF32, "rocm sigmoid_gate_bf16: gate must be f32"); + const int64_t n = out.Numel(); + if (n == 0) return; + hipStream_t s = AsStream(q); + if (attn.dtype == DType::kBF16) { + SigmoidGateK<<>>(out.Ptr<__hip_bfloat16>(), + attn.Ptr<__hip_bfloat16>(), + gate.Ptr(), n); + } else { + SigmoidGateK<<>>(out.Ptr<__hip_bfloat16>(), + attn.Ptr(), + gate.Ptr(), n); + } + Check(hipGetLastError(), "sigmoid_gate launch"); +} + +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, + const Tensor& cos_sin, const RmsNormArgs& na, + const RopeArgs& ra) { + const int64_t t = q_out.shape[0], hq = q_out.shape[1], dh = q_out.shape[2]; + const int64_t hkv = k_out.shape[1]; + if (t == 0) return; + const int64_t items = t * (hq + hkv); + hipStream_t s = AsStream(q); + auto launch = [&](auto src_tag, auto qk_tag, auto gate_tag) { + using Tsrc = decltype(src_tag); + using Tqk = decltype(qk_tag); + using Tgate = decltype(gate_tag); + AttnQkNormRopeGateK<<>>( + q_out.Ptr(), k_out.Ptr(), gate_out.Ptr(), qgate.Ptr(), + kf.Ptr(), q_norm.Ptr(), k_norm.Ptr(), cos_sin.Ptr(), + t, hq, hkv, dh, qgate.stride[0], kf.stride[0], ra.rotary_dim, na.eps, na.gemma); + }; + // Validated combos (mirror the CUDA lane): (f32 src, f32 out); (bf16 src, + // bf16 out + bf16 gate); (bf16 src, bf16 out + f32 gate). + if (qgate.dtype == DType::kBF16) { + if (gate_out.dtype == DType::kF32) { + launch(__hip_bfloat16{}, __hip_bfloat16{}, float{}); + } else { + launch(__hip_bfloat16{}, __hip_bfloat16{}, __hip_bfloat16{}); + } + } else { + launch(float{}, float{}, float{}); + } + Check(hipGetLastError(), "attn_qk_norm_rope_gate launch"); +} + +} // namespace vt::rocm diff --git a/src/vt/rocm/rocm_ops.hip b/src/vt/rocm/rocm_ops.hip index e2820c99d..48ccc5615 100644 --- a/src/vt/rocm/rocm_ops.hip +++ b/src/vt/rocm/rocm_ops.hip @@ -67,6 +67,17 @@ void GdnPrefillKernelRocm(Queue& q, Tensor& out, const Tensor& q_in, const Tenso void GdnDecodeKernelRocm(Queue& q, Tensor& out, const Tensor& q_in, const Tensor& k, const Tensor& v, const Tensor& g, const Tensor& beta, Tensor& state, const Tensor* state_idx, const GdnArgs& args); +// BACKEND-ROCM-GDN-KERNELS family 5 (rocm_gdn_fused.hip): gated RMSNorm, +// sigmoid gate, and the fused full-attention preamble. +void RmsNormGatedKernelRocm(Queue& q, Tensor& out, const Tensor& x, const Tensor& gate, + const Tensor& w, const RmsNormGatedArgs& args); +void SigmoidGateBf16KernelRocm(Queue& q, Tensor& out, const Tensor& attn, + const Tensor& gate); +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, + const Tensor& cos_sin, const RmsNormArgs& na, + const RopeArgs& ra); namespace { @@ -143,6 +154,15 @@ struct Registrar { reinterpret_cast(static_cast(&GdnPrefillKernelRocm))); RegisterOp(OpId::kGdnDecode, DeviceType::kROCM, reinterpret_cast(static_cast(&GdnDecodeKernelRocm))); + RegisterOp(OpId::kRmsNormGated, DeviceType::kROCM, + reinterpret_cast( + static_cast(&RmsNormGatedKernelRocm))); + RegisterOp(OpId::kSigmoidGateBf16, DeviceType::kROCM, + reinterpret_cast( + static_cast(&SigmoidGateBf16KernelRocm))); + RegisterOp(OpId::kAttnQkNormRopeGate, DeviceType::kROCM, + reinterpret_cast( + static_cast(&AttnQkNormRopeGateKernelRocm))); } } registrar; diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index 625d920f4..1de8dad36 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -1128,6 +1128,15 @@ std::vector Bf16Bits(const std::vector& src) { return out; } +// Rank-3 padded-row view [T, H, D] over a [T, row_stride] f32 buffer — the +// merged-qkvz slice shape the GDN/attention glue ops consume in the model. +Tensor T3PaddedF32(void* p, Device d, int64_t t, int64_t h, int64_t w, + int64_t row_stride) { + Tensor t3 = Tensor::Contiguous(p, DType::kF32, d, {t, h, w}); + t3.stride[0] = row_stride; + return t3; +} + // --- GDN cases (BACKEND-ROCM-GDN-KERNELS) ------------------------------------- @@ -1598,6 +1607,192 @@ TEST_CASE("GDN prefill/decode recurrence matches the CPU oracle within NMSE <= 5 } +TEST_CASE("RmsNormGated and SigmoidGate match the CPU oracle") { + // §5. RmsNormGated: NMSE (rms reduction + gate activation), both gate + // activations, and BOTH gate layouts — contiguous rank-2 and the padded-row + // rank-3 [T,Hv,D] merged-qkvz view. SigmoidGateBf16 is a single multiply + // with an RNE store both sides apply: bit-exact. + const int64_t T = 5, HV = 3, D = 32; + const int64_t rows = T * HV; + const int64_t gate_outer = HV * D + 8; // padded token stride (rank-3 arm) + const size_t xn = static_cast(rows * D); + const std::vector x = RandomVec(xn, 831); + const std::vector gate = RandomVec(static_cast(T * gate_outer), 832); + const std::vector w = RandomVec(static_cast(D), 833, 0.2f, 1.0f); + + for (bool sig : {false, true}) { + for (bool rank3 : {false, true}) { + CAPTURE(sig); + CAPTURE(rank3); + vt::RmsNormGatedArgs args; + args.sigmoid_gate = sig; + // rank3 arm: x/gate/out are [T,Hv,D] (gate padded-row); rank2 arm: all + // [rows,D] contiguous (gate buffer's leading rows*D elements). + std::vector ref(xn, 0.0f); + { + vt::Backend& cpu = vt::GetBackend(DeviceType::kCPU); + Queue cq = cpu.CreateQueue(); + const Device cd{DeviceType::kCPU, 0}; + std::vector cx = x, cg = gate, cw = w; + Tensor tx = rank3 ? Tensor::Contiguous(cx.data(), DType::kF32, cd, {T, HV, D}) + : T2(cx.data(), cd, rows, D); + Tensor tg = rank3 ? T3PaddedF32(cg.data(), cd, T, HV, D, gate_outer) + : T2(cg.data(), cd, rows, D); + Tensor tw = T1(cw.data(), cd, D); + Tensor tout = rank3 ? Tensor::Contiguous(ref.data(), DType::kF32, cd, {T, HV, D}) + : T2(ref.data(), cd, rows, D); + vt::RmsNormGated(cq, tout, tx, tg, tw, args); + cpu.DestroyQueue(cq); + } + for (DeviceType dt : RegisteredDevices()) { + if (!OpAvailable(vt::OpId::kRmsNormGated, dt)) continue; + CAPTURE(DeviceName(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + const Device d{dt, 0}; + DevBuf dx(dev, q, xn), dg(dev, q, gate.size()), dw(dev, q, D), dout(dev, q, xn); + dx.Upload(x); + dg.Upload(gate); + dw.Upload(w); + Tensor tx = rank3 ? Tensor::Contiguous(dx.ptr(), DType::kF32, d, {T, HV, D}) + : T2(dx.ptr(), d, rows, D); + Tensor tg = rank3 ? T3PaddedF32(dg.ptr(), d, T, HV, D, gate_outer) + : T2(dg.ptr(), d, rows, D); + Tensor tw = T1(dw.ptr(), d, D); + Tensor tout = rank3 ? Tensor::Contiguous(dout.ptr(), DType::kF32, d, {T, HV, D}) + : T2(dout.ptr(), d, rows, D); + vt::RmsNormGated(q, tout, tx, tg, tw, args); + CHECK(Nmse(ref, dout.Download()) <= kNmseTol); + dev.DestroyQueue(q); + } + } + } + + // SigmoidGateBf16: out bf16, gate f32, attn bf16 OR f32 (the FA-2 prefill + // combo) — single multiply with the same RNE store on both sides: bit-exact. + const size_t sn = 256; + const std::vector attn_f = RandomVec(sn, 841); + const std::vector gate_f = RandomVec(sn, 842); + const std::vector attn_bf = Bf16Bits(attn_f); + for (bool attn_is_bf16 : {true, false}) { + CAPTURE(attn_is_bf16); + std::vector ref_sg(sn, 0); + { + vt::Backend& cpu = vt::GetBackend(DeviceType::kCPU); + Queue cq = cpu.CreateQueue(); + const Device cd{DeviceType::kCPU, 0}; + std::vector ca = attn_bf; + std::vector caf = attn_f, cg = gate_f; + Tensor ta = attn_is_bf16 + ? Tensor::Contiguous(ca.data(), DType::kBF16, cd, {static_cast(sn)}) + : Tensor::Contiguous(caf.data(), DType::kF32, cd, {static_cast(sn)}); + Tensor tg = T1(cg.data(), cd, static_cast(sn)); + Tensor tout = Tensor::Contiguous(ref_sg.data(), DType::kBF16, cd, {static_cast(sn)}); + vt::SigmoidGateBf16(cq, tout, ta, tg); + cpu.DestroyQueue(cq); + } + for (DeviceType dt : RegisteredDevices()) { + if (!OpAvailable(vt::OpId::kSigmoidGateBf16, dt)) continue; + CAPTURE(DeviceName(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + const Device d{dt, 0}; + DevBufBytes da(dev, q, sn * (attn_is_bf16 ? 2 : 4)); + DevBuf dg(dev, q, sn); + DevBufBytes dout(dev, q, sn * 2); + if (attn_is_bf16) { + da.Upload(attn_bf.data()); + } else { + da.Upload(attn_f.data()); + } + dg.Upload(gate_f); + Tensor ta = Tensor::Contiguous(da.ptr(), attn_is_bf16 ? DType::kBF16 : DType::kF32, d, + {static_cast(sn)}); + Tensor tg = T1(dg.ptr(), d, static_cast(sn)); + Tensor tout = Tensor::Contiguous(dout.ptr(), DType::kBF16, d, {static_cast(sn)}); + vt::SigmoidGateBf16(q, tout, ta, tg); + std::vector got(sn); + dout.Download(got.data()); + CHECK(got == ref_sg); + dev.DestroyQueue(q); + } + } +} + + +TEST_CASE("AttnQkNormRopeGate matches the CPU oracle within NMSE <= 5e-4") { + // Fused full-attention preamble: split q|gate + (gemma) qk-RMSNorm(Dh) + + // partial NeoX RoPE-from-cache + gate passthrough. Padded qgate/kf token + // strides; plain + gemma norm variants. All arithmetic: NMSE except the + // gate passthrough (pure movement). + const int64_t T = 4; + const int64_t HQ = 3, HKV = 2, DH = 32, ROT = 16; + const int64_t qg_outer = HQ * 2 * DH + 7, kf_outer = HKV * DH + 5; + const std::vector qgate = RandomVec(static_cast(T * qg_outer), 881, -0.5f, 0.5f); + const std::vector kf = RandomVec(static_cast(T * kf_outer), 882, -0.5f, 0.5f); + const std::vector qn = RandomVec(static_cast(DH), 883, 0.2f, 1.0f); + const std::vector kn = RandomVec(static_cast(DH), 884, 0.2f, 1.0f); + const std::vector cs = RandomVec(static_cast(T * ROT), 885, -1.0f, 1.0f); + for (bool gemma : {false, true}) { + CAPTURE(gemma); + vt::RmsNormArgs na; + na.eps = 1e-6f; + na.gemma = gemma; + vt::RopeArgs ra; + ra.rotary_dim = ROT; + std::vector ref_qo(static_cast(T * HQ * DH)), + ref_ko(static_cast(T * HKV * DH)), ref_go(static_cast(T * HQ * DH)); + { + vt::Backend& cpu = vt::GetBackend(DeviceType::kCPU); + Queue cq = cpu.CreateQueue(); + const Device cd{DeviceType::kCPU, 0}; + std::vector cqg = qgate, ckf = kf, cqn = qn, ckn = kn, ccs = cs; + Tensor tqg = Tensor::Contiguous(cqg.data(), DType::kF32, cd, {T, HQ * 2 * DH}); + tqg.stride[0] = qg_outer; // padded token rows (merged-projection view) + Tensor tkf = Tensor::Contiguous(ckf.data(), DType::kF32, cd, {T, HKV * DH}); + tkf.stride[0] = kf_outer; + Tensor tqn = T1(cqn.data(), cd, DH); + Tensor tkn = T1(ckn.data(), cd, DH); + Tensor tcs = T2(ccs.data(), cd, T, ROT); + Tensor tqo = Tensor::Contiguous(ref_qo.data(), DType::kF32, cd, {T, HQ, DH}); + Tensor tko = Tensor::Contiguous(ref_ko.data(), DType::kF32, cd, {T, HKV, DH}); + Tensor tgo = Tensor::Contiguous(ref_go.data(), DType::kF32, cd, {T, HQ, DH}); + vt::AttnQkNormRopeGate(cq, tqo, tko, tgo, tqg, tkf, tqn, tkn, tcs, na, ra); + cpu.DestroyQueue(cq); + } + for (DeviceType dt : RegisteredDevices()) { + if (!OpAvailable(vt::OpId::kAttnQkNormRopeGate, dt)) continue; + CAPTURE(DeviceName(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + const Device d{dt, 0}; + DevBuf dqg(dev, q, qgate.size()), dkf(dev, q, kf.size()), dqn(dev, q, DH), + dkn(dev, q, DH), dcs(dev, q, cs.size()), dqo(dev, q, ref_qo.size()), + dko(dev, q, ref_ko.size()), dgo(dev, q, ref_go.size()); + dqg.Upload(qgate); + dkf.Upload(kf); + dqn.Upload(qn); + dkn.Upload(kn); + dcs.Upload(cs); + Tensor tqg = Tensor::Contiguous(dqg.ptr(), DType::kF32, d, {T, HQ * 2 * DH}); + tqg.stride[0] = qg_outer; + Tensor tkf = Tensor::Contiguous(dkf.ptr(), DType::kF32, d, {T, HKV * DH}); + tkf.stride[0] = kf_outer; + Tensor tqn = T1(dqn.ptr(), d, DH); + Tensor tkn = T1(dkn.ptr(), d, DH); + Tensor tcs = T2(dcs.ptr(), d, T, ROT); + Tensor tqo = Tensor::Contiguous(dqo.ptr(), DType::kF32, d, {T, HQ, DH}); + Tensor tko = Tensor::Contiguous(dko.ptr(), DType::kF32, d, {T, HKV, DH}); + Tensor tgo = Tensor::Contiguous(dgo.ptr(), DType::kF32, d, {T, HQ, DH}); + vt::AttnQkNormRopeGate(q, tqo, tko, tgo, tqg, tkf, tqn, tkn, tcs, na, ra); + CHECK(Nmse(ref_qo, dqo.Download()) <= kNmseTol); + CHECK(Nmse(ref_ko, dko.Download()) <= kNmseTol); + CHECK(Nmse(ref_go, dgo.Download()) <= kNmseTol); // gate passthrough: exact movement + dev.DestroyQueue(q); + } + } +} + TEST_CASE("reference tier: an op with no native kernel matches the CPU oracle (unified only)") { constexpr int64_t kRows = 7, kCols = 48; constexpr size_t kN = kRows * kCols; From 07e9cc393cf3701e85262547ec926a389e893f0e Mon Sep 17 00:00:00 2001 From: Vikash Loomba Date: Mon, 10 Aug 2026 22:20:26 -0700 Subject: [PATCH 7/7] test(rocm): AttnQkNormRopeGate at the real Qwen3.5-0.8B dims (Dh=256, rot=64 partial) (#41) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The divergence investigation (per-layer drift table on #41) localized the open-prompt divergence to the full-attention block, so the preamble op got a real-dims arm: Dh=256 with partial_rotary 0.25 (rot=64 — a 192-dim pass-through tail the synthetic 32/16 arm never exercised), Hq=8/Hkv=2, gemma and plain. PASSES 44/44 at NMSE <= 5e-4 CPU-vs-ROCm — op 67 is cleared at the real config; the remaining attention-path suspects (the MRoPE cos_sin cache and paged attention, both pre-GDN-slice ops) are recorded on #41. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: pi:kimi-k3 [pi] --- tests/vt/test_backend_cross_device.cpp | 65 ++++++++++++++++++++++++++ 1 file changed, 65 insertions(+) diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index 1de8dad36..bbb2013a5 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -1726,6 +1726,71 @@ TEST_CASE("AttnQkNormRopeGate matches the CPU oracle within NMSE <= 5e-4") { // strides; plain + gemma norm variants. All arithmetic: NMSE except the // gate passthrough (pure movement). const int64_t T = 4; + // Real Qwen3.5-0.8B attention dims first: Dh=256, rot=64 (partial_rotary + // 0.25), Hq=8, Hkv=2 — the config the model actually runs; the synthetic + // 32/16 arm below does not exercise the 192-dim pass-through tail. + { + const int64_t HQr = 8, HKVr = 2, DHr = 256, ROTr = 64; + const int64_t qgo = HQr * 2 * DHr + 7, kfo = HKVr * DHr + 5; + const std::vector qg = RandomVec(static_cast(T * qgo), 981, -0.5f, 0.5f); + const std::vector kfv = RandomVec(static_cast(T * kfo), 982, -0.5f, 0.5f); + const std::vector qnr = RandomVec(static_cast(DHr), 983, 0.2f, 1.0f); + const std::vector knr = RandomVec(static_cast(DHr), 984, 0.2f, 1.0f); + const std::vector csr = RandomVec(static_cast(T * ROTr), 985, -1.0f, 1.0f); + for (bool gemma : {false, true}) { + CAPTURE(gemma); + vt::RmsNormArgs na2; na2.eps = 1e-6f; na2.gemma = gemma; + vt::RopeArgs ra2; ra2.rotary_dim = static_cast(ROTr); + std::vector rq(static_cast(T * HQr * DHr)); + std::vector rk(static_cast(T * HKVr * DHr)); + std::vector rg(static_cast(T * HQr * DHr)); + { + vt::Backend& cpu = vt::GetBackend(DeviceType::kCPU); + Queue cq = cpu.CreateQueue(); + const Device cd{DeviceType::kCPU, 0}; + std::vector a = qg, b = kfv, e = qnr, f = knr, g = csr; + Tensor tqg = Tensor::Contiguous(a.data(), DType::kF32, cd, {T, HQr * 2 * DHr}); + tqg.stride[0] = qgo; + Tensor tkf = Tensor::Contiguous(b.data(), DType::kF32, cd, {T, HKVr * DHr}); + tkf.stride[0] = kfo; + Tensor tqn = T1(e.data(), cd, DHr); + Tensor tkn = T1(f.data(), cd, DHr); + Tensor tcs = T2(g.data(), cd, T, ROTr); + Tensor tqo = Tensor::Contiguous(rq.data(), DType::kF32, cd, {T, HQr, DHr}); + Tensor tko = Tensor::Contiguous(rk.data(), DType::kF32, cd, {T, HKVr, DHr}); + Tensor tgo = Tensor::Contiguous(rg.data(), DType::kF32, cd, {T, HQr, DHr}); + vt::AttnQkNormRopeGate(cq, tqo, tko, tgo, tqg, tkf, tqn, tkn, tcs, na2, ra2); + cpu.DestroyQueue(cq); + } + for (DeviceType dt : RegisteredDevices()) { + if (!OpAvailable(vt::OpId::kAttnQkNormRopeGate, dt)) continue; + CAPTURE(DeviceName(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + const Device d{dt, 0}; + DevBuf dqg(dev, q, qg.size()), dkf(dev, q, kfv.size()), dqn(dev, q, DHr), + dkn(dev, q, DHr), dcs(dev, q, csr.size()), dqo(dev, q, rq.size()), + dko(dev, q, rk.size()), dgo(dev, q, rg.size()); + dqg.Upload(qg); dkf.Upload(kfv); dqn.Upload(qnr); dkn.Upload(knr); dcs.Upload(csr); + Tensor tqg = Tensor::Contiguous(dqg.ptr(), DType::kF32, d, {T, HQr * 2 * DHr}); + tqg.stride[0] = qgo; + Tensor tkf = Tensor::Contiguous(dkf.ptr(), DType::kF32, d, {T, HKVr * DHr}); + tkf.stride[0] = kfo; + Tensor tqn = T1(dqn.ptr(), d, DHr); + Tensor tkn = T1(dkn.ptr(), d, DHr); + Tensor tcs = T2(dcs.ptr(), d, T, ROTr); + Tensor tqo = Tensor::Contiguous(dqo.ptr(), DType::kF32, d, {T, HQr, DHr}); + Tensor tko = Tensor::Contiguous(dko.ptr(), DType::kF32, d, {T, HKVr, DHr}); + Tensor tgo = Tensor::Contiguous(dgo.ptr(), DType::kF32, d, {T, HQr, DHr}); + vt::AttnQkNormRopeGate(q, tqo, tko, tgo, tqg, tkf, tqn, tkn, tcs, na2, ra2); + CHECK(Nmse(rq, dqo.Download()) <= kNmseTol); + CHECK(Nmse(rk, dko.Download()) <= kNmseTol); + CHECK(Nmse(rg, dgo.Download()) <= kNmseTol); + dev.DestroyQueue(q); + } + } + } + const int64_t HQ = 3, HKV = 2, DH = 32, ROT = 16; const int64_t qg_outer = HQ * 2 * DH + 7, kf_outer = HKV * DH + 5; const std::vector qgate = RandomVec(static_cast(T * qg_outer), 881, -0.5f, 0.5f);