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 <int ncols>
+__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 <int ncols>
+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<ncols>, 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<char> 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 <bool pt>
+// (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 <bool pt, bool tc = false>
__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<false>, 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<true, true>, 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,