Skip to content

cuda: speed up Q1_0 extraction with byte permutes #73

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
khosravipasha merged 2 commits into from
Jul 17, 2026
+55 −49
Merged

cuda: speed up Q1_0 extraction with byte permutes #73

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
727e7ae
cuda: speed up Q1_0 extraction with byte permutes
bri-prism Jul 15, 2026
e6fdfcc
cuda: use unsigned halfword for Q1_0 byte-perm selectors
bri-prism Jul 16, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
31 changes: 12 additions & 19 deletions ggml/src/ggml-cuda/mmq.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -350,29 +350,23 @@ template <int mmq_y, bool need_check> 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 uint16_t * qxi = (const uint16_t *) bxi->qs + kqsx * 2;

Comment on lines 352 to 354
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)
}
}
Expand Down Expand Up @@ -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);

73 changes: 43 additions & 30 deletions ggml/src/ggml-cuda/vecdotq.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -675,6 +675,33 @@ 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 uint16_t q) {
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
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));
#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) {

Expand All @@ -684,44 +711,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 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;

// 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(
Expand Down
Loading