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, - ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, + ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, ggml_cuda_mmq_write_back_mma); #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<<>>( x, nullptr, vy, ne00, s01, s02, s03, ne0, ne1, ne2, 0, gate); break; + case MMQ_Q8_1_DS_LAYOUT_D1: + quantize_mmq_q8_1<<>>( + x, nullptr, vy, ne00, s01, s02, s03, ne0, ne1, ne2, 0, gate); + break; case MMQ_Q8_1_DS_LAYOUT_DS4: quantize_mmq_q8_1<<>>( 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<<>>( 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<<>>( + 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<<>>( 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 <<>>(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 + <<>>(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 <<>>(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<<>>( 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<<>>( + 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<<>>( x, ids_src1_inv, vy, ne00, /*s01=*/0, /*s02=*/stride_token, /*s03=*/0, ne0, /*ne1=*/(int) nrows_dst, /*ne2=*/1, n_expert_used);