perf(GFX1100-TG200): T14 row-split greedy argmax arm

.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<Q4_K> 2.46 ms/tok (25%), wvSplitKSml 2.32 ms/tok (24%),
-KQuantGemvMmvqK<Q6_K> 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.

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

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<int64_t>(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<int>(gitem % blocks_per_row);
+  const int lane = static_cast<int>(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<int>(j) < arg)) {
+      best = v; arg = static_cast<int>(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<int>(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<int>(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<int64_t>(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,39 @@ 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 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';
+      }();
+      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<void**>(&part_arg),
+                               static_cast<size_t>(rows) * kBpr * sizeof(int64_t), st),
+                "argmax_split scratch");
+          Check(hipMallocAsync(reinterpret_cast<void**>(&part_val),
+                               static_cast<size_t>(rows) * kBpr * sizeof(float), st),
+                "argmax_split scratch");
+          part_rows = rows;
+        }
+        ArgmaxSplitPhaseAK<<<rows * kBpr, kBlock, 0, st>>>(
+            part_arg, part_val, logits.Ptr<float>(), rows, cols, kBpr);
+        ArgmaxSplitPhaseBK<<<static_cast<unsigned>(rows), kBlock, 0, st>>>(
+            out.Ptr<int64_t>(), part_arg, part_val, kBpr);
+        return;
+      }
       ArgmaxK<float><<<static_cast<unsigned>(rows), kBlock, 0, st>>>(
           out.Ptr<int64_t>(), logits.Ptr<float>(), rows, cols);
-    else if (logits.dtype == DType::kBF16)
+    } else if (logits.dtype == DType::kBF16)
       ArgmaxK<__hip_bfloat16><<<static_cast<unsigned>(rows), kBlock, 0, st>>>(
           out.Ptr<int64_t>(), logits.Ptr<__hip_bfloat16>(), rows, cols);
     else