From 2db9ba1464f3de0aceb2b5289963e69fc369cb66 Mon Sep 17 00:00:00 2001
From: Radoslav Gerganov <rgerganov@gmail.com>
Date: Fri, 18 Apr 2025 10:13:42 +0300
Subject: [PATCH 1/6] rpc : add RPC_CMD_HELLO (#12955)

Add RPC_CMD_HELLO for getting the version of the protocol implemend by
the server. Follow the semantic versioning rules at https://semver.org

Hopefully this bring better user experience when we make breaking
changes at the protocol level and avoid issues like #12465
---
 examples/rpc/rpc-server.cpp    |  5 +++-
 ggml/include/ggml-rpc.h        |  3 ++
 ggml/src/ggml-rpc/ggml-rpc.cpp | 54 +++++++++++++++++++++++++++++++++-
 3 files changed, 60 insertions(+), 2 deletions(-)

diff --git a/examples/rpc/rpc-server.cpp b/examples/rpc/rpc-server.cpp
index f8f2cb90e..9db554257 100644
--- a/examples/rpc/rpc-server.cpp
+++ b/examples/rpc/rpc-server.cpp
@@ -297,7 +297,10 @@ int main(int argc, char * argv[]) {
         }
         cache_dir = cache_dir_str.c_str();
     }
-    printf("Starting RPC server\n");
+    printf("Starting RPC server v%d.%d.%d\n",
+           RPC_PROTO_MAJOR_VERSION,
+           RPC_PROTO_MINOR_VERSION,
+           RPC_PROTO_PATCH_VERSION);
     printf("  endpoint       : %s\n", endpoint.c_str());
     printf("  local cache    : %s\n", cache_dir ? cache_dir : "n/a");
     printf("  backend memory : %zu MB\n", free_mem / (1024 * 1024));
diff --git a/ggml/include/ggml-rpc.h b/ggml/include/ggml-rpc.h
index 4e0d210f8..c8b6097f7 100644
--- a/ggml/include/ggml-rpc.h
+++ b/ggml/include/ggml-rpc.h
@@ -7,6 +7,9 @@
 extern "C" {
 #endif
 
+#define RPC_PROTO_MAJOR_VERSION    1
+#define RPC_PROTO_MINOR_VERSION    0
+#define RPC_PROTO_PATCH_VERSION    0
 #define GGML_RPC_MAX_SERVERS       16
 
 // backend API
diff --git a/ggml/src/ggml-rpc/ggml-rpc.cpp b/ggml/src/ggml-rpc/ggml-rpc.cpp
index 3189ae85d..a0667b7d7 100644
--- a/ggml/src/ggml-rpc/ggml-rpc.cpp
+++ b/ggml/src/ggml-rpc/ggml-rpc.cpp
@@ -92,12 +92,19 @@ enum rpc_cmd {
     RPC_CMD_GET_DEVICE_MEMORY,
     RPC_CMD_INIT_TENSOR,
     RPC_CMD_GET_ALLOC_SIZE,
+    RPC_CMD_HELLO,
     RPC_CMD_COUNT,
 };
 
 // Try RPC_CMD_SET_TENSOR_HASH first when data size is larger than this threshold
 const size_t HASH_THRESHOLD = 10 * 1024 * 1024;
 
+struct rpc_msg_hello_rsp {
+    uint8_t major;
+    uint8_t minor;
+    uint8_t patch;
+};
+
 struct rpc_msg_get_alloc_size_req {
     rpc_tensor tensor;
 };
@@ -400,6 +407,20 @@ static bool send_rpc_cmd(const std::shared_ptr<socket_t> & sock, enum rpc_cmd cm
 
 // RPC client-side implementation
 
+static bool check_server_version(const std::shared_ptr<socket_t> & sock) {
+    rpc_msg_hello_rsp response;
+    bool status = send_rpc_cmd(sock, RPC_CMD_HELLO, nullptr, 0, &response, sizeof(response));
+    GGML_ASSERT(status);
+    if (response.major != RPC_PROTO_MAJOR_VERSION || response.minor > RPC_PROTO_MINOR_VERSION) {
+        fprintf(stderr, "RPC server version mismatch: %d.%d.%d\n", response.major, response.minor, response.patch);
+        return false;
+    }
+    if (response.minor != RPC_PROTO_MINOR_VERSION || response.patch != RPC_PROTO_PATCH_VERSION) {
+        fprintf(stderr, "WARNING: RPC server version mismatch: %d.%d.%d\n", response.major, response.minor, response.patch);
+    }
+    return true;
+}
+
 static std::shared_ptr<socket_t> get_socket(const std::string & endpoint) {
     static std::mutex mutex;
     std::lock_guard<std::mutex> lock(mutex);
@@ -433,6 +454,9 @@ static std::shared_ptr<socket_t> get_socket(const std::string & endpoint) {
     if (sock == nullptr) {
         return nullptr;
     }
+    if (!check_server_version(sock)) {
+        return nullptr;
+    }
     GGML_PRINT_DEBUG("[%s] connected to %s, sockfd=%d\n", __func__, endpoint.c_str(), sock->fd);
     sockets[endpoint] = sock;
     return sock;
@@ -818,6 +842,7 @@ public:
     }
     ~rpc_server();
 
+    void hello(rpc_msg_hello_rsp & response);
     void alloc_buffer(const rpc_msg_alloc_buffer_req & request, rpc_msg_alloc_buffer_rsp & response);
     void get_alignment(rpc_msg_get_alignment_rsp & response);
     void get_max_size(rpc_msg_get_max_size_rsp & response);
@@ -846,6 +871,13 @@ private:
     std::unordered_set<ggml_backend_buffer_t> buffers;
 };
 
+void rpc_server::hello(rpc_msg_hello_rsp & response) {
+    response.major = RPC_PROTO_MAJOR_VERSION;
+    response.minor = RPC_PROTO_MINOR_VERSION;
+    response.patch = RPC_PROTO_PATCH_VERSION;
+    GGML_PRINT_DEBUG("[%s] version: %d.%d.%d\n", __func__, response.major, response.minor, response.patch);
+}
+
 bool rpc_server::get_alloc_size(const rpc_msg_get_alloc_size_req & request, rpc_msg_get_alloc_size_rsp & response) {
     ggml_backend_buffer_type_t buft;
     struct ggml_init_params params {
@@ -1271,8 +1303,24 @@ rpc_server::~rpc_server() {
 static void rpc_serve_client(ggml_backend_t backend, const char * cache_dir,
                              sockfd_t sockfd, size_t free_mem, size_t total_mem) {
     rpc_server server(backend, cache_dir);
+    uint8_t cmd;
+    if (!recv_data(sockfd, &cmd, 1)) {
+        return;
+    }
+    // the first command sent by the client must be HELLO
+    if (cmd != RPC_CMD_HELLO) {
+        fprintf(stderr, "Expected HELLO command, update client\n");
+        return;
+    }
+    if (!recv_msg(sockfd, nullptr, 0)) {
+        return;
+    }
+    rpc_msg_hello_rsp response;
+    server.hello(response);
+    if (!send_msg(sockfd, &response, sizeof(response))) {
+        return;
+    }
     while (true) {
-        uint8_t cmd;
         if (!recv_data(sockfd, &cmd, 1)) {
             break;
         }
@@ -1282,6 +1330,10 @@ static void rpc_serve_client(ggml_backend_t backend, const char * cache_dir,
             break;
         }
         switch (cmd) {
+            case RPC_CMD_HELLO: {
+                // HELLO command is handled above
+                return;
+            }
             case RPC_CMD_ALLOC_BUFFER: {
                 rpc_msg_alloc_buffer_req request;
                 if (!recv_msg(sockfd, &request, sizeof(request))) {

From b9154ecff93ff54dc554411eb844a2a654be49f2 Mon Sep 17 00:00:00 2001
From: Xuan-Son Nguyen <son@huggingface.co>
Date: Fri, 18 Apr 2025 10:04:51 +0200
Subject: [PATCH 2/6] mtmd : add methods to access `mtmd_image_tokens` (#12906)

* mtmd : add more api around mtmd_image_tokens

* mtmd : ability to calc image hash

* shared_ptr for mtmd_image_tokens

* move hash to user-define ID (fixed)

* fix prompt_modified

* rm redundant data member
---
 examples/llava/gemma3-cli.cpp | 11 +++--
 examples/llava/mtmd.cpp       | 88 ++++++++++++++++++++++++-----------
 examples/llava/mtmd.h         | 37 ++++++++++-----
 3 files changed, 92 insertions(+), 44 deletions(-)

diff --git a/examples/llava/gemma3-cli.cpp b/examples/llava/gemma3-cli.cpp
index 693b738a8..3d5664750 100644
--- a/examples/llava/gemma3-cli.cpp
+++ b/examples/llava/gemma3-cli.cpp
@@ -184,18 +184,19 @@ static int eval_message(gemma3_context & ctx, common_chat_msg & msg, std::vector
     text.text          = formatted_chat.prompt;
     text.add_special   = add_bos;
     text.parse_special = true;
-    mtmd_input_chunks_ptr chunks(mtmd_tokenize(ctx.ctx_vision.get(), text, bitmaps));
-    if (chunks == nullptr) {
-        LOG_ERR("Unable to tokenize prompt\n");
+    mtmd_input_chunks chunks;
+    int32_t res = mtmd_tokenize(ctx.ctx_vision.get(), chunks, text, bitmaps);
+    if (res != 0) {
+        LOG_ERR("Unable to tokenize prompt, res = %d\n", res);
         return 1;
     }
 
-    if (mtmd_helper_eval(ctx.ctx_vision.get(), ctx.lctx, chunks.get(), ctx.n_past, 0, ctx.n_batch)) {
+    if (mtmd_helper_eval(ctx.ctx_vision.get(), ctx.lctx, chunks, ctx.n_past, 0, ctx.n_batch)) {
         LOG_ERR("Unable to eval prompt\n");
         return 1;
     }
 
-    ctx.n_past += mtmd_helper_get_n_tokens(chunks.get());
+    ctx.n_past += mtmd_helper_get_n_tokens(chunks);
 
     return 0;
 }
diff --git a/examples/llava/mtmd.cpp b/examples/llava/mtmd.cpp
index 114c274bc..3fd5bebc6 100644
--- a/examples/llava/mtmd.cpp
+++ b/examples/llava/mtmd.cpp
@@ -16,6 +16,7 @@ struct mtmd_context {
     struct clip_ctx * ctx_clip;
     const struct llama_model * text_model;
     std::vector<float> image_embd_v; // image embedding vector
+
     bool print_timings;
     int n_threads;
     std::string image_marker;
@@ -24,7 +25,11 @@ struct mtmd_context {
 
     mtmd_context(const char * mmproj_fname,
                    const llama_model * text_model,
-                   const mtmd_context_params & ctx_params) : print_timings(ctx_params.print_timings), n_threads(ctx_params.n_threads), image_marker(ctx_params.image_marker) {
+                   const mtmd_context_params & ctx_params) :
+        print_timings(ctx_params.print_timings),
+        n_threads    (ctx_params.n_threads),
+        image_marker (ctx_params.image_marker)
+    {
         clip_context_params ctx_clip_params;
         ctx_clip_params.use_gpu   = ctx_params.use_gpu;
         ctx_clip_params.verbosity = ctx_params.verbosity;
@@ -49,6 +54,7 @@ struct mtmd_image_tokens {
     uint32_t ny; // number of tokens in y direction
     uint32_t n_tokens() const { return nx * ny; }
     clip_image_f32_batch batch_f32; // preprocessed image patches
+    std::string id; // optional user-defined ID, useful for KV cache tracking
 };
 
 mtmd_context * mtmd_init_from_file(const char * mmproj_fname,
@@ -88,10 +94,10 @@ static std::vector<llama_token> mtmd_tokenize_text_internal(
     return result;
 }
 
-mtmd_input_chunks * mtmd_tokenize(mtmd_context * ctx,
-                                const mtmd_input_text & text,
-                                const std::vector<mtmd_bitmap> & bitmaps) {
-    mtmd_input_chunks * output = new mtmd_input_chunks;
+int32_t mtmd_tokenize(mtmd_context * ctx,
+                        std::vector<mtmd_input_chunk> & output,
+                        const mtmd_input_text & text,
+                        const std::vector<mtmd_bitmap> & bitmaps) {
     auto vocab = llama_model_get_vocab(ctx->text_model);
 
     std::string prompt_modified(text.text);
@@ -105,9 +111,9 @@ mtmd_input_chunks * mtmd_tokenize(mtmd_context * ctx,
         string_replace_all(prompt_modified, ctx->image_marker, marker_modified);
     }
 
-    std::vector<std::string> parts = string_split_str(text.text, ctx->image_marker);
-    output->clear();
-    output->reserve(parts.size());
+    std::vector<std::string> parts = string_split_str(prompt_modified, ctx->image_marker);
+    output.clear();
+    output.reserve(parts.size());
 
     size_t i_img = 0;
 
@@ -123,14 +129,14 @@ mtmd_input_chunks * mtmd_tokenize(mtmd_context * ctx,
             std::move(tokens),
             {},
         };
-        output->emplace_back(std::move(chunk));
+        output.emplace_back(std::move(chunk));
 
         if (&parts.back() != &part) {
             // add image token to middle of 2 parts
 
             if (i_img >= bitmaps.size()) {
                 LOG_ERR("%s: error: not enough images for %d parts\n", __func__, (int)parts.size());
-                return nullptr;
+                return 1;
             }
 
             // shim layer
@@ -145,34 +151,48 @@ mtmd_input_chunks * mtmd_tokenize(mtmd_context * ctx,
             bool ok = clip_image_preprocess(ctx->ctx_clip, img_u8.get(), &batch_f32);
             if (!ok) {
                 LOG_ERR("Unable to preprocess image\n");
-                return nullptr;
+                return 2;
             }
 
-            mtmd_image_tokens * image_tokens = new mtmd_image_tokens;
+            mtmd_image_tokens_ptr image_tokens(new mtmd_image_tokens);
             image_tokens->nx = clip_n_patches(ctx->ctx_clip); // TODO @ngxson : use clip_n_patches_by_image
             image_tokens->ny = 1; // TODO
             image_tokens->batch_f32 = std::move(batch_f32);
+            image_tokens->id = bitmaps[i_img].id; // optional
 
             mtmd_input_chunk chunk{
                 MTMD_INPUT_CHUNK_TYPE_IMAGE,
                 {},
-                image_tokens,
+                std::move(image_tokens),
             };
-            output->emplace_back(std::move(chunk));
+            output.emplace_back(std::move(chunk));
             i_img++;
         }
     }
 
-    return output;
+    return 0;
 }
 
-void mtmd_input_chunks_free(mtmd_input_chunks * chunks) {
-    for (auto & chunk : *chunks) {
-        if (chunk.type == MTMD_INPUT_CHUNK_TYPE_IMAGE && chunk.tokens_image) {
-            delete chunk.tokens_image;
-        }
+void mtmd_image_tokens_free(mtmd_image_tokens * image_tokens) {
+    if (image_tokens) {
+        delete image_tokens;
     }
-    delete chunks;
+}
+
+size_t mtmd_image_tokens_get_n_tokens(const mtmd_image_tokens * image_tokens) {
+    return image_tokens->n_tokens();
+}
+
+size_t mtmd_image_tokens_get_nx(const mtmd_image_tokens * image_tokens) {
+    return image_tokens->nx;
+}
+
+size_t mtmd_image_tokens_get_ny(const mtmd_image_tokens * image_tokens) {
+    return image_tokens->ny;
+}
+
+std::string mtmd_image_tokens_get_id(const mtmd_image_tokens * image_tokens) {
+    return image_tokens->id;
 }
 
 int32_t mtmd_encode(mtmd_context * ctx, const mtmd_image_tokens * image_tokens) {
@@ -190,9 +210,9 @@ float * mtmd_get_output_embd(mtmd_context * ctx) {
     return ctx->image_embd_v.data();
 }
 
-size_t mtmd_helper_get_n_tokens(mtmd_input_chunks * chunks) {
+size_t mtmd_helper_get_n_tokens(mtmd_input_chunks & chunks) {
     size_t n_tokens = 0;
-    for (auto & chunk : *chunks) {
+    for (auto & chunk : chunks) {
         if (chunk.type == MTMD_INPUT_CHUNK_TYPE_TEXT) {
             n_tokens += chunk.tokens_text.size();
         } else if (chunk.type == MTMD_INPUT_CHUNK_TYPE_IMAGE) {
@@ -241,7 +261,7 @@ struct decode_embd_batch {
 
 int32_t mtmd_helper_eval(mtmd_context * ctx,
         llama_context * lctx,
-        mtmd_input_chunks * chunks,
+        mtmd_input_chunks & chunks,
         llama_pos pos0,
         llama_seq_id seq_id,
         int32_t n_batch) {
@@ -249,8 +269,8 @@ int32_t mtmd_helper_eval(mtmd_context * ctx,
     llama_pos n_past = pos0;
     llama_batch text_batch = llama_batch_init(n_batch, 0, 1);
 
-    for (auto & chunk : *chunks) {
-        bool is_last = &chunk == &chunks->back();
+    for (auto & chunk : chunks) {
+        bool is_last = &chunk == &chunks.back();
         if (chunk.type == MTMD_INPUT_CHUNK_TYPE_TEXT) {
             // TODO @ngxson : may need to split into smaller batches
             text_batch.n_tokens = chunk.tokens_text.size();
@@ -279,7 +299,7 @@ int32_t mtmd_helper_eval(mtmd_context * ctx,
             if (ctx->print_timings) {
                 LOG_INF("encoding image...\n");
             }
-            ret = mtmd_encode(ctx, chunk.tokens_image);
+            ret = mtmd_encode(ctx, chunk.tokens_image.get());
             if (ret != 0) {
                 LOG_ERR("failed to encode image\n");
                 llama_batch_free(text_batch);
@@ -289,7 +309,7 @@ int32_t mtmd_helper_eval(mtmd_context * ctx,
                 LOG_INF("image encoded in %" PRId64 " ms\n", ggml_time_ms() - t0);
             }
 
-            int32_t n_tokens = chunk.tokens_image->n_tokens();
+            int32_t n_tokens = mtmd_image_tokens_get_n_tokens(chunk.tokens_image.get());
             float * embd = mtmd_get_output_embd(ctx);
             decode_embd_batch batch_img(embd, n_tokens, n_past, 0);
             int64_t t1 = ggml_time_ms();
@@ -339,3 +359,15 @@ int32_t mtmd_helper_bitmap_init_from_file(const char * fname, mtmd_bitmap & outp
     std::memcpy(output.data.data(), data, output.nx * output.ny * 3);
     return 0;
 }
+
+bool mtmd_decode_use_non_causal(mtmd_context * ctx) {
+    projector_type proj_type = clip_get_projector_type(ctx->ctx_clip);
+    if (proj_type == PROJECTOR_TYPE_GEMMA3) {
+        return true;
+    }
+    return false;
+}
+
+void mtmd_image_tokens_deleter::operator()(mtmd_image_tokens * val) {
+    mtmd_image_tokens_free(val);
+}
diff --git a/examples/llava/mtmd.h b/examples/llava/mtmd.h
index 598f6947b..78be192dd 100644
--- a/examples/llava/mtmd.h
+++ b/examples/llava/mtmd.h
@@ -39,12 +39,18 @@ struct mtmd_bitmap {
     uint32_t nx;
     uint32_t ny;
     std::vector<unsigned char> data;
+    std::string id; // optional user-defined id, for ex: can be set to image hash, useful for KV cache tracking
 };
 
+struct mtmd_image_tokens_deleter {
+    void operator()(mtmd_image_tokens * val); // forward declaration
+};
+using mtmd_image_tokens_ptr = std::unique_ptr<mtmd_image_tokens, mtmd_image_tokens_deleter>;
+
 struct mtmd_input_chunk {
     mtmd_input_chunk_type type;
     std::vector<llama_token> tokens_text;
-    mtmd_image_tokens * tokens_image = nullptr;
+    mtmd_image_tokens_ptr tokens_image;
 };
 
 using mtmd_input_chunks = std::vector<mtmd_input_chunk>;
@@ -82,12 +88,21 @@ MTMD_API void mtmd_free(mtmd_context * ctx);
 //   3. "<end_of_image>\ndescribe it in detail."
 // number of bitmaps must be equal to the number of image markers in the prompt
 // this function is thread-safe (shared ctx)
-MTMD_API mtmd_input_chunks * mtmd_tokenize(mtmd_context * ctx,
+// return values:
+//   0 on success
+//   1 on number of images not matching the number of markers
+//   2 on image preprocessing error
+MTMD_API int32_t mtmd_tokenize(mtmd_context * ctx,
+                                std::vector<mtmd_input_chunk> & output,
                                 const mtmd_input_text & text,
                                 const std::vector<mtmd_bitmap> & bitmaps);
 
-// free image chunk data
-MTMD_API void mtmd_input_chunks_free(mtmd_input_chunks * chunks);
+// access mtmd_image_tokens
+MTMD_API size_t      mtmd_image_tokens_get_n_tokens(const mtmd_image_tokens * image_tokens);
+MTMD_API size_t      mtmd_image_tokens_get_nx(const mtmd_image_tokens * image_tokens);
+MTMD_API size_t      mtmd_image_tokens_get_ny(const mtmd_image_tokens * image_tokens);
+MTMD_API std::string mtmd_image_tokens_get_id(const mtmd_image_tokens * image_tokens);
+MTMD_API void        mtmd_image_tokens_free(mtmd_image_tokens * image_tokens);
 
 // returns 0 on success
 MTMD_API int32_t mtmd_encode(mtmd_context * ctx,
@@ -96,12 +111,17 @@ MTMD_API int32_t mtmd_encode(mtmd_context * ctx,
 // get output embeddings from the last encode pass
 MTMD_API float * mtmd_get_output_embd(mtmd_context * ctx);
 
+// whether we need to set non-causal mask before llama_decode
+MTMD_API bool mtmd_decode_use_non_causal(mtmd_context * ctx);
+
+
+
 //
 // helper functions (can be implemented based on other functions)
 //
 
 // helper to count the total number of tokens from a list of chunks, useful to keep track of n_past
-MTMD_API size_t mtmd_helper_get_n_tokens(mtmd_input_chunks * chunks);
+MTMD_API size_t mtmd_helper_get_n_tokens(mtmd_input_chunks & chunks);
 
 // helper function that automatically:
 // 1. run llama_decode() on text chunks
@@ -110,7 +130,7 @@ MTMD_API size_t mtmd_helper_get_n_tokens(mtmd_input_chunks * chunks);
 // otherwise, returns 0 on success
 MTMD_API int32_t mtmd_helper_eval(mtmd_context * ctx,
                                 llama_context * lctx,
-                                mtmd_input_chunks * chunks,
+                                mtmd_input_chunks & chunks,
                                 llama_pos pos0,
                                 llama_seq_id seq_id,
                                 int32_t n_batch);
@@ -132,11 +152,6 @@ struct mtmd_context_deleter {
 };
 using mtmd_context_ptr = std::unique_ptr<mtmd_context, mtmd_context_deleter>;
 
-struct mtmd_input_chunks_deleter {
-    void operator()(mtmd_input_chunks * val) { mtmd_input_chunks_free(val); }
-};
-using mtmd_input_chunks_ptr = std::unique_ptr<mtmd_input_chunks, mtmd_input_chunks_deleter>;
-
 #else
 
 static_assert(false && "C header is not yet supported by this library");

From 8d6600576318dfc6b091fca744b0fd36a5e5255f Mon Sep 17 00:00:00 2001
From: Akarshan Biswas <akarshan@menlo.ai>
Date: Fri, 18 Apr 2025 19:27:56 +0530
Subject: [PATCH 3/6] SYCL: Refactor and enable FP16 in binary broadcast OPs
 (#12975)

* SYCL: refactor move to a separate file

* Fix binbcast

* Remove duplicates

* fix include formatting

* fix typo
---
 ggml/src/ggml-sycl/backend.hpp      |   1 +
 ggml/src/ggml-sycl/binbcast.cpp     | 350 ++++++++++++++++++++++++++++
 ggml/src/ggml-sycl/binbcast.hpp     |  39 ++++
 ggml/src/ggml-sycl/common.hpp       | 281 ----------------------
 ggml/src/ggml-sycl/element_wise.cpp |  47 ----
 ggml/src/ggml-sycl/element_wise.hpp |  32 +--
 ggml/src/ggml-sycl/ggml-sycl.cpp    |  15 +-
 7 files changed, 393 insertions(+), 372 deletions(-)
 create mode 100644 ggml/src/ggml-sycl/binbcast.cpp
 create mode 100644 ggml/src/ggml-sycl/binbcast.hpp

diff --git a/ggml/src/ggml-sycl/backend.hpp b/ggml/src/ggml-sycl/backend.hpp
index 73d807cab..de814ef91 100644
--- a/ggml/src/ggml-sycl/backend.hpp
+++ b/ggml/src/ggml-sycl/backend.hpp
@@ -13,6 +13,7 @@
 #ifndef GGML_SYCL_BACKEND_HPP
 #define GGML_SYCL_BACKEND_HPP
 
+#include "binbcast.hpp"
 #include "concat.hpp"
 #include "common.hpp"
 #include "conv.hpp"
diff --git a/ggml/src/ggml-sycl/binbcast.cpp b/ggml/src/ggml-sycl/binbcast.cpp
new file mode 100644
index 000000000..0a9d3a927
--- /dev/null
+++ b/ggml/src/ggml-sycl/binbcast.cpp
@@ -0,0 +1,350 @@
+#include "binbcast.hpp"
+
+#include <cstddef>
+#include <cstdint>
+#include <sycl/sycl.hpp>
+
+#include "ggml.h"
+
+template<float (*bin_op)(const float, const float), typename src0_t, typename src1_t, typename dst_t>
+static void k_bin_bcast(const src0_t * src0, const src1_t * src1, dst_t * dst,
+        int ne0, int ne1, int ne2, int ne3,
+        int ne10, int ne11, int ne12, int ne13,
+        /*int s0, */ int s1,  int s2,  int s3,
+        /*int s00,*/ int s01, int s02, int s03,
+        /*int s10,*/ int s11, int s12, int s13,
+        const sycl::nd_item<3> &item_ct1) {
+    const int i0s = item_ct1.get_local_range(2) * item_ct1.get_group(2) +
+                    item_ct1.get_local_id(2);
+    const int i1 = (item_ct1.get_local_range(1) * item_ct1.get_group(1) +
+                    item_ct1.get_local_id(1));
+    const int i2 = (item_ct1.get_local_range(0) * item_ct1.get_group(0) +
+                    item_ct1.get_local_id(0)) /
+                   ne3;
+    const int i3 = (item_ct1.get_local_range(0) * item_ct1.get_group(0) +
+                    item_ct1.get_local_id(0)) %
+                   ne3;
+
+    if (i0s >= ne0 || i1 >= ne1 || i2 >= ne2 || i3 >= ne3) {
+        return;
+    }
+
+    const int i11 = i1 % ne11;
+    const int i12 = i2 % ne12;
+    const int i13 = i3 % ne13;
+
+    const size_t i_src0 =  i3*s03 +  i2*s02 +  i1*s01;
+    const size_t i_src1 = i13*s13 + i12*s12 + i11*s11;
+    const size_t i_dst  =  i3*s3  +  i2*s2  +  i1*s1;
+
+    const src0_t * src0_row = src0 + i_src0;
+    const src1_t * src1_row = src1 + i_src1;
+    dst_t * dst_row = dst + i_dst;
+
+    for (int i0 = i0s; i0 < ne0;
+         i0 += item_ct1.get_local_range(2) * item_ct1.get_group_range(2)) {
+        const int i10 = i0 % ne10;
+        dst_row[i0] = (dst_t)bin_op(src0 ? (float)src0_row[i0] : 0.0f, (float)src1_row[i10]);
+    }
+}
+
+template<float (*bin_op)(const float, const float), typename src0_t, typename src1_t, typename dst_t>
+static void k_bin_bcast_unravel(const src0_t * src0, const src1_t * src1, dst_t * dst,
+        int ne0, int ne1, int ne2, int ne3,
+        int ne10, int ne11, int ne12, int ne13,
+        /*int s0, */ int s1,  int s2,  int s3,
+        /*int s00,*/ int s01, int s02, int s03,
+        /*int s10,*/ int s11, int s12, int s13,
+        const sycl::nd_item<3> &item_ct1) {
+
+    const int i = item_ct1.get_local_range(2) * item_ct1.get_group(2) +
+                  item_ct1.get_local_id(2);
+
+    const int i3 = i/(ne2*ne1*ne0);
+    const int i2 = (i/(ne1*ne0)) % ne2;
+    const int i1 = (i/ne0) % ne1;
+    const int i0 = i % ne0;
+
+    if (i0 >= ne0 || i1 >= ne1 || i2 >= ne2 || i3 >= ne3) {
+        return;
+    }
+
+    const int i11 = i1 % ne11;
+    const int i12 = i2 % ne12;
+    const int i13 = i3 % ne13;
+
+    const size_t i_src0 =  i3*s03 +  i2*s02 +  i1*s01;
+    const size_t i_src1 = i13*s13 + i12*s12 + i11*s11;
+    const size_t i_dst  =  i3*s3  +  i2*s2  +  i1*s1;
+
+    const src0_t * src0_row = src0 + i_src0;
+    const src1_t * src1_row = src1 + i_src1;
+    dst_t * dst_row = dst + i_dst;
+
+    const int i10 = i0 % ne10;
+    dst_row[i0] = (dst_t)bin_op(src0 ? (float)src0_row[i0] : 0.0f, (float)src1_row[i10]);
+}
+
+
+template<float (*bin_op)(const float, const float)>
+struct bin_bcast_sycl {
+    template <typename src0_t, typename src1_t, typename dst_t>
+    void operator()(const src0_t * src0_dd, const src1_t * src1_dd, dst_t * dst_dd, const int64_t ne00,
+                    const int64_t ne01, const int64_t ne02, const int64_t ne03, const int64_t ne10, const int64_t ne11,
+                    const int64_t ne12, const int64_t ne13, const int64_t ne0, const int64_t ne1, const int64_t ne2,
+                    const int64_t ne3, const size_t nb00, const size_t nb01, const size_t nb02, const size_t nb03,
+                    const size_t nb10, const size_t nb11, const size_t nb12, const size_t nb13, const size_t nb0,
+                    const size_t nb1, const size_t nb2, const size_t nb3, const bool src0_is_contiguous,
+                    const bool src1_is_contiguous, const bool dst_is_contiguous, queue_ptr stream) {
+        int nr0 = ne10 / ne0;
+        int nr1 = ne11/ne1;
+        int nr2 = ne12/ne2;
+        int nr3 = ne13/ne3;
+
+        int nr[4] = { nr0, nr1, nr2, nr3 };
+
+        // collapse dimensions until first broadcast dimension
+        int64_t cne[] = {ne0, ne1, ne2, ne3};
+        int64_t cne0[] = {ne00, ne01, ne02, ne03};
+        int64_t cne1[] = {ne10, ne11, ne12, ne13};
+        size_t cnb[] = {nb0, nb1, nb2, nb3};
+        size_t cnb0[] = {nb00, nb01, nb02, nb03};
+        size_t cnb1[] = {nb10, nb11, nb12, nb13};
+        auto collapse = [](int64_t cne[]) {
+            cne[0] *= cne[1];
+            cne[1] = cne[2];
+            cne[2] = cne[3];
+            cne[3] = 1;
+        };
+
+        auto collapse_nb = [](size_t cnb[], int64_t cne[]) {
+            cnb[1] *= cne[1];
+            cnb[2] *= cne[2];
+            cnb[3] *= cne[3];
+        };
+
+        if (src0_is_contiguous && src1_is_contiguous && dst_is_contiguous) {
+            for (int i = 0; i < 4; i++) {
+                if (nr[i] != 1) {
+                    break;
+                }
+                if (i > 0) {
+                    collapse_nb(cnb, cne);
+                    collapse_nb(cnb0, cne0);
+                    collapse_nb(cnb1, cne1);
+                    collapse(cne);
+                    collapse(cne0);
+                    collapse(cne1);
+                }
+            }
+        }
+        {
+            int64_t ne0 = cne[0];
+            int64_t ne1 = cne[1];
+            int64_t ne2 = cne[2];
+            int64_t ne3 = cne[3];
+
+            int64_t ne10 = cne1[0];
+            int64_t ne11 = cne1[1];
+            int64_t ne12 = cne1[2];
+            int64_t ne13 = cne1[3];
+
+            size_t nb0 = cnb[0];
+            size_t nb1 = cnb[1];
+            size_t nb2 = cnb[2];
+            size_t nb3 = cnb[3];
+
+            size_t nb00 = cnb0[0];
+            size_t nb01 = cnb0[1];
+            size_t nb02 = cnb0[2];
+            size_t nb03 = cnb0[3];
+
+            size_t nb10 = cnb1[0];
+            size_t nb11 = cnb1[1];
+            size_t nb12 = cnb1[2];
+            size_t nb13 = cnb1[3];
+
+            size_t s0 = nb0 / sizeof(dst_t);
+            size_t s1 = nb1 / sizeof(dst_t);
+            size_t s2 = nb2 / sizeof(dst_t);
+            size_t s3 = nb3 / sizeof(dst_t);
+
+            size_t s10 = nb10 / sizeof(src1_t);
+            size_t s11 = nb11 / sizeof(src1_t);
+            size_t s12 = nb12 / sizeof(src1_t);
+            size_t s13 = nb13 / sizeof(src1_t);
+
+            size_t s00 = nb00 / sizeof(src0_t);
+            size_t s01 = nb01 / sizeof(src0_t);
+            size_t s02 = nb02 / sizeof(src0_t);
+            size_t s03 = nb03 / sizeof(src0_t);
+
+            GGML_UNUSED(s00);
+
+            GGML_ASSERT(nb0 % sizeof(dst_t) == 0);
+            GGML_ASSERT(nb1 % sizeof(dst_t) == 0);
+            GGML_ASSERT(nb2 % sizeof(dst_t) == 0);
+            GGML_ASSERT(nb3 % sizeof(dst_t) == 0);
+
+            GGML_ASSERT(nb00 % sizeof(src0_t) == 0);
+            GGML_ASSERT(nb01 % sizeof(src0_t) == 0);
+            GGML_ASSERT(nb02 % sizeof(src0_t) == 0);
+            GGML_ASSERT(nb03 % sizeof(src0_t) == 0);
+
+            GGML_ASSERT(nb10 % sizeof(src1_t) == 0);
+            GGML_ASSERT(nb11 % sizeof(src1_t) == 0);
+            GGML_ASSERT(nb12 % sizeof(src1_t) == 0);
+            GGML_ASSERT(nb13 % sizeof(src1_t) == 0);
+
+            GGML_ASSERT(s0 == 1);
+            GGML_ASSERT(s10 == 1);
+
+            const int block_size = 128;
+
+            int64_t hne0 = std::max(ne0/2LL, 1LL);
+
+            sycl::range<3> block_dims(1, 1, 1);
+            block_dims[2] = std::min<unsigned int>(hne0, block_size);
+            block_dims[1] = std::min<unsigned int>(
+                ne1, block_size / (unsigned int)block_dims[2]);
+            block_dims[0] = std::min(
+                std::min<unsigned int>(
+                    ne2 * ne3, block_size / (unsigned int)block_dims[2] /
+                                   (unsigned int)block_dims[1]),
+                64U);
+
+            sycl::range<3> block_nums(
+                (ne2 * ne3 + block_dims[0] - 1) / block_dims[0],
+                (ne1 + block_dims[1] - 1) / block_dims[1],
+                (hne0 + block_dims[2] - 1) / block_dims[2]);
+
+            if (block_nums[0] > 65535) {
+                // this is the maximum number of blocks in z direction, fallback to 1D grid kernel
+                int block_num = (ne0*ne1*ne2*ne3 + block_size - 1) / block_size;
+                {
+                    dpct::has_capability_or_fail(stream->get_device(),
+                                                 {sycl::aspect::fp16});
+
+                    stream->parallel_for(
+                        sycl::nd_range<3>(sycl::range<3>(1, 1, block_num) *
+                                              sycl::range<3>(1, 1, block_size),
+                                          sycl::range<3>(1, 1, block_size)),
+                        [=](sycl::nd_item<3> item_ct1) {
+                            k_bin_bcast_unravel<bin_op>(
+                                src0_dd, src1_dd, dst_dd, ne0, ne1, ne2, ne3,
+                                ne10, ne11, ne12, ne13, s1, s2, s3, s01, s02,
+                                s03, s11, s12, s13, item_ct1);
+                        });
+                }
+            } else {
+                /*
+                DPCT1049:16: The work-group size passed to the SYCL kernel may
+                exceed the limit. To get the device limit, query
+                info::device::max_work_group_size. Adjust the work-group size if
+                needed.
+                */
+                dpct::has_capability_or_fail(stream->get_device(),
+                                             {sycl::aspect::fp16});
+
+                stream->parallel_for(
+                    sycl::nd_range<3>(block_nums * block_dims, block_dims),
+                    [=](sycl::nd_item<3> item_ct1) {
+                        k_bin_bcast<bin_op>(src0_dd, src1_dd, dst_dd, ne0, ne1,
+                                            ne2, ne3, ne10, ne11, ne12, ne13,
+                                            s1, s2, s3, s01, s02, s03, s11, s12, s13,
+                                            item_ct1);
+                    });
+            }
+        }
+    }
+};
+
+template <class op>
+inline void ggml_sycl_op_bin_bcast(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1,
+                                   ggml_tensor * dst) {
+    dpct::queue_ptr main_stream = ctx.stream();
+    GGML_TENSOR_BINARY_OP_LOCALS
+
+    if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
+        op()((const float *) src0->data, (const float *) src1->data, (float *) dst->data, ne00, ne01, ne02, ne03, ne10,
+             ne11, ne12, ne13, ne0, ne1, ne2, ne3, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb0, nb1, nb2, nb3,
+             ggml_is_contiguous(src0), ggml_is_contiguous(src1), ggml_is_contiguous(dst), main_stream);
+    } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F16) {
+        op()((const sycl::half *) src0->data, (const sycl::half *) src1->data, (sycl::half *) dst->data, ne00, ne01,
+             ne02, ne03, ne10, ne11, ne12, ne13, ne0, ne1, ne2, ne3, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13,
+             nb0, nb1, nb2, nb3, ggml_is_contiguous(src0), ggml_is_contiguous(src1), ggml_is_contiguous(dst),
+             main_stream);
+    } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F16) {
+        op()((const sycl::half *) src0->data, (const float *) src1->data, (sycl::half *) dst->data, ne00, ne01, ne02,
+             ne03, ne10, ne11, ne12, ne13, ne0, ne1, ne2, ne3, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb0, nb1,
+             nb2, nb3, ggml_is_contiguous(src0), ggml_is_contiguous(src1), ggml_is_contiguous(dst), main_stream);
+    } else if (src0->type == GGML_TYPE_I32 && src1->type == GGML_TYPE_I32 && dst->type == GGML_TYPE_I32) {
+        op()((const int32_t *) src0->data, (const int32_t *) src1->data, (int32_t *) dst->data, ne00, ne01, ne02, ne03,
+             ne10, ne11, ne12, ne13, ne0, ne1, ne2, ne3, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb0, nb1, nb2,
+             nb3, ggml_is_contiguous(src0), ggml_is_contiguous(src1), ggml_is_contiguous(dst), main_stream);
+    } else if (src0->type == GGML_TYPE_I16 && src1->type == GGML_TYPE_I16 && dst->type == GGML_TYPE_I16) {
+        op()((const int16_t *) src0->data, (const int16_t *) src1->data, (int16_t *) dst->data, ne00, ne01, ne02, ne03,
+             ne10, ne11, ne12, ne13, ne0, ne1, ne2, ne3, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb0, nb1, nb2,
+             nb3, ggml_is_contiguous(src0), ggml_is_contiguous(src1), ggml_is_contiguous(dst), main_stream);
+    } else {
+        fprintf(stderr, "%s: unsupported types: dst: %s, src0: %s, src1: %s\n", __func__, ggml_type_name(dst->type),
+                ggml_type_name(src0->type), ggml_type_name(src1->type));
+        GGML_ABORT("fatal error");
+    }
+}
+
+inline void ggml_sycl_op_add(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
+
+    ggml_sycl_op_bin_bcast<bin_bcast_sycl<op_add>>(ctx, dst->src[0], dst->src[1], dst);
+}
+
+inline void ggml_sycl_op_sub(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
+
+    ggml_sycl_op_bin_bcast<bin_bcast_sycl<op_sub>>(ctx, dst->src[0], dst->src[1], dst);
+}
+
+inline void ggml_sycl_op_mul(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
+
+    ggml_sycl_op_bin_bcast<bin_bcast_sycl<op_mul>>(ctx, dst->src[0], dst->src[1], dst);
+}
+
+inline void ggml_sycl_op_div(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
+
+    ggml_sycl_op_bin_bcast<bin_bcast_sycl<op_div>>(ctx, dst->src[0], dst->src[1], dst);
+}
+
+inline void ggml_sycl_op_repeat(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
+    ggml_sycl_op_bin_bcast<bin_bcast_sycl<op_repeat>>(ctx, dst, dst->src[0], dst);
+}
+
+
+void ggml_sycl_add(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    GGML_SYCL_DEBUG("call %s\n", __func__);
+    ggml_sycl_op_add(ctx, dst);
+    GGML_SYCL_DEBUG("call %s done\n", __func__);
+}
+
+void ggml_sycl_sub(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    GGML_SYCL_DEBUG("call %s\n", __func__);
+    ggml_sycl_op_sub(ctx, dst);
+    GGML_SYCL_DEBUG("call %s done\n", __func__);
+}
+
+void ggml_sycl_mul(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    GGML_SYCL_DEBUG("call %s\n", __func__);
+    ggml_sycl_op_mul(ctx, dst);
+    GGML_SYCL_DEBUG("call %s done\n", __func__);
+}
+
+void ggml_sycl_div(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    GGML_SYCL_DEBUG("call %s\n", __func__);
+    ggml_sycl_op_div(ctx, dst);
+    GGML_SYCL_DEBUG("call %s done\n", __func__);
+}
+
+void ggml_sycl_repeat(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    GGML_SYCL_DEBUG("call %s\n", __func__);
+    ggml_sycl_op_repeat(ctx, dst);
+    GGML_SYCL_DEBUG("call %s done\n", __func__);
+}
+
diff --git a/ggml/src/ggml-sycl/binbcast.hpp b/ggml/src/ggml-sycl/binbcast.hpp
new file mode 100644
index 000000000..9cce0f053
--- /dev/null
+++ b/ggml/src/ggml-sycl/binbcast.hpp
@@ -0,0 +1,39 @@
+#ifndef GGML_SYCL_BINBCAST_HPP
+#define GGML_SYCL_BINBCAST_HPP
+#include "common.hpp"
+
+
+static __dpct_inline__ float op_repeat(const float a, const float b) {
+    return b;
+    GGML_UNUSED(a);
+}
+
+static __dpct_inline__ float op_add(const float a, const float b) {
+    return a + b;
+}
+
+static __dpct_inline__ float op_sub(const float a, const float b) {
+    return a - b;
+}
+
+static __dpct_inline__ float op_mul(const float a, const float b) {
+    return a * b;
+}
+
+static __dpct_inline__ float op_div(const float a, const float b) {
+    return a / b;
+}
+
+void ggml_sycl_add(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+
+void ggml_sycl_sub(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+
+void ggml_sycl_mul(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+
+void ggml_sycl_div(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+
+void ggml_sycl_repeat(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+
+
+#endif //GGML_SYCL_BINBCAST_HPP
+
diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp
index 3e1ceeaa4..96becabc8 100644
--- a/ggml/src/ggml-sycl/common.hpp
+++ b/ggml/src/ggml-sycl/common.hpp
@@ -494,286 +494,5 @@ static __dpct_inline__ Tp* get_pointer(sycl::local_accessor<Tp, dim> acc) {
 
 int64_t downsample_sycl_global_range(int64_t accumulate_block_num, int64_t block_size);
 
-template<float (*bin_op)(const float, const float), typename src0_t, typename src1_t, typename dst_t>
-static void k_bin_bcast(const src0_t * src0, const src1_t * src1, dst_t * dst,
-        int ne0, int ne1, int ne2, int ne3,
-        int ne10, int ne11, int ne12, int ne13,
-        /*int s0, */ int s1,  int s2,  int s3,
-        /*int s00,*/ int s01, int s02, int s03,
-        /*int s10,*/ int s11, int s12, int s13,
-        const sycl::nd_item<3> &item_ct1) {
-    const int i0s = item_ct1.get_local_range(2) * item_ct1.get_group(2) +
-                    item_ct1.get_local_id(2);
-    const int i1 = (item_ct1.get_local_range(1) * item_ct1.get_group(1) +
-                    item_ct1.get_local_id(1));
-    const int i2 = (item_ct1.get_local_range(0) * item_ct1.get_group(0) +
-                    item_ct1.get_local_id(0)) /
-                   ne3;
-    const int i3 = (item_ct1.get_local_range(0) * item_ct1.get_group(0) +
-                    item_ct1.get_local_id(0)) %
-                   ne3;
-
-    if (i0s >= ne0 || i1 >= ne1 || i2 >= ne2 || i3 >= ne3) {
-        return;
-    }
-
-    const int i11 = i1 % ne11;
-    const int i12 = i2 % ne12;
-    const int i13 = i3 % ne13;
-
-    const size_t i_src0 =  i3*s03 +  i2*s02 +  i1*s01;
-    const size_t i_src1 = i13*s13 + i12*s12 + i11*s11;
-    const size_t i_dst  =  i3*s3  +  i2*s2  +  i1*s1;
-
-    const src0_t * src0_row = src0 + i_src0;
-    const src1_t * src1_row = src1 + i_src1;
-    dst_t * dst_row = dst + i_dst;
-
-    for (int i0 = i0s; i0 < ne0;
-         i0 += item_ct1.get_local_range(2) * item_ct1.get_group_range(2)) {
-        const int i10 = i0 % ne10;
-        dst_row[i0] = (dst_t)bin_op(src0 ? (float)src0_row[i0] : 0.0f, (float)src1_row[i10]);
-    }
-}
-
-template<float (*bin_op)(const float, const float), typename src0_t, typename src1_t, typename dst_t>
-static void k_bin_bcast_unravel(const src0_t * src0, const src1_t * src1, dst_t * dst,
-        int ne0, int ne1, int ne2, int ne3,
-        int ne10, int ne11, int ne12, int ne13,
-        /*int s0, */ int s1,  int s2,  int s3,
-        /*int s00,*/ int s01, int s02, int s03,
-        /*int s10,*/ int s11, int s12, int s13,
-        const sycl::nd_item<3> &item_ct1) {
-
-    const int i = item_ct1.get_local_range(2) * item_ct1.get_group(2) +
-                  item_ct1.get_local_id(2);
-
-    const int i3 = i/(ne2*ne1*ne0);
-    const int i2 = (i/(ne1*ne0)) % ne2;
-    const int i1 = (i/ne0) % ne1;
-    const int i0 = i % ne0;
-
-    if (i0 >= ne0 || i1 >= ne1 || i2 >= ne2 || i3 >= ne3) {
-        return;
-    }
-
-    const int i11 = i1 % ne11;
-    const int i12 = i2 % ne12;
-    const int i13 = i3 % ne13;
-
-    const size_t i_src0 =  i3*s03 +  i2*s02 +  i1*s01;
-    const size_t i_src1 = i13*s13 + i12*s12 + i11*s11;
-    const size_t i_dst  =  i3*s3  +  i2*s2  +  i1*s1;
-
-    const src0_t * src0_row = src0 + i_src0;
-    const src1_t * src1_row = src1 + i_src1;
-    dst_t * dst_row = dst + i_dst;
-
-    const int i10 = i0 % ne10;
-    dst_row[i0] = (dst_t)bin_op(src0 ? (float)src0_row[i0] : 0.0f, (float)src1_row[i10]);
-}
-
-
-template<float (*bin_op)(const float, const float)>
-struct bin_bcast_sycl {
-    template <typename src0_t, typename src1_t, typename dst_t>
-    void operator()(ggml_backend_sycl_context & ctx,
-                    const struct ggml_tensor *src0,
-                    const struct ggml_tensor *src1, struct ggml_tensor *dst,
-                    const src0_t *src0_dd, const src1_t *src1_dd, dst_t *dst_dd,
-                    queue_ptr stream) {
-
-        GGML_TENSOR_BINARY_OP_LOCALS
-
-        int nr0 = ne10/ne0;
-        int nr1 = ne11/ne1;
-        int nr2 = ne12/ne2;
-        int nr3 = ne13/ne3;
-
-        int nr[4] = { nr0, nr1, nr2, nr3 };
-
-        // collapse dimensions until first broadcast dimension
-        int64_t cne[] = {ne0, ne1, ne2, ne3};
-        int64_t cne0[] = {ne00, ne01, ne02, ne03};
-        int64_t cne1[] = {ne10, ne11, ne12, ne13};
-        size_t cnb[] = {nb0, nb1, nb2, nb3};
-        size_t cnb0[] = {nb00, nb01, nb02, nb03};
-        size_t cnb1[] = {nb10, nb11, nb12, nb13};
-        auto collapse = [](int64_t cne[]) {
-            cne[0] *= cne[1];
-            cne[1] = cne[2];
-            cne[2] = cne[3];
-            cne[3] = 1;
-        };
-
-        auto collapse_nb = [](size_t cnb[], int64_t cne[]) {
-            cnb[1] *= cne[1];
-            cnb[2] *= cne[2];
-            cnb[3] *= cne[3];
-        };
-
-        if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && ggml_is_contiguous(dst)) {
-            for (int i = 0; i < 4; i++) {
-                if (nr[i] != 1) {
-                    break;
-                }
-                if (i > 0) {
-                    collapse_nb(cnb, cne);
-                    collapse_nb(cnb0, cne0);
-                    collapse_nb(cnb1, cne1);
-                    collapse(cne);
-                    collapse(cne0);
-                    collapse(cne1);
-                }
-            }
-        }
-        {
-            int64_t ne0 = cne[0];
-            int64_t ne1 = cne[1];
-            int64_t ne2 = cne[2];
-            int64_t ne3 = cne[3];
-
-            int64_t ne10 = cne1[0];
-            int64_t ne11 = cne1[1];
-            int64_t ne12 = cne1[2];
-            int64_t ne13 = cne1[3];
-
-            size_t nb0 = cnb[0];
-            size_t nb1 = cnb[1];
-            size_t nb2 = cnb[2];
-            size_t nb3 = cnb[3];
-
-            size_t nb00 = cnb0[0];
-            size_t nb01 = cnb0[1];
-            size_t nb02 = cnb0[2];
-            size_t nb03 = cnb0[3];
-
-            size_t nb10 = cnb1[0];
-            size_t nb11 = cnb1[1];
-            size_t nb12 = cnb1[2];
-            size_t nb13 = cnb1[3];
-
-            size_t s0 = nb0 / sizeof(dst_t);
-            size_t s1 = nb1 / sizeof(dst_t);
-            size_t s2 = nb2 / sizeof(dst_t);
-            size_t s3 = nb3 / sizeof(dst_t);
-
-            size_t s10 = nb10 / sizeof(src1_t);
-            size_t s11 = nb11 / sizeof(src1_t);
-            size_t s12 = nb12 / sizeof(src1_t);
-            size_t s13 = nb13 / sizeof(src1_t);
-
-            size_t s00 = nb00 / sizeof(src0_t);
-            size_t s01 = nb01 / sizeof(src0_t);
-            size_t s02 = nb02 / sizeof(src0_t);
-            size_t s03 = nb03 / sizeof(src0_t);
-
-            GGML_UNUSED(s00);
-
-            GGML_ASSERT(nb0 % sizeof(dst_t) == 0);
-            GGML_ASSERT(nb1 % sizeof(dst_t) == 0);
-            GGML_ASSERT(nb2 % sizeof(dst_t) == 0);
-            GGML_ASSERT(nb3 % sizeof(dst_t) == 0);
-
-            GGML_ASSERT(nb00 % sizeof(src0_t) == 0);
-            GGML_ASSERT(nb01 % sizeof(src0_t) == 0);
-            GGML_ASSERT(nb02 % sizeof(src0_t) == 0);
-            GGML_ASSERT(nb03 % sizeof(src0_t) == 0);
-
-            GGML_ASSERT(nb10 % sizeof(src1_t) == 0);
-            GGML_ASSERT(nb11 % sizeof(src1_t) == 0);
-            GGML_ASSERT(nb12 % sizeof(src1_t) == 0);
-            GGML_ASSERT(nb13 % sizeof(src1_t) == 0);
-
-            GGML_ASSERT(s0 == 1);
-            GGML_ASSERT(s10 == 1);
-
-            const int block_size = 128;
-
-            int64_t hne0 = std::max(ne0/2LL, 1LL);
-
-            sycl::range<3> block_dims(1, 1, 1);
-            block_dims[2] = std::min<unsigned int>(hne0, block_size);
-            block_dims[1] = std::min<unsigned int>(
-                ne1, block_size / (unsigned int)block_dims[2]);
-            block_dims[0] = std::min(
-                std::min<unsigned int>(
-                    ne2 * ne3, block_size / (unsigned int)block_dims[2] /
-                                   (unsigned int)block_dims[1]),
-                64U);
-
-            sycl::range<3> block_nums(
-                (ne2 * ne3 + block_dims[0] - 1) / block_dims[0],
-                (ne1 + block_dims[1] - 1) / block_dims[1],
-                (hne0 + block_dims[2] - 1) / block_dims[2]);
-
-            if (block_nums[0] > 65535) {
-                // this is the maximum number of blocks in z direction, fallback to 1D grid kernel
-                int block_num = (ne0*ne1*ne2*ne3 + block_size - 1) / block_size;
-                {
-                    dpct::has_capability_or_fail(stream->get_device(),
-                                                 {sycl::aspect::fp16});
-
-                    stream->parallel_for(
-                        sycl::nd_range<3>(sycl::range<3>(1, 1, block_num) *
-                                              sycl::range<3>(1, 1, block_size),
-                                          sycl::range<3>(1, 1, block_size)),
-                        [=](sycl::nd_item<3> item_ct1) {
-                            k_bin_bcast_unravel<bin_op>(
-                                src0_dd, src1_dd, dst_dd, ne0, ne1, ne2, ne3,
-                                ne10, ne11, ne12, ne13, s1, s2, s3, s01, s02,
-                                s03, s11, s12, s13, item_ct1);
-                        });
-                }
-            } else {
-                /*
-                DPCT1049:16: The work-group size passed to the SYCL kernel may
-                exceed the limit. To get the device limit, query
-                info::device::max_work_group_size. Adjust the work-group size if
-                needed.
-                */
-                dpct::has_capability_or_fail(stream->get_device(),
-                                             {sycl::aspect::fp16});
-
-                stream->parallel_for(
-                    sycl::nd_range<3>(block_nums * block_dims, block_dims),
-                    [=](sycl::nd_item<3> item_ct1) {
-                        k_bin_bcast<bin_op>(src0_dd, src1_dd, dst_dd, ne0, ne1,
-                                            ne2, ne3, ne10, ne11, ne12, ne13,
-                                            s1, s2, s3, s01, s02, s03, s11, s12, s13,
-                                            item_ct1);
-                    });
-            }
-        }
-        GGML_UNUSED(ctx);
-    }
-};
-
-template <class op>
-inline void ggml_sycl_op_bin_bcast(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
-                                   const ggml_tensor *src1, ggml_tensor *dst) {
-    dpct::queue_ptr main_stream = ctx.stream();
-
-    if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
-        op()(ctx, src0, src1, dst, (const float *)src0->data, (const float *)src1->data, (float *)dst->data, main_stream);
-    } else if (src0->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F16) {
-        op()(ctx, src0, src1, dst, (const sycl::half *)src0->data, (const float *)src1->data,
-             (sycl::half *)dst->data, main_stream);
-    } else if (src0->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F32) {
-        op()(ctx, src0, src1, dst, (const sycl::half *)src0->data, (const float *)src1->data, (float *)dst->data,
-             main_stream);
-    } else if (src0->type == GGML_TYPE_I32 && dst->type == GGML_TYPE_I32) {
-        op()(ctx, src0, src1, dst, (const int32_t *)src0->data, (const int32_t *)src1->data, (int32_t *)dst->data,
-             main_stream);
-    } else if (src0->type == GGML_TYPE_I16 && dst->type == GGML_TYPE_I16) {
-        op()(ctx, src0, src1, dst, (const int16_t *)src0->data, (const int16_t *)src1->data, (int16_t *)dst->data,
-             main_stream);
-    } else {
-        fprintf(stderr, "%s: unsupported types: dst: %s, src0: %s, src1: %s\n", __func__,
-            ggml_type_name(dst->type), ggml_type_name(src0->type), ggml_type_name(src1->type));
-        GGML_ABORT("fatal error");
-    }
-}
-
 bool gpu_has_xmx(sycl::device &dev);
 #endif // GGML_SYCL_COMMON_HPP
diff --git a/ggml/src/ggml-sycl/element_wise.cpp b/ggml/src/ggml-sycl/element_wise.cpp
index b36aa7a9d..fc25d98dd 100644
--- a/ggml/src/ggml-sycl/element_wise.cpp
+++ b/ggml/src/ggml-sycl/element_wise.cpp
@@ -1261,27 +1261,6 @@ inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor *dst)
 }
 
 
-inline void ggml_sycl_op_add(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
-
-    ggml_sycl_op_bin_bcast<bin_bcast_sycl<op_add>>(ctx, dst->src[0], dst->src[1], dst);
-}
-
-inline void ggml_sycl_op_sub(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
-
-    ggml_sycl_op_bin_bcast<bin_bcast_sycl<op_sub>>(ctx, dst->src[0], dst->src[1], dst);
-}
-
-inline void ggml_sycl_op_mul(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
-
-    ggml_sycl_op_bin_bcast<bin_bcast_sycl<op_mul>>(ctx, dst->src[0], dst->src[1], dst);
-}
-
-inline void ggml_sycl_op_div(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
-
-    ggml_sycl_op_bin_bcast<bin_bcast_sycl<op_div>>(ctx, dst->src[0], dst->src[1], dst);
-}
-
-
 void ggml_sycl_sqrt(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
     GGML_SYCL_DEBUG("call %s: DST Tensor type: %s\n", __func__, ggml_type_name(dst->type));
     ggml_sycl_op_sqrt(ctx, dst);
@@ -1409,29 +1388,3 @@ void ggml_sycl_clamp(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
     GGML_SYCL_DEBUG("call %s done\n", __func__);
 }
 
-
-
-
-void ggml_sycl_add(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
-    GGML_SYCL_DEBUG("call %s\n", __func__);
-    ggml_sycl_op_add(ctx, dst);
-    GGML_SYCL_DEBUG("call %s done\n", __func__);
-}
-
-void ggml_sycl_sub(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
-    GGML_SYCL_DEBUG("call %s\n", __func__);
-    ggml_sycl_op_sub(ctx, dst);
-    GGML_SYCL_DEBUG("call %s done\n", __func__);
-}
-
-void ggml_sycl_mul(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
-    GGML_SYCL_DEBUG("call %s\n", __func__);
-    ggml_sycl_op_mul(ctx, dst);
-    GGML_SYCL_DEBUG("call %s done\n", __func__);
-}
-
-void ggml_sycl_div(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
-    GGML_SYCL_DEBUG("call %s\n", __func__);
-    ggml_sycl_op_div(ctx, dst);
-    GGML_SYCL_DEBUG("call %s done\n", __func__);
-}
diff --git a/ggml/src/ggml-sycl/element_wise.hpp b/ggml/src/ggml-sycl/element_wise.hpp
index 76d91ba10..e623cb56f 100644
--- a/ggml/src/ggml-sycl/element_wise.hpp
+++ b/ggml/src/ggml-sycl/element_wise.hpp
@@ -10,27 +10,6 @@ T neg_infinity() {
     return -std::numeric_limits<T>::infinity();
 }
 
-static __dpct_inline__ float op_repeat(const float a, const float b) {
-    return b;
-    GGML_UNUSED(a);
-}
-
-static __dpct_inline__ float op_add(const float a, const float b) {
-    return a + b;
-}
-
-static __dpct_inline__ float op_sub(const float a, const float b) {
-    return a - b;
-}
-
-static __dpct_inline__ float op_mul(const float a, const float b) {
-    return a * b;
-}
-
-static __dpct_inline__ float op_div(const float a, const float b) {
-    return a / b;
-}
-
 template<typename T>
 struct typed_data {
     const T * src;
@@ -87,14 +66,5 @@ void ggml_sycl_pad(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
 
 void ggml_sycl_clamp(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
 
-// ---------
-
-void ggml_sycl_add(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
-
-void ggml_sycl_sub(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
-
-void ggml_sycl_mul(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
-
-void ggml_sycl_div(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
-
 #endif // GGML_SYCL_ELEMENTWISE_HPP
+
diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index 4d2fda0bf..1de34c962 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -1967,11 +1967,6 @@ catch (sycl::exception const &exc) {
   std::exit(1);
 }
 
-static void ggml_sycl_op_repeat(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
-    ggml_sycl_op_bin_bcast<bin_bcast_sycl<op_repeat>>(ctx, dst, dst->src[0], dst);
-}
-
-
 inline void ggml_sycl_op_mul_mat_sycl(
     ggml_backend_sycl_context & ctx,
     const ggml_tensor *src0, const ggml_tensor *src1, ggml_tensor *dst,
@@ -2600,12 +2595,6 @@ catch (sycl::exception const &exc) {
 }
 
 
-static void ggml_sycl_repeat(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
-    GGML_SYCL_DEBUG("call %s\n", __func__);
-    ggml_sycl_op_repeat(ctx, dst);
-    GGML_SYCL_DEBUG("call %s done\n", __func__);
-}
-
 static void ggml_sycl_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
     GGML_SYCL_DEBUG("call %s\n", __func__);
     ggml_sycl_op_get_rows(ctx, dst);
@@ -3972,7 +3961,6 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g
         case GGML_OP_ARGMAX:
         case GGML_OP_NONE:
         case GGML_OP_RESHAPE:
-        case GGML_OP_REPEAT:
         case GGML_OP_VIEW:
         case GGML_OP_PERMUTE:
         case GGML_OP_TRANSPOSE:
@@ -3982,7 +3970,8 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g
         case GGML_OP_SUB:
         case GGML_OP_MUL:
         case GGML_OP_DIV:
-            return (op->src[0]->type == GGML_TYPE_F32);
+        case GGML_OP_REPEAT:
+            return true;
         case GGML_OP_SQR:
         case GGML_OP_SQRT:
         case GGML_OP_SIN:

From 35370ba94593c3abc9550328377d461fc4f822b7 Mon Sep 17 00:00:00 2001
From: Xuan-Son Nguyen <son@huggingface.co>
Date: Fri, 18 Apr 2025 19:58:12 +0200
Subject: [PATCH 4/6] server : use std::move whenever possible (#12936)

* server : use std::move whenever possible

* use r-value ref

* Apply suggestions from code review

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>

* make task creation scoped

* restore std::move

* fix task_id not set correctly

* apply changes from suggestion

Co-authored-by: ggerganov <ggerganov@gmail.com>

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
---
 examples/server/server.cpp | 245 ++++++++++++++++++++-----------------
 1 file changed, 132 insertions(+), 113 deletions(-)

diff --git a/examples/server/server.cpp b/examples/server/server.cpp
index d87bda1a0..c580ec123 100644
--- a/examples/server/server.cpp
+++ b/examples/server/server.cpp
@@ -1552,29 +1552,30 @@ struct server_queue {
     std::condition_variable condition_tasks;
 
     // callback functions
-    std::function<void(server_task)> callback_new_task;
-    std::function<void(void)>        callback_update_slots;
+    std::function<void(server_task &&)> callback_new_task;
+    std::function<void(void)>           callback_update_slots;
 
     // Add a new task to the end of the queue
-    int post(server_task task, bool front = false) {
+    int post(server_task && task, bool front = false) {
         std::unique_lock<std::mutex> lock(mutex_tasks);
         GGML_ASSERT(task.id != -1);
         // if this is cancel task make sure to clean up pending tasks
         if (task.type == SERVER_TASK_TYPE_CANCEL) {
             cleanup_pending_task(task.id_target);
         }
-        QUE_DBG("new task, id = %d, front = %d\n", task.id, front);
+        const int task_id = task.id;
+        QUE_DBG("new task, id = %d, front = %d\n", task_id, front);
         if (front) {
             queue_tasks.push_front(std::move(task));
         } else {
             queue_tasks.push_back(std::move(task));
         }
         condition_tasks.notify_one();
-        return task.id;
+        return task_id;
     }
 
     // multi-task version of post()
-    int post(std::vector<server_task> & tasks, bool front = false) {
+    int post(std::vector<server_task> && tasks, bool front = false) {
         std::unique_lock<std::mutex> lock(mutex_tasks);
         for (auto & task : tasks) {
             if (task.id == -1) {
@@ -1596,7 +1597,7 @@ struct server_queue {
     }
 
     // Add a new task, but defer until one slot is available
-    void defer(server_task task) {
+    void defer(server_task && task) {
         std::unique_lock<std::mutex> lock(mutex_tasks);
         QUE_DBG("defer task, id = %d\n", task.id);
         queue_tasks_deferred.push_back(std::move(task));
@@ -1611,7 +1612,7 @@ struct server_queue {
     }
 
     // Register function to process a new task
-    void on_new_task(std::function<void(server_task)> callback) {
+    void on_new_task(std::function<void(server_task &&)> callback) {
         callback_new_task = std::move(callback);
     }
 
@@ -1660,7 +1661,7 @@ struct server_queue {
                     lock.unlock();
                     break;
                 }
-                server_task task = queue_tasks.front();
+                server_task task = std::move(queue_tasks.front());
                 queue_tasks.pop_front();
                 lock.unlock();
 
@@ -2004,7 +2005,7 @@ struct server_context {
 
             slot.reset();
 
-            slots.push_back(slot);
+            slots.push_back(std::move(slot));
         }
 
         default_generation_settings_for_props = slots[0].to_json();
@@ -2105,7 +2106,7 @@ struct server_context {
         return true;
     }
 
-    bool launch_slot_with_task(server_slot & slot, const server_task & task) {
+    bool launch_slot_with_task(server_slot & slot, server_task && task) {
         slot.reset();
         slot.id_task       = task.id;
         slot.index         = task.index;
@@ -2113,10 +2114,10 @@ struct server_context {
         slot.params        = std::move(task.params);
         slot.prompt_tokens = std::move(task.prompt_tokens);
 
-        if (!are_lora_equal(task.params.lora, slot.lora)) {
+        if (!are_lora_equal(slot.params.lora, slot.lora)) {
             // if lora is changed, we cannot reuse cached tokens
             slot.cache_tokens.clear();
-            slot.lora = task.params.lora;
+            slot.lora = slot.params.lora;
         }
 
         bool can_detokenize = can_be_detokenized(ctx, slot.prompt_tokens);
@@ -2547,10 +2548,10 @@ struct server_context {
             server_task task(SERVER_TASK_TYPE_CANCEL);
             task.id_target = id_task;
             queue_results.remove_waiting_task_id(id_task);
-            cancel_tasks.push_back(task);
+            cancel_tasks.push_back(std::move(task));
         }
         // push to beginning of the queue, so it has highest priority
-        queue_tasks.post(cancel_tasks, true);
+        queue_tasks.post(std::move(cancel_tasks), true);
     }
 
     // receive the results from task(s)
@@ -2637,7 +2638,7 @@ struct server_context {
     // Functions to process the task
     //
 
-    void process_single_task(server_task task) {
+    void process_single_task(server_task && task) {
         switch (task.type) {
             case SERVER_TASK_TYPE_COMPLETION:
             case SERVER_TASK_TYPE_INFILL:
@@ -2651,17 +2652,17 @@ struct server_context {
                     if (slot == nullptr) {
                         // if no slot is available, we defer this task for processing later
                         SRV_DBG("no slot is available, defer task, id_task = %d\n", task.id);
-                        queue_tasks.defer(task);
+                        queue_tasks.defer(std::move(task));
                         break;
                     }
                     if (slot->is_processing()) {
                         // if requested slot is unavailable, we defer this task for processing later
                         SRV_DBG("requested slot is unavailable, defer task, id_task = %d\n", task.id);
-                        queue_tasks.defer(task);
+                        queue_tasks.defer(std::move(task));
                         break;
                     }
 
-                    if (!launch_slot_with_task(*slot, task)) {
+                    if (!launch_slot_with_task(*slot, std::move(task))) {
                         SRV_ERR("failed to launch slot with task, id_task = %d\n", task.id);
                         break;
                     }
@@ -2740,7 +2741,7 @@ struct server_context {
                     if (slot->is_processing()) {
                         // if requested slot is unavailable, we defer this task for processing later
                         SRV_DBG("requested slot is unavailable, defer task, id_task = %d\n", task.id);
-                        queue_tasks.defer(task);
+                        queue_tasks.defer(std::move(task));
                         break;
                     }
 
@@ -2776,7 +2777,7 @@ struct server_context {
                     if (slot->is_processing()) {
                         // if requested slot is unavailable, we defer this task for processing later
                         SRV_DBG("requested slot is unavailable, defer task, id_task = %d\n", task.id);
-                        queue_tasks.defer(task);
+                        queue_tasks.defer(std::move(task));
                         break;
                     }
 
@@ -2819,7 +2820,7 @@ struct server_context {
                     if (slot->is_processing()) {
                         // if requested slot is unavailable, we defer this task for processing later
                         SRV_DBG("requested slot is unavailable, defer task, id_task = %d\n", task.id);
-                        queue_tasks.defer(task);
+                        queue_tasks.defer(std::move(task));
                         break;
                     }
 
@@ -2871,7 +2872,7 @@ struct server_context {
 
             server_task task(SERVER_TASK_TYPE_NEXT_RESPONSE);
             task.id = queue_tasks.get_new_id();
-            queue_tasks.post(task);
+            queue_tasks.post(std::move(task));
         }
 
         // apply context-shift if needed
@@ -3633,14 +3634,17 @@ int main(int argc, char ** argv) {
         }
 
         // request slots data using task queue
-        server_task task(SERVER_TASK_TYPE_METRICS);
-        task.id = ctx_server.queue_tasks.get_new_id();
-        ctx_server.queue_results.add_waiting_task_id(task.id);
-        ctx_server.queue_tasks.post(task, true); // high-priority task
+        int task_id = ctx_server.queue_tasks.get_new_id();
+        {
+            server_task task(SERVER_TASK_TYPE_METRICS);
+            task.id = task_id;
+            ctx_server.queue_results.add_waiting_task_id(task_id);
+            ctx_server.queue_tasks.post(std::move(task), true); // high-priority task
+        }
 
         // get the result
-        server_task_result_ptr result = ctx_server.queue_results.recv(task.id);
-        ctx_server.queue_results.remove_waiting_task_id(task.id);
+        server_task_result_ptr result = ctx_server.queue_results.recv(task_id);
+        ctx_server.queue_results.remove_waiting_task_id(task_id);
 
         if (result->is_error()) {
             res_error(res, result->to_json());
@@ -3669,16 +3673,17 @@ int main(int argc, char ** argv) {
         }
 
         // request slots data using task queue
-        server_task task(SERVER_TASK_TYPE_METRICS);
-        task.id = ctx_server.queue_tasks.get_new_id();
-        task.metrics_reset_bucket = true;
-
-        ctx_server.queue_results.add_waiting_task_id(task.id);
-        ctx_server.queue_tasks.post(task, true); // high-priority task
+        int task_id = ctx_server.queue_tasks.get_new_id();
+        {
+            server_task task(SERVER_TASK_TYPE_METRICS);
+            task.id = task_id;
+            ctx_server.queue_results.add_waiting_task_id(task_id);
+            ctx_server.queue_tasks.post(std::move(task), true); // high-priority task
+        }
 
         // get the result
-        server_task_result_ptr result = ctx_server.queue_results.recv(task.id);
-        ctx_server.queue_results.remove_waiting_task_id(task.id);
+        server_task_result_ptr result = ctx_server.queue_results.recv(task_id);
+        ctx_server.queue_results.remove_waiting_task_id(task_id);
 
         if (result->is_error()) {
             res_error(res, result->to_json());
@@ -3775,17 +3780,20 @@ int main(int argc, char ** argv) {
         }
         std::string filepath = params.slot_save_path + filename;
 
-        server_task task(SERVER_TASK_TYPE_SLOT_SAVE);
-        task.id = ctx_server.queue_tasks.get_new_id();
-        task.slot_action.slot_id  = id_slot;
-        task.slot_action.filename = filename;
-        task.slot_action.filepath = filepath;
+        int task_id = ctx_server.queue_tasks.get_new_id();
+        {
+            server_task task(SERVER_TASK_TYPE_SLOT_SAVE);
+            task.id = task_id;
+            task.slot_action.slot_id  = id_slot;
+            task.slot_action.filename = filename;
+            task.slot_action.filepath = filepath;
 
-        ctx_server.queue_results.add_waiting_task_id(task.id);
-        ctx_server.queue_tasks.post(task);
+            ctx_server.queue_results.add_waiting_task_id(task_id);
+            ctx_server.queue_tasks.post(std::move(task));
+        }
 
-        server_task_result_ptr result = ctx_server.queue_results.recv(task.id);
-        ctx_server.queue_results.remove_waiting_task_id(task.id);
+        server_task_result_ptr result = ctx_server.queue_results.recv(task_id);
+        ctx_server.queue_results.remove_waiting_task_id(task_id);
 
         if (result->is_error()) {
             res_error(res, result->to_json());
@@ -3804,17 +3812,20 @@ int main(int argc, char ** argv) {
         }
         std::string filepath = params.slot_save_path + filename;
 
-        server_task task(SERVER_TASK_TYPE_SLOT_RESTORE);
-        task.id = ctx_server.queue_tasks.get_new_id();
-        task.slot_action.slot_id  = id_slot;
-        task.slot_action.filename = filename;
-        task.slot_action.filepath = filepath;
+        int task_id = ctx_server.queue_tasks.get_new_id();
+        {
+            server_task task(SERVER_TASK_TYPE_SLOT_RESTORE);
+            task.id = task_id;
+            task.slot_action.slot_id  = id_slot;
+            task.slot_action.filename = filename;
+            task.slot_action.filepath = filepath;
 
-        ctx_server.queue_results.add_waiting_task_id(task.id);
-        ctx_server.queue_tasks.post(task);
+            ctx_server.queue_results.add_waiting_task_id(task_id);
+            ctx_server.queue_tasks.post(std::move(task));
+        }
 
-        server_task_result_ptr result = ctx_server.queue_results.recv(task.id);
-        ctx_server.queue_results.remove_waiting_task_id(task.id);
+        server_task_result_ptr result = ctx_server.queue_results.recv(task_id);
+        ctx_server.queue_results.remove_waiting_task_id(task_id);
 
         if (result->is_error()) {
             res_error(res, result->to_json());
@@ -3826,15 +3837,18 @@ int main(int argc, char ** argv) {
     };
 
     const auto handle_slots_erase = [&ctx_server, &res_error, &res_ok](const httplib::Request & /* req */, httplib::Response & res, int id_slot) {
-        server_task task(SERVER_TASK_TYPE_SLOT_ERASE);
-        task.id = ctx_server.queue_tasks.get_new_id();
-        task.slot_action.slot_id = id_slot;
+        int task_id = ctx_server.queue_tasks.get_new_id();
+        {
+            server_task task(SERVER_TASK_TYPE_SLOT_ERASE);
+            task.id = task_id;
+            task.slot_action.slot_id = id_slot;
 
-        ctx_server.queue_results.add_waiting_task_id(task.id);
-        ctx_server.queue_tasks.post(task);
+            ctx_server.queue_results.add_waiting_task_id(task_id);
+            ctx_server.queue_tasks.post(std::move(task));
+        }
 
-        server_task_result_ptr result = ctx_server.queue_results.recv(task.id);
-        ctx_server.queue_results.remove_waiting_task_id(task.id);
+        server_task_result_ptr result = ctx_server.queue_results.recv(task_id);
+        ctx_server.queue_results.remove_waiting_task_id(task_id);
 
         if (result->is_error()) {
             res_error(res, result->to_json());
@@ -3938,9 +3952,10 @@ int main(int argc, char ** argv) {
         }
 
         auto completion_id = gen_chatcmplid();
-        std::vector<server_task> tasks;
-
+        std::unordered_set<int> task_ids;
         try {
+            std::vector<server_task> tasks;
+
             const auto & prompt = data.at("prompt");
             // TODO: this log can become very long, put it behind a flag or think about a more compact format
             //SRV_DBG("Prompt: %s\n", prompt.is_string() ? prompt.get<std::string>().c_str() : prompt.dump(2).c_str());
@@ -3955,9 +3970,9 @@ int main(int argc, char ** argv) {
 
                 task.prompt_tokens    = std::move(tokenized_prompts[i]);
                 task.params           = server_task::params_from_json_cmpl(
-                                            ctx_server.ctx,
-                                            ctx_server.params_base,
-                                            data);
+                        ctx_server.ctx,
+                        ctx_server.params_base,
+                        data);
                 task.id_selected_slot = json_value(data, "id_slot", -1);
 
                 // OAI-compat
@@ -3965,18 +3980,18 @@ int main(int argc, char ** argv) {
                 task.params.oaicompat_cmpl_id         = completion_id;
                 // oaicompat_model is already populated by params_from_json_cmpl
 
-                tasks.push_back(task);
+                tasks.push_back(std::move(task));
             }
+
+            task_ids = server_task::get_list_id(tasks);
+            ctx_server.queue_results.add_waiting_tasks(tasks);
+            ctx_server.queue_tasks.post(std::move(tasks));
         } catch (const std::exception & e) {
             res_error(res, format_error_response(e.what(), ERROR_TYPE_INVALID_REQUEST));
             return;
         }
 
-        ctx_server.queue_results.add_waiting_tasks(tasks);
-        ctx_server.queue_tasks.post(tasks);
-
         bool stream = json_value(data, "stream", false);
-        const auto task_ids = server_task::get_list_id(tasks);
 
         if (!stream) {
             ctx_server.receive_multi_results(task_ids, [&](std::vector<server_task_result_ptr> & results) {
@@ -4268,6 +4283,7 @@ int main(int argc, char ** argv) {
         // create and queue the task
         json responses = json::array();
         bool error = false;
+        std::unordered_set<int> task_ids;
         {
             std::vector<server_task> tasks;
             for (size_t i = 0; i < tokenized_prompts.size(); i++) {
@@ -4280,28 +4296,27 @@ int main(int argc, char ** argv) {
                 // OAI-compat
                 task.params.oaicompat = oaicompat;
 
-                tasks.push_back(task);
+                tasks.push_back(std::move(task));
             }
 
+            task_ids = server_task::get_list_id(tasks);
             ctx_server.queue_results.add_waiting_tasks(tasks);
-            ctx_server.queue_tasks.post(tasks);
-
-            // get the result
-            std::unordered_set<int> task_ids = server_task::get_list_id(tasks);
-
-            ctx_server.receive_multi_results(task_ids, [&](std::vector<server_task_result_ptr> & results) {
-                for (auto & res : results) {
-                    GGML_ASSERT(dynamic_cast<server_task_result_embd*>(res.get()) != nullptr);
-                    responses.push_back(res->to_json());
-                }
-            }, [&](const json & error_data) {
-                res_error(res, error_data);
-                error = true;
-            }, req.is_connection_closed);
-
-            ctx_server.queue_results.remove_waiting_task_ids(task_ids);
+            ctx_server.queue_tasks.post(std::move(tasks));
         }
 
+        // get the result
+        ctx_server.receive_multi_results(task_ids, [&](std::vector<server_task_result_ptr> & results) {
+            for (auto & res : results) {
+                GGML_ASSERT(dynamic_cast<server_task_result_embd*>(res.get()) != nullptr);
+                responses.push_back(res->to_json());
+            }
+        }, [&](const json & error_data) {
+            res_error(res, error_data);
+            error = true;
+        }, req.is_connection_closed);
+
+        ctx_server.queue_results.remove_waiting_task_ids(task_ids);
+
         if (error) {
             return;
         }
@@ -4367,6 +4382,7 @@ int main(int argc, char ** argv) {
         // create and queue the task
         json responses = json::array();
         bool error = false;
+        std::unordered_set<int> task_ids;
         {
             std::vector<server_task> tasks;
             std::vector<llama_tokens> tokenized_docs = tokenize_input_prompts(ctx_server.vocab, documents, /* add_special */ false, true);
@@ -4376,26 +4392,24 @@ int main(int argc, char ** argv) {
                 task.id            = ctx_server.queue_tasks.get_new_id();
                 task.index         = i;
                 task.prompt_tokens = format_rerank(ctx_server.vocab, tokenized_query, tokenized_docs[i]);
-                tasks.push_back(task);
+                tasks.push_back(std::move(task));
             }
 
+            task_ids = server_task::get_list_id(tasks);
             ctx_server.queue_results.add_waiting_tasks(tasks);
-            ctx_server.queue_tasks.post(tasks);
-
-            // get the result
-            std::unordered_set<int> task_ids = server_task::get_list_id(tasks);
-
-            ctx_server.receive_multi_results(task_ids, [&](std::vector<server_task_result_ptr> & results) {
-                for (auto & res : results) {
-                    GGML_ASSERT(dynamic_cast<server_task_result_rerank*>(res.get()) != nullptr);
-                    responses.push_back(res->to_json());
-                }
-            }, [&](const json & error_data) {
-                res_error(res, error_data);
-                error = true;
-            }, req.is_connection_closed);
+            ctx_server.queue_tasks.post(std::move(tasks));
         }
 
+        ctx_server.receive_multi_results(task_ids, [&](std::vector<server_task_result_ptr> & results) {
+            for (auto & res : results) {
+                GGML_ASSERT(dynamic_cast<server_task_result_rerank*>(res.get()) != nullptr);
+                responses.push_back(res->to_json());
+            }
+        }, [&](const json & error_data) {
+            res_error(res, error_data);
+            error = true;
+        }, req.is_connection_closed);
+
         if (error) {
             return;
         }
@@ -4431,14 +4445,19 @@ int main(int argc, char ** argv) {
             res_error(res, format_error_response("Request body must be an array", ERROR_TYPE_INVALID_REQUEST));
             return;
         }
-        server_task task(SERVER_TASK_TYPE_SET_LORA);
-        task.id = ctx_server.queue_tasks.get_new_id();
-        task.set_lora = parse_lora_request(ctx_server.params_base.lora_adapters, body);
-        ctx_server.queue_results.add_waiting_task_id(task.id);
-        ctx_server.queue_tasks.post(task);
 
-        server_task_result_ptr result = ctx_server.queue_results.recv(task.id);
-        ctx_server.queue_results.remove_waiting_task_id(task.id);
+        int task_id = ctx_server.queue_tasks.get_new_id();
+        {
+            server_task task(SERVER_TASK_TYPE_SET_LORA);
+            task.id = task_id;
+            task.set_lora = parse_lora_request(ctx_server.params_base.lora_adapters, body);
+            ctx_server.queue_results.add_waiting_task_id(task_id);
+            ctx_server.queue_tasks.post(std::move(task));
+        }
+
+        // get the result
+        server_task_result_ptr result = ctx_server.queue_results.recv(task_id);
+        ctx_server.queue_results.remove_waiting_task_id(task_id);
 
         if (result->is_error()) {
             res_error(res, result->to_json());
@@ -4582,8 +4601,8 @@ int main(int argc, char ** argv) {
         common_chat_templates_source(ctx_server.chat_templates.get()),
         common_chat_format_example(ctx_server.chat_templates.get(), ctx_server.params_base.use_jinja).c_str());
 
-    ctx_server.queue_tasks.on_new_task([&ctx_server](const server_task & task) {
-        ctx_server.process_single_task(task);
+    ctx_server.queue_tasks.on_new_task([&ctx_server](server_task && task) {
+        ctx_server.process_single_task(std::move(task));
     });
 
     ctx_server.queue_tasks.on_update_slots([&ctx_server]() {

From aff9d107b066c9d464f7ab1324280acfdcebf569 Mon Sep 17 00:00:00 2001
From: Chris Thompson <christopherthompson81@gmail.com>
Date: Fri, 18 Apr 2025 12:30:41 -0600
Subject: [PATCH 5/6] gguf-py : GGUF Editor GUI - Python + Qt6 (#12930)

---
 gguf-py/README.md                             |    7 +
 gguf-py/gguf/scripts/__init__.py              |    1 +
 gguf-py/gguf/scripts/gguf_editor_gui.py       | 1610 +++++++++++++++++
 gguf-py/pyproject.toml                        |    7 +-
 requirements/requirements-all.txt             |    2 +
 requirements/requirements-gguf_editor_gui.txt |    3 +
 6 files changed, 1629 insertions(+), 1 deletion(-)
 create mode 100755 gguf-py/gguf/scripts/gguf_editor_gui.py
 create mode 100644 requirements/requirements-gguf_editor_gui.txt

diff --git a/gguf-py/README.md b/gguf-py/README.md
index dd4ab7bde..ca7e09c68 100644
--- a/gguf-py/README.md
+++ b/gguf-py/README.md
@@ -11,6 +11,11 @@ as an example for its usage.
 pip install gguf
 ```
 
+Optionally, you can install gguf with the extra 'gui' to enable the visual GGUF editor.
+```sh
+pip install gguf[gui]
+```
+
 ## API Examples/Simple Tools
 
 [examples/writer.py](https://github.com/ggml-org/llama.cpp/blob/master/gguf-py/examples/writer.py) — Generates `example.gguf` in the current directory to demonstrate generating a GGUF file. Note that this file cannot be used as a model.
@@ -25,6 +30,8 @@ pip install gguf
 
 [gguf/scripts/gguf_new_metadata.py](https://github.com/ggml-org/llama.cpp/blob/master/gguf-py/gguf/scripts/gguf_new_metadata.py) — Copies a GGUF file with added/modified/removed metadata values.
 
+[gguf/scripts/gguf_editor_gui.py](https://github.com/ggml-org/llama.cpp/blob/master/gguf-py/gguf/scripts/gguf_editor_gui.py) — Allows for viewing, editing, adding, or removing metadata values within a GGUF file as well as viewing its tensors with a Qt interface.
+
 ## Development
 Maintainers who participate in development of this package are advised to install it in editable mode:
 
diff --git a/gguf-py/gguf/scripts/__init__.py b/gguf-py/gguf/scripts/__init__.py
index e77f2e9c9..72cc73e70 100644
--- a/gguf-py/gguf/scripts/__init__.py
+++ b/gguf-py/gguf/scripts/__init__.py
@@ -4,3 +4,4 @@ from .gguf_convert_endian import main as gguf_convert_endian_entrypoint
 from .gguf_dump import main as gguf_dump_entrypoint
 from .gguf_set_metadata import main as gguf_set_metadata_entrypoint
 from .gguf_new_metadata import main as gguf_new_metadata_entrypoint
+from .gguf_editor_gui import main as gguf_editor_gui_entrypoint
diff --git a/gguf-py/gguf/scripts/gguf_editor_gui.py b/gguf-py/gguf/scripts/gguf_editor_gui.py
new file mode 100755
index 000000000..9dab6ca27
--- /dev/null
+++ b/gguf-py/gguf/scripts/gguf_editor_gui.py
@@ -0,0 +1,1610 @@
+#!/usr/bin/env python3
+from __future__ import annotations
+
+import logging
+import argparse
+import os
+import sys
+import numpy
+import enum
+from pathlib import Path
+from typing import Any, Optional, Tuple, Type
+import warnings
+
+import numpy as np
+from PySide6.QtWidgets import (
+    QApplication, QMainWindow, QWidget, QVBoxLayout, QHBoxLayout,
+    QPushButton, QLabel, QLineEdit, QFileDialog, QTableWidget,
+    QTableWidgetItem, QComboBox, QMessageBox, QTabWidget,
+    QTextEdit, QFormLayout,
+    QHeaderView, QDialog, QDialogButtonBox
+)
+from PySide6.QtCore import Qt
+
+# Necessary to load the local gguf package
+if "NO_LOCAL_GGUF" not in os.environ and (Path(__file__).parent.parent.parent.parent / 'gguf-py').exists():
+    sys.path.insert(0, str(Path(__file__).parent.parent.parent))
+
+import gguf
+from gguf import GGUFReader, GGUFWriter, GGUFValueType, ReaderField
+from gguf.constants import TokenType, RopeScalingType, PoolingType, GGMLQuantizationType
+
+logger = logging.getLogger("gguf-editor-gui")
+
+# Map of key names to enum types for automatic enum interpretation
+KEY_TO_ENUM_TYPE = {
+    gguf.Keys.Tokenizer.TOKEN_TYPE: TokenType,
+    gguf.Keys.Rope.SCALING_TYPE: RopeScalingType,
+    gguf.Keys.LLM.POOLING_TYPE: PoolingType,
+    gguf.Keys.General.FILE_TYPE: GGMLQuantizationType,
+}
+
+# Define the tokenizer keys that should be edited together
+TOKENIZER_LINKED_KEYS = [
+    gguf.Keys.Tokenizer.LIST,
+    gguf.Keys.Tokenizer.TOKEN_TYPE,
+    gguf.Keys.Tokenizer.SCORES
+]
+
+
+class TokenizerEditorDialog(QDialog):
+    def __init__(self, tokens, token_types, scores, parent=None):
+        super().__init__(parent)
+        self.setWindowTitle("Edit Tokenizer Data")
+        self.resize(900, 600)
+
+        self.tokens = tokens.copy() if tokens else []
+        self.token_types = token_types.copy() if token_types else []
+        self.scores = scores.copy() if scores else []
+
+        # Ensure all arrays have the same length
+        max_len = max(len(self.tokens), len(self.token_types), len(self.scores))
+        if len(self.tokens) < max_len:
+            self.tokens.extend([""] * (max_len - len(self.tokens)))
+        if len(self.token_types) < max_len:
+            self.token_types.extend([0] * (max_len - len(self.token_types)))
+        if len(self.scores) < max_len:
+            self.scores.extend([0.0] * (max_len - len(self.scores)))
+
+        layout = QVBoxLayout(self)
+
+        # Add filter controls
+        filter_layout = QHBoxLayout()
+        filter_layout.addWidget(QLabel("Filter:"))
+        self.filter_edit = QLineEdit()
+        self.filter_edit.setPlaceholderText("Type to filter tokens...")
+        self.filter_edit.textChanged.connect(self.apply_filter)
+        filter_layout.addWidget(self.filter_edit)
+
+        # Add page controls
+        self.page_size = 100  # Show 100 items per page
+        self.current_page = 0
+        self.total_pages = max(1, (len(self.tokens) + self.page_size - 1) // self.page_size)
+
+        self.page_label = QLabel(f"Page 1 of {self.total_pages}")
+        filter_layout.addWidget(self.page_label)
+
+        prev_page = QPushButton("Previous")
+        prev_page.clicked.connect(self.previous_page)
+        filter_layout.addWidget(prev_page)
+
+        next_page = QPushButton("Next")
+        next_page.clicked.connect(self.next_page)
+        filter_layout.addWidget(next_page)
+
+        layout.addLayout(filter_layout)
+
+        # Tokenizer data table
+        self.tokens_table = QTableWidget()
+        self.tokens_table.setColumnCount(4)
+        self.tokens_table.setHorizontalHeaderLabels(["Index", "Token", "Type", "Score"])
+        self.tokens_table.horizontalHeader().setSectionResizeMode(0, QHeaderView.ResizeMode.ResizeToContents)
+        self.tokens_table.horizontalHeader().setSectionResizeMode(1, QHeaderView.ResizeMode.Stretch)
+        self.tokens_table.horizontalHeader().setSectionResizeMode(2, QHeaderView.ResizeMode.ResizeToContents)
+        self.tokens_table.horizontalHeader().setSectionResizeMode(3, QHeaderView.ResizeMode.ResizeToContents)
+
+        layout.addWidget(self.tokens_table)
+
+        # Controls
+        controls_layout = QHBoxLayout()
+
+        add_button = QPushButton("Add Token")
+        add_button.clicked.connect(self.add_token)
+        controls_layout.addWidget(add_button)
+
+        remove_button = QPushButton("Remove Selected")
+        remove_button.clicked.connect(self.remove_selected)
+        controls_layout.addWidget(remove_button)
+
+        controls_layout.addStretch()
+
+        layout.addLayout(controls_layout)
+
+        # Buttons
+        buttons = QDialogButtonBox(QDialogButtonBox.StandardButton.Ok | QDialogButtonBox.StandardButton.Cancel)
+        buttons.accepted.connect(self.accept)
+        buttons.rejected.connect(self.reject)
+        layout.addWidget(buttons)
+
+        # Initialize the filtered values
+        self.filtered_indices = list(range(len(self.tokens)))
+
+        # Load data for the first page
+        self.load_page()
+
+    def apply_filter(self):
+        """Filter the tokens based on the search text."""
+        filter_text = self.filter_edit.text().lower()
+
+        if not filter_text:
+            # No filter, show all values
+            self.filtered_indices = list(range(len(self.tokens)))
+        else:
+            # Apply filter
+            self.filtered_indices = []
+            for i, token in enumerate(self.tokens):
+                if filter_text in str(token).lower():
+                    self.filtered_indices.append(i)
+
+        # Reset to first page and reload
+        self.total_pages = max(1, (len(self.filtered_indices) + self.page_size - 1) // self.page_size)
+        self.current_page = 0
+        self.page_label.setText(f"Page 1 of {self.total_pages}")
+        self.load_page()
+
+    def previous_page(self):
+        """Go to the previous page of results."""
+        if self.current_page > 0:
+            self.current_page -= 1
+            self.page_label.setText(f"Page {self.current_page + 1} of {self.total_pages}")
+            self.load_page()
+
+    def next_page(self):
+        """Go to the next page of results."""
+        if self.current_page < self.total_pages - 1:
+            self.current_page += 1
+            self.page_label.setText(f"Page {self.current_page + 1} of {self.total_pages}")
+            self.load_page()
+
+    def load_page(self):
+        """Load the current page of tokenizer data."""
+        self.tokens_table.setRowCount(0)  # Clear the table
+
+        # Calculate start and end indices for the current page
+        start_idx = self.current_page * self.page_size
+        end_idx = min(start_idx + self.page_size, len(self.filtered_indices))
+
+        # Pre-allocate rows for better performance
+        self.tokens_table.setRowCount(end_idx - start_idx)
+
+        for row, i in enumerate(range(start_idx, end_idx)):
+            orig_idx = self.filtered_indices[i]
+
+            # Index
+            index_item = QTableWidgetItem(str(orig_idx))
+            index_item.setData(Qt.ItemDataRole.UserRole, orig_idx)  # Store original index
+            index_item.setFlags(index_item.flags() & ~Qt.ItemFlag.ItemIsEditable)
+            self.tokens_table.setItem(row, 0, index_item)
+
+            # Token
+            token_item = QTableWidgetItem(str(self.tokens[orig_idx]))
+            self.tokens_table.setItem(row, 1, token_item)
+
+            # Token Type
+            token_type = self.token_types[orig_idx] if orig_idx < len(self.token_types) else 0
+            try:
+                enum_val = TokenType(token_type)
+                display_text = f"{enum_val.name} ({token_type})"
+            except (ValueError, KeyError):
+                display_text = f"Unknown ({token_type})"
+
+            type_item = QTableWidgetItem(display_text)
+            type_item.setData(Qt.ItemDataRole.UserRole, token_type)
+
+            # Make type cell editable with a double-click handler
+            type_item.setFlags(type_item.flags() & ~Qt.ItemFlag.ItemIsEditable)
+            self.tokens_table.setItem(row, 2, type_item)
+
+            # Score
+            score = self.scores[orig_idx] if orig_idx < len(self.scores) else 0.0
+            score_item = QTableWidgetItem(str(score))
+            self.tokens_table.setItem(row, 3, score_item)
+
+        # Connect double-click handler for token type cells
+        self.tokens_table.cellDoubleClicked.connect(self.handle_cell_double_click)
+
+    def handle_cell_double_click(self, row, column):
+        """Handle double-click on a cell, specifically for token type editing."""
+        if column == 2:  # Token Type column
+            orig_item = self.tokens_table.item(row, 0)
+            if orig_item:
+                orig_idx = orig_item.data(Qt.ItemDataRole.UserRole)
+                self.edit_token_type(row, orig_idx)
+
+    def edit_token_type(self, row, orig_idx):
+        """Edit a token type using a dialog with a dropdown of all enum options."""
+        current_value = self.token_types[orig_idx] if orig_idx < len(self.token_types) else 0
+
+        # Create a dialog with enum options
+        dialog = QDialog(self)
+        dialog.setWindowTitle("Select Token Type")
+        layout = QVBoxLayout(dialog)
+
+        combo = QComboBox()
+        for enum_val in TokenType:
+            combo.addItem(f"{enum_val.name} ({enum_val.value})", enum_val.value)
+
+        # Set current value
+        try:
+            if isinstance(current_value, int):
+                enum_val = TokenType(current_value)
+                combo.setCurrentText(f"{enum_val.name} ({current_value})")
+        except (ValueError, KeyError):
+            pass
+
+        layout.addWidget(combo)
+
+        buttons = QDialogButtonBox(QDialogButtonBox.StandardButton.Ok | QDialogButtonBox.StandardButton.Cancel)
+        buttons.accepted.connect(dialog.accept)
+        buttons.rejected.connect(dialog.reject)
+        layout.addWidget(buttons)
+
+        if dialog.exec() == QDialog.DialogCode.Accepted:
+            # Get the selected value
+            new_value = combo.currentData()
+            enum_val = TokenType(new_value)
+            display_text = f"{enum_val.name} ({new_value})"
+
+            # Update the display
+            type_item = self.tokens_table.item(row, 2)
+            if type_item:
+                type_item.setText(display_text)
+                type_item.setData(Qt.ItemDataRole.UserRole, new_value)
+
+            # Update the actual value
+            self.token_types[orig_idx] = new_value
+
+    def add_token(self):
+        """Add a new token to the end of the list."""
+        # Add to the end of the arrays
+        self.tokens.append("")
+        self.token_types.append(0)  # Default to normal token
+        self.scores.append(0.0)
+
+        orig_idx = len(self.tokens) - 1
+
+        # Add to filtered indices if it matches the current filter
+        filter_text = self.filter_edit.text().lower()
+        if not filter_text or filter_text in "":
+            self.filtered_indices.append(orig_idx)
+
+        # Update pagination
+        self.total_pages = max(1, (len(self.filtered_indices) + self.page_size - 1) // self.page_size)
+
+        # Go to the last page to show the new item
+        self.current_page = self.total_pages - 1
+        self.page_label.setText(f"Page {self.current_page + 1} of {self.total_pages}")
+
+        # Reload the page
+        self.load_page()
+
+    def remove_selected(self):
+        """Remove selected tokens from all arrays."""
+        selected_rows = []
+        for item in self.tokens_table.selectedItems():
+            row = item.row()
+            if row not in selected_rows:
+                selected_rows.append(row)
+
+        if not selected_rows:
+            return
+
+        # Get original indices in descending order to avoid index shifting
+        orig_indices = []
+        for row in selected_rows:
+            orig_item = self.tokens_table.item(row, 0)
+            if orig_item:
+                orig_indices.append(orig_item.data(Qt.ItemDataRole.UserRole))
+        orig_indices.sort(reverse=True)
+
+        # Remove from all arrays
+        for idx in orig_indices:
+            if idx < len(self.tokens):
+                del self.tokens[idx]
+            if idx < len(self.token_types):
+                del self.token_types[idx]
+            if idx < len(self.scores):
+                del self.scores[idx]
+
+        # Rebuild filtered_indices
+        self.filtered_indices = []
+        filter_text = self.filter_edit.text().lower()
+
+        for i, token in enumerate(self.tokens):
+            if not filter_text or filter_text in str(token).lower():
+                self.filtered_indices.append(i)
+
+        # Update pagination
+        self.total_pages = max(1, (len(self.filtered_indices) + self.page_size - 1) // self.page_size)
+        self.current_page = min(self.current_page, self.total_pages - 1)
+        self.page_label.setText(f"Page {self.current_page + 1} of {self.total_pages}")
+
+        # Reload the page
+        self.load_page()
+
+    def get_data(self):
+        """Return the edited tokenizer data."""
+        return self.tokens, self.token_types, self.scores
+
+
+class ArrayEditorDialog(QDialog):
+    def __init__(self, array_values, element_type, key=None, parent=None):
+        super().__init__(parent)
+        self.setWindowTitle("Edit Array Values")
+        self.resize(700, 500)
+
+        self.array_values = array_values
+        self.element_type = element_type
+        self.key = key
+
+        # Get enum type for this array if applicable
+        self.enum_type = None
+        if key in KEY_TO_ENUM_TYPE and element_type == GGUFValueType.INT32:
+            self.enum_type = KEY_TO_ENUM_TYPE[key]
+
+        layout = QVBoxLayout(self)
+
+        # Add enum type information if applicable
+        if self.enum_type is not None:
+            enum_info_layout = QHBoxLayout()
+            enum_label = QLabel(f"Editing {self.enum_type.__name__} values:")
+            enum_info_layout.addWidget(enum_label)
+
+            # Add a legend for the enum values
+            enum_values = ", ".join([f"{e.name}={e.value}" for e in self.enum_type])
+            enum_values_label = QLabel(f"Available values: {enum_values}")
+            enum_values_label.setWordWrap(True)
+            enum_info_layout.addWidget(enum_values_label, 1)
+
+            layout.addLayout(enum_info_layout)
+
+        # Add search/filter controls
+        filter_layout = QHBoxLayout()
+        filter_layout.addWidget(QLabel("Filter:"))
+        self.filter_edit = QLineEdit()
+        self.filter_edit.setPlaceholderText("Type to filter values...")
+        self.filter_edit.textChanged.connect(self.apply_filter)
+        filter_layout.addWidget(self.filter_edit)
+
+        # Add page controls for large arrays
+        self.page_size = 100  # Show 100 items per page
+        self.current_page = 0
+        self.total_pages = max(1, (len(array_values) + self.page_size - 1) // self.page_size)
+
+        self.page_label = QLabel(f"Page 1 of {self.total_pages}")
+        filter_layout.addWidget(self.page_label)
+
+        prev_page = QPushButton("Previous")
+        prev_page.clicked.connect(self.previous_page)
+        filter_layout.addWidget(prev_page)
+
+        next_page = QPushButton("Next")
+        next_page.clicked.connect(self.next_page)
+        filter_layout.addWidget(next_page)
+
+        layout.addLayout(filter_layout)
+
+        # Array items table
+        self.items_table = QTableWidget()
+
+        # Set up columns based on whether we have an enum type
+        if self.enum_type is not None:
+            self.items_table.setColumnCount(3)
+            self.items_table.setHorizontalHeaderLabels(["Index", "Value", "Actions"])
+            self.items_table.horizontalHeader().setSectionResizeMode(0, QHeaderView.ResizeMode.ResizeToContents)
+            self.items_table.horizontalHeader().setSectionResizeMode(1, QHeaderView.ResizeMode.Stretch)
+            self.items_table.horizontalHeader().setSectionResizeMode(2, QHeaderView.ResizeMode.ResizeToContents)
+        else:
+            self.items_table.setColumnCount(2)
+            self.items_table.setHorizontalHeaderLabels(["Index", "Value"])
+            self.items_table.horizontalHeader().setSectionResizeMode(0, QHeaderView.ResizeMode.ResizeToContents)
+            self.items_table.horizontalHeader().setSectionResizeMode(1, QHeaderView.ResizeMode.Stretch)
+
+        layout.addWidget(self.items_table)
+
+        # Controls
+        controls_layout = QHBoxLayout()
+
+        add_button = QPushButton("Add Item")
+        add_button.clicked.connect(self.add_item)
+        controls_layout.addWidget(add_button)
+
+        remove_button = QPushButton("Remove Selected")
+        remove_button.clicked.connect(self.remove_selected)
+        controls_layout.addWidget(remove_button)
+
+        # Add bulk edit button for enum arrays
+        if self.enum_type is not None:
+            bulk_edit_button = QPushButton("Bulk Edit Selected")
+            bulk_edit_button.clicked.connect(self.bulk_edit_selected)
+            controls_layout.addWidget(bulk_edit_button)
+
+        controls_layout.addStretch()
+
+        layout.addLayout(controls_layout)
+
+        # Buttons
+        buttons = QDialogButtonBox(QDialogButtonBox.StandardButton.Ok | QDialogButtonBox.StandardButton.Cancel)
+        buttons.accepted.connect(self.accept)
+        buttons.rejected.connect(self.reject)
+        layout.addWidget(buttons)
+
+        # Initialize the filtered values
+        self.filtered_indices = list(range(len(self.array_values)))
+
+        # Load array values for the first page
+        self.load_page()
+
+    def apply_filter(self):
+        """Filter the array values based on the search text."""
+        filter_text = self.filter_edit.text().lower()
+
+        if not filter_text:
+            # No filter, show all values
+            self.filtered_indices = list(range(len(self.array_values)))
+        else:
+            # Apply filter
+            self.filtered_indices = []
+            for i, value in enumerate(self.array_values):
+                # For enum values, search in both name and value
+                if self.enum_type is not None and isinstance(value, int):
+                    try:
+                        enum_val = self.enum_type(value)
+                        display_text = f"{enum_val.name} ({value})".lower()
+                        if filter_text in display_text:
+                            self.filtered_indices.append(i)
+                    except (ValueError, KeyError):
+                        # If not a valid enum value, just check the raw value
+                        if filter_text in str(value).lower():
+                            self.filtered_indices.append(i)
+                else:
+                    # For non-enum values, just check the string representation
+                    if filter_text in str(value).lower():
+                        self.filtered_indices.append(i)
+
+        # Reset to first page and reload
+        self.total_pages = max(1, (len(self.filtered_indices) + self.page_size - 1) // self.page_size)
+        self.current_page = 0
+        self.page_label.setText(f"Page 1 of {self.total_pages}")
+        self.load_page()
+
+    def previous_page(self):
+        """Go to the previous page of results."""
+        if self.current_page > 0:
+            self.current_page -= 1
+            self.page_label.setText(f"Page {self.current_page + 1} of {self.total_pages}")
+            self.load_page()
+
+    def next_page(self):
+        """Go to the next page of results."""
+        if self.current_page < self.total_pages - 1:
+            self.current_page += 1
+            self.page_label.setText(f"Page {self.current_page + 1} of {self.total_pages}")
+            self.load_page()
+
+    def load_page(self):
+        """Load the current page of array values."""
+        self.items_table.setRowCount(0)  # Clear the table
+
+        # Calculate start and end indices for the current page
+        start_idx = self.current_page * self.page_size
+        end_idx = min(start_idx + self.page_size, len(self.filtered_indices))
+
+        # Pre-allocate rows for better performance
+        self.items_table.setRowCount(end_idx - start_idx)
+
+        for row, i in enumerate(range(start_idx, end_idx)):
+            orig_idx = self.filtered_indices[i]
+            value = self.array_values[orig_idx]
+
+            # Index
+            index_item = QTableWidgetItem(str(orig_idx))
+            index_item.setData(Qt.ItemDataRole.UserRole, orig_idx)  # Store original index
+            index_item.setFlags(index_item.flags() & ~Qt.ItemFlag.ItemIsEditable)
+            self.items_table.setItem(row, 0, index_item)
+
+            # Value
+            if self.enum_type is not None:
+                # Display enum value and name
+                try:
+                    if isinstance(value, (int, numpy.signedinteger)):
+                        enum_val = self.enum_type(value)
+                        display_text = f"{enum_val.name} ({value})"
+                    else:
+                        display_text = str(value)
+                except (ValueError, KeyError):
+                    display_text = f"Unknown ({value})"
+
+                # Store the enum value in the item
+                value_item = QTableWidgetItem(display_text)
+                value_item.setData(Qt.ItemDataRole.UserRole, value)
+                value_item.setFlags(value_item.flags() & ~Qt.ItemFlag.ItemIsEditable)
+                self.items_table.setItem(row, 1, value_item)
+
+                # Add an edit button in a separate column
+                edit_button = QPushButton("Edit")
+                edit_button.setProperty("row", row)
+                edit_button.clicked.connect(self.edit_array_enum_value)
+
+                # Create a widget to hold the button
+                button_widget = QWidget()
+                button_layout = QHBoxLayout(button_widget)
+                button_layout.setContentsMargins(2, 2, 2, 2)
+                button_layout.addWidget(edit_button)
+                button_layout.addStretch()
+
+                self.items_table.setCellWidget(row, 2, button_widget)
+            else:
+                value_item = QTableWidgetItem(str(value))
+                self.items_table.setItem(row, 1, value_item)
+
+    def edit_array_enum_value(self):
+        """Handle editing an enum value in the array editor."""
+        button = self.sender()
+        row = button.property("row")
+
+        # Get the original index from the table item
+        orig_item = self.items_table.item(row, 0)
+        new_item = self.items_table.item(row, 1)
+        if orig_item and new_item and self.enum_type and self.edit_enum_value(row, self.enum_type):
+            orig_idx = orig_item.data(Qt.ItemDataRole.UserRole)
+            new_value = new_item.data(Qt.ItemDataRole.UserRole)
+            # Update the stored value in the array
+            if isinstance(new_value, (int, float, str, bool)):
+                self.array_values[orig_idx] = new_value
+
+    def bulk_edit_selected(self):
+        """Edit multiple enum values at once."""
+        if not self.enum_type:
+            return
+
+        selected_rows = set()
+        for item in self.items_table.selectedItems():
+            selected_rows.add(item.row())
+
+        if not selected_rows:
+            QMessageBox.information(self, "No Selection", "Please select at least one row to edit.")
+            return
+
+        # Create a dialog with enum options
+        dialog = QDialog(self)
+        dialog.setWindowTitle(f"Bulk Edit {self.enum_type.__name__} Values")
+        layout = QVBoxLayout(dialog)
+
+        layout.addWidget(QLabel(f"Set {len(selected_rows)} selected items to:"))
+
+        combo = QComboBox()
+        for enum_val in self.enum_type:
+            combo.addItem(f"{enum_val.name} ({enum_val.value})", enum_val.value)
+
+        layout.addWidget(combo)
+
+        buttons = QDialogButtonBox(QDialogButtonBox.StandardButton.Ok | QDialogButtonBox.StandardButton.Cancel)
+        buttons.accepted.connect(dialog.accept)
+        buttons.rejected.connect(dialog.reject)
+        layout.addWidget(buttons)
+
+        if dialog.exec() == QDialog.DialogCode.Accepted:
+            # Get the selected value
+            new_value = combo.currentData()
+            enum_val = self.enum_type(new_value)
+            display_text = f"{enum_val.name} ({new_value})"
+
+            # Update all selected rows
+            for row in selected_rows:
+                orig_item = self.items_table.item(row, 0)
+                new_item = self.items_table.item(row, 1)
+                if orig_item and new_item:
+                    orig_idx = orig_item.data(Qt.ItemDataRole.UserRole)
+                    self.array_values[orig_idx] = new_value
+
+                    # Update the display
+                    new_item.setText(display_text)
+                    new_item.setData(Qt.ItemDataRole.UserRole, new_value)
+
+    def add_item(self):
+        # Add to the end of the array
+        orig_idx = len(self.array_values)
+
+        # Add default value based on type
+        if self.enum_type is not None:
+            # Default to first enum value
+            default_value = list(self.enum_type)[0].value
+            self.array_values.append(default_value)
+        else:
+            if self.element_type == GGUFValueType.STRING:
+                self.array_values.append("")
+            else:
+                self.array_values.append(0)
+
+        # Add to filtered indices if it matches the current filter
+        self.filtered_indices.append(orig_idx)
+
+        # Update pagination
+        self.total_pages = max(1, (len(self.filtered_indices) + self.page_size - 1) // self.page_size)
+
+        # Go to the last page to show the new item
+        self.current_page = self.total_pages - 1
+        self.page_label.setText(f"Page {self.current_page + 1} of {self.total_pages}")
+
+        # Reload the page
+        self.load_page()
+
+    def remove_selected(self):
+        selected_rows = []
+        for item in self.items_table.selectedItems():
+            row = item.row()
+            if row not in selected_rows:
+                selected_rows.append(row)
+
+        if not selected_rows:
+            return
+
+        # Get original indices in descending order to avoid index shifting
+        orig_indices = list()
+        for row in selected_rows:
+            orig_item = self.items_table.item(row, 0)
+            if orig_item:
+                orig_indices.append(orig_item.data(Qt.ItemDataRole.UserRole))
+        orig_indices.sort(reverse=True)
+
+        # Remove from array_values
+        for idx in orig_indices:
+            del self.array_values[idx]
+
+        # Rebuild filtered_indices
+        self.filtered_indices = []
+        filter_text = self.filter_edit.text().lower()
+
+        for i, value in enumerate(self.array_values):
+            if not filter_text:
+                self.filtered_indices.append(i)
+            else:
+                # Apply filter
+                if self.enum_type is not None and isinstance(value, int):
+                    try:
+                        enum_val = self.enum_type(value)
+                        display_text = f"{enum_val.name} ({value})".lower()
+                        if filter_text in display_text:
+                            self.filtered_indices.append(i)
+                    except (ValueError, KeyError):
+                        if filter_text in str(value).lower():
+                            self.filtered_indices.append(i)
+                else:
+                    if filter_text in str(value).lower():
+                        self.filtered_indices.append(i)
+
+        # Update pagination
+        self.total_pages = max(1, (len(self.filtered_indices) + self.page_size - 1) // self.page_size)
+        self.current_page = min(self.current_page, self.total_pages - 1)
+        self.page_label.setText(f"Page {self.current_page + 1} of {self.total_pages}")
+
+        # Reload the page
+        self.load_page()
+
+    def edit_enum_value(self, row: int, enum_type: Type[enum.Enum]):
+        """Edit an enum value using a dialog with a dropdown of all enum options."""
+        # Get the original index from the table item
+        orig_item = self.items_table.item(row, 0)
+        if orig_item:
+            orig_idx = orig_item.data(Qt.ItemDataRole.UserRole)
+        else:
+            return
+        current_value = self.array_values[orig_idx]
+
+        # Create a dialog with enum options
+        dialog = QDialog(self)
+        dialog.setWindowTitle(f"Select {enum_type.__name__} Value")
+        layout = QVBoxLayout(dialog)
+
+        # Add description
+        description = QLabel(f"Select a {enum_type.__name__} value:")
+        layout.addWidget(description)
+
+        # Use a combo box for quick selection
+        combo = QComboBox()
+        for enum_val in enum_type:
+            combo.addItem(f"{enum_val.name} ({enum_val.value})", enum_val.value)
+
+        # Set current value
+        try:
+            if isinstance(current_value, int):
+                enum_val = enum_type(current_value)
+                combo.setCurrentText(f"{enum_val.name} ({current_value})")
+        except (ValueError, KeyError):
+            pass
+
+        layout.addWidget(combo)
+
+        buttons = QDialogButtonBox(QDialogButtonBox.StandardButton.Ok | QDialogButtonBox.StandardButton.Cancel)
+        buttons.accepted.connect(dialog.accept)
+        buttons.rejected.connect(dialog.reject)
+        layout.addWidget(buttons)
+
+        if dialog.exec() == QDialog.DialogCode.Accepted:
+            # Update the value display and stored data
+            new_value = combo.currentData()
+            enum_val = enum_type(new_value)
+            display_text = f"{enum_val.name} ({new_value})"
+
+            new_item = self.items_table.item(row, 1)
+            if new_item:
+                new_item.setText(display_text)
+                new_item.setData(Qt.ItemDataRole.UserRole, new_value)
+
+            # Update the actual array value
+            self.array_values[orig_idx] = new_value
+            return True
+        return False
+
+    def get_array_values(self):
+        # The array_values list is kept up-to-date as edits are made
+        return self.array_values
+
+
+class AddMetadataDialog(QDialog):
+    def __init__(self, parent=None):
+        super().__init__(parent)
+        self.setWindowTitle("Add Metadata")
+        self.resize(400, 200)
+
+        layout = QVBoxLayout(self)
+
+        form_layout = QFormLayout()
+
+        self.key_edit = QLineEdit()
+        form_layout.addRow("Key:", self.key_edit)
+
+        self.type_combo = QComboBox()
+        for value_type in GGUFValueType:
+            if value_type != GGUFValueType.ARRAY:  # Skip array type for simplicity
+                self.type_combo.addItem(value_type.name, value_type)
+        form_layout.addRow("Type:", self.type_combo)
+
+        self.value_edit = QTextEdit()
+        form_layout.addRow("Value:", self.value_edit)
+
+        layout.addLayout(form_layout)
+
+        buttons = QDialogButtonBox(QDialogButtonBox.StandardButton.Ok | QDialogButtonBox.StandardButton.Cancel)
+        buttons.accepted.connect(self.accept)
+        buttons.rejected.connect(self.reject)
+        layout.addWidget(buttons)
+
+    def get_data(self) -> Tuple[str, GGUFValueType, Any]:
+        key = self.key_edit.text()
+        value_type = self.type_combo.currentData()
+        value_text = self.value_edit.toPlainText()
+
+        # Convert value based on type
+        if value_type == GGUFValueType.UINT8:
+            value = np.uint8(int(value_text))
+        elif value_type == GGUFValueType.INT8:
+            value = np.int8(int(value_text))
+        elif value_type == GGUFValueType.UINT16:
+            value = np.uint16(int(value_text))
+        elif value_type == GGUFValueType.INT16:
+            value = np.int16(int(value_text))
+        elif value_type == GGUFValueType.UINT32:
+            value = np.uint32(int(value_text))
+        elif value_type == GGUFValueType.INT32:
+            value = np.int32(int(value_text))
+        elif value_type == GGUFValueType.FLOAT32:
+            value = np.float32(float(value_text))
+        elif value_type == GGUFValueType.BOOL:
+            value = value_text.lower() in ('true', 'yes', '1')
+        elif value_type == GGUFValueType.STRING:
+            value = value_text
+        else:
+            value = value_text
+
+        return key, value_type, value
+
+
+class GGUFEditorWindow(QMainWindow):
+    def __init__(self):
+        super().__init__()
+
+        self.setWindowTitle("GGUF Editor")
+        self.resize(1000, 800)
+
+        self.current_file = None
+        self.reader = None
+        self.modified = False
+        self.metadata_changes = {}  # Store changes to apply when saving
+        self.metadata_to_remove = set()  # Store keys to remove when saving
+
+        self.setup_ui()
+
+    def setup_ui(self):
+        central_widget = QWidget()
+        self.setCentralWidget(central_widget)
+
+        main_layout = QVBoxLayout(central_widget)
+
+        # File controls
+        file_layout = QHBoxLayout()
+
+        self.file_path_edit = QLineEdit()
+        self.file_path_edit.setReadOnly(True)
+        file_layout.addWidget(self.file_path_edit)
+
+        open_button = QPushButton("Open GGUF")
+        open_button.clicked.connect(self.open_file)
+        file_layout.addWidget(open_button)
+
+        save_button = QPushButton("Save As...")
+        save_button.clicked.connect(self.save_file)
+        file_layout.addWidget(save_button)
+
+        main_layout.addLayout(file_layout)
+
+        # Tabs for different views
+        self.tabs = QTabWidget()
+
+        # Metadata tab
+        self.metadata_tab = QWidget()
+        metadata_layout = QVBoxLayout(self.metadata_tab)
+
+        # Metadata table
+        self.metadata_table = QTableWidget()
+        self.metadata_table.setColumnCount(4)
+        self.metadata_table.setHorizontalHeaderLabels(["Key", "Type", "Value", "Actions"])
+        self.metadata_table.horizontalHeader().setSectionResizeMode(0, QHeaderView.ResizeMode.Stretch)
+        self.metadata_table.horizontalHeader().setSectionResizeMode(1, QHeaderView.ResizeMode.ResizeToContents)
+        self.metadata_table.horizontalHeader().setSectionResizeMode(2, QHeaderView.ResizeMode.Stretch)
+        self.metadata_table.horizontalHeader().setSectionResizeMode(3, QHeaderView.ResizeMode.ResizeToContents)
+        metadata_layout.addWidget(self.metadata_table)
+
+        # Metadata controls
+        metadata_controls = QHBoxLayout()
+
+        add_metadata_button = QPushButton("Add Metadata")
+        add_metadata_button.clicked.connect(self.add_metadata)
+        metadata_controls.addWidget(add_metadata_button)
+
+        metadata_controls.addStretch()
+
+        metadata_layout.addLayout(metadata_controls)
+
+        # Tensors tab
+        self.tensors_tab = QWidget()
+        tensors_layout = QVBoxLayout(self.tensors_tab)
+
+        self.tensors_table = QTableWidget()
+        self.tensors_table.setColumnCount(5)
+        self.tensors_table.setHorizontalHeaderLabels(["Name", "Type", "Shape", "Elements", "Size (bytes)"])
+        self.tensors_table.horizontalHeader().setSectionResizeMode(0, QHeaderView.ResizeMode.Stretch)
+        self.tensors_table.horizontalHeader().setSectionResizeMode(1, QHeaderView.ResizeMode.ResizeToContents)
+        self.tensors_table.horizontalHeader().setSectionResizeMode(2, QHeaderView.ResizeMode.ResizeToContents)
+        self.tensors_table.horizontalHeader().setSectionResizeMode(3, QHeaderView.ResizeMode.ResizeToContents)
+        self.tensors_table.horizontalHeader().setSectionResizeMode(4, QHeaderView.ResizeMode.ResizeToContents)
+        tensors_layout.addWidget(self.tensors_table)
+
+        # Add tabs to tab widget
+        self.tabs.addTab(self.metadata_tab, "Metadata")
+        self.tabs.addTab(self.tensors_tab, "Tensors")
+
+        main_layout.addWidget(self.tabs)
+
+        # Status bar
+        self.statusBar().showMessage("Ready")
+
+    def load_file(self, file_path):
+        """Load a GGUF file by path"""
+        try:
+            self.statusBar().showMessage(f"Loading {file_path}...")
+            QApplication.processEvents()
+
+            self.reader = GGUFReader(file_path, 'r')
+            self.current_file = file_path
+            self.file_path_edit.setText(file_path)
+
+            self.load_metadata()
+            self.load_tensors()
+
+            self.metadata_changes = {}
+            self.metadata_to_remove = set()
+            self.modified = False
+
+            self.statusBar().showMessage(f"Loaded {file_path}")
+            return True
+        except Exception as e:
+            QMessageBox.critical(self, "Error", f"Failed to open file: {str(e)}")
+            self.statusBar().showMessage("Error loading file")
+            return False
+
+    def open_file(self):
+        file_path, _ = QFileDialog.getOpenFileName(
+            self, "Open GGUF File", "", "GGUF Files (*.gguf);;All Files (*)"
+        )
+
+        if not file_path:
+            return
+
+        self.load_file(file_path)
+
+    def load_metadata(self):
+        self.metadata_table.setRowCount(0)
+
+        if not self.reader:
+            return
+
+        # Disconnect to prevent triggering during loading
+        with warnings.catch_warnings():
+            warnings.filterwarnings('ignore')
+            self.metadata_table.itemChanged.disconnect(self.on_metadata_changed)
+
+        for i, (key, field) in enumerate(self.reader.fields.items()):
+            self.metadata_table.insertRow(i)
+
+            # Key
+            key_item = QTableWidgetItem(key)
+            key_item.setFlags(key_item.flags() & ~Qt.ItemFlag.ItemIsEditable)
+            self.metadata_table.setItem(i, 0, key_item)
+
+            # Type
+            if not field.types:
+                type_str = "N/A"
+            elif field.types[0] == GGUFValueType.ARRAY:
+                nest_count = len(field.types) - 1
+                element_type = field.types[-1].name
+                # Check if this is an enum array
+                enum_type = self.get_enum_for_key(key)
+                if enum_type is not None and field.types[-1] == GGUFValueType.INT32:
+                    element_type = enum_type.__name__
+                type_str = '[' * nest_count + element_type + ']' * nest_count
+            else:
+                type_str = str(field.types[0].name)
+                # Check if this is an enum field
+                enum_type = self.get_enum_for_key(key)
+                if enum_type is not None and field.types[0] == GGUFValueType.INT32:
+                    type_str = enum_type.__name__
+
+            type_item = QTableWidgetItem(type_str)
+            type_item.setFlags(type_item.flags() & ~Qt.ItemFlag.ItemIsEditable)
+            self.metadata_table.setItem(i, 1, type_item)
+
+            # Value
+            value_str = self.format_field_value(field)
+            value_item = QTableWidgetItem(value_str)
+
+            # Make only simple values editable
+            if len(field.types) == 1 and field.types[0] != GGUFValueType.ARRAY:
+                value_item.setFlags(value_item.flags() | Qt.ItemFlag.ItemIsEditable)
+            else:
+                value_item.setFlags(value_item.flags() & ~Qt.ItemFlag.ItemIsEditable)
+
+            self.metadata_table.setItem(i, 2, value_item)
+
+            # Actions
+            actions_widget = QWidget()
+            actions_layout = QHBoxLayout(actions_widget)
+            actions_layout.setContentsMargins(2, 2, 2, 2)
+
+            # Add Edit button for arrays and enum fields
+            if field.types and field.types[0] == GGUFValueType.ARRAY:
+                edit_button = QPushButton("Edit")
+                edit_button.setProperty("row", i)
+                edit_button.setProperty("key", key)
+                edit_button.clicked.connect(self.edit_array_metadata)
+                actions_layout.addWidget(edit_button)
+
+                # Add special label for tokenizer linked fields
+                if key in TOKENIZER_LINKED_KEYS:
+                    edit_button.setText("Edit Tokenizer")
+                    edit_button.setToolTip("Edit all tokenizer data together")
+            elif len(field.types) == 1 and self.get_enum_for_key(key) is not None:
+                edit_button = QPushButton("Edit")
+                edit_button.setProperty("row", i)
+                edit_button.setProperty("key", key)
+                edit_button.clicked.connect(self.edit_metadata_enum)
+                actions_layout.addWidget(edit_button)
+
+            remove_button = QPushButton("Remove")
+            remove_button.setProperty("row", i)
+            remove_button.setProperty("key", key)
+            remove_button.clicked.connect(self.remove_metadata)
+            actions_layout.addWidget(remove_button)
+
+            self.metadata_table.setCellWidget(i, 3, actions_widget)
+
+        # Reconnect after loading
+        self.metadata_table.itemChanged.connect(self.on_metadata_changed)
+
+    def extract_array_values(self, field: ReaderField) -> list:
+        """Extract all values from an array field."""
+        if not field.types or field.types[0] != GGUFValueType.ARRAY:
+            return []
+
+        curr_type = field.types[1]
+        array_values = []
+        total_elements = len(field.data)
+
+        if curr_type == GGUFValueType.STRING:
+            for element_pos in range(total_elements):
+                value_string = str(bytes(field.parts[-1 - (total_elements - element_pos - 1) * 2]), encoding='utf-8')
+                array_values.append(value_string)
+        elif self.reader and curr_type in self.reader.gguf_scalar_to_np:
+            for element_pos in range(total_elements):
+                array_values.append(field.parts[-1 - (total_elements - element_pos - 1)][0])
+
+        return array_values
+
+    def get_enum_for_key(self, key: str) -> Optional[Type[enum.Enum]]:
+        """Get the enum type for a given key if it exists."""
+        return KEY_TO_ENUM_TYPE.get(key)
+
+    def format_enum_value(self, value: Any, enum_type: Type[enum.Enum]) -> str:
+        """Format a value as an enum if possible."""
+        try:
+            if isinstance(value, (int, str)):
+                enum_value = enum_type(value)
+                return f"{enum_value.name} ({value})"
+        except (ValueError, KeyError):
+            pass
+        return str(value)
+
+    def format_field_value(self, field: ReaderField) -> str:
+        if not field.types:
+            return "N/A"
+
+        if len(field.types) == 1:
+            curr_type = field.types[0]
+            if curr_type == GGUFValueType.STRING:
+                return str(bytes(field.parts[-1]), encoding='utf-8')
+            elif self.reader and curr_type in self.reader.gguf_scalar_to_np:
+                value = field.parts[-1][0]
+                # Check if this field has an enum type
+                enum_type = self.get_enum_for_key(field.name)
+                if enum_type is not None:
+                    return self.format_enum_value(value, enum_type)
+                return str(value)
+
+        if field.types[0] == GGUFValueType.ARRAY:
+            array_values = self.extract_array_values(field)
+            render_element = min(5, len(array_values))
+
+            # Get enum type for this array if applicable
+            enum_type = self.get_enum_for_key(field.name)
+
+            if enum_type is not None:
+                array_elements = []
+                for i in range(render_element):
+                    array_elements.append(self.format_enum_value(array_values[i], enum_type))
+            else:
+                array_elements = [str(array_values[i]) for i in range(render_element)]
+
+            return f"[ {', '.join(array_elements).strip()}{', ...' if len(array_values) > len(array_elements) else ''} ]"
+
+        return "Complex value"
+
+    def load_tensors(self):
+        self.tensors_table.setRowCount(0)
+
+        if not self.reader:
+            return
+
+        for i, tensor in enumerate(self.reader.tensors):
+            self.tensors_table.insertRow(i)
+
+            # Name
+            name_item = QTableWidgetItem(tensor.name)
+            name_item.setFlags(name_item.flags() & ~Qt.ItemFlag.ItemIsEditable)
+            self.tensors_table.setItem(i, 0, name_item)
+
+            # Type
+            type_item = QTableWidgetItem(tensor.tensor_type.name)
+            type_item.setFlags(type_item.flags() & ~Qt.ItemFlag.ItemIsEditable)
+            self.tensors_table.setItem(i, 1, type_item)
+
+            # Shape
+            shape_str = " × ".join(str(d) for d in tensor.shape)
+            shape_item = QTableWidgetItem(shape_str)
+            shape_item.setFlags(shape_item.flags() & ~Qt.ItemFlag.ItemIsEditable)
+            self.tensors_table.setItem(i, 2, shape_item)
+
+            # Elements
+            elements_item = QTableWidgetItem(str(tensor.n_elements))
+            elements_item.setFlags(elements_item.flags() & ~Qt.ItemFlag.ItemIsEditable)
+            self.tensors_table.setItem(i, 3, elements_item)
+
+            # Size
+            size_item = QTableWidgetItem(f"{tensor.n_bytes:,}")
+            size_item.setFlags(size_item.flags() & ~Qt.ItemFlag.ItemIsEditable)
+            self.tensors_table.setItem(i, 4, size_item)
+
+    def on_metadata_changed(self, item):
+        if item.column() != 2:  # Only handle value column changes
+            return
+
+        row = item.row()
+        orig_item = self.metadata_table.item(row, 0)
+        key = None
+        if orig_item:
+            key = orig_item.text()
+        new_value = item.text()
+
+        field = None
+        if self.reader and key:
+            field = self.reader.get_field(key)
+        if not field or not field.types or not key:
+            return
+
+        value_type = field.types[0]
+
+        # Check if this is an enum field
+        enum_type = self.get_enum_for_key(key)
+        if enum_type is not None and value_type == GGUFValueType.INT32:
+            # Try to parse the enum value from the text
+            try:
+                # Check if it's a name
+                try:
+                    enum_val = enum_type[new_value]
+                    converted_value = enum_val.value
+                except (KeyError, AttributeError):
+                    # Check if it's a number or "NAME (value)" format
+                    if '(' in new_value and ')' in new_value:
+                        # Extract the value from "NAME (value)" format
+                        value_part = new_value.split('(')[1].split(')')[0].strip()
+                        converted_value = int(value_part)
+                    else:
+                        # Try to convert directly to int
+                        converted_value = int(new_value)
+
+                # Validate that it's a valid enum value
+                enum_type(converted_value)
+
+                # Store the change
+                self.metadata_changes[key] = (value_type, converted_value)
+                self.modified = True
+
+                # Update display with formatted enum value
+                formatted_value = self.format_enum_value(converted_value, enum_type)
+                item.setText(formatted_value)
+
+                self.statusBar().showMessage(f"Changed {key} to {formatted_value}")
+                return
+            except (ValueError, KeyError) as e:
+                QMessageBox.warning(
+                    self,
+                    f"Invalid Enum Value ({e})",
+                    f"'{new_value}' is not a valid {enum_type.__name__} value.\n"
+                    f"Valid values are: {', '.join(v.name for v in enum_type)}")
+
+                # Revert to original value
+                original_value = self.format_field_value(field)
+                item.setText(original_value)
+                return
+
+        try:
+            # Convert the string value to the appropriate type
+            if value_type == GGUFValueType.UINT8:
+                converted_value = np.uint8(int(new_value))
+            elif value_type == GGUFValueType.INT8:
+                converted_value = np.int8(int(new_value))
+            elif value_type == GGUFValueType.UINT16:
+                converted_value = np.uint16(int(new_value))
+            elif value_type == GGUFValueType.INT16:
+                converted_value = np.int16(int(new_value))
+            elif value_type == GGUFValueType.UINT32:
+                converted_value = np.uint32(int(new_value))
+            elif value_type == GGUFValueType.INT32:
+                converted_value = np.int32(int(new_value))
+            elif value_type == GGUFValueType.FLOAT32:
+                converted_value = np.float32(float(new_value))
+            elif value_type == GGUFValueType.BOOL:
+                converted_value = new_value.lower() in ('true', 'yes', '1')
+            elif value_type == GGUFValueType.STRING:
+                converted_value = new_value
+            else:
+                # Unsupported type for editing
+                return
+
+            # Store the change
+            self.metadata_changes[key] = (value_type, converted_value)
+            self.modified = True
+
+            self.statusBar().showMessage(f"Changed {key} to {new_value}")
+        except ValueError:
+            QMessageBox.warning(self, "Invalid Value", f"The value '{new_value}' is not valid for type {value_type.name}")
+
+            # Revert to original value
+            original_value = self.format_field_value(field)
+            item.setText(original_value)
+
+    def remove_metadata(self):
+        button = self.sender()
+        key = button.property("key")
+        row = button.property("row")
+
+        reply = QMessageBox.question(
+            self, "Confirm Removal",
+            f"Are you sure you want to remove the metadata key '{key}'?",
+            QMessageBox.StandardButton.Yes | QMessageBox.StandardButton.No, QMessageBox.StandardButton.No
+        )
+
+        if reply == QMessageBox.StandardButton.Yes:
+            self.metadata_table.removeRow(row)
+            self.metadata_to_remove.add(key)
+
+            # If we previously had changes for this key, remove them
+            if key in self.metadata_changes:
+                del self.metadata_changes[key]
+
+            self.modified = True
+            self.statusBar().showMessage(f"Marked {key} for removal")
+
+    def edit_metadata_enum(self):
+        """Edit an enum metadata field."""
+        button = self.sender()
+        key = button.property("key")
+        row = button.property("row")
+
+        field = None
+        if self.reader:
+            field = self.reader.get_field(key)
+        if not field or not field.types:
+            return
+
+        enum_type = self.get_enum_for_key(key)
+        if enum_type is None:
+            return
+
+        # Get current value
+        current_value = field.contents()
+
+        # Create a dialog with enum options
+        dialog = QDialog(self)
+        dialog.setWindowTitle(f"Select {enum_type.__name__} Value")
+        layout = QVBoxLayout(dialog)
+
+        combo = QComboBox()
+        for enum_val in enum_type:
+            combo.addItem(f"{enum_val.name} ({enum_val.value})", enum_val.value)
+
+        # Set current value
+        try:
+            if isinstance(current_value, (int, str)):
+                enum_val = enum_type(current_value)
+                combo.setCurrentText(f"{enum_val.name} ({current_value})")
+        except (ValueError, KeyError):
+            pass
+
+        layout.addWidget(combo)
+
+        buttons = QDialogButtonBox(QDialogButtonBox.StandardButton.Ok | QDialogButtonBox.StandardButton.Cancel)
+        buttons.accepted.connect(dialog.accept)
+        buttons.rejected.connect(dialog.reject)
+        layout.addWidget(buttons)
+
+        if dialog.exec() == QDialog.DialogCode.Accepted:
+            # Get the selected value
+            new_value = combo.currentData()
+            enum_val = enum_type(new_value)
+
+            # Store the change
+            self.metadata_changes[key] = (field.types[0], new_value)
+            self.modified = True
+
+            # Update display
+            display_text = f"{enum_val.name} ({new_value})"
+            target_item = self.metadata_table.item(row, 2)
+            if target_item:
+                target_item.setText(display_text)
+
+            self.statusBar().showMessage(f"Changed {key} to {display_text}")
+
+    def edit_array_metadata(self):
+        button = self.sender()
+        key = button.property("key")
+        row = button.property("row")
+
+        # Check if this is one of the linked tokenizer keys
+        if key in TOKENIZER_LINKED_KEYS:
+            self.edit_tokenizer_metadata(key)
+            return
+
+        field = None
+        if self.reader:
+            field = self.reader.get_field(key)
+        if not field or not field.types or field.types[0] != GGUFValueType.ARRAY:
+            return
+
+        # Get array element type
+        element_type = field.types[1]
+
+        # Extract array values
+        array_values = self.extract_array_values(field)
+
+        # Open array editor dialog
+        dialog = ArrayEditorDialog(array_values, element_type, key, self)
+        if dialog.exec() == QDialog.DialogCode.Accepted:
+            new_values = dialog.get_array_values()
+
+            # Store the change
+            self.metadata_changes[key] = (GGUFValueType.ARRAY, (element_type, new_values))
+            self.modified = True
+
+            # Update display
+            enum_type = self.get_enum_for_key(key)
+            if enum_type is not None and element_type == GGUFValueType.INT32:
+                value_str = f"[ {', '.join(self.format_enum_value(v, enum_type) for v in new_values[:5])}{', ...' if len(new_values) > 5 else ''} ]"
+            else:
+                value_str = f"[ {', '.join(str(v) for v in new_values[:5])}{', ...' if len(new_values) > 5 else ''} ]"
+            target_item = self.metadata_table.item(row, 2)
+            if target_item:
+                target_item.setText(value_str)
+
+            self.statusBar().showMessage(f"Updated array values for {key}")
+
+    def edit_tokenizer_metadata(self, trigger_key):
+        """Edit the linked tokenizer metadata arrays together."""
+        if not self.reader:
+            return
+
+        # Get all three fields
+        tokens_field = self.reader.get_field(gguf.Keys.Tokenizer.LIST)
+        token_types_field = self.reader.get_field(gguf.Keys.Tokenizer.TOKEN_TYPE)
+        scores_field = self.reader.get_field(gguf.Keys.Tokenizer.SCORES)
+
+        # Extract values from each field
+        tokens = self.extract_array_values(tokens_field) if tokens_field else []
+        token_types = self.extract_array_values(token_types_field) if token_types_field else []
+        scores = self.extract_array_values(scores_field) if scores_field else []
+
+        # Apply any pending changes
+        if gguf.Keys.Tokenizer.LIST in self.metadata_changes:
+            _, (_, tokens) = self.metadata_changes[gguf.Keys.Tokenizer.LIST]
+        if gguf.Keys.Tokenizer.TOKEN_TYPE in self.metadata_changes:
+            _, (_, token_types) = self.metadata_changes[gguf.Keys.Tokenizer.TOKEN_TYPE]
+        if gguf.Keys.Tokenizer.SCORES in self.metadata_changes:
+            _, (_, scores) = self.metadata_changes[gguf.Keys.Tokenizer.SCORES]
+
+        # Open the tokenizer editor dialog
+        dialog = TokenizerEditorDialog(tokens, token_types, scores, self)
+        if dialog.exec() == QDialog.DialogCode.Accepted:
+            new_tokens, new_token_types, new_scores = dialog.get_data()
+
+            # Store changes for all three arrays
+            if tokens_field:
+                self.metadata_changes[gguf.Keys.Tokenizer.LIST] = (
+                    GGUFValueType.ARRAY,
+                    (tokens_field.types[1], new_tokens)
+                )
+
+            if token_types_field:
+                self.metadata_changes[gguf.Keys.Tokenizer.TOKEN_TYPE] = (
+                    GGUFValueType.ARRAY,
+                    (token_types_field.types[1], new_token_types)
+                )
+
+            if scores_field:
+                self.metadata_changes[gguf.Keys.Tokenizer.SCORES] = (
+                    GGUFValueType.ARRAY,
+                    (scores_field.types[1], new_scores)
+                )
+
+            self.modified = True
+
+            # Update display for all three fields
+            self.update_tokenizer_display(gguf.Keys.Tokenizer.LIST, new_tokens)
+            self.update_tokenizer_display(gguf.Keys.Tokenizer.TOKEN_TYPE, new_token_types)
+            self.update_tokenizer_display(gguf.Keys.Tokenizer.SCORES, new_scores)
+
+            self.statusBar().showMessage("Updated tokenizer data")
+
+    def update_tokenizer_display(self, key, values):
+        """Update the display of a tokenizer field in the metadata table."""
+        for row in range(self.metadata_table.rowCount()):
+            key_item = self.metadata_table.item(row, 0)
+            if key_item and key_item.text() == key:
+                value_str = f"[ {', '.join(str(v) for v in values[:5])}{', ...' if len(values) > 5 else ''} ]"
+                value_item = self.metadata_table.item(row, 2)
+                if value_item:
+                    value_item.setText(value_str)
+                break
+
+    def add_metadata(self):
+        dialog = AddMetadataDialog(self)
+        if dialog.exec() == QDialog.DialogCode.Accepted:
+            key, value_type, value = dialog.get_data()
+
+            if not key:
+                QMessageBox.warning(self, "Invalid Key", "Key cannot be empty")
+                return
+
+            # Check if key already exists
+            for row in range(self.metadata_table.rowCount()):
+                orig_item = self.metadata_table.item(row, 0)
+                if orig_item and orig_item.text() == key:
+                    QMessageBox.warning(self, "Duplicate Key", f"Key '{key}' already exists")
+                    return
+
+            # Add to table
+            row = self.metadata_table.rowCount()
+            self.metadata_table.insertRow(row)
+
+            # Key
+            key_item = QTableWidgetItem(key)
+            key_item.setFlags(key_item.flags() & ~Qt.ItemFlag.ItemIsEditable)
+            self.metadata_table.setItem(row, 0, key_item)
+
+            # Type
+            type_item = QTableWidgetItem(value_type.name)
+            type_item.setFlags(type_item.flags() & ~Qt.ItemFlag.ItemIsEditable)
+            self.metadata_table.setItem(row, 1, type_item)
+
+            # Value
+            value_item = QTableWidgetItem(str(value))
+            value_item.setFlags(value_item.flags() | Qt.ItemFlag.ItemIsEditable)
+            self.metadata_table.setItem(row, 2, value_item)
+
+            # Actions
+            actions_widget = QWidget()
+            actions_layout = QHBoxLayout(actions_widget)
+            actions_layout.setContentsMargins(2, 2, 2, 2)
+
+            remove_button = QPushButton("Remove")
+            remove_button.setProperty("row", row)
+            remove_button.setProperty("key", key)
+            remove_button.clicked.connect(self.remove_metadata)
+            actions_layout.addWidget(remove_button)
+
+            self.metadata_table.setCellWidget(row, 3, actions_widget)
+
+            # Store the change
+            self.metadata_changes[key] = (value_type, value)
+            self.modified = True
+
+            self.statusBar().showMessage(f"Added new metadata key {key}")
+
+    def save_file(self):
+        if not self.reader:
+            QMessageBox.warning(self, "No File Open", "Please open a GGUF file first")
+            return
+
+        if not self.modified and not self.metadata_changes and not self.metadata_to_remove:
+            QMessageBox.information(self, "No Changes", "No changes to save")
+            return
+
+        file_path, _ = QFileDialog.getSaveFileName(
+            self, "Save GGUF File As", "", "GGUF Files (*.gguf);;All Files (*)"
+        )
+
+        if not file_path:
+            return
+
+        try:
+            self.statusBar().showMessage(f"Saving to {file_path}...")
+            QApplication.processEvents()
+
+            # Get architecture and endianness from the original file
+            arch = 'unknown'
+            field = self.reader.get_field(gguf.Keys.General.ARCHITECTURE)
+            if field:
+                arch = field.contents()
+
+            # Create writer
+            writer = GGUFWriter(file_path, arch=arch, endianess=self.reader.endianess)
+
+            # Get alignment if present
+            alignment = None
+            field = self.reader.get_field(gguf.Keys.General.ALIGNMENT)
+            if field:
+                alignment = field.contents()
+                if alignment is not None:
+                    writer.data_alignment = alignment
+
+            # Copy metadata with changes
+            for field in self.reader.fields.values():
+                # Skip virtual fields and fields written by GGUFWriter
+                if field.name == gguf.Keys.General.ARCHITECTURE or field.name.startswith('GGUF.'):
+                    continue
+
+                # Skip fields marked for removal
+                if field.name in self.metadata_to_remove:
+                    continue
+
+                # Apply changes if any
+                if field.name in self.metadata_changes:
+                    value_type, value = self.metadata_changes[field.name]
+                    if value_type == GGUFValueType.ARRAY:
+                        # Handle array values
+                        element_type, array_values = value
+                        writer.add_array(field.name, array_values)
+                    else:
+                        writer.add_key_value(field.name, value, value_type)
+                else:
+                    # Copy original value
+                    value = field.contents()
+                    if value is not None and field.types:
+                        writer.add_key_value(field.name, value, field.types[0])
+
+            # Add new metadata
+            for key, (value_type, value) in self.metadata_changes.items():
+                # Skip if the key already existed (we handled it above)
+                if self.reader.get_field(key) is not None:
+                    continue
+
+                writer.add_key_value(key, value, value_type)
+
+            # Add tensors (including data)
+            for tensor in self.reader.tensors:
+                writer.add_tensor(tensor.name, tensor.data, raw_shape=tensor.data.shape, raw_dtype=tensor.tensor_type)
+
+            # Write header and metadata
+            writer.open_output_file(Path(file_path))
+            writer.write_header_to_file()
+            writer.write_kv_data_to_file()
+
+            # Write tensor data using the optimized method
+            writer.write_tensors_to_file(progress=False)
+
+            writer.close()
+
+            self.statusBar().showMessage(f"Saved to {file_path}")
+
+            # Ask if user wants to open the new file
+            reply = QMessageBox.question(
+                self, "Open Saved File",
+                "Would you like to open the newly saved file?",
+                QMessageBox.StandardButton.Yes | QMessageBox.StandardButton.No, QMessageBox.StandardButton.Yes
+            )
+
+            if reply == QMessageBox.StandardButton.Yes:
+                self.reader = GGUFReader(file_path, 'r')
+                self.current_file = file_path
+                self.file_path_edit.setText(file_path)
+
+                self.load_metadata()
+                self.load_tensors()
+
+                self.metadata_changes = {}
+                self.metadata_to_remove = set()
+                self.modified = False
+
+        except Exception as e:
+            QMessageBox.critical(self, "Error", f"Failed to save file: {str(e)}")
+            self.statusBar().showMessage("Error saving file")
+
+
+def main() -> None:
+    parser = argparse.ArgumentParser(description="GUI GGUF Editor")
+    parser.add_argument("model_path", nargs="?", help="path to GGUF model file to load at startup")
+    parser.add_argument("--verbose", action="store_true", help="increase output verbosity")
+
+    args = parser.parse_args()
+
+    logging.basicConfig(level=logging.DEBUG if args.verbose else logging.INFO)
+
+    app = QApplication(sys.argv)
+    window = GGUFEditorWindow()
+    window.show()
+
+    # Load model if specified
+    if args.model_path:
+        if os.path.isfile(args.model_path) and args.model_path.endswith('.gguf'):
+            window.load_file(args.model_path)
+        else:
+            logger.error(f"Invalid model path: {args.model_path}")
+            QMessageBox.warning(
+                window,
+                "Invalid Model Path",
+                f"The specified file does not exist or is not a GGUF file: {args.model_path}")
+
+    sys.exit(app.exec())
+
+
+if __name__ == '__main__':
+    main()
diff --git a/gguf-py/pyproject.toml b/gguf-py/pyproject.toml
index d214e8720..9745d3ea8 100644
--- a/gguf-py/pyproject.toml
+++ b/gguf-py/pyproject.toml
@@ -1,6 +1,6 @@
 [tool.poetry]
 name = "gguf"
-version = "0.16.0"
+version = "0.16.1"
 description = "Read and write ML models in GGUF for GGML"
 authors = ["GGML <ggml@ggml.ai>"]
 packages = [
@@ -23,10 +23,14 @@ numpy = ">=1.17"
 tqdm = ">=4.27"
 pyyaml = ">=5.1"
 sentencepiece = ">=0.1.98,<=0.2.0"
+PySide6 = { version = "^6.9", optional = true }
 
 [tool.poetry.dev-dependencies]
 pytest = "^5.2"
 
+[tool.poetry.extras]
+gui = ["PySide6"]
+
 [build-system]
 requires = ["poetry-core>=1.0.0"]
 build-backend = "poetry.core.masonry.api"
@@ -36,3 +40,4 @@ gguf-convert-endian = "gguf.scripts:gguf_convert_endian_entrypoint"
 gguf-dump = "gguf.scripts:gguf_dump_entrypoint"
 gguf-set-metadata = "gguf.scripts:gguf_set_metadata_entrypoint"
 gguf-new-metadata = "gguf.scripts:gguf_new_metadata_entrypoint"
+gguf-editor-gui = "gguf.scripts:gguf_editor_gui_entrypoint"
diff --git a/requirements/requirements-all.txt b/requirements/requirements-all.txt
index 439db8886..eba0a59f6 100644
--- a/requirements/requirements-all.txt
+++ b/requirements/requirements-all.txt
@@ -11,3 +11,5 @@
 -r ./requirements-convert_legacy_llama.txt
 -r ./requirements-convert_llama_ggml_to_gguf.txt
 -r ./requirements-tool_bench.txt
+
+-r ./requirements-gguf_editor_gui.txt
diff --git a/requirements/requirements-gguf_editor_gui.txt b/requirements/requirements-gguf_editor_gui.txt
new file mode 100644
index 000000000..920dc7cf9
--- /dev/null
+++ b/requirements/requirements-gguf_editor_gui.txt
@@ -0,0 +1,3 @@
+numpy~=1.26.4
+PySide6~=6.9.0
+gguf>=0.16.0

From 6408210082cc0a61b992b487be7e2ff2efbb9e36 Mon Sep 17 00:00:00 2001
From: Daniel Tang <danielzgtg.opensource@gmail.com>
Date: Fri, 18 Apr 2025 16:02:55 -0400
Subject: [PATCH 6/6] main : Fix Ctrl+D/newline handling (#12951)

This restores the behavior from #491. This does not affect Ctrl+D's ability to
terminate --multiline-input lines (#1040).

This also actually implements #587: "If the user wants the text to end in a
newline, this should be accomplished by explicitly adding a newline by using
\ followed by return, then returning control by pressing return again."

Fixes #12949
---
 examples/main/main.cpp | 21 ++++++++++++++++-----
 1 file changed, 16 insertions(+), 5 deletions(-)

diff --git a/examples/main/main.cpp b/examples/main/main.cpp
index fd7410a64..c59b941bf 100644
--- a/examples/main/main.cpp
+++ b/examples/main/main.cpp
@@ -865,9 +865,22 @@ int main(int argc, char ** argv) {
                 console::set_display(console::reset);
                 display = true;
 
-                // Add tokens to embd only if the input buffer is non-empty
-                // Entering a empty line lets the user pass control back
-                if (buffer.length() > 1) {
+                if (buffer.empty()) { // Ctrl+D on empty line exits
+                    LOG("EOF by user\n");
+                    break;
+                }
+
+                if (buffer.back() == '\n') {
+                    // Implement #587:
+                    // If the user wants the text to end in a newline,
+                    // this should be accomplished by explicitly adding a newline by using \ followed by return,
+                    // then returning control by pressing return again.
+                    buffer.pop_back();
+                }
+
+                if (buffer.empty()) { // Enter key on empty line lets the user pass control back
+                    LOG_DBG("empty line, passing control back\n");
+                } else { // Add tokens to embd only if the input buffer is non-empty
                     // append input suffix if any
                     if (!params.input_suffix.empty() && !params.conversation_mode) {
                         LOG_DBG("appending input suffix: '%s'\n", params.input_suffix.c_str());
@@ -915,8 +928,6 @@ int main(int argc, char ** argv) {
 
                     n_remain -= line_inp.size();
                     LOG_DBG("n_remain: %d\n", n_remain);
-                } else {
-                    LOG_DBG("empty line, passing control back\n");
                 }
 
                 input_echo = false; // do not echo this again