diff --git a/common/chat-auto-parser-generator.cpp b/common/chat-auto-parser-generator.cpp index bf44091d6..3f036bb5b 100644 --- a/common/chat-auto-parser-generator.cpp +++ b/common/chat-auto-parser-generator.cpp @@ -65,7 +65,7 @@ common_chat_params peg_generator::generate_parser(const common_chat_template & data.grammar = build_grammar([&](const common_grammar_builder & builder) { foreach_function(inputs.tools, [&](const json & tool) { const auto & function = tool.at("function"); - auto schema = function.at("parameters"); + auto schema = function.contains("parameters") ? function.at("parameters") : json::object(); builder.resolve_refs(schema); }); parser.build_grammar(builder, data.grammar_lazy); @@ -221,7 +221,7 @@ common_peg_parser analyze_tools::build_tool_parser_tag_json(parser_build_context foreach_function(inputs.tools, [&](const json & tool) { const auto & func = tool.at("function"); std::string name = func.at("name"); - const auto & schema = func.at("parameters"); + const auto & schema = func.contains("parameters") ? func.at("parameters") : json::object(); // Build call_id parser based on position (if supported) common_peg_parser call_id_section = p.eps(); @@ -282,19 +282,11 @@ common_peg_parser analyze_tools::build_tool_parser_tag_tagged(parser_build_conte common_peg_parser tool_choice = p.choice(); foreach_function(inputs.tools, [&](const json & tool) { - const auto & func = tool.at("function"); - std::string name = func.at("name"); - const auto & params = func.at("parameters"); - - if (!params.contains("properties") || !params.at("properties").is_object()) { - return; - } - - const auto & properties = params.at("properties"); + const auto & func = tool.at("function"); + std::string name = func.at("name"); + const auto & params = func.contains("parameters") ? func.at("parameters") : json::object(); + const auto & properties = params.contains("properties") ? params.at("properties") : json::object(); std::set required; - if (params.contains("required") && params.at("required").is_array()) { - params.at("required").get_to(required); - } // Build parser for each argument, separating required and optional std::vector required_parsers; @@ -311,17 +303,18 @@ common_peg_parser analyze_tools::build_tool_parser_tag_tagged(parser_build_conte } } - auto arg = p.tool_arg( - p.tool_arg_open(arguments.name_prefix + p.tool_arg_name(p.literal(param_name)) + - arguments.name_suffix) + - arguments.value_prefix + - (type == "string" ? p.tool_arg_string_value(p.schema(p.until(arguments.value_suffix), - "tool-" + name + "-arg-" + param_name + "-schema", - param_schema, true)) : - p.tool_arg_json_value(p.schema( - p.json(), "tool-" + name + "-arg-" + param_name + "-schema", param_schema, false)) + - p.space()) + - p.tool_arg_close(p.literal(arguments.value_suffix))); + auto arg = + p.tool_arg(p.tool_arg_open(arguments.name_prefix + p.tool_arg_name(p.literal(param_name)) + + arguments.name_suffix) + + arguments.value_prefix + + (type == "string" ? + p.tool_arg_string_value(p.schema(p.until(arguments.value_suffix), + "tool-" + name + "-arg-" + param_name + "-schema", + param_schema, true)) : + p.tool_arg_json_value(p.schema( + p.json(), "tool-" + name + "-arg-" + param_name + "-schema", param_schema, false)) + + p.space()) + + p.tool_arg_close(p.literal(arguments.value_suffix))); auto named_arg = p.rule("tool-" + name + "-arg-" + param_name, arg); if (is_required) { diff --git a/common/common.cpp b/common/common.cpp index 4ea4eaea1..f807f3b04 100644 --- a/common/common.cpp +++ b/common/common.cpp @@ -663,14 +663,53 @@ bool string_parse_kv_override(const char * data, std::vector= start_char && c <= end_char) { + matched = true; + break; + } + class_start += 3; + } else { + if (*class_start == c) { + matched = true; + break; + } + class_start++; + } + } + + return negated ? !matched : matched; +} + +// simple glob: * matches non-/ chars, ** matches anything including /, [] matches character class static inline bool glob_match(const char * pattern, const char * str) { if (*pattern == '\0') { return *str == '\0'; } if (pattern[0] == '*' && pattern[1] == '*') { const char * p = pattern + 2; - if (*p == '/') p++; if (glob_match(p, str)) return true; if (*str != '\0') return glob_match(pattern, str + 1); return false; @@ -685,6 +724,26 @@ static inline bool glob_match(const char * pattern, const char * str) { if (*pattern == '?' && *str != '\0' && *str != '/') { return glob_match(pattern + 1, str + 1); } + if (*pattern == '[') { + const char * class_end = pattern + 1; + // If first character after '[' is ']' or '-', treat it as literal + if (*class_end == ']' || *class_end == '-') { + class_end++; + } + while (*class_end != '\0' && *class_end != ']') { + class_end++; + } + if (*class_end == ']') { + if (*str == '\0') return false; + bool matched = glob_class_match(*str, pattern + 1, class_end); + return matched && glob_match(class_end + 1, str + 1); + } else { + if (*str == '[') { + return glob_match(pattern + 1, str + 1); + } + return false; + } + } if (*pattern == *str) { return glob_match(pattern + 1, str + 1); } diff --git a/convert_hf_to_gguf.py b/convert_hf_to_gguf.py index bcf98cfae..51f0d1ab1 100755 --- a/convert_hf_to_gguf.py +++ b/convert_hf_to_gguf.py @@ -31,10 +31,10 @@ import gguf from gguf.vocab import MistralTokenizerType, MistralVocab try: - from mistral_common.tokens.tokenizers.base import TokenizerVersion # type: ignore[import-not-found] - from mistral_common.tokens.tokenizers.multimodal import DATASET_MEAN as _MISTRAL_COMMON_DATASET_MEAN, DATASET_STD as _MISTRAL_COMMON_DATASET_STD # type: ignore[import-not-found] - from mistral_common.tokens.tokenizers.tekken import Tekkenizer # type: ignore[import-not-found] - from mistral_common.tokens.tokenizers.sentencepiece import ( # type: ignore[import-not-found] + from mistral_common.tokens.tokenizers.base import TokenizerVersion # type: ignore[import-not-found, ty:unresolved-import] + from mistral_common.tokens.tokenizers.multimodal import DATASET_MEAN as _MISTRAL_COMMON_DATASET_MEAN, DATASET_STD as _MISTRAL_COMMON_DATASET_STD # type: ignore[import-not-found, ty:unresolved-import] + from mistral_common.tokens.tokenizers.tekken import Tekkenizer # type: ignore[import-not-found, ty:unresolved-import] + from mistral_common.tokens.tokenizers.sentencepiece import ( # type: ignore[import-not-found, ty:unresolved-import] SentencePieceTokenizer, ) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index a389767fb..6d4ea5d07 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -2351,7 +2351,8 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor * static_assert(MMVQ_MAX_BATCH_SIZE == MMVF_MAX_BATCH_SIZE); if (ne2 <= MMVQ_MAX_BATCH_SIZE) { if (ggml_is_quantized(src0->type)) { - if (ne2 <= MMVQ_MMID_MAX_BATCH_SIZE) { + const int mmvq_mmid_max = get_mmvq_mmid_max_batch(src0->type, cc); + if (ne2 <= mmvq_mmid_max) { ggml_cuda_mul_mat_vec_q(ctx, src0, src1, ids, dst); return; } @@ -2958,14 +2959,18 @@ static bool ggml_cuda_graph_check_compability(ggml_cgraph * cgraph) { } // [TAG_MUL_MAT_ID_CUDA_GRAPHS] - if (node->op == GGML_OP_MUL_MAT_ID && (!ggml_is_quantized(node->src[0]->type) || node->ne[2] > MMVQ_MMID_MAX_BATCH_SIZE)) { - // under these conditions, the mul_mat_id operation will need to synchronize the stream, so we cannot use CUDA graphs - // TODO: figure out a way to enable for larger batch sizes, without hurting performance - // ref: https://github.com/ggml-org/llama.cpp/pull/18958 - use_cuda_graph = false; + if (node->op == GGML_OP_MUL_MAT_ID) { + const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; + const int mmvq_mmid_max = get_mmvq_mmid_max_batch(node->src[0]->type, cc); + if (!ggml_is_quantized(node->src[0]->type) || node->ne[2] > mmvq_mmid_max) { + // under these conditions, the mul_mat_id operation will need to synchronize the stream, so we cannot use CUDA graphs + // TODO: figure out a way to enable for larger batch sizes, without hurting performance + // ref: https://github.com/ggml-org/llama.cpp/pull/18958 + use_cuda_graph = false; #ifndef NDEBUG - GGML_LOG_DEBUG("%s: disabling CUDA graphs due to unsupported node type\n", __func__); + GGML_LOG_DEBUG("%s: disabling CUDA graphs due to unsupported node type\n", __func__); #endif + } } if (node->op == GGML_OP_ADD && node->src[1] && node->src[1]->ne[1] > 1) { diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 66bd8beea..8d80d1dd9 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -97,6 +97,194 @@ static __host__ mmvq_parameter_table_id get_device_table_id(int cc) { return MMVQ_PARAMETERS_GENERIC; } +// Per-architecture maximum batch size for which MMVQ should be used for MUL_MAT_ID. +// Returns a value <= MMVQ_MAX_BATCH_SIZE. Default is MMVQ_MAX_BATCH_SIZE. +// Check https://github.com/ggml-org/llama.cpp/pull/20905#issuecomment-4145835627 for details + +static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_pascal_older(ggml_type type) { + switch (type) { + case GGML_TYPE_IQ1_S: return 6; + case GGML_TYPE_IQ1_M: return 6; + case GGML_TYPE_IQ2_S: return 4; + case GGML_TYPE_IQ2_XS: return 5; + case GGML_TYPE_IQ2_XXS: return 5; + case GGML_TYPE_IQ3_S: return 4; + case GGML_TYPE_IQ3_XXS: return 4; + case GGML_TYPE_IQ4_NL: return 6; + case GGML_TYPE_IQ4_XS: return 5; + case GGML_TYPE_MXFP4: return 4; + case GGML_TYPE_Q2_K: return 4; + case GGML_TYPE_Q3_K: return 4; + case GGML_TYPE_Q4_0: return 6; + case GGML_TYPE_Q4_1: return 6; + case GGML_TYPE_Q4_K: return 5; + case GGML_TYPE_Q5_0: return 6; + case GGML_TYPE_Q5_1: return 6; + case GGML_TYPE_Q5_K: return 5; + case GGML_TYPE_Q6_K: return 4; + case GGML_TYPE_Q8_0: return 4; + default: return MMVQ_MAX_BATCH_SIZE; + } +} + +static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_turing_plus(ggml_type type) { + switch (type) { + case GGML_TYPE_IQ2_S: return 7; + case GGML_TYPE_IQ3_S: return 6; + case GGML_TYPE_IQ3_XXS: return 7; + case GGML_TYPE_MXFP4: return 7; + case GGML_TYPE_Q2_K: return 7; + case GGML_TYPE_Q3_K: return 5; + default: return MMVQ_MAX_BATCH_SIZE; + } +} + +static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_gcn(ggml_type type) { + switch (type) { + case GGML_TYPE_IQ1_S: return 5; + case GGML_TYPE_IQ1_M: return 5; + case GGML_TYPE_IQ2_S: return 4; + case GGML_TYPE_IQ2_XS: return 4; + case GGML_TYPE_IQ2_XXS: return 4; + case GGML_TYPE_IQ3_S: return 4; + case GGML_TYPE_IQ3_XXS: return 4; + case GGML_TYPE_IQ4_NL: return 6; + case GGML_TYPE_IQ4_XS: return 4; + case GGML_TYPE_Q2_K: return 4; + case GGML_TYPE_Q3_K: return 4; + case GGML_TYPE_Q4_0: return 5; + case GGML_TYPE_Q4_1: return 5; + case GGML_TYPE_Q4_K: return 4; + case GGML_TYPE_Q5_K: return 4; + case GGML_TYPE_Q6_K: return 4; + case GGML_TYPE_Q8_0: return 4; + default: return MMVQ_MAX_BATCH_SIZE; + } +} + +static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_cdna(ggml_type type) { + switch (type) { + case GGML_TYPE_IQ2_S: return 5; + case GGML_TYPE_IQ2_XS: return 5; + case GGML_TYPE_IQ2_XXS: return 5; + case GGML_TYPE_IQ3_S: return 4; + case GGML_TYPE_IQ3_XXS: return 5; + default: return MMVQ_MAX_BATCH_SIZE; + } +} + +static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_rdna1_rdna2(ggml_type type) { + switch (type) { + case GGML_TYPE_IQ2_S: return 4; + case GGML_TYPE_IQ2_XS: return 4; + case GGML_TYPE_IQ2_XXS: return 4; + case GGML_TYPE_IQ3_S: return 4; + case GGML_TYPE_IQ3_XXS: return 4; + case GGML_TYPE_Q2_K: return 7; + case GGML_TYPE_Q3_K: return 4; + case GGML_TYPE_Q4_K: return 5; + case GGML_TYPE_Q5_K: return 6; + case GGML_TYPE_Q6_K: return 5; + default: return MMVQ_MAX_BATCH_SIZE; + } +} + +static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_rdna3(ggml_type type) { + switch (type) { + case GGML_TYPE_IQ1_S: return 6; + case GGML_TYPE_IQ1_M: return 6; + case GGML_TYPE_IQ2_S: return 4; + case GGML_TYPE_IQ2_XS: return 4; + case GGML_TYPE_IQ2_XXS: return 4; + case GGML_TYPE_IQ3_S: return 4; + case GGML_TYPE_IQ3_XXS: return 4; + case GGML_TYPE_IQ4_NL: return 6; + case GGML_TYPE_IQ4_XS: return 6; + case GGML_TYPE_Q4_K: return 4; + case GGML_TYPE_Q5_K: return 4; + case GGML_TYPE_Q6_K: return 4; + default: return MMVQ_MAX_BATCH_SIZE; + } +} + +static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_rdna4(ggml_type type) { + switch (type) { + case GGML_TYPE_IQ1_S: return 7; + case GGML_TYPE_IQ1_M: return 7; + case GGML_TYPE_IQ2_S: return 4; + case GGML_TYPE_IQ2_XS: return 4; + case GGML_TYPE_IQ2_XXS: return 4; + case GGML_TYPE_IQ3_S: return 4; + case GGML_TYPE_IQ3_XXS: return 4; + case GGML_TYPE_IQ4_NL: return 7; + case GGML_TYPE_IQ4_XS: return 5; + case GGML_TYPE_MXFP4: return 5; + case GGML_TYPE_Q3_K: return 4; + case GGML_TYPE_Q4_0: return 7; + case GGML_TYPE_Q4_1: return 7; + case GGML_TYPE_Q4_K: return 4; + case GGML_TYPE_Q5_0: return 7; + case GGML_TYPE_Q5_1: return 7; + case GGML_TYPE_Q5_K: return 5; + case GGML_TYPE_Q6_K: return 5; + case GGML_TYPE_Q8_0: return 7; + default: return MMVQ_MAX_BATCH_SIZE; + } +} + +// Host function: returns the max batch size for the current arch+type at runtime. +int get_mmvq_mmid_max_batch(ggml_type type, int cc) { + // NVIDIA: Volta, Ada Lovelace, and Blackwell always use MMVQ for MUL_MAT_ID. + if (cc == GGML_CUDA_CC_VOLTA || cc >= GGML_CUDA_CC_ADA_LOVELACE) { + return MMVQ_MAX_BATCH_SIZE; + } + if (cc >= GGML_CUDA_CC_TURING) { + return get_mmvq_mmid_max_batch_turing_plus(type); + } + if (GGML_CUDA_CC_IS_NVIDIA(cc)) { + return get_mmvq_mmid_max_batch_pascal_older(type); + } + // AMD + if (GGML_CUDA_CC_IS_RDNA4(cc)) { + return get_mmvq_mmid_max_batch_rdna4(type); + } + if (GGML_CUDA_CC_IS_RDNA3(cc)) { + return get_mmvq_mmid_max_batch_rdna3(type); + } + if (GGML_CUDA_CC_IS_RDNA1(cc) || GGML_CUDA_CC_IS_RDNA2(cc)) { + return get_mmvq_mmid_max_batch_rdna1_rdna2(type); + } + if (GGML_CUDA_CC_IS_CDNA(cc)) { + return get_mmvq_mmid_max_batch_cdna(type); + } + if (GGML_CUDA_CC_IS_GCN(cc)) { + return get_mmvq_mmid_max_batch_gcn(type); + } + return MMVQ_MAX_BATCH_SIZE; +} + +// Device constexpr: returns the max batch size for the current arch+type at compile time. +template +static constexpr __device__ int get_mmvq_mmid_max_batch_for_device() { +#if defined(RDNA4) + return get_mmvq_mmid_max_batch_rdna4(type); +#elif defined(RDNA3) + return get_mmvq_mmid_max_batch_rdna3(type); +#elif defined(RDNA2) || defined(RDNA1) + return get_mmvq_mmid_max_batch_rdna1_rdna2(type); +#elif defined(CDNA) + return get_mmvq_mmid_max_batch_cdna(type); +#elif defined(GCN) + return get_mmvq_mmid_max_batch_gcn(type); +#elif defined(__CUDA_ARCH__) && (__CUDA_ARCH__ == GGML_CUDA_CC_VOLTA || __CUDA_ARCH__ >= GGML_CUDA_CC_ADA_LOVELACE) + return MMVQ_MAX_BATCH_SIZE; +#elif defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= GGML_CUDA_CC_TURING + return get_mmvq_mmid_max_batch_turing_plus(type); +#else + return get_mmvq_mmid_max_batch_pascal_older(type); +#endif +} + static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_dst, mmvq_parameter_table_id table_id) { if (table_id == MMVQ_PARAMETERS_GENERIC) { switch (ncols_dst) { @@ -195,7 +383,7 @@ static constexpr __host__ __device__ int calc_rows_per_block(int ncols_dst, int return 1; } -template +template __launch_bounds__(calc_nwarps(type, ncols_dst, get_device_table_id())*ggml_cuda_get_physical_warp_size(), 1) static __global__ void mul_mat_vec_q( const void * __restrict__ vx, const void * __restrict__ vy, const int32_t * __restrict__ ids, const ggml_cuda_mm_fusion_args_device fusion, float * __restrict__ dst, @@ -222,22 +410,13 @@ static __global__ void mul_mat_vec_q( const uint32_t channel_dst = blockIdx.y; - uint32_t token_idx = 0; uint32_t channel_x; uint32_t channel_y; uint32_t sample_dst; - if constexpr (is_multi_token_id) { - // Multi-token MUL_MAT_ID path, adding these in the normal path causes a perf regression for n_tokens=1 case - token_idx = blockIdx.z; - channel_x = ids[channel_dst + token_idx * ids_stride]; - channel_y = fastmodulo(channel_dst, nchannels_y); - sample_dst = 0; - } else { - channel_x = ncols_dst == 1 && ids ? ids[channel_dst] : fastdiv(channel_dst, channel_ratio); - channel_y = ncols_dst == 1 && ids ? fastmodulo(channel_dst, nchannels_y) : channel_dst; - sample_dst = blockIdx.z; - } + channel_x = ncols_dst == 1 && ids ? ids[channel_dst] : fastdiv(channel_dst, channel_ratio); + channel_y = ncols_dst == 1 && ids ? fastmodulo(channel_dst, nchannels_y) : channel_dst; + sample_dst = blockIdx.z; const uint32_t sample_x = fastdiv(sample_dst, sample_ratio); const uint32_t sample_y = sample_dst; @@ -294,9 +473,6 @@ static __global__ void mul_mat_vec_q( float tmp_gate[ncols_dst][rows_per_cuda_block] = {{0.0f}}; const block_q8_1 * y = ((const block_q8_1 *) vy) + sample_y*stride_sample_y + channel_y*stride_channel_y; - if constexpr (is_multi_token_id) { - y += token_idx*stride_col_y; - } const int kbx_offset = sample_x*stride_sample_x + channel_x*stride_channel_x + row0*stride_row_x; for (int kbx = tid / (qi/vdr); kbx < blocks_per_row_x; kbx += blocks_per_iter) { @@ -350,10 +526,6 @@ static __global__ void mul_mat_vec_q( dst += sample_dst*stride_sample_dst + channel_dst*stride_channel_dst + row0; - if constexpr (is_multi_token_id) { - dst += token_idx*stride_col_dst; - } - // sum up partial sums and write back result #pragma unroll for (int j = 0; j < ncols_dst; ++j) { @@ -413,6 +585,69 @@ static __global__ void mul_mat_vec_q( } } +// Dedicated MoE multi-token kernel. +// Grid: (ceil(nrows_x / c_rows_per_block), nchannels_dst) +// Block: (warp_size, ncols_dst) - each warp handles one token independently. +// No shared memory reduction needed since each warp works alone. +template +__launch_bounds__(get_mmvq_mmid_max_batch_for_device()*ggml_cuda_get_physical_warp_size(), 1) +static __global__ void mul_mat_vec_q_moe( + const void * __restrict__ vx, const void * __restrict__ vy, const int32_t * __restrict__ ids, + float * __restrict__ dst, + const uint32_t ncols_x, const uint3 nchannels_y, const uint32_t nrows_x, + const uint32_t stride_row_x, const uint32_t stride_col_y, const uint32_t stride_col_dst, + const uint32_t stride_channel_x, const uint32_t stride_channel_y, const uint32_t stride_channel_dst, + const uint32_t ncols_dst, const uint32_t ids_stride) { + + constexpr int qk = ggml_cuda_type_traits::qk; + constexpr int qi = ggml_cuda_type_traits::qi; + constexpr int vdr = get_vdr_mmvq(type); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + + constexpr vec_dot_q_cuda_t vec_dot_q_cuda = get_vec_dot_q_cuda(type); + + const uint32_t token_idx = threadIdx.y; + const int row0 = c_rows_per_block*blockIdx.x; + const int blocks_per_row_x = ncols_x / qk; + constexpr int blocks_per_iter = vdr * warp_size / qi; + + const uint32_t channel_dst = blockIdx.y; + + if (token_idx >= ncols_dst) { + return; + } + + const uint32_t channel_x = ids[channel_dst + token_idx * ids_stride]; + const uint32_t channel_y = fastmodulo(channel_dst, nchannels_y); + + const block_q8_1 * y = ((const block_q8_1 *) vy) + channel_y*stride_channel_y + token_idx*stride_col_y; + const int kbx_offset = channel_x*stride_channel_x + row0*stride_row_x; + + // partial sum for each thread + float tmp[c_rows_per_block] = {0.0f}; + + for (int kbx = threadIdx.x / (qi/vdr); kbx < blocks_per_row_x; kbx += blocks_per_iter) { + const int kby = kbx * (qk/QK8_1); + const int kqs = vdr * (threadIdx.x % (qi/vdr)); + +#pragma unroll + for (int i = 0; i < c_rows_per_block; ++i) { + tmp[i] += vec_dot_q_cuda(vx, &y[kby], kbx_offset + i*stride_row_x + kbx, kqs); + } + } + + // Warp-level reduction only - no shared memory needed +#pragma unroll + for (int i = 0; i < c_rows_per_block; ++i) { + tmp[i] = warp_reduce_sum(tmp[i]); + } + + // Write results + if (threadIdx.x < c_rows_per_block && (c_rows_per_block == 1 || uint32_t(row0 + threadIdx.x) < nrows_x)) { + dst[channel_dst*stride_channel_dst + token_idx*stride_col_dst + row0 + threadIdx.x] = tmp[threadIdx.x]; + } +} + template static std::pair calc_launch_params( const int ncols_dst, const int nrows_x, const int nchannels_dst, const int nsamples_or_ntokens, @@ -425,7 +660,7 @@ static std::pair calc_launch_params( return {block_nums, block_dims}; } -template +template static void mul_mat_vec_q_switch_fusion( const void * vx, const void * vy, const int32_t * ids, const ggml_cuda_mm_fusion_args_device fusion, float * dst, const uint32_t ncols_x, const uint3 nchannels_y, const uint32_t stride_row_x, const uint32_t stride_col_y, @@ -438,7 +673,7 @@ static void mul_mat_vec_q_switch_fusion( const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr; if constexpr (c_ncols_dst == 1) { if (has_fusion) { - mul_mat_vec_q<<>> + mul_mat_vec_q<<>> (vx, vy, ids, fusion, dst, ncols_x, nchannels_y, stride_row_x, stride_col_y, stride_col_dst, channel_ratio, stride_channel_x, stride_channel_y, stride_channel_dst, sample_ratio, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride); @@ -448,12 +683,33 @@ static void mul_mat_vec_q_switch_fusion( GGML_ASSERT(!has_fusion && "fusion only supported for ncols_dst=1"); - mul_mat_vec_q<<>> + mul_mat_vec_q<<>> (vx, vy, ids, fusion, dst, ncols_x, nchannels_y, stride_row_x, stride_col_y, stride_col_dst, channel_ratio, stride_channel_x, stride_channel_y, stride_channel_dst, sample_ratio, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride); } +template +static void mul_mat_vec_q_moe_launch( + const void * vx, const void * vy, const int32_t * ids, float * dst, + const uint32_t ncols_x, const uint3 nchannels_y, const uint32_t nrows_x, + const uint32_t stride_row_x, const uint32_t stride_col_y, const uint32_t stride_col_dst, + const uint32_t stride_channel_x, const uint32_t stride_channel_y, const uint32_t stride_channel_dst, + const uint32_t ncols_dst, const uint32_t ids_stride, + const int warp_size, const int nchannels_dst, cudaStream_t stream) { + + constexpr int rows_per_block = 2; // 2 gives best perf based on tuning + const int64_t nblocks_rows = (nrows_x + rows_per_block - 1) / rows_per_block; + const dim3 block_nums(nblocks_rows, nchannels_dst); + const dim3 block_dims(warp_size, ncols_dst); + + mul_mat_vec_q_moe<<>>( + vx, vy, ids, dst, ncols_x, nchannels_y, nrows_x, + stride_row_x, stride_col_y, stride_col_dst, + stride_channel_x, stride_channel_y, stride_channel_dst, + ncols_dst, ids_stride); +} + template static void mul_mat_vec_q_switch_ncols_dst( const void * vx, const void * vy, const int32_t * ids, const ggml_cuda_mm_fusion_args_device fusion, float * dst, @@ -472,20 +728,62 @@ static void mul_mat_vec_q_switch_ncols_dst( const uint3 sample_ratio_fd = init_fastdiv_values(nsamples_dst / nsamples_x); const int device = ggml_cuda_get_device(); + const int cc = ggml_cuda_info().devices[device].cc; const int warp_size = ggml_cuda_info().devices[device].warp_size; - const mmvq_parameter_table_id table_id = get_device_table_id(ggml_cuda_info().devices[device].cc); + const mmvq_parameter_table_id table_id = get_device_table_id(cc); const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr; const bool has_ids = ids != nullptr; + const auto should_use_small_k = [&](int c_ncols_dst) { + // When K is small, increase rows_per_block to match nwarps so each warp has more work to do + // Trigger when the full thread block covers all K blocks in a single loop iteration and few threads remain idle. + constexpr int qk = ggml_cuda_type_traits::qk; + constexpr int qi = ggml_cuda_type_traits::qi; + constexpr int vdr = get_vdr_mmvq(type); + const int blocks_per_row_x = ncols_x / qk; + const int blocks_per_iter_1warp = vdr * warp_size / qi; + const int nwarps = calc_nwarps(type, c_ncols_dst, table_id); + bool use = nwarps > 1 && blocks_per_row_x < nwarps * blocks_per_iter_1warp; + + constexpr std::array iq_slow_turing = { + GGML_TYPE_IQ3_XXS, + GGML_TYPE_IQ3_S, + }; + constexpr std::array iq_slow_other = { + GGML_TYPE_IQ1_S, GGML_TYPE_IQ1_M, GGML_TYPE_IQ2_XXS, GGML_TYPE_IQ2_XS, + GGML_TYPE_IQ2_S, GGML_TYPE_IQ3_XXS, GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_XS, + }; + constexpr std::array slow_pascal = { + GGML_TYPE_IQ3_S, + GGML_TYPE_Q2_K, + GGML_TYPE_Q3_K, + }; + + const bool is_nvidia_turing_plus = GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_TURING; + const bool is_nvidia_pascal_older = GGML_CUDA_CC_IS_NVIDIA(cc) && cc < GGML_CUDA_CC_VOLTA; + + if (is_nvidia_turing_plus) { + if (ncols_dst == 1 && + std::find(iq_slow_turing.begin(), iq_slow_turing.end(), type) != iq_slow_turing.end()) { + use = false; + } + } else if ((ncols_dst == 1 && std::find(iq_slow_other.begin(), iq_slow_other.end(), type) != iq_slow_other.end()) || + (is_nvidia_pascal_older && std::find(slow_pascal.begin(), slow_pascal.end(), type) != slow_pascal.end()) || + GGML_CUDA_CC_IS_RDNA(cc)) { + use = false; + } + + return use; + }; + if (has_ids && ncols_dst > 1) { - // Multi-token MUL_MAT_ID path only - single-token goes through regular path below - constexpr int c_ncols_dst = 1; - std::pair dims = calc_launch_params(c_ncols_dst, nrows_x, nchannels_dst, ncols_dst, warp_size, table_id); - mul_mat_vec_q_switch_fusion(vx, vy, ids, fusion, dst, ncols_x, nchannels_y_fd, stride_row_x, stride_col_y, stride_col_dst, - channel_ratio_fd, stride_channel_x, stride_channel_y, stride_channel_dst, - sample_ratio_fd, stride_sample_x, stride_sample_y, stride_sample_dst, - dims.first, dims.second, 0, ids_stride, stream); + // Multi-token MUL_MAT_ID path - dedicated MoE kernel + mul_mat_vec_q_moe_launch( + vx, vy, ids, dst, ncols_x, nchannels_y_fd, nrows_x, + stride_row_x, stride_col_y, stride_col_dst, + stride_channel_x, stride_channel_y, stride_channel_dst, + ncols_dst, ids_stride, warp_size, nchannels_dst, stream); return; } @@ -493,31 +791,24 @@ static void mul_mat_vec_q_switch_ncols_dst( case 1: { constexpr int c_ncols_dst = 1; - // When K is small, increase rows_per_block to match nwarps so each warp has more work to do - // Trigger when the full thread block covers all K blocks in a single loop iteration and few threads remain idle. - constexpr int qk = ggml_cuda_type_traits::qk; - constexpr int qi = ggml_cuda_type_traits::qi; - constexpr int vdr = get_vdr_mmvq(type); - const int blocks_per_row_x = ncols_x / qk; - const int blocks_per_iter_1warp = vdr * warp_size / qi; - const int nwarps = calc_nwarps(type, c_ncols_dst, table_id); - const bool use_small_k = nwarps > 1 && blocks_per_row_x < nwarps * blocks_per_iter_1warp; + bool use_small_k = should_use_small_k(c_ncols_dst); + if (use_small_k) { - std::pair dims = calc_launch_params(c_ncols_dst, nrows_x, nchannels_dst, nsamples_dst, - warp_size, table_id, true); - mul_mat_vec_q_switch_fusion( + std::pair dims = calc_launch_params(c_ncols_dst, nrows_x, nchannels_dst, + nsamples_dst, warp_size, table_id, true); + mul_mat_vec_q_switch_fusion( vx, vy, ids, fusion, dst, ncols_x, nchannels_y_fd, stride_row_x, stride_col_y, stride_col_dst, - channel_ratio_fd, stride_channel_x, stride_channel_y, stride_channel_dst, - sample_ratio_fd, stride_sample_x, stride_sample_y, stride_sample_dst, - dims.first, dims.second, 0, ids_stride, stream); + channel_ratio_fd, stride_channel_x, stride_channel_y, stride_channel_dst, sample_ratio_fd, + stride_sample_x, stride_sample_y, stride_sample_dst, dims.first, dims.second, 0, ids_stride, + stream); } else { - std::pair dims = calc_launch_params(c_ncols_dst, nrows_x, nchannels_dst, nsamples_dst, - warp_size, table_id); + std::pair dims = calc_launch_params(c_ncols_dst, nrows_x, nchannels_dst, + nsamples_dst, warp_size, table_id); mul_mat_vec_q_switch_fusion( vx, vy, ids, fusion, dst, ncols_x, nchannels_y_fd, stride_row_x, stride_col_y, stride_col_dst, - channel_ratio_fd, stride_channel_x, stride_channel_y, stride_channel_dst, - sample_ratio_fd, stride_sample_x, stride_sample_y, stride_sample_dst, - dims.first, dims.second, 0, ids_stride, stream); + channel_ratio_fd, stride_channel_x, stride_channel_y, stride_channel_dst, sample_ratio_fd, + stride_sample_x, stride_sample_y, stride_sample_dst, dims.first, dims.second, 0, ids_stride, + stream); } } break; case 2: { diff --git a/ggml/src/ggml-cuda/mmvq.cuh b/ggml/src/ggml-cuda/mmvq.cuh index 8a154631f..6bf0a8e86 100644 --- a/ggml/src/ggml-cuda/mmvq.cuh +++ b/ggml/src/ggml-cuda/mmvq.cuh @@ -1,7 +1,10 @@ #include "common.cuh" #define MMVQ_MAX_BATCH_SIZE 8 // Max. batch size for which to use MMVQ kernels. -#define MMVQ_MMID_MAX_BATCH_SIZE 4 // Max. batch size for which to use MMVQ kernels for MUL_MAT_ID + +// Returns the maximum batch size for which MMVQ should be used for MUL_MAT_ID, +// based on the quantization type and GPU architecture (compute capability). +int get_mmvq_mmid_max_batch(ggml_type type, int cc); void ggml_cuda_mul_mat_vec_q(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst, const ggml_cuda_mm_fusion_args_host * fusion = nullptr); diff --git a/gguf-py/gguf/vocab.py b/gguf-py/gguf/vocab.py index e4ab5e1e4..5cd729dfa 100644 --- a/gguf-py/gguf/vocab.py +++ b/gguf-py/gguf/vocab.py @@ -14,12 +14,12 @@ except ImportError: SentencePieceProcessor: Any = None try: - from mistral_common.tokens.tokenizers.mistral import MistralTokenizer # type: ignore[import-not-found] - from mistral_common.tokens.tokenizers.tekken import Tekkenizer # type: ignore[import-not-found] - from mistral_common.tokens.tokenizers.utils import ( # type: ignore[import-not-found] + from mistral_common.tokens.tokenizers.mistral import MistralTokenizer # type: ignore[import-not-found, ty:unresolved-import] + from mistral_common.tokens.tokenizers.tekken import Tekkenizer # type: ignore[import-not-found, ty:unresolved-import] + from mistral_common.tokens.tokenizers.utils import ( # type: ignore[import-not-found, ty:unresolved-import] _filter_valid_tokenizer_files, ) - from mistral_common.tokens.tokenizers.sentencepiece import ( # type: ignore[import-not-found] + from mistral_common.tokens.tokenizers.sentencepiece import ( # type: ignore[import-not-found, ty:unresolved-import] SentencePieceTokenizer, ) except ImportError: @@ -32,7 +32,7 @@ else: _mistral_common_installed = True try: - from mistral_common.tokens.tokenizers.utils import ( # type: ignore[import-not-found] + from mistral_common.tokens.tokenizers.utils import ( # type: ignore[import-not-found, ty:unresolved-import] get_one_valid_tokenizer_file, ) except ImportError: diff --git a/src/llama-arch.cpp b/src/llama-arch.cpp index a90518763..bd6d28e5f 100644 --- a/src/llama-arch.cpp +++ b/src/llama-arch.cpp @@ -557,6 +557,8 @@ static std::set llm_get_tensor_names(llm_arch arch) { LLM_TENSOR_OUTPUT_NORM, LLM_TENSOR_OUTPUT, LLM_TENSOR_ROPE_FREQS, + LLM_TENSOR_ROPE_FACTORS_LONG, + LLM_TENSOR_ROPE_FACTORS_SHORT, LLM_TENSOR_ATTN_NORM, LLM_TENSOR_ATTN_Q, LLM_TENSOR_ATTN_K, diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp index 13212b10d..a8c2675a3 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp @@ -1162,6 +1162,12 @@ struct ggml_tensor * llama_model_loader::create_tensor( if (overrides->buft == ggml_backend_cpu_buffer_type()) { // when overriding to a CPU buffer, consider the extra buffer types buft = select_weight_buft(hparams, t_meta, op, buft_list_cpu); + if (use_mmap) { + static std::once_flag once; + std::call_once(once, [] { + LLAMA_LOG_WARN("llama_model_loader: tensor overrides to CPU are used with mmap enabled - consider using --no-mmap for better performance\n"); + }); + } } else { buft = overrides->buft; } diff --git a/tools/server/public/index.html.gz b/tools/server/public/index.html.gz index 5e7326a58..1a69dde32 100644 Binary files a/tools/server/public/index.html.gz and b/tools/server/public/index.html.gz differ diff --git a/tools/server/server-cors-proxy.h b/tools/server/server-cors-proxy.h index c412d4c25..ffdd073b9 100644 --- a/tools/server/server-cors-proxy.h +++ b/tools/server/server-cors-proxy.h @@ -32,13 +32,22 @@ static server_http_res_ptr proxy_request(const server_http_req & req, std::strin SRV_INF("proxying %s request to %s://%s:%i%s\n", method.c_str(), parsed_url.scheme.c_str(), parsed_url.host.c_str(), parsed_url.port, parsed_url.path.c_str()); + std::map headers; + for (auto [key, value] : req.headers) { + auto new_key = key; + if (string_starts_with(new_key, "X-Proxy-Header-")) { + string_replace_all(new_key, "X-Proxy-Header-", ""); + } + headers[new_key] = value; + } + auto proxy = std::make_unique( method, parsed_url.scheme, parsed_url.host, parsed_url.port, parsed_url.path, - req.headers, + headers, req.body, req.should_stop, 600, // timeout_read (default to 10 minutes) diff --git a/tools/server/server-http.h b/tools/server/server-http.h index 3621064cd..f8a174c44 100644 --- a/tools/server/server-http.h +++ b/tools/server/server-http.h @@ -35,7 +35,7 @@ using server_http_res_ptr = std::unique_ptr; struct server_http_req { std::map params; // path_params + query_params - std::map headers; // reserved for future use + std::map headers; // used by MCP proxy std::string path; std::string query_string; // query parameters string (e.g. "action=save") std::string body; diff --git a/tools/server/tests/utils.py b/tools/server/tests/utils.py index 8f62210bc..a9a7e3c4f 100644 --- a/tools/server/tests/utils.py +++ b/tools/server/tests/utils.py @@ -116,7 +116,7 @@ class ServerProcess: self.server_port = int(os.environ["PORT"]) self.external_server = "DEBUG_EXTERNAL" in os.environ - def start(self, timeout_seconds: int | None = DEFAULT_HTTP_TIMEOUT) -> None: + def start(self, timeout_seconds: int = DEFAULT_HTTP_TIMEOUT) -> None: if self.external_server: print(f"[external_server]: Assuming external server running on {self.server_host}:{self.server_port}") return diff --git a/tools/server/webui/src/lib/components/app/models/ModelId.svelte b/tools/server/webui/src/lib/components/app/models/ModelId.svelte index 5fda49342..272b697c4 100644 --- a/tools/server/webui/src/lib/components/app/models/ModelId.svelte +++ b/tools/server/webui/src/lib/components/app/models/ModelId.svelte @@ -18,7 +18,8 @@ showRaw = undefined, aliases, tags, - class: className = '' + class: className = '', + ...rest }: Props = $props(); const badgeClass = @@ -36,9 +37,9 @@ {#if resolvedShowRaw} - + {:else} - + {#if showOrgName && parsed.orgName && !(aliases && aliases.length > 0)}{parsed.orgName}/{/if}{displayName} diff --git a/tools/server/webui/src/lib/components/app/models/ModelsSelector.svelte b/tools/server/webui/src/lib/components/app/models/ModelsSelector.svelte index 0ed6e5c1d..ccda78690 100644 --- a/tools/server/webui/src/lib/components/app/models/ModelsSelector.svelte +++ b/tools/server/webui/src/lib/components/app/models/ModelsSelector.svelte @@ -271,50 +271,49 @@ {#if isRouter} { - e.preventDefault(); - e.stopPropagation(); - }} - > - + {#if updating || isLoadingModel} + + {:else} + + {/if} - - + + + {#snippet child({ props })} + + {/snippet} diff --git a/tools/server/webui/src/lib/services/mcp.service.ts b/tools/server/webui/src/lib/services/mcp.service.ts index b38d18676..cf401bc30 100644 --- a/tools/server/webui/src/lib/services/mcp.service.ts +++ b/tools/server/webui/src/lib/services/mcp.service.ts @@ -39,7 +39,13 @@ import type { MCPResourceContent, MCPReadResourceResult } from '$lib/types'; -import { buildProxiedUrl, throwIfAborted, isAbortError, createBase64DataUrl } from '$lib/utils'; +import { + buildProxiedUrl, + buildProxiedHeaders, + throwIfAborted, + isAbortError, + createBase64DataUrl +} from '$lib/utils'; interface ToolResultContentItem { type: string; @@ -118,7 +124,7 @@ export class MCPService { const requestInit: RequestInit = {}; if (config.headers) { - requestInit.headers = config.headers; + requestInit.headers = buildProxiedHeaders(config.headers); } if (config.credentials) { diff --git a/tools/server/webui/src/lib/stores/chat.svelte.ts b/tools/server/webui/src/lib/stores/chat.svelte.ts index 3af91be83..5f3812ed3 100644 --- a/tools/server/webui/src/lib/stores/chat.svelte.ts +++ b/tools/server/webui/src/lib/stores/chat.svelte.ts @@ -28,6 +28,7 @@ import { filterByLeafNodeId, findDescendantMessages, findLeafNode, + findMessageById, isAbortError } from '$lib/utils'; import { @@ -416,7 +417,7 @@ class ChatStore { if (!activeConv) return false; try { const allMessages = await conversationsStore.getConversationMessages(activeConv.id); - const systemMessage = allMessages.find((m) => m.id === messageId); + const systemMessage = findMessageById(allMessages, messageId); if (!systemMessage || systemMessage.role !== MessageRole.SYSTEM) return false; const rootMessage = allMessages.find((m) => m.type === 'root' && m.parent === null); if (!rootMessage) return false; @@ -878,7 +879,7 @@ class ChatStore { const msg = conversationsStore.activeMessages[idx]; if (msg.role !== MessageRole.ASSISTANT) return; const allMessages = await conversationsStore.getConversationMessages(activeConv.id); - const parentMessage = allMessages.find((m) => m.id === msg.parent); + const parentMessage = findMessageById(allMessages, msg.parent); if (!parentMessage) return; this.setChatLoading(activeConv.id, true); this.clearChatStreaming(activeConv.id); @@ -928,7 +929,7 @@ class ChatStore { if (!activeConv) return { totalCount: 0, userMessages: 0, assistantMessages: 0, messageTypes: [] }; const allMessages = await conversationsStore.getConversationMessages(activeConv.id); - const messageToDelete = allMessages.find((m) => m.id === messageId); + const messageToDelete = findMessageById(allMessages, messageId); // For system messages, don't count descendants as they will be preserved (reparented to root) if (messageToDelete?.role === MessageRole.SYSTEM) { @@ -975,7 +976,7 @@ class ChatStore { if (!activeConv) return; try { const allMessages = await conversationsStore.getConversationMessages(activeConv.id); - const messageToDelete = allMessages.find((m) => m.id === messageId); + const messageToDelete = findMessageById(allMessages, messageId); if (!messageToDelete) return; @@ -1024,7 +1025,7 @@ class ChatStore { this.clearChatStreaming(activeConv.id); const allMessages = await conversationsStore.getConversationMessages(activeConv.id); - const dbMessage = allMessages.find((m) => m.id === messageId); + const dbMessage = findMessageById(allMessages, messageId); if (!dbMessage) { this.setChatLoading(activeConv.id, false); @@ -1280,7 +1281,10 @@ class ChatStore { let messageIdForResponse: string; - if (msg.children.length === 0) { + const dbMsg = findMessageById(allMessages, msg.id); + const hasChildren = dbMsg ? dbMsg.children.length > 0 : msg.children.length > 0; + + if (!hasChildren) { // No responses after this message — update in place instead of branching const updates: Partial = { content: newContent, diff --git a/tools/server/webui/src/lib/utils/branching.ts b/tools/server/webui/src/lib/utils/branching.ts index e60fb206f..4e117b3c2 100644 --- a/tools/server/webui/src/lib/utils/branching.ts +++ b/tools/server/webui/src/lib/utils/branching.ts @@ -17,6 +17,17 @@ import { MessageRole } from '$lib/enums'; +/** + * Finds a message by its ID in the given messages array. + */ +export function findMessageById( + messages: readonly DatabaseMessage[], + id: string | null | undefined +): DatabaseMessage | undefined { + if (!id) return undefined; + return messages.find((m) => m.id === id); +} + /** * Filters messages to get the conversation path from root to a specific leaf node. * If the leafNodeId doesn't exist, returns the path with the latest timestamp. diff --git a/tools/server/webui/src/lib/utils/cors-proxy.ts b/tools/server/webui/src/lib/utils/cors-proxy.ts index 39c368ba0..2eb10b16d 100644 --- a/tools/server/webui/src/lib/utils/cors-proxy.ts +++ b/tools/server/webui/src/lib/utils/cors-proxy.ts @@ -19,6 +19,21 @@ export function buildProxiedUrl(targetUrl: string): URL { return proxyUrl; } +/** + * Wrap original headers for proxying through the CORS proxy. This avoids issues with duplicated llama.cpp-specific and target headers when using the CORS proxy. + * @param headers - The original headers to be proxied to target + * @returns List of "wrapped" headers to be sent to the CORS proxy + */ +export function buildProxiedHeaders(headers: Record): Record { + const proxiedHeaders: Record = {}; + + for (const [key, value] of Object.entries(headers)) { + proxiedHeaders[`X-Proxy-Header-${key}`] = value; + } + + return proxiedHeaders; +} + /** * Get a proxied URL string for use in fetch requests. * @param targetUrl - The original URL to proxy diff --git a/tools/server/webui/src/lib/utils/index.ts b/tools/server/webui/src/lib/utils/index.ts index 2caaf9ac3..455d4f2c3 100644 --- a/tools/server/webui/src/lib/utils/index.ts +++ b/tools/server/webui/src/lib/utils/index.ts @@ -22,6 +22,7 @@ export { default as autoResizeTextarea } from './autoresize-textarea'; // Branching utilities export { filterByLeafNodeId, + findMessageById, findLeafNode, findDescendantMessages, getMessageSiblings, @@ -38,7 +39,7 @@ export { highlightCode, detectIncompleteCodeBlock, type IncompleteCodeBlock } fr export { setConfigValue, getConfigValue, configToParameterRecord } from './config-helpers'; // CORS Proxy -export { buildProxiedUrl, getProxiedUrlString } from './cors-proxy'; +export { buildProxiedUrl, getProxiedUrlString, buildProxiedHeaders } from './cors-proxy'; // Conversation utilities export { createMessageCountMap, getMessageCount } from './conversation-utils';