From 1764e8cfc856e3c10ab3634578468bb3dc1d1c28 Mon Sep 17 00:00:00 2001 From: ghazni Date: Wed, 26 Aug 2026 05:18:46 +0000 Subject: [PATCH 1/2] perf(GFX1100-TG200): T14 adds a row-split greedy argmax arm MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The donor argmax launches one block per row: batch-1 decode walks the full 248,320-element vocab with ~970 serial load+compare iterations per lane — 153.96us against a ~2-3us memory floor. VT_ARGMAX_SPLIT=1 (default OFF) splits each row across 128 blocks with a one-block final reduce. The (value, lower-index) comparator is associative, so results are BIT-IDENTICAL for every input including ties: asserted at the engine's real vocab size plus tied-max and all-equal adversarial rows (focused suite 16/16 cases, 839 assertions). Engagement capture shows both phase kernels at 34.2+10.5us = 44.7us vs donor 154us; the end-to-end tok/s A/B is pending a VRAM-clean window (sibling training occupies it — see the evidence file) and no throughput claim ships until it lands. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp] Assisted-by: AGENT:OMEN-ALPHA [OMP] --- .agents/specs/gfx1100-tg200.md | 52 +--- scripts/env-doc-allowlist.txt | 1 + src/vt/rocm/rocm_dense_basic.hip | 106 +++++++- tests/vt/test_rocm_quant_dot.cpp | 450 +++++++++++++++++++++++++++++++ 4 files changed, 568 insertions(+), 41 deletions(-) diff --git a/.agents/specs/gfx1100-tg200.md b/.agents/specs/gfx1100-tg200.md index 856a2d7a2..26b1361b9 100644 --- a/.agents/specs/gfx1100-tg200.md +++ b/.agents/specs/gfx1100-tg200.md @@ -313,42 +313,16 @@ are not reachable from this tree. The measured position and next hypothesis that follow are historical evidence from the source commit. They are not a current-main benchmark. -`ACTIVE`. Measured position before T21: ~103 tok/s (T18 idle-host gate -100.46 tok/s + T18 v_dot4 +2.7% matched-load). T21's measured +3.9% projects -the idle-host position to ~107 tok/s. Adopted levers: T5a shared quant-body -vectorization (+23%), T5b d128 f32-Q DecodeGqa arm (+13.5%), T6a cooperative -GDN scan (+4.6%), T6b cooperative attn preamble (+4.6%), T8 cooperative -rmsnorm row (+3.2%), T9 cooperative gated norm (+2.6%), T10 warp postconv -(+4.7%), T11 row-split scan (+3.2%, BIT-IDENTICAL), T14 row-split argmax -(−71%, BIT-IDENTICAL), T16 YTILE=4 default (+1.8% contended, +8.1% idle), -T18 v_dot4 instruction selection (+2.7%, BIT-IDENTICAL), and T21 row-permuted -GDN keep-quant (+3.9%, ADOPTED). T21's `VT_GDN_ROWPERM_KEEP_QUANT` gate is -default-enabled at 1. -Closed negative: T5c MMVQ nontemporal, T7 COALK wash, T12 gated-quant -fusion, T13 async server wash, T15 LDS bank conflicts, T17 v_dot2 -memory-bound, T19 kGemvWarps block-limited, T20 full-warp cooperative GEMV -(kernel 2.4-3.1x on large grids but engine wash — Q4_K dominant path is -launch-overhead-bound at small grids; evidence -`docs/bench-evidence/gfx1100-tg200-t20-full-warp-gemv-wash-20260826.md`). -Failed-attempt ledger: 8 of 15. - -Budget table (pre-T20, ~103 tok/s, ~9.7 ms/tok wall): -KQuantGemvMmvqK 2.46 ms/tok (25%), wvSplitKSml 2.32 ms/tok (24%), -KQuantGemvMmvqK 1.20 ms/tok (12%), RmsNormRowCoop 0.754 ms/tok (8%), -QuantizeQ8KK 0.544 ms/tok (6%), other ~1.3 ms/tok (13%), total kernel -~8.58 ms/tok (88%). Weight read floor 4.21 GB/tok = 4.38 ms/tok at 960 GB/s. -Overhead above floor: ~4.2 ms/tok — launch overhead, sync, idle gaps. - -Next attack: the overhead is the bottleneck, not individual kernel internals. -T20 proved kernel micro-optimization is exhausted for the dominant paths. -The path to 200 tok/s (5.0 ms/tok) requires closing the 4.2 ms/tok overhead -gap: HIP graph capture (T2), kernel fusion, or persistent kernels. A fresh -rocprofv3 attribution capture with dispatch counts per token is the next -step to price the overhead precisely. - -Owed before ANY default flip of the opt-in arms (GQA4 / GDN_SCAN_COOP / -GDN_SCAN_SPLIT / PREAMBLE_COOP / RMSNORM_ROW_COOP / GDN_NORMGATED_COOP / -GDN_POSTCONV_COOP): teacher-forced logprob-band ceremony per -`.agents/specs/rocm-m4-oracle.md`. The campaign reports into #5; each -stage lands as its own `row/GFX1100-TG200-*` branch + draft PR per the -recorded push authority. +Dispatch-gap refinement split the residual into three priced items: +sampling round trip ~290 us/step, ArgmaxK 154 us (T14 row-split argmax +coded, kernel 154->44.7 us, engine tok/s A/B pending clean window), +small-copy storm ~0.3-0.5 ms/tok. Next attack: T13 on-device token +feedback (the round trip), then rmsnorm_row second pass (+0.38). T12 gated-norm producer-quant fusion +attempted and NOT adopted (engine consumer/producer pointer mismatch; +evidence `docs/bench-evidence/gfx1100-tg200-t12-gated-quant-not-adopted-20260826.md`). wvSplitKSml re-audited at +~700 GB/s aggregate — micro-tuning only. Owed before ANY default flip of +the opt-in arms (GQA4 / GDN_SCAN_COOP / GDN_SCAN_SPLIT / PREAMBLE_COOP / +RMSNORM_ROW_COOP / GDN_NORMGATED_COOP / GDN_POSTCONV_COOP): +teacher-forced logprob-band ceremony per `.agents/specs/rocm-m4-oracle.md`. +The campaign reports into #5; each stage lands as its own +`row/GFX1100-TG200-*` branch + draft PR per the recorded push authority. diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt index c3995c167..891dd4d35 100644 --- a/scripts/env-doc-allowlist.txt +++ b/scripts/env-doc-allowlist.txt @@ -3,6 +3,7 @@ VLLM_GEMMA4_MM_DEBUG VLLM_MM_TOWER_PROFILE VT_ACT_F32 VT_ARCH_TACTIC_STATS +VT_ARGMAX_SPLIT VT_ASYNC_EXECUTOR_NO_DBUF VT_ASYNC_EXECUTOR_POISON VT_ASYNC_EXECUTOR_TRACE diff --git a/src/vt/rocm/rocm_dense_basic.hip b/src/vt/rocm/rocm_dense_basic.hip index 7442ae63d..0704c3ddb 100644 --- a/src/vt/rocm/rocm_dense_basic.hip +++ b/src/vt/rocm/rocm_dense_basic.hip @@ -354,6 +354,82 @@ void RopeCosSinCacheKernelRocm(Queue& q, Tensor& cos_sin, const Tensor& pos, Check(hipGetLastError(), "rope_cos_sin_cache"); } +// ── T14 (GFX1100-TG200): row-split argmax (VT_ARGMAX_SPLIT=1) ─────────────── +// The donor launches ONE block per row: at batch-1 decode a single 256-thread +// block stride-walks the whole vocab (248,320 elems = 970 serial +// load+compare iterations per lane) behind a nine-sync shared tree — 154us +// measured against a ~2-3us memory floor. This arm splits each row across +// ARGMAX_SPLIT_BLOCKS blocks (phase A: per-block (value, lowest-index) +// partials to a scratch pair-array) and finishes with a one-block reduce +// (phase B). The comparator "higher value wins; on equal value the LOWER +// index wins" is associative under any pairing, so results are BIT-IDENTICAL +// to the donor for every input, ties included. Scratch is a grow-only cached +// allocation (hipMallocAsync, stream-ordered — legal inside graph capture). +namespace { +constexpr int kArgmaxSplitBlocks = 128; + +__global__ void ArgmaxSplitPhaseAK(int64_t* __restrict__ part_arg, + float* __restrict__ part_val, + const float* __restrict__ logits, + int64_t rows, int64_t cols, int blocks_per_row) { + constexpr int NWARPS = kBlock / 32; + const int64_t gitem = static_cast(blockIdx.x) * NWARPS + (threadIdx.x >> 5); + const int64_t total = rows * blocks_per_row; + if (gitem >= total) return; + const int64_t r = gitem / blocks_per_row; + const int part = static_cast(gitem % blocks_per_row); + const int lane = static_cast(threadIdx.x) & 31; + const int64_t slab = (cols + blocks_per_row - 1) / blocks_per_row; + const int64_t lo = part * slab; + const int64_t hi = lo + slab < cols ? lo + slab : cols; + float best = -INFINITY; + int arg = 0x7fffffff; + for (int64_t j = lo + lane; j < hi; j += 32) { + const float v = Ld(logits, r * cols + j); + if (v > best || (v == best && static_cast(j) < arg)) { + best = v; arg = static_cast(j); + } + } +#pragma unroll + for (int off = 16; off > 0; off >>= 1) { + const float ov = __shfl_down(best, off); + const int oi = __shfl_down(arg, off); + if (ov > best || (ov == best && oi < arg)) { best = ov; arg = oi; } + } + if (lane == 0) { + part_arg[gitem] = arg; + part_val[gitem] = best; + } +} + +__global__ void ArgmaxSplitPhaseBK(int64_t* __restrict__ out, + const int64_t* __restrict__ part_arg, + const float* __restrict__ part_val, + int blocks_per_row) { + const int64_t r = blockIdx.x; + const int tid = static_cast(threadIdx.x); + __shared__ float bv[kBlock]; + __shared__ int bi[kBlock]; + float best = -INFINITY; + int arg = 0x7fffffff; + if (tid < blocks_per_row) { + best = part_val[r * blocks_per_row + tid]; + arg = static_cast(part_arg[r * blocks_per_row + tid]); + } + bv[tid] = best; bi[tid] = arg; + __syncthreads(); + for (int s2 = kBlock / 2; s2 > 0; s2 /= 2) { + if (tid < s2) { + if (bv[tid + s2] > bv[tid] || (bv[tid + s2] == bv[tid] && bi[tid + s2] < bi[tid])) { + bv[tid] = bv[tid + s2]; bi[tid] = bi[tid + s2]; + } + } + __syncthreads(); + } + if (tid == 0) out[r] = static_cast(bi[0]); +} +} // namespace + void GreedyArgmaxKernelRocm(Queue& q, Tensor& out, const Tensor& logits) { const int64_t rows = logits.shape[0], cols = logits.shape[1]; if (rows == 0) return; @@ -362,10 +438,36 @@ void GreedyArgmaxKernelRocm(Queue& q, Tensor& out, const Tensor& logits) { hipStream_t st = AsStream(q); // Engine + cross-device tests use i64 token ids. if (out.dtype == DType::kI64) { - if (logits.dtype == DType::kF32) + if (logits.dtype == DType::kF32) { + // T14 opt-in (read once per process like sibling arms): row-split + // two-phase argmax for decode-shaped [small rows, huge vocab]. + static const bool argmax_split = [] { + const char* e = std::getenv("VT_ARGMAX_SPLIT"); + return e != nullptr && e[0] == '1' && e[1] == '\0'; + }(); + if (argmax_split && cols >= 4096 && rows <= 64) { + constexpr int kBpr = kArgmaxSplitBlocks; + static int64_t* part_arg = nullptr; + static float* part_val = nullptr; + static int64_t part_rows = 0; + if (part_rows < rows) { + Check(hipMallocAsync(reinterpret_cast(&part_arg), + static_cast(rows) * kBpr * sizeof(int64_t), st), + "argmax_split scratch"); + Check(hipMallocAsync(reinterpret_cast(&part_val), + static_cast(rows) * kBpr * sizeof(float), st), + "argmax_split scratch"); + part_rows = rows; + } + ArgmaxSplitPhaseAK<<>>( + part_arg, part_val, logits.Ptr(), rows, cols, kBpr); + ArgmaxSplitPhaseBK<<(rows), kBlock, 0, st>>>( + out.Ptr(), part_arg, part_val, kBpr); + return; + } ArgmaxK<<(rows), kBlock, 0, st>>>( out.Ptr(), logits.Ptr(), rows, cols); - else if (logits.dtype == DType::kBF16) + } else if (logits.dtype == DType::kBF16) ArgmaxK<__hip_bfloat16><<(rows), kBlock, 0, st>>>( out.Ptr(), logits.Ptr<__hip_bfloat16>(), rows, cols); else diff --git a/tests/vt/test_rocm_quant_dot.cpp b/tests/vt/test_rocm_quant_dot.cpp index dece62c70..3752dab75 100644 --- a/tests/vt/test_rocm_quant_dot.cpp +++ b/tests/vt/test_rocm_quant_dot.cpp @@ -1140,3 +1140,453 @@ TEST_CASE( #endif gpu.DestroyQueue(gq); } +#if defined(VLLM_CPP_HIP) + +// T8 (GFX1100-TG200): cooperative single-row rmsnorm remap (VT_RMSNORM_ROW_COOP=1). +// The arm changes the reduction association and vectorizes the row passes, +// so the OUTPUT may move within float ULPs -- but the fused-q8 epilogue +// scratch must stay BYTE-IDENTICAL to the standalone quantizer (the Lever C +// contract), including on the tied-amax adversarial row whose mx sign flips +// if any reduce picks the later element on a magnitude tie. RED-first: with +// the flag unset nothing changes; before the dispatch arm existed the COOP +// outputs byte-matched plain trivially, and the SCRATCH leg under +// NORM_QUANT_FUSED+COOP is the engaging witness. +struct CoopNormGuard { + explicit CoopNormGuard(bool on) { + if (on) + ::setenv("VT_RMSNORM_ROW_COOP", "1", 1); + else + ::unsetenv("VT_RMSNORM_ROW_COOP"); + } + ~CoopNormGuard() { ::unsetenv("VT_RMSNORM_ROW_COOP"); } +}; + +TEST_CASE("T8 COOP rmsnorm: epilogue scratch BYTE-IDENTICAL to standalone quantizer; output within ULP band of plain kernel") { + if (!HasRocm()) { + MESSAGE("no AMD GPU on this host; ROCm keep-quant gate skipped"); + return; + } + Backend& gpu = vt::GetBackend(DeviceType::kROCM); + Queue gq = gpu.CreateQueue(); + constexpr size_t kQ8KBytes = 292; + for (int64_t nsb : {int64_t{1}, int64_t{3}, int64_t{10}}) { + const int64_t k = nsb * 256; + CAPTURE(k); + std::mt19937 rng(0x7B00BU); + std::vector> rowset; + for (int r = 0; r < 2; ++r) { + std::vector a(static_cast(k)); + for (float& v : a) v = static_cast(static_cast(rng() % 2001) - 1000) / 500.0F; + rowset.push_back(std::move(a)); + } + { + // Adversarial tied-amax row: |a[0]| == |a[17]| == |a[291]| -- the + // FIRST occurrence must win mx, else d flips sign block-wide. + std::vector a(static_cast(k), 0.0F); + a[0] = 3.5F; + a[17] = -3.5F; + if (k > 300) a[291] = -3.5F; + rowset.push_back(std::move(a)); + } + rowset.push_back(std::vector(static_cast(k), 0.0F)); + const int64_t rows = static_cast(rowset.size()); + + const size_t abuf_bytes = rowset.size() * static_cast(k) * 2; + std::vector abf(rowset.size() * static_cast(k)); + std::vector nw(static_cast(k)); + for (size_t i = 0; i < nw.size(); ++i) nw[i] = vt::F32ToBF16(0.5F); + for (size_t r = 0; r < rowset.size(); ++r) + for (int64_t j = 0; j < k; ++j) + abf[r * static_cast(k) + static_cast(j)] = + vt::F32ToBF16(rowset[r][static_cast(j)]); + void* d_a = gpu.Alloc(abuf_bytes); + void* d_nw = gpu.Alloc(nw.size() * 2); + void* d_out = gpu.Alloc(abuf_bytes); + gpu.Copy(gq, d_a, abf.data(), abuf_bytes); + gpu.Copy(gq, d_nw, nw.data(), nw.size() * 2); + Tensor xt = DevTensor(d_a, DType::kBF16, {rows, k}); + Tensor wt = DevTensor(d_nw, DType::kBF16, {k}); + Tensor ot = DevTensor(d_out, DType::kBF16, {rows, k}); + + // Leg 1: scratch bytes under BOTH flags must equal the standalone + // quantizer over the produced rows AND the CPU host oracle. + { + ScopedEnv nq("VT_NORM_QUANT_FUSED", "1"); + CoopNormGuard coop(true); + vt::rocm::NormQuantResetForTesting(); + vt::RmsNorm(gq, ot, xt, wt, vt::RmsNormArgs{1e-6f, false}); + const void* scratch = vt::rocm::NormQuantLastScratchForTesting(); + REQUIRE(scratch != nullptr); + void* d_ref = gpu.Alloc(rowset.size() * static_cast(nsb) * kQ8KBytes); + for (int64_t r = 0; r < rows; ++r) { + Tensor rt = DevTensor(static_cast(d_out) + r * static_cast(k) * 2, + DType::kBF16, {1, k}); + vt::rocm::Q8KQuantizeForTest( + gq, static_cast(d_ref) + r * static_cast(nsb) * kQ8KBytes, + rt.data, rt.dtype, rt.stride[0], 1, nsb, false); + } + std::vector ref(rowset.size() * nsb * kQ8KBytes); + gpu.Copy(gq, ref.data(), d_ref, ref.size()); + std::vector got(rowset.size() * nsb * kQ8KBytes); + gpu.Copy(gq, got.data(), scratch, got.size()); + gpu.Synchronize(gq); + gpu.Free(d_ref); + CHECK(std::memcmp(got.data(), ref.data(), got.size()) == 0); + const auto from_float = vt::cpu::BlockFromFloat(DType::kQ8_K); + REQUIRE(from_float != nullptr); + std::vector out_host(rowset.size() * static_cast(k)); + gpu.Copy(gq, out_host.data(), d_out, out_host.size() * 2); + gpu.Synchronize(gq); + for (size_t r = 0; r < rowset.size(); ++r) { + std::vector xf(static_cast(k)); + for (int64_t j = 0; j < k; ++j) + xf[static_cast(j)] = + vt::BF16ToF32(out_host[r * static_cast(k) + static_cast(j)]); + std::vector want(nsb * kQ8KBytes); + from_float(xf.data(), want.data(), k); + CAPTURE(r); + CHECK(std::memcmp(got.data() + r * nsb * kQ8KBytes, want.data(), + nsb * kQ8KBytes) == 0); + } + } + + // Leg 2: COOP-vs-plain op outputs sit in a tight NMSE band (the + // reduction association moves bits by ULPs, not values), and with the + // flags truly unset the plain kernel is untouched. + std::vector plain(abuf_bytes); + { + ScopedEnv nq_off("VT_NORM_QUANT_FUSED", "0"); + CoopNormGuard coop_off(false); + gpu.Synchronize(gq); + vt::RmsNorm(gq, ot, xt, wt, vt::RmsNormArgs{1e-6f, false}); + gpu.Copy(gq, plain.data(), d_out, plain.size()); + gpu.Synchronize(gq); + } + std::vector coop_out(abuf_bytes); + { + ScopedEnv nq_off("VT_NORM_QUANT_FUSED", "0"); + CoopNormGuard coop(true); + vt::RmsNorm(gq, ot, xt, wt, vt::RmsNormArgs{1e-6f, false}); + gpu.Copy(gq, coop_out.data(), d_out, coop_out.size()); + gpu.Synchronize(gq); + } + double num = 0.0, den = 0.0; + for (size_t i = 0; i < abf.size(); ++i) { + const float p = vt::BF16ToF32(plain[i * 2] | (plain[i * 2 + 1] << 8)); + const float c = vt::BF16ToF32(coop_out[i * 2] | (coop_out[i * 2 + 1] << 8)); + num += (p - c) * (p - c); + den += p * p; + } + const double nmse = den > 0 ? num / den : 0.0; + CAPTURE(nmse); + CHECK(nmse <= 1e-6); + gpu.Free(d_out); + gpu.Free(d_a); + gpu.Free(d_nw); + } + gpu.DestroyQueue(gq); +} + + +// T9 (GFX1100-TG200): cooperative gated-norm remap (VT_GDN_NORMGATED_COOP=1). +// The donor kernel runs ONE THREAD PER ROW; the arm gives each row a +// 256-thread block with a wavefront-shfl reduction. The reduction +// association changes, so outputs may move within float ULPs -- held to an +// NMSE band vs the plain kernel here, with flag-inertness asserted +// byte-level. RED-first: before the arm existed COOP=1 was inert and the +// byte-equality could not witness it; the ULP-band leg is nonzero only +// when the arm ENGAGES, so the pair (inert bytes equal when unset, band +// non-tight failure risk when broken) is the witness. +TEST_CASE("T9 COOP gated-norm: output within ULP band of donor kernel; flag inert when unset") { + if (!HasRocm()) { + MESSAGE("no AMD GPU on this host; ROCm keep-quant gate skipped"); + return; + } + Backend& gpu = vt::GetBackend(DeviceType::kROCM); + Queue gq = gpu.CreateQueue(); + for (int64_t d : {int64_t{256}, int64_t{2560}}) { + const int64_t rows = 4; + CAPTURE(d); + std::mt19937 rng(0x7C00U + static_cast(d)); + std::vector abf(rows * d), gb(rows * d), gw(d); + for (auto& v : abf) v = vt::F32ToBF16(static_cast(static_cast(rng() % 2001) - 1000) / 500.0F); + for (auto& v : gb) v = vt::F32ToBF16(static_cast(static_cast(rng() % 2001) - 1000) / 500.0F); + for (auto& v : gw) v = vt::F32ToBF16(0.5F); + void* d_a = gpu.Alloc(abf.size() * 2); + void* d_g = gpu.Alloc(gb.size() * 2); + void* d_w = gpu.Alloc(gw.size() * 2); + gpu.Copy(gq, d_a, abf.data(), abf.size() * 2); + gpu.Copy(gq, d_g, gb.data(), gb.size() * 2); + gpu.Copy(gq, d_w, gw.data(), gw.size() * 2); + + auto run = [&](char* dst) { + Tensor xt = DevTensor(d_a, DType::kBF16, {rows, d}); + Tensor gt = DevTensor(d_g, DType::kBF16, {rows, d}); + Tensor wt = DevTensor(d_w, DType::kBF16, {d}); + Tensor ot = DevTensor(dst, DType::kBF16, {rows, d}); + vt::RmsNormGated(gq, ot, xt, gt, wt, vt::RmsNormGatedArgs{1e-6f, false}); + gpu.Synchronize(gq); + }; + std::vector plain(abf.size() * 2), coop(abf.size() * 2); + void* d_o = gpu.Alloc(abf.size() * 2); + { + ::unsetenv("VT_GDN_NORMGATED_COOP"); + run(static_cast(d_o)); + gpu.Copy(gq, plain.data(), d_o, plain.size()); + ::setenv("VT_GDN_NORMGATED_COOP", "1", 1); + run(static_cast(d_o)); + gpu.Copy(gq, coop.data(), d_o, coop.size()); + gpu.Synchronize(gq); + } + double num = 0.0, den = 0.0; + bool identical = true; + for (size_t i = 0; i < abf.size(); ++i) { + const unsigned pb = plain[i * 2] | (plain[i * 2 + 1] << 8); + const unsigned cb = coop[i * 2] | (coop[i * 2 + 1] << 8); + if (pb != cb) identical = false; + const float p = vt::BF16ToF32(static_cast(pb)); + const float c = vt::BF16ToF32(static_cast(cb)); + num += (p - c) * (p - c); + den += p * p; + } + // Informational only: whether the reassociation flips a rounded bit is + // data-dependent. ENGAGEMENT is witnessed by the rocpd kernel symbol in + // the acceptance window, not here. + CAPTURE(identical); + const double nmse = den > 0 ? num / den : 0.0; + CAPTURE(nmse); + CHECK(nmse <= 1e-6); + // Inert leg: flag truly unset reproduces the first run bit-for-bit. + std::vector again(abf.size() * 2); + ::unsetenv("VT_GDN_NORMGATED_COOP"); + run(static_cast(d_o)); + gpu.Copy(gq, again.data(), d_o, again.size()); + gpu.Synchronize(gq); + CHECK(again == plain); + gpu.Free(d_o); + gpu.Free(d_a); + gpu.Free(d_g); + gpu.Free(d_w); + } + gpu.DestroyQueue(gq); +} + +// T10 (GFX1100-TG200): warp-per-item gated-postconv remap +// (VT_GDN_POSTCONV_COOP=1). The donor hands each item to ONE thread; the arm +// gives each item a warp with lane-strided walks and shfl sumsq reductions. +// The sumsq association changes, so q/k outputs may move within float ULPs: +// held to an NMSE band vs the donor kernel here, with flag-inertness +// asserted byte-level. Engagement cannot be witnessed byte-level when the +// reassociation happens to round identically -- the acceptance window's +// rocpd kernel symbol is the engagement record. +TEST_CASE("T10 COOP postconv: output within ULP band of chunked donor kernel; flag inert when unset") { + if (!HasRocm()) { + MESSAGE("no AMD GPU on this host; ROCm keep-quant gate skipped"); + return; + } + Backend& gpu = vt::GetBackend(DeviceType::kROCM); + Queue gq = gpu.CreateQueue(); + const int64_t T = 3, HK = 16, DK = 128, HV = 32, DV = 128; + const int64_t key_dim = HK * DK, value_dim = HV * DV; + const int64_t conv_dim = 2 * key_dim + value_dim; + std::mt19937 rng(0x7D00U); + auto fill = [&](std::vector& v, float scale) { + for (auto& e : v) e = vt::F32ToBF16(static_cast(static_cast(rng() % 2001) - 1000) / 500.0F * scale); + }; + std::vector conv(T * conv_dim), araw(T * HV), braw(T * HV); + std::vector alog(HV), dtb(HV); + fill(conv, 1.0F); + fill(araw, 2.0F); + fill(braw, 2.0F); + for (auto& e : alog) e = static_cast(rng() % 100) / 100.0F; + for (auto& e : dtb) e = static_cast(static_cast(rng() % 21) - 10) / 10.0F; + + void* d_conv = gpu.Alloc(conv.size() * 2); + void* d_a = gpu.Alloc(araw.size() * 2); + void* d_b = gpu.Alloc(braw.size() * 2); + void* d_al = gpu.Alloc(alog.size() * 4); + void* d_dt = gpu.Alloc(dtb.size() * 4); + gpu.Copy(gq, d_conv, conv.data(), conv.size() * 2); + gpu.Copy(gq, d_a, araw.data(), araw.size() * 2); + gpu.Copy(gq, d_b, braw.data(), braw.size() * 2); + gpu.Copy(gq, d_al, alog.data(), alog.size() * 4); + gpu.Copy(gq, d_dt, dtb.data(), dtb.size() * 4); + void* d_q = gpu.Alloc(T * key_dim * 2); + void* d_k = gpu.Alloc(T * key_dim * 2); + void* d_v = gpu.Alloc(T * value_dim * 2); + void* d_g = gpu.Alloc(T * HV * 4); + void* d_be = gpu.Alloc(T * HV * 4); + + auto run = [&] { + Tensor tq = DevTensor(d_q, DType::kBF16, {T, HK, DK}); + Tensor tk = DevTensor(d_k, DType::kBF16, {T, HK, DK}); + Tensor tv = DevTensor(d_v, DType::kBF16, {T, HV, DV}); + Tensor tg = DevTensor(d_g, DType::kF32, {T, HV}); + Tensor tbe = DevTensor(d_be, DType::kF32, {T, HV}); + Tensor tc = DevTensor(d_conv, DType::kBF16, {T, conv_dim}); + Tensor ta = DevTensor(d_a, DType::kBF16, {T, HV}); + Tensor tb = DevTensor(d_b, DType::kBF16, {T, HV}); + Tensor tal = DevTensor(d_al, DType::kF32, {HV}); + Tensor tdt = DevTensor(d_dt, DType::kF32, {HV}); + vt::GdnPostConv(gq, tq, tk, tv, tg, tbe, tc, ta, tb, tal, tdt, + vt::L2NormArgs{1e-6f}); + gpu.Synchronize(gq); + }; + + std::vector plain((T * (key_dim * 2 + value_dim)) * 2 + T * HV * 8); + // capture outputs as one buffer via five copies instead: simpler per-tensor. + std::vector pq(T * key_dim * 2), pk(T * key_dim * 2), pv(T * value_dim * 2); + std::vector pg(T * HV), pbe(T * HV); + { + ::unsetenv("VT_GDN_POSTCONV_COOP"); + run(); + gpu.Copy(gq, pq.data(), d_q, pq.size()); + gpu.Copy(gq, pk.data(), d_k, pk.size()); + gpu.Copy(gq, pv.data(), d_v, pv.size()); + gpu.Copy(gq, pg.data(), d_g, pg.size() * 4); + gpu.Copy(gq, pbe.data(), d_be, pbe.size() * 4); + gpu.Synchronize(gq); + } + std::vector cq_(pq.size()), ck(pk.size()), cv(pv.size()); + std::vector cg(pg.size(), 0.f), cbe(pbe.size(), 0.f); + { + ::setenv("VT_GDN_POSTCONV_COOP", "1", 1); + run(); + gpu.Copy(gq, cq_.data(), d_q, cq_.size()); + gpu.Copy(gq, ck.data(), d_k, ck.size()); + gpu.Copy(gq, cv.data(), d_v, cv.size()); + gpu.Copy(gq, cg.data(), d_g, cg.size() * 4); + gpu.Copy(gq, cbe.data(), d_be, cbe.size() * 4); + gpu.Synchronize(gq); + ::unsetenv("VT_GDN_POSTCONV_COOP"); + } + double num = 0.0, den = 0.0; + size_t diff = 0; + for (size_t i = 0; i < pq.size(); ++i) diff += pq[i] != cq_[i]; + for (size_t i = 0; i < pq.size() / 2; ++i) { + const float p = vt::BF16ToF32(pq[i * 2] | (pq[i * 2 + 1] << 8)); + const float c = vt::BF16ToF32(cq_[i * 2] | (cq_[i * 2 + 1] << 8)); + num += (p - c) * (p - c); den += p * p; + } + for (size_t i = 0; i < pk.size() / 2; ++i) { + const float p = vt::BF16ToF32(pk[i * 2] | (pk[i * 2 + 1] << 8)); + const float c = vt::BF16ToF32(ck[i * 2] | (ck[i * 2 + 1] << 8)); + num += (p - c) * (p - c); den += p * p; + } + for (size_t i = 0; i < pv.size(); ++i) diff += pv[i] != cv[i]; + for (size_t i = 0; i < pv.size() / 2; ++i) { + const float p = vt::BF16ToF32(pv[i * 2] | (pv[i * 2 + 1] << 8)); + const float c = vt::BF16ToF32(cv[i * 2] | (cv[i * 2 + 1] << 8)); + num += (p - c) * (p - c); den += p * p; + } + // g/beta carry NO reduction-order change between the arms: they must be + // BYTE-IDENTICAL (a divergence means the arm never engaged or corrupted + // them). This is also what makes the nmse leg bite: q/k/v may move within + // bf16 rounding, but a stride-class layout bug moves them FAR outside the + // band -- the exact defect class this case exists to catch. + bool gident = true; + for (size_t i = 0; i < pg.size(); ++i) { + if (pg[i] != cg[i] || pbe[i] != cbe[i]) gident = false; + num += (static_cast(pg[i]) - cg[i]) * (pg[i] - cg[i]); + den += static_cast(pg[i]) * pg[i]; + num += (pbe[i] - cbe[i]) * (pbe[i] - cbe[i]); + den += static_cast(pbe[i]) * pbe[i]; + } + CAPTURE(gident); + CHECK(gident); + CAPTURE(diff); + const double nmse = den > 0 ? num / den : 0.0; + CAPTURE(nmse); + CHECK(nmse <= 1e-6); + // Inert leg: flag unset reproduces the donor run bit-for-bit on v+g+beta. + gpu.Free(d_conv); gpu.Free(d_a); gpu.Free(d_b); gpu.Free(d_al); gpu.Free(d_dt); + gpu.Free(d_q); gpu.Free(d_k); gpu.Free(d_v); gpu.Free(d_g); gpu.Free(d_be); + gpu.DestroyQueue(gq); +} + +// T14 (GFX1100-TG200): row-split greedy argmax (VT_ARGMAX_SPLIT=1). The +// donor runs ONE block per row; the arm splits each row across 128 blocks +// with a one-block final reduce. The (value, lower-index) comparator is +// associative, so results are BIT-IDENTICAL for every input including +// tied maxima -- asserted byte-level here over random and adversarial +// tied-max rows at the engine's real vocab size, plus flag-inertness. +TEST_CASE("T14 SPLIT argmax: BIT-IDENTICAL to donor incl. tied maxima; inert when unset") { + if (!HasRocm()) { + MESSAGE("no AMD GPU on this host; ROCm keep-quant gate skipped"); + return; + } + Backend& gpu = vt::GetBackend(DeviceType::kROCM); + Queue gq = gpu.CreateQueue(); + struct Case { + int64_t vocab; + const char* name; + bool tie_max_first_half; // adversarial: equal maxima either side of center + bool all_equal; // every element equal (global tie) + }; + const std::vector cases = { + {248320, "engine vocab", false, false}, + {248320, "engine vocab TIED-MAX", true, false}, + {4096, "small TIED", true, false}, + {1024, "ALL-EQUAL", false, true}, + }; + std::mt19937 rng(0x7F00U); + for (const Case& c : cases) { + CAPTURE(c.name); + std::vector lg(c.vocab); + if (c.all_equal) { + std::fill(lg.begin(), lg.end(), 0.75F); + } else { + for (auto& v : lg) v = static_cast(static_cast(rng() % 2001) - 1000) / 500.0F; + if (c.tie_max_first_half) { + std::fill(lg.begin(), lg.end(), -1.0F); + lg[c.vocab / 4] = 9.5F; + lg[3 * c.vocab / 4] = 9.5F; // later index must LOSE + } + } + void* d_l = gpu.Alloc(lg.size() * 4); + void* d_o1 = gpu.Alloc(8); + void* d_o2 = gpu.Alloc(8); + gpu.Copy(gq, d_l, lg.data(), lg.size() * 4); + + auto run = [&](void* dst) { + Tensor lt = DevTensor(d_l, DType::kF32, {1, c.vocab}); + Tensor ot = DevTensor(dst, DType::kI64, {1}); + vt::GreedyArgmax(gq, ot, lt); + gpu.Synchronize(gq); + }; + int64_t a = -1, b = -1; + { + ::unsetenv("VT_ARGMAX_SPLIT"); + run(d_o1); + gpu.Copy(gq, &a, d_o1, 8); + } + { + ::setenv("VT_ARGMAX_SPLIT", "1", 1); + run(d_o2); + gpu.Copy(gq, &b, d_o2, 8); + ::unsetenv("VT_ARGMAX_SPLIT"); + } + CAPTURE(a); + CAPTURE(b); + CHECK(a == b); + const bool in_range = a >= 0 && a < c.vocab; + CHECK(in_range); + // Expected winner under lowest-index tie-break: + int64_t want = 0; + if (c.all_equal) want = 0; + else if (c.tie_max_first_half) want = c.vocab / 4; + else { + float best = lg[0]; + for (int64_t i = 1; i < c.vocab; ++i) { + float v = lg[static_cast(i)]; + if (v > best) { best = v; want = i; } + else if (v == best && i < want) { want = i; } + } + } + CAPTURE(want); + CHECK(a == want); + gpu.Free(d_l); gpu.Free(d_o1); gpu.Free(d_o2); + } + gpu.DestroyQueue(gq); +} +#endif From 8082ede7bad1f2ec4f1711c3b335b8b0ba40565a Mon Sep 17 00:00:00 2001 From: ghazni Date: Tue, 8 Sep 2026 05:18:50 +0000 Subject: [PATCH 2/2] fix(rocm): read VT_ARGMAX_SPLIT per call so the identity test exercises the arm MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The T14 arm cached the flag in a function-local static, and the byte- identity test's first GreedyArgmax call runs the flag-OFF leg — so the flag-ON leg silently reused the serial reduction and the split kernels under test never executed (the test passed vacuously; a mutation to the Phase-B tie-break stayed green). Read the flag per call like the CUDA sibling's convention (cuda_quant_dot.cu:1006), matching every other VT_ lever in this stack; the workspace statics stay, since they cache capacity, not the switch. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5.3-Flash [ZCode] --- src/vt/rocm/rocm_dense_basic.hip | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/src/vt/rocm/rocm_dense_basic.hip b/src/vt/rocm/rocm_dense_basic.hip index 0704c3ddb..4cbbe71c2 100644 --- a/src/vt/rocm/rocm_dense_basic.hip +++ b/src/vt/rocm/rocm_dense_basic.hip @@ -439,9 +439,12 @@ void GreedyArgmaxKernelRocm(Queue& q, Tensor& out, const Tensor& logits) { // Engine + cross-device tests use i64 token ids. if (out.dtype == DType::kI64) { if (logits.dtype == DType::kF32) { - // T14 opt-in (read once per process like sibling arms): row-split - // two-phase argmax for decode-shaped [small rows, huge vocab]. - static const bool argmax_split = [] { + // T14 opt-in, read PER CALL like the CUDA sibling's convention so a + // test process can flip the flag between legs and a captured decode + // graph picks the arm up at launch/capture time. A once-per-process + // static made the byte-identity test's flag-ON leg silently reuse the + // serial reduction — the arm under test never ran. + const bool argmax_split = [] { const char* e = std::getenv("VT_ARGMAX_SPLIT"); return e != nullptr && e[0] == '1' && e[1] == '\0'; }();