ggml : add ggml_backend_op_alloc_size_may_expand, use it in RPC (#27960 )

some backends (Metal, SYCL, WebGPU) require additional memory for fleeting data for certain ops, which is reflected in their get_alloc_size implementations. add ggml_backend_op_alloc_size_may_expand() to the backend utils, listing these ops, and assert in ggml_backend_buft_get_alloc_size that a backend expanding the alloc size of a compute op only does so for ops listed in the helper. use the helper in the RPC backend to decide whether to query the remote server for the actual alloc size, instead of a hardcoded list. Assisted-by: pi:llama.cpp/Qwen3.8-27B
rpc: fix apple rdma error spew on teardown (#27908 )
2026-08-31 04:27:33 +02:00 · 2026-08-30 09:17:47 +03:00 · 2026-08-30 09:16:26 +03:00 · 2026-08-30 09:06:29 +03:00 · 2026-08-30 09:02:22 +03:00 · 2026-08-30 09:01:51 +03:00
20 changed files with 2053 additions and 1178 deletions
@@ -424,6 +424,10 @@ extern "C" {
    // Compare the output of two backends
    GGML_API bool ggml_backend_compare_graph_backend(ggml_backend_t backend1, ggml_backend_t backend2, struct ggml_cgraph * graph, ggml_backend_eval_callback callback, void * user_data, struct ggml_tensor const * const * test_nodes, size_t num_test_nodes);

+    // returns true for ops that may require additional memory for fleeting data on some backends,
+    // i.e. the backend's get_alloc_size may return more than ggml_nbytes for the output tensor
+    GGML_API bool ggml_backend_op_alloc_size_may_expand(enum ggml_op op);
+
    // Tensor initialization
    GGML_API enum ggml_status ggml_backend_tensor_alloc(ggml_backend_buffer_t buffer, struct ggml_tensor * tensor, void * addr);
    GGML_API enum ggml_status ggml_backend_view_init(struct ggml_tensor * tensor);
@@ -6,7 +6,7 @@
 extern "C" {
 #endif

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

@@ -865,12 +865,7 @@ 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: {
-                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};
-                }
+                split_state = {GGML_BACKEND_SPLIT_AXIS_MIRRORED, {0}, {1}, 1};
            } break;
            case GGML_OP_DUP: {
                split_state = handle_generic(src_ss, /*scalar_only =*/ true);
@@ -2288,14 +2283,6 @@ 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
@@ -65,6 +65,13 @@ size_t ggml_backend_buft_get_alloc_size(ggml_backend_buffer_type_t buft, const s
    if (buft->iface.get_alloc_size) {
        size_t size = buft->iface.get_alloc_size(buft, tensor);
        assert(size >= ggml_nbytes(tensor));
+
+        // [TAG_ALLOC_SIZE_EXPAND]
+        // if you hit this assert, update ggml_backend_op_alloc_size_may_expand() accordingly
+        GGML_ASSERT(size <= ggml_nbytes(tensor) ||
+                    ggml_op_is_empty(tensor->op) ||
+                    ggml_backend_op_alloc_size_may_expand(tensor->op));
+
        return size;
    }
    return ggml_nbytes(tensor);
@@ -2101,6 +2108,22 @@ ggml_backend_t ggml_backend_sched_get_tensor_backend(ggml_backend_sched_t sched,

 // utils

+// [TAG_ALLOC_SIZE_EXPAND]
+// returns true for ops that may require additional memory for fleeting data on some backends,
+// i.e. the backend's get_alloc_size may return more than ggml_nbytes for the output tensor
+bool ggml_backend_op_alloc_size_may_expand(enum ggml_op op) {
+    switch (op) {
+        case GGML_OP_FLASH_ATTN_EXT:
+        case GGML_OP_MUL_MAT_ID:
+        case GGML_OP_CUMSUM:
+        case GGML_OP_ARGSORT:
+        case GGML_OP_TOP_K:
+            return true;
+        default:
+            return false;
+    }
+}
+
 enum ggml_status ggml_backend_view_init(struct ggml_tensor * tensor) {
    GGML_ASSERT(tensor);
    GGML_ASSERT(tensor->buffer == NULL);
@@ -69,30 +69,15 @@ 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         physical_idx = 0;
+    int         virtual_idx  = 0;
+    int         domain_id    = 0;
+    std::string domain_name;
    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
@@ -361,7 +346,6 @@ 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;
@@ -376,9 +360,6 @@ 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;
@@ -392,12 +373,12 @@ struct ggml_hexagon_session {

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

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

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

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

    void enqueue_op(const htp_opnode & node);
@@ -430,14 +411,38 @@ 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_hexagon_session * sess) {
-        this->sess = sess;
-        this->name = name;
+    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_hexagon_session * sess;
-    std::string            name;
+    ggml_backend_hexagon_device_context * dev_ctx;
+    std::string                           name;
 };

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

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

 static void ggml_backend_hexagon_buffer_free_buffer(ggml_backend_buffer_t buffer) {
@@ -1494,24 +1500,26 @@ 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 sess = static_cast<ggml_backend_hexagon_buffer_type_context *>(buffer_type->context)->sess;
+    auto dev_ctx = static_cast<ggml_backend_hexagon_buffer_type_context *>(buffer_type->context)->dev_ctx;
+    auto sess    = dev_ctx->session();
    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", sess->c_name(), exc.what());
+        GGML_LOG_ERROR("ggml-hex: %s failed to allocate device buffer context: %s\n", dev_ctx->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 sess = static_cast<ggml_backend_hexagon_buffer_type_context *>(buffer_type->context)->sess;
+    auto dev_ctx = static_cast<ggml_backend_hexagon_buffer_type_context *>(buffer_type->context)->dev_ctx;
+    auto sess    = dev_ctx->session();
    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", sess->c_name(), exc.what());
+        GGML_LOG_ERROR("ggml-hex: %s failed to allocate host buffer context: %s\n", dev_ctx->c_name(), exc.what());
        return nullptr;
    }
 }
@@ -1536,7 +1544,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->sess->max_bufsize;
+    return context->dev_ctx->max_bufsize;
 }

 static bool ggml_backend_hexagon_buffer_type_is_host(ggml_backend_buffer_type_t buft) {
@@ -1567,6 +1575,22 @@ 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;
 }
@@ -2811,8 +2835,7 @@ 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(int dev_id) noexcept(false) {
-    const auto & config = opt_device_configs[dev_id];
+void ggml_hexagon_session::allocate(const ggml_hexagon_device_config & config) noexcept(false) {
    int phys_idx = config.physical_idx;
    int virt_idx = config.virtual_idx;

@@ -2823,21 +2846,31 @@ void ggml_hexagon_session::allocate(int dev_id) noexcept(false) {

    this->phys_idx   = phys_idx;
    this->virt_idx   = virt_idx;
-    this->domain_id  = get_domain_id(phys_idx);
+    this->domain_id  = config.domain_id;
    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());

-    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)");
+    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");
    }

-    std::string dom_name = get_domain_name(phys_idx);
+    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)");
+        }
+    }

    // Create new session if virtual_idx > 0
    if (virt_idx > 0) {
@@ -2849,7 +2882,8 @@ void ggml_hexagon_session::allocate(int dev_id) noexcept(false) {

        int err = remote_session_control(FASTRPC_RESERVE_NEW_SESSION, (void *) &n, sizeof(n));
        if (err != AEE_SUCCESS) {
-            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);
+            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);
            throw std::runtime_error("ggml-hex: remote_session_control(new-sess) failed (see log for details)");
        }

@@ -2857,9 +2891,20 @@ void ggml_hexagon_session::allocate(int dev_id) noexcept(false) {
        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;

-    // Get session URI
+        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);
+        }
+    }

    char session_uri[256];
    {
@@ -2877,31 +2922,18 @@ void ggml_hexagon_session::allocate(int dev_id) noexcept(false) {

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

-            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)");
+            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);
        }
    }

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

@@ -2991,7 +3023,7 @@ void ggml_hexagon_session::allocate(int dev_id) noexcept(false) {
    this->op_batch = new ggml_hexagon_opbatch(this, opt_opbatch, this->max_vmem);

    // Start dspqueue/opbatch processing
-    err = htp_iface_start(this->handle, dev_id, this->queue_id, opt_nhvx, opt_nhmx, this->max_vmem);
+    err = htp_iface_start(this->handle, this->session_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)");
@@ -3054,33 +3086,23 @@ void ggml_hexagon_session::release() noexcept(true) {
    this->cloned_buffers.clear();
 }

-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;
-
+ggml_hexagon_session::ggml_hexagon_session(const ggml_hexagon_device_config & config, ggml_backend_dev_t dev) noexcept(false) {
    op_batch = nullptr;
    op_queue = nullptr;
    fence_seq = ((uintptr_t)this) & 0xFFFF;

    try {
-        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);
+        allocate(config);
    } 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
@@ -3957,11 +3979,13 @@ 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 && t->buffer->buft == &sess->host_buffer_type;
+    return t && t->buffer && ggml_backend_buft_is_host(t->buffer->buft);
+    GGML_UNUSED(sess);
 }

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

 static bool ggml_hexagon_supported_mul_mat(const struct ggml_hexagon_session * sess, const struct ggml_tensor * dst) {
@@ -5269,7 +5293,8 @@ 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 sess = static_cast<ggml_hexagon_session *>(dev->context);
+    auto dev_ctx = static_cast<ggml_backend_hexagon_device_context *>(dev->context);
+    auto sess    = dev_ctx->session();

    return new ggml_backend{
        /* .guid      = */ ggml_backend_hexagon_guid(),
@@ -5282,8 +5307,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 sess = static_cast<ggml_hexagon_session *>(dev->context);
-    return sess->c_name();
+    auto dev_ctx = static_cast<ggml_backend_hexagon_device_context *>(dev->context);
+    return dev_ctx->c_name();

    GGML_UNUSED(dev);
 }
@@ -5321,16 +5346,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 sess = static_cast<ggml_hexagon_session *>(dev->context);
-    return &sess->buffer_type;
+    auto dev_ctx = static_cast<ggml_backend_hexagon_device_context *>(dev->context);
+    return &dev_ctx->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 sess = static_cast<ggml_hexagon_session *>(dev->context);
-    return &sess->host_buffer_type;
+    auto dev_ctx = static_cast<ggml_backend_hexagon_device_context *>(dev->context);
+    return &dev_ctx->host_buffer_type;
 }

 static bool ggml_hexagon_supported_cpy(const struct ggml_hexagon_session * sess, const struct ggml_tensor * op) {
@@ -5421,7 +5446,8 @@ 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 sess = static_cast<ggml_hexagon_session *>(dev->context);
+    auto dev_ctx = static_cast<ggml_backend_hexagon_device_context *>(dev->context);
+    auto sess    = dev_ctx->session();

    // reject ops that match the filter
    if (opt_opfilter && std::regex_match(ggml_op_desc(op), *opt_opfilter)) {
@@ -5493,6 +5519,7 @@ 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;
@@ -5505,6 +5532,7 @@ 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;
@@ -5590,17 +5618,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 sess = static_cast<ggml_hexagon_session *>(dev->context);
+    auto dev_ctx = static_cast<ggml_backend_hexagon_device_context *>(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 == &sess->host_buffer_type) || (buft == &sess->buffer_type);
+    bool supp = (buft == &dev_ctx->host_buffer_type) || (buft == &dev_ctx->buffer_type);
 #endif

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

@@ -5629,16 +5657,11 @@ 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 / sessions
+    // Create devices
    for (size_t i = 0; i < opt_ndev; 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;
-        }
+        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]);
    }

 }
@@ -5646,10 +5669,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 / sessions
+    // Release devices
    for (size_t i = 0; i < opt_ndev; i++) {
-        auto sess = static_cast<ggml_hexagon_session *>(devices[i].context);
-        delete sess;
+        auto dev_ctx = static_cast<ggml_backend_hexagon_device_context *>(devices[i].context);
+        delete dev_ctx;
    }
 }

@@ -5818,6 +5841,85 @@ 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,
@@ -5983,6 +6085,9 @@ 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,6 +73,7 @@ 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
@@ -99,6 +100,7 @@ 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
@@ -206,6 +208,13 @@ 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) {
@@ -367,6 +376,7 @@ 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,6 +116,8 @@ 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,10 +1,27 @@
 #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,6 +47,10 @@ 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,6 +3,7 @@
 #import "ggml-impl.h"
 #import "ggml-backend-impl.h"
 #import "ggml-metal-impl.h"
+#import "ggml-metal-common.h"

 #include <Foundation/Foundation.h>

@@ -788,6 +789,10 @@ 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];
 }

@@ -1410,6 +1415,30 @@ 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;
@@ -1713,9 +1742,15 @@ 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 has_simdgroup_reduction && op->src[0]->type != GGML_TYPE_NVFP4;
+            return ggml_metal_supports_mul_mat_op(
+                    has_simdgroup_reduction, op, false,
+                    ggml_metal_op_mul_mat_id_use_mm(op, has_simdgroup_mm));
        case GGML_OP_SET:
        case GGML_OP_CPY:
        case GGML_OP_DUP:
@@ -2362,10 +2362,6 @@ 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) &&
@@ -2468,12 +2464,7 @@ 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_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) {
+    } else if (ggml_metal_op_mul_mat_use_mm(op, props_dev->has_simdgroup_mm)) {
        //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
@@ -2622,13 +2613,7 @@ int ggml_metal_op_mul_mat_id(ggml_metal_op_t ctx, int idx) {
    const uint32_t r2 = 1;
    const uint32_t r3 = 1;

-    // 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)) {
+    if (ggml_metal_op_mul_mat_id_use_mm(op, props_dev->has_simdgroup_mm)) {
        // 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,6 +720,218 @@ 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 } },
@@ -812,6 +1024,204 @@ 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,10 +34,14 @@ 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,6 +8,7 @@
 #include <cstdlib>
 #include <cstring>
 #include <string>
+#include <dlfcn.h>
 #include <poll.h>
 #include <sys/socket.h>
 #include <unistd.h>
@@ -114,16 +115,9 @@ struct apple_rdma::impl {

    ~impl() {
        broken = true;
-        // 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);
-        }
+        // 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);
        if (send_mr) ibv_dereg_mr(send_mr);
        if (recv_mr) ibv_dereg_mr(recv_mr);
        free(send_mem);
@@ -184,11 +178,28 @@ 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,7 +2402,138 @@ 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,
@@ -2410,98 +2541,107 @@ 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 * k), cgh);
-        sycl::local_accessor<int, 1> shared_idx(sycl::range<1>(block_size * k), 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);

        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;

-                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;
-                    }
-                }
+                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);
            });
    });
 }
@@ -2902,7 +3042,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(src0_dd, dst_dd, ncols, nrows, k, main_stream);
+    top_k_f32_sycl(ctx, 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.53.1"
+HTTPLIB_VERSION = "refs/tags/v0.54.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) {
+    if (k_rot && k_rot->buffer) {
        kv_self->set_input_k_rot(k_rot);
    }
 }
@@ -8,8 +8,8 @@
 #ifndef CPPHTTPLIB_HTTPLIB_H
 #define CPPHTTPLIB_HTTPLIB_H

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

 #ifdef _WIN32
 #if defined(_WIN32_WINNT) && _WIN32_WINNT < 0x0A00
@@ -134,6 +134,16 @@
 #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
@@ -1429,9 +1439,16 @@ public:
  DataSink &operator=(DataSink &&) = delete;

  std::function<bool(const char *data, size_t data_len)> write;
-  std::function<bool()> is_writable;
-  std::function<void()> done;
-  std::function<void(const Headers &trailer)> done_with_trailer;
+
+  // 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::ostream os;

 private:
@@ -1516,7 +1533,10 @@ 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());
-      if (n == 0) { break; }
+      // 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 (!sink.write(buf, n)) { return false; }
      length -= n;
    }
@@ -1723,6 +1743,14 @@ 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;
@@ -1788,6 +1816,11 @@ 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 {
@@ -1827,6 +1860,7 @@ enum class Error {
  InvalidRangeHeader,
  UnsupportedContentEncoding,
  WebSocketHandshake,
+  UserCallbackException,

  // For internal use only
  SSLPeerCouldBeClosed_,
@@ -2020,6 +2054,10 @@ 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,
@@ -2107,6 +2145,17 @@ 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);
@@ -2174,6 +2223,10 @@ 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(
@@ -2202,6 +2255,35 @@ 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_;
@@ -2215,6 +2297,11 @@ 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;
@@ -2226,9 +2313,21 @@ 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,
@@ -2259,6 +2358,10 @@ 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,
@@ -2292,6 +2395,10 @@ 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;
@@ -2313,6 +2420,7 @@ private:
  Handlers delete_handlers_;
  HandlersForContentReader delete_handlers_for_content_reader_;
  Handlers options_handlers_;
+  CustomHandlers custom_handlers_;

  struct WebSocketHandlerEntry {
    std::unique_ptr<detail::MatcherBase> matcher;
@@ -3500,6 +3608,16 @@ 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,
@@ -3520,6 +3638,9 @@ 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);
@@ -3534,11 +3655,13 @@ 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);

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

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

@@ -4318,6 +4441,11 @@ 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_;
Author	SHA1	Message	Date
Georgi GerganovandGitHub	73f56d105b	ggml : add ggml_backend_op_alloc_size_may_expand, use it in RPC (#27960 ) some backends (Metal, SYCL, WebGPU) require additional memory for fleeting data for certain ops, which is reflected in their get_alloc_size implementations. add ggml_backend_op_alloc_size_may_expand() to the backend utils, listing these ops, and assert in ggml_backend_buft_get_alloc_size that a backend expanding the alloc size of a compute op only does so for ops listed in the helper. use the helper in the RPC backend to decide whether to query the remote server for the actual alloc size, instead of a hardcoded list. Assisted-by: pi:llama.cpp/Qwen3.8-27B	2026-08-30 09:17:47 +03:00
Ryan CandGitHub	742347b2e7	rpc: fix apple rdma error spew on teardown (#27908 )	2026-08-30 09:16:26 +03:00
Nils GladitzandGitHub	093adb242e	metal: add fa-vec tunings for M3 Ultra (#27999 )	2026-08-30 09:06:29 +03:00
Daya AdiantoandGitHub	b8b743c3c1	metal : Add fa-vec tuning for M3 Pro (#27963 ) Related issue: #27668	2026-08-30 09:02:22 +03:00
Alessandro de Oliveira Faria (A.K.A.CABELO)andGitHub	dc7aecf70d	vendor : update cpp-httplib to 0.54.0 (#27919 ) * vendor : update cpp-httplib to 0.54.0 * vendor : update cpp-httplib to 0.54.0 and 0.54.1	2026-08-30 09:01:51 +03:00
Ryan CandGitHub	2bf0415152	rpc : fix pre-rdma macOS versions (#27815 )	2026-08-30 08:59:25 +03:00
Max Krasnyansky GitHub Marco Colombo Alexander Lu Ehsan Bateni	9e54e687cb	hexagon: support for device discovery and create sessions on demand (#27785 ) * hex-devices: add support for lazy session allocation and cleanup dev interfaces Co-authored-by: Marco Colombo <mcolombo@qti.qualcomm.com> * hex-devices: support for runtime discovery of available NPU cores Co-authored-by: Alexander Lu <alexlu@qti.qualcomm.com> Co-authored-by: Ehsan Bateni <ebateni@qti.qualcomm.com> * hex-devices: reject non-existing devices early during init --------- Co-authored-by: Marco Colombo <mcolombo@qti.qualcomm.com> Co-authored-by: Alexander Lu <alexlu@qti.qualcomm.com> Co-authored-by: Ehsan Bateni <ebateni@qti.qualcomm.com>	2026-08-30 08:57:55 +03:00
TitaniumtownandGitHub	370cb12e8b	sycl: split long rows in TOP_K instead of one work-group per row (#27847 )	2026-08-30 08:57:08 +03:00
QuintinShawandGitHub	d882575cc8	metal : fix null-pipeline crash for F16 src1 mul_mat/mul_mat_id (#25648 ) * metal : fail closed on mul_mat shapes with missing F16 kernels * metal : abort on nil pipeline in encoder_set_pipeline * metal : address review comments * metal : share mul_mat mm dispatch with supports_op	2026-08-30 08:56:35 +03:00
fairydreaming GitHub Stanisław Szymczyk AesSedai	bdf3955159	memory : copy Hadamard matrix to k_rot tensor only if it has buffer assigned to prevent crashes during context shift of unquantized K cache (#27967 ) Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com> Co-authored-by: AesSedai <7980540+AesSedai@users.noreply.github.com>	2026-08-30 07:47:15 +02:00