From 727e7ae446f14bfd9b72622d3969b4f0a2f07331 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Wed, 15 Jul 2026 12:49:12 +0100 Subject: [PATCH 1/2] cuda: speed up Q1_0 extraction with byte permutes --- ggml/src/ggml-cuda/mmq.cuh | 31 ++++++--------- ggml/src/ggml-cuda/vecdotq.cuh | 72 ++++++++++++++++++++-------------- 2 files changed, 54 insertions(+), 49 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index f730c6f4de66..a5c7064386ab 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -350,29 +350,23 @@ template static __device__ __forceinline__ void loa } 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; const int dst_offset = kbx*(scale_entries_per_block*QI8_0) + kqsx*QI8_0; #pragma unroll - for (int j = 0; j < 8; ++j) { + for (int j = 0; j < 2; ++j) { + const int4 v = unpack_q1_0_bytes(qxi[j]); + #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) - x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + dst_offset + j] = unpacked_bytes[j]; + x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + dst_offset + j*4+0] = v.x; + x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + dst_offset + j*4+1] = v.y; + x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + dst_offset + j*4+2] = v.z; + x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + dst_offset + j*4+3] = v.w; #else - x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j] = unpacked_bytes[j]; + x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+0] = v.x; + x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+1] = v.y; + x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+2] = v.z; + x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j*4+3] = v.w; #endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) } } @@ -4279,4 +4273,3 @@ void ggml_cuda_op_mul_mat_q( const int64_t src1_padded_row_size, cudaStream_t stream); bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t n_experts); - diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index 24cf8edb7de6..d638d823355c 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -675,6 +675,32 @@ static __device__ __forceinline__ float vec_dot_q6_K_q8_1_impl_mmq( return d6 * sumf_d; } +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 int s0 = __byte_perm(0x01FF, 0x01FF, n0 >> 0); + const int s1 = __byte_perm(0x01FF, 0x01FF, n1 >> 0); + const int s2 = __byte_perm(0x01FF, 0x01FF, n0 >> 16); + const int s3 = __byte_perm(0x01FF, 0x01FF, n1 >> 16); + + return make_int4(__byte_perm(s0, s1, 0x5410), __byte_perm(s0, s1, 0x7632), __byte_perm(s2, s3, 0x5410), + __byte_perm(s2, s3, 0x7632)); +#else + int values[4]; +#pragma unroll + for (int j = 0; j < 4; ++j) { + const int bits4 = (q >> (4 * j)) & 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; + values[j] = (b0 & 0xFF) | ((b1 & 0xFF) << 8) | ((b2 & 0xFF) << 16) | ((b3 & 0xFF) << 24); + } + return make_int4(values[0], values[1], values[2], values[3]); +#endif +} + static __device__ __forceinline__ float vec_dot_q1_0_q8_1( const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { @@ -684,44 +710,30 @@ static __device__ __forceinline__ float vec_dot_q1_0_q8_1( // Q8_1: 32 elements per block with individual scales // iqs selects which of the 4 chunks of 32 elements to process (0-3) - const float d1 = bq1_0->d; + const float d1 = bq1_0->d; + const int16_t * qs = (const int16_t *) bq1_0->qs + iqs * 2; // Process only the chunk specified by iqs const block_q8_1 * bq8_1_chunk = bq8_1 + iqs; - // Load 32 bits (4 bytes) for this chunk from Q1_0 - const int offset = iqs * 4; - const int v = bq1_0->qs[offset + 0] | (bq1_0->qs[offset + 1] << 8) | - (bq1_0->qs[offset + 2] << 16) | (bq1_0->qs[offset + 3] << 24); - - // Unpack 32 bits into 32 raw UNSIGNED {0,1} lanes -- no per-element sign - // materialization. Symbol = 2*bit - 1, so sum(symbol*act) = 2*sum(bit*act) - // - sum(act); that affine correction is applied once at the end instead - // (matches the deferred-correction pattern vec_dot_q4_0_q8_1_impl uses). - int vi_bytes[8]; -#pragma unroll - for (int j = 0; j < 8; ++j) { - const int shift = j * 4; - const int bits4 = (v >> shift) & 0x0F; - const int b0 = (bits4 >> 0) & 1; - const int b1 = (bits4 >> 1) & 1; - const int b2 = (bits4 >> 2) & 1; - const int b3 = (bits4 >> 3) & 1; - vi_bytes[j] = b0 | (b1 << 8) | (b2 << 16) | (b3 << 24); - } - - // Compute dot product for this 32-element chunk int sumi = 0; #pragma unroll - for (int j = 0; j < 8; ++j) { - const int u = get_int_b4(bq8_1_chunk->qs, j); - sumi = ggml_cuda_dp4a(vi_bytes[j], u, sumi); + for (int j = 0; j < 2; ++j) { + const int4 v = unpack_q1_0_bytes(qs[j]); + + const int u0 = get_int_b4(bq8_1_chunk->qs, j * 4 + 0); + const int u1 = get_int_b4(bq8_1_chunk->qs, j * 4 + 1); + const int u2 = get_int_b4(bq8_1_chunk->qs, j * 4 + 2); + const int u3 = get_int_b4(bq8_1_chunk->qs, j * 4 + 3); + + sumi = ggml_cuda_dp4a(v.x, u0, sumi); + sumi = ggml_cuda_dp4a(v.y, u1, sumi); + sumi = ggml_cuda_dp4a(v.z, u2, sumi); + sumi = ggml_cuda_dp4a(v.w, u3, sumi); } - // ds.x = d8 (per-block activation scale), ds.y = sum(act) in real units - // (see quantize_q8_1: y[ib].ds = make_half2(d, sum)). - const float2 ds8f = __half22float2(bq8_1_chunk->ds); - return d1 * (2.0f * sumi * ds8f.x - ds8f.y); + const float d8 = __low2float(bq8_1_chunk->ds); + return d1 * d8 * sumi; } static __device__ __forceinline__ float vec_dot_q2_0_q8_1( From e6fdfcc10545d6c9ce801875a8b62daa2dab65b8 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Thu, 16 Jul 2026 23:45:46 +0100 Subject: [PATCH 2/2] cuda: use unsigned halfword for Q1_0 byte-perm selectors 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. --- ggml/src/ggml-cuda/mmq.cuh | 2 +- ggml/src/ggml-cuda/vecdotq.cuh | 17 +++++++++-------- 2 files changed, 10 insertions(+), 9 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index a5c7064386ab..08d8e98dc2f7 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -350,7 +350,7 @@ template static __device__ __forceinline__ void loa } const block_q1_0 * bxi = (const block_q1_0 *) x + kbx0 + i*stride + kbx; - const int16_t * qxi = (const int16_t *) bxi->qs + kqsx * 2; + const uint16_t * qxi = (const uint16_t *) bxi->qs + kqsx * 2; const int dst_offset = kbx*(scale_entries_per_block*QI8_0) + kqsx*QI8_0; #pragma unroll diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index d638d823355c..b3bccb711e10 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -675,14 +675,15 @@ static __device__ __forceinline__ float vec_dot_q6_K_q8_1_impl_mmq( return d6 * sumf_d; } -static __device__ __forceinline__ int4 unpack_q1_0_bytes(const int16_t q) { +static __device__ __forceinline__ int4 unpack_q1_0_bytes(const uint16_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 int s0 = __byte_perm(0x01FF, 0x01FF, n0 >> 0); - const int s1 = __byte_perm(0x01FF, 0x01FF, n1 >> 0); - const int s2 = __byte_perm(0x01FF, 0x01FF, n0 >> 16); - const int s3 = __byte_perm(0x01FF, 0x01FF, n1 >> 16); + const uint32_t q32 = q; + const int n0 = __byte_perm(0x11100100, 0x11100100, q32 >> 0); + const int n1 = __byte_perm(0x11100100, 0x11100100, q32 >> 2); + const int s0 = __byte_perm(0x01FF, 0x01FF, n0 >> 0); + const int s1 = __byte_perm(0x01FF, 0x01FF, n1 >> 0); + const int s2 = __byte_perm(0x01FF, 0x01FF, n0 >> 16); + const int s3 = __byte_perm(0x01FF, 0x01FF, n1 >> 16); return make_int4(__byte_perm(s0, s1, 0x5410), __byte_perm(s0, s1, 0x7632), __byte_perm(s2, s3, 0x5410), __byte_perm(s2, s3, 0x7632)); @@ -711,7 +712,7 @@ static __device__ __forceinline__ float vec_dot_q1_0_q8_1( // iqs selects which of the 4 chunks of 32 elements to process (0-3) const float d1 = bq1_0->d; - const int16_t * qs = (const int16_t *) bq1_0->qs + iqs * 2; + const uint16_t * qs = (const uint16_t *) bq1_0->qs + iqs * 2; // Process only the chunk specified by iqs const block_q8_1 * bq8_1_chunk = bq8_1 + iqs;