diff --git a/ggml/src/ggml-cuda/mmq-vec-dot.cuh b/ggml/src/ggml-cuda/mmq-vec-dot.cuh
index be091d27..65ca3f3d 100644
--- a/ggml/src/ggml-cuda/mmq-vec-dot.cuh
+++ b/ggml/src/ggml-cuda/mmq-vec-dot.cuh
@@ -179,7 +179,7 @@ static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma(
float dB;
const int j = j0 + tile_C::get_j(0);
- if (ds_layout == MMQ_Q8_1_DS_LAYOUT_D4) {
+ if (ds_layout == MMQ_Q8_1_DS_LAYOUT_D4 || ds_layout == MMQ_Q8_1_DS_LAYOUT_D1) {
dB = y_df[j*MMQ_TILE_Y_K + k01/QI8_1];
} else {
dB = __low2float(y_ds[j*MMQ_TILE_Y_K + k01/QI8_1]);
@@ -244,6 +244,39 @@ static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma(
}
}
+ if constexpr (ds_layout == MMQ_Q8_1_DS_LAYOUT_D1) {
+ // x and y both have one scale per MMQ_TILE_NE_K ints (128 values), so the k01 mmas accumulate in int32
+ // and the result is rescaled once.
+#pragma unroll
+ for (int j0 = 0; j0 < J; j0 += ntx*tile_C::J) {
+ tile_C C[ntx];
+#pragma unroll
+ for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += QI8_0) {
+ tile_B B;
+ load_generic(B, y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K);
+#pragma unroll
+ for (int n = 0; n < ntx; ++n) {
+ mma(C[n], A[n][k01/QI8_0], B);
+ }
+ }
+
+ float dB[tile_C::ne/2];
+#pragma unroll
+ for (int l = 0; l < tile_C::ne/2; ++l) {
+ dB[l] = y_df[(j0 + tile_C::get_j(l))*MMQ_TILE_Y_K];
+ }
+
+#pragma unroll
+ for (int n = 0; n < ntx; ++n) {
+#pragma unroll
+ for (int l = 0; l < tile_C::ne; ++l) {
+ sum[(j0/tile_C::J + n)*tile_C::ne + l] += C[n].x[l]*dA[n][l/2][0]*dB[l%2];
+ }
+ }
+ }
+ return;
+ }
+
#pragma unroll
for (int j0 = 0; j0 < J; j0 += ntx*tile_C::J) {
#pragma unroll
@@ -257,7 +290,7 @@ static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma(
for (int l = 0; l < tile_C::ne/2; ++l) {
const int j = j0 + tile_C::get_j(l);
- if (ds_layout == MMQ_Q8_1_DS_LAYOUT_D4) {
+ if (ds_layout == MMQ_Q8_1_DS_LAYOUT_D4 || ds_layout == MMQ_Q8_1_DS_LAYOUT_D1) {
dB[l] = y_df[j*MMQ_TILE_Y_K + k01/QI8_1];
} else {
dB[l] = __low2float(y_ds[j*MMQ_TILE_Y_K + k01/QI8_1]);
diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh
index 47de164c..99b05d86 100644
--- a/ggml/src/ggml-cuda/mmq.cuh
+++ b/ggml/src/ggml-cuda/mmq.cuh
@@ -8,6 +8,13 @@
#define MMQ_DP4A_MAX_BATCH_SIZE 64 // Max. batch size to use for dp4a MMQ kernels when FP16 tensor cores are available.
#define MMQ_PTQ1_0_MAX_BATCH_SIZE (1 << 30) // MMQ at every batch unless GGML_CUDA_PTQ1_0_MMQ_MAX_BATCH lowers it
+
+// PTQ1_0 MMQ quantizes src1 with one int8 scale per 128 values (one PTQ1_0 block) instead of one per 32, so the
+// tensor cores accumulate a whole block in int32 and the float rescale runs once per 128 instead of once per 32.
+// Changes results versus per-32 activation scales; set to 0 for the per-32 layout.
+#ifndef GGML_CUDA_PTQ1_0_MMQ_Y128
+#define GGML_CUDA_PTQ1_0_MMQ_Y128 1
+#endif
#define MMQ_ITER_K 256
#define MMQ_ITER_K_FP4 512
#define MMQ_NWARPS 8
@@ -19,6 +26,7 @@ typedef void (*ggml_cuda_mmq_write_back_t)(const float * __restrict__ sum, const
enum mmq_q8_1_ds_layout {
MMQ_Q8_1_DS_LAYOUT_D4,
+ MMQ_Q8_1_DS_LAYOUT_D1, // like D4, but one scale per 128 values, written to all four d4 slots
MMQ_Q8_1_DS_LAYOUT_DS4,
MMQ_Q8_1_DS_LAYOUT_D2S6,
};
@@ -64,10 +72,11 @@ static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) {
case GGML_TYPE_Q1_0:
case GGML_TYPE_Q2_0:
case GGML_TYPE_PQ2_0:
+ return MMQ_Q8_1_DS_LAYOUT_D4;
#if !defined(GGML_USE_HIP)
case GGML_TYPE_PTQ1_0:
+ return GGML_CUDA_PTQ1_0_MMQ_Y128 ? MMQ_Q8_1_DS_LAYOUT_D1 : MMQ_Q8_1_DS_LAYOUT_D4;
#endif
- return MMQ_Q8_1_DS_LAYOUT_D4;
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q4_1:
return MMQ_Q8_1_DS_LAYOUT_DS4;
@@ -756,7 +765,8 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func
case GGML_TYPE_PTQ1_0:
return ggml_cuda_mmq_util_funcs(
-1, ggml_cuda_mmq_load_tiles_ptq1_0<type, J, fallback>,
- ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma<type, J, fallback, MMQ_Q8_1_DS_LAYOUT_D4>,
+ ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma<type, J, fallback,
+ GGML_CUDA_PTQ1_0_MMQ_Y128 ? MMQ_Q8_1_DS_LAYOUT_D1 : MMQ_Q8_1_DS_LAYOUT_D4>,
ggml_cuda_mmq_write_back_mma<type, J, fallback>);
#endif
case GGML_TYPE_Q4_0:
diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu
index 8cbee011..ad4306b7 100644
--- a/ggml/src/ggml-cuda/quantize.cu
+++ b/ggml/src/ggml-cuda/quantize.cu
@@ -497,7 +497,7 @@ static __global__ void quantize_mmq_q8_1(
const float * __restrict__ norm_weight = nullptr,
const float * __restrict__ norm_scale = nullptr) {
- constexpr int vals_per_scale = ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6 ? 64 : 32;
+ constexpr int vals_per_scale = ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6 ? 64 : ds_layout == MMQ_Q8_1_DS_LAYOUT_D1 ? 128 : 32;
constexpr int vals_per_sum = ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6 ? 16 : 32;
const int64_t i0 = ((int64_t)blockDim.x*blockIdx.y + threadIdx.x)*4;
@@ -559,7 +559,7 @@ static __global__ void quantize_mmq_q8_1(
}
float sum;
- if (ds_layout != MMQ_Q8_1_DS_LAYOUT_D4) {
+ if (ds_layout != MMQ_Q8_1_DS_LAYOUT_D4 && ds_layout != MMQ_Q8_1_DS_LAYOUT_D1) {
sum = xi.x + xi.y + xi.z + xi.w;
// Calculate sums across vals_per_sum/4 threads.
@@ -627,6 +627,10 @@ void quantize_mmq_q8_1_swiglu_cuda(
quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D4, false, true><<<num_blocks, block_size, 0, stream>>>(
x, nullptr, vy, ne00, s01, s02, s03, ne0, ne1, ne2, 0, gate);
break;
+ case MMQ_Q8_1_DS_LAYOUT_D1:
+ quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D1, false, true><<<num_blocks, block_size, 0, stream>>>(
+ x, nullptr, vy, ne00, s01, s02, s03, ne0, ne1, ne2, 0, gate);
+ break;
case MMQ_Q8_1_DS_LAYOUT_DS4:
quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_DS4, false, true><<<num_blocks, block_size, 0, stream>>>(
x, nullptr, vy, ne00, s01, s02, s03, ne0, ne1, ne2, 0, gate);
@@ -655,6 +659,10 @@ void quantize_mmq_q8_1_rms_cuda(
quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D4, false, false, true><<<num_blocks, block_size, 0, stream>>>(
x, nullptr, vy, ne00, s01, s02, s03, ne0, ne1, ne2, 0, nullptr, weight, scale);
break;
+ case MMQ_Q8_1_DS_LAYOUT_D1:
+ quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D1, false, false, true><<<num_blocks, block_size, 0, stream>>>(
+ x, nullptr, vy, ne00, s01, s02, s03, ne0, ne1, ne2, 0, nullptr, weight, scale);
+ break;
case MMQ_Q8_1_DS_LAYOUT_DS4:
quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_DS4, false, false, true><<<num_blocks, block_size, 0, stream>>>(
x, nullptr, vy, ne00, s01, s02, s03, ne0, ne1, ne2, 0, nullptr, weight, scale);
@@ -719,6 +727,10 @@ void quantize_mmq_q8_1_cuda(
quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D4, false>
<<<num_blocks, block_size, 0, stream>>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0);
break;
+ case MMQ_Q8_1_DS_LAYOUT_D1:
+ quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D1, false>
+ <<<num_blocks, block_size, 0, stream>>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0);
+ break;
case MMQ_Q8_1_DS_LAYOUT_DS4:
quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_DS4, false>
<<<num_blocks, block_size, 0, stream>>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, /*n_expert_used=*/0);
@@ -749,6 +761,10 @@ void quantize_scatter_mmq_q8_1_cuda(
quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D4, true><<<num_blocks, block_size, 0, stream>>>(
x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used);
break;
+ case MMQ_Q8_1_DS_LAYOUT_D1:
+ quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D1, true><<<num_blocks, block_size, 0, stream>>>(
+ x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used);
+ break;
case MMQ_Q8_1_DS_LAYOUT_DS4:
quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_DS4, true><<<num_blocks, block_size, 0, stream>>>(
x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used);