cuda: speed up Q1_0 extraction with byte permutes#73
Conversation
khosravipasha
left a comment
khosravipasha
left a comment
There was a problem hiding this comment.
Looks good, I was thinknig of grabbing the changes from main later but this branch probably hard to rebase in future so good to merge now.
Do we need to test more on the accuray side?
There was a problem hiding this comment.
Pull request overview
This PR optimizes CUDA Q1_0 unpacking by replacing scalar per-bit expansion with a __byte_perm-based helper, and reuses that helper in both the MMVQ (vecdot) and MMQ (tiling/prefill) paths while keeping scalar fallbacks for HIP/MUSA.
Changes:
- Add
unpack_q1_0_bytes()to expand 16 packed Q1_0 bits into 4xdp4a-ready int32s via__byte_perm(CUDA) or scalar expansion (HIP/MUSA). - Update
vec_dot_q1_0_q8_1to use the new helper and switch to a signed (+1/-1) dot product formulation. - Update
load_tiles_q1_0to populate the existing MMQ shared-memory layout using the same helper.
Reviewed changes
Copilot reviewed 2 out of 2 changed files in this pull request and generated 3 comments.
| File | Description |
|---|---|
| ggml/src/ggml-cuda/vecdotq.cuh | Introduces and uses unpack_q1_0_bytes() to speed up Q1_0 extraction in MMVQ vecdot. |
| ggml/src/ggml-cuda/mmq.cuh | Reuses the same unpack helper when loading Q1_0 tiles for MMQ kernels. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
| static __device__ __forceinline__ int4 unpack_q1_0_bytes(const int16_t q) { | ||
| #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) | ||
| const int n0 = __byte_perm(0x11100100, 0x11100100, q >> 0); | ||
| const int n1 = __byte_perm(0x11100100, 0x11100100, q >> 2); |
| const float d1 = bq1_0->d; | ||
| const int16_t * qs = (const int16_t *) bq1_0->qs + iqs * 2; | ||
|
|
| const block_q1_0 * bxi = (const block_q1_0 *) x + kbx0 + i*stride + kbx; | ||
| const int qs_offset = 4*kqsx; | ||
| const int qs0 = bxi->qs[qs_offset + 0] | (bxi->qs[qs_offset + 1] << 8) | | ||
| (bxi->qs[qs_offset + 2] << 16) | (bxi->qs[qs_offset + 3] << 24); | ||
|
|
||
| int unpacked_bytes[8]; | ||
| #pragma unroll | ||
| for (int j = 0; j < 8; ++j) { | ||
| const int shift = j * 4; | ||
| const int bits4 = (qs0 >> shift) & 0x0F; | ||
| const int b0 = (bits4 & 0x01) ? 1 : -1; | ||
| const int b1 = (bits4 & 0x02) ? 1 : -1; | ||
| const int b2 = (bits4 & 0x04) ? 1 : -1; | ||
| const int b3 = (bits4 & 0x08) ? 1 : -1; | ||
| unpacked_bytes[j] = (b0 & 0xFF) | ((b1 & 0xFF) << 8) | ((b2 & 0xFF) << 16) | ((b3 & 0xFF) << 24); | ||
| } | ||
| const int16_t * qxi = (const int16_t *) bxi->qs + kqsx * 2; | ||
|
|
unpack_q1_0_bytes() built the __byte_perm selectors by right-shifting a signed int16_t, so a packed halfword with the top bit set sign-extends through the shift and corrupts the selector nibbles. Take the packed word as uint16_t (widened through uint32_t before the shift) at the helper and at both the MMVQ and MMQ call sites.
|
This was also merged in llama.cpp for Q1_0, tested it there myself and was safe so should be good to mrege. |
|
@bri-prism the ones is llama.cpp does not have your changes on top though, are they importatn enough to send a follow up to llama.cpp? |
What
__byte_permoperations instead of scalar per-bit expansion.Why
Scalar Q1_0 expansion adds measurable overhead before the existing int8 dot product. On an H100 with a large Q1_0 model, replacing that expansion increased TG128 from 99.32 to 110.38 tok/s and PP512 from 2767.55 to 2843.85 tok/s in an interleaved A/B.
This adapts the extraction approach from ggml-org#25628 to this fork's pre-split MMQ layout.
How
unpack_q1_0_bytesloads 16 packed bits at a time, expands them into signed int8 lanes with byte permutations, and returns four packed values ready fordp4a.vec_dot_q1_0_q8_1uses the helper directly for MMVQ.load_tiles_q1_0writes the same values into the existing MMQ shared-memory layout, so the downstream matrix path is unchanged.HIP and MUSA select the original scalar signed expansion at compile time because
__byte_permis CUDA-specific.Test plan
sm_90.test-backend-opsMUL_MAT correctness: 43/43 supported cases passed against the CPU reference.