MLXCEL 환경 변수¶
이 문서는 mlxcel에서 사용하는 MLXCEL_* shell 환경 변수를 개발자 관점에서 정리합니다. 안정적인 운영 설정, 서비스 기본값, 빌드 제어, 저수준 진단 스위치를 함께 다룹니다.
일반 배포에서는 같은 기능의 CLI 플래그가 있으면 플래그를 우선 사용하십시오. 환경 변수는 프로세스 전체에 적용되며 일부 저수준 스위치는 최초 사용 시 한 번 읽힌 뒤 캐시됩니다. mlxcel 또는 mlxcel-server를 시작하기 전에 설정해야 합니다.
규칙¶
- 같은 옵션을 제어하는 CLI 플래그와 환경 변수가 함께 있으면, 별도 설명이 없는 한 CLI 플래그가 우선합니다.
- llama.cpp 호환 설정 일부는
LLAMA_ARG_*별칭도 사용합니다. 이 문서는 mlxcel 전용MLXCEL_*이름에 초점을 둡니다. - boolean 파싱은 모든 내부 knob에서 완전히 같지는 않습니다.
- 사용자에게 노출된 서버 옵션은 대체로
true/false,1/0,yes/no,on/off를 허용합니다. - 많은 진단 스위치는 presence-based라서 값이 설정되어 있으면 활성화됩니다.
- "falsy disables"라고 적힌 항목은
0,false,off,no를 비활성화 값으로 취급합니다. - advanced, diagnostic, test-only로 표시한 변수는 안정적인 공개 API가 아닙니다. 벤치마크, rollback, kernel 개발, CI 용도입니다.
런타임과 크기 산정¶
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLXCEL_DEVICE | gpu, metal, cpu | GPU 선호 | 런타임 디바이스 override. cpu는 CPU 실행을 요청합니다. 잘못된 값은 경고 후 GPU 선호로 폴백하며, GPU 백엔드가 없으면 런타임은 CPU로 폴백합니다. 시작 로그는 둘 중 무엇이 런타임을 CPU에 올렸는지 밝힙니다. cpu 요청은 GPU 백엔드 사용 가능 여부와 나란히 요청으로 보고되므로, GPU 호스트의 CPU 로그를 백엔드 부재로 오해하지 않습니다. |
MLXCEL_WIRED_LIMIT | max, 0, none, bytes, NGB, NMB | Apple GPU에서 max | Apple Silicon GPU wired-memory 한도. 0 / none은 비활성화하고 명시 크기는 wired memory를 제한합니다. |
MLXCEL_WIRE_WEIGHTS | 1/true/on/yes | 미설정(off) | macOS 전용입니다. 로드가 끝난 뒤 평가된 weight buffer를 1 GiB 단위로 mlock하고, 시스템 wired 총량이 RAM에서 max(7 GiB, RAM의 15%)를 뺀 값에 닿기 전에 멈춥니다. 기본값인 eager MLXCEL_WEIGHT_MATERIALIZATION이 필요합니다. wired 총량은 /metrics의 mlxcel_weights_wired_bytes로 보고하고 실패는 모두 로그에 남깁니다. soak 측정 전이라 기본값은 off이며 macOS 메모리 가드레일을 참고하십시오. |
MLXCEL_MEMORY_LIMIT | 0, none, bytes, NGB, NMB | 미설정 | MLX allocator soft cap. 0이 아닌 값은 mlxcel inspect와 --estimate-memory의 권위 있는 사용 가능 메모리 예산으로도 사용됩니다. |
MLXCEL_HEADROOM_FACTOR | 양의 f64 | 1.20 | 통합 메모리 추정기(mlxcel inspect, --estimate-memory, paged KV auto 예산, 프롬프트 캐시 용량 기본값)에서 weights + KV에 곱하는 allocator overhead 배수입니다. activation 예약분(MLXCEL_ACTIVATION_MULT)과 ROCm 빌드의 backend in-flight 예약분(MLX_ROCM_MAX_INFLIGHT_MB)은 별도 항이어서 이 배수가 적용되지 않습니다. 0 < value <= 1.0 값은 headroom 항을 비활성화하고, 잘못됐거나 0 이하인 값은 경고 후 기본값으로 폴백합니다. |
MLXCEL_ACTIVATION_MULT | 양의 유한 f64 | 2.0 | 통합 메모리 추정기의 activation working set(batch x min(context, prefill chunk) x (hidden + intermediate) x 2 bytes)에 곱하는 배수이며, 사용처는 MLXCEL_HEADROOM_FACTOR와 같습니다. 이 항에는 MLXCEL_HEADROOM_FACTOR가 적용되지 않습니다. 보정용 override이며 잘못됐거나 유한하지 않거나 0 이하인 값은 경고 후 2.0을 사용합니다. |
MLXCEL_CACHE_DIR | 디렉터리 경로 | $HOME/.cache/mlxcel | 기본 model store($MLXCEL_CACHE_DIR/models/<owner>/<name>)와 토크나이저 language-analysis 파일을 포함한 기타 mlxcel 디스크 캐시의 루트입니다. |
MLXCEL_MODELS_DIR | 디렉터리 경로 | 미설정 | model-store 루트 직접 재정의. repo-id 다운로드와 조회는 추가 models/ 컴포넌트 없이 $MLXCEL_MODELS_DIR/<owner>/<name>을 사용합니다. --models-dir가 우선합니다. |
MLXCEL_DEFAULT_ORG | HuggingFace owner | mlx-community | 모델 인자에 slash가 없을 때 사용할 기본 owner입니다. 예: Qwen3-4B-4bit -> mlx-community/Qwen3-4B-4bit. |
MLXCEL_SERVER_DECODE_STORAGE | auto, dense, paged | auto | 연속 배칭 decode storage backend. --decode-storage-backend 플래그가 우선합니다. |
MLXCEL_SURGERY | YAML 파일 경로 | 미설정 | surgery 기능이 빌드되어 있고 --surgery가 없을 때 사용할 런타임 모델 변환 파일입니다. |
MLXCEL_CACHE_LIMIT | 0, none, bytes, NGB, NMB | ROCm 빌드에서 2GB, 그 외 미설정 | MLX 자체 buffer cache의 상한(set_cache_limit)입니다. 미설정이면 backend 기본값을 따르는데, Metal과 CUDA에는 상한이 없고 ROCm에서는 2 GiB입니다. ROCm allocator의 자체 기본값으로는 cache가 allocator 메모리 한도까지 커지기 때문입니다. 0이나 none은 상한을 없애고 MLX 기본 cache 동작을 그대로 둡니다. CUDA에서는 decode loop의 주기적 cache clear가 기본적으로 꺼져 있으므로 메모리 증가를 제어하는 주된 수단이고, 아래 MLXCEL_CACHE_CLEAR_INTERVAL을 함께 참조하십시오. |
MLXCEL_CACHE_CLEAR_INTERVAL | 0(비활성화), 양의 정수(토큰 주기) | CUDA에서 0, Metal/CPU에서 256 | decode loop와 batch scheduler가 주기적으로 호출하는 clear_memory_cache의 주기(생성 토큰 수 기준)입니다. 기본값은 backend에 따라 다릅니다. CUDA에서는 주기적 clear가 memory pool을 흔들어 CUDA-graph 재사용을 깨뜨리므로 기본적으로 꺼져 있고(대신 MLXCEL_CACHE_LIMIT로 상한을 두십시오), Metal/CPU에서는 Python mlx-lm과 동일하게 256토큰 주기로 저렴하게 정리합니다. |
MLXCEL_PREFILL_CHUNK | 부호 없는 정수 토큰 수, 0이면 비활성화 | 2048 | CLI, 서버, 벤치마크 모든 경로의 prefill 청크 크기입니다. 프롬프트를 이 토큰 수만큼씩 나눠 모델에 넣으므로, sliding-window cache가 청크 사이에서 잘리고 prefill 최고 메모리는 프롬프트 전체가 아닌 청크 크기를 따라갑니다. 0이면 프롬프트 전체를 한 번에 prefill합니다. 서버에서 --prefill-chunk-size를 주지 않은 상태로 이 변수를 설정하면, 적응형 기본값(다른 시퀀스가 디코딩 중이 아니면 2048, 디코드 배치가 돌고 있으면 512) 대신 두 상태 모두 이 값으로 고정되며 0은 서버 청크 분할도 끕니다. 파싱할 수 없는 값은 경고를 남기고 기본값을 유지합니다. |
MLXCEL_MAX_BATCH_PREFILL_TOKENS | 부호 없는 정수 토큰 수, 0이면 비활성화 | 유도값 (2 * max_batch_prefill * min(prefill_chunk_size, 512), 기본 설정에서 4096) | 서버 배치 prefill 한 창의 패딩 포함 토큰 예산으로, prefill 배치가 점유할 수 있는 순간 메모리를 제한합니다. 512 상한은 2048 토큰 청크 기본값 때문에 예산이 네 배로 커지지 않게 합니다. 0이면 상한을 두지 않습니다. --max-batch-prefill-tokens가 우선합니다. |
MLXCEL_PREFILL_GRANT_INTERVAL | 부호 없는 정수 디코드 틱 수, 0이면 비활성화 | 16 | 연속 배칭 스케줄러의 prefill 공정성 grant. --prefill-chunk-size보다 긴 프롬프트는 청크 사이에 대기하는데, N 디코드 틱마다 청크를 하나 배정해 긴 프롬프트 하나 때문에 배치 전체가 멈추지 않게 합니다. 0이면 grant를 주지 않습니다. |
MLXCEL_KV_ADMISSION_WATERMARK | 0.0부터 0.5까지의 비율 | 0.01 | 행이 디코딩 중일 때 prefill admission이 paged KV block 예산에서 비워 두는 비율로, 디코딩 중인 행이 preemption 없이 늘어날 공간을 확보합니다. 그 공간을 만들기 위해 admission은 대기하기 전에 오래 쓰지 않은 프롬프트 캐시 접두사를 축출하고 우선순위가 엄격히 낮은 행을 preempt합니다. 디코드 배치가 돌고 있고 block 예산이 있을 때만 적용되며 0이면 비활성화됩니다. 범위를 벗어나거나 숫자가 아닌 값은 시작 시 거부됩니다. --kv-admission-watermark가 우선합니다. 연속 배칭을 참조하십시오. |
MLXCEL_REQUEST_MEMORY_BUDGET | 양의 bytes, NK/NKB, NM/NMB, NG/NGB | 런타임 상한(device 권장 working set, wired 상한, MLX allocator 상한) 중 가장 작은 양수 값 | Metal의 텍스트 전용 Nemotron-H에 적용하는 요청 admission 상한입니다. 이 경우 스케줄러는 활성 slot 하나만 운영하고 요청마다 보수적인 메모리 추정치로 admission을 판단합니다. 이 변수는 런타임 상한을 낮출 수만 있습니다. 명시한 값이 0, none, 빈 값, 유한하지 않은 값, overflow 값이면 서버가 시작하지 않습니다. 예산을 넘는 요청은 prefill 전에(스트리밍 요청이면 SSE 헤더 전에) HTTP 503 request_memory_budget_exceeded를 받습니다. |
MLXCEL_PREFILL_DEQUANT_MIN_M | 0/off/false/no이면 비활성화, 양의 정수 행 수 | M1 세대 Apple Silicon에서 1024, 다른 Apple 세대와 CUDA, ROCm에서는 off | scale이 입력과 같은 dtype(f16 또는 bf16)인 affine 4-bit projection을 quantized_matmul 대신 dequantize와 dense matmul로 실행하기 시작하는 입력 행 수입니다. 출력 폭이 충분히 넓어 두 경로가 같은 바이트를 내는 projection만 해당합니다. M1 Ultra에서는 1024행 prefill이 모델에 따라 최대 약 15% 빨라지고, prefill 최고 메모리는 0.1~0.6 GB 늘어납니다. |
MLXCEL_MIXED_STEP | 1/true/yes/on이면 활성화, 미설정이거나 그 외 값이면 비활성화 | off | 실험적. 연속 배칭 스케줄러의 prefill/decode 혼합 스텝 프로토타입으로, 한 틱이 두 종류의 작업을 함께 처리합니다. 설계와 기각된 대안은 ADR 0005를 참고하세요. |
MLXCEL_ATTENTION_CHUNK_BUDGET_MB | 부호 없는 정수 MiB, 0이면 비활성화 | 1024 | 고급. CUDA 어텐션 폴백의 score matrix 바이트 예산. fused 커널에 도달할 수 없는 구성에서 폴백은 이 예산 안에 들어가도록 score matrix를 분할합니다. |
MLXCEL_ALLOWED_ORIGINS | 쉼표로 구분한 origin 목록 | 미설정 (제한 없음) | CORS를 지정한 origin으로 제한합니다. 예: https://app.example.com,https://admin.example.com. --allowed-origins 플래그가 우선합니다. |
MLXCEL_WEIGHT_MATERIALIZATION | eager, lazy (대소문자 무시) | eager | 생성 모델의 체크포인트 텐서를 언제 읽고 업로드하는지 정합니다. eager는 모델 로드 끝에 이를 평가하므로 Model loaded in ...이 상주 가중치 기준으로 보고되고 첫 토큰이 가중치 I/O를 치르지 않으며, 서버는 첫 요청이 아니라 시작 시점에 비용을 냅니다. 모달리티 타워 텐서(vision_tower, audio_tower, multi_modal_projector 등)는 지연 상태로 남아 첫 이미지나 오디오 요청에서 로드되므로 텍스트 전용 실행은 이들을 들고 있지 않습니다. lazy는 각 가중치를 처음 도달하는 평가 안에서 materialize하는 이전 동작으로 되돌립니다. tensor-parallel 및 pipeline-parallel 로드와 임베딩·재순위화 로드는 항상 첫 사용 시 materialize합니다. 잘못된 값은 경고 후 eager를 사용합니다. |
MLXCEL_DENSE_GEMV_ROWS | 정수 행 상한, 0이면 비활성화 | 6 | Apple GPU generation 13, 14(M1, M2)에서 이 행 수 이하의 dense matmul을 steel GEMM tile 대신 GEMV kernel로 배치 실행합니다. 0이면 upstream dispatch로 되돌립니다. 프로세스당 한 번 읽습니다. |
관련 운영 문서: 설정, CLI 명령어, 메모리 추정, 런타임 모델 변환.
MLX 런타임 변수¶
다음 non-MLXCEL_* 값은 운영자가 명시하지 않았을 때 CLI가 초기화합니다.
| 변수 | 값 | 자동 기본값 | 설명 |
|---|---|---|---|
MLX_MAX_OPS_PER_BUFFER | 양의 정수 | M1-M4에서 1000, CUDA compute capability 12.1(GB10)의 laguna checkpoint에서 100, ROCm에서 2000(backend 자체 기본값), 그 외 MLX 기본값 | Metal command buffer 하나, CUDA graph 하나, ROCm command batch 하나에 담았다가 commit하는 op 수입니다. CUDA에서는 MLX_MAX_MB_PER_BUFFER와 함께 작동해 어느 한쪽 예산을 넘으면 MLX가 commit하고, ROCm에서는 batch가 MLX_ROCM_MAX_INFLIGHT_MB의 4분의 1을 할당했을 때도 commit합니다. GB10 기본값은 family allowlist로 두는데, 두 예산을 함께 올리면 측정한 모든 workload에서 Laguna는 빨라졌지만 다른 family는 느려졌기 때문입니다. peak memory도 늘어납니다. 명시한 값이 항상 우선합니다. |
MLX_MAX_MB_PER_BUFFER | 양의 정수(바이트가 아니라 원소 수를 20비트 오른쪽 시프트한 값) | Metal: MLX 기본값이며 MLXCEL_DECODE_MB_PER_BUFFER가 decode 단계에서만 올림. CUDA: compute capability 12.1의 laguna 체크포인트에서 1000, 그 밖에는 MLX 기본값 | command buffer 하나나 CUDA graph 하나의 입력 크기 예산. 직접 설정하면 prefill과 decode 모두에 그 값을 고정하고 decode 전용 상향을 끕니다. 명시한 값이 항상 우선합니다. |
MLX_CUDA_SDPA_CACHE_SIZE | unsigned integer | CUDA에서 2000 (MLX 자체 기본값은 256) | MLX cuDNN SDPA execution-plan cache의 LRU 용량. 서로 다른 attention shape마다 miss가 나고, 누적 miss가 용량의 두 배를 넘으면 MLX가 Cache thrashing으로 중단하므로 프롬프트 길이가 다양한 장기 실행 서버에는 큰 예산이 필요합니다. 미리 할당하는 것이 아니라 LRU 상한입니다. 명시한 값이 항상 우선합니다. |
MLX_CUDA_GRAPH_CACHE_SIZE | 양의 정수 | CUDA에서 2000 | 장시간 shape-diverse decode를 위해 MLX graph-cache LRU 용량을 높입니다. 명시한 값이 항상 우선합니다. |
MLX_USE_CUDA_GRAPHS | 0이면 비활성화, 그 외 MLX 지원 값이면 활성화 | MLX 기본값, CUDA Gemma 4는 자동 0 | MLX CUDA graph capture를 제어합니다. mlxcel은 알려진 incremental-decode hazard를 피하기 위해 첫 평가 전에 Gemma 4에서 끄며, 운영자가 어느 방향으로든 명시한 값은 항상 우선합니다. |
MLX_ENABLE_TF32 | 0(full f32), 1(TF32급) | 1(MLX 기본값이며 mlxcel은 덮어쓰지 않음) | CUDA와 Metal에서 f32 GEMM 정밀도를 정하는 MLX 자체 switch입니다. 기본값에서는 f32 GEMM 입력이 10-bit mantissa를 거치는데, CUDA Ampere 이후에서는 cuBLAS TF32 연산, Apple GPU generation 17에서는 NAX kernel을 사용합니다. bf16, f16, 양자화 matmul에는 영향이 없습니다. f32 checkpoint에서 f32와 정확히 같은 출력이 필요하면 0으로 설정하십시오. 저장소의 테스트 바이너리는 이 값을 0으로 고정합니다. |
MLX_PTX_CACHE_DIR | 디렉터리 경로 | CUDA에서 $MLXCEL_CACHE_DIR/cuda-ptx/<MLX-pin> | 영속 NVRTC PTX 캐시입니다. mlxcel은 오래된 커널이 MLX 업그레이드를 넘어 남지 않도록 고정 MLX commit별 디렉터리를 만듭니다. 명시 경로가 항상 우선합니다. |
MLX_BF16_NATIVE | 0이면 legacy 변환, 그 외 값이면 native bf16 유지 | native bf16 | 저수준 bridge 호환 switch입니다. native bf16 경로를 쓸 수 없는 하드웨어/구성에서 0으로 설정하면 bf16 tensor를 로드 시 f32로 확장하므로 메모리 사용량이 늘어납니다. |
MLXCEL_CCCL_DIR | 존재하는 디렉터리 경로 | bundled/build-tree header | CCCL header의 runtime NVRTC include override입니다. 재배치 가능한 CUDA package와 embedder용입니다. |
MLXCEL_CUTLASS_DIR | 존재하는 디렉터리 경로 | bundled/build-tree header | CUTLASS/CuTe header의 runtime NVRTC include override입니다. |
MLXCEL_QUIET_JIT | 설정하면 숨김 | 미설정 | cold-cache CUDA kernel compile 안내를 한 번 출력하는 동작만 숨기며 JIT compile 자체를 끄지는 않습니다. |
외부 모델 저장소¶
Qwen3.8-Flash-Next(qwen4_exp) 체크포인트가 n-gram table과 routed expert를 어디에 두는지 제어합니다. 외부 가중치 저장소를 참고하십시오.
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLXCEL_PLE_IO | direct, mmap, ram | mmap | n-gram table이 외부에 있는 체크포인트(text_config.ple_storage)의 I/O 경로입니다. 그 밖의 값은 로드 오류입니다. mmap은 row 파일을 매핑하고 선택한 row를 page cache를 거쳐 복사합니다. direct는 작은 thread pool로 row를 unbuffered read하고(macOS의 F_NOCACHE, 가능한 Linux의 O_DIRECT) 인접한 row를 합칩니다. ram은 table 전체를 anonymous memory로 한 번 읽고 row cache를 두지 않습니다. ple_storage가 없는 체크포인트에서는 무시하며, 절대 경로 manifest나 모델 디렉터리 밖으로 풀리는 manifest·row 파일은 로드에서 거부합니다. |
MLXCEL_PLE_CACHE_BYTES | 부호 없는 정수 bytes, 0이면 비활성화 | 536870912 (0.5 GiB) | 외부 n-gram table의 row cache 예산입니다. 미설정이면 ple_storage나 manifest의 양수 cache_rows에 row stride를 곱해 cache 크기를 정합니다. ram 모드에서는 쓰지 않습니다. |
MLXCEL_EXPERT_STREAM | 1/on/true, 0/off/false, auto | auto | Metal에서 Qwen3.8-Flash-Next 체크포인트의 routed expert를 SSD에서 스트리밍합니다. routed expert는 상주 slot arena에 두고 token이 그 expert로 라우팅될 때 SSD에서 읽습니다. auto는 routed expert가 Metal working set의 60%를 넘을 때만 켜고, 스트리밍할 수 없는 layout이면 상주 경로로 물러납니다. 1은 강제하며 그런 layout은 로드 오류로 만듭니다. 출력은 exact하게 유지됩니다. 다른 MoE 계열과 CUDA, ROCm은 지원 대상이 아닙니다. |
MLXCEL_EXPERT_ARENA_BYTES | 부호 없는 정수 bytes | 사용 가능한 메모리에서 유도 | MLXCEL_EXPERT_STREAM의 arena 크기이며, layer마다 num_experts개까지 64~num_experts slot으로 나눕니다. 미설정이면 macOS가 지금 내줄 수 있는 메모리에서 상주 non-expert weight와 6 GB headroom을 뺀 값을 expert 크기로 상한을 두고 씁니다. |
다운로더¶
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLXCEL_NO_PROGRESS | 비어 있지 않은 임의 값 | 미설정 | 대화형 다운로드 progress bar를 숨깁니다. NO_COLOR와 CI=true도 progress bar를 숨깁니다. |
MLXCEL_MODELS_DIR | 디렉터리 경로 | 미설정 | download, run/generate/serve/inspect의 repo-id 자동 해석, list, rm이 사용하는 store root입니다. CLI --models-dir가 우선합니다. |
MLXCEL_DEFAULT_ORG | HuggingFace owner | mlx-community | bare 모델 이름 앞에 붙일 owner입니다. |
MLXCEL_ALLOW_INSECURE_ENDPOINT | 비어 있지 않은 임의 값 | 미설정 | HTTPS가 아닌 HF_ENDPOINT로 Hugging Face token을 보내는 것을 허용합니다. 감사된 내부 mirror가 아니면 설정하지 마십시오. |
MLXCEL_EXTRA_CA_CERTS | 파일 경로 (PEM 번들, 최대 1 MiB) | 미설정 | 다운로드 클라이언트가 검증에 쓰는 운영체제 신뢰 저장소에 사용자 CA 인증서를 더하며, 서버의 원격 미디어 fetcher도 같은 클라이언트 설정을 씁니다. TLS를 검사하는 프록시나 사설 레지스트리 뒤에 있는 호스트에서 PEM 번들을 지정하세요. |
프롬프트 캐시¶
대응하는 CLI 플래그가 없을 때 적용됩니다. 동작과 튜닝은 프롬프트 캐시를 참조하십시오.
| 변수 | 값 | 기본값 | 대응 플래그 |
|---|---|---|---|
MLXCEL_PROMPT_CACHE_ENABLED | boolean | true | --prompt-cache-enabled |
MLXCEL_PROMPT_CACHE_CAPACITY_BYTES | 부호 없는 bytes | 모델 기반 산정, 최소 2147483648 | --prompt-cache-capacity-bytes |
MLXCEL_PROMPT_CACHE_MAX_ENTRIES | unsigned integer | 1024 | --prompt-cache-max-entries |
MLXCEL_PROMPT_CACHE_TTL | unsigned seconds | 3600 | --prompt-cache-ttl |
MLXCEL_PROMPT_CACHE_MIN_PREFIX | unsigned tokens | 32 | --prompt-cache-min-prefix |
MLXCEL_PROMPT_CACHE_SNAPSHOT_CAPACITY_BYTES | 부호 없는 bytes | 모델 기반 산정, 최소 536870912 | --prompt-cache-snapshot-capacity-bytes. prefix cache와 별도로 계산하는 session snapshot store의 byte 예산입니다. |
MLXCEL_PROMPT_CACHE_SNAPSHOT_MAX_ENTRIES | unsigned integer | 4096 | --prompt-cache-snapshot-max-entries |
MLXCEL_PROMPT_CACHE_SNAPSHOT_TTL | unsigned seconds | 7200 | --prompt-cache-snapshot-ttl |
APC_ENABLED | boolean | true | --apc-enabled. 호환되지 않는 recurrent/hybrid 모델에서는 자동 비활성화됩니다. |
APC_BLOCK_SIZE | unsigned tokens | 16 | --apc-block-size |
APC_NUM_BLOCKS | unsigned integer | 파생값 | --apc-num-blocks |
APC_HASH | sha256, blake3 | sha256 | --apc-hash |
CLI 플래그가 없고 MLXCEL_PROMPT_CACHE_ENABLED와 llama.cpp 호환 LLAMA_ARG_CACHE_REUSE가 모두 있으면 MLXCEL_PROMPT_CACHE_ENABLED가 우선합니다.
플래그와 변수 어느 쪽도 용량을 정하지 않으면 두 store의 크기는 로드한 모델을 기준으로 정해집니다. snapshot 전용 family(Qwen 3.5, Jamba, Falcon-H1, LFM2, Nemotron-H, Mamba, Granite 4 hybrid, PLaMo 2 같은 recurrent hybrid와 Gemma 3, Gemma 4, Llama 4, AFMoE 같은 attention-cache family)의 snapshot store는 max(536870912, min(6 * entry, left / 4))를 받습니다. 여기서 entry는 대표 길이(--ctx-size, 없으면 8192)에서의 항목 하나 크기이고 left는 모델 자체의 추정 footprint를 빼고 남은 메모리이며, 항목 하나 전체가 left / 2 안에 들어가면 그 크기까지 올라갑니다. dense KV family의 KV store도 같은 방식으로 정하되 2 GiB fallback 아래로는 내려가지 않습니다. store 전체보다 큰 항목이 나오면 store마다 한 번, 올려야 할 플래그 이름을 담은 경고를 남깁니다.
| 변수 | 값 | 기본값 | 대응 플래그 |
|---|---|---|---|
MLXCEL_ENABLE_VLM_PREFIX_CACHE | boolean | false | --enable-vlm-prefix-cache. 같은 이미지를 이어가는 멀티턴 대화를 위한 실험적 VLM 프롬프트 프리픽스 캐시 공유입니다. |
MLXCEL_DISABLE_BOUNDARY_SNAPSHOT | 정확히 1이면 비활성화, true·on을 포함한 그 외 값은 활성 상태 유지 | on | 진단용 킬 스위치. prefill이 히스토리 경계에서 스냅샷을 기증하지 않게 합니다. 이 기증이 스냅샷 전용 계열(Gemma 4, Qwen 3.5, SSM 하이브리드)에서 멀티턴 프롬프트 캐시가 적중하게 만드는 장치라 기본값이 on입니다. |
MLXCEL_DISABLE_CACHE_WARMUP | 정확히 1이면 비활성화, 그 외 값은 활성 상태 유지 | on | 진단용 킬 스위치. 대화의 히스토리 경계 스냅샷을 직전 assistant 답변까지 확장해 다음 턴이 새 user 메시지만 prefill하게 만드는 백그라운드 예열을 끕니다. MLXCEL_DISABLE_BOUNDARY_SNAPSHOT과는 의도적으로 분리돼 있습니다. |
MLXCEL_APC_TRACE | 정확히 1이면 활성화, true·on을 포함한 그 외 값은 off 유지 | off | 진단용. 스케줄러의 프롬프트 캐시 결정마다 추적 줄을 하나씩 남깁니다. 요청이 어느 스냅샷에 매칭됐는지, 발산한 스냅샷이 왜 거절됐는지 포함합니다. 첫 사용 시 한 번만 읽으므로 프로세스 시작 후 설정해도 효과가 없습니다. |
투기적 디코딩¶
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLXCEL_DRAFT_KIND | dflash, mtp, full-model, mlxcel serve와 mlxcel-server에서는 prompt-lookup도 가능 | auto / 없음 | CLI 플래그와 LLAMA_ARG_DRAFT_KIND가 없을 때 --draft-kind의 alias로 쓰입니다. full-model은 --draft-model / --model-draft가 가리키는 일반 작은 체크포인트로 드래프팅합니다. mlxcel generate는 prompt lookup을 --prompt-lookup으로 받으므로 여기서 prompt-lookup을 거부합니다. |
MLXCEL_DRAFT_BLOCK_SIZE | unsigned integer | 시작 시 결정 | CLI 플래그와 LLAMA_ARG_DRAFT_BLOCK_SIZE가 없을 때 쓰는 --draft-block-size 별칭입니다. 셋 다 없으면 초안 모델이 선언한 폭, 측정 기반 장치별 기본값(GB10에서 affine 양자화 대상 모델과 짝지은 DFlash는 4), MTP와 full-model 4, DFlash 16 순으로 정합니다. 시작 시 speculative= 줄이 폭과 그 폭을 만든 규칙을 기록합니다. 작업자가 소유하는 값이므로 바꾸려면 재시작해야 합니다. |
MLXCEL_DFLASH_PER_ROW_ACCEPTS | 0/false/no/off이면 비활성화, 미설정이나 그 밖의 값이면 활성화 | on | batched Qwen 3.5 DFlash round loop에서 행별 accept commit을 씁니다. 각 행이 자기 accept 수를 commit하고, round는 가장 짧은 행에 맞춰 잘리지 않습니다. attention cache는 행별 연속 prefix로 compact하고, query와 key는 각 행의 논리 위치로 회전하며 남은 stale gap은 mask합니다. 각 linear-attention layer는 받아들인 prefix 위에서 convolution과 recurrent state를 다시 재생합니다. 0을 주면 A/B 비교용으로 clamp가 돌아오며, 대상의 attention cache가 quantized, paged, head-trimmed이면 행별 compact가 불가능하므로 자동으로 clamp를 씁니다. |
MLXCEL_ENABLE_MTP_B1 | truthy / falsy | adaptive | B=1 MTP 결정을 고정하고 adaptive profiling을 건너뜁니다. 0, false, no, off는 classic decode를, 그 외 값은 MTP를 강제합니다. 미설정이면 MLXCEL_MTP_ADAPTIVE가 pairing별로 profile하고 기존 hardware rule은 모호한 결과의 fallback으로만 사용합니다. |
MLXCEL_ENABLE_MTP_BATCH | truthy 값 | off | Advanced. parity/debug 테스트를 위해 Gemma 4 MTP batched burst 경로를 강제합니다. row-wise geometry(Gemma 4 31B, CUDA의 12B)에서는 batched burst가 decode-exact이며, 프롬프트가 sliding window 안에 들어가는 최대 4행 window에 대해 CUDA에서 실행됩니다. 그 밖의 window와 Metal, ROCm은 classic decode로 내려갑니다. window는 함께 대기열에 들어왔고 seed가 같으며 token bias가 없는 요청끼리만, 그리고 --max-batch-prefill 1일 때만 만들어집니다. |
MLXCEL_ENABLE_MTP_BATCH_RAGGED | truthy 값 | off | Advanced. Batched MTP window에서 서로 다른 prompt 길이를 허용합니다. MLXCEL_ENABLE_MTP_BATCH가 필요하며 target sliding window를 넘는 prompt는 거부합니다. |
MLXCEL_MTP_DISABLE_DIVERGENT_FIX | 1 | off | Advanced. batched Gemma 4 MTP에서 divergent accept 이후 적용되는 행별 compaction을 되돌립니다. 진단 전용입니다. |
MLXCEL_MTP_BLOCK_CONTROLLER | proxy, requested, 또는 미설정 시 throughput comparator | throughput | --draft-block-size가 drafter의 설정 depth를 넘을 때 B=1 MTP verify width를 결정할 컨트롤러를 고릅니다. 기본값은 설정 depth와 요청 상한 사이에서 32라운드 측정 창을 번갈아 돌리고 라운드 시간당 emit 토큰이 더 많은 쪽을 유지하며, 도전자가 현 우위를 밀어내려면 2% 앞서야 하고 진 쪽은 점점 늘어나는 backoff로 재도전합니다. 근거는 generator에 남으므로 서버 프로세스는 요청 사이에 이를 유지합니다. proxy는 fully-accepted-prefix 게이트로 되돌리는데, acceptance 0.585에서는 설정 prefix가 온전히 수용되는 일이 드물어 Gemma 4 12B 조합을 최적값보다 약 5% 아래에 묶어 둡니다. requested는 첫 라운드부터 요청 width로 draft하며 유해 상한 보호가 꺼지므로, 배포 설정이 아니라 width 스윕 측정 모드입니다. prefer_requested_block_size를 설정하는 drafter(Qwen 3.5 MTP)는 두 컨트롤러를 모두 우회하고, batched B>1 루프는 proxy 게이트에 그대로 머뭅니다. generator당 한 번 읽습니다. |
MLXCEL_MTP_ADAPTIVE | 0/false/no/off면 비활성화, 그 외 값이나 미설정이면 활성화 | on | Adaptive B=1 MTP 정책. Server가 pairing별 qualifying B=1 burst 4개를 profile해 enable/decline verdict를 저장하고 측정이 모호할 때만 static hardware table을 사용합니다. |
MLXCEL_MTP_TICK_SLICE | 0/false/no/off면 비활성화, 그 외 값이나 미설정이면 활성화 | on | 틱 협조형 B=1 MTP 서빙. B=1 MTP 요청이 슬라이스 사이에 스케줄러를 양보하므로 투기적 버스트가 큐 앞단을 막지 않습니다. |
MLXCEL_MTP_SLICE_GRANT_ROUNDS | 음이 아닌 정수 | 8 | 틱 슬라이스 투기 슬롯을 한 번 점유할 때의 grant 예산이며, 실행한 슬라이스 수로 셉니다. 예산을 다 쓰면 슬롯이 대기 중인 다른 요청으로 넘어갑니다. 0이면 회전 없이 활성 요청이 슬롯을 계속 씁니다. |
MLXCEL_MTP_QUANTIZE_DRAFTER | 0/false/no/off면 비활성화, 그 외 값이거나 미설정이면 활성화 | on | Qwen 3.5 MTP drafter의 projection을 로드 시점에 양자화해 bf16 drafter를 약 810 MiB에서 약 228 MiB로 줄입니다. 타깃이 모든 제안을 검증하므로 정확성이 아니라 acceptance를 거는 거래이고, acceptance는 측정 가능한 수준으로 움직이지 않았습니다. 프로세스당 한 번 읽습니다. |
MLXCEL_MTP_ALLOW_INEXACT | 1/true/yes/on이면 활성화, 미설정이나 그 밖의 값이면 비활성화 | off | 시작 시 exactness probe가 multi-token verify block이 single-token chain과 byte 단위로 같지 않다고 보고해도 Qwen 3.5 또는 Gemma 4 MTP를 켭니다. 이를 받아들이면 temperature 0 투기적 출력이 비투기적 디코드와 달라질 수 있습니다. 프로세스당 한 번 읽습니다. |
MLXCEL_MTP_PROBE_PROMPT_LEN | >= 2인 정수 | 8 | 고급. Qwen 3.5 블록 대 단일 토큰 시퀀스 정확성 검사가 비교 전에 프리필하는 프롬프트 길이. 다른 체크포인트에서 길이 스윕을 반복하기 위한 값이며, 늘리면 작업자 시작 시 검사마다 프리필이 한 번씩 늘어납니다. |
MLXCEL_GLM_MTP_HIDDEN_TAP | pre(pre_norm, pre-norm도 가능). 미설정이나 그 밖의 값이면 기본값 | 최종 정규화 뒤 | GLM-4.7-Flash(glm4_moe_lite_mtp) MTP 초안 모델에 넣을 대상 모델 은닉 상태. pre는 승인율 측정용이며 승인율만 바꾸고 출력 토큰은 바꾸지 않습니다. 프로세스당 한 번 읽습니다. |
MLXCEL_PRINT_TOKEN_IDS | presence enables | 미설정 | 진단. 오프라인 mlxcel generate가 프롬프트 ID와 생성 토큰 ID를 [prompt ids (N): ...], [token ids (N): ...] 형식으로 stderr에 출력하여, 일반 실행과 --draft-model 실행을 ID로 비교할 수 있게 합니다. |
MLXCEL_QMV_WIDE | 값이 설정돼 있으면 고정 | 미설정 | 고급. MLX의 qmv_wide 커널 선택을 직접 고정합니다. 명시 설정이 exactness gate의 자체 재시도보다 우선하므로 재시도는 건너뛰고, decline 메시지가 고정한 설정에서는 바이트 동일성 계약에 도달할 수 없었다고 알려줍니다. |
MLXCEL_GDN_CHAIN_PARITY | 0이면 비활성화, 그 외 값이거나 미설정이면 활성화 | on | 고급, 진단용. Qwen 3.5 MTP의 verify·rollback-replay 경로가 쓰는 Metal chain-parity gated-delta 커널을 제어합니다. 0으로 두면 바이트 동일성 계약을 포기하며, 이 커널 자체의 비용을 A/B로 귀속시키려고 존재합니다. Metal 전용입니다. |
MLXCEL_MTP_TREE_BRANCHES | 음이 아닌 정수 | 0 | 고급. tree 라운드가 verify position을 쓸 수 있는 2위 후보 수입니다. 0은 linear tree를 제안하며 이는 정확히 chain 동작이자 유일하게 바이트 동일한 선택입니다. |
MLXCEL_MTP_TREE | truthy 값 | off | Target이 지원한다고 보고할 때 MTP tree-round 구성을 활성화합니다. 기본은 꺼져 있고 chain 경로가 측정된 production baseline입니다. |
MLXCEL_MTP_TREE_MARGIN | float | gate 없음 | 고급. 어떤 스텝을 branch 예산을 쓸 만큼 불확실하다고 볼 상위 두 logit 격차입니다. |
MLXCEL_MTP_DRAFT_PROFILE | 1/true/yes/on이면 활성화, 미설정이거나 그 외 값이면 비활성화 | off | 진단용. MTP drafter step의 각 구성 요소를 다음 것이 시작되기 전에 evaluate해서 MTP drafter step profile 로그가 그래프 빌드 시간이 아니라 실제 GPU 시간을 보고하게 합니다. 추가 동기화가 step 비용을 부풀리므로 비용으로는 비profile 수치를, 귀속에만 profile 수치를 인용하십시오. |
MLXCEL_MTP_GAP_LOG | truthy 값 | off | 진단용. verify position마다 position, token, top1, top2, gap을 출력합니다. exactness 비용을 near-tie 조건부로 만들 수 있는지 판단할 때 씁니다. |
지원되는 사용자 경로는 투기적 디코딩을 참조하십시오.
루프 감지 (Loop Detection)¶
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLXCEL_LOOP_DETECTION | off/0/none/false/disabled, on/default/true/enabled, 또는 MIN,MAX,COUNT | 미설정(계열별 자동 활성화 정책 적용) | tail N-gram 반복("loop") 감지를 위한 전역 override입니다. 미설정이면 계열별 자동 활성화 정책이 그대로 적용되며 Gemma 4 계열은 기본적으로 켜져 있습니다. off는 모든 요청에서 강제로 비활성화하고 on은 권장 threshold(min=1, max=20, count=4)를 강제하며 1,20,4 같은 명시적 triple은 min_pattern_size, max_pattern_size, min_count를 직접 지정합니다. API 요청 본문의 max_pattern_size / min_pattern_size / min_count가 있으면 이 변수보다 항상 우선합니다. 값이 잘못되면 경고 후 무시합니다. |
루프 감지는 생성 스트림 끝에서 짧은 토큰 패턴이 반복될 때 생성을 조기 종료하는 기능으로, vLLM이 SamplingParams에서 제공하는 것과 같은 방식의 대응책입니다. 요청 단위 필드는 REST API 레퍼런스를 참조하십시오.
블록 디퓨전¶
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLXCEL_DIFFUSION_DEBUG_CANVAS | 1 | off | Test-only. 구현 간 parity 테스트를 위해 난수 캔버스 초기화를 고정된 결정적 패턴으로 대체합니다. 일반 생성에는 쓸모 있는 출력이 나오지 않습니다. |
생성 메커니즘과 --diffusion-* CLI 플래그는 블록 디퓨전을 참조하십시오.
KV 캐시와 TurboQuant¶
가능하면 --cache-type-k, --cache-type-v, --kv-cache-mode, --turbo-boundary-v, 서버 batch KV quantization 플래그를 우선 사용하십시오.
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLXCEL_KV_CACHE_BUDGET | auto, bytes, none, 또는 0 | auto | 통합 paged KV 캐시 풀의 블록 예산입니다. auto는 memory-aware cap을 도출하고 none/0은 opt-out합니다. --kv-cache-budget 별칭이며 paged 디코드 스토리지 백엔드의 풀 기반 FP16 모델에 적용됩니다. |
MLXCEL_KV_BOUNDARY_V_LAYERS | integer count | 2 | Turbo4 계열에서 더 높은 precision으로 유지할 첫/마지막 V layer 개수입니다. 0은 비활성화합니다. --turbo-boundary-v가 cache construction 전에 이 값을 씁니다. |
MLXCEL_TURBO_BOUNDARY_V | integer count | alias | MLXCEL_KV_BOUNDARY_V_LAYERS 호환 별칭입니다. 둘 다 있으면 primary 이름이 우선합니다. |
MLXCEL_KV_SKIP_LAST_LAYER | boolean | true | 연속 배칭 KV quantization의 --kv-skip-last-layer에 대한 mlxcel-native 폴백입니다. |
MLXCEL_SPARSE_V_THRESHOLD | non-negative float | 1e-6 | Sparse-V alive threshold. 0은 Sparse-V를 비활성화하고 잘못된 값은 경고 후 기본값을 사용합니다. |
MLXCEL_SPARSE_V_COUNT | 출력 파일 경로 | 미설정(off) | Diagnostic. 비어 있지 않은 경로를 지정하면 단일 토큰 Turbo4Asym decode step마다 call_idx,kv_tokens,skipped,total 행을 해당 CSV에 추가하며, 모든 head에서 MLXCEL_SPARSE_V_THRESHOLD 미만인 post-softmax attention weight 개수를 셉니다. graph 전용 side computation이라 미설정 시 비용이 없습니다. benchmarks/sparse_v_skip/의 측정 자산을 참조하십시오. |
MLXCEL_SPARSE_V_KERNEL | falsy disables | macOS에서 enabled | fused Sparse-V/dequant Metal kernel 허용 여부입니다. 0, false, off, no로 graph 폴백을 강제합니다. |
MLXCEL_TURBO4_DEQUANT_SDPA | falsy disables | on | symmetric Turbo4의 dequant-first SDPA 경로를 제어합니다. |
MLXCEL_TURBO4_DELEGATED_DEQUANT_SDPA | falsy disables | on | Turbo4Delegated의 기본 dequant-first SDPA 경로를 제어합니다. |
MLXCEL_TURBO4_DELEGATED_FUSED | truthy enables | off | Advanced. dequant-first SDPA를 끈 상태에서 비교하기 위한 이전 custom fused delegated-kernel 경로를 활성화합니다. |
MLXCEL_TURBO4_DELEGATED_FP16_FAST_PATH | truthy enables | off | Advanced. delegated mode에서 packed sidecar를 유지하면서 unified FP16 V working set을 둡니다. |
MLXCEL_TURBO4_DELEGATED_FP16_SIDECARS | predecode, eager, lazy, on-demand | predecode | delegated FP16 fast path의 sidecar 유지 정책입니다. |
MLXCEL_ENABLE_DIRECT_PREFILL_CACHE_STORE | presence enables | off | Advanced. 적용 가능한 경우 incoming prefill tensor를 초기 KV cache buffer로 직접 설치합니다. |
MLXCEL_KV_INPLACE_WRITE | 0/false/off/no이면 비활성화 | on | single-token FP16 디코드 단계가 K/V 행을 slice_update 대신 dense cache buffer에 제자리로 씁니다. slice_update는 이전 파이프라인 단계가 아직 buffer를 잡고 있는 동안 매 단계 cache 전체를 복사합니다. 제자리 쓰기는 현재 buffer를 cache 자신의 마지막 쓰기가 만든 경우에만 쓰므로, 프롬프트 캐시 snapshot과 공유하는 buffer는 덮어쓰지 않습니다. 서버의 paged pool도 시퀀스 하나가 소유한 block의 행에 같은 제자리 쓰기를 사용합니다. M1 Ultra 측정에서 command-r7b 4-bit 디코드가 context 16에서 +2.5%, 2048에서 +8.2%였습니다. |
MLXCEL_DIAG_SKIP_DECODE_KV_WRITE | 존재하면 활성화 | off | 진단용이며 설정 중에는 출력이 유효하지 않습니다. single-token FP16 디코드 단계가 K/V 행 쓰기를 건너뛰므로, 디코드 벤치마크로 단계별 cache 쓰기를 없앴을 때의 상한을 읽을 수 있습니다. cache가 늘지 않고 position도 진행하지 않으므로 생성 텍스트에는 의미가 없습니다. 처음 사용할 때 경고를 한 번 출력합니다. |
MLXCEL_TURBO4_ASYM_DEQUANT_SDPA | falsy면 비활성화 | on | 비대칭 Turbo4Asym(FP16 K + 4비트 V)의 dequant-first SDPA 경로. falsy 값이면 손실이 있는 sparse-V 경로를 씁니다. |
MLXCEL_MLA_ABSORBED | 1/true/on/yes면 활성화 | off | DeepSeek 계열의 matrix-absorbed MLA 디코드. 압축 해제한 head별 K/V 대신 압축 latent (ckv, kpe)를 캐시하므로 긴 컨텍스트에서 KV 사용량이 크게 줄어듭니다. |
MLXCEL_MLA_SPLIT_KV | 1/true/on/yes면 활성화 | off | 고급. latent 구간을 청크로 나누고 각 부분 softmax 상태를 merge 커널로 합칩니다. MLXCEL_MLA_ABSORBED가 필요합니다. |
MLXCEL_CASCADE_ATTENTION | 1/true/on/yes면 활성화, 0/false/off/no는 kill switch | off | 2단계 cascade 디코드. 여러 동시 시퀀스가 공유하는 페이지 단위 프롬프트 prefix를 한 번만 어텐션한 뒤 요청별 suffix 상태와 합쳐, 공유 시스템 프롬프트에서 생기는 중복 계산을 없앱니다. |
MLXCEL_CASCADE_MIN_SHARED_PAGES | 음이 아닌 정수 | 16 | cascade를 쓰기 전에 서브그룹이 공유해야 하는 온전한 페이지 수. 기본 블록 크기 32에서 16페이지는 512토큰입니다. 이 값에 못 미치면 평면 launch를 씁니다. |
MLXCEL_CASCADE_MIN_MEMBERS | 2 이상의 정수 | 2 | cascade를 쓰기 전에 구간을 공유해야 하는 시퀀스 수. 2 미만이면 없앨 중복 계산이 없습니다. |
MLXCEL_NATIVE_WINDOW_SDPA | auto, 1/on, 0/off | auto | 고급, A/B selector. 반환된 key가 window를 넘어선 뒤 sliding window 위의 causal multi-query prefill이 Metal의 MLX SDPA에 도달하는 방식을 고릅니다. MLX 네이티브 window_size 인자(mask_mode = "causal")를 쓰거나, array-mask 경로가 늘 쓰던 명시적 full-width windowed-causal mask를 씁니다. auto는 MLX가 해당 shape의 fused kernel을 보고할 때만 native 경로를 쓰고, 1은 graph fallback을 거치더라도 native 진입점을 강제하며, 0은 legacy mask 경로를 유지합니다. 정적으로 자격을 갖춘 Metal prefill shape에만 적용되며(softcap 없음, attention sink 없음, 추가 mask array 없음, query block이 8행 초과이고 key 수 이하, 지원 head dim, f16/bf16/f32), decode, 작은 verify block, CPU, CUDA, ROCm은 기존 경로를 유지합니다. 프로세스당 한 번 읽습니다. |
MLXCEL_LOG_WINDOW_SDPA | 1/sampled, all | off | causal window prefill마다 [mlxcel][window-sdpa] 줄을 stderr로 출력하며, 선택한 경로(native 또는 legacy), 그 이유, 양쪽 arm에 대한 MLX의 fused-kernel 판정, q/k shape를 담습니다. sampled(1과 같음)는 처음 16회 호출과 이후 256회마다 출력합니다. 프로세스당 한 번 읽습니다. |
비디오와 로컬 미디어¶
비디오 입력을 지원하는 VLM 요청 처리에 적용됩니다. 비전-언어 모델을 참조하십시오.
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLXCEL_ALLOW_PRIVATE_MEDIA_URLS | 1이면 활성화 | 미설정(거부 동작) | 서버가 원격 미디어를 가져올 때 수행하는 네트워크 주소 검사를 전부 통과로 바꿉니다. 내부 오브젝트 스토어에서 미디어를 실제로 서비스하는 배포를 위한 것입니다. 사설망에 닿는 fetch proxy는 운영자가 명시적으로 결정해야 하므로 기본값은 off입니다. 서버를 참조하십시오. |
MLXCEL_VIDEO_DIR_ALLOWLIST | 쉼표로 구분한 디렉터리 | 미설정 | 로컬 video_url 파일 경로는 정규화한 이 디렉터리 중 하나 아래로 해석될 때만 허용됩니다. 디렉터리는 owner만 쓸 수 있게 두십시오. group이나 world가 쓸 수 있는 항목은 시작 시 경고합니다. 네이티브 비디오 family와 video-frames fallback 모두 서버 시작 시(router가 로드하는 모델마다 한 번, in-process mlxcel run에서도 한 번) 목록을 해석하므로, 아직 없는 항목은 경고와 함께 제외됩니다. 변수를 바꾸거나 목록의 디렉터리를 만든 뒤에는 다시 시작하십시오. data URI와 HTTP(S) 비디오 payload는 ffmpeg가 읽기 전에 비공개 0700 디렉터리 안의 새 0600 파일로 기록됩니다. |
MLXCEL_VIDEO_MAX_FRAMES | unsigned integer (최소 2) | 16 | 고유 비디오 경로가 없는 체크포인트에 클립을 순서가 있는 정지 이미지로 보낼 때 남기는 프레임 수. 2보다 작은 값은 2로 올립니다. mlxcel serve, mlxcel-server, mlxcel generate의 --video-max-frames와 같습니다. 프레임은 --max-images에 포함됩니다. |
MLXCEL_VIDEO_FPS | > 0인 float | 2.0 | 요청의 video_url.fps가 없을 때 그 대체 경로의 샘플링 속도. 서버 전용(--video-fps)이며 CLI는 --fps를 씁니다. |
MLXCEL_VIDEO_MAX_PIXELS | unsigned integer | 16777216 | source video의 width * height가 cap을 넘으면 거부합니다. |
MLXCEL_VIDEO_MAX_DURATION_SEC | float seconds | 600 | source video가 cap보다 길면 거부합니다. |
MLXCEL_VIDEO_MAX_PNG_FRAME_BYTES | unsigned bytes | 268435456 | ffmpeg PNG stream splitter의 frame당 cap입니다. |
MLXCEL_MEDIA_NAT64_PREFIXES | 쉼표로 구분한 IPv6 prefix addr/len, len은 32, 40, 48, 56, 64, 96 중 하나 | 미설정(prefix 없음) | 미디어 주소 정책이 해독하는 NAT64 prefix입니다. 목록에 있는 prefix 안의 주소는 내장한 IPv4 주소로 판정하므로(RFC 6052), 합성된 공개 대상은 가져오고 합성된 사설 대상은 계속 거부합니다. well-known 64:ff9b::/96은 항상 해독합니다. 목록의 prefix는 global unicast 2000::/3 또는 RFC 8215 local-use block 64:ff9b:1::/48 안에 있어야 하며, 파싱되지 않거나 다른 길이를 쓰거나 그 범위를 벗어난 항목은 이 변수 이름을 밝히며 시작을 실패시킵니다. 시작 시 한 번 읽고, MLXCEL_ALLOW_PRIVATE_MEDIA_URLS=1은 여전히 모든 주소 검사를 끕니다. |
MLXCEL_KIMI_K3_MAX_MEDIA_TOKENS(기본값 32768)는 Kimi K3 요청 하나의 이미지가 합쳐서 쓸 수 있는 media token 수를 제한합니다. 체크포인트의 크기 조정 규칙에서는 이미지 한 장이 17,000토큰 가까이 될 수 있고 타워의 어텐션이 이미지별 patch 수의 제곱에 비례하기 때문입니다.
서버 이미지 admission 한도에는 llama.cpp 호환 별칭 LLAMA_ARG_MAX_IMAGE_WIDTH, LLAMA_ARG_MAX_IMAGE_HEIGHT도 있으며 둘 다 기본값은 16384입니다. serve CLI 레퍼런스를 참조하십시오.
Disaggregated router 보안¶
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLXCEL_DECODE_ALLOWLIST | 쉼표로 구분한 numeric IP:port | 미설정(경고 후 permissive) | Prefill worker가 KV handoff를 위해 연결할 수 있는 router-selectable decode node 전체 집합입니다. 잘못된 항목과 hostname은 건너뜁니다. 신뢰 cluster에서 forged decode_target이 prompt-derived KV state를 다른 곳으로 보내지 못하게 설정하십시오. |
MLXCEL_ROUTER_STATS_VERBOSE | 1/true/yes/on이면 활성화 | off | GET /router/stats에 raw peer address를 포함합니다. 기본값은 address를 가리고 안정적인 router node id만 표시합니다. |
MLXCEL_CLUSTER_SECRET_FILE | cluster secret을 담은 파일 경로 | 미설정 | distributed pipeline role과 disaggregated role을 위한 cluster별 pre-shared secret입니다. 파일은 최소 16바이트를 담아야 하며(끝의 newline 하나는 무시), openssl rand -hex 32 > cluster.secret으로 만들고 모든 호스트에 복사하십시오. 우선순위는 --cluster-secret-file, 이 변수, MLXCEL_CLUSTER_SECRET, cluster TOML의 [cluster] secret_file 순서이며, 두 환경 변수를 모두 설정하면 오류입니다. 모든 pipeline role과 disaggregated role은 secret 없이는 시작을 거부하고, 단일 노드 서빙은 이 값을 읽지 않습니다. |
MLXCEL_CLUSTER_SECRET | secret 값(최소 16바이트) | 미설정 | secret을 환경 변수로 주입하는 container와 service unit을 위한 MLXCEL_CLUSTER_SECRET_FILE의 inline 형태입니다. 우선순위와 fail-closed 동작은 같고 값은 절대 로그에 남지 않습니다. 환경 변수는 같은 사용자의 다른 프로세스에 보이므로 공유 호스트에서는 파일 형태를 권장하며, 두 환경 변수를 모두 설정하면 오류입니다. |
OpenXLA 백엔드¶
루트 Cargo manifest는 xla-backend, xla-iree, 그리고 xla-diagnostics 계열 feature를 선언하며, 전부 기본값이 off입니다. 따라서 기본 Metal/CUDA 바이너리는 OpenXLA 경로를 아예 컴파일하지 않고 이 이름들도 읽지 않습니다. --features xla-backend(seam과 StableHLO emitter까지. 네이티브 툴체인 불필요)나 --features xla-iree(실제 IREE 런타임까지. 빌드 시점에 IREE_DIST, IREE_CUDA_HOME, IREE_MACOS_HOME, IREE_ROCM_HOME 중 하나가 필요)로 빌드하면 읽습니다. OpenXLA 백엔드 소스 상태를 참조하십시오.
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLXCEL_BACKEND | xla | MLX(미설정) | OpenXLA/StableHLO 백엔드를 선택합니다. xla-backend 없이 빌드한 기본 바이너리에서는 동작하지 않습니다. |
MLXCEL_XLA_DEVICE | local-task, local-sync, metal, cuda, hip | Apple Silicon은 metal, 그 외 local-task | OpenXLA 엔진의 IREE HAL device입니다. local-task / local-sync는 CPU driver입니다. cuda와 hip은 자동으로 선택되지 않으며, 해당 driver로 빌드한 런타임이 필요하므로 직접 지정해야 합니다. |
MLXCEL_XLA_LOCAL_TASK_WORKERS | 1부터 256까지의 정수 | 미설정(현재 NUMA node의 물리 core당 worker 하나, 최대 64) | local-task(CPU) device의 worker 수입니다. shim은 IREE 전역 task flag 대신 8 MiB worker stack으로 그 device를 직접 만드는데, 기본 128 KiB stack은 CUDA나 ROCm 런타임 라이브러리를 로드한 프로세스의 static TLS를 담지 못하기 때문입니다. device가 local-task일 때만 읽으며, 0, 음수, 정수가 아닌 값, 256 초과 값은 clamp되지 않고 이 변수 이름을 밝히는 오류로 로드를 실패시킵니다. |
MLXCEL_XLA_ROCM_TARGET | gfx* chip 하나. 예: gfx1151, gfx1100, gfx942 | 미설정 | MLXCEL_XLA_DEVICE=hip일 때 필수이며 iree-compile에 --iree-rocm-target으로 전달됩니다. 미설정, 빈 값, 쉼표로 구분한 값, gfx*가 아닌 값은 iree-compile 실행 전에 거부됩니다. scripts/iree/setup-rocm.sh --env가 감지한 호스트 target을 export합니다. |
MLXCEL_XLA_PRECISION | f32, f16, bf16 | GPU device(metal, cuda, hip)에서 f16, CPU device에서 f32 | StableHLO emitter의 matmul contraction precision입니다. norm, softmax, RoPE는 항상 f32로 유지됩니다. metal device에서 bf16은 로드 시 거부됩니다(Metal GPU target은 bf16 연산을 지원하지 않음). 인식할 수 없는 값은 [compile] 진단과 함께 거부됩니다. |
MLXCEL_XLA_QUANT | packed | 미설정 | MLX 4/8-bit 양자화 가중치를 로드 시 dequantize하는 대신 graph 안에 packed 상태로 두고 device 쪽에서 dequantize합니다. cuda, hip, local-task에서 지원하고, packed prefill이 fault를 내는 metal device에서는 로드 시 거부됩니다. token-exact하지만 기본은 off인데, IREE가 in-graph dequant와 matmul을 별도 dispatch로 낮추기 때문에 host에서 dequantize한 가중치의 디코드가 더 빠르기 때문입니다. packed는 host upload 바이트를 줄입니다. |
MLXCEL_XLA_VISION_BACKEND | auto, iree, host | auto | CLI와 서버가 공유하는 prepared-prefill 경로의 vision tower/projector 선택자입니다. auto는 상주 IREE 모듈이 있으면 그것을 쓰고 없으면 host(MLX) 경로로 폴백하며, iree는 폴백하지 않고 시작에 실패하고, host는 vision을 MLX에 둡니다. Qwen2-VL과 Qwen2.5-VL은 host vision 경로가 없으므로 xla-iree 빌드가 필요하고 host를 거부합니다. |
MLXCEL_XLA_CONTEXT_CAPACITY | 1..=2147483647 범위의 정수 토큰 수 | 미설정(CLI는 요청 크기, 서버는 유도 bucket) | 설정하면 OpenXLA prefill/decode 쌍과 모든 KV 버퍼에 정적 시퀀스 용량 하나를 고정하고, effective_prompt_len + max_new_tokens가 이를 넘는 요청은 native 실행 전에 거부합니다. 미설정이면 CLI가 요청에서 shape를 정합니다. 프롬프트와 생성 예산을 담는 [256, 512, 1024, 2048, 4096, 8192, 16384, 32768]의 가장 작은 rung을 고르고 체크포인트의 max_position_embeddings로 상한을 둡니다. 무제한 -n -1은 512토큰 예산으로 shape를 잡고 허용 예산을 shape가 남기는 공간에 맞춰 자릅니다. 용량을 키우면 KV 메모리는 선형으로, prefill attention mask는 제곱으로 늘고 별도 컴파일 artifact가 생깁니다. StableHLO 동적 시퀀스 shape는 쓰지 않습니다. |
MLXCEL_XLA_CONTEXT_BUCKETS | 쉼표로 구분한 정수, 각각 1..=2147483647 범위 | 유도값(256, min(4096, 체크포인트 window), worst-case 이미지 floor가 더 크면 그 값) | 서버 OpenXLA 연속 배칭 worker 전용(xla-iree)으로, 컴파일하는 정적 context shape를 정렬·중복 제거해 지정합니다. 유도보다 트래픽을 잘 아는 운영자를 위한 값입니다. 요청은 자기를 받아들이는 가장 작은 bucket으로 가므로 짧은 요청은 256 shape를 유지합니다. MLXCEL_XLA_CONTEXT_CAPACITY가 설정되면 무시되며 bucket은 하나가 됩니다. 빈 목록, 0, 정수가 아닌 값은 worker 시작을 중단시킵니다. CLI 경로는 읽지 않습니다. |
MLXCEL_XLA_IREE_COMPILE | iree-compile 경로 | 빌드 시점 고정값 | 번들 그래프를 vmfb로 낮출 때 쓰는 iree-compile 바이너리를 런타임에 덮어씁니다. CUDA와 ROCm 소스 런타임 빌드는 IREE_CUDA_COMPILE 또는 IREE_ROCM_COMPILE에서 짝이 맞는 컴파일러를 빌드 시점에 고정합니다. |
MLXCEL_XLA_LOAD_TIMING | 설정하면 활성화 | off | 진단용. OpenXLA 언어 모델 로드와 Gemma3n audio 런타임 로드의 각 단계마다 mlxcel-xla-load: phase=<label> elapsed_ms=<n> rss_kib=<n> 줄을 stderr로 출력합니다. RSS는 /proc/self/status에서 읽으므로 Linux가 아닌 곳에서는 unavailable로 표시됩니다. |
MLXCEL_PHI4MM_AUDIO_PRECISION | f32, bf16 | f32 | 남아 있는 XLA 소스의 Phi-4 multimodal audio encoder precision입니다. 잘못된 값이면 graph 준비가 실패합니다. |
MLXCEL_XLA_GEMMA3N_TRACE_LAYER_START | 0 이상의 layer index | 미설정 | Gemma 3n host/XLA trace 비교를 시작할 진단 layer입니다. |
MLXCEL_XLA_REFERENCE_EXPECT_UNEXPANDED_IDS | test/reference policy 값 | 미설정 | 내부 XLA reference-admission contract switch이며 기본 서버 설정이 아닙니다. |
MLXCEL_XLA_IREE_RUN_MODULE | 실행 파일 경로 | IREE 도구에서 탐색 | 남아 있는 Gemma 3n probe에서 쓸 iree-run-module test/tool override입니다. |
MLXCEL_XLA_IREE_TEST_TARGET | local 또는 probe target | local | ignored retained-backend 통합 probe의 target selector입니다. |
서버 오디오 admission¶
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLXCEL_AUDIO_QUEUE_DEPTH | 부호 없는 정수 | 8 | 오디오 워커 커맨드 큐의 상한. 큐가 가득 차면 새 오디오 요청은 구조화된 503을 받습니다. --audio-queue-depth 플래그가 우선합니다. |
MLXCEL_AUDIO_REQUEST_TIMEOUT_SECS | 부호 없는 정수 초 | 120 | 요청별 응답 타임아웃. 멈췄거나 비정상적으로 느린 오디오 요청이 큐 슬롯을 붙들지 않고 놓아줍니다. --audio-request-timeout-secs 플래그가 우선합니다. |
서버 저장소와 요청 한도¶
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLXCEL_DECODE_TIMEOUT | 부호 없는 정수 초 | 600 | 요청별 디코드 워치독. 생성이 토큰을 더 내놓지 못하는 요청을 취소합니다. --decode-timeout 플래그가 우선합니다. --timeout / LLAMA_ARG_TIMEOUT은 다른 설정으로, llama-server의 3600초 기본값을 쓰는 HTTP 소켓 read/write 타임아웃입니다. 디코드 워치독을 함께 설정하지 않고 --timeout만 설정하면 시작 시 마이그레이션 경고를 남깁니다. 0은 모든 요청을 만료시키는 대신 내장 300초 가드로 물러나며 경고합니다. |
MLXCEL_ENABLE_SETTINGS_ENDPOINT | true/false, 1/0, yes/no, on/off | 미설정(엔드포인트 없음) | --settings의 환경 변수 쌍입니다. /v1/settings와 /settings의 인증된 GET과 PATCH를 켭니다. 플래그 파싱 뒤에 적용되므로 명령줄의 명시적 --settings가 우선합니다. 해석할 수 없는 값은 경고와 함께 무시합니다. API 키 없이 loopback이 아닌 TCP 리스너에 엔드포인트를 노출하는 것은 시작 시점에 거부합니다. |
MLXCEL_CONVERSATION_STORE_MAX_BYTES | 부호 없는 정수 바이트 | 67108864 | conversation store가 보유하는 대략적인 JSON 바이트 수를 제한합니다. --conversation-store-max-bytes 플래그가 우선합니다. |
MLXCEL_RESPONSES_STORE_MAX_BYTES | 부호 없는 정수 바이트 | 268435456 | Responses API store가 보유하는 대략적인 JSON 바이트 수를 제한합니다. --responses-store-max-bytes 플래그가 우선합니다. |
MLXCEL_RESPONSES_BACKGROUND_MAX_JOBS | 부호 없는 정수 | 64 | 동시에 활성인 background: true Responses API job 수의 상한입니다. job마다 생성 thread를 하나씩 점유하며 0은 background 모드를 끕니다. --responses-background-max-jobs 플래그가 우선합니다. |
MLXCEL_RESPONSES_BACKGROUND_EPHEMERAL_TTL_SECS | 1 이상의 부호 없는 정수 초 | 600 | store: false인 background job의 최종 객체를 polling으로 읽을 수 있는 시간입니다. 0은 시작 실패입니다. --responses-background-ephemeral-ttl-secs 플래그가 우선합니다. |
MLXCEL_ANTHROPIC_PIN_BILLING_STAMP | 0이면 비활성화(정확히 0만), 그 밖의 값이나 미설정은 pinning 유지 | 미설정(pinning 켜짐) | Anthropic Messages의 system 텍스트에 들어 있는 Claude Code x-anthropic-billing-header에서 요청마다 바뀌는 cch와 세션마다 바뀌는 cc_version 일부를 같은 길이의 f 문자로 바꿔 렌더링합니다. 덕분에 다음 턴이 이전 프롬프트의 캐시된 prefix를 다시 씁니다. 정확히 그 header로 시작하는 최상위 system 텍스트만 대상입니다. OnceLock으로 한 번 읽습니다. 서버를 참고하십시오. |
MLXCEL_REASONING_ALIAS_FIELD | reasoning, none | reasoning | --reasoning-alias-field의 네이티브 환경 변수 짝입니다. reasoning은 모든 Chat Completions reasoning_content 값을 동일한 reasoning 필드로 복제하고, none은 그 복제 별칭만 억제합니다. 모델이 사고하는지 여부, 기존 reasoning_content 필드, Responses API reasoning 이벤트, Anthropic thinking 블록은 바뀌지 않습니다. |
MLXCEL_REASONING_ECHO_MAX_BYTES | 0 이상의 정수 bytes, 0이면 비활성화 | 16777216 (16 MiB) | 프롬프트 캐시가 켜져 있을 때 생성한 reasoning을 content만 담은 assistant history 턴에 다시 넣는 store의 byte 예산입니다(/v1/chat/completions, /v1/responses, /v1/messages). 덕분에 template이 이전 턴을 생성 당시 모습대로 다시 렌더링하고 후속 요청도 프롬프트 캐시에 hit합니다. store는 항목 수를 4096개로 제한하고, 예산의 8분의 1보다 큰 trace 하나는 받지 않으며, 가장 오래 쓰지 않은 항목부터 축출합니다. 파싱할 수 없는 값은 경고 후 기본값을 유지합니다. 서버 시작 시 읽습니다. |
임베딩 및 재순위화¶
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLXCEL_EMBEDDING_MODEL | 로컬 경로 또는 HuggingFace 저장소 식별자 | 미설정 | 주 모델 옆에 불러올 임베딩 체크포인트입니다. LLAMA_ARG_EMBEDDING_MODEL은 호환 별칭이며 --embedding-model이 우선합니다. |
MLXCEL_RERANKER_MODEL | 로컬 경로 또는 HuggingFace 저장소 식별자 | 미설정 | 주 모델 옆에 불러올 재순위화 모델입니다. LLAMA_ARG_RERANKER_MODEL은 호환 별칭이며 --reranker-model이 우선합니다. |
MLXCEL_IMAGE_MODEL | 로컬 경로 또는 owner/name repo-id | 미설정 | -m의 채팅 모델 옆에서 /v1/images/generations와 /v1/images/edits로 서비스하는 이미지 생성 체크포인트이며, -m과 같은 방식으로 해석합니다. -m이 없으면(그리고 LLAMA_ARG_MODEL이나 --hf-repo도 없으면) 이미지 전용 서빙을 선택합니다. 체크포인트가 서빙 모델이 되고, 채팅과 토크나이저 경로는 501로 응답하며, /v1/models는 이 모델을 image 기능으로 나열합니다. --image-model이 우선합니다. 로드 실패는 시작 오류입니다. |
MLXCEL_EMBEDDING_BATCH_SIZE | 양의 정수 | 16 | 임베딩 순전파 한 번에 넣는 텍스트 수입니다. --embedding-batch-size가 우선합니다. |
MLXCEL_EMBEDDING_MAX_LENGTH | 양의 정수 토큰 수 | 체크포인트에서 결정, 최대 8192 | 자동으로 구한 임베딩 입력 상한을 더 낮춥니다. --embedding-max-length가 우선합니다. |
MLXCEL_EMBEDDING_QUEUE_DEPTH | 양의 정수 | 8 | 각 임베딩·재순위화 작업자 명령 큐의 상한입니다. --embedding-queue-depth가 우선합니다. |
MLXCEL_EMBEDDING_REQUEST_TIMEOUT_SECS | 부호 없는 정수 초 | 120 | 임베딩과 재순위화 요청의 응답 제한 시간입니다. 0이면 기본값으로 돌아가고 명령행 옵션이 우선합니다. |
MLXCEL_RERANK_BATCH_SIZE | 부호 없는 정수 | 0 | 재순위화 순전파 한 번에 넣는 쌍 수입니다. 0은 종류별 기본값인 텍스트 8, 멀티모달 2를 선택합니다. |
MLXCEL_EMBEDDING_POOLING | cls, mean, max, lasttoken | 체크포인트·계열 | --pooling이 없을 때 사용하는 진단용 풀링 덮어쓰기입니다. 명령행 옵션이 이 변수와 체크포인트 설정보다 우선합니다. |
모델 로딩 방식, 큐 오류, 정규화 제어는 임베딩 및 재순위화 API를 참고하십시오.
빌드 타임 변수¶
mlxcel-core build script가 읽습니다. CUDA 빌드는 CUDA_HOME, MLX_CUDA_ARCHITECTURES 같은 non-MLXCEL_* 변수도 사용합니다. 소스에서 빌드하기를 참조하십시오.
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLXCEL_BUILD_METAL | 1/0, on/off, true/false, yes/no | macOS 빌드에서 on | CMake MLX_BUILD_METAL 설정을 override합니다. 잘못된 값은 빌드를 실패시킵니다. |
MLXCEL_BUILD_ACCELERATE | 1/0, on/off, true/false, yes/no | macOS 빌드에서 on | CMake MLX_BUILD_ACCELERATE 설정을 override합니다. 잘못된 값은 빌드를 실패시킵니다. |
MLXCEL_CXX_MARCH | -march= 값 또는 none | native | 릴리스 빌드에서 C++ 브리지의 ISA 기준선. 빌드 호스트가 아닌 다른 머신에서 돌릴 바이너리라면 x86-64-v3 같은 이식성 있는 기준선을 지정하고, none이면 -march를 아예 넣지 않습니다. |
MLXCEL_GIT_COMMIT | 임의 문자열, 예: commit hash | 미설정 | build script가 아니라 mlxcel crate가 compile time에(option_env!) 읽습니다. cargo build 환경에 있던 값이 WebUI bootstrap 응답(GET /ui-api/v1/bootstrap)의 server.build.git_commit으로 embed됩니다. 미설정이면 null로 보고합니다. 저장소는 이 값을 설정하지 않으며 packager가 설정할 수 있습니다. 런타임에 설정해도 효과가 없습니다. |
ROCm 빌드(--features rocm)는 ROCm CMake 패키지, hipcc, rocminfo 위치로 ROCM_PATH(기본값 /opt/rocm)를, gfx 타겟으로 MLX_ROCM_ARCHITECTURES(기본값은 rocminfo로 감지)를 읽습니다. 소스에서 빌드하기를 참조하십시오. ROCm backend가 런타임에 읽는 변수는 ROCm 백엔드 변수에 정리되어 있습니다.
네이티브 Windows 빌드(실험 단계)는 읽는 변수가 다릅니다. Windows CUDA 빌드는 CUDA_PATH(CUDA_HOME이 있으면 먼저 씀)를 읽고 cuDNN을 CUDA 툴킷 디렉터리에 풀지 않은 경우에만 cuDNN 루트인 CUDNN_PATH를 읽습니다. Windows ROCm 빌드는 ROCM_PATH를 읽고 이 값이 없거나 비어 있으면 HIP_PATH를 읽습니다(기본값이 없어 둘 다 없으면 두 이름을 밝히며 빌드가 실패함). gfx 타겟은 rocminfo 대신 SDK의 amdgpu-arch.exe로 감지하며 MLX_ROCM_ARCHITECTURES로 덮어쓸 수 있습니다. Windows에서 이 경로들은 빌드할 때만 필요합니다. 실험적 지원을 참조하십시오.
--features xla-iree OpenXLA 빌드는 IREE_DIST를 읽거나, 소스 빌드 런타임 하나와 짝이 맞는 컴파일러를 읽습니다. IREE_CUDA_HOME과 IREE_CUDA_COMPILE, IREE_MACOS_HOME과 IREE_MACOS_COMPILE, IREE_ROCM_HOME과 IREE_ROCM_COMPILE이 각각 짝이고, 이 트리는 scripts/iree/setup-cuda.sh, scripts/iree/setup-macos.sh, scripts/iree/setup-rocm.sh(또는 make iree-rocm)가 만듭니다. OpenXLA 백엔드 소스 상태를 참조하십시오.
MLX_BUILD_CUDA는 Cargo cuda feature가 선택하는 CMake 옵션이지 런타임 switch가 아닙니다. MLXCEL_MLX_COMMIT, MLXCEL_XLA_IREE_DIST는 build script가 생성하여 해당 crate에 bake하는 compile-time 값입니다. MLX_PIN_CMAKE_FILE은 MLX pin parser fixture test에서만 사용합니다.
ROCm 백엔드 변수¶
--features rocm 빌드의 ROCm backend(src/lib/mlxcelverse/overlays/rocm/mlx/backend/rocm/)가 읽는 변수입니다. 튜닝, A/B 비교, bisect용 switch이며 안정된 공개 API가 아니고, 배포 설정으로 권장하는 것도 없습니다. 행에 따로 적지 않았다면 해당 코드 경로가 처음 실행될 때 프로세스당 한 번 읽으므로 프로세스 시작 전에 설정하십시오. 아래 표는 운영자가 다룰 만한 일부만 담았고, backend가 읽는 나머지 kernel 선택 switch의 이름과 기본값은 ROCm overlay 소스를 따릅니다. 빌드 옵션과 현재 ROCm 지원 상태는 소스에서 빌드하기에 있습니다.
메모리와 디바이스¶
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLX_ROCM_MAX_INFLIGHT_MB | 0 이상의 정수 MiB, 그 밖의 값은 경고 후 기본값 유지 | 1024 | commit했지만 끝나지 않은 command batch들이 할당할 수 있는 메모리입니다. batch는 이 예산의 4분의 1을 할당하면 commit되고, 그보다 많은 양이 진행 중이면 host가 가장 오래된 batch를 기다리므로 prefill의 임시 버퍼가 한꺼번에 살아 있지 않습니다(Meta-Llama-3.1-8B-Instruct-4bit 최고 메모리가 20.60 GB에서 6.14 GB로 줄어듦). 0이면 상한을 끕니다. 메모리 추정기는 이 예산을 별도의 Backend in-flight 항(mlxcel inspect --json의 backend_inflight_bytes)으로 더합니다. |
MLX_ROCM_GPU_WATCHDOG_SECS | 0 이상의 정수 초, 그 밖의 값은 경고 후 off 유지 | 0(off) | host가 GPU 작업 하나를 기다리며 막혀 있을 수 있는 최대 시간이며, 넘으면 오류로 실패합니다. GPU fault는 이 값 없이도 약 1초 안에 보고되고, watchdog은 끝나지 않는 kernel을 위한 것이므로 서버가 받는 가장 긴 prefill보다 길게 잡아야 합니다. watchdog으로 실패해도 프로세스는 다시 시작해야 합니다. |
MLX_ROCM_FFT_CACHE_SIZE | 2147483647 이하의 양의 정수, 그 밖의 값은 경고 후 기본값 유지 | 128 | plan cache가 유지하는 hipFFT plan 수입니다. 값을 줄이면 메모리 대신 plan 재생성 비용을 치릅니다. |
MLX_ROCM_FORCE_WARP_SIZE | 32 또는 64, 그 밖의 값은 stderr 한 줄을 남기고 무시 | 디바이스의 wavefront 크기 | host 쪽 launch 폭을 덮어씁니다. 디버그 전용이며 값이 틀리면 출력이 깨집니다. |
MLX_HSACO_CACHE_DIR | 디렉터리 경로 | 시스템 임시 디렉터리 아래 mlx/<version>/hsaco | JIT 컴파일한 kernel의 디스크 cache 루트입니다. GPU 아키텍처마다 하위 디렉터리를 하나씩 둡니다. |
ROCM_HOME, ROCM_PATH, HIP_PATH | 디렉터리 경로, 이 순서로 확인하며 HIP_PATH는 Windows 전용, 빈 값은 미설정으로 취급 | Linux는 /opt/rocm이 있으면 그 경로, Windows는 ROCm 런타임 디렉터리 안이나 옆의 include, 그다음 빌드 시점에 기록한 SDK include(아직 있을 때) | 런타임 JIT(hiprtc)가 헤더를 가져오는 ROCm 설치 위치입니다. 아무것도 설정되지 않고 대체 경로도 없으면 첫 JIT 컴파일이 실패합니다. 실행 시 설정한 값이 대체 경로보다 우선하고 그 헤더가 GPU 커널로 컴파일되므로 신뢰하는 계정만 쓸 수 있는 디렉터리를 가리켜야 합니다. AMD HIP SDK 설치 프로그램은 HIP_PATH를 시스템 전체에 설정하는데 그러면 런타임 번들의 include보다 먼저 쓰이므로 mlxcel 실행 시 지우거나 번들 루트를 가리키게 하십시오. |
MLXCEL_ROCM_RUNTIME_DIR | 절대 디렉터리 경로(Windows 전용), 빈 값은 미설정으로 취급 | 미설정: 실행 파일 디렉터리에 HIP 런타임(ROCm 10은 amdhip64_7.dll)이 있으면 그 디렉터리, 없으면 빌드 시점에 기록한 SDK bin | 모든 ROCm DLL을 로드하는 디렉터리로, 보통 풀어 놓은 mlxcel ROCm 런타임 번들의 bin 폴더입니다. 설정되어 있으면 이것만 후보이므로 상대 경로이거나 HIP 런타임이 없는 디렉터리면 다른 후보로 넘어가지 않고 첫 ROCm 호출에서 메시지와 함께 멈춥니다. 고른 디렉터리에 없는 DLL도 다른 설치에서 가져오지 않습니다. 고른 디렉터리는 [mlx-rocm] ROCm runtime: <dir> (<source>)로 한 번 출력합니다. 디렉터리 안의 코드가 mlxcel 프로세스에서 실행되므로 신뢰하는 계정만 쓸 수 있어야 합니다. Windows ROCm 런타임 번들을 참조하십시오. |
양자화 matmul과 MoE¶
| 변수 | 값 | 기본값 | 설명 |
|---|---|---|---|
MLX_ROCM_GATHER_QMV_EXPERT_BATCHED | 0이면 off, 1이면 on, 그 밖의 값은 기본값 유지. 호출마다 읽음 | on | 정렬된 MoE prefill용 expert-batched gather_qmm kernel입니다(group 64의 bf16/f16 affine 4-bit 또는 8-bit, group 32의 bf16 mxfp4). 0이면 모든 방식을 행 단위 gather kernel로 보내므로, MoE 회귀를 A/B 비교하거나 bisect할 때 씁니다. |
MLX_ROCM_WMMA_QMM | 미설정, 0, 1 | 미설정(디바이스에 네이티브 WMMA가 있고 CU가 적은 iGPU가 아니면 fused WMMA kernel) | quantized_matmul이 두 행 이상인 bf16 affine GEMM을 실행하는 방식입니다. 0은 WMMA kernel을 쓰지 않고, 1은 shape가 맞으면 항상 쓰며 MLX_ROCM_WMMA_QMM_MAX_M을 무시합니다. |
MLX_ROCM_WMMA_QMM_MAX_M | 1부터 2147483647까지의 정수, 그 밖의 값은 경고 후 기본값 유지 | RDNA 3.5(gfx1150~gfx1152)에서 128, 그 외 상한 없음 | 이 행 수부터는 해당 GEMM을 WMMA kernel 대신 dequantize와 hipBLASLt로 보냅니다. |
MLX_ROCM_QMM_DEQUANT_CACHE_SIZE | 0부터 2147483647까지의 정수 | 0(off) | dequantize-plus-GEMM 경로가 재사용하려고 보관하는 dequantize된 가중치 수입니다. 이 경로에 오는 shape는 cache hit가 거의 없어 기본은 off이며, hit가 나는 workload는 양수 크기로 켤 수 있습니다. |
MLX_ROCM_QMM_DEQUANT_CACHE_MAX_BYTES | bytes | 268435456(256 MiB) | 그 cache를 켰을 때의 byte 상한입니다. |
MLX_NO_HIPBLASLT | 존재 여부 | hipBLASLt 사용 | 모든 GEMM을 hipBLASLt 대신 rocBLAS 경로로 보냅니다. hipBLASLt 수치 문제가 의심될 때 쓰는 진단용입니다. |
ROCm에서 달라지는 기본값¶
아래 변수의 전체 설명은 이 페이지의 다른 표에 있으며, ROCm 빌드에서는 기본값만 다릅니다.
| 변수 | ROCm 빌드의 기본값 | 그 외 |
|---|---|---|
MLXCEL_CACHE_LIMIT | 2GB | 미설정 |
MLXCEL_FUSED_ADD_RMSNORM | on | Metal과 CUDA에서 off |
MLXCEL_FUSED_ROPE_APPEND | on | Metal과 CUDA에서 off |
MLXCEL_FUSED_MOE_SGY | 2 | Metal과 CUDA에서 8 |
MLXCEL_FUSED_MOE_MAX_DFF | 8192(CUDA와 같은 값) | Metal에서 4096 |
MLX_MAX_OPS_PER_BUFFER | 2000 | MLX 런타임 변수 참조 |
하드웨어와 커널 진단¶
프로파일링, rollback, 실험용 변수입니다. 일반 배포 설정으로 의존하지 마십시오.
| 변수 | 값 | 기본값 | 용도 |
|---|---|---|---|
MLXCEL_METAL4_ATTENTION | 0/false/no/off면 비활성화, 미설정이거나 그 외 값이면 활성화 | 하드웨어가 지원하면 on | 고급, 진단용 킬 스위치. 해당 하드웨어에서 M5 뉴럴 액셀러레이터 fused attention 경로를 끄고 일반 SDPA 경로가 돌게 합니다. 끄는 쪽으로만 동작합니다. 하드웨어 검사가 거짓인 곳에서 켤 수는 없으므로 M1에서 설정해도 무효입니다. M5가 이전 Apple GPU 세대와 수치로 어긋날 때 가장 먼저 의심할 경로입니다. 프로세스당 한 번 읽고, Metal이 아닌 빌드에서는 무효입니다. |
MLXCEL_NO_PADDED_PREFILL | 존재하면 비활성화 | 자동 | CLI와 mlxcel-server에서 M5+/Neural Accelerator prefill tile 정렬을 끕니다. MLXCEL_FORCE_PADDED_PREFILL보다 우선합니다. |
MLXCEL_FORCE_PADDED_PREFILL | 존재하면 활성화 | off | 하드웨어와 관계없이 CLI와 서버에서 32토큰 prefill tile 정렬을 강제하므로, M5 Neural Accelerator 없이도 padded prefill과 그 cache trim을 테스트할 수 있습니다. M5가 아니면 순수한 오버헤드이며 검증 전용입니다. |
MLXCEL_FORCE_PADDED_PREFILL_MASK | presence enables | off | 명시적인 padded prefill mask 경로를 강제합니다. |
MLXCEL_LOG_NA_ATTENTION | sampled, all, truthy | off | Neural Accelerator attention dispatch 결정을 로그로 남깁니다. |
MLXCEL_ENABLE_FUSED_CAUSAL_PREFILL_ATTENTION | presence enables | off | 지원되는 경우 실험적 Llama-family fused causal prefill 경로를 활성화합니다. |
MLXCEL_TRACE_ARCH | presence enables | off | 프로세스당 한 번, device init 시점에 CUDA 아키텍처 상황을 stderr로 출력합니다. 실행 중인 GPU의 compute capability, 이 바이너리가 컴파일된 MLX_CUDA_ARCHITECTURES 목록, 그리고 해당 장치가 사전 컴파일된 cubin으로 서비스되는지 JIT 컴파일된 PTX로 서비스되는지를 보여줍니다. CUDA 빌드에서는 dispatcher가 첫 호출에서 고르는 quantized-matmul 경로(qmm_sm90, qmm_sm80, qmm_naive, qmv, fp_qmv)도 함께 보고하는데, MLX가 달리 노출하지 않는 정보입니다. 진단 전용이라 동기화도 추가 그래프 작업도 없으며, ROCm 빌드에서는 대신 장치의 gfx 타겟과 컴파일된 gfx 목록을 출력하고, Metal이나 CPU 전용 빌드에서는 compute capability를 확인할 수 없다고 보고합니다. 아키텍처 불일치 자체는 이 변수 없이도 보고됩니다. 컴파일된 아키텍처가 호스트 GPU를 포함하지 않는 바이너리는 양쪽을 모두 명시한 오류와 함께 시작을 거부하며, MLXCEL_DEVICE=cpu로 우회할 수 있습니다. |
MLXCEL_V4_DENSE | 존재하면 dense pooled attention 강제 | off (index_topk 이후 sparse) | DeepSeek-V4 진단용. index_topk 이후에도 모든 sparse 레이어를 선택 행 gather 대신 dense local+pooled concat에 묶어 두어, 그 레이어를 sparse 경로가 근사하는 full-attention 오라클로 만듭니다. index_topk 아래에서는 이미 그렇게 동작하므로 long-context 동작만 달라집니다. indexer의 compressor는 그대로 돌면서 pooled 캐시를 전진시키므로, 변수를 제거하면 올바르게 채워진 pool로 sparse 선택이 재개됩니다. 서빙 모드가 아닙니다. concat이 Np pooled 행 전부를 안고 가므로, sparse 경로가 피하려고 만들어진 바로 그 방식대로 메모리가 컨텍스트에 비례해 늘어납니다. |
MLXCEL_V4_FLAT_INDEX | 존재하면 flat scan 강제 | off (크기 게이트가 열리면 hierarchy) | DeepSeek-V4 진단용. sparse attention은 그대로 둔 채 pooled 행 선택만 flat O(Np) scan으로 강제해 2단 HiSA hierarchy를 끕니다. 두 경로가 같은 pooled 행을 같은 방식으로 점수화하므로 한 프롬프트를 양쪽으로 돌려 보면 "hierarchy가 잘못된 블록을 남겼다"와 "sparse combine이 틀렸다"를 갈라낼 수 있습니다. hierarchy는 index_block * index_keep(1024) pooled 행, 비율 4에서 대략 4096 프롬프트 토큰을 넘어야 개입하므로 달리 이분하기 어려운 지점입니다. 메모리를 씁니다. flat scan은 hierarchy가 피하려고 존재하는 [B, H, L, Np] 점수 텐서를 실체화합니다. |
MLXCEL_ENABLE_FUSED_QKV_SPLIT_ROPE | presence enables | off | 실험적 fused QKV projection/split/RoPE 경로를 활성화합니다. |
MLXCEL_GEMMA4_ENABLE_FUSED_QKV | presence enables | off | Gemma 4 fused-QKV projection 실험을 활성화합니다. |
MLXCEL_GEMMA4_NO_KV_SHARED_TAIL_TRIM | 1이면 비활성 | 트림 켬 | 첫 KV-shared 레이어에서 히든 상태, 레이어별 입력, 마스크를 샘플링된 행으로 자르지 않고 KV-shared 꼬리 구간의 프리필을 전체 쿼리 폭으로 유지합니다. 트림 경로는 전체 폭 경로와 수치적으로 동등하므로, 이 변수는 한 바이너리에서 teacher-forced 추적의 두 갈래를 실행하려고 둡니다. *_last_* 프리필 항목만 트림하며 num_kv_shared_layers > 0일 때만 적용됩니다. |
MLXCEL_CUDA_DISABLE_GEMV | 1이면 활성, 미설정이나 다른 값은 비활성 | off | CUDA 전용. MLX가 자체 gemv 커널로 보내던 dense matmul(M == 1이면서 B가 전치된 경우, 또는 N == 1이고 바이어스 없음)을 cuBLASLt로 보내는 진단 스위치입니다. bf16/f16 모델의 B=1 디코드가 바뀌고, 양자화 matmul과 gather_mm은 각자의 커널을 유지합니다. 프로세스마다 한 번 읽으며, 튜닝이 아니라 gemv와 cuBLASLt 비교 측정용입니다. |
MLXCEL_DSA_DENSE | presence enables | off | A/B 진단을 위해 DeepSeek-V3.2/GLM-MoE-DSA가 DSA indexer를 건너뛰고 dense full attention을 사용하게 합니다. |
MLXCEL_MINIMAX_M3_DENSE | presence enables | off | A/B 진단을 위해 MiniMax-M3가 block-sparse indexer를 건너뛰고 dense attention을 사용하게 합니다. |
MLXCEL_FUSED_QK_NORM | 1/true/on/yes는 enable, 그 외 값이나 미설정은 disable | off(opt-in) | Qwen3와 Qwen3-MoE decode용 fused 단일 토큰 QKV projection + Q/K RMSNorm + RoPE kernel입니다. graph 경로와 RMS < 5e-3 이내로 일치하지만 지금까지 측정한 어떤 backend(M1 Ultra, M5 Max, GB10/CUDA)에서도 decode 속도 향상이 확인되지 않아 opt-in으로 남아 있습니다. |
MLXCEL_FUSED_XIELU | 0/false/off/no면 비활성화, 그 밖의 값 또는 미설정이면 활성화 | on | Apertus MLP activation용 단일 실행 fused xIELU kernel(Metal, ROCm)입니다. Apple Silicon에서 greedy temp-0 decode가 elementwise 경로와 byte 단위로 같고 M1 Ultra와 M5 Max에서 더 빠르게 측정됐으며, ROCm port도 gfx1151에서 elementwise 경로와 byte 단위로 같습니다. CUDA와 CPU에서는 FFI가 동등한 elementwise graph로 폴백합니다. Apertus 전용입니다. elementwise 경로를 강제하려면 0으로 설정하십시오. |
MLXCEL_DISABLE_COMPILED_SWITCH_QGEGLU | presence disables | 지원 시 compiled path | Gemma 4 compiled Switch-QGeGLU decode 경로를 rollback합니다. |
MLXCEL_ENABLE_SOFTCAP_GQA_DECODE_GROUPED | 0이 아닌 값은 활성화, 0은 비활성화 | on | grouped softcap-GQA decode. q_len == 1이고 mask가 없으면 K와 V를 [B, H_kv, S, D]로 두고 matmul 안에서 n_rep을 broadcast합니다. fallback은 먼저 확장하므로 decode 단계마다 살아 있는 cache 전체의 n_rep 배 복사본을 씁니다. 이 경로에 도달하는 것은 softcap 계열뿐이라 실제로 달라지는 것은 Gemma 2(attn_logit_softcapping)입니다. greedy 출력은 어느 쪽이든 바이트 단위로 같습니다. |
MLXCEL_DISABLE_SOFTCAP_GQA_DECODE_GROUPED | 1 disables, 0 enables | 미설정 | grouped softcap-GQA decode의 legacy rollback/override입니다. |
MLXCEL_DISABLE_SINGLE_QUERY_MASKLESS | truthy disables | maskless path on | single-query maskless attention 경로를 비활성화합니다. |
MLXCEL_EXPERIMENTAL_BOOL_CAUSAL_MASK | truthy enables | off | 실험적 boolean causal-mask 경로를 활성화합니다. |
MLXCEL_PIPELINE_GRANULARITY | off, layer, block:N | off | pipeline 실험용 layer-boundary async-eval hint를 삽입합니다. |
MLXCEL_AUTOTUNE | 1/true/on/yes면 튜닝, cache/read/readonly/read-only면 캐시된 tactic만 사용, 0/false/off/no나 미설정이면 비활성화 | off | shape 버킷 단위 커널 autotuner. 튜닝 모드는 버킷마다 후보 tactic을 측정해 승자를 캐시하고, 읽기 전용 모드는 다른 곳에서 만든 캐시를 측정 없이 적용합니다. |
MLXCEL_BENCH_LLC_BYTES | 양의 정수 (바이트) | 자동 추정값 | cold-L2 벤치마크 모드가 입력 rotation 크기를 정할 때 쓰는 last-level cache 크기를 덮어씁니다. |
MLXCEL_FUSED_MOE | 0/false/off/no이면 비활성화, 그 밖의 값 또는 미설정이면 활성화 | on | Metal, CUDA, ROCm(HIP port, wave32 gfx1151에서만 검증)의 fused single-token decode-MoE kernel입니다. MLXCEL_FUSED_MOE_MAX_DFF보다 넓은 expert는 gather_qmm에 남습니다. fused 경로는 모든 checkpoint에서 byte 단위로 같지는 않으므로, 새 MoE port를 reference 구현과 비교할 때는 끄십시오. |
MLXCEL_FUSED_ADD_NORM | 0/false/off/no이면 비활성화, 그 밖의 값 또는 미설정이면 활성화 | on | decode 중 parallel-residual block의 (attn + mlp) + x와 이를 소비하는 LayerNorm을 하나의 kernel(Metal, ROCm)로 합치며, 결과는 각 backend의 fused하지 않은 쌍과 byte 단위로 같습니다. CUDA와 CPU, 6656보다 넓은 row, 혼합 dtype에서는 fused하지 않은 쌍으로 폴백합니다. Cohere2에서 사용합니다. |
MLXCEL_DECODE_MB_PER_BUFFER | 0/off/false/no면 비활성화, 양의 정수 | M5 이전 Apple Silicon에서 1000, 그 밖에는 꺼짐 | 파이프라인 decode 단계(generate 반복문과 서버 lookahead decode)에만 적용하는 Metal command-buffer 입력 예산이며, prefill과 동기 단계는 장치 기본값을 유지합니다. decode에서는 토큰마다 전체 가중치를 읽으므로 기본 예산에서는 한두 계층마다 buffer를 commit하고 그 경계마다 GPU가 쉽니다. M1 Ultra에서 측정한 decode 이득은 모델에 따라 변화 없음부터 약 20%까지입니다. 같은 값을 prefill에 쓰면 긴 프롬프트에서 peak memory가 대략 두 배가 되므로 prefill은 기본값을 유지합니다. MLX_MAX_MB_PER_BUFFER를 설정하면 무시합니다. |
MLXCEL_DEBUG_KERNEL_BACKEND | presence enables | off | 진단. mlxcel 커스텀 커널이 선택한 GPU 백엔드(metal, cuda, rocm, none)와 그 백엔드에 커스텀 커널 포트가 있는지, MLX 그래프 대체 경로를 쓰는지를 한 번 출력합니다. |
MLXCEL_SDPA_FALLBACK_MAX_QUERIES | 음이 아닌 정수 | 32 | CUDA 전용. 더 긴 key 시퀀스 위에서 query 행이 2에서 N개인 masked SDPA 호출(투기적 검증 형태)은 cuDNN을 건너뛰고 MLX ops 대체 경로를 탑니다. cuDNN은 key 길이가 바뀔 때마다 execution plan을 다시 만들기 때문입니다. 0은 상류 dispatch로 되돌립니다. |
MLXCEL_SDPA_PLAN_BUCKET_MAX_QUERIES | 음이 아닌 정수 | 0 (꺼짐) | CUDA 전용, opt-in. query 행이 2에서 N개인 array-masked 검증 호출의 cuDNN plan-cache key를 bucket으로 묶어 plan 하나가 여러 라운드를 처리하게 합니다. 보통 길이의 문맥에서는 대부분의 폭에서 ops 대체 경로보다 느리게 측정되어 꺼져 있으며, 32로 켭니다. |
MLXCEL_SDPA_PLAN_BUCKET_MAX_MB | 음이 아닌 정수 (MiB) | 64 | CUDA 전용. K와 V가 캐시 버퍼 하나의 slice가 아닐 때 bucket 경로가 만드는 zero-padded K/V 복사본의 텐서별 상한. 0은 복사 경로를 끕니다. |
MLXCEL_SDPA_PLAN_DEBUG | 1이면 활성화 | off | CUDA 전용, 진단. cuDNN SDPA 호출마다 plan 재사용을 결정하는 필드와 plan을 새로 만들었는지 재사용했는지를 한 줄씩 기록합니다. 계층별 호출마다 출력하므로 켜 둔 채로 두지 마십시오. |
MLXCEL_FUSED_MOE_SGY | 1-32 | 8(Metal, CUDA) / 2(ROCm) | fused decode-MoE kernel의 threadgroup당 simdgroup(Metal), warp(CUDA), wavefront(ROCm) 수입니다. 출력은 이 값에 의존하지 않습니다. 하드웨어별로 조정하십시오. |
MLXCEL_FUSED_MOE_MAX_DFF | 양의 정수 | 4096(Metal) / 8192(CUDA, ROCm) | fused 경로의 expert-intermediate(Dff) 상한입니다. 이를 넘으면 caller가 gather_qmm으로 폴백합니다. |
MLXCEL_FUSED_MOE_RELU2 | 설정하면 활성화 | off | nemotron 계열 expert를 위한 squared-ReLU fused routed-expert 경로. |
MLXCEL_FUSED_MOE_PARITY_CHECK | 미설정이거나 0이면 비활성화, 0이 아닌 정수면 활성화 | off | 진단용. fused decode-MoE 커널의 호출별 parity/결정성 probe. 체크포인트의 fused 출력이 레퍼런스를 따라가는지 확인할 때 씁니다. |
MLXCEL_FUSED_MOE_PARITY_THRESHOLD | 실수 | 0.05 | MLXCEL_FUSED_MOE_PARITY_CHECK probe의 정규화 RMS 편차 임계값. |
MLXCEL_INKLING_MOE_KERNEL | 0이면 비활성화, 그 밖의 값 또는 미설정이면 활성화 | on(Metal) | port가 있는 backend에서 fused Inkling MoE decode kernel을 켭니다. route kernel은 gate matmul 이후 routing chain(sigmoid, correction bias, top-k, logsigmoid softmax, scaling, shared-expert gamma broadcast)을 대체하고, down-combine kernel은 routed affine q4/g64 down projection을 top-k 가중 expert 합과 합쳐 배치 8 이하를 처리합니다. port가 없는 backend는 graph 경로를 타므로 0은 rollback일 뿐입니다. |
MLXCEL_FUSED_ADD_RMSNORM | 0/false/off/no이면 비활성화, 1/true/on/yes이면 활성화 | off(Metal, CUDA) / on(ROCm) | residual-add와 RMSNorm을 합친 decode kernel로, hidden state를 두 번 읽던 것을 한 번으로 줄입니다. |
MLXCEL_FUSED_ROPE_APPEND | 0/false/off/no이면 비활성화, 1/true/on/yes이면 활성화 | off(Metal, CUDA) / on(ROCm) | q/k RoPE와 KV-append-layout을 합친 decode kernel입니다. |
MLXCEL_SPARSE_PAGED_ATTENTION | 0/false/off/no면 비활성화, 그 외 값이나 미설정이면 활성화 | on | page indirection을 이용한 fused sparse 디코드. 요청의 live 페이지가 할당량의 일부에 불과할 때 커널이 선택된 행만 gather합니다. |
MLXCEL_SPARSE_PAGED_MIN_SPARSITY | 음이 아닌 정수 | 8 | fused sparse 디코드를 dispatch하기 전에 launch가 넘어야 하는 최소 live_len / selected_rows 비율. |
MLXCEL_SPARSE_PAGED_DUMP | 1/true/on/yes면 활성화 | off | 진단용. fused sparse 디코드에서 요청별로 선택된 pool 행을 출력합니다. 디바이스와 동기화하므로 운영 스위치가 아니라 디버깅 도구입니다. |
MLXCEL_PAGED_DECODE_SPLITS | 양의 정수 | 미설정 (메모리 예산 상한) | fused paged-attention 디코드 커널의 NumSplits 토큰 분할 launch 형태를 메모리 예산에서 유도하지 않고 고정합니다. |
MLXCEL_PAGED_ATTENTION_NATIVE | 1/true/on/yes는 native 강제, 0/false/off/no는 gather 강제, 미설정은 adaptive 선택에 맡김 | adaptive | A/B 측정을 위해 paged-attention decode dispatch의 한쪽 경로를 강제합니다. fused kernel은 Metal, CUDA, ROCm에서 실행되며 ROCm은 CUDA 선택 규칙을 따릅니다. |
MLXCEL_PAGED_SLAB_BLOCKS | 양의 정수(블록), 값 그대로 사용. 0은 풀 기본값 고정 | --ctx-size와 --parallel에서 유도 | paged 풀의 slab 크기(블록 단위)입니다. fused 커널은 한쪽당 연속 버퍼 하나를 읽으므로, 행이 첫 slab을 넘어선 레이어는 gather 경로로 되돌아갑니다. 서버는 확정된 값을 시작 시 Paged KV slab size: N blocks per layer로 기록합니다. --ctx-size가 함의하는 것보다 긴 컨텍스트를 서비스한다면 값을 올리십시오. 0은 과거의 32블록 풀 기본값을 고정하며, 그 부수 효과로 slab 하나를 넘어서면 fused 경로에 도달할 수 없게 됩니다. 음이 아닌 정수가 아닌 값은 경고와 함께 무시되므로, 변수를 설정하지 않은 것과 똑같이 유도값이 적용됩니다. |
MLXCEL_PAGED_V2_MIN_KV_TOKENS | 음이 아닌 정수(토큰) | 4096 | 단일 요청 launch가 paged decode v2로 dispatch되기 위해 넘어야 할 가시 KV 토큰 수입니다. 이 바닥값은 유일하게 측정된 손실(배치 1, 1024 토큰)을 제외합니다. 파싱 실패 시 기본값으로 되돌아갑니다. |
MLXCEL_PAGED_V2_MIN_KV_TOKENS_PER_REQUEST | 음이 아닌 정수(토큰) | 512 | 다중 요청 launch에서 요청당 넘어야 할 가시 KV 토큰 수입니다. 측정된 최소 이득보다 낮게 잡아, 명목 1K를 살짝 밑도는 프롬프트가 근거 없이 거절되지 않게 합니다. |
MLXCEL_PAGED_ATTENTION_V2 | 1/true/on/yes면 활성화, 미설정이거나 그 외 값이면 비활성화 | off | 진단용. 라이브러리 전용 v1 진입점이 fused v2 커널을 먼저 시도하게 합니다. 서버의 프로덕션 v2 디코드는 제어하지 않으며, 그쪽은 위의 토큰 바닥값이 관장합니다. |
MLXCEL_PAGED_DECODE_V2_CHUNK | 양의 정수(페이지) | 미설정 | 진단용. 모든 paged decode v2 launch의 pages_per_chunk를 고정해 plan의 점유율 휴리스틱과 autotune 값을 모두 덮어씁니다. plan이 값을 가능 범위로 clamp합니다. |
MLXCEL_PAGED_DECODE_V2_TARGET_CTAS | 양의 정수 | 유도값(Apple: gpu_core_count * 8, 최소 64, 그 외: 512) | 진단용. v2 plan이 chunk 크기를 이진 탐색할 때 겨누는 점유율 목표를 옮깁니다. 유도된 디바이스 병렬도가 틀린 호스트에서 씁니다. |
MLXCEL_QMV_MULTIROW | 0이면 비활성화, 그 외 값이나 미설정이면 활성화 | on | CUDA 전용. weight를 분할 상환하는 multirow qmv 경로. broadcast weight와 작은 행 창에서 weight를 한 번 읽어 여러 행에 씁니다. |
MLXCEL_QMV_MULTIROW_MAX_ROWS | 1-8 | 8 | CUDA 전용. multirow qmv 경로의 행 창 상한으로, dispatch 창을 좁힙니다. |
MLXCEL_QMM_TILE_M | 양의 정수 | 자동 | CUDA 전용. quantized matmul 커널의 M 타일을 shape별 선택 대신 고정합니다. |
MLXCEL_QMM_TILE_N, MLXCEL_QMM_TILE_K | 양의 정수 | 128, max(64, group_size) | CUDA 전용, 실험적. quantized-matmul CTA의 나머지 tile axis를 고정합니다. 잘못된 layout은 NVRTC compile을 실패시킬 수 있으므로 kernel tuning에만 사용하십시오. |
MLXCEL_QMM_FOLD_BATCH | 0이면 비활성화, 그 밖의 값 또는 미설정이면 활성화 | on | CUDA 전용. 하나의 2D 양자화 가중치를 공유하는 배치를 QuantizedMatmul::eval_gpu에서 행으로 fold하므로, batched decode가 배치 항목마다가 아니라 가중치를 한 번 읽습니다. MLXCEL_QMM_FOLD_MIN_ROWS 미만이면 같은 qmv kernel이 fold된 행 위에서 돌고, 이상이면 GEMM 하나가 가중치를 한 번 읽습니다. 첫 사용 시 한 번 읽으며 0이면 배치별 dispatch로 되돌립니다. |
MLXCEL_QMM_FOLD_MIN_ROWS | 2부터 8까지의 정수, 그 밖의 값은 기본값 유지 | compute capability 12.1에서 5, 그 밖에서 8 | CUDA 전용. 큰 가중치(MLXCEL_QMM_FOLD_MIN_ELEMENTS 참조)의 fold된 M == 1 decode 배치가 qmv를 떠나 GEMM kernel로 가는 행 수입니다. 이 폭부터는 배치된 행이 혼자 디코드한 같은 행과 bit 단위로 같지 않습니다. verify block과 prefill 조각(M > 1)은 항상 8을 유지합니다. 첫 사용 시 한 번 읽습니다. |
MLXCEL_QMM_FOLD_MIN_ELEMENTS | 양의 정수 | 16777216 (2^24) | CUDA 전용. MLXCEL_QMM_FOLD_MIN_ROWS가 적용되는 가중치 크기 N * K입니다. 더 작은 가중치는 compute capability 12.1에서 8행까지, 그 밖에서 7행까지 qmv를 유지합니다. 첫 사용 시 한 번 읽습니다. |
MLXCEL_QMM_NAIVE_TILE_N | 64 또는 128, 그 밖의 값이나 미설정이면 shared-memory 예산에 선택을 맡김 | 미설정 | CUDA 전용, 진단용. qmm_naive kernel의 N tile을 고정합니다. dispatch마다 읽고 캐시하지 않으므로 한 프로세스에서 두 폭을 모두 돌려 bit 단위로 비교할 수 있으며, 폭마다 자체 module을 컴파일합니다. |
MLXCEL_TRACE_QMM_TILE | 설정하면 활성화 | off | CUDA 전용, 진단용. 선택된 qmm_naive tile을 컴파일된 kernel의 레지스터 수, spill, 달성 점유율과 함께 kernel 종류마다 한 번 stderr로 출력합니다. |
MLX_DEBUG_KERNELS | 1이면 활성화 | off | 진단용. custom Turbo kernel launch를 포함한 MLX runtime-JIT kernel source dump를 요청합니다. |
MLXCEL_GATHER_QMM_GROUPED | 0이면 비활성화, 그 외 값이나 미설정이면 활성화 | on | CUDA 전용. GatherQMM::eval_gpu의 정렬 MoE prefill 빠른 경로. |
MLXCEL_GATHER_QMM_GROUPED_MIN_ROWS | 양의 정수 | 8 | CUDA 전용. MLXCEL_GATHER_QMM_GROUPED의 분할 상환 임계값으로, 정렬 배치가 이 행 수에 도달해야 빠른 경로가 켜집니다. |
MLXCEL_SSM_KERNEL | 0이면 비활성화, 그 밖의 값 또는 미설정이면 활성화 | on | hybrid SSM 모델(granite-4.0-h, falcon-h1, plamo-2, nemotron-h)용 fused single-token SSM decode kernel(ssm_update_kernel)로, port가 있는 모든 backend(Metal, CUDA, ROCm)에서 쓰입니다. port가 없는 backend나 기본 디바이스가 CPU인 경우는 graph 경로를 탑니다. graph 경로를 강제하려면 0으로 설정하십시오. |
MLXCEL_SSM_CUDA_KERNEL | 0이면 비활성화, 그 밖의 값 또는 미설정이면 활성화 | on | MLXCEL_SSM_KERNEL의 alias입니다. 둘 중 하나라도 0이면 모든 backend에서 graph 경로를 강제합니다. |
MLXCEL_MAMBA1_SCAN_KERNEL | 0이면 비활성화, 그 밖의 값 또는 미설정이면 기본값 유지 | on(Metal, ROCm, CUDA에서는 Jamba) | Jamba, Mamba, Falcon-Mamba의 Mamba1 selective scan용 fused kernel로, prefill과 decode 모두에 쓰입니다. Metal과 ROCm kernel은 state를 float32로 유지하므로 결과가 graph scan과 조금 다르고, CUDA에서 Jamba가 쓰는 port는 graph scan과 같은 방식으로 반올림해 bit 단위로 같습니다. M1 Ultra(Jamba reasoning 3B 4-bit) 측정에서 prefill은 +45%~+54%, decode는 약 +10%였습니다. graph scan을 강제하려면 0으로 설정하십시오. 호출마다 확인합니다. |
MLXCEL_ONE_BIT_KERNEL | 0/false/off/no이면 비활성화, 미설정이나 그 밖의 값은 kernel 유지 | on(Metal) | affine 1-bit checkpoint(bits: 1)를 mlxcel 자체 fused Metal qmv(16행 미만)와 tiled qmm kernel로 실행합니다. 0이면 dequantize graph를 강제하는데, 이 graph는 parity 기준이자 CUDA와 ROCm이 항상 쓰는 경로입니다. 프로세스당 한 번 읽습니다. |
MLXCEL_SDPA_DETERMINISTIC | 1(0이 아닌 정수)이면 활성화, true 같은 단어는 off로 읽음 | off(0) | CUDA 전용. MLX의 single-row decode SDPA를 sdpa_vector kernel로 보내고 stream-K cuDNN plan을 배제해 bit 단위로 재현되게 하며, decode 비용은 0~13%입니다. mlxcel-engine-parity --expect-identical에 필요합니다. 서버의 paged decode kernel에는 영향이 없습니다. |
MLXCEL_SDPA_VECTOR_LARGE_D | 0/false/off/no면 비활성화, 그 외 값이나 미설정이면 활성화 | on | CUDA 전용. CUDA supports_sdpa_vector 검사가 큰 head dimension을 받아들이도록 합니다. |
MLXCEL_COMPILED_QGELU_MLP | 0/false/off/no면 비활성화, 그 외 값이나 미설정이면 활성화 | on | affine quantized GeGLU MLP(gate, up, gelu, down)를 mx::compile 그래프 하나로 컴파일합니다. |
MLXCEL_SHAPELESS_COMPILE_AUDIT | 0이 아닌 임의 값이면 활성화 | off | 감사 모드이며 서빙 설정이 아닙니다. C++ 브리지의 모든 shapeless mlx::core::compile 지점을 eager oracle로 감싸, 호출마다 컴파일하지 않은 함수도 실행하고 지점별 호출, 입력 signature, 불일치를 shapeless_compile_audit_report용으로 기록합니다. 첫 compiled 함수가 만들어지기 전에 설정해야 합니다. scripts/audit_shapeless_compile.sh가 구동합니다. |
MLXCEL_SAMPLING_GUMBEL | 0/false/off/no이면 비활성화, 그 밖의 값 또는 미설정이면 활성화 | on | Metal, CUDA, ROCm의 softmax-free Gumbel-max categorical sampling kernel입니다. 명시적 softmax 없이 logits에서 바로 샘플링합니다. |
MLXCEL_SAMPLING_REJECTION | 0/false/off/no이면 비활성화, 그 밖의 값 또는 미설정이면 활성화 | on | Metal, CUDA, ROCm에서 dual-pivot rejection으로 정렬 없이 top-k, top-p, min-p sampling을 수행해 전체 vocabulary 정렬을 피합니다. top-k와 top-p를 함께 쓰면 vocabulary 32768(ROCm은 152064)까지 이 kernel로 보냅니다. |
MLXCEL_NVFP4_DENSE_REPACK | 1/true/on/yes이면 dense fallback 강제 | off(direct transcode) | 직접 ModelOpt-triplet NVFP4 transcode 대신 dense fallback을 선택합니다. 네이티브 NVFP4 kernel이 없는 ROCm에서는 dense affine 4-bit 변환이 이미 기본이므로 이 변수는 아무것도 바꾸지 않습니다. |
MLXCEL_NVFP4_NATIVE_REPACK | 1/true/on/yes이면 dense fallback 내부에서 native NVFP4 선택 | 기본 native | dense NVFP4 fallback 경로의 호환성 selector입니다. ROCm에서는 ModelOpt NVFP4를 항상 affine 4-bit로 변환하므로 무시되며, 로드 로그에 그렇게 표시됩니다. |
MLXCEL_DISABLE_FUSED_GLOBAL_SCALE | 1/true/on/yes면 비활성화 | off (fold 켜짐) | NVFP4 global-scale fold를 끄고 scale을 별도 단계로 적용합니다. |
MLXCEL_CUDA_F16_NORMALIZE | 1/true/on/yes면 활성화, 0/false/off/no면 비활성화 | sm80 미만에서 on, sm80 이상에서 off (opt-in) | CUDA 전용. 단일 dtype decode 그래프를 위해 weight를 로드 시점에 bf16에서 f16으로 정규화합니다. 하드웨어가 갈리므로 기본값도 갈립니다. Ampere 이상은 bf16 ALU가 있어 f16으로 얻는 것이 없으므로 bf16을 유지합니다. Volta와 Turing(sm70, sm75)은 bf16 ALU도, bf16 tensor-core MMA atom도, cuBLAS bf16 GEMM도 없어 모든 bf16 피연산자가 실행 전에 변환되므로, 그 변환을 로드 시점으로 옮기고 양자화 체크포인트까지 대상에 넣습니다. 양자화 체크포인트는 packed plane이 u32로 남고 bf16 side-data만 움직이기 때문입니다. pre-Ampere에서 bf16을 유지하려면 0으로 설정하십시오. |
MLXCEL_CUDA_F16_FRAGILE | 1/true/on/yes면 활성화, 미설정이면 제외 유지 | off (opt-in) | sm_80 미만 CUDA 전용. f16-fragile 예외 목록에 오른 계열을 그래도 변환합니다. sm_80 미만에서 그 목록에 남은 것은 apertus 하나이며, xIELU가 입력을 제곱하는 탓에 f16에서는 채점한 모든 위치가 NaN이 됩니다. 활성화 값이 다르다고 볼 근거가 있는 체크포인트이거나, 목록이 아직 모르는 계열을 직접 측정할 때만 쓰십시오. |
MLXCEL_CUDA_RUNTIME_DIR | 절대 디렉터리 경로(Windows CUDA 빌드 전용), 빈 값은 미설정으로 취급 | 미설정: 실행 파일 디렉터리에 CUDA 런타임 DLL이 있으면 그 디렉터리, 없으면 빌드 시점에 기록한 toolchain과 cuDNN bin\<arch> | 모든 CUDA, cuDNN, NVRTC DLL을 로드하는 디렉터리로, 보통 풀어 놓은 mlxcel CUDA 런타임 번들의 bin 폴더입니다. 설정되어 있으면 이것만 후보이므로 상대 경로이거나 런타임 DLL이 없는 디렉터리면 폴백하지 않고 첫 CUDA 호출에서 메시지와 함께 멈춥니다. 고른 디렉터리는 [mlx-cuda] CUDA runtime: <dir> (<source>)로 한 번 출력합니다. 디렉터리 안의 코드가 mlxcel 프로세스에서 실행되므로 신뢰하는 계정만 쓸 수 있게 두십시오. |
MLXCEL_CUDA_PINNED_LOAD | 0/false/off면 비활성화, 그 밖의 값이나 미설정이면 기본값 유지 | on | CUDA와 ROCm(ROCm에서는 MLXCEL_ROCM_PINNED_LOAD, MLXCEL_CUDA_PINNED_LOAD도 fallback으로 인정). 체크포인트 텐서를 텐서마다 버퍼를 잡는 대신 재사용하는 64 MiB pinned host staging slot 네 개로 업로드하므로, 한 slot으로의 읽기가 이전 slot의 복사와 겹칩니다. 256 MiB pinned 메모리는 첫 업로드에서 할당해 프로세스가 유지하며, 할당에 실패하면 경고 한 번을 남기고 pageable 경로를 씁니다. 0이면 pageable 경로를 강제하며, 업로드되는 바이트는 어느 쪽이든 같습니다. |
MLXCEL_ROCM_PINNED_LOAD | 0/false/off면 비활성화, 그 밖의 값이나 미설정이면 기본값 유지 | on | MLXCEL_CUDA_PINNED_LOAD의 ROCm 표기이며, MLXCEL_CUDA_PINNED_LOAD도 fallback으로 인정합니다. pinned host staging 동작과 pageable fallback은 같습니다. |
MLXCEL_CUDA_GRAPH_CAPTURE_BARRIER | 0/false/off/no면 비활성화, 1/true/on/yes, 미설정 또는 그 밖의 값이면 활성화 | on | CUDA 전용. MLX CUDA graph 안에 capture된 모든 child graph(cuBLASLt matmul, cuDNN convolution 또는 attention, 기타 stream capture)를 barrier로 만들어, 앞의 열린 branch를 모두 합류시키고 뒤에 추가되는 모든 node를 gate합니다. MLX가 buffer에서 유도한 데이터 의존성에 더해 적용됩니다. upstream은 선언한 buffer로만 child graph의 순서를 정하는데, GB10에서는 관련 없는 kernel이 겹치면서 신규 프로세스의 약 40%에서 cuBLASLt f32 GEMM이 깨졌습니다. upstream 순서와 A/B로 비교하려면 0으로 설정하십시오. |
MLXCEL_CUDA_GRAPHS_DOT_VERBOSE | 1/true/on/yes면 활성화, 0/false/off/no, 미설정 또는 그 밖의 값이면 off 유지 | off | CUDA 전용. MLX의 MLX_SAVE_CUDA_GRAPHS_DOT_FILE이 설정된 상태에서 1이면 각 graph를 cudaGraphDebugDotFlagsVerbose로 dump해 kernel과 memset node 파라미터(grid, shared memory, memset 대상)가 node 종류 옆에 나타납니다. 동시 실행 node가 어떤 buffer를 건드리는지 보여주는 정보입니다. |
MLXCEL_QMV_ROWS_PER_WARP | 1, 2, 4. 그 외 값은 아키텍처별 기본값 | sm_70에서 2, 그 외 1 | CUDA 전용. 고전적인 M*B == 1 decode에서 warp 하나가 담당하는 qmv 출력 행 수입니다. 기본 커널은 warp 하나에 출력 행 하나를 주고 그 warp가 activation 벡터 전체를 읽게 하는데, 이것이 Ampere 미만의 decode가 weight 대역폭에 묶이지 않는 이유입니다. warp당 R행은 activation 타일을 한 번 읽어 R개 weight 행에 모두 적용하므로 activation 트래픽이 R분의 1이 되고 weight 트래픽은 그대로입니다. temperature 0에서 출력은 R = 1과 바이트 단위로 같습니다. 1은 기본 커널로 가는 rollback입니다. 프로세스당 한 번 읽습니다. |
MLXCEL_QMV_MIN_BLOCKS | 1, 3, 4. 그 외 값은 아키텍처별 기본값 | sm_70에서 3, 그 외 1 | CUDA 전용. 행 단위로 블록화한 qmv 커널의 __launch_bounds__에 넘기는 resident block 하한이며, ptxas에 레지스터 상한을 지시하는 수단입니다. 256스레드 블록에서 3은 스레드당 85개, 4는 64개로 제한합니다. 예산은 느슨하게도 빡빡하게도 잘못 잡을 수 있어, 행 수가 아니라 이 값이 튜닝을 담당합니다. 1은 제약 없는 경우이자 rollback입니다. 프로세스당 한 번 읽습니다. |
MLXCEL_KEEP_BF16 | 아무 값이나 설정하면 활성화 | off | 측정 도구이며 지원되는 구성이 아닙니다. 모든 백엔드의 모든 체크포인트에서 로드 시점 bf16에서 f16으로의 변환을 건너뛰어, bf16 모델이 전부 bf16으로 로드되게 합니다. dtype 정책을 그것이 대체한 혼합 구성과 A/B로 비교하려고 존재합니다. bf16이 f16보다 낫다는 증거가 아닙니다. 강제 bf16 부분집합이 없는 계열에서는 같은 하드웨어에서 두 dtype이 오차 범위 안에 들어옵니다. 비용을 만드는 것은 한 체크포인트가 두 dtype을 함께 갖는 상태입니다. MLX가 f16과 bf16의 경계마다 f32로 승격하기 때문입니다. 로드마다 읽습니다. |
MLXCEL_BAILING_LINEAR_CHUNKED_PREFILL | 0/false/off/no면 비활성화, 그 외 값이나 미설정이면 활성화 | on | bailing_moe_linear(Ling / Ring) gated linear attention의 청크 closed-form prefill로, 64토큰 청크 단위로 평가합니다. |
MLXCEL_PROFILE_NEMOTRON_MOE | 설정하면 활성화 | off | 진단용. fused_moe_forward 안의 router, routed expert, combine, shared expert 경계에서 강제로 평가한 뒤 경계마다 타이밍 한 줄을 출력합니다. |
로깅, 프로파일링, 캡처¶
대부분의 스위치는 동기화 또는 추가 graph 작업을 강제하므로 throughput 측정값을 바꿀 수 있습니다.
| 변수 | 값 | 기본값 | 용도 |
|---|---|---|---|
MLXCEL_TRACE_DTYPE | 존재하면 활성화 | off | fused Mamba2 forward 안의 중간 dtype(input projection 뒤, SSM update 뒤)을 출력하며, 출력할 때마다 graph를 평가합니다. |
MLXCEL_FORCE_SYNC | 존재하면 활성화 | off | 모든 디코드 단계를 동기식으로 실행합니다. 서버 BatchScheduler의 lookahead 디코드 파이프라인과, mlxcel generate(--prompt-lookup의 일반 라운드 포함), 추론 세션, mlxcel-bench-decode, parity harness arm이 쓰는 raw-completion client의 B=1 파이프라인을 끕니다. greedy 출력은 어느 쪽이든 같습니다. |
MLXCEL_PROFILE_TTFT | 존재하면 활성화 | off | mlxcel generate --profile 생성마다, 그리고 mlxcel-bench-decode 측정 pass마다 [TTFT] 한 줄을 출력해 첫 토큰 이전 구간을 setup, build(lazy graph 구성, host 전용), sample, eval(모든 가중치 materialize, 모든 JIT 모듈 로드, 모든 CUDA graph 인스턴스화까지 포함하는 단 한 번의 blocking 평가), post로 나누고, 보고된 prefill 시간과 귀속되지 않은 잔여 시간을 함께 보여줍니다. sampler graph 구성은 첫 평가 안에서 실행되므로 sample은 0이고 eval이 둘을 모두 담습니다. |
MLXCEL_PROFILE_BLOCKS | presence enables | off | 구현된 모델 계열에서 per-block timing을 출력합니다. |
MLXCEL_PROFILE_FORWARD | presence enables | off | 구현된 모델에서 forward profiling을 활성화합니다. |
MLXCEL_PROFILE_QWEN3_MOE_DETAIL | presence enables | off | Qwen3 MoE 내부를 프로파일링합니다. |
MLXCEL_PROFILE_MOE_INNER | presence enables | off | Gemma 4 MoE sub-operation을 프로파일링합니다. |
MLXCEL_PROFILE_PER_LAYER | presence enables | off | Gemma 4 per-layer timing을 출력합니다. |
MLXCEL_PROFILE_LAYER_BUILD | presence enables | off | Gemma 4 layer-build timing을 추가합니다. |
MLXCEL_PROFILE_LAYER_SUBOPS | presence enables | off | Gemma 4 per-suboperation timing을 추가합니다. |
MLXCEL_METAL_CAPTURE_PATH | 경로 | 미설정 | 일반 mlxcel generate의 측정 구간(한 토큰 warmup 이후)과 mlxcel-bench-decode의 측정 pass를 Metal capture로 감쌉니다. MTL_CAPTURE_ENABLED=1이 필요합니다. |
MLXCEL_BENCH_PHASE_MARKS | 정확히 1(trimmed) | off | mlxcel-bench-decode 전용. host 시계 기준 [phase] <label> monotonic_ns=... boottime_ns=... mark를 stderr로 출력해 kernel trace(예: rocprofv3)를 측정된 decode 구간으로 잘라낼 수 있게 합니다. scripts/rocm_decode_profile.py가 이 mark를 읽습니다. |
MLXCEL_DEBUG_GEMMA4_LOAD | presence enables | off | Gemma 4 safetensors 로딩 진단을 출력합니다. |
MLXCEL_NO_PRECISION_WARNING | presence suppresses | warning on | Apple Silicon에서 bf16 precision/performance 안내를 숨깁니다. |
MLXCEL_LANG_BIAS_COUNTERS | 0이 아닌 비어 있지 않은 값이면 활성화 | off | bias 적용 전 suppression counter인 mlxcel_lang_bias_tokens_suppressed_total과 mlxcel_lang_bias_byte_fragment_suppressions_total을 수집합니다. bias 적용 전 argmax를 읽으려면 token bias(logit_bias, ignore_eos, --lang-bias)가 있는 요청의 디코드 단계마다 device 동기화가 필요하므로, 이 변수가 설정된 동안 그런 요청은 batched fused sampler와 lookahead 파이프라인을 쓰지 않습니다. 프로세스당 한 번 읽습니다. |
MLXCEL_VOICECHAT_PROFILE_STAGES | 0이 아닌 비어 있지 않은 값이면 활성화 | off | Nemotron VoiceChat mlxcel generate --stream --profile 요약에 하위 단계 timer(예: perception.layers, tts.backbone, codec.decoder_istft)를 추가합니다. timer가 경계마다 평가를 강제하므로 이 실행의 frame 합계는 real-time factor가 아닙니다. |
MLXCEL_VOICECHAT_PROFILE_FRAMES | 파일 경로 | 미설정 | 긴 입력의 정상 상태를 분석할 수 있도록 cold frame을 포함한 모든 VoiceChat frame의 시간을 JSON 배열로 기록합니다. |
MLXCEL_APC_TRACE | 1이면 활성화 (정확히 일치할 때만) | off | 프롬프트 캐시 trace 로깅으로, 사유별 캐시 reject 횟수와 hit/miss 판단을 기록합니다. 설정 여부만 보는 다른 스위치와 달리 1과 정확히 일치해야 합니다. |
MLXCEL_API_KEY | secret 문자열 | 미설정 | API key 소스가 아닙니다. mlxcel에서 이 값으로 인증하는 것은 없습니다(--api-key 또는 LLAMA_API_KEY 사용). 설정하면 그 값이 log secret으로 등록되어 HF_TOKEN, LLAMA_API_KEY처럼 모든 log sink에서 가려집니다. |
MLXCEL_PROFILE_PIPELINE, MLXCEL_PROFILE_PIPELINE_DETAIL, MLXCEL_EXPORT_DECODE_DOT, MLXCEL_CAPTURE_DECODE, MLXCEL_TRACE_ASTYPE, MLXCEL_TRACE_ENABLE_THINKING, MLXCEL_TRACE_BATCH_ROWS, MLXCEL_TRACE_IMAGE, MLXCEL_TRACE_IMAGE_PROMPT, MLXCEL_TRACE_LORA, MLXCEL_TRACE_LORA_MODE, MLXCEL_TRACE_TP_SIZE, MLXCEL_LOGIT_TRACE_SCALAR_REFERENCE은 현재 코드가 읽지 않습니다. 디코드의 Metal trace는 MLXCEL_METAL_CAPTURE_PATH로, 첫 토큰 이전 시간 분석은 MLXCEL_PROFILE_TTFT로 얻으십시오.
테스트와 CI 변수¶
저장소 자동화용이며 일반 사용자 실행에는 필요하지 않습니다. 테스트를 참조하십시오.
| 변수 | 용도 |
|---|---|
MLXCEL_CI_PP_MODEL | pipeline-parallel CI 통합 테스트가 사용할 모델 경로입니다. |
MLXCEL_SKIP_HEAVY_TESTS | 일부 heavy test를 건너뜁니다. |
MLXCEL_BENCH_DATE | Turbo KV benchmark test metadata override입니다. |
MLXCEL_BENCH_MACHINE | Turbo KV benchmark test metadata override입니다. |
MLXCEL_ALLOW_PARALLEL_CUDA_TESTS | CUDA 테스트 직렬화 가드를 경고로 낮춥니다. 이 가드는 cuda feature에서 mlxcel-core 스위트가 두 개 이상의 테스트 스레드로 실행되려 할 때 실패합니다. |
MLXCEL_ALLOW_CONCURRENT_GPU_TESTS | GPU 배타성 가드를 경고로 낮춥니다. 이 가드는 다른 mlxcel-core 테스트 바이너리가 이미 실행 중일 때 실패합니다. 한 디바이스에서 두 스위트가 서로 간섭하기 때문입니다. |
MLXCEL_TEST_VIDEO | ffmpeg 기반 비디오 테스트의 정상 skip을 하드 실패로 바꿉니다. 정확히 문자열 1일 때만 켜집니다. make verify-test-video가 ffmpeg 설치 후 설정합니다. |
MLXCEL_REQUIRE_PINNED_CHECKPOINTS | 고정 체크포인트 계약 테스트의 정상 skip을 하드 실패로 바꿉니다. 정확히 문자열 1일 때만 켜집니다. 고정된 Muse Glimmer 체크포인트를 가진 머신을 위한 것으로, 그 머신에서 무조건 skip하면 그 체크포인트의 유일한 커버리지가 조용히 사라지기 때문입니다. |
MLXCEL_REQUIRE_MODELS | 실제 체크포인트가 없을 때 skip 대신 테스트 실패로 만듭니다. tests/common/mod.rs와 src/ 아래 real-model parity suite에 적용되며 0이 아닌 값이면 활성화됩니다. 실제 체크포인트를 쓰는 WebUI Makefile gate는 MLXCEL_REQUIRE_MODELS=1 없이는 실행되지 않습니다. |
MLXCEL_AUDIO_PROBE_DIR | 지정 디렉터리에 audio-encoder parity probe를 기록합니다. |
MLXCEL_SPEC_DIST_SAMPLES | speculative-decoding distribution test의 sample 수를 override합니다. |
MLXCEL_TEST_GUMBEL_SAMPLES | Gumbel goodness-of-fit test의 sample 수를 override합니다. |
MLXCEL_TEST_REJECTION_SAMPLES | rejection-sampling test의 sample 수를 override합니다. |
MLXCEL_DUMP_CONFIG, MLXCEL_DUMP_OUT, MLXCEL_DUMP_DIR | ignored OpenXLA StableHLO dump test의 입출력입니다. |
MLXCEL_MOE_N | ignored OpenXLA MoE probe의 row 수이며 기본값은 4입니다. |
MLXCEL_FREEZE_GOLDENS, MLXCEL_FREEZE_MOE | 정확히 1일 때 retained OpenXLA test가 추적 중인 golden asset을 다시 쓰도록 허용합니다. 일반 test run에서는 설정하지 마십시오. |
환경 변수가 아닌 이름¶
트리에 MLXCEL_처럼 보이지만 의도적으로 제외한 문자열이 있습니다.
MLXCEL_BENCH는 shell script 내부 local variable이며 사용자 환경 설정이 아닙니다.MLXCEL_KEY는 실제 key인MLXCEL_PROMPT_CACHE_ENABLED를 담는 Rust 내부 상수 이름입니다.__MLXCEL_MINICPMO_IMAGE__는 내부 prompt sentinel입니다.MLXCEL_CUDA_ARCHITECTURES와MLXCEL_ROCM_ARCHITECTURES는build.rs가 rustc에 넘기는 컴파일 시점 값이고,MLXCEL_CUDA_INCLUDE_DIR,MLXCEL_ROCM_INCLUDE_DIR,MLXCEL_ROCM_BIN_DIR는 MLX CMake 빌드가 설정하는 컴파일 정의입니다. 어느 것도 실행 시점에 셸에서 읽지 않으며, 대상 선택에는MLX_CUDA_ARCHITECTURES와MLX_ROCM_ARCHITECTURES를 사용합니다.- CMake 파일의
MLXCEL_TURBO_KERNEL_*이름은 CMake list variable이지 process environment variable이 아닙니다. MLXCEL_BRIDGE_METAL_BACKEND,MLXCEL_BRIDGE_CUDA_BACKEND,MLXCEL_BRIDGE_ROCM_BACKEND,MLXCEL_BRIDGE_GPU_BACKEND는build.rs가 선택하는 C++ preprocessor definition이며 shell 환경 설정이 아닙니다. MLX GPU helper를 호출하는 브리지 C++ 코드는MLXCEL_BRIDGE_GPU_BACKEND뒤에 두어 CPU 전용 빌드도 링크되게 합니다.LLAMA_ARG_LANG_BIAS_POLICY,LLAMA_ARG_LANG_BIAS_CONFIG는 지원하지 않는다고 명시하는 경고문에만 등장합니다.--lang-bias-policy,--lang-bias-config를 사용하십시오.MLX_VLM_DRAFT_KIND는 오래된 오류 안내에만 등장하고 실제로 읽지 않습니다.--draft-kind또는 문서화된LLAMA_ARG_DRAFT_KIND/MLXCEL_DRAFT_KIND별칭을 사용하십시오.MLX_DISABLE_COMPILE은 graph compilation(compile)을 끄는 MLX 자체 switch입니다. mlxcel의 Rust 코드는 이를 파싱하지 않고, ROCm backend를 포함한 MLX가 읽습니다.MLX_API,MLX_GET_TYPE,MLX_MIXED_BINARY_OP,MLX_MTL_PRAGMA_UNROLL,MLX_CUDA_SM90A_ENABLED,MLX_CCCL_DIR는 shell 설정이 아니라 C/C++ compile-time macro입니다. 이름이 비슷한 runtime header override는MLXCEL_CCCL_DIR입니다.- 과거 investigation 문서에는
MLXCEL_DISABLE_GEN_STREAM,MLXCEL_ENABLE_COMPILED_SWITCH_QGEGLU처럼 현재 revision의 executable code가 읽지 않는 옛 실험 이름이 남아 있을 수 있습니다.