rpc: allow -sm tensor

2026-08-30 20:17:36 +02:00 · 2026-08-31 00:02:46 +08:00
18 changed files with 1175 additions and 2023 deletions
@@ -6,7 +6,7 @@
 extern "C" {
 #endif

-#define RPC_PROTO_MAJOR_VERSION    6
+#define RPC_PROTO_MAJOR_VERSION    7
 #define RPC_PROTO_MINOR_VERSION    0
 #define RPC_PROTO_PATCH_VERSION    0

@@ -865,7 +865,12 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
        ggml_backend_meta_split_state split_state;
        switch (tensor->op) {
            case GGML_OP_NONE: {
-                split_state = {GGML_BACKEND_SPLIT_AXIS_MIRRORED, {0}, {1}, 1};
+                if (tensor->view_src != nullptr) {
+                    // full-tensor view created with ggml_view_tensor, transparent for the split state
+                    split_state = ggml_backend_meta_get_split_state(stc, tensor->view_src, assume_sync);
+                } else {
+                    split_state = {GGML_BACKEND_SPLIT_AXIS_MIRRORED, {0}, {1}, 1};
+                }
            } break;
            case GGML_OP_DUP: {
                split_state = handle_generic(src_ss, /*scalar_only =*/ true);
@@ -2283,6 +2288,14 @@ static enum ggml_status ggml_backend_meta_graph_compute(ggml_backend_t backend,
                cgraph_ij->uid = ggml_graph_next_uid();
            }
        }
+
+        // Aux graph contents are rewritten on every compute but are identical across calls while the subgraphs are reused,
+        // so they can get stable uids on rebuild. Only safe without a comm backend, where the fallback usage is deterministic.
+        if (backend_ctx->comm_ctx == nullptr) {
+            for (ggml_cgraph * cgraph_aux : backend_ctx->cgraphs_aux) {
+                cgraph_aux->uid = ggml_graph_next_uid();
+            }
+        }
    }

    size_t iga = 0; // i graph aux
@@ -69,15 +69,30 @@ using u32vec  = std::vector<uint32_t>;
 #define GGML_HEXAGON_FENCE_SLOT_SIZE       128

 struct ggml_hexagon_device_config {
-    int         physical_idx = 0;
-    int         virtual_idx  = 0;
-    int         domain_id    = 0;
-    std::string domain_name;
+    int physical_idx = 0;
+    int virtual_idx  = 0;
    std::string name;
 };

 static ggml_hexagon_device_config opt_device_configs[GGML_HEXAGON_MAX_SESSIONS];

+static int get_domain_id(int physical_idx) {
+    switch (physical_idx) {
+        case 0:  return 3;  // CDSP0 (all devices)
+        case 1:  return 4;  // CDSP1 (IQ9, IQ10)
+        case 2:  return 18; // CDSP2 (IQ10)
+        case 3:  return 19; // CDSP3 (IQ10)
+        default: return CDSP_DOMAIN_ID + physical_idx;
+    }
+}
+
+static std::string get_domain_name(int physical_idx) {
+    if (physical_idx == 0) {
+        return CDSP_DOMAIN_NAME;
+    }
+    return std::string("cdsp") + std::to_string(physical_idx);
+}
+
 static int    opt_arch    = 0; // autodetect
 static size_t opt_ndev    = 1;
 static size_t opt_nhvx    = 0; // use all
@@ -346,6 +361,7 @@ struct ggml_hexagon_session {
    uint32_t         session_id;
    uint32_t         domain_id;
    uint64_t         queue_id;
+    int              dev_id;
    int              phys_idx;
    int              virt_idx;
    bool             valid_session;
@@ -360,6 +376,9 @@ struct ggml_hexagon_session {
    std::unordered_map<int, std::unique_ptr<ggml_hexagon_shared_buffer>> cloned_buffers;
    std::unordered_set<ggml_hexagon_session *>                           sync_peers;

+    ggml_backend_buffer_type buffer_type        = {};
+    ggml_backend_buffer_type host_buffer_type   = {};
+
    uint32_t n_threads   = 0;
    uint32_t n_hvx       = 0;
    uint32_t n_hmx       = 0;
@@ -373,12 +392,12 @@ struct ggml_hexagon_session {

    mutable std::unordered_set<const ggml_tensor *> needs_repack;

-    ggml_hexagon_session(const ggml_hexagon_device_config & config, ggml_backend_dev_t dev = nullptr) noexcept(false);
+    ggml_hexagon_session(int dev_id, ggml_backend_dev_t dev) noexcept(false);
    ~ggml_hexagon_session() noexcept(true);

    const char* c_name() const { return name.c_str(); }

-    void allocate(const ggml_hexagon_device_config & config) noexcept(false);
+    void allocate(int dev_id) noexcept(false);
    void release() noexcept(true);

    void enqueue_op(const htp_opnode & node);
@@ -411,38 +430,14 @@ struct ggml_hexagon_session {

 // ** backend buffers

-struct ggml_backend_hexagon_device_context {
-    int                        dev_id;
-    ggml_hexagon_device_config config;
-    ggml_backend_dev_t         dev = nullptr;
-    size_t                     max_bufsize = 0;
-
-    ggml_backend_buffer_type buffer_type      = {};
-    ggml_backend_buffer_type host_buffer_type = {};
-
-    std::unique_ptr<ggml_hexagon_session> sess;
-
-    ggml_backend_hexagon_device_context(int dev_id, const ggml_hexagon_device_config & config, ggml_backend_dev_t dev);
-    ~ggml_backend_hexagon_device_context();
-
-    const char * c_name() const { return config.name.c_str(); }
-
-    ggml_hexagon_session * session() {
-        if (!sess) {
-            sess = std::make_unique<ggml_hexagon_session>(config, dev);
-        }
-        return sess.get();
-    }
-};
-
 struct ggml_backend_hexagon_buffer_type_context {
-    ggml_backend_hexagon_buffer_type_context(const std::string & name, ggml_backend_hexagon_device_context * dev_ctx) {
-        this->dev_ctx = dev_ctx;
-        this->name    = name;
+    ggml_backend_hexagon_buffer_type_context(const std::string & name, ggml_hexagon_session * sess) {
+        this->sess = sess;
+        this->name = name;
    }

-    ggml_backend_hexagon_device_context * dev_ctx;
-    std::string                           name;
+    ggml_hexagon_session * sess;
+    std::string            name;
 };

 struct ggml_hexagon_rpcmem_block {
@@ -581,8 +576,7 @@ struct ggml_hexagon_shared_buffer {
 };

 static ggml_hexagon_session * ggml_backend_hexagon_buffer_get_sess(ggml_backend_buffer_t buffer) {
-    auto sbuf = static_cast<ggml_hexagon_shared_buffer *>(buffer->context);
-    return sbuf->sess;
+    return static_cast<ggml_backend_hexagon_buffer_type_context *>(buffer->buft->context)->sess;
 }

 static void ggml_backend_hexagon_buffer_free_buffer(ggml_backend_buffer_t buffer) {
@@ -1500,26 +1494,24 @@ static const char * ggml_backend_hexagon_buffer_type_name(ggml_backend_buffer_ty

 static ggml_backend_buffer_t ggml_backend_hexagon_buffer_type_alloc_buffer(
            ggml_backend_buffer_type_t buffer_type, size_t size) {
-    auto dev_ctx = static_cast<ggml_backend_hexagon_buffer_type_context *>(buffer_type->context)->dev_ctx;
-    auto sess    = dev_ctx->session();
+    auto sess = static_cast<ggml_backend_hexagon_buffer_type_context *>(buffer_type->context)->sess;
    try {
        ggml_hexagon_shared_buffer * sbuf = new ggml_hexagon_shared_buffer(sess, size, false, GGML_HEXAGON_FENCE_BUFFER_SIZE);
        return ggml_backend_buffer_init(buffer_type, ggml_backend_hexagon_buffer_interface, sbuf, size);
    } catch (const std::exception & exc) {
-        GGML_LOG_ERROR("ggml-hex: %s failed to allocate device buffer context: %s\n", dev_ctx->c_name(), exc.what());
+        GGML_LOG_ERROR("ggml-hex: %s failed to allocate device buffer context: %s\n", sess->c_name(), exc.what());
        return nullptr;
    }
 }

 static ggml_backend_buffer_t ggml_backend_hexagon_host_buffer_type_alloc_buffer(
            ggml_backend_buffer_type_t buffer_type, size_t size) {
-    auto dev_ctx = static_cast<ggml_backend_hexagon_buffer_type_context *>(buffer_type->context)->dev_ctx;
-    auto sess    = dev_ctx->session();
+    auto sess = static_cast<ggml_backend_hexagon_buffer_type_context *>(buffer_type->context)->sess;
    try {
        ggml_hexagon_shared_buffer * sbuf = new ggml_hexagon_shared_buffer(sess, size, false, GGML_HEXAGON_FENCE_BUFFER_SIZE);
        return ggml_backend_buffer_init(buffer_type, ggml_backend_hexagon_host_buffer_interface, sbuf, size);
    } catch (const std::exception & exc) {
-        GGML_LOG_ERROR("ggml-hex: %s failed to allocate host buffer context: %s\n", dev_ctx->c_name(), exc.what());
+        GGML_LOG_ERROR("ggml-hex: %s failed to allocate host buffer context: %s\n", sess->c_name(), exc.what());
        return nullptr;
    }
 }
@@ -1544,7 +1536,7 @@ static size_t ggml_backend_hexagon_buffer_type_get_alloc_size(ggml_backend_buffe

 static size_t ggml_backend_hexagon_buffer_type_get_max_size(ggml_backend_buffer_type_t buft) {
    auto * context = static_cast<ggml_backend_hexagon_buffer_type_context *>(buft->context);
-    return context->dev_ctx->max_bufsize;
+    return context->sess->max_bufsize;
 }

 static bool ggml_backend_hexagon_buffer_type_is_host(ggml_backend_buffer_type_t buft) {
@@ -1575,22 +1567,6 @@ static ggml_backend_buffer_type_i ggml_backend_hexagon_host_buffer_type_interfac
    /* .is_host          = */ ggml_backend_hexagon_host_buffer_type_is_host,
 };

-ggml_backend_hexagon_device_context::ggml_backend_hexagon_device_context(int dev_id, const ggml_hexagon_device_config & config, ggml_backend_dev_t dev)
-    : dev_id(dev_id), config(config), dev(dev), max_bufsize(opt_mbuf) {
-    buffer_type.device  = dev;
-    buffer_type.iface   = ggml_backend_hexagon_buffer_type_interface;
-    buffer_type.context = new ggml_backend_hexagon_buffer_type_context(config.name, this);
-
-    host_buffer_type.device  = dev;
-    host_buffer_type.iface   = ggml_backend_hexagon_host_buffer_type_interface;
-    host_buffer_type.context = new ggml_backend_hexagon_buffer_type_context(config.name + "-HOST", this);
-}
-
-ggml_backend_hexagon_device_context::~ggml_backend_hexagon_device_context() {
-    delete static_cast<ggml_backend_hexagon_buffer_type_context *>(buffer_type.context);
-    delete static_cast<ggml_backend_hexagon_buffer_type_context *>(host_buffer_type.context);
-}
-
 static bool ggml_backend_buffer_is_hexagon(const struct ggml_backend_buffer * b) {
    return b->buft->iface.get_alignment == ggml_backend_hexagon_buffer_type_get_alignment;
 }
@@ -2835,7 +2811,8 @@ static size_t ggml_hexagon_measure_max_vmem(ggml_hexagon_session *sess) {
    return vmem - step; // backoff to account for overhead from internal mappings
 }

-void ggml_hexagon_session::allocate(const ggml_hexagon_device_config & config) noexcept(false) {
+void ggml_hexagon_session::allocate(int dev_id) noexcept(false) {
+    const auto & config = opt_device_configs[dev_id];
    int phys_idx = config.physical_idx;
    int virt_idx = config.virtual_idx;

@@ -2846,31 +2823,21 @@ void ggml_hexagon_session::allocate(const ggml_hexagon_device_config & config) n

    this->phys_idx   = phys_idx;
    this->virt_idx   = virt_idx;
-    this->domain_id  = config.domain_id;
+    this->domain_id  = get_domain_id(phys_idx);
    this->session_id = 0;
+    this->dev_id     = dev_id;
    this->name       = config.name;
    this->op_pending = 0;

    GGML_LOG_DEBUG("ggml-hex: %s allocating new session\n", this->name.c_str());

-    if (config.domain_id < 0 || config.domain_name.empty()) {
-        GGML_LOG_ERROR("ggml-hex: %s: invalid physical CDSP core %d\n", config.name.c_str(), config.physical_idx);
-        throw std::runtime_error("ggml-hex: invalid physical CDSP core");
+    domain * my_domain = htpdrv_get_domain(this->domain_id);
+    if (my_domain == NULL) {
+        GGML_LOG_ERROR("ggml-hex: unable to get domain struct for CDSP (domain_id %d)\n", this->domain_id);
+        throw std::runtime_error("ggml-hex: failed to get CDSP domain (see log for details)");
    }

-    const std::string & dom_name = config.domain_name;
-
-    // Enable Unsigned PD for all domains
-    {
-        struct remote_rpc_control_unsigned_module u;
-        u.domain = -1;
-        u.enable = 1;
-        int err  = remote_session_control(DSPRPC_CONTROL_UNSIGNED_MODULE, (void *) &u, sizeof(u));
-        if (err != AEE_SUCCESS) {
-            GGML_LOG_ERROR("ggml-hex: %s failed to enable unsigned PD : error 0x%x\n", this->c_name(), err);
-            throw std::runtime_error("ggml-hex: remote_session_control(unsign) failed (see log for details)");
-        }
-    }
+    std::string dom_name = get_domain_name(phys_idx);

    // Create new session if virtual_idx > 0
    if (virt_idx > 0) {
@@ -2882,8 +2849,7 @@ void ggml_hexagon_session::allocate(const ggml_hexagon_device_config & config) n

        int err = remote_session_control(FASTRPC_RESERVE_NEW_SESSION, (void *) &n, sizeof(n));
        if (err != AEE_SUCCESS) {
-            GGML_LOG_ERROR("ggml-hex: %s failed to reserve new session (physical %d, virtual %d) : error 0x%x\n",
-                           this->c_name(), phys_idx, virt_idx, err);
+            GGML_LOG_ERROR("ggml-hex: failed to reserve new session %d (physical %d, virtual %d) : error 0x%x\n", dev_id, phys_idx, virt_idx, err);
            throw std::runtime_error("ggml-hex: remote_session_control(new-sess) failed (see log for details)");
        }

@@ -2891,21 +2857,10 @@ void ggml_hexagon_session::allocate(const ggml_hexagon_device_config & config) n
        this->session_id    = n.session_id;
        this->domain_id     = n.effective_domain_id;
        this->valid_session = true;
-    } else {
-        struct remote_rpc_effective_domain_id eff = {};
-        eff.domain_name     = const_cast<char *>(dom_name.c_str());
-        eff.domain_name_len = dom_name.size();
-        eff.session_id      = 0;
-
-        int err = remote_session_control(FASTRPC_GET_EFFECTIVE_DOMAIN_ID, (void *) &eff, sizeof(eff));
-        if (err == AEE_SUCCESS) {
-            this->domain_id = eff.effective_domain_id;
-        } else {
-            GGML_LOG_DEBUG("ggml-hex: %s FASTRPC_GET_EFFECTIVE_DOMAIN_ID returned 0x%x, using domain_id %d\n",
-                           this->name.c_str(), err, this->domain_id);
-        }
    }

+    // Get session URI
+
    char session_uri[256];
    {
        char htp_uri[256];
@@ -2922,18 +2877,31 @@ void ggml_hexagon_session::allocate(const ggml_hexagon_device_config & config) n

        int err = remote_session_control(FASTRPC_GET_URI, (void *) &u, sizeof(u));
        if (err != AEE_SUCCESS) {
-            snprintf(session_uri, sizeof(session_uri), "%s&_dom=%s&_session=%u",
-                     htp_uri, dom_name.c_str(), this->session_id);
+            // fallback to single session uris
+            int htp_URI_domain_len = strlen(htp_uri) + MAX_DOMAIN_NAMELEN;

-            GGML_LOG_WARN("ggml-hex: %s failed to get URI (physical %d, virtual %d) : error 0x%x. Falling back to single session URI: %s\n",
-                          this->c_name(), phys_idx, virt_idx, err, session_uri);
+            snprintf(session_uri, htp_URI_domain_len, "%s%s", htp_uri, my_domain->uri);
+
+            GGML_LOG_WARN("ggml-hex: failed to get URI for session %d (physical %d, virtual %d) : error 0x%x. Falling back to single session URI: %s\n", dev_id, phys_idx, virt_idx, err, session_uri);
+        }
+    }
+
+    // Enable Unsigned PD
+    {
+        struct remote_rpc_control_unsigned_module u;
+        u.domain = this->domain_id;
+        u.enable = 1;
+        int err  = remote_session_control(DSPRPC_CONTROL_UNSIGNED_MODULE, (void *) &u, sizeof(u));
+        if (err != AEE_SUCCESS) {
+            GGML_LOG_ERROR("ggml-hex: failed to enable unsigned PD for session %d : error 0x%x\n", dev_id, err);
+            throw std::runtime_error("ggml-hex: remote_session_control(unsign) failed (see log for details)");
        }
    }

    // Open session
    int err = htp_iface_open(session_uri, &this->handle);
    if (err != AEE_SUCCESS) {
-        GGML_LOG_ERROR("ggml-hex: %s failed to open session : error 0x%x\n", this->c_name(), err);
+        GGML_LOG_ERROR("ggml-hex: failed to open session %d : error 0x%x\n", dev_id, err);
        throw std::runtime_error("ggml-hex: failed to open session (see log for details)");
    }

@@ -3023,7 +2991,7 @@ void ggml_hexagon_session::allocate(const ggml_hexagon_device_config & config) n
    this->op_batch = new ggml_hexagon_opbatch(this, opt_opbatch, this->max_vmem);

    // Start dspqueue/opbatch processing
-    err = htp_iface_start(this->handle, this->session_id, this->queue_id, opt_nhvx, opt_nhmx, this->max_vmem);
+    err = htp_iface_start(this->handle, dev_id, this->queue_id, opt_nhvx, opt_nhmx, this->max_vmem);
    if (err != 0) {
        GGML_LOG_ERROR("ggml-hex: %s failed to start session: 0x%08x\n", this->c_name(), (unsigned) err);
        throw std::runtime_error("ggml-hex: iface start failed (see log for details)");
@@ -3086,23 +3054,33 @@ void ggml_hexagon_session::release() noexcept(true) {
    this->cloned_buffers.clear();
 }

-ggml_hexagon_session::ggml_hexagon_session(const ggml_hexagon_device_config & config, ggml_backend_dev_t dev) noexcept(false) {
+ggml_hexagon_session::ggml_hexagon_session(int dev_id, ggml_backend_dev_t dev) noexcept(false) {
+    buffer_type.device      = dev;
+    host_buffer_type.device = dev;
+
    op_batch = nullptr;
    op_queue = nullptr;
    fence_seq = ((uintptr_t)this) & 0xFFFF;

    try {
-        allocate(config);
+        allocate(dev_id);
+
+        buffer_type.iface   = ggml_backend_hexagon_buffer_type_interface;
+        buffer_type.context = new ggml_backend_hexagon_buffer_type_context(this->name, this);
+
+        host_buffer_type.iface   = ggml_backend_hexagon_host_buffer_type_interface;
+        host_buffer_type.context = new ggml_backend_hexagon_buffer_type_context(this->name + "-HOST", this);
    } catch (const std::exception & exc) {
        release();
        throw;
    }
-
-    GGML_UNUSED(dev);
 }

 ggml_hexagon_session::~ggml_hexagon_session() noexcept(true) {
    release();
+
+    delete static_cast<ggml_backend_hexagon_buffer_type_context *>(buffer_type.context);
+    delete static_cast<ggml_backend_hexagon_buffer_type_context *>(host_buffer_type.context);
 }

 // ** backend interface
@@ -3979,13 +3957,11 @@ static void ggml_hexagon_precompute_fused_mmnx_params(
 }

 static bool ggml_hexagon_tensor_is_host(const struct ggml_hexagon_session * sess, const struct ggml_tensor * t) {
-    return t && t->buffer && ggml_backend_buft_is_host(t->buffer->buft);
-    GGML_UNUSED(sess);
+    return t && t->buffer && t->buffer->buft == &sess->host_buffer_type;
 }

 static bool ggml_hexagon_tensor_is_non_host(const struct ggml_hexagon_session * sess, const struct ggml_tensor * t) {
-    return t && t->buffer && !ggml_backend_buft_is_host(t->buffer->buft);
-    GGML_UNUSED(sess);
+    return t && t->buffer && t->buffer->buft != &sess->host_buffer_type;
 }

 static bool ggml_hexagon_supported_mul_mat(const struct ggml_hexagon_session * sess, const struct ggml_tensor * dst) {
@@ -5293,8 +5269,7 @@ bool ggml_backend_is_hexagon(ggml_backend_t backend) {
 // device interface

 static ggml_backend_t ggml_backend_hexagon_device_init(ggml_backend_dev_t dev, const char * params) {
-    auto dev_ctx = static_cast<ggml_backend_hexagon_device_context *>(dev->context);
-    auto sess    = dev_ctx->session();
+    auto sess = static_cast<ggml_hexagon_session *>(dev->context);

    return new ggml_backend{
        /* .guid      = */ ggml_backend_hexagon_guid(),
@@ -5307,8 +5282,8 @@ static ggml_backend_t ggml_backend_hexagon_device_init(ggml_backend_dev_t dev, c
 }

 static const char * ggml_backend_hexagon_device_get_name(ggml_backend_dev_t dev) {
-    auto dev_ctx = static_cast<ggml_backend_hexagon_device_context *>(dev->context);
-    return dev_ctx->c_name();
+    auto sess = static_cast<ggml_hexagon_session *>(dev->context);
+    return sess->c_name();

    GGML_UNUSED(dev);
 }
@@ -5346,16 +5321,16 @@ static void ggml_backend_hexagon_device_get_props(ggml_backend_dev_t dev, struct
 }

 static ggml_backend_buffer_type_t ggml_backend_hexagon_device_get_buffer_type(ggml_backend_dev_t dev) {
-    auto dev_ctx = static_cast<ggml_backend_hexagon_device_context *>(dev->context);
-    return &dev_ctx->buffer_type;
+    auto sess = static_cast<ggml_hexagon_session *>(dev->context);
+    return &sess->buffer_type;
 }

 static ggml_backend_buffer_type_t ggml_backend_hexagon_device_get_host_buffer_type(ggml_backend_dev_t dev) {
    if (!opt_hostbuf) {
        return NULL;
    }
-    auto dev_ctx = static_cast<ggml_backend_hexagon_device_context *>(dev->context);
-    return &dev_ctx->host_buffer_type;
+    auto sess = static_cast<ggml_hexagon_session *>(dev->context);
+    return &sess->host_buffer_type;
 }

 static bool ggml_hexagon_supported_cpy(const struct ggml_hexagon_session * sess, const struct ggml_tensor * op) {
@@ -5446,8 +5421,7 @@ static bool ggml_hexagon_supported_fill(const struct ggml_hexagon_session * sess
 }

 static bool ggml_backend_hexagon_device_supports_op(ggml_backend_dev_t dev, const struct ggml_tensor * op) {
-    auto dev_ctx = static_cast<ggml_backend_hexagon_device_context *>(dev->context);
-    auto sess    = dev_ctx->session();
+    auto sess = static_cast<ggml_hexagon_session *>(dev->context);

    // reject ops that match the filter
    if (opt_opfilter && std::regex_match(ggml_op_desc(op), *opt_opfilter)) {
@@ -5519,7 +5493,6 @@ static bool ggml_backend_hexagon_device_supports_op(ggml_backend_dev_t dev, cons
                    supp = ggml_hexagon_supported_unary(sess, op);
                    break;
                default:
-                    supp = false;
                    break;
            }
            break;
@@ -5532,7 +5505,6 @@ static bool ggml_backend_hexagon_device_supports_op(ggml_backend_dev_t dev, cons
                    supp = ggml_hexagon_supported_activations(sess, op);
                    break;
                default:
-                    supp = false;
                    break;
            }
            break;
@@ -5618,17 +5590,17 @@ static bool ggml_backend_hexagon_device_supports_op(ggml_backend_dev_t dev, cons
 }

 static bool ggml_backend_hexagon_device_supports_buft(ggml_backend_dev_t dev, ggml_backend_buffer_type_t buft) {
-    auto dev_ctx = static_cast<ggml_backend_hexagon_device_context *>(dev->context);
+    auto sess = static_cast<ggml_hexagon_session *>(dev->context);

    // Technically we can clone hexagon buffers from any session but for some reason the output is garbled with layer-split,
    // tensor-split works correctly, so it needs mode debugging and investigation. For now accept only our own buffers.
 #if 0
    bool supp = (buft->iface.get_alignment == ggml_backend_hexagon_buffer_type_get_alignment);
 #else
-    bool supp = (buft == &dev_ctx->host_buffer_type) || (buft == &dev_ctx->buffer_type);
+    bool supp = (buft == &sess->host_buffer_type) || (buft == &sess->buffer_type);
 #endif

-    HEX_VERBOSE("ggml-hex: %s device-supports-buft %s %s\n", dev_ctx->c_name(), ggml_backend_buft_name(buft), supp ? "yes" : "no");
+    HEX_VERBOSE("ggml-hex: %s device-supports-buft %s %s\n", sess->name.c_str(), ggml_backend_buft_name(buft), supp ? "yes" : "no");
    return supp;
 }

@@ -5657,11 +5629,16 @@ ggml_hexagon_registry::ggml_hexagon_registry(ggml_backend_reg_t reg) {

    GGML_LOG_INFO("ggml-hex: Hexagon Arch version v%d\n", opt_arch);

-    // Create devices
+    // Create devices / sessions
    for (size_t i = 0; i < opt_ndev; i++) {
-        devices[i].iface   = ggml_backend_hexagon_device_i;
-        devices[i].reg     = reg;
-        devices[i].context = new ggml_backend_hexagon_device_context(i, opt_device_configs[i], &devices[i]);
+        devices[i].iface = ggml_backend_hexagon_device_i;
+        devices[i].reg   = reg;
+        try {
+            devices[i].context = new ggml_hexagon_session(i, &devices[i]);
+        } catch (const std::exception & exc) {
+            GGML_LOG_ERROR("ggml-hex: failed to create device/session %zu\n", i);
+            devices[i].context = nullptr;
+        }
    }

 }
@@ -5669,10 +5646,10 @@ ggml_hexagon_registry::ggml_hexagon_registry(ggml_backend_reg_t reg) {
 ggml_hexagon_registry::~ggml_hexagon_registry() {
    GGML_LOG_INFO("ggml-hex: releasing registry\n");

-    // Release devices
+    // Release devices / sessions
    for (size_t i = 0; i < opt_ndev; i++) {
-        auto dev_ctx = static_cast<ggml_backend_hexagon_device_context *>(devices[i].context);
-        delete dev_ctx;
+        auto sess = static_cast<ggml_hexagon_session *>(devices[i].context);
+        delete sess;
    }
 }

@@ -5841,85 +5818,6 @@ template<typename T, int BASE=10> std::string vec_to_str(std::vector<T> v) {
    return str;
 }

-// Enumerate NPU (aka CDSP) domains via FASTRPC_GET_DOMAINS if supported,
-// and populate domain_id and domain_name for all configured devices.
-static void ggml_hexagon_discover_devices() {
-    std::unordered_map<int, fastrpc_domain> cdsp_map;
-    bool discovery_supported = false;
-
-    system_req_payload domain_info = {};
-    domain_info.id              = FASTRPC_GET_DOMAINS;
-    domain_info.sys.domains     = nullptr;
-    domain_info.sys.max_domains = 0;
-    domain_info.sys.flags       = DOMAINS_LIST_FLAGS_SET_TYPE(0, FASTRPC_NSP);
-
-    int err = remote_system_request(&domain_info);
-    if (err == AEE_SUCCESS && domain_info.sys.num_domains > 0) {
-        std::vector<fastrpc_domain> domains(domain_info.sys.num_domains);
-        domain_info.sys.domains     = domains.data();
-        domain_info.sys.max_domains = (int) domains.size();
-
-        err = remote_system_request(&domain_info);
-        if (err == AEE_SUCCESS) {
-            discovery_supported = true;
-            const int n_domains = std::min(domain_info.sys.num_domains, (int) domains.size());
-            for (int i = 0; i < n_domains; i++) {
-                GGML_LOG_INFO("ggml-hex: FASTRPC_GET_DOMAINS[%d]: type %d id %d name '%s' status %d instance-id %d\n",
-                              i, (int) domains[i].type, domains[i].id, domains[i].name, domains[i].status, domains[i].instance_id);
-                if (domains[i].type != FASTRPC_NSP) {
-                    GGML_LOG_DEBUG("ggml-hex:   skipping non-CDSP domain (type=%d)\n", (int) domains[i].type);
-                    continue;
-                }
-                if (!domains[i].status) {
-                    GGML_LOG_WARN("ggml-hex:   skipping CDSP domain id=%d (status=down)\n", domains[i].id);
-                    continue;
-                }
-                cdsp_map[domains[i].instance_id] = domains[i];
-                GGML_LOG_INFO("ggml-hex: using CDSP domain: instance-id %d id %d name '%s'\n",
-                              domains[i].instance_id, domains[i].id, domains[i].name);
-            }
-        } else {
-            GGML_LOG_WARN("ggml-hex: FASTRPC_GET_DOMAINS fetch failed (0x%x), using static CDSP domains\n", (unsigned) err);
-        }
-    } else if (err != AEE_SUCCESS) {
-        GGML_LOG_DEBUG("ggml-hex: FASTRPC_GET_DOMAINS query failed (0x%x), using static CDSP domains\n", (unsigned) err);
-    }
-
-    // Populate domain IDs and names for all configured devices
-    for (size_t i = 0; i < opt_ndev; i++) {
-        auto & cfg = opt_device_configs[i];
-        if (discovery_supported) {
-            auto it = cdsp_map.find(cfg.physical_idx);
-            if (it != cdsp_map.end()) {
-                cfg.domain_id   = it->second.id;
-                cfg.domain_name = it->second.name;
-            } else {
-                GGML_LOG_ERROR("ggml-hex: physical CDSP core %d not found on device (%zu CDSP core(s) available)\n",
-                               cfg.physical_idx, cdsp_map.size());
-                cfg.domain_id   = -1;
-                cfg.domain_name = "";
-            }
-        } else {
-            switch (cfg.physical_idx) {
-                case 0:
-                    cfg.domain_id   = 3;
-                    cfg.domain_name = CDSP_DOMAIN_NAME;
-                    break;
-                case 1:
-                    cfg.domain_id   = 4;
-                    cfg.domain_name = "cdsp1";
-                    break;
-                default:
-                    GGML_LOG_ERROR("ggml-hex: physical CDSP core %d not supported without dynamic discovery\n",
-                                   cfg.physical_idx);
-                    cfg.domain_id   = -1;
-                    cfg.domain_name = "";
-                    break;
-            }
-        }
-    }
-}
-
 static void ggml_hexagon_init(ggml_backend_reg * reg) {
    // Basic sanity checks to make sure definitions match
    static_assert((unsigned int) HTP_TYPE_Q4_0 == (unsigned int) GGML_TYPE_Q4_0,
@@ -6085,9 +5983,6 @@ static void ggml_hexagon_init(ggml_backend_reg * reg) {
    }
 #endif

-    // Resolve domain info for all configured devices
-    ggml_hexagon_discover_devices();
-
    if (str_profile) {
        opt_pmu_evt = [&]() -> std::vector<uint32_t> {
            auto v  = str_to_vec<uint32_t>(str_profile);
@@ -73,7 +73,6 @@ typedef int (*remote_handle64_close_pfn_t)(remote_handle h);
 typedef int (*remote_handle_control_pfn_t)(uint32_t req, void* data, uint32_t datalen);
 typedef int (*remote_handle64_control_pfn_t)(remote_handle64 h, uint32_t req, void* data, uint32_t datalen);
 typedef int (*remote_session_control_pfn_t)(uint32_t req, void *data, uint32_t datalen);
-typedef int (*remote_system_request_pfn_t)(system_req_payload * req);

 //
 // Driver API pfns
@@ -100,7 +99,6 @@ remote_handle64_close_pfn_t   remote_handle64_close_pfn   = nullptr;
 remote_handle_control_pfn_t   remote_handle_control_pfn   = nullptr;
 remote_handle64_control_pfn_t remote_handle64_control_pfn = nullptr;
 remote_session_control_pfn_t  remote_session_control_pfn  = nullptr;
-remote_system_request_pfn_t   remote_system_request_pfn   = nullptr;

 //
 // Driver API
@@ -208,13 +206,6 @@ HTPDRV_API int remote_session_control(uint32_t req, void * data, uint32_t datale
    return remote_session_control_pfn(req, data, datalen);
 }

-HTPDRV_API int remote_system_request(system_req_payload * req) {
-    if (!remote_system_request_pfn) {
-        return AEE_EUNSUPPORTEDAPI;
-    }
-    return remote_system_request_pfn(req);
-}
-
 #ifdef _WIN32

 static std::string wstr_to_str(std::wstring_view wstr) {
@@ -376,7 +367,6 @@ int htpdrv_init() {
    dlsym(handle.get(), remote_handle64_control_pfn_t, remote_handle64_control_pfn, remote_handle64_control, false);
    dlsym(handle.get(), remote_session_control_pfn_t, remote_session_control_pfn, remote_session_control, false);
    dlsym(handle.get(), remote_handle64_close_pfn_t, remote_handle64_close_pfn, remote_handle64_close, false);
-    dlsym(handle.get(), remote_system_request_pfn_t, remote_system_request_pfn, remote_system_request, true);

    lib_cdsp_rpc_handle = std::move(handle);
    initialized         = true;
@@ -116,8 +116,6 @@ HTPDRV_API domain * htpdrv_get_domain(int domain_id);
 */
 HTPDRV_API int htpdrv_get_arch(int domain, int * arch);

-HTPDRV_API int remote_system_request(system_req_payload * req);
-
 #ifdef __cplusplus
 }
 #endif
@@ -1,27 +1,10 @@
 #include "ggml-metal-common.h"

-#include "ggml.h"
 #include "ggml-impl.h"
 #include "ggml-backend-impl.h"

 #include <vector>

-bool ggml_metal_op_mul_mat_use_mm(const struct ggml_tensor * op, bool has_simdgroup_mm) {
-    const int64_t ne00 = op->src[0]->ne[0];
-    const int64_t ne11 = op->src[1]->ne[1];
-
-    return !ggml_is_transposed(op->src[0]) &&
-           !ggml_is_transposed(op->src[1]) &&
-           has_simdgroup_mm && ne00 >= 64 && ne11 > 8;
-}
-
-bool ggml_metal_op_mul_mat_id_use_mm(const struct ggml_tensor * op, bool has_simdgroup_mm) {
-    const int64_t ne00 = op->src[0]->ne[0];
-    const int64_t ne21 = op->src[2]->ne[1];
-
-    return has_simdgroup_mm && ne00 >= 64 && ne21 >= 32;
-}
-
 // represents a memory range (i.e. an interval from a starting address p0 to an ending address p1 in a given buffer pb)
 // the type indicates whether it is a source range (i.e. ops read data from it) or a destination range (i.e. ops write data to it)
 struct ggml_mem_range {
@@ -47,10 +47,6 @@ bool ggml_mem_ranges_check(ggml_mem_ranges_t mrs, const struct ggml_tensor * ten
 //       if it proves to work well, we can start using it for other backends in the future
 void ggml_graph_optimize(struct ggml_cgraph * gf);

-// mat-mat vs mat-vec dispatch; used by both supports_op and ggml_metal_op_mul_mat*
-bool ggml_metal_op_mul_mat_use_mm   (const struct ggml_tensor * op, bool has_simdgroup_mm);
-bool ggml_metal_op_mul_mat_id_use_mm(const struct ggml_tensor * op, bool has_simdgroup_mm);
-
 #ifdef __cplusplus
 }
 #endif
@@ -3,7 +3,6 @@
 #import "ggml-impl.h"
 #import "ggml-backend-impl.h"
 #import "ggml-metal-impl.h"
-#import "ggml-metal-common.h"

 #include <Foundation/Foundation.h>

@@ -789,10 +788,6 @@ void ggml_metal_encoder_debug_group_pop (ggml_metal_encoder_t encoder) {
 }

 void ggml_metal_encoder_set_pipeline(ggml_metal_encoder_t encoder, struct ggml_metal_pipeline_with_params pipeline) {
-    if (!pipeline.pipeline) {
-        GGML_ABORT("%s: nil Metal pipeline (missing kernel; see compile_pipeline log above)\n", __func__);
-    }
-
    [encoder->obj setComputePipelineState:pipeline.pipeline->obj];
 }

@@ -1415,30 +1410,6 @@ void ggml_metal_device_get_memory(ggml_metal_device_t dev, size_t * free, size_t
    }
 }

-static bool ggml_metal_supports_mul_mat_op(
-        bool has_simdgroup_reduction,
-        const struct ggml_tensor * op,
-        bool src0_f16_has_mv,
-        bool mm_path) {
-    if (!has_simdgroup_reduction || op->src[0]->type == GGML_TYPE_NVFP4) {
-        return false;
-    }
-
-    if (op->src[1]->type != GGML_TYPE_F16) {
-        return true;
-    }
-
-    if (op->src[0]->type == GGML_TYPE_BF16) {
-        return false;
-    }
-
-    if (src0_f16_has_mv && op->src[0]->type == GGML_TYPE_F16) {
-        return true;
-    }
-
-    return mm_path;
-}
-
 bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_tensor * op) {
    const bool has_simdgroup_mm        = dev->props.has_simdgroup_mm;
    const bool has_simdgroup_reduction = dev->props.has_simdgroup_reduction;
@@ -1742,15 +1713,9 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
        case GGML_OP_GATED_DELTA_NET:
            return has_simdgroup_reduction && op->src[2]->ne[0] % 32 == 0;
        case GGML_OP_SOLVE_TRI:
-            return has_simdgroup_reduction && op->src[0]->type == GGML_TYPE_F32;
        case GGML_OP_MUL_MAT:
-            return ggml_metal_supports_mul_mat_op(
-                    has_simdgroup_reduction, op, true,
-                    ggml_metal_op_mul_mat_use_mm(op, has_simdgroup_mm));
        case GGML_OP_MUL_MAT_ID:
-            return ggml_metal_supports_mul_mat_op(
-                    has_simdgroup_reduction, op, false,
-                    ggml_metal_op_mul_mat_id_use_mm(op, has_simdgroup_mm));
+            return has_simdgroup_reduction && op->src[0]->type != GGML_TYPE_NVFP4;
        case GGML_OP_SET:
        case GGML_OP_CPY:
        case GGML_OP_DUP:
@@ -2362,6 +2362,10 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) {
    const int16_t r2 = ne12/ne02;
    const int16_t r3 = ne13/ne03;

+    // find the break-even point where the matrix-matrix kernel becomes more efficient compared
+    // to the matrix-vector kernel
+    const int ne11_mm_min = 8;
+
    // first try to use small-batch mat-mv kernels
    // these should be efficient for BS [2, ~8]
    if (op->src[1]->type == GGML_TYPE_F32 && (ne00%128 == 0) &&
@@ -2464,7 +2468,12 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) {
        ggml_metal_encoder_set_buffer  (enc, ggml_metal_get_buffer_id(op),         3);

        ggml_metal_encoder_dispatch_threadgroups(enc, ((ne01 + r0ptg - 1)/r0ptg), ((ne11 + r1ptg - 1)/r1ptg), ne12*ne13, 32, nsg, 1);
-    } else if (ggml_metal_op_mul_mat_use_mm(op, props_dev->has_simdgroup_mm)) {
+    } else if (
+        !ggml_is_transposed(op->src[0]) &&
+        !ggml_is_transposed(op->src[1]) &&
+        // for now the matrix-matrix multiplication kernel only works on A14+/M1+ SoCs
+        // AMD GPU and older A-chips will reuse matrix-vector multiplication kernel
+        props_dev->has_simdgroup_mm && ne00 >= 64 && ne11 > ne11_mm_min) {
        //GGML_LOG_INFO("matrix: ne00 = %6d, ne01 = %6d, ne02 = %6d, ne11 = %6d, ne12 = %6d\n", ne00, ne01, ne02, ne11, ne12);

        // some Metal matrix data types require aligned pointers
@@ -2613,7 +2622,13 @@ int ggml_metal_op_mul_mat_id(ggml_metal_op_t ctx, int idx) {
    const uint32_t r2 = 1;
    const uint32_t r3 = 1;

-    if (ggml_metal_op_mul_mat_id_use_mm(op, props_dev->has_simdgroup_mm)) {
+    // find the break-even point where the matrix-matrix kernel becomes more efficient compared
+    // to the matrix-vector kernel
+    // ne20 = n_used_experts
+    // ne21 = n_rows (batch size)
+    const int ne21_mm_id_min = 32;
+
+    if (props_dev->has_simdgroup_mm && ne00 >= 64 && (ne21 >= ne21_mm_id_min)) {
        // some Metal matrix data types require aligned pointers
        // ref: https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf (Table 2.5)
        //switch (op->src[0]->type) {
@@ -720,218 +720,6 @@ constexpr fa_vec_entry_t fa_vec_tuned_table[] = {
    { { GGML_METAL_DEVICE_M2_ULTRA, GGML_TYPE_Q8_0, 576, 512, -1, 0 }, { 1, 4 } },
    { { GGML_METAL_DEVICE_M2_ULTRA, GGML_TYPE_Q8_0, 576, 512, -1, 1 }, { 1, 4 } },

-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 32, 32, 1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 32, 32, 1, 3 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 32, 32, 2, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 32, 32, 2, 3 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 32, 32, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 32, 32, 3, 3 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 64, 64, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 64, 64, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 64, 64, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 64, 64, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 96, 96, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 96, 96, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 96, 96, 1, 4 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 96, 96, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 96, 96, 2, 4 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 128, 128, -1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 128, 128, 1, 2 }, { 1, 1 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 128, 128, 1, 4 }, { 1, 1 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 128, 128, 2, 2 }, { 1, 1 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 128, 128, 2, 4 }, { 1, 1 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 192, 192, -1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 192, 192, 1, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 192, 192, 1, 4 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 192, 128, -1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 192, 128, 1, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 192, 128, 1, 4 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 192, 128, 2, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 256, 256, -1, 0 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 256, 256, 3, 0 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 256, 256, -1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 320, 256, 3, 0 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 320, 256, -1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 512, 512, 3, 0 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 32, 32, -1, 1 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 32, 32, 1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 32, 32, 2, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 32, 32, 2, 4 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 32, 32, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 64, 64, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 64, 64, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 64, 64, 3, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 64, 64, 3, 3 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 96, 96, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 96, 96, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 96, 96, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 128, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 128, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 128, 128, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 128, 128, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 128, 128, 2, 4 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 128, 128, 3, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 192, 192, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 192, 192, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 192, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 192, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 192, 128, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 192, 128, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 192, 128, 2, 4 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 192, 128, 3, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 256, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 256, 256, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 320, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 320, 256, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 512, 512, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 512, 512, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 576, 512, 2, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 576, 512, 3, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 576, 512, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 576, 512, 1, 1 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 576, 512, 1, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 576, 512, 1, 3 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_0, 576, 512, 1, 4 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 32, 32, -1, 1 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 32, 32, 1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 32, 32, 2, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 32, 32, 2, 4 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 32, 32, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 64, 64, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 64, 64, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 64, 64, 3, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 64, 64, 3, 3 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 96, 96, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 96, 96, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 96, 96, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 128, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 128, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 128, 128, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 128, 128, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 128, 128, 3, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 192, 192, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 192, 192, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 192, 192, 2, 3 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 192, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 192, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 192, 128, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 192, 128, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 192, 128, 3, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 256, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 256, 256, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 320, 256, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 576, 512, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q4_1, 576, 512, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 32, 32, -1, 1 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 32, 32, 1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 32, 32, 2, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 32, 32, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 64, 64, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 64, 64, -1, 1 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 64, 64, 1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 64, 64, 2, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 64, 64, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 96, 96, -1, 1 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 96, 96, 1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 96, 96, 2, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 96, 96, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 128, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 128, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 128, 128, 2, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 128, 128, 2, 3 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 128, 128, 3, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 128, 128, 3, 3 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 192, 192, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 192, 192, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 192, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 192, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 256, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 256, 256, -1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 256, 256, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 256, 256, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 256, 256, 3, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 320, 256, -1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 320, 256, 1, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 320, 256, 2, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 320, 256, 3, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 512, 512, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 512, 512, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 576, 512, 2, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 576, 512, 2, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 576, 512, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 576, 512, 2, 3 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_0, 576, 512, 2, 4 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 32, 32, -1, 1 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 32, 32, 1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 32, 32, 2, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 32, 32, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 64, 64, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 64, 64, -1, 1 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 64, 64, 1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 64, 64, 2, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 64, 64, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 96, 96, -1, 1 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 96, 96, 1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 96, 96, 2, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 96, 96, 2, 4 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 96, 96, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 96, 96, 3, 4 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 128, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 128, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 128, 128, 1, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 128, 128, 2, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 128, 128, 2, 3 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 128, 128, 3, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 128, 128, 3, 3 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 192, 192, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 192, 192, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 192, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 192, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 256, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 256, 256, -1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 256, 256, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 256, 256, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 256, 256, 3, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 320, 256, -1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 320, 256, 1, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 320, 256, 2, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 320, 256, 3, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 512, 512, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 512, 512, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 576, 512, 2, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 576, 512, 2, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 576, 512, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 576, 512, 2, 3 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q5_1, 576, 512, 2, 4 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 32, 32, -1, 1 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 32, 32, 1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 32, 32, 2, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 32, 32, 2, 4 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 32, 32, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 64, 64, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 64, 64, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 64, 64, 3, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 64, 64, 3, 3 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 96, 96, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 128, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 128, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 128, 128, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 128, 128, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 192, 192, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 192, 192, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 192, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 192, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 192, 128, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 192, 128, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 256, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 256, 256, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 320, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 320, 256, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 512, 512, -1, 0 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 512, 512, -1, 1 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 576, 512, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_Q8_0, 576, 512, -1, 1 }, { 1, 4 } },
-
    { { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 32, 32, 1, 3 }, { 4, 4 } },
    { { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 32, 32, 2, 1 }, { 2, 4 } },
    { { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 32, 32, 2, 3 }, { 4, 4 } },
@@ -1024,204 +812,6 @@ constexpr fa_vec_entry_t fa_vec_tuned_table[] = {
    { { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 576, 512, 1, 1 }, { 1, 2 } },
    { { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 576, 512, 1, 4 }, { 1, 2 } },

-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 32, 32, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 32, 32, 3, 3 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 64, 64, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 64, 64, 1, 1 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 64, 64, 1, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 64, 64, 2, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 96, 96, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 96, 96, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 96, 96, 1, 3 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 96, 96, 1, 4 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 96, 96, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 128, 128, -1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 128, 128, 1, 1 }, { 1, 1 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 128, 128, 1, 2 }, { 1, 1 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 128, 128, 1, 4 }, { 1, 1 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 192, 192, -1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 192, 192, 1, 1 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 192, 192, 1, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 192, 128, -1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 192, 128, 1, 1 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 192, 128, 1, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 256, 256, -1, 0 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 256, 256, -1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 256, 256, 1, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 320, 256, -1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 512, 512, 3, 3 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_F16, 512, 512, 3, 4 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 32, 32, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 32, 32, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 32, 32, 3, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 64, 64, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 64, 64, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 64, 64, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 96, 96, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 96, 96, 3, 3 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 96, 96, 3, 4 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 128, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 128, 128, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 192, 192, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 192, 192, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 192, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 192, 128, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 192, 128, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 256, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 256, 256, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 320, 256, 2, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 320, 256, 3, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 320, 256, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 512, 512, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 512, 512, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 576, 512, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 576, 512, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 576, 512, 1, 1 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 576, 512, 1, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 576, 512, 1, 3 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_0, 576, 512, 1, 4 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 32, 32, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 32, 32, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 32, 32, 3, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 64, 64, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 64, 64, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 64, 64, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 64, 64, 2, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 96, 96, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 96, 96, 3, 3 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 96, 96, 3, 4 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 128, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 128, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 128, 128, 1, 2 }, { 1, 1 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 128, 128, 1, 4 }, { 1, 1 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 128, 128, 2, 2 }, { 1, 1 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 128, 128, 3, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 192, 192, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 192, 192, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 192, 192, 1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 192, 192, 2, 3 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 192, 192, 2, 4 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 192, 128, 2, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 192, 128, 3, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 192, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 192, 128, 3, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 256, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 256, 256, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 320, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 320, 256, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 576, 512, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q4_1, 576, 512, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 32, 32, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 32, 32, 3, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 32, 32, 3, 3 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 32, 32, 3, 4 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 64, 64, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 64, 64, -1, 1 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 64, 64, 1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 64, 64, 2, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 64, 64, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 96, 96, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 96, 96, 1, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 96, 96, 1, 3 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 128, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 128, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 128, 128, 1, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 128, 128, 1, 3 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 192, 192, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 192, 192, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 192, 192, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 192, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 192, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 256, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 256, 256, -1, 1 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 256, 256, 1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 256, 256, 2, 3 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 256, 256, 2, 4 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 320, 256, 1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 512, 512, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 512, 512, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 576, 512, 2, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 576, 512, 3, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 576, 512, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 576, 512, 1, 1 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 576, 512, 1, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 576, 512, 1, 3 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_0, 576, 512, 1, 4 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 32, 32, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 32, 32, 3, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 32, 32, 3, 3 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 32, 32, 3, 4 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 64, 64, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 64, 64, -1, 1 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 64, 64, 1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 64, 64, 2, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 64, 64, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 96, 96, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 96, 96, 1, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 96, 96, 1, 3 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 128, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 128, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 128, 128, 1, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 128, 128, 1, 3 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 192, 192, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 192, 192, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 192, 192, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 192, 192, 3, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 192, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 192, 128, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 256, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 256, 256, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 256, 256, 2, 3 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 256, 256, 2, 4 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 320, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 320, 256, 1, 1 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 320, 256, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 320, 256, 1, 3 }, { 2, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 512, 512, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 512, 512, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 512, 512, 1, 3 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 576, 512, 2, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 576, 512, 3, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 576, 512, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 576, 512, 1, 1 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 576, 512, 1, 2 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 576, 512, 1, 3 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q5_1, 576, 512, 1, 4 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 32, 32, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 32, 32, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 32, 32, 3, 2 }, { 4, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 64, 64, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 64, 64, -1, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 64, 64, 1, 2 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 96, 96, 2, 3 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 96, 96, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 96, 96, 3, 3 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 96, 96, 3, 4 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 128, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 128, 128, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 128, 128, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 128, 128, 3, 3 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 128, 128, 3, 4 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 192, 192, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 192, 192, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 192, 128, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 192, 128, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 192, 128, 2, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 192, 128, 2, 3 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 192, 128, 2, 4 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 192, 128, 3, 1 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 192, 128, 3, 3 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 192, 128, 3, 4 }, { 2, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 256, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 256, 256, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 320, 256, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 320, 256, -1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 512, 512, -1, 0 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 512, 512, -1, 1 }, { 1, 2 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 512, 512, 1, 1 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 576, 512, -1, 0 }, { 1, 4 } },
-    { { GGML_METAL_DEVICE_M3_ULTRA, GGML_TYPE_Q8_0, 576, 512, -1, 1 }, { 1, 4 } },
-
    { { GGML_METAL_DEVICE_M4, GGML_TYPE_F16, 32, 32, 1, 1 }, { 2, 4 } },
    { { GGML_METAL_DEVICE_M4, GGML_TYPE_F16, 32, 32, 1, 3 }, { 4, 4 } },
    { { GGML_METAL_DEVICE_M4, GGML_TYPE_F16, 32, 32, 2, 1 }, { 2, 4 } },
@@ -34,14 +34,10 @@ if (GGML_RPC_RDMA)
        find_library(RDMA_LIB ${RDMA_LIB_NAME} REQUIRED)
    endif()
    target_compile_definitions(ggml-rpc PRIVATE GGML_RPC_RDMA)
+    target_link_libraries(ggml-rpc PRIVATE ${RDMA_LIB})
    if (APPLE)
-        # librdma.dylib only exists on macOS 26.2 and later. Link it weakly so a build made
-        # where it exists still loads where it does not; checked at runtime before use.
-        target_link_options(ggml-rpc PRIVATE "LINKER:-weak_library,${RDMA_LIB}")
        target_compile_definitions(ggml-rpc PRIVATE GGML_RPC_RDMA_APPLE)
        target_sources(ggml-rpc PRIVATE transport-apple.cpp)
-    else()
-        target_link_libraries(ggml-rpc PRIVATE ${RDMA_LIB})
    endif()
    message(STATUS "  RDMA transport enabled (${RDMA_DESC})")
 else()
@@ -8,7 +8,6 @@
 #include <cstdlib>
 #include <cstring>
 #include <string>
-#include <dlfcn.h>
 #include <poll.h>
 #include <sys/socket.h>
 #include <unistd.h>
@@ -115,9 +114,16 @@ struct apple_rdma::impl {

    ~impl() {
        broken = true;
-        // destroy the QP first: it can still write to the rings until it is gone.
-        // no IBV_QPS_ERR before it - Apple's provider then fails every region unmap.
-        if (qp) ibv_destroy_qp(qp);
+        // the QP must be destroyed before the memory it can still write to is
+        // deregistered and freed: ERR only starts flushing the posted WQEs
+        if (qp) {
+            struct ibv_qp_attr a = {};
+            a.qp_state = IBV_QPS_ERR;
+            ibv_modify_qp(qp, &a, IBV_QP_STATE);
+            struct ibv_wc wc[RDMA_NBUF * 2];
+            while (ibv_poll_cq(cq, RDMA_NBUF * 2, wc) > 0) {}
+            ibv_destroy_qp(qp);
+        }
        if (send_mr) ibv_dereg_mr(send_mr);
        if (recv_mr) ibv_dereg_mr(recv_mr);
        free(send_mem);
@@ -178,28 +184,11 @@ static uint8_t rdma_first_active_port(struct ibv_context * ctx, struct ibv_port_
    return 0;
 }

-// librdma.dylib is weak-linked, so its symbols are null when it is absent. Nothing may
-// call one before this has returned true.
-static bool rdma_library_present() {
-    static const bool present = [] {
-        void * handle = dlopen("/usr/lib/librdma.dylib", RTLD_LAZY);
-        if (handle == nullptr) {
-            return false;
-        }
-        dlclose(handle);
-        return true;
-    }();
-    return present;
-}
-
 // Called before the endpoints are exchanged: pick the local device facing this
 // peer, create a UC QP and register the frame rings. RDMA is point-to-point, so
 // the device is the one whose GID equals the bootstrap connection's local
 // address, i.e. the one cabled to the peer.
 std::unique_ptr<apple_rdma> apple_rdma::probe(int fd, const uint8_t * target_gid, uint8_t * caps) {
-    if (!rdma_library_present()) {
-        return nullptr;
-    }
    int ndev = 0;
    ibv_device ** devs = ibv_get_device_list(&ndev);
    if (!devs) return nullptr;
@@ -2402,138 +2402,7 @@ static void argsort_f32_i32_sycl(const float *x, int *dst, const int ncols,
    }
 }

-// Scan and block merge, shared by every launch shape below so a partitioned row uses the
-// same insertion order as an unpartitioned one.
-//
-// src_map  != nullptr: report src_map[col] instead of col, so a merge pass can carry the
-//                      original column index through.
-// out_vals != nullptr: also emit the k winning values, for a later merge pass.
-// swap01:              emit in the output order the single-pass path uses.
-static void top_k_scan_merge_f32(
-    const float *   src_vals,
-    const int32_t * src_map,
-    const int       begin,
-    const int       end,
-    const int       k,
-    const int       block_size,
-    float *         shared_vals,
-    int *           shared_idx,
-    float *         out_vals,
-    int32_t *       out_idx,
-    const bool      swap01,
-    const sycl::nd_item<1> & item_ct1
-) {
-    const int tid = item_ct1.get_local_id(0);
-
-    // The running top-k lives in SLM (shared local memory) rather than a private array:
-    // an array indexed by a runtime position cannot be register-allocated, so a private
-    // one lands in scratch, i.e. device memory, and insertion is this kernel's dominant
-    // cost.
-    //
-    // Lane-strided (lv[i * block_size]) rather than lane-blocked (lv[i]) so a given i is
-    // contiguous across lanes; a k-strided layout would put every lane of a shift step in
-    // the same SLM bank.
-    float * lv = shared_vals + tid;
-    int * li = shared_idx + tid;
-
-    for (int i = 0; i < k; i++) {
-        lv[i * block_size] = -FLT_MAX;
-        li[i * block_size] = -1;
-    }
-
-    // The k-th best, cached in a register. The reject test is taken for the large
-    // majority of elements scanned, and in that case touches no memory.
-    float kth = -FLT_MAX;
-
-    for (int col = begin + tid; col < end; col += block_size) {
-        float val = src_vals[col];
-
-        if (val > kth) {
-            int pos = k - 1;
-            while (pos > 0 && val > lv[(pos - 1) * block_size]) {
-                pos--;
-            }
-
-            for (int i = k - 1; i > pos; i--) {
-                lv[i * block_size] = lv[(i - 1) * block_size];
-                li[i * block_size] = li[(i - 1) * block_size];
-            }
-            lv[pos * block_size] = val;
-            li[pos * block_size] = src_map ? src_map[col] : col;
-
-            kth = lv[(k - 1) * block_size];
-        }
-    }
-
-    item_ct1.barrier(sycl::access::fence_space::local_space);
-
-    if (tid != 0) {
-        return;
-    }
-
-    // Same treatment for the merge accumulator, past the per-lane region.
-    float * fv = shared_vals + (size_t) k * block_size;
-    int * fi = shared_idx + (size_t) k * block_size;
-
-    for (int i = 0; i < k; i++) {
-        fv[i] = -FLT_MAX;
-        fi[i] = -1;
-    }
-
-    float fkth = -FLT_MAX;
-
-    // Candidates are visited in the same (t, i) order as before, so tie-breaking is
-    // unchanged.
-    for (int t = 0; t < block_size; t++) {
-        for (int i = 0; i < k; i++) {
-            float val = shared_vals[i * block_size + t];
-
-            if (val <= fkth) {
-                // Lane t's list is sorted descending, so once one of its entries loses
-                // to the k-th best, every later entry loses too. fkth only rises, so
-                // that stays true for the rest of the merge. This turns the merge from
-                // block_size*k steps into roughly block_size plus the candidates
-                // accepted.
-                break;
-            }
-
-            int idx = shared_idx[i * block_size + t];
-
-            int pos = k - 1;
-            while (pos > 0 && val > fv[pos - 1]) {
-                pos--;
-            }
-
-            for (int j = k - 1; j > pos; j--) {
-                fv[j] = fv[j - 1];
-                fi[j] = fi[j - 1];
-            }
-            fv[pos] = val;
-            fi[pos] = idx;
-
-            fkth = fv[k - 1];
-        }
-    }
-
-    if (out_vals) {
-        for (int i = 0; i < k; i++) {
-            out_vals[i] = fv[i];
-        }
-    }
-
-    for (int i = 0; i < k; i++) {
-        out_idx[i] = fi[i];
-    }
-
-    if (swap01 && k > 1) {
-        int32_t temp = out_idx[0];
-        out_idx[0] = out_idx[1];
-        out_idx[1] = temp;
-    }
-}
-
 static void top_k_f32_sycl(
-    ggml_backend_sycl_context & ctx,
    const float * src,
    int32_t * dst_indices,
    const int64_t ncols,
@@ -2541,107 +2410,98 @@ static void top_k_f32_sycl(
    const int k,
    dpct::queue_ptr main_stream
 ) {
-    // A row is scanned by exactly one work-group, so a vocabulary-sized row leaves the
-    // rest of the device idle. What the scan is short of is memory requests in flight,
-    // not bandwidth or per-request latency, so lanes in flight is the lever: split the
-    // row across independent work-groups, have each emit its partition's top-k, and
-    // merge those nsplit*k candidates in a second launch.
-    //
-    // split_block trades parallelism against SLM residency. Its cost is
-    // (split_block + 1) * k * 8 bytes of SLM per group, so at the k <= 32 ceiling 128
-    // lanes need about 33 KB, which leaves a single resident group per Xe-core. Revisit
-    // if the supported k ever grows.
-    constexpr int split_block = 128;
-    constexpr int max_splits  = 128;
-    constexpr int min_cols    = 8192;
-
-    int nsplit = 1;
-    if (ncols >= min_cols) {
-        // A partition is then always >= split_block = 128 columns, hence always more than
-        // the k <= 32 ceiling, so no pass is ever padded with -FLT_MAX sentinels.
-        const int64_t want = ncols / split_block;
-        nsplit = (int) (want > max_splits ? max_splits : want);
-    }
-
-    if (nsplit > 1) {
-        const int nchunk = (int) ((ncols + nsplit - 1) / nsplit);
-        const size_t ncand = (size_t) nrows * nsplit * k;
-
-        ggml_sycl_pool_alloc<float>   part_vals(ctx.pool(), ncand);
-        ggml_sycl_pool_alloc<int32_t> part_idx(ctx.pool(), ncand);
-
-        float *   pv = part_vals.get();
-        int32_t * pi = part_idx.get();
-
-        const sycl::range<1> block_dims(split_block);
-
-        main_stream->submit([&](sycl::handler &cgh) {
-            sycl::local_accessor<float, 1> shared_vals(sycl::range<1>((split_block + 1) * k), cgh);
-            sycl::local_accessor<int, 1> shared_idx(sycl::range<1>((split_block + 1) * k), cgh);
-
-            cgh.parallel_for(
-                sycl::nd_range<1>(sycl::range<1>(nrows * nsplit) * block_dims, block_dims),
-                [=](sycl::nd_item<1> item_ct1) {
-                    const int grp = item_ct1.get_group(0);
-                    const int row = grp / nsplit;
-                    const int part = grp % nsplit;
-
-                    const int begin = part * nchunk;
-                    int end = begin + nchunk;
-                    if (end > (int) ncols) {
-                        end = (int) ncols;
-                    }
-
-                    top_k_scan_merge_f32(
-                        src + (int64_t) row * ncols, nullptr, begin, end, k, split_block,
-                        shared_vals.get_multi_ptr<sycl::access::decorated::no>().get(),
-                        shared_idx.get_multi_ptr<sycl::access::decorated::no>().get(),
-                        pv + (size_t) grp * k, pi + (size_t) grp * k, false, item_ct1);
-                });
-        });
-
-        main_stream->submit([&](sycl::handler &cgh) {
-            sycl::local_accessor<float, 1> shared_vals(sycl::range<1>((split_block + 1) * k), cgh);
-            sycl::local_accessor<int, 1> shared_idx(sycl::range<1>((split_block + 1) * k), cgh);
-
-            cgh.parallel_for(
-                sycl::nd_range<1>(sycl::range<1>(nrows) * block_dims, block_dims),
-                [=](sycl::nd_item<1> item_ct1) {
-                    const int row = item_ct1.get_group(0);
-                    const size_t off = (size_t) row * nsplit * k;
-
-                    top_k_scan_merge_f32(
-                        pv + off, pi + off, 0, nsplit * k, k, split_block,
-                        shared_vals.get_multi_ptr<sycl::access::decorated::no>().get(),
-                        shared_idx.get_multi_ptr<sycl::access::decorated::no>().get(),
-                        nullptr, dst_indices + (int64_t) row * k, true, item_ct1);
-                });
-        });
-
-        return;
-    }
-
    const int block_size = 128;

    const sycl::range<1> block_dims(block_size);
    const sycl::range<1> grid_dims(nrows);

    main_stream->submit([&](sycl::handler &cgh) {
-        sycl::local_accessor<float, 1> shared_vals(sycl::range<1>((block_size + 1) * k), cgh);
-        sycl::local_accessor<int, 1> shared_idx(sycl::range<1>((block_size + 1) * k), cgh);
+        sycl::local_accessor<float, 1> shared_vals(sycl::range<1>(block_size * k), cgh);
+        sycl::local_accessor<int, 1> shared_idx(sycl::range<1>(block_size * k), cgh);

        cgh.parallel_for(
            sycl::nd_range<1>(grid_dims * block_dims, block_dims),
            [=](sycl::nd_item<1> item_ct1) {
                const int row = item_ct1.get_group(0);
+                const int tid = item_ct1.get_local_id(0);

                if (row >= nrows) return;

-                top_k_scan_merge_f32(
-                    src + (int64_t) row * ncols, nullptr, 0, (int) ncols, k, block_size,
-                    shared_vals.get_multi_ptr<sycl::access::decorated::no>().get(),
-                    shared_idx.get_multi_ptr<sycl::access::decorated::no>().get(),
-                    nullptr, dst_indices + (int64_t) row * k, true, item_ct1);
+                const float * src_row = src + row * ncols;
+                int32_t * dst_idx_row = dst_indices + row * k;
+
+                float local_vals[32];
+                int local_idx[32];
+
+                for (int i = 0; i < k; i++) {
+                    local_vals[i] = -FLT_MAX;
+                    local_idx[i] = -1;
+                }
+
+                for (int col = tid; col < ncols; col += block_size) {
+                    float val = src_row[col];
+
+                    if (val > local_vals[k-1]) {
+                        int pos = k - 1;
+                        while (pos > 0 && val > local_vals[pos - 1]) {
+                            pos--;
+                        }
+
+                        for (int i = k - 1; i > pos; i--) {
+                            local_vals[i] = local_vals[i - 1];
+                            local_idx[i] = local_idx[i - 1];
+                        }
+                        local_vals[pos] = val;
+                        local_idx[pos] = col;
+                    }
+                }
+
+                for (int i = 0; i < k; i++) {
+                    shared_vals[tid * k + i] = local_vals[i];
+                    shared_idx[tid * k + i] = local_idx[i];
+                }
+                item_ct1.barrier(sycl::access::fence_space::local_space);
+
+                if (tid == 0) {
+                    float final_vals[32];
+                    int final_idx[32];
+
+                    for (int i = 0; i < k; i++) {
+                        final_vals[i] = -FLT_MAX;
+                        final_idx[i] = -1;
+                    }
+
+                    for (int t = 0; t < block_size; t++) {
+                        for (int i = 0; i < k; i++) {
+                            float val = shared_vals[t * k + i];
+                            int idx = shared_idx[t * k + i];
+
+                            if (val > final_vals[k-1]) {
+                                int pos = k - 1;
+                                while (pos > 0 && val > final_vals[pos - 1]) {
+                                    pos--;
+                                }
+
+                                for (int j = k - 1; j > pos; j--) {
+                                    final_vals[j] = final_vals[j - 1];
+                                    final_idx[j] = final_idx[j - 1];
+                                }
+                                final_vals[pos] = val;
+                                final_idx[pos] = idx;
+                            }
+                        }
+                    }
+
+                    for (int i = 0; i < k; i++) {
+                        dst_idx_row[i] = final_idx[i];
+                    }
+
+                    if (k > 1) {
+                        int32_t temp = dst_idx_row[0];
+                        dst_idx_row[0] = dst_idx_row[1];
+                        dst_idx_row[1] = temp;
+                    }
+                }
            });
    });
 }
@@ -3042,7 +2902,7 @@ static void ggml_sycl_op_top_k(ggml_backend_sycl_context & ctx, ggml_tensor * ds
    GGML_ASSERT(k > 0 && k <= 32);
    GGML_ASSERT(k <= ncols);

-    top_k_f32_sycl(ctx, src0_dd, dst_dd, ncols, nrows, k, main_stream);
+    top_k_f32_sycl(src0_dd, dst_dd, ncols, nrows, k, main_stream);
 }

 inline void ggml_sycl_op_argmax(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
@@ -5,7 +5,7 @@ import os
 import sys
 import subprocess

-HTTPLIB_VERSION = "refs/tags/v0.54.1"
+HTTPLIB_VERSION = "refs/tags/v0.53.1"

 # used by examples/gguf-hash, these repos have no release tag, so we pin a commit
 XXHASH_COMMIT      = "9f465f1ea932d6ad9a26cd77496311ffa544cd68"
@@ -2034,7 +2034,7 @@ void llm_graph_input_k_shift::set_input(const llama_ubatch * ubatch) {
        kv_self->set_input_k_shift(k_shift);
    }

-    if (k_rot && k_rot->buffer) {
+    if (k_rot) {
        kv_self->set_input_k_rot(k_rot);
    }
 }
@@ -8,8 +8,8 @@
 #ifndef CPPHTTPLIB_HTTPLIB_H
 #define CPPHTTPLIB_HTTPLIB_H

-#define CPPHTTPLIB_VERSION "0.54.1"
-#define CPPHTTPLIB_VERSION_NUM "0x003601"
+#define CPPHTTPLIB_VERSION "0.53.1"
+#define CPPHTTPLIB_VERSION_NUM "0x003501"

 #ifdef _WIN32
 #if defined(_WIN32_WINNT) && _WIN32_WINNT < 0x0A00
@@ -134,16 +134,6 @@
 #define CPPHTTPLIB_FORM_URL_ENCODED_PAYLOAD_MAX_LENGTH 8192
 #endif

-#ifndef CPPHTTPLIB_STATIC_FILE_COMPRESSION_MIN_LENGTH
-// 1400 rather than a round number: a body that already fits in one 1500-byte
-// MTU gains nothing from being made smaller.
-#define CPPHTTPLIB_STATIC_FILE_COMPRESSION_MIN_LENGTH 1400
-#endif
-
-#ifndef CPPHTTPLIB_STATIC_FILE_COMPRESSION_MAX_LENGTH
-#define CPPHTTPLIB_STATIC_FILE_COMPRESSION_MAX_LENGTH (4 * 1024 * 1024) // 4MB
-#endif
-
 #ifndef CPPHTTPLIB_RANGE_MAX_COUNT
 #define CPPHTTPLIB_RANGE_MAX_COUNT 1024
 #endif
@@ -1439,16 +1429,9 @@ public:
  DataSink &operator=(DataSink &&) = delete;

  std::function<bool(const char *data, size_t data_len)> write;
-
-  // Only `write` is mandatory. The rest are defaulted so that a provider
-  // calling one on a writer that does not set it gets sensible behaviour
-  // rather than std::bad_function_call thrown from a worker thread. Capturing
-  // `this` is safe: DataSink is neither copyable nor movable.
-  std::function<bool()> is_writable = []() { return true; };
-  std::function<void()> done = []() {};
-  std::function<void(const Headers &trailer)> done_with_trailer =
-      [this](const Headers & /*trailer*/) { done(); };
-
+  std::function<bool()> is_writable;
+  std::function<void()> done;
+  std::function<void(const Headers &trailer)> done_with_trailer;
  std::ostream os;

 private:
@@ -1533,10 +1516,7 @@ make_file_body(const std::string &filepath) {
      auto to_read = (std::min)(sizeof(buf), length);
      f.read(buf, static_cast<std::streamsize>(to_read));
      auto n = static_cast<size_t>(f.gcount());
-      // The file is shorter than the size make_file_body() measured, which the
-      // caller has already committed to as Content-Length. The body cannot be
-      // completed, so fail as every other error here does.
-      if (n == 0) { return false; }
+      if (n == 0) { break; }
      if (!sink.write(buf, n)) { return false; }
      length -= n;
    }
@@ -1743,14 +1723,6 @@ struct Request {
 #endif
 };

-namespace detail {
-
-// Declared up here, away from the rest of the compression helpers, because
-// `Response` stores one.
-enum class EncodingType { None = 0, Gzip, Brotli, Zstd };
-
-} // namespace detail
-
 struct Response {
  std::string version;
  int status = -1;
@@ -1816,11 +1788,6 @@ struct Response {
  bool content_provider_success_ = false;
  std::string file_content_path_;
  std::string file_content_content_type_;
-
-  // Content coding chosen for a file-backed content provider, decided once
-  // where the file is opened so that the ETag and the body cannot disagree.
-  // `EncodingType::None` for every other kind of response.
-  detail::EncodingType file_content_encoding_ = detail::EncodingType::None;
 };

 enum class Error {
@@ -1860,7 +1827,6 @@ enum class Error {
  InvalidRangeHeader,
  UnsupportedContentEncoding,
  WebSocketHandshake,
-  UserCallbackException,

  // For internal use only
  SSLPeerCouldBeClosed_,
@@ -2054,10 +2020,6 @@ private:

 int close_socket(socket_t sock) noexcept;

-bool is_accept_resource_error();
-
-bool is_accept_transient_error();
-
 ssize_t write_headers(Stream &strm, const Headers &headers);

 bool set_socket_opt_time(socket_t sock, int level, int optname, time_t sec,
@@ -2145,17 +2107,6 @@ public:
  Server &Delete(const std::string &pattern, HandlerWithContentReader handler);
  Server &Options(const std::string &pattern, Handler handler);

-  // Register a handler for an HTTP method outside the built-in set (e.g. the
-  // WebDAV methods from RFC 4918). Registering a method here is what makes the
-  // server accept it; an unregistered method is still rejected with 400.
-  // `method` must be a valid HTTP method token and must not be one of the
-  // built-in methods, which have their own registration functions above. A
-  // rejected registration makes is_valid() return false, so listen() fails.
-  Server &CustomRoute(const std::string &method, const std::string &pattern,
-                      Handler handler);
-  Server &CustomRoute(const std::string &method, const std::string &pattern,
-                      HandlerWithContentReader handler);
-
  Server &WebSocket(const std::string &pattern, WebSocketHandler handler);
  Server &WebSocket(const std::string &pattern, WebSocketHandler handler,
                    SubProtocolSelector sub_protocol_selector);
@@ -2223,10 +2174,6 @@ public:

  Server &set_payload_max_length(size_t length);

-  Server &set_static_file_compression(bool on);
-  Server &set_static_file_compression_min_length(size_t length);
-  Server &set_static_file_compression_max_length(size_t length);
-
  Server &set_websocket_ping_interval(time_t sec);
  template <class Rep, class Period>
  Server &set_websocket_ping_interval(
@@ -2255,35 +2202,6 @@ protected:
                       const std::function<void(Request &)> &setup_request,
                       bool *websocket_upgraded = nullptr);

-  // Runs the per-connection serving loop and stops an exception thrown by a
-  // user callback from escaping the worker thread.
-  //
-  // process_request() wraps only routing() in a try/catch. Content providers,
-  // the post-routing, error, logging and expect-100 handlers and WebSocket
-  // handlers all run outside it, and the task queue calls the job without a
-  // catch, so an exception from any of those would terminate the process.
-  //
-  // No 500 is possible here: by the time a content provider runs, the status
-  // line and headers are already on the wire. Report it through the error
-  // logger and drop the connection, which is what the peer observes either
-  // way. Other connections are unaffected.
-  template <typename Serve> bool serve_guarded(Serve &&serve) const {
-#ifdef CPPHTTPLIB_NO_EXCEPTIONS
-    return serve();
-#else
-    try {
-      return serve();
-    } catch (...) {
-      // The error logger is a user callback too, so it must not be able to
-      // throw the guard back open.
-      try {
-        output_error_log(Error::UserCallbackException, nullptr);
-      } catch (...) {}
-      return false;
-    }
-#endif
-  }
-
  std::atomic<socket_t> svr_sock_{INVALID_SOCKET};

  std::vector<std::string> trusted_proxies_;
@@ -2297,11 +2215,6 @@ protected:
  time_t idle_interval_sec_ = CPPHTTPLIB_IDLE_INTERVAL_SECOND;
  time_t idle_interval_usec_ = CPPHTTPLIB_IDLE_INTERVAL_USECOND;
  size_t payload_max_length_ = CPPHTTPLIB_PAYLOAD_MAX_LENGTH;
-  bool static_file_compression_ = false;
-  size_t static_file_compression_min_length_ =
-      CPPHTTPLIB_STATIC_FILE_COMPRESSION_MIN_LENGTH;
-  size_t static_file_compression_max_length_ =
-      CPPHTTPLIB_STATIC_FILE_COMPRESSION_MAX_LENGTH;
  time_t websocket_ping_interval_sec_ =
      CPPHTTPLIB_WEBSOCKET_PING_INTERVAL_SECOND;
  int websocket_max_missed_pongs_ = CPPHTTPLIB_WEBSOCKET_MAX_MISSED_PONGS;
@@ -2313,21 +2226,9 @@ private:
      std::vector<std::pair<std::unique_ptr<detail::MatcherBase>,
                            HandlerWithContentReader>>;

-  // Both handler tables for one custom method live in a single entry, so that
-  // routing() needs only one map lookup per request to reach either of them.
-  struct CustomHandlerEntry {
-    Handlers handlers;
-    HandlersForContentReader handlers_for_content_reader;
-  };
-  using CustomHandlers = std::map<std::string, CustomHandlerEntry>;
-
  static std::unique_ptr<detail::MatcherBase>
  make_matcher(const std::string &pattern);

-  static const std::set<std::string> &builtin_methods();
-  CustomHandlerEntry *custom_entry_for_registration(const std::string &method);
-  const CustomHandlerEntry *find_custom_entry(const std::string &method) const;
-
  template <typename H>
  Server &add_handler(
      std::vector<std::pair<std::unique_ptr<detail::MatcherBase>, H>> &handlers,
@@ -2358,10 +2259,6 @@ private:
      const HandlersForContentReader &handlers) const;

  bool parse_request_line(const char *s, Request &req) const;
-  detail::EncodingType static_file_encoding(const Request &req,
-                                            const std::string &content_type,
-                                            size_t length) const;
-  bool apply_static_file_compression(const Request &req, Response &res) const;
  void apply_ranges(const Request &req, Response &res,
                    std::string &content_type, std::string &boundary) const;
  bool write_response(Stream &strm, bool close_connection, Request &req,
@@ -2395,10 +2292,6 @@ private:
  std::atomic<bool> is_running_{false};
  std::atomic<bool> is_decommissioned{false};

-  // Set when CustomRoute() refuses a registration. Written before listen(),
-  // read by is_valid() on the same thread, so it needs no synchronization.
-  bool has_invalid_registration_ = false;
-
  struct MountPointEntry {
    std::string mount_point;
    std::string base_dir;
@@ -2420,7 +2313,6 @@ private:
  Handlers delete_handlers_;
  HandlersForContentReader delete_handlers_for_content_reader_;
  Handlers options_handlers_;
-  CustomHandlers custom_handlers_;

  struct WebSocketHandlerEntry {
    std::unique_ptr<detail::MatcherBase> matcher;
@@ -3608,16 +3500,6 @@ void split(const char *b, const char *e, char d,
 void split(const char *b, const char *e, char d, size_t m,
           std::function<void(const char *, const char *)> fn);

-bool split_find(const char *b, const char *e, char d,
-                std::function<bool(const char *, const char *)> fn);
-
-bool has_header_token(const Headers &headers, const std::string &key,
-                      const std::string &token);
-
-std::string websocket_accept_key(const std::string &client_key);
-
-bool is_websocket_upgrade(const Request &req);
-
 bool process_client_socket(
    socket_t sock, time_t read_timeout_sec, time_t read_timeout_usec,
    time_t write_timeout_sec, time_t write_timeout_usec,
@@ -3638,9 +3520,6 @@ socket_t create_client_socket(const std::string &host, const std::string &ip,
 const char *get_header_value(const Headers &headers, const std::string &key,
                             const char *def, size_t id);

-std::string get_combined_header_value(const Headers &headers,
-                                      const std::string &key);
-
 std::string params_to_query_str(const Params &params);

 void parse_query_text(const char *data, std::size_t size, Params &params);
@@ -3655,13 +3534,11 @@ bool parse_range_header(const std::string &s, Ranges &ranges);
 bool parse_accept_header(const std::string &s,
                         std::vector<std::string> &content_types);

-void parse_disposition_params(const std::string &s, Params &params);
-
 ssize_t send_socket(socket_t sock, const void *ptr, size_t size, int flags);

 ssize_t read_socket(socket_t sock, void *ptr, size_t size, int flags);

-EncodingType encoding_type(const Request &req, const std::string &content_type);
+enum class EncodingType { None = 0, Gzip, Brotli, Zstd };

 EncodingType encoding_type(const Request &req, const Response &res);

@@ -4441,11 +4318,6 @@ private:
  int unacked_pings_ = 0;
  std::atomic<bool> closed_{false};
  std::mutex write_mutex_;
-  // Owned by whichever thread is parsing frames off strm_. Only one thread
-  // may do so: read_websocket_frame() reads a payload until it has the whole
-  // declared length, so a second parser stealing bytes silently corrupts the
-  // message the first one is assembling.
-  std::mutex read_mutex_;
  std::thread ping_thread_;
  std::mutex ping_mutex_;
  std::condition_variable ping_cv_;