diff --git a/.agents/specs/rocm-q8k-cooperative-quantizer.md b/.agents/specs/rocm-q8k-cooperative-quantizer.md new file mode 100644 index 000000000..eac67822f --- /dev/null +++ b/.agents/specs/rocm-q8k-cooperative-quantizer.md @@ -0,0 +1,763 @@ +# Cooperative ROCm Q8_K activation quantizer + +Issue: [#1876](https://github.com/mudler/vllm.cpp/issues/1876) + +Row: `BACKEND-ROCM` + +## Now + +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`. + +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 + +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 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 + +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. + +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 + +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. + +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 architecture-scoped recorded default. +- Any other value refuses by name instead of choosing an arm silently. + +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. + +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 + +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 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. + +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. + +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. + +## 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`. +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 +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. 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. + +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. +- 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. +- 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. + +## 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. 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 + +- 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. +- 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. + +## Owed + +- [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. + +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 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. + +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. + +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 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 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 `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`. diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index c42a89280..c4d2523ca 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -220,6 +220,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 2c76dfb86..b62a5f7be 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -510,6 +510,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/src/vt/rocm/rocm_grouped_gemm.hip b/src/vt/rocm/rocm_grouped_gemm.hip index 4e617d43a..5a2b67995 100644 --- a/src/vt/rocm/rocm_grouped_gemm.hip +++ b/src/vt/rocm/rocm_grouped_gemm.hip @@ -24,10 +24,12 @@ #include #include #include +#include #include "vt/ops.h" #include "vt/rocm/rocm_device_bind.h" #include "vt/rocm/rocm_f16_codec.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" @@ -44,6 +46,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; } @@ -106,9 +109,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; @@ -137,6 +142,67 @@ __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, + 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); + + __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) { @@ -566,6 +632,103 @@ 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; + +// 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; + 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]{}; +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, + Q8KKernelWitnessCount* kernel_witness) { + 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, 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, kernel_witness); + 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); + // 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 // 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 @@ -626,6 +789,42 @@ 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, + 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* 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]; @@ -659,9 +858,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); @@ -738,9 +936,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 4c596bfd5..80199a4bc 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -27,6 +27,7 @@ #include #include +#include #include #include #include @@ -38,9 +39,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 { @@ -2558,8 +2562,479 @@ TEST_CASE("keep-quant Q6_K GEMM runs at the production launch geometry") { 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); +void* Q8KSetKernelExecutionWitnessForTest(void* device_counts); } // 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); +} + +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 architecture-scoped") { + 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)); + 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()); + 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); + ScopedQ8KKernelExecutionWitness kernel_witness(kernel_witness_device.ptr()); + + 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, true}}; + for (const ArmCase& arm : arms) { + CAPTURE(arm.value == nullptr ? std::string("unset") : std::string(arm.value)); + { + 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); +} + 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