콘텐츠로 이동

소스에서 빌드하기

이 페이지에서는 macOS(Apple Silicon), Linux(CUDA, AMD ROCm), 그리고 Windows(NVIDIA CUDA, AMD ROCm)에서 소스로 mlxcel을 빌드하는 방법을 설명합니다. ROCm 지원은 두 운영체제 모두 0.8부터이며, Windows 빌드와 마찬가지로 미리 빌드한 아카이브가 없습니다. Windows CUDA 빌드는 아직 실험 단계입니다. macOS Apple Silicon에서 가장 빠르게 패키지로 설치하려면 빠른 시작의 Homebrew 경로를 사용하십시오. Linux 빌드는 CUDA 전용이라 Homebrew tap에 포함되지 않으므로 GitHub 릴리스 페이지에서 mlxcel-linux-*-cuda13 아카이브를 받거나 이 페이지의 방법대로 소스에서 빌드하십시오. v0.8.0 릴리스는 프리릴리스이고 첨부된 아카이브가 없으며 releases/latest는 이전 안정 릴리스를 가리키므로, v0.8.0을 실행하려면 git clone --branch v0.8.0 https://github.com/lablup/mlxcel.git으로 태그를 받아 빌드하십시오. 소스 빌드는 로컬 컴파일러 플래그, 아직 릴리스되지 않은 변경사항, 하드웨어별 CUDA 설정이 필요할 때 유용합니다.

목차


사전 요구사항

macOS (Apple Silicon)

mlxcel은 Apple Silicon(M1, M2, M3, M4, M5 시리즈)에서 동작합니다. Intel Mac은 지원하지 않습니다. Apple의 MLX 프레임워크가 Apple Silicon 통합 메모리 아키텍처를 필요로 하기 때문입니다.

요구사항 세부 내용
하드웨어 Apple Silicon Mac (M1 / M2 / M3 / M4 / M5 시리즈)
Rust 2024 에디션. v0.8.0 소스 트리는 Rust 1.97.1을 고정 사용
Xcode Command Line Tools 시스템 헤더 및 링커에 필요
Metal Toolchain MLX Metal GPU 셰이더 컴파일에 필요

Rust 툴체인이 없다면 rustup.rs에서 설치하십시오.

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

Xcode Command Line Tools를 설치합니다.

xcode-select --install

Metal Toolchain을 설치합니다. MLX 셰이더 컴파일에 필요하며 Xcode와는 별도의 구성요소입니다.

xcodebuild -downloadComponent MetalToolchain

Linux (CUDA)

Linux 빌드는 CUDA 백엔드를 통해 NVIDIA GPU를 대상으로 합니다. Grace Hopper 및 유사한 ARM 기반 NVIDIA 플랫폼용 aarch64와 x86_64 시스템 모두 지원됩니다.

요구사항 세부 내용
아키텍처 aarch64 또는 x86_64
Rust 2024 에디션. v0.8.0 소스 트리는 Rust 1.97.1을 고정 사용
CMake 3.24 이상 (MLX C++ 라이브러리 빌드에 필요)
CUDA Toolkit 12.0 이상, nvcc 포함
cuDNN 9 이상
시스템 라이브러리 libopenblas-dev, liblapack-dev, liblapacke-dev

Ubuntu 또는 Debian에서 시스템 의존성을 설치합니다.

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

참고: 설치된 CUDA 버전에 맞게 cuDNN 패키지명을 조정하십시오. 예를 들어 CUDA 12.x에서는 libcudnn9-dev-cuda-12를 사용합니다.

Linux (AMD ROCm)

AMD GPU 지원은 Linux에서 0.8부터이며 소스 빌드만 지원합니다. 릴리스 아티팩트가 없고, 선택 사항인 ROCm CI 작업은 저장소가 ROCM_CI_ENABLED를 설정하지 않으면 꺼져 있습니다. 이 백엔드는 상류 MLX에 포함되어 있지 않으며, mlxcel은 NripeshN/mlx의 rocm-support 브랜치(MIT)를 src/lib/mlxcelverse/overlays/rocm/에 가져와 Metal과 CUDA 빌드가 쓰는 것과 같은 고정 MLX 커밋 위에 적용합니다. 포크 커밋은 src/lib/mlxcelverse/overlays/rocm/UPSTREAM에 기록되어 있고, mlxcel이 그 위에 더한 변경은 같은 디렉터리의 LOCAL_FIXES.md에 정리되어 있습니다.

검증한 구성은 AMD Ryzen AI MAX+ 395와 Radeon 8060S(gfx1151, RDNA 3.5), 96 GiB VRAM carve-out, Debian 13, ROCm 10.0.0 패키지(HIP 7.15)입니다. 다른 RDNA 3, 3.5, 4 제품도 빌드될 것으로 예상합니다. CDNA 제품(예: MI300)은 컴파일되지만 튜닝은 없고, 32-lane wavefront에 걸쳐 reduce하는 fused 커널은 이 제품에서 MLX 그래프 대체 경로를 탑니다(현재 상태의 Wave64 행 참조).

요구사항 세부사항
Rust, CMake, C++ Rust 2024 edition, CMake, C++20 컴파일러
시스템 라이브러리 pkg-config, OpenSSL 헤더, libopenblas-dev, liblapack-dev, liblapacke-dev
ROCm CMake 패키지 hip, rocblas, rocthrust, rocprim, hiprand, rocwmma, hipblaslt, hipfft, hiprtc와 hipcc, rocminfo
장치 접근 빌드하고 실행하는 사용자가 video와 render 그룹에 속해야 합니다
sudo apt-get install -y \
    build-essential cmake git pkg-config libssl-dev \
    libopenblas-dev liblapack-dev liblapacke-dev
# ROCm 자체는 AMD 저장소(repo.radeon.com)에서 설치합니다.
rocminfo | grep -E '^\s+Name:\s+gfx'   # GPU 타겟이 나와야 합니다

빌드 방법

macOS 빌드

# 저장소 클론
git clone https://github.com/lablup/mlxcel.git
cd mlxcel

# Metal 및 Accelerate 프레임워크 지원으로 빌드
cargo build --release --features metal,accelerate

metal 피처는 MLX C++ 라이브러리를 통해 Apple의 Metal GPU 가속을 활성화합니다. accelerate 피처는 모델 로딩 및 CPU 폴백 경로에서 사용되는 최적화된 CPU 선형 대수 연산을 위해 Apple의 Accelerate 프레임워크를 활성화합니다.

macOS에서는 feature 없이 cargo build --release만 실행해도 MLX가 Metal과 Accelerate를 켭니다. 위 명령의 명시적 feature는 의도한 백엔드를 문서로 남기는 역할을 합니다. MTP qmv_wide 재시도 같은 런타임 제어는 feature 지정 여부와 관계없이 실제로 빌드된 백엔드를 따르며, MLXCEL_BUILD_METAL=OFF를 주면 Cargo metal feature를 선택했더라도 Metal 백엔드와 그 제어가 꺼집니다.

체크포인트를 로드하지 않고 이 런타임 제어를 확인하려면 다음을 실행합니다.

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

이 probe는 QMV 선택과 decode command-buffer override를 확인한 뒤 이전 값으로 되돌리며, Metal 장치가 없으면 건너뛰었다고 보고합니다.

lablup/mlxcel#1988의 영향을 받는 이전 plain macOS 빌드라면 서버 명령 앞에 MLXCEL_QMV_WIDE=0을 붙여 새 프로세스를 시작하십시오. 시작 시점에 더 좁은 Metal projection 커널을 고정하고 exactness gate는 그대로 유지합니다. 처리량이 줄 수 있고 Metal에만 적용되며, 로드한 모델의 probe는 여전히 통과해야 합니다.

Linux 빌드

# 저장소 클론
git clone https://github.com/lablup/mlxcel.git
cd mlxcel

# CUDA 지원으로 빌드
cargo build --release --features cuda

첫 빌드는 MLX C++ 라이브러리와 해당 CUDA 커널이 Cargo 빌드 프로세스의 일부로 소스에서 컴파일되기 때문에 몇 분 정도 소요됩니다.

Linux ROCm 빌드

cargo build --release --features rocm

rocm은 cuda나 metal과 함께 쓸 수 없으며, 함께 지정하면 충돌을 밝히며 빌드가 실패합니다. 첫 빌드는 hipcc로 MLX 장치 코드를 컴파일하므로 Rust 빌드에 몇 분이 더 걸립니다. 타겟 선택과 현재 상태는 ROCm 빌드 옵션을 참조하십시오.

증분 빌드는 헤더 의존성을 추적합니다. hipcc가 HIP 오브젝트마다 의존성 파일을 남기므로 src/lib/mlxcelverse/overlays/rocm/ 아래 헤더를 고치면 다음 빌드는 그 헤더를 직접 또는 간접으로 include하는 .hip 파일만 다시 컴파일하고, 바뀐 것이 없으면 아무것도 다시 컴파일하지 않습니다. overlay는 CMake의 configure_file로 빌드 트리에 복사되는데, 이 함수는 내용이 다를 때만 사본을 다시 쓰기 때문에 overlay 헤더를 touch만 해서는 아무것도 다시 빌드되지 않습니다. 의존성 파일에는 ROCm과 시스템 헤더도 들어가지만, 장치 컴파일에서만 include하는 헤더(rocWMMA)와 컴파일러 자체는 빠져 있습니다.

그래도 HIP를 깨끗하게 다시 빌드해야 한다면(예: hipcc나 rocWMMA를 올린 뒤, 또는 bisect 중 오래된 오브젝트를 배제하려고) 사용하는 빌드 프로파일의 HIP 오브젝트를 지우고 overlay 파일 하나를 touch해서 Cargo가 빌드 스크립트를 다시 실행하게 합니다.

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

release는 실제로 빌드하는 프로파일 디렉터리(debug, test-fast 등)로 바꾸십시오. touch를 빼면 Cargo가 바뀐 입력이 없다고 보고 빌드 스크립트를 건너뛰어 이전 커널을 그대로 링크합니다. cargo clean -p mlxcel-core --release로도 다시 빌드할 수 있지만, 이 경우 MLX C++ 라이브러리 전체를 다시 빌드합니다.


Cargo feature 플래그

두 바이너리가 같은 루트 패키지에서 빌드되므로 하나의 feature 조합이 mlxcel과 mlxcel-server에 함께 적용됩니다. cargo build --features <a,b>로 넘기십시오. 배포 빌드는 플랫폼 백엔드 플래그만 켜고, 나머지는 opt-in seam이거나 테스트용 발판입니다.

Feature 기본값 효과
webui on 번들 관리용 WebUI 자산을 두 바이너리에 내장하고 --webui를 활성화합니다. 이 feature 없이 빌드한 바이너리는 시작 시 --webui를 거부합니다. WebUI를 참조하십시오.
surgery on weight-load surgery. 로드 시점의 scale, add, prune, replace, interpolate weight-space 편집을 위해 --surgery <config.yaml>과 MLXCEL_SURGERY를 노출합니다. surgery 설정을 주지 않으면 로드 경로는 이 feature 없이 빌드한 것과 바이트 단위로 같습니다.
metal off Apple Silicon Metal GPU 백엔드. macOS 표준입니다.
accelerate off Apple Accelerate CPU BLAS 백엔드. macOS 표준입니다.
cuda off NVIDIA CUDA GPU 백엔드. NVIDIA 호스트에 필요하며, 켜지 않으면 CPU 전용 빌드가 됩니다. Windows에서도 네이티브로 빌드됩니다(실험 단계, Windows (NVIDIA CUDA) 참조).
rocm off Linux와 Windows용 AMD ROCm/HIP GPU 백엔드로, src/lib/mlxcelverse/overlays/rocm/의 ROCm overlay로 빌드합니다. 0.8부터 지원하며 소스 빌드만 지원합니다. cuda, metal과 함께 쓸 수 없습니다. Windows 빌드는 Windows (AMD ROCm)을 참조하십시오.
experimental-backend off MLX가 아닌 compute backend seam 자리를 예약합니다. 커널을 싣지 않고 런타임 dispatch도 늘리지 않습니다.
xla-backend off OpenXLA/StableHLO seam과 그래프 emitter까지. 실행 엔진은 없고 네이티브 툴체인도 필요 없습니다.
xla-iree off xla-backend에 실제 IREE 런타임을 더합니다. 빌드 시점에 IREE_DIST, 또는 IREE_CUDA_HOME, IREE_ROCM_HOME, IREE_MACOS_HOME 중 하나와 짝이 되는 compile 변수(IREE_CUDA_COMPILE, IREE_ROCM_COMPILE, IREE_MACOS_COMPILE)가 필요합니다. IREE_ROCM_HOME은 gfx1151, gfx1100, gfx942용으로 소스 빌드한 HIP 지원 IREE 런타임입니다.
xla-diagnostics-cpu, xla-diagnostics, xla-diagnostics-rocm off eager MLX 경로와 상주 IREE 그래프를 비교하는 진단 전용 빌드입니다(예: Qwen2.5-VL vision oracle). xla-diagnostics-cpu는 xla-iree에 mlxcel-xla 진단 기능을 더한 것이고, xla-diagnostics는 여기에 cuda를, xla-diagnostics-rocm은 AMD GPU oracle 실행을 위해 rocm을 더합니다. 배포 빌드용이 아닙니다.
test-utils off 테스트 전용 헬퍼. distributed_integration, pipeline_e2e, paged_handoff_parity 통합 테스트를 빌드하는 데 필요하며 바이너리에는 필요 없습니다.

default = ["surgery", "webui"]이므로 그냥 cargo build하면 surgery와 WebUI만 켜지고, 실제 빌드는 그 위에 플랫폼 백엔드를 얹습니다. --no-default-features로 빌드하면 둘 다 빠지므로 남길 것을 다시 지정하십시오:

# Metal + Accelerate, surgery 크레이트와 WebUI 없음.
cargo build --release --no-default-features --features metal,accelerate

# Metal + Accelerate와 surgery, WebUI 없음.
cargo build --release --no-default-features --features metal,accelerate,surgery

XLA feature가 무엇을 컴파일하고 빌드 시점에 무엇을 요구하는지, 그리고 어떤 모델 경로(Qwen2.5-VL vision tower 포함)가 IREE에서 실행되는지는 OpenXLA 백엔드를 참조하십시오.


CUDA 빌드 옵션

GPU 아키텍처 자동 감지

기본적으로 빌드 스크립트는 nvidia-smi를 통해 빌드 머신에 있는 GPU의 컴퓨트 캐퍼빌리티(SM 버전)를 확인하고 해당 타겟에 맞춰 CUDA 커널을 컴파일합니다. 이렇게 하면 가장 작은 바이너리와 하드웨어에 최적화된 커널 특수화가 만들어집니다. 감지한 값은 릴리스 워크플로와 같은 표기를 따릅니다. Hopper에는 CUDA의 아키텍처 전용 a 접미사를 붙이고(감지된 90은 90a가 됩니다), Blackwell(100, 120, 121)은 접미사 없이 둡니다. 감지에 실패하면 90a로 빌드합니다.

Blackwell에서 접미사 없는 표기로도 MLX의 하드웨어 NVFP4/MXFP4 변환기를 사용합니다. MLX CMake 빌드가 이 변환기를 쓰는 번역 단위 하나에만 아키텍처 전용 이미지를 추가하므로, 다른 커널의 코드를 바꾸지 않고 block-float 양자화가 하드웨어 경로를 탑니다. 목록에 121a를 넣지 마십시오. 장치의 모든 커널 머신 코드가 바뀌며, 121f는 컴파일에 실패합니다.

수동 아키텍처 지정

MLX_CUDA_ARCHITECTURES 환경 변수를 사용하여 하나 이상의 SM 버전을 직접 지정할 수 있습니다. 다음 상황에서 필요합니다.

  • GPU가 없는 머신에서 빌드하는 경우 (헤드리스 빌드 서버, CI 파이프라인)
  • 빌드 머신과 다른 GPU에서 실행할 바이너리를 만드는 경우
  • 단일 바이너리에서 여러 GPU 세대를 지원해야 하는 경우
# Hopper (H100, H200). 전용 Hopper 양자화 커널에는 `a` 접미사가 필요하며
# 그냥 `90`으로 빌드하면 이 커널이 빠집니다.
MLX_CUDA_ARCHITECTURES=90a cargo build --release --features cuda

# Blackwell GB10 (DGX Spark). 릴리스 워크플로와 같이 접미사 없이 빌드합니다.
MLX_CUDA_ARCHITECTURES=121 cargo build --release --features cuda

# 단일 바이너리에 여러 아키텍처 포함 (바이너리 크기 증가, 빌드 시간 증가)
MLX_CUDA_ARCHITECTURES="90a;121" cargo build --release --features cuda

사용자 지정 CUDA 경로

CUDA Toolkit이 기본 위치(/usr/local/cuda)가 아닌 곳에 설치된 경우, CUDA_HOME을 올바른 경로로 설정하십시오.

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

두 변수를 함께 사용할 수도 있습니다.

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

GPU 아키텍처 호환성 표

모든 GPU 세대가 mlxcel의 모든 기능을 지원하는 것은 아닙니다. 아래 표는 호환성을 요약한 것입니다.

GPU 패밀리 SM 버전 4비트 양자화 모델 fp16 / bf16 모델 비고
Turing (RTX 20xx) 75 가능(범용) 가능 텐서코어 양자화 커널 없음. 범용 qmm_naive/qmv 폴백 사용(느리고 fp16, 미검증)
Ampere (RTX 30xx, A100) 80 / 86 가능 가능 텐서코어 qmm_sm80(SM 80 이상)과 qmv 디코드 경로
Ada Lovelace (RTX 40xx) 89 가능 가능 텐서코어 qmm_sm80과 qmv
Hopper (H100, H200) 90 가능(WGMMA) 가능 전용 qmm_sm90 WGMMA 커널. 90a 빌드라야 컴파일됨
Blackwell (B100, B200) 100 가능 가능 텐서코어 qmm_sm80과 qmv
Blackwell GB10 (DGX Spark) 120 / 121 가능 가능 qmm_sm80과 qmv. 흔치 않은 일부 shape은 폴백하거나 실패

어느 양자화 커널이 어디서 도나요? MLX CUDA 백엔드의 양자화 matmul은 컴퓨트 캐퍼빌리티에 따라 여러 커널로 분기합니다. 전용 Hopper WGMMA 커널(qmm_sm90)은 SM 90에서 동작하고 바이너리가 90a를 타겟으로 빌드됐을 때만 컴파일됩니다. mlxcel은 MLX_CUDA_ARCHITECTURES를 항상 명시적으로 넘기므로 빌드 스크립트가 Hopper에 a 접미사를 붙입니다(위 참고). 텐서코어 qmm_sm80 커널은 SM 80 이상(Ampere, Ada, Hopper, Blackwell)을 담당합니다. 범용 qmm_naive/qmv 경로는 아키텍처 게이트가 없어 지원되는 모든 GPU에서 돕니다. 4비트 추론은 라인업 전반에서 동작하며 Hopper가 가장 빠른 전용 경로를 씁니다. CUDA 모델 테스트의 GB10(SM 121) 항목에서도 수십 개의 4비트 모델이 돌아갑니다. 비양자화 bf16/fp16 모델은 모든 GPU에서 동작하지만 메모리를 약 4배 더 쓰고 흔치 않은 일부 matmul shape은 최적화 커널이 없어 naive로 떨어지거나 실패합니다.


ROCm 빌드 옵션

HIP 아키텍처 선택

빌드는 빌드 호스트의 rocminfo가 보고하는 gfx 타겟에 맞춰 MLX 장치 코드를 컴파일합니다. GPU가 없는 호스트에서 빌드하거나 여러 타겟을 지정하려면 MLX_ROCM_ARCHITECTURES를 설정하십시오.

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

모든 항목은 HIP 타겟 ID여야 합니다. gfx 뒤에 소문자 16진수(gfx1151), 제네릭 타겟(gfx11-generic), 또는 기능 접미사가 붙은 타겟(gfx90a:xnack+) 형식입니다. native, 대문자 이름, 옵션 문자열 같은 그 밖의 값은 CMake configure 단계에서 Not a HIP target ID와 함께 중단됩니다. 각 항목이 --offload-arch 컴파일러 인자가 되기 때문입니다.

변수와 rocminfo 어느 쪽에서도 타겟을 얻지 못하면 빌드는 이 변수를 밝히며 실패합니다. 선택한 목록은 바이너리에 기록됩니다. mlxcel generate는 시작 시 HIP architecture <gfx target>; compiled for [<list>] 줄을 출력하고, MLXCEL_TRACE_ARCH=1은 어느 진입점에서든 같은 정보를 출력합니다. 장치의 타겟이 컴파일된 목록에 없으면 런타임 초기화에서 둘을 모두 밝히며 거부합니다. ROCm이 /opt/rocm 아래에 없으면 ROCM_PATH를 설정하십시오. 빌드는 그 경로에서 CMake 패키지, hipcc, rocminfo를 읽습니다.

ROCm에서 실행

바이너리에 $ROCM_PATH/lib가 rpath로 들어가므로 LD_LIBRARY_PATH는 필요하지 않습니다. 서버는 시작 시 Runtime device: GPU를 기록하며, GPU 프로세스는 rocm-smi --showpids에 나타납니다. UMA 호스트에서는 GPU가 운영체제와 다른 GPU 프로세스와 메모리를 공유하므로, 큰 모델을 로드하기 전에 rocm-smi --showpids로 다른 사용자를 확인하십시오.

expert-batched gather_qmm prefill 커널 대신 행 단위 MoE 커널을 쓰려면(예: 처리량 변화를 A/B 비교하거나 MoE 회귀를 bisect할 때) MLX_ROCM_GATHER_QMV_EXPERT_BATCHED=0을 설정하십시오. ROCm 백엔드가 실행 중에 읽는 변수는 환경 변수 레퍼런스의 ROCm 백엔드 변수 절에 정리되어 있으며, 안정된 인터페이스가 아니라 튜닝과 bisect용 스위치입니다.

메모리 사용량

두 가지 기본값이 할당기 메모리 사용량을 가중치 크기 가까이로 유지합니다(#2062). gfx1151에서 pp512/tg128로 측정하면 Meta-Llama-3.1-8B-Instruct-4bit의 MLX 최대 사용량은 이 기본값을 쓸 때 6.14 GB, 쓰지 않을 때 20.60 GB이고, decode 처리량은 같으며 prefill 차이는 2% 이내입니다.

  • In-flight 상한, MLX_ROCM_MAX_INFLIGHT_MB(기본값 1024). prefill은 연산마다 임시 버퍼를 할당하고, 각 버퍼는 해당 command batch가 GPU에서 끝나야 해제됩니다. 백엔드는 이 예산의 4분의 1을 할당한 batch를 commit하고, 진행 중인 양이 예산을 넘으면 가장 오래된 batch가 끝날 때까지 기다리므로 임시 버퍼가 대략 이 MiB 수 안에 머뭅니다. 0이면 상한을 두지 않습니다.
  • 버퍼 캐시 상한, MLXCEL_CACHE_LIMIT(ROCm 기본값 2 GiB). 해제된 버퍼는 재사용을 위해 캐시되는데, ROCm 할당기는 요청 크기와 정확히 같은 버퍼만 재사용합니다. ROCm 빌드는 MLXCEL_CACHE_LIMIT을 따로 지정하지 않으면 이 캐시를 2 GiB로 제한하며, 0이나 none을 주면 제한을 없앱니다.

로드 전 메모리 추정은 가중치, KV 캐시, headroom, activation 예약분에 더해 in-flight 예산을 mlxcel inspect의 별도 Backend in-flight 줄(--json 출력에서는 backend_inflight_bytes)로 계산합니다(#2155). 메모리 추정을 참조하십시오.

현재 상태

영역 ROCm 상태
affine 4비트 / 8비트 체크포인트 dense와 MoE 모두 네이티브로 실행합니다.
mxfp8, mxfp4 체크포인트 gather_qmm을 통한 MoE expert를 포함해 네이티브로 실행하며, 로드 로그에 모드마다 한 번 그 사실을 남깁니다(Quantization mode mxfp4: running on native ROCm kernels, no load-time conversion ...). 벤더 FP8 블록 체크포인트(quant_method: fp8, 128x128 블록)는 로드 시 mxfp8로 재양자화한 뒤 네이티브로 실행합니다.
NVFP4 체크포인트 네이티브 커널이 없습니다. ModelOpt NVFP4 체크포인트는 환경 변수 없이 로드 시 affine 4비트로 변환하고, 로드 로그에 경로와 이유를 남깁니다. 변환할 수 없는 레이어가 있으면 그 이름을 밝히며 로드가 실패합니다. ROCm에서는 MLXCEL_NVFP4_NATIVE_REPACK을 무시합니다. MLX 네이티브 NVFP4 export(config.json에 "mode": "nvfp4"가 있는 것, 예: mlx-community/*-nvfp4)는 로드 시 변환 경로가 없어 로드 단계에서 거부되므로 affine export를 사용하십시오.
affine MoE 모델(예: Qwen3-30B-A3B) 네이티브로 실행합니다. 단일 토큰 decode는 Metal, CUDA와 같은 fused 2-커널 MoE 경로(gate-up 다음 down)를 HIP 포트로 실행하며, MLXCEL_FUSED_MOE=0을 주면 gather_qmm 그래프 경로를 씁니다. MLXCEL_FUSED_MOE_MAX_DFF(CUDA와 같은 기본값 8192)보다 넓은 expert는 gather_qmm에 남습니다.
모델 전용 fused 커널(BitNet bitlinear_matmul, Apertus xIELU, Cohere2 add3 + LayerNorm, Mamba1 selective scan, Nemotron-H의 opt-in MLXCEL_FUSED_MOE_RELU2) HIP 포트로 실행합니다.
fused residual-add RMSNorm, RoPE + KV append(MLXCEL_FUSED_ADD_RMSNORM, MLXCEL_FUSED_ROPE_APPEND) ROCm에서는 기본으로 켜져 있고 Metal과 CUDA에서는 꺼져 있습니다. 두 HIP 포트는 대체하는 그래프와 바이트 단위로 같은 출력을 냅니다. 그래프를 쓰려면 각 변수를 0으로 설정하십시오.
paged attention(배치 paged decode, MLA split-KV, sparse paged decode) v1 decode, v2 partial, merge 커널이 HIP 포트로 실행되므로 이 경로들은 Metal, CUDA와 마찬가지로 gather 후 SDPA 대신 fused 커널을 씁니다. MLXCEL_PAGED_ATTENTION_NATIVE=0이면 gather 경로를 씁니다.
샘플링 decode(temperature > 0) fused Gumbel-max 샘플러와 rejection 샘플러가 HIP 커널로 실행됩니다. top-k와 top-p를 함께 쓰면 ROCm에서는 어휘 크기 152064까지(Metal과 CUDA에서는 32768까지) rejection 커널을 씁니다. MLXCEL_SAMPLING_GUMBEL=0, MLXCEL_SAMPLING_REJECTION=0이면 그래프를 씁니다.
하이브리드 SSM decode(granite-4.0-h, falcon-h1, plamo-2, Nemotron-H) 단일 토큰 Mamba2 SSM 업데이트가 SSD 그래프 대신 fused HIP 커널로 실행됩니다. MLXCEL_SSM_KERNEL=0이면 그래프를 씁니다.
FFT, Hadamard, SearchSorted 지원합니다. FFT는 hipFFT에서 실행하며 MLX_ROCM_FFT_CACHE_SIZE(CUDA와 같은 기본값 128)가 유지할 변환 plan 수를 제한합니다. 양의 정수가 아닌 값은 경고와 함께 무시합니다.
mlxcel-server chat completions dense와 affine MoE 체크포인트에서 스트리밍과 비스트리밍 모두 동작합니다.
오디오(음성 인식, 음성 합성) 동작합니다.
UMA 호스트의 메모리 추정 정확합니다. 추정기는 호스트 RAM보다 먼저 ROCm 할당기의 상한(검증 호스트에서는 96 GiB carve-out 중 76.80 GiB)을 읽고, 메모리 사용량에서 설명한 Backend in-flight 예약분을 더합니다.
GPU 오류 오류로 보고합니다. HIP가 거부한 launch는 그 launch를 낸 평가만 실패시키고 장치는 계속 쓸 수 있습니다. 비동기 오류(범위 밖 접근)는 기다리던 평가를 1초 안팎에 실패시키는데, queue 오류 뒤에는 HIP 런타임이 이후 호출을 모두 거부하므로 프로세스를 다시 시작해야 합니다. 끝나지 않는 커널은 MLX_ROCM_GPU_WATCHDOG_SECS(기본값 0, 꺼짐)를 설정하지 않으면 감지하지 못하며, 설정하면 그 시간을 넘긴 호스트 대기를 실패시킵니다.
진단 정보 AMD 벤더, 장치 이름, gfx 타겟, 장치 메모리를 보고합니다. 컴파일된 gfx 목록이 장치를 포함하지 않는 바이너리는 시작을 거부합니다.
Wave64 장치(CDNA: MI200 gfx90a, MI300 gfx942) 실행해 보지 않았습니다. lane 간 reduce를 하는 fused HIP 커널은 32-lane wavefront에서만 실행해 봤기 때문에, mlxcel은 장치가 32-lane wavefront를 보고할 때만 이 커널을 선택합니다. 64-lane 장치에서는 해당 경로가 MLX 그래프 대체 경로를 쓰고, 프로세스마다 한 번 stderr에 그 사실을 남깁니다. BitLinear에는 그래프 대체 경로가 없어 wave64 장치에서는 BitNet 체크포인트를 로드 단계에서 거부합니다.
ROCm 빌드의 CPU 장치(MLXCEL_DEVICE=cpu) 실행되지만 느립니다. 서빙 모드가 아니라 정확성 비교 기준이자 gfx 빌드가 맞지 않을 때의 우회 수단입니다. 이 장치의 BLAS 작업은 OpenBLAS 스레드 하나로 실행합니다.
모델 카탈로그 mlxcel arch --json과 WebUI는 모든 계열의 ROCm 지원을 partial로 보고합니다.
Windows 0.8부터 지원하며 소스 빌드만 되고 단일 노드만 지원합니다. Windows (AMD ROCm)을 참조하십시오.
다중 GPU, 분산 추론 지원하지 않습니다.

참고용 수치(greedy decode, 짧은 프롬프트, 검증 호스트 기준)로 Qwen3-0.6B-4bit는 약 250 tok/s, gpt-oss-20b-MXFP4-Q4는 약 8 tok/s입니다.


실험적 지원

Windows 빌드는 소스 빌드가 되고 각각 검증 호스트 한 대에서 실행을 확인했으며, 둘 다 미리 빌드한 아카이브를 제공하지 않습니다. Windows CUDA는 실험 단계이고, Windows ROCm은 0.8부터 지원하며 아래 구성에서 검증했습니다. 검증 구성을 벗어나면 문제가 생길 수 있으니 발견한 문제는 이슈 트래커에 알려 주십시오.

플랫폼 Feature 검증 구성 범위
Windows (NVIDIA CUDA) cuda NVIDIA RTX Spark N1X(aarch64-pc-windows-msvc, sm_121), CUDA 13.4.3, cuDNN 9.27 단일 노드. 빌드할 때 쓴 CUDA 툴킷 디렉터리나 런타임 번들로 실행합니다.
Windows (AMD ROCm) rocm Radeon 8060S(x86_64-pc-windows-msvc, gfx1151), ROCm 10.1.0 단일 노드. 개발용 SDK로도, 다른 위치로 옮길 수 있는 런타임 번들로도 실행합니다.

NVIDIA GPU가 있는 Windows 머신에서 공식 지원하는 방법은 여전히 WSL2에서 Linux CUDA 빌드를 쓰는 것입니다.

Windows 빌드 공통 사항

두 Windows 빌드 모두 MSVC 툴체인으로 네이티브 mlxcel.exe와 mlxcel-server.exe를 만듭니다. nvcc와 ROCm의 clang이 MSVC ABI로 컴파일하므로 GNU 툴체인으로는 빌드되지 않습니다.

요구사항 상세
Visual Studio 2022 이상, "Desktop development with C++" 워크로드. 함께 설치되는 CMake와 Ninja로 충분합니다.
Rust 머신에 맞는 MSVC 타깃(x86_64-pc-windows-msvc, Windows on Arm의 CUDA 빌드는 aarch64-pc-windows-msvc). -gnu 타깃은 쓰지 않습니다. 버전은 저장소의 rust-toolchain.toml이 정합니다.
셸 타깃 아키텍처용 Visual Studio 개발자 프롬프트("x64 Native Tools" 또는 "ARM64 Native Tools Command Prompt"). 여기서 실행해야 cl.exe, CMake, Ninja가 PATH에 잡힙니다.

CMAKE_GENERATOR를 지정하지 않으면 MLX는 Ninja generator로 빌드합니다. Visual Studio generator는 NVIDIA MSBuild 통합 없이는 .cu 파일을 컴파일하지 못하고 HIP 컴파일 명령은 아예 실행하지 못합니다. target\ 아래 MLX 의존성 트리는 경로가 깊으므로 CMake, nvcc, hipcc가 경로를 열지 못한다고 하면 Win32 long path를 켜거나 C:\t처럼 짧은 CARGO_TARGET_DIR를 지정하십시오.

빌드는 MLX CMake가 CPU 백엔드용으로 내려받은 OpenBLAS DLL을 실행 파일 옆에 복사합니다. 바이너리를 옮길 때 이 DLL도 mlxcel.exe 옆에 두십시오.

두 Windows 빌드에 공통으로 적용되는 제한은 다음과 같습니다.

영역 상태
단일 GPU 생성과 mlxcel-server 검증 구성에서 동작합니다.
텐서·파이프라인 병렬, 분리형 서빙, --pp-auto, --pp-peer 시작 시 해당 플래그를 밝히는 오류로 거부합니다.
Unix domain socket 리슨(--host *.sock) 거부합니다. TCP로 리슨하십시오.
router 모드의 모델 다운로드와 캐시 삭제(WebUI 모델 라이브러리) 거부합니다. 디렉터리 fd 기준 파일시스템 연산에 의존하는데 Windows 구현이 없습니다.
CPU 디바이스(MLXCEL_DEVICE=cpu, 또는 GPU feature 없는 빌드) 빌드와 로드는 되지만 양자화 matmul이 MLX의 스칼라 CPU 루프로 돌아갑니다. CUDA 호스트에서 gemma-4-e4b-it-qat-4bit로 24토큰 generate를 실행했을 때 30분이 지나도 끝나지 않았습니다.
미리 빌드한 아카이브, CI 아카이브는 없습니다. CI는 GPU 없이 컴파일만 하며, 호스팅되는 Windows build (compile-only) 작업이 Windows 빌드 스크립트, build_support/, overlay의 런타임 디렉터리 소스, 번들 패키저, MLX 핀이 바뀔 때 실행되어 mlxcel-mlx-pin 단위 테스트와 tests/ Python 스위트, ROCm 10.1.0 gfx1151 TheRock SDK에 대한 cargo check -p mlxcel-core --features rocm을 확인합니다. Windows에서 cuda feature를 빌드하는 CI 작업은 없습니다.

Windows (NVIDIA CUDA)

검증 구성은 aarch64-pc-windows-msvc, NVIDIA RTX Spark N1X(Blackwell, sm_121, 드라이버 616.62), Visual Studio 2026(MSVC 14.51), CUDA 13.4.3, cuDNN 9.27입니다. 빌드는 x64와 arm64 툴킷 구조를 똑같이 처리하지만 x86_64-pc-windows-msvc GPU 호스트에서 실행한 적은 아직 없습니다.

Windows 빌드 공통 사항에 더해 다음이 필요합니다.

  • CUDA Toolkit 13.x와 CUDA 13용 cuDNN 9. CUDA 12.x는 Windows에서 DLL을 bin\<arch>\가 아니라 bin\에 두기 때문에 아직 처리하지 않습니다. x64에서는 CUDA 설치 프로그램을 써도 됩니다. cuDNN은 redistributable 아카이브 구조여야 하며 cuDNN 설치 프로그램이 만드는 lib\13.x\x64 구조는 찾지 않습니다.
  • Windows on Arm용으로 NVIDIA는 redistributable 아카이브만 배포합니다(developer.download.nvidia.com/compute/cuda/redist/와 .../cudnn/redist/의 windows-arm64 항목). cuda_cudart, cuda_nvcc, cuda_crt, cuda_nvrtc, cccl, libcublas, libcufft, libcusolver, libcusparse, libnvjitlink, libnvvm 아카이브와 cuDNN 아카이브를 한 디렉터리에 풀면 되고 관리자 권한은 필요 없습니다.
  • 툴킷을 지원하는 NVIDIA 디스플레이 드라이버.

빌드가 기대하는 툴킷 구조는 bin\nvcc.exe, include\, lib\<arch>\, bin\<arch>\이고 <arch>는 x64 또는 arm64입니다. cuDNN은 같은 디렉터리에 있어도 되고 같은 구조의 별도 루트에 있어도 됩니다.

개발자 프롬프트에서 빌드합니다.

set CUDA_PATH=C:\path\to\cuda
rem cuDNN을 CUDA_PATH에 풀지 않은 경우에만:
set CUDNN_PATH=C:\path\to\cudnn
cargo build --release --features cuda

CUDA_HOME이 있으면 CUDA_PATH보다 먼저 씁니다. GPU 아키텍처는 Linux와 마찬가지로 nvidia-smi로 감지하고 MLX_CUDA_ARCHITECTURES로 덮어쓸 수 있습니다(CUDA 빌드 옵션 참조).

CUDA와 cuDNN DLL은 빌드 시점에 기록한 툴킷의 bin\<arch> 디렉터리에서 delay-load하고 NVRTC는 기록해 둔 include\에서 CUDA 헤더를 찾습니다. 그래서 툴킷이 PATH에 없고 CUDA_PATH가 설정되지 않아도 첫 JIT 컴파일까지 포함해 실행됩니다. 디스플레이 드라이버의 nvcuda.dll은 드라이버에서 가져옵니다. 번들 없이 실행하면 그 툴킷 디렉터리가 있어야 합니다. 다른 머신에서 실행하려면 런타임 번들을 만드십시오.

Windows CUDA 런타임 번들

런타임 번들에는 CUDA, cuDNN, NVRTC DLL과 JIT 컴파일에 쓰는 헤더가 들어 있어서 툴킷 없이 mlxcel을 실행할 수 있습니다. 번들은 scripts/package_windows_cuda_runtime.py(Python 3, 표준 라이브러리만 사용)로 만듭니다.

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

바이너리는 런타임 디렉터리 하나를 MLXCEL_CUDA_RUNTIME_DIR(절대 경로), cuBLASLt DLL이 있는 실행 파일 디렉터리, 빌드 시점에 기록한 툴킷 순서로 고르고, 고른 결과를 [mlx-cuda] CUDA runtime: C:\mlxcel\bin (executable directory)처럼 출력합니다. 번들에서 필수 DLL이 없으면 기대한 전체 경로를 알리고 멈춥니다. cuFFT와 cuSOLVER는 이 검사 대상이 아니며 처음 호출할 때 로드합니다.

  • RTX Spark N1X(드라이버 617.74, CUDA 13.4, cuDNN 9.27)에서 툴킷 이름을 바꾼 채 시험했습니다. 콜드 JIT 캐시의 mlxcel generate와 mlxcel-server 채팅 요청이 번들 폴더와 MLXCEL_CUDA_RUNTIME_DIR 양쪽에서 동작했습니다. x64 GPU 호스트와 디스플레이 드라이버만 있는 머신에서는 시험하지 않았습니다.
  • 번들에는 Visual C++ 런타임이 없습니다. openblas.dll은 일반 import이므로 번들 밖에서 실행하는 바이너리 옆에 복사본을 두십시오.
  • 툴킷 아카이브를 한 디렉터리에 풀면 LICENSE 파일이 서로 덮어씁니다. cuda_cudart와 cudnn 아카이브의 LICENSE는 --cuda-license와 --cudnn-license로, NVIDIA/cccl 릴리스 태그의 LICENSE는 --cccl-license로 넘기십시오.
  • 번들은 자기 머신에서 쓰는 용도입니다. CUDA EULA가 JIT에 쓰는 헤더를 모두 재배포 가능으로 나열하지는 않으므로, 공개 배포 전에는 NVIDIA에 조건을 확인하십시오.

Windows (AMD ROCm)

MSVC로 빌드한 rocm feature 네이티브 mlxcel.exe와 mlxcel-server.exe이며, 0.8부터 지원하고 소스 빌드만 되며 단일 노드입니다. Linux 빌드와 같은 ROCm overlay를 쓰고, 위의 Windows 빌드가 공유하는 단일 노드 제약을 그대로 받습니다.

검증 구성은 x86_64-pc-windows-msvc, 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), Visual Studio 18(MSVC 14.51)입니다. 이 호스트에서 mlx-community/Qwen3-0.6B-4bit는 약 115 tok/s, mlx-community/Qwen3-1.7B-bf16은 약 30 tok/s로 디코딩합니다. 빌드는 x86_64-pc-windows-msvc만 받습니다.

Windows 빌드 공통 사항에 더해 다음이 필요합니다.

  • Windows용 ROCm 10.x SDK. ROCm 10에는 HIP SDK 설치 프로그램이 없으므로 https://stable.repo.amd.com/rocm/core/tarball/에서 GPU 계열에 맞는 TheRock tarball(gfx1151이면 therock-dist-windows-gfx1151-10.1.0.tar.gz)을 받거나 AMD ROCm pip wheel의 SDK를 씁니다. 빌드에는 bin\hipcc.exe, lib\의 import library, 타깃 감지용 amdgpu-arch.exe(10.1 tarball은 lib\llvm\bin\에 둠)가 필요합니다.
  • GPU를 지원하는 AMD 디스플레이 드라이버.

1단계: SDK 설치

Windows에 기본으로 들어 있는 tar로 tarball을 풀고 C:\rocm\10.1.0\bin\hipcc.exe가 있는지 확인합니다. 아카이브에 최상위 폴더가 있으면 그 내용을 한 단계 위로 옮기십시오.

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

SDK 디렉터리는 관리자만(또는 본인만) 쓸 수 있어야 합니다. 바이너리가 그 bin\에서 DLL을 로드하고 include\의 헤더로 GPU 커널을 컴파일하므로 거기에 쓸 수 있는 계정은 모든 mlxcel 프로세스 안에서 코드를 실행할 수 있습니다. C:\ 바로 아래 만든 폴더는 드라이브 루트로부터 인증된 모든 사용자의 수정 권한을 상속하므로 계정이 둘 이상인 머신에서는 관리자 프롬프트에서 그 권한을 지우십시오.

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

Administrators와 SYSTEM은 모든 권한, Users는 읽기와 실행 권한만 남습니다. 적용 후 icacls C:\rocm\10.1.0\bin 출력에 Authenticated Users 항목이 없어야 합니다.

2단계: 빌드

"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

ROCM_PATH가 없으면 HIP_PATH를 읽습니다. 경로는 절대 경로여야 하며 상대 경로이거나 SDK bin에 amdhip64_*.dll이 없으면 빌드가 실패합니다. HIP 타깃은 amdgpu-arch로 감지합니다(Windows SDK에는 rocminfo가 없음). 감지를 건너뛰거나 설치되지 않은 GPU용으로 빌드하려면 MLX_ROCM_ARCHITECTURES=gfx1151처럼 지정하고 여러 개면 세미콜론으로 구분합니다. 바이너리는 빌드한 타깃의 GPU에서만 실행됩니다.

검증 호스트에서 전체 릴리스 빌드는 약 18분 걸렸습니다. 이 호스트는 GPU carve-out 뒤 Windows가 보는 RAM이 약 32 GB인데, 메모리가 부족한 상태에서 빌드 한 번이 링크 단계에서 실패했습니다(link.exe 종료 코드 0xC0000142). 병렬도를 낮추니 완료되었습니다.

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

3단계: 빌드 확인

ROCM_PATH는 빌드할 때만 필요합니다. 새 프롬프트에서 바이너리를 실행해 봅니다.

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

stderr 첫 줄들에 런타임 디렉터리와 GPU가 나옵니다. 예를 들면 [mlx-rocm] ROCm runtime: C:\rocm\10.1.0\bin (build-time ROCm SDK)와 [mlx-rocm] bound HIP device 0: gfx1151 (...)입니다. 첫 실행에서는 hiprtc로 fused 커널을 컴파일해 ~/.cache/mlxcel/hsaco 또는 MLXCEL_CACHE_DIR 아래에 캐시합니다(MLX_HSACO_CACHE_DIR로 위치를 바꿀 수 있으며, 다른 사용자가 쓸 수 있는 디렉터리는 거부합니다).

이렇게 만든 개발 빌드는 빌드 시점에 기록한 SDK bin에서 ROCm DLL을 로드하므로 SDK가 같은 경로에 있는 머신에서만 실행됩니다. 다른 머신에서 실행하려면 런타임 번들을 만드십시오.

Windows ROCm 런타임 번들

런타임 번들에는 mlxcel에 필요한 ROCm DLL, 커널 데이터, 헤더가 들어 있어서 실행할 머신에는 AMD 디스플레이 드라이버만 있으면 됩니다. 번들은 scripts/package_windows_rocm_runtime.py(Python 3, 표준 라이브러리만 사용)가 SDK로부터 만듭니다. --arch는 바이너리를 빌드한 GPU 타깃마다 한 번씩 넘깁니다.

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

dist\mlxcel-rocm-runtime-10.1.0-win-x64.zip이 만들어지며 ROCm 10.1.0의 gfx1151 기준 압축 76 MiB, 해제 225 MiB입니다. 바이너리까지 들어간 바로 실행 가능한 폴더가 필요하면 zip 대신 폴더로 만듭니다.

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

스크립트는 담을 파일이 SDK에 하나라도 없으면 실패하므로 이후 ROCm 릴리스에서 DLL 이름이 바뀌면 패키징 단계에서 드러납니다. rocBLAS, hipBLASLt, rocFFT가 자기 DLL 기준 ..\.kpack\에서 디바이스 코드를 읽기 때문에 번들은 SDK의 폴더 구조를 그대로 유지합니다.

폴더 내용
bin\ mlxcel이 로드하는 DLL 10개(amdhip64_7, amd_comgr, rocm_kpack, hiprtc0716, hiprtc-builtins0716, rocblas, libhipblaslt, origami, hipfft, rocfft)와 요청한 타깃으로 줄인 rocBLAS, hipBLASLt 커널 데이터. 이 폴더가 런타임 디렉터리입니다.
.kpack\ 디바이스 코드 아카이브. <DLL 폴더>\..\.kpack\<라이브러리>_<타깃>.kpack으로 엽니다.
include\hip\ hiprtc 커널 컴파일러가 쓰는 헤더.
licenses\ 번들에 담은 구성요소의 SDK 라이선스 파일. ROCm 10.1.0 tarball에는 HIP 런타임, hipRTC, rocm_kpack의 라이선스 파일이 없습니다.

실행할 머신에서는 번들을 C:\mlxcel처럼 짧은 경로에 풀고 두 가지 배치 중 하나를 고릅니다.

  • mlxcel.exe, mlxcel-server.exe, libopenblas.dll을 풀어 놓은 bin\에 넣고(위의 --add-binary 폴더에는 이미 들어 있음) C:\mlxcel\bin\mlxcel.exe로 실행합니다.
  • 바이너리는 다른 곳에 두고 MLXCEL_ROCM_RUNTIME_DIR=C:\mlxcel\bin을 설정합니다.

바이너리는 프로세스마다 런타임 디렉터리 하나를 골라 [mlx-rocm] ROCm runtime: C:\mlxcel\bin (executable directory)처럼 한 번 출력합니다. 후보를 확인하는 순서는 다음과 같습니다.

  1. MLXCEL_ROCM_RUNTIME_DIR. 절대 경로여야 하며 설정되어 있으면 이것만 후보입니다. 상대 경로이거나 HIP 런타임이 없는 폴더면 문제를 밝히는 메시지와 함께 멈춥니다.
  2. mlxcel.exe 또는 mlxcel-server.exe가 있는 폴더. 바이너리가 import하는 HIP 런타임(ROCm 10은 amdhip64_7.dll)이 있을 때만 후보가 됩니다.
  3. 빌드 시점에 기록한 SDK bin. HIP 런타임이 있을 때만 후보가 됩니다.

ROCm DLL은 모두 고른 폴더에서 가져옵니다. 거기에 없는 DLL은 기대한 전체 경로를 밝히며 멈추고 다른 후보나 작업 디렉터리, PATH, AMD 디스플레이 드라이버가 System32에 설치하는 이전 버전 amdhip64_7.dll과 amd_comgr.dll에서 가져오지 않습니다. 서로 다른 ROCm 릴리스의 DLL은 함께 동작한다는 보장이 없기 때문입니다(SDK의 HIP 런타임과 드라이버의 amd_comgr.dll을 같이 쓰면 디바이스를 하나도 찾지 못함). HIP 런타임이 이름으로 여는 amd_comgr*.dll이나 hiprtc-builtins*.dll이 폴더에 없을 때도 ROCm DLL을 로드하기 전에 멈춥니다.

설치 규칙은 다음과 같습니다.

  • 애플리케이션마다 런타임 폴더를 따로 두십시오. 예를 들어 llama.cpp의 Windows ROCm 릴리스는 ROCm 10.0의 amdhip64_7.dll, amd_comgr.dll, rocm_kpack.dll을 실행 파일 옆에 둡니다. 두 프로그램을 한 폴더에 풀면 이 파일들이 덮어쓰여 두 릴리스가 섞입니다. 나란히 둬야 한다면 MLXCEL_ROCM_RUNTIME_DIR이 mlxcel 쪽 폴더를 가리키게 하십시오.
  • 런타임 폴더, 그 상위 폴더(.kpack\과 include\가 있는 곳), mlxcel.exe가 있는 폴더는 1단계에서 설명한 이유로 관리자만(또는 본인만) 쓸 수 있어야 합니다. C:\Program Files는 이미 그렇게 보호되고 C:\ 바로 아래 만든 폴더라면 위의 icacls 명령을 적용하십시오.
  • 런타임 폴더 경로는 130자 이하로 두십시오. rocBLAS와 hipBLASLt는 그 아래 최대 128자 깊이의 커널 파일을 여는데 Windows 경로 길이 제한(260자) 때문에 133자 위치에서는 첫 행렬 곱에서 access violation으로 종료되었습니다. 고른 폴더가 이보다 길면 바이너리가 경고를 출력합니다.
  • ROCM_HOME, ROCM_PATH, HIP_PATH가 설정되어 있으면 JIT 헤더로 번들의 include\보다 먼저 씁니다. AMD HIP SDK 설치 프로그램은 HIP_PATH를 시스템 전체에 설정하므로 그런 머신에서는 mlxcel 실행 시 set HIP_PATH=로 지우거나 번들 루트를 가리키게 해서 다른 ROCm 버전의 헤더로 커널을 컴파일하지 않게 하십시오.
  • 번들은 Radeon 8060S(gfx1151)의 Adrenalin 32.0.31041.1004에서 시험했습니다. 다른 드라이버 버전과 GPU는 시험하지 않았습니다.

Windows ROCm 빌드는 host-shadow 메모리 모델을 씁니다. Windows는 APU에서도 GPU 할당을 CPU 주소 공간에 매핑하지 않으므로 CPU가 GPU 버퍼를 읽거나 쓸 때는 pinned host 사본을 거치고 다음 커널이 그 버퍼를 읽기 전에 되돌려 씁니다. 통합 GPU의 allocator 한도는 HIP가 보고하는 디바이스 메모리의 80%이며(검증 호스트에서 107.87 GiB 중 86.29 GiB), Linux와 달리 시스템 RAM으로 늘어나지 않습니다. FFT(hipFFT, rocFFT)는 번들에 들어 있지만 Windows에서 실행해 보지 않았습니다.


빌드 결과물

빌드가 성공하면 target/release/ 아래에 두 개의 바이너리가 생성됩니다(Windows에서는 target\release\, 또는 설정한 CARGO_TARGET_DIR 아래에 mlxcel.exe, mlxcel-server.exe가 생기고 그 옆에 OpenBLAS DLL이 놓입니다).

바이너리 경로 설명
mlxcel target/release/mlxcel run, generate, embed, rerank, serve, list, arch, inspect, download, detect, rm, tune, split-mtp(마지막은 기본 surgery feature 필요)의 13개 명령을 모두 제공하는 메인 CLI
mlxcel-server target/release/mlxcel-server llama-server 드롭인 대체용으로 설계된 독립 실행형 HTTP 서버. 서버 전용 배포에서도 모델을 받을 수 있도록 download 서브명령어도 제공합니다.

시스템 전체에서 사용할 수 있도록 PATH에 포함된 디렉터리(예: /usr/local/bin)에 바이너리를 복사할 수 있습니다.

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

문제 해결

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

이 오류는 Metal 셰이더 컴파일러가 설치되어 있지 않음을 의미합니다. Metal Toolchain은 Xcode와 별개로 설치해야 하는 선택적 구성요소입니다. 다음 명령으로 설치하십시오.

xcodebuild -downloadComponent MetalToolchain

설치 후 cargo build를 다시 실행합니다.

Cannot find CUDA library directory (Linux)

빌드 스크립트가 기본 경로(/usr/local/cuda)에서 CUDA를 찾지 못하는 경우입니다. CUDA_HOME을 실제 CUDA 설치 디렉터리로 설정하십시오.

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

영구적으로 적용하려면 쉘 프로파일(~/.bashrc, ~/.zshrc 등)에 export 줄을 추가하십시오.

빌드 중 nvidia-smi: command not found (Linux)

빌드 스크립트는 nvidia-smi를 사용하여 GPU의 컴퓨트 캐퍼빌리티를 자동으로 감지합니다. 헤드리스 빌드 서버, 컨테이너화된 CI 환경, 또는 GPU 드라이버가 설치되지 않은 머신에서는 nvidia-smi를 사용할 수 없을 수 있습니다.

자동 감지를 우회하려면 타겟 아키텍처를 수동으로 지정하십시오.

# GPU의 SM 버전을 지정하세요. SM 90 이상은 `a`를 붙입니다(예: 90a).
MLX_CUDA_ARCHITECTURES=90a cargo build --release --features cuda

error[E0463]: can't find crate for 'std' 또는 링커 오류 (Linux)

rust-toolchain.toml에 고정된 툴체인(v0.8.0에서는 Rust 1.97.1)을 설치하십시오. rustup이 있으면 checkout에서 Cargo를 실행할 때 이 버전이 자동 선택됩니다.

rustup toolchain install 1.97.1 --component rustfmt --component clippy
rustc --version  # 이 checkout에서는 1.97.1이어야 함

빌드 시간이 너무 오래 걸리는 경우 (Linux)

첫 빌드에서는 이 현상이 정상입니다. MLX C++ 라이브러리, 해당 CUDA 커널, 지원 라이브러리(fmt, gguflib)가 모두 소스에서 컴파일됩니다. 이후 증분 빌드는 훨씬 빠릅니다. 빠른 NVMe 드라이브를 장착한 최신 워크스테이션에서 전체 CUDA 첫 빌드는 일반적으로 10~20분 정도 소요됩니다.

CMake, nvcc, hipcc가 경로를 열지 못하는 경우 (Windows)

target\ 아래 MLX 의존성 트리는 Windows 경로 길이 제한(260자)을 넘을 만큼 깊습니다. Win32 long path를 켜거나 짧은 디렉터리에 빌드하십시오.

set CARGO_TARGET_DIR=C:\t\mlxcel

mlxcel: cannot find the ROCm runtime (Windows)

바이너리 옆에 amdhip64_7.dll이 없고 빌드할 때 쓴 SDK도 그 경로에 없는 경우입니다. 바이너리를 풀어 놓은 런타임 번들의 bin\ 폴더에 넣거나 MLXCEL_ROCM_RUNTIME_DIR을 그 폴더로 설정하십시오. MLXCEL_ROCM_RUNTIME_DIR이 절대 경로여야 한다거나 amdhip64_7.dll이 없다는 메시지는 변수가 상대 경로나 잘못된 폴더로 설정되었다는 뜻이며 이때 다른 후보로 넘어가지 않습니다.

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

고른 런타임 폴더에 파일이 빠진 경우로, 손으로 파일을 복사했거나 일부를 지웠을 때 생깁니다. 번들을 다시 풀거나, SDK에 파일이 없으면 실패하는 scripts\package_windows_rocm_runtime.py로 번들을 다시 만드십시오. 빠진 파일을 다른 ROCm 설치나 System32의 DLL로 채우지 마십시오.