rocm: vectorized HIP PTQ1_0 vec_dot - #211
Conversation
The model file stores weights as densely packed base-3 digits (trits). To use them, the GPU must unpack each trit and multiply-accumulate it with the activation. The old AMD code did this one trit at a time in a long serial chain (128 slow steps per block). The new code unpacks 4 trits at once using the GPU's byte-shuffle and 4-wide dot-product instructions. Same math, ~3x less instruction work, which is why decode got ~88% faster. A wrong guess about a shuffle instruction's argument order was caught by on-device probes and unit tests before it could do harm; the corrected convention is now documented in a code comment.
Ankk98
changed the title
bri-prism
left a comment
bri-prism
left a comment
There was a problem hiding this comment.
Agent review: posted by the maintainer's coding agent at their request.
No findings in this source pass. The four-byte trit extraction, element mapping and integer activation-sum correction are consistent with the previous scalar HIP path.
HIP compilation and device numerical tests were not run here. Before merging, retain a comparison against the scalar path that covers all packed-byte values and the qh tail on the supported AMD targets.
Reviewed commit: e435472d5391720bda720fbbbebf8d75938f8fca.
| const int q = (int) __builtin_amdgcn_perm(w_hi, w_lo, 0x07050301); | ||
| sumi[e >> 5] += ggml_cuda_dp4a(q, u, 0); | ||
| sumu[e >> 5] += ggml_cuda_dp4a(0x01010101, u, 0); |
There was a problem hiding this comment.
Addressed the comment, extended tests/test-ptq1_0-cuda-dot.cpp with:
- a host transcription of the HIP vector idiom, unsigned {0,1,2} digits, bias folded into activation sums, qh tail in the unsigned scheme;
- a hip_perm emulator encoding the device-proven rule, including the non-obvious part I verified while writing it: selectors are byte-granular ((s >> 8*i) & 7), not nibbles, nibble indexing contradicts the [b0,0,b1,0] comments and fails everything, byte indexing passes everything;
- per-trial vec-vs-reference integer checks plus an exhaustive sweep (all 256 byte values × 6 packed words + 2 qh bytes × 4 chunks = 8192 exact checks).
Results: 20000 exact integer checks, 0 mismatches; 8192 exhaustive, 0 mismatches; exit 0. Device side stays covered by test-backend-ops on gfx1151.
| // Four packed bytes advance in the low bytes of 16-bit lanes, so one | ||
| // 32-bit multiply steps four trit streams at once (3*255 < 2^16, no | ||
| // cross-lane carry). Digits come out as {0,1,2}; the -1 folds into the | ||
| // activation sums, which a per-byte subtract cannot do borrow-free. | ||
| // NB: __builtin_amdgcn_perm picks sel bit2=0 from the SECOND arg, the | ||
| // opposite of CUDA __byte_perm, so the operands are swapped vs _multi. |
Transcribe the vector idiom (unsigned digits, bias fold, swapped-operand perm, qh tail) and assert it against the scalar reference over all 256 byte values at every word position. Unwrap the vecdotq comment to one sentence per line. Assisted-by: Opencode with Muse Spark
bri-prism
left a comment
bri-prism
left a comment
There was a problem hiding this comment.
Re-reviewed at 85a6c5f. The vectorization is correct on my read, the test is a real gate, and there is one evidentiary gap plus a sequencing hazard with #215.
The arithmetic checks out. I traced it by hand rather than trusting the commit message. Folding the bias as sum(q*a) - sum(a) with unsigned {0,1,2} digits is algebraically equivalent to the signed form. The element mapping is preserved: the scalar path's e = t*16 + m with m = 4w + j gives the same index the vector path computes as e = t*16 + 4w + j. The 0x07050301 selector places digit(b0..b3) in result bytes 0..3, which matches the little-endian order get_int_b4 loads the activations in. All three loops together touch each of the 128 elements exactly once, 80 plus 40 plus 8. No overflow is possible: sumu peaks near 32127 and sumi near 232*127, both far inside int32. The (e & 31) >> 2 group index stays 4-element aligned in both main loops. No findings there.
The byte-granular selector discovery is the valuable part of this PR and I am glad it is written down. Selectors being (s >> 8*i) & 7 rather than nibble-indexed, with sel bit2 = 0 choosing the second operand, is the kind of thing that costs someone a day.
Where I think the evidence is thinner than it looks.
hip_perm in the test is a host emulation of __builtin_amdgcn_perm, and vec_dot_ptq1_0_q8_1_vec is a transcription of the kernel rather than the kernel. That means the suite establishes that the algorithm is self-consistent given a belief about the intrinsic. It cannot establish the belief.
This matters more than usual here because of how the emulation was arrived at. You describe trying nibble indexing, finding it fails everything, then adopting byte indexing because it passes everything. That is fitting the emulator until the algorithm agrees with the reference. It is a reasonable way to find the rule, but it means the emulator and the kernel now share whatever assumption makes them agree, and the host test can no longer discriminate between "this matches the hardware" and "these two copies match each other". If the real intrinsic differs from hip_perm in some case neither of us has thought of, every one of those 28192 checks still passes and the device still returns wrong numbers.
The one artifact that would settle it is the device run, and that is the one reported without numbers. The host results are given precisely, 20000 exact integer checks and 8192 exhaustive, 0 mismatches, exit 0. The device side is given as "stays covered by test-backend-ops on gfx1151", which is an assertion rather than a result. Could you post the actual test-backend-ops output for the PTQ1_0 cases at 85a6c5f, with the case count? My earlier note asked for validation on the supported AMD targets, and gfx1151 output would close it outright.
Two smaller things.
The exhaustive sweep sets all four bytes of a packed word to the same value via v * 0x01010101u, so a within-word byte transposition such as b0 against b1 would pass that loop unnoticed. The 5000 random trials do use independent bytes and compare against ref_dequant, so the case is covered overall and I am not asking for a change. It is only that the exhaustive loop alone proves less than its name suggests.
The test carries its own copy of the kernel body. A later edit to the HIP branch in vecdotq.cuh will not be caught unless someone remembers to edit the transcription too. A short comment at the top of the HIP branch pointing at tests/test-ptq1_0-cuda-dot.cpp would make that link visible to whoever touches it next.
Sequencing with #215. Both PRs modify vec_dot_ptq1_0_q8_1, on opposite branches of the same #if. This PR leaves the #else calling vec_dot_ptq1_0_q8_1_multi, which is correct. #215 replaces that same #else with return 0.0f, which I have flagged there as a regression that silently zeroes PTQ1_0 on MUSA. Neither branch contains the other, and the two versions of this one function differ by about fifty lines, so they will conflict. Whoever resolves that conflict could reinstate the MUSA stub without noticing it is a defect, since it looks like an intentional part of the other change. Worth agreeing an order before either lands.
Method note: source reading and hand-tracing at 85a6c5f. I did not build, and I have no AMD hardware, so nothing above is a device result.
|
Device code is byte-identical between e435472 and 85a6c5f (the delta is comment + host test only). So, this run covers both.
test-backend-ops test -b ROCm0 -j 16 -p 'type_a=ptq1_0' → 120/120 pass (skipped not supported combos excluded from the count):
For the conflicts among the branches, we can merge the #211 first (it's hand-verified and device-proven, #215 is still open against prism), then #215 can rebase onto it. We need to inform the branch owner and make sure that the conflict resolver must preserve the vec_dot_ptq1_0_q8_1_multi call in the #else, the return 0.0f variant must not survive the merge. We can loop in the #215 owner i.e. @professorpalmer @bri-prism Is this ok? |
bri-prism
left a comment
bri-prism
left a comment
There was a problem hiding this comment.
Agent review: posted by the maintainer's coding agent at their request.
Approving at 3a82de9. Thanks for posting the device output. The 120/120 PTQ1_0 test-backend-ops pass on gfx1151 checks __builtin_amdgcn_perm against the CPU reference on real hardware, which is the evidence I was asking for. The only change since 85a6c5f is the pointer comment, so that run covers this head.
On sequencing: yes, merging this first and then rebasing #215 onto it works for me, with one condition for whoever resolves the conflict. #215 (and #221, which carries it) now widens this branch to #if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA). This PR puts __builtin_amdgcn_perm inside that branch, and it is an AMD-only builtin, so keeping #215's condition over this body would break the MUSA build. After the rebase, the vectorized path needs to stay under GGML_USE_HIP only, with MUSA on the scalar loop, and the #else must not end up returning 0.0f for any backend that can reach it. @professorpalmer, flagging this for the #215 rebase.
Method note: source reading at 3a82de9 plus your posted device output. I did not build or run it.
|
Understood, and thank you for spelling out the condition. After #211 is on
#221 carries the same function, so the same split applies there on its next rebase. I will not rebase either branch onto this one before it lands. |
Overview
Through this branch I am trying to optimize Bonsai 2 27B to run efficiently on strix-halo through rocm path.
The model file stores weights as densely packed base-3 digits (trits). To use them, the GPU must unpack each trit and multiply-accumulate it with the activation. The old AMD code did this one trit at a time in a long serial chain (128 slow steps per block). The new code unpacks 4 trits at once using the GPU's byte-shuffle and 4-wide dot-product instructions. Same math, ~3x less instruction work, which is why decode got ~88% faster.
A wrong guess about a shuffle instruction's argument order was caught by on-device probes and unit tests before it could do harm; the corrected convention is now documented in a code comment.
Additional information
Bench summary (AC + high-perf,
-ngl 99 -fa on -ub 2048, cooled between runs)Test Hardware
Notes
Requirements