diff --git a/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh b/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh index 899309c9..29078046 100644 --- a/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh +++ b/ggml/src/ggml-cuda/mmvq-ptq1_0.cuh @@ -442,6 +442,302 @@ static void mul_mat_vec_ptq1_0_pt_launch( vx, vy, fusion, dst, ncols_x, nrows_x, stride_row_x, stride_col_y, stride_col_dst, rows_per_cta, bpr_fd); } + +// Trits t (0, 1, 2) of four bytes of five trits. For a byte b let H_s = floor(3^s * b / 256), H_0 = 0: +// trit s is H_s - 3*H_(s-1). In a 16-bit lane 3^s * b never overflows (3^5 * 255 < 65536), so H_s of +// four bytes is one multiply per lane pair and one byte permute, and every byte of H_s - 3*H_(s-1) +// is 0, 1 or 2, so the word-wide subtraction never borrows across bytes. + +// H_s of the four bytes held as two 16-bit-lane words already multiplied by 3^s +static __device__ __forceinline__ uint32_t ptq1_0_high_bytes(const uint32_t wlo, const uint32_t whi) { + return __byte_perm(wlo, whi, 0x7531); +} + +// four trits from H_s and H_(s-1) of the same four bytes +static __device__ __forceinline__ uint32_t ptq1_0_trits(const uint32_t h, const uint32_t h_prev) { + return h - 3*h_prev; +} + +// --------------------------------------------------------------------------- +// Tensor-core PTQ1_0 mat-vec for 2 to 8 activation columns. +// +// The dp4a kernel above pays one dp4a per column for every four decoded weights, so on an +// RTX 3060 (integer-ALU bound) a 4-column verify costs ~1.8x a single column and 5+ columns +// fall to the MMQ tile path. Here a warp decodes the trits of 16 rows once and dots them +// against up to 8 columns with mma.m16n8k32, so the cost of a verify barely depends on its width. +// +// Same bits as the dp4a kernel: each 32-element sub-block sum is an exact integer either way +// (sum (t - 1)*a is a mma of the trits t accumulated onto a mma of all -1 bytes), the fp32 fold +// acc = fma(d8_k, (float) sumi_k, acc) over k = 0..3 and the scale by d are the same +// operations, and the per-row sum over K blocks keeps the dp4a kernel's order: the slot of +// residue l (mod 32) adds the blocks l, l+32, ... in that order, then the slots are combined in +// the pairing of warp_reduce_sum. Warp w of the CTA owns the blocks kb = w (mod 4) and keeps +// those residues' running sums in registers. +// +// A fragment of m16n8k32 (row major): lane (g = lane/4, q = lane%4) holds k = 4q..4q+3 and +// 16+4q..16+4q+3 of rows g and g+8. Byte group q of qs[0..15] gives elements 16t + 4q + b +// (t = 0..4). The rest is shared out across the quad: lanes 0 and 1 decode byte group q of +// qs[16..23] (elements 80 + 8t + 4q + b), lane 2 decodes qh, and the digits the other lanes +// need come over by shuffle. +// --------------------------------------------------------------------------- + +#define PTQ1_0_TC_THREADS 256 +#define PTQ1_0_TC_MAX_COLS 8 + +// c += a*b with a the trits (u8) and b the activations (s8) +static __device__ __forceinline__ void ptq1_0_tc_mma_u8(int (&c)[4], const uint32_t (&a)[4], const int b0, const int b1) { +#if defined(AMPERE_MMA_AVAILABLE) + asm volatile("mma.sync.aligned.m16n8k32.row.col.s32.u8.s8.s32 {%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9}, {%0, %1, %2, %3};" + : "+r"(c[0]), "+r"(c[1]), "+r"(c[2]), "+r"(c[3]) + : "r"(a[0]), "r"(a[1]), "r"(a[2]), "r"(a[3]), "r"(b0), "r"(b1)); +#else + GGML_UNUSED(c); GGML_UNUSED(a); GGML_UNUSED(b0); GGML_UNUSED(b1); + NO_DEVICE_CODE; +#endif +} + +// c = -(sum of the activations of each column), the same for every row +static __device__ __forceinline__ void ptq1_0_tc_mma_neg_sum(int (&c)[4], const int b0, const int b1) { +#if defined(AMPERE_MMA_AVAILABLE) + const int m1 = -1; // four s8 -1 + asm volatile("mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32 {%0, %1, %2, %3}, {%4, %4, %4, %4}, {%5, %6}, {%7, %7, %7, %7};" + : "=r"(c[0]), "=r"(c[1]), "=r"(c[2]), "=r"(c[3]) + : "r"(m1), "r"(b0), "r"(b1), "r"(0)); +#else + GGML_UNUSED(c); GGML_UNUSED(b0); GGML_UNUSED(b1); + NO_DEVICE_CODE; +#endif +} + +// the five trit words of one 4-byte group, digit s of byte b in byte b of t[s-1] +static __device__ __forceinline__ void ptq1_0_tc_trits5(const uint32_t packed, uint32_t (&t)[5]) { + uint32_t wlo = __byte_perm(packed, 0, 0x4140); + uint32_t whi = __byte_perm(packed, 0, 0x4342); + uint32_t h_prev = 0; +#pragma unroll + for (int s = 0; s < 5; ++s) { + wlo *= 3; + whi *= 3; + const uint32_t h = ptq1_0_high_bytes(wlo, whi); + t[s] = ptq1_0_trits(h, h_prev); + h_prev = h; + } +} + +// The A fragments of one 128-weight block of one row from its words xq = qs[4q..4q+3], +// xo = qs[16+4q..] (q < 2) or the qh/d word (q >= 2): frag[k][0] holds k-offsets 4q..4q+3 of +// sub-block k, frag[k][1] the offsets 16+4q..16+4q+3. Every lane of the warp must call this. +static __device__ __forceinline__ void ptq1_0_tc_row_frags(const uint32_t xq, const uint32_t xo, const int q, uint32_t (&frag)[4][2]) { + uint32_t t1[5]; + uint32_t t2[5]; + ptq1_0_tc_trits5(xq, t1); + ptq1_0_tc_trits5(q < 2 ? xo : (xo & 0xFFFF), t2); // qh[0], qh[1]: four trits each + + const uint32_t from_low1 = __shfl_xor_sync(0xFFFFFFFF, t2[1], 2); // lanes 2, 3: digit 2 of qs[16+4(q-2) ..] + const uint32_t from_low3 = __shfl_xor_sync(0xFFFFFFFF, t2[3], 2); // lanes 2, 3: digit 4 + const uint32_t qh12 = __byte_perm(t2[0], t2[1], 0x5410); // lane 2: elements 120..123 + const uint32_t qh34 = __shfl_xor_sync(0xFFFFFFFF, __byte_perm(t2[2], t2[3], 0x5410), 1); // lane 3: 124..127 + + const bool low = q < 2; + frag[0][0] = t1[0]; // elements 0 + 4q + frag[0][1] = t1[1]; // elements 16 + 4q + frag[1][0] = t1[2]; // elements 32 + 4q + frag[1][1] = t1[3]; // elements 48 + 4q + frag[2][0] = t1[4]; // elements 64 + 4q + frag[2][1] = low ? t2[0] : from_low1; // elements 80 + 4q + frag[3][0] = low ? t2[2] : from_low3; // elements 96 + 4q + frag[3][1] = low ? t2[4] : (q == 2 ? qh12 : qh34); // elements 112 + 4q +} + +template +__launch_bounds__(PTQ1_0_TC_THREADS, 2) +static __global__ void mul_mat_vec_ptq1_0_tc( + const void * vx_ptr, const void * vy_ptr, float * dst_ptr, + const int ncols_x, const int nrows_x, const int stride_row_x, const int stride_col_y, const int stride_col_dst) { + static_assert(ncols >= 1 && ncols <= PTQ1_0_TC_MAX_COLS, "1 to 8 columns"); + constexpr int nwarps = PTQ1_0_TC_THREADS / WARP_SIZE; + static_assert(WARP_SIZE % nwarps == 0, "a warp owns the K blocks of fixed residues mod 32"); + constexpr int nslots = WARP_SIZE / nwarps; // residues mod 32 per warp + + const void * GGML_CUDA_RESTRICT vx = vx_ptr; + const void * GGML_CUDA_RESTRICT vy = vy_ptr; + float * GGML_CUDA_RESTRICT dst = dst_ptr; + + __shared__ float slots[ncols*16][WARP_SIZE + 1]; + + const int bpr = ncols_x / QK_PTQ1_0; + const int row0 = 16 * blockIdx.x; + const int warp = threadIdx.x / WARP_SIZE; + const int lane = threadIdx.x % WARP_SIZE; + const int g = lane / 4; + const int q = lane % 4; + + // rows g and g+8 of the tile; tail rows are clamped and their results not written + const uint32_t * x0 = (const uint32_t *) ((const block_ptq1_0 *) vx + (int64_t) min(row0 + g, nrows_x - 1)*stride_row_x); + const uint32_t * x1 = (const uint32_t *) ((const block_ptq1_0 *) vx + (int64_t) min(row0 + g + 8, nrows_x - 1)*stride_row_x); + constexpr int xw = sizeof(block_ptq1_0) / sizeof(uint32_t); + const int xo_word = q < 2 ? 4 + q : 6; + + // B fragment column of this lane (zero beyond ncols), and the two C columns it folds + // activations in the tensor-core layout (quantize_row_q8_1_ptq1_0_tc_cuda): block kb of a column is 144 + // bytes, the 32 at 32q hold this lane's B words of the four sub-blocks, the (d, s) pairs follow at 128 + const bool b_ok = g < ncols; + const char * yb = (const char *) ((const block_q8_1 *) vy + (b_ok ? g : 0)*stride_col_y) + 32*q; + const int c0 = 2*q; + const int c1 = 2*q + 1; + const char * yd0 = (const char *) ((const block_q8_1 *) vy + (c0 < ncols ? c0 : 0)*stride_col_y) + 128; + const char * yd1 = (const char *) ((const block_q8_1 *) vy + (c1 < ncols ? c1 : 0)*stride_col_y) + 128; + + ggml_cuda_pdl_sync(); + + float slot[nslots][4]; +#pragma unroll + for (int i = 0; i < nslots; ++i) { +#pragma unroll + for (int e = 0; e < 4; ++e) { + slot[i][e] = 0.0f; + } + } + + auto block = [&](const int kb, float (&p)[4]) { + const uint32_t * b0w = x0 + kb*xw; + const uint32_t * b1w = x1 + kb*xw; + const uint32_t xq0 = b0w[q], xo0 = b0w[xo_word], xd0 = b0w[6]; + const uint32_t xq1 = b1w[q], xo1 = b1w[xo_word], xd1 = b1w[6]; + // b[2k], b[2k+1]: k-offsets 4q.. and 16+4q.. of sub-block k + const int4 bv0 = b_ok ? *(const int4 *) (yb + kb*144) : make_int4(0, 0, 0, 0); + const int4 bv1 = b_ok ? *(const int4 *) (yb + kb*144 + 16) : make_int4(0, 0, 0, 0); + const int b[8] = { bv0.x, bv0.y, bv0.z, bv0.w, bv1.x, bv1.y, bv1.z, bv1.w }; + const int4 ds0 = *(const int4 *) (yd0 + kb*144); + const int4 ds1 = *(const int4 *) (yd1 + kb*144); + + uint32_t f0[4][2]; + uint32_t f1[4][2]; + ptq1_0_tc_row_frags(xq0, xo0, q, f0); + ptq1_0_tc_row_frags(xq1, xo1, q, f1); + + // C elements: (row g, col c0), (row g, col c1), (row g+8, col c0), (row g+8, col c1) + float acc[4] = { 0.0f, 0.0f, 0.0f, 0.0f }; +#pragma unroll + for (int k = 0; k < 4; ++k) { + int c[4]; + ptq1_0_tc_mma_neg_sum(c, b[2*k], b[2*k + 1]); + const uint32_t a[4] = { f0[k][0], f1[k][0], f0[k][1], f1[k][1] }; + ptq1_0_tc_mma_u8(c, a, b[2*k], b[2*k + 1]); + + const float d8_0 = __low2float(((const half2 *) &ds0)[k]); + const float d8_1 = __low2float(((const half2 *) &ds1)[k]); + acc[0] = __fmaf_rn(d8_0, (float) c[0], acc[0]); + acc[1] = __fmaf_rn(d8_1, (float) c[1], acc[1]); + acc[2] = __fmaf_rn(d8_0, (float) c[2], acc[2]); + acc[3] = __fmaf_rn(d8_1, (float) c[3], acc[3]); + } + const float d0 = __half2float(__ushort_as_half((unsigned short) (xd0 >> 16))); + const float d1 = __half2float(__ushort_as_half((unsigned short) (xd1 >> 16))); + p[0] = __fmul_rn(d0, acc[0]); + p[1] = __fmul_rn(d0, acc[1]); + p[2] = __fmul_rn(d1, acc[2]); + p[3] = __fmul_rn(d1, acc[3]); + }; + + const int bpr_full = bpr - bpr % WARP_SIZE; + for (int kb0 = 0; kb0 < bpr_full; kb0 += WARP_SIZE) { +#pragma unroll + for (int i = 0; i < nslots; ++i) { + float p[4]; + block(kb0 + nwarps*i + warp, p); +#pragma unroll + for (int e = 0; e < 4; ++e) { + slot[i][e] += p[e]; + } + } + } +#pragma unroll + for (int i = 0; i < nslots; ++i) { + const int kb = bpr_full + nwarps*i + warp; + if (kb < bpr) { // uniform across the warp + float p[4]; + block(kb, p); +#pragma unroll + for (int e = 0; e < 4; ++e) { + slot[i][e] += p[e]; + } + } + } + +#pragma unroll + for (int i = 0; i < nslots; ++i) { + if (c0 < ncols) { + slots[c0*16 + g ][nwarps*i + warp] = slot[i][0]; + slots[c0*16 + g + 8][nwarps*i + warp] = slot[i][2]; + } + if (c1 < ncols) { + slots[c1*16 + g ][nwarps*i + warp] = slot[i][1]; + slots[c1*16 + g + 8][nwarps*i + warp] = slot[i][3]; + } + } + __syncthreads(); + + // one thread per (column, row): the 32 slots combined in warp_reduce_sum's pairing + // (distance 16, then 8, 4, 2, 1), which is the result that butterfly leaves in lane 0 + for (int w = threadIdx.x; w < ncols*16; w += PTQ1_0_TC_THREADS) { + float v[WARP_SIZE]; +#pragma unroll + for (int l = 0; l < WARP_SIZE; ++l) { + v[l] = slots[w][l]; + } +#pragma unroll + for (int off = WARP_SIZE/2; off > 0; off >>= 1) { +#pragma unroll + for (int l = 0; l < off; ++l) { + v[l] += v[l + off]; + } + } + const int j = w / 16; + const int row = row0 + w % 16; + if (row < nrows_x) { + dst[j*stride_col_dst + row] = v[0]; + } + } +} + +// GGML_CUDA_PTQ1_TC=0 keeps the dp4a kernel for 2 to 4 columns (and MMQ above), for A/B +static bool ptq1_0_tc_enabled(const int cc) { + static const bool env_on = [] { + const char * e = getenv("GGML_CUDA_PTQ1_TC"); + return e == nullptr || atoi(e) != 0; + }(); + return env_on && GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_AMPERE; +} + +template +static void mul_mat_vec_ptq1_0_tc_launch( + const void * vx, const void * vy, float * dst, + const int ncols_x, const int nrows_x, const int stride_row_x, const int stride_col_y, const int stride_col_dst, + cudaStream_t stream) { + const dim3 block_nums((nrows_x + 15) / 16, 1, 1); + const dim3 block_dims(PTQ1_0_TC_THREADS, 1, 1); + const ggml_cuda_kernel_launch_params lp = ggml_cuda_kernel_launch_params(block_nums, block_dims, 0, stream); + ggml_cuda_kernel_launch(mul_mat_vec_ptq1_0_tc, lp, vx, vy, dst, ncols_x, nrows_x, stride_row_x, stride_col_y, stride_col_dst); +} + +// plain 2D PTQ1_0 x 2..8 columns without fusion, activations quantized by quantize_row_q8_1_ptq1_0_tc_cuda +static void mul_mat_vec_ptq1_0_tc_cuda( + const void * vx, const void * vy, float * dst, + const int ncols_x, const int nrows_x, const int ncols_dst, + const int stride_row_x, const int stride_col_y, const int stride_col_dst, cudaStream_t stream) { + switch (ncols_dst) { + case 2: mul_mat_vec_ptq1_0_tc_launch<2>(vx, vy, dst, ncols_x, nrows_x, stride_row_x, stride_col_y, stride_col_dst, stream); break; + case 3: mul_mat_vec_ptq1_0_tc_launch<3>(vx, vy, dst, ncols_x, nrows_x, stride_row_x, stride_col_y, stride_col_dst, stream); break; + case 4: mul_mat_vec_ptq1_0_tc_launch<4>(vx, vy, dst, ncols_x, nrows_x, stride_row_x, stride_col_y, stride_col_dst, stream); break; + case 5: mul_mat_vec_ptq1_0_tc_launch<5>(vx, vy, dst, ncols_x, nrows_x, stride_row_x, stride_col_y, stride_col_dst, stream); break; + case 6: mul_mat_vec_ptq1_0_tc_launch<6>(vx, vy, dst, ncols_x, nrows_x, stride_row_x, stride_col_y, stride_col_dst, stream); break; + case 7: mul_mat_vec_ptq1_0_tc_launch<7>(vx, vy, dst, ncols_x, nrows_x, stride_row_x, stride_col_y, stride_col_dst, stream); break; + case 8: mul_mat_vec_ptq1_0_tc_launch<8>(vx, vy, dst, ncols_x, nrows_x, stride_row_x, stride_col_y, stride_col_dst, stream); break; + default: GGML_ABORT("PTQ1_0 tensor-core mat-vec takes 2 to 8 columns"); + } +} + // true when the dedicated kernel handles this call (plain 2D, K a multiple of 128, 1 to 4 columns) static bool mul_mat_vec_ptq1_0_pt_switch( const void * vx, const void * vy, const ggml_cuda_mm_fusion_args_device & fusion, float * dst, diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 3678c788..8dfbea4b 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -301,7 +301,7 @@ bool ggml_cuda_should_use_mmvq(enum ggml_type type, int cc, int64_t ne11) { // MMQ tile path up to 4 columns; from 5 on the branch-free MMQ tile loader is faster // (RTX 3060, K = 5120 shapes: a batch of 8 in 64.4 ms through MMQ against 112.7 ms // through the 8-column mat-vec, llama-bench pp8), so 5 and above take that path - return ne11 <= PTQ1_0_PT_MAX_COLS; + return ne11 <= (ptq1_0_tc_enabled(cc) ? PTQ1_0_TC_MAX_COLS : PTQ1_0_PT_MAX_COLS); } #endif // k-quants cost more to decode and mvq redoes that per column, so MMQ wins sooner. @@ -1504,6 +1504,21 @@ void ggml_cuda_mul_mat_vec_q( const int64_t ne10_padded = GGML_PAD(ne10, MATRIX_ROW_PADDING); ggml_cuda_pool_alloc src1_q8_1(ctx.pool(), ne13*ne12 * ne11*ne10_padded * sizeof(block_q8_1)/QK8_1); + +#if !defined(GGML_USE_HIP) + // PTQ1_0 with 2 to 8 columns, plain 2D, no fusion: the tensor-core mat-vec, which takes its own activation layout + const bool has_fusion = fusion_local.x_bias || fusion_local.gate || fusion_local.gate_bias; + if (src0->type == GGML_TYPE_PTQ1_0 && !ids && !has_fusion && ne11 >= 3 && ne11 <= PTQ1_0_TC_MAX_COLS && + ne02 == 1 && ne03 == 1 && ne12 == 1 && ne13 == 1 && ne00 % QK_PTQ1_0 == 0 && + ptq1_0_tc_enabled(ggml_cuda_info().devices[ggml_cuda_get_device()].cc)) { + quantize_row_q8_1_ptq1_0_tc_cuda(src1_d, src1_q8_1.get(), ne10, src1->nb[1] / ts_src1, src1->nb[2] / ts_src1, + src1->nb[3] / ts_src1, ne10_padded, ne11, ne12, ne13, stream); + mul_mat_vec_ptq1_0_tc_cuda(src0->data, src1_q8_1.get(), dst_d, ne00, ne01, ne11, src0->nb[1] / ts_src0, + ne10_padded / QK8_1, dst->nb[1] / ts_dst, stream); + return; + } +#endif + { const int64_t s11 = src1->nb[1] / ts_src1; const int64_t s12 = src1->nb[2] / ts_src1; diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index 5c661286..8cbee011 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -53,8 +53,9 @@ static __device__ __forceinline__ float nvfp4_native_scale_error( #endif // defined(BLACKWELL_MMA_AVAILABLE) // pt: write the planar-transposed layout consumed by the PTQ1_0 mat-vec path -// (see mmvq-ptq1_0.cuh) instead of block_q8_1; same quantization, same bytes per row -template +// (see mmvq-ptq1_0.cuh) instead of block_q8_1; same quantization, same bytes per row. +// tc (with pt): the block-major layout of the PTQ1_0 tensor-core mat-vec instead +template __launch_bounds__(CUDA_QUANTIZE_BLOCK_SIZE, 1) static __global__ void quantize_q8_1( const float * x_ptr, void * vy_ptr, @@ -102,6 +103,17 @@ static __global__ void quantize_q8_1( const int64_t nblk = ne0 / QK_PTQ1_0; const int64_t kb = i0 / QK_PTQ1_0; const int e = i0 % QK_PTQ1_0; + if constexpr (tc) { + // element e = 32k + 16h + 4q + b of block kb at kb*144 + 32q + 8k + 4h + b, the four (d, s) after the 128 quants + const int k = e / 32, h = (e / 16) % 2, qq = (e / 4) % 4, b = e % 4; + ycol[kb*144 + 32*qq + 8*k + 4*h + b] = q; + if (iqs > 0) { + return; + } + half2 * ds = (half2 *) (ycol + kb*144 + 128) + k; + *ds = make_half2(d, sum); + return; + } ycol[((e / 16)*nblk + kb) * 16 + (e % 16)] = q; if (iqs > 0) { @@ -677,6 +689,20 @@ void quantize_row_q8_1_cuda( ggml_cuda_kernel_launch(quantize_q8_1, launch_params, x, vy, ne00, s01, s02, s03, ne0, ne1, ne2_fastdiv); } +void quantize_row_q8_1_ptq1_0_tc_cuda( + const float * x, void * vy, const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, + const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, cudaStream_t stream) { + GGML_ASSERT(ne0 % QK_PTQ1_0 == 0); + + const uint3 ne2_fastdiv = init_fastdiv_values(ne2); + + const int64_t block_num_x = (ne0 + CUDA_QUANTIZE_BLOCK_SIZE - 1) / CUDA_QUANTIZE_BLOCK_SIZE; + const dim3 num_blocks(block_num_x, ne1, ne2*ne3); + const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE, 1, 1); + const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(num_blocks, block_size, 0, stream); + ggml_cuda_kernel_launch(quantize_q8_1, launch_params, x, vy, ne00, s01, s02, s03, ne0, ne1, ne2_fastdiv); +} + void quantize_mmq_q8_1_cuda( const float * x, const int32_t * ids, void * vy, const ggml_type type_src0, const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, diff --git a/ggml/src/ggml-cuda/quantize.cuh b/ggml/src/ggml-cuda/quantize.cuh index 335ff587..6d668298 100644 --- a/ggml/src/ggml-cuda/quantize.cuh +++ b/ggml/src/ggml-cuda/quantize.cuh @@ -21,6 +21,11 @@ void quantize_row_q8_1_cuda( ggml_type type_src0, int64_t ne00, int64_t s01, int64_t s02, int64_t s03, int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne3, cudaStream_t stream); +// activations in the layout of the PTQ1_0 tensor-core mat-vec (mmvq-ptq1_0.cuh), same quantization as quantize_row_q8_1_cuda +void quantize_row_q8_1_ptq1_0_tc_cuda( + const float * x, void * vy, int64_t ne00, int64_t s01, int64_t s02, int64_t s03, + int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne3, cudaStream_t stream); + void quantize_mmq_q8_1_cuda( const float * x, const int32_t * ids, void * vy, ggml_type type_src0, int64_t ne00, int64_t s01, int64_t s02, int64_t s03,