* MLX: reduce Qwen3.8 Flash Next memory use
For NVFP4 creation, retain source precision for sensitive QSA projections
and non-expert MTP weights. Keep routed experts and PLE at NVFP4, and
promote other eligible weights to MXFP8.
Keep NVFP4 PLE tables host-backed, transferring only selected packed rows
and their scales to MLX. Honor scalar and per-expert NVFP4 global scales
in MoE, and per-shard global scales in host-backed PLE.
Use compressed QSA block selections in fused indexed Metal attention for
long prefill batches, retaining the graph fallback. Co-locate component
loading and forward methods, with explicit quantization and cache ownership.
* mlx: skip materializing host-backed tensors
Let models exclude tensors already owned by host storage from MLX loading.
Avoid allocating and discarding the full Qwen PLE tables at startup while
preserving the resident fallback.
* address comments
NVIDIA's libcusolver looks for its dependencies only in its own
directory, so libcublas and libnvJitLink resolved to the host's copies
and MLX failed to load when those didn't match the bundled ones. Preload
the bundled copies by full path before loading MLX.
MLX's argpartition sorts the whole row, and a row longer than one block
of the GPU sort takes a multi-pass merge sort. TopKIndices sorts each
2048-value chunk in a single block, then only the chunks' top k, and
returns the same indices in the same order.
The sampler's top_k and top_logprobs and the DFlash2 candidate selector
use it in place of a full-vocabulary argpartition.
* MLX, MLX-C: version bump
The MLX-C pin now provides its own bindings for gather_qmm global_scale
and fast::gated_delta_update, so those two carried patches drop out.
Carry 0001-mlx-c-sdpa-window-size instead: sliding-window attention
added a window_size argument to fast::scaled_dot_product_attention
(ml-explore/mlx#4734) that MLX-C has not bound yet, to be regenerated
from MLX-C once the upstream binding PR merges. The SDPA wrappers gain
a window size argument; all call sites default to the full context.
* mlxrunner: dispatch sliding-window prefills to the native SDPA window
Sliding layers no longer materialize a causal∩sliding mask for prefill
chunks past the window: resolve() dispatches {causal, window} natively
when the rotating cache's L>1 concat has linearized the ring and the
mask is otherwise pure causal, via a new MaskApplier.NativeWindow hint.
Vision relaxations, padding, model arrays, and decode keep the
materialized path. No model changes; sinks compose natively.
Preserve source tool JSON alongside decoded API values so Flint can retain
numeric types and schema order in prompts. Existing API marshaling is unchanged.
Assemble ModelOpt expert banks during safetensors import, and reject direct
Flint NVFP4 creation before conversion. Other models retain their existing
import and quantization paths.
Add optional learned sinks and device causal masks to shared MLX attention.
Existing callers retain the no-sink and host-mask paths; Flint opts into the
new behavior.
* Deduplicate MLX CUDA runtime payloads
Move CUDA libraries shared by the GGML and MLX backends into the sibling CUDA payload, and teach in-process and subprocess MLX loading to resolve that packaged dependency directory.
Keep MLX-private libraries and JIT headers in the MLX overlay, include the full OpenBLAS Fortran runtime closure, and verify those dependencies in release packaging.
* address comments
* ci: shorten Windows Vulkan build path
Keep the nested shader-generator compiler probes below Windows path limits in checkouts with longer repository names.
* mlx: fix Nemotron bugs and support global scales
Accept published Nemotron layer names while retaining legacy aliases, and
fix layer-count inference and detached MTP quantization. Stabilize Mamba
softplus, preserve FP32 vision normalization, and compile Mamba gate and
norm scaling.
Support per-expert F32 global scales and preserve activation-scale
companions in ModelOpt and compressed-tensors NVFP4 imports. Normalize
reciprocal globals and reject invalid scales. Keep activation scales
separate from weight-only inference.
* mlx: add Nemotron 3.5 Super
* minor adjustments
* mlx: port the Metal custom kernels to CUDA
mamba2_scan, depthwise_conv_silu and gated_delta (plus their _states/_bias
variants) had a Metal source only, so on CUDA each fell back to graph ops.
Deviations from the Metal sources:
- The scan returns early past D: Metal dispatches an exact thread count,
CUDA rounds the grid up to whole blocks.
- T is dereferenced; CUDA passes inputs by address, Metal scalarizes
zero-dimensional ones.
- simd_sum becomes a shfl butterfly, and the threadgroup q/k caches become
per-warp registers: each warp normalizes the row its own lanes consume,
costing a redundant norm but no shared memory and no barriers.
- The gates and SiLU call MLX's Sigmoid/LogAddExp/Exp instead of restating
their formulas. Metal cannot: its custom-kernel preamble is metal::utils()
alone, while CUDA embeds the op headers for nvrtc.
The parity tests now run on Metal and CUDA hosts. The fused gated-delta
comparison takes a tolerance: the kernel and the graph path reduce in a
different order, so the bf16 output can land one ULP away.
* address comments
* Add Kev decision scoring to MLX
Load Kev models with their FP32 decision head and release metadata.
Support System One scoring with shared-state caching and Kev-specific limits.
* address comments
* mlx: compact restored recurrent state after speculative rollback
Per-token recurrent captures can leave a restored state pointing into a larger buffer. If a request ends before another model step, the live cache retains that buffer even though it accounts for only the selected state.
Track when a restored state may share that backing buffer. At request close, compact marked states in one batch before retaining them in the prefix cache. Clear the marker when a normal model step replaces the state, and cover restoration and buffer detachment in tests.
* address comments
* Reduce model lookup and MLX decision request overhead
Avoid decoding unrelated manifests when resolving existing model names.
Reuse small Metal scratch buffers between decision requests, releasing
larger working sets and clearing scratch after scoring errors.
* address comments
* mlxrunner: restore System One scoring
Carry the MLX implementation and its tests from the original System One
commit. The API/compiler and llama-server implementation already exist on
the release branch, so retain those versions. Adapt the optional model test
to the renamed decision package and its explicit rendering/context setup.
(cherry picked from commit 26e546ddb3)
* Enable MLX System One and add regression coverage
* mlxrunner: add Laya support for System One
Import Laya checkpoints and implement bidirectional decision scoring.
* mlxrunner: release cached buffers after System One scoring
* decision: match Tev1 and Nimble publisher prompt formats
Select Tev1 prompts using renderer metadata and fix JSON formatting, with regression coverage for GGUF and MLX.
* mlx: support Clef and Clef Flash decision models
Add model import, quantization safeguards, and text/image scoring with the Clef joint decision head.
* support vision capabilities for decision models
* Make Clef scoring cancellable between prefill chunks
Preserve joint-head inputs and test cancellation, cleanup, and recovery.
* mlx: separate decision preparation and readout from forward
Let serial and batched scoring share model-owned preparation and readout
without duplicating architecture-specific logic in the runner.
* mlx: add Strands Decider support
* implement caching
---------
Co-authored-by: ParthSareen <parth.sareen@ollama.com>
* mlx: speed up Qwen 3.8 prompt processing
Use MLX's gated-delta kernel for long scans and fold dense MLP global scales into SwiGLU.
* address comments
The MLX runner is the only Go inference runner left and is no longer
experimental, so its packages leave x/. The bindings become a top-level
mlx package beside the carried patches in mlx/compat, mirroring how
llama/ holds the llama.cpp integration, and the runner becomes mlxrunner
with the architectures nested under the package they implement.
Subpackages move with their parent unless listed.
x/mlxrunner/mlx mlx
x/internal/mlxthread mlx/mlxthread
x/internal/mlxthreadtest mlx/mlxthread/mlxthreadtest
x/internal/mlxtest mlx/mlxtest
x/quant mlx/quant
mlx/compat/*.patch mlx/compat/mlx-c (MLX patches go in mlx/compat/mlx)
x/mlxrunner mlxrunner
x/models/nn mlxrunner/nn
x/models/<arch> mlxrunner/model/<arch>
x/mlxrunner/imports.go mlxrunner/model/architectures (new package)
x/create create
x/safetensors fs/safetensors
x/tokenizer mlxrunner/tokenizer
Every package keeps its name, so the Go changes are the import path
rewrites the moves force, and the CMake, Dockerfile, CI cache keys, drift
check and Darwin payload script follow the new paths. Four edits are not
paths: the runner's blank architecture imports become the package
mlxrunner/model/architectures, so the list to extend for a new model sits
beside the architecture directories; a depguard rule keeps the two test
harnesses out of non-test code, as the x/internal placement used to; the
CI change filter's two entries for the long-deleted x/imagegen/mlx now
name the bindings' CMake project and the carried patches, so a change to
either builds the payload; and the tokenizer parity test reads its
fixtures from its own testdata instead of walking out of x/.
x/server and x/imagegen/manifest stay for the next two commits.
Daniel HiltgenandGitHub
Jesse Gross