llama : change starcoder2 rope type

.github/workflows/build.yml

@@ -145,28 +145,6 @@ jobs:
           cd build
           ctest -L main --verbose
 
-  ubuntu-22-cmake-vulkan:
-    runs-on: ubuntu-22.04
-
-    steps:
-      - name: Clone
-        id: checkout
-        uses: actions/checkout@v3
-
-      - name: Dependencies
-        id: depends
-        run: |
-          sudo apt-get update
-          sudo apt-get install build-essential libvulkan-dev
-
-      - name: Build
-        id: cmake_build
-        run: |
-          mkdir build
-          cd build
-          cmake -DLLAMA_VULKAN=ON ..
-          cmake --build . --config Release -j $(nproc)
-
   ubuntu-22-cmake-sycl:
     runs-on: ubuntu-22.04
 

convert-hf-to-gguf.py

@@ -96,9 +96,11 @@ class Model:
         if (n_head_kv := self.hparams.get("num_key_value_heads")) is not None:
             self.gguf_writer.add_head_count_kv(n_head_kv)
 
+        if (rope_theta := self.hparams.get("rope_theta")) is not None:
+            self.gguf_writer.add_rope_freq_base(rope_theta)
         if (f_rms_eps := self.hparams.get("rms_norm_eps")) is not None:
             self.gguf_writer.add_layer_norm_rms_eps(f_rms_eps)
-        if (f_norm_eps := self.find_hparam(["layer_norm_eps", "layer_norm_epsilon"], optional=True)) is not None:
+        if (f_norm_eps := self.find_hparam(["layer_norm_eps", "layer_norm_epsilon", "norm_epsilon"], optional=True)) is not None:
             self.gguf_writer.add_layer_norm_eps(f_norm_eps)
         if (n_experts := self.hparams.get("num_local_experts")) is not None:
             self.gguf_writer.add_expert_count(n_experts)
@@ -281,6 +283,8 @@ class Model:
             return gguf.MODEL_ARCH.NOMIC_BERT
         if arch == "GemmaForCausalLM":
             return gguf.MODEL_ARCH.GEMMA
+        if arch == "Starcoder2ForCausalLM":
+            return gguf.MODEL_ARCH.STARCODER2
 
         raise NotImplementedError(f'Architecture "{arch}" not supported!')
 

examples/server/README.md

@@ -18,7 +18,6 @@ The project is under active development, and we are [looking for feedback and co
 
 - `--threads N`, `-t N`: Set the number of threads to use during generation.
 - `-tb N, --threads-batch N`: Set the number of threads to use during batch and prompt processing. If not specified, the number of threads will be set to the number of threads used for generation.
-- `--threads-http N`: number of threads in the http server pool to process requests (default: `std::thread::hardware_concurrency()`)
 - `-m FNAME`, `--model FNAME`: Specify the path to the LLaMA model file (e.g., `models/7B/ggml-model.gguf`).
 - `-a ALIAS`, `--alias ALIAS`: Set an alias for the model. The alias will be returned in API responses.
 - `-c N`, `--ctx-size N`: Set the size of the prompt context. The default is 512, but LLaMA models were built with a context of 2048, which will provide better results for longer input/inference. The size may differ in other models, for example, baichuan models were build with a context of 4096.

examples/server/server.cpp

@@ -33,7 +33,8 @@
 
 using json = nlohmann::json;
 
-struct server_params {
+struct server_params
+{
     std::string hostname = "127.0.0.1";
     std::vector<std::string> api_keys;
     std::string public_path = "examples/server/public";
@@ -43,56 +44,108 @@ struct server_params {
     int32_t write_timeout = 600;
     bool slots_endpoint = true;
     bool metrics_endpoint = false;
-    int n_threads_http = -1;
 };
 
 bool server_verbose = false;
 bool server_log_json = true;
 
-enum stop_type {
+static size_t common_part(const std::vector<llama_token> &a, const std::vector<llama_token> &b)
+{
+    size_t i;
+    for (i = 0; i < a.size() && i < b.size() && a[i] == b[i]; i++)
+    {
+    }
+    return i;
+}
+
+enum stop_type
+{
     STOP_FULL,
     STOP_PARTIAL,
 };
 
-// TODO: can become bool if we can't find use of more states
-enum slot_state {
-    IDLE,
-    PROCESSING,
-};
+static bool ends_with(const std::string &str, const std::string &suffix)
+{
+    return str.size() >= suffix.size() &&
+           0 == str.compare(str.size() - suffix.size(), suffix.size(), suffix);
+}
 
-enum slot_command {
-    NONE,
-    LOAD_PROMPT,
-    RELEASE,
-};
+static size_t find_partial_stop_string(const std::string &stop,
+                                       const std::string &text)
+{
+    if (!text.empty() && !stop.empty())
+    {
+        const char text_last_char = text.back();
+        for (int64_t char_index = stop.size() - 1; char_index >= 0; char_index--)
+        {
+            if (stop[char_index] == text_last_char)
+            {
+                const std::string current_partial = stop.substr(0, char_index + 1);
+                if (ends_with(text, current_partial))
+                {
+                    return text.size() - char_index - 1;
+                }
+            }
+        }
+    }
+    return std::string::npos;
+}
 
-struct slot_params {
-    bool stream       = true;
-    bool cache_prompt = false; // remember the prompt to avoid reprocessing all prompt
+// TODO: reuse llama_detokenize
+template <class Iter>
+static std::string tokens_to_str(llama_context *ctx, Iter begin, Iter end)
+{
+    std::string ret;
+    for (; begin != end; ++begin)
+    {
+        ret += llama_token_to_piece(ctx, *begin);
+    }
+    return ret;
+}
 
-    uint32_t seed      = -1; // RNG seed
-    int32_t  n_keep    =  0; // number of tokens to keep from initial prompt
-    int32_t  n_predict = -1; // new tokens to predict
+// format incomplete utf-8 multibyte character for output
+static std::string tokens_to_output_formatted_string(const llama_context *ctx, const llama_token token)
+{
+    std::string out = token == -1 ? "" : llama_token_to_piece(ctx, token);
+    // if the size is 1 and first bit is 1, meaning it's a partial character
+    //   (size > 1 meaning it's already a known token)
+    if (out.size() == 1 && (out[0] & 0x80) == 0x80)
+    {
+        std::stringstream ss;
+        ss << std::hex << (out[0] & 0xff);
+        std::string res(ss.str());
+        out = "byte: \\x" + res;
+    }
+    return out;
+}
 
-    std::vector<std::string> antiprompt;
+// convert a vector of completion_token_output to json
+static json probs_vector_to_json(const llama_context *ctx, const std::vector<completion_token_output> &probs)
+{
+    json out = json::array();
+    for (const auto &prob : probs)
+    {
+        json probs_for_token = json::array();
+        for (const auto &p : prob.probs)
+        {
+            std::string tok_str = tokens_to_output_formatted_string(ctx, p.tok);
+            probs_for_token.push_back(json
+            {
+                {"tok_str", tok_str},
+                {"prob",    p.prob},
+            });
+        }
+        std::string tok_str = tokens_to_output_formatted_string(ctx, prob.tok);
+        out.push_back(json{
+            {"content", tok_str},
+            {"probs",   probs_for_token},
+        });
+    }
+    return out;
+}
 
-    json input_prefix;
-    json input_suffix;
-};
-
-struct slot_image {
-    int32_t id;
-
-    bool request_encode_image = false;
-    float * image_embedding = nullptr;
-    int32_t image_tokens = 0;
-
-    clip_image_u8 * img_data;
-
-    std::string prefix_prompt; // before of this image
-};
-
-struct server_slot {
+struct llama_client_slot
+{
     int id;
     int task_id = -1;
 
@@ -112,8 +165,8 @@ struct server_slot {
     int32_t i_batch     = -1;
     int32_t n_predict   = -1;
 
-    int32_t n_prompt_tokens           = 0;
-    int32_t n_prompt_tokens_processed = 0;
+    int32_t num_prompt_tokens           = 0;
+    int32_t num_prompt_tokens_processed = 0;
 
     json prompt;
     std::string generated_text;
@@ -148,8 +201,8 @@ struct server_slot {
     std::vector<slot_image> images;
 
     // stats
-    size_t n_sent_text = 0; // number of sent text character
-    size_t n_sent_token_probs = 0;
+    size_t sent_count = 0;
+    size_t sent_token_probs_index = 0;
 
     int64_t t_start_process_prompt;
     int64_t t_start_genereration;
@@ -161,7 +214,7 @@ struct server_slot {
     int multitask_id = -1;
 
     void reset() {
-        n_prompt_tokens        = 0;
+        num_prompt_tokens      = 0;
         generated_text         = "";
         truncated              = false;
         stopped_eos            = false;
@@ -169,15 +222,16 @@ struct server_slot {
         stopped_limit          = false;
         stopping_word          = "";
         n_past                 = 0;
-        n_sent_text            = 0;
-        n_sent_token_probs     = 0;
+        sent_count             = 0;
+        sent_token_probs_index = 0;
         infill                 = false;
         ga_i                   = 0;
         n_past_se              = 0;
 
         generated_token_probs.clear();
 
-        for (slot_image & img : images) {
+        for (slot_image & img : images)
+        {
             free(img.image_embedding);
             if (img.img_data) {
                 clip_image_u8_free(img.img_data);
@@ -189,15 +243,19 @@ struct server_slot {
     }
 
     bool has_budget(gpt_params &global_params) {
-        if (params.n_predict == -1 && global_params.n_predict == -1) {
+        if (params.n_predict == -1 && global_params.n_predict == -1)
+        {
             return true; // limitless
         }
 
         n_remaining = -1;
 
-        if (params.n_predict != -1) {
+        if (params.n_predict != -1)
+        {
             n_remaining = params.n_predict - n_decoded;
-        } else if (global_params.n_predict != -1) {
+        }
+        else if (global_params.n_predict != -1)
+        {
             n_remaining = global_params.n_predict - n_decoded;
         }
 
@@ -213,7 +271,8 @@ struct server_slot {
     }
 
     void add_token_string(const completion_token_output &token) {
-        if (command == RELEASE) {
+        if (command == RELEASE)
+        {
             return;
         }
         cache_tokens.push_back(token.tok);
@@ -231,10 +290,10 @@ struct server_slot {
     json get_formated_timings() {
         return json
         {
-            {"prompt_n",               n_prompt_tokens_processed},
+            {"prompt_n",               num_prompt_tokens_processed},
             {"prompt_ms",              t_prompt_processing},
-            {"prompt_per_token_ms",    t_prompt_processing / n_prompt_tokens_processed},
-            {"prompt_per_second",      1e3 / t_prompt_processing * n_prompt_tokens_processed},
+            {"prompt_per_token_ms",    t_prompt_processing / num_prompt_tokens_processed},
+            {"prompt_per_second",      1e3 / t_prompt_processing * num_prompt_tokens_processed},
 
             {"predicted_n",            n_decoded},
             {"predicted_ms",           t_token_generation},
@@ -245,18 +304,18 @@ struct server_slot {
 
     void print_timings() const {
        char buffer[512];
-        double t_token = t_prompt_processing / n_prompt_tokens_processed;
-        double n_tokens_second = 1e3 / t_prompt_processing * n_prompt_tokens_processed;
+        double t_token = t_prompt_processing / num_prompt_tokens_processed;
+        double n_tokens_second = 1e3 / t_prompt_processing * num_prompt_tokens_processed;
         sprintf(buffer, "prompt eval time     = %10.2f ms / %5d tokens (%8.2f ms per token, %8.2f tokens per second)",
-                t_prompt_processing, n_prompt_tokens_processed,
+                t_prompt_processing, num_prompt_tokens_processed,
                 t_token, n_tokens_second);
         LOG_INFO(buffer, {
-            {"slot_id",                   id},
-            {"task_id",                   task_id},
-            {"t_prompt_processing",       t_prompt_processing},
-            {"n_prompt_tokens_processed", n_prompt_tokens_processed},
-            {"t_token",                   t_token},
-            {"n_tokens_second",           n_tokens_second},
+            {"slot_id",                     id},
+            {"task_id",                     task_id},
+            {"t_prompt_processing",         t_prompt_processing},
+            {"num_prompt_tokens_processed", num_prompt_tokens_processed},
+            {"t_token",                     t_token},
+            {"n_tokens_second",             n_tokens_second},
         });
 
         t_token = t_token_generation / n_decoded;
@@ -284,7 +343,7 @@ struct server_slot {
     }
 };
 
-struct server_metrics {
+struct llama_metrics {
     uint64_t n_prompt_tokens_processed_total = 0;
     uint64_t n_tokens_predicted_total        = 0;
 
@@ -295,16 +354,18 @@ struct server_metrics {
     uint64_t t_tokens_generation      = 0;
 
 
-    void on_prompt_eval(const server_slot &slot) {
-        n_prompt_tokens_processed_total += slot.n_prompt_tokens_processed;
-        n_prompt_tokens_processed       += slot.n_prompt_tokens_processed;
-        t_prompt_processing             += slot.t_prompt_processing;
+    void on_prompt_eval(const llama_client_slot &slot) {
+        n_prompt_tokens_processed_total += slot.num_prompt_tokens_processed;
+
+        n_prompt_tokens_processed += slot.num_prompt_tokens_processed;
+        t_prompt_processing       += slot.t_prompt_processing;
     }
 
-    void on_prediction(const server_slot &slot) {
+    void on_prediction(const llama_client_slot &slot) {
         n_tokens_predicted_total += slot.n_decoded;
-        n_tokens_predicted       += slot.n_decoded;
-        t_tokens_generation      += slot.t_token_generation;
+
+        n_tokens_predicted  += slot.n_decoded;
+        t_tokens_generation += slot.t_token_generation;
     }
 
     void reset_bucket() {
@@ -343,13 +404,13 @@ struct llama_server_context
     std::string name_assistant;
 
     // slots / clients
-    std::vector<server_slot> slots;
+    std::vector<llama_client_slot> slots;
     json default_generation_settings_for_props;
 
-    llama_server_queue    queue_tasks;
+    llama_server_queue queue_tasks;
     llama_server_response queue_results;
 
-    server_metrics metrics;
+    llama_metrics metrics;
 
     ~llama_server_context()
     {
@@ -426,7 +487,7 @@ struct llama_server_context
         LOG_INFO("initializing slots", {{"n_slots", params.n_parallel}});
         for (int i = 0; i < params.n_parallel; i++)
         {
-            server_slot slot;
+            llama_client_slot slot;
 
             slot.id = i;
             slot.n_ctx = n_ctx_slot;
@@ -518,11 +579,11 @@ struct llama_server_context
         return prompt_tokens;
     }
 
-    server_slot* get_slot(int id) {
+    llama_client_slot* get_slot(int id) {
         int64_t t_last = ggml_time_us();
-        server_slot *last_used = nullptr;
+        llama_client_slot *last_used = nullptr;
 
-        for (server_slot & slot : slots)
+        for (llama_client_slot & slot : slots)
         {
             if (slot.id == id && slot.available())
             {
@@ -539,7 +600,7 @@ struct llama_server_context
         return last_used;
     }
 
-    bool launch_slot_with_data(server_slot* &slot, json data) {
+    bool launch_slot_with_data(llama_client_slot* &slot, json data) {
         slot_params default_params;
         llama_sampling_params default_sparams;
 
@@ -827,7 +888,7 @@ struct llama_server_context
         clean_kv_cache = false;
     }
 
-    void system_prompt_update() {
+    void update_system_prompt() {
         kv_cache_clear();
         system_tokens.clear();
 
@@ -872,9 +933,9 @@ struct llama_server_context
         system_need_update = false;
     }
 
-    void system_prompt_notify() {
+    void notify_system_prompt_changed() {
         // release all slots
-        for (server_slot &slot : slots)
+        for (llama_client_slot &slot : slots)
         {
             slot.release();
         }
@@ -882,17 +943,17 @@ struct llama_server_context
         system_need_update = true;
     }
 
-    void system_prompt_process(const json &sys_props) {
+    void process_system_prompt_data(const json &sys_props) {
         system_prompt  = sys_props.value("prompt", "");
         name_user      = sys_props.value("anti_prompt", "");
         name_assistant = sys_props.value("assistant_name", "");
 
 
-        system_prompt_notify();
+        notify_system_prompt_changed();
     }
 
     static size_t find_stopping_strings(const std::string &text, const size_t last_token_size,
-                                        const stop_type type, server_slot &slot)
+                                        const stop_type type, llama_client_slot &slot)
     {
         size_t stop_pos = std::string::npos;
 
@@ -914,8 +975,8 @@ struct llama_server_context
             {
                 if (type == STOP_FULL)
                 {
-                    slot.stopped_word   = true;
-                    slot.stopping_word  = word;
+                    slot.stopped_word = true;
+                    slot.stopping_word = word;
                     slot.has_next_token = false;
                 }
                 stop_pos = pos;
@@ -925,7 +986,7 @@ struct llama_server_context
         return stop_pos;
     }
 
-    bool process_token(completion_token_output &result, server_slot &slot) {
+    bool process_token(completion_token_output &result, llama_client_slot &slot) {
         // remember which tokens were sampled - used for repetition penalties during sampling
         const std::string token_str = llama_token_to_piece(ctx, result.tok);
         slot.sampled = result.tok;
@@ -971,7 +1032,7 @@ struct llama_server_context
 
         if (!incomplete)
         {
-            size_t pos = std::min(slot.n_sent_text, slot.generated_text.size());
+            size_t pos = std::min(slot.sent_count, slot.generated_text.size());
             const std::string str_test = slot.generated_text.substr(pos);
             bool is_stop_full = false;
             size_t stop_pos = find_stopping_strings(str_test, token_str.size(), STOP_FULL, slot);
@@ -981,7 +1042,7 @@ struct llama_server_context
                 slot.generated_text.erase(
                     slot.generated_text.begin() + pos + stop_pos,
                     slot.generated_text.end());
-                pos = std::min(slot.n_sent_text, slot.generated_text.size());
+                pos = std::min(slot.sent_count, slot.generated_text.size());
             }
             else
             {
@@ -994,7 +1055,7 @@ struct llama_server_context
             {
                 // no send the stop word in the response
                 result.text_to_send = slot.generated_text.substr(pos, std::string::npos);
-                slot.n_sent_text += result.text_to_send.size();
+                slot.sent_count += result.text_to_send.size();
                 // add the token to slot queue and cache
             }
             slot.add_token_string(result);
@@ -1038,7 +1099,7 @@ struct llama_server_context
         return slot.has_next_token; // continue
     }
 
-    bool process_images(server_slot &slot) const
+    bool process_images(llama_client_slot &slot) const
     {
         for (slot_image &img : slot.images)
         {
@@ -1071,7 +1132,7 @@ struct llama_server_context
         queue_results.send(res);
     }
 
-    json get_formated_generation(server_slot &slot)
+    json get_formated_generation(llama_client_slot &slot)
     {
         const auto eos_bias = slot.sparams.logit_bias.find(llama_token_eos(model));
         const bool ignore_eos = eos_bias != slot.sparams.logit_bias.end() &&
@@ -1118,7 +1179,7 @@ struct llama_server_context
         };
     }
 
-    void send_partial_response(server_slot &slot, completion_token_output tkn)
+    void send_partial_response(llama_client_slot &slot, completion_token_output tkn)
     {
         task_result res;
         res.id = slot.task_id;
@@ -1138,13 +1199,13 @@ struct llama_server_context
         {
             std::vector<completion_token_output> probs_output = {};
             const std::vector<llama_token> to_send_toks = llama_tokenize(ctx, tkn.text_to_send, false);
-            size_t probs_pos      = std::min(slot.n_sent_token_probs,                       slot.generated_token_probs.size());
-            size_t probs_stop_pos = std::min(slot.n_sent_token_probs + to_send_toks.size(), slot.generated_token_probs.size());
+            size_t probs_pos      = std::min(slot.sent_token_probs_index,                       slot.generated_token_probs.size());
+            size_t probs_stop_pos = std::min(slot.sent_token_probs_index + to_send_toks.size(), slot.generated_token_probs.size());
             if (probs_pos < probs_stop_pos)
             {
                 probs_output = std::vector<completion_token_output>(slot.generated_token_probs.begin() + probs_pos, slot.generated_token_probs.begin() + probs_stop_pos);
             }
-            slot.n_sent_token_probs = probs_stop_pos;
+            slot.sent_token_probs_index = probs_stop_pos;
             res.result_json["completion_probabilities"] = probs_vector_to_json(ctx, probs_output);
         }
 
@@ -1157,7 +1218,7 @@ struct llama_server_context
         queue_results.send(res);
     }
 
-    void send_final_response(server_slot &slot)
+    void send_final_response(llama_client_slot &slot)
     {
         task_result res;
         res.id = slot.task_id;
@@ -1172,7 +1233,7 @@ struct llama_server_context
             {"stop",                true},
             {"model",               params.model_alias},
             {"tokens_predicted",    slot.n_decoded},
-            {"tokens_evaluated",    slot.n_prompt_tokens},
+            {"tokens_evaluated",    slot.num_prompt_tokens},
             {"generation_settings", get_formated_generation(slot)},
             {"prompt",              slot.prompt},
             {"truncated",           slot.truncated},
@@ -1210,7 +1271,7 @@ struct llama_server_context
         queue_results.send(res);
     }
 
-    void send_embedding(server_slot &slot)
+    void send_embedding(llama_client_slot &slot)
     {
         task_result res;
         res.id = slot.task_id;
@@ -1221,7 +1282,9 @@ struct llama_server_context
         const int n_embd = llama_n_embd(model);
         if (!params.embedding)
         {
-            LOG_WARNING("embedding disabled", {{"params.embedding", params.embedding}});
+            LOG_WARNING("embedding disabled", {
+                                                  {"params.embedding", params.embedding},
+                                              });
             res.result_json = json
             {
                 {"embedding", std::vector<float>(n_embd, 0.0f)},
@@ -1233,7 +1296,7 @@ struct llama_server_context
             std::vector<float> embedding(data, data + n_embd);
             res.result_json = json
             {
-                {"embedding", embedding},
+                {"embedding", embedding },
             };
         }
         queue_results.send(res);
@@ -1282,7 +1345,7 @@ struct llama_server_context
     }
 
     // for multiple images processing
-    bool ingest_images(server_slot &slot, int n_batch)
+    bool ingest_images(llama_client_slot &slot, int n_batch)
     {
         int image_idx = 0;
 
@@ -1321,17 +1384,7 @@ struct llama_server_context
                 }
 
                 const int n_embd = llama_n_embd(model);
-                llama_batch batch_img = {
-                    n_eval,
-                    nullptr,
-                    (img.image_embedding + i * n_embd),
-                    nullptr,
-                    nullptr,
-                    nullptr,
-                    nullptr,
-                    slot.n_past,
-                    1, 0
-                };
+                llama_batch batch_img = { n_eval, nullptr, (img.image_embedding + i * n_embd), nullptr, nullptr, nullptr, nullptr, slot.n_past, 1, 0, };
                 if (llama_decode(ctx, batch_img))
                 {
                     LOG_TEE("%s : failed to eval image\n", __func__);
@@ -1401,7 +1454,7 @@ struct llama_server_context
         switch (task.type)
         {
             case TASK_TYPE_COMPLETION: {
-                server_slot *slot = get_slot(json_value(task.data, "slot_id", -1));
+                llama_client_slot *slot = get_slot(json_value(task.data, "slot_id", -1));
                 if (slot == nullptr)
                 {
                     // if no slot is available, we defer this task for processing later
@@ -1416,10 +1469,10 @@ struct llama_server_context
                         send_error(task, "system prompt can only be updated when all slots are idle");
                         break;
                     }
-                    system_prompt_process(task.data["system_prompt"]);
+                    process_system_prompt_data(task.data["system_prompt"]);
 
                     // reset cache_tokens for all slots
-                    for (server_slot &slot : slots)
+                    for (llama_client_slot &slot : slots)
                     {
                         slot.cache_tokens.clear();
                         slot.n_past    = 0;
@@ -1459,20 +1512,20 @@ struct llama_server_context
                 int n_idle_slots       = 0;
                 int n_processing_slots = 0;
 
-                for (server_slot &slot: slots) {
+                for (llama_client_slot &slot: slots) {
                     json slot_data = get_formated_generation(slot);
                     slot_data["id"] = slot.id;
                     slot_data["task_id"] = slot.task_id;
                     slot_data["state"] = slot.state;
                     slot_data["prompt"] = slot.prompt;
                     slot_data["next_token"] = {
-                            {"has_next_token",       slot.has_next_token},
-                            {"n_remain",             slot.n_remaining},
+                            {"has_next_token", slot.has_next_token},
+                            {"n_remain", slot.n_remaining},
                             {"num_tokens_predicted", slot.n_decoded},
-                            {"stopped_eos",          slot.stopped_eos},
-                            {"stopped_word",         slot.stopped_word},
-                            {"stopped_limit",        slot.stopped_limit},
-                            {"stopping_word",        slot.stopping_word},
+                            {"stopped_eos", slot.stopped_eos},
+                            {"stopped_word", slot.stopped_word},
+                            {"stopped_limit", slot.stopped_limit},
+                            {"stopping_word", slot.stopping_word},
                     };
                     if (slot_data["state"] == IDLE) {
                         n_idle_slots++;
@@ -1510,10 +1563,10 @@ struct llama_server_context
                         { "n_tokens_predicted",              metrics.n_tokens_predicted},
                         { "t_tokens_generation",             metrics.t_tokens_generation},
 
-                        { "kv_cache_tokens_count",           llama_get_kv_cache_token_count(ctx)},
-                        { "kv_cache_used_cells",             llama_get_kv_cache_used_cells(ctx)},
+                        { "kv_cache_tokens_count",          llama_get_kv_cache_token_count(ctx)},
+                        { "kv_cache_used_cells",            llama_get_kv_cache_used_cells(ctx)},
 
-                        { "slots",                           slots_data },
+                        { "slots",                          slots_data },
                 };
                 metrics.reset_bucket();
                 queue_results.send(res);
@@ -1544,7 +1597,7 @@ struct llama_server_context
         if (system_need_update)
         {
             LOG_INFO("updating system prompt", {});
-            system_prompt_update();
+            update_system_prompt();
         }
 
         llama_batch_clear(batch);
@@ -1565,7 +1618,7 @@ struct llama_server_context
         task.target_id = -1;
         queue_tasks.post(task);
 
-        for (server_slot &slot : slots)
+        for (llama_client_slot &slot : slots)
         {
             if (slot.ga_n == 1)
             {
@@ -1701,50 +1754,45 @@ struct llama_server_context
                         prompt_tokens = tokenize(slot.prompt, system_prompt.empty() && add_bos_token);  // add BOS if there isn't system prompt
                     }
 
-                    slot.n_prompt_tokens = prompt_tokens.size();
+                    slot.num_prompt_tokens = prompt_tokens.size();
 
                     if (slot.params.n_keep < 0)
                     {
-                        slot.params.n_keep = slot.n_prompt_tokens;
+                        slot.params.n_keep = slot.num_prompt_tokens;
                     }
                     slot.params.n_keep = std::min(slot.n_ctx - 4, slot.params.n_keep);
 
                     // if input prompt is too big, truncate it
-                    if (slot.n_prompt_tokens >= slot.n_ctx)
+                    if (slot.num_prompt_tokens >= slot.n_ctx)
                     {
                         const int n_left = slot.n_ctx - slot.params.n_keep;
                         const int n_block_size = n_left / 2;
-                        const int erased_blocks = (slot.n_prompt_tokens - slot.params.n_keep - n_block_size) / n_block_size;
+                        const int erased_blocks = (slot.num_prompt_tokens - slot.params.n_keep - n_block_size) / n_block_size;
 
-                        std::vector<llama_token> new_tokens(
-                            prompt_tokens.begin(),
-                            prompt_tokens.begin() + slot.params.n_keep);
-                        new_tokens.insert(
-                            new_tokens.end(),
-                            prompt_tokens.begin() + slot.params.n_keep + erased_blocks * n_block_size,
-                            prompt_tokens.end());
+                        std::vector<llama_token> new_tokens(prompt_tokens.begin(), prompt_tokens.begin() + slot.params.n_keep);
+                        new_tokens.insert(new_tokens.end(), prompt_tokens.begin() + slot.params.n_keep + erased_blocks * n_block_size, prompt_tokens.end());
 
                         LOG_VERBOSE("input truncated", {
-                            {"n_ctx",      slot.n_ctx},
-                            {"n_keep",     slot.params.n_keep},
-                            {"n_left",     n_left},
+                            {"n_ctx",  slot.n_ctx},
+                            {"n_keep", slot.params.n_keep},
+                            {"n_left", n_left},
                             {"new_tokens", tokens_to_str(ctx, new_tokens.cbegin(), new_tokens.cend())},
                         });
                         slot.truncated = true;
                         prompt_tokens = new_tokens;
 
-                        slot.n_prompt_tokens = prompt_tokens.size();
-                        GGML_ASSERT(slot.n_prompt_tokens < slot.n_ctx);
+                        slot.num_prompt_tokens = prompt_tokens.size();
+                        GGML_ASSERT(slot.num_prompt_tokens < slot.n_ctx);
                     }
 
                     if (!slot.params.cache_prompt)
                     {
                         llama_sampling_reset(slot.ctx_sampling);
 
-                        slot.n_past    = 0;
+                        slot.n_past = 0;
                         slot.n_past_se = 0;
-                        slot.ga_i      = 0;
-                        slot.n_prompt_tokens_processed = slot.n_prompt_tokens;
+                        slot.ga_i = 0;
+                        slot.num_prompt_tokens_processed = slot.num_prompt_tokens;
                     }
                     else
                     {
@@ -1763,7 +1811,7 @@ struct llama_server_context
                             slot.n_past -= 1;
                         }
 
-                        slot.n_prompt_tokens_processed = slot.n_prompt_tokens - slot.n_past;
+                        slot.num_prompt_tokens_processed = slot.num_prompt_tokens - slot.n_past;
 
                         if (slot.ga_n != 1)
                         {
@@ -1788,13 +1836,13 @@ struct llama_server_context
                             { "slot_id", slot.id },
                             { "task_id", slot.task_id },
                             { "n_past",  slot.n_past },
-                            { "n_prompt_tokens_processed", slot.n_prompt_tokens_processed }
+                            { "num_prompt_tokens_processed", slot.num_prompt_tokens_processed }
                         });
                     }
 
                     slot.cache_tokens = prompt_tokens;
 
-                    if (slot.n_past == slot.n_prompt_tokens && slot.n_past > 0)
+                    if (slot.n_past == slot.num_prompt_tokens && slot.n_past > 0)
                     {
                         // we have to evaluate at least 1 token to generate logits.
                         LOG_INFO("we have to evaluate at least 1 token to generate logits", {
@@ -1850,8 +1898,8 @@ struct llama_server_context
                     if (has_images && !ingest_images(slot, n_batch))
                     {
                         LOG_ERROR("failed processing images", {
-                            {"slot_id", slot.id},
-                            {"task_id", slot.task_id},
+                            "slot_id", slot.id,
+                            "task_id", slot.task_id,
                         });
                         // FIXME @phymbert: to be properly tested
                         //  early returning without changing the slot state will block the slot for ever
@@ -2001,6 +2049,10 @@ struct llama_server_context
         LOG_VERBOSE("slots updated", {});
         return true;
     }
+
+    void run_on_all_tasks_finished() {
+        update_slots();
+    }
 };
 
 static void server_print_usage(const char *argv0, const gpt_params &params,
@@ -2013,7 +2065,6 @@ static void server_print_usage(const char *argv0, const gpt_params &params,
     printf("  -v, --verbose             verbose output (default: %s)\n", server_verbose ? "enabled" : "disabled");
     printf("  -t N, --threads N         number of threads to use during computation (default: %d)\n", params.n_threads);
     printf("  -tb N, --threads-batch N  number of threads to use during batch and prompt processing (default: same as --threads)\n");
-    printf("  --threads-http N          number of threads in the http server pool to process requests (default: hardware concurrency)\n");
     printf("  -c N, --ctx-size N        size of the prompt context (default: %d)\n", params.n_ctx);
     printf("  --rope-scaling {none,linear,yarn}\n");
     printf("                            RoPE frequency scaling method, defaults to linear unless specified by the model\n");
@@ -2082,8 +2133,8 @@ static void server_print_usage(const char *argv0, const gpt_params &params,
     printf("  --override-kv KEY=TYPE:VALUE\n");
     printf("                            advanced option to override model metadata by key. may be specified multiple times.\n");
     printf("                            types: int, float, bool. example: --override-kv tokenizer.ggml.add_bos_token=bool:false\n");
-    printf("  -gan N, --grp-attn-n N    set the group attention factor to extend context size through self-extend(default: 1=disabled), used together with group attention width `--grp-attn-w`\n");
-    printf("  -gaw N, --grp-attn-w N    set the group attention width to extend context size through self-extend(default: 512), used together with group attention factor `--grp-attn-n`\n");
+    printf("  -gan N, --grp-attn-n N    set the group attention factor to extend context size through self-extend(default: 1=disabled), used together with group attention width `--grp-attn-w`");
+    printf("  -gaw N, --grp-attn-w N    set the group attention width to extend context size through self-extend(default: 512), used together with group attention factor `--grp-attn-n`");
     printf("  --chat-template JINJA_TEMPLATE\n");
     printf("                            set custom jinja chat template (default: template taken from model's metadata)\n");
     printf("                            Note: only commonly used templates are accepted, since we don't have jinja parser\n");
@@ -2300,15 +2351,6 @@ static void server_params_parse(int argc, char **argv, server_params &sparams,
             }
             params.n_threads_batch = std::stoi(argv[i]);
         }
-        else if (arg == "--threads-http")
-        {
-            if (++i >= argc)
-            {
-                invalid_param = true;
-                break;
-            }
-            sparams.n_threads_http = std::stoi(argv[i]);
-        }
         else if (arg == "-b" || arg == "--batch-size")
         {
             if (++i >= argc)
@@ -2519,7 +2561,7 @@ static void server_params_parse(int argc, char **argv, server_params &sparams,
                 std::istreambuf_iterator<char>(),
                 std::back_inserter(systm_content)
             );
-            llama.system_prompt_process(json::parse(systm_content));
+            llama.process_system_prompt_data(json::parse(systm_content));
         }
         else if (arg == "-ctk" || arg == "--cache-type-k") {
             params.cache_type_k = argv[++i];
@@ -2650,7 +2692,7 @@ static void server_params_parse(int argc, char **argv, server_params &sparams,
 
 /* llama.cpp completion api semantics */
 static json format_partial_response(
-    llama_server_context &llama, server_slot *slot, const std::string &content, const std::vector<completion_token_output> &probs
+    llama_server_context &llama, llama_client_slot *slot, const std::string &content, const std::vector<completion_token_output> &probs
 ) {
     json res = json
     {
@@ -2706,7 +2748,14 @@ static void log_server_request(const httplib::Request &req, const httplib::Respo
     });
 }
 
-static void append_to_generated_text_from_generated_token_probs(llama_server_context &llama, server_slot *slot)
+struct token_translator
+{
+    llama_context * ctx;
+    std::string operator()(llama_token tok)                    const { return llama_token_to_piece(ctx, tok); }
+    std::string operator()(const completion_token_output &cto) const { return (*this)(cto.tok); }
+};
+
+static void append_to_generated_text_from_generated_token_probs(llama_server_context &llama, llama_client_slot *slot)
 {
     auto & gtps = slot->generated_token_probs;
     auto translator = token_translator{llama.ctx};
@@ -3460,11 +3509,6 @@ int main(int argc, char **argv)
     }*/
     //);
 
-    if (sparams.n_threads_http > 0) {
-        log_data["n_threads_http"] =  std::to_string(sparams.n_threads_http);
-        svr.new_task_queue = [&sparams] { return new httplib::ThreadPool(sparams.n_threads_http); };
-    }
-
     LOG_INFO("HTTP server listening", log_data);
     // run the HTTP server in a thread - see comment below
     std::thread t([&]()
@@ -3482,8 +3526,8 @@ int main(int argc, char **argv)
         &llama_server_context::process_single_task, &llama, std::placeholders::_1));
     llama.queue_tasks.on_finish_multitask(std::bind(
         &llama_server_context::on_finish_multitask, &llama, std::placeholders::_1));
-    llama.queue_tasks.on_run_slots(std::bind(
-        &llama_server_context::update_slots, &llama));
+    llama.queue_tasks.on_all_tasks_finished(std::bind(
+        &llama_server_context::run_on_all_tasks_finished, &llama));
     llama.queue_results.on_multitask_update(std::bind(
         &llama_server_queue::update_multitask,
         &llama.queue_tasks,

examples/server/utils.hpp

@@ -37,6 +37,10 @@ extern bool server_log_json;
 #define LOG_WARNING(MSG, ...) server_log("WARN", __func__, __LINE__, MSG, __VA_ARGS__)
 #define LOG_INFO(   MSG, ...) server_log("INFO", __func__, __LINE__, MSG, __VA_ARGS__)
 
+//
+// parallel
+//
+
 enum server_state {
     SERVER_STATE_LOADING_MODEL,  // Server is starting up, model not fully loaded yet
     SERVER_STATE_READY,          // Server is ready and model is loaded
@@ -74,8 +78,51 @@ struct task_multi {
     std::vector<task_result> results{};
 };
 
+// TODO: can become bool if we can't find use of more states
+enum slot_state
+{
+    IDLE,
+    PROCESSING,
+};
+
+enum slot_command
+{
+    NONE,
+    LOAD_PROMPT,
+    RELEASE,
+};
+
+struct slot_params
+{
+    bool stream       = true;
+    bool cache_prompt = false; // remember the prompt to avoid reprocessing all prompt
+
+    uint32_t seed      = -1; // RNG seed
+    int32_t  n_keep    =  0; // number of tokens to keep from initial prompt
+    int32_t  n_predict = -1; // new tokens to predict
+
+    std::vector<std::string> antiprompt;
+
+    json input_prefix;
+    json input_suffix;
+};
+
+struct slot_image
+{
+    int32_t id;
+
+    bool request_encode_image = false;
+    float * image_embedding = nullptr;
+    int32_t image_tokens = 0;
+
+    clip_image_u8 * img_data;
+
+    std::string prefix_prompt; // before of this image
+};
+
 // completion token output with probabilities
-struct completion_token_output {
+struct completion_token_output
+{
     struct token_prob
     {
         llama_token tok;
@@ -87,13 +134,8 @@ struct completion_token_output {
     std::string text_to_send;
 };
 
-struct token_translator {
-    llama_context * ctx;
-    std::string operator()(llama_token tok)                    const { return llama_token_to_piece(ctx, tok); }
-    std::string operator()(const completion_token_output &cto) const { return (*this)(cto.tok); }
-};
-
-static inline void server_log(const char *level, const char *function, int line, const char *message, const nlohmann::ordered_json &extra) {
+static inline void server_log(const char *level, const char *function, int line, const char *message, const nlohmann::ordered_json &extra)
+{
     std::stringstream ss_tid;
     ss_tid << std::this_thread::get_id();
     json log = nlohmann::ordered_json{
@@ -141,7 +183,8 @@ static inline void server_log(const char *level, const char *function, int line,
 //
 
 template <typename T>
-static T json_value(const json &body, const std::string &key, const T &default_value) {
+static T json_value(const json &body, const std::string &key, const T &default_value)
+{
     // Fallback null to default value
     return body.contains(key) && !body.at(key).is_null()
         ? body.value(key, default_value)
@@ -157,7 +200,8 @@ inline bool verify_custom_template(const std::string & tmpl) {
 }
 
 // Format given chat. If tmpl is empty, we take the template from model metadata
-inline std::string format_chat(const struct llama_model * model, const std::string & tmpl, const std::vector<json> & messages) {
+inline std::string format_chat(const struct llama_model * model, const std::string & tmpl, const std::vector<json> & messages)
+{
     size_t alloc_size = 0;
     // vector holding all allocated string to be passed to llama_chat_apply_template
     std::vector<std::string> str(messages.size() * 2);
@@ -206,7 +250,7 @@ struct llama_server_queue {
     // callback functions
     std::function<void(task_server&)> callback_new_task;
     std::function<void(task_multi&)> callback_finish_multitask;
-    std::function<void(void)> callback_run_slots;
+    std::function<void(void)> callback_all_task_finished;
 
     // Add a new task to the end of the queue
     int post(task_server task) {
@@ -239,14 +283,14 @@ struct llama_server_queue {
         callback_new_task = callback;
     }
 
-    // Register function to process a multitask when it is finished
+    // Register function to process a multitask
     void on_finish_multitask(std::function<void(task_multi&)> callback) {
         callback_finish_multitask = callback;
     }
 
-    // Register the function to be called when all slots data is ready to be processed
-    void on_run_slots(std::function<void(void)> callback) {
-        callback_run_slots = callback;
+    // Register the function to be called when the batch of tasks is finished
+    void on_all_tasks_finished(std::function<void(void)> callback) {
+        callback_all_task_finished = callback;
     }
 
     // Call when the state of one slot is changed
@@ -268,13 +312,7 @@ struct llama_server_queue {
         condition_tasks.notify_all();
     }
 
-    /**
-     * Main loop consists of these steps:
-     * - Wait until a new task arrives
-     * - Process the task (i.e. maybe copy data into slot)
-     * - Check if multitask is finished
-     * - Run all slots
-     */
+    // Start the main loop.
     void start_loop() {
         running = true;
         while (true) {
@@ -293,8 +331,8 @@ struct llama_server_queue {
                     LOG_VERBOSE("callback_new_task", {{"task_id", task.id}});
                     callback_new_task(task);
                 }
-                LOG_VERBOSE("update_multitasks", {});
-                // check if we have any finished multitasks
+                LOG_VERBOSE("callback_all_task_finished", {});
+                // process and update all the multitasks
                 auto queue_iterator = queue_multitasks.begin();
                 while (queue_iterator != queue_multitasks.end())
                 {
@@ -311,9 +349,8 @@ struct llama_server_queue {
                         ++queue_iterator;
                     }
                 }
-                // all tasks in the current loop is processed, slots data is now ready
-                LOG_VERBOSE("callback_run_slots", {});
-                callback_run_slots();
+                // all tasks in the current loop is finished
+                callback_all_task_finished();
             }
             LOG_VERBOSE("wait for new task", {});
             // wait for new task
@@ -371,14 +408,12 @@ struct llama_server_response {
     std::mutex mutex_results;
     std::condition_variable condition_results;
 
-    // add the task_id to the list of tasks waiting for response
     void add_waiting_task_id(int task_id) {
         LOG_VERBOSE("waiting for task id", {{"task_id", task_id}});
         std::unique_lock<std::mutex> lock(mutex_results);
         waiting_task_ids.insert(task_id);
     }
 
-    // when the request is finished, we can remove task associated with it
     void remove_waiting_task_id(int task_id) {
         LOG_VERBOSE("remove waiting for task id", {{"task_id", task_id}});
         std::unique_lock<std::mutex> lock(mutex_results);
@@ -539,96 +574,3 @@ static std::string gen_chatcmplid()
     chatcmplid << "chatcmpl-" << random_string();
     return chatcmplid.str();
 }
-
-//
-// other common utils
-//
-
-static size_t common_part(const std::vector<llama_token> &a, const std::vector<llama_token> &b)
-{
-    size_t i;
-    for (i = 0; i < a.size() && i < b.size() && a[i] == b[i]; i++)
-    {
-    }
-    return i;
-}
-
-static bool ends_with(const std::string &str, const std::string &suffix)
-{
-    return str.size() >= suffix.size() &&
-           0 == str.compare(str.size() - suffix.size(), suffix.size(), suffix);
-}
-
-static size_t find_partial_stop_string(const std::string &stop,
-                                       const std::string &text)
-{
-    if (!text.empty() && !stop.empty())
-    {
-        const char text_last_char = text.back();
-        for (int64_t char_index = stop.size() - 1; char_index >= 0; char_index--)
-        {
-            if (stop[char_index] == text_last_char)
-            {
-                const std::string current_partial = stop.substr(0, char_index + 1);
-                if (ends_with(text, current_partial))
-                {
-                    return text.size() - char_index - 1;
-                }
-            }
-        }
-    }
-    return std::string::npos;
-}
-
-// TODO: reuse llama_detokenize
-template <class Iter>
-static std::string tokens_to_str(llama_context *ctx, Iter begin, Iter end)
-{
-    std::string ret;
-    for (; begin != end; ++begin)
-    {
-        ret += llama_token_to_piece(ctx, *begin);
-    }
-    return ret;
-}
-
-// format incomplete utf-8 multibyte character for output
-static std::string tokens_to_output_formatted_string(const llama_context *ctx, const llama_token token)
-{
-    std::string out = token == -1 ? "" : llama_token_to_piece(ctx, token);
-    // if the size is 1 and first bit is 1, meaning it's a partial character
-    //   (size > 1 meaning it's already a known token)
-    if (out.size() == 1 && (out[0] & 0x80) == 0x80)
-    {
-        std::stringstream ss;
-        ss << std::hex << (out[0] & 0xff);
-        std::string res(ss.str());
-        out = "byte: \\x" + res;
-    }
-    return out;
-}
-
-// convert a vector of completion_token_output to json
-static json probs_vector_to_json(const llama_context *ctx, const std::vector<completion_token_output> &probs)
-{
-    json out = json::array();
-    for (const auto &prob : probs)
-    {
-        json probs_for_token = json::array();
-        for (const auto &p : prob.probs)
-        {
-            std::string tok_str = tokens_to_output_formatted_string(ctx, p.tok);
-            probs_for_token.push_back(json
-            {
-                {"tok_str", tok_str},
-                {"prob",    p.prob},
-            });
-        }
-        std::string tok_str = tokens_to_output_formatted_string(ctx, prob.tok);
-        out.push_back(json{
-            {"content", tok_str},
-            {"probs",   probs_for_token},
-        });
-    }
-    return out;
-}

ggml-sycl.cpp

@@ -12726,7 +12726,6 @@ static void ggml_sycl_op_mul_mat(const ggml_tensor *src0,
 
     GGML_ASSERT(dst->backend != GGML_BACKEND_TYPE_GPU_SPLIT);
     GGML_ASSERT(src1->backend != GGML_BACKEND_TYPE_GPU_SPLIT);
-    GGML_ASSERT(src1->type == GGML_TYPE_F32 || (src1->ne[2] == 1 && src1->ne[3] == 1));
 
     GGML_ASSERT(ne12 >= ne02 && ne12 % ne02 == 0);
 
@@ -13270,23 +13269,31 @@ static void k_compute_batched_ptrs(const sycl::half *src0_as_f16,
     int64_t i03 = i13 / r3;
     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;
+    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/2 + i13*nb13/2;
+    ptrs_dst[0*ne23 + i12 + i13*ne12] = (      char *)         dst + i12*nbd2   + i13*nbd3;
 }
 
-static void ggml_sycl_mul_mat_batched_sycl(const ggml_tensor *src0,
-                                             const ggml_tensor *src1,
-                                             ggml_tensor *dst) try {
+static void ggml_sycl_mul_mat_mat_batched_sycl(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(src0->backend != GGML_BACKEND_TYPE_GPU_SPLIT);
     GGML_ASSERT(src0->type == GGML_TYPE_F16);
+    GGML_ASSERT(src1->type == GGML_TYPE_F32);
 
-    GGML_TENSOR_BINARY_OP_LOCALS
+    GGML_TENSOR_LOCALS(int64_t, ne0, src0, ne);
 
-    const int64_t ne_dst  = ggml_nelements(dst);
+    GGML_TENSOR_LOCALS(int64_t, nb0, src0, nb);
+
+    GGML_TENSOR_LOCALS(int64_t, ne1, src1, ne);
+
+    GGML_TENSOR_LOCALS(int64_t, nb1, src1, nb);
+
+    const int64_t ne1 = ggml_nelements(src1);
+    const int64_t ne  = ggml_nelements(dst);
 
     SYCL_CHECK(ggml_sycl_set_device(g_main_device));
     dpct::queue_ptr main_stream = g_syclStreams[g_main_device_index][0];
@@ -13305,16 +13312,11 @@ static void ggml_sycl_mul_mat_batched_sycl(const ggml_tensor *src0,
     float * dst_ddf = (float *) dst_extra->data_device[g_main_device_index];
 
     // convert src1 to fp16
-    sycl_pool_alloc<sycl::half> src1_f16_alloc;
-    if (src1->type != GGML_TYPE_F16) {
-      const to_fp16_sycl_t to_fp16_sycl = ggml_get_to_fp16_sycl(src1->type);
-      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);
-    }
-    sycl::half *src1_f16 = src1->type == GGML_TYPE_F16 ? (sycl::half *)src1_ddf
-                                                       : src1_f16_alloc.get();
+    const to_fp16_sycl_t to_fp16_sycl = ggml_get_to_fp16_sycl(src1->type);
+    GGML_ASSERT(to_fp16_sycl != nullptr);
+
+    sycl_pool_alloc<sycl::half> src1_as_f16(ne1);
+    to_fp16_sycl(src1_ddf, src1_as_f16.get(), ne1, main_stream);
 
     sycl_pool_alloc<sycl::half> dst_f16;
     char * dst_t;
@@ -13335,12 +13337,20 @@ static void ggml_sycl_mul_mat_batched_sycl(const ggml_tensor *src0,
     const void * alpha = &alpha_f16;
     const void * beta  = &beta_f16;
 
-    // TODO: Renable (dst->op_params[0] =! GGML_PREC_DEFAULT) pathway
-    // once oneMKL open source supports half, half, float, float: datatypes
-    dst_t = (char *) dst_f16.alloc(ne_dst);
+    if (dst->op_params[0] == GGML_PREC_DEFAULT) {
+        dst_t = (char *) dst_f16.alloc(ne);
 
-    nbd2 /= sizeof(float) / sizeof(sycl::half);
-    nbd3 /= sizeof(float) / sizeof(sycl::half);
+        nbd2 /= sizeof(float) / sizeof(sycl::half);
+        nbd3 /= sizeof(float) / sizeof(sycl::half);
+    } else {
+        dst_t = (char *) dst_ddf;
+
+        cu_compute_type = dpct::library_data_t::real_float;
+        cu_data_type = dpct::library_data_t::real_float;
+
+        alpha = &alpha_f32;
+        beta  = &beta_f32;
+    }
 
     GGML_ASSERT(ne12 % ne02 == 0);
     GGML_ASSERT(ne13 % ne03 == 0);
@@ -13376,10 +13386,10 @@ static void ggml_sycl_mul_mat_batched_sycl(const ggml_tensor *src0,
             *g_sycl_handles[g_main_device_index], oneapi::mkl::transpose::trans,
             oneapi::mkl::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,
+            nb01 / sizeof(sycl::half), src0->nb[2] / sizeof(sycl::half),
+            (const char *)src1_as_f16.get(), dpct::library_data_t::real_half,
+            nb11 / sizeof(float), src1->nb[2] / sizeof(float), beta,
+            (char *)dst_t, cu_data_type, ne01, dst->nb[2] / sizeof(float),
             ne12 * ne13, cu_compute_type)));
     } else {
         // use syclGemmBatchedEx
@@ -13399,35 +13409,44 @@ static void ggml_sycl_mul_mat_batched_sycl(const ggml_tensor *src0,
                                          {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;
+                const sycl::half *src1_as_f16_get_ct1 = src1_as_f16.get();
+                const void **ptrs_src_get_ct3 = ptrs_src.get();
+                void **ptrs_dst_get_ct4 = ptrs_dst.get();
+
                 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);
+                                         src0_as_f16, src1_as_f16_get_ct1,
+                                         dst_t, ptrs_src_get_ct3,
+                                         ptrs_dst_get_ct4, ne12, ne13, ne23,
+                                         nb02, nb03, nb12, nb13, nbd2, nbd3, r2,
+                                         r3, item_ct1);
                                  });
             });
         }
+        /*
+        DPCT1010:95: SYCL uses exceptions to report errors and does not use the
+        error codes. The call was replaced with 0. You need to rewrite this
+        code.
+        */
+        SYCL_CHECK(0);
+
         SYCL_CHECK(CHECK_TRY_ERROR(dpct::gemm_batch(
             *g_sycl_handles[g_main_device_index], oneapi::mkl::transpose::trans,
             oneapi::mkl::transpose::nontrans, ne01, ne11, ne10, alpha,
             (const void **)(ptrs_src.get() + 0 * ne23),
-            dpct::library_data_t::real_half, nb01 / nb00,
+            dpct::library_data_t::real_half, nb01 / sizeof(sycl::half),
             (const void **)(ptrs_src.get() + 1 * ne23),
-            dpct::library_data_t::real_half, nb11 / nb10, beta,
+            dpct::library_data_t::real_half, nb11 / sizeof(float), beta,
             (void **)(ptrs_dst.get() + 0 * ne23), cu_data_type, ne01, ne23,
             cu_compute_type)));
     }
 #endif
 
-    const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(GGML_TYPE_F16);
-    to_fp32_sycl(dst_f16.get(), dst_ddf, ne_dst, main_stream);
+    if (dst->op_params[0] == GGML_PREC_DEFAULT) {
+        const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(GGML_TYPE_F16);
+        to_fp32_sycl(dst_f16.get(), dst_ddf, ne, main_stream);
+    }
 }
 catch (sycl::exception const &exc) {
   std::cerr << exc.what() << "Exception caught at file:" << __FILE__
@@ -13472,10 +13491,10 @@ static void ggml_sycl_mul_mat(const ggml_tensor * src0, const ggml_tensor * src1
         // KQV single-batch
         // GGML_SYCL_DEBUG("ggml_sycl_mul_mat_vec_nc\n");
         ggml_sycl_mul_mat_vec_nc(src0, src1, dst);
-    } else if (!split && all_on_device && use_xmx && src0->type == GGML_TYPE_F16 && !ggml_is_transposed(src0) && !ggml_is_transposed(src1)) {
+    } else if (!split && all_on_device && use_xmx && src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F32 && !ggml_is_transposed(src0) && !ggml_is_transposed(src1)) {
         // KQ + KQV multi-batch
-        // GGML_SYCL_DEBUG("ggml_sycl_mul_mat_batched_sycl\n");
-        ggml_sycl_mul_mat_batched_sycl(src0, src1, dst);
+        // GGML_SYCL_DEBUG("ggml_sycl_mul_mat_mat_batched_sycl\n");
+        ggml_sycl_mul_mat_mat_batched_sycl(src0, src1, dst);
     } else if (src0->type == GGML_TYPE_F32) {
         // GGML_SYCL_DEBUG("ggml_sycl_op_mul_mat\n");
         ggml_sycl_op_mul_mat(src0, src1, dst, ggml_sycl_op_mul_mat_sycl, false);

gguf-py/gguf/constants.py

@@ -112,6 +112,7 @@ class MODEL_ARCH(IntEnum):
     INTERNLM2  = auto()
     MINICPM    = auto()
     GEMMA      = auto()
+    STARCODER2 = auto()
 
 
 class MODEL_TENSOR(IntEnum):
@@ -169,6 +170,7 @@ MODEL_ARCH_NAMES: dict[MODEL_ARCH, str] = {
     MODEL_ARCH.INTERNLM2:      "internlm2",
     MODEL_ARCH.MINICPM:        "minicpm",
     MODEL_ARCH.GEMMA:          "gemma",
+    MODEL_ARCH.STARCODER2:     "starcoder2",
 }
 
 TENSOR_NAMES: dict[MODEL_TENSOR, str] = {
@@ -526,6 +528,21 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
         MODEL_TENSOR.FFN_UP,
         MODEL_TENSOR.FFN_NORM,
     ],
+    MODEL_ARCH.STARCODER2: [
+        MODEL_TENSOR.TOKEN_EMBD,
+        MODEL_TENSOR.OUTPUT_NORM,
+        MODEL_TENSOR.OUTPUT,
+        MODEL_TENSOR.ROPE_FREQS,
+        MODEL_TENSOR.ATTN_NORM,
+        MODEL_TENSOR.ATTN_Q,
+        MODEL_TENSOR.ATTN_K,
+        MODEL_TENSOR.ATTN_V,
+        MODEL_TENSOR.ATTN_OUT,
+        MODEL_TENSOR.ATTN_ROT_EMBD,
+        MODEL_TENSOR.FFN_NORM,
+        MODEL_TENSOR.FFN_DOWN,
+        MODEL_TENSOR.FFN_UP,
+    ],
     # TODO
 }
 
@@ -554,6 +571,10 @@ MODEL_TENSOR_SKIP: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
         MODEL_TENSOR.ROPE_FREQS,
         MODEL_TENSOR.ATTN_ROT_EMBD,
     ],
+    MODEL_ARCH.STARCODER2: [
+        MODEL_TENSOR.ROPE_FREQS,
+        MODEL_TENSOR.ATTN_ROT_EMBD,
+    ],
 }
 
 #

gguf-py/gguf/tensor_mapping.py

@@ -210,6 +210,7 @@ class TensorNameMap:
             "model.layers.layers.{bid}.mlp.up_proj",                  # plamo
             "model.layers.{bid}.feed_forward.w3",                     # internlm2
             "encoder.layers.{bid}.mlp.fc11",                          # nomic-bert
+            "model.layers.{bid}.mlp.c_fc",                            # starcoder2
         ),
 
         MODEL_TENSOR.FFN_UP_EXP: (
@@ -256,6 +257,7 @@ class TensorNameMap:
             "model.layers.layers.{bid}.mlp.down_proj",                # plamo
             "model.layers.{bid}.feed_forward.w2",                     # internlm2
             "encoder.layers.{bid}.mlp.fc2",                           # nomic-bert
+            "model.layers.{bid}.mlp.c_proj",                          # starcoder2
         ),
 
         MODEL_TENSOR.FFN_DOWN_EXP: (

llama.cpp

@@ -211,6 +211,7 @@ enum llm_arch {
     LLM_ARCH_INTERNLM2,
     LLM_ARCH_MINICPM,
     LLM_ARCH_GEMMA,
+    LLM_ARCH_STARCODER2,
     LLM_ARCH_UNKNOWN,
 };
 
@@ -238,6 +239,7 @@ static std::map<llm_arch, const char *> LLM_ARCH_NAMES = {
     { LLM_ARCH_INTERNLM2,       "internlm2"  },
     { LLM_ARCH_MINICPM,         "minicpm"    },
     { LLM_ARCH_GEMMA,           "gemma"      },
+    { LLM_ARCH_STARCODER2,      "starcoder2" },
 };
 
 enum llm_kv {
@@ -779,6 +781,24 @@ static std::map<llm_arch, std::map<llm_tensor, std::string>> LLM_TENSOR_NAMES =
             { LLM_TENSOR_FFN_UP,          "blk.%d.ffn_up" },
         },
     },
+    {
+        LLM_ARCH_STARCODER2,
+        {
+            { LLM_TENSOR_TOKEN_EMBD,      "token_embd" },
+            { LLM_TENSOR_OUTPUT_NORM,     "output_norm" },
+            { LLM_TENSOR_OUTPUT,          "output" },
+            { LLM_TENSOR_ROPE_FREQS,      "rope_freqs" },
+            { LLM_TENSOR_ATTN_NORM,       "blk.%d.attn_norm" },
+            { LLM_TENSOR_ATTN_Q,          "blk.%d.attn_q" },
+            { LLM_TENSOR_ATTN_K,          "blk.%d.attn_k" },
+            { LLM_TENSOR_ATTN_V,          "blk.%d.attn_v" },
+            { LLM_TENSOR_ATTN_OUT,        "blk.%d.attn_output" },
+            { LLM_TENSOR_ATTN_ROT_EMBD,   "blk.%d.attn_rot_embd" },
+            { LLM_TENSOR_FFN_NORM,        "blk.%d.ffn_norm" },
+            { LLM_TENSOR_FFN_DOWN,        "blk.%d.ffn_down" },
+            { LLM_TENSOR_FFN_UP,          "blk.%d.ffn_up" },
+        },
+    },
     {
         LLM_ARCH_UNKNOWN,
         {
@@ -3321,6 +3341,16 @@ static void llm_load_hparams(
                     default: model.type = e_model::MODEL_UNKNOWN;
                }
             } break;
+        case LLM_ARCH_STARCODER2:
+            {
+                ml.get_key(LLM_KV_ATTENTION_LAYERNORM_EPS, hparams.f_norm_eps);
+                switch (hparams.n_layer) {
+                    case 30: model.type = e_model::MODEL_3B; break;
+                    case 32: model.type = e_model::MODEL_7B; break;
+                    case 40: model.type = e_model::MODEL_15B; break;
+                    default: model.type = e_model::MODEL_UNKNOWN;
+                }
+            } break;
         default: (void)0;
     }
 
@@ -4491,6 +4521,56 @@ static bool llm_load_tensors(
                         layer.ffn_down = ml.create_tensor(ctx_split, tn(LLM_TENSOR_FFN_DOWN, "weight", i), {  n_ff, n_embd});
                     }
                 } break;
+            case LLM_ARCH_STARCODER2:
+                {
+                    model.tok_embd = ml.create_tensor(ctx_input, tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab});
+
+                    // output
+                    {
+                        model.output_norm   = ml.create_tensor(ctx_output, tn(LLM_TENSOR_OUTPUT_NORM, "weight"), {n_embd});
+                        model.output_norm_b = ml.create_tensor(ctx_output, tn(LLM_TENSOR_OUTPUT_NORM, "bias"),   {n_embd});
+
+                        model.output = ml.create_tensor(ctx_output_split, tn(LLM_TENSOR_OUTPUT, "weight"), {n_embd, n_vocab}, false);
+                        // if output is NULL, init from the input tok embed
+                        if (model.output == NULL) {
+                            model.output = ml.create_tensor(ctx_output, tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab});
+                            ml.n_created--; // artificial tensor
+                            ml.size_data += ggml_nbytes(model.output);
+                        }
+
+                    }
+
+                    for (int i = 0; i < n_layer; ++i) {
+                        ggml_context * ctx_layer = ctx_for_layer(i);
+                        ggml_context * ctx_split = ctx_for_layer_split(i);
+
+                        auto & layer = model.layers[i];
+
+                        layer.attn_norm   = ml.create_tensor(ctx_layer, tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd});
+                        layer.attn_norm_b = ml.create_tensor(ctx_layer, tn(LLM_TENSOR_ATTN_NORM, "bias", i),   {n_embd});
+
+                        layer.wq = ml.create_tensor(ctx_split, tn(LLM_TENSOR_ATTN_Q,   "weight", i), {n_embd, n_embd});
+                        layer.wk = ml.create_tensor(ctx_split, tn(LLM_TENSOR_ATTN_K,   "weight", i), {n_embd, n_embd_gqa});
+                        layer.wv = ml.create_tensor(ctx_split, tn(LLM_TENSOR_ATTN_V,   "weight", i), {n_embd, n_embd_gqa});
+                        layer.wo = ml.create_tensor(ctx_split, tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_embd, n_embd});
+
+                        // optional bias tensors
+                        layer.bq = ml.create_tensor(ctx_layer, tn(LLM_TENSOR_ATTN_Q,   "bias", i), {n_embd});
+                        layer.bk = ml.create_tensor(ctx_layer, tn(LLM_TENSOR_ATTN_K,   "bias", i), {n_embd_gqa});
+                        layer.bv = ml.create_tensor(ctx_layer, tn(LLM_TENSOR_ATTN_V,   "bias", i), {n_embd_gqa});
+                        layer.bo = ml.create_tensor(ctx_layer, tn(LLM_TENSOR_ATTN_OUT, "bias", i), {n_embd});
+
+                        layer.ffn_norm   = ml.create_tensor(ctx_layer, tn(LLM_TENSOR_FFN_NORM, "weight", i), {n_embd});
+                        layer.ffn_norm_b = ml.create_tensor(ctx_layer, tn(LLM_TENSOR_FFN_NORM, "bias", i),   {n_embd});
+
+                        layer.ffn_down = ml.create_tensor(ctx_split, tn(LLM_TENSOR_FFN_DOWN, "weight", i), {  n_ff, n_embd});
+                        layer.ffn_up   = ml.create_tensor(ctx_split, tn(LLM_TENSOR_FFN_UP,   "weight", i), {n_embd,   n_ff});
+
+                        // optional bias tensors
+                        layer.ffn_down_b = ml.create_tensor(ctx_layer, tn(LLM_TENSOR_FFN_DOWN, "bias", i), {n_embd});
+                        layer.ffn_up_b   = ml.create_tensor(ctx_layer, tn(LLM_TENSOR_FFN_UP ,  "bias", i), {  n_ff});
+                    }
+                } break;
             default:
                 throw std::runtime_error("unknown architecture");
         }
@@ -7560,6 +7640,120 @@ struct llm_build_context {
 
         return gf;
     }
+
+    struct ggml_cgraph * build_starcoder2() {
+        struct ggml_cgraph * gf = ggml_new_graph_custom(ctx0, LLAMA_MAX_NODES, false);
+
+        const int64_t n_embd_head = hparams.n_embd_head_v;
+        GGML_ASSERT(n_embd_head == hparams.n_embd_head_k);
+        GGML_ASSERT(n_embd_head == hparams.n_rot);
+
+        struct ggml_tensor * cur;
+        struct ggml_tensor * inpL;
+
+        inpL = llm_build_inp_embd(ctx0, hparams, batch, model.tok_embd, lctx.inp_tokens, lctx.inp_embd, cb);
+        cb(inpL, "inp_embd", -1);
+
+        // inp_pos - contains the positions
+        struct ggml_tensor * inp_pos = ggml_view_1d(ctx0, lctx.inp_pos, n_tokens, 0);
+        cb(inp_pos, "inp_pos", -1);
+
+        // KQ_mask (mask for 1 head, it will be broadcasted to all heads)
+        struct ggml_tensor * KQ_mask = ggml_view_2d(ctx0, lctx.inp_KQ_mask, n_kv, n_tokens, n_kv*ggml_type_size(lctx.inp_KQ_mask->type), 0);
+        cb(KQ_mask, "KQ_mask", -1);
+
+        for (int il = 0; il < n_layer; ++il) {
+            struct ggml_tensor * inpSA = inpL;
+
+            // norm
+            cur = llm_build_norm(ctx0, inpL, hparams,
+                    model.layers[il].attn_norm, model.layers[il].attn_norm_b,
+                    LLM_NORM, cb, il);
+            cb(cur, "attn_norm", il);
+
+            // self-attention
+            {
+                // compute Q and K and RoPE them
+                struct ggml_tensor * Qcur = ggml_mul_mat(ctx0, model.layers[il].wq, cur);
+                cb(Qcur, "Qcur", il);
+                if (model.layers[il].bq) {
+                    Qcur = ggml_add(ctx0, Qcur, model.layers[il].bq);
+                    cb(Qcur, "Qcur", il);
+                }
+
+                struct ggml_tensor * Kcur = ggml_mul_mat(ctx0, model.layers[il].wk, cur);
+                cb(Kcur, "Kcur", il);
+                if (model.layers[il].bk) {
+                    Kcur = ggml_add(ctx0, Kcur, model.layers[il].bk);
+                    cb(Kcur, "Kcur", il);
+                }
+
+                struct ggml_tensor * Vcur = ggml_mul_mat(ctx0, model.layers[il].wv, cur);
+                cb(Vcur, "Vcur", il);
+                if (model.layers[il].bv) {
+                    Vcur = ggml_add(ctx0, Vcur, model.layers[il].bv);
+                    cb(Vcur, "Vcur", il);
+                }
+
+                Qcur = ggml_rope_custom(
+                    ctx0, ggml_reshape_3d(ctx0, Qcur, n_embd_head, n_head, n_tokens), inp_pos,
+                    n_rot, rope_type, 0, n_orig_ctx, freq_base, freq_scale,
+                    ext_factor, attn_factor, beta_fast, beta_slow
+                );
+                cb(Qcur, "Qcur", il);
+
+                Kcur = ggml_rope_custom(
+                    ctx0, ggml_reshape_3d(ctx0, Kcur, n_embd_head, n_head_kv, n_tokens), inp_pos,
+                    n_rot, rope_type, 0, n_orig_ctx, freq_base, freq_scale,
+                    ext_factor, attn_factor, beta_fast, beta_slow
+                );
+                cb(Kcur, "Kcur", il);
+
+                cur = llm_build_kv(ctx0, model, hparams, kv_self, gf,
+                        model.layers[il].wo, model.layers[il].bo,
+                        Kcur, Vcur, Qcur, KQ_mask, nullptr, n_ctx, n_tokens, kv_head, n_kv, 1.0f/sqrtf(float(n_embd_head)), cb, il);
+                cb(cur, "kqv_out", il);
+            }
+
+            struct ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpSA);
+            cb(ffn_inp, "ffn_inp", il);
+
+            // feed-forward network
+
+            cur = llm_build_norm(ctx0, ffn_inp, hparams,
+                    model.layers[il].ffn_norm, model.layers[il].ffn_norm_b,
+                    LLM_NORM, cb, il);
+            cb(cur, "ffn_norm", il);
+
+            cur = llm_build_ffn(ctx0, cur,
+                        model.layers[il].ffn_up,   model.layers[il].ffn_up_b,
+                        NULL,                      NULL,
+                        model.layers[il].ffn_down, model.layers[il].ffn_down_b,
+                        NULL,
+                        LLM_FFN_GELU, LLM_FFN_SEQ, cb, il);
+            cb(cur, "ffn_out", il);
+            cur = ggml_add(ctx0, cur, ffn_inp);
+            cb(cur, "l_out", il);
+
+            // input for next layer
+            inpL = cur;
+        }
+
+        cur = inpL;
+
+        cur = llm_build_norm(ctx0, cur, hparams,
+                model.output_norm, model.output_norm_b,
+                LLM_NORM, cb, -1);
+        cb(cur, "result_norm", -1);
+
+        // lm_head
+        cur = ggml_mul_mat(ctx0, model.output, cur);
+        cb(cur, "result_output", -1);
+
+        ggml_build_forward_expand(gf, cur);
+
+        return gf;
+    }
 };
 
 static struct ggml_cgraph * llama_build_graph_defrag(llama_context & lctx, const std::vector<uint32_t> & ids) {
@@ -7706,6 +7900,10 @@ static struct ggml_cgraph * llama_build_graph(
             {
                 result = llm.build_gemma();
             } break;
+        case LLM_ARCH_STARCODER2:
+            {
+                result = llm.build_starcoder2();
+            } break;
         default:
             GGML_ASSERT(false);
     }
@@ -12086,6 +12284,7 @@ enum llama_rope_type llama_rope_type(const struct llama_model * model) {
         case LLM_ARCH_QWEN2:
         case LLM_ARCH_PHI2:
         case LLM_ARCH_GEMMA:
+        case LLM_ARCH_STARCODER2:
             return LLAMA_ROPE_TYPE_NEOX;
 
         // all model arches should be listed explicitly here