diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 896ba0df..bc7fadb0 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -515,6 +515,13 @@ struct ggml_cuda_type_traits { static constexpr int qi = QI1_S; }; +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = 32; + static constexpr int qr = 2; + static constexpr int qi = 4; +}; + template<> struct ggml_cuda_type_traits { static constexpr int qk = QK_K; diff --git a/ggml/src/ggml-cuda/iqk_mmvq.cu b/ggml/src/ggml-cuda/iqk_mmvq.cu index b08ea718..2340b54a 100644 --- a/ggml/src/ggml-cuda/iqk_mmvq.cu +++ b/ggml/src/ggml-cuda/iqk_mmvq.cu @@ -50,13 +50,6 @@ struct ggml_cuda_type_traits { static constexpr int qi = QI5_XS; }; -template<> -struct ggml_cuda_type_traits { - static constexpr int qk = 32; - static constexpr int qr = 2; - static constexpr int qi = 4; -}; - // Reminder: // constexpr int qk = ggml_cuda_type_traits::qk; diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 2f7a9bfd..60c2037f 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -85,6 +85,9 @@ void ggml_cuda_op_mul_mat_q( case GGML_TYPE_IQ1_S: mul_mat_q_case(ctx, args, stream); break; + case GGML_TYPE_IQ1_S_R4: + mul_mat_q_case(ctx, args, stream); + break; case GGML_TYPE_IQ4_XS: mul_mat_q_case(ctx, args, stream); break; @@ -150,6 +153,7 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11) { case GGML_TYPE_IQ3_XXS: case GGML_TYPE_IQ3_S: case GGML_TYPE_IQ1_S: + case GGML_TYPE_IQ1_S_R4: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_KS: @@ -174,6 +178,9 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11) { if (int8_mma_available(cc)) { return true; } + if (type == GGML_TYPE_IQ1_S_R4) { + return false; + } if (cc < MIN_CC_DP4A) { return false; diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 7a51c514..c98fa561 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -79,6 +79,7 @@ static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) { case GGML_TYPE_IQ3_S: return MMQ_Q8_1_DS_LAYOUT_D4; case GGML_TYPE_IQ1_S: + case GGML_TYPE_IQ1_S_R4: return MMQ_Q8_1_DS_LAYOUT_DS4; case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ4_NL: @@ -186,6 +187,7 @@ static constexpr __host__ __device__ tile_x_sizes mmq_get_dp4a_tile_x_sizes(ggml case GGML_TYPE_IQ3_XXS : return MMQ_DP4A_TXS_Q8_0; case GGML_TYPE_IQ3_S : return MMQ_DP4A_TXS_Q8_0; case GGML_TYPE_IQ1_S : return MMQ_DP4A_TXS_Q8_0; + case GGML_TYPE_IQ1_S_R4: return MMQ_DP4A_TXS_Q8_0; case GGML_TYPE_IQ4_XS : return MMQ_DP4A_TXS_Q8_0; case GGML_TYPE_IQ4_NL : return MMQ_DP4A_TXS_Q8_0; case GGML_TYPE_IQ4_KS : return MMQ_DP4A_TXS_Q8_0; @@ -231,6 +233,7 @@ static constexpr __host__ __device__ int mmq_get_mma_tile_x_k(ggml_type type) { case GGML_TYPE_IQ3_XXS : return MMQ_MMA_TILE_X_K_Q8_0; case GGML_TYPE_IQ3_S : return MMQ_MMA_TILE_X_K_Q8_0; case GGML_TYPE_IQ1_S : return MMQ_MMA_TILE_X_K_Q8_0; + case GGML_TYPE_IQ1_S_R4: return MMQ_MMA_TILE_X_K_Q8_0; case GGML_TYPE_IQ4_XS : return MMQ_MMA_TILE_X_K_Q8_0; case GGML_TYPE_IQ4_NL : return MMQ_MMA_TILE_X_K_Q8_0; case GGML_TYPE_IQ4_KS : return MMQ_MMA_TILE_X_K_Q8_0; @@ -318,6 +321,74 @@ template static __device__ __forceinlin } } +template static __device__ __forceinline__ void load_tiles_iq1_s_r4( + const char * __restrict__ x, int * __restrict__ x_tile, const int & kbx0, const int & i_max, const int & stride) { + +#ifdef INT8_MMA_AVAILABLE + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + 2*WARP_SIZE); +#else + constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q4_0, mmq_y); + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + txs.qs); +#endif // INT8_MMA_AVAILABLE + + const int kbx = threadIdx.x / 4; + const int kqsx = threadIdx.x % 4; + + int32_t grid32[2]; + +#pragma unroll + for (int i0 = 0; i0 < mmq_y; i0 += nwarps) { + int i = i0 + threadIdx.y; + + if (need_check) { + i = min(i, i_max); + } + const int i4 = i/4; + const int ir = i%4; + + const block_iq1_s_r4 * bxi = (const block_iq1_s_r4 *)(x + 4*i4*stride + 4*sizeof(half)) + kbx0 + kbx; + + grid32[0] = iq1s_grid_gpu[bxi->qs[4*kqsx+ir] | (((bxi->qh[ir] >> 3*kqsx) & 7) << 8)]; + grid32[1] = ((grid32[0] >> 4) & 0x0f0f0f0f) << 3; + grid32[0] = (grid32[0] & 0x0f0f0f0f) << 3; + const int shift = bxi->qh[ir] & 0x8000 ? 0x09090909 : 0x07070707; + +#ifdef INT8_MMA_AVAILABLE + x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + 8*kbx + 2*kqsx + 0] = __vsubss4(grid32[0], shift); + x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + 8*kbx + 2*kqsx + 1] = __vsubss4(grid32[1], shift); +#else + // TODO + //x_qs[i*(WARP_SIZE + 1) + threadIdx.x] = qs0; +#endif // INT8_MMA_AVAILABLE + } + + const int blocks_per_tile_x_row = WARP_SIZE / 4; + const int kbxd = threadIdx.x % blocks_per_tile_x_row; + +#pragma unroll + for (int i0 = 0; i0 < mmq_y; i0 += nwarps * 4) { + int i = i0 + threadIdx.y * 4 + threadIdx.x / blocks_per_tile_x_row; + + if (need_check) { + i = min(i, i_max); + } + const int i4 = i/4; + const int ir = i%4; + + const half * dptr = (const half *)(x + 4*i4*stride); + const block_iq1_s_r4 * bxi = (const block_iq1_s_r4 *)(dptr + 4) + kbx0 + kbxd; + +#ifdef INT8_MMA_AVAILABLE + x_df[i*MMQ_MMA_TILE_X_K_Q8_0 + kbxd] = 0.125f * __half2float(dptr[ir]) * (((bxi->qh[ir] >> 11) & 14) + 1); +#else + // TODO + //x_df[i*(WARP_SIZE/QI4_0) + i/QI4_0 + kbxd] = bxi->d; +#endif // INT8_MMA_AVAILABLE + } +} + template static __device__ __forceinline__ void vec_dot_q4_0_q8_1_dp4a( const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int & k00) { @@ -3132,6 +3203,14 @@ struct mmq_type_traits { static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q8_1_q8_1_dp4a; }; +template +struct mmq_type_traits { + static constexpr int vdr = VDR_Q4_0_Q8_1_MMQ; + static constexpr load_tiles_mmq_t load_tiles = load_tiles_iq1_s_r4; + static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_0_q8_1_mma; + static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q4_0_q8_1_dp4a; +}; + template struct mmq_type_traits { static constexpr int vdr = VDR_IQ4_NL_Q8_1_MMQ; @@ -3656,6 +3735,7 @@ extern DECL_MMQ_CASE(GGML_TYPE_IQ4_K); extern DECL_MMQ_CASE(GGML_TYPE_IQ5_K); extern DECL_MMQ_CASE(GGML_TYPE_IQ5_KS); extern DECL_MMQ_CASE(GGML_TYPE_IQ6_K); +extern DECL_MMQ_CASE(GGML_TYPE_IQ1_S_R4); // ------------------------------------------------------------------------------------------------------------------------- diff --git a/ggml/src/ggml-cuda/template-instances/mmq-instance-iq1_s_r4.cu b/ggml/src/ggml-cuda/template-instances/mmq-instance-iq1_s_r4.cu new file mode 100644 index 00000000..7bc5cf1a --- /dev/null +++ b/ggml/src/ggml-cuda/template-instances/mmq-instance-iq1_s_r4.cu @@ -0,0 +1,5 @@ +// This file has been autogenerated by generate_cu_files.py, do not edit manually. + +#include "../mmq.cuh" + +DECL_MMQ_CASE(GGML_TYPE_IQ1_S_R4);