diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cu index 06dfae700..b91875d64 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cu @@ -20,6 +20,15 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, con case GGML_TYPE_Q8_0: mul_mat_q_case(ctx, args, stream); break; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + mul_mat_q_case(ctx, args, stream); + break; + case GGML_TYPE_Q2_0_ROCMFP2: + mul_mat_q_case(ctx, args, stream); + break; + case GGML_TYPE_Q3_0_ROCMFPX: + mul_mat_q_case(ctx, args, stream); + break; case GGML_TYPE_MXFP4: case GGML_TYPE_NVFP4: #ifndef GGML_CUDA_BLACKWELL_CONSUMER @@ -316,6 +325,12 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t case GGML_TYPE_IQ4_NL: mmq_supported = true; break; + case GGML_TYPE_Q2_0_ROCMFP2: + case GGML_TYPE_Q3_0_ROCMFPX: + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + // ROCmFPX MMQ variants are implemented for gfx1151 only. + mmq_supported = GGML_CUDA_CC_IS_RDNA3_5(cc); + break; default: mmq_supported = false; break; diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh index d8df400fb..2e2cc03b3 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh @@ -66,6 +66,11 @@ static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) { return MMQ_Q8_1_DS_LAYOUT_DS4; case GGML_TYPE_Q8_0: return MMQ_Q8_1_DS_LAYOUT_D4; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + return MMQ_Q8_1_DS_LAYOUT_D4; + case GGML_TYPE_Q2_0_ROCMFP2: + case GGML_TYPE_Q3_0_ROCMFPX: + return MMQ_Q8_1_DS_LAYOUT_D4; case GGML_TYPE_MXFP4: return MMQ_Q8_1_DS_LAYOUT_D4; case GGML_TYPE_NVFP4: @@ -212,6 +217,7 @@ static constexpr __device__ int get_mmq_y_device() { #define MMQ_DP4A_TXS_Q4_1 tile_x_sizes{mmq_y*MMQ_TILE_NE_K + mmq_y, mmq_y*MMQ_TILE_NE_K/QI4_1 + mmq_y/QI4_1, 0} #define MMQ_DP4A_TXS_Q8_0 tile_x_sizes{mmq_y*MMQ_TILE_NE_K*2 + mmq_y, mmq_y*MMQ_TILE_NE_K*2/QI8_0 + mmq_y/(QI8_0/2), 0} #define MMQ_DP4A_TXS_Q8_0_16 tile_x_sizes{mmq_y*MMQ_TILE_NE_K*2 + mmq_y, mmq_y*MMQ_TILE_NE_K*4/QI8_0 + mmq_y/(QI8_0/4), 0} +#define MMQ_DP4A_TXS_ROCMFPX tile_x_sizes{mmq_y*MMQ_TILE_NE_K*2 + mmq_y, mmq_y*MMQ_TILE_NE_K/2 + mmq_y/2, 0} #define MMQ_DP4A_TXS_Q8_1 tile_x_sizes{mmq_y*MMQ_TILE_NE_K*2 + mmq_y, mmq_y*MMQ_TILE_NE_K*2/QI8_1 + mmq_y/(QI8_1/2), 0} #define MMQ_DP4A_TXS_Q2_K tile_x_sizes{mmq_y*MMQ_TILE_NE_K*2 + mmq_y, mmq_y*MMQ_TILE_NE_K + mmq_y, 0} #define MMQ_DP4A_TXS_Q3_K tile_x_sizes{mmq_y*MMQ_TILE_NE_K*2 + mmq_y, mmq_y, mmq_y*MMQ_TILE_NE_K/8 + mmq_y/8} @@ -226,6 +232,11 @@ static constexpr __host__ __device__ tile_x_sizes mmq_get_dp4a_tile_x_sizes(ggml case GGML_TYPE_Q5_0: return MMQ_DP4A_TXS_Q8_0; case GGML_TYPE_Q5_1: return MMQ_DP4A_TXS_Q8_1; case GGML_TYPE_Q8_0: return MMQ_DP4A_TXS_Q8_0; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + return MMQ_DP4A_TXS_Q8_0; + case GGML_TYPE_Q2_0_ROCMFP2: + case GGML_TYPE_Q3_0_ROCMFPX: + return MMQ_DP4A_TXS_ROCMFPX; case GGML_TYPE_MXFP4: return MMQ_DP4A_TXS_Q8_1; case GGML_TYPE_NVFP4: return MMQ_DP4A_TXS_Q8_0_16; case GGML_TYPE_Q2_K: return MMQ_DP4A_TXS_Q2_K; @@ -248,6 +259,7 @@ static constexpr __host__ __device__ tile_x_sizes mmq_get_dp4a_tile_x_sizes(ggml #define MMQ_MMA_TILE_X_K_Q8_0 (2*MMQ_TILE_NE_K + 2*MMQ_TILE_NE_K/QI8_0 + 4) #define MMQ_MMA_TILE_X_K_FP4 (2*MMQ_TILE_NE_K + 8 + 4) // MXFP4 #define MMQ_MMA_TILE_X_K_NVFP4 (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K/2 + 4) // NVFP4 +#define MMQ_MMA_TILE_X_K_ROCMFPX (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K/2 + 4) #define MMQ_MMA_TILE_X_K_Q8_1 (2*MMQ_TILE_NE_K + 2*MMQ_TILE_NE_K/QI8_0 + 4) #define MMQ_MMA_TILE_X_K_Q2_K (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K + 4) #define MMQ_MMA_TILE_X_K_Q3_K (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K/2 + 4) @@ -261,6 +273,7 @@ static_assert(MMQ_MMA_TILE_X_K_Q6_K % 8 == 4, "Wrong padding."); static_assert(MMQ_MMA_TILE_X_K_FP4 % 8 == 4, "Wrong padding."); static_assert(MMQ_MMA_TILE_X_K_FP4 == MMQ_MMA_TILE_X_K_Q8_1, "Wrong tile size for MXFP4"); static_assert(MMQ_MMA_TILE_X_K_NVFP4 % 8 == 4, "Wrong padding."); +static_assert(MMQ_MMA_TILE_X_K_ROCMFPX % 8 == 4, "Wrong padding."); static constexpr __host__ __device__ int mmq_get_mma_tile_x_k(ggml_type type) { @@ -270,6 +283,11 @@ static constexpr __host__ __device__ int mmq_get_mma_tile_x_k(ggml_type type) { case GGML_TYPE_Q5_0: return MMQ_MMA_TILE_X_K_Q8_0; case GGML_TYPE_Q5_1: return MMQ_MMA_TILE_X_K_Q8_1; case GGML_TYPE_Q8_0: return MMQ_MMA_TILE_X_K_Q8_0; + case GGML_TYPE_Q4_0_ROCMFP4_FAST: + return MMQ_MMA_TILE_X_K_Q8_0; + case GGML_TYPE_Q2_0_ROCMFP2: + case GGML_TYPE_Q3_0_ROCMFPX: + return MMQ_MMA_TILE_X_K_ROCMFPX; // tile sizes are the same for Q8_1 and FP4 for blackwell case GGML_TYPE_MXFP4: return MMQ_MMA_TILE_X_K_Q8_1; case GGML_TYPE_NVFP4: return MMQ_MMA_TILE_X_K_NVFP4; @@ -780,6 +798,269 @@ template static __device__ __forceinline__ void loa } } +// ROCmFP4_FAST has one UE4M3 scale per 32-weight block. Its expanded int8 +// tile is therefore layout-compatible with Q8_0 MMQ after Codebook10 decode. +template static __device__ __forceinline__ void load_tiles_rocmfp4_fast( + const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + constexpr int nwarps = mmq_get_nwarps_device(); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + +#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + 2*MMQ_TILE_NE_K); +#else + constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q4_0_ROCMFP4_FAST, mmq_y); + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + txs.qs); +#endif + + constexpr int words_per_block = QK_ROCMFP4 / 4; + constexpr int threads_per_row = MMQ_ITER_K / 8; + constexpr int nrows = warp_size / threads_per_row; + const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; + +#pragma unroll + for (int i0 = 0; i0 < mmq_y; i0 += nrows*nwarps) { + int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); + if (need_check) { + i = min(i, i_max); + } + +#pragma unroll + for (int part = 0; part < 2; ++part) { + const int word = part*MMQ_TILE_NE_K + txi; + const int block = word / words_per_block; + const int word_in_block = word % words_per_block; + const block_rocmfp4_fast * bxi = (const block_rocmfp4_fast *) x + kbx0 + i*stride + block; + const int2 values = rocmfp4_get_int_from_codebook_16( + rocmfp4_get_qs_i32(bxi->qs, word_in_block % 4), kvalues_rocmfp4); + const int packed = word_in_block < 4 ? values.x : values.y; + +#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + word] = packed; +#else + x_qs[i*(2*MMQ_TILE_NE_K + 1) + word] = packed; +#endif + } + + if (txi < 2*MMQ_TILE_NE_K/QI8_0) { + const block_rocmfp4_fast * bxi = (const block_rocmfp4_fast *) x + kbx0 + i*stride + txi; + const float scale = rocmfp4_ue4m3_to_fp32_half_finite(bxi->e); + +#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) + x_df[i*MMQ_MMA_TILE_X_K_Q8_0 + txi] = scale; +#else + x_df[i*(2*MMQ_TILE_NE_K/QI8_0) + i/(QI8_0/2) + txi] = scale; +#endif + } + } +} + +// ROCmFP2/FP3 use two UE4M3 scales per 32-value block. Expand each packed +// block into the same int8 tile shape as Q8_0, but retain one scale per +// 16-value half-block so the DP4A reduction remains numerically exact. +template static __device__ __forceinline__ void load_tiles_rocmfp2( + const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + constexpr int nwarps = mmq_get_nwarps_device(); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + +#if defined(AMD_WMMA_AVAILABLE) + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + 2*MMQ_TILE_NE_K); +#else + constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q2_0_ROCMFP2, mmq_y); + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + txs.qs); +#endif + + constexpr int words_per_block = QK_ROCMFP2 / 4; + constexpr int threads_per_row = MMQ_ITER_K / 8; + constexpr int nrows = warp_size / threads_per_row; + const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; + +#pragma unroll + for (int i0 = 0; i0 < mmq_y; i0 += nrows*nwarps) { + int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); + if (need_check) { + i = min(i, i_max); + } + +#pragma unroll + for (int part = 0; part < 2; ++part) { + const int word = part*MMQ_TILE_NE_K + txi; + const int block = word / words_per_block; + const int word_in_block = word % words_per_block; + const block_rocmfp2 * bxi = (const block_rocmfp2 *) x + kbx0 + i*stride + block; +#if defined(AMD_WMMA_AVAILABLE) + x_qs[i*MMQ_MMA_TILE_X_K_ROCMFPX + word] = +#else + x_qs[i*(2*MMQ_TILE_NE_K + 1) + word] = +#endif + rocmfpx_pack4_fp2_vec_cuda(bxi->qs, 4*word_in_block); + } + + if (txi < MMQ_TILE_NE_K/2) { + const int half = txi; + const int block = half / 2; + const block_rocmfp2 * bxi = (const block_rocmfp2 *) x + kbx0 + i*stride + block; +#if defined(AMD_WMMA_AVAILABLE) + x_df[i*MMQ_MMA_TILE_X_K_ROCMFPX + half] = rocmfpx_ue4m3_to_fp32_finite(bxi->e[half & 1]); +#else + x_df[i*(MMQ_TILE_NE_K/2) + i/2 + half] = rocmfpx_ue4m3_to_fp32_finite(bxi->e[half & 1]); +#endif + } + } +} + +template static __device__ __forceinline__ void load_tiles_rocmfp3( + const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { + constexpr int nwarps = mmq_get_nwarps_device(); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + +#if defined(AMD_WMMA_AVAILABLE) + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + 2*MMQ_TILE_NE_K); +#else + constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q3_0_ROCMFPX, mmq_y); + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + txs.qs); +#endif + + constexpr int words_per_block = QK_ROCMFP3 / 4; + constexpr int threads_per_row = MMQ_ITER_K / 8; + constexpr int nrows = warp_size / threads_per_row; + const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x; + +#pragma unroll + for (int i0 = 0; i0 < mmq_y; i0 += nrows*nwarps) { + int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row); + if (need_check) { + i = min(i, i_max); + } + +#pragma unroll + for (int part = 0; part < 2; ++part) { + const int word = part*MMQ_TILE_NE_K + txi; + const int block = word / words_per_block; + const int word_in_block = word % words_per_block; + const block_rocmfp3 * bxi = (const block_rocmfp3 *) x + kbx0 + i*stride + block; +#if defined(AMD_WMMA_AVAILABLE) + x_qs[i*MMQ_MMA_TILE_X_K_ROCMFPX + word] = +#else + x_qs[i*(2*MMQ_TILE_NE_K + 1) + word] = +#endif + rocmfpx_pack4_fp3_vec_cuda(bxi->qs, 4*word_in_block); + } + + if (txi < MMQ_TILE_NE_K/2) { + const int half = txi; + const int block = half / 2; + const block_rocmfp3 * bxi = (const block_rocmfp3 *) x + kbx0 + i*stride + block; +#if defined(AMD_WMMA_AVAILABLE) + x_df[i*MMQ_MMA_TILE_X_K_ROCMFPX + half] = rocmfpx_ue4m3_to_fp32_finite(bxi->e[half & 1]); +#else + x_df[i*(MMQ_TILE_NE_K/2) + i/2 + half] = rocmfpx_ue4m3_to_fp32_finite(bxi->e[half & 1]); +#endif + } + } +} + +template +static __device__ __forceinline__ void vec_dot_rocmfpx_q8_1_dp4a( + const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { + constexpr int nwarps = mmq_get_nwarps_device(); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q2_0_ROCMFP2, mmq_y); + + const int * x_qs = (const int *) x; + const float * x_df = (const float *) x_qs + txs.qs; + const int * y_qs = (const int *) y + 4; + const float * y_df = (const float *) y; + +#pragma unroll + for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += VDR_ROCMFP2_Q8_1_MMQ) { + const int k0 = k00 + k01; + +#pragma unroll + for (int j0 = 0; j0 < mmq_x; j0 += nwarps) { + const int j = j0 + threadIdx.y; + const float dB = y_df[j*MMQ_TILE_Y_K + (k0/QI8_1) % (MMQ_TILE_NE_K/QI8_1)]; + +#pragma unroll + for (int i0 = 0; i0 < mmq_y; i0 += warp_size) { + const int i = i0 + threadIdx.x; + int dot = 0; +#pragma unroll + for (int l = 0; l < VDR_ROCMFP2_Q8_1_MMQ; ++l) { + dot = ggml_cuda_dp4a(x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0 + l], + y_qs[j*MMQ_TILE_Y_K + (k01 + l) % MMQ_TILE_NE_K], dot); + } + const float dA = x_df[i*(MMQ_TILE_NE_K/2) + i/2 + k0/VDR_ROCMFP2_Q8_1_MMQ]; + sum[j0/nwarps*mmq_y/warp_size + i0/warp_size] += (float) dot * dA * dB; + } + } + } +} + +template +static __device__ __forceinline__ void vec_dot_rocmfpx_q8_1_mma( + const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { +#if defined(AMD_WMMA_AVAILABLE) + constexpr data_layout input_layout = get_input_data_layout(); + typedef tile<16, 4, int, input_layout> tile_A; + typedef tile<16, 4, int, input_layout> tile_B; + typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C; + + constexpr int granularity = mmq_get_granularity_device(mmq_x); + constexpr int rows_per_warp = granularity; + constexpr int ntx = rows_per_warp/tile_C::I; + + y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K); + + const int * x_qs = (const int *) x; + const float * x_df = (const float *) x_qs + 2*MMQ_TILE_NE_K; + const int * y_qs = (const int *) y + 4; + const float * y_df = (const float *) y; + + const int i0 = (threadIdx.y / ntx) * rows_per_warp; + + for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += 4) { + const int k0 = k00 + k01; + + tile_A A[ntx]; +#pragma unroll + for (int n = 0; n < ntx; ++n) { + load_generic(A[n], x_qs + (i0 + n*tile_A::I)*MMQ_MMA_TILE_X_K_ROCMFPX + k0, + MMQ_MMA_TILE_X_K_ROCMFPX); + } + +#pragma unroll + for (int j0 = 0; j0 < mmq_x; j0 += ntx*tile_C::J) { + tile_B B; + load_generic(B, y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); + + const int j = j0 + tile_C::get_j(0); + const float dB = y_df[j*MMQ_TILE_Y_K + k01/QI8_1]; + +#pragma unroll + for (int n = 0; n < ntx; ++n) { + tile_C C; + mma(C, A[n], B); + +#pragma unroll + for (int l = 0; l < tile_C::ne; ++l) { + const int i = i0 + n*tile_C::I + tile_C::get_i(l); + const float dA = x_df[i*MMQ_MMA_TILE_X_K_ROCMFPX + k0/4]; + sum[(j0/tile_C::J + n)*tile_C::ne + l] += C.x[l]*dA*dB; + } + } + } + } +#else + vec_dot_rocmfpx_q8_1_dp4a(x, y, sum, k00); +#endif +} + template static __device__ __forceinline__ void load_tiles_mxfp4( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { constexpr int nwarps = mmq_get_nwarps_device(); @@ -3392,6 +3673,30 @@ struct mmq_type_traits { static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q8_0_q8_1_dp4a; }; +template +struct mmq_type_traits { + static constexpr int vdr = VDR_ROCMFP4_FAST_Q8_1_MMQ; + static constexpr load_tiles_mmq_t load_tiles = load_tiles_rocmfp4_fast; + 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_q8_0_q8_1_dp4a; +}; + +template +struct mmq_type_traits { + static constexpr int vdr = VDR_ROCMFP2_Q8_1_MMQ; + static constexpr load_tiles_mmq_t load_tiles = load_tiles_rocmfp2; + static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_rocmfpx_q8_1_mma; + static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_rocmfpx_q8_1_dp4a; +}; + +template +struct mmq_type_traits { + static constexpr int vdr = VDR_ROCMFP3_Q8_1_MMQ; + static constexpr load_tiles_mmq_t load_tiles = load_tiles_rocmfp3; + static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_rocmfpx_q8_1_mma; + static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_rocmfpx_q8_1_dp4a; +}; + template struct mmq_type_traits { static constexpr int vdr = VDR_MXFP4_Q8_1_MMQ; @@ -3534,9 +3839,17 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( int * tile_y = data_mul_mat_q + mmq_x; int * tile_x = tile_y + GGML_PAD(mmq_x*MMQ_TILE_Y_K, nwarps*warp_size); -#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) +#if defined(AMD_WMMA_AVAILABLE) constexpr vec_dot_mmq_t vec_dot = mmq_type_traits::vec_dot_mma; constexpr mmq_write_back_t write_back = mmq_write_back_mma; +#elif defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) + constexpr bool rocmfpx_dp4a_layout = type == GGML_TYPE_Q2_0_ROCMFP2 || type == GGML_TYPE_Q3_0_ROCMFPX; + constexpr vec_dot_mmq_t vec_dot = rocmfpx_dp4a_layout ? + mmq_type_traits::vec_dot_dp4a : + mmq_type_traits::vec_dot_mma; + constexpr mmq_write_back_t write_back = rocmfpx_dp4a_layout ? + mmq_write_back_dp4a : + mmq_write_back_mma; #else constexpr vec_dot_mmq_t vec_dot = mmq_type_traits::vec_dot_dp4a; constexpr mmq_write_back_t write_back = mmq_write_back_dp4a; @@ -4051,7 +4364,11 @@ static size_t mmq_get_nbytes_shared(const int mmq_x, const int mmq_y, const int const tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(type, mmq_y); const int mmq_tile_x_k = mmq_get_mma_tile_x_k(type); const size_t nbs_ids = mmq_x*sizeof(int); - const size_t nbs_x = (turing_mma_available(cc) || amd_mfma_available(cc) || amd_wmma_available(cc)) ? mmq_y*mmq_tile_x_k*sizeof(int) : txs.qs*sizeof(int) + txs.dm*sizeof(half2) + txs.sc*sizeof(int); + constexpr bool rocmfpx_type = type == GGML_TYPE_Q2_0_ROCMFP2 || type == GGML_TYPE_Q3_0_ROCMFPX; + const bool use_mma_layout = rocmfpx_type ? amd_wmma_available(cc) : + (turing_mma_available(cc) || amd_mfma_available(cc) || amd_wmma_available(cc)); + const size_t nbs_x = use_mma_layout ? mmq_y*mmq_tile_x_k*sizeof(int) : + txs.qs*sizeof(int) + txs.dm*sizeof(half2) + txs.sc*sizeof(int); const size_t nbs_y = mmq_x * (sizeof(block_q8_1_mmq)); return nbs_ids + nbs_x + GGML_PAD(nbs_y, nwarps*warp_size*sizeof(int)); } @@ -4242,6 +4559,9 @@ extern DECL_MMQ_CASE(GGML_TYPE_Q4_1); extern DECL_MMQ_CASE(GGML_TYPE_Q5_0); extern DECL_MMQ_CASE(GGML_TYPE_Q5_1); extern DECL_MMQ_CASE(GGML_TYPE_Q8_0); +extern DECL_MMQ_CASE(GGML_TYPE_Q4_0_ROCMFP4_FAST); +extern DECL_MMQ_CASE(GGML_TYPE_Q2_0_ROCMFP2); +extern DECL_MMQ_CASE(GGML_TYPE_Q3_0_ROCMFPX); extern DECL_MMQ_CASE(GGML_TYPE_MXFP4); extern DECL_MMQ_CASE(GGML_TYPE_NVFP4); extern DECL_MMQ_CASE(GGML_TYPE_Q2_K); diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/generate_cu_files.py b/server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/generate_cu_files.py index 17cee853e..22a7d57f4 100755 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/generate_cu_files.py +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/generate_cu_files.py @@ -33,6 +33,7 @@ TYPES_MMQ = [ "GGML_TYPE_Q4_0", "GGML_TYPE_Q4_1", "GGML_TYPE_Q5_0", "GGML_TYPE_Q5_1", "GGML_TYPE_Q8_0", + "GGML_TYPE_Q4_0_ROCMFP4_FAST", "GGML_TYPE_Q2_0_ROCMFP2", "GGML_TYPE_Q3_0_ROCMFPX", "GGML_TYPE_Q2_K", "GGML_TYPE_Q3_K", "GGML_TYPE_Q4_K", "GGML_TYPE_Q5_K", "GGML_TYPE_Q6_K", "GGML_TYPE_IQ2_XXS", "GGML_TYPE_IQ2_XS", "GGML_TYPE_IQ2_S", "GGML_TYPE_IQ3_XXS", "GGML_TYPE_IQ3_S", "GGML_TYPE_IQ1_S", "GGML_TYPE_IQ4_NL", "GGML_TYPE_IQ4_XS", "GGML_TYPE_MXFP4", "GGML_TYPE_NVFP4" diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/mmq-instance-q2_0_rocmfp2.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/mmq-instance-q2_0_rocmfp2.cu new file mode 100644 index 000000000..6fb657de1 --- /dev/null +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/mmq-instance-q2_0_rocmfp2.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_Q2_0_ROCMFP2); diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/mmq-instance-q3_0_rocmfpx.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/mmq-instance-q3_0_rocmfpx.cu new file mode 100644 index 000000000..85338fc41 --- /dev/null +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/mmq-instance-q3_0_rocmfpx.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_Q3_0_ROCMFPX); diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/mmq-instance-q4_0_rocmfp4_fast.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/mmq-instance-q4_0_rocmfp4_fast.cu new file mode 100644 index 000000000..df63d09ae --- /dev/null +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/mmq-instance-q4_0_rocmfp4_fast.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_Q4_0_ROCMFP4_FAST);