mirror of
https://github.com/LostRuins/koboldcpp.git
synced 2026-08-14 19:05:22 +00:00
Merge branch 'upstream' into concedo_experimental
# Conflicts: # .devops/intel.Dockerfile # .github/workflows/python-type-check.yml # embd_res/templates/Qwen3.5-4B.jinja # examples/model-conversion/scripts/causal/compare-logits.py # examples/model-conversion/scripts/utils/check-nmse.py # examples/model-conversion/scripts/utils/compare_tokens.py # examples/model-conversion/scripts/utils/semantic_check.py # examples/sycl/build.sh # examples/sycl/run-llama2.sh # ggml/src/ggml-hexagon/htp/flash-attn-ops.c # ggml/src/ggml-hexagon/htp/hex-dma.h # ggml/src/ggml-hexagon/htp/rope-ops.c # scripts/gen-unicode-data.py # tests/test-chat.cpp
This commit is contained in:
Concedo
commit
a3a5897d93
20 changed files with 565 additions and 152 deletions
|
|
@ -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<std::string> 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<common_peg_parser> 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) {
|
||||
|
|
|
|||
|
|
@ -663,14 +663,53 @@ bool string_parse_kv_override(const char * data, std::vector<llama_model_kv_over
|
|||
return true;
|
||||
}
|
||||
|
||||
// simple glob: * matches non-/ chars, ** matches anything including /
|
||||
static inline bool glob_class_match(const char c, const char * pattern, const char * class_end) {
|
||||
const char * class_start = pattern;
|
||||
bool negated = false;
|
||||
|
||||
if (*class_start == '!') {
|
||||
negated = true;
|
||||
class_start++;
|
||||
}
|
||||
|
||||
// If first character after negation is ']' or '-', treat it as literal
|
||||
if (*class_start == ']' || *class_start == '-') {
|
||||
if (class_start < class_end && *class_start == c) {
|
||||
return !negated;
|
||||
}
|
||||
class_start++;
|
||||
}
|
||||
|
||||
bool matched = false;
|
||||
|
||||
while (class_start < class_end) {
|
||||
if (class_start + 2 < class_end && class_start[1] == '-' && class_start[2] != ']') {
|
||||
char start_char = *class_start;
|
||||
char end_char = class_start[2];
|
||||
if (c >= 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);
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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,
|
||||
)
|
||||
|
||||
|
|
|
|||
|
|
@ -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) {
|
||||
|
|
|
|||
|
|
@ -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 <ggml_type type>
|
||||
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 <ggml_type type, int ncols_dst, bool has_fusion, bool is_multi_token_id = false, bool small_k = false>
|
||||
template <ggml_type type, int ncols_dst, bool has_fusion, bool small_k = false>
|
||||
__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 <ggml_type type, int c_rows_per_block>
|
||||
__launch_bounds__(get_mmvq_mmid_max_batch_for_device<type>()*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<type>::qk;
|
||||
constexpr int qi = ggml_cuda_type_traits<type>::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<warp_size>(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<ggml_type type>
|
||||
static std::pair<dim3, dim3> 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<dim3, dim3> calc_launch_params(
|
|||
return {block_nums, block_dims};
|
||||
}
|
||||
|
||||
template<ggml_type type, int c_ncols_dst, bool is_multi_token_id = false, bool small_k = false>
|
||||
template<ggml_type type, int c_ncols_dst, bool small_k = false>
|
||||
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<type, c_ncols_dst, true, is_multi_token_id, small_k><<<block_nums, block_dims, nbytes_shared, stream>>>
|
||||
mul_mat_vec_q<type, c_ncols_dst, true, small_k><<<block_nums, block_dims, nbytes_shared, stream>>>
|
||||
(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<type, c_ncols_dst, false, is_multi_token_id, small_k><<<block_nums, block_dims, nbytes_shared, stream>>>
|
||||
mul_mat_vec_q<type, c_ncols_dst, false, small_k><<<block_nums, block_dims, nbytes_shared, stream>>>
|
||||
(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 <ggml_type type>
|
||||
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<type, rows_per_block><<<block_nums, block_dims, 0, stream>>>(
|
||||
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 <ggml_type type>
|
||||
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<type>::qk;
|
||||
constexpr int qi = ggml_cuda_type_traits<type>::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<ggml_type, 2> iq_slow_turing = {
|
||||
GGML_TYPE_IQ3_XXS,
|
||||
GGML_TYPE_IQ3_S,
|
||||
};
|
||||
constexpr std::array<ggml_type, 8> 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<ggml_type, 3> 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<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst, ncols_dst, warp_size, table_id);
|
||||
mul_mat_vec_q_switch_fusion<type, c_ncols_dst, true>(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<type>(
|
||||
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<type>::qk;
|
||||
constexpr int qi = ggml_cuda_type_traits<type>::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<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst, nsamples_dst,
|
||||
warp_size, table_id, true);
|
||||
mul_mat_vec_q_switch_fusion<type, c_ncols_dst, false, true>(
|
||||
std::pair<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst,
|
||||
nsamples_dst, warp_size, table_id, true);
|
||||
mul_mat_vec_q_switch_fusion<type, c_ncols_dst, true>(
|
||||
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<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst, nsamples_dst,
|
||||
warp_size, table_id);
|
||||
std::pair<dim3, dim3> dims = calc_launch_params<type>(c_ncols_dst, nrows_x, nchannels_dst,
|
||||
nsamples_dst, warp_size, table_id);
|
||||
mul_mat_vec_q_switch_fusion<type, c_ncols_dst>(
|
||||
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: {
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
|
|
|
|||
|
|
@ -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:
|
||||
|
|
|
|||
|
|
@ -557,6 +557,8 @@ static std::set<llm_tensor> 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,
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
}
|
||||
|
|
|
|||
Binary file not shown.
|
|
@ -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<std::string, std::string> 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<server_http_proxy>(
|
||||
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)
|
||||
|
|
|
|||
|
|
@ -35,7 +35,7 @@ using server_http_res_ptr = std::unique_ptr<server_http_res>;
|
|||
|
||||
struct server_http_req {
|
||||
std::map<std::string, std::string> params; // path_params + query_params
|
||||
std::map<std::string, std::string> headers; // reserved for future use
|
||||
std::map<std::string, std::string> headers; // used by MCP proxy
|
||||
std::string path;
|
||||
std::string query_string; // query parameters string (e.g. "action=save")
|
||||
std::string body;
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -18,7 +18,8 @@
|
|||
showRaw = undefined,
|
||||
aliases,
|
||||
tags,
|
||||
class: className = ''
|
||||
class: className = '',
|
||||
...rest
|
||||
}: Props = $props();
|
||||
|
||||
const badgeClass =
|
||||
|
|
@ -36,9 +37,9 @@
|
|||
</script>
|
||||
|
||||
{#if resolvedShowRaw}
|
||||
<TruncatedText class="font-medium {className}" showTooltip={false} text={modelId} />
|
||||
<TruncatedText class="font-medium {className}" showTooltip={false} text={modelId} {...rest} />
|
||||
{:else}
|
||||
<span class="flex min-w-0 flex-wrap items-center gap-1 {className}">
|
||||
<span class="flex min-w-0 flex-wrap items-center gap-1 {className}" {...rest}>
|
||||
<span class="min-w-0 truncate font-medium">
|
||||
{#if showOrgName && parsed.orgName && !(aliases && aliases.length > 0)}{parsed.orgName}/{/if}{displayName}
|
||||
</span>
|
||||
|
|
|
|||
|
|
@ -271,50 +271,49 @@
|
|||
{#if isRouter}
|
||||
<DropdownMenu.Root bind:open={isOpen} onOpenChange={handleOpenChange}>
|
||||
<DropdownMenu.Trigger
|
||||
disabled={disabled || updating}
|
||||
onclick={(e) => {
|
||||
e.preventDefault();
|
||||
e.stopPropagation();
|
||||
}}
|
||||
>
|
||||
<button
|
||||
type="button"
|
||||
class={cn(
|
||||
`inline-grid cursor-pointer grid-cols-[1fr_auto_1fr] items-center gap-1.5 rounded-sm bg-muted-foreground/10 px-1.5 py-1 text-xs transition hover:text-foreground focus:outline-none focus-visible:ring-2 focus-visible:ring-ring focus-visible:ring-offset-2 disabled:cursor-not-allowed disabled:opacity-60`,
|
||||
!isCurrentModelInCache
|
||||
? 'bg-red-400/10 !text-red-400 hover:bg-red-400/20 hover:text-red-400'
|
||||
: forceForegroundText
|
||||
class={cn(
|
||||
`inline-grid cursor-pointer grid-cols-[1fr_auto_1fr] items-center gap-1.5 rounded-sm bg-muted-foreground/10 px-1.5 py-1 text-xs transition hover:text-foreground focus:outline-none focus-visible:ring-2 focus-visible:ring-ring focus-visible:ring-offset-2 disabled:cursor-not-allowed disabled:opacity-60`,
|
||||
!isCurrentModelInCache
|
||||
? 'bg-red-400/10 !text-red-400 hover:bg-red-400/20 hover:text-red-400'
|
||||
: forceForegroundText
|
||||
? 'text-foreground'
|
||||
: isHighlightedCurrentModelActive
|
||||
? 'text-foreground'
|
||||
: isHighlightedCurrentModelActive
|
||||
? 'text-foreground'
|
||||
: 'text-muted-foreground',
|
||||
isOpen ? 'text-foreground' : ''
|
||||
)}
|
||||
style="max-width: min(calc(100cqw - 9rem), 20rem)"
|
||||
disabled={disabled || updating}
|
||||
>
|
||||
<Package class="h-3.5 w-3.5" />
|
||||
: 'text-muted-foreground',
|
||||
isOpen ? 'text-foreground' : ''
|
||||
)}
|
||||
style="max-width: min(calc(100cqw - 9rem), 20rem)"
|
||||
disabled={disabled || updating}
|
||||
>
|
||||
<Package class="h-3.5 w-3.5" />
|
||||
|
||||
{#if selectedOption}
|
||||
<Tooltip.Root>
|
||||
<Tooltip.Trigger class="min-w-0 overflow-hidden">
|
||||
<ModelId modelId={selectedOption.model} class="min-w-0" showOrgName />
|
||||
</Tooltip.Trigger>
|
||||
{#if selectedOption}
|
||||
<Tooltip.Root>
|
||||
<Tooltip.Trigger>
|
||||
<!-- prevent another nested button element -->
|
||||
{#snippet child({ props })}
|
||||
<ModelId
|
||||
modelId={selectedOption.model}
|
||||
class="min-w-0 overflow-hidden"
|
||||
showOrgName
|
||||
{...props}
|
||||
/>
|
||||
{/snippet}
|
||||
</Tooltip.Trigger>
|
||||
|
||||
<Tooltip.Content>
|
||||
<p class="font-mono">{selectedOption.model}</p>
|
||||
</Tooltip.Content>
|
||||
</Tooltip.Root>
|
||||
{:else}
|
||||
<span class="min-w-0 font-medium">Select model</span>
|
||||
{/if}
|
||||
<Tooltip.Content>
|
||||
<p class="font-mono">{selectedOption.model}</p>
|
||||
</Tooltip.Content>
|
||||
</Tooltip.Root>
|
||||
{:else}
|
||||
<span class="min-w-0 font-medium">Select model</span>
|
||||
{/if}
|
||||
|
||||
{#if updating || isLoadingModel}
|
||||
<Loader2 class="h-3 w-3.5 animate-spin" />
|
||||
{:else}
|
||||
<ChevronDown class="h-3 w-3.5" />
|
||||
{/if}
|
||||
</button>
|
||||
{#if updating || isLoadingModel}
|
||||
<Loader2 class="h-3 w-3.5 animate-spin" />
|
||||
{:else}
|
||||
<ChevronDown class="h-3 w-3.5" />
|
||||
{/if}
|
||||
</DropdownMenu.Trigger>
|
||||
|
||||
<DropdownMenu.Content
|
||||
|
|
@ -407,8 +406,16 @@
|
|||
|
||||
{#if selectedOption}
|
||||
<Tooltip.Root>
|
||||
<Tooltip.Trigger class="min-w-0 overflow-hidden">
|
||||
<ModelId modelId={selectedOption.model} class="min-w-0" showOrgName />
|
||||
<Tooltip.Trigger>
|
||||
<!-- prevent another nested button element -->
|
||||
{#snippet child({ props })}
|
||||
<ModelId
|
||||
modelId={selectedOption.model}
|
||||
class="min-w-0 overflow-hidden"
|
||||
showOrgName
|
||||
{...props}
|
||||
/>
|
||||
{/snippet}
|
||||
</Tooltip.Trigger>
|
||||
|
||||
<Tooltip.Content>
|
||||
|
|
|
|||
|
|
@ -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) {
|
||||
|
|
|
|||
|
|
@ -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<DatabaseMessage> = {
|
||||
content: newContent,
|
||||
|
|
|
|||
|
|
@ -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.
|
||||
|
|
|
|||
|
|
@ -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<string, string>): Record<string, string> {
|
||||
const proxiedHeaders: Record<string, string> = {};
|
||||
|
||||
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
|
||||
|
|
|
|||
|
|
@ -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';
|
||||
|
|
|
|||
Loading…
Add table
Add a link
Reference in a new issue