diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 69d7ae8067..ba5f673c73 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -937,6 +937,9 @@ static __device__ __forceinline__ uint2 fast_div_modulo(uint32_t n, const uint3 typedef void (*dequantize_kernel_t)(const void * vx, const int64_t ib, const int iqs, float2 & v); +template +using dequantize_kq_t = void (*)(const void * vx, const int64_t ib, dst_t * y, const int tid); + static __device__ __forceinline__ float get_alibi_slope( const float max_bias, const uint32_t h, const uint32_t n_head_log2, const float m0, const float m1 ) { diff --git a/ggml/src/ggml-cuda/convert.cu b/ggml/src/ggml-cuda/convert.cu index f04a2d5a2c..946e02af50 100644 --- a/ggml/src/ggml-cuda/convert.cu +++ b/ggml/src/ggml-cuda/convert.cu @@ -140,358 +140,107 @@ static __global__ void dequantize_block_q4_1(const void * __restrict__ vx, dst_t template static __global__ void dequantize_block_q2_K(const void * __restrict__ vx, dst_t * __restrict__ yy) { + const int64_t i = blockIdx.x; - const int64_t i = blockIdx.x; - const block_q2_K * x = (const block_q2_K *) vx; - - const int64_t tid = threadIdx.x; - const int64_t n = tid/32; - const int64_t l = tid - 32*n; - const int64_t is = 8*n + l/16; - - const uint8_t q = x[i].qs[32*n + l]; - dst_t * y = yy + i*QK_K + 128*n; - - float dall = __low2half(x[i].dm); - float dmin = __high2half(x[i].dm); - y[l+ 0] = ggml_cuda_cast(dall * (x[i].scales[is+0] & 0xF) * ((q >> 0) & 3) - dmin * (x[i].scales[is+0] >> 4)); - y[l+32] = ggml_cuda_cast(dall * (x[i].scales[is+2] & 0xF) * ((q >> 2) & 3) - dmin * (x[i].scales[is+2] >> 4)); - y[l+64] = ggml_cuda_cast(dall * (x[i].scales[is+4] & 0xF) * ((q >> 4) & 3) - dmin * (x[i].scales[is+4] >> 4)); - y[l+96] = ggml_cuda_cast(dall * (x[i].scales[is+6] & 0xF) * ((q >> 6) & 3) - dmin * (x[i].scales[is+6] >> 4)); + dequantize_q2_K(vx, i, yy + i*QK_K, threadIdx.x); } template static __global__ void dequantize_block_q3_K(const void * __restrict__ vx, dst_t * __restrict__ yy) { - const int64_t i = blockIdx.x; - const block_q3_K * x = (const block_q3_K *) vx; - const int64_t r = threadIdx.x/4; - const int64_t tid = r/2; - const int64_t is0 = r%2; - const int64_t l0 = 16*is0 + 4*(threadIdx.x%4); - const int64_t n = tid / 4; - const int64_t j = tid - 4*n; - - uint8_t m = 1 << (4*n + j); - int64_t is = 8*n + 2*j + is0; - int shift = 2*j; - - int8_t us = is < 4 ? (x[i].scales[is-0] & 0xF) | (((x[i].scales[is+8] >> 0) & 3) << 4) : - is < 8 ? (x[i].scales[is-0] & 0xF) | (((x[i].scales[is+4] >> 2) & 3) << 4) : - is < 12 ? (x[i].scales[is-8] >> 4) | (((x[i].scales[is+0] >> 4) & 3) << 4) : - (x[i].scales[is-8] >> 4) | (((x[i].scales[is-4] >> 6) & 3) << 4); - float d_all = x[i].d; - float dl = d_all * (us - 32); - - dst_t * y = yy + i*QK_K + 128*n + 32*j; - const uint8_t * q = x[i].qs + 32*n; - const uint8_t * hm = x[i].hmask; - - for (int l = l0; l < l0+4; ++l) { - y[l] = ggml_cuda_cast(dl * ((int8_t)((q[l] >> shift) & 3) - ((hm[l] & m) ? 0 : 4))); - } -} - -static inline __device__ void get_scale_min_k4(int j, const uint8_t * q, uint8_t & d, uint8_t & m) { - if (j < 4) { - d = q[j] & 63; m = q[j + 4] & 63; - } else { - d = (q[j+4] & 0xF) | ((q[j-4] >> 6) << 4); - m = (q[j+4] >> 4) | ((q[j-0] >> 6) << 4); - } + dequantize_q3_K(vx, i, yy + i*QK_K, threadIdx.x); } template static __global__ void dequantize_block_q4_K(const void * __restrict__ vx, dst_t * __restrict__ yy) { - const block_q4_K * x = (const block_q4_K *) vx; - const int64_t i = blockIdx.x; - // assume 32 threads - const int64_t tid = threadIdx.x; - const int64_t il = tid/8; - const int64_t ir = tid%8; - const int64_t is = 2*il; - const int64_t n = 4; - - dst_t * y = yy + i*QK_K + 64*il + n*ir; - - const float dall = __low2half(x[i].dm); - const float dmin = __high2half(x[i].dm); - - const uint8_t * q = x[i].qs + 32*il + n*ir; - - uint8_t sc, m; - get_scale_min_k4(is + 0, x[i].scales, sc, m); - const float d1 = dall * sc; const float m1 = dmin * m; - get_scale_min_k4(is + 1, x[i].scales, sc, m); - const float d2 = dall * sc; const float m2 = dmin * m; - for (int l = 0; l < n; ++l) { - y[l + 0] = ggml_cuda_cast(d1 * (q[l] & 0xF) - m1); - y[l +32] = ggml_cuda_cast(d2 * (q[l] >> 4) - m2); - } + dequantize_q4_K(vx, i, yy + i*QK_K, threadIdx.x); } template static __global__ void dequantize_block_q5_K(const void * __restrict__ vx, dst_t * __restrict__ yy) { - const block_q5_K * x = (const block_q5_K *) vx; - const int64_t i = blockIdx.x; - // assume 64 threads - this is very slightly better than the one below - const int64_t tid = threadIdx.x; - const int64_t il = tid/16; // il is in 0...3 - const int64_t ir = tid%16; // ir is in 0...15 - const int64_t is = 2*il; // is is in 0...6 - - dst_t * y = yy + i*QK_K + 64*il + 2*ir; - - const float dall = __low2half(x[i].dm); - const float dmin = __high2half(x[i].dm); - - const uint8_t * ql = x[i].qs + 32*il + 2*ir; - const uint8_t * qh = x[i].qh + 2*ir; - - uint8_t sc, m; - get_scale_min_k4(is + 0, x[i].scales, sc, m); - const float d1 = dall * sc; const float m1 = dmin * m; - get_scale_min_k4(is + 1, x[i].scales, sc, m); - const float d2 = dall * sc; const float m2 = dmin * m; - - uint8_t hm = 1 << (2*il); - y[ 0] = ggml_cuda_cast(d1 * ((ql[ 0] & 0xF) + (qh[ 0] & hm ? 16 : 0)) - m1); - y[ 1] = ggml_cuda_cast(d1 * ((ql[ 1] & 0xF) + (qh[ 1] & hm ? 16 : 0)) - m1); - hm <<= 1; - y[32] = ggml_cuda_cast(d2 * ((ql[ 0] >> 4) + (qh[ 0] & hm ? 16 : 0)) - m2); - y[33] = ggml_cuda_cast(d2 * ((ql[ 1] >> 4) + (qh[ 1] & hm ? 16 : 0)) - m2); + dequantize_q5_K(vx, i, yy + i*QK_K, threadIdx.x); } template static __global__ void dequantize_block_q6_K(const void * __restrict__ vx, dst_t * __restrict__ yy) { - const block_q6_K * x = (const block_q6_K *) vx; - const int64_t i = blockIdx.x; - // assume 64 threads - this is very slightly better than the one below - const int64_t tid = threadIdx.x; - const int64_t ip = tid/32; // ip is 0 or 1 - const int64_t il = tid - 32*ip; // 0...32 - const int64_t is = 8*ip + il/16; - - dst_t * y = yy + i*QK_K + 128*ip + il; - - const float d = x[i].d; - - const uint8_t * ql = x[i].ql + 64*ip + il; - const uint8_t qh = x[i].qh[32*ip + il]; - const int8_t * sc = x[i].scales + is; - - y[ 0] = ggml_cuda_cast(d * sc[0] * ((int8_t)((ql[ 0] & 0xF) | (((qh >> 0) & 3) << 4)) - 32)); - y[32] = ggml_cuda_cast(d * sc[2] * ((int8_t)((ql[32] & 0xF) | (((qh >> 2) & 3) << 4)) - 32)); - y[64] = ggml_cuda_cast(d * sc[4] * ((int8_t)((ql[ 0] >> 4) | (((qh >> 4) & 3) << 4)) - 32)); - y[96] = ggml_cuda_cast(d * sc[6] * ((int8_t)((ql[32] >> 4) | (((qh >> 6) & 3) << 4)) - 32)); + dequantize_q6_K(vx, i, yy + i*QK_K, threadIdx.x); } template static __global__ void dequantize_block_iq2_xxs(const void * __restrict__ vx, dst_t * __restrict__ yy) { + const int64_t i = blockIdx.x; - const int64_t i = blockIdx.x; - const block_iq2_xxs * x = (const block_iq2_xxs *) vx; - - const int64_t tid = threadIdx.x; - const int64_t il = tid/8; // 0...3 - const int64_t ib = tid%8; // 0...7 - dst_t * y = yy + i*QK_K + 32*ib + 8*il; - const uint16_t * q2 = x[i].qs + 4*ib; - const uint8_t * aux8 = (const uint8_t *)q2; - const uint8_t * grid = (const uint8_t *)(iq2xxs_grid + aux8[il]); - const uint32_t aux32 = q2[2] | (q2[3] << 16); - const float d = (float)x[i].d * (0.5f + (aux32 >> 28)) * 0.25f; - const uint8_t signs = ksigns_iq2xs[(aux32 >> 7*il) & 127]; - for (int j = 0; j < 8; ++j) { - y[j] = ggml_cuda_cast(d * grid[j] * (signs & kmask_iq2xs[j] ? -1.f : 1.f)); - } + dequantize_iq2_xxs(vx, i, yy + i*QK_K, threadIdx.x); } template static __global__ void dequantize_block_iq2_xs(const void * __restrict__ vx, dst_t * __restrict__ yy) { + const int64_t i = blockIdx.x; - const int64_t i = blockIdx.x; - const block_iq2_xs * x = (const block_iq2_xs *) vx; - - const int64_t tid = threadIdx.x; - const int64_t il = tid/8; // 0...3 - const int64_t ib = tid%8; // 0...7 - dst_t * y = yy + i*QK_K + 32*ib + 8*il; - const uint16_t * q2 = x[i].qs + 4*ib; - const uint8_t * grid = (const uint8_t *)(iq2xs_grid + (q2[il] & 511)); - const float d = (float)x[i].d * (0.5f + ((x[i].scales[ib] >> 4*(il/2)) & 0xf)) * 0.25f; - const uint8_t signs = ksigns_iq2xs[q2[il] >> 9]; - for (int j = 0; j < 8; ++j) { - y[j] = ggml_cuda_cast(d * grid[j] * (signs & kmask_iq2xs[j] ? -1.f : 1.f)); - } + dequantize_iq2_xs(vx, i, yy + i*QK_K, threadIdx.x); } template static __global__ void dequantize_block_iq2_s(const void * __restrict__ vx, dst_t * __restrict__ yy) { + const int64_t i = blockIdx.x; - const int64_t i = blockIdx.x; - const block_iq2_s * x = (const block_iq2_s *) vx; - - const int64_t tid = threadIdx.x; - const int64_t il = tid/8; // 0...3 - const int64_t ib = tid%8; // 0...7 - dst_t * y = yy + i*QK_K + 32*ib + 8*il; - const uint8_t * grid = (const uint8_t *)(iq2s_grid + (x[i].qs[4*ib+il] | ((x[i].qh[ib] << (8-2*il)) & 0x300))); - const float d = (float)x[i].d * (0.5f + ((x[i].scales[ib] >> 4*(il/2)) & 0xf)) * 0.25f; - const uint8_t signs = x[i].qs[QK_K/8+4*ib+il]; - for (int j = 0; j < 8; ++j) { - y[j] = ggml_cuda_cast(d * grid[j] * (signs & kmask_iq2xs[j] ? -1.f : 1.f)); - } + dequantize_iq2_s(vx, i, yy + i*QK_K, threadIdx.x); } template static __global__ void dequantize_block_iq3_xxs(const void * __restrict__ vx, dst_t * __restrict__ yy) { + const int64_t i = blockIdx.x; - const int64_t i = blockIdx.x; - const block_iq3_xxs * x = (const block_iq3_xxs *) vx; - - const int64_t tid = threadIdx.x; - const int64_t il = tid/8; // 0...3 - const int64_t ib = tid%8; // 0...7 - dst_t * y = yy + i*QK_K + 32*ib + 8*il; - const uint8_t * q3 = x[i].qs + 8*ib; - const uint16_t * gas = (const uint16_t *)(x[i].qs + QK_K/4) + 2*ib; - const uint8_t * grid1 = (const uint8_t *)(iq3xxs_grid + q3[2*il+0]); - const uint8_t * grid2 = (const uint8_t *)(iq3xxs_grid + q3[2*il+1]); - const uint32_t aux32 = gas[0] | (gas[1] << 16); - const float d = (float)x[i].d * (0.5f + (aux32 >> 28)) * 0.5f; - const uint8_t signs = ksigns_iq2xs[(aux32 >> 7*il) & 127]; - for (int j = 0; j < 4; ++j) { - y[j+0] = ggml_cuda_cast(d * grid1[j] * (signs & kmask_iq2xs[j+0] ? -1.f : 1.f)); - y[j+4] = ggml_cuda_cast(d * grid2[j] * (signs & kmask_iq2xs[j+4] ? -1.f : 1.f)); - } + dequantize_iq3_xxs(vx, i, yy + i*QK_K, threadIdx.x); } template static __global__ void dequantize_block_iq3_s(const void * __restrict__ vx, dst_t * __restrict__ yy) { + const int64_t i = blockIdx.x; - const int64_t i = blockIdx.x; - const block_iq3_s * x = (const block_iq3_s *) vx; - - const int64_t tid = threadIdx.x; - const int64_t il = tid/8; // 0...3 - const int64_t ib = tid%8; // 0...7 - dst_t * y = yy + i*QK_K + 32*ib + 8*il; - const uint8_t * qs = x[i].qs + 8*ib; - const uint8_t * grid1 = (const uint8_t *)(iq3s_grid + (qs[2*il+0] | ((x[i].qh[ib] << (8-2*il)) & 256))); - const uint8_t * grid2 = (const uint8_t *)(iq3s_grid + (qs[2*il+1] | ((x[i].qh[ib] << (7-2*il)) & 256))); - const float d = (float)x[i].d * (1 + 2*((x[i].scales[ib/2] >> 4*(ib%2)) & 0xf)); - const uint8_t signs = x[i].signs[4*ib + il]; - for (int j = 0; j < 4; ++j) { - y[j+0] = ggml_cuda_cast(d * grid1[j] * (signs & kmask_iq2xs[j+0] ? -1.f : 1.f)); - y[j+4] = ggml_cuda_cast(d * grid2[j] * (signs & kmask_iq2xs[j+4] ? -1.f : 1.f)); - } + dequantize_iq3_s(vx, i, yy + i*QK_K, threadIdx.x); } template static __global__ void dequantize_block_iq1_s(const void * __restrict__ vx, dst_t * __restrict__ yy) { + const int64_t i = blockIdx.x; - const int64_t i = blockIdx.x; - const block_iq1_s * x = (const block_iq1_s *) vx; - - const int64_t tid = threadIdx.x; - const int64_t il = tid/8; // 0...3 - const int64_t ib = tid%8; // 0...7 - dst_t * y = yy + i*QK_K + 32*ib + 8*il; - const float delta = x[i].qh[ib] & 0x8000 ? -1 - IQ1S_DELTA : -1 + IQ1S_DELTA; - const float d = (float)x[i].d * (2*((x[i].qh[ib] >> 12) & 7) + 1); - uint32_t grid32[2]; const int8_t * q = (const int8_t *)grid32; - grid32[0] = iq1s_grid_gpu[x[i].qs[4*ib+il] | (((x[i].qh[ib] >> 3*il) & 7) << 8)]; - grid32[1] = (grid32[0] >> 4) & 0x0f0f0f0f; - grid32[0] &= 0x0f0f0f0f; - for (int j = 0; j < 8; ++j) { - y[j] = ggml_cuda_cast(d * (q[j] + delta)); - } + dequantize_iq1_s(vx, i, yy + i*QK_K, threadIdx.x); } template static __global__ void dequantize_block_iq1_m(const void * __restrict__ vx, dst_t * __restrict__ yy) { + const int64_t i = blockIdx.x; - const int64_t i = blockIdx.x; - const block_iq1_m * x = (const block_iq1_m *) vx; - - const int64_t tid = threadIdx.x; - const int64_t il = tid/8; // 0...3 - const int64_t ib = tid%8; // 0...7 - dst_t * y = yy + i*QK_K + 32*ib + 8*il; - const uint16_t * sc = (const uint16_t *)x[i].scales; - iq1m_scale_t scale; - scale.u16 = (sc[0] >> 12) | ((sc[1] >> 8) & 0x00f0) | ((sc[2] >> 4) & 0x0f00) | (sc[3] & 0xf000); - const int64_t ib16 = 2*ib + il/2; // sc[ib16/4] >> 3*(ib16%4) -> sc[ib/2] >> 3*((2*ib+il/2)%4); - const float d = (float)scale.f16 * (2*((sc[ib16/4] >> 3*(ib16%4)) & 0x7) + 1); - const float delta = x[i].qh[2*ib+il/2] & (0x08 << 4*(il%2)) ? -1 - IQ1M_DELTA : -1 + IQ1M_DELTA; - uint32_t grid32[2]; const int8_t * q = (const int8_t *)grid32; - grid32[0] = iq1s_grid_gpu[x[i].qs[4*ib+il] | (((x[i].qh[2*ib+il/2] >> 4*(il%2)) & 7) << 8)]; - grid32[1] = (grid32[0] >> 4) & 0x0f0f0f0f; - grid32[0] &= 0x0f0f0f0f; - for (int j = 0; j < 8; ++j) { - y[j] = ggml_cuda_cast(d * (q[j] + delta)); - } + dequantize_iq1_m(vx, i, yy + i*QK_K, threadIdx.x); } template static __global__ void dequantize_block_iq4_nl(const void * __restrict__ vx, dst_t * __restrict__ yy) { + const int64_t i = blockIdx.x; - const int64_t i = blockIdx.x; - const block_iq4_nl * x = (const block_iq4_nl *) vx + i*(QK_K/QK4_NL); - - const int64_t tid = threadIdx.x; - const int64_t il = tid/8; // 0...3 - const int64_t ib = tid%8; // 0...7 - dst_t * y = yy + i*QK_K + 32*ib + 4*il; - const uint8_t * q4 = x[ib].qs + 4*il; - const float d = (float)x[ib].d; - for (int j = 0; j < 4; ++j) { - y[j+ 0] = ggml_cuda_cast(d * kvalues_iq4nl[q4[j] & 0xf]); - y[j+16] = ggml_cuda_cast(d * kvalues_iq4nl[q4[j] >> 4]); - } + dequantize_iq4_nl(vx, i, yy + i*QK_K, threadIdx.x); } template static __global__ void dequantize_block_iq4_xs(const void * __restrict__ vx, dst_t * __restrict__ yy) { - const int64_t i = blockIdx.x; - const block_iq4_xs * x = (const block_iq4_xs *)vx; + const int64_t i = blockIdx.x; - const int64_t tid = threadIdx.x; - const int64_t il = tid/8; // 0...3 - const int64_t ib = tid%8; // 0...7 - dst_t * y = yy + i*QK_K + 32*ib + 4*il; - const uint8_t * q4 = x[i].qs + 16*ib + 4*il; - const float d = (float)x[i].d * ((((x[i].scales_l[ib/2] >> 4*(ib%2)) & 0xf) | (((x[i].scales_h >> 2*ib) & 3) << 4)) - 32); - for (int j = 0; j < 4; ++j) { - y[j+ 0] = ggml_cuda_cast(d * kvalues_iq4nl[q4[j] & 0xf]); - y[j+16] = ggml_cuda_cast(d * kvalues_iq4nl[q4[j] >> 4]); - } + dequantize_iq4_xs(vx, i, yy + i*QK_K, threadIdx.x); } template static __global__ void dequantize_block_mxfp4(const void * __restrict__ vx, dst_t * __restrict__ yy) { + const int64_t i = blockIdx.x; - const int64_t i = blockIdx.x; - const block_mxfp4 * x = (const block_mxfp4 *) vx + i*(QK_K/QK_MXFP4); - - const int64_t tid = threadIdx.x; - const int64_t il = tid/8; // 0...3 - const int64_t ib = tid%8; // 0...7 - dst_t * y = yy + i*QK_K + 32*ib + 4*il; - const uint8_t * q4 = x[ib].qs + 4*il; - const float d = ggml_cuda_e8m0_to_fp32(x[ib].e); - for (int j = 0; j < 4; ++j) { - y[j+ 0] = ggml_cuda_cast(d * kvalues_mxfp4[q4[j] & 0xf]*0.5f); - y[j+16] = ggml_cuda_cast(d * kvalues_mxfp4[q4[j] >> 4]*0.5f); - } + dequantize_mxfp4(vx, i, yy + i*QK_K, threadIdx.x); } template diff --git a/ggml/src/ggml-cuda/dequantize.cuh b/ggml/src/ggml-cuda/dequantize.cuh index 9ae1342fc0..8ab5ad8e74 100644 --- a/ggml/src/ggml-cuda/dequantize.cuh +++ b/ggml/src/ggml-cuda/dequantize.cuh @@ -1,4 +1,5 @@ #include "common.cuh" +#include "convert.cuh" static __device__ __forceinline__ void dequantize_q1_0(const void * vx, const int64_t ib, const int iqs, float2 & v){ const block_q1_0 * x = (const block_q1_0 *) vx; @@ -97,3 +98,335 @@ static __device__ __forceinline__ void dequantize_q8_0(const void * vx, const in v.x *= d; v.y *= d; } + +//================================== k-quants + +// Each call dequantizes one super-block of QK_K values into y using the +// thread layout of the caller: 32 threads for q4_K, 64 threads otherwise. + +template +static __device__ __forceinline__ void dequantize_q2_K(const void * vx, const int64_t ib, dst_t * yy, const int tid) { + const block_q2_K * x = (const block_q2_K *) vx; + + const int64_t n = tid/32; + const int64_t l = tid - 32*n; + const int64_t is = 8*n + l/16; + + const uint8_t q = x[ib].qs[32*n + l]; + dst_t * y = yy + 128*n; + + float dall = __low2half(x[ib].dm); + float dmin = __high2half(x[ib].dm); + y[l+ 0] = ggml_cuda_cast(dall * (x[ib].scales[is+0] & 0xF) * ((q >> 0) & 3) - dmin * (x[ib].scales[is+0] >> 4)); + y[l+32] = ggml_cuda_cast(dall * (x[ib].scales[is+2] & 0xF) * ((q >> 2) & 3) - dmin * (x[ib].scales[is+2] >> 4)); + y[l+64] = ggml_cuda_cast(dall * (x[ib].scales[is+4] & 0xF) * ((q >> 4) & 3) - dmin * (x[ib].scales[is+4] >> 4)); + y[l+96] = ggml_cuda_cast(dall * (x[ib].scales[is+6] & 0xF) * ((q >> 6) & 3) - dmin * (x[ib].scales[is+6] >> 4)); +} + +template +static __device__ __forceinline__ void dequantize_q3_K(const void * vx, const int64_t ib, dst_t * yy, const int tid) { + const block_q3_K * x = (const block_q3_K *) vx; + + const int64_t r = tid/4; + const int64_t t = r/2; + const int64_t is0 = r%2; + const int64_t l0 = 16*is0 + 4*(tid%4); + const int64_t n = t / 4; + const int64_t j = t - 4*n; + + uint8_t m = 1 << (4*n + j); + int64_t is = 8*n + 2*j + is0; + int shift = 2*j; + + int8_t us = is < 4 ? (x[ib].scales[is-0] & 0xF) | (((x[ib].scales[is+8] >> 0) & 3) << 4) : + is < 8 ? (x[ib].scales[is-0] & 0xF) | (((x[ib].scales[is+4] >> 2) & 3) << 4) : + is < 12 ? (x[ib].scales[is-8] >> 4) | (((x[ib].scales[is+0] >> 4) & 3) << 4) : + (x[ib].scales[is-8] >> 4) | (((x[ib].scales[is-4] >> 6) & 3) << 4); + float d_all = x[ib].d; + float dl = d_all * (us - 32); + + dst_t * y = yy + 128*n + 32*j; + const uint8_t * q = x[ib].qs + 32*n; + const uint8_t * hm = x[ib].hmask; + + for (int l = l0; l < l0+4; ++l) { + y[l] = ggml_cuda_cast(dl * ((int8_t)((q[l] >> shift) & 3) - ((hm[l] & m) ? 0 : 4))); + } +} + +static inline __device__ void get_scale_min_k4(int j, const uint8_t * q, uint8_t & d, uint8_t & m) { + if (j < 4) { + d = q[j] & 63; m = q[j + 4] & 63; + } else { + d = (q[j+4] & 0xF) | ((q[j-4] >> 6) << 4); + m = (q[j+4] >> 4) | ((q[j-0] >> 6) << 4); + } +} + +template +static __device__ __forceinline__ void dequantize_q4_K(const void * vx, const int64_t ib, dst_t * yy, const int tid) { + const block_q4_K * x = (const block_q4_K *) vx; + + // assume 32 threads + const int64_t il = tid/8; + const int64_t ir = tid%8; + const int64_t is = 2*il; + const int64_t n = 4; + + dst_t * y = yy + 64*il + n*ir; + + const float dall = __low2half(x[ib].dm); + const float dmin = __high2half(x[ib].dm); + + const uint8_t * q = x[ib].qs + 32*il + n*ir; + + uint8_t sc, m; + get_scale_min_k4(is + 0, x[ib].scales, sc, m); + const float d1 = dall * sc; const float m1 = dmin * m; + get_scale_min_k4(is + 1, x[ib].scales, sc, m); + const float d2 = dall * sc; const float m2 = dmin * m; + for (int l = 0; l < n; ++l) { + y[l + 0] = ggml_cuda_cast(d1 * (q[l] & 0xF) - m1); + y[l +32] = ggml_cuda_cast(d2 * (q[l] >> 4) - m2); + } +} + +template +static __device__ __forceinline__ void dequantize_q5_K(const void * vx, const int64_t ib, dst_t * yy, const int tid) { + const block_q5_K * x = (const block_q5_K *) vx; + + // assume 64 threads - this is very slightly better than the one below + const int64_t il = tid/16; // il is in 0...3 + const int64_t ir = tid%16; // ir is in 0...15 + const int64_t is = 2*il; // is is in 0...6 + + dst_t * y = yy + 64*il + 2*ir; + + const float dall = __low2half(x[ib].dm); + const float dmin = __high2half(x[ib].dm); + + const uint8_t * ql = x[ib].qs + 32*il + 2*ir; + const uint8_t * qh = x[ib].qh + 2*ir; + + uint8_t sc, m; + get_scale_min_k4(is + 0, x[ib].scales, sc, m); + const float d1 = dall * sc; const float m1 = dmin * m; + get_scale_min_k4(is + 1, x[ib].scales, sc, m); + const float d2 = dall * sc; const float m2 = dmin * m; + + uint8_t hm = 1 << (2*il); + y[ 0] = ggml_cuda_cast(d1 * ((ql[ 0] & 0xF) + (qh[ 0] & hm ? 16 : 0)) - m1); + y[ 1] = ggml_cuda_cast(d1 * ((ql[ 1] & 0xF) + (qh[ 1] & hm ? 16 : 0)) - m1); + hm <<= 1; + y[32] = ggml_cuda_cast(d2 * ((ql[ 0] >> 4) + (qh[ 0] & hm ? 16 : 0)) - m2); + y[33] = ggml_cuda_cast(d2 * ((ql[ 1] >> 4) + (qh[ 1] & hm ? 16 : 0)) - m2); +} + +template +static __device__ __forceinline__ void dequantize_q6_K(const void * vx, const int64_t ib, dst_t * yy, const int tid) { + const block_q6_K * x = (const block_q6_K *) vx; + + // assume 64 threads - this is very slightly better than the one below + const int64_t ip = tid/32; // ip is 0 or 1 + const int64_t il = tid - 32*ip; // 0...32 + const int64_t is = 8*ip + il/16; + + dst_t * y = yy + 128*ip + il; + + const float d = x[ib].d; + + const uint8_t * ql = x[ib].ql + 64*ip + il; + const uint8_t qh = x[ib].qh[32*ip + il]; + const int8_t * sc = x[ib].scales + is; + + y[ 0] = ggml_cuda_cast(d * sc[0] * ((int8_t)((ql[ 0] & 0xF) | (((qh >> 0) & 3) << 4)) - 32)); + y[32] = ggml_cuda_cast(d * sc[2] * ((int8_t)((ql[32] & 0xF) | (((qh >> 2) & 3) << 4)) - 32)); + y[64] = ggml_cuda_cast(d * sc[4] * ((int8_t)((ql[ 0] >> 4) | (((qh >> 4) & 3) << 4)) - 32)); + y[96] = ggml_cuda_cast(d * sc[6] * ((int8_t)((ql[32] >> 4) | (((qh >> 6) & 3) << 4)) - 32)); +} + +//================================== i-quants + +// Each call dequantizes one super-block of QK_K values into y with 32 +// threads; iq4_nl packs QK_K/QK4_NL sub-blocks per super-block. + +template +static __device__ __forceinline__ void dequantize_iq2_xxs(const void * vx, const int64_t ibs, dst_t * yy, const int tid) { + + const block_iq2_xxs * x = (const block_iq2_xxs *) vx; + + const int64_t il = tid/8; // 0...3 + const int64_t ib = tid%8; // 0...7 + dst_t * y = yy + 32*ib + 8*il; + const uint16_t * q2 = x[ibs].qs + 4*ib; + const uint8_t * aux8 = (const uint8_t *)q2; + const uint8_t * grid = (const uint8_t *)(iq2xxs_grid + aux8[il]); + const uint32_t aux32 = q2[2] | (q2[3] << 16); + const float d = (float)x[ibs].d * (0.5f + (aux32 >> 28)) * 0.25f; + const uint8_t signs = ksigns_iq2xs[(aux32 >> 7*il) & 127]; + for (int j = 0; j < 8; ++j) { + y[j] = ggml_cuda_cast(d * grid[j] * (signs & kmask_iq2xs[j] ? -1.f : 1.f)); + } +} + +template +static __device__ __forceinline__ void dequantize_iq2_xs(const void * vx, const int64_t ibs, dst_t * yy, const int tid) { + + const block_iq2_xs * x = (const block_iq2_xs *) vx; + + const int64_t il = tid/8; // 0...3 + const int64_t ib = tid%8; // 0...7 + dst_t * y = yy + 32*ib + 8*il; + const uint16_t * q2 = x[ibs].qs + 4*ib; + const uint8_t * grid = (const uint8_t *)(iq2xs_grid + (q2[il] & 511)); + const float d = (float)x[ibs].d * (0.5f + ((x[ibs].scales[ib] >> 4*(il/2)) & 0xf)) * 0.25f; + const uint8_t signs = ksigns_iq2xs[q2[il] >> 9]; + for (int j = 0; j < 8; ++j) { + y[j] = ggml_cuda_cast(d * grid[j] * (signs & kmask_iq2xs[j] ? -1.f : 1.f)); + } +} + +template +static __device__ __forceinline__ void dequantize_iq2_s(const void * vx, const int64_t ibs, dst_t * yy, const int tid) { + + const block_iq2_s * x = (const block_iq2_s *) vx; + + const int64_t il = tid/8; // 0...3 + const int64_t ib = tid%8; // 0...7 + dst_t * y = yy + 32*ib + 8*il; + const uint8_t * grid = (const uint8_t *)(iq2s_grid + (x[ibs].qs[4*ib+il] | ((x[ibs].qh[ib] << (8-2*il)) & 0x300))); + const float d = (float)x[ibs].d * (0.5f + ((x[ibs].scales[ib] >> 4*(il/2)) & 0xf)) * 0.25f; + const uint8_t signs = x[ibs].qs[QK_K/8+4*ib+il]; + for (int j = 0; j < 8; ++j) { + y[j] = ggml_cuda_cast(d * grid[j] * (signs & kmask_iq2xs[j] ? -1.f : 1.f)); + } +} + +template +static __device__ __forceinline__ void dequantize_iq3_xxs(const void * vx, const int64_t ibs, dst_t * yy, const int tid) { + + const block_iq3_xxs * x = (const block_iq3_xxs *) vx; + + const int64_t il = tid/8; // 0...3 + const int64_t ib = tid%8; // 0...7 + dst_t * y = yy + 32*ib + 8*il; + const uint8_t * q3 = x[ibs].qs + 8*ib; + const uint16_t * gas = (const uint16_t *)(x[ibs].qs + QK_K/4) + 2*ib; + const uint8_t * grid1 = (const uint8_t *)(iq3xxs_grid + q3[2*il+0]); + const uint8_t * grid2 = (const uint8_t *)(iq3xxs_grid + q3[2*il+1]); + const uint32_t aux32 = gas[0] | (gas[1] << 16); + const float d = (float)x[ibs].d * (0.5f + (aux32 >> 28)) * 0.5f; + const uint8_t signs = ksigns_iq2xs[(aux32 >> 7*il) & 127]; + for (int j = 0; j < 4; ++j) { + y[j+0] = ggml_cuda_cast(d * grid1[j] * (signs & kmask_iq2xs[j+0] ? -1.f : 1.f)); + y[j+4] = ggml_cuda_cast(d * grid2[j] * (signs & kmask_iq2xs[j+4] ? -1.f : 1.f)); + } +} + +template +static __device__ __forceinline__ void dequantize_iq3_s(const void * vx, const int64_t ibs, dst_t * yy, const int tid) { + + const block_iq3_s * x = (const block_iq3_s *) vx; + + const int64_t il = tid/8; // 0...3 + const int64_t ib = tid%8; // 0...7 + dst_t * y = yy + 32*ib + 8*il; + const uint8_t * qs = x[ibs].qs + 8*ib; + const uint8_t * grid1 = (const uint8_t *)(iq3s_grid + (qs[2*il+0] | ((x[ibs].qh[ib] << (8-2*il)) & 256))); + const uint8_t * grid2 = (const uint8_t *)(iq3s_grid + (qs[2*il+1] | ((x[ibs].qh[ib] << (7-2*il)) & 256))); + const float d = (float)x[ibs].d * (1 + 2*((x[ibs].scales[ib/2] >> 4*(ib%2)) & 0xf)); + const uint8_t signs = x[ibs].signs[4*ib + il]; + for (int j = 0; j < 4; ++j) { + y[j+0] = ggml_cuda_cast(d * grid1[j] * (signs & kmask_iq2xs[j+0] ? -1.f : 1.f)); + y[j+4] = ggml_cuda_cast(d * grid2[j] * (signs & kmask_iq2xs[j+4] ? -1.f : 1.f)); + } +} + +template +static __device__ __forceinline__ void dequantize_iq1_s(const void * vx, const int64_t ibs, dst_t * yy, const int tid) { + + const block_iq1_s * x = (const block_iq1_s *) vx; + + const int64_t il = tid/8; // 0...3 + const int64_t ib = tid%8; // 0...7 + dst_t * y = yy + 32*ib + 8*il; + const float delta = x[ibs].qh[ib] & 0x8000 ? -1 - IQ1S_DELTA : -1 + IQ1S_DELTA; + const float d = (float)x[ibs].d * (2*((x[ibs].qh[ib] >> 12) & 7) + 1); + uint32_t grid32[2]; const int8_t * q = (const int8_t *)grid32; + grid32[0] = iq1s_grid_gpu[x[ibs].qs[4*ib+il] | (((x[ibs].qh[ib] >> 3*il) & 7) << 8)]; + grid32[1] = (grid32[0] >> 4) & 0x0f0f0f0f; + grid32[0] &= 0x0f0f0f0f; + for (int j = 0; j < 8; ++j) { + y[j] = ggml_cuda_cast(d * (q[j] + delta)); + } +} + +template +static __device__ __forceinline__ void dequantize_iq1_m(const void * vx, const int64_t ibs, dst_t * yy, const int tid) { + + const block_iq1_m * x = (const block_iq1_m *) vx; + + const int64_t il = tid/8; // 0...3 + const int64_t ib = tid%8; // 0...7 + dst_t * y = yy + 32*ib + 8*il; + const uint16_t * sc = (const uint16_t *)x[ibs].scales; + iq1m_scale_t scale; + scale.u16 = (sc[0] >> 12) | ((sc[1] >> 8) & 0x00f0) | ((sc[2] >> 4) & 0x0f00) | (sc[3] & 0xf000); + const int64_t ib16 = 2*ib + il/2; // sc[ib16/4] >> 3*(ib16%4) -> sc[ib/2] >> 3*((2*ib+il/2)%4); + const float d = (float)scale.f16 * (2*((sc[ib16/4] >> 3*(ib16%4)) & 0x7) + 1); + const float delta = x[ibs].qh[2*ib+il/2] & (0x08 << 4*(il%2)) ? -1 - IQ1M_DELTA : -1 + IQ1M_DELTA; + uint32_t grid32[2]; const int8_t * q = (const int8_t *)grid32; + grid32[0] = iq1s_grid_gpu[x[ibs].qs[4*ib+il] | (((x[ibs].qh[2*ib+il/2] >> 4*(il%2)) & 7) << 8)]; + grid32[1] = (grid32[0] >> 4) & 0x0f0f0f0f; + grid32[0] &= 0x0f0f0f0f; + for (int j = 0; j < 8; ++j) { + y[j] = ggml_cuda_cast(d * (q[j] + delta)); + } +} + +template +static __device__ __forceinline__ void dequantize_iq4_nl(const void * vx, const int64_t ibs, dst_t * yy, const int tid) { + + const block_iq4_nl * x = (const block_iq4_nl *) vx + ibs*(QK_K/QK4_NL); + + const int64_t il = tid/8; // 0...3 + const int64_t ib = tid%8; // 0...7 + dst_t * y = yy + 32*ib + 4*il; + const uint8_t * q4 = x[ib].qs + 4*il; + const float d = (float)x[ib].d; + for (int j = 0; j < 4; ++j) { + y[j+ 0] = ggml_cuda_cast(d * kvalues_iq4nl[q4[j] & 0xf]); + y[j+16] = ggml_cuda_cast(d * kvalues_iq4nl[q4[j] >> 4]); + } +} + +template +static __device__ __forceinline__ void dequantize_iq4_xs(const void * vx, const int64_t ibs, dst_t * yy, const int tid) { + const block_iq4_xs * x = (const block_iq4_xs *)vx; + + const int64_t il = tid/8; // 0...3 + const int64_t ib = tid%8; // 0...7 + dst_t * y = yy + 32*ib + 4*il; + const uint8_t * q4 = x[ibs].qs + 16*ib + 4*il; + const float d = (float)x[ibs].d * ((((x[ibs].scales_l[ib/2] >> 4*(ib%2)) & 0xf) | (((x[ibs].scales_h >> 2*ib) & 3) << 4)) - 32); + for (int j = 0; j < 4; ++j) { + y[j+ 0] = ggml_cuda_cast(d * kvalues_iq4nl[q4[j] & 0xf]); + y[j+16] = ggml_cuda_cast(d * kvalues_iq4nl[q4[j] >> 4]); + } +} + +template +static __device__ __forceinline__ void dequantize_mxfp4(const void * vx, const int64_t ibs, dst_t * yy, const int tid) { + + const block_mxfp4 * x = (const block_mxfp4 *) vx + ibs*(QK_K/QK_MXFP4); + + const int64_t il = tid/8; // 0...3 + const int64_t ib = tid%8; // 0...7 + dst_t * y = yy + 32*ib + 4*il; + const uint8_t * q4 = x[ib].qs + 4*il; + const float d = ggml_cuda_e8m0_to_fp32(x[ib].e); + for (int j = 0; j < 4; ++j) { + y[j+ 0] = ggml_cuda_cast(d * kvalues_mxfp4[q4[j] & 0xf]*0.5f); + y[j+16] = ggml_cuda_cast(d * kvalues_mxfp4[q4[j] >> 4]*0.5f); + } +} diff --git a/ggml/src/ggml-cuda/getrows.cu b/ggml/src/ggml-cuda/getrows.cu index dcf9ca7a70..a9ec4f697c 100644 --- a/ggml/src/ggml-cuda/getrows.cu +++ b/ggml/src/ggml-cuda/getrows.cu @@ -40,6 +40,35 @@ static __global__ void k_get_rows( } } +template dequantize_kq> +static __global__ void k_get_rows_kq( + const void * __restrict__ src0, const int32_t * __restrict__ src1, dst_t * __restrict__ dst, + const int64_t ne00, /*const int64_t ne01, const int64_t ne02, const int64_t ne03,*/ + /*const int64_t ne10,*/ const int64_t ne11, const uint3 ne12_fdv, /*const int64_t ne13,*/ + /*const size_t s0,*/ const size_t s1, const size_t s2, const size_t s3, + /*const size_t nb00,*/ const size_t nb01, const size_t nb02, const size_t nb03, + const size_t s10, const size_t s11, const size_t s12/*, const size_t s13*/) { + + ggml_cuda_pdl_sync(); + const int64_t nsb = ne00/QK_K; // super-blocks per row + for (int64_t z = blockIdx.z; z < ne11*(int64_t)ne12_fdv.z; z += gridDim.z) { + // The x and y dimensions of the grid are swapped because the maximum allowed grid size for x is higher. + const int i10 = blockIdx.x; + const uint2 dm = fast_div_modulo((uint32_t)z, ne12_fdv); + const int i11 = dm.x; + const int i12 = dm.y; + + const int i01 = src1[i10*s10 + i11*s11 + i12*s12]; + + dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3; + const void * src0_row = (const char *) src0 + i01*nb01 + i11*nb02 + i12*nb03; + + for (int64_t ib = blockIdx.y; ib < nsb; ib += gridDim.y) { + dequantize_kq(src0_row, ib, dst_row + ib*QK_K, threadIdx.x); + } + } +} + template static __global__ void k_get_rows_float( const src0_t * src0_ptr, const int32_t * src1_ptr, dst_t * dst_ptr, @@ -164,6 +193,43 @@ static void get_rows_cuda_q( s10, s11, s12/*, s13*/); } +template dequantize_kq> +static void get_rows_cuda_kq( + const void * src0_d, const int32_t * src1_d, dst_t * dst_d, + const int64_t ne00, const size_t nb01, const size_t nb02, const size_t nb03, + const int64_t ne10, const int64_t ne11, const int64_t ne12, const size_t nb10, const size_t nb11, const size_t nb12, + const size_t nb1, const size_t nb2, const size_t nb3, + cudaStream_t stream) { + GGML_ASSERT(ne00 % QK_K == 0); + const int64_t nsb = ne00/QK_K; + + const dim3 block_dims(block_dim, 1, 1); + const dim3 block_nums(ne10, MIN(nsb, UINT16_MAX), MIN(ne11*ne12, UINT16_MAX)); + + // strides in elements + // const size_t s0 = nb0 / sizeof(dst_t); + const size_t s1 = nb1 / sizeof(dst_t); + const size_t s2 = nb2 / sizeof(dst_t); + const size_t s3 = nb3 / sizeof(dst_t); + + const size_t s10 = nb10 / sizeof(int32_t); + const size_t s11 = nb11 / sizeof(int32_t); + const size_t s12 = nb12 / sizeof(int32_t); + // const size_t s13 = nb13 / sizeof(int32_t); + + GGML_ASSERT(ne12 > 0); + GGML_ASSERT(ne11 <= std::numeric_limits::max() / ne12); + const uint3 ne12_fdv = init_fastdiv_values(ne12); + + k_get_rows_kq<<>>( + src0_d, src1_d, dst_d, + ne00, /*ne01, ne02, ne03,*/ + /*ne10,*/ ne11, ne12_fdv, /*ne13,*/ + /* s0,*/ s1, s2, s3, + /* nb00,*/ nb01, nb02, nb03, + s10, s11, s12/*, s13*/); +} + template static void get_rows_cuda_float( const src0_t * src0_d, const int32_t * src1_d, dst_t * dst_d, @@ -274,8 +340,67 @@ static void ggml_cuda_get_rows_switch_src0_type( get_rows_cuda_q(src0_d, src1_d, dst_d, ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); break; + case GGML_TYPE_Q2_K: + get_rows_cuda_kq<64, dst_t, dequantize_q2_K>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_Q3_K: + get_rows_cuda_kq<64, dst_t, dequantize_q3_K>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_Q4_K: + get_rows_cuda_kq<32, dst_t, dequantize_q4_K>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_Q5_K: + get_rows_cuda_kq<64, dst_t, dequantize_q5_K>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_Q6_K: + get_rows_cuda_kq<64, dst_t, dequantize_q6_K>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_IQ2_XXS: + get_rows_cuda_kq<32, dst_t, dequantize_iq2_xxs>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_IQ2_XS: + get_rows_cuda_kq<32, dst_t, dequantize_iq2_xs>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_IQ2_S: + get_rows_cuda_kq<32, dst_t, dequantize_iq2_s>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_IQ3_XXS: + get_rows_cuda_kq<32, dst_t, dequantize_iq3_xxs>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_IQ3_S: + get_rows_cuda_kq<32, dst_t, dequantize_iq3_s>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_IQ1_S: + get_rows_cuda_kq<32, dst_t, dequantize_iq1_s>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_IQ1_M: + get_rows_cuda_kq<32, dst_t, dequantize_iq1_m>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_IQ4_NL: + get_rows_cuda_kq<32, dst_t, dequantize_iq4_nl>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_IQ4_XS: + get_rows_cuda_kq<32, dst_t, dequantize_iq4_xs>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; + case GGML_TYPE_MXFP4: + get_rows_cuda_kq<32, dst_t, dequantize_mxfp4>(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; default: - // TODO: k-quants GGML_ABORT("%s: unsupported src0 type: %s\n", __func__, ggml_type_name(src0_type)); break; } diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 40e57ff0b3..e73a7b8906 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -4845,7 +4845,25 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q2_K: + case GGML_TYPE_Q3_K: + case GGML_TYPE_Q4_K: + case GGML_TYPE_Q5_K: + case GGML_TYPE_Q6_K: + case GGML_TYPE_IQ2_XXS: + case GGML_TYPE_IQ2_XS: + case GGML_TYPE_IQ2_S: + case GGML_TYPE_IQ3_XXS: + case GGML_TYPE_IQ3_S: + case GGML_TYPE_IQ1_S: + case GGML_TYPE_IQ1_M: + case GGML_TYPE_IQ4_XS: return true; + case GGML_TYPE_IQ4_NL: + case GGML_TYPE_MXFP4: + // 32-value sub-blocks, the row size does not guarantee + // the QK_K super-blocks the get_rows kernel iterates on + return op->src[0]->ne[0] % QK_K == 0; default: return false; }