Repository navigation
HIP follow-up: native-pack fixes, type-aware PLE rows, and clean test provisioning - #160
Closed
samuelishida wants to merge 9 commits into
Closed
samuelishida wants to merge 9 commits into
samuelishida wants to merge 9 commits into
Conversation
bartowski's IQ2_XS keeps token_embd.weight as Q2_K (GGML type 10), which the native embedding path refused because it consulted the MMVQ set. Add Q2_K to a separate dequant/gather set instead of the MMVQ set: the flat dequantizers, iq_embed_rows and iq_row_bytes accept it, iq_supported/iq_mmvq keep refusing it, and NativeEmbed::load gates on the dequant set (the WIP rule: a dequantizer that cannot run must fail loudly, and one that can must not be hidden behind an unrelated capability check). The decoder is llama.cpp's Q2_K dequantizer with its 64 threads folded onto 32, the same form as the other block types in this file. tests/hip/q2k_embed_parity.cpp checks the decoder against a hand-computed reference (including that the two gather rows are reordered, not replayed) and, with two file arguments, against real rows of the IQ2_XS shard compared to gguf.dequantize. Verified on the RX 7900 XTX: 0 mismatches on synthetic rows and 0 on model rows 0 and 248319.
The PLE table `per_layer_token_embd.weight` is IQ4_NL in the original pack, Q4_0 in Bartowski's IQ2_XS GGUF and Q5_1 in the AD IQ4_XS-M64 checkpoint. The reader hardcoded IQ4_NL and 90-byte rows, so any other encoding was refused at open (IQ2_XS) or read with the wrong geometry. Decode and row geometry now follow the GGUF tensor type: - `ple_row_bytes(type)`: Q4_0/IQ4_NL 90 B, Q5_1 120 B; anything else is refused. - `ple_dequant_row(type, ...)`: codebook for IQ4_NL, `(q - 8) * d` for Q4_0, and the checkpoint's own Q5_1 layout (its qh bits 16..31 are the second half; the generic ggml Q5_1 helpers in this tree shift differently and are not reusable here). - the PleTable row buffers, the direct-reader stride, the row cache and the batched gather all take the row size from the table type, with overflow and short-read checks kept fail-closed. Synthetic PLE parity fixtures stay on the default 90-byte row; a real Q4_0 table opens and decodes through both the mmap and direct paths. Verified on the RX 7900 XTX with the IQ2_XS shard: table opens, 320001536 rows, prefill and decode run. ple_reader_selftest stays green.
…path A generate-mode run on a native (IQ) pack never zeroed its session state. The canonical path does it from `put_input` at position 0 (`session_zero`, with the embedding as `R`), but a native pack leaves that loop before the first `put_input` - it runs verify windows instead - and neither the batched prompt path nor the verifier zeroes anything. The state buffer is `cudaMalloc`'d, so the GDN recurrence, the QSA KV/indexer state, the PLE history and `R` all began at whatever the allocator last held: finite on a fresh allocation, overflow or NaN after reuse. The symptom was prompt-independent: every prompt decoded to the same token, the final residual held non-finite values up to 3.4e38 (94 of 10240 on a 5-token prompt), and it happened from the first token, with or without a prefill chunk. `--serve` zeroes exactly this state when `resume == 0`; generate mode always starts from an empty sequence, so it must too. Verified against llama.cpp (pinned 3cf03257f, the same GGUF) on the RX 7900 XTX: "The capital of France is" now decodes to " Paris. The capital of Germany" - the first four tokens are identical to llama.cpp's greedy tokens (11751 13 561 6511). The final residual is finite.
…e path A native pack decodes through verify windows only, and the three oracle dumps were all wired to the paths it does not take: `--dump-logits` wrote only the header (rows came from the ordinary loop), `--dump-final-r` wrote no frames at all, and `--dump-layers`'s ladder lives in `session_loop`, which a native pack never runs - and the two dump flags also kept the device residency table from being staged, so `--spec` was refused later with "--spec needs the device residency table". Together that gated the branch's own bisection tooling out of its only native configuration: the exact hole this review fell into, where a native-pack run could not produce a single comparable artifact and the root cause of a degenerate output survived until llama.cpp was built as an oracle. The residency table moves out from under the dump flags (it is the expert tier's state, the token graph keeps its own capture gate), and the verify windows now produce the oracles themselves: - `--dump-logits`: after each window, the accepted rows of the head's own logits buffer are copied back and written. Row t of a window at p is the distribution after the tokens at p..p+t - what the ordinary loop's row at p+t is; the rejected rows are conditioned on drafts the ordinary path would never feed and are skipped. `--check-logits` scans the same rows. - `--dump-layers`: `Verifier::set_layer_dump` arms a residual ladder captured into every window graph - one device-to-host copy per layer, row 0 the embedding broadcast, the same convention `session_loop` uses. The rows are the fused-layout residual (layer l's pending hyper-connection write is not folded in; the next fused read folds it) - bitwise what the default native decode path holds, which is the right oracle for tracing where two engine runs diverge, but not the canonical ladder: that needs the unfused canonical path. Refused on a layer split. - `--dump-final-r`: `final_R(t)` per accepted row, the same frames the ordinary loop writes. Two honesty fixes on the way. The logits header used to state a count computed from `n_prompt - 1 + max_new`, which is the truth only when the token loop ran every position: a batched prompt (the native packs' required configuration) or an early EOS stop writes fewer rows, and the header then described a different file from the one written - the same defect class `tools/logits_identical.py` already refuses. The close-time header is now rewritten with the rows the file actually holds. And a window that runs past the generation cap leaves its later rows conditioned on drafts the run never emitted; those rows are dropped so the file cannot read as "the text continued". `--dump-halves` under `--spec` is refused outright: its copies live in the captured canonical block path and the file would be frames of zeros - evidence that looks like evidence but is not. Measured on IQ2_XS (native pack, --spec 4, 8 tokens): 8 logits rows whose argmax replays the emitted tokens token for token, against llama.cpp (b3119-b6b, same GGUF) the top-1 agrees 8/8 with mean Spearman 0.733 over the reference's top-100 (the Q5_K head is quantized; the reference is f16); 8 ladder frames at positions [4, 12) with first frame position printed; final-R frames finite, max |R| 10.8.
The shipped table is calibrated for hipBLASLt 100100 and the loader refuses a mismatched version (solution IDs are scoped to the version), so on ROCm 7.2.1 the tuned dense-prefill path silently falls back to hipBLASEx. Recalibrated on the same RX 7900 XTX with the branch's own tuner over the same 30 (dtype, N, K, ldy, T) cases. 26 of the 30 cases produced a validated, timed solution and are written as-is. For 4 cases - (1,2560,1), (4,10240,4), (512,2560,512) and (2560,2560,2560), each at T=4096 only - every heuristic solution produced non-finite output on 100202 while the same shapes calibrated fine on 100100, so the tuner's gates refuse to write rows for them; the engine's lookup falls back to hipblasGemmEx for a missing row, which is what the shipped 100100 table had those shapes on anyway. The hip_prefill_hipblaslt_gemm parity test passes with this table (STRATA_HIPBLASLT_TUNING=tools/hip/gfx1100-hipblaslt-100202.txt): all cases PASS with relative_l2 <= 1e-4 and max_abs <= 5e-3, no padding writes.
Three suite-hygiene fixes, all with the same shape: a test whose inputs cannot
exist in the repo must not red-run for every consumer, and a test that CAN run
must not sit skipped behind an environment variable nobody sets.
- ple_parity: the test hard-failed for every fresh clone, because its block
fixtures (over 100 MB of dequantized f32 weights - GitHub's own per-file
limit) and its artifact-matched oracle include ("GENERATED by
tests/python/gen_ple_oracle.py", itself never committed, nor was the capture
tool ple_layer_xcheck.cpp it depends on) were never in the tree. The CMake
now adds the test only when the fixtures AND the artifact paths are
provisioned (STRATA_PLE_TEST_GGUF / STRATA_PLE_TEST_PACK), and the configure
step says exactly what is missing; a provisioned-but-wrong run still
hard-fails inside the test itself. The test also takes the pack from
$STRATA_PLE_PACK, mirroring how it already took the GGUF. The missing
supply chain is documented in tests/python/README.md instead of being
reinvented badly: a reviewer-side reconstruction of the ggml capture would
be a fixture that looks authoritative without being it, which is the exact
failure mode the project's own comments call out.
- hip_prefill_hipblaslt_gemm skipped unless STRATA_HIPBLASLT_TUNING was set,
even when a table for the runtime's own hipBLASLt version ships in the repo.
It now falls back to tools/hip/<arch>-hipblaslt-<version>.txt (env still
wins), so on ROCm 7.2.1 with the committed 100202 table the test runs
unconditionally and passes (0.3 s, all cases within tolerances).
- docs/AMD_HIP.md documents the native-pack oracle dumps (their spec-window
row semantics, the fused-layout caveat on the ladder) and
docs/AMD_HIP_PERFORMANCE.md the 100202 recalibration: 26 of 30 cases
validated; four shapes at T=4096 produce non-finite output in every
hipBLASLt solution on 100202 and are intentionally absent, with the runtime
falling back to hipblasGemmEx for missing rows.
Full suite on both configured builds: build-hip 30/30, build-hip-mmq 31/31,
nothing skipped.
- .gitignore: ignore the local .pi/ harness metadata directory. - src/program/generate.cpp: remove the unused `argmax()` helper (pre-existing warning in both HIP builds); use `sizeof(float)` instead of the literal 4 in the verify-window logits device-to-host copy. - tests/hip/q2k_embed_parity.cpp: clarify the support-set contract check so it reads as an assertion, not a riddle; add trailing newline; wrap hipFree calls in HIP_CHECK to silence the new `[[nodiscard]]` warnings. No functional change; both HIP builds remain warning-free for the touched files and ctest is still 30/30 and 31/31.
The tensor-core prompt-attention kernel is already refused at runtime on AMD (line 695), but the file unconditionally included <math_constants.h> for CUDART_INF_F, which is not present on a ROCm-only build. Replace it with the portable INFINITY macro from <cmath> so the HIP build succeeds without a CUDA toolkit installed.
Trim the all-caps, rant-like asides in generate.cpp, verify.cpp and verify.hpp so the diagnostic-dump and session-zero explanations read as factual notes rather than bug-diaries. Shorten the CMake and hipblaslt-test provisioning comments similarly.
Owner
|
Thanks @samuelishida, great work validating native packs on ROCm 7.2.1. We'll take the
Splitting it into 3-4 PRs (session zeroing, PLE/Q2_K types, diagnostics, the hipBLASLt table) would make the gate much quicker. |
Owner
Contributor
Author
|
You're welcome! This was one of the most exciting inference projects I've seen recently, so I'm more than happy to contribute. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Closed in favor of the split PRs requested during review:
CUDART_INF_F→INFINITY(fixes HIP backend build failure #157)--spec+ple_parityprovisioningThe two review fixes (restore verify loop order; keep host_res gate for --no-capture/--no-token-graph) are in #169.