From 6c7fd67b647a76846d1691cd181011dff4549d02 Mon Sep 17 00:00:00 2001 From: piDack <104877312+piDack@users.noreply.github.com> Date: Wed, 7 May 2025 15:23:11 +0800 Subject: [PATCH 01/10] llama : support tie embedding for chatglm models (#13328) --- src/llama-model.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/src/llama-model.cpp b/src/llama-model.cpp index 8d25070f4..774e343fb 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -3510,7 +3510,11 @@ bool llama_model::load_tensors(llama_model_loader & ml) { // output output_norm = create_tensor(tn(LLM_TENSOR_OUTPUT_NORM, "weight"), {n_embd}, 0); - output = create_tensor(tn(LLM_TENSOR_OUTPUT, "weight"), {n_embd, n_vocab}, 0); + output = create_tensor(tn(LLM_TENSOR_OUTPUT, "weight"), {n_embd, n_vocab}, TENSOR_NOT_REQUIRED); + // if output is NULL, init from the input tok embed + if (output == NULL) { + output = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, TENSOR_DUPLICATED); + } for (int i = 0; i < n_layer; ++i) { auto & layer = layers[i]; From 4773d7a02ffdb05ba9e673ff21ce95351836e33a Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Wed, 7 May 2025 10:28:02 +0300 Subject: [PATCH 02/10] examples : remove infill (#13283) ggml-ci --- Makefile | 5 - common/arg.cpp | 12 +- common/common.h | 1 - examples/CMakeLists.txt | 1 - examples/infill/CMakeLists.txt | 5 - examples/infill/README.md | 47 --- examples/infill/infill.cpp | 590 --------------------------------- 7 files changed, 6 insertions(+), 655 deletions(-) delete mode 100644 examples/infill/CMakeLists.txt delete mode 100644 examples/infill/README.md delete mode 100644 examples/infill/infill.cpp diff --git a/Makefile b/Makefile index 68b6fe59a..958ad8f2f 100644 --- a/Makefile +++ b/Makefile @@ -1187,11 +1187,6 @@ llama-cli: tools/main/main.cpp \ @echo '==== Run ./llama-cli -h for help. ====' @echo -llama-infill: examples/infill/infill.cpp \ - $(OBJ_ALL) - $(CXX) $(CXXFLAGS) -c $< -o $(call GET_OBJ_FILE, $<) - $(CXX) $(CXXFLAGS) $(filter-out %.h $<,$^) $(call GET_OBJ_FILE, $<) -o $@ $(LDFLAGS) - llama-run: tools/run/run.cpp \ $(OBJ_ALL) $(CXX) $(CXXFLAGS) -c $< -o $(call GET_OBJ_FILE, $<) diff --git a/common/arg.cpp b/common/arg.cpp index 490991b23..5e07e8a69 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -1283,7 +1283,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex [](common_params & params) { params.use_color = true; } - ).set_examples({LLAMA_EXAMPLE_MAIN, LLAMA_EXAMPLE_INFILL, LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_LOOKUP})); + ).set_examples({LLAMA_EXAMPLE_MAIN, LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_LOOKUP})); add_opt(common_arg( {"-t", "--threads"}, "N", string_format("number of threads to use during generation (default: %d)", params.cpuparams.n_threads), @@ -1416,7 +1416,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex add_opt(common_arg( {"-n", "--predict", "--n-predict"}, "N", string_format( - ex == LLAMA_EXAMPLE_MAIN || ex == LLAMA_EXAMPLE_INFILL + ex == LLAMA_EXAMPLE_MAIN ? "number of tokens to predict (default: %d, -1 = infinity, -2 = until context filled)" : "number of tokens to predict (default: %d, -1 = infinity)", params.n_predict), @@ -1655,7 +1655,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex params.input_prefix = value; params.enable_chat_template = false; } - ).set_examples({LLAMA_EXAMPLE_MAIN, LLAMA_EXAMPLE_INFILL})); + ).set_examples({LLAMA_EXAMPLE_MAIN})); add_opt(common_arg( {"--in-suffix"}, "STRING", "string to suffix after user inputs with (default: empty)", @@ -1663,7 +1663,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex params.input_suffix = value; params.enable_chat_template = false; } - ).set_examples({LLAMA_EXAMPLE_MAIN, LLAMA_EXAMPLE_INFILL})); + ).set_examples({LLAMA_EXAMPLE_MAIN})); add_opt(common_arg( {"--no-warmup"}, "skip warming up the model with an empty run", @@ -1680,7 +1680,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex [](common_params & params) { params.spm_infill = true; } - ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_INFILL})); + ).set_examples({LLAMA_EXAMPLE_SERVER})); add_opt(common_arg( {"--samplers"}, "SAMPLERS", string_format("samplers that will be used for generation in the order, separated by \';\'\n(default: %s)", sampler_type_names.c_str()), @@ -2892,7 +2892,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex [](common_params & params) { params.simple_io = true; } - ).set_examples({LLAMA_EXAMPLE_MAIN, LLAMA_EXAMPLE_INFILL})); + ).set_examples({LLAMA_EXAMPLE_MAIN})); add_opt(common_arg( {"--positive-file"}, "FNAME", string_format("positive prompts file, one prompt per line (default: '%s')", params.cvector_positive_file.c_str()), diff --git a/common/common.h b/common/common.h index dfd6e2093..400f674b2 100644 --- a/common/common.h +++ b/common/common.h @@ -66,7 +66,6 @@ enum llama_example { LLAMA_EXAMPLE_COMMON, LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_MAIN, - LLAMA_EXAMPLE_INFILL, LLAMA_EXAMPLE_EMBEDDING, LLAMA_EXAMPLE_PERPLEXITY, LLAMA_EXAMPLE_RETRIEVAL, diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index eca0d0b09..4ca9230c5 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -21,7 +21,6 @@ else() add_subdirectory(gguf-hash) add_subdirectory(gguf) add_subdirectory(gritlm) - add_subdirectory(infill) add_subdirectory(lookahead) add_subdirectory(lookup) add_subdirectory(parallel) diff --git a/examples/infill/CMakeLists.txt b/examples/infill/CMakeLists.txt deleted file mode 100644 index fb26628d8..000000000 --- a/examples/infill/CMakeLists.txt +++ /dev/null @@ -1,5 +0,0 @@ -set(TARGET llama-infill) -add_executable(${TARGET} infill.cpp) -install(TARGETS ${TARGET} RUNTIME) -target_link_libraries(${TARGET} PRIVATE common llama ${CMAKE_THREAD_LIBS_INIT}) -target_compile_features(${TARGET} PRIVATE cxx_std_17) diff --git a/examples/infill/README.md b/examples/infill/README.md deleted file mode 100644 index df4d976f2..000000000 --- a/examples/infill/README.md +++ /dev/null @@ -1,47 +0,0 @@ -# llama.cpp/example/infill - -This example shows how to use the infill mode with Code Llama models supporting infill mode. -Currently the 7B and 13B models support infill mode. - -Infill supports most of the options available in the main example. - -For further information have a look at the main README.md in llama.cpp/example/main/README.md - -## Common Options - -In this section, we cover the most commonly used options for running the `infill` program with the LLaMA models: - -- `-m FNAME, --model FNAME`: Specify the path to the LLaMA model file (e.g., `models/7B/ggml-model.bin`). -- `-i, --interactive`: Run the program in interactive mode, allowing you to provide input directly and receive real-time responses. -- `-n N, --n-predict N`: Set the number of tokens to predict when generating text. Adjusting this value can influence the length of the generated text. -- `-c N, --ctx-size N`: Set the size of the prompt context. The default is 4096, but if a LLaMA model was built with a longer context, increasing this value will provide better results for longer input/inference. -- `--spm-infill`: Use Suffix/Prefix/Middle pattern for infill (instead of Prefix/Suffix/Middle) as some models prefer this. - -## Input Prompts - -The `infill` program provides several ways to interact with the LLaMA models using input prompts: - -- `--in-prefix PROMPT_BEFORE_CURSOR`: Provide the prefix directly as a command-line option. -- `--in-suffix PROMPT_AFTER_CURSOR`: Provide the suffix directly as a command-line option. -- `--interactive-first`: Run the program in interactive mode and wait for input right away. (More on this below.) - -## Interaction - -The `infill` program offers a seamless way to interact with LLaMA models, allowing users to receive real-time infill suggestions. The interactive mode can be triggered using `--interactive`, and `--interactive-first` - -### Interaction Options - -- `-i, --interactive`: Run the program in interactive mode, allowing users to get real time code suggestions from model. -- `--interactive-first`: Run the program in interactive mode and immediately wait for user input before starting the text generation. -- `--color`: Enable colorized output to differentiate visually distinguishing between prompts, user input, and generated text. - -### Example - -Download a model that supports infill, for example CodeLlama: -```console -scripts/hf.sh --repo TheBloke/CodeLlama-13B-GGUF --file codellama-13b.Q5_K_S.gguf --outdir models -``` - -```bash -./llama-infill -t 10 -ngl 0 -m models/codellama-13b.Q5_K_S.gguf -c 4096 --temp 0.7 --repeat_penalty 1.1 -n 20 --in-prefix "def helloworld():\n print(\"hell" --in-suffix "\n print(\"goodbye world\")\n " -``` diff --git a/examples/infill/infill.cpp b/examples/infill/infill.cpp deleted file mode 100644 index 4e2f7b727..000000000 --- a/examples/infill/infill.cpp +++ /dev/null @@ -1,590 +0,0 @@ -#include "arg.h" -#include "common.h" -#include "console.h" -#include "sampling.h" -#include "log.h" -#include "llama.h" - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#if defined (__unix__) || (defined (__APPLE__) && defined (__MACH__)) -#include -#include -#elif defined (_WIN32) -#define WIN32_LEAN_AND_MEAN -#ifndef NOMINMAX -#define NOMINMAX -#endif -#include -#include -#endif - -#if defined(_MSC_VER) -#pragma warning(disable: 4244 4267) // possible loss of data -#endif - -static llama_context ** g_ctx; -static llama_model ** g_model; -static common_sampler ** g_smpl; -static common_params * g_params; -static std::vector * g_input_tokens; -static std::ostringstream * g_output_ss; -static std::vector * g_output_tokens; - -static bool is_interacting = false; - -#if defined (__unix__) || (defined (__APPLE__) && defined (__MACH__)) || defined (_WIN32) -static void sigint_handler(int signo) { - if (signo == SIGINT) { - if (!is_interacting) { - is_interacting = true; - } else { - console::cleanup(); - LOG("\n"); - common_perf_print(*g_ctx, *g_smpl); - - // make sure all logs are flushed - LOG("Interrupted by user\n"); - common_log_pause(common_log_main()); - - _exit(130); - } - } -} -#endif - -int main(int argc, char ** argv) { - common_params params; - g_params = ¶ms; - - if (!common_params_parse(argc, argv, params, LLAMA_EXAMPLE_INFILL)) { - return 1; - } - - common_init(); - - auto & sparams = params.sampling; - - console::init(params.simple_io, params.use_color); - atexit([]() { console::cleanup(); }); - - if (params.logits_all) { - LOG_ERR("\n************\n"); - LOG_ERR("%s: please use the 'perplexity' tool for perplexity calculations\n", __func__); - LOG_ERR("************\n\n"); - - return 0; - } - - if (params.embedding) { - LOG_ERR("\n************\n"); - LOG_ERR("%s: please use the 'embedding' tool for embedding calculations\n", __func__); - LOG_ERR("************\n\n"); - - return 0; - } - - if (params.n_ctx != 0 && params.n_ctx < 8) { - LOG_WRN("%s: minimum context size is 8, using minimum size.\n", __func__); - params.n_ctx = 8; - } - - if (!params.interactive_first && (params.input_prefix.empty() && params.input_suffix.empty())) { - LOG_ERR("\n************\n"); - LOG_ERR("%s: please use '--interactive_first' or specify '--in_prefix' and/or '--in_suffix'\n", __func__); - LOG_ERR("************\n\n"); - - return 0; - } - - if (params.rope_freq_base != 0.0) { - LOG_WRN("%s: changing RoPE frequency base to %g.\n", __func__, params.rope_freq_base); - } - - if (params.rope_freq_scale != 0.0) { - LOG_WRN("%s: scaling RoPE frequency by %g.\n", __func__, params.rope_freq_scale); - } - - LOG_INF("%s: llama backend init\n", __func__); - llama_backend_init(); - llama_numa_init(params.numa); - - llama_model * model = nullptr; - llama_context * ctx = nullptr; - common_sampler * smpl = nullptr; - - g_model = &model; - g_ctx = &ctx; - g_smpl = &smpl; - - // load the model and apply lora adapter, if any - LOG_INF("%s: load the model and apply lora adapter, if any\n", __func__); - common_init_result llama_init = common_init_from_params(params); - - model = llama_init.model.get(); - ctx = llama_init.context.get(); - - if (model == NULL) { - LOG_ERR("%s: unable to load model\n", __func__); - return 1; - } - - const llama_vocab * vocab = llama_model_get_vocab(model); - - const int n_ctx_train = llama_model_n_ctx_train(model); - const int n_ctx = llama_n_ctx(ctx); - LOG_DBG("n_ctx: %d\n", n_ctx); - - if (n_ctx > n_ctx_train) { - LOG_WRN("%s: model was trained on only %d context tokens (%d specified)\n", __func__, n_ctx_train, n_ctx); - } - - // print system information - { - LOG_INF("\n"); - LOG_INF("%s\n", common_params_get_system_info(params).c_str()); - } - const bool add_bos = llama_vocab_get_add_bos(vocab); - GGML_ASSERT(!llama_vocab_get_add_eos(vocab)); - - std::vector embd_inp; - std::vector embd_end; - std::vector inp_pfx = common_tokenize(ctx, params.input_prefix, false); - std::vector inp_sfx = common_tokenize(ctx, params.input_suffix, false); - - GGML_ASSERT(llama_vocab_fim_pre(vocab) >= 0); - GGML_ASSERT(llama_vocab_fim_suf(vocab) >= 0); - - inp_pfx.insert(inp_pfx.begin(), llama_vocab_fim_pre(vocab)); - inp_sfx.insert(inp_sfx.begin(), llama_vocab_fim_suf(vocab)); - - embd_inp = params.spm_infill ? inp_sfx : inp_pfx; - embd_end = params.spm_infill ? inp_pfx : inp_sfx; - if (add_bos) { - embd_inp.insert(embd_inp.begin(), llama_vocab_bos(vocab)); - } - embd_inp.insert(embd_inp.end(), embd_end.begin(), embd_end.end()); - - const llama_token middle_token = llama_vocab_fim_mid(vocab); - if (middle_token >= 0) { - embd_inp.push_back(middle_token); - } - - LOG_DBG("add_bos: %d\n", add_bos); - LOG_DBG("prefix: \"%s\"\n", params.input_prefix.c_str()); - LOG_DBG("suffix: \"%s\"\n", params.input_suffix.c_str()); - LOG_DBG("tokens: %s\n", string_from(ctx, embd_inp).c_str()); - - // Should not run without any tokens - if (embd_inp.empty()) { - embd_inp.push_back(llama_vocab_bos(vocab)); - LOG_WRN("embd_inp was considered empty and bos was added: %s\n", string_from(ctx, embd_inp).c_str()); - } - - if ((int) embd_inp.size() > n_ctx - 4) { - LOG_ERR("%s: prompt is too long (%d tokens, max %d)\n", __func__, (int) embd_inp.size(), n_ctx - 4); - return 1; - } - - // number of tokens to keep when resetting context - if (params.n_keep < 0 || params.n_keep > (int) embd_inp.size()) { - params.n_keep = (int)embd_inp.size(); - } - - LOG_INF("inp_pfx: %s\n", string_from(ctx, inp_pfx).c_str()); - LOG_INF("inp_sfx: %s\n", string_from(ctx, inp_sfx).c_str()); - - // enable interactive mode if interactive start is specified - if (params.interactive_first) { - params.interactive = true; - } - - if (params.verbose_prompt) { - LOG_INF("\n"); - LOG_INF("%s: prompt: '%s'\n", __func__, params.prompt.c_str()); - LOG_INF("%s: number of tokens in prompt = %zu\n", __func__, embd_inp.size()); - for (int i = 0; i < (int) embd_inp.size(); i++) { - LOG_INF("%6d -> '%s'\n", embd_inp[i], common_token_to_piece(ctx, embd_inp[i]).c_str()); - } - - if (params.n_keep > 0) { - LOG_INF("%s: static prompt based on n_keep: '", __func__); - for (int i = 0; i < params.n_keep; i++) { - LOG_CNT("%s", common_token_to_piece(ctx, embd_inp[i]).c_str()); - } - LOG_CNT("'\n"); - } - LOG_INF("\n"); - } - - if (params.interactive) { -#if defined (__unix__) || (defined (__APPLE__) && defined (__MACH__)) - struct sigaction sigint_action; - sigint_action.sa_handler = sigint_handler; - sigemptyset (&sigint_action.sa_mask); - sigint_action.sa_flags = 0; - sigaction(SIGINT, &sigint_action, NULL); -#elif defined (_WIN32) - auto console_ctrl_handler = +[](DWORD ctrl_type) -> BOOL { - return (ctrl_type == CTRL_C_EVENT) ? (sigint_handler(SIGINT), true) : false; - }; - SetConsoleCtrlHandler(reinterpret_cast(console_ctrl_handler), true); -#endif - - LOG_INF("%s: interactive mode on.\n", __func__); - - if (params.input_prefix_bos) { - LOG_INF("Input prefix with BOS\n"); - } - - if (!params.input_prefix.empty()) { - LOG_INF("Input prefix: '%s'\n", params.input_prefix.c_str()); - } - - if (!params.input_suffix.empty()) { - LOG_INF("Input suffix: '%s'\n", params.input_suffix.c_str()); - } - } - smpl = common_sampler_init(model, sparams); - - LOG_INF("sampler seed: %u\n", common_sampler_get_seed(smpl)); - LOG_INF("sampler params: \n%s\n", sparams.print().c_str()); - LOG_INF("sampler chain: %s\n", common_sampler_print(smpl).c_str()); - - LOG_INF("generate: n_ctx = %d, n_batch = %d, n_predict = %d, n_keep = %d\n", n_ctx, params.n_batch, params.n_predict, params.n_keep); - - LOG_INF("\n"); - LOG_INF("\n##### Infill mode #####\n\n"); - if (params.interactive) { - const char *control_message; - if (params.multiline_input) { - control_message = " - To return control to LLaMA, end your input with '\\'.\n" - " - To return control without starting a new line, end your input with '/'.\n"; - } else { - control_message = " - Press Return to return control to LLaMA.\n" - " - To return control without starting a new line, end your input with '/'.\n" - " - If you want to submit another line, end your input with '\\'.\n"; - } - LOG_INF("== Running in interactive mode. ==\n"); -#if defined (__unix__) || (defined (__APPLE__) && defined (__MACH__)) || defined (_WIN32) - LOG_INF( " - Press Ctrl+C to interject at any time.\n"); -#endif - LOG_INF( "%s\n", control_message); - - is_interacting = params.interactive_first; - } - - bool input_echo = true; - - int n_past = 0; - int n_remain = params.n_predict; - int n_consumed = 0; - - std::vector input_tokens; g_input_tokens = &input_tokens; - std::vector output_tokens; g_output_tokens = &output_tokens; - std::ostringstream output_ss; g_output_ss = &output_ss; - - // the first thing we will do is to output the prompt, so set color accordingly - console::set_display(console::prompt); - - std::vector embd; - - while (n_remain != 0 || params.interactive) { - // predict - if (!embd.empty()) { - // Note: n_ctx - 4 here is to match the logic for commandline prompt handling via - // --prompt or --file which uses the same value. - int max_embd_size = n_ctx - 4; - - // Ensure the input doesn't exceed the context size by truncating embd if necessary. - if ((int) embd.size() > max_embd_size) { - const int skipped_tokens = (int) embd.size() - max_embd_size; - embd.resize(max_embd_size); - - console::set_display(console::error); - LOG_WRN("<>", skipped_tokens, skipped_tokens != 1 ? "s" : ""); - console::set_display(console::reset); - } - - // infinite text generation via context swapping - // if we run out of context: - // - take the n_keep first tokens from the original prompt (via n_past) - // - take half of the last (n_ctx - n_keep) tokens and recompute the logits in batches - if (n_past + (int) embd.size() > n_ctx) { - if (params.n_predict == -2) { - LOG_DBG("\n\n%s: context full and n_predict == -%d => stopping\n", __func__, params.n_predict); - break; - } - - const int n_left = n_past - params.n_keep - 1; - const int n_discard = n_left/2; - - LOG_DBG("context full, swapping: n_past = %d, n_left = %d, n_ctx = %d, n_keep = %d, n_discard = %d\n", - n_past, n_left, n_ctx, params.n_keep, n_discard); - - llama_kv_self_seq_rm (ctx, 0, params.n_keep + 1 , params.n_keep + n_discard + 1); - llama_kv_self_seq_add(ctx, 0, params.n_keep + 1 + n_discard, n_past, -n_discard); - - n_past -= n_discard; - - LOG_DBG("after swap: n_past = %d\n", n_past); - - LOG_DBG("embd: %s\n", string_from(ctx, embd).c_str()); - - } - - // evaluate tokens in batches - // embd is typically prepared beforehand to fit within a batch, but not always - for (int i = 0; i < (int) embd.size(); i += params.n_batch) { - int n_eval = (int) embd.size() - i; - if (n_eval > params.n_batch) { - n_eval = params.n_batch; - } - - LOG_DBG("eval: %s\n", string_from(ctx, embd).c_str()); - - if (llama_decode(ctx, llama_batch_get_one(&embd[i], n_eval))) { - LOG_ERR("%s : failed to eval\n", __func__); - return 1; - } - - n_past += n_eval; - - LOG_DBG("n_past = %d\n", n_past); - } - - } - - embd.clear(); - - if ((int) embd_inp.size() <= n_consumed && !is_interacting) { - const llama_token id = common_sampler_sample(smpl, ctx, -1); - - common_sampler_accept(smpl, id, true); - - // LOG_DBG("last: %s\n", string_from(ctx, smpl->prev.to_vector()).c_str()); - - embd.push_back(id); - - // echo this to console - input_echo = true; - - // decrement remaining sampling budget - --n_remain; - - LOG_DBG("n_remain: %d\n", n_remain); - } else { - // some user input remains from prompt or interaction, forward it to processing - LOG_DBG("embd_inp.size(): %d, n_consumed: %d\n", (int) embd_inp.size(), n_consumed); - while ((int) embd_inp.size() > n_consumed) { - embd.push_back(embd_inp[n_consumed]); - - // push the prompt in the sampling context in order to apply repetition penalties later - // for the prompt, we don't apply grammar rules - common_sampler_accept(smpl, embd_inp[n_consumed], false); - - ++n_consumed; - if ((int) embd.size() >= params.n_batch) { - break; - } - } - } - - // display text - if (input_echo) { - for (auto id : embd) { - const std::string token_str = common_token_to_piece(ctx, id); - LOG("%s", token_str.c_str()); - - if (embd.size() > 1) { - input_tokens.push_back(id); - } else { - output_tokens.push_back(id); - output_ss << token_str; - } - } - } - // reset color to default if we there is no pending user input - if (input_echo && (int) embd_inp.size() == n_consumed) { - console::set_display(console::reset); - } - - // if not currently processing queued inputs; - if ((int) embd_inp.size() <= n_consumed) { - // deal with eot token in infill mode - if ((common_sampler_last(smpl) == llama_vocab_eot(vocab) || is_interacting) && params.interactive){ - if (is_interacting && !params.interactive_first) { - // print an eot token - LOG("%s", common_token_to_piece(ctx, llama_vocab_eot(vocab)).c_str()); - } - LOG("\n"); - console::set_display(console::user_input); - std::string buffer; - std::string line; - bool another_line=true; - // set a new prefix via stdin - do { - another_line = console::readline(line, params.multiline_input); - buffer += line; - } while (another_line); - // check if we got an empty line, if so we use the old input - if (!buffer.empty() && !(buffer.length() == 1 && buffer[0] == '\n')) { - params.input_prefix = buffer; - } - buffer.clear(); - // set a new suffix via stdin - do { - another_line = console::readline(line, params.multiline_input); - buffer += line; - } while (another_line); - // check if we got an empty line - if (!buffer.empty() && !(buffer.length() == 1 && buffer[0] == '\n')) { - params.input_suffix = buffer; - } - buffer.clear(); - // done taking input, reset color - console::set_display(console::reset); - - if (params.escape) { - //process escape sequences, for the initial prompt this is done in common.cpp when we load the params, but for the interactive mode we need to do it here - string_process_escapes(params.input_prefix); - string_process_escapes(params.input_suffix); - } - - // tokenize new prefix and suffix - std::vector inp_pfx = common_tokenize(ctx, params.input_prefix, false); - std::vector inp_sfx = common_tokenize(ctx, params.input_suffix, false); - - inp_pfx.insert(inp_pfx.begin(), llama_vocab_fim_pre(vocab)); - inp_sfx.insert(inp_sfx.begin(), llama_vocab_fim_suf(vocab)); - - embd_inp = params.spm_infill ? inp_sfx : inp_pfx; - embd_end = params.spm_infill ? inp_pfx : inp_sfx; - if (add_bos) { - embd_inp.insert(embd_inp.begin(), llama_vocab_bos(vocab)); - } - embd_inp.insert(embd_inp.end(), embd_end.begin(), embd_end.end()); - - if (middle_token >= 0) { - embd_inp.push_back(middle_token); - } - - embd.clear(); - n_remain = params.n_predict; - n_past = 0; - n_consumed = 0; - is_interacting = false; - } - // deal with end of generation tokens in interactive mode - else if (llama_vocab_is_eog(vocab, common_sampler_last(smpl))) { - LOG_DBG("found EOS token\n"); - - if (params.interactive) { - - is_interacting = true; - LOG("\n"); - console::set_display(console::user_input); - } - } - - if (n_past > 0 && is_interacting && !params.interactive) { - LOG_DBG("waiting for user input\n"); - - if (params.input_prefix_bos) { - LOG_DBG("adding input prefix BOS token\n"); - embd_inp.push_back(llama_vocab_bos(vocab)); - } - - std::string buffer; - if (!params.input_prefix.empty()) { - LOG_DBG("appending input prefix: '%s'\n", params.input_prefix.c_str()); - buffer += params.input_prefix; - LOG("%s", buffer.c_str()); - } - - std::string line; - bool another_line = true; - do { - another_line = console::readline(line, params.multiline_input); - buffer += line; - } while (another_line); - - // done taking input, reset color - console::set_display(console::reset); - - // Add tokens to embd only if the input buffer is non-empty - // Entering a empty line lets the user pass control back - if (buffer.length() > 1) { - // append input suffix if any - if (!params.input_suffix.empty()) { - LOG_DBG("appending input suffix: '%s'\n", params.input_suffix.c_str()); - buffer += params.input_suffix; - LOG("%s", params.input_suffix.c_str()); - } - - LOG_DBG("buffer: '%s'\n", buffer.c_str()); - - const size_t original_size = embd_inp.size(); - - const auto line_inp = common_tokenize(ctx, buffer, false); - LOG_DBG("input tokens: %s\n", string_from(ctx, line_inp).c_str()); - - embd_inp.insert(embd_inp.end(), line_inp.begin(), line_inp.end()); - - for (size_t i = original_size; i < embd_inp.size(); ++i) { - const llama_token token = embd_inp[i]; - output_tokens.push_back(token); - output_ss << common_token_to_piece(ctx, token); - } - - n_remain -= line_inp.size(); - LOG_DBG("n_remain: %d\n", n_remain); - } else { - LOG_DBG("empty line, passing control back\n"); - } - - input_echo = false; // do not echo this again - } - - if (n_past > 0) { - if (is_interacting) { - common_sampler_reset(smpl); - } - is_interacting = false; - } - } - - // end of generation - if (!embd.empty() && llama_vocab_is_eog(vocab, embd.back()) && !params.interactive) { - break; - } - - // In interactive mode, respect the maximum number of tokens and drop back to user input when reached. - // We skip this logic when n_predict == -1 (infinite) or -2 (stop at context size). - if (params.interactive && n_remain <= 0 && params.n_predict >= 0) { - n_remain = params.n_predict; - is_interacting = true; - } - } - if (!params.interactive && n_remain <= 0) { - LOG("%s", common_token_to_piece(ctx, llama_vocab_eot(vocab)).c_str()); - } - - LOG("\n"); - common_perf_print(ctx, smpl); - - common_sampler_free(smpl); - llama_backend_free(); - - return 0; -} From 1f73301b63668d61b5f3109489050a27dc3f65be Mon Sep 17 00:00:00 2001 From: R0CKSTAR Date: Wed, 7 May 2025 15:48:23 +0800 Subject: [PATCH 03/10] cuda : remove nrows_x in mul_mat_q_process_tile (#13325) Signed-off-by: Xiaodong Ye --- ggml/src/ggml-cuda/mmq.cuh | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index b8143a7b2..80baf459c 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -2522,7 +2522,7 @@ template static __device__ __forceinline__ void mul_mat_q_process_tile( const char * __restrict__ x, const int offset_x, const int * __restrict__ y, const int * __restrict__ ids_dst, float * __restrict__ dst, float * __restrict__ tmp_fixup, - const int nrows_x, const int stride_row_x, const int ncols_y, const int stride_col_dst, + const int stride_row_x, const int ncols_y, const int stride_col_dst, const int tile_x_max_i, const int tile_y_max_j, const int kb0_start, const int kb0_stop) { constexpr int qk = ggml_cuda_type_traits::qk; @@ -2689,7 +2689,7 @@ static __global__ void mul_mat_q( constexpr bool fixup = false; mul_mat_q_process_tile - (x, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, nrows_x, stride_row_x, ncols_y, stride_col_dst, + (x, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, stride_row_x, ncols_y, stride_col_dst, tile_x_max_i, tile_y_max_j, 0, ncols_x/qk); return; } @@ -2767,7 +2767,7 @@ static __global__ void mul_mat_q( constexpr bool fixup = false; // All but (potentially) the last iterations write their data to dst rather than the fixup buffer. mul_mat_q_process_tile - (x, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, nrows_x, stride_row_x, ncols_y, stride_col_dst, + (x, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, stride_row_x, ncols_y, stride_col_dst, tile_x_max_i, tile_y_max_j, kb0_start, kb0_stop); kbc += blocks_per_ne00; @@ -2834,7 +2834,7 @@ static __global__ void mul_mat_q( constexpr bool fixup = true; // Last index writes its data to fixup buffer to avoid data races with other blocks. mul_mat_q_process_tile - (x, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, nrows_x, stride_row_x, ncols_y, stride_col_dst, + (x, offset_x, y + offset_y, ids_dst_shared, dst + offset_dst, tmp_fixup, stride_row_x, ncols_y, stride_col_dst, tile_x_max_i, tile_y_max_j, kb0_start, kb0_stop); } From 39e73ae0d69f882d7e29cecc6dd8f5052fca6731 Mon Sep 17 00:00:00 2001 From: Ycros <18012+ycros@users.noreply.github.com> Date: Wed, 7 May 2025 18:23:28 +1000 Subject: [PATCH 04/10] common : Add a warning when we can't match samplers from a string or char. (#13330) --- common/sampling.cpp | 17 +++++++++++------ 1 file changed, 11 insertions(+), 6 deletions(-) diff --git a/common/sampling.cpp b/common/sampling.cpp index bbaec5b80..28705e24c 100644 --- a/common/sampling.cpp +++ b/common/sampling.cpp @@ -1,6 +1,7 @@ #include "sampling.h" #include "common.h" +#include "log.h" #include #include @@ -534,14 +535,16 @@ std::vector common_sampler_types_from_names(const std::vect auto sampler = sampler_canonical_name_map.find(name); if (sampler != sampler_canonical_name_map.end()) { samplers.push_back(sampler->second); - } else { - if (allow_alt_names) { - sampler = sampler_alt_name_map.find(name); - if (sampler != sampler_alt_name_map.end()) { - samplers.push_back(sampler->second); - } + continue; + } + if (allow_alt_names) { + sampler = sampler_alt_name_map.find(name); + if (sampler != sampler_alt_name_map.end()) { + samplers.push_back(sampler->second); + continue; } } + LOG_WRN("%s: unable to match sampler by name '%s'\n", __func__, name.c_str()); } return samplers; @@ -568,6 +571,8 @@ std::vector common_sampler_types_from_chars(const std::stri const auto sampler = sampler_name_map.find(c); if (sampler != sampler_name_map.end()) { samplers.push_back(sampler->second); + } else { + LOG_WRN("%s: unable to match sampler by char '%c'\n", __func__, c); } } From bc4e1128f78be0fbb4e2fa630adb6a04b969ac68 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Sigbj=C3=B8rn=20Skj=C3=A6ret?= Date: Wed, 7 May 2025 12:49:27 +0200 Subject: [PATCH 05/10] llama : deci : support ffn-free with attention (#13296) --- src/llama-model.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/llama-model.cpp b/src/llama-model.cpp index 774e343fb..7b66e2a03 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -4792,7 +4792,7 @@ struct llm_build_deci : public llm_graph_context { } // FFN-free layer of Llama-3_1-Nemotron-Ultra-253B - if (n_head == 0 && n_ff == 0) { + if (n_ff == 0) { continue; } From bba9d945c14b93c5264b7956e00736601ca6f89a Mon Sep 17 00:00:00 2001 From: Jared Tweed Date: Fri, 2 May 2025 02:41:35 -0700 Subject: [PATCH 06/10] cmake : removed stdc++fs (whisper/3097) * removed stdc++fs * kept line, but removed stdc++fs --- ggml/src/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/CMakeLists.txt b/ggml/src/CMakeLists.txt index 43d9fc4fe..ddea5ad38 100644 --- a/ggml/src/CMakeLists.txt +++ b/ggml/src/CMakeLists.txt @@ -214,7 +214,7 @@ add_library(ggml target_link_libraries(ggml PUBLIC ggml-base) if (CMAKE_SYSTEM_NAME MATCHES "Linux") - target_link_libraries(ggml PRIVATE dl stdc++fs) + target_link_libraries(ggml PRIVATE dl) endif() function(ggml_add_backend_library backend) From 13b0a04597a4581cad4d9027a848f450c623801d Mon Sep 17 00:00:00 2001 From: Daniel Bevenius Date: Mon, 5 May 2025 13:09:35 +0200 Subject: [PATCH 07/10] whisper: remove MSVC warnings pragmas (whisper/3090) * ggml : remove MSVC warnings pragmas This commit removes the MSVC-specific pragmas as these are now handled in ggml/CMakeLists.txt. * whisper : remove MSVC warning pragmas This commit removes the MSVC-specific pragmas. These are now handled in the ggml/CMakeLists.txt file. --- ggml/CMakeLists.txt | 2 ++ ggml/src/ggml-cpu/ggml-cpu-aarch64.cpp | 2 -- ggml/src/ggml-cpu/ggml-cpu-quants.c | 6 ------ ggml/src/ggml-cpu/ggml-cpu.c | 13 ------------- ggml/src/ggml-cpu/ops.cpp | 13 ------------- ggml/src/ggml-cpu/vec.cpp | 6 ------ ggml/src/ggml-cuda/common.cuh | 4 ---- ggml/src/ggml-quants.c | 6 ------ ggml/src/ggml-sycl/common.hpp | 4 ---- 9 files changed, 2 insertions(+), 54 deletions(-) diff --git a/ggml/CMakeLists.txt b/ggml/CMakeLists.txt index b463cbd9b..a8300e16d 100644 --- a/ggml/CMakeLists.txt +++ b/ggml/CMakeLists.txt @@ -366,6 +366,8 @@ if (MSVC) /wd4005 # Macro redefinition /wd4244 # Conversion from one type to another type, possible loss of data /wd4267 # Conversion from 'size_t' to a smaller type, possible loss of data + /wd4996 # Disable POSIX deprecation warnings + /wd4702 # Unreachable code warnings ) function(disable_msvc_warnings target_name) if(TARGET ${target_name}) diff --git a/ggml/src/ggml-cpu/ggml-cpu-aarch64.cpp b/ggml/src/ggml-cpu/ggml-cpu-aarch64.cpp index 175cba329..8ff6d64a4 100644 --- a/ggml/src/ggml-cpu/ggml-cpu-aarch64.cpp +++ b/ggml/src/ggml-cpu/ggml-cpu-aarch64.cpp @@ -72,8 +72,6 @@ static_assert(sizeof(block_iq4_nlx4) == 4 * sizeof(ggml_half) + QK4_NL * 2, "wro #if defined(__GNUC__) #pragma GCC diagnostic ignored "-Woverlength-strings" -#elif defined(_MSC_VER) -#pragma warning(disable: 4244 4267) // possible loss of data #endif #define UNUSED GGML_UNUSED diff --git a/ggml/src/ggml-cpu/ggml-cpu-quants.c b/ggml/src/ggml-cpu/ggml-cpu-quants.c index 987e15316..ccd0651eb 100644 --- a/ggml/src/ggml-cpu/ggml-cpu-quants.c +++ b/ggml/src/ggml-cpu/ggml-cpu-quants.c @@ -20,12 +20,6 @@ #define GROUP_MAX_EPS_IQ1_M 1e-7f #define GROUP_MAX_EPS_IQ1_S 1e-12f -#if defined(_MSC_VER) -// disable "possible loss of data" to avoid warnings for hundreds of casts -// we should just be careful :) -#pragma warning(disable: 4244 4267) -#endif - #define UNUSED GGML_UNUSED // some compilers don't provide _mm256_set_m128i, e.g. gcc 7 diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 64405449e..a30e67f22 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -50,19 +50,6 @@ #include "llamafile/sgemm.h" #endif -#if defined(_MSC_VER) -// disable "possible loss of data" to avoid hundreds of casts -// we should just be careful :) -#pragma warning(disable: 4244 4267) - -// disable POSIX deprecation warnings -// these functions are never going away, anyway -#pragma warning(disable: 4996) - -// unreachable code because of multiple instances of code after GGML_ABORT -#pragma warning(disable: 4702) -#endif - // Note: once we move threading into a separate C++ file // will use std::hardware_destructive_interference_size instead of hardcoding it here // and we'll use C++ attribute syntax. diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index 7413192b7..955fec59a 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -8,19 +8,6 @@ #include -#if defined(_MSC_VER) -// disable "possible loss of data" to avoid hundreds of casts -// we should just be careful :) -#pragma warning(disable: 4244 4267) - -// disable POSIX deprecation warnings -// these functions are never going away, anyway -#pragma warning(disable: 4996) - -// unreachable code because of multiple instances of code after GGML_ABORT -#pragma warning(disable: 4702) -#endif - // ggml_compute_forward_dup static void ggml_compute_forward_dup_same_cont( diff --git a/ggml/src/ggml-cpu/vec.cpp b/ggml/src/ggml-cpu/vec.cpp index dfe2218e3..02d406182 100644 --- a/ggml/src/ggml-cpu/vec.cpp +++ b/ggml/src/ggml-cpu/vec.cpp @@ -2,12 +2,6 @@ #include -#if defined(_MSC_VER) -// disable "possible loss of data" to avoid hundreds of casts -// we should just be careful :) -#pragma warning(disable: 4244 4267) -#endif - // precomputed gelu table for f16 (128 KB) ggml_fp16_t ggml_table_gelu_f16[1 << 16]; diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 2ea014e64..919217dfa 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -130,10 +130,6 @@ static int ggml_cuda_highest_compiled_arch(const int arch) { #define MATRIX_ROW_PADDING 512 // last row of quant. matrices is a multiple of this to avoid out-of-bounds memory accesses -#if defined(_MSC_VER) -#pragma warning(disable: 4244 4267) // possible loss of data -#endif - #define GGML_CUDA_MAX_STREAMS 8 [[noreturn]] diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index ac918a60d..84ec6dfe3 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -19,12 +19,6 @@ #define GROUP_MAX_EPS_IQ1_M 1e-7f #define GROUP_MAX_EPS_IQ1_S 1e-12f -#if defined(_MSC_VER) -// disable "possible loss of data" to avoid warnings for hundreds of casts -// we should just be careful :) -#pragma warning(disable: 4244 4267) -#endif - #define UNUSED GGML_UNUSED // reference implementation for deterministic creation of model files diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp index c3d9d1864..c71cc89c0 100644 --- a/ggml/src/ggml-sycl/common.hpp +++ b/ggml/src/ggml-sycl/common.hpp @@ -80,10 +80,6 @@ extern int g_ggml_sycl_disable_optimize; // max batch size to use MMQ kernels when tensor cores are available #define MMQ_MAX_BATCH_SIZE 32 -#if defined(_MSC_VER) -#pragma warning(disable : 4244 4267) // possible loss of data -#endif - // dmmv = dequantize_mul_mat_vec #ifndef GGML_SYCL_DMMV_X #define GGML_SYCL_DMMV_X 32 From d879433824ac3c16b3b5f00075895d0ee9688e34 Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Wed, 7 May 2025 16:39:36 +0300 Subject: [PATCH 08/10] sync : ggml ggml-ci --- scripts/sync-ggml.last | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/scripts/sync-ggml.last b/scripts/sync-ggml.last index 3e93dbf8a..1f7c650c2 100644 --- a/scripts/sync-ggml.last +++ b/scripts/sync-ggml.last @@ -1 +1 @@ -0482de9c63b9134eb462c7732888c0ee0dbc2755 +b59bddafe278877dfa22a80e53a637513862babb From 814f795e063c257f33b921eab4073484238a151a Mon Sep 17 00:00:00 2001 From: Diego Devesa Date: Wed, 7 May 2025 16:36:33 +0200 Subject: [PATCH 09/10] docker : disable arm64 and intel images (#13356) --- .github/workflows/docker.yml | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/.github/workflows/docker.yml b/.github/workflows/docker.yml index 114cbf83e..f8e072fb8 100644 --- a/.github/workflows/docker.yml +++ b/.github/workflows/docker.yml @@ -36,10 +36,14 @@ jobs: matrix: config: # Multi-stage build - - { tag: "cpu", dockerfile: ".devops/cpu.Dockerfile", platforms: "linux/amd64,linux/arm64", full: true, light: true, server: true, free_disk_space: false } + # Note: the arm64 images are failing, which prevents the amd64 images from being built + # https://github.com/ggml-org/llama.cpp/issues/11888 + #- { tag: "cpu", dockerfile: ".devops/cpu.Dockerfile", platforms: "linux/amd64,linux/arm64", full: true, light: true, server: true, free_disk_space: false } + - { tag: "cpu", dockerfile: ".devops/cpu.Dockerfile", platforms: "linux/amd64", full: true, light: true, server: true, free_disk_space: false } - { tag: "cuda", dockerfile: ".devops/cuda.Dockerfile", platforms: "linux/amd64", full: true, light: true, server: true, free_disk_space: false } - { tag: "musa", dockerfile: ".devops/musa.Dockerfile", platforms: "linux/amd64", full: true, light: true, server: true, free_disk_space: true } - - { tag: "intel", dockerfile: ".devops/intel.Dockerfile", platforms: "linux/amd64", full: true, light: true, server: true, free_disk_space: false } + # Note: the intel images are failing due to an out of disk space error + # - { tag: "intel", dockerfile: ".devops/intel.Dockerfile", platforms: "linux/amd64", full: true, light: true, server: true, free_disk_space: false } - { tag: "vulkan", dockerfile: ".devops/vulkan.Dockerfile", platforms: "linux/amd64", full: true, light: true, server: true, free_disk_space: false } # Note: the rocm images are failing due to a compiler error and are disabled until this is fixed to allow the workflow to complete #- {tag: "rocm", dockerfile: ".devops/rocm.Dockerfile", platforms: "linux/amd64,linux/arm64", full: true, light: true, server: true, free_disk_space: true } From 8733e0cf6eefc7c7752297cc22d0836706f4222c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Alberto=20Cabrera=20P=C3=A9rez?= Date: Thu, 8 May 2025 10:08:01 +0100 Subject: [PATCH 10/10] sycl: addressing non-contiguous src1 mul_mats (nc and batched) (#13343) * sycl: fixed non-contiguous src1 mul_mats (nc and batched) * Fixed wrong static_cast inside kernel --- ggml/src/ggml-sycl/common.hpp | 17 ++-- ggml/src/ggml-sycl/convert.cpp | 66 ++++++++----- ggml/src/ggml-sycl/convert.hpp | 21 ++-- ggml/src/ggml-sycl/ggml-sycl.cpp | 159 +++++++++++++++---------------- 4 files changed, 138 insertions(+), 125 deletions(-) diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp index c71cc89c0..69aad938e 100644 --- a/ggml/src/ggml-sycl/common.hpp +++ b/ggml/src/ggml-sycl/common.hpp @@ -114,17 +114,12 @@ static void crash() { GGML_ABORT("SYCL error"); } -#define SYCL_CHECK(err) \ - do { \ - auto err_ = (err); \ - if (err_ != 0) \ - ggml_sycl_error( \ - #err, \ - __func__, \ - __FILE__, \ - __LINE__, \ - "Meet error in this line code!"); \ - } while (0) +#define SYCL_CHECK(err) \ + do { \ + auto err_ = (err); \ + if (err_ != 0) \ + ggml_sycl_error(#err, __func__, __FILE__, __LINE__, "Exception caught in this line of code."); \ + } while (0) #if DPCT_COMPAT_RT_VERSION >= 11100 #define GGML_SYCL_ASSUME(x) __builtin_assume(x) diff --git a/ggml/src/ggml-sycl/convert.cpp b/ggml/src/ggml-sycl/convert.cpp index 76ac6a4dd..b2f8a6569 100644 --- a/ggml/src/ggml-sycl/convert.cpp +++ b/ggml/src/ggml-sycl/convert.cpp @@ -437,41 +437,52 @@ static void dequantize_row_iq4_nl_sycl(const void *vx, dst_t *y, const int64_t k } template -static void convert_unary(const void * __restrict__ vx, dst_t * __restrict__ y, const int64_t k, - const sycl::nd_item<3> &item_ct1) { +static void convert_unary_nc(const void * __restrict__ vx, dst_t * __restrict__ y, const int64_t ne00, const int64_t ne01, + const int64_t ne02, const int64_t s01, const int64_t s02, const int64_t s03, + const sycl::nd_item<3> & item_ct1) { + const int64_t work_group_size = item_ct1.get_local_range(2); - const int64_t global_id = item_ct1.get_local_id(2) + work_group_size * item_ct1.get_group(2); + const int64_t global_id = item_ct1.get_local_id(2) + work_group_size * item_ct1.get_group(2); + + const int64_t i01 = item_ct1.get_group(1); + const int64_t i02 = item_ct1.get_group(0) % ne02; + const int64_t i03 = item_ct1.get_group(0) / ne02; // make each work-item deal with more elements since sycl global range can not exceed max int - const src_t * x = (const src_t *) vx; - for (int64_t i = global_id; i < k; i += work_group_size * item_ct1.get_group_range(2)) { - y[i] = x[i]; + const src_t * x = static_cast(vx); + const int64_t ix = i03 * s03 + i02 * s02 + i01 * s01; + const int64_t iy = ((i03 * ne02 + i02) * ne01 + i01) * ne00; + +#pragma unroll + for (int64_t i00 = global_id; i00 < ne00; i00 += work_group_size * item_ct1.get_group_range(2)) { + y[iy + i00] = static_cast(x[ix + i00]); } } template -static void convert_unary_sycl(const void *__restrict__ vx, - dst_t *__restrict__ y, const int64_t k, - dpct::queue_ptr stream) { - const int64_t num_blocks = (k + SYCL_DEQUANTIZE_BLOCK_SIZE - 1) / SYCL_DEQUANTIZE_BLOCK_SIZE; +static void convert_unary_nc_sycl(const void * __restrict__ vx, dst_t * __restrict__ y, + const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03, + const int64_t s01, const int64_t s02, const int64_t s03, dpct::queue_ptr queue) { + dpct::has_capability_or_fail(queue->get_device(), { sycl::aspect::fp16 }); + + sycl::range<3> global_size(ne02 * ne03, ne01, ceil_div(ne00, SYCL_DEQUANTIZE_BLOCK_SIZE)); // decrease global range when it exceeds the max int - int64_t local_size = downsample_sycl_global_range(num_blocks, SYCL_DEQUANTIZE_BLOCK_SIZE); - sycl::range<3> block_nums(1, 1, num_blocks); - sycl::range<3> local_range(1, 1, local_size); - { - dpct::has_capability_or_fail(stream->get_device(), - {sycl::aspect::fp16}); + // TODO: Downsample logic is separated from the kernel, a rewrite is desirable + int64_t downsized_workgroup = downsample_sycl_global_range(global_size[0], SYCL_DEQUANTIZE_BLOCK_SIZE); + sycl::range<3> workgroup_size(1, 1, downsized_workgroup); - stream->parallel_for( - sycl::nd_range<3>(block_nums * local_range, local_range), - [=](sycl::nd_item<3> item_ct1) { - convert_unary(vx, y, k, item_ct1); - }); - } + queue->parallel_for(sycl::nd_range<3>(global_size * workgroup_size, workgroup_size), [=](sycl::nd_item<3> item_ct1) { + convert_unary_nc(vx, y, ne00, ne01, ne02, s01, s02, s03, item_ct1); + }); } -to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor *dst) { +template +static void convert_unary_sycl(const void * vx, dst_t * y, const int64_t k, dpct::queue_ptr queue) { + convert_unary_nc_sycl(vx, y, k, 1, 1, 1, k, k, k, queue); +} + +to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) { switch (type) { case GGML_TYPE_Q4_0: if (dst->src[0]->extra && @@ -574,3 +585,12 @@ to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor *dst) { return nullptr; } } + +to_fp16_nc_sycl_t get_to_fp16_nc_sycl(ggml_type type) { + switch (type) { + case GGML_TYPE_F32: + return convert_unary_nc_sycl; + default: + return nullptr; + } +} diff --git a/ggml/src/ggml-sycl/convert.hpp b/ggml/src/ggml-sycl/convert.hpp index 355dae22b..f8cb573e3 100644 --- a/ggml/src/ggml-sycl/convert.hpp +++ b/ggml/src/ggml-sycl/convert.hpp @@ -1,6 +1,6 @@ // // MIT license -// Copyright (C) 2024 Intel Corporation +// Copyright (C) 2025 Intel Corporation // SPDX-License-Identifier: MIT // @@ -16,12 +16,19 @@ #include "common.hpp" template -using to_t_sycl_t = void (*)(const void *__restrict__ x, T *__restrict__ y, - int64_t k, dpct::queue_ptr stream); -typedef to_t_sycl_t to_fp32_sycl_t; +using to_t_sycl_t = void (*)(const void * __restrict__ x, T * __restrict__ y, int64_t k, dpct::queue_ptr stream); +typedef to_t_sycl_t to_fp32_sycl_t; typedef to_t_sycl_t to_fp16_sycl_t; -to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor *dst); -to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor *dst); +to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst); +to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor * dst); -#endif // GGML_SYCL_CONVERT_HPP +// Nc = Non-contiguous +template +using to_t_nc_sycl_t = void (*)(const void * x, T * y, int64_t ne00, int64_t ne01, int64_t ne02, int64_t ne03, + int64_t s01, int64_t s02, int64_t s03, dpct::queue_ptr queue); + +typedef to_t_nc_sycl_t to_fp16_nc_sycl_t; +to_fp16_nc_sycl_t get_to_fp16_nc_sycl(ggml_type type); + +#endif // GGML_SYCL_CONVERT_HPP diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index ea5d10f40..68a26fa48 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -2694,35 +2694,31 @@ catch (sycl::exception const &exc) { std::exit(1); } -static void k_compute_batched_ptrs(const sycl::half *src0_as_f16, - const sycl::half *src1_as_f16, char *dst, - const void **ptrs_src, void **ptrs_dst, - int64_t ne12, int64_t ne13, int64_t ne23, - size_t nb02, size_t nb03, size_t nb12, - size_t nb13, size_t nbd2, size_t nbd3, - int64_t r2, int64_t r3, - const sycl::nd_item<3> &item_ct1) { - int64_t i13 = item_ct1.get_group(2) * item_ct1.get_local_range(2) + - item_ct1.get_local_id(2); - int64_t i12 = item_ct1.get_group(1) * item_ct1.get_local_range(1) + - item_ct1.get_local_id(1); +static void k_compute_batched_ptrs(const sycl::half * src0_as_f16, const sycl::half * src1_as_f16, char * dst, + const void ** ptrs_src, void ** ptrs_dst, int64_t ne12, int64_t ne13, int64_t ne23, + size_t nb02, size_t nb03, size_t nb12, size_t nb13, size_t nbd2, size_t nbd3, + int64_t r2, int64_t r3, const sycl::nd_item<3> & item_ct1) { + const int64_t i13 = item_ct1.get_group(2) * item_ct1.get_local_range(2) + item_ct1.get_local_id(2); + const int64_t i12 = item_ct1.get_group(1) * item_ct1.get_local_range(1) + item_ct1.get_local_id(1); if (i13 >= ne13 || i12 >= ne12) { return; } - int64_t i03 = i13 / r3; - int64_t i02 = i12 / r2; + const int64_t i03 = i13 / r3; + const int64_t i02 = i12 / r2; - ptrs_src[0*ne23 + i12 + i13*ne12] = (const char *) src0_as_f16 + i02*nb02 + i03*nb03; - ptrs_src[1*ne23 + i12 + i13*ne12] = (const char *) src1_as_f16 + i12*nb12 + i13*nb13; - ptrs_dst[0*ne23 + i12 + i13*ne12] = ( char *) dst + i12*nbd2 + i13*nbd3; + const uint8_t * src0_bytes = reinterpret_cast(src0_as_f16); + const uint8_t * src1_bytes = reinterpret_cast(src1_as_f16); + uint8_t * dst_bytes = reinterpret_cast(dst); + + ptrs_src[0 * ne23 + i12 + i13 * ne12] = src0_bytes + i02 * nb02 + i03 * nb03; + ptrs_src[1 * ne23 + i12 + i13 * ne12] = src1_bytes + i12 * nb12 + i13 * nb13; + ptrs_dst[0 * ne23 + i12 + i13 * ne12] = dst_bytes + i12 * nbd2 + i13 * nbd3; } -static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, - const ggml_tensor *src0, - const ggml_tensor *src1, - ggml_tensor *dst) try { +static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, + const ggml_tensor * src1, ggml_tensor * dst) try { GGML_ASSERT(!ggml_is_transposed(src0)); GGML_ASSERT(!ggml_is_transposed(src1)); GGML_ASSERT(!ggml_backend_buffer_is_sycl_split(src0->buffer)); @@ -2730,102 +2726,100 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, GGML_TENSOR_BINARY_OP_LOCALS + // TODO: see https://github.com/ggml-org/llama.cpp/pull/13155 + // Batched mul_mat requires a rewrite to support both oneDNN and non-contiguous dst + GGML_ASSERT(ggml_is_contiguous(dst)); SYCL_CHECK(ggml_sycl_set_device(ctx.device)); - queue_ptr main_stream = ctx.stream();; + queue_ptr queue = ctx.stream(); - void * src0_ddq = src0->data; - sycl::half *src0_as_f16 = (sycl::half *)src0_ddq; - float * src1_ddf = (float *) src1->data; - float * dst_ddf = (float *) dst->data; + dpct::has_capability_or_fail(queue->get_device(), { sycl::aspect::fp16 }); + + const sycl::half * src0_f16 = static_cast(src0->data); + float * dst_ddf = static_cast(dst->data); + + const sycl::half * src1_f16 = static_cast(src1->data); + const size_t type_size_src1 = ggml_type_size(src1->type); + GGML_ASSERT(nb10 == type_size_src1); + + // SRC1 strides + int64_t s11 = nb11 / type_size_src1; + int64_t s12 = nb12 / type_size_src1; + int64_t s13 = nb13 / type_size_src1; + ggml_sycl_pool_alloc src1_f16_alloc(ctx.pool()); // convert src1 to fp16 - ggml_sycl_pool_alloc src1_f16_alloc(ctx.pool()); if (src1->type != GGML_TYPE_F16) { - const to_fp16_sycl_t to_fp16_sycl = ggml_get_to_fp16_sycl(src1->type, dst); + const to_fp16_nc_sycl_t to_fp16_nc_sycl = get_to_fp16_nc_sycl(src1->type); + GGML_ASSERT(to_fp16_nc_sycl != nullptr); const int64_t ne_src1 = ggml_nelements(src1); src1_f16_alloc.alloc(ne_src1); - GGML_ASSERT(to_fp16_sycl != nullptr); - to_fp16_sycl(src1_ddf, src1_f16_alloc.get(), ne_src1, main_stream); + to_fp16_nc_sycl(src1_f16, src1_f16_alloc.get(), ne10, ne11, ne12, ne13, s11, s12, s13, queue); + + src1_f16 = src1_f16_alloc.get(); + s11 = ne10; + s12 = ne11 * s11; + s13 = ne12 * s12; } - sycl::half *src1_f16 = src1->type == GGML_TYPE_F16 ? (sycl::half *)src1_ddf - : src1_f16_alloc.get(); - char * dst_t; + ggml_sycl_pool_alloc dst_f16(ctx.pool()); + char * dst_t = reinterpret_cast(dst_ddf); - dpct::library_data_t cu_compute_type = dpct::library_data_t::real_float; - dpct::library_data_t cu_data_type = dpct::library_data_t::real_float; + dpct::library_data_t mkl_compute_type = dpct::library_data_t::real_float; + dpct::library_data_t mkl_data_type = dpct::library_data_t::real_float; // dst strides size_t nbd2 = dst->nb[2]; size_t nbd3 = dst->nb[3]; const float alpha_f32 = 1.0f; - const float beta_f32 = 0.0f; + const float beta_f32 = 0.0f; const void * alpha = &alpha_f32; const void * beta = &beta_f32; - dst_t = (char *) dst_ddf; - GGML_ASSERT(ne12 % ne02 == 0); GGML_ASSERT(ne13 % ne03 == 0); // broadcast factors - const int64_t r2 = ne12/ne02; - const int64_t r3 = ne13/ne03; + const int64_t r2 = ne12 / ne02; + const int64_t r3 = ne13 / ne03; if (r2 == 1 && r3 == 1 && ggml_is_contiguous_2(src0) && ggml_is_contiguous_2(src1)) { // there is no broadcast and src0, src1 are contiguous across dims 2, 3 - SYCL_CHECK(CHECK_TRY_ERROR(dpct::gemm_batch( - *main_stream, oneapi::math::transpose::trans, oneapi::math::transpose::nontrans, ne01, ne11, ne10, alpha, - (const char *) src0_as_f16, dpct::library_data_t::real_half, nb01 / nb00, nb02 / nb00, - (const char *) src1_f16, dpct::library_data_t::real_half, nb11 / nb10, nb12 / nb10, beta, (char *) dst_t, - cu_data_type, ne01, nb2 / nb0, ne12 * ne13, cu_compute_type))); + SYCL_CHECK(CHECK_TRY_ERROR(dpct::gemm_batch(*queue, oneapi::math::transpose::trans, + oneapi::math::transpose::nontrans, ne01, ne11, ne10, alpha, + src0_f16, dpct::library_data_t::real_half, nb01 / nb00, nb02 / nb00, + src1_f16, dpct::library_data_t::real_half, s11, s12, beta, dst_t, + mkl_data_type, ne0, ne1 * ne0, ne12 * ne13, mkl_compute_type))); } else { - const int ne23 = ne12*ne13; + const int ne23 = ne12 * ne13; - ggml_sycl_pool_alloc ptrs_src(ctx.pool(), 2*ne23); - ggml_sycl_pool_alloc< void *> ptrs_dst(ctx.pool(), 1*ne23); + ggml_sycl_pool_alloc ptrs_src(ctx.pool(), 2 * ne23); + ggml_sycl_pool_alloc ptrs_dst(ctx.pool(), 1 * ne23); ggml_sycl_pool_alloc> matrix_info(ctx.host_pool(), 1); sycl::range<3> block_dims(1, ne12, ne13); - /* - DPCT1049:47: The work-group size passed to the SYCL kernel may exceed - the limit. To get the device limit, query - info::device::max_work_group_size. Adjust the work-group size if needed. - */ - { - dpct::has_capability_or_fail(main_stream->get_device(), - {sycl::aspect::fp16}); - - main_stream->submit([&](sycl::handler &cgh) { - const void **ptrs_src_get = ptrs_src.get(); - void **ptrs_dst_get = ptrs_dst.get(); - size_t nb12_scaled = src1->type == GGML_TYPE_F16 ? nb12 : nb12 / 2; - size_t nb13_scaled = src1->type == GGML_TYPE_F16 ? nb13 : nb13 / 2; - cgh.parallel_for(sycl::nd_range<3>(block_dims, block_dims), - [=](sycl::nd_item<3> item_ct1) { - k_compute_batched_ptrs( - src0_as_f16, src1_f16, - dst_t, ptrs_src_get, - ptrs_dst_get, ne12, ne13, ne23, - nb02, nb03, nb12_scaled, nb13_scaled, - nbd2, nbd3, r2, r3, item_ct1); - }); + queue->submit([&](sycl::handler & cgh) { + const void ** ptrs_src_get = ptrs_src.get(); + void ** ptrs_dst_get = ptrs_dst.get(); + size_t nb12_scaled = src1->type == GGML_TYPE_F16 ? nb12 : s12 * sizeof(sycl::half); + size_t nb13_scaled = src1->type == GGML_TYPE_F16 ? nb13 : s13 * sizeof(sycl::half); + cgh.parallel_for(sycl::nd_range<3>(block_dims, block_dims), [=](sycl::nd_item<3> item_ct1) { + k_compute_batched_ptrs(src0_f16, src1_f16, dst_t, ptrs_src_get, ptrs_dst_get, ne12, ne13, ne23, nb02, + nb03, nb12_scaled, nb13_scaled, nbd2, nbd3, r2, r3, item_ct1); }); - } + }); + SYCL_CHECK(CHECK_TRY_ERROR(dpct::gemm_batch( - *main_stream, oneapi::math::transpose::trans, oneapi::math::transpose::nontrans, ne01, ne11, ne10, alpha, + *queue, oneapi::math::transpose::trans, oneapi::math::transpose::nontrans, ne01, ne11, ne10, alpha, (const void **) (ptrs_src.get() + 0 * ne23), dpct::library_data_t::real_half, nb01 / nb00, - (const void **) (ptrs_src.get() + 1 * ne23), dpct::library_data_t::real_half, nb11 / nb10, beta, - (void **) (ptrs_dst.get() + 0 * ne23), cu_data_type, ne01, ne23, cu_compute_type, matrix_info.get()))); + (const void **) (ptrs_src.get() + 1 * ne23), dpct::library_data_t::real_half, s11, beta, + (void **) (ptrs_dst.get() + 0 * ne23), mkl_data_type, ne0, ne23, mkl_compute_type, matrix_info.get()))); } -} -catch (sycl::exception const &exc) { - std::cerr << exc.what() << "Exception caught at file:" << __FILE__ - << ", line:" << __LINE__ << std::endl; - std::exit(1); +} catch (const sycl::exception & exc) { + std::cerr << exc.what() << "Exception caught at file:" << __FILE__ << ", line:" << __LINE__ << std::endl; + std::exit(1); } inline bool ggml_sycl_supports_mmq(enum ggml_type type) { @@ -2966,7 +2960,7 @@ static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor // The kernel from the if path is faster for that specific case, but does not support all mul mats. ggml_sycl_mul_mat_batched_sycl(ctx, src0, src1, dst); } - } else if (!split && src0->type == GGML_TYPE_F16 && !ggml_is_contiguous(src0) && !ggml_is_transposed(src1) && src1->ne[1] == 1) { + } else if (!split && src0->type == GGML_TYPE_F16 && !ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && !ggml_is_transposed(src1) && src1->ne[1] == 1) { // KQV single-batch ggml_sycl_mul_mat_vec_nc(ctx, src0, src1, dst); } else if (!split && src0->type == GGML_TYPE_F16 && !ggml_is_transposed(src0) && !ggml_is_transposed(src1) && src1->ne[2]*src1->ne[3] > 1) { @@ -3873,9 +3867,6 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g if (a->ne[3] != b->ne[3]) { return false; } - if (!ggml_is_contiguous(b)) { - return false; - } ggml_type a_type = a->type; if (a_type == GGML_TYPE_IQ4_NL || a_type == GGML_TYPE_IQ4_XS || a_type == GGML_TYPE_IQ3_XXS || a_type == GGML_TYPE_IQ3_S ||