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) diff --git a/CMakeLists.txt b/CMakeLists.txt index 3f0f009d5..9481ee57e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1312,6 +1312,9 @@ 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_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( @@ -1326,6 +1329,9 @@ 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_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 dbd2d1455..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 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 + 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 89e172f76..ae9577dab 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -32,6 +32,19 @@ 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)); 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 +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_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_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_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..c79093a18 100644 --- a/src/vt/rocm/rocm_ops.hip +++ b/src/vt/rocm/rocm_ops.hip @@ -38,6 +38,27 @@ 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); +// 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); +// 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 { @@ -95,6 +116,21 @@ 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))); + RegisterOp(OpId::kCausalConv1dFwd, DeviceType::kROCM, + reinterpret_cast( + static_cast(&CausalConv1dFwdKernelRocm))); + 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 29e4b4abd..cae719410 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -1064,6 +1064,403 @@ 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("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("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;