merge: pr/343 -- GDN gated-delta recurrence (kGdnPrefill/kGdnDecode), family 4 (#343)

rocm_gdn_scan.hip: the gated-delta recurrence itself, the op the whole slice
exists to reach, in both its prefill and decode forms.

Fourth of the stacked gfx1100 GDN slice. New ROCm-only TU under VLLM_CPP_HIP,
additive registrations, additive cross-device cases; no existing backend's path
changes. Conflicts resolved by union (additive vs additive).

Refs #343, #41.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: Claude:claude-opus-5 [ClaudeCode]

Author: mudler

CMakeLists.txt

@@ -1328,6 +1328,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(
@@ -1346,6 +1347,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()

docs/FEATURES.md

@@ -307,7 +307,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` |

docs/USAGE.md

@@ -43,8 +43,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

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 <hip/hip_bf16.h>
+#include <hip/hip_fp16.h>
+#include <hip/hip_runtime.h>
+
+#include <cstdint>
+#include <cstddef>
+#include <stdexcept>
+#include <string>
+
+#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<hipStream_t>(q.handle);
+}
+inline unsigned GridFor(int64_t n) {
+  if (n <= 0) return 1;
+  const int64_t g = (n + kBlock - 1) / kBlock;
+  return static_cast<unsigned>(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 <typename Tin, typename Tout, typename TState>
+__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 <typename Tin, typename Tout>
+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<unsigned>(hv_n), static_cast<unsigned>(n));
+  const size_t shmem = 2 * static_cast<size_t>(dk) * sizeof(float);
+  if (state.dtype == DType::kF16) {
+    GdnScanK<Tin, Tout, __half><<<grid, kBlock, shmem, s>>>(
+        out.Ptr<Tout>(), q_in.Ptr<Tin>(), k.Ptr<Tin>(), v.Ptr<Tin>(),
+        g.Ptr<float>(), beta.Ptr<float>(), state.Ptr<__half>(), qsl, state_idx,
+        state.shape[0], hk_n, dk, hv_n, dv, args.scale);
+  } else if (state.dtype == DType::kBF16) {
+    GdnScanK<Tin, Tout, __hip_bfloat16><<<grid, kBlock, shmem, s>>>(
+        out.Ptr<Tout>(), q_in.Ptr<Tin>(), k.Ptr<Tin>(), v.Ptr<Tin>(),
+        g.Ptr<float>(), beta.Ptr<float>(), state.Ptr<__hip_bfloat16>(), qsl,
+        state_idx, state.shape[0], hk_n, dk, hv_n, dv, args.scale);
+  } else {
+    GdnScanK<Tin, Tout, float><<<grid, kBlock, shmem, s>>>(
+        out.Ptr<Tout>(), q_in.Ptr<Tin>(), k.Ptr<Tin>(), v.Ptr<Tin>(),
+        g.Ptr<float>(), beta.Ptr<float>(), state.Ptr<float>(), qsl, state_idx,
+        state.shape[0], hk_n, dk, hv_n, dv, args.scale);
+  }
+  Check(hipGetLastError(), "gdn scan launch");
+}
+
+template <typename Tin>
+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<Tin, __hip_bfloat16>(s, out, q_in, k, v, g, beta, state, qsl,
+                                            state_idx, n, args);
+  else
+    LaunchGdnScanState<Tin, float>(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<int32_t>(), nullptr, n, args);
+  } else {
+    LaunchGdnScanIn<float>(s, out, q_in, k, v, g, beta, state, qsl.Ptr<int32_t>(),
+                           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<int32_t>() : 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<float>(s, out, q_in, k, v, g, beta, state, nullptr, si, n, args);
+  }
+}
+
+}  // namespace vt::rocm

src/vt/rocm/rocm_ops.hip

@@ -73,6 +73,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 {
 
@@ -173,6 +181,10 @@ struct Registrar {
     RegisterOp(OpId::kGdnPostConv, DeviceType::kROCM,
                reinterpret_cast<void*>(
                    static_cast<GdnPostConvFn>(&GdnPostConvKernelRocm)));
+    RegisterOp(OpId::kGdnPrefill, DeviceType::kROCM,
+               reinterpret_cast<void*>(static_cast<GdnPrefillFn>(&GdnPrefillKernelRocm)));
+    RegisterOp(OpId::kGdnDecode, DeviceType::kROCM,
+               reinterpret_cast<void*>(static_cast<GdnDecodeFn>(&GdnDecodeKernelRocm)));
   }
 } registrar;
 

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<int32_t> qsl = {0, 4, 5};
+  const int64_t N = 2, T = 5;
+  const size_t qkn = static_cast<size_t>(T * HK * DK), vn = static_cast<size_t>(T * HV * DV);
+  const size_t gbn = static_cast<size_t>(T * HV), stn = static_cast<size_t>(N * HV * DV * DK);
+  const std::vector<float> qin = RandomVec(qkn, 851, -0.5f, 0.5f);
+  const std::vector<float> kin = RandomVec(qkn, 852, -0.5f, 0.5f);
+  const std::vector<float> vin = RandomVec(vn, 853, -0.5f, 0.5f);
+  const std::vector<float> gin = RandomVec(gbn, 854, -0.3f, -0.01f);  // log-decay < 0
+  const std::vector<float> bin = RandomVec(gbn, 855, 0.0f, 0.5f);
+
+  std::vector<float> 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<float> hq = qin, hk_ = kin, hv_ = vin, hg = gin, hb = bin;
+    std::vector<int32_t> 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<float>(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<size_t>(B * HK * DK), dvn = static_cast<size_t>(B * HV * DV);
+  const size_t dgbn = static_cast<size_t>(B * HV);
+  const std::vector<int32_t> sidx = {2, -1, 0};
+  const std::vector<float> dq_in = RandomVec(dqkn, 861, -0.5f, 0.5f);
+  const std::vector<float> dk_in = RandomVec(dqkn, 862, -0.5f, 0.5f);
+  const std::vector<float> dv_in = RandomVec(dvn, 863, -0.5f, 0.5f);
+  const std::vector<float> dg_in = RandomVec(dgbn, 864, -0.3f, -0.01f);
+  const std::vector<float> db_in = RandomVec(dgbn, 865, 0.0f, 0.5f);
+  const size_t dstn = static_cast<size_t>(SLOTS * HV * DV * DK);
+  const std::vector<float> 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<float> dst_arm(dst0.begin(), dst0.begin() + st_rows * HV * DV * DK);
+    std::vector<float> 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<float> hq = dq_in, hk_ = dk_in, hv_ = dv_in, hg = dg_in, hb = db_in;
+      std::vector<int32_t> 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<size_t>(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<float>(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;