diff --git a/.agents/issues/KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3/ISSUE-LOCAL-01M2F0PQWGSCXG0N4951NF9DPZ.md b/.agents/issues/KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3/ISSUE-LOCAL-01M2F0PQWGSCXG0N4951NF9DPZ.md new file mode 100644 index 000000000..45e887273 --- /dev/null +++ b/.agents/issues/KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3/ISSUE-LOCAL-01M2F0PQWGSCXG0N4951NF9DPZ.md @@ -0,0 +1,19 @@ +ID: ISSUE-LOCAL-01M2F0PQWGSCXG0N4951NF9DPZ +Title: Enable the existing quantized WMMA prefill kernels on gfx1100 +Row: KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3 +State: CLOSED +Kind: feature +GitHub: - +Mirror: PENDING +Availability: FULL +Created: 2026-09-13 +Updated: 2026-09-13 +Closed: 2026-09-13 + +## Problem + +Current main compiles and dispatches its generic rocWMMA Q4_K and Q6_K prefill kernels only on gfx1200/gfx1201. The installed rocWMMA 2.2.1 implements the same 16x16x16 signed-int8 operation on gfx1100. Verify admission-only reuse on physical gfx1100, preserve the attention architecture guard, and prove production reachability, numerical correctness, and same-binary prefill performance before accepting the default. The user selects RDNA3 WMMA as the next work and explicitly deprioritizes the unrelated #2773 characterization campaign. + +## Resolution + +13 September 2026: The landing change admits physical gfx1100 through the existing Q4_K and Q6_K WMMA dispatch. Independent review and operator verification pass the architecture, 240 original matrix, and public 1024-logit gates. Eight exact native prompt arrays and all generated IDs in six native processes match both task-pinned oracles. Native traces contain 152 WMMA calls per prefill and zero during decode. The observed prefill ratio is 1.2891 under dynamic clocks. The report and linked archive at docs/bench-evidence/rocm-rdna3-quant-wmma/README.md retain every value, ratio, failed attempt, and measurement limit. ISSUE-LOCAL-01M2F4WCD6ZK5VH5S8TF83APD6 remains open for whole-model floors and accepted clock attribution. diff --git a/.agents/issues/KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3/ISSUE-LOCAL-01M2F51K16TDAC95DM1VEKBBCK.md b/.agents/issues/KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3/ISSUE-LOCAL-01M2F51K16TDAC95DM1VEKBBCK.md new file mode 100644 index 000000000..74be88b95 --- /dev/null +++ b/.agents/issues/KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3/ISSUE-LOCAL-01M2F51K16TDAC95DM1VEKBBCK.md @@ -0,0 +1,19 @@ +ID: ISSUE-LOCAL-01M2F51K16TDAC95DM1VEKBBCK +Title: Classify the RDNA3 WMMA evidence in the permitted per-run layout +Row: KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3 +State: CLOSED +Kind: bug +GitHub: - +Mirror: PENDING +Availability: FULL +Created: 2026-09-13 +Updated: 2026-09-13 +Closed: 2026-09-13 + +## Problem + +The row stores operator and review receipts in nested directories. check-pr-size.py rejects those paths because BENCH_EVIDENCE_RUN permits one run directory followed by a filename. Flatten the receipt filenames, preserve original bytes and source locations, update local links and manifests, and validate the exact integrated diff without changing the checker. + +## Resolution + +13 September 2026: The landing change flattens all RDNA3 WMMA receipts into the permitted per-run directory. evidence-path-map.json preserves original filenames, byte hashes, and scoped manifest adaptations. Local links and copy manifests resolve the flat paths. Source locations and original measurements remain unchanged. The checker itself is unchanged. Classification of the final integrated diff is recorded with the final records verification receipt. diff --git a/.agents/issues/KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3/ISSUE-LOCAL-01M2GNY58NHBK3D4JQQ738M6GR.md b/.agents/issues/KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3/ISSUE-LOCAL-01M2GNY58NHBK3D4JQQ738M6GR.md new file mode 100644 index 000000000..c791eeecd --- /dev/null +++ b/.agents/issues/KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3/ISSUE-LOCAL-01M2GNY58NHBK3D4JQQ738M6GR.md @@ -0,0 +1,19 @@ +ID: ISSUE-LOCAL-01M2GNY58NHBK3D4JQQ738M6GR +Title: Move RDNA3 WMMA raw evidence out of the pull request diff +Row: KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3 +State: CLOSED +Kind: documentation +GitHub: - +Mirror: PENDING +Availability: FULL +Created: 2026-09-14 +Updated: 2026-09-14 +Closed: 2026-09-14 + +## Problem + +The RDNA3 WMMA pull request includes 236 evidence files and 42,704 evidence lines for a two-file runtime admission change. The developer requested a compact report with the complete raw evidence preserved in a separately accessible, checksum-verified archive. Keep the same PR, implementation, tests, model results, and all qualifications. + +## Resolution + +14 September 2026: The linked release archive preserves all 236 original evidence files from d6e40c91f634a041c873c7a04516d55c4d05772a. The operator verified anonymous download, the archive checksum, and every original byte before replacement. docs/bench-evidence/rocm-rdna3-quant-wmma/README.md retains the source revision, archive identity, extraction procedure, correctness results, and performance qualifications. The implementation, tests, and validation harness remain byte-identical. Altered-archive, altered-file, and missing-file integrity checks fail as intended and pass after restoration. diff --git a/.agents/issues/_owed/ISSUE-LOCAL-01M2F4WCD6ZK5VH5S8TF83APD6.md b/.agents/issues/_owed/ISSUE-LOCAL-01M2F4WCD6ZK5VH5S8TF83APD6.md new file mode 100644 index 000000000..957519173 --- /dev/null +++ b/.agents/issues/_owed/ISSUE-LOCAL-01M2F4WCD6ZK5VH5S8TF83APD6.md @@ -0,0 +1,21 @@ +ID: ISSUE-LOCAL-01M2F4WCD6ZK5VH5S8TF83APD6 +Title: Resolve the remaining gfx1100 Qwen3.5-4B Q4_K_M performance gaps +Row: - +State: OPEN +Kind: performance +GitHub: - +Mirror: PENDING +Availability: FULL +Created: 2026-09-13 +Updated: 2026-09-13 +Closed: - + +## Problem + +The completed RDNA3 WMMA experiment preserves exact tokens and improves its recorded same-binary prefill metric. Native whole-workload throughput and decode remain below the task-pinned primary and llama.cpp comparisons. Tail latency, sampled memory against llama.cpp, matching oracle traces, comparable timing windows, and accepted clock attribution remain unresolved. Establish the missing measurements and close every below-floor axis on the identical retained workload. Do not treat the architecture-admission result as full-model parity or as a performance ceiling. + +The next traceable candidates are Gated DeltaNet prefill and scalar remainder work at input lengths 183 and 174. Native `src/vt/rocm/rocm_gdn_scan.hip:55-104` runs a serial token recurrence for prefill and decode, and the native trace contains `GdnScanK`. At the task-pinned primary revision, `vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py:240-324` selects a chunk-prefill backend whose native method calls `fla_chunk_gated_delta_rule`. These source differences are hypotheses. Matching generated oracle traces and accepted clock windows must establish the executed paths and their costs before ranking them. + +## Resolution + +- diff --git a/.agents/specs/kernel-quant-ciq-gemm-rocm.md b/.agents/specs/kernel-quant-ciq-gemm-rocm.md index 0308ce9f0..3ca3b2df0 100644 --- a/.agents/specs/kernel-quant-ciq-gemm-rocm.md +++ b/.agents/specs/kernel-quant-ciq-gemm-rocm.md @@ -247,8 +247,8 @@ scale layout through the tile op. ## Owed -- gfx1100 (RDNA3) WMMA tile: `VikashLoomba`'s audit on #2109, a separate row - and spec. +- gfx1100 (RDNA3) WMMA admission: measured and completed by this change in + [KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3](rocm-rdna3-quant-wmma.md). - gfx1151 (RDNA3.5) WMMA tile: needs Strix Halo hardware to verify; a separate row and spec. - hipBLASLt per-superblock scale support: unmeasured; recorded as an open diff --git a/.agents/specs/rocm-rdna3-quant-wmma.md b/.agents/specs/rocm-rdna3-quant-wmma.md new file mode 100644 index 000000000..78c742228 --- /dev/null +++ b/.agents/specs/rocm-rdna3-quant-wmma.md @@ -0,0 +1,227 @@ +# Reuse quantized WMMA prefill on gfx1100 + +Row: `KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3`. +Parent: `KERNEL-QUANT-CIQ-GEMM-ROCM`, which retains its RDNA4 work. +Issue: `ISSUE-LOCAL-01M2F0PQWGSCXG0N4951NF9DPZ`. +Base: `cef9f821632eb33156bb7d24a0a85eeab7f7484a`. +Integration: one pull request, selected by the developer on 13 September 2026. +Commit this specification before implementation. + +## Now + +`DONE` for architecture admission, 13 September 2026. +Implementation `c3fe98ba6c55ce71e75746e1b944a27640464e0f` passes G1-G3 and +fresh mutation review. Both executing tile bodies remain unchanged. +The operator independently reproduces the hardware, public, and 240 original +primary fixture gates. G5 retains its explicitly recorded baseline resource skips. +G4 executes both task-pinned oracles and six native processes with exact input +and generated token IDs. Native traces prove prefill-only WMMA dispatch. +The observed prefill ratio is 1.2891 with WMMA enabled. +Full-model floors remain failing, and accepted clock attribution remains pending +under the dedicated issue in `## Owed`. +[Evidence report and archive](../../docs/bench-evidence/rocm-rdna3-quant-wmma/README.md) +retain commands, results, mutations, and measurement limits. + +## Scope + +Enable the existing 16 by 16 quantized prefill tiles on gfx1100 through the +ordinary `vt::MatmulBTQuant` dispatch. Preserve Q4_K and Q6_K block decoding, +Q8_K activation quantization, output dtypes, scales, and remainder handling. +Preserve the existing `VT_ROCM_QUANT_WMMA=0` scalar control. +Keep admission bounded to gfx1100 and the already admitted gfx1200/gfx1201. +Do not enable other gfx11 devices without their own evidence. + +This is an architecture-admission change, not a new model or arithmetic port. +The user selected it independently of the broader Qwen state characterization. +Do not add graph activation, attention changes, new formats, grouped expert +tiles, kernel redesign, or a general numerical-characterization dependency. + +## Sources and hypothesis + +The parent [kernel spec](kernel-quant-ciq-gemm-rocm.md) records the existing +implementation and its oracle hierarchy. Current local source uses generic +rocWMMA fragments in `src/vt/rocm/rocm_grouped_gemm.hip::KQuantGemmKWmmaQ6K` +and `src/vt/rocm/rocm_grouped_gemm.hip::KQuantGemmKWmmaQ4K`. +Each operation loads fragments through `load_matrix_sync`, computes with +`mma_sync`, and stores a row-major tile through `store_matrix_sync`. +The scalar epilogue reads that stored tile, not the hardware register layout. + +Installed rocWMMA 2.2.1 defines the signed-int8 gfx11 specialization in +`internal/wmma_impl.hpp:602-634`. It selects +`__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32` and an int32 accumulator. +Its `internal/config.hpp:222-226` requires wave32 and block dimension 16. +The existing kernel uses both. The installed compiler is HIP 7.15.26333, +Clang 23 at revision `8f497e0992fb7513f7f78a6f6b6f1056c375e961`. +Record the compiler and header hashes with actual build evidence. + +Pinned llama.cpp `b10451`, revision `10bf611e5`, confirms the distinct RDNA3 +operand layout in `ggml/src/ggml-cuda/mma.cuh::get_input_data_layout` and its +RDNA3 integer WMMA specialization. rocWMMA owns that difference here. +The source is available at the exact revision through the upstream repository. +The developer explicitly selected current upstream main/master for this +experiment on 13 September 2026, instead of the previously installed oracles. +The fetched task snapshots are vLLM +`39545e475d3627287ff69c25465dc0bd405f67e1`, llama.cpp +`093a2f86c3e37c54fa3e1f9efb17b304f3433abd`, and GGUF plugin +`d4c1f0d082fc7cd4350da56689109a01c1f29d6c`. +The plugin's current main is unchanged from the repository pin. +Build and execute isolated runtimes at these revisions before accepting their +results. Recheck the executing source chain there. This task-specific developer +direction does not claim reconciliation of unrelated repository goldens. + +## Design + +1. Add an additive quant-specific architecture policy header. Its predicate + accepts gfx1100 plus the current gfx1200/gfx1201 predicate, with valid + feature suffixes and rejection of malformed or unmeasured names. +2. Admit `__gfx1100__` at the quantized translation unit's rocWMMA include + and compile guard. Use the quant-specific predicate at its runtime gate. +3. Keep `GcnArchNameIsGfx12PrefillWmma` unchanged. Attention uses that predicate + but its gfx1100 device body is still excluded. Widening it could launch + an empty attention kernel. +4. Reuse both tile bodies. Change a body only if a concrete failing test proves + an architecture-specific correction is required. Report that finding before + expanding this design. Do not change the scalar control or output dtype. + +## Tests and gates + +G1: Add a CPU architecture-policy test before admission. It must reject +unmeasured architectures and retain the attention rejection of gfx1100. +The existing Q4_K/Q6_K physical tests must run on gfx1100 independently of +the new production predicate. Before admission, dispatch assertions must fail. +Preserve their existing fixtures and tolerances, F32/BF16 outputs, aligned +shapes, joint tails, and asymmetric bottom/right tails. Include a partial +four-wave block because the kernels contain block-wide barriers. + +G2: Build the actual HIP translation unit for gfx1100 and inspect emitted +ISA for signed-int8 WMMA. Execute registered-operation tests on the local +7900 XTX under the configured GPU mutex. Kernel output, not just a host +dispatch counter, must be checked. Re-run scalar-control outputs in separate +processes because the environment selection is cached. + +Port the applicable plugin `tests/test_kernels.py::test_mmq` coverage with its +original Q4_K/Q6_K sample tensors, token counts 7/83/128/2048, hidden widths +256/1024, seed zero, and F16/BF16/F32 inputs. Preserve the upstream tolerances. +Native outputs support F32/BF16 only; explicitly narrow a native F32 result +for the upstream F16-output comparison and document that harness adaptation. +Reuse the sample verification/export helpers under `tools/rocm_quant_gather` +where useful. Synthetic local correctness tests remain tighter independent +guards; they do not replace the upstream fixture comparison. + +G3: Prove production reachability through the public load/completion path on +a Q4_K/Q6_K fixture or the recorded Qwen3.5-4B Q4_K_M checkpoint. Use prompt +lengths that enter prefill tiles, including a nonmultiple of 16. Record completed +tokens and both format dispatch counters or matching profiler kernel traces. +Deleting each launch or restoring gfx12-only admission must fail this gate. + +G4: After correctness, compare the same binary with the environment unset and +with `VT_ROCM_QUANT_WMMA=0`. Use identical weights, prompts, token counts, +batching, and sampling. Record prefill, decode, latency, and memory on an idle +host with interleaved repeats. Run the applicable pinned primary workload and +the quant-matched llama.cpp floor before accepting a performance conclusion. +The retained checkpoint is `Qwen3.5-4B-Q4_K_M-unsloth-e87f1764.gguf`, SHA256 +`00fe7986ff5f6b463e62455821146049db6f9313603938a70800d1fb69ef11a4`, from +`unsloth/Qwen3.5-4B-GGUF@e87f176479d0855a907a41277aca2f8ee7a09523`. +Verify its hash before use. A missing oracle resource stays explicitly pending. + +G5: Run the full declared local gate: the HIP `rocm|cross_device` CTest set, +the production reachability case, CPU architecture tests, and the full +`scripts/agent-preflight.sh --staged`. Qualify skips and pair baseline failures. +A fresh reviewer inspects the immutable head, mutates admission and both launch +sites in a scratch worktree, and restores source hashes after each mutation. +The operator independently reruns the focused hardware and production gates. + +## Risks and stop conditions + +The rocWMMA implementation can compile but use a wrong fragment layout. +Output comparisons and ISA inspection distinguish that failure from admission. +Partial blocks can expose barrier assumptions. Tail tests must complete. +An instruction-path win need not improve full-model prefill. Measure both. +Retain the existing behavior on every excluded device and on decode. +If a new kernel algorithm is required, report `NEEDS_DECISION` with evidence. +Do not silently broaden the row or claim an unrun gate passed. + +## Evidence packaging + +`ISSUE-LOCAL-01M2GNY58NHBK3D4JQQ738M6GR` owns the developer-requested +packaging correction of 14 September 2026. Keep this correction in the existing +pull request. Preserve every implementation, test, and validation-harness byte. + +Archive all 236 evidence files from +`d6e40c91f634a041c873c7a04516d55c4d05772a` with their original relative paths. +Publish the archive and checksums as assets of a dedicated release on the +developer's fork. Pin that release to the source revision. Verify anonymous +download, the archive checksum, and every extracted file against the source tree +before replacing the raw files in the diff. + +Keep one concise evidence report with the download link, source revision, +SHA256, extraction instructions, correctness results, model results, review +verdicts, and all material performance and gate qualifications. Update this +specification and its issue references to that report. Detailed values, failed +attempts, commands, and original manifests remain in the archive. + +The packaging gate checks archive completeness and bytes, repeatable archive +creation, safe member paths, and detection of missing or altered files. It also +checks retained links, canonical records, diff classification, commit style, +and equality of product, test, and harness files with the reviewed head. +A fresh reviewer verifies the immutable packaging commit and downloaded archive. +The operator independently repeats the packaging gate. Reuse the recorded +hardware results when their executing files remain byte-identical. Stop if +publication or preservation cannot be verified. Do not weaken a checker or +reclassify pending measurements as passing. + +On 14 September 2026, the published archive preserves all 236 original files. +The [retained report](../../docs/bench-evidence/rocm-rdna3-quant-wmma/README.md#retrieve-the-complete-evidence) +pins the source revision, download, checksum, and extraction procedure. +The operator verified anonymous download and every original byte before replacement. +The original detailed reports and their qualifications remain unchanged in the archive. + +## Owed + +`ISSUE-LOCAL-01M2F4WCD6ZK5VH5S8TF83APD6` owns the remaining full-model +performance gaps on the retained gfx1100 workload. It includes every below-floor +axis, matching oracle traces, comparable timing windows, and accepted clock +attribution. The [evidence report and archive](../../docs/bench-evidence/rocm-rdna3-quant-wmma/README.md) +retain all measured values, ratios, and limits. This debt is separate from +architecture admission and from the deferred Qwen state characterization. + +Other formats, grouped expert tiles, and other RDNA3 devices keep their prior +owners. The parent RDNA4 row and its open cooperative-tile work remain separate. + +## Row inventory + +| ID | Upstream source | Local anchor | Tests and evidence | Spec | State | Owner | Issue | +|---|---|---|---|---|---|---|---| +| `KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3` | Pinned llama.cpp RDNA3 integer WMMA; rocWMMA 2.2.1 gfx11 fragments | `KQuantGemmKWmmaQ4K`, `KQuantGemmKWmmaQ6K` | [G1-G5 receipts](../../docs/bench-evidence/rocm-rdna3-quant-wmma/README.md), independent review and operator verification | [This spec](rocm-rdna3-quant-wmma.md) | `DONE` | RDNA3 helper; coordinating operator | `ISSUE-LOCAL-01M2F0PQWGSCXG0N4951NF9DPZ` | + +## Outcome + +The existing generic rocWMMA tile bodies execute correctly on physical gfx1100. +Architecture admission is sufficient. No arithmetic or fragment-layout repair +is needed. The Q4_K and Q6_K bodies retain their original byte hashes. +The original 240 plugin cases preserve every input mode and tolerance. +The public gate matches 1024 logits and eight completion tokens against scalar. +Independent mutations detect missing admission, missing launches, scalar-control +changes, and finite corrupted output. The operator repeats the applicable gates. + +The full-model workload uses eight requests with 183 and 174 input tokens in +alternating order. Exact input arrays and all 128 generated IDs match the primary +and llama.cpp captures. All six native process outputs match both oracles. +The enabled trace contains 152 WMMA calls per prefill and zero during decode. +The scalar trace contains no WMMA calls. + +Enabled and disabled median prefill rates are 224.12 and 173.86 tokens/s. +Median first-token latencies are 386.51 and 617.09 ms. +Mean per-stream decode rates are 53.33 and 53.15 tokens/s. +Sampled whole-device memory peaks are 4.864 and 4.857 GB. +These dynamic-clock observations motivate retaining the enabled default. +They do not establish clock-attributed performance or full-model parity. +The dedicated performance issue retains every below-floor axis and missing +measurement obligation. No apparent performance limit is accepted. + +The default architecture predicate adds only measured gfx1100 to gfx1200/gfx1201. +Other gfx11 devices remain excluded because they lack their own physical evidence. +The attention predicate remains unchanged because its gfx1100 body is not enabled. +The scalar environment override remains available for reproducible controls. +Output dtypes, quantization formats, and tile arithmetic retain their prior values +because this row proves admission without changing those contracts. diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 18cade43a..dbbc67a4b 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -347,7 +347,7 @@ expand at load instead of reaching a device provider that cannot execute them. | CPU (x86, Arm i8mm; A76 assembly correct/default, llama speed gate open, and the closed 20-core floor ran a SUPERSEDED fork denominator rather than the stock `b10451` pin, re-take owed #1003) | ✅ `CPU_ATTN` registered (#1371/#1392, [spec](../.agents/specs/attn-validate-configuration.md)) | ◐ | ☐ | ✅ | | Metal (Apple Silicon) | ✅ builds under Apple Clang with project warnings promoted to errors, the Qwen3.5 MoE loader included; its layout-refusal path uses the same messages and behavior on every platform (#1054) | ☐ | ☐ | ✅ | | Vulkan | ◐ **EXL3 is now NATIVE** ([#2530](https://github.com/mudler/vllm.cpp/issues/2530), [spec](../.agents/specs/backend-vulkan-exl3.md)): `kCastF16` and `kExl3Gemm` were the only two ops an EXL3 checkpoint ran on the CPU reference tier on a Vulkan queue, MEASURED at two fallback notices and now at ZERO. `kExl3Gemm` is TRANSCRIBED from the portable CPU reference, not ported from the CUDA kernel, whose 90 KiB shared-memory budget alone exceeds Vulkan's 16 KiB guarantee before one reaches `mma.sync`, `ldmatrix`, `cp.async` or a grid-wide barrier Vulkan has at no version -- so the gate is BYTE equality with the CPU arm on all three codebooks and every width, not a tolerance. Verified on llvmpipe with NO GPU and NO lease, which is what makes it a CI gate rather than a lease gate; no real Vulkan GPU and no checkpoint run, and NO speed number is claimed on any axis. `kExl3MoeMlp` and `kExl3HadR128` are owed | ☐ | ☐ | ✅ | -| ROCm | W0: 5 gfx archs; dense/GDN all-native; 0.8B dispatch fixed. **M4:** Qwen3-0.6B/3.5-0.8B 16/16 (#41). **M3:** `ROCM_ATTN` registered (#1056/#1065, [spec](../.agents/specs/rocm-attn-backend.md)). **EXL3 is now NATIVE on `gfx1151`** ([#2433](https://github.com/mudler/vllm.cpp/issues/2433), [spec](../.agents/specs/backend-rocm-exl3.md)): `kCastF16` and `kExl3Gemm` were the only two ops an EXL3 checkpoint still ran on the CPU reference tier, and a 3.0bpw Llama-3.2-1B now completes with ZERO reference-tier hits at 8.27 tok/s warm against 0.83 for the same tree with the two registrations disabled -- byte-identical greedy output either way. `kExl3Gemm` is TRANSCRIBED from the portable CPU reference, not ported from the CUDA kernel, whose 90 KiB shared-memory budget alone does not fit AMD's 64 KiB LDS; the gate is therefore byte equality with the CPU arm rather than a tolerance. No BF16 denominator and no clock attribution: the control HUNG the GPU in the same lease ([#2511](https://github.com/mudler/vllm.cpp/issues/2511)). `kExl3MoeMlp` is owed. CPU parity open (#269). Native grouped BF16 expert kernels pass component gates on gfx1100; full Qwen3 MoE token parity remains open ([spec](../.agents/specs/rocm-bf16-moe.md)) | Native ([ROCm guide](ROCM.md), device-specific wrapped-call-safe derived recount): full GDN, MoE combine/gate, keep-quant GEMM, EXL3 trellis GEMM; ctest-green gfx1151/1103/1100/1201/1200 (#41). **APU managed allocation is now MEASURED and NARROWED** ([#2511](https://github.com/mudler/vllm.cpp/issues/2511)): `hipMallocManaged` on a part reporting `PageableMemoryAccess = 0` faulted the GPU in 17 of 21 legs on `gfx1151`, against 0 of 21 for plain `hipMalloc`, and the plain arm is the first thing that has completed a gate-sized Qwen3.8-27B Q4_K run on that board. The managed branch now requires a device that can take a recoverable page fault, and `UnifiedMemory()` follows the allocator -- so gfx1151 and gfx1103 lose the CPU reference tier and an op with no native ROCm kernel refuses by name there, naming the attribute and `VT_ROCM_MANAGED_ALLOC=1`. | ✅ | ✅ | +| ROCm | W0: 5 gfx archs; dense/GDN all-native; 0.8B dispatch fixed. **M4:** Qwen3-0.6B/3.5-0.8B 16/16 (#41). **M3:** `ROCM_ATTN` registered (#1056/#1065, [spec](../.agents/specs/rocm-attn-backend.md)). **EXL3 is now NATIVE on `gfx1151`** ([#2433](https://github.com/mudler/vllm.cpp/issues/2433), [spec](../.agents/specs/backend-rocm-exl3.md)): `kCastF16` and `kExl3Gemm` were the only two ops an EXL3 checkpoint still ran on the CPU reference tier, and a 3.0bpw Llama-3.2-1B now completes with ZERO reference-tier hits at 8.27 tok/s warm against 0.83 for the same tree with the two registrations disabled -- byte-identical greedy output either way. `kExl3Gemm` is TRANSCRIBED from the portable CPU reference, not ported from the CUDA kernel, whose 90 KiB shared-memory budget alone does not fit AMD's 64 KiB LDS; the gate is therefore byte equality with the CPU arm rather than a tolerance. No BF16 denominator and no clock attribution: the control HUNG the GPU in the same lease ([#2511](https://github.com/mudler/vllm.cpp/issues/2511)). `kExl3MoeMlp` is owed. CPU parity open (#269). Native grouped BF16 expert kernels pass component gates on gfx1100; full Qwen3 MoE token parity remains open ([spec](../.agents/specs/rocm-bf16-moe.md)). Q4_K and Q6_K prefill uses rocWMMA on gfx1100 and retains gfx1200/gfx1201 admission. The gfx1100 public gate matches its scalar control. Full-model performance remains pending ([spec](../.agents/specs/rocm-rdna3-quant-wmma.md)) | Native ([ROCm guide](ROCM.md), device-specific wrapped-call-safe derived recount): full GDN, MoE combine/gate, keep-quant GEMM, EXL3 trellis GEMM; ctest-green gfx1151/1103/1100/1201/1200 (#41). **APU managed allocation is now MEASURED and NARROWED** ([#2511](https://github.com/mudler/vllm.cpp/issues/2511)): `hipMallocManaged` on a part reporting `PageableMemoryAccess = 0` faulted the GPU in 17 of 21 legs on `gfx1151`, against 0 of 21 for plain `hipMalloc`, and the plain arm is the first thing that has completed a gate-sized Qwen3.8-27B Q4_K run on that board. The managed branch now requires a device that can take a recoverable page fault, and `UnifiedMemory()` follows the allocator -- so gfx1151 and gfx1103 lose the CPU reference tier and an op with no native ROCm kernel refuses by name there, naming the attribute and `VT_ROCM_MANAGED_ALLOC=1`. | ✅ | ✅ | | XPU / TPU | ☐ | ✅ | ◐ | ☐ | | Tenstorrent Blackhole | ◐ `ACTIVE`, OPT-125m 6/6; Qwen3-0.6B wired; Mistral-7B-v0.3 16/16 on P150 ([spec](../.agents/specs/tenstorrent-mistral.md)). 16x16 rerun and residual-RMS owed ([spec](../.agents/specs/tenstorrent-backend.md)) | ✅ | ☐ | ☐ | | Tenstorrent host-free decode | ◐ DEFAULT since #1604 (`0` opts out): no per-step host readback; 2.1x default-leg tok/s; both golden pairs re-adjudicated, both paged gates 16/16. Capture DEFAULT for Qwen3-dense since the #1625 flip (`0` opts out) — 27.7 vs 12.9 tok/s at c=1, 2.15x over the opt-out, [details](benchmarks/tt-capture-default-decode.md); Qwen3.5 dense (GDN) joined the default since [#2907](https://github.com/mudler/vllm.cpp/issues/2907), the decode-side GDN ops running device-resident under capture with [#2812](https://github.com/mudler/vllm.cpp/issues/2812)'s captured fatal fixed; 4B near-tie pair owed ([#2811](https://github.com/mudler/vllm.cpp/issues/2811)); async off (#1627) | ☐ | ☐ | ☐ | diff --git a/docs/bench-evidence/rocm-rdna3-quant-wmma/README.md b/docs/bench-evidence/rocm-rdna3-quant-wmma/README.md new file mode 100644 index 000000000..d3d05aa89 --- /dev/null +++ b/docs/bench-evidence/rocm-rdna3-quant-wmma/README.md @@ -0,0 +1,109 @@ +# gfx1100 quantized WMMA evidence + +Row: `KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3`. +[Specification](../../../.agents/specs/rocm-rdna3-quant-wmma.md). +Architecture admission is `DONE`. Whole-model performance floors remain `FAILING`. +This packaging change preserves the reviewed implementation, tests, and validation harness. + +## Retrieve the complete evidence + +[Release](https://github.com/VikashLoomba/vllm.cpp/releases/tag/rdna3-wmma-evidence-d6e40c91f634) · +[Download archive](https://github.com/VikashLoomba/vllm.cpp/releases/download/rdna3-wmma-evidence-d6e40c91f634/rdna3-wmma-evidence-d6e40c91f634.tar.gz). +Source revision: `d6e40c91f634a041c873c7a04516d55c4d05772a`. +The 511,480-byte archive preserves all 236 original files under this evidence path. +It adds a root `SHA256SUMS` manifest. Archive SHA256: +`996ed227235b9e85196055b9005e2adf7a1d8c86a61ef610249fcfd0be36eb0d`. + +Run these commands with Python 3.12 or later. Extraction uses a new private directory. + +```sh +( +set -eu +archive=rdna3-wmma-evidence-d6e40c91f634.tar.gz +curl -q --fail --location --output "$archive" \ + "https://github.com/VikashLoomba/vllm.cpp/releases/download/rdna3-wmma-evidence-d6e40c91f634/$archive" +printf '%s %s\n' \ + 996ed227235b9e85196055b9005e2adf7a1d8c86a61ef610249fcfd0be36eb0d "$archive" | sha256sum --check --strict +dest=$(mktemp -d) +python3 - "$archive" "$dest" <<'PYARCHIVE' +import pathlib, sys, tarfile +with tarfile.open(sys.argv[1], "r:gz") as archive: + for member in archive.getmembers(): + path = pathlib.PurePosixPath(member.name) + assert not path.is_absolute() and ".." not in path.parts + assert member.isfile() or member.isdir() + archive.extractall(sys.argv[2], filter="data") +PYARCHIVE +cd "$dest" +sha256sum --check --strict SHA256SUMS +printf 'Extracted evidence to: %s\n' "$PWD" +) +``` + +Open `docs/bench-evidence/rocm-rdna3-quant-wmma/README.md` and `model-summary.md` inside the extraction for the full reports. +Their original cross-links outside the evidence directory require the source checkout at the revision above. +The archive retains commands, manifests, failed attempts, compiler identities, and receipt qualifications unchanged. +Large raw matrices, model logits, and full profiler traces were already external at their sealed source paths. + +## Correctness, review, and scope + +Implementation: `c3fe98ba6c55ce71e75746e1b944a27640464e0f`. +Measurements used physical RX 7900 XTX `gfx1100` on 13 September 2026. +Every GPU invocation held `/home/vikash/gpu.lock` and selected device 0. +Native tooling used HIP 7.15, Clang 23, and rocWMMA 2.2.1. The archive seals exact toolchain hashes. + +| Gate | Result | Archived receipts | +|---|---|---| +| G1 architecture policy | Four assertions fail before admission. Green: 16 cases, 109 assertions. Attention admission stays unchanged. | `arch-red.log`, `arch-green.log` | +| G1-G2 physical tiles | Red: four cases fail ten dispatch assertions. Green and separate scalar control: 46 assertions each. | `hardware-red.log`, `hardware-green.log`, `hardware-scalar.log` | +| G2 arithmetic and ISA | Both tile bodies remain byte-identical. Eight signed-int8 WMMA instructions span Q4_K/Q6_K and F32/BF16. | `tile-bodies.json`, `gfx1100-wmma-isa.txt` | +| G2 original primary fixtures | All 240 cases pass against the dense reference and plugin outputs with original modes and tolerances. | `original-mmq-summary.json`, `original-mmq-comparison.json` | +| G3 production reachability | Prompts of 16 and 37 tokens enter both formats. All 1024 logits and eight completion tokens match scalar. | `public-red.log`, `public-green.log` | +| G5 fresh review | `PASS`, no findings. Mutated admission, launches, scalar override, and finite corruption fail their gates. Source hashes restore exactly. | `review-review-report.json`, `production-mutations.json` | +| G5 operator verification | Independent architecture, public, and 240-case gates pass. HIP: 29 registered tests, zero failures, five baseline skips. | `operator-receipt.json`, `model-operator-mmq-comparison.json` | + +Physical cases retain F32/BF16 outputs, joint and asymmetric tails, and the partial four-wave block `M=32,N=48,K=512`. +Original fixtures retain `M=7,83,128,2048`, `K=256,1024`, seed zero, and F16/BF16/F32 inputs. +The native harness narrows F32 to F16 for upstream F16 output comparisons. Native Q8_K activation and output contracts remain unchanged. +Source chain: plugin `tests/test_kernels.py:145-196` → `ops.py:200` → `csrc/gguf/gguf_kernel.hip:221` → `mmq_hip.cuh:490,591`. + +HIP skips require three model fixtures and two-visible-device coverage. Only device 0 was exposed. +Historical full preflights preserve 12 skips. The operator's original role failure and its scoped branch-rename resolution remain recorded. +Seven NumPy suites subsequently pass in isolation. Five adherence checkpoint subcases remain unavailable without `VT_LTX25_ADHERENCE_MODEL`. +The archived final-records preflight retains its onboarding fixture failure and the passing isolated rerun of all 39 cases. +The later source-head readiness run has zero failed checks and five argument-dependent skips, with exit 1 under readiness policy. +PR classification passed separately. ARM, CPU, CUDA, and Triton argument-dependent checks concern unchanged architecture paths. +These qualified results do not constitute an all-green readiness claim. + +## Model identity and observed results + +Checkpoint: `Qwen3.5-4B-Q4_K_M-unsloth-e87f1764.gguf`, 2,740,937,888 bytes. +Source: `unsloth/Qwen3.5-4B-GGUF@e87f176479d0855a907a41277aca2f8ee7a09523`. +SHA256: `00fe7986ff5f6b463e62455821146049db6f9313603938a70800d1fb69ef11a4`. +Task snapshots: vLLM `39545e475d3627287ff69c25465dc0bd405f67e1`, +llama.cpp `093a2f86c3e37c54fa3e1f9efb17b304f3433abd`, +GGUF plugin `d4c1f0d082fc7cd4350da56689109a01c1f29d6c`. +These developer-selected snapshots do not advance the repository parity pin. + +Eight requests alternate 183/174 input tokens, totaling 1428 inputs and 128 generated IDs. +Concurrency, temperature, and seed are 1, 0, and 0. Each request generates 16 tokens without chat templates or EOS-logit masks. +All eight input arrays match the primary. All six native process outputs match both oracles exactly. +Native `rocprofv3` captures contain 152 WMMA calls per prefill, 1216 total, and zero during decode or with WMMA disabled. +Both traced completion arrays match the primary. Native traces do not establish cross-engine invocation parity. + +| Observed axis | WMMA enabled | WMMA disabled | Enabled/scalar | +|---|---:|---:|---:| +| Prefill input tokens per summed first-token latency | 224.12 tokens/s | 173.86 tokens/s | 1.2891 | +| Median first-token latency | 386.51 ms | 617.09 ms | 0.6263 | +| Mean per-stream decode rate | 53.33 tokens/s | 53.15 tokens/s | 1.0034 | +| Sampled whole-device memory peak | 4.8638 GB | 4.8567 GB | 1.0015 | + +Native medians cover three complete processes per mode in `on1,off1,off2,on2,on3,off3` order, including first-use effects. +Each oracle loads once and captures all four two-request legs. Primary compilation caches were warmed by its failed first attempt. +Clocks vary dynamically. Monitor windows include initialization and teardown, so accepted clock attribution remains `PENDING`. +Engine cache capacities differ. The primary reserves approximately 17 GiB, preventing equal-capacity memory conclusions. +The llama.cpp client includes extra logit, memory, and finite-value instrumentation. Its decode-call-only timings are a separate scope. +Memory values are sampled device-wide or process-tree peaks. Small decode or memory differences do not establish stable changes. +Native/primary total whole-run throughput is 0.9443 and decode is 0.8880. Native/llama.cpp ratios are 0.2439 and 0.5743. +The [open performance issue](../../../.agents/issues/_owed/ISSUE-LOCAL-01M2F4WCD6ZK5VH5S8TF83APD6.md) retains every below-floor axis, comparable timing windows, matching oracle traces, and accepted clock reproduction. +The archived `model-summary.md` retains every value and ratio. Architecture admission claims no full-model parity or performance ceiling. diff --git a/include/vt/rocm/rocm_quant_wmma_arch.h b/include/vt/rocm/rocm_quant_wmma_arch.h new file mode 100644 index 000000000..d55d84974 --- /dev/null +++ b/include/vt/rocm/rocm_quant_wmma_arch.h @@ -0,0 +1,31 @@ +// Quantized prefill admission stays separate from the attention policy. +#pragma once + +#include + +namespace vt::rocm { + +// rocWMMA owns the different gfx11/gfx12 operand layouts. Admit only the +// measured gfx1100 device and existing gfx1200/gfx1201 devices. Attention +// keeps its gfx12 policy because its gfx11 device body is still excluded. +constexpr bool GcnArchNameHasQuantWmma(std::string_view gcn_arch) { + const auto colon = gcn_arch.find(':'); + const auto stem = gcn_arch.substr(0, colon); + // Require an exact stem followed by HIP target features with explicit signs. + if (stem != "gfx1100" && stem != "gfx1200" && stem != "gfx1201") return false; + if (colon == std::string_view::npos) return true; + auto features = gcn_arch.substr(colon + 1); + bool xnack = false, sramecc = false; + while (!features.empty()) { + const auto next = features.find(':'); + const auto feature = features.substr(0, next); + if ((feature == "xnack+" || feature == "xnack-") && !xnack) xnack = true; + else if ((feature == "sramecc+" || feature == "sramecc-") && !sramecc) sramecc = true; + else return false; + if (next == std::string_view::npos) return true; + features.remove_prefix(next + 1); + } + return false; +} + +} // namespace vt::rocm diff --git a/src/vt/rocm/rocm_grouped_gemm.hip b/src/vt/rocm/rocm_grouped_gemm.hip index f1573a214..208626efe 100644 --- a/src/vt/rocm/rocm_grouped_gemm.hip +++ b/src/vt/rocm/rocm_grouped_gemm.hip @@ -15,11 +15,9 @@ #include #include #include -// rocWMMA int8 tile arm (KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA4, issue #2109), -// gfx1200/gfx1201 (RDNA4) only — same gate rocm_paged_attn.hip already uses -// for its own WMMA arm, and the same reason: a fat build also emits other -// gfx12 targets (for example gfx1036) that must not see this include. -#if defined(__gfx1200__) || defined(__gfx1201__) +// The quantized fragment API supports measured gfx1100 and gfx1200/gfx1201. +// Keep this device guard independent of the attention WMMA architecture gate. +#if defined(__gfx1100__) || defined(__gfx1200__) || defined(__gfx1201__) #include #define VT_ROCM_QUANT_WMMA_OK 1 #endif @@ -38,12 +36,12 @@ #include "vt/grow_only_stream_scratch.h" #include "vt/ops.h" #include "vt/rocm/rocm_act_quant.h" -#include "vt/rocm/rocm_arch.h" #include "vt/rocm/rocm_device_bind.h" #include "vt/rocm/rocm_f16_codec.h" #include "vt/rocm/rocm_mmvq_policy.h" #include "vt/rocm/rocm_norm_quant_bridge.h" #include "vt/rocm/rocm_quant_iq_tables.h" +#include "vt/rocm/rocm_quant_wmma_arch.h" #include "vt/rocm/rocm_runtime.h" // Block layouts — the single source of truth (ggml-common.h mirrors). @@ -461,12 +459,8 @@ __global__ void KQuantGemmKWmmaQ6K(OutT* __restrict__ out, const uint8_t* __rest const BlockQ8_K* __restrict__ act, int64_t m, int64_t n, int64_t nsb, int64_t n_tiles) { #if !defined(VT_ROCM_QUANT_WMMA_OK) - // Compiled for a target other than gfx1200/gfx1201, or for the host pass: - // never launched there (`Gfx12QuantWmmaHostOk` gates every launch site - // below on the actual device's runtime gcnArchName, not on this macro — - // the macro is only visible to a device compilation pass for gfx1200/ - // gfx1201 itself, exactly the trap issue #785 already found and fixed for - // `rocm_paged_attn.hip`'s own WMMA arm). + // Non-admitted device and host passes retain a stub. QuantWmmaHostOk + // checks the actual device architecture before any launch of this body. (void)out; (void)weight; (void)act; (void)m; (void)n; (void)nsb; (void)n_tiles; return; #else @@ -761,17 +755,9 @@ __global__ void KQuantGemmKWmmaQ4K(OutT* __restrict__ out, const uint8_t* __rest #endif // VT_ROCM_QUANT_WMMA_OK } -// Per-device, once: launch the Q6_K WMMA tile only when this device's -// gcnArchName is gfx1200/gfx1201 (prefix match). NEVER gated by -// VT_ROCM_QUANT_WMMA_OK at a host call site — that macro is defined only on -// the device compilation pass for gfx1200/gfx1201 itself, so a host-side -// `#if` on it always compiles away the launch (issue #785's exact trap, -// documented in `.agents/specs/rocm-sharedk-wmma-host-dispatch.md`). Reuses -// the same HIP-free, unit-tested prefix check `rocm_paged_attn.hip`'s own -// WMMA dispatch already established (`vt::rocm::GcnArchNameIsGfx12PrefillWmma`, -// `include/vt/rocm/rocm_arch.h`); the name is attention-flavored but the test -// it names is pure architecture matching, shared rather than duplicated here. -bool Gfx12QuantWmmaHostOk(int device_index) { +// Cache each device's quant-specific runtime admission. The compile guard is +// device-only, so testing it on the host would compile away every launch. +bool QuantWmmaHostOk(int device_index) { struct Slot { std::once_flag once; bool ok = false; @@ -785,7 +771,7 @@ bool Gfx12QuantWmmaHostOk(int device_index) { return; } slots[static_cast(device_index)].ok = - vt::rocm::GcnArchNameIsGfx12PrefillWmma(prop.gcnArchName); + vt::rocm::GcnArchNameHasQuantWmma(prop.gcnArchName); }); return slots[static_cast(device_index)].ok; } @@ -1441,7 +1427,7 @@ std::atomic g_kq_wmma_dispatches{0}; // site is reached independently, not infer Q4_K's from Q6_K's. std::atomic g_kq_wmma_q4k_dispatches{0}; -// VT_ROCM_QUANT_WMMA=0 forces the scalar arm even on gfx1200/gfx1201 — the +// VT_ROCM_QUANT_WMMA=0 forces the scalar arm even on admitted WMMA devices — the // A/B escape hatch every WMMA arm in this tree ships with // (VT_ATTN_PREFILL_SHAREDK_WMMA is the paged-attention precedent). Default // on: the gate below still requires the runtime architecture match and @@ -1803,7 +1789,7 @@ void MatmulBTQuantKernelRocmGdn(Queue& q, Tensor& out, const Tensor& a, const Te g_mmvq_route_baseline.fetch_add(1, std::memory_order_relaxed); const int64_t grid = (m * n + kWarpsPerBlock - 1) / kWarpsPerBlock; - // RDNA4 WMMA tile arm: Q6_K and Q4_K prefill only (m > 1, since + // Quantized WMMA tile arm: Q6_K and Q4_K prefill only (m > 1, since // KQuantDecodeCoopWarps above already owns m == 1). Q5_K is `## Owed` — // same 32-wide scale/min shape as Q4_K plus a high-bit plane, not yet // ported. Every other format and every other architecture falls through @@ -1829,7 +1815,7 @@ void MatmulBTQuantKernelRocmGdn(Queue& q, Tensor& out, const Tensor& a, const Te // much as the whole GEMM to fill a typically thin remainder strip, for // a large aligned corner and wide N. const bool wmma_ok = (fmt == 2 || fmt == 0) && m >= 16 && n >= 16 && - QuantWmmaEnabled() && Gfx12QuantWmmaHostOk(q.device.index); + QuantWmmaEnabled() && QuantWmmaHostOk(q.device.index); auto launch = [&](auto ot) { using OutT = decltype(ot); auto* o = static_cast(out.data); diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 35ee143bd..b98062b85 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -2976,6 +2976,8 @@ vllm_cpp_add_test(test_cuda_embedding_quant vt/test_cuda_embedding_quant.cpp) target_include_directories(test_cuda_embedding_quant PRIVATE ${CMAKE_SOURCE_DIR}/src) # BACKEND-ROCM-QUANT-GATHER: the production gate starts at the public loader. vllm_cpp_add_test(test_capi_rocm_embedding_quant capi/test_rocm_embedding_quant.cpp) +# Quantized WMMA reaches both formats through the public completion API. +vllm_cpp_add_test(test_capi_rocm_quant_wmma capi/test_rocm_quant_wmma.cpp) vllm_cpp_add_test(test_rocm_embedding_quant vt/test_rocm_embedding_quant.cpp) target_include_directories(test_rocm_embedding_quant PRIVATE ${CMAKE_SOURCE_DIR}/src) if(VLLM_CPP_HIP) @@ -2987,6 +2989,8 @@ endif() add_executable(rocm_quant_gather_capture ../tools/rocm_quant_gather/capture.cpp) target_link_libraries(rocm_quant_gather_capture PRIVATE vllm) target_include_directories(rocm_quant_gather_capture PRIVATE ${CMAKE_SOURCE_DIR}/src ${CMAKE_SOURCE_DIR}/tests) +add_executable(rocm_quant_wmma_capture ../tools/rocm_quant_wmma/capture.cpp) +target_link_libraries(rocm_quant_wmma_capture PRIVATE vllm) # CPU-ELEM-GEMM: the specialized + SIMD elementwise (bf16/f16/f32) GEMM. The # gate is BYTE-IDENTITY against an independent scalar reference across every # dtype combination, both weight orientations, ragged K/N, row-strided diff --git a/tests/capi/rocm_quant_gather_fixture.h b/tests/capi/rocm_quant_gather_fixture.h index defabe86a..f908fe3d5 100644 --- a/tests/capi/rocm_quant_gather_fixture.h +++ b/tests/capi/rocm_quant_gather_fixture.h @@ -119,7 +119,9 @@ inline std::string DenseEmbedding(const Format& format, return dense; } -inline std::string BuildModel(const Format& format, bool dense_embedding) { +inline std::string BuildModel(const Format& format, bool dense_embedding, + const Format* gate_up_format = nullptr, + const Format* down_format = nullptr) { using namespace gguf_test; GgufModelBuilder builder; builder.AddKv(StrKv("general.architecture", "qwen35")); @@ -166,7 +168,14 @@ inline std::string BuildModel(const Format& format, bool dense_embedding) { dense_embedding ? 30 : format.ggml, dense_embedding ? DenseEmbedding(format, packed) : packed); std::mt19937 rng(kSeed); - const auto projection = [&](const char* name, uint64_t input, uint64_t output) { + const auto projection = [&](const char* name, uint64_t input, uint64_t output, + const Format* quant = nullptr) { + if (quant != nullptr) { + builder.AddTensor(name, {input, output}, quant->ggml, + PackedTable(*quant, static_cast(output), + static_cast(input))); + return; + } std::string bytes; bytes.reserve(static_cast(input * output * 2)); for (uint64_t i = 0; i < input * output; ++i) { @@ -188,9 +197,9 @@ inline std::string BuildModel(const Format& format, bool dense_embedding) { projection("blk.0.attn_k.weight", 256, 64); projection("blk.0.attn_v.weight", 256, 64); projection("blk.0.attn_output.weight", 256, 256); - projection("blk.0.ffn_gate.weight", 256, 256); - projection("blk.0.ffn_up.weight", 256, 256); - projection("blk.0.ffn_down.weight", 256, 256); + projection("blk.0.ffn_gate.weight", 256, 256, gate_up_format); + projection("blk.0.ffn_up.weight", 256, 256, gate_up_format); + projection("blk.0.ffn_down.weight", 256, 256, down_format); norm("output_norm.weight", 256); norm("blk.0.attn_norm.weight", 256); norm("blk.0.post_attention_norm.weight", 256); diff --git a/tests/capi/test_rocm_quant_wmma.cpp b/tests/capi/test_rocm_quant_wmma.cpp new file mode 100644 index 000000000..d9eae519a --- /dev/null +++ b/tests/capi/test_rocm_quant_wmma.cpp @@ -0,0 +1,182 @@ +// KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA3: public Q4_K/Q6_K prefill reachability. +// A separate scalar process supplies complete logits for the identical GGUF. +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include "capi/rocm_quant_gather_fixture.h" +#include "vllm.h" +#include "vt/backend.h" +#include "vt/ops.h" + +#if defined(VLLM_CPP_HIP) && !defined(VLLM_CPP_CUDA) && defined(__unix__) +#include +#include +extern char** environ; + +namespace vt::rocm { +std::string DeviceArchName(int); +uint64_t KQuantWmmaDispatchCount(); +uint64_t KQuantWmmaQ4KDispatchCount(); +} // namespace vt::rocm + +namespace { +using Json = nlohmann::json; +struct Completion { + std::vector logits; + int callbacks = 0; + bool invalid = false; +}; + +void Require(bool value, const std::string& message) { + if (!value) throw std::runtime_error(message); +} + +void Capture(const int32_t*, int32_t, float* logits, int32_t vocab, void* data) { + auto& capture = *static_cast(data); + if (vocab != 128 || capture.callbacks >= 4) { + capture.invalid = true; + return; + } + for (int i = 0; i < vocab; ++i) { + if (!std::isfinite(logits[i])) capture.invalid = true; + capture.logits.push_back(logits[i]); + } + ++capture.callbacks; +} + +Json Run(bool expect_wmma) { + const rocm_gather_test::Format q4{vt::DType::kQ4_K, 12, "Q4_K"}; + const rocm_gather_test::Format q6{vt::DType::kQ6_K, 14, "Q6_K"}; + gguf_test::TempFile model_file(rocm_gather_test::BuildModel(q4, true, &q4, &q6)); + Json result = Json::array(); + for (const int tokens : {16, 37}) { + auto params = vllm_model_params_default(); + params.model_path = model_file.path().c_str(); + params.device = 0; + params.block_size = 16; + params.num_blocks = 16; + params.max_model_len = 64; + params.max_num_seqs = 1; + vllm_engine* raw = nullptr; + Require(vllm_engine_load(¶ms, &raw) == VLLM_OK, + "public load: " + std::string(vllm_last_error())); + std::unique_ptr engine(raw, vllm_engine_free); + std::vector prompt(static_cast(tokens)); + for (int i = 0; i < tokens; ++i) prompt[static_cast(i)] = (i * 19 + 1) % 128; + Completion capture; + capture.logits.reserve(4 * 128); + std::vector ids(4); + auto sampling = vllm_sampling_params_default(); + sampling.temperature = 0.0F; + sampling.max_tokens = 4; + sampling.ignore_eos = 1; + sampling.has_seed = 1; + sampling.seed = rocm_gather_test::kSeed; + sampling.logits_processor = Capture; + sampling.logits_processor_user_data = &capture; + int32_t count = 0; + const auto q4_before = vt::rocm::KQuantWmmaQ4KDispatchCount(); + const auto q6_before = vt::rocm::KQuantWmmaDispatchCount(); + const auto reference_before = vt::GetReferenceTierHits(); + Require(vllm_complete_tokens(engine.get(), prompt.data(), tokens, &sampling, + ids.data(), 4, &count, nullptr) == VLLM_OK, + "public completion: " + std::string(vllm_last_error())); + const auto q4_calls = vt::rocm::KQuantWmmaQ4KDispatchCount() - q4_before; + const auto q6_calls = vt::rocm::KQuantWmmaDispatchCount() - q6_before; + Require(count == 4 && capture.callbacks == 4 && !capture.invalid, + "completion must return four tokens and finite full logits"); + Require(vt::GetReferenceTierHits() == reference_before, "unexpected CPU reference fallback"); + Require((q4_calls > 0) == expect_wmma, "Q4_K public WMMA dispatch mismatch"); + Require((q6_calls > 0) == expect_wmma, "Q6_K public WMMA dispatch mismatch"); + std::printf("PUBLIC prompt=%d q4_wmma=%llu q6_wmma=%llu tokens=", tokens, + static_cast(q4_calls), + static_cast(q6_calls)); + for (const auto id : ids) std::printf(" %d", id); + std::puts(""); + result.push_back({{"prompt_tokens", tokens}, {"ids", ids}, {"logits", capture.logits}}); + } + return result; +} + +void ScalarProcess(const char* executable, const std::string& output) { + // posix_spawn executes a fresh process because QuantWmmaEnabled caches the + // environment on first use. The parent's GPU mutex also covers this child. + std::vector variables; + for (char** entry = environ; *entry != nullptr; ++entry) { + const std::string value(*entry); + if (value.rfind("VT_ROCM_QUANT_WMMA=", 0) != 0) variables.push_back(value); + } + variables.emplace_back("VT_ROCM_QUANT_WMMA=0"); + std::vector environment; + for (auto& value : variables) environment.push_back(value.data()); + environment.push_back(nullptr); + char* args[] = {const_cast(executable), const_cast("--scalar"), + const_cast(output.c_str()), nullptr}; + pid_t child = 0; + Require(posix_spawnp(&child, executable, nullptr, nullptr, args, environment.data()) == 0, + "cannot start scalar control"); + int status = 0; + Require(waitpid(child, &status, 0) == child && WIFEXITED(status) && WEXITSTATUS(status) == 0, + "scalar control failed"); +} +} // namespace +#endif + +int main(int argc, char** argv) { +#if defined(VLLM_CPP_HIP) && !defined(VLLM_CPP_CUDA) && defined(__unix__) + try { + const std::string arch = vt::rocm::DeviceArchName(0); + const auto stem = arch.substr(0, arch.find(':')); + if (stem != "gfx1100" && stem != "gfx1200" && stem != "gfx1201") { + std::printf("SKIPPED: physical gfx1100 or gfx1200/gfx1201 required, got %s\n", arch.c_str()); + return 77; + } + if (argc == 3 && std::string(argv[1]) == "--scalar") { + const auto report = Run(false); + std::ofstream file(argv[2]); + file << report.dump(2) << '\n'; + Require(static_cast(file), "cannot write scalar result"); + return 0; + } + Require(argc == 1, "use no arguments or --scalar OUTPUT"); + const char* knob = std::getenv("VT_ROCM_QUANT_WMMA"); + Require(knob == nullptr || std::string(knob) != "0", "run the public gate with WMMA enabled"); + gguf_test::TempFile scalar_file(""); + ScalarProcess(argv[0], scalar_file.path()); + std::ifstream stream(scalar_file.path()); + const auto expected = Json::parse(stream); + const auto actual = Run(true); + Require(actual.size() == expected.size(), "scalar prompt count differs"); + float max_error = 0; + for (size_t p = 0; p < actual.size(); ++p) { + Require(actual[p]["ids"] == expected[p]["ids"], "scalar completion tokens differ"); + const auto got = actual[p]["logits"].get>(); + const auto ref = expected[p]["logits"].get>(); + Require(got.size() == ref.size(), "scalar logits shape differs"); + for (size_t i = 0; i < got.size(); ++i) { + const float error = std::abs(got[i] - ref[i]); + max_error = std::max(max_error, error); + Require(error <= 0.02F + 0.01F * std::abs(ref[i]), "public logits differ from scalar control"); + } + } + std::printf("PUBLIC WMMA PASS: 1024 logits, token-exact, max_abs_error=%g\n", max_error); + return 0; + } catch (const std::exception& error) { + std::fprintf(stderr, "PUBLIC WMMA FAILED: %s\n", error.what()); + return 1; + } +#else + (void)argc; + (void)argv; + std::puts("SKIPPED: the public AUTO selection gate requires a HIP-only Unix build"); + return 77; +#endif +} diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index 033c87091..1ea50e7a6 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -3675,13 +3675,16 @@ TEST_CASE("ROCm Q6_K decode spreads one row's superblocks over several warps") { } } -// KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA4 (issue #2109), W1: the RDNA4 WMMA int8 -// tile arm of the Q6_K prefill GEMM. Runs ONLY on gfx1200/gfx1201 — every -// other ROCm target keeps the scalar arm the case above already covers, so -// this returns early rather than skip-reporting on hardware it does not -// target (`GcnArchNameIsGfx12PrefillWmma` is the same host gate the kernel's -// own dispatch decision uses, per `include/vt/rocm/rocm_arch.h`). -TEST_CASE("keep-quant Q6_K WMMA tile arm matches the CPU oracle on RDNA4") { +// Q4_K/Q6_K WMMA runs on physical gfx1100, gfx1200, and gfx1201. +// Preserve the original RDNA4 fixtures, including six tiles in a four-wave +// block and both output dtypes. Scalar controls run in a separate process. +namespace { +bool ExpectQuantWmma() { + const char* value = std::getenv("VT_ROCM_QUANT_WMMA"); + return value == nullptr || std::strcmp(value, "0") != 0; +} +} // namespace +TEST_CASE("keep-quant Q6_K WMMA tile arm matches the CPU oracle on gfx1100 and RDNA4") { const bool rocm_registered = [] { for (DeviceType dt : RegisteredDevices()) if (dt == DeviceType::kROCM) return true; @@ -3691,7 +3694,10 @@ TEST_CASE("keep-quant Q6_K WMMA tile arm matches the CPU oracle on RDNA4") { REQUIRE(OpAvailable(vt::OpId::kMatmulBTQuant, DeviceType::kROCM)); const std::string actual_arch = vt::rocm::DeviceArchName(0); - if (!vt::rocm::GcnArchNameIsGfx12PrefillWmma(actual_arch)) return; + // Test admission is independent of the production policy: narrowing that + // policy must fail dispatch assertions instead of skipping this case. + const auto stem = actual_arch.substr(0, actual_arch.find(':')); + if (stem != "gfx1100" && stem != "gfx1200" && stem != "gfx1201") return; // Tile-aligned M and N (both multiples of 16, and not equal, so the grid // exercises a non-square m_tiles x n_tiles), K spanning more than one @@ -3767,7 +3773,7 @@ TEST_CASE("keep-quant Q6_K WMMA tile arm matches the CPU oracle on RDNA4") { // site's call in a scratch copy leaves this counter flat and reds this // case, which the NMSE checks above cannot do on their own — the scalar // fallback would still pass them. - CHECK(wmma_after > wmma_before); + CHECK((wmma_after > wmma_before) == ExpectQuantWmma()); } rocm.DestroyQueue(q); } @@ -3777,7 +3783,7 @@ TEST_CASE("keep-quant Q6_K WMMA tile arm matches the CPU oracle on RDNA4") { // carries a second per-sub-block correction (`dmin * sumi`) Q6_K has no // equivalent of, so this is not just the Q6_K case with a different dtype — // it exercises a materially different code path in `KQuantGemmKWmmaQ4K`. -TEST_CASE("keep-quant Q4_K WMMA tile arm matches the CPU oracle on RDNA4") { +TEST_CASE("keep-quant Q4_K WMMA tile arm matches the CPU oracle on gfx1100 and RDNA4") { const bool rocm_registered = [] { for (DeviceType dt : RegisteredDevices()) if (dt == DeviceType::kROCM) return true; @@ -3787,7 +3793,10 @@ TEST_CASE("keep-quant Q4_K WMMA tile arm matches the CPU oracle on RDNA4") { REQUIRE(OpAvailable(vt::OpId::kMatmulBTQuant, DeviceType::kROCM)); const std::string actual_arch = vt::rocm::DeviceArchName(0); - if (!vt::rocm::GcnArchNameIsGfx12PrefillWmma(actual_arch)) return; + // Test admission is independent of the production policy: narrowing that + // policy must fail dispatch assertions instead of skipping this case. + const auto stem = actual_arch.substr(0, actual_arch.find(':')); + if (stem != "gfx1100" && stem != "gfx1200" && stem != "gfx1201") return; constexpr int64_t M = 32, N = 48, K = 512; constexpr int64_t kBlockBytes = 144; // sizeof(BlockQ4_K) @@ -3858,7 +3867,7 @@ TEST_CASE("keep-quant Q4_K WMMA tile arm matches the CPU oracle on RDNA4") { CHECK(Nmse(ref, gotf) <= kNmseTol); } const uint64_t wmma_after = vt::rocm::KQuantWmmaQ4KDispatchCount(); - CHECK(wmma_after > wmma_before); + CHECK((wmma_after > wmma_before) == ExpectQuantWmma()); } rocm.DestroyQueue(q); } @@ -3881,7 +3890,10 @@ TEST_CASE("keep-quant GEMM matches the CPU oracle when M and N are not multiples REQUIRE(OpAvailable(vt::OpId::kMatmulBTQuant, DeviceType::kROCM)); const std::string actual_arch = vt::rocm::DeviceArchName(0); - if (!vt::rocm::GcnArchNameIsGfx12PrefillWmma(actual_arch)) return; + // Test admission is independent of the production policy: narrowing that + // policy must fail dispatch assertions instead of skipping this case. + const auto stem = actual_arch.substr(0, actual_arch.find(':')); + if (stem != "gfx1100" && stem != "gfx1200" && stem != "gfx1201") return; constexpr int64_t M = 37, N = 50, K = 512; struct Fmt { @@ -3959,7 +3971,7 @@ TEST_CASE("keep-quant GEMM matches the CPU oracle when M and N are not multiples // silent full fallback to scalar would pass the NMSE check above just // as well, which is exactly why #2109's own real-model measurement // needed a hand-trimmed prompt before this fix. - CHECK(wmma_after > wmma_before); + CHECK((wmma_after > wmma_before) == ExpectQuantWmma()); rocm.DestroyQueue(q); } } @@ -3988,7 +4000,10 @@ TEST_CASE("keep-quant GEMM matches the CPU oracle when only one of M/N is misali REQUIRE(OpAvailable(vt::OpId::kMatmulBTQuant, DeviceType::kROCM)); const std::string actual_arch = vt::rocm::DeviceArchName(0); - if (!vt::rocm::GcnArchNameIsGfx12PrefillWmma(actual_arch)) return; + // Test admission is independent of the production policy: narrowing that + // policy must fail dispatch assertions instead of skipping this case. + const auto stem = actual_arch.substr(0, actual_arch.find(':')); + if (stem != "gfx1100" && stem != "gfx1200" && stem != "gfx1201") return; constexpr int64_t K = 512; struct Fmt { @@ -4074,7 +4089,7 @@ TEST_CASE("keep-quant GEMM matches the CPU oracle when only one of M/N is misali : vt::rocm::KQuantWmmaQ4KDispatchCount(); // Reachability: the WMMA arm must still fire for its aligned corner // even though one dimension is a remainder-only split. - CHECK(wmma_after > wmma_before); + CHECK((wmma_after > wmma_before) == ExpectQuantWmma()); rocm.DestroyQueue(q); } } diff --git a/tests/vt/test_rocm_arch.cpp b/tests/vt/test_rocm_arch.cpp index 68734511c..09022c77d 100644 --- a/tests/vt/test_rocm_arch.cpp +++ b/tests/vt/test_rocm_arch.cpp @@ -16,6 +16,7 @@ #include #include "vt/rocm/rocm_arch.h" +#include "vt/rocm/rocm_quant_wmma_arch.h" #include "vt/rocm/rocm_skinny_gemm_arch.h" using vt::rocm::CapabilityFromGcnArch; @@ -110,6 +111,27 @@ TEST_CASE("SharedK WMMA host gate is gfx1200/gfx1201 prefix, not substring") { static_assert(!GcnArchNameIsGfx12PrefillWmma("gfx12010")); } +TEST_CASE("quantized WMMA admits only measured device names and HIP features") { + using vt::rocm::GcnArchNameHasQuantWmma; + for (const char* name : {"gfx1100", "gfx1200", "gfx1201", "gfx1100:xnack-", + "gfx1100:sramecc+:xnack-", "gfx1200:sramecc-", + "gfx1201:xnack+:sramecc+"}) { + CAPTURE(std::string(name)); + CHECK(GcnArchNameHasQuantWmma(name)); + } + for (const char* name : {"", "gfx110", "gfx11000", "gfx1100garbage", + "agfx1100", "gfx1100:", "gfx1100:xnack", "gfx1100:foo+", + "gfx1100:xnack-:", "gfx1100::xnack-", "gfx1201garbage", + "gfx1100:xnack-:xnack+", "gfx1200:sramecc+:sramecc+", + "gfx1101", "gfx1102", "gfx1103", "gfx1150", "gfx1151", + "gfx1202", "gfx1210", "gfx942", "gfx1030"}) { + CAPTURE(std::string(name)); + CHECK_FALSE(GcnArchNameHasQuantWmma(name)); + } + CHECK_FALSE(vt::rocm::GcnArchNameIsGfx12PrefillWmma("gfx1100")); + CHECK_FALSE(vt::rocm::GcnArchNameIsGfx12PrefillWmma("gfx1100:xnack-")); +} + TEST_CASE("the parse is constexpr, so a wrong answer is a compile error") { // Not decoration: it is what lets the capability be asserted without a device. static_assert(CapabilityFromGcnArch("gfx1100")->first == 11); diff --git a/tools/rocm_quant_wmma/capture.cpp b/tools/rocm_quant_wmma/capture.cpp new file mode 100644 index 000000000..29a88d22b --- /dev/null +++ b/tools/rocm_quant_wmma/capture.cpp @@ -0,0 +1,133 @@ +// Native adapter for plugin d4c1f0d tests/test_kernels.py::test_mmq. +// Input generation and the unchanged upstream tolerances live in primary.py. +#include +#include +#include +#include +#include +#include +#include +#include +#include "vt/backend.h" +#include "vt/ops.h" + +#ifdef VLLM_CPP_HIP +namespace vt::rocm { +std::string DeviceArchName(int); +uint64_t KQuantWmmaDispatchCount(); +uint64_t KQuantWmmaQ4KDispatchCount(); +} // namespace vt::rocm +#endif + +namespace { +using Json = nlohmann::json; +using Path = std::filesystem::path; +std::string Read(const Path& path) { + std::ifstream file(path, std::ios::binary); + if (!file) throw std::runtime_error("cannot read " + path.string()); + return {std::istreambuf_iterator(file), std::istreambuf_iterator()}; +} +void Write(const Path& path, const std::string& bytes) { + std::ofstream file(path, std::ios::binary); + file.write(bytes.data(), static_cast(bytes.size())); + if (!file) throw std::runtime_error("cannot write " + path.string()); +} + +#ifdef VLLM_CPP_HIP +constexpr vt::Device kDevice{vt::DeviceType::kROCM, 0}; +class Storage { + public: + Storage() : backend(vt::GetBackend(kDevice.type)), queue(backend.CreateQueue()) {} + ~Storage() { + for (auto* allocation : allocations) backend.Free(allocation); + backend.DestroyQueue(queue); + } + void* Allocate(size_t size) { + auto* pointer = backend.Alloc(size); + allocations.push_back(pointer); + return pointer; + } + vt::Tensor Upload(const std::string& bytes, vt::DType dtype, int64_t rows, int64_t width) { + auto* pointer = Allocate(bytes.size()); + backend.Copy(queue, pointer, bytes.data(), bytes.size()); + return vt::Tensor::Contiguous(pointer, dtype, kDevice, {rows, width}); + } + vt::Backend& backend; + vt::Queue queue; + std::vector allocations; +}; + +void Capture(const Path& input, const Path& output) { + const auto arch = vt::rocm::DeviceArchName(0); + const auto stem = arch.substr(0, arch.find(':')); + if (stem != "gfx1100" && stem != "gfx1200" && stem != "gfx1201") + throw std::runtime_error("capture requires an admitted physical WMMA device"); + const auto manifest = Json::parse(Read(input)); + std::filesystem::create_directories(output); + Json report = {{"architecture", arch}, {"cases", Json::array()}}; + for (const auto& entry : manifest.at("cases")) { + const auto name = entry.at("name").get(); + const int type = entry.at("type").get(); + if (type != 12 && type != 14) throw std::runtime_error("only Q4_K/Q6_K are in scope"); + const auto dtype = type == 12 ? vt::DType::kQ4_K : vt::DType::kQ6_K; + const int64_t m = entry.at("tokens").get(); + const int64_t n = entry.at("rows").get(); + const int64_t k = entry.at("width").get(); + const auto input_dtype = entry.at("dtype").get(); + const auto act_dtype = input_dtype == "f16" ? vt::DType::kF16 : + input_dtype == "bf16" ? vt::DType::kBF16 : vt::DType::kF32; + if (input_dtype != "f16" && input_dtype != "bf16" && input_dtype != "f32") + throw std::runtime_error("invalid activation dtype"); + // f32 is the harness adaptation for F16 output, absent from the native ABI. + // BF16 uses its native two-byte output. primary.py narrows only F16 cases. + const auto out_dtype = act_dtype == vt::DType::kBF16 ? vt::DType::kBF16 : vt::DType::kF32; + const auto packed = Read(input.parent_path() / entry.at("packed").get()); + const auto activation = Read(input.parent_path() / entry.at("activation").get()); + if (m <= 0 || n <= 0 || k <= 0 || k % 256 != 0 || + packed.size() != static_cast(n) * vt::RowSizeBytes(dtype, k) || + activation.size() != static_cast(m * k) * vt::SizeOf(act_dtype)) + throw std::runtime_error("fixture geometry or bytes differ"); + Storage storage; + auto a = storage.Upload(activation, act_dtype, m, k); + auto b = storage.Upload(packed, dtype, n, k); + std::string result(static_cast(m * n) * vt::SizeOf(out_dtype), '\0'); + auto out = vt::Tensor::Contiguous(storage.Allocate(result.size()), out_dtype, kDevice, {m, n}); + const auto count = type == 12 ? vt::rocm::KQuantWmmaQ4KDispatchCount : vt::rocm::KQuantWmmaDispatchCount; + const auto before = count(); + const auto reference_before = vt::GetReferenceTierHits(); + vt::MatmulBTQuant(storage.queue, out, a, b); + storage.backend.Copy(storage.queue, result.data(), out.data, result.size()); + storage.backend.Synchronize(storage.queue); + const auto calls = count() - before; + const char* knob = std::getenv("VT_ROCM_QUANT_WMMA"); + const bool expected = m >= 16 && n >= 16 && (knob == nullptr || std::string(knob) != "0"); + if (calls != static_cast(expected) || vt::GetReferenceTierHits() != reference_before) + throw std::runtime_error("native WMMA dispatch or reference fallback mismatch: " + name); + const auto result_name = name + ".bin"; + Write(output / result_name, result); + report["cases"].push_back({{"name", name}, {"dtype", vt::Name(out_dtype)}, + {"wmma_calls", calls}, {"result", result_name}}); + std::printf("NATIVE %s wmmas=%llu\n", name.c_str(), static_cast(calls)); + } + if (report["cases"].empty()) throw std::runtime_error("empty input manifest"); + Write(output / "report.json", report.dump(2) + "\n"); +} +#endif +} // namespace + +int main(int argc, char** argv) { + try { +#ifdef VLLM_CPP_HIP + if (argc != 3) throw std::runtime_error("use MANIFEST OUTPUT"); + Capture(argv[1], argv[2]); + return 0; +#else + (void)argc; + (void)argv; + throw std::runtime_error("native capture requires a HIP build"); +#endif + } catch (const std::exception& error) { + std::fprintf(stderr, "CAPTURE FAILED: %s\n", error.what()); + return 1; + } +} diff --git a/tools/rocm_quant_wmma/primary.py b/tools/rocm_quant_wmma/primary.py new file mode 100644 index 000000000..05d1eac12 --- /dev/null +++ b/tools/rocm_quant_wmma/primary.py @@ -0,0 +1,151 @@ +#!/usr/bin/env python3 +"""Port plugin d4c1f0d tests/test_kernels.py::test_mmq for Q4_K and Q6_K. + +Preserve every original tensor, tokens 7/83/128/2048, widths 256/1024, +F16/BF16/F32 inputs, seed zero, and the original dtype-specific tolerances. +The sole dtype adaptation narrows native F32 output to F16 after capture, +because the native MatmulBTQuant contract has no F16 output. +Run capture inside the task's pinned primary runtime under the GPU mutex. +""" +from __future__ import annotations + +import argparse +import hashlib +import json +from pathlib import Path +import random + +PIN = "d4c1f0d082fc7cd4350da56689109a01c1f29d6c" +ATOL = {"f16": 1, "bf16": 1.5, "f32": 1.2} +RTOL = {"f16": 1e-1, "bf16": 1e4, "f32": 2e1} + + +def seal(path: Path) -> dict: + return {"bytes": path.stat().st_size, "sha256": hashlib.sha256(path.read_bytes()).hexdigest()} + + +def raw(tensor) -> bytes: + import torch + return tensor.contiguous().view(torch.uint8).cpu().numpy().tobytes() + + +def capture(fixture_manifest: Path, fixtures: Path, output: Path) -> None: + import gguf + import numpy as np + import torch + import vllm_gguf_plugin.ops as ops + + source = json.loads(fixture_manifest.read_text()) + expected_tensors = [entry for entry in source["cases"] if entry["type"] in (12, 14)] + if len(expected_tensors) != 20: + raise ValueError("the original Q4_K/Q6_K corpus must contain all 20 tensors") + output.mkdir(parents=True, exist_ok=True) + manifest = {"version": 1, "plugin_pin": PIN, "fixture_manifest": seal(fixture_manifest), + "torch": torch.__version__, "torch_git": torch.version.git_version, + "hip": torch.version.hip, "architecture": torch.cuda.get_device_properties(0).gcnArchName, + "ops_source": seal(Path(ops.__file__)), "fixtures": {}, "cases": []} + for quant_name in ("Q4_K", "Q6_K"): + quant_type = gguf.GGMLQuantizationType[quant_name] + for hidden in (256, 1024): + filename = f"Quant_{quant_name}_{hidden}.gguf" + fixture = fixtures / filename + expected = source["fixtures"][filename] + if seal(fixture) != expected: + raise ValueError(f"original fixture hash or size differs: {filename}") + manifest["fixtures"][filename] = expected + tensors = gguf.GGUFReader(fixture).tensors + selected = [entry for entry in expected_tensors if entry["source"] == filename] + if [tensor.name for tensor in tensors] != [entry["tensor"] for entry in selected]: + raise ValueError("original tensor names or ordering differs") + for name, dtype in (("f16", torch.half), ("bf16", torch.bfloat16), ("f32", torch.float32)): + for tokens in (7, 83, 128, 2048): + # Exact upstream seed_everything and rand device/dtype. + random.seed(0) + np.random.seed(0) + torch.manual_seed(0) + x = torch.rand((tokens, hidden), dtype=dtype, device="cuda") + activation_name = f"{quant_name}-{hidden}-{name}-{tokens}.activation" + (output / activation_name).write_bytes(raw(x)) + for index, tensor in enumerate(tensors): + case_name = f"{quant_name}-{hidden}-{index}-{name}-{tokens}" + packed_name = f"{quant_name}-{hidden}-{index}.packed" + packed_path = output / packed_name + packed_path.write_bytes(tensor.data.tobytes()) + with torch.inference_mode(): + weight = torch.tensor(gguf.dequantize(tensor.data, quant_type), device="cuda").to(dtype) + reference = x @ weight.T + packed = torch.tensor(tensor.data, device="cuda") + actual = ops.ggml_mul_mat_a8(packed, x, int(quant_type), packed.shape[0]).to(dtype) + torch.testing.assert_close(actual, reference, atol=ATOL[name], rtol=RTOL[name]) + reference_name = case_name + ".reference" + primary_name = case_name + ".primary" + (output / reference_name).write_bytes(raw(reference)) + (output / primary_name).write_bytes(raw(actual)) + manifest["cases"].append({"name": case_name, "type": int(quant_type), + "source": filename, "tensor": tensor.name, + "tokens": tokens, "rows": int(packed.shape[0]), "width": hidden, + "dtype": name, "seed": 0, "packed": packed_name, + "packed_seal": seal(packed_path), "activation": activation_name, + "activation_seal": seal(output / activation_name), + "reference": reference_name, "primary": primary_name, + "reference_seal": seal(output / reference_name), + "primary_seal": seal(output / primary_name), + "native_extension": ops._cuda_gemm_kernel_available("ggml_mul_mat_a8", int(quant_type)), + "atol": ATOL[name], "rtol": RTOL[name]}) + print(f"PRIMARY {case_name} PASS", flush=True) + if len(manifest["cases"]) != 240: + raise ValueError("original coverage is not 240 tensor cases") + (output / "manifest.json").write_text(json.dumps(manifest, indent=2) + "\n") + + +def compare(manifest_path: Path, native: Path) -> None: + import torch + + manifest = json.loads(manifest_path.read_text()) + report = json.loads((native / "report.json").read_text()) + actual_cases = {entry["name"]: entry for entry in report["cases"]} + if len(actual_cases) != len(report["cases"]) or set(actual_cases) != {entry["name"] for entry in manifest["cases"]}: + raise ValueError("native case coverage differs from the complete primary manifest") + dtypes = {"f16": torch.half, "bf16": torch.bfloat16, "f32": torch.float32} + comparisons = [] + for case in manifest["cases"]: + entry = actual_cases[case["name"]] + for name in ("packed", "activation", "reference", "primary"): + if seal(manifest_path.parent / case[name]) != case[name + "_seal"]: + raise ValueError(f"changed {name} bytes: {case['name']}") + dtype = dtypes[case["dtype"]] + expected_native_dtype = "bf16" if case["dtype"] == "bf16" else "f32" + if entry["dtype"] != expected_native_dtype: + raise ValueError("native output dtype differs from the explicit adaptation") + result_path = native / entry["result"] + got = torch.frombuffer(bytearray(result_path.read_bytes()), dtype=dtypes[entry["dtype"]]).clone().to(dtype) + if got.numel() != case["tokens"] * case["rows"]: + raise ValueError("native output shape differs") + for field in ("reference", "primary"): + ref = torch.frombuffer(bytearray((manifest_path.parent / case[field]).read_bytes()), dtype=dtype).clone() + torch.testing.assert_close(got, ref, atol=ATOL[case["dtype"]], rtol=RTOL[case["dtype"]]) + comparisons.append({"name": case["name"], "verdict": "PASS", "result": seal(result_path), + "atol": ATOL[case["dtype"]], "rtol": RTOL[case["dtype"]]}) + print(f"NATIVE PARITY {case['name']} PASS", flush=True) + (native / "comparison.json").write_text(json.dumps({"verdict": "PASS", "cases": comparisons}, indent=2) + "\n") + + +def main() -> None: + parser = argparse.ArgumentParser(description=__doc__) + commands = parser.add_subparsers(dest="command", required=True) + run = commands.add_parser("capture") + run.add_argument("fixture_manifest", type=Path) + run.add_argument("fixtures", type=Path) + run.add_argument("output", type=Path) + check = commands.add_parser("compare") + check.add_argument("manifest", type=Path) + check.add_argument("native", type=Path) + args = parser.parse_args() + if args.command == "capture": + capture(args.fixture_manifest, args.fixtures, args.output) + else: + compare(args.manifest, args.native) + + +if __name__ == "__main__": + main()