Skip to content

Building from Source

This page covers building mlxcel from source on macOS (Apple Silicon), Linux (CUDA and AMD ROCm), and Windows (NVIDIA CUDA and AMD ROCm). ROCm support arrived in 0.8 on both operating systems and, like the Windows builds, has no prebuilt archives. The Windows CUDA build is still experimental. For the fastest package-based install on macOS Apple Silicon, use the Homebrew path in Quickstart. Linux builds are CUDA-specific and are not in the Homebrew tap; download the mlxcel-linux-*-cuda13 archives from the GitHub release page or build from source as described here. The v0.8.0 release is a prerelease with no archives attached, and releases/latest points to the previous stable release, so to run v0.8.0 clone the tag with git clone --branch v0.8.0 https://github.com/lablup/mlxcel.git and build it. Building from source remains useful when you need local compiler flags, unreleased changes, or hardware-specific CUDA settings.

Table of Contents


Prerequisites

macOS (Apple Silicon)

mlxcel targets Apple Silicon hardware (M1, M2, M3, M4, and M5 series). Intel-based Macs are not supported because the project depends on Apple's MLX framework, which requires Apple Silicon's unified memory architecture.

Requirement Details
Hardware Apple Silicon Mac (M1 / M2 / M3 / M4 / M5 series)
Rust 2024 edition; the v0.8.0 source tree pins Rust 1.97.1
Xcode Command Line Tools Required for system headers and linker
Metal Toolchain Required to compile MLX Metal GPU shaders

Install the Rust toolchain from rustup.rs if you do not have it:

curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh

Install Xcode Command Line Tools:

xcode-select --install

Install the Metal Toolchain (required for MLX shader compilation, which is separate from Xcode):

xcodebuild -downloadComponent MetalToolchain

Linux (CUDA)

Linux builds target NVIDIA GPUs via the CUDA backend. Both aarch64 (for Grace Hopper and similar ARM-based NVIDIA platforms) and x86_64 systems are supported.

Requirement Details
Architecture aarch64 or x86_64
Rust 2024 edition; the v0.8.0 source tree pins Rust 1.97.1
CMake 3.24 or later (for building the MLX C++ library)
CUDA Toolkit 12.0 or later, including nvcc
cuDNN 9 or later
System libraries libopenblas-dev, liblapack-dev, liblapacke-dev

Install system dependencies on Ubuntu or Debian:

sudo apt-get update
sudo apt-get install -y \
    cmake \
    libcudnn9-dev-cuda-13 \
    libopenblas-dev \
    liblapack-dev \
    liblapacke-dev

Note: Adjust the cuDNN package name to match your installed CUDA version (e.g., libcudnn9-dev-cuda-12 for CUDA 12.x).

Linux (AMD ROCm)

AMD GPU support on Linux arrived in 0.8 and is a source build only: there is no release artifact, and the optional ROCm CI job is disabled unless the repository sets ROCM_CI_ENABLED. The backend is not part of upstream MLX; mlxcel vendors it from the rocm-support branch of NripeshN/mlx (MIT) into src/lib/mlxcelverse/overlays/rocm/ and applies it on top of the same pinned MLX commit the Metal and CUDA builds use. src/lib/mlxcelverse/overlays/rocm/UPSTREAM records the fork commit, and LOCAL_FIXES.md in the same directory lists the changes mlxcel carries on top of it.

The tested configuration is an AMD Ryzen AI MAX+ 395 with a Radeon 8060S (gfx1151, RDNA 3.5) and a 96 GiB VRAM carve-out, on Debian 13 with ROCm 10.0.0 packages (HIP 7.15). Other RDNA 3, 3.5, and 4 parts are expected to build. CDNA parts (for example MI300) compile but carry no tuning, and the fused kernels that reduce across a 32-lane wavefront take their MLX graph fallbacks on them (see the Wave64 row in Current Status).

Requirement Details
Rust, CMake, C++ Rust 2024 edition, CMake, and a C++20 compiler
System libraries pkg-config, OpenSSL headers, libopenblas-dev, liblapack-dev, liblapacke-dev
ROCm CMake packages hip, rocblas, rocthrust, rocprim, hiprand, rocwmma, hipblaslt, hipfft, and hiprtc, plus hipcc and rocminfo
Device access The build and run user must be in the video and render groups
sudo apt-get install -y \
    build-essential cmake git pkg-config libssl-dev \
    libopenblas-dev liblapack-dev liblapacke-dev
# ROCm itself comes from AMD's repository (repo.radeon.com).
rocminfo | grep -E '^\s+Name:\s+gfx'   # should list your GPU target

Build Instructions

macOS Build

# Clone the repository
git clone https://github.com/lablup/mlxcel.git
cd mlxcel

# Build with Metal and Accelerate framework support
cargo build --release --features metal,accelerate

The metal feature enables Apple's Metal GPU acceleration through the MLX C++ library. The accelerate feature enables Apple's Accelerate framework for optimized CPU linear algebra operations, which is used during model loading and on CPU fallback paths.

On macOS, MLX also enables Metal and Accelerate for a plain cargo build --release; the explicit features above document the intended backend. Runtime controls such as the MTP qmv_wide retry follow the backend actually built, including plain builds. MLXCEL_BUILD_METAL=OFF disables the Metal backend and its controls even when the Cargo metal feature is selected.

To check those runtime controls without loading a checkpoint, run:

cargo run -p mlxcel-core --profile test-fast --example metal_runtime_switch_probe

The probe checks QMV selection and the decode command-buffer override, then restores their previous values. It reports a skip when no Metal device is available.

For an older plain macOS build affected by #1988, start a fresh process with MLXCEL_QMV_WIDE=0 before the server command. This pins the narrower Metal projection kernel at startup while keeping the exactness gate enabled; it may reduce throughput, applies only to Metal, and still requires the loaded model's probe to pass.

Linux Build

# Clone the repository
git clone https://github.com/lablup/mlxcel.git
cd mlxcel

# Build with CUDA support
cargo build --release --features cuda

The first build will take several minutes because the MLX C++ library and its CUDA kernels are compiled from source as part of the Cargo build process.

Linux ROCm Build

cargo build --release --features rocm

rocm cannot be combined with cuda or metal; the build fails with a message naming the conflict. The first build compiles the MLX device code with hipcc, which adds a few minutes to the Rust build. See ROCm Build Options for target selection and the current status.

Incremental builds track header dependencies: hipcc writes a dependency file next to each HIP object, so after you change a header under src/lib/mlxcelverse/overlays/rocm/, the next build recompiles the .hip files that include it, directly or transitively, and a build with no change recompiles none. The overlay reaches the build tree through CMake's configure_file, which rewrites a copy only when its content differs, so a bare touch of an overlay header rebuilds nothing. The dependency files also list ROCm and system headers, but not headers that only the device compilation includes (rocWMMA's) or the compiler itself.

To force a clean HIP rebuild (for example after upgrading hipcc or rocWMMA, or to rule out a stale object while bisecting), delete the HIP objects of the build profile you use and touch an overlay file so that Cargo reruns the build script:

rm -rf target/release/build/mlxcel-core-*/out/build/_deps/mlx-build/mlx/backend/rocm/hip_objs
touch src/lib/mlxcelverse/overlays/rocm/mlx/backend/rocm/CMakeLists.txt
cargo build --release --features rocm

Replace release with the profile directory you build (debug, test-fast, and so on). Without the touch, Cargo sees no changed input, skips the build script, and keeps linking the old kernels. cargo clean -p mlxcel-core --release also forces the rebuild, but rebuilds the whole MLX C++ library too.


Cargo Feature Flags

Both binaries build from the same root package, so one feature set applies to mlxcel and mlxcel-server alike. Pass them with cargo build --features <a,b>. Shipping builds enable the platform backend flags and nothing else; the rest are opt-in seams or test scaffolding.

Feature Default Effect
webui on Embeds the bundled administrative WebUI assets into both binaries and enables --webui. A binary built without it refuses --webui at startup. See WebUI.
surgery on Weight-load surgery. Exposes --surgery <config.yaml> and MLXCEL_SURGERY for scale, add, prune, replace and interpolate weight-space edits at load time. With no surgery config supplied, the load path is byte-for-byte identical to a build without the feature.
metal off Apple Silicon Metal GPU backend. Standard on macOS.
accelerate off Apple Accelerate CPU BLAS backend. Standard on macOS.
cuda off NVIDIA CUDA GPU backend. Required on NVIDIA hosts; a plain build is CPU-only. Also builds natively on Windows (experimental, see Windows (NVIDIA CUDA)).
rocm off AMD ROCm/HIP GPU backend on Linux and Windows, built from the ROCm overlay in src/lib/mlxcelverse/overlays/rocm/. Supported since 0.8, source build only. Mutually exclusive with cuda and metal. See Windows (AMD ROCm) for the Windows build.
experimental-backend off Reserves the non-MLX compute-backend seam slot. Ships no kernels and adds no runtime dispatch.
xla-backend off OpenXLA/StableHLO seam and graph emitter, with no execution engine. Needs no native toolchain.
xla-iree off xla-backend plus a real IREE runtime. Needs IREE_DIST, or one of IREE_CUDA_HOME, IREE_ROCM_HOME, or IREE_MACOS_HOME with its matching compile variable (IREE_CUDA_COMPILE, IREE_ROCM_COMPILE, IREE_MACOS_COMPILE), at build time. IREE_ROCM_HOME is a source-built HIP-enabled IREE runtime for gfx1151, gfx1100, or gfx942.
xla-diagnostics-cpu, xla-diagnostics, xla-diagnostics-rocm off Diagnostics-only builds that compare the eager MLX path against the resident IREE graphs (for example the Qwen2.5-VL vision oracle). xla-diagnostics-cpu is xla-iree plus the mlxcel-xla diagnostics; xla-diagnostics adds cuda, and xla-diagnostics-rocm adds rocm for AMD GPU oracle runs. Not for shipping builds.
test-utils off Test-only helpers, required to build the distributed_integration, pipeline_e2e and paged_handoff_parity integration tests. Not needed for the binaries.

default = ["surgery", "webui"], so a plain cargo build enables surgery and the WebUI only, and a real build adds a platform backend on top. Build with --no-default-features to drop both, and add back the one you want to keep:

# Metal + Accelerate, no surgery crate and no WebUI.
cargo build --release --no-default-features --features metal,accelerate

# Metal + Accelerate with surgery, without the WebUI.
cargo build --release --no-default-features --features metal,accelerate,surgery

See The OpenXLA Backend for what the XLA features compile, what they need at build time, and which model paths (including the Qwen2.5-VL vision tower) run on IREE.


CUDA Build Options

GPU Architecture Auto-Detection

By default, the build script queries nvidia-smi to determine the compute capability (SM version) of the GPU present on the build machine and compiles CUDA kernels for that target. This produces the smallest possible binary and the fastest kernel specialization for your hardware. The detected value is spelled the way the release workflow spells it: Hopper gets CUDA's architecture-specific a suffix (a detected 90 becomes 90a), and Blackwell (100, 120, 121) stays plain. If detection fails, the build falls back to 90a.

On Blackwell the plain spelling still gets MLX's hardware NVFP4/MXFP4 converters: the MLX CMake build adds the architecture-specific image to the one translation unit that uses them, so block-float quantization takes the hardware path without changing the code of any other kernel. Do not put 121a in the list: it changes the machine code of every kernel on the device, and 121f fails to compile.

Manual Architecture Override

Use the MLX_CUDA_ARCHITECTURES environment variable to specify one or more SM versions. This is necessary when:

  • Building on a machine without a GPU (headless build servers, CI pipelines)
  • Creating a binary that runs on a different GPU than the build machine
  • Targeting multiple GPU generations in a single binary
# Hopper (H100, H200). The `a` suffix is required for the dedicated Hopper
# quantized kernel; a plain `90` builds without it.
MLX_CUDA_ARCHITECTURES=90a cargo build --release --features cuda

# Blackwell GB10 (DGX Spark). Plain, as the release workflow builds it.
MLX_CUDA_ARCHITECTURES=121 cargo build --release --features cuda

# Multiple architectures in a single binary (larger binary, longer build)
MLX_CUDA_ARCHITECTURES="90a;121" cargo build --release --features cuda

Custom CUDA Path

If your CUDA Toolkit is not installed in the default location (/usr/local/cuda), set CUDA_HOME to point to the correct path:

CUDA_HOME=/opt/cuda-13.0 cargo build --release --features cuda

You can combine both variables:

CUDA_HOME=/opt/cuda-13.0 MLX_CUDA_ARCHITECTURES=90a cargo build --release --features cuda

GPU Architecture Compatibility Table

Not all GPU generations support all mlxcel features. The table below summarizes compatibility:

GPU Family SM Version 4-bit Quantized Models fp16 / bf16 Models Notes
Turing (RTX 20xx) 75 Yes (generic) Yes No tensor-core quant kernel; uses the generic qmm_naive/qmv fallback (slower, fp16, untested)
Ampere (RTX 30xx, A100) 80 / 86 Yes Yes Tensor-core qmm_sm80 (SM ≥ 80) plus the qmv decode path
Ada Lovelace (RTX 40xx) 89 Yes Yes Tensor-core qmm_sm80 plus qmv
Hopper (H100, H200) 90 Yes (WGMMA) Yes Dedicated qmm_sm90 WGMMA kernel; requires a 90a build to be compiled in
Blackwell (B100, B200) 100 Yes Yes Tensor-core qmm_sm80 plus qmv
Blackwell GB10 (DGX Spark) 120 / 121 Yes Yes qmm_sm80 plus qmv; a few uncommon shapes fall back or fail

Which quantized kernel runs where? Quantized matmul in the MLX CUDA backend dispatches across several kernels by compute capability. The dedicated Hopper WGMMA kernel (qmm_sm90) runs on SM 90 and is compiled only when the binary targets 90a; because mlxcel always passes MLX_CUDA_ARCHITECTURES explicitly, the build script appends that a suffix for Hopper (see above). The tensor-core qmm_sm80 kernel covers SM ≥ 80 (Ampere, Ada, Hopper, Blackwell). The generic qmm_naive/qmv path has no architecture gate and runs on any supported GPU. So 4-bit inference works across the lineup, with Hopper taking the fastest dedicated path; the GB10 (SM 121) entries in the CUDA model tests run dozens of 4-bit models. Non-quantized bf16/fp16 models work on every GPU but use roughly 4x more memory, and a few uncommon matmul shapes have no optimized kernel and fall back to the naive path or fail.


ROCm Build Options

HIP Architecture Selection

The build compiles MLX device code for the gfx targets that rocminfo reports on the build host. Set MLX_ROCM_ARCHITECTURES to choose them explicitly, for example on a host without a GPU or for several targets:

MLX_ROCM_ARCHITECTURES=gfx1151 cargo build --release --features rocm
MLX_ROCM_ARCHITECTURES="gfx1100;gfx1151" cargo build --release --features rocm

Every entry must be a HIP target ID: gfx followed by lowercase hex digits (gfx1151), a generic target (gfx11-generic), or a target with feature suffixes (gfx90a:xnack+). Anything else, such as native, an uppercase name, or an option string, stops the CMake configure step with Not a HIP target ID, because each entry becomes an --offload-arch compiler argument.

If neither the variable nor rocminfo yields a target, the build fails and names the variable. The chosen list is recorded in the binary. mlxcel generate prints a HIP architecture <gfx target>; compiled for [<list>] line at startup, MLXCEL_TRACE_ARCH=1 prints the same picture from any entry point, and a device whose target is not in the compiled list is refused at runtime initialization with both named. Set ROCM_PATH when ROCm is not installed under /opt/rocm; the build reads its CMake packages, hipcc, and rocminfo from there.

Running on ROCm

The binaries carry $ROCM_PATH/lib as an rpath, so LD_LIBRARY_PATH is not needed. The server logs Runtime device: GPU at startup, and a GPU process appears in rocm-smi --showpids. On a UMA host the GPU shares memory with the operating system and with other GPU processes, so check rocm-smi --showpids for other tenants before loading a large model.

To take the per-row MoE kernels instead of the expert-batched gather_qmm prefill kernel (for example to A/B a throughput change or to bisect a MoE regression), set MLX_ROCM_GATHER_QMV_EXPERT_BATCHED=0. The variables the ROCm backend reads at run time are listed in ROCm Backend Variables. They are tuning and bisection switches rather than a stable interface.

Memory Footprint

Two defaults keep the allocator's footprint close to the weights (#2062). Measured on gfx1151 at pp512/tg128, the MLX peak for Meta-Llama-3.1-8B-Instruct-4bit is 6.14 GB with them and 20.60 GB without them, with decode throughput unchanged and prefill within 2%.

  • In-flight bound, MLX_ROCM_MAX_INFLIGHT_MB (default 1024). A prefill allocates a transient for every operation, and each one is released only when its command batch finishes on the GPU. The backend commits a batch once it has allocated a quarter of this budget and waits for the oldest batch while more than the budget is in flight, so the transients stay within roughly this many MiB. 0 removes the bound.
  • Buffer-cache bound, MLXCEL_CACHE_LIMIT (default 2 GiB on ROCm). Freed buffers are cached for reuse, and the ROCm allocator reuses only a buffer of exactly the requested size. A ROCm build caps that cache at 2 GiB unless MLXCEL_CACHE_LIMIT says otherwise; 0 or none removes the cap.

The pre-load memory estimate counts the in-flight budget as its own Backend in-flight line in mlxcel inspect (backend_inflight_bytes in --json output), on top of the weights, KV cache, headroom, and activation reserve (#2155). See Memory Estimation.

Current Status

Area Status on ROCm
Affine 4-bit / 8-bit checkpoints Run natively, dense and MoE.
mxfp8 and mxfp4 checkpoints Run natively, including MoE experts through gather_qmm. The load log says so once per mode (Quantization mode mxfp4: running on native ROCm kernels, no load-time conversion ...). Vendor FP8 block checkpoints (quant_method: fp8, 128x128 blocks) are requantized to mxfp8 at load and run natively.
NVFP4 checkpoints No native kernel. ModelOpt NVFP4 checkpoints are converted to affine 4-bit at load with no environment variables, and the load log names the route and the reason; a layer that cannot be converted fails the load with its name. MLXCEL_NVFP4_NATIVE_REPACK is ignored on ROCm. MLX-native NVFP4 exports ("mode": "nvfp4" in config.json, such as mlx-community/*-nvfp4) have no load-time conversion and are refused at load; use an affine export instead.
Affine MoE models (for example Qwen3-30B-A3B) Run natively. Single-token decode takes the fused two-kernel MoE path (gate-up, then down) through its HIP port, as on Metal and CUDA; MLXCEL_FUSED_MOE=0 selects the gather_qmm graph path instead. Experts wider than MLXCEL_FUSED_MOE_MAX_DFF (8192 by default, as on CUDA) stay on gather_qmm.
Model-specific fused kernels (BitNet bitlinear_matmul, Apertus xIELU, Cohere2 add3 + LayerNorm, Mamba1 selective scan, Nemotron-H's opt-in MLXCEL_FUSED_MOE_RELU2) Run as HIP ports.
Fused residual-add RMSNorm and RoPE + KV append (MLXCEL_FUSED_ADD_RMSNORM, MLXCEL_FUSED_ROPE_APPEND) On by default on ROCm, off on Metal and CUDA. Both HIP ports produce output byte-identical to the graph they replace. Set either variable to 0 to use the graph.
Paged attention (batched paged decode, MLA split-KV, sparse paged decode) The v1 decode, v2 partial, and merge kernels run as HIP ports, so these paths take the fused kernels instead of gather-then-SDPA, as on Metal and CUDA. MLXCEL_PAGED_ATTENTION_NATIVE=0 forces the gather path.
Sampled decode (temperature > 0) The fused Gumbel-max sampler and the rejection sampler run as HIP kernels. Top-k combined with top-p takes the rejection kernel up to a vocabulary of 152064 on ROCm (32768 on Metal and CUDA). MLXCEL_SAMPLING_GUMBEL=0 and MLXCEL_SAMPLING_REJECTION=0 force the graph.
Hybrid SSM decode (granite-4.0-h, falcon-h1, plamo-2, Nemotron-H) The single-token Mamba2 SSM update runs the fused HIP kernel instead of the SSD graph. MLXCEL_SSM_KERNEL=0 forces the graph.
FFT, Hadamard, and SearchSorted Supported. FFT runs on hipFFT; MLX_ROCM_FFT_CACHE_SIZE (default 128, as on CUDA) bounds how many transform plans stay alive, and a value that is not a positive integer is ignored with a warning.
mlxcel-server chat completions Work for dense and affine MoE checkpoints, streaming and non-streaming.
Audio (speech to text, text to speech) Works.
Memory estimation on UMA hosts Correct. The estimator reads the ROCm allocator's cap (76.80 GiB of the 96 GiB carve-out on the tested host) before it would reach host RAM, and adds the Backend in-flight reserve described in Memory Footprint.
GPU faults Reported as errors. A launch HIP rejects fails the evaluation that issued it and the device stays usable. An asynchronous fault (an out-of-bounds access) fails the waiting evaluation within about a second; after a queue fault the HIP runtime rejects every later call, so the process has to restart. A kernel that never finishes is not detected unless MLX_ROCM_GPU_WATCHDOG_SECS (default 0, off) is set, which fails any single host wait that outlives it.
Diagnostics Report the AMD vendor, device name, gfx target, and device memory. A binary whose compiled gfx list does not cover the device refuses to start.
Wave64 devices (CDNA: MI200 gfx90a, MI300 gfx942) Not run. The fused HIP kernels that reduce across lanes have run only on 32-lane wavefronts, so mlxcel selects them only when the device reports a 32-lane wavefront; on a 64-lane device those paths use their MLX graph fallbacks, and stderr says so once per process. BitLinear has no graph fallback, so BitNet checkpoints are refused at load on a wave64 device.
CPU device on a ROCm build (MLXCEL_DEVICE=cpu) Runs, slowly: a correctness reference and an escape hatch for a mismatched gfx build, not a serving mode. BLAS work on this device runs on one OpenBLAS thread.
Model catalog mlxcel arch --json and the WebUI report every family's ROCm support as partial.
Windows Supported since 0.8, source build only, single node. See Windows (AMD ROCm).
Multiple GPUs, distributed inference Not supported.

For orientation only (greedy decode, short prompt, tested host): Qwen3-0.6B-4bit runs at about 250 tok/s and gpt-oss-20b-MXFP4-Q4 at about 8 tok/s.


Experimental Support

The Windows builds build from source and have run on one validated host each, and neither ships a prebuilt archive. Windows CUDA is experimental; Windows ROCm has been supported since 0.8 and is validated on the configuration below. Expect rough edges outside the validated configuration, and report problems on the issue tracker.

Platform Feature Validated configuration Scope
Windows (NVIDIA CUDA) cuda NVIDIA RTX Spark N1X (aarch64-pc-windows-msvc, sm_121), CUDA 13.4.3, cuDNN 9.27 Single node. Runs against the CUDA toolkit directory it was built with, or from a runtime bundle.
Windows (AMD ROCm) rocm Radeon 8060S (x86_64-pc-windows-msvc, gfx1151), ROCm 10.1.0 Single node. Runs from a developer SDK or from a relocatable runtime bundle.

WSL2 with the Linux CUDA build remains the supported way to run mlxcel on a Windows machine with an NVIDIA GPU.

Windows Builds in Common

Both Windows builds produce native mlxcel.exe and mlxcel-server.exe with the MSVC toolchain. The GNU toolchain does not work, because nvcc and ROCm's clang compile for the MSVC ABI.

Requirement Details
Visual Studio 2022 or newer with the "Desktop development with C++" workload. Its bundled CMake and Ninja are enough.
Rust The MSVC target of the machine (x86_64-pc-windows-msvc, or aarch64-pc-windows-msvc for CUDA on Windows on Arm), not a -gnu target. The checkout's rust-toolchain.toml selects the version.
Shell A Visual Studio developer prompt for the target architecture ("x64 Native Tools" or "ARM64 Native Tools Command Prompt"), so cl.exe, CMake, and Ninja are on PATH.

The build uses the Ninja generator for MLX unless CMAKE_GENERATOR is set; the Visual Studio generator cannot compile .cu files without NVIDIA's MSBuild integration, and cannot run the HIP compile commands at all. The MLX dependency tree under target\ is deep, so if CMake, nvcc, or hipcc reports a path it cannot open, enable Win32 long paths or set a short CARGO_TARGET_DIR such as C:\t.

The build copies the OpenBLAS DLL that MLX's CMake downloads for its CPU backend next to the executables. Keep it beside mlxcel.exe when you move the binaries.

Both Windows builds share these limits:

Area Status
Single-GPU generation and mlxcel-server Works on the validated configuration.
Tensor and pipeline parallelism, disaggregated serving, --pp-auto, --pp-peer Refused at startup with an error naming the flag.
Unix domain socket listen (--host *.sock) Refused; listen on TCP.
Router-mode model downloads and cache removal (WebUI model library) Refused: they depend on directory-fd-anchored filesystem operations, which have no Windows implementation.
CPU device (MLXCEL_DEVICE=cpu, or a build without a GPU feature) Builds and loads, but quantized matmul runs MLX's scalar CPU loops. On the CUDA host a 24-token generate with gemma-4-e4b-it-qat-4bit had not finished after 30 minutes.
Prebuilt archives, CI No archives. CI is compile-only, with no GPU: the hosted Windows build (compile-only) job runs when the Windows build scripts, build_support/, the overlays' runtime-directory sources, the bundle packagers or the MLX pin change, and it covers the mlxcel-mlx-pin unit tests, the tests/ Python suites, and cargo check -p mlxcel-core --features rocm against the ROCm 10.1.0 gfx1151 TheRock SDK. No CI job builds the cuda feature on Windows.

Windows (NVIDIA CUDA)

Validated on aarch64-pc-windows-msvc with an NVIDIA RTX Spark N1X (Blackwell, sm_121, driver 616.62), Visual Studio 2026 (MSVC 14.51), CUDA 13.4.3, and cuDNN 9.27. The build resolves x64 and arm64 toolkit layouts alike, but no x86_64-pc-windows-msvc GPU host has run it yet.

Prerequisites, in addition to Windows Builds in Common:

  • CUDA Toolkit 13.x and cuDNN 9 built for CUDA 13. CUDA 12.x is not handled, because its Windows layout keeps the DLLs in bin\ rather than bin\<arch>\. On x64 the CUDA installer works. cuDNN must use the redistributable archive layout; the cuDNN installer's lib\13.x\x64 layout is not searched.
  • On Windows on Arm, NVIDIA ships only the redistributable archives (the windows-arm64 entries under developer.download.nvidia.com/compute/cuda/redist/ and .../cudnn/redist/). Unpack the cuda_cudart, cuda_nvcc, cuda_crt, cuda_nvrtc, cccl, libcublas, libcufft, libcusolver, libcusparse, libnvjitlink, and libnvvm archives and the cuDNN archive into one directory. No administrator rights are needed.
  • An NVIDIA display driver that supports the toolkit.

The build expects the toolkit layout bin\nvcc.exe, include\, lib\<arch>\, and bin\<arch>\, where <arch> is x64 or arm64. cuDNN may live in the same directory or in its own root with the same layout.

Build from the developer prompt:

set CUDA_PATH=C:\path\to\cuda
rem Only when cuDNN is not unpacked into CUDA_PATH:
set CUDNN_PATH=C:\path\to\cudnn
cargo build --release --features cuda

CUDA_HOME is honored ahead of CUDA_PATH. The GPU architecture is detected with nvidia-smi as on Linux, and MLX_CUDA_ARCHITECTURES overrides it (see CUDA Build Options).

The CUDA and cuDNN DLLs are delay-loaded from the toolkit's bin\<arch> directory recorded at build time, and NVRTC finds the CUDA headers in the recorded include\. The binaries therefore run, including the first JIT compile, without the toolkit on PATH and without CUDA_PATH set. The display driver's nvcuda.dll comes from the driver. Without a bundle, the binaries depend on that toolkit directory at run time. To run on another machine, build a runtime bundle.

Windows CUDA runtime bundle

A runtime bundle carries the CUDA, cuDNN and NVRTC DLLs and the headers the JIT compiles against, so mlxcel runs without the toolkit. Build one with scripts/package_windows_cuda_runtime.py (Python 3, standard library):

python scripts\package_windows_cuda_runtime.py --cuda-path C:\path\to\cuda --stage-dir dist\mlxcel --add-binary target\release\mlxcel.exe --add-binary target\release\mlxcel-server.exe

The binaries use one runtime directory, chosen in this order: MLXCEL_CUDA_RUNTIME_DIR (an absolute path), the executable's directory when it holds the cuBLASLt DLL, then the toolkit recorded at build time. They print the choice, for example [mlx-cuda] CUDA runtime: C:\mlxcel\bin (executable directory). In a bundle, a missing required DLL stops the binary with the full path it expected. cuFFT and cuSOLVER are not in that check and load when first called.

  • This was validated on an RTX Spark N1X (driver 617.74, CUDA 13.4, cuDNN 9.27) with the toolkit renamed: mlxcel generate from a cold JIT cache and a mlxcel-server chat request ran from the bundle folder and through MLXCEL_CUDA_RUNTIME_DIR. x64 GPU hosts and a machine with only the display driver have not been tried.
  • The bundle does not carry the Visual C++ runtime. openblas.dll is a normal import, so keep a copy beside a binary that runs outside the bundle.
  • With toolkit archives unpacked into one directory, the LICENSE files overwrite each other. Pass the cuda_cudart and cudnn archives' LICENSE with --cuda-license and --cudnn-license, and the NVIDIA/cccl release tag's LICENSE with --cccl-license.
  • The bundle is meant for your own machines. The CUDA EULA does not list every header the JIT includes as redistributable, so confirm the terms with NVIDIA before publishing one.

Windows (AMD ROCm)

A native mlxcel.exe and mlxcel-server.exe with the rocm feature, built with MSVC and supported since 0.8, source build only and single node. It shares the ROCm overlay with the Linux build and the single-node limits of the Windows builds above.

Validated on x86_64-pc-windows-msvc with an AMD Ryzen AI MAX+ 395 (Radeon 8060S, gfx1151, BIOS UMA carve-out), Windows 11, Adrenalin 32.0.31041.1004, ROCm Core SDK 10.1.0 (TheRock tarball), and Visual Studio 18 (MSVC 14.51). On that host mlx-community/Qwen3-0.6B-4bit decodes at about 115 tok/s and mlx-community/Qwen3-1.7B-bf16 at about 30 tok/s. The build accepts only x86_64-pc-windows-msvc.

Prerequisites, in addition to Windows Builds in Common:

  • A ROCm 10.x SDK for Windows. ROCm 10 has no HIP SDK installer; use the TheRock tarball for your GPU family from https://stable.repo.amd.com/rocm/core/tarball/ (for gfx1151, therock-dist-windows-gfx1151-10.1.0.tar.gz) or the SDK from AMD's ROCm pip wheels. The build needs bin\hipcc.exe, the import libraries in lib\, and amdgpu-arch.exe for target detection (the 10.1 tarball keeps it in lib\llvm\bin\).
  • An AMD display driver that supports the GPU.

Step 1: Install the SDK

Extract the tarball with the tar that ships with Windows, then check that C:\rocm\10.1.0\bin\hipcc.exe exists (if the archive has a top-level folder, move its contents up one level):

mkdir C:\rocm\10.1.0
tar -xzf therock-dist-windows-gfx1151-10.1.0.tar.gz -C C:\rocm\10.1.0

The SDK directory must be writable only by administrators (or only by you). The binaries load DLLs from its bin\ and compile GPU kernels from the headers in its include\, so any account that can write there can run code in every mlxcel process. A folder created directly under C:\ inherits Modify access for all authenticated users from the drive root; on a machine with more than one account, remove that access from an administrator prompt:

icacls C:\rocm /inheritance:r /grant:r *S-1-5-32-544:(OI)(CI)F *S-1-5-18:(OI)(CI)F *S-1-5-32-545:(OI)(CI)RX

This leaves Administrators and SYSTEM with full control and Users with read and execute. Afterwards icacls C:\rocm\10.1.0\bin lists no Authenticated Users entry.

Step 2: Build

From the "x64 Native Tools Command Prompt":

set ROCM_PATH=C:\rocm\10.1.0
set CARGO_TARGET_DIR=C:\t\mlxcel
cargo build --release --features rocm

HIP_PATH is read when ROCM_PATH is unset. The path must be absolute; the build fails on a relative one, and also when the SDK bin holds no amdhip64_*.dll. The HIP targets are detected with amdgpu-arch (Windows SDKs ship no rocminfo). Set MLX_ROCM_ARCHITECTURES=gfx1151 (semicolon-separated for several) to skip detection or to build for a GPU that is not installed; the binary runs only on the targets it was built for.

A full release build took about 18 minutes on the validated host. Windows sees about 32 GB of RAM there after the GPU carve-out, and one build failed at link time (link.exe exit 0xC0000142) while memory was low. Limiting parallelism completed it:

set CMAKE_BUILD_PARALLEL_LEVEL=8
cargo build --release --features rocm -j 8

Step 3: Check the build

ROCM_PATH is needed only at build time. Run the binary from a new prompt:

C:\t\mlxcel\release\mlxcel.exe generate -m mlx-community/Qwen3-0.6B-4bit -p "What is the capital of France?" -n 32 --temp 0

The first lines of stderr name the runtime directory and the GPU, for example [mlx-rocm] ROCm runtime: C:\rocm\10.1.0\bin (build-time ROCm SDK) and [mlx-rocm] bound HIP device 0: gfx1151 (...). The first run also compiles fused kernels with hiprtc and caches them under ~/.cache/mlxcel/hsaco or MLXCEL_CACHE_DIR (MLX_HSACO_CACHE_DIR moves the cache, and a directory another user can write is refused).

A developer build like this loads the ROCm DLLs from the SDK bin recorded at build time, so it runs only on a machine with the SDK at that path. To run it elsewhere, build a runtime bundle.

Windows ROCm Runtime Bundle

The runtime bundle carries the ROCm DLLs, kernel data, and headers mlxcel needs, so the target machine needs only the AMD display driver. scripts/package_windows_rocm_runtime.py (Python 3, standard library) builds it from an SDK. Pass --arch once per GPU target the binaries were built for:

python scripts\package_windows_rocm_runtime.py --rocm-path C:\rocm\10.1.0 --arch gfx1151 --out dist

This writes dist\mlxcel-rocm-runtime-10.1.0-win-x64.zip (76 MiB, 225 MiB extracted, for gfx1151 with ROCm 10.1.0). To produce a ready-to-run folder with the binaries in it, stage the tree instead of zipping it:

python scripts\package_windows_rocm_runtime.py --rocm-path C:\rocm\10.1.0 --arch gfx1151 --stage-dir dist\mlxcel --add-binary C:\t\mlxcel\release\mlxcel.exe --add-binary C:\t\mlxcel\release\mlxcel-server.exe --add-binary C:\t\mlxcel\release\libopenblas.dll

The script fails when the SDK lacks any file it packs, so a DLL renamed in a later ROCm release is caught when packaging. The bundle keeps the SDK's folder layout, because rocBLAS, hipBLASLt, and rocFFT read their device code from ..\.kpack\ relative to their own DLL:

Folder Contents
bin\ The ten DLLs mlxcel loads (amdhip64_7, amd_comgr, rocm_kpack, hiprtc0716, hiprtc-builtins0716, rocblas, libhipblaslt, origami, hipfft, rocfft) and the rocBLAS and hipBLASLt kernel data, reduced to the requested targets. This is the runtime directory.
.kpack\ Device-code archives, opened as <DLL folder>\..\.kpack\<library>_<target>.kpack.
include\hip\ Headers the hiprtc kernel compiler needs.
licenses\ SDK license files of the shipped components. The ROCm 10.1.0 tarball has none for the HIP runtime, hipRTC, or rocm_kpack.

To install on the target machine, extract the bundle to a short path such as C:\mlxcel and use one of two layouts:

  • Put mlxcel.exe, mlxcel-server.exe, and libopenblas.dll in the extracted bin\ (the --add-binary folder above already has them) and run C:\mlxcel\bin\mlxcel.exe.
  • Keep the binaries elsewhere and set MLXCEL_ROCM_RUNTIME_DIR=C:\mlxcel\bin.

The binaries pick one runtime directory per process and print it once, for example [mlx-rocm] ROCm runtime: C:\mlxcel\bin (executable directory). The candidates are tried in this order:

  1. MLXCEL_ROCM_RUNTIME_DIR, which must be an absolute path. When it is set it is the only candidate: a relative path, or a folder without the HIP runtime, stops the binary with a message naming the problem.
  2. The folder of mlxcel.exe or mlxcel-server.exe, when it holds the HIP runtime the binary imports (amdhip64_7.dll for ROCm 10).
  3. The SDK bin recorded at build time, when it holds the HIP runtime.

Every ROCm DLL comes from the chosen folder. A DLL missing there stops the binary with the full path it expected, and is never taken from another candidate, from the working directory, from PATH, or from the older amdhip64_7.dll and amd_comgr.dll the AMD display driver installs in System32; DLLs from two ROCm releases are not guaranteed to work together (the SDK's HIP runtime next to the driver's amd_comgr.dll finds no device). The binary also stops before loading any ROCm DLL when the folder holds no amd_comgr*.dll or no hiprtc-builtins*.dll, which the HIP runtime opens by name.

Install rules:

  • Give each application its own runtime folder. llama.cpp's Windows ROCm release, for example, ships ROCm 10.0 amdhip64_7.dll, amd_comgr.dll, and rocm_kpack.dll beside its executables; extracting both into one folder replaces those files and mixes two releases. Point MLXCEL_ROCM_RUNTIME_DIR at mlxcel's folder when the two must sit side by side.
  • The runtime folder, its parent (which holds .kpack\ and include\), and the folder holding mlxcel.exe must be writable only by administrators (or only by you), for the reasons given in Step 1. C:\Program Files has that protection; for a folder directly under C:\, apply the icacls command above.
  • Keep the runtime folder path at most 130 characters long. rocBLAS and hipBLASLt open kernel files up to 128 characters below it, and with the 260-character Windows path limit a 133-character location made the first matrix multiply crash with an access violation. The binaries print a warning when the chosen folder is longer.
  • ROCM_HOME, ROCM_PATH, and HIP_PATH, when set, take precedence over the bundle's include\ for the JIT headers. The AMD HIP SDK installer sets HIP_PATH system-wide; on such a machine clear it for mlxcel (set HIP_PATH=) or point it at the bundle root, so kernels are not compiled against another ROCm version's headers.
  • The bundle was tested with Adrenalin 32.0.31041.1004 on a Radeon 8060S (gfx1151). Other driver versions and GPUs are untested.

Windows ROCm builds use a host-shadow memory model: Windows does not map GPU allocations into the CPU address space, even on an APU, so CPU access to a GPU buffer goes through a pinned host copy that is flushed back before the next kernel reads it. On an integrated GPU the allocator limit is 80% of the device memory HIP reports (86.29 GiB of 107.87 GiB on the validated host); unlike Linux, system RAM does not raise it. FFT (hipFFT and rocFFT) ships in the bundle but has not been run on Windows.


Build Output

A successful build produces two binaries under target/release/ (on Windows, mlxcel.exe and mlxcel-server.exe under target\release\, or under CARGO_TARGET_DIR when it is set, with the OpenBLAS DLL beside them):

Binary Path Description
mlxcel target/release/mlxcel Main CLI with all 13 commands: run, generate, embed, rerank, serve, list, arch, inspect, download, detect, rm, tune, and split-mtp (the last one with the default surgery feature)
mlxcel-server target/release/mlxcel-server Standalone HTTP server, designed as a llama-server drop-in replacement; also includes the download subcommand

You can copy these binaries to a directory on your PATH (e.g., /usr/local/bin) for system-wide access:

sudo cp target/release/mlxcel target/release/mlxcel-server /usr/local/bin/

Troubleshooting

cannot execute tool 'metal' due to missing Metal Toolchain (macOS)

This error means the Metal shader compiler is not installed. The Metal Toolchain is a separate optional component from Xcode. Install it with:

xcodebuild -downloadComponent MetalToolchain

After installation, re-run cargo build.

Cannot find CUDA library directory (Linux)

The build script cannot locate CUDA in the default path (/usr/local/cuda). Set CUDA_HOME to your actual CUDA installation directory:

export CUDA_HOME=/usr/local/cuda-13
cargo build --release --features cuda

To make this permanent, add the export line to your shell profile (~/.bashrc, ~/.zshrc, etc.).

nvidia-smi: command not found during build (Linux)

The build script uses nvidia-smi to detect your GPU's compute capability automatically. On headless build servers, containerized CI environments, or machines without GPU drivers installed, nvidia-smi may not be available.

Set the target architecture manually to bypass auto-detection:

# Use your GPU's SM version; append `a` for SM 90 and above (e.g. 90a).
MLX_CUDA_ARCHITECTURES=90a cargo build --release --features cuda

error[E0463]: can't find crate for 'std' or linker errors (Linux)

Install the toolchain pinned by rust-toolchain.toml (Rust 1.97.1 in v0.8.0). Running Cargo in the checkout selects it automatically when rustup is installed:

rustup toolchain install 1.97.1 --component rustfmt --component clippy
rustc --version  # should show 1.97.1 in this checkout

Build takes a very long time (Linux)

This is expected on first build. The MLX C++ library, its CUDA kernels, and supporting libraries (fmt, gguflib) are all compiled from source. Subsequent incremental builds are much faster. On a modern workstation with a fast NVMe drive, the first full CUDA build typically takes 10-20 minutes.

CMake, nvcc, or hipcc cannot open a path (Windows)

The MLX dependency tree under target\ is deep enough to exceed the 260-character Windows path limit. Enable Win32 long paths, or build into a short directory:

set CARGO_TARGET_DIR=C:\t\mlxcel

mlxcel: cannot find the ROCm runtime (Windows)

The binary found no amdhip64_7.dll beside itself and the SDK it was built against is not at its build path. Put the binary in the bin\ folder of an extracted runtime bundle, or set MLXCEL_ROCM_RUNTIME_DIR to that folder. A message that MLXCEL_ROCM_RUNTIME_DIR must be absolute or holds no amdhip64_7.dll means the variable is set to a relative path or to the wrong folder; it is not ignored in favor of another candidate.

holds no amd_comgr*.dll or cannot load <name>.dll from the ROCm runtime directory (Windows)

The chosen runtime folder is incomplete, usually because files were copied by hand or a file was deleted. Extract the bundle again, or rebuild it with scripts\package_windows_rocm_runtime.py, which fails when the SDK lacks a file. Do not fill the gap with DLLs from another ROCm installation or from System32.