From 87b46882875fc868d152baa658d9cd161a0e40b5 Mon Sep 17 00:00:00 2001 From: Vikash Loomba Date: Sat, 29 Aug 2026 08:17:19 -0700 Subject: [PATCH 1/7] spec(BACKEND-ROCM): decompose the Q8_K activation quantizer Issue #1876 needs a committed correctness and measurement contract before the ROCm kernel changes. This specification pins the byte oracle, both production routes, the same-binary lever, and the gfx1100 adoption gate. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:gpt-5 [codex] --- .agents/issue-index.md | 1 + .../specs/rocm-q8k-cooperative-quantizer.md | 575 ++++++++++++++++++ 2 files changed, 576 insertions(+) create mode 100644 .agents/specs/rocm-q8k-cooperative-quantizer.md diff --git a/.agents/issue-index.md b/.agents/issue-index.md index df1e9bba7..d093698c6 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -876,3 +876,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#2220](https://github.com/mudler/vllm.cpp/issues/2220) | `LTX25-ORACLE-ABSOLUTE` | **A CUDA toolkit staged off CIFS compiles but cannot be LINKED against, and the precondition that should catch it checks the one link that works.** `/workspace` is CIFS and stores no symlink, so a staged toolkit carries only `libcudart.so.13.3.29` and `libcublasLt.so.13.6.0.2`. `ltx25-oracle-absolute-render.sh` rebuilt the links with `b=${f%%.so.*}; ln -sf "$f" "$b.so"; ln -sf "$f" "$b.so.${f#*.so.}"` -- and `${f#*.so.}` strips the SHORTEST prefix, so it expands to `13.3.29` rather than `13`. The second `ln` therefore links `libcudart.so.13.3.29` to ITSELF and **`libcudart.so.13`, the SONAME, is never created**. That is the name `ld` resolves versioned undefined symbols against, so CMake reports `Found CUDAToolkit`, every CUDA TU compiles, and the job dies 21 minutes later with 38 `undefined reference to ...@libcudart.so.13` / `@libcublasLt.so.13` and `ninja: build stopped`. `need_ok` tested `[ -f .../libcublasLt.so ]`, which is exactly the link the loop DID create, so it passed on an unlinkable toolkit -- a precondition that cannot fail. LATENT, not new: the staging branch is a FALLBACK, and every earlier lease found `/usr/local/cuda` 13.0.88 and never took it; `dgx:gpu0` went `unhealthy ... worker_lost` for 3h20m on 2026-08-28 and returned without a toolkit, which exercised it for the first time. The two runs A/B in their own configure logs: `20260827T220845Z` `/usr/local/cuda` 13.0.88 built in 1192 s, `20260828T224529Z` `/root/cudatk` 13.3.73 failed at link. Fixed in flow: take the MAJOR (`v=${f#*.so.}; ${v%%.*}`), prefer `ldconfig -n` which reads each object's own `DT_SONAME`, and assert `.so` resolves AND `.so.` exists for both libraries BEFORE the build. Red-before/green-after on a replica of the CIFS layout: old loop creates no `.so.13`, new logic creates both, and the guard FAILS on the old layout, PASSES on the new, and FAILS on the real NAS source. `rc` job `1ad519b1-4e75-41d7-9386-9932076390f1`, exit 34. Also recorded in [`environment.md`](environment.md) as a lease-environment fact, because it will bite the next row | bug | | [#2244](https://github.com/mudler/vllm.cpp/issues/2244) | `BACKEND-TENSTORRENT-QWEN35` | **The staging upload still pays tt-metal's full creation path on every step.** W4 (#2107) landed levers 1+2 — bulk bf16 staging and single-slot resolution, 0.104 → 0.177 tok/s (+70%), staging `Numel()` share 27.09% → 1.76% — and re-attributed the residual: ~23% of the staging chain is tt-metal per-upload internal work (a fresh `MeshBuffer` allocation, cluster/chip discovery, CQ completion handling) and ~19.2% is CPU threadpool spin, with lever 3 (batch per-layer staging) explicitly not taken. Every upload still goes through `UploadRowsBf16` (`src/vt/tenstorrent/tenstorrent_ops.cpp:469`), which builds a new `ttnn::Tensor` via `from_span`, so identical geometry pays the creation path every step. The lever the W4 record named: allocate the device buffer once per staging slot — lifecycle tied to the slot structures under the #1486 never-destroy rule — and write the host bytes through the mesh command queue (`MeshCommandQueue::enqueue_write`/`enqueue_write_shard`), making the upload allocation-free. The tt-metal-internal half is a proof obligation, not an assumption: read the pinned tt-metal's mesh write path and trace the executed chain before declaring any part of the lever unreachable. `StagingStats` gains route counters for the new path; the capture-unsafe host-write refusals keep their semantics; the f32-conversion arms keep their declared dtypes. Invariant: staging stays bit-identical — the sacred golden pair 16/16 and the full TT suite green; this wave changes speed, never tokens. Evidence owed: same-method before/after profile on the P150 (identical leg, lock discipline) plus a fresh benchmark-record entry; a wall that does not move is a reported result, not a failure — the attribution shifts or the lever is named unreachable with the trace that proves it. Owned by `BACKEND-TENSTORRENT-QWEN35`, next wave after W3 (#2201, landed via #2217) | feature | | [#2277](https://github.com/mudler/vllm.cpp/issues/2277) | `MODEL-MM-glm5-next-glm5-next-for-conditional-generation` | **The published GLM-5.3-Flash GGUF is `tokenizer.ggml.pre = "glm4"`, and our pre-tokenizer table refuses that name — this is where the loader stops once [#2268](https://github.com/mudler/vllm.cpp/issues/2268) is fixed.** Measured 2026-08-29 on one tree and one build directory, three legs of one probe object driven through `LoadedEngine::FromModelDir` on `device = kCPU` at `/mnt/nas_share/rc/ckpt/GLM-5.3-Flash-UD-Q2_K_XL/`, headers only: the baseline reader stops at `attention.key_length_mla - attention.key_length is -256 but rope.dimension_count is 0`; with the MLA convention fixed but `attention.linear_head_count` still required it stops at `missing metadata key glm5next.attention.linear_head_count`, one key along; with both fixed it stops at `tokenizer: unsupported tokenizer.ggml.pre "glm4"`, past config resolution entirely. `src/vllm/tokenizer/tokenizer.cpp::FromGguf` maps seven pre names — `qwen35`, `qwen2`, `llama-bpe`, the four GPT-4o names, `deepseek-llm`, the three DeepSeek-V3 names and `laguna` — and refuses the rest by name. `glm4` is what every GLM-4 / GLM-5 GGUF carries; shard 1's KV block states `tokenizer.ggml.model = gpt2`, `tokenizer.ggml.pre = glm4`, 154880 tokens and 321649 merges. **The splitting rule is free and the BOS is not.** llama.cpp maps `glm4` and `chatglm-bpe` to `LLAMA_VOCAB_PRE_TYPE_CHATGLM4` (`b10451:src/llama-vocab.cpp:2256-2258`), whose regex at `:398` is BYTE-IDENTICAL to `LLAMA_VOCAB_PRE_TYPE_LLAMA3`'s at `:289`, so `SplitPattern::kLlama3` is EXACT here rather than the "close approximation" that [#347](https://github.com/mudler/vllm.cpp/issues/347) and [#1924](https://github.com/mudler/vllm.cpp/issues/1924) each had to undo — compare the two byte strings in the fix rather than trusting this sentence. But the same branch sets `special_bos_id = LLAMA_TOKEN_NULL` (`:2259`) while the artifact states `tokenizer.ggml.bos_token_id = 154822`, so llama.cpp DISCARDS a BOS id the file carries; a port that reads it and prepends it emits one token no reference run emits, on every request, and a token gate built from our own tokenizer could not see it because both sides would agree. Scope: map both names onto the CHATGLM4 rule with the byte comparison recorded rather than asserted, mirror the `special_bos_id` suppression with a case that fails if a BOS is prepended, and gate through `FromModelDir` on a `pre = "glm4"` fixture so the refusal that moves is the production one. Recorded as O20 in [`specs/glm5-next-flash.md`](specs/glm5-next-flash.md), which carries the paired measurement | bug | +| [#1876](https://github.com/mudler/vllm.cpp/issues/1876) | `BACKEND-ROCM` | The byte-identical cooperative Q8_K quantizer must route dense and grouped ROCm keep-quant paths, gated and measured on gfx1100. gfx1201 remains external | perf | diff --git a/.agents/specs/rocm-q8k-cooperative-quantizer.md b/.agents/specs/rocm-q8k-cooperative-quantizer.md new file mode 100644 index 000000000..aa4c2b6e0 --- /dev/null +++ b/.agents/specs/rocm-q8k-cooperative-quantizer.md @@ -0,0 +1,575 @@ +# Cooperative ROCm Q8_K activation quantizer + +Issue: [#1876](https://github.com/mudler/vllm.cpp/issues/1876) + +Row: `BACKEND-ROCM` + +## Now + +The specification is `READY`. Product implementation and tests are `PENDING`. +The first implementation must keep the cooperative arm off by default. + +This commit adds only this specification and the issue ownership row. It makes +no matrix or lifecycle change. + +## Issue ownership + +Issue #1876 owns this `BACKEND-ROCM` performance slice. The issue-index row +names `BACKEND-ROCM` as the owner and `perf` as the kind. + +The issue stays open until the accepted implementation and its evidence land. +A negative result also leaves the issue open. + +## Git integration + +Use one pull request for the specification and implementation. This choice is +the repository default and is recorded in the local developer preferences. + +The specification commit must precede all product and test commits. A fresh +implementer must start from this committed specification. + +## Scope + +The implementation has these in-scope changes: + +- Add one cooperative Q8_K activation quantizer for the ROCm backend. +- Add one shared ROCm Q8_K launcher for dense and grouped keep-quant calls. +- Add the same-binary `VT_ROCM_Q8K_BLOCK` selection lever. +- Add direct scratch-byte tests and production-route witnesses. +- Gate correctness and performance on local `gfx1100` hardware. +- Record a negative result if the candidate fails an acceptance rule. + +The implementation must not change these surfaces: + +- The keep-quant GEMM dot bodies. +- Issue #1910's cooperative Q6_K GEMM. +- Graph enablement or graph policy. +- Producer fusion or quantization inside a GEMM prologue. +- CUDA kernels, launch policy, or environment variables. +- Quantization formats or the `BlockQ8_K` layout. +- Area matrices, lifecycle state, or other keyed records. +- Public documentation before a real-checkpoint gate succeeds. +- Product code from pull request #1936. + +## Reference hierarchy + +### vLLM primary reference + +The primary pin is `5559679229bc961848b121ccdeaa8fa5d79bec98`. +That tree has no executable GGUF Q8_K activation quantizer. + +The gap check searched `vllm/model_executor/layers/quantization/**`, +`vllm/model_executor/**`, and `csrc/**`. Searches for `Q8_K`, `q8_K`, and +`quantize_q8` found no implementation. Therefore vLLM supplies no byte oracle +or launch decomposition for this slice. + +This absence does not change mirror priority. If a future vLLM pin implements +this behavior, reconcile the row against vLLM before landing. + +### Behavioral byte oracle + +The byte oracle is the local CPU path: + +- `include/vt/quant.h:69::BlockFromFloat` exposes the activation encoder. +- `src/vt/cpu/cpu_quant_act.cpp:88::QuantizeRowQ8_K` defines the arithmetic. +- `src/vt/cpu/cpu_quant_blocks.h:129::BlockQ8_K` defines the 292-byte layout. + +The CPU encoder mirrors llama.cpp +`ggml/src/ggml-quants.c::quantize_row_q8_K_ref`. The block layout mirrors +llama.cpp `ggml/src/ggml-common.h::block_q8_K`. + +The pinned llama.cpp source places the reference encoder at +`ggml/src/ggml-quants.c:2768::quantize_row_q8_K_ref` and the block at +`ggml/src/ggml-common.h:371::block_q8_K`. The current local CPU comment still +names the older line coordinate `ggml-quants.c:2696`. The persistent anchor is +the symbol at pin `10bf611e533d81f739128304991c5e133c6aebd8`. + +### Local decomposition donor + +The primary implementation donor is the current CUDA path: + +- `src/vt/cuda/cuda_quant_dot.cu:238::QuantizeQ8KPreqKernel` maps one thread to + each activation element. +- `src/vt/cuda/cuda_quant_dot.cu:1643::LaunchQuantizeQ8K` selects one quantizer and + is shared by dense and grouped consumers. +- `src/vt/cuda/cuda_quant_dot.cu:176::QuantizeQ8KKernel` retains the serial control + arm. + +`QuantizeQ8KPreqKernel` uses a 256-thread block for one row superblock. Its +reduction carries `(abs, value, index)` and selects the lowest original index +on an exact absolute-value tie. + +This change ports that decomposition to HIP. It does not port the CUDA default +or the CUDA environment variable. + +### llama.cpp secondary oracle + +The secondary pin is llama.cpp +`10bf611e533d81f739128304991c5e133c6aebd8`, tag `b10451`. The pin and its +gateability record are in `.agents/oracles/llama-cpp.md`. + +The relevant dense and grouped executing chain is: + +1. `ggml/src/ggml-cuda/ggml-cuda.cu:1812::ggml_cuda_mul_mat` selects MMVQ for a + dense quantized matrix-vector operation. +2. `ggml/src/ggml-cuda/ggml-cuda.cu:1899::ggml_cuda_mul_mat_id` selects the same + MMVQ path for eligible grouped expert operations. +3. `ggml/src/ggml-cuda/ggml-cuda.cu:1783::ggml_cuda_should_fuse_mul_mat_vec_q` + governs fused eligible nodes that also call the same MMVQ entry point. +4. `ggml/src/ggml-cuda/mmvq.cu:1153::ggml_cuda_mul_mat_vec_q` allocates Q8_1 + activation scratch. +5. The same function calls + `ggml/src/ggml-cuda/quantize.cu:558::quantize_row_q8_1_cuda`. +6. That launcher starts `ggml/src/ggml-cuda/quantize.cu:54::quantize_q8_1` with + `CUDA_QUANTIZE_BLOCK_SIZE == 256`. +7. `ggml/src/ggml-cuda/mmvq.cu:1000::mul_mat_vec_q_switch_type` consumes the Q8_1 + scratch in the quantized matrix-vector kernel. + +The Q8_1 kernel supplies decomposition and performance context only. Q8_1 has +32-element blocks, half precision metadata, and no Q8_K `bsums` field. It is +not a byte oracle for Q8_K. + +### Local production chain + +The dense path is: + +1. `src/vllm/model_executor/model_loader/gguf_keep_quant.cpp:83::GgufQuantComputeAvailable` + keeps eligible weights quantized when the ROCm provider exists. +2. Model linear calls reach `src/vt/ops.cpp:149::MatmulBT`. +3. A block weight routes to `src/vt/ops.cpp:191::MatmulBTQuant`. +4. `src/vt/rocm/rocm_ops.hip:212` registers + `src/vt/rocm/rocm_grouped_gemm.hip:633::MatmulBTQuantKernelRocm`. +5. `MatmulBTQuantKernelRocm` allocates Q8_K scratch and launches + `src/vt/rocm/rocm_grouped_gemm.hip:148::QuantizeQ8KK` at line 666. +6. `KQuantGemmK` or `KQuantGemmKCoopQ6K` consumes that scratch. + +The grouped path is: + +1. Model expert calls reach `src/vt/ops.cpp:221::MatmulBTQuantGrouped`. +2. `src/vt/rocm/rocm_ops.hip:215` registers + `src/vt/rocm/rocm_grouped_gemm.hip:699::MatmulBTQuantGroupedKernelRocm`. +3. `MatmulBTQuantGroupedKernelRocm` allocates Q8_K scratch and launches + `QuantizeQ8KK` directly at line 744. +4. `src/vt/rocm/rocm_grouped_gemm.hip:384::GroupedKQ8K` consumes that + scratch for each routed expert. + +Both launch sites must move to one shared ROCm launcher. No direct Q8_K launch +may remain in either production consumer. + +## Observed gap + +`src/vt/rocm/rocm_grouped_gemm.hip::QuantizeQ8KK` maps one thread to one +256-element superblock. The thread scans all elements, quantizes all elements, +and computes all 16 sums. + +The amax loop calls `DLoadAct` to compute `fabsf`. It calls `DLoadAct` again +when an element wins. The candidate must load each element once and retain the +value used by the reduction. + +Both `MatmulBTQuantKernelRocm` and `MatmulBTQuantGroupedKernelRocm` launch the +serial kernel directly. A dense-only change cannot close issue #1876. + +The unmerged T27 commit +`05455b6a97f1fe60615af105b63b8611ac681873` is historical evidence. Its +8-thread candidate measured 91.532 to 93.417 tok/s on `gfx1100`. The 2.06% +change had 5 of 5 byte-identical winning pairs. + +T27 was dense-only, had no direct scratch-byte suite, and stayed off by +default. Do not cherry-pick it. Its numbers do not establish a current result. + +## Design + +### Cooperative kernel + +Port the CUDA block-per-superblock decomposition to HIP. Launch one 256-thread +block for each `(row, superblock)` pair. + +Each thread must do these operations: + +1. Load its source element once through `DLoadAct`. +2. Store its absolute value, signed value, and original index for reduction. +3. Participate in a lowest-index signed-amax reduction. +4. Write one `qs` byte after the scale is known. + +The first 16 threads each compute one signed sum over 16 consecutive `qs` +bytes. One thread writes the f32 delta. + +The implementation can use shared memory or wave operations. The result and +the original index rule are binding. A wave-specific shortcut must retain a +correct block-wide reduction on `gfx1100`. + +### Shared launcher + +Add one private launcher beside the ROCm kernels. The launcher accepts the +scratch pointer, source pointer, activation dtype, row stride, row count, +superblock count, stream, and production route. + +The route is `dense` or `grouped`. The launcher selects the legacy or +cooperative kernel, increments its route witness, and performs the launch. + +Both production consumers must call this launcher. The launcher must preserve +the existing scratch allocation, stream ordering, and error check. + +### Same-binary selection + +`VT_ROCM_Q8K_BLOCK` accepts only `0` and `1`: + +- `0` selects the current serial `QuantizeQ8KK` arm. +- `1` selects the cooperative block arm. +- An unset value selects the recorded default. +- Any other value refuses by name instead of choosing an arm silently. + +The first implementation default is `0`. Read the selection at launch time so +one test process can exercise both values. + +A final default of `1` is allowed only by the acceptance rules in this spec. +If the default becomes `1`, value `0` remains the legacy same-binary arm. + +### Test-only seams + +Add private test-only functions in the ROCm translation unit. Tests may +forward-declare them, following the existing `KQuantCoopDispatchCount` +convention. + +Extend `tests/vt/test_backend_cross_device.cpp`. Keep the existing test target +and do not add a CMake registration for this slice. + +The seam must include these capabilities: + +- Launch either arm explicitly into caller-provided Q8_K scratch. +- Reset and read dense legacy, dense candidate, grouped legacy, and grouped + candidate counters. +- Resolve the environment policy without changing the public ABI. + +The explicit-arm hook must not read `VT_ROCM_Q8K_BLOCK`. This separation lets +the direct byte test compare both kernels regardless of the default. + +No declaration may enter `include/vllm.h` or another public header. + +## Arithmetic invariants + +Each Q8_K superblock contains exactly 256 source elements and 292 output bytes. +The output fields are one f32 `d`, 256 signed `qs` bytes, and 16 signed `bsums`. + +The implementation must preserve these rules: + +- Decode f32, f16, and bf16 source bytes through the existing `DLoadAct` rules. +- Treat `a_rs` as an element stride and never read row padding as data. +- Select the first source element with the largest absolute value. +- Let the lowest original index win an exact absolute-value tie. +- Preserve the selected element's sign in `mx`. +- Set `iscale` to `-127.0f / mx`. +- Use `DNearestInt(iscale * x)` for every quant. +- Clamp only the upper side with `v < 127 ? v : 127`. +- Set each `bsums[g]` to the exact integer sum of `qs[16*g..16*g+15]`. +- Set `d` to `1.0f / iscale`. +- For an all-zero superblock, write zero to all 292 bytes. + +The tie rule is not optional. Opposite-sign tied maxima produce a different +scale sign when the wrong index wins. + +## Red-first tests + +Write and run the direct tests before product code. Preserve the failing output +under `/tmp`. + +The direct test must compare all 292 bytes of every `BlockQ8_K` against both +references: + +1. The legacy ROCm GPU arm. +2. The independent CPU `vt::cpu::BlockFromFloat(DType::kQ8_K)` arm. + +Run the comparison for f32, f16, and bf16 sources. Convert f16 and bf16 source +values independently to f32 before calling the CPU encoder. + +Cover these input classes: + +- Deterministic random finite data. +- All-zero data with sentinel-filled output storage. +- Opposite-sign exact tied maxima in both sign orders. +- Ties that cross a wave and a reduction boundary. +- Nontrivial row stride with sentinel padding. +- More than one activation row. +- `nsb` values 1, 2, 3, 10, and 16. + +The test must name the first mismatching byte, block, row, dtype, and input +class. A field-level equality check is insufficient. + +The production tests must call `vt::MatmulBTQuant` and +`vt::MatmulBTQuantGrouped`. For each environment arm, they must assert the +matching route counter changes and the other route counter stays unchanged. + +Also test unset, `0`, `1`, and one invalid value. Test the explicit-arm hook +independently from the environment policy. + +The record-first specification commit has no product behavior to make red. Do +not fabricate a red result for this commit. + +## Production reachability + +Output equivalence cannot prove that production used the cooperative arm. +Route counters are the causal witnesses. + +The dense witness must enter through `vt::MatmulBTQuant`. It must observe the +selected arm inside `MatmulBTQuantKernelRocm` through the shared launcher. + +The grouped witness must enter through `vt::MatmulBTQuantGrouped`. It must +observe the selected arm inside `MatmulBTQuantGroupedKernelRocm` through the +same launcher. + +The direct explicit-arm test proves kernel arithmetic. It does not replace +either production witness. + +The profiler must also contain the selected kernel name and the matching route +counter totals. A counter-only result cannot establish GPU engagement. + +## Build and correctness gates + +Use a fresh Release build directory. Do not copy a configured build tree. + +Configure with these binding values: + +```sh +BUILD_DIR="$(mktemp -d /tmp/vllmcpp-rocm-q8k-release.XXXXXX)" +cmake -S . -B "$BUILD_DIR" -G Ninja \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_CXX_COMPILER=/usr/bin/c++ \ + -DCMAKE_HIP_COMPILER=/opt/rocm/lib/llvm/bin/clang++ \ + -DROCM_PATH=/opt/rocm \ + -DVLLM_CPP_HIP=ON \ + -DVLLM_CPP_HIP_ARCHITECTURES=gfx1100 +cmake --build "$BUILD_DIR" -j 4 +``` + +Every HIP runtime command must use this environment: + +```sh +LD_LIBRARY_PATH=/opt/rocm/lib:/opt/rocm/lib/llvm/lib +``` + +Every GPU command must hold this lock for the complete command: + +```sh +flock /home/vikash/gpu.lock +``` + +Run the new focused case under the lock. Then run the full applicable Release +test gate under the same lock. + +Use these complete test commands: + +```sh +flock /home/vikash/gpu.lock env \ + LD_LIBRARY_PATH=/opt/rocm/lib:/opt/rocm/lib/llvm/lib \ + "$BUILD_DIR/tests/test_backend_cross_device" "--test-case=*ROCm Q8_K*" +flock /home/vikash/gpu.lock env \ + LD_LIBRARY_PATH=/opt/rocm/lib:/opt/rocm/lib/llvm/lib \ + ctest --test-dir "$BUILD_DIR" --output-on-failure +``` + +Run `scripts/agent-preflight.sh` after the focused and full tests pass. Run the +staged quiet preflight immediately before the implementation commit. + +If a command writes through `tee`, enable `set -o pipefail` and record the +producer's exit status. A log file must not replace the pipeline return code. + +## Real-checkpoint correctness gate + +Use this exact model artifact: + +- Path: `/home/vikash/models/Qwen3.5-4B-Q4_K_M-unsloth-e87f1764.gguf`. +- Size: 2,740,937,888 bytes. +- SHA-256: `00fe7986ff5f6b463e62455821146049db6f9313603938a70800d1fb69ef11a4`. +- Source: `unsloth/Qwen3.5-4B-GGUF`. +- Revision: `e87f176479d0855a907a41277aca2f8ee7a09523`. + +Use `tools/tg200-prompt.txt` with SHA-256 +`e2b801cc6a5739cd317c2f77adfb67040667de524ab60ca64aac39f79c846bba`. +The blob exists on the local `remotes/pr/1936` ref but is absent at this base. +Import only that prompt artifact before measurement and verify its hash. + +Do not import product code from pull request #1936. If the prompt artifact is +still absent, the real-checkpoint gate stays `PENDING`. + +The correctness run uses batch 1, greedy decode, `--max-tokens 256`, +`--temperature 0`, and `--seed 0`. Both arms must return byte-identical +completion bytes and identical token counts. + +## Performance protocol + +Build one Release binary and record its SHA-256. Compare only these two arms: + +- Legacy: `VT_ROCM_Q8K_BLOCK=0`. +- Candidate: `VT_ROCM_Q8K_BLOCK=1`. + +Do not compare two binaries. Do not compare an unset arm against an explicit +arm. + +Run one discarded warmup completion for each arm. Then run five measured pairs +in alternating order: `0,1`, `1,0`, `0,1`, `1,0`, and `0,1`. + +Keep the model, prompt, batch, token count, sampling values, device, binary, +and other environment values identical. Hash every measured completion. + +Hold `/home/vikash/gpu.lock` around the complete series. Record the lock +holder, GPU identity, ROCm version, boot identifier, host load, GPU processes, +free memory, utilization, temperature, power, and clock state before and after +each leg. + +Record generation seconds, engine tok/s, peak device memory, and peak host +resident set size for every leg. Record absolute values and candidate-to-legacy +ratios. + +The candidate wins the engine gate only if all these rules hold: + +- All five candidate legs beat their paired legacy legs. +- Median paired engine tok/s improves by at least 2.0%. +- The exact one-sided sign test is 5 of 5, with `p = 0.03125` under equal odds. +- Completion bytes and token counts match in all ten measured legs. +- Device memory and host resident memory do not increase beyond instrument + resolution. +- No leg has invalid contention, clock, thermal, or power evidence. + +The 2.0% threshold is an advance decision. It matches the historical effect +size without accepting a smaller unratified change as signal. + +## Profiler protocol + +Use `rocprofv3 --kernel-trace --stats` for both arms. Use the same binary, +model, prompt, device, sampling values, and environment. + +After a discarded profiler warmup, capture two valid pairs of +`--max-tokens 4` and `--max-tokens 36` for each arm. Subtract each 4-token +capture from its paired 36-token capture, then divide by 32 decode tokens. + +For each arm, record these values: + +- Quantizer calls per token. +- Quantizer microseconds per call. +- Quantizer kernel milliseconds per token. +- Total kernel milliseconds per token. +- Engine tok/s. +- Exact binary SHA-256. +- Full environment and contention evidence. +- Dense and grouped route counter values. +- The selected kernel name and arm engagement. + +Repeat the profile pair if either capture has invalid contention evidence. +Use one tool and one workload for both arms. + +The candidate profile must reduce quantizer kernel milliseconds per token by +more than the within-arm repeat spread. Total kernel milliseconds per token +must not regress. + +llama.cpp at clean pin `b10451` is a secondary floor when it builds and runs on +the same artifact. Record its clean tree, binary hash, ignored tensors, and +exact workload. + +Do not compare llama.cpp Q8_1 bytes with Q8_K bytes. The floor never changes +the local CPU oracle's priority. + +## Negative mutation plan + +A fresh reviewer must use a scratch copy of the immutable implementation head. +The reviewer must rebuild after each mutation and after each restoration. + +Run these mutations separately: + +1. Corrupt the exact-tie selector so a higher index can win. +2. Bypass the shared launcher in the dense production consumer. +3. Bypass the shared launcher in the grouped production consumer. +4. Force the legacy arm when the candidate arm is selected. +5. Corrupt one `qs` byte equality guarantee. +6. Corrupt one `bsums` equality guarantee. +7. Corrupt the `d` byte equality guarantee. + +Each mutation must fail the focused test for the intended reason. The reviewer +must restore the scratch tree byte-for-byte after each mutation. + +After each restoration, rebuild the affected target and rerun the focused +test. The final scratch tree must match the reviewed head exactly. + +## Evidence surfaces + +Keep implementation and review evidence under `/tmp` until the operator +imports the accepted evidence into the owning change. + +Record these items: + +- Base SHA, implementation SHA, and tree SHA. +- Exact configure, build, test, benchmark, and profile commands. +- Command exit status and output path. +- Red-first failure and focused green result. +- Full Release gate result. +- All mutation failures and restoration checks. +- Model and prompt hashes. +- Binary and relevant source hashes. +- Route counters and profiler kernel names. +- Ten measured A/B legs and their raw completion hashes. +- The 4-token and 36-token profiler outputs for both arms. +- Hardware, ROCm, clock, thermal, power, memory, boot, and contention state. + +Do not edit `docs/FEATURES.md`, `docs/BENCHMARKS.md`, or another public document +until the real-checkpoint correctness and performance gates succeed. + +## Risks + +- A block-wide reduction can select the wrong signed maximum on an exact tie. +- HIP wave width can invalidate a warp-only reduction that worked in CUDA. +- Shared memory and barriers can cost more than the serial loop on `gfx1100`. +- Dense routing can work while grouped routing still launches the legacy arm. +- A default flip can hide a dead candidate if the environment policy is wrong. +- Profiler totals can mix prefill and decode without the 4-token subtraction. +- The prompt artifact can be absent even when its Git blob is available. +- A `gfx1100` result does not predict `gfx1200` or `gfx1201` behavior. + +## Owed + +- `gfx1200` runtime validation is `PENDING` external hardware. +- `gfx1201` runtime validation is `PENDING` external hardware, including + validation from @bakon11. +- Public documentation is owed only after successful real-checkpoint gates. +- A llama.cpp floor measurement is owed if a clean pinned build cannot run in + this implementation flow. + +Do not infer either external architecture result from `gfx1100` source review, +compilation, or measurement. + +## Stop conditions + +Stop and keep the default off if any direct 292-byte comparison differs. + +Stop if either production path does not reach the shared launcher. Output +equivalence does not waive this condition. + +Stop if the profiler does not show the candidate kernel and matching arm +engagement. + +Stop if the candidate does not beat the ratified engine threshold. Record the +negative result and restore the candidate product code instead of landing dead +speculative code. + +Stop with `NEEDS_CONTEXT` if an unavailable input changes the correctness or +measurement contract. Do not replace an unavailable input with a new artifact. + +Stop with `NEEDS_DECISION` if a proposed scope change exceeds issue #1876. + +## Acceptance and default decision + +Correctness is accepted only when every direct block is byte-identical to both +the legacy GPU arm and the CPU encoder. Both production counters must prove +route engagement. + +Performance is accepted only when the five-pair engine rule and the profiler +rule both pass on `gfx1100`. All completions must stay byte-identical. + +Keep `VT_ROCM_Q8K_BLOCK` default-off if correctness, reachability, engine speed, +or profile engagement fails. A source-level expectation cannot override a +failed gate. + +A final default-on decision is allowed only after the specification records +exact-byte correctness, dense engagement, grouped engagement, the accepted +five-pair win, and the quantizer profile reduction. + +If the default becomes on, `VT_ROCM_Q8K_BLOCK=0` remains the permanent legacy +A/B arm. All performance and default claims remain scoped to `gfx1100`. From 212cce260592bd3f4c8c6368699e3872252cadf5 Mon Sep 17 00:00:00 2001 From: Vikash Loomba Date: Sat, 29 Aug 2026 08:42:34 -0700 Subject: [PATCH 2/7] spec(BACKEND-ROCM): bind the reviewed Q8_K gate prompt Pull request #2270 owns the reviewed gate prompt, so this branch must not duplicate it. The gate now permits local validation against the immutable reviewed worktree. Publication remains ordered after a maintainer lands the prompt. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:gpt-5 [codex] --- .../specs/rocm-q8k-cooperative-quantizer.md | 32 +++++++++++++++---- 1 file changed, 26 insertions(+), 6 deletions(-) diff --git a/.agents/specs/rocm-q8k-cooperative-quantizer.md b/.agents/specs/rocm-q8k-cooperative-quantizer.md index aa4c2b6e0..84c14768b 100644 --- a/.agents/specs/rocm-q8k-cooperative-quantizer.md +++ b/.agents/specs/rocm-q8k-cooperative-quantizer.md @@ -28,6 +28,11 @@ the repository default and is recorded in the local developer preferences. The specification commit must precede all product and test commits. A fresh implementer must start from this committed specification. +Local correctness and performance evidence can proceed against the reviewed +external prompt source in the real-checkpoint gate. Publication and merge of +#1876 remain ordered after records-only pull request #2270 or an equivalent +maintainer-landed evidence commit that tracks the same reviewed prompt artifact. + ## Scope The implementation has these in-scope changes: @@ -385,11 +390,21 @@ Use this exact model artifact: Use `tools/tg200-prompt.txt` with SHA-256 `e2b801cc6a5739cd317c2f77adfb67040667de524ab60ca64aac39f79c846bba`. -The blob exists on the local `remotes/pr/1936` ref but is absent at this base. -Import only that prompt artifact before measurement and verify its hash. - -Do not import product code from pull request #1936. If the prompt artifact is -still absent, the real-checkpoint gate stays `PENDING`. +Use these prompt sources in order: + +1. After #2270 lands, use the tracked `tools/tg200-prompt.txt` and verify its + SHA-256. +2. Until #2270 lands, local validation can use only + `/home/vikash/vllm.cpp-rocm-launch-evidence/tools/tg200-prompt.txt`. Before + each validation, assert that the worktree has HEAD + `88b1b1bc80c7c7024d64b9ab10626a93ff279a95`, tree + `e2fb82f523d9572a5a3e26437d912c5ea0a5f76c`, an empty + `git status --porcelain=v1`, the tracked path `tools/tg200-prompt.txt`, and + the required SHA-256. + +If neither verified source exists, the real-checkpoint gate stays `PENDING`. +Never substitute prompt text copied from prose. Do not import the prompt into +this branch. No product file from pull request #1936 is imported. The correctness run uses batch 1, greedy decode, `--max-tokens 256`, `--temperature 0`, and `--seed 0`. Both arms must return byte-identical @@ -503,6 +518,8 @@ Record these items: - Full Release gate result. - All mutation failures and restoration checks. - Model and prompt hashes. +- The prompt source path, its tracked commit and tree, and its clean-worktree + assertion when local validation uses the reviewed #2270 worktree. - Binary and relevant source hashes. - Route counters and profiler kernel names. - Ten measured A/B legs and their raw completion hashes. @@ -520,7 +537,7 @@ until the real-checkpoint correctness and performance gates succeed. - Dense routing can work while grouped routing still launches the legacy arm. - A default flip can hide a dead candidate if the environment policy is wrong. - Profiler totals can mix prefill and decode without the 4-token subtraction. -- The prompt artifact can be absent even when its Git blob is available. +- The reviewed prompt source can be absent or fail a provenance assertion. - A `gfx1100` result does not predict `gfx1200` or `gfx1201` behavior. ## Owed @@ -549,6 +566,9 @@ Stop if the candidate does not beat the ratified engine threshold. Record the negative result and restore the candidate product code instead of landing dead speculative code. +Keep the real-checkpoint gate `PENDING` if neither ordered prompt source passes +all provenance assertions. Do not reconstruct the prompt from prose. + Stop with `NEEDS_CONTEXT` if an unavailable input changes the correctness or measurement contract. Do not replace an unavailable input with a new artifact. From 6ae9f62f7af90188d73af520abd0b656a15f6a1c Mon Sep 17 00:00:00 2001 From: Vikash Loomba Date: Sat, 29 Aug 2026 09:08:34 -0700 Subject: [PATCH 3/7] spec(BACKEND-ROCM): bound the Q8_K default to gfx1100 Only gfx1100 has accepted measurement authority. A global unset default could therefore enable the cooperative arm on unvalidated devices. The specification now makes unset selection device-scoped and requires resolver-keyed policy tests before product code. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:gpt-5 [codex] --- .../specs/rocm-q8k-cooperative-quantizer.md | 90 ++++++++++++++----- 1 file changed, 70 insertions(+), 20 deletions(-) diff --git a/.agents/specs/rocm-q8k-cooperative-quantizer.md b/.agents/specs/rocm-q8k-cooperative-quantizer.md index 84c14768b..78d8834ad 100644 --- a/.agents/specs/rocm-q8k-cooperative-quantizer.md +++ b/.agents/specs/rocm-q8k-cooperative-quantizer.md @@ -7,7 +7,8 @@ Row: `BACKEND-ROCM` ## Now The specification is `READY`. Product implementation and tests are `PENDING`. -The first implementation must keep the cooperative arm off by default. +The first implementation must keep the cooperative arm off by default on every +architecture. This commit adds only this specification and the issue ownership row. It makes no matrix or lifecycle change. @@ -215,20 +216,36 @@ cooperative kernel, increments its route witness, and performs the launch. Both production consumers must call this launcher. The launcher must preserve the existing scratch allocation, stream ordering, and error check. +The launcher must resolve default eligibility from the queue's actual device. +It must key any cached architecture result by both the device and resolver. +It must never use one process-global cached eligibility boolean. Follow the +device-aware pattern established by issue #1183, or use an equivalent +per-device, thread-safe cache. The launch hot path must not take a process-wide +mutex. + ### Same-binary selection `VT_ROCM_Q8K_BLOCK` accepts only `0` and `1`: - `0` selects the current serial `QuantizeQ8KK` arm. - `1` selects the cooperative block arm. -- An unset value selects the recorded default. +- An unset value selects the architecture-scoped recorded default. - Any other value refuses by name instead of choosing an arm silently. -The first implementation default is `0`. Read the selection at launch time so -one test process can exercise both values. +The first implementation's unset default is the legacy arm on every +architecture. Read the environment selection at launch time so one test +process can exercise every value. + +If and only if `gfx1100` acceptance succeeds, the final unset policy is allowed +to select the candidate on a queue whose actual device resolves to `gfx1100`. +The unset policy must select legacy on `gfx1200`, `gfx1201`, an unknown +architecture, or an architecture resolution failure. Pending validation on +those architectures never authorizes a global candidate default. -A final default of `1` is allowed only by the acceptance rules in this spec. -If the default becomes `1`, value `0` remains the legacy same-binary arm. +Explicit `0` remains the legacy same-binary arm. Explicit `1` remains a +diagnostic opt-in on any ROCm architecture, even when architecture resolution +fails. Outside measured `gfx1100` hardware, explicit `1` carries no +correctness, performance, or default claim. ### Test-only seams @@ -244,7 +261,9 @@ The seam must include these capabilities: - Launch either arm explicitly into caller-provided Q8_K scratch. - Reset and read dense legacy, dense candidate, grouped legacy, and grouped candidate counters. -- Resolve the environment policy without changing the public ABI. +- Resolve the environment and architecture policy without changing the public + ABI or querying the process environment from the pure policy function. +- Supply a synthetic architecture resolver and device to the policy caller. The explicit-arm hook must not read `VT_ROCM_Q8K_BLOCK`. This separation lets the direct byte test compare both kernels regardless of the default. @@ -304,8 +323,23 @@ The production tests must call `vt::MatmulBTQuant` and `vt::MatmulBTQuantGrouped`. For each environment arm, they must assert the matching route counter changes and the other route counter stays unchanged. -Also test unset, `0`, `1`, and one invalid value. Test the explicit-arm hook -independently from the environment policy. +Add a HIP-free pure policy matrix for unset, `0`, `1`, and one invalid value. +Run every value against synthetic `gfx1100`, `gfx1200`, `gfx1201`, unknown, +and resolution-failure results. Unset must select legacy in every case for the +first implementation. After accepted `gfx1100` default enablement, only the +synthetic `gfx1100` case can change to candidate. Explicit `0` must select +legacy, explicit `1` must select candidate, and invalid must refuse for every +synthetic result. + +Exercise device hops and resolver changes in one process. The matrix must fail +if one device's eligibility is reused for another device or resolver. Test the +explicit-arm hook independently from the environment policy. + +Use an actual resolved `gfx1100` queue for the production-route evidence. With +the environment unset, both production entry points must increment the arm +required by the current recorded default. The first implementation must prove +legacy engagement. An accepted default change must rerun this evidence and +prove candidate engagement. The record-first specification commit has no product behavior to make red. Do not fabricate a red result for this commit. @@ -494,9 +528,11 @@ Run these mutations separately: 2. Bypass the shared launcher in the dense production consumer. 3. Bypass the shared launcher in the grouped production consumer. 4. Force the legacy arm when the candidate arm is selected. -5. Corrupt one `qs` byte equality guarantee. -6. Corrupt one `bsums` equality guarantee. -7. Corrupt the `d` byte equality guarantee. +5. Make the unset policy select the candidate globally after one eligible + device resolves. +6. Corrupt one `qs` byte equality guarantee. +7. Corrupt one `bsums` equality guarantee. +8. Corrupt the `d` byte equality guarantee. Each mutation must fail the focused test for the intended reason. The reviewer must restore the scratch tree byte-for-byte after each mutation. @@ -536,6 +572,9 @@ until the real-checkpoint correctness and performance gates succeed. - Shared memory and barriers can cost more than the serial loop on `gfx1100`. - Dense routing can work while grouped routing still launches the legacy arm. - A default flip can hide a dead candidate if the environment policy is wrong. +- A process-global architecture result can enable an unvalidated device after + one `gfx1100` launch. +- A process-wide cache mutex can serialize the quantizer launch hot path. - Profiler totals can mix prefill and decode without the 4-token subtraction. - The reviewed prompt source can be absent or fail a provenance assertion. - A `gfx1100` result does not predict `gfx1200` or `gfx1201` behavior. @@ -546,6 +585,9 @@ until the real-checkpoint correctness and performance gates succeed. - `gfx1201` runtime validation is `PENDING` external hardware, including validation from @bakon11. - Public documentation is owed only after successful real-checkpoint gates. + If the `gfx1100` unset default changes, the documentation must state the + architecture-scoped unset policy and the diagnostic-only scope of explicit + `1` on other ROCm architectures. - A llama.cpp floor measurement is owed if a clean pinned build cannot run in this implementation flow. @@ -562,6 +604,10 @@ equivalence does not waive this condition. Stop if the profiler does not show the candidate kernel and matching arm engagement. +Stop if the policy matrix lets unset select the candidate on `gfx1200`, +`gfx1201`, unknown, or resolution failure. Stop if policy eligibility uses one +process-global cached boolean or a process-wide mutex on the launch hot path. + Stop if the candidate does not beat the ratified engine threshold. Record the negative result and restore the candidate product code instead of landing dead speculative code. @@ -583,13 +629,17 @@ route engagement. Performance is accepted only when the five-pair engine rule and the profiler rule both pass on `gfx1100`. All completions must stay byte-identical. -Keep `VT_ROCM_Q8K_BLOCK` default-off if correctness, reachability, engine speed, -or profile engagement fails. A source-level expectation cannot override a -failed gate. +Keep the unset `VT_ROCM_Q8K_BLOCK` policy on the legacy arm for every +architecture if correctness, reachability, engine speed, or profile engagement +fails. A source-level expectation cannot override a failed gate. -A final default-on decision is allowed only after the specification records -exact-byte correctness, dense engagement, grouped engagement, the accepted -five-pair win, and the quantizer profile reduction. +A final unset-candidate decision is allowed only for a resolved `gfx1100` +device. It is allowed only after the specification records exact-byte +correctness, actual-device dense engagement, actual-device grouped engagement, +the accepted five-pair win, and the quantizer profile reduction on `gfx1100`. -If the default becomes on, `VT_ROCM_Q8K_BLOCK=0` remains the permanent legacy -A/B arm. All performance and default claims remain scoped to `gfx1100`. +If the `gfx1100` unset default becomes candidate, `VT_ROCM_Q8K_BLOCK=0` +remains the permanent legacy A/B arm. Unset remains legacy on `gfx1200`, +`gfx1201`, unknown, and resolution failure until each architecture gets its own +ratified acceptance authority. All performance and default claims remain +scoped to `gfx1100`. From 70d8d317efb88f28e08762599b49546ccc220597 Mon Sep 17 00:00:00 2001 From: Vikash Loomba Date: Sat, 29 Aug 2026 10:51:07 -0700 Subject: [PATCH 4/7] perf(BACKEND-ROCM): cooperate across each Q8_K activation block The legacy ROCm quantizer serializes each 256-element activation block. Use one 256-thread block while preserving the Q8_K byte encoding. Keep the architecture-scoped unset policy on the legacy arm until the operator accepts the real-checkpoint performance gates. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:gpt-5 [codex] --- scripts/env-doc-allowlist.txt | 1 + src/vt/rocm/rocm_grouped_gemm.hip | 190 ++++++++++- tests/vt/test_backend_cross_device.cpp | 418 +++++++++++++++++++++++++ 3 files changed, 603 insertions(+), 6 deletions(-) diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt index e17b4c2d8..efbd511cb 100644 --- a/scripts/env-doc-allowlist.txt +++ b/scripts/env-doc-allowlist.txt @@ -175,6 +175,7 @@ VT_ROCM_GDN_POSTCONV_CHUNK VT_ROCM_GEMM_COMPUTE VT_ROCM_GEMV VT_ROCM_HIPBLASLT +VT_ROCM_Q8K_BLOCK VT_ROCM_SKINNY VT_SILU_FP4_FAST VT_SPEC_TEST_SELECT_SPIN_MS diff --git a/src/vt/rocm/rocm_grouped_gemm.hip b/src/vt/rocm/rocm_grouped_gemm.hip index 71e47c93e..81570d725 100644 --- a/src/vt/rocm/rocm_grouped_gemm.hip +++ b/src/vt/rocm/rocm_grouped_gemm.hip @@ -18,14 +18,17 @@ #include #include +#include #include #include #include #include #include +#include #include "vt/ops.h" #include "vt/rocm/rocm_device_bind.h" +#include "vt/rocm/rocm_runtime.h" // Block layouts — the single source of truth (ggml-common.h mirrors). #include "vt/cpu/cpu_quant_blocks.h" @@ -177,6 +180,64 @@ __global__ void QuantizeQ8KK(BlockQ8_K* __restrict__ scratch, const void* __rest y.d = 1.0f / iscale; } +// Q8_K (block-per-256-superblock): cuda_quant_dot.cu +// QuantizeQ8KPreqKernel. One thread loads and writes one activation element. +// The reduction retains the signed value and the lowest original index, which +// preserves the legacy first-occurrence rule for opposite-sign absolute ties. +__global__ void QuantizeQ8KCooperativeK(BlockQ8_K* __restrict__ scratch, + const void* __restrict__ a, ActDT adt, + int64_t a_rs, int64_t m, int64_t nsb) { + const int64_t sb = static_cast(blockIdx.x); + const int64_t row = static_cast(blockIdx.y); + if (sb >= nsb || row >= m) return; + const int tid = static_cast(threadIdx.x); + const int64_t elem0 = row * a_rs + sb * kQK_K; + const float value = DLoadAct(a, adt, elem0 + tid); + + __shared__ float reduction_abs[kQK_K]; + __shared__ float reduction_value[kQK_K]; + __shared__ int reduction_index[kQK_K]; + reduction_abs[tid] = fabsf(value); + reduction_value[tid] = value; + reduction_index[tid] = tid; + __syncthreads(); + +#pragma unroll + for (int stride = kQK_K / 2; stride > 0; stride >>= 1) { + if (tid < stride) { + const float other_abs = reduction_abs[tid + stride]; + const int other_index = reduction_index[tid + stride]; + if (other_abs > reduction_abs[tid] || + (other_abs == reduction_abs[tid] && other_index < reduction_index[tid])) { + reduction_abs[tid] = other_abs; + reduction_value[tid] = reduction_value[tid + stride]; + reduction_index[tid] = other_index; + } + } + __syncthreads(); + } + + BlockQ8_K& output = scratch[row * nsb + sb]; + if (reduction_abs[0] == 0.0f) { + if (tid == 0) output.d = 0.0f; + output.qs[tid] = 0; + if (tid < kQK_K / 16) output.bsums[tid] = 0; + return; + } + + const float iscale = -127.0f / reduction_value[0]; + const int quant = DNearestInt(iscale * value); + output.qs[tid] = static_cast(quant < 127 ? quant : 127); + __syncthreads(); + if (tid < kQK_K / 16) { + int sum = 0; +#pragma unroll + for (int ii = 0; ii < 16; ++ii) sum += output.qs[tid * 16 + ii]; + output.bsums[tid] = static_cast(sum); + } + if (tid == 0) output.d = 1.0f / iscale; +} + // ---- dot superblocks (1:1 ports) ---- // Q8_0 x Q8_0: cuda_quant_dot.cu QuantDotGemmQ8_0 — dp4a int core. __device__ inline float DotQ8_0(const BlockQ8_0* wb, const BlockQ8_0* ab) { @@ -593,6 +654,97 @@ void* EnsureQuantScratch(size_t need, hipStream_t s) { return sc.buf; } +enum class Q8KQuantArm : uint8_t { kLegacy = 0, kCandidate = 1 }; +enum class Q8KProductionRoute : uint8_t { kDense = 0, kGrouped = 1 }; +using Q8KArchResolver = std::string (*)(int) noexcept; + +// This constant changes only after the real-checkpoint and profiler gates +// accept gfx1100. The first implementation remains legacy on every device. +constexpr bool kQ8KGfx1100DefaultAccepted = false; + +struct Q8KArchCacheEntry { + int device_index; + Q8KArchResolver resolve; + std::string arch; +}; + +class Q8KArchCache { + public: + std::string Resolve(int device_index, Q8KArchResolver resolve) { + for (const Q8KArchCacheEntry& entry : entries_) { + if (entry.device_index == device_index && entry.resolve == resolve) return entry.arch; + } + const std::string arch = resolve == nullptr ? std::string{} : resolve(device_index); + entries_.push_back(Q8KArchCacheEntry{device_index, resolve, arch}); + return arch; + } + + private: + std::vector entries_; +}; + +std::string ResolveQ8KArch(int device_index, Q8KArchResolver resolve) { + // Match rocm_skinny_gemm_arch.h: cache each device and resolver on the + // calling thread. The launch hot path takes no process-wide mutex. + static thread_local Q8KArchCache cache; + return cache.Resolve(device_index, resolve); +} + +bool Q8KArchIsGfx1100(const std::string& arch) { + constexpr char kStem[] = "gfx1100"; + constexpr size_t kStemLength = sizeof(kStem) - 1; + if (arch.size() < kStemLength || arch.compare(0, kStemLength, kStem) != 0) return false; + if (arch.size() == kStemLength) return true; + const char next = arch[kStemLength]; + return next < '0' || next > '9'; +} + +Q8KQuantArm SelectQ8KQuantArm(const char* env_value, bool gfx1100_default_accepted, + int device_index, Q8KArchResolver resolve) { + if (env_value != nullptr) { + if (std::strcmp(env_value, "0") == 0) return Q8KQuantArm::kLegacy; + if (std::strcmp(env_value, "1") == 0) return Q8KQuantArm::kCandidate; + throw std::runtime_error(std::string("vt rocm: VT_ROCM_Q8K_BLOCK=") + env_value + + " must be 0 or 1"); + } + const std::string arch = ResolveQ8KArch(device_index, resolve); + return gfx1100_default_accepted && Q8KArchIsGfx1100(arch) + ? Q8KQuantArm::kCandidate + : Q8KQuantArm::kLegacy; +} + +std::atomic g_q8k_route_dispatches[2][2]{}; + +void LaunchQ8KQuantArm(BlockQ8_K* scratch, const void* act, ActDT dtype, + int64_t row_stride, int64_t rows, int64_t nsb, hipStream_t stream, + Q8KQuantArm arm, const Q8KProductionRoute* route) { + if (route != nullptr) { + g_q8k_route_dispatches[static_cast(*route)][static_cast(arm)].fetch_add( + 1, std::memory_order_relaxed); + } + if (arm == Q8KQuantArm::kCandidate) { + const dim3 grid(static_cast(nsb), static_cast(rows), 1); + QuantizeQ8KCooperativeK<<>>( + scratch, act, dtype, row_stride, rows, nsb); + Check(hipGetLastError(), "q8_K cooperative quant"); + return; + } + constexpr int kThreads = 128; + const int64_t blocks = (rows * nsb + kThreads - 1) / kThreads; + QuantizeQ8KK<<(blocks), kThreads, 0, stream>>>( + scratch, act, dtype, row_stride, rows, nsb); + Check(hipGetLastError(), "q8_K legacy quant"); +} + +void LaunchQ8KQuantizer(BlockQ8_K* scratch, const void* act, ActDT dtype, + int64_t row_stride, int64_t rows, int64_t nsb, + hipStream_t stream, Q8KProductionRoute route, int device_index) { + const Q8KQuantArm arm = SelectQ8KQuantArm(std::getenv("VT_ROCM_Q8K_BLOCK"), + kQ8KGfx1100DefaultAccepted, device_index, + vt::rocm::DeviceArchName); + LaunchQ8KQuantArm(scratch, act, dtype, row_stride, rows, nsb, stream, arm, &route); +} + // How many dense K-quant dispatches took the cooperative arm. Host-side and // relaxed: nothing reads it on the decode path, and the ROCm decode gate reads // it to prove the call site below is REACHED rather than merely compiled @@ -630,6 +782,34 @@ uint64_t KQuantCoopDispatchCount() { return g_kq_coop_dispatches.load(std::memory_order_relaxed); } +void Q8KQuantizeForTest(Queue& q, void* scratch, const void* act, DType dtype, + int64_t row_stride, int64_t rows, int64_t nsb, bool candidate) { + EnsureQueueDevice(q); + if (dtype != DType::kF32 && dtype != DType::kF16 && dtype != DType::kBF16) { + throw std::runtime_error("vt rocm: Q8KQuantizeForTest requires f32, f16, or bf16"); + } + LaunchQ8KQuantArm(static_cast(scratch), act, ActDtOf(dtype), row_stride, rows, + nsb, static_cast(q.handle), + candidate ? Q8KQuantArm::kCandidate : Q8KQuantArm::kLegacy, nullptr); +} + +bool Q8KCandidateSelectedForTest(const char* env_value, bool gfx1100_default_accepted, + int device_index, + std::string (*resolve)(int) noexcept) { + return SelectQ8KQuantArm(env_value, gfx1100_default_accepted, device_index, resolve) == + Q8KQuantArm::kCandidate; +} + +void Q8KResetRouteDispatchCountsForTest() { + for (auto& route : g_q8k_route_dispatches) + for (std::atomic& count : route) count.store(0, std::memory_order_relaxed); +} + +uint64_t Q8KRouteDispatchCountForTest(bool grouped, bool candidate) { + return g_q8k_route_dispatches[grouped ? 1 : 0][candidate ? 1 : 0].load( + std::memory_order_relaxed); +} + void MatmulBTQuantKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tensor& b) { EnsureQueueDevice(q); const int64_t m = a.shape[0], k = a.shape[1], n = b.shape[0]; @@ -663,9 +843,8 @@ void MatmulBTQuantKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tenso const size_t w_row_bytes = static_cast(nsb) * w_block_bytes; BlockQ8_K* qact = static_cast(EnsureQuantScratch( static_cast(m) * nsb * sizeof(BlockQ8_K), s)); - QuantizeQ8KK<<((m * nsb + 127) / 128), 128, 0, s>>>( - qact, a.data, ActDtOf(a.dtype), a.stride[0], m, nsb); - Check(hipGetLastError(), "q8_K quant"); + LaunchQ8KQuantizer(qact, a.data, ActDtOf(a.dtype), a.stride[0], m, nsb, s, + Q8KProductionRoute::kDense, q.device.index); const int64_t grid = (m * n + kWarpsPerBlock - 1) / kWarpsPerBlock; const int fmt = b.dtype == DType::kQ6_K ? 2 : b.dtype == DType::kQ5_K ? 1 : 0; const int coop = KQuantDecodeCoopWarps(b.dtype, m, nsb); @@ -741,9 +920,8 @@ void MatmulBTQuantGroupedKernelRocm(Queue& q, Tensor& out, const Tensor& act, const size_t w_row_bytes = static_cast(nsb) * w_block_bytes; BlockQ8_K* qact = static_cast(EnsureQuantScratch( static_cast(Pa) * nsb * sizeof(BlockQ8_K), s)); - QuantizeQ8KK<<((Pa * nsb + 127) / 128), 128, 0, s>>>( - qact, act.data, ActDtOf(act.dtype), act.stride[0], Pa, nsb); - Check(hipGetLastError(), "q8_K quant"); + LaunchQ8KQuantizer(qact, act.data, ActDtOf(act.dtype), act.stride[0], Pa, nsb, s, + Q8KProductionRoute::kGrouped, q.device.index); const int64_t grid = (P * n + kWarpsPerBlock - 1) / kWarpsPerBlock; const int fmt = weight.dtype == DType::kQ6_K ? 2 : weight.dtype == DType::kQ5_K ? 1 : 0; auto launch = [&](auto ot) { diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index 5ebee07b4..7a8687bf5 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -25,6 +25,7 @@ // a partial backend is a supported, tested state (src/vt/ops.cpp:104-111). #include +#include #include #include #include @@ -36,9 +37,12 @@ #include "support/test_env.h" // SetEnv/UnsetEnv — MSVC has no setenv (#603) #include "vt/backend.h" +#include "../../src/vt/cpu/cpu_quant_blocks.h" #include "vt/op_provider.h" #include "vt/ops.h" +#include "vt/quant.h" #include "vt/recipes.h" +#include "vt/rocm/rocm_runtime.h" namespace { @@ -2322,8 +2326,422 @@ TEST_CASE("non-grouped keep-quant GEMM (Q8_0/Q4_K/Q5_K/Q6_K) matches the CPU ora namespace vt::rocm { int KQuantDecodeCoopWarps(vt::DType wdt, int64_t m, int64_t nsb); uint64_t KQuantCoopDispatchCount(); +void Q8KQuantizeForTest(vt::Queue& q, void* scratch, const void* act, vt::DType dtype, + int64_t row_stride, int64_t rows, int64_t nsb, bool candidate); +bool Q8KCandidateSelectedForTest(const char* env_value, bool gfx1100_default_accepted, + int device_index, + std::string (*resolve)(int) noexcept); +void Q8KResetRouteDispatchCountsForTest(); +uint64_t Q8KRouteDispatchCountForTest(bool grouped, bool candidate); } // namespace vt::rocm +namespace { + +constexpr size_t kQ8KBlockBytes = sizeof(vt::cpu::BlockQ8_K); +static_assert(kQ8KBlockBytes == 292); + +std::string Q8KMatrixResolver(int device_index) noexcept { + switch (device_index) { + case 0: return "gfx1100"; + case 1: return "gfx1200"; + case 2: return "gfx1201"; + case 3: return "unknown"; + default: return {}; + } +} + +std::string Q8KAlternateResolver(int device_index) noexcept { + if (device_index == 0) return "gfx1200"; + if (device_index == 1) return "gfx1100"; + return {}; +} + +class ScopedQ8KEnv { + public: + explicit ScopedQ8KEnv(const char* value) { + if (const char* old = std::getenv("VT_ROCM_Q8K_BLOCK")) { + had_value_ = true; + old_value_ = old; + } + if (value == nullptr) vllm_test::UnsetEnv("VT_ROCM_Q8K_BLOCK"); + else vllm_test::SetEnv("VT_ROCM_Q8K_BLOCK", value); + } + + ~ScopedQ8KEnv() { + if (had_value_) vllm_test::SetEnv("VT_ROCM_Q8K_BLOCK", old_value_); + else vllm_test::UnsetEnv("VT_ROCM_Q8K_BLOCK"); + } + + ScopedQ8KEnv(const ScopedQ8KEnv&) = delete; + ScopedQ8KEnv& operator=(const ScopedQ8KEnv&) = delete; + + private: + bool had_value_ = false; + std::string old_value_; +}; + +size_t Q8KSourceElementBytes(DType dtype) { + return dtype == DType::kF32 ? sizeof(float) : sizeof(uint16_t); +} + +const char* Q8KDTypeName(DType dtype) { + if (dtype == DType::kF32) return "f32"; + if (dtype == DType::kF16) return "f16"; + return "bf16"; +} + +void StoreQ8KSource(std::vector& bytes, DType dtype, size_t index, float value) { + const size_t offset = index * Q8KSourceElementBytes(dtype); + if (dtype == DType::kF32) { + std::memcpy(bytes.data() + offset, &value, sizeof(value)); + } else { + const uint16_t bits = dtype == DType::kF16 ? vt::F32ToF16(value) : vt::F32ToBF16(value); + std::memcpy(bytes.data() + offset, &bits, sizeof(bits)); + } +} + +float LoadQ8KSource(const std::vector& bytes, DType dtype, size_t index) { + const size_t offset = index * Q8KSourceElementBytes(dtype); + if (dtype == DType::kF32) { + float value = 0.0f; + std::memcpy(&value, bytes.data() + offset, sizeof(value)); + return value; + } + uint16_t bits = 0; + std::memcpy(&bits, bytes.data() + offset, sizeof(bits)); + return dtype == DType::kF16 ? vt::F16ToF32(bits) : vt::BF16ToF32(bits); +} + +enum class Q8KInputKind { + kRandom, + kZero, + kPositiveNegativeWaveTie, + kNegativePositiveWaveTie, + kPositiveNegativeReductionTie, + kNegativePositiveReductionTie, +}; + +struct Q8KInputCase { + Q8KInputKind kind; + const char* name; +}; + +constexpr std::array kQ8KInputCases{{ + {Q8KInputKind::kRandom, "random"}, + {Q8KInputKind::kZero, "all-zero-sentinel"}, + {Q8KInputKind::kPositiveNegativeWaveTie, "positive-negative-wave-tie"}, + {Q8KInputKind::kNegativePositiveWaveTie, "negative-positive-wave-tie"}, + {Q8KInputKind::kPositiveNegativeReductionTie, "positive-negative-reduction-tie"}, + {Q8KInputKind::kNegativePositiveReductionTie, "negative-positive-reduction-tie"}, +}}; + +void FillQ8KInput(std::vector& source, DType dtype, int64_t row_stride, + int64_t rows, int64_t nsb, Q8KInputKind kind) { + std::mt19937 rng(1876u + static_cast(nsb) * 17u + + static_cast(kind) * 101u); + std::uniform_real_distribution random_value(-3.0f, 3.0f); + const int64_t k = nsb * vt::cpu::kQK_K; + for (int64_t row = 0; row < rows; ++row) { + for (int64_t col = 0; col < k; ++col) { + float value = 0.0f; + if (kind == Q8KInputKind::kRandom) { + value = random_value(rng); + } else if (kind != Q8KInputKind::kZero) { + const int bounded = static_cast((col * 29 + row * 11) % 101) - 50; + value = static_cast(bounded) / 64.0f; + } + StoreQ8KSource(source, dtype, static_cast(row * row_stride + col), value); + } + if (kind == Q8KInputKind::kRandom || kind == Q8KInputKind::kZero) continue; + for (int64_t sb = 0; sb < nsb; ++sb) { + int first = 31; + int second = 32; + float first_value = 7.0f; + float second_value = -7.0f; + if (kind == Q8KInputKind::kNegativePositiveWaveTie) { + first_value = -7.0f; + second_value = 7.0f; + } else if (kind == Q8KInputKind::kPositiveNegativeReductionTie) { + first = 127; + second = 128; + } else if (kind == Q8KInputKind::kNegativePositiveReductionTie) { + first = 127; + second = 128; + first_value = -7.0f; + second_value = 7.0f; + } + const int64_t block_start = row * row_stride + sb * vt::cpu::kQK_K; + StoreQ8KSource(source, dtype, static_cast(block_start + first), first_value); + StoreQ8KSource(source, dtype, static_cast(block_start + second), second_value); + } + } +} + +void CheckQ8KBlockBytes(const std::vector& got, + const std::vector& expected, const char* reference, + DType dtype, const char* input_class, int64_t row, int64_t sb, + int64_t nsb) { + const size_t block = static_cast(row * nsb + sb); + const size_t offset = block * kQ8KBlockBytes; + size_t mismatch = kQ8KBlockBytes; + for (size_t byte = 0; byte < kQ8KBlockBytes; ++byte) { + if (got[offset + byte] != expected[offset + byte]) { + mismatch = byte; + break; + } + } + CAPTURE(std::string(Q8KDTypeName(dtype))); + CAPTURE(std::string(input_class)); + CAPTURE(std::string(reference)); + CAPTURE(row); + CAPTURE(sb); + CAPTURE(mismatch); + CHECK_MESSAGE(mismatch == kQ8KBlockBytes, + "ROCm Q8_K byte mismatch: dtype/class/reference/row/sb/byte are captured"); +} + +void CheckOnlyQ8KRouteCounter(bool grouped, bool candidate) { + for (bool observed_grouped : {false, true}) { + for (bool observed_candidate : {false, true}) { + const uint64_t count = + vt::rocm::Q8KRouteDispatchCountForTest(observed_grouped, observed_candidate); + CAPTURE(grouped); + CAPTURE(candidate); + CAPTURE(observed_grouped); + CAPTURE(observed_candidate); + CHECK(count == ((grouped == observed_grouped && candidate == observed_candidate) ? 1u + : 0u)); + } + } +} + +void CheckNoQ8KRouteCounters() { + for (bool grouped : {false, true}) + for (bool candidate : {false, true}) + CHECK(vt::rocm::Q8KRouteDispatchCountForTest(grouped, candidate) == 0); +} + +} // namespace + +TEST_CASE("ROCm Q8_K policy is resolver-keyed and default-off") { + for (bool default_accepted : {false, true}) { + for (int device = 0; device < 5; ++device) { + CAPTURE(default_accepted); + CAPTURE(device); + const bool unset_expected = default_accepted && device == 0; + CHECK(vt::rocm::Q8KCandidateSelectedForTest( + nullptr, default_accepted, device, Q8KMatrixResolver) == unset_expected); + CHECK_FALSE(vt::rocm::Q8KCandidateSelectedForTest( + "0", default_accepted, device, Q8KMatrixResolver)); + CHECK(vt::rocm::Q8KCandidateSelectedForTest( + "1", default_accepted, device, Q8KMatrixResolver)); + CHECK_THROWS_WITH_AS( + vt::rocm::Q8KCandidateSelectedForTest( + "invalid", default_accepted, device, Q8KMatrixResolver), + doctest::Contains("VT_ROCM_Q8K_BLOCK=invalid"), std::runtime_error); + } + } + + // One resolver says device 0 is gfx1100 and device 1 is gfx1200. The next + // resolver reverses those answers. A cache keyed only by device fails here. + CHECK(vt::rocm::Q8KCandidateSelectedForTest(nullptr, true, 0, Q8KMatrixResolver)); + CHECK_FALSE(vt::rocm::Q8KCandidateSelectedForTest(nullptr, true, 1, Q8KMatrixResolver)); + CHECK_FALSE(vt::rocm::Q8KCandidateSelectedForTest(nullptr, true, 0, + Q8KAlternateResolver)); + CHECK(vt::rocm::Q8KCandidateSelectedForTest(nullptr, true, 1, + Q8KAlternateResolver)); +} + +TEST_CASE("ROCm Q8_K cooperative quantizer matches both 292-byte oracles") { + const bool rocm_registered = [] { + for (DeviceType dtype : RegisteredDevices()) + if (dtype == DeviceType::kROCM) return true; + return false; + }(); + if (!rocm_registered) return; + + vt::Backend& rocm = vt::GetBackend(DeviceType::kROCM); + Queue q = rocm.CreateQueue(); + constexpr int64_t kRows = 3; + constexpr int64_t kPadding = 37; + constexpr std::array kSuperblocks{{1, 2, 3, 10, 16}}; + constexpr std::array kDTypes{{DType::kF32, DType::kF16, DType::kBF16}}; + const vt::cpu::FromFloatFn cpu_oracle = vt::cpu::BlockFromFloat(DType::kQ8_K); + REQUIRE(cpu_oracle != nullptr); + vt::rocm::Q8KResetRouteDispatchCountsForTest(); + + for (DType dtype : kDTypes) { + for (const Q8KInputCase& input_case : kQ8KInputCases) { + for (int64_t nsb : kSuperblocks) { + CAPTURE(std::string(Q8KDTypeName(dtype))); + CAPTURE(std::string(input_case.name)); + CAPTURE(nsb); + const int64_t k = nsb * vt::cpu::kQK_K; + const int64_t row_stride = k + kPadding; + const size_t source_bytes = static_cast(kRows * row_stride) * + Q8KSourceElementBytes(dtype); + const size_t scratch_bytes = + static_cast(kRows * nsb) * kQ8KBlockBytes; + std::vector source(source_bytes, 0xD3); + FillQ8KInput(source, dtype, row_stride, kRows, nsb, input_case.kind); + + std::vector cpu_bytes(scratch_bytes, 0xA5); + std::vector cpu_row(static_cast(k)); + for (int64_t row = 0; row < kRows; ++row) { + for (int64_t col = 0; col < k; ++col) { + cpu_row[static_cast(col)] = LoadQ8KSource( + source, dtype, static_cast(row * row_stride + col)); + } + cpu_oracle(cpu_row.data(), + cpu_bytes.data() + static_cast(row * nsb) * kQ8KBlockBytes, k); + } + + DevBufBytes device_source(rocm, q, source_bytes); + DevBufBytes device_legacy(rocm, q, scratch_bytes); + DevBufBytes device_candidate(rocm, q, scratch_bytes); + std::vector sentinel(scratch_bytes, 0xA5); + device_source.Upload(source.data()); + device_legacy.Upload(sentinel.data()); + device_candidate.Upload(sentinel.data()); + vt::rocm::Q8KQuantizeForTest(q, device_legacy.ptr(), device_source.ptr(), dtype, + row_stride, kRows, nsb, false); + vt::rocm::Q8KQuantizeForTest(q, device_candidate.ptr(), device_source.ptr(), dtype, + row_stride, kRows, nsb, true); + std::vector legacy(scratch_bytes); + std::vector candidate(scratch_bytes); + device_legacy.Download(legacy.data()); + device_candidate.Download(candidate.data()); + + for (int64_t row = 0; row < kRows; ++row) { + for (int64_t sb = 0; sb < nsb; ++sb) { + CheckQ8KBlockBytes(candidate, legacy, "legacy-gpu", dtype, input_case.name, + row, sb, nsb); + CheckQ8KBlockBytes(candidate, cpu_bytes, "cpu-block-from-float", dtype, + input_case.name, row, sb, nsb); + } + } + } + } + } + CheckNoQ8KRouteCounters(); + rocm.DestroyQueue(q); +} + +TEST_CASE("ROCm Q8_K dense and grouped production routes select one actual arm") { + const bool rocm_registered = [] { + for (DeviceType dtype : RegisteredDevices()) + if (dtype == DeviceType::kROCM) return true; + return false; + }(); + if (!rocm_registered) return; + REQUIRE(OpAvailable(vt::OpId::kMatmulBTQuant, DeviceType::kROCM)); + REQUIRE(OpAvailable(vt::OpId::kMatmulBTQuantGrouped, DeviceType::kROCM)); + + const std::string actual_arch = vt::rocm::DeviceArchName(0); + if (actual_arch.rfind("gfx1100", 0) != 0) return; + REQUIRE(actual_arch.rfind("gfx1100", 0) == 0); + + constexpr int64_t kM = 2; + constexpr int64_t kP = 2; + constexpr int64_t kN = 4; + constexpr int64_t kE = 3; + constexpr int64_t kK = vt::cpu::kQK_K; + constexpr size_t kQ4BlockBytes = sizeof(vt::cpu::BlockQ4_K); + std::vector dense_act = RandomVec(static_cast(kM * kK), 18760, -0.5f, 0.5f); + std::vector grouped_act = RandomVec(static_cast(kP * kK), 18761, -0.5f, 0.5f); + std::vector dense_weight(static_cast(kN) * kQ4BlockBytes); + std::vector grouped_weight(static_cast(kE * kN) * kQ4BlockBytes); + std::mt19937 rng(18762); + for (uint8_t& byte : dense_weight) byte = static_cast(rng() & 0xff); + for (uint8_t& byte : grouped_weight) byte = static_cast(rng() & 0xff); + auto set_q4_deltas = [](std::vector& bytes) { + for (size_t offset = 0; offset < bytes.size(); offset += kQ4BlockBytes) { + const uint16_t d = vt::F32ToF16(0.0125f); + const uint16_t dmin = vt::F32ToF16(0.0075f); + std::memcpy(bytes.data() + offset, &d, sizeof(d)); + std::memcpy(bytes.data() + offset + sizeof(d), &dmin, sizeof(dmin)); + } + }; + set_q4_deltas(dense_weight); + set_q4_deltas(grouped_weight); + const std::vector expert_ids{2, 0}; + + vt::Backend& rocm = vt::GetBackend(DeviceType::kROCM); + Queue q = rocm.CreateQueue(); + const Device device{DeviceType::kROCM, 0}; + DevBuf dense_act_device(rocm, q, dense_act.size()); + DevBuf grouped_act_device(rocm, q, grouped_act.size()); + DevBufBytes dense_weight_device(rocm, q, dense_weight.size()); + DevBufBytes grouped_weight_device(rocm, q, grouped_weight.size()); + DevBufI32 expert_ids_device(rocm, q, expert_ids.size()); + DevBuf dense_out(rocm, q, static_cast(kM * kN)); + DevBuf grouped_out(rocm, q, static_cast(kP * kN)); + dense_act_device.Upload(dense_act); + grouped_act_device.Upload(grouped_act); + dense_weight_device.Upload(dense_weight.data()); + grouped_weight_device.Upload(grouped_weight.data()); + expert_ids_device.Upload(expert_ids); + + Tensor dense_act_tensor = T2(dense_act_device.ptr(), device, kM, kK); + Tensor dense_weight_tensor = Tensor::Contiguous( + dense_weight_device.ptr(), DType::kQ4_K, device, {kN, kK}); + Tensor dense_out_tensor = T2(dense_out.ptr(), device, kM, kN); + Tensor grouped_act_tensor = T2(grouped_act_device.ptr(), device, kP, kK); + Tensor grouped_weight_tensor = Tensor::Contiguous( + grouped_weight_device.ptr(), DType::kQ4_K, device, {kE * kN, kK}); + Tensor expert_ids_tensor = TI32(expert_ids_device.ptr(), device, kP); + Tensor grouped_out_tensor = T2(grouped_out.ptr(), device, kP, kN); + + auto run_dense = [&] { + vt::MatmulBTQuant(q, dense_out_tensor, dense_act_tensor, dense_weight_tensor); + rocm.Synchronize(q); + }; + auto run_grouped = [&] { + vt::MatmulBTQuantGrouped(q, grouped_out_tensor, grouped_act_tensor, + grouped_weight_tensor, expert_ids_tensor); + rocm.Synchronize(q); + }; + + struct ArmCase { + const char* value; + bool candidate; + }; + const ArmCase arms[] = {{"0", false}, {"1", true}, {nullptr, false}}; + for (const ArmCase& arm : arms) { + CAPTURE(arm.value == nullptr ? std::string("unset") : std::string(arm.value)); + { + ScopedQ8KEnv env(arm.value); + vt::rocm::Q8KResetRouteDispatchCountsForTest(); + run_dense(); + CheckOnlyQ8KRouteCounter(false, arm.candidate); + } + { + ScopedQ8KEnv env(arm.value); + vt::rocm::Q8KResetRouteDispatchCountsForTest(); + run_grouped(); + CheckOnlyQ8KRouteCounter(true, arm.candidate); + } + } + + { + ScopedQ8KEnv env("invalid-production"); + vt::rocm::Q8KResetRouteDispatchCountsForTest(); + CHECK_THROWS_WITH_AS(run_dense(), + doctest::Contains("VT_ROCM_Q8K_BLOCK=invalid-production"), + std::runtime_error); + CheckNoQ8KRouteCounters(); + } + { + ScopedQ8KEnv env("invalid-production"); + vt::rocm::Q8KResetRouteDispatchCountsForTest(); + CHECK_THROWS_WITH_AS(run_grouped(), + doctest::Contains("VT_ROCM_Q8K_BLOCK=invalid-production"), + std::runtime_error); + CheckNoQ8KRouteCounters(); + } + rocm.DestroyQueue(q); +} + TEST_CASE("ROCm Q6_K decode spreads one row's superblocks over several warps") { // Issue #1910: KQuantGemmK strides ONE warp's 32 lanes over nsb = K/256 // superblocks, so a 4096-wide projection (nsb = 16) leaves lanes 16..31 with From be70d25bbc67e3ce4d242c44d1bd4b47cdd52328 Mon Sep 17 00:00:00 2001 From: Vikash Loomba Date: Sat, 29 Aug 2026 13:24:57 -0700 Subject: [PATCH 5/7] test(BACKEND-ROCM): witness the selected Q8_K kernel on device Host route counters proved selection before launch, so a legacy kernel could still satisfy the candidate route assertion. Record one arm-specific count inside the launched kernel and combine it with both production route checks. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:gpt-5 [codex] --- src/vt/rocm/rocm_grouped_gemm.hip | 34 +++++++++++---- tests/vt/test_backend_cross_device.cpp | 57 ++++++++++++++++++++++++++ 2 files changed, 84 insertions(+), 7 deletions(-) diff --git a/src/vt/rocm/rocm_grouped_gemm.hip b/src/vt/rocm/rocm_grouped_gemm.hip index 81570d725..c87f35f68 100644 --- a/src/vt/rocm/rocm_grouped_gemm.hip +++ b/src/vt/rocm/rocm_grouped_gemm.hip @@ -45,6 +45,7 @@ using vt::cpu::kQK8_0; using vt::cpu::kQK_K; enum class ActDT : int { kF32 = 0, kF16 = 1, kBF16 = 2 }; +using Q8KKernelWitnessCount = unsigned long long; inline ActDT ActDtOf(DType dt) { return dt == DType::kF32 ? ActDT::kF32 : dt == DType::kF16 ? ActDT::kF16 : ActDT::kBF16; } @@ -149,9 +150,11 @@ __global__ void QuantizeQ8_0K(BlockQ8_0* __restrict__ scratch, const void* __res // Q8_K (thread-per-256-superblock): cuda_quant_dot.cu QuantizeQ8KKernel. __global__ void QuantizeQ8KK(BlockQ8_K* __restrict__ scratch, const void* __restrict__ a, - ActDT adt, int64_t a_rs, int64_t m, int64_t nsb) { + ActDT adt, int64_t a_rs, int64_t m, int64_t nsb, + Q8KKernelWitnessCount* kernel_witness) { const int64_t t = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; if (t >= m * nsb) return; + if (t == 0 && kernel_witness != nullptr) atomicAdd(&kernel_witness[0], 1ULL); const int64_t i = t / nsb; const int64_t sb = t % nsb; const int64_t elem0 = i * a_rs + sb * kQK_K; @@ -186,11 +189,14 @@ __global__ void QuantizeQ8KK(BlockQ8_K* __restrict__ scratch, const void* __rest // preserves the legacy first-occurrence rule for opposite-sign absolute ties. __global__ void QuantizeQ8KCooperativeK(BlockQ8_K* __restrict__ scratch, const void* __restrict__ a, ActDT adt, - int64_t a_rs, int64_t m, int64_t nsb) { + int64_t a_rs, int64_t m, int64_t nsb, + Q8KKernelWitnessCount* kernel_witness) { const int64_t sb = static_cast(blockIdx.x); const int64_t row = static_cast(blockIdx.y); if (sb >= nsb || row >= m) return; const int tid = static_cast(threadIdx.x); + if (sb == 0 && row == 0 && tid == 0 && kernel_witness != nullptr) + atomicAdd(&kernel_witness[1], 1ULL); const int64_t elem0 = row * a_rs + sb * kQK_K; const float value = DLoadAct(a, adt, elem0 + tid); @@ -714,10 +720,12 @@ Q8KQuantArm SelectQ8KQuantArm(const char* env_value, bool gfx1100_default_accept } std::atomic g_q8k_route_dispatches[2][2]{}; +thread_local Q8KKernelWitnessCount* g_q8k_kernel_execution_witness_for_test = nullptr; void LaunchQ8KQuantArm(BlockQ8_K* scratch, const void* act, ActDT dtype, int64_t row_stride, int64_t rows, int64_t nsb, hipStream_t stream, - Q8KQuantArm arm, const Q8KProductionRoute* route) { + Q8KQuantArm arm, const Q8KProductionRoute* route, + Q8KKernelWitnessCount* kernel_witness) { if (route != nullptr) { g_q8k_route_dispatches[static_cast(*route)][static_cast(arm)].fetch_add( 1, std::memory_order_relaxed); @@ -725,14 +733,14 @@ void LaunchQ8KQuantArm(BlockQ8_K* scratch, const void* act, ActDT dtype, if (arm == Q8KQuantArm::kCandidate) { const dim3 grid(static_cast(nsb), static_cast(rows), 1); QuantizeQ8KCooperativeK<<>>( - scratch, act, dtype, row_stride, rows, nsb); + scratch, act, dtype, row_stride, rows, nsb, kernel_witness); Check(hipGetLastError(), "q8_K cooperative quant"); return; } constexpr int kThreads = 128; const int64_t blocks = (rows * nsb + kThreads - 1) / kThreads; QuantizeQ8KK<<(blocks), kThreads, 0, stream>>>( - scratch, act, dtype, row_stride, rows, nsb); + scratch, act, dtype, row_stride, rows, nsb, kernel_witness); Check(hipGetLastError(), "q8_K legacy quant"); } @@ -742,7 +750,11 @@ void LaunchQ8KQuantizer(BlockQ8_K* scratch, const void* act, ActDT dtype, const Q8KQuantArm arm = SelectQ8KQuantArm(std::getenv("VT_ROCM_Q8K_BLOCK"), kQ8KGfx1100DefaultAccepted, device_index, vt::rocm::DeviceArchName); - LaunchQ8KQuantArm(scratch, act, dtype, row_stride, rows, nsb, stream, arm, &route); + // The ordinary path passes null. Its instrumentation cost is one pointer + // argument and one kernel-leader null check, with no allocation, host copy, + // synchronization, or atomic operation. + LaunchQ8KQuantArm(scratch, act, dtype, row_stride, rows, nsb, stream, arm, &route, + g_q8k_kernel_execution_witness_for_test); } // How many dense K-quant dispatches took the cooperative arm. Host-side and @@ -790,7 +802,8 @@ void Q8KQuantizeForTest(Queue& q, void* scratch, const void* act, DType dtype, } LaunchQ8KQuantArm(static_cast(scratch), act, ActDtOf(dtype), row_stride, rows, nsb, static_cast(q.handle), - candidate ? Q8KQuantArm::kCandidate : Q8KQuantArm::kLegacy, nullptr); + candidate ? Q8KQuantArm::kCandidate : Q8KQuantArm::kLegacy, nullptr, + nullptr); } bool Q8KCandidateSelectedForTest(const char* env_value, bool gfx1100_default_accepted, @@ -810,6 +823,13 @@ uint64_t Q8KRouteDispatchCountForTest(bool grouped, bool candidate) { std::memory_order_relaxed); } +void* Q8KSetKernelExecutionWitnessForTest(void* device_counts) { + Q8KKernelWitnessCount* previous = g_q8k_kernel_execution_witness_for_test; + g_q8k_kernel_execution_witness_for_test = + static_cast(device_counts); + return previous; +} + void MatmulBTQuantKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tensor& b) { EnsureQueueDevice(q); const int64_t m = a.shape[0], k = a.shape[1], n = b.shape[0]; diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index 7a8687bf5..b95c3cb4f 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -2333,6 +2333,7 @@ bool Q8KCandidateSelectedForTest(const char* env_value, bool gfx1100_default_acc std::string (*resolve)(int) noexcept); void Q8KResetRouteDispatchCountsForTest(); uint64_t Q8KRouteDispatchCountForTest(bool grouped, bool candidate); +void* Q8KSetKernelExecutionWitnessForTest(void* device_counts); } // namespace vt::rocm namespace { @@ -2521,6 +2522,52 @@ void CheckNoQ8KRouteCounters() { CHECK(vt::rocm::Q8KRouteDispatchCountForTest(grouped, candidate) == 0); } +constexpr size_t kQ8KKernelWitnessArmCount = 2; +using Q8KKernelWitnessCounts = std::array; + +class ScopedQ8KKernelExecutionWitness { + public: + explicit ScopedQ8KKernelExecutionWitness(void* device_counts) + : previous_(vt::rocm::Q8KSetKernelExecutionWitnessForTest(device_counts)) {} + + ~ScopedQ8KKernelExecutionWitness() { + vt::rocm::Q8KSetKernelExecutionWitnessForTest(previous_); + } + + ScopedQ8KKernelExecutionWitness(const ScopedQ8KKernelExecutionWitness&) = delete; + ScopedQ8KKernelExecutionWitness& operator=(const ScopedQ8KKernelExecutionWitness&) = + delete; + + private: + void* previous_ = nullptr; +}; + +void ResetQ8KKernelExecutionWitness(DevBufBytes& device_counts) { + const Q8KKernelWitnessCounts zero{}; + device_counts.Upload(zero.data()); +} + +Q8KKernelWitnessCounts ReadQ8KKernelExecutionWitness(DevBufBytes& device_counts) { + Q8KKernelWitnessCounts counts{}; + device_counts.Download(counts.data()); + return counts; +} + +void CheckOnlyQ8KKernelExecutionWitness(DevBufBytes& device_counts, bool candidate) { + const Q8KKernelWitnessCounts counts = ReadQ8KKernelExecutionWitness(device_counts); + for (bool observed_candidate : {false, true}) { + CAPTURE(candidate); + CAPTURE(observed_candidate); + CHECK(counts[observed_candidate ? 1 : 0] == + (candidate == observed_candidate ? 1u : 0u)); + } +} + +void CheckNoQ8KKernelExecutionWitness(DevBufBytes& device_counts) { + const Q8KKernelWitnessCounts counts = ReadQ8KKernelExecutionWitness(device_counts); + for (uint64_t count : counts) CHECK(count == 0); +} + } // namespace TEST_CASE("ROCm Q8_K policy is resolver-keyed and default-off") { @@ -2676,6 +2723,7 @@ TEST_CASE("ROCm Q8_K dense and grouped production routes select one actual arm") DevBufI32 expert_ids_device(rocm, q, expert_ids.size()); DevBuf dense_out(rocm, q, static_cast(kM * kN)); DevBuf grouped_out(rocm, q, static_cast(kP * kN)); + DevBufBytes kernel_witness_device(rocm, q, sizeof(Q8KKernelWitnessCounts)); dense_act_device.Upload(dense_act); grouped_act_device.Upload(grouped_act); dense_weight_device.Upload(dense_weight.data()); @@ -2691,6 +2739,7 @@ TEST_CASE("ROCm Q8_K dense and grouped production routes select one actual arm") grouped_weight_device.ptr(), DType::kQ4_K, device, {kE * kN, kK}); Tensor expert_ids_tensor = TI32(expert_ids_device.ptr(), device, kP); Tensor grouped_out_tensor = T2(grouped_out.ptr(), device, kP, kN); + ScopedQ8KKernelExecutionWitness kernel_witness(kernel_witness_device.ptr()); auto run_dense = [&] { vt::MatmulBTQuant(q, dense_out_tensor, dense_act_tensor, dense_weight_tensor); @@ -2712,32 +2761,40 @@ TEST_CASE("ROCm Q8_K dense and grouped production routes select one actual arm") { ScopedQ8KEnv env(arm.value); vt::rocm::Q8KResetRouteDispatchCountsForTest(); + ResetQ8KKernelExecutionWitness(kernel_witness_device); run_dense(); CheckOnlyQ8KRouteCounter(false, arm.candidate); + CheckOnlyQ8KKernelExecutionWitness(kernel_witness_device, arm.candidate); } { ScopedQ8KEnv env(arm.value); vt::rocm::Q8KResetRouteDispatchCountsForTest(); + ResetQ8KKernelExecutionWitness(kernel_witness_device); run_grouped(); CheckOnlyQ8KRouteCounter(true, arm.candidate); + CheckOnlyQ8KKernelExecutionWitness(kernel_witness_device, arm.candidate); } } { ScopedQ8KEnv env("invalid-production"); vt::rocm::Q8KResetRouteDispatchCountsForTest(); + ResetQ8KKernelExecutionWitness(kernel_witness_device); CHECK_THROWS_WITH_AS(run_dense(), doctest::Contains("VT_ROCM_Q8K_BLOCK=invalid-production"), std::runtime_error); CheckNoQ8KRouteCounters(); + CheckNoQ8KKernelExecutionWitness(kernel_witness_device); } { ScopedQ8KEnv env("invalid-production"); vt::rocm::Q8KResetRouteDispatchCountsForTest(); + ResetQ8KKernelExecutionWitness(kernel_witness_device); CHECK_THROWS_WITH_AS(run_grouped(), doctest::Contains("VT_ROCM_Q8K_BLOCK=invalid-production"), std::runtime_error); CheckNoQ8KRouteCounters(); + CheckNoQ8KKernelExecutionWitness(kernel_witness_device); } rocm.DestroyQueue(q); } From 17f4568d11add943919524d5ed42c4fc5cf0eb80 Mon Sep 17 00:00:00 2001 From: Vikash Loomba Date: Sat, 29 Aug 2026 16:13:25 -0700 Subject: [PATCH 6/7] perf(BACKEND-ROCM): select cooperative Q8_K quantization on gfx1100 The accepted gfx1100 correctness, route, engine, and profiler gates justify the architecture-scoped default. Other architectures remain legacy, and strict controls retain same-binary diagnosis and rollback. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:gpt-5 [codex] --- .../specs/rocm-q8k-cooperative-quantizer.md | 138 ++++++++++++++++-- docs/ENVIRONMENT.md | 1 + docs/ROCM.md | 16 ++ docs/USAGE.md | 3 + scripts/env-doc-allowlist.txt | 1 - src/vt/rocm/rocm_grouped_gemm.hip | 6 +- tests/vt/test_backend_cross_device.cpp | 4 +- 7 files changed, 154 insertions(+), 15 deletions(-) diff --git a/.agents/specs/rocm-q8k-cooperative-quantizer.md b/.agents/specs/rocm-q8k-cooperative-quantizer.md index 78d8834ad..6d88b206b 100644 --- a/.agents/specs/rocm-q8k-cooperative-quantizer.md +++ b/.agents/specs/rocm-q8k-cooperative-quantizer.md @@ -6,12 +6,15 @@ Row: `BACKEND-ROCM` ## Now -The specification is `READY`. Product implementation and tests are `PENDING`. -The first implementation must keep the cooperative arm off by default on every -architecture. +The local implementation, gfx1100-only unset default, acceptance evidence, and +fresh review are accepted. Review of immutable head +`be70d25bbc67e3ce4d242c44d1bd4b47cdd52328` found no findings and returned +`PASS`. -This commit adds only this specification and the issue ownership row. It makes -no matrix or lifecycle change. +Issue #1876 and publication remain open and `PENDING` pull request #2270 or an +equivalent maintainer-landed evidence commit. Runtime and default validation on +gfx1200 and gfx1201 remain `PENDING` external hardware. Nothing in this outcome +claims upstream landing. The owning `BACKEND-ROCM` row remains `ACTIVE`. ## Issue ownership @@ -565,6 +568,124 @@ Record these items: Do not edit `docs/FEATURES.md`, `docs/BENCHMARKS.md`, or another public document until the real-checkpoint correctness and performance gates succeed. +## Outcome + +The accepted implementation is head +`be70d25bbc67e3ce4d242c44d1bd4b47cdd52328`, tree +`6cffd44c485374fcca729615ff93373042b6553c`. Its CLI SHA-256 is +`289a76a00fbeaa4af9ee5d8d74dad5a501b4935d59e57d350aa53be512759b0b`, and +its `libvllm.so.0.0.3` SHA-256 is +`13d46bf5f2a6055635e4b4fb3f2c583b5a653d8f3846f0f5d27e5c728c9fc336`. +The review record is `/tmp/rocm-q8k-review2.out`, SHA-256 +`f26e6997e95d95c54d13fd243944baefe09e8901f050daa87d823155ad93d04f`; +it records findings `NONE` and verdict `PASS`. The operator focused gate passed +3/3 cases and 3561/3561 assertions. The full gate retained the same 25 failures +as its immutable base, and deterministic preflight ended directly with `All +gates green.` + +### Correctness and reachability + +The direct gate compared all 292 bytes of every candidate `BlockQ8_K` with +both the legacy GPU encoder and `vt::cpu::BlockFromFloat(DType::kQ8_K)`. It +covered f32, f16, and bf16 inputs; deterministic random values; zero output; +opposite-sign first-absolute-maximum ties in both orders across wave and +reduction boundaries; padded stride; three rows; and +`nsb={1,2,3,10,16}`. Every comparison was byte-identical. + +The public dense and grouped operations proved exclusive host-route identity +and arm-distinct device writes from inside the kernels. Explicit `0`, explicit +`1`, and unset were exercised on the actual gfx1100 device; invalid values left +both witness classes at zero. Five independent reviewer mutations made actual +kernel identity, legacy identity, candidate identity, dense reachability, and +grouped reachability fail. Each byte-for-byte restoration returned the focused +gate to 3561/3561 assertions. + +The accepted real checkpoint is +`Qwen3.5-4B-Q4_K_M-unsloth-e87f1764.gguf`, 2,740,937,888 bytes, from +`unsloth/Qwen3.5-4B-GGUF@e87f176479d0855a907a41277aca2f8ee7a09523`, with +SHA-256 `00fe7986ff5f6b463e62455821146049db6f9313603938a70800d1fb69ef11a4`. +The prompt SHA-256 is +`e2b801cc6a5739cd317c2f77adfb67040667de524ab60ca64aac39f79c846bba`. +Its clean tracked source was head +`88b1b1bc80c7c7024d64b9ab10626a93ff279a95`, tree +`e2fb82f523d9572a5a3e26437d912c5ea0a5f76c`. All ten measured A/B legs +produced 256 tokens and one completion SHA-256, +`769bf8eebae5390db7b6aec5b9ab8e84caa9bc4124f659d77b7240f4494ed245`. + +### Engine and profiler evidence + +The binding same-binary A/B is `/tmp/rocm-q8k-ab2-be70d25`; its summary and +provenance SHA-256 values are +`1a92664c1f95d6710cd83bdd65344d24057e9e5ac4792d2a05b9ecbbe98a9ef2` and +`ec3837756fbcd22c1b76a7219e84bae8a68aabd24e524d3cc2b87946fe04cb9e`. +The candidate won all five pairs. Legacy and candidate medians were 30.943 and +40.316 tok/s, the median paired improvement was 31.377%, and the exact +one-sided sign test was `p=0.03125`. + +The binding two-repeat `rocprofv3 --kernel-trace --stats` subtraction is +`/tmp/rocm-q8k-prof-be70d25`; its summary and provenance SHA-256 values are +`a938c56225460a2f0f50bb58cbffe8591af416401638e29c0d97a2c8b0488d40` and +`692073e885ceb43f5f328802304a5ca24748e8ff084373d709005875f7e2b48c`. +Both arms issued 129 selected quantizer calls per decode token. Median +quantizer time fell from 7.8811 to 0.3426 ms/token, a 95.65% reduction, and +median total kernel time fell from 25.7287 to 18.1923 ms/token. The 7.5385 +ms/token reduction exceeded the 0.0080 ms/token legacy and 0.0043 ms/token +candidate repeat spreads. Each capture contained only the selected quantizer +kernel name, and every 4-token output was the prefix of its matching 36-token +output. + +GDB route counters were `[516,0,0,0]` for explicit `0` and `[0,516,0,0]` for +explicit `1`, ordered as dense legacy, dense candidate, grouped legacy, grouped +candidate. The 516 selected dense dispatches matched the profiler's selected +quantizer calls. The focused production gate supplies the separate grouped +device witness because the dense checkpoint does not exercise grouped routing. + +### Memory, hardware, and rejected evidence + +Peak VRAM maxima were identical at 7,523,020,800 bytes. The raw paired +candidate-minus-legacy deltas were `-8192,+8192,-4096,-8192,-4096` bytes, +inside the 12-16 KiB within-arm sampling spread. Overlapping host RSS samples +also resolved no arm increase. The cooperative arm adds no allocation, +synchronization, or copy, and its default-null witness performs no atomic. +Memory therefore showed no increase beyond repeat resolution. + +The accepted hardware was ROCm device 0, an AMD Radeon RX 7900 XTX resolving +to gfx1100 at PCI `0000:03:00.0` and `/sys/class/drm/card1/device`. Every GPU +run held `/home/vikash/gpu.lock` with `HIP_VISIBLE_DEVICES=0`, +`ROCR_VISIBLE_DEVICES=0`, and the fixed ROCm library path. No external KFD +process appeared. Measured temperature was 62-64 C, sampled peak package power +was 327 W, and no contention, clock, thermal, or power evidence invalidated a +leg. + +The earlier `/tmp/rocm-q8k-ab-be70d25` campaign is rejected. It sampled DRM +card0 while ROCm device 0 resolved to card1; its `INVALID.md` records the +disposition, and no leg was reused. The historical T27 commit +`05455b6a97f1fe60615af105b63b8611ac681873` is also not acceptance evidence: +its 8-thread candidate was dense-only, had no direct scratch-byte suite, and +remained off by default. Its 91.532 to 93.417 tok/s result, 2.06% change, and +five byte-identical winning pairs cannot establish the current result. + +### Default rationale and remaining work + +Unset selects the cooperative arm only when the queue-device resolver returns +gfx1100, including a valid feature-suffix spelling. Exact-byte correctness, +both production-route witnesses, the five-pair engine win, and the profiler +reduction ratify that value. Unset remains legacy for gfx1200, gfx1201, unknown +architectures, and resolution failure because none has its own runtime/default +acceptance. Explicit `VT_ROCM_Q8K_BLOCK=0` remains the permanent legacy A/B +control. Explicit `1` remains a diagnostic candidate override regardless of +architecture so future validation can compare the two arms without another +binary; it makes no correctness, performance, or default claim outside +validated gfx1100. Every other value is refused so a misspelling cannot change +the quantizer silently. + +The clean pinned llama.cpp `b10451` floor remains secondary and cannot alter +the Q8_K byte oracle or this default decision. Gfx1200 and gfx1201 runtime and +default validation remain `PENDING` external hardware, including gfx1201 +validation from @bakon11. Issue #1876, publication, and merge remain ordered +after pull request #2270 or equivalent maintainer-landed prompt evidence. The +broader `BACKEND-ROCM` row therefore remains `ACTIVE`. + ## Risks - A block-wide reduction can select the wrong signed maximum on an exact tie. @@ -584,10 +705,9 @@ until the real-checkpoint correctness and performance gates succeed. - `gfx1200` runtime validation is `PENDING` external hardware. - `gfx1201` runtime validation is `PENDING` external hardware, including validation from @bakon11. -- Public documentation is owed only after successful real-checkpoint gates. - If the `gfx1100` unset default changes, the documentation must state the - architecture-scoped unset policy and the diagnostic-only scope of explicit - `1` on other ROCm architectures. +- The architecture-scoped public control documentation is complete locally; + its publication remains ordered after pull request #2270 or equivalent + maintainer-landed prompt evidence. - A llama.cpp floor measurement is owed if a clean pinned build cannot run in this implementation flow. diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 34f489705..42070cdbd 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -219,6 +219,7 @@ portable/reference path. In normal operation leave them unset. | `VT_QWEN3VL_ATTN_EAGER` | off (flash-tiled attention on) | Set to `1` to force the naive per-key block-reduction attention (`vt::Attention`) in the Qwen3-VL / Qwen3.6-27B vision tower instead of the default flash-tiled kernel. Rollback / A-B knob; token-identical to the default path | | `VT_DEVICE_KV_CACHE` | on (CUDA) | Host-side KV cache instead of the on-device one | | `VT_GPU_SAMPLE` | on (CUDA) | Host-side sampling instead of on-GPU sampling | +| `VT_ROCM_Q8K_BLOCK` | unset (cooperative only on queue-device-resolved `gfx1100`) | Exact `0` selects the permanent legacy A/B arm. Exact `1` forces the cooperative candidate diagnostically. Explicit `1` outside validated `gfx1100` is not a support claim or a default claim. Unset keeps the legacy arm on `gfx1200`, `gfx1201`, unknown architectures, and architecture-resolution failure. Every other value is refused. See [Q8_K activation quantization](ROCM.md#select-q8_k-activation-quantization) | | `VT_GDN_PACKED_DECODE` | on (CUDA GDN) | Unpacked GDN decode path | | `VT_GDN_DECODE_BV` | `32` (CUDA GDN decode experiment) | Exact `16` selects the byte-identical 16-value fused-recurrence tile; unset and every other spelling keep the 32-value schedule. Experimental opt-in; no release or cross-hardware default change | | `VT_GDN_DECODE_SWIZZLE` | `0` (CUDA GDN decode experiment) | Exact `1` enables the shared-memory bank swizzle only for the `BV=16`, `Dv=Dk=128`, eight-lane production geometry; all other values and shapes keep the incumbent layout | diff --git a/docs/ROCM.md b/docs/ROCM.md index ac137d578..062ef6214 100644 --- a/docs/ROCM.md +++ b/docs/ROCM.md @@ -49,6 +49,22 @@ VT_OP_PROVIDER_STATS=1 ./build-hip/examples/vllm-cli \ Run the same command with `--device cpu` to compare greedy tokens. `VT_OP_PROVIDER_STATS=1` reports native operations and CPU fallbacks. +## Select Q8_K activation quantization + +`VT_ROCM_Q8K_BLOCK` selects the Q8_K activation quantizer used by ROCm dense +and grouped keep-quant operations. When it is unset, a queue device that +resolves to `gfx1100`, including a valid feature-suffix spelling, uses the +cooperative arm. Unset uses the legacy arm on `gfx1200`, `gfx1201`, an unknown +architecture, or architecture-resolution failure. This policy does not imply +runtime validation on `gfx1200` or `gfx1201`. + +The parser is strict. Set `VT_ROCM_Q8K_BLOCK=0` for the permanent legacy A/B +arm, or set `VT_ROCM_Q8K_BLOCK=1` to force the cooperative diagnostic arm +regardless of architecture. Explicit `1` is diagnostic-only outside validated +`gfx1100`; it does not make a correctness, performance, or default claim for +another architecture. Any value other than `0` or `1` refuses the operation +instead of choosing an arm silently. + ## Understand fallback behavior An integrated GPU can use the CPU reference tier when the backend reports diff --git a/docs/USAGE.md b/docs/USAGE.md index f0fb662f0..964818a52 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -480,6 +480,9 @@ skips with that refusal quoted. - Read the matching model or task guide before you add model-specific flags. - If startup fails, use the exact error text to find the refused file, option, operation, or checkpoint arm in the focused guides. +- On ROCm, see [Q8_K activation quantization](ROCM.md#select-q8_k-activation-quantization) + to inspect or override the architecture-scoped quantizer during + troubleshooting. - On ROCm, GGUF mixture-of-experts checkpoints compute on the quantized expert blocks (Q8_0, Q4_K, Q5_K, Q6_K) instead of being dequantized to bf16 at load time. diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt index efbd511cb..e17b4c2d8 100644 --- a/scripts/env-doc-allowlist.txt +++ b/scripts/env-doc-allowlist.txt @@ -175,7 +175,6 @@ VT_ROCM_GDN_POSTCONV_CHUNK VT_ROCM_GEMM_COMPUTE VT_ROCM_GEMV VT_ROCM_HIPBLASLT -VT_ROCM_Q8K_BLOCK VT_ROCM_SKINNY VT_SILU_FP4_FAST VT_SPEC_TEST_SELECT_SPIN_MS diff --git a/src/vt/rocm/rocm_grouped_gemm.hip b/src/vt/rocm/rocm_grouped_gemm.hip index c87f35f68..928f4aa42 100644 --- a/src/vt/rocm/rocm_grouped_gemm.hip +++ b/src/vt/rocm/rocm_grouped_gemm.hip @@ -664,9 +664,9 @@ enum class Q8KQuantArm : uint8_t { kLegacy = 0, kCandidate = 1 }; enum class Q8KProductionRoute : uint8_t { kDense = 0, kGrouped = 1 }; using Q8KArchResolver = std::string (*)(int) noexcept; -// This constant changes only after the real-checkpoint and profiler gates -// accept gfx1100. The first implementation remains legacy on every device. -constexpr bool kQ8KGfx1100DefaultAccepted = false; +// The accepted real-checkpoint and profiler gates apply only to gfx1100. +// Explicit arms and every other architecture remain governed by the selector. +constexpr bool kQ8KGfx1100DefaultAccepted = true; struct Q8KArchCacheEntry { int device_index; diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index b95c3cb4f..726d2e7bf 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -2570,7 +2570,7 @@ void CheckNoQ8KKernelExecutionWitness(DevBufBytes& device_counts) { } // namespace -TEST_CASE("ROCm Q8_K policy is resolver-keyed and default-off") { +TEST_CASE("ROCm Q8_K policy is resolver-keyed and architecture-scoped") { for (bool default_accepted : {false, true}) { for (int device = 0; device < 5; ++device) { CAPTURE(default_accepted); @@ -2755,7 +2755,7 @@ TEST_CASE("ROCm Q8_K dense and grouped production routes select one actual arm") const char* value; bool candidate; }; - const ArmCase arms[] = {{"0", false}, {"1", true}, {nullptr, false}}; + const ArmCase arms[] = {{"0", false}, {"1", true}, {nullptr, true}}; for (const ArmCase& arm : arms) { CAPTURE(arm.value == nullptr ? std::string("unset") : std::string(arm.value)); { From 39398fece25282d358250de09f60ed7ef740e25d Mon Sep 17 00:00:00 2001 From: Vikash Loomba Date: Mon, 31 Aug 2026 19:51:47 -0700 Subject: [PATCH 7/7] record(BACKEND-ROCM): retain issue 1876 under the owning spec Pull request #2270 supplies the tracked prompt prerequisite. Issue #1876 remains open for gfx1200 and gfx1201 acceptance. The authenticated account cannot add the live `Row:` line. The spec's `## Owed` link supplies `BACKEND-ROCM` ownership. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:gpt-5 [codex] --- .../specs/rocm-q8k-cooperative-quantizer.md | 74 +++++++++---------- 1 file changed, 36 insertions(+), 38 deletions(-) diff --git a/.agents/specs/rocm-q8k-cooperative-quantizer.md b/.agents/specs/rocm-q8k-cooperative-quantizer.md index 6d88b206b..eac67822f 100644 --- a/.agents/specs/rocm-q8k-cooperative-quantizer.md +++ b/.agents/specs/rocm-q8k-cooperative-quantizer.md @@ -11,18 +11,22 @@ fresh review are accepted. Review of immutable head `be70d25bbc67e3ce4d242c44d1bd4b47cdd52328` found no findings and returned `PASS`. -Issue #1876 and publication remain open and `PENDING` pull request #2270 or an -equivalent maintainer-landed evidence commit. Runtime and default validation on -gfx1200 and gfx1201 remain `PENDING` external hardware. Nothing in this outcome -claims upstream landing. The owning `BACKEND-ROCM` row remains `ACTIVE`. +Pull request [#2270](https://github.com/mudler/vllm.cpp/pull/2270) landed as +`5575f689f`. Its tracked `tools/tg200-prompt.txt` satisfies the publication +ordering prerequisite. The Q8_K implementation in this branch is not yet +upstream. Runtime and default validation on gfx1200 and gfx1201 remain +`PENDING` external hardware. The owning `BACKEND-ROCM` row remains `ACTIVE`. ## Issue ownership -Issue #1876 owns this `BACKEND-ROCM` performance slice. The issue-index row -names `BACKEND-ROCM` as the owner and `perf` as the kind. +Live issue #1876 remains `OPEN` and lacks a canonical `Row:` line. Under the +current derived policy, the full issue link in this spec's `## Owed` section +supplies `BACKEND-ROCM` ownership. The retired issue index and its historical +`perf` classification do not supply current ownership. -The issue stays open until the accepted implementation and its evidence land. -A negative result also leaves the issue open. +The proposed pull request lands the accepted gfx1100 slice and addresses #1876 +without closing it. Issue #1876 remains open because gfx1200 and gfx1201 +runtime and default acceptance are still owed. ## Git integration @@ -32,10 +36,10 @@ the repository default and is recorded in the local developer preferences. The specification commit must precede all product and test commits. A fresh implementer must start from this committed specification. -Local correctness and performance evidence can proceed against the reviewed -external prompt source in the real-checkpoint gate. Publication and merge of -#1876 remain ordered after records-only pull request #2270 or an equivalent -maintainer-landed evidence commit that tracks the same reviewed prompt artifact. +The accepted pre-landing evidence used the reviewed external prompt source in +the real-checkpoint gate. Pull request #2270 later tracked the same reviewed +prompt artifact as `tools/tg200-prompt.txt`. Commit `5575f689f` satisfies the +publication ordering prerequisite for issue #1876. ## Scope @@ -427,21 +431,17 @@ Use this exact model artifact: Use `tools/tg200-prompt.txt` with SHA-256 `e2b801cc6a5739cd317c2f77adfb67040667de524ab60ca64aac39f79c846bba`. -Use these prompt sources in order: - -1. After #2270 lands, use the tracked `tools/tg200-prompt.txt` and verify its - SHA-256. -2. Until #2270 lands, local validation can use only - `/home/vikash/vllm.cpp-rocm-launch-evidence/tools/tg200-prompt.txt`. Before - each validation, assert that the worktree has HEAD - `88b1b1bc80c7c7024d64b9ab10626a93ff279a95`, tree - `e2fb82f523d9572a5a3e26437d912c5ea0a5f76c`, an empty - `git status --porcelain=v1`, the tracked path `tools/tg200-prompt.txt`, and - the required SHA-256. - -If neither verified source exists, the real-checkpoint gate stays `PENDING`. -Never substitute prompt text copied from prose. Do not import the prompt into -this branch. No product file from pull request #1936 is imported. +Pull request #2270 landed as `5575f689f`, so use the tracked +`tools/tg200-prompt.txt` and verify its SHA-256. The accepted pre-landing +validation used only +`/home/vikash/vllm.cpp-rocm-launch-evidence/tools/tg200-prompt.txt`. Its +worktree had HEAD `88b1b1bc80c7c7024d64b9ab10626a93ff279a95`, tree +`e2fb82f523d9572a5a3e26437d912c5ea0a5f76c`, an empty +`git status --porcelain=v1`, the tracked prompt path, and the required SHA-256. + +If the tracked source is absent or its hash differs, a real-checkpoint rerun +stays `PENDING`. Never substitute prompt text copied from prose or another +worktree. No product file from pull request #1936 was imported. The correctness run uses batch 1, greedy decode, `--max-tokens 256`, `--temperature 0`, and `--seed 0`. Both arms must return byte-identical @@ -557,8 +557,8 @@ Record these items: - Full Release gate result. - All mutation failures and restoration checks. - Model and prompt hashes. -- The prompt source path, its tracked commit and tree, and its clean-worktree - assertion when local validation uses the reviewed #2270 worktree. +- The pre-landing prompt source path, its tracked commit and tree, and its + clean-worktree assertion. - Binary and relevant source hashes. - Route counters and profiler kernel names. - Ten measured A/B legs and their raw completion hashes. @@ -682,9 +682,10 @@ the quantizer silently. The clean pinned llama.cpp `b10451` floor remains secondary and cannot alter the Q8_K byte oracle or this default decision. Gfx1200 and gfx1201 runtime and default validation remain `PENDING` external hardware, including gfx1201 -validation from @bakon11. Issue #1876, publication, and merge remain ordered -after pull request #2270 or equivalent maintainer-landed prompt evidence. The -broader `BACKEND-ROCM` row therefore remains `ACTIVE`. +validation from @bakon11. Pull request #2270 landed as `5575f689f`, so the +tracked prompt and publication ordering prerequisite are satisfied. The Q8_K +implementation in this pull request is not yet upstream. The broader +`BACKEND-ROCM` row therefore remains `ACTIVE`. ## Risks @@ -702,12 +703,9 @@ broader `BACKEND-ROCM` row therefore remains `ACTIVE`. ## Owed -- `gfx1200` runtime validation is `PENDING` external hardware. -- `gfx1201` runtime validation is `PENDING` external hardware, including - validation from @bakon11. -- The architecture-scoped public control documentation is complete locally; - its publication remains ordered after pull request #2270 or equivalent - maintainer-landed prompt evidence. +- [Issue #1876](https://github.com/mudler/vllm.cpp/issues/1876) owns gfx1200 and + gfx1201 runtime and default validation. Both remain `PENDING` external + hardware, and the gfx1201 scope includes validation from @bakon11. - A llama.cpp floor measurement is owed if a clean pinned build cannot run in this implementation flow.