Layr-Labs/mlx-c
diff:
ignored:
+580
-3
+299
-0
This is an overview of the changes in Layr-Labs/mlx-c,
a fork of ml-explore/mlx-c.
MLX C is the C boundary between the MLX core and every other language binding. Layr-Labs
runs an LLM/VLM serving stack on Apple silicon — continuous batching, paged KV caches,
MoE and quantized kernels — built as mlx-swift-lm
on mlx-swift, which links a forked
mlx core through this fork. Every change here is a
carry-patch: a thin C wrapper over a function the forked core added, so Swift can reach
it. The upstream API and ABI are left untouched; the fork only adds entry points, and pins
the core it builds against to the Layr-Labs revision that provides them.
Downstream consumers are named per section so a change here can be traced to the Swift symbol that needs it.
Core pin
+2
-2
CMakeLists.txt fetches the MLX core from Layr-Labs/mlx at an exact commit instead of
an upstream release tag. The C functions below wrap core functions that only exist on the
fork, so the pin is what makes them link; it is bumped together with the header changes
each time the core gains a new entry point.
diff --git ml-explore/mlx-c/CMakeLists.txt Layr-Labs/mlx-c/CMakeLists.txt
index fcf4ea77bf2e571fbe596e224baf9045a3e43527..d9165d5867d10ca7ce1b62940f4360b98be777f1 100644
--- ml-explore/mlx-c/CMakeLists.txt
+++ Layr-Labs/mlx-c/CMakeLists.txt
@@ -34,8 +34,8 @@ find_package(MLX REQUIRED)
else()
FetchContent_Declare(
mlx
- GIT_REPOSITORY "https://github.com/ml-explore/mlx.git"
- GIT_TAG v0.31.2)
+ GIT_REPOSITORY "https://github.com/Layr-Labs/mlx.git"
+ GIT_TAG 76a96e291fa9b1db946f8d5b7d1091110d672cac)
FetchContent_MakeAvailable(mlx)
endif()
Memory: allocator observability and footprint policy
+111
-1
The Swift admission code decides whether a request fits by asking the allocator, and the upstream C API only offered independent counters. The fork adds, in order of arrival:
mlx_get_num_resources/mlx_get_resource_limit— the live Metal buffer count and its hard ceiling, the quantity behind the[metal::malloc] Resource limit exceededcrash (consumed asMemory.numResources/Memory.resourceLimitin mlx-swift).mlx_get_memory_snapshot— one coherent active/cache/peak observation instead of three separate calls that can straddle an allocation.mlx_get_allocation_footprint_policyand its checked_bound/_maximum_extrahelpers — the allocator’s immutable rounding geometry (alignment, rounding threshold, minimum allocation, power-of-two cut-off, cache page size) so callers can compute the real footprint of a requested size without allocating; bound helpers return 1 on an invalid policy or overflow without raising through the error callback.mlx_get_allocation_size_upper_bound— the same bound for the current policy.
The generator override in python/mlxhooks.py keeps mlx_get_memory_snapshot intact
when the bindings are regenerated from the core headers.
diff --git ml-explore/mlx-c/mlx/c/memory.cpp Layr-Labs/mlx-c/mlx/c/memory.cpp
index fb38921d5e928cfbecf3037006d251ebe870bb54..34e21abe6eee2bdd4f8a1b9683052a599451c915 100644
--- ml-explore/mlx-c/mlx/c/memory.cpp
+++ Layr-Labs/mlx-c/mlx/c/memory.cpp
@@ -4,9 +4,9 @@ /* This file is auto-generated. Do not edit manually. */
/* */
#include "mlx/c/memory.h"
+#include "mlx/memory.h"
#include "mlx/c/error.h"
#include "mlx/c/private/mlx.h"
-#include "mlx/memory.h"
extern "C" int mlx_clear_cache(void) {
try {
@@ -44,9 +44,40 @@ return 1;
}
return 0;
}
+extern "C" int mlx_get_memory_snapshot(size_t* active, size_t* cache, size_t* peak) {
+ try {
+ auto snapshot = mlx::core::get_memory_snapshot();
+ *active = snapshot.active_memory;
+ *cache = snapshot.cache_memory;
+ *peak = snapshot.peak_memory;
+ } catch (std::exception& e) {
+ mlx_error(e.what());
+ return 1;
+ }
+ return 0;
+}
+
+extern "C" int mlx_get_num_resources(size_t* res) {
+ try {
+ *res = mlx::core::get_num_resources();
+ } catch (std::exception& e) {
+ mlx_error(e.what());
+ return 1;
+ }
+ return 0;
+}
extern "C" int mlx_get_peak_memory(size_t* res) {
try {
*res = mlx::core::get_peak_memory();
+ } catch (std::exception& e) {
+ mlx_error(e.what());
+ return 1;
+ }
+ return 0;
+}
+extern "C" int mlx_get_resource_limit(size_t* res) {
+ try {
+ *res = mlx::core::get_resource_limit();
} catch (std::exception& e) {
mlx_error(e.what());
return 1;
@@ -89,3 +120,37 @@ return 1;
}
return 0;
}
+
+extern "C" int mlx_get_allocation_size_upper_bound(size_t* res, size_t size) {
+ try {
+ *res = mlx::core::get_allocation_size_upper_bound(size);
+ } catch (std::exception& e) {
+ mlx_error(e.what());
+ return 1;
+ }
+ return 0;
+}
+
+extern "C" int mlx_get_allocation_footprint_policy(mlx_allocation_footprint_policy* res) {
+ if (!res) { return 1; }
+ auto p = mlx::core::get_allocation_footprint_policy();
+ *res = {p.alignment, p.rounding_threshold, p.minimum_allocation,
+ p.power_of_two_below, p.cache_page_size};
+ return 0;
+}
+
+static mlx::core::AllocationFootprintPolicy core_policy(
+ const mlx_allocation_footprint_policy& p) noexcept {
+ return {p.alignment, p.rounding_threshold, p.minimum_allocation,
+ p.power_of_two_below, p.cache_page_size};
+}
+
+extern "C" int mlx_allocation_footprint_policy_bound(size_t* res,
+ const mlx_allocation_footprint_policy* policy, size_t size) {
+ return res && policy && core_policy(*policy).upper_bound(size, *res) ? 0 : 1;
+}
+
+extern "C" int mlx_allocation_footprint_policy_maximum_extra(size_t* res,
+ const mlx_allocation_footprint_policy* policy) {
+ return res && policy && core_policy(*policy).maximum_extra_bytes(*res) ? 0 : 1;
+}
diff --git ml-explore/mlx-c/mlx/c/memory.h Layr-Labs/mlx-c/mlx/c/memory.h
index bae9e08ec392561a46ae50f995bd7a9294648d01..ab52b1a9416762a586b501efc3e16e8d17f14a4c 100644
--- ml-explore/mlx-c/mlx/c/memory.h
+++ Layr-Labs/mlx-c/mlx/c/memory.h
@@ -30,9 +30,31 @@ /**@{*/
int mlx_clear_cache(void);
int mlx_get_active_memory(size_t* res);
+int mlx_get_allocation_size_upper_bound(size_t* res, size_t size);
+
+/* Immutable allocator geometry. Bound functions return 1 on invalid policy or
+ * overflow without invoking the error callback or allocating an exception. */
+typedef struct mlx_allocation_footprint_policy {
+ size_t alignment;
+ size_t rounding_threshold;
+ size_t minimum_allocation;
+ size_t power_of_two_below;
+ size_t cache_page_size;
+} mlx_allocation_footprint_policy;
+
+int mlx_get_allocation_footprint_policy(mlx_allocation_footprint_policy* res);
+int mlx_allocation_footprint_policy_bound(size_t* res,
+ const mlx_allocation_footprint_policy* policy, size_t size);
+int mlx_allocation_footprint_policy_maximum_extra(size_t* res,
+ const mlx_allocation_footprint_policy* policy);
+
int mlx_get_cache_memory(size_t* res);
int mlx_get_memory_limit(size_t* res);
+int mlx_get_memory_snapshot(size_t* active, size_t* cache, size_t* peak);
+
+int mlx_get_num_resources(size_t* res);
int mlx_get_peak_memory(size_t* res);
+int mlx_get_resource_limit(size_t* res);
int mlx_reset_peak_memory(void);
int mlx_set_cache_limit(size_t* res, size_t limit);
int mlx_set_memory_limit(size_t* res, size_t limit);
diff --git ml-explore/mlx-c/python/mlxhooks.py Layr-Labs/mlx-c/python/mlxhooks.py
index 4fc3698667953535b9e1d37c2fecc39965e2e879..590f3c9d41c9b20bfaacc361429236aff5c6ecc2 100644
--- ml-explore/mlx-c/python/mlxhooks.py
+++ Layr-Labs/mlx-c/python/mlxhooks.py
@@ -506,3 +506,26 @@ }"""
)
pass
return True
+
+
+def mlx_get_memory_snapshot(f, implementation):
+ if not implementation:
+ print(
+ "int mlx_get_memory_snapshot(size_t* active, size_t* cache, size_t* peak);"
+ )
+ else:
+ print(
+ '''\
+extern "C" int mlx_get_memory_snapshot(size_t* active, size_t* cache, size_t* peak) {
+ try {
+ auto snapshot = mlx::core::get_memory_snapshot();
+ *active = snapshot.active_memory;
+ *cache = snapshot.cache_memory;
+ *peak = snapshot.peak_memory;
+ } catch (std::exception& e) {
+ mlx_error(e.what());
+ return 1;
+ }
+ return 0;
+}'''
+ )
Array: buffer metadata and constant-cache identity
+61
-0
Two non-evaluating inspections the quantized-inference path in mlx-swift relies on:
mlx_array_get_buffer_info— backing metadata (allocated bytes, data offset, element count, row-contiguity, uniqueness) for an already evaluated array. Unavailable backing is reported asavailable=falserather than forcing an eval or a stream sync; views may report shared backing and the fields confer no ownership._mlx_array_constant_cache_identity— the backing descriptor’s identity plus a transform-safety predicate, so a Swift constant cache can reuse an unchanged constant’s exact cast across calls without mistaking a mutated array for its former contents. It refuses compile/autodiff tracing and retained graphs, and never evaluates, hashes, or reads data back. Marked internal: the caller must retain the source descriptor while the identity is cached.
diff --git ml-explore/mlx-c/mlx/c/array.cpp Layr-Labs/mlx-c/mlx/c/array.cpp
index 7c7342d3c2318abf2020acf56d8f7c2e6c2fb364..a74d1aa00118fb3aacfcbd0e49d9402b0bdefd13 100644
--- ml-explore/mlx-c/mlx/c/array.cpp
+++ Layr-Labs/mlx-c/mlx/c/array.cpp
@@ -6,6 +6,7 @@ #include "mlx/c/array.h"
#include "mlx/c/error.h"
#include "mlx/c/private/mlx.h"
#include "mlx/c/string.h"
+#include "mlx/transforms_impl.h"
extern "C" size_t mlx_dtype_size(mlx_dtype dtype) {
return mlx_dtype_to_cpp(dtype).size();
@@ -659,3 +660,47 @@ return 1;
}
return 0;
}
+
+extern "C" int mlx_array_get_buffer_info(
+ bool* available, size_t* allocated_bytes, size_t* data_offset,
+ size_t* data_elements, bool* row_contiguous, bool* unique,
+ const mlx_array arr) {
+ *available = false;
+ *allocated_bytes = 0;
+ *data_offset = 0;
+ *data_elements = 0;
+ *row_contiguous = false;
+ *unique = false;
+ try {
+ const auto& value = mlx_array_get_(arr);
+ if (!value.is_available() || !value.data_shared_ptr()) {
+ return 0;
+ }
+ *allocated_bytes = value.buffer().ptr() ? value.buffer_size() : 0;
+ *data_offset = value.offset();
+ *data_elements = value.data_size();
+ *row_contiguous = value.flags().row_contiguous;
+ *unique = value.is_donatable();
+ *available = true;
+ } catch (std::exception& e) {
+ mlx_error(e.what());
+ return 1;
+ }
+ return 0;
+}
+
+// Non-evaluating identity for bounded immutable-constant caches. The caller
+// retains an array snapshot for every cached identity to prevent address ABA.
+extern "C" int _mlx_array_constant_cache_identity(
+ uintptr_t* identity, bool* can_cache, const mlx_array arr) {
+ try {
+ const auto& value = mlx_array_get_(arr);
+ *identity = value.id();
+ *can_cache = !mlx::core::detail::in_tracing() &&
+ !mlx::core::detail::retain_graph() && !value.is_tracer();
+ } catch (std::exception& e) {
+ mlx_error(e.what());
+ return 1;
+ }
+ return 0;
+}
diff --git ml-explore/mlx-c/mlx/c/array.h Layr-Labs/mlx-c/mlx/c/array.h
index a3b382bb2a5d4a338c459730f99d1b13db0acd7e..47e12646e7b6cca62bdc95ba0a0ee48ccc482283 100644
--- ml-explore/mlx-c/mlx/c/array.h
+++ Layr-Labs/mlx-c/mlx/c/array.h
@@ -387,6 +387,22 @@ * Internal function: use at your own risk.
*/
int _mlx_array_is_available(bool* res, const mlx_array arr);
+/** Non-blocking backing metadata. Unavailable arrays return available=false.
+ * Views may report shared backing; these fields do not confer ownership. */
+int mlx_array_get_buffer_info(
+ bool* available, size_t* allocated_bytes, size_t* data_offset,
+ size_t* data_elements, bool* row_contiguous, bool* unique,
+ const mlx_array arr);
+
+
+/**
+ * Backing descriptor identity and whether a constant cast may be retained.
+ * Does not evaluate the array. Retain a snapshot while using the identity.
+ * Internal function: use at your own risk.
+ */
+int _mlx_array_constant_cache_identity(
+ uintptr_t* identity, bool* can_cache, const mlx_array arr);
+
/**
* Wait on the array to be available. After this `_mlx_array_is_available`
* returns `true`. Internal function: use at your own risk.
Fast: mutable-input Metal kernels and the SDPA signature
+33
-0
mlx_fast_metal_kernel_new_mutable— a second constructor for custom Metal kernels that declares which caller-owned input buffers the kernel writes (the core’s mutable-input contract, used for in-place paged KV updates). The existing constructor and its read-only ABI are unchanged.- The scaled-dot-product-attention binding passes the
force_fusedargument MLX 0.32.2 added to the core call, asfalse, so the C signature and pre-0.32.2 behaviour are preserved.
diff --git ml-explore/mlx-c/mlx/c/fast.cpp Layr-Labs/mlx-c/mlx/c/fast.cpp
index ac17d54951b771090cd01f73a1095682dbad74a3..052de3c4598478fbbbb303c837cc14174ac13d5b 100644
--- ml-explore/mlx-c/mlx/c/fast.cpp
+++ Layr-Labs/mlx-c/mlx/c/fast.cpp
@@ -475,6 +475,28 @@ }
return {nullptr};
}
+extern "C" mlx_fast_metal_kernel mlx_fast_metal_kernel_new_mutable(
+ const char* name,
+ const mlx_vector_string input_names,
+ const mlx_vector_string output_names,
+ const char* source,
+ const char* header,
+ bool ensure_row_contiguous,
+ bool atomic_outputs,
+ const mlx_vector_string mutable_input_names) {
+ try {
+ return mlx_fast_metal_kernel({new mlx_fast_metal_kernel_cpp_(
+ mlx::core::fast::metal_kernel_with_mutable_inputs(
+ name, mlx_vector_string_get_(input_names),
+ mlx_vector_string_get_(output_names), source,
+ mlx_vector_string_get_(mutable_input_names), header,
+ ensure_row_contiguous, atomic_outputs))});
+ } catch (std::exception& e) {
+ mlx_error(e.what());
+ }
+ return {nullptr};
+}
+
inline mlx::core::fast::CustomKernelFunction& mlx_fast_metal_kernel_get_(
mlx_fast_metal_kernel d) {
if (!d.ctx) {
@@ -624,6 +646,7 @@ (mask_arr.ctx ? std::make_optional(mlx_array_get_(mask_arr))
: std::nullopt),
(sinks.ctx ? std::make_optional(mlx_array_get_(sinks))
: std::nullopt),
+ false,
mlx_stream_get_(s)));
} catch (std::exception& e) {
mlx_error(e.what());
diff --git ml-explore/mlx-c/mlx/c/fast.h Layr-Labs/mlx-c/mlx/c/fast.h
index c825d00e52ac440cf8417537bdb46a3b9b4dc363..02b458d56c702c310314ab4032f16eab3256a9d5 100644
--- ml-explore/mlx-c/mlx/c/fast.h
+++ Layr-Labs/mlx-c/mlx/c/fast.h
@@ -153,6 +153,16 @@ bool atomic_outputs);
void mlx_fast_metal_kernel_free(mlx_fast_metal_kernel cls);
+mlx_fast_metal_kernel mlx_fast_metal_kernel_new_mutable(
+ const char* name,
+ const mlx_vector_string input_names,
+ const mlx_vector_string output_names,
+ const char* source,
+ const char* header,
+ bool ensure_row_contiguous,
+ bool atomic_outputs,
+ const mlx_vector_string mutable_input_names);
+
int mlx_fast_metal_kernel_apply(
mlx_vector_array* outputs,
mlx_fast_metal_kernel cls,
Stream: cross-thread default streams
+33
-0
MLX 0.32 made the default stream’s Metal command encoder thread-local, so a graph built
on one thread and evaluated on another — Swift async/await continuations, actor
executors — aborts with There is no Stream(gpu, N) in current thread.
mlx_thread_unsafe_cpu_stream_new / mlx_thread_unsafe_gpu_stream_new expose the
core’s new_thread_unsafe_stream, whose encoder is registered globally, letting
mlx-swift restore a single process-wide default stream. “Thread-unsafe” means the caller
serializes submission; it is not locked for concurrent multi-thread use.
diff --git ml-explore/mlx-c/mlx/c/stream.cpp Layr-Labs/mlx-c/mlx/c/stream.cpp
index 2d17997b386fd11ff5f8d66b0b0e007ab439e460..f5eb067653afdf4d8956573223d29d2286206b39 100644
--- ml-explore/mlx-c/mlx/c/stream.cpp
+++ Layr-Labs/mlx-c/mlx/c/stream.cpp
@@ -116,3 +116,21 @@ mlx_error(e.what());
return mlx_stream_new_();
}
}
+extern "C" mlx_stream mlx_thread_unsafe_cpu_stream_new(void) {
+ try {
+ return mlx_stream_new_(
+ mlx::core::new_thread_unsafe_stream(mlx::core::Device::DeviceType::cpu));
+ } catch (std::exception& e) {
+ mlx_error(e.what());
+ return mlx_stream_new_();
+ }
+}
+extern "C" mlx_stream mlx_thread_unsafe_gpu_stream_new(void) {
+ try {
+ return mlx_stream_new_(
+ mlx::core::new_thread_unsafe_stream(mlx::core::Device::DeviceType::gpu));
+ } catch (std::exception& e) {
+ mlx_error(e.what());
+ return mlx_stream_new_();
+ }
+}
diff --git ml-explore/mlx-c/mlx/c/stream.h Layr-Labs/mlx-c/mlx/c/stream.h
index d5865b80616abe468f58acf39f129f024020d105..7906e65afd3ce7a596b722d254636815c6253466 100644
--- ml-explore/mlx-c/mlx/c/stream.h
+++ Layr-Labs/mlx-c/mlx/c/stream.h
@@ -79,6 +79,21 @@ * Returns the current default GPU stream.
*/
mlx_stream mlx_default_gpu_stream_new(void);
+/**
+ * Returns a new CPU stream usable from ANY thread.
+ *
+ * Unlike the per-thread default stream, this stream's command encoder is
+ * registered globally (mlx::core::new_thread_unsafe_stream), so it stays valid
+ * when work hops threads (e.g. Swift async/await continuations). The caller is
+ * responsible for not submitting to it concurrently from multiple threads.
+ */
+mlx_stream mlx_thread_unsafe_cpu_stream_new(void);
+
+/**
+ * Returns a new GPU stream usable from ANY thread (see the CPU variant).
+ */
+mlx_stream mlx_thread_unsafe_gpu_stream_new(void);
+
/**@}*/
#ifdef __cplusplus
Fork tooling
+340
-0
The gate that keeps this page honest: check_forkdiff.py fails CI when base.hash
is not the merge-base with upstream, or when a file the fork changes is not described
by a section above. FORKDIFF.md explains the day-to-day; the README links the page.
diff --git ml-explore/mlx-c/FORKDIFF.md Layr-Labs/mlx-c/FORKDIFF.md
new file mode 100644
index 0000000000000000000000000000000000000000..657d19e31395411107b98356b6af1563da0ca1b1
--- /dev/null
+++ Layr-Labs/mlx-c/FORKDIFF.md
@@ -0,0 +1,79 @@
+# Fork diff: what this fork changes, and keeping that page honest
+
+This repository is a fork of [`ml-explore/mlx-c`](https://github.com/ml-explore/mlx-c).
+Everything it changes relative to upstream is published as a browsable page:
+
+**https://layr-labs.github.io/mlx-c/**
+
+The page is rendered by [`protolambda/forkdiff`](https://github.com/protolambda/forkdiff)
+from [`fork.yaml`](fork.yaml) at the repo root, in the style of
+[op-geth's go-ethereum fork diff](https://op-geth.optimism.io/). `fork.yaml`
+groups the changed files into sections with a paragraph each, names the exact
+upstream commit the fork is rebased onto (`base.hash`), and lists files that
+are not code (`ignore`). Every push to `main` re-renders and redeploys it
+(`.github/workflows/forkdiff-pages.yml`).
+
+**One-time setup.** GitHub Pages has to be switched on by a repo admin:
+Settings → Pages → Source: **GitHub Actions**. The workflow token cannot do
+this itself. Until it's done the deploy workflow still builds the page (kept
+as a run artifact, `forkdiff-page`) and exits with a notice instead of failing.
+
+## The gate
+
+A fork-diff page is only useful while it is true, and two things make it go
+stale silently. `scripts/check_forkdiff.py` runs on every PR
+(`.github/workflows/forkdiff-check.yml`) and fails on both:
+
+| Drift | Check | Why it matters |
+|-------|-------|----------------|
+| **Rebase onto newer upstream** | `base.hash` must equal `git merge-base HEAD upstream/main` and be an ancestor of upstream `main` | After a rebase the old base makes upstream's own commits look like fork changes: files that are not ours. |
+| **New fork change nobody described** | every path in `git diff --name-only base.hash HEAD` must match a section glob or a global `ignore` | An undescribed file is a change the page can't explain. |
+| **Section describing code we no longer carry** | every glob must match at least one changed path | Stale sections are as misleading as missing ones. |
+
+The check also renders the page, so a `fork.yaml` that forkdiff itself rejects
+cannot merge. The deploy workflow runs the same check before publishing, so a
+stale analysis is never served.
+
+## Day to day
+
+**Adding or changing fork files in a PR.** If the check lists uncovered files,
+add each to the section that explains it in `fork.yaml` (or start a new
+section with a short description). Files that are not code go under the
+top-level `ignore`. Keep globs specific: a `mlx/**` catch-all would swallow
+upstream's changes after a bad rebase and defeat the gate.
+
+**Rebasing onto newer upstream.** The gate will fail with the new merge-base
+in its message:
+
+```
+git fetch https://github.com/ml-explore/mlx-c.git main:refs/remotes/upstream/main
+git merge-base HEAD refs/remotes/upstream/main # → new base.hash
+```
+
+Set `base.hash` to that value, then run the check locally and fix what it
+reports — usually files upstream absorbed (stale globs to delete) and files
+that moved (globs to rename):
+
+```
+python3 -m pip install pyyaml
+python3 scripts/check_forkdiff.py --upstream-ref refs/remotes/upstream/main
+```
+
+**Previewing the page locally** (Go 1.21+):
+
+```
+go run github.com/protolambda/forkdiff@v0.1.1 -repo . -fork fork.yaml -out tmp/index.html
+open tmp/index.html
+```
+
+## Design notes
+
+- `base.hash` is a full 40-hex commit id, never a branch name: a symbolic base
+ would move underneath the page and the gate alike.
+- Section `ignore` lists count as coverage (forkdiff still lists those files,
+ grayed out); the top-level `ignore` is for things that aren't code at all.
+- The glob semantics are forkdiff's: `*` and `?` stop at `/`, `**` spans
+ directories (and may match none), `[!x]` negates a class.
+- The check is pure git + PyYAML so the same command runs locally and in CI.
+ It is shared verbatim across the Layr-Labs MLX forks (`mlx`, `mlx-c`,
+ `mlx-swift`, `mlx-swift-lm`).
diff --git ml-explore/mlx-c/README.md Layr-Labs/mlx-c/README.md
index d57abc96e88b8a25ba4cc33f6475d85728c7baf0..e3bf56463961f9021fc9de2fa9aacf7528255b7f 100644
--- ml-explore/mlx-c/README.md
+++ Layr-Labs/mlx-c/README.md
@@ -1,5 +1,7 @@
# MLX C
+> **This is a fork.** `Layr-Labs/mlx-c` tracks [`ml-explore/mlx-c`](https://github.com/ml-explore/mlx-c) and adds the C entry points Layr-Labs' Swift inference stack needs from its forked MLX core. Everything changed relative to upstream is published as a fork diff at **https://layr-labs.github.io/mlx-c/**, described in [`fork.yaml`](fork.yaml) and kept honest by CI — see [FORKDIFF.md](FORKDIFF.md).
+
MLX C is a C API for [MLX](https://github.com/ml-explore/mlx).
MLX is an array framework for machine learning on Apple silicon. MLX C expands
diff --git ml-explore/mlx-c/scripts/check_forkdiff.py Layr-Labs/mlx-c/scripts/check_forkdiff.py
new file mode 100644
index 0000000000000000000000000000000000000000..8e10e811d614f83c414a1d323ddc64c7e6a3193b
--- /dev/null
+++ Layr-Labs/mlx-c/scripts/check_forkdiff.py
@@ -0,0 +1,259 @@
+#!/usr/bin/env python3
+"""Gate: ``fork.yaml`` must describe this fork as it is *now*.
+
+The fork-diff page (rendered with protolambda/forkdiff and published on GitHub
+Pages) is only useful while it is true. Two things make it go stale silently:
+
+1. **A rebase onto newer upstream.** ``base.hash`` still points at the old
+ upstream commit, so the page shows upstream's own changes as if the fork
+ made them. This script requires ``base.hash`` to equal
+ ``git merge-base HEAD <upstream>``; a rebase moves the merge-base, and the
+ gate stays red until the hash is bumped.
+2. **A new fork change nobody described.** Every path in
+ ``git diff --name-only base.hash HEAD`` must match a glob in some section
+ (or a global ``ignore``), and every glob must still match something — a
+ section describing code the fork no longer carries is as misleading as a
+ missing one.
+
+Both are pure git + YAML, so the same check runs locally::
+
+ python3 scripts/check_forkdiff.py # coverage only
+ python3 scripts/check_forkdiff.py --upstream-ref upstream/main
+
+Exit status is non-zero on any violation. ``--review-status-out`` writes the
+JSON the CI comment synthesizer consumes (same shape as history-check).
+"""
+
+from __future__ import annotations
+
+import argparse
+import json
+import re
+import subprocess
+import sys
+from pathlib import Path
+from typing import Iterable
+
+import yaml
+
+GLOB_CLASS_RE = re.compile(r"\[([^\]]*)\]")
+
+
+def glob_to_regex(glob: str) -> re.Pattern[str]:
+ """Translate a forkdiff glob into a regex over the repo-relative path.
+
+ Semantics follow the doublestar rules forkdiff uses: ``*`` and ``?`` never
+ cross a ``/``; ``**`` matches any number of directories (``a/**/b`` also
+ matches ``a/b``); ``[...]`` character classes pass through, with a leading
+ ``!`` meaning negation.
+ """
+ out: list[str] = []
+ i = 0
+ while i < len(glob):
+ c = glob[i]
+ if c == "*":
+ if glob.startswith("**", i):
+ if glob.startswith("**/", i):
+ out.append("(?:.*/)?")
+ i += 3
+ continue
+ out.append(".*")
+ i += 2
+ continue
+ out.append("[^/]*")
+ elif c == "?":
+ out.append("[^/]")
+ elif c == "[":
+ end = glob.find("]", i + 1)
+ if end == -1:
+ out.append(re.escape(c))
+ else:
+ cls = glob[i + 1 : end]
+ if cls.startswith("!"):
+ cls = "^" + cls[1:]
+ out.append("[" + cls + "]")
+ i = end + 1
+ continue
+ else:
+ out.append(re.escape(c))
+ i += 1
+ return re.compile("^" + "".join(out) + "$")
+
+
+def collect_globs(node: dict, path: str = "def") -> list[tuple[str, str]]:
+ """Every glob in the section tree as ``(section path, glob)``.
+
+ A section's ``ignore`` list counts as coverage too: forkdiff still lists
+ those files under the section (grayed out), so they are described.
+ """
+ found: list[tuple[str, str]] = []
+ title = node.get("title") or "(untitled)"
+ here = f"{path} › {title}" if path != "def" else title
+ for g in node.get("globs") or []:
+ found.append((here, str(g)))
+ for g in node.get("ignore") or []:
+ found.append((here, str(g)))
+ for child in node.get("sub") or []:
+ found.extend(collect_globs(child, here))
+ return found
+
+
+def git(*args: str, cwd: Path) -> str:
+ result = subprocess.run(
+ ["git", *args], cwd=cwd, capture_output=True, text=True, check=False
+ )
+ if result.returncode != 0:
+ raise RuntimeError(f"git {' '.join(args)} failed: {result.stderr.strip()}")
+ return result.stdout.strip()
+
+
+def changed_paths(repo: Path, base: str, head: str) -> list[str]:
+ out = git("diff", "--name-only", f"{base}..{head}", cwd=repo)
+ return [line for line in out.splitlines() if line]
+
+
+def check_coverage(
+ paths: Iterable[str], globs: list[tuple[str, str]]
+) -> tuple[list[str], list[tuple[str, str]], dict[str, int]]:
+ """Return ``(uncovered paths, stale globs, matches per glob)``."""
+ compiled = [(section, g, glob_to_regex(g)) for section, g in globs]
+ hits: dict[str, int] = {g: 0 for _, g in globs}
+ uncovered: list[str] = []
+ for p in paths:
+ matched = False
+ for _, g, rx in compiled:
+ if rx.match(p):
+ hits[g] += 1
+ matched = True
+ if not matched:
+ uncovered.append(p)
+ stale = [(section, g) for section, g in globs if hits[g] == 0]
+ return uncovered, stale, hits
+
+
+def check_base(repo: Path, base: str, head: str, upstream_ref: str | None) -> list[str]:
+ problems: list[str] = []
+ if not re.fullmatch(r"[0-9a-f]{40}", base):
+ problems.append(
+ f"base.hash must be a full 40-hex commit id, got {base!r} — a short or "
+ "symbolic ref would silently move under the page."
+ )
+ return problems
+ try:
+ git("cat-file", "-e", f"{base}^{{commit}}", cwd=repo)
+ except RuntimeError:
+ problems.append(
+ f"base.hash {base[:12]} is not present in this repository. The fork must "
+ "sit on top of it (fetch upstream if the clone is shallow)."
+ )
+ return problems
+ if upstream_ref is None:
+ return problems
+ try:
+ git("merge-base", "--is-ancestor", base, upstream_ref, cwd=repo)
+ except RuntimeError:
+ problems.append(
+ f"base.hash {base[:12]} is not an ancestor of {upstream_ref} — it must name "
+ "a commit on upstream main, not a fork commit."
+ )
+ merge_base = git("merge-base", head, upstream_ref, cwd=repo)
+ if merge_base != base:
+ problems.append(
+ f"base.hash {base[:12]} != merge-base({head}, {upstream_ref}) = "
+ f"{merge_base[:12]}. The fork was rebased onto newer upstream; update "
+ "fork.yaml's base.hash and re-describe the sections (see docs/forkdiff.md)."
+ )
+ return problems
+
+
+def review_status(problems: list[str], detail: str) -> list[dict]:
+ if not problems:
+ return []
+ return [
+ {
+ "source": "fork diff analysis",
+ "results": [
+ {
+ "kind": "action_required",
+ "title": "fork.yaml no longer describes the fork",
+ "summary": problems[0]
+ if len(problems) == 1
+ else f"{len(problems)} issues: {problems[0]}",
+ "detail": detail,
+ "how_to_fix": (
+ "See docs/forkdiff.md. After a rebase: set base.hash to "
+ "`git merge-base HEAD upstream/main`. For new files: add them to "
+ "the section that explains them (or to a global `ignore` if they "
+ "are not code). Then run `python3 scripts/check_forkdiff.py "
+ "--upstream-ref upstream/main` locally."
+ ),
+ }
+ ],
+ }
+ ]
+
+
+def main(argv: list[str] | None = None) -> int:
+ ap = argparse.ArgumentParser(description=__doc__.split("\n\n")[0])
+ ap.add_argument("--repo", default=".", help="path to the fork checkout")
+ ap.add_argument("--fork", default="fork.yaml", help="fork page definition")
+ ap.add_argument("--head", default="HEAD", help="the fork revision to describe")
+ ap.add_argument(
+ "--upstream-ref",
+ default=None,
+ help="a ref holding upstream main (e.g. refs/remotes/upstream/main); "
+ "enables the merge-base check",
+ )
+ ap.add_argument("--review-status-out", default=None, help="write review-status JSON here")
+ args = ap.parse_args(argv)
+
+ repo = Path(args.repo).resolve()
+ fork_path = repo / args.fork
+ spec = yaml.safe_load(fork_path.read_text(encoding="utf-8")) or {}
+ base = str((spec.get("base") or {}).get("hash") or "").strip()
+ problems: list[str] = []
+ lines: list[str] = []
+
+ problems += check_base(repo, base, args.head, args.upstream_ref)
+
+ globs = collect_globs(spec.get("def") or {})
+ globs += [("(global ignore)", str(g)) for g in spec.get("ignore") or []]
+
+ paths: list[str] = []
+ if re.fullmatch(r"[0-9a-f]{40}", base):
+ try:
+ paths = changed_paths(repo, base, args.head)
+ except RuntimeError as exc:
+ problems.append(str(exc))
+
+ uncovered, stale, hits = check_coverage(paths, globs)
+ if uncovered:
+ problems.append(
+ f"{len(uncovered)} changed file(s) are not described by any fork.yaml section"
+ )
+ lines.append("Uncovered files (add each to the section that explains it):")
+ lines += [f" - {p}" for p in uncovered]
+ if stale:
+ problems.append(f"{len(stale)} glob(s) match nothing the fork changes")
+ lines.append("Stale globs (the fork no longer changes anything they name):")
+ lines += [f" - {g} [{section}]" for section, g in stale]
+
+ covered = len(paths) - len(uncovered)
+ print(f"fork.yaml: base {base[:12]} head {args.head} changed files {len(paths)} "
+ f"described {covered} sections+ignores {len(globs)}")
+ if args.upstream_ref is None:
+ print(" (no --upstream-ref: merge-base check skipped)")
+ for line in lines:
+ print(line)
+ for p in problems:
+ print(f"::error::{p}")
+
+ if args.review_status_out:
+ Path(args.review_status_out).write_text(
+ json.dumps(review_status(problems, "\n".join(lines))), encoding="utf-8"
+ )
+ return 1 if problems else 0
+
+
+if __name__ == "__main__":
+ sys.exit(main())
Ignored changes
+299
-0
diff --git ml-explore/mlx-c/.github/workflows/forkdiff-check.yml Layr-Labs/mlx-c/.github/workflows/forkdiff-check.yml
new file mode 100644
index 0000000000000000000000000000000000000000..e09dc8969643eb0a83092fa562a541c39b5928a0
--- /dev/null
+++ Layr-Labs/mlx-c/.github/workflows/forkdiff-check.yml
@@ -0,0 +1,76 @@
+name: Fork Diff Check
+
+# Fails a PR whose `fork.yaml` no longer describes this fork.
+#
+# The fork-diff page (https://layr-labs.github.io/mlx-c/, rendered by
+# protolambda/forkdiff from `fork.yaml`, deployed by forkdiff-pages.yml) is
+# only useful while it is true, and two things make it go stale silently:
+#
+# 1. A rebase onto newer upstream. `base.hash` keeps pointing at the old
+# upstream commit, so the page shows upstream's own changes as the fork's.
+# `scripts/check_forkdiff.py` requires base.hash == merge-base(HEAD,
+# upstream/main); a rebase moves the merge-base and the gate stays red
+# until the hash — and the sections — are brought up to date.
+# 2. A fork change nobody described. Every file in the base..HEAD diff must
+# match a section glob (or a global ignore), and every glob must still
+# match something.
+#
+# The page is also rendered here so a fork.yaml that forkdiff itself rejects
+# cannot merge.
+
+on:
+ pull_request:
+ push:
+ branches: [main]
+ workflow_dispatch:
+
+permissions:
+ contents: read
+
+jobs:
+ analysis-up-to-date:
+ name: fork.yaml describes the fork
+ runs-on: ubuntu-latest
+ timeout-minutes: 20
+ steps:
+ - uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2
+ with:
+ fetch-depth: 0 # full history: merge-base with upstream, and forkdiff diffs against base.hash
+
+ - name: Fetch upstream main
+ run: |
+ git fetch --no-tags --quiet https://github.com/ml-explore/mlx-c.git \
+ main:refs/remotes/upstream/main
+ echo "upstream/main = $(git rev-parse --short refs/remotes/upstream/main)"
+ echo "merge-base = $(git merge-base HEAD refs/remotes/upstream/main | cut -c1-12)"
+
+ - uses: actions/setup-python@a309ff8b426b58ec0e2a45f0f869d46889d02405 # v6.2.0
+ with:
+ python-version: "3.12"
+
+ - name: Verify fork.yaml against the real diff
+ run: |
+ python3 -m pip install --quiet pyyaml
+ python3 scripts/check_forkdiff.py --upstream-ref refs/remotes/upstream/main
+
+ - uses: actions/setup-go@40f1582b2485089dde7abd97c1529aa768e1baff # v5.6.0
+ with:
+ go-version: "1.24"
+ cache: false
+
+ - name: Render the fork-diff page
+ # fork.yaml names refs/heads/main; on a PR the checkout is a detached
+ # merge commit, so point the local branch at what we are checking.
+ run: |
+ git update-ref refs/heads/main HEAD
+ mkdir -p tmp/pages
+ go run github.com/protolambda/forkdiff@v0.1.1 \
+ -repo . -fork fork.yaml -out tmp/pages/index.html
+ echo "rendered $(wc -c < tmp/pages/index.html) bytes"
+
+ - name: Upload rendered page
+ uses: actions/upload-artifact@043fb46d1a93c77aae656e7c1c64a875d1fc6a0a # v7
+ with:
+ name: forkdiff-page
+ path: tmp/pages/index.html
+ retention-days: 7
diff --git ml-explore/mlx-c/.github/workflows/forkdiff-pages.yml Layr-Labs/mlx-c/.github/workflows/forkdiff-pages.yml
new file mode 100644
index 0000000000000000000000000000000000000000..7e4447313221bd445d8a515eaec6c1925fe8fe62
--- /dev/null
+++ Layr-Labs/mlx-c/.github/workflows/forkdiff-pages.yml
@@ -0,0 +1,99 @@
+name: Deploy Fork Diff
+
+# Renders `fork.yaml` with protolambda/forkdiff and publishes the result as the
+# repository's GitHub Pages site — the same setup ethereum-optimism/op-geth
+# uses for its go-ethereum fork diff. Runs the analysis gate first, so a page
+# that lies about the fork is never published.
+
+on:
+ push:
+ branches: [main]
+ workflow_dispatch:
+
+permissions:
+ contents: read
+ pages: write
+ id-token: write
+
+concurrency:
+ group: "pages"
+ cancel-in-progress: true
+
+jobs:
+ deploy:
+ name: Render and deploy
+ environment:
+ name: github-pages
+ url: ${{ steps.deployment.outputs.page_url }}
+ runs-on: ubuntu-latest
+ timeout-minutes: 20
+ steps:
+ - uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6.0.2
+ with:
+ fetch-depth: 0 # forkdiff diffs against base.hash, deep in history
+
+ - name: Fetch upstream main
+ run: |
+ git fetch --no-tags --quiet https://github.com/ml-explore/mlx-c.git \
+ main:refs/remotes/upstream/main
+
+ - uses: actions/setup-python@a309ff8b426b58ec0e2a45f0f869d46889d02405 # v6.2.0
+ with:
+ python-version: "3.12"
+
+ - name: Refuse to publish a stale analysis
+ run: |
+ python3 -m pip install --quiet pyyaml
+ python3 scripts/check_forkdiff.py --upstream-ref refs/remotes/upstream/main
+
+ - uses: actions/setup-go@40f1582b2485089dde7abd97c1529aa768e1baff # v5.6.0
+ with:
+ go-version: "1.24"
+ cache: false
+
+ - name: Build forkdiff
+ run: |
+ mkdir -p tmp/pages
+ go run github.com/protolambda/forkdiff@v0.1.1 \
+ -repo . -fork fork.yaml -out tmp/pages/index.html
+ touch tmp/pages/.nojekyll
+
+ - name: Is GitHub Pages enabled?
+ # Pages must be switched on once by a repo admin (Settings → Pages →
+ # Source: "GitHub Actions"). The workflow token cannot do that itself
+ # (`configure-pages` enablement needs an admin PAT), so until then the
+ # page is built and kept as an artifact but not deployed — a notice,
+ # not a red run.
+ id: pages
+ env:
+ GH_TOKEN: ${{ github.token }}
+ run: |
+ if gh api "repos/${GITHUB_REPOSITORY}/pages" --silent 2>/dev/null; then
+ echo "enabled=true" >> "$GITHUB_OUTPUT"
+ else
+ echo "enabled=false" >> "$GITHUB_OUTPUT"
+ echo "::notice::GitHub Pages is not enabled for ${GITHUB_REPOSITORY}; built the page but skipped the deploy. Enable it once under Settings → Pages (Source: GitHub Actions) and re-run this workflow."
+ fi
+
+ - name: Keep the rendered page as an artifact
+ if: steps.pages.outputs.enabled != 'true'
+ uses: actions/upload-artifact@043fb46d1a93c77aae656e7c1c64a875d1fc6a0a # v7
+ with:
+ name: forkdiff-page
+ path: tmp/pages/index.html
+ retention-days: 30
+
+ - name: Setup Pages
+ if: steps.pages.outputs.enabled == 'true'
+ uses: actions/configure-pages@983d7736d9b0ae728b81ab479565c72886d7745b # v5.0.0
+
+ - name: Upload artifact
+ if: steps.pages.outputs.enabled == 'true'
+ uses: actions/upload-pages-artifact@fc324d3547104276b827a68afc52ff2a11cc49c9 # v5.0.0
+ with:
+ path: tmp/pages
+
+ - name: Deploy to GitHub Pages
+ if: steps.pages.outputs.enabled == 'true'
+ id: deployment
+ uses: actions/deploy-pages@368f82528645a54fb793d4d04e342629a3f51346 # v5.0.1
diff --git ml-explore/mlx-c/fork.yaml Layr-Labs/mlx-c/fork.yaml
new file mode 100644
index 0000000000000000000000000000000000000000..7b45caa39b06b19e950e95f8e41b24f704ecdbc3
--- /dev/null
+++ Layr-Labs/mlx-c/fork.yaml
@@ -0,0 +1,124 @@
+title: "Layr-Labs/mlx-c - MLX C fork diff overview"
+footer: |
+ Fork-diff overview of [`Layr-Labs/mlx-c`](https://github.com/Layr-Labs/mlx-c), a fork of
+ [`ml-explore/mlx-c`](https://github.com/ml-explore/mlx-c) · the C bridge in Layr-Labs'
+ Apple-silicon inference stack ([`mlx`](https://github.com/Layr-Labs/mlx) →
+ [`mlx-c`](https://github.com/Layr-Labs/mlx-c) → [`mlx-swift`](https://github.com/Layr-Labs/mlx-swift) →
+ [`mlx-swift-lm`](https://github.com/Layr-Labs/mlx-swift-lm)) · created with
+ [Forkdiff](https://github.com/protolambda/forkdiff)
+base:
+ name: ml-explore/mlx-c
+ url: https://github.com/ml-explore/mlx-c
+ # The upstream commit this fork is rebased onto. CI (`scripts/check_forkdiff.py`)
+ # requires this to equal `git merge-base HEAD upstream/main`, so a rebase onto
+ # newer upstream fails the gate until the hash — and the sections below — are
+ # brought up to date. Upstream main as of 2026-04-23 (bindings for MLX 0.31.2).
+ hash: fba4470b89073180056c9ea46c443051375f7399
+fork:
+ name: Layr-Labs/mlx-c
+ url: https://github.com/Layr-Labs/mlx-c
+ ref: refs/heads/main
+def:
+ title: "Layr-Labs/mlx-c"
+ description: |
+ This is an overview of the changes in [`Layr-Labs/mlx-c`](https://github.com/Layr-Labs/mlx-c),
+ a fork of [`ml-explore/mlx-c`](https://github.com/ml-explore/mlx-c).
+
+ MLX C is the C boundary between the MLX core and every other language binding. Layr-Labs
+ runs an LLM/VLM serving stack on Apple silicon — continuous batching, paged KV caches,
+ MoE and quantized kernels — built as [`mlx-swift-lm`](https://github.com/Layr-Labs/mlx-swift-lm)
+ on [`mlx-swift`](https://github.com/Layr-Labs/mlx-swift), which links a forked
+ [`mlx`](https://github.com/Layr-Labs/mlx) core through this fork. Every change here is a
+ **carry-patch**: a thin C wrapper over a function the forked core added, so Swift can reach
+ it. The upstream API and ABI are left untouched; the fork only adds entry points, and pins
+ the core it builds against to the Layr-Labs revision that provides them.
+
+ Downstream consumers are named per section so a change here can be traced to the Swift
+ symbol that needs it.
+ sub:
+ - title: "Core pin"
+ description: |
+ `CMakeLists.txt` fetches the MLX core from `Layr-Labs/mlx` at an exact commit instead of
+ an upstream release tag. The C functions below wrap core functions that only exist on the
+ fork, so the pin is what makes them link; it is bumped together with the header changes
+ each time the core gains a new entry point.
+ globs:
+ - "CMakeLists.txt"
+ - title: "Memory: allocator observability and footprint policy"
+ description: |
+ The Swift admission code decides whether a request fits by asking the allocator, and the
+ upstream C API only offered independent counters. The fork adds, in order of arrival:
+
+ - `mlx_get_num_resources` / `mlx_get_resource_limit` — the live Metal buffer *count* and
+ its hard ceiling, the quantity behind the `[metal::malloc] Resource limit exceeded`
+ crash (consumed as `Memory.numResources` / `Memory.resourceLimit` in mlx-swift).
+ - `mlx_get_memory_snapshot` — one coherent active/cache/peak observation instead of three
+ separate calls that can straddle an allocation.
+ - `mlx_get_allocation_footprint_policy` and its checked `_bound` / `_maximum_extra`
+ helpers — the allocator's immutable rounding geometry (alignment, rounding threshold,
+ minimum allocation, power-of-two cut-off, cache page size) so callers can compute the
+ real footprint of a requested size without allocating; bound helpers return 1 on an
+ invalid policy or overflow without raising through the error callback.
+ - `mlx_get_allocation_size_upper_bound` — the same bound for the current policy.
+
+ The generator override in `python/mlxhooks.py` keeps `mlx_get_memory_snapshot` intact
+ when the bindings are regenerated from the core headers.
+ globs:
+ - "mlx/c/memory.h"
+ - "mlx/c/memory.cpp"
+ - "python/mlxhooks.py"
+ - title: "Array: buffer metadata and constant-cache identity"
+ description: |
+ Two non-evaluating inspections the quantized-inference path in mlx-swift relies on:
+
+ - `mlx_array_get_buffer_info` — backing metadata (allocated bytes, data offset, element
+ count, row-contiguity, uniqueness) for an *already evaluated* array. Unavailable backing
+ is reported as `available=false` rather than forcing an eval or a stream sync; views may
+ report shared backing and the fields confer no ownership.
+ - `_mlx_array_constant_cache_identity` — the backing descriptor's identity plus a
+ transform-safety predicate, so a Swift constant cache can reuse an unchanged constant's
+ exact cast across calls without mistaking a mutated array for its former contents. It
+ refuses compile/autodiff tracing and retained graphs, and never evaluates, hashes, or
+ reads data back. Marked internal: the caller must retain the source descriptor while
+ the identity is cached.
+ globs:
+ - "mlx/c/array.h"
+ - "mlx/c/array.cpp"
+ - title: "Fast: mutable-input Metal kernels and the SDPA signature"
+ description: |
+ - `mlx_fast_metal_kernel_new_mutable` — a second constructor for custom Metal kernels
+ that declares which caller-owned input buffers the kernel writes (the core's
+ mutable-input contract, used for in-place paged KV updates). The existing constructor
+ and its read-only ABI are unchanged.
+ - The scaled-dot-product-attention binding passes the `force_fused` argument MLX 0.32.2
+ added to the core call, as `false`, so the C signature and pre-0.32.2 behaviour are
+ preserved.
+ globs:
+ - "mlx/c/fast.h"
+ - "mlx/c/fast.cpp"
+ - title: "Stream: cross-thread default streams"
+ description: |
+ MLX 0.32 made the default stream's Metal command encoder thread-local, so a graph built
+ on one thread and evaluated on another — Swift `async/await` continuations, actor
+ executors — aborts with `There is no Stream(gpu, N) in current thread`.
+ `mlx_thread_unsafe_cpu_stream_new` / `mlx_thread_unsafe_gpu_stream_new` expose the
+ core's `new_thread_unsafe_stream`, whose encoder is registered globally, letting
+ mlx-swift restore a single process-wide default stream. "Thread-unsafe" means the caller
+ serializes submission; it is not locked for concurrent multi-thread use.
+ globs:
+ - "mlx/c/stream.h"
+ - "mlx/c/stream.cpp"
+ - title: "Fork tooling"
+ description: |
+ The gate that keeps this page honest: `check_forkdiff.py` fails CI when `base.hash`
+ is not the merge-base with upstream, or when a file the fork changes is not described
+ by a section above. `FORKDIFF.md` explains the day-to-day; the README links the page.
+ globs:
+ - "scripts/check_forkdiff.py"
+ - "FORKDIFF.md"
+ - "README.md"
+
+# ignored globally, does not count towards line count
+ignore:
+ - "fork.yaml"
+ - ".github/**"