server: metrics: add llamacpp:prompt_seconds_total and llamacpp:tokens_predicted_seconds_total, reset bucket only on /metrics. Fix values cast to int. Add Process-Start-Time-Unix header. (#5937)

Closes #5850

common/log.h

@@ -297,7 +297,7 @@ inline std::string log_filename_generator_impl(LogTriState multilog, const std::
 #ifndef _MSC_VER
     #define LOG(...) LOG_IMPL(__VA_ARGS__, "")
 #else
-    #define LOG(str, ...) LOG_IMPL("%s" str, "", __VA_ARGS__, "")
+    #define LOG(str, ...) LOG_IMPL("%s" str, "", ##__VA_ARGS__, "")
 #endif
 
 // Main TEE macro.
@@ -311,7 +311,7 @@ inline std::string log_filename_generator_impl(LogTriState multilog, const std::
 #ifndef _MSC_VER
     #define LOG_TEE(...) LOG_TEE_IMPL(__VA_ARGS__, "")
 #else
-    #define LOG_TEE(str, ...) LOG_TEE_IMPL("%s" str, "", __VA_ARGS__, "")
+    #define LOG_TEE(str, ...) LOG_TEE_IMPL("%s" str, "", ##__VA_ARGS__, "")
 #endif
 
 // LOG macro variants with auto endline.
@@ -319,8 +319,8 @@ inline std::string log_filename_generator_impl(LogTriState multilog, const std::
     #define LOGLN(...) LOG_IMPL(__VA_ARGS__, "\n")
     #define LOG_TEELN(...) LOG_TEE_IMPL(__VA_ARGS__, "\n")
 #else
-    #define LOGLN(str, ...) LOG_IMPL("%s" str, "", __VA_ARGS__, "\n")
-    #define LOG_TEELN(str, ...) LOG_TEE_IMPL("%s" str, "", __VA_ARGS__, "\n")
+    #define LOGLN(str, ...) LOG_IMPL("%s" str, "", ##__VA_ARGS__, "\n")
+    #define LOG_TEELN(str, ...) LOG_TEE_IMPL("%s" str, "", ##__VA_ARGS__, "\n")
 #endif
 
 // INTERNAL, DO NOT USE

examples/llama-bench/llama-bench.cpp

@@ -173,6 +173,7 @@ struct cmd_params {
     std::vector<bool> no_kv_offload;
     std::vector<std::vector<float>> tensor_split;
     std::vector<bool> use_mmap;
+    std::vector<bool> embeddings;
     int reps;
     bool verbose;
     output_formats output_format;
@@ -192,6 +193,7 @@ static const cmd_params cmd_params_defaults = {
     /* no_kv_offload */ {false},
     /* tensor_split  */ {std::vector<float>(llama_max_devices(), 0.0f)},
     /* use_mmap      */ {true},
+    /* embeddings    */ {false},
     /* reps          */ 5,
     /* verbose       */ false,
     /* output_format */ MARKDOWN
@@ -214,6 +216,7 @@ static void print_usage(int /* argc */, char ** argv) {
     printf("  -mg, --main-gpu <i>                 (default: %s)\n", join(cmd_params_defaults.main_gpu, ",").c_str());
     printf("  -nkvo, --no-kv-offload <0|1>        (default: %s)\n", join(cmd_params_defaults.no_kv_offload, ",").c_str());
     printf("  -mmp, --mmap <0|1>                  (default: %s)\n", join(cmd_params_defaults.use_mmap, ",").c_str());
+    printf("  -embd, --embeddings <0|1>           (default: %s)\n", join(cmd_params_defaults.embeddings, ",").c_str());
     printf("  -ts, --tensor_split <ts0/ts1/..>    (default: 0)\n");
     printf("  -r, --repetitions <n>               (default: %d)\n", cmd_params_defaults.reps);
     printf("  -o, --output <csv|json|md|sql>      (default: %s)\n", output_format_str(cmd_params_defaults.output_format));
@@ -382,6 +385,13 @@ static cmd_params parse_cmd_params(int argc, char ** argv) {
             }
             auto p = split<bool>(argv[i], split_delim);
             params.use_mmap.insert(params.use_mmap.end(), p.begin(), p.end());
+        } else if (arg == "-embd" || arg == "--embeddings") {
+            if (++i >= argc) {
+                invalid_param = true;
+                break;
+            }
+            auto p = split<bool>(argv[i], split_delim);
+            params.embeddings.insert(params.embeddings.end(), p.begin(), p.end());
         } else if (arg == "-ts" || arg == "--tensor-split") {
             if (++i >= argc) {
                 invalid_param = true;
@@ -453,6 +463,7 @@ static cmd_params parse_cmd_params(int argc, char ** argv) {
     if (params.no_kv_offload.empty()){ params.no_kv_offload = cmd_params_defaults.no_kv_offload; }
     if (params.tensor_split.empty()) { params.tensor_split = cmd_params_defaults.tensor_split; }
     if (params.use_mmap.empty())     { params.use_mmap = cmd_params_defaults.use_mmap; }
+    if (params.embeddings.empty())   { params.embeddings = cmd_params_defaults.embeddings; }
     if (params.n_threads.empty())    { params.n_threads = cmd_params_defaults.n_threads; }
 
     return params;
@@ -472,6 +483,7 @@ struct cmd_params_instance {
     bool no_kv_offload;
     std::vector<float> tensor_split;
     bool use_mmap;
+    bool embeddings;
 
     llama_model_params to_llama_mparams() const {
         llama_model_params mparams = llama_model_default_params();
@@ -502,6 +514,7 @@ struct cmd_params_instance {
         cparams.type_k = type_k;
         cparams.type_v = type_v;
         cparams.offload_kqv = !no_kv_offload;
+        cparams.embeddings = embeddings;
 
         return cparams;
     }
@@ -517,6 +530,7 @@ static std::vector<cmd_params_instance> get_cmd_params_instances(const cmd_param
     for (const auto & mg : params.main_gpu)
     for (const auto & ts : params.tensor_split)
     for (const auto & mmp : params.use_mmap)
+    for (const auto & embd : params.embeddings)
     for (const auto & nb : params.n_batch)
     for (const auto & tk : params.type_k)
     for (const auto & tv : params.type_v)
@@ -540,6 +554,7 @@ static std::vector<cmd_params_instance> get_cmd_params_instances(const cmd_param
                 /* .no_kv_offload= */ nkvo,
                 /* .tensor_split = */ ts,
                 /* .use_mmap     = */ mmp,
+                /* .embeddings   = */ embd,
             };
             instances.push_back(instance);
         }
@@ -562,6 +577,7 @@ static std::vector<cmd_params_instance> get_cmd_params_instances(const cmd_param
                 /* .no_kv_offload= */ nkvo,
                 /* .tensor_split = */ ts,
                 /* .use_mmap     = */ mmp,
+                /* .embeddings   = */ embd,
             };
             instances.push_back(instance);
         }
@@ -597,6 +613,7 @@ struct test {
     bool no_kv_offload;
     std::vector<float> tensor_split;
     bool use_mmap;
+    bool embeddings;
     int n_prompt;
     int n_gen;
     std::string test_time;
@@ -619,6 +636,7 @@ struct test {
         no_kv_offload = inst.no_kv_offload;
         tensor_split = inst.tensor_split;
         use_mmap = inst.use_mmap;
+        embeddings = inst.embeddings;
         n_prompt = inst.n_prompt;
         n_gen = inst.n_gen;
         // RFC 3339 date-time format
@@ -690,7 +708,7 @@ struct test {
             "n_batch", "n_threads", "type_k", "type_v",
             "n_gpu_layers", "split_mode",
             "main_gpu", "no_kv_offload",
-            "tensor_split", "use_mmap",
+            "tensor_split", "use_mmap", "embeddings",
             "n_prompt", "n_gen", "test_time",
             "avg_ns", "stddev_ns",
             "avg_ts", "stddev_ts"
@@ -710,7 +728,7 @@ struct test {
         }
         if (field == "cuda" || field == "opencl"  || field == "vulkan" || field == "kompute" || field == "metal" ||
             field == "gpu_blas" || field == "blas" || field == "sycl" ||field == "f16_kv" || field == "no_kv_offload" ||
-            field == "use_mmap") {
+            field == "use_mmap" || field == "embeddings") {
             return BOOL;
         }
         if (field == "avg_ts" || field == "stddev_ts") {
@@ -744,7 +762,7 @@ struct test {
             std::to_string(n_batch), std::to_string(n_threads), ggml_type_name(type_k), ggml_type_name(type_v),
             std::to_string(n_gpu_layers), split_mode_str(split_mode),
             std::to_string(main_gpu), std::to_string(no_kv_offload),
-            tensor_split_str, std::to_string(use_mmap),
+            tensor_split_str, std::to_string(use_mmap), std::to_string(embeddings),
             std::to_string(n_prompt), std::to_string(n_gen), test_time,
             std::to_string(avg_ns()), std::to_string(stdev_ns()),
             std::to_string(avg_ts()), std::to_string(stdev_ts())
@@ -914,6 +932,9 @@ struct markdown_printer : public printer {
         if (field == "use_mmap") {
             return "mmap";
         }
+        if (field == "embeddings") {
+            return "embd";
+        }
         if (field == "tensor_split") {
             return "ts";
         }
@@ -957,6 +978,9 @@ struct markdown_printer : public printer {
         if (params.use_mmap.size() > 1 || params.use_mmap != cmd_params_defaults.use_mmap) {
             fields.emplace_back("use_mmap");
         }
+        if (params.embeddings.size() > 1 || params.embeddings != cmd_params_defaults.embeddings) {
+            fields.emplace_back("embeddings");
+        }
         fields.emplace_back("test");
         fields.emplace_back("t/s");
 

examples/server/server.cpp

@@ -335,8 +335,12 @@ struct server_slot {
 };
 
 struct server_metrics {
+    const int64_t t_start = ggml_time_us();
+
     uint64_t n_prompt_tokens_processed_total = 0;
+    uint64_t t_prompt_processing_total       = 0;
     uint64_t n_tokens_predicted_total        = 0;
+    uint64_t t_tokens_generation_total       = 0;
 
     uint64_t n_prompt_tokens_processed = 0;
     uint64_t t_prompt_processing       = 0;
@@ -348,12 +352,14 @@ struct server_metrics {
         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;
+        t_prompt_processing_total       += slot.t_prompt_processing;
     }
 
     void on_prediction(const server_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_total   += slot.n_decoded;
+        n_tokens_predicted         += slot.n_decoded;
+        t_tokens_generation        += slot.t_token_generation;
+        t_tokens_generation_total  += slot.t_token_generation;
     }
 
     void reset_bucket() {
@@ -1123,7 +1129,7 @@ struct server_context {
             });
         }
 
-        if (!slot.cache_tokens.empty() && result.tok == llama_token_eos(model)) {
+        if (result.tok == llama_token_eos(model)) {
             slot.stopped_eos    = true;
             slot.has_next_token = false;
 
@@ -1502,9 +1508,12 @@ struct server_context {
                         { "idle",                            n_idle_slots       },
                         { "processing",                      n_processing_slots },
                         { "deferred",                        queue_tasks.queue_tasks_deferred.size() },
+                        { "t_start",                         metrics.t_start},
 
                         { "n_prompt_tokens_processed_total", metrics.n_prompt_tokens_processed_total},
+                        { "t_tokens_generation_total",       metrics.t_tokens_generation_total},
                         { "n_tokens_predicted_total",        metrics.n_tokens_predicted_total},
+                        { "t_prompt_processing_total",       metrics.t_prompt_processing_total},
 
                         { "n_prompt_tokens_processed",       metrics.n_prompt_tokens_processed},
                         { "t_prompt_processing",             metrics.t_prompt_processing},
@@ -1517,7 +1526,9 @@ struct server_context {
                         { "slots",                           slots_data },
                     };
 
-                    metrics.reset_bucket();
+                    if (json_value(task.data, "reset_bucket", false)) {
+                        metrics.reset_bucket();
+                    }
                     queue_results.send(res);
                 } break;
         }
@@ -2709,6 +2720,7 @@ int main(int argc, char ** argv) {
             task.id_multi  = -1;
             task.id_target = -1;
             task.type = SERVER_TASK_TYPE_METRICS;
+            task.data.push_back({{"reset_bucket", true}});
 
             ctx_server.queue_results.add_waiting_task_id(task.id);
             ctx_server.queue_tasks.post(task);
@@ -2732,20 +2744,28 @@ int main(int argc, char ** argv) {
                 {"counter", {{
                         {"name",  "prompt_tokens_total"},
                         {"help",  "Number of prompt tokens processed."},
-                        {"value",  data["n_prompt_tokens_processed_total"]}
+                        {"value",  (uint64_t) data["n_prompt_tokens_processed_total"]}
+                }, {
+                        {"name",  "prompt_seconds_total"},
+                        {"help",  "Prompt process time"},
+                        {"value",  (uint64_t) data["t_prompt_processing_total"] / 1.e3}
                 }, {
                         {"name",  "tokens_predicted_total"},
                         {"help",  "Number of generation tokens processed."},
-                        {"value",  data["n_tokens_predicted_total"]}
+                        {"value",  (uint64_t) data["n_tokens_predicted_total"]}
+                }, {
+                        {"name",  "tokens_predicted_seconds_total"},
+                        {"help",  "Predict process time"},
+                        {"value",  (uint64_t) data["t_tokens_generation_total"] / 1.e3}
                 }}},
                 {"gauge", {{
                         {"name",  "prompt_tokens_seconds"},
                         {"help",  "Average prompt throughput in tokens/s."},
-                        {"value",  n_prompt_tokens_processed ? 1e3 / t_prompt_processing * n_prompt_tokens_processed : 0}
+                        {"value",  n_prompt_tokens_processed ? 1.e3 / t_prompt_processing * n_prompt_tokens_processed : 0.}
                 },{
                         {"name",  "predicted_tokens_seconds"},
                         {"help",  "Average generation throughput in tokens/s."},
-                        {"value",  n_tokens_predicted ? 1e3 / t_tokens_generation * n_tokens_predicted : 0}
+                        {"value",  n_tokens_predicted ? 1.e3 / t_tokens_generation * n_tokens_predicted : 0.}
                 },{
                         {"name",  "kv_cache_usage_ratio"},
                         {"help",  "KV-cache usage. 1 means 100 percent usage."},
@@ -2753,15 +2773,15 @@ int main(int argc, char ** argv) {
                 },{
                         {"name",  "kv_cache_tokens"},
                         {"help",  "KV-cache tokens."},
-                        {"value",  data["kv_cache_tokens_count"]}
+                        {"value",  (uint64_t) data["kv_cache_tokens_count"]}
                 },{
                         {"name",  "requests_processing"},
                         {"help",  "Number of request processing."},
-                        {"value",  data["processing"]}
+                        {"value",  (uint64_t) data["processing"]}
                 },{
                         {"name",  "requests_deferred"},
                         {"help",  "Number of request deferred."},
-                        {"value",  data["deferred"]}
+                        {"value",  (uint64_t) data["deferred"]}
                 }}}
             };
 
@@ -2775,13 +2795,16 @@ int main(int argc, char ** argv) {
                     const std::string name = metric_def["name"];
                     const std::string help = metric_def["help"];
 
-                    auto value = json_value(metric_def, "value", 0);
+                    auto value = json_value(metric_def, "value", 0.);
                     prometheus << "# HELP llamacpp:" << name << " " << help  << "\n"
                                << "# TYPE llamacpp:" << name << " " << type  << "\n"
                                << "llamacpp:"        << name << " " << value << "\n";
                 }
             }
 
+            const int64_t t_start = data["t_start"];
+            res.set_header("Process-Start-Time-Unix", std::to_string(t_start));
+
             res.set_content(prometheus.str(), "text/plain; version=0.0.4");
             res.status = 200; // HTTP OK
         });

examples/server/tests/features/server.feature

@@ -29,6 +29,7 @@ Feature: llama.cpp server
     And   a completion request with no api error
     Then  <n_predicted> tokens are predicted matching <re_content>
     And   prometheus metrics are exposed
+    And   metric llamacpp:tokens_predicted is <n_predicted>
 
     Examples: Prompts
       | prompt                           | n_predict | re_content                       | n_predicted |

examples/server/tests/features/steps/steps.py

@@ -586,14 +586,24 @@ async def step_prometheus_metrics_exported(context):
             metric_exported = False
             if context.debug:
                 print(f"/metrics answer:\n{metrics_raw}\n")
+            context.metrics = {}
             for metric in parser.text_string_to_metric_families(metrics_raw):
                 match metric.name:
                     case "llamacpp:kv_cache_usage_ratio":
                         assert len(metric.samples) > 0
                         metric_exported = True
+                context.metrics[metric.name] = metric
+            assert int(metrics_response.headers["Process-Start-Time-Unix"]) > 0, "no header process start time"
             assert metric_exported, "No metrics exported"
 
 
+@step(u'metric {metric_name} is {metric_value:d}')
+def step_assert_metric_value(context, metric_name, metric_value):
+    if metric_name not in context.metrics:
+        assert False, f"no metric {metric_name} in {context.metrics.keys()}"
+    assert context.metrics[metric_name].samples[0].value == metric_value, f"metric: {context.metrics[metric_name]}"
+
+
 @step(u'available models')
 def step_available_models(context):
     # openai client always expects an api_key
@@ -879,7 +889,6 @@ def assert_n_tokens_predicted(completion_response, expected_predicted_n=None, re
                                                      f' {n_predicted} <> {expected_predicted_n}')
 
 
-
 async def gather_tasks_results(context):
     n_tasks = len(context.concurrent_tasks)
     if context.debug:

ggml-sycl.cpp

@@ -3559,31 +3559,12 @@ class sycl_gpu_mgr {
         int work_group_size = 0;
         std::string gpus_list = "";
 
-        /*
-        Use all GPU with same top max compute units
-        */
         sycl_gpu_mgr() {
             detect_sycl_gpu_list_with_max_cu();
             get_allow_gpus();
             create_context_with_gpus();
         }
 
-        /*
-        Use the assigned GPU as only one
-        */
-        sycl_gpu_mgr(int main_gpu_id) {
-            sycl::device device = dpct::dev_mgr::instance().get_device(main_gpu_id);
-            dpct::device_info prop;
-            dpct::get_device_info(prop, device);
-            gpus.push_back(main_gpu_id);
-            devices.push_back(device);
-            work_group_size = prop.get_max_work_group_size();
-            max_compute_units = prop.get_max_compute_units();
-
-            get_allow_gpus();
-            create_context_with_gpus();
-        }
-
         void create_context_with_gpus() {
             sycl::context ctx = sycl::context(devices);
             assert(gpus.size() > 0);
@@ -3599,7 +3580,7 @@ class sycl_gpu_mgr {
                 gpus_list += std::to_string(gpus[i]);
                 gpus_list += ",";
             }
-            if (gpus_list.length() > 1) {
+            if (gpus_list.length() > 2) {
                 gpus_list.pop_back();
             }
         }
@@ -3648,8 +3629,8 @@ class sycl_gpu_mgr {
                 if (gpus[i] == id)
                     return i;
             }
-            printf("miss to get device index by id=%d\n", id);
-            GGML_ASSERT(false);
+            assert(false);
+            return -1;
         }
 
         int get_next_index(int id) {
@@ -3658,7 +3639,8 @@ class sycl_gpu_mgr {
                 if (gpus[i] == id)
                     return i;
             }
-            GGML_ASSERT(false);
+            assert(false);
+            return -1;
         }
 };
 
@@ -3667,7 +3649,6 @@ static int g_device_count = -1;
 static int g_all_sycl_device_count = -1;
 static int g_main_device = -1;
 static int g_main_device_id = -1;
-static bool g_ggml_backend_sycl_buffer_type_initialized = false;
 
 static std::array<float, GGML_SYCL_MAX_DEVICES> g_default_tensor_split = {};
 
@@ -13244,7 +13225,7 @@ void ggml_backend_sycl_print_sycl_devices() {
 }
 
 void print_gpu_device_list() {
-    fprintf(stderr, "detect %d SYCL GPUs: [%s] with top Max compute units:%d\n",
+    fprintf(stderr, "detect %d SYCL GPUs: [%s] with Max compute units:%d\n",
             g_sycl_gpu_mgr->get_gpu_count(),
             g_sycl_gpu_mgr->gpus_list.c_str(),
             g_sycl_gpu_mgr->max_compute_units);
@@ -13283,15 +13264,6 @@ void ggml_init_sycl() try {
 #else
         fprintf(stderr, "%s: GGML_SYCL_F16: no\n", __func__);
 #endif
-
-/* NOT REMOVE, keep it for next optimize for XMX.
-#if defined(SYCL_USE_XMX)
-        fprintf(stderr, "%s: SYCL_USE_XMX: yes\n", __func__);
-#else
-        fprintf(stderr, "%s: SYCL_USE_XMX: no\n", __func__);
-#endif
-*/
-
         if (CHECK_TRY_ERROR(g_all_sycl_device_count =
                             dpct::dev_mgr::instance().device_count()) != 0) {
             initialized = true;
@@ -13300,61 +13272,68 @@ void ggml_init_sycl() try {
         }
         GGML_ASSERT(g_all_sycl_device_count <= GGML_SYCL_MAX_DEVICES);
         ggml_backend_sycl_print_sycl_devices();
+
         if (!g_sycl_gpu_mgr) g_sycl_gpu_mgr = new sycl_gpu_mgr();
+
+        g_device_count = g_sycl_gpu_mgr->get_gpu_count();
+        g_work_group_size = g_sycl_gpu_mgr->work_group_size;
+
         print_gpu_device_list();
-        initialized = true;
-        g_sycl_loaded = true;
-    }
 
+        int64_t total_vram = 0;
 
-
-    g_device_count = g_sycl_gpu_mgr->get_gpu_count();
-    g_work_group_size = g_sycl_gpu_mgr->work_group_size;
-
-    int64_t total_vram = 0;
-
-
-    for (int id = 0; id < GGML_SYCL_MAX_DEVICES; ++id) {
-        g_device_caps[id].vmm = 0;
-        g_device_caps[id].device_id = -1;
-        g_device_caps[id].cc = 0;
-        g_tensor_split[id] = 0;
-        g_default_tensor_split[id] = 0;
-    }
-
-    for (int i = 0; i < g_device_count; ++i) {
-        int device_id = g_sycl_gpu_mgr->gpus[i];
-        g_device_caps[i].vmm = 0;
-
-        dpct::device_info prop;
-        SYCL_CHECK(CHECK_TRY_ERROR(dpct::get_device_info(
-            prop, dpct::dev_mgr::instance().get_device(device_id))));
-
-        g_default_tensor_split[i] = total_vram;
-        total_vram += prop.get_global_mem_size();
-
-        g_device_caps[i].cc =
-            100 * prop.get_major_version() + 10 * prop.get_minor_version();
-    }
-
-    for (int i = 0; i < g_device_count; ++i) {
-        g_default_tensor_split[i] /= total_vram;
-    }
-
-    for (int i = 0; i < g_device_count; ++i) {
-        SYCL_CHECK(ggml_sycl_set_device(i));
-
-        // create sycl streams
-        for (int is = 0; is < MAX_STREAMS; ++is) {
-            SYCL_CHECK(CHECK_TRY_ERROR(
-                g_syclStreams[i][is] =
-                    dpct::get_current_device().create_queue(
-                        g_sycl_gpu_mgr->get_co_ctx(), dpct::get_current_device())));
+/* NOT REMOVE, keep it for next optimize for XMX.
+#if defined(SYCL_USE_XMX)
+        fprintf(stderr, "%s: SYCL_USE_XMX: yes\n", __func__);
+#else
+        fprintf(stderr, "%s: SYCL_USE_XMX: no\n", __func__);
+#endif
+*/
+        for (int id = 0; id < GGML_SYCL_MAX_DEVICES; ++id) {
+            g_device_caps[id].vmm = 0;
+            g_device_caps[id].device_id = -1;
+            g_device_caps[id].cc = 0;
+            g_tensor_split[id] = 0;
+            g_default_tensor_split[id] = 0;
         }
 
-        const dpct::queue_ptr stream = g_syclStreams[i][0];
-        // create sycl handle
-        SYCL_CHECK(CHECK_TRY_ERROR(g_sycl_handles[i] = stream));
+        for (int i = 0; i < g_device_count; ++i) {
+            int device_id = g_sycl_gpu_mgr->gpus[i];
+            g_device_caps[i].vmm = 0;
+
+            dpct::device_info prop;
+            SYCL_CHECK(CHECK_TRY_ERROR(dpct::get_device_info(
+                prop, dpct::dev_mgr::instance().get_device(device_id))));
+
+            g_default_tensor_split[i] = total_vram;
+            total_vram += prop.get_global_mem_size();
+
+            g_device_caps[i].cc =
+                100 * prop.get_major_version() + 10 * prop.get_minor_version();
+        }
+
+        for (int i = 0; i < g_device_count; ++i) {
+            g_default_tensor_split[i] /= total_vram;
+        }
+
+        for (int i = 0; i < g_device_count; ++i) {
+            SYCL_CHECK(ggml_sycl_set_device(i));
+
+            // create sycl streams
+            for (int is = 0; is < MAX_STREAMS; ++is) {
+                SYCL_CHECK(CHECK_TRY_ERROR(
+                    g_syclStreams[i][is] =
+                        dpct::get_current_device().create_queue(
+                            g_sycl_gpu_mgr->get_co_ctx(), dpct::get_current_device())));
+            }
+
+            const dpct::queue_ptr stream = g_syclStreams[i][0];
+            // create sycl handle
+            SYCL_CHECK(CHECK_TRY_ERROR(g_sycl_handles[i] = stream));
+        }
+
+        initialized = true;
+        g_sycl_loaded = true;
     }
 }
 catch (sycl::exception const &exc) {
@@ -16753,24 +16732,22 @@ static ggml_backend_buffer_type_i ggml_backend_sycl_buffer_type_interface = {
     /* .is_host          = */ nullptr,
 };
 
-ggml_backend_buffer_type_t ggml_backend_sycl_buffer_type(int device_index) {
-    if (device_index>=g_device_count or device_index<0) {
-        printf("ggml_backend_sycl_buffer_type error: device_index:%d is out of range [0, %d], miss to call ggml_backend_sycl_set_single_device()\n",
-            device_index, g_device_count-1);
-        GGML_ASSERT(device_index<g_device_count);
-    }
+ggml_backend_buffer_type_t ggml_backend_sycl_buffer_type(int device) {
     static struct ggml_backend_buffer_type ggml_backend_sycl_buffer_types[GGML_SYCL_MAX_DEVICES];
 
-    if (!g_ggml_backend_sycl_buffer_type_initialized) {
+    static bool ggml_backend_sycl_buffer_type_initialized = false;
+
+    if (!ggml_backend_sycl_buffer_type_initialized) {
         for (int i = 0; i < g_device_count; i++) {
             ggml_backend_sycl_buffer_types[i] = {
                 /* .iface    = */ ggml_backend_sycl_buffer_type_interface,
                 /* .context  = */ new ggml_backend_sycl_buffer_type_context{i, GGML_SYCL_NAME + std::to_string(g_sycl_gpu_mgr->gpus[i])},
             };
         }
-        g_ggml_backend_sycl_buffer_type_initialized = true;
+        ggml_backend_sycl_buffer_type_initialized = true;
     }
-    return &ggml_backend_sycl_buffer_types[device_index];
+
+    return &ggml_backend_sycl_buffer_types[device];
 }
 
 // sycl split buffer type
@@ -17519,17 +17496,6 @@ GGML_API GGML_CALL int ggml_backend_sycl_get_device_index(int device_id) {
     return g_sycl_gpu_mgr->get_index(device_id);
 }
 
-GGML_API GGML_CALL void ggml_backend_sycl_set_single_device(int main_gpu_id) {
-    GGML_ASSERT(main_gpu_id<g_all_sycl_device_count);
-    printf("ggml_backend_sycl_set_single_device: use single device: %d\n", main_gpu_id);
-    if (g_sycl_gpu_mgr) {
-        delete g_sycl_gpu_mgr;
-    }
-    g_sycl_gpu_mgr = new sycl_gpu_mgr(main_gpu_id);
-    ggml_init_sycl();
-    g_ggml_backend_sycl_buffer_type_initialized = false;
-}
-
 extern "C" int ggml_backend_sycl_reg_devices();
 
 int ggml_backend_sycl_reg_devices() {

ggml-sycl.h

@@ -28,7 +28,6 @@ GGML_API GGML_CALL int   ggml_backend_sycl_get_device_count();
 GGML_API GGML_CALL ggml_backend_buffer_type_t ggml_backend_sycl_split_buffer_type(const float * tensor_split);
 GGML_API GGML_CALL void ggml_backend_sycl_get_device_memory(int device, size_t *free, size_t *total);
 GGML_API GGML_CALL int ggml_backend_sycl_get_device_index(int device_id);
-GGML_API GGML_CALL void ggml_backend_sycl_set_single_device(int main_gpu);
 
 #ifdef  __cplusplus
 }

llama.cpp

@@ -3750,14 +3750,6 @@ static bool llm_load_tensors(
     model.main_gpu     = main_gpu;
     model.n_gpu_layers = n_gpu_layers;
 
-#ifdef GGML_USE_SYCL
-    if (split_mode == LLAMA_SPLIT_MODE_NONE) {
-        ggml_backend_sycl_set_single_device(main_gpu);
-        //SYCL use device index (0, 1, 2), instead if device id.
-        main_gpu = ggml_backend_sycl_get_device_index(main_gpu);
-    }
-#endif
-
     const int64_t n_layer     = hparams.n_layer;
     const int64_t i_gpu_start = std::max((int64_t) hparams.n_layer - n_gpu_layers, (int64_t) 0);
 
@@ -3896,7 +3888,13 @@ static bool llm_load_tensors(
                     {
                         model.output_norm = ml.create_tensor(ctx_output,       tn(LLM_TENSOR_OUTPUT_NORM, "weight"), {n_embd});
                         if (model.arch != LLM_ARCH_MINICPM){
-                            model.output = ml.create_tensor(ctx_output_split, tn(LLM_TENSOR_OUTPUT,      "weight"), {n_embd, n_vocab});
+                            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);
+                            }
                         }
                     }
 
@@ -12268,13 +12266,13 @@ struct llama_context * llama_new_context_with_model(
                 ctx->backends.push_back(backend);
             } else {
                 // LLAMA_SPLIT_LAYER requires a backend for each GPU
-
+                int id_list[GGML_SYCL_MAX_DEVICES];
+                ggml_sycl_get_gpu_list(id_list, GGML_SYCL_MAX_DEVICES);
                 for (int i = 0; i < ggml_backend_sycl_get_device_count(); ++i) {
+                    int device_id = id_list[i];
                     ggml_backend_t backend = ggml_backend_sycl_init(i);
                     if (backend == nullptr) {
-                        int id_list[GGML_SYCL_MAX_DEVICES];
-                        ggml_sycl_get_gpu_list(id_list, GGML_SYCL_MAX_DEVICES);
-                        LLAMA_LOG_ERROR("%s: failed to initialize SYCL%d (index %d)backend\n", __func__, id_list[i], i);
+                        LLAMA_LOG_ERROR("%s: failed to initialize SYCL%d (index %d)backend\n", __func__, device_id, i);
                         llama_free(ctx);
                         return nullptr;
                     }