patches/0030-neo-pipeline.patch

From: opencoti
Subject: [PATCH 0030] F5 M3 — NEO: asymmetric GPU/CPU attention pipelining

Sits on top of 0020 (M2 head-residency split). When --neo-pipeline != off and
the headinfer split is active for the layer, build_attn_mha builds TWO
ggml_flash_attn_ext ops on disjoint Q-head ranges instead of M2's
concat-then-stream-back path: cur_g consumes the GPU half of K/V on the GPU
backend; cur_c consumes the CPU half of K/V on the CPU backend (auto-routed
by the scheduler from each FA op's K/V buffer-type). Output halves are
ggml_concat'd on the head axis — same head-axis concat M2 used, but on
O(n_tokens) FA outputs instead of O(n_kv) K/V.

A small orchestrator (ggml/include/ggml-neo-pipeline.h + ggml/src/
ggml-neo-pipeline.cpp) lets the graph builder register each (cur_g, cur_c)
pair; the CUDA backend's graph_compute hook routes registered cur_g to a
non-default CUDA stream (curr_stream_no=1) so concurrent CPU FA on cur_c can
overlap with stream-1 GPU FA. A cudaEvent is recorded on stream 1 after the
GPU FA dispatch; when a downstream split's graph reads cur_g (typically the
concat join), the hook inserts a cudaStreamWaitEvent on the consuming stream
so the read sees completed data. CUDA graph capture is disabled for
NEO-engaged splits; off-path graphs keep the fast cuda_graph path.

Default off in every layer: cparams.neo_pipeline_mode == 0 means
build_attn_mha_neo is never called and the M3-A structural split never
engages — the M2 concat-then-stream-back path runs unchanged, byte-identical
to the 0020-shipped binary. Adapter exposes --neo-pipeline off|on|auto via
common/arg.cpp; @opencoti/llamafile typed config field
neoPipelineMode + OPENCOTI_LLAMAFILE_NEO_PIPELINE env mirror the M2
template.

Solidpc RTX 3090 + Qwen2.5-Coder-0.5B IQ4_XS bench (perf/llamafile/
neo-pipeline.bench.ts) at ctx=4096 / predict=128 / frac=0.5:

  B baseline (no split)        326.69 tps
  M m2-only  (split, M2 path)  186.96 tps
  N neo-on   (split, M3 path)  187.04 tps  (+0.04% vs M)  ← C2 FAIL
  O off-id   (--neo-pipeline off) 151.65 tps

  C1 PASS  N completion === B completion (correctness intact)
  C2 FAIL  N tps - M tps < 5% (NEO overlap mechanism is correct but the
           per-step GPU FA on half-heads finishes in much less wall time
           than the CPU FA on half-heads, so even perfect overlap is
           bounded by the slower CPU half — for this model + ctx,
           M3-B's stream-swap has no observable win over M3-A's
           structural split alone)
  C3 PASS  N GPU KV (24 MiB) ≈ M GPU KV (24 MiB) — residency preserved
  C4 PASS  O completion === M completion (--neo-pipeline off byte-identical)

Per the M3-D gate in docs/features/advanced_kv.md (F5 M3): C2 failed but
the orchestrator is correct and the off-path is byte-identical to M2,
so the orchestrator file ships as a wired-but-no-observable-win
mechanism (kept for M4 ScoutAttention reuse) rather than as a
dormant stub. The M3-A structural win (per-step CPU→GPU stream-back
eliminated) is the actual deliverable. M3-C EWMA load-aware controller
is deferred — without an observable M3-B overlap win it would be
cosmetic.

Surgical-hook count stays 18 — all 0030 changes are in vendored-source
files, captured here via snapshot-diff (NOT git diff HEAD; bug-121).

diff --git a/llama.cpp/BUILD.mk b/llama.cpp/BUILD.mk
--- a/llama.cpp/BUILD.mk
+++ b/llama.cpp/BUILD.mk
@@ -27,6 +27,7 @@ GGML_SRCS_CPP := \
 	llama.cpp/ggml/src/ggml-backend-meta.cpp \
 	llama.cpp/ggml/src/ggml-backend-reg.cpp \
 	llama.cpp/ggml/src/ggml-backend.cpp \
+	llama.cpp/ggml/src/ggml-neo-pipeline.cpp \
 	llama.cpp/ggml/src/ggml-opt.cpp \
 	llama.cpp/ggml/src/ggml-threading.cpp \
 	llama.cpp/ggml/src/ggml.cpp \
diff --git a/llama.cpp/common/arg.cpp b/llama.cpp/common/arg.cpp
--- a/llama.cpp/common/arg.cpp
+++ b/llama.cpp/common/arg.cpp
@@ -1429,6 +1429,17 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
             params.headinfer_gpu_heads_frac = std::stof(value);
         }
     ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_HEADINFER_GPU_HEADS_FRAC"));
+    // opencoti F5 M3 neo — see docs/features/advanced_kv.md
+    add_opt(common_arg(
+        {"--neo-pipeline"}, "MODE",
+        "neo (asymmetric GPU/CPU attention pipelining): off (default), on (engage when headinfer split is active for the layer), auto (reserved, same as on). Requires --headinfer-gpu-heads-frac < 1.0 to do anything.",
+        [](common_params & params, const std::string & value) {
+            if      (value == "off")  params.neo_pipeline_mode = 0;
+            else if (value == "on")   params.neo_pipeline_mode = 1;
+            else if (value == "auto") params.neo_pipeline_mode = 2;
+            else throw std::invalid_argument("--neo-pipeline must be off|on|auto");
+        }
+    ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_NEO_PIPELINE"));
     add_opt(common_arg(
         {"--chunks"}, "N",
         string_format("max number of chunks to process (default: %d, -1 = all)", params.n_chunks),
diff --git a/llama.cpp/common/common.cpp b/llama.cpp/common/common.cpp
--- a/llama.cpp/common/common.cpp
+++ b/llama.cpp/common/common.cpp
@@ -1634,6 +1634,8 @@ struct llama_context_params common_context_params_to_llama(const common_params &
     cparams.slot_shrink_idle_ms = params.slot_shrink_idle_ms < 0 ? 0u : (uint32_t) params.slot_shrink_idle_ms;
     // opencoti F5 M2 headinfer — see docs/features/advanced_kv.md
     cparams.headinfer_gpu_heads_frac = params.headinfer_gpu_heads_frac;
+    // opencoti F5 M3 neo — see docs/features/advanced_kv.md
+    cparams.neo_pipeline_mode = params.neo_pipeline_mode;
     cparams.n_rs_seq          = params.speculative.need_n_rs_seq();
     cparams.n_batch           = params.n_batch;
     cparams.n_ubatch          = params.n_ubatch;
diff --git a/llama.cpp/common/common.h b/llama.cpp/common/common.h
--- a/llama.cpp/common/common.h
+++ b/llama.cpp/common/common.h
@@ -440,6 +440,11 @@ struct common_params {
     // memory and streamed back per step. 1.0 = off (no split). Only active for
     // GPU-offloaded layers with flash-attention (non-transposed V cache).
     float   headinfer_gpu_heads_frac = 1.0f;
+    // opencoti F5 M3 neo — see docs/features/advanced_kv.md
+    // 0 = off (default; M2 concat path runs), 1 = on (engage when split
+    // active for the layer), 2 = auto (reserved for future load-based
+    // engagement; same as `on` today).
+    uint32_t neo_pipeline_mode    =     0;
     int32_t n_batch               =  2048; // logical batch size for prompt processing (must be >=32 to use BLAS)
     int32_t n_ubatch              =   512; // physical batch size for prompt processing (must be >=32 to use BLAS)
     int32_t n_keep                =     0; // number of tokens to keep from initial prompt
diff --git a/llama.cpp/ggml/include/ggml-neo-pipeline.h b/llama.cpp/ggml/include/ggml-neo-pipeline.h
new file mode 100644
--- /dev/null
+++ b/llama.cpp/ggml/include/ggml-neo-pipeline.h
@@ -0,0 +1,78 @@
+// opencoti F5 M3 NEO — asymmetric GPU/CPU attention pipelining.
+//
+// This header declares the orchestrator surface that lets the M3 two-flash-attn
+// structural split overlap its GPU and CPU halves. The graph builder registers
+// each (cur_g, cur_c) pair after constructing it; the CUDA backend consults the
+// registry at compute time to (a) dispatch GPU FA on a non-default CUDA stream
+// so the scheduler's stream-0 sync no longer blocks the CPU FA's input copy
+// and (b) wait on the recorded stream-1 event when a downstream split consumes
+// cur_g (notably the concat that joins both halves), keeping correctness intact.
+//
+// Default-off: ggml_neo_pipeline_set_enabled(false) makes register_pair a no-op
+// and the CUDA hooks fast-skip. The off-path is byte-identical to the M3-A
+// shipped path; M3-A in turn is byte-identical to M2's concat path when
+// --neo-pipeline off (cparams.neo_pipeline_mode == 0). See
+// docs/features/advanced_kv.md (F5 M3 NEO).
+#pragma once
+
+#include "ggml.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+// Global on/off switch. Off by default; the graph builder flips this on for
+// the duration of a graph build when cparams.neo_pipeline_mode != 0 and at
+// least one layer has an active headinfer split. While off, all functions in
+// this file are O(1) no-ops.
+GGML_API void ggml_neo_pipeline_set_enabled(bool enabled);
+GGML_API bool ggml_neo_pipeline_enabled(void);
+
+// Registers a NEO FA pair. Called from build_attn_mha_neo after constructing
+// the two flash-attention ops on disjoint head ranges. cur_g lives on the GPU
+// backend (CUDA-allocated K/V), cur_c on the CPU backend (CPU-allocated K/V).
+// The CUDA backend later consults this registry to route cur_g to a non-default
+// stream so cur_c can dispatch concurrently. When disabled, this is a no-op.
+GGML_API void ggml_neo_pipeline_register_pair(const struct ggml_tensor * cur_g,
+                                              const struct ggml_tensor * cur_c);
+
+// Membership probe. Returns true iff t was the cur_g argument of a prior
+// register_pair call in this graph. O(1) fast-path; the CUDA backend calls
+// this once per node when scanning a graph it is about to dispatch.
+GGML_API bool ggml_neo_pipeline_is_gpu_partner(const struct ggml_tensor * t);
+
+// CUDA event slot for cur_g. The CUDA backend creates a cudaEvent (with
+// cudaEventCreateWithFlags + cudaEventDisableTiming) the first time it
+// dispatches cur_g on a non-default stream and stores the opaque handle via
+// set_event. A later split whose graph reads cur_g calls get_event and
+// inserts a cudaStreamWaitEvent so the consuming stream waits for the
+// stream-1 work. Both functions are no-ops when NEO is disabled or t isn't
+// registered (get_event returns NULL). The void* is treated as cudaEvent_t
+// at the CUDA call site; the orchestrator does not know the underlying type.
+GGML_API void * ggml_neo_pipeline_get_event(const struct ggml_tensor * cur_g);
+GGML_API void   ggml_neo_pipeline_set_event(const struct ggml_tensor * cur_g, void * event);
+
+// Clears all per-graph registrations. Called by the CUDA backend at the end
+// of each graph dispatch so registrations don't leak across decode steps.
+// (The graph builder reinstalls them on the next build.) Notably does NOT
+// destroy stored events — those are destroyed by the CUDA backend, which
+// owns the cudaEvent lifecycle (call destroy_all_events at shutdown).
+GGML_API void ggml_neo_pipeline_clear_registrations(void);
+
+// Walks every (cur_g, event) entry, invoking destroyer(event) for each
+// non-null event, then clears the event map. Called at session/backend
+// shutdown so cudaEvents are properly released. destroyer is typically a
+// thin lambda wrapping cudaEventDestroy at the CUDA call site.
+GGML_API void ggml_neo_pipeline_destroy_all_events(void (*destroyer)(void * event));
+
+// Counter of NEO pair dispatches observed by the CUDA backend in the current
+// process. Surfaced for the M3-E bench's residency / engagement checks.
+GGML_API int  ggml_neo_pipeline_dispatch_count(void);
+
+// Increments the dispatch counter. Called by the CUDA backend each time it
+// observes and routes a registered cur_g to a non-default stream.
+GGML_API void ggml_neo_pipeline_inc_dispatch_count(void);
+
+#ifdef __cplusplus
+}
+#endif
diff --git a/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu b/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu
--- a/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -2,6 +2,10 @@
 #include "ggml-impl.h"
 #include "ggml-backend-impl.h"
 
+// opencoti F5 M3-B NEO orchestrator (registry of cur_g/cur_c pairs for the
+// asymmetric attention pipelining; consulted in graph_compute below).
+#include "ggml-neo-pipeline.h"
+
 #include "ggml-cuda/allreduce.cuh"
 #include "ggml-cuda/common.cuh"
 #include "ggml-cuda/acc.cuh"
@@ -4485,6 +4489,58 @@ static enum ggml_status GGML_CALL ggml_backend_cuda_graph_compute(ggml_backend_t
 
     ggml_cuda_set_device(cuda_ctx->device);
 
+    // opencoti F5 M3-B NEO: scan the graph for registered cur_g flash-attn
+    // outputs and for downstream consumers of prior cur_g. The orchestrator
+    // is dormant unless the graph builder called set_enabled(true) (which
+    // only happens when --neo-pipeline != off AND at least one layer has
+    // an active headinfer split). When dormant this loop body is one
+    // atomic-load fast-skip.
+    //
+    // When NEO is engaged for this split:
+    //   - If the split computes a cur_g, swap curr_stream_no to 1 so the
+    //     kernels dispatch on the alternate CUDA stream; the scheduler's
+    //     stream-0 sync in compute_splits will NOT block on this work, so
+    //     the subsequent split's q_c host copy starts immediately and the
+    //     CPU FA on cur_c can overlap with our stream-1 work. We record an
+    //     event on stream 1 at exit so downstream consumers know when
+    //     cur_g is done.
+    //   - If the split reads a prior cur_g (typically the concat join),
+    //     stream-wait on the recorded event before kernels read it.
+    //
+    // CUDA graph capture is disabled when NEO is active for this split:
+    // capturing kernel dispatch on a non-default stream interacts with the
+    // graph_key cache and the existing concurrent_event re-ordering pass
+    // in ways we haven't validated. Off-path graphs keep the cuda_graph
+    // fast path; only NEO-engaged splits pay this cost.
+    bool          neo_active           = false;
+    ggml_tensor * neo_cur_g_to_record  = nullptr;
+    if (ggml_neo_pipeline_enabled()) {
+        for (int i = 0; i < cgraph->n_nodes; ++i) {
+            ggml_tensor * node = cgraph->nodes[i];
+            if (ggml_neo_pipeline_is_gpu_partner(node)) {
+                neo_active          = true;
+                neo_cur_g_to_record = node;
+                break;
+            }
+            for (int s = 0; s < GGML_MAX_SRC; ++s) {
+                if (node->src[s] != nullptr && ggml_neo_pipeline_is_gpu_partner(node->src[s])) {
+                    void * ev = ggml_neo_pipeline_get_event(node->src[s]);
+                    if (ev != nullptr) {
+                        CUDA_CHECK(cudaStreamWaitEvent(cuda_ctx->stream(), (cudaEvent_t) ev, 0));
+                    }
+                }
+            }
+        }
+    }
+    if (neo_active) {
+        // Materialize stream 1 (creates it on first use) and switch the
+        // context to dispatch all subsequent kernels there. Both calls are
+        // cheap: stream() is a vector lookup + lazy create; the curr_stream_no
+        // write is a single int.
+        (void) cuda_ctx->stream(cuda_ctx->device, 1);
+        cuda_ctx->curr_stream_no = 1;
+    }
+
     bool use_cuda_graph             = false;
     bool cuda_graph_update_required = false;
     const void * graph_key = nullptr;
@@ -4524,6 +4580,16 @@ static enum ggml_status GGML_CALL ggml_backend_cuda_graph_compute(ggml_backend_t
     }
 #endif // USE_CUDA_GRAPH
 
+    // opencoti F5 M3-B NEO: force-disable cuda_graph capture for NEO-active
+    // splits. Capturing on a non-default stream is structurally fine in CUDA
+    // but interacts with the existing graph_key cache + concurrent_event
+    // re-ordering pass in ways we haven't validated. Off-path graphs keep
+    // the fast cuda_graph path.
+    if (neo_active) {
+        use_cuda_graph             = false;
+        cuda_graph_update_required = false;
+    }
+
     if (use_cuda_graph && cuda_graph_update_required) {
         // Start CUDA graph capture
         {
@@ -4536,6 +4602,23 @@ static enum ggml_status GGML_CALL ggml_backend_cuda_graph_compute(ggml_backend_t
 
     ggml_cuda_graph_evaluate_and_capture(cuda_ctx, cgraph, use_cuda_graph, cuda_graph_update_required, graph_key);
 
+    // opencoti F5 M3-B NEO exit: record an event on stream 1 marking cur_g's
+    // completion so the downstream concat (consumer split) can stream-wait
+    // on it from stream 0. Restore curr_stream_no=0 so the next graph_compute
+    // call starts on the default stream as normal.
+    if (neo_active && neo_cur_g_to_record != nullptr) {
+        void * ev = ggml_neo_pipeline_get_event(neo_cur_g_to_record);
+        if (ev == nullptr) {
+            cudaEvent_t new_ev;
+            CUDA_CHECK(cudaEventCreateWithFlags(&new_ev, cudaEventDisableTiming));
+            ev = (void *) new_ev;
+            ggml_neo_pipeline_set_event(neo_cur_g_to_record, ev);
+        }
+        CUDA_CHECK(cudaEventRecord((cudaEvent_t) ev, cuda_ctx->stream(cuda_ctx->device, 1)));
+        ggml_neo_pipeline_inc_dispatch_count();
+        cuda_ctx->curr_stream_no = 0;
+    }
+
     return GGML_STATUS_SUCCESS;
 }
 
diff --git a/llama.cpp/ggml/src/ggml-neo-pipeline.cpp b/llama.cpp/ggml/src/ggml-neo-pipeline.cpp
new file mode 100644
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-neo-pipeline.cpp
@@ -0,0 +1,126 @@
+// opencoti F5 M3 NEO orchestrator implementation.
+//
+// Backend-agnostic registry that lets the CUDA backend recognize the M3-built
+// (cur_g, cur_c) flash-attention pairs and route cur_g to a non-default CUDA
+// stream so cur_c can dispatch on the CPU backend concurrently.
+//
+// The orchestrator stores tensor pointers and opaque cudaEvent handles. It
+// does not call any CUDA API directly — the CUDA backend creates events with
+// cudaEventCreateWithFlags and hands them to set_event; on shutdown the CUDA
+// backend supplies a destroyer callback that wraps cudaEventDestroy.
+//
+// All maps are protected by a single mutex. Lookups are O(1) (unordered_map).
+// While the global enabled flag is false, every public function fast-skips
+// without touching the maps, so the off-path is one atomic-load + branch.
+//
+// See ggml/include/ggml-neo-pipeline.h for the public API contract and
+// docs/features/advanced_kv.md (F5 M3) for the surrounding design.
+
+#include "ggml-neo-pipeline.h"
+
+#include <atomic>
+#include <mutex>
+#include <unordered_map>
+#include <unordered_set>
+
+namespace {
+
+std::atomic<bool> g_enabled{false};
+std::atomic<int>  g_dispatch_count{0};
+
+std::mutex g_mu;
+
+// Registered cur_g tensors. Membership probe is O(1).
+std::unordered_set<const ggml_tensor *> g_gpu_partners;
+
+// Map from cur_g → cudaEvent_t handle (opaque void *). NULL slot means the
+// CUDA backend hasn't created an event for this tensor yet — get_event
+// returns nullptr in that case so the caller knows to create-and-set.
+std::unordered_map<const ggml_tensor *, void *> g_events;
+
+} // namespace
+
+void ggml_neo_pipeline_set_enabled(bool enabled) {
+    g_enabled.store(enabled, std::memory_order_release);
+}
+
+bool ggml_neo_pipeline_enabled(void) {
+    return g_enabled.load(std::memory_order_acquire);
+}
+
+void ggml_neo_pipeline_register_pair(const ggml_tensor * cur_g, const ggml_tensor * cur_c) {
+    (void) cur_c; // cpu partner is informational; the orchestrator only acts on cur_g
+    if (!ggml_neo_pipeline_enabled()) {
+        return;
+    }
+    if (cur_g == nullptr) {
+        return;
+    }
+    std::lock_guard<std::mutex> lock(g_mu);
+    g_gpu_partners.insert(cur_g);
+}
+
+bool ggml_neo_pipeline_is_gpu_partner(const ggml_tensor * t) {
+    if (!ggml_neo_pipeline_enabled()) {
+        return false;
+    }
+    if (t == nullptr) {
+        return false;
+    }
+    std::lock_guard<std::mutex> lock(g_mu);
+    return g_gpu_partners.find(t) != g_gpu_partners.end();
+}
+
+void * ggml_neo_pipeline_get_event(const ggml_tensor * cur_g) {
+    if (!ggml_neo_pipeline_enabled()) {
+        return nullptr;
+    }
+    if (cur_g == nullptr) {
+        return nullptr;
+    }
+    std::lock_guard<std::mutex> lock(g_mu);
+    const auto it = g_events.find(cur_g);
+    if (it == g_events.end()) {
+        return nullptr;
+    }
+    return it->second;
+}
+
+void ggml_neo_pipeline_set_event(const ggml_tensor * cur_g, void * event) {
+    if (!ggml_neo_pipeline_enabled()) {
+        return;
+    }
+    if (cur_g == nullptr) {
+        return;
+    }
+    std::lock_guard<std::mutex> lock(g_mu);
+    g_events[cur_g] = event;
+}
+
+void ggml_neo_pipeline_clear_registrations(void) {
+    std::lock_guard<std::mutex> lock(g_mu);
+    g_gpu_partners.clear();
+    // Note: g_events deliberately NOT cleared here. Events are reusable
+    // across decode steps as long as the underlying tensor pointers stay
+    // stable. They're cleaned up by destroy_all_events on shutdown.
+}
+
+void ggml_neo_pipeline_destroy_all_events(void (*destroyer)(void * event)) {
+    std::lock_guard<std::mutex> lock(g_mu);
+    if (destroyer != nullptr) {
+        for (auto & kv : g_events) {
+            if (kv.second != nullptr) {
+                destroyer(kv.second);
+            }
+        }
+    }
+    g_events.clear();
+}
+
+int ggml_neo_pipeline_dispatch_count(void) {
+    return g_dispatch_count.load(std::memory_order_acquire);
+}
+
+void ggml_neo_pipeline_inc_dispatch_count(void) {
+    g_dispatch_count.fetch_add(1, std::memory_order_relaxed);
+}
diff --git a/llama.cpp/include/llama.h b/llama.cpp/include/llama.h
--- a/llama.cpp/include/llama.h
+++ b/llama.cpp/include/llama.h
@@ -358,6 +358,10 @@ extern "C" {
         // streamed back per step). 1.0 = off. Only effective for offloaded
         // layers with flash-attention (non-transposed V cache).
         float    headinfer_gpu_heads_frac;
+        // opencoti F5 M3 neo — see docs/features/advanced_kv.md
+        // 0 = off (default), 1 = on, 2 = auto. Effective only when the
+        // M2 head split is active for a given layer.
+        uint32_t neo_pipeline_mode;
         uint32_t n_rs_seq;          // number of recurrent-state snapshots per seq for rollback (0 = no rollback) [EXPERIMENTAL]
         int32_t  n_threads;         // number of threads to use for generation
         int32_t  n_threads_batch;   // number of threads to use for batch processing
diff --git a/llama.cpp/src/llama-context.cpp b/llama.cpp/src/llama-context.cpp
--- a/llama.cpp/src/llama-context.cpp
+++ b/llama.cpp/src/llama-context.cpp
@@ -64,6 +64,8 @@ llama_context::llama_context(
     cparams.slot_shrink_idle_ms = params.slot_shrink_idle_ms;
     // opencoti F5 M2 headinfer — see docs/features/advanced_kv.md
     cparams.headinfer_gpu_heads_frac = params.headinfer_gpu_heads_frac;
+    // opencoti F5 M3 neo — see docs/features/advanced_kv.md
+    cparams.neo_pipeline_mode = params.neo_pipeline_mode;
 
     cparams.n_threads        = params.n_threads;
     cparams.n_threads_batch  = params.n_threads_batch;
@@ -3356,6 +3358,7 @@ llama_context_params llama_context_default_params() {
         /*.slot_initial_ctx            =*/ 0,
         /*.slot_shrink_idle_ms         =*/ 0,
         /*.headinfer_gpu_heads_frac    =*/ 1.0f,
+        /*.neo_pipeline_mode            =*/ 0,  // opencoti F5 M3 neo — off by default
         /*.n_rs_seq                    =*/ 0,
         /*.n_threads                   =*/ GGML_DEFAULT_N_THREADS, // TODO: better default
         /*.n_threads_batch             =*/ GGML_DEFAULT_N_THREADS,
diff --git a/llama.cpp/src/llama-cparams.h b/llama.cpp/src/llama-cparams.h
--- a/llama.cpp/src/llama-cparams.h
+++ b/llama.cpp/src/llama-cparams.h
@@ -23,6 +23,15 @@ struct llama_cparams {
     uint32_t slot_shrink_idle_ms;
     // opencoti F5 M2 headinfer — see docs/features/advanced_kv.md
     float    headinfer_gpu_heads_frac;
+    // opencoti F5 M3 neo — asymmetric GPU/CPU pipelining of attention.
+    // 0 = off (M2 concat path runs; binary equivalent to upstream when
+    //         headinfer_gpu_heads_frac == 1.0).
+    // 1 = on  (engage when headinfer split is active for the layer:
+    //         build two flash_attn_ext ops on disjoint head ranges, run
+    //         them on GPU and CPU backends respectively, concat outputs).
+    // 2 = auto (same engagement as `on` for now — reserved for future
+    //         load-based engagement heuristics).
+    uint32_t neo_pipeline_mode;
     uint32_t n_rs_seq;        // number of recurrent-state snapshots per seq for rollback
     int32_t  n_threads;       // number of threads to use for generation
     int32_t  n_threads_batch; // number of threads to use for batch processing
diff --git a/llama.cpp/src/llama-graph.cpp b/llama.cpp/src/llama-graph.cpp
--- a/llama.cpp/src/llama-graph.cpp
+++ b/llama.cpp/src/llama-graph.cpp
@@ -11,6 +11,10 @@
 #include "llama-memory-hybrid-iswa.h"
 #include "llama-memory-recurrent.h"
 
+// opencoti F5 M3 NEO orchestrator (registers cur_g/cur_c pairs so the CUDA
+// backend can route GPU FA to a non-default stream and overlap with CPU FA).
+#include "ggml-neo-pipeline.h"
+
 #include <cassert>
 #include <cmath>
 #include <cstring>
@@ -2085,6 +2089,140 @@ ggml_tensor * llm_graph_context::build_attn_mha(
     return cur;
 }
 
+// opencoti F5 M3 neo — head-axis-split attention. Builds two flash_attn_ext
+// ops on disjoint head ranges:
+//   GPU side: q[..., :gpu_heads*gqa, ...] × k_g × v_g  (auto-routes to GPU
+//             via the scheduler because k_g/v_g are views on the layer's
+//             GPU buffer-type tensor — see ggml-backend.cpp:797-852).
+//   CPU side: q[..., gpu_heads*gqa:, ...] × k_c × v_c  (auto-routes to CPU
+//             because k_c/v_c are views on the CPU buffer-type tensor that
+//             M2's headinfer allocates alongside the GPU half — see
+//             llama-kv-cache.cpp:get_k_cpu/get_v_cpu).
+// Outputs are concatenated on the head axis and fed into the standard
+// reshape_2d → Wo flow. This eliminates M2's per-step CPU→GPU stream-back
+// (O(n_kv) bytes per layer) — only the small O(n_tokens) CPU FA output
+// is bounced to GPU for the concat. Pre-M3-B, the two FA ops run serially
+// under the scheduler's default dispatch; M3-B's orchestrator will overlap
+// them concurrently.
+ggml_tensor * llm_graph_context::build_attn_mha_neo(
+        ggml_tensor * q,
+        ggml_tensor * k_g,
+        ggml_tensor * v_g,
+        ggml_tensor * k_c,
+        ggml_tensor * v_c,
+            uint32_t  gpu_heads,
+        ggml_tensor * kq_b,
+        ggml_tensor * kq_mask,
+        ggml_tensor * sinks,
+        ggml_tensor * v_mla,
+              float   kq_scale,
+                int   il) const {
+    // Preconditions enforced by the M2 split gate that produced k_c/v_c.
+    GGML_ASSERT(cparams.flash_attn && "NEO path requires flash-attn");
+    GGML_ASSERT(kq_b == nullptr     && "NEO path does not support KQ bias");
+    GGML_ASSERT(v_mla == nullptr    && "NEO path does not support MLA");
+    GGML_ASSERT(gpu_heads > 0       && "NEO path requires an active head split");
+    GGML_ASSERT(k_g && v_g && k_c && v_c);
+
+    // K/V layout is identical between the GPU and CPU halves (non-transposed
+    // V per the M2 gate). Stream axis matches across both — both are views
+    // on tensors allocated with the same n_stream at construction.
+    GGML_ASSERT(k_g->ne[3] == k_c->ne[3]);
+    GGML_ASSERT(v_g->ne[3] == v_c->ne[3]);
+    GGML_ASSERT(k_g->nb[1] <= k_g->nb[2] && "headinfer split requires non-transposed V");
+
+    // The full Q has all n_head_q heads. Heads are split on KV-head-group
+    // boundaries (M2's GQA clamp ensures gpu_heads is a whole number of
+    // groups). gqa_ratio = n_head_q / n_head_kv; n_head_kv = gpu_heads +
+    // cpu_heads. Q-heads serving the same KV head are CONTIGUOUS in Q's
+    // dim 1 under standard GGUF, so the slice is a clean view (no extra
+    // permute beyond the one already in build_attn_mha).
+    const uint32_t n_head_q   = (uint32_t) q->ne[1];
+    const uint32_t n_head_kv  = (uint32_t) (k_g->ne[1] + k_c->ne[1]);
+    GGML_ASSERT(n_head_q % n_head_kv == 0 && "GQA ratio must divide n_head_q");
+    const uint32_t gqa_ratio  = n_head_q / n_head_kv;
+    const uint32_t gpu_q_heads = gpu_heads * gqa_ratio;
+    const uint32_t cpu_q_heads = (n_head_kv - gpu_heads) * gqa_ratio;
+    GGML_ASSERT(gpu_q_heads + cpu_q_heads == n_head_q);
+
+    // Mirror build_attn_mha's n_stream reshape + permute on Q, but slice Q
+    // on the head axis BEFORE the permute so the dim-1 slice stays
+    // contiguous in the input layout.
+    const auto n_stream = k_g->ne[3];
+    ggml_tensor * q_4 = ggml_view_4d(ctx0, q,
+            q->ne[0], q->ne[1], q->ne[2]/n_stream, n_stream,
+            q->nb[1], q->nb[2], q->nb[3]/n_stream, 0);
+
+    ggml_tensor * q_g = ggml_view_4d(ctx0, q_4,
+            q_4->ne[0], gpu_q_heads, q_4->ne[2], q_4->ne[3],
+            q_4->nb[1], q_4->nb[2], q_4->nb[3],
+            /*offset=*/ 0);
+    ggml_tensor * q_c = ggml_view_4d(ctx0, q_4,
+            q_4->ne[0], cpu_q_heads, q_4->ne[2], q_4->ne[3],
+            q_4->nb[1], q_4->nb[2], q_4->nb[3],
+            /*offset=*/ (size_t) gpu_q_heads * q_4->nb[1]);
+
+    q_g = ggml_permute(ctx0, q_g, 0, 2, 1, 3);
+    q_c = ggml_permute(ctx0, q_c, 0, 2, 1, 3);
+    ggml_tensor * k_g_p = ggml_permute(ctx0, k_g, 0, 2, 1, 3);
+    ggml_tensor * v_g_p = ggml_permute(ctx0, v_g, 0, 2, 1, 3);
+    ggml_tensor * k_c_p = ggml_permute(ctx0, k_c, 0, 2, 1, 3);
+    ggml_tensor * v_c_p = ggml_permute(ctx0, v_c, 0, 2, 1, 3);
+
+    // F32 K → F16 cast mirrors build_attn_mha lines 1881-1887 (the embedding
+    // model non-causal path). v_trans is asserted false above, so no extra
+    // transpose is needed; V already has the same dim-1 head layout as K.
+    if (k_g_p->type == GGML_TYPE_F32) k_g_p = ggml_cast(ctx0, k_g_p, GGML_TYPE_F16);
+    if (k_c_p->type == GGML_TYPE_F32) k_c_p = ggml_cast(ctx0, k_c_p, GGML_TYPE_F16);
+    if (v_g_p->type == GGML_TYPE_F32) v_g_p = ggml_cast(ctx0, v_g_p, GGML_TYPE_F16);
+    if (v_c_p->type == GGML_TYPE_F32) v_c_p = ggml_cast(ctx0, v_c_p, GGML_TYPE_F16);
+
+    // Sinks (if present) carry one F32 entry per Q-head; slice them to
+    // match each half. The flash_attn_ext_add_sinks op asserts
+    // sinks->ne[0] == q->src[0]->ne[2] (the q-head axis after permute).
+    ggml_tensor * sinks_g = nullptr;
+    ggml_tensor * sinks_c = nullptr;
+    if (sinks) {
+        GGML_ASSERT((uint32_t) sinks->ne[0] == n_head_q);
+        sinks_g = ggml_view_1d(ctx0, sinks, gpu_q_heads, 0);
+        sinks_c = ggml_view_1d(ctx0, sinks, cpu_q_heads,
+                               (size_t) gpu_q_heads * sinks->nb[0]);
+    }
+
+    const float alibi_bias  = hparams.f_max_alibi_bias;
+    const float logit_cap   = hparams.attn_soft_cap ? hparams.f_attn_logit_softcapping : 0.0f;
+
+    ggml_tensor * cur_g = ggml_flash_attn_ext(ctx0, q_g, k_g_p, v_g_p, kq_mask, kq_scale, alibi_bias, logit_cap);
+    cb(cur_g, LLAMA_TENSOR_NAME_FATTN "_g", il);
+    ggml_flash_attn_ext_add_sinks(cur_g, sinks_g);
+    ggml_flash_attn_ext_set_prec (cur_g, GGML_PREC_F32);
+
+    ggml_tensor * cur_c = ggml_flash_attn_ext(ctx0, q_c, k_c_p, v_c_p, kq_mask, kq_scale, alibi_bias, logit_cap);
+    cb(cur_c, LLAMA_TENSOR_NAME_FATTN "_c", il);
+    ggml_flash_attn_ext_add_sinks(cur_c, sinks_c);
+    ggml_flash_attn_ext_set_prec (cur_c, GGML_PREC_F32);
+
+    // opencoti F5 M3-B: enable + register the pair so the CUDA backend swaps
+    // cur_g to a non-default stream at dispatch time, letting cur_c (on CPU)
+    // overlap. set_enabled(true) is idempotent and only flipped here — the
+    // call site already gates on neo_pipeline_mode != 0 && headinfer_split_active,
+    // so reaching this point means NEO is engaged for this layer.
+    ggml_neo_pipeline_set_enabled(true);
+    ggml_neo_pipeline_register_pair(cur_g, cur_c);
+
+    // FA output shape is [v_dim, n_q_heads_split, n_tokens, n_stream].
+    // Concat on dim 1 (q-head axis) to reassemble the full Q-head set.
+    // The scheduler synchronizes here — cur_c gets lifted to GPU so the
+    // concat runs on whichever backend hosts cur_g (typically GPU).
+    ggml_tensor * cur = ggml_concat(ctx0, cur_g, cur_c, 1);
+    cb(cur, LLAMA_TENSOR_NAME_FATTN, il);
+
+    // Match build_attn_mha's flash-attn final reshape (line 1913).
+    cur = ggml_reshape_2d(ctx0, cur, cur->ne[0]*cur->ne[1], cur->ne[2]*cur->ne[3]);
+
+    return cur;
+}
+
 llm_graph_input_attn_no_cache * llm_graph_context::build_attn_inp_no_cache() const {
     auto inp = std::make_unique<llm_graph_input_attn_no_cache>(hparams, cparams);
 
@@ -2237,10 +2375,29 @@ ggml_tensor * llm_graph_context::build_attn(
     const auto & kq_mask = inp->get_kq_mask();
 
     ggml_tensor * q = q_cur;
-    ggml_tensor * k = mctx_cur->get_k(ctx0, il);
-    ggml_tensor * v = mctx_cur->get_v(ctx0, il);
 
-    ggml_tensor * cur = build_attn_mha(q, k, v, kq_b, kq_mask, sinks, v_mla, kq_scale, il);
+    // opencoti F5 M3 neo — head-axis-split attention. When the M2 head
+    // residency split is active for this layer AND --neo-pipeline is on,
+    // build two flash_attn_ext ops on disjoint head ranges instead of the
+    // M2 concat-reassembly path (which streams k_cpu/v_cpu back to GPU each
+    // step and runs a single full-head FA op). The scheduler routes each
+    // FA op to the backend matching its K/V buffer-type. v_mla must be
+    // null (build_attn for KV asserts this above at :2102), kq_b too;
+    // the M2 gate ensures non-transposed V. See docs/features/advanced_kv.md.
+    ggml_tensor * cur;
+    if (cparams.neo_pipeline_mode != 0 && mctx_cur->headinfer_split_active(il)) {
+        ggml_tensor * k_g = mctx_cur->get_k_gpu(ctx0, il);
+        ggml_tensor * v_g = mctx_cur->get_v_gpu(ctx0, il);
+        ggml_tensor * k_c = mctx_cur->get_k_cpu(ctx0, il);
+        ggml_tensor * v_c = mctx_cur->get_v_cpu(ctx0, il);
+        cur = build_attn_mha_neo(q, k_g, v_g, k_c, v_c,
+                                 mctx_cur->headinfer_gpu_heads(il),
+                                 kq_b, kq_mask, sinks, v_mla, kq_scale, il);
+    } else {
+        ggml_tensor * k = mctx_cur->get_k(ctx0, il);
+        ggml_tensor * v = mctx_cur->get_v(ctx0, il);
+        cur = build_attn_mha(q, k, v, kq_b, kq_mask, sinks, v_mla, kq_scale, il);
+    }
     cb(cur, "kqv_out", il);
 
     if (inp->self_v_rot) {
diff --git a/llama.cpp/src/llama-graph.h b/llama.cpp/src/llama-graph.h
--- a/llama.cpp/src/llama-graph.h
+++ b/llama.cpp/src/llama-graph.h
@@ -908,6 +908,27 @@ struct llm_graph_context {
                   float   kq_scale,
                     int   il) const;
 
+    // opencoti F5 M3 neo — head-axis split flash-attn. Builds two
+    // ggml_flash_attn_ext ops on disjoint head ranges (GPU vs CPU subsets,
+    // routed automatically by the scheduler via the K/V tensors' buffer
+    // types) and concatenates their outputs on the head axis before the
+    // reshape/Wo projection. Flash-attention only (assert kq_b == nullptr,
+    // v_mla == nullptr, !v_trans for V — the M2 split gate already enforces
+    // this). Output shape matches build_attn_mha's flash-attn branch.
+    ggml_tensor * build_attn_mha_neo(
+            ggml_tensor * q,       // [n_embd_head_q, n_head_q, n_tokens]
+            ggml_tensor * k_g,     // GPU-resident K subset (gpu_heads on dim 1)
+            ggml_tensor * v_g,     // GPU-resident V subset (gpu_heads on dim 1)
+            ggml_tensor * k_c,     // CPU-resident K subset (cpu_heads on dim 1)
+            ggml_tensor * v_c,     // CPU-resident V subset (cpu_heads on dim 1)
+                uint32_t  gpu_heads,
+            ggml_tensor * kq_b,    // must be nullptr
+            ggml_tensor * kq_mask,
+            ggml_tensor * sinks,   // [n_head_q] or nullptr
+            ggml_tensor * v_mla,   // must be nullptr
+                  float   kq_scale,
+                    int   il) const;
+
     llm_graph_input_attn_no_cache * build_attn_inp_no_cache() const;
 
     ggml_tensor * build_attn(
diff --git a/llama.cpp/src/llama-kv-cache.cpp b/llama.cpp/src/llama-kv-cache.cpp
--- a/llama.cpp/src/llama-kv-cache.cpp
+++ b/llama.cpp/src/llama-kv-cache.cpp
@@ -1810,6 +1810,116 @@ ggml_tensor * llama_kv_cache::get_v(ggml_context * ctx, int32_t il, uint32_t n_k
             ggml_row_size(v->type, kv_size*n_embd_v_gqa)*sinfo.s0);
 }
 
+// opencoti F5 M3 neo — head-split-aware accessors. These factor the per-half
+// view machinery that get_k/get_v perform internally before their concat,
+// exposing the GPU-only and CPU-only sub-tensors as independent ggml views.
+// M3 NEO uses them to build two flash_attn_ext ops (one per backend, scheduler
+// auto-routes by buffer type) and merge their outputs, eliminating M2's
+// per-step O(n_kv) CPU→GPU stream-back. Dimensions and strides are identical
+// to the corresponding halves inside M2's get_k/get_v.
+
+bool llama_kv_cache::headinfer_split_active(int32_t il) const {
+    const auto it = map_layer_ids.find(il);
+    if (it == map_layer_ids.end()) {
+        return false;
+    }
+    const auto & layer = layers[it->second];
+    return layer.gpu_heads > 0 && layer.k_cpu != nullptr;
+}
+
+uint32_t llama_kv_cache::headinfer_gpu_heads(int32_t il) const {
+    const auto it = map_layer_ids.find(il);
+    if (it == map_layer_ids.end()) {
+        return 0;
+    }
+    return layers[it->second].gpu_heads;
+}
+
+ggml_tensor * llama_kv_cache::get_k_gpu(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const {
+    const int32_t ikv = map_layer_ids.at(il);
+    const auto & layer = layers[ikv];
+    GGML_ASSERT(layer.gpu_heads > 0 && "get_k_gpu called without an active headinfer split");
+
+    auto * k = layer.k;
+    const uint64_t kv_size = get_size();
+    const uint32_t ns      = sinfo.s1 - sinfo.s0 + 1;
+    const uint32_t n_embd_head_k = hparams.n_embd_head_k(il);
+    const uint32_t n_head_gpu    = layer.gpu_heads;
+    const uint64_t n_embd_gpu    = (uint64_t) n_head_gpu * n_embd_head_k;
+
+    return ggml_view_4d(ctx, k,
+            n_embd_head_k, n_head_gpu, n_kv, ns,
+            ggml_row_size(k->type, n_embd_head_k),
+            ggml_row_size(k->type, n_embd_gpu),
+            ggml_row_size(k->type, n_embd_gpu*kv_size),
+            ggml_row_size(k->type, n_embd_gpu*kv_size)*sinfo.s0);
+}
+
+ggml_tensor * llama_kv_cache::get_k_cpu(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const {
+    const int32_t ikv = map_layer_ids.at(il);
+    const auto & layer = layers[ikv];
+    GGML_ASSERT(layer.gpu_heads > 0 && layer.k_cpu && "get_k_cpu called without an active headinfer split");
+
+    auto * kc = layer.k_cpu;
+    const uint64_t kv_size = get_size();
+    const uint32_t ns      = sinfo.s1 - sinfo.s0 + 1;
+    const uint32_t n_embd_head_k = hparams.n_embd_head_k(il);
+    const uint32_t n_head_cpu    = hparams.n_head_kv(il) - layer.gpu_heads;
+    const uint64_t n_embd_cpu    = (uint64_t) n_head_cpu * n_embd_head_k;
+
+    return ggml_view_4d(ctx, kc,
+            n_embd_head_k, n_head_cpu, n_kv, ns,
+            ggml_row_size(kc->type, n_embd_head_k),
+            ggml_row_size(kc->type, n_embd_cpu),
+            ggml_row_size(kc->type, n_embd_cpu*kv_size),
+            ggml_row_size(kc->type, n_embd_cpu*kv_size)*sinfo.s0);
+}
+
+ggml_tensor * llama_kv_cache::get_v_gpu(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const {
+    const int32_t ikv = map_layer_ids.at(il);
+    const auto & layer = layers[ikv];
+    GGML_ASSERT(layer.gpu_heads > 0 && "get_v_gpu called without an active headinfer split");
+    // M2's gate is `!v_trans && n_stream == 1` — when split is active the
+    // V layout is always non-transposed (matches get_v's non-split branch
+    // shape modulo head count). Same head-dim-1 stride pattern as K.
+    GGML_ASSERT(!v_trans && "headinfer split requires non-transposed V");
+
+    auto * v = layer.v;
+    const uint64_t kv_size = get_size();
+    const uint32_t ns      = sinfo.s1 - sinfo.s0 + 1;
+    const uint32_t n_embd_head_v = hparams.n_embd_head_v(il);
+    const uint32_t n_head_gpu    = layer.gpu_heads;
+    const uint64_t n_embd_gpu    = (uint64_t) n_head_gpu * n_embd_head_v;
+
+    return ggml_view_4d(ctx, v,
+            n_embd_head_v, n_head_gpu, n_kv, ns,
+            ggml_row_size(v->type, n_embd_head_v),
+            ggml_row_size(v->type, n_embd_gpu),
+            ggml_row_size(v->type, n_embd_gpu*kv_size),
+            ggml_row_size(v->type, n_embd_gpu*kv_size)*sinfo.s0);
+}
+
+ggml_tensor * llama_kv_cache::get_v_cpu(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const {
+    const int32_t ikv = map_layer_ids.at(il);
+    const auto & layer = layers[ikv];
+    GGML_ASSERT(layer.gpu_heads > 0 && layer.v_cpu && "get_v_cpu called without an active headinfer split");
+    GGML_ASSERT(!v_trans && "headinfer split requires non-transposed V");
+
+    auto * vc = layer.v_cpu;
+    const uint64_t kv_size = get_size();
+    const uint32_t ns      = sinfo.s1 - sinfo.s0 + 1;
+    const uint32_t n_embd_head_v = hparams.n_embd_head_v(il);
+    const uint32_t n_head_cpu    = hparams.n_head_kv(il) - layer.gpu_heads;
+    const uint64_t n_embd_cpu    = (uint64_t) n_head_cpu * n_embd_head_v;
+
+    return ggml_view_4d(ctx, vc,
+            n_embd_head_v, n_head_cpu, n_kv, ns,
+            ggml_row_size(vc->type, n_embd_head_v),
+            ggml_row_size(vc->type, n_embd_cpu),
+            ggml_row_size(vc->type, n_embd_cpu*kv_size),
+            ggml_row_size(vc->type, n_embd_cpu*kv_size)*sinfo.s0);
+}
+
 ggml_tensor * llama_kv_cache::cpy_k(ggml_context * ctx, ggml_tensor * k_cur, ggml_tensor * k_idxs, int32_t il, const slot_info & sinfo) const {
     GGML_UNUSED(sinfo);
 
@@ -3107,6 +3217,31 @@ ggml_tensor * llama_kv_cache_context::get_v(ggml_context * ctx, int32_t il) cons
     return kv->get_v(ctx, il, n_kv, sinfos[i_cur]);
 }
 
+// opencoti F5 M3 neo — per-batch wrappers around the head-split accessors.
+ggml_tensor * llama_kv_cache_context::get_k_gpu(ggml_context * ctx, int32_t il) const {
+    return kv->get_k_gpu(ctx, il, n_kv, sinfos[i_cur]);
+}
+
+ggml_tensor * llama_kv_cache_context::get_v_gpu(ggml_context * ctx, int32_t il) const {
+    return kv->get_v_gpu(ctx, il, n_kv, sinfos[i_cur]);
+}
+
+ggml_tensor * llama_kv_cache_context::get_k_cpu(ggml_context * ctx, int32_t il) const {
+    return kv->get_k_cpu(ctx, il, n_kv, sinfos[i_cur]);
+}
+
+ggml_tensor * llama_kv_cache_context::get_v_cpu(ggml_context * ctx, int32_t il) const {
+    return kv->get_v_cpu(ctx, il, n_kv, sinfos[i_cur]);
+}
+
+bool llama_kv_cache_context::headinfer_split_active(int32_t il) const {
+    return kv->headinfer_split_active(il);
+}
+
+uint32_t llama_kv_cache_context::headinfer_gpu_heads(int32_t il) const {
+    return kv->headinfer_gpu_heads(il);
+}
+
 ggml_tensor * llama_kv_cache_context::cpy_k(ggml_context * ctx, ggml_tensor * k_cur, ggml_tensor * k_idxs, int32_t il) const {
     return kv->cpy_k(ctx, k_cur, k_idxs, il, sinfos[i_cur]);
 }
diff --git a/llama.cpp/src/llama-kv-cache.h b/llama.cpp/src/llama-kv-cache.h
--- a/llama.cpp/src/llama-kv-cache.h
+++ b/llama.cpp/src/llama-kv-cache.h
@@ -217,6 +217,23 @@ public:
     ggml_tensor * get_k(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const;
     ggml_tensor * get_v(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const;
 
+    // opencoti F5 M3 neo — head-split-aware K/V accessors. Return separate
+    // views onto the GPU-resident and host-resident head subsets without the
+    // concat-reassembly that get_k/get_v perform. Built on the SAME view
+    // machinery (see llama-kv-cache.cpp get_k for the dimensions); callers
+    // are expected to gate via headinfer_split_active(il) and to build their
+    // own combiner (M3 NEO builds two flash_attn_ext ops + a head-axis concat
+    // on the small attention outputs, eliminating M2's O(n_kv) stream-back).
+    // For non-split layers these return the equivalent of the un-split view
+    // for the "_gpu" half and nullptr for "_cpu" — callers must check
+    // headinfer_split_active first.
+    ggml_tensor * get_k_gpu(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const;
+    ggml_tensor * get_v_gpu(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const;
+    ggml_tensor * get_k_cpu(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const;
+    ggml_tensor * get_v_cpu(ggml_context * ctx, int32_t il, uint32_t n_kv, const slot_info & sinfo) const;
+    bool          headinfer_split_active(int32_t il) const;
+    uint32_t      headinfer_gpu_heads   (int32_t il) const;
+
     // store k_cur and v_cur in the cache based on the provided head location
     ggml_tensor * cpy_k(ggml_context * ctx, ggml_tensor * k_cur, ggml_tensor * k_idxs, int32_t il, const slot_info & sinfo) const;
     ggml_tensor * cpy_v(ggml_context * ctx, ggml_tensor * v_cur, ggml_tensor * v_idxs, int32_t il, const slot_info & sinfo) const;
@@ -444,6 +461,16 @@ public:
     ggml_tensor * get_k(ggml_context * ctx, int32_t il) const;
     ggml_tensor * get_v(ggml_context * ctx, int32_t il) const;
 
+    // opencoti F5 M3 neo — head-split-aware K/V accessors. Per-batch wrappers
+    // that delegate to the underlying llama_kv_cache with the current n_kv /
+    // slot info. See llama-kv-cache.h for semantics.
+    ggml_tensor * get_k_gpu(ggml_context * ctx, int32_t il) const;
+    ggml_tensor * get_v_gpu(ggml_context * ctx, int32_t il) const;
+    ggml_tensor * get_k_cpu(ggml_context * ctx, int32_t il) const;
+    ggml_tensor * get_v_cpu(ggml_context * ctx, int32_t il) const;
+    bool          headinfer_split_active(int32_t il) const;
+    uint32_t      headinfer_gpu_heads   (int32_t il) const;
+
     // store k_cur and v_cur in the cache based on the provided head location
     // note: the heads in k_cur and v_cur should be laid out contiguously in memory
     //   - k_cur  [n_embd_head_k, n_head_k, n_tokens]
diff --git a/llamafile/build-functions.sh b/llamafile/build-functions.sh
--- a/llamafile/build-functions.sh
+++ b/llamafile/build-functions.sh
@@ -218,10 +218,14 @@ compile_ggml_core() {
     local llama_cpp_dir="$1"
     local build_dir="$2"
 
+    # opencoti F5 M3-B: ggml-neo-pipeline.cpp added so the CUDA backend's
+    # graph_compute hook can call the orchestrator API (registered cur_g
+    # detection + event slot get/set + dispatch-count increment).
     local ggml_core_sources="\
         $llama_cpp_dir/ggml/src/ggml.c \
         $llama_cpp_dir/ggml/src/ggml-alloc.c \
         $llama_cpp_dir/ggml/src/ggml-backend.cpp \
+        $llama_cpp_dir/ggml/src/ggml-neo-pipeline.cpp \
         $llama_cpp_dir/ggml/src/ggml-backend-meta.cpp \
         $llama_cpp_dir/ggml/src/ggml-quants.c \
         $llama_cpp_dir/ggml/src/ggml-threading.cpp"