wmantly/llama.cpp-xrip
1
0 Fork 0
Code Issues Pull Requests Actions 31 Packages Projects Releases Wiki Activity

cuda: drop SM75 Q4_K decode reuse, superseded by upstream branchless unpack

Browse Source 
Same story as the packed scale loads. The Q4_K reuse path decoded a block once
and applied it across ncols_dst columns, which beat the old branchy generic
path because that path re-ran the scale unpack per column. Upstream's
branchless rewrite hoists the unpack out of the loop, and the generic path is
now the faster of the two.

Measured on a CMP 50HX (cc 7.5), test-backend-ops perf, medians of 5 reps after
a discarded warmup, two independent runs:

  reuse disabled vs enabled, q4_K n=2..4:  -0.52% and -0.54% (faster without)
  n=5..8, identical code in both builds:   -0.02%  (noise check)
  widening the gate to n<=8 instead:       +0.80%  (worse)

14 of 15 cases at n=2..4 favour removal, so the effect is consistent rather
than drift.

The IQ4_XS reuse path is deliberately kept: upstream's change only touched
Q4_K/Q5_K scale unpacking, IQ4_XS decodes via kvalues_iq4nl table lookups and
its generic path did not improve. Measured with GGML_SM75_IQ4_REUSE=0, removing
it costs +1.14% at n=2..4, so it stays.
This commit is contained in:
xrip committed 2026-09-21 20:52:02 +03:00
1 parent 12877b6224
commit 6f65493dde
2 files changed
-80
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

ggml/src/ggml-cuda/mmvq.cu
-24
View File
@@ -745,30 +745,6 @@ static __global__ void mul_mat_vec_q(
}
#endif
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_TURING
if constexpr (type == GGML_TYPE_Q4_K && ncols_dst >= 2 && ncols_dst <= 4) {
#pragma unroll
for (int i = 0; i < rows_per_cuda_block; ++i) {
const int bx = kbx_offset + i * stride_row_x + kbx;
const sm75_q4_K_decoded w = sm75_q4_K_decode(vx, bx, kqs);
#pragma unroll
for (int j = 0; j < ncols_dst; ++j) {
tmp[j][i] += sm75_q4_K_dot(w, &y[j * stride_col_y + kby], kqs);
}
if constexpr (has_fusion) {
if (use_gate) {
const sm75_q4_K_decoded gate = sm75_q4_K_decode(vgate, bx, kqs);
#pragma unroll
for (int j = 0; j < ncols_dst; ++j) {
tmp_gate[j][i] += sm75_q4_K_dot(gate, &y[j * stride_col_y + kby], kqs);
}
}
}
}
continue;
}
#endif
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK
// start the next iterations' weight loads early
if constexpr (mmvq_should_prefetch(type)) {
ggml/src/ggml-cuda/vecdotq.cuh
-56
View File
@@ -916,62 +916,6 @@ static __device__ __forceinline__ float vec_dot_q3_K_q8_1(
return vec_dot_q3_K_q8_1_impl_mmvq(vl, vh, u, bq3_K->scales, scale_offset, d, d8);
}
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_TURING
struct sm75_q4_K_decoded {
int values[4];
uint16_t scales;
uint16_t mins;
float2 dm;
};
static __device__ __forceinline__ sm75_q4_K_decoded sm75_q4_K_decode(const void * vx, int kbx, int iqs) {
static_assert(QR4_K == 2 && QI8_1 == 8 && VDR_Q4_K_Q8_1_MMVQ == 2, "Q4_K reuse expects two words per lane");
const block_q4_K * b = (const block_q4_K *) vx + kbx;
const int offset = QR4_K * ((iqs / 2) / (QI8_1 / 2));
const int * q = (const int *) (b->qs + 16 * offset + 4 * ((iqs / 2) % 4));
sm75_q4_K_decoded w;
#pragma unroll
for (int i = 0; i < QR4_K; ++i) {
w.values[2 * i] = (q[0] >> (4 * i)) & 0x0f0f0f0f;
w.values[2 * i + 1] = (q[4] >> (4 * i)) & 0x0f0f0f0f;
}
const uint32_t * scales = (const uint32_t *) b->scales;
const int j = offset / 2;
const int shift = 16 * (j & 1);
const uint16_t s0 = scales[0] >> shift;
const uint16_t s1 = scales[1] >> shift;
if (j < 2) {
w.scales = s0 & 0x3f3f;
w.mins = s1 & 0x3f3f;
} else {
const uint16_t s2 = scales[2] >> shift;
w.scales = (s2 & 0x0f0f) | ((s0 & 0xc0c0) >> 2);
w.mins = ((s2 >> 4) & 0x0f0f) | ((s1 & 0xc0c0) >> 2);
}
w.dm = __half22float2(b->dm);
return w;
}
static __device__ __forceinline__ float sm75_q4_K_dot(const sm75_q4_K_decoded & w, const block_q8_1 * y, int iqs) {
const int offset = QR4_K * ((iqs / 2) / (QI8_1 / 2));
float sum_d = 0.0f;
float sum_m = 0.0f;
#pragma unroll
for (int i = 0; i < QR4_K; ++i) {
const block_q8_1 * b = y + offset + i;
const int * q = (const int *) b->qs + ((iqs / 2) % 4);
const int dot = ggml_cuda_dp4a(w.values[2 * i + 1], q[4], ggml_cuda_dp4a(w.values[2 * i], q[0], 0));
const int sum = ggml_cuda_dp4a(0x01010101, q[4], ggml_cuda_dp4a(0x01010101, q[0], 0));
const int scale = (w.scales >> (8 * i)) & 0xff;
const int min = (w.mins >> (8 * i)) & 0xff;
const float d = __low2float(b->ds);
sum_d += d * (dot * scale);
sum_m += d * (sum * min);
}
return w.dm.x * sum_d - w.dm.y * sum_m;
}
#endif
static __device__ __forceinline__ float vec_dot_q4_K_q8_1(
const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {