[
EnvVarEntry {
name: "RLX_ALLOW_THROTTLE",
group: "debug",
summary: "Skip thermal gate for one-off benches (prefer `just throttle`)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_ARCH",
group: "misc",
summary: "Trainer architecture to build: mlp | cnn",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_ARENA_ALIGN",
group: "cpu",
summary: "Arena buffer alignment in bytes (default 64)",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_ARENA_NO_REUSE",
group: "cpu",
summary: "Give every tensor its own arena slot; disables slot reuse",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_ARENA_WRITE_TRACE",
group: "cpu",
summary: "Log the arena byte range each CPU thunk dirties (finds slot aliasing/overruns)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_ATTN_DEBUG",
group: "misc",
summary: "Print CPU attention thunk selection and shapes",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_BENCHMARK_PICK",
group: "device",
summary: "Micro-benchmark N runs to pick the fastest device (needs inputs)",
kind: EnvKind::U64,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Device,
},
EnvVarEntry {
name: "RLX_BENCH_CSV",
group: "misc",
summary: "Append each bench row to this CSV file",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_BENCH_DISPATCH_ONLY",
group: "misc",
summary: "Bench dispatch overhead only, skipping the compute",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_BENCH_ITERS",
group: "misc",
summary: "Iteration count for the op benchmarks",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_CACHE_PARAM_INVARIANT",
group: "compile",
summary: "Hoist param-invariant subgraph into prepare-once graph",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_CHECK",
group: "debug",
summary: "Static-check breadth: off | all | strict",
kind: EnvKind::String,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_COREML_F16",
group: "coreml",
summary: "Store CoreML activations and dequant output in f16 (~half bandwidth)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("coreml"),
},
EnvVarEntry {
name: "RLX_COREML_FLEXIBLE_INPUTS",
group: "coreml",
summary: "Declare CoreML inputs as flexible-shape",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("coreml"),
},
EnvVarEntry {
name: "RLX_COREML_HOST_DEQUANT",
group: "coreml",
summary: "Force CoreML hybrid host dequant segments",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("coreml"),
},
EnvVarEntry {
name: "RLX_COREML_NATIVE_FLEX",
group: "coreml",
summary: "Compile dynamic graphs once with CoreML ShapeRange instead of per shape",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("coreml"),
},
EnvVarEntry {
name: "RLX_COREML_NATIVE_SCAN",
group: "coreml",
summary: "Lower Scan natively on CoreML; set 0 to host it",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("coreml"),
},
EnvVarEntry {
name: "RLX_COREML_Q1_MODE",
group: "coreml",
summary: "Q1_0 on-device lowering mode: lut (default) | f32 (legacy unfold)",
kind: EnvKind::String,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("coreml"),
},
EnvVarEntry {
name: "RLX_COREML_SEG_REPORT",
group: "coreml",
summary: "Report the CoreML/host segmentation plan",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("coreml"),
},
EnvVarEntry {
name: "RLX_COREML_UNITS",
group: "coreml",
summary: "CoreML compute units: cpu | gpu | all | ane",
kind: EnvKind::String,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("coreml"),
},
EnvVarEntry {
name: "RLX_COREML_VERIFY",
group: "coreml",
summary: "MIL program verifier: `strict` rejects a malformed program at build time",
kind: EnvKind::String,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("coreml"),
},
EnvVarEntry {
name: "RLX_CPU_MATMUL_F64_ACCUM",
group: "cpu",
summary: "Accumulate the CPU matmul K-reduction in f64 (precision/validated mode; default-off, vendor-BLAS fast path untouched)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_CPU_DUMP_FLAT",
group: "misc",
summary: "Also print this flat element index when dumping a diverging node",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_CPU_DUMP_NODES",
group: "misc",
summary: "Print per-node max|x| and nonzero counts during execution",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_CPU_DUMP_NODES_LIMIT",
group: "misc",
summary: "Cap how many nodes the CPU dump writes (default 2000)",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_CUDA_ARENA_DEBUG",
group: "cuda",
summary: "Log CUDA arena allocation and slot reuse",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_ARENA_NO_REUSE",
group: "cuda",
summary: "Give every CUDA tensor its own arena slot; disables slot reuse",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_ARENA_POOL",
group: "cuda",
summary: "Pool CUDA arena allocations across compiles",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_ARENA_POOL_CHUNK_BYTES",
group: "cuda",
summary: "Max bytes per pooled CUDA arena chunk",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_ARENA_POOL_MAX",
group: "cuda",
summary: "Max buffers retained in the CUDA arena pool",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_COMPILE_MODE",
group: "cuda",
summary: "CUDA compile mode: jit (default) | aot",
kind: EnvKind::String,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_COMPILE_TIMING",
group: "cuda",
summary: "Report NVRTC compile time per kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_CONV_BWD_CUDNN",
group: "cuda",
summary: "Allow cuDNN for grouped/degenerate Conv2d backward shapes",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_CONV_BWD_HOST",
group: "cuda",
summary: "Force CUDA Conv2d backward through CPU host-fallback (parity/debug)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_CONV_FWD_CUDNN",
group: "cuda",
summary: "Force CUDA Conv2d forward through cuDNN (override host/default routing)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_CONV_FWD_HOST",
group: "cuda",
summary: "Force CUDA Conv2d forward through CPU host-fallback (bisect)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_CONV_STABLE_BWD",
group: "cuda",
summary: "Prefer cuDNN IMPLICIT_GEMM (ALGO_1) for conv backward (default on; set 0 to opt out)",
kind: EnvKind::BoolOr { default: true },
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_CONV_TF32",
group: "cuda",
summary: "Enable TF32 tensor-core math for cuDNN conv (default is FMA)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_DUMP_INTERMEDIATE",
group: "cuda",
summary: "Dump intermediate node buffers (see RLX_CUDA_DUMP_NODES_LIMIT)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_DUMP_NODES",
group: "cuda",
summary: "Dump output buffers per node for Metal/CPU cross-diff",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_DUMP_NODES_LIMIT",
group: "cuda",
summary: "Cap how many nodes RLX_CUDA_DUMP_INTERMEDIATE writes",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_EXEC_MODE",
group: "cuda",
summary: "CUDA execution mode: stream (default) | graph | multistream:N",
kind: EnvKind::String,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_GDN_HOST",
group: "cuda",
summary: "Force GatedDeltaNet through the host D2H-CPU-H2D path",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_GGUF_FUSED_M1",
group: "cuda",
summary: "Fused m=1 GGUF decode GEMV; set 0 to keep planning scratch",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_GGUF_HOST",
group: "cuda",
summary: "Dequantize GGUF weights on the host instead of on device",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_IM2COL_HOST",
group: "cuda",
summary: "Force the host im2col fallback instead of the GPU kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_INDEXING_FULL_ARENA",
group: "cuda",
summary: "Give indexing ops the full arena instead of a window",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_INDEXING_HOST",
group: "cuda",
summary: "Force ND indexing (GatherND/GatherElements/Scatter*) to the CPU host route",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_LOG_CONV_PATH",
group: "cuda",
summary: "Log which convolution implementation was selected",
kind: EnvKind::Path,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_LOG_FALLBACK",
group: "cuda",
summary: "Log every op that falls back off its native CUDA path",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_MATMUL_PRECISE",
group: "cuda",
summary: "Compensated k-reduction for matmul, improving precision at long K",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_NONDET_CONV",
group: "cuda",
summary: "Allow non-deterministic cuDNN conv backward algos (atomicAdd; faster, noisy)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_NO_CUBLAS",
group: "cuda",
summary: "Skip both cuBLAS tiers so dense GEMM runs rlx's own tiled `matmul` kernel (A/B vs the vendor library; also how the dispatch tuner reaches the tile it tunes)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_NO_CUBLASLT",
group: "cuda",
summary: "Disable cuBLASLt and use plain cuBLAS",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_NO_CUDNN",
group: "cuda",
summary: "Skip cuDNN entirely (im2col / custom kernels only; silence missing-lib warning)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_NO_PACKED_BSHD_ATTN",
group: "cuda",
summary: "Disable the packed BSHD attention kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_NO_TF32",
group: "cuda",
summary: "Disable TF32 tensor cores (costs ~1e-4 relative error otherwise)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_NO_ZERO_ARENA",
group: "cuda",
summary: "Skip zeroing the CUDA arena on allocation",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_PARITY",
group: "cuda",
summary: "Tighten CUDA numerics for cross-backend parity (implies no TF32)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_PATH_TRACE",
group: "cuda",
summary: "Trace which implementation path each op takes",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_PINNED_IO",
group: "cuda",
summary: "Pinned host staging for faster D2H; on by default, set 0 to disable",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_PTX_CACHE",
group: "cuda",
summary: "Directory for the compiled PTX cache (overrides XDG_CACHE_HOME)",
kind: EnvKind::Path,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_Q4K_GEMV_WARP",
group: "cuda",
summary: "Warp-per-row Q4_K decode GEMV; full occupancy at small k, unlike the block-per-row coop kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_TRACE_FAB",
group: "cuda",
summary: "Trace fused-attention-block construction",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_WMMA",
group: "cuda",
summary: "Opt in to the WMMA tensor-core matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDNN_DIR",
group: "cuda",
summary: "Directory containing libcudnn.so* to preload (pip/conda wheel failsafe)",
kind: EnvKind::Path,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_DBG_BINF",
group: "debug",
summary: "Print binary-op shape inference during ONNX import",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_DBG_CONV",
group: "debug",
summary: "Print convolution lowering details",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_DBG_CUSTOM",
group: "debug",
summary: "Log host custom-op staging (onnx.* dtype bridge) on GPU backends",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_DBG_SHAPES",
group: "debug",
summary: "Print inferred shapes during ONNX import",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_DBG_STEP",
group: "debug",
summary: "Print each wgpu execution step as it runs",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_DEBUG_NANS",
group: "debug",
summary: "Runtime NaN/Inf localizer (1 or abort)",
kind: EnvKind::Enum(&["", "1", "abort"]),
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_DECODE_BUCKET_PEAK_BYTES",
group: "misc",
summary: "Override the estimated compile peak bytes for bucketed decode",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_DECODE_BUCKET_RESIDENT_BYTES",
group: "misc",
summary: "Override the assumed resident bytes per decode bucket",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_DECODE_ONESHOT_PEAK_BYTES",
group: "misc",
summary: "Override the estimated compile peak bytes for one-shot decode",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_DECOMPOSE_FUSION_REGIONS",
group: "misc",
summary: "Decompose fusion regions back to primitives before lowering",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_DECOMPOSE_SPLINE_BWD",
group: "misc",
summary: "Force the KAN spline VJP to the decomposed exp/mul/reduce chain instead of the fused SplineActivationBackwardX/Coeff ops (debug/parity).",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_DEQUANT_CACHE",
group: "misc",
summary: "Cache dequantized weights; set 0 to disable caching entirely",
kind: EnvKind::Path,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_DEVICE",
group: "device",
summary: "Default device hint for resolved runs (cpu, metal, mlx, cuda, gpu, …)",
kind: EnvKind::String,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Device,
},
EnvVarEntry {
name: "RLX_DEVICES",
group: "device",
summary: "Allow-list of devices for DevicePolicy::from_env",
kind: EnvKind::String,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Device,
},
EnvVarEntry {
name: "RLX_DEVICE_CHAIN",
group: "device",
summary: "Fallback order when a preferred device fails (e.g. cuda,gpu,cpu)",
kind: EnvKind::String,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Device,
},
EnvVarEntry {
name: "RLX_DIRECT_CONV",
group: "misc",
summary: "Opt in to the direct (non-im2col) convolution path",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_DISABLE_CONV_BIAS_ACT_FUSION",
group: "compile",
summary: "Skip Conv+Bias+Act fusion",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_DISABLE_MPS",
group: "misc",
summary: "Disable MPSMatrixMultiplication; use rlx's own MSL sgemm",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_DISABLE_MPSGRAPH",
group: "metal",
summary: "Force Metal thunk path instead of MPSGraph regions",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_DISABLE_MPSGRAPH_EXECUTABLE",
group: "metal",
summary: "Disable the precompiled MPSGraphExecutable path",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_DISABLE_NOMIC_FUSION",
group: "misc",
summary: "Disable the Nomic-embed fusion pattern for an A/B",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_DISPATCH_REPORT",
group: "debug",
summary: "Print legalize/dispatch report during compile (1 = on)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_DUMP_SCHED",
group: "misc",
summary: "Dump the encode schedule",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_ENABLE_FUSE_TRANSFORMER_LAYER",
group: "misc",
summary: "Enable whole-transformer-layer fusion",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_ENV_DEPRECATIONS",
group: "debug",
summary: "Emit one-shot messages when deprecated RLX_* aliases are used",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_F32_DUMP",
group: "misc",
summary: "Directory the trainer writes f32 weight dumps to",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_FAST_CONV",
group: "cpu",
summary: "CPU Conv2d im2col+BLAS path (default on; set 0 for scalar nested loops)",
kind: EnvKind::BoolOr { default: true },
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_FFT_CPU_PARALLEL",
group: "fft",
summary: "Rayon-parallelize the CPU FFT batch loop; set 0 to serialize",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_FFT_CUFFT",
group: "fft",
summary: "Use cuFFT; set 0 to fall back to rlx's own FFT kernels",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_FFT_FAST",
group: "gpu",
summary: "Enable native on-chip GPU FFT when compiled (0 disables)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_FFT_FORCE_MIXED",
group: "fft",
summary: "Force the precompiled mixed-radix FFT instead of the generated kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_FFT_FUSE_DEBUG",
group: "fft",
summary: "Print FFT fusion decisions",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_FFT_FUSE_REAL",
group: "fft",
summary: "Fuse the real-to-complex FFT prologue; set 0 to disable",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_FFT_GEN",
group: "fft",
summary: "Use the per-size generated FFT kernel; set 0 for the mixed-radix fallback",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_FFT_MULTIROW",
group: "fft",
summary: "Opt in to the multi-row small-n FFT path (default off)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_FFT_NATIVE",
group: "fft",
summary: "Native Stockham FFT path; set 0 to disable it",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_FFT_RADIX",
group: "fft",
summary: "Cap the FFT radix for an A/B: 2 | 4 (default 8)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_FFT_RADIX4",
group: "fft",
summary: "Radix-4 CPU FFT for pure powers of four; set 0 to disable",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_FK_BATCH_SINGLE_KERNEL",
group: "misc",
summary: "Emit one fused kernel for the batched FK path",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_FOLD_LIVENESS_DEBUG",
group: "compile",
summary: "Report how many Conv3D epilogue folds extended a live range (absorbed / read-through)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_FORCE_DEVICE",
group: "misc",
summary: "Hard-pin the backend, overriding the cost model's device choice",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_FUSED_GPU_CONIC_SCANLINE",
group: "misc",
summary: "Experimental GPU conic scanline splat rasterizer",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_FUSE_ATTN_THRESHOLD",
group: "compile",
summary: "Sequence length above which attention fusion engages",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_FUSE_BATCH_PREPROCESS",
group: "compile",
summary: "Merge parallel region slices into BatchElementwiseRegion; 0 disables",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_FUSE_REGION_PROLOGUE",
group: "compile",
summary: "Fold ResizeNearest2x into the ElementwiseRegion prologue; 0 disables",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_DISABLE_SCCP",
group: "compile",
summary: "Disable sparse conditional constant propagation (A/B the pass)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_FUSION_REPORT",
group: "compile",
summary: "Print fusion pass before/after report",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_GDN_UNFUSE_FOR_AD",
group: "compile",
summary: "Unroll Op::GatedDeltaNet's time loop for autodiff instead of emitting the fused backward (~32x slower)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_GGUF_MATMUL_LEGACY",
group: "misc",
summary: "Use the legacy GGUF matmul dispatch instead of gguf_matmul_bt_dispatch",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_GGUF_TRACE",
group: "misc",
summary: "Trace GGUF header and tensor parsing",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_GLSLANG",
group: "vulkan",
summary: "Path to `glslangValidator` for GLSL->SPIR-V (naga lacks memoryBarrierShared)",
kind: EnvKind::Path,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_GPU_HANDLE_HOST_MIRROR",
group: "misc",
summary: "Mirror GPU handle outputs to the host on every read",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_GPU_TUNING_CACHE",
group: "misc",
summary: "Path to the persisted GPU dispatch tuning cache (measured kernel-variant + tile winners per arch/op/shape-bucket); unset = under $XDG_CACHE_HOME",
kind: EnvKind::Path,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_HIGHER_ORDER_NO_FUSE",
group: "misc",
summary: "Disable elementwise fusion for higher-order derivatives",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_HYBRID_K",
group: "misc",
summary: "Override the k dimension in the hybrid dequant-matmul parity test",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_HYBRID_M",
group: "misc",
summary: "Override the hybrid-split m dimension (default 96)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_ICB_TRACE",
group: "misc",
summary: "Trace indirect-command-buffer construction",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_IMP_DBG",
group: "misc",
summary: "Print per-op decisions during ONNX import",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_INDEXING_FULL_ARENA",
group: "gpu",
summary: "Force full-arena mirror for indexing host-fallback (bisect; slow on discrete GPUs)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_IQ_TEST_DIR",
group: "misc",
summary: "Directory holding the IQ/TQ test GGUFs",
kind: EnvKind::Path,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_IR_DUMP",
group: "misc",
summary: "Write a full pipeline IR dump to this path prefix or directory",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_KEEP_ELEMENTWISE_REGIONS",
group: "misc",
summary: "Keep fused ElementwiseRegion nodes through lowering",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_KERNEL_DISPATCH",
group: "compile",
summary: "common|native|force_common|force_native kernel dispatch policy",
kind: EnvKind::Enum(&["common", "native", "force_common", "force_native"]),
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_KV_CACHE_DBG",
group: "misc",
summary: "Log KV cache eviction decisions",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_KV_CACHE_MAX_RESIDENT",
group: "misc",
summary: "Byte ceiling above which KV cache buckets become evictable",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_KV_CACHE_NO_EVICT",
group: "misc",
summary: "Never evict KV cache buckets",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_LINT_NUMERICS",
group: "compile",
summary: "Static provable NaN/Inf lint during compile",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_LM_HEAD_PARALLEL",
group: "misc",
summary: "Scan vocab rows in parallel via Rayon; set 0 to serialize",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_LOG_EPOCH_LOSS",
group: "misc",
summary: "Print the training loss after each epoch",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_LSTM_DEBUG",
group: "misc",
summary: "Print LSTM lowering details during ONNX import",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_METAL_ATTN_BWD_GPU",
group: "metal",
summary: "Run attention backward on the GPU; set 0 to host it",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_ATTN_TRACE",
group: "metal",
summary: "Trace Metal attention kernel selection",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CMDBUF_TRACE",
group: "metal",
summary: "Trace Metal command-buffer status transitions",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONCAT_HOST",
group: "metal",
summary: "Opt in to the host-staged Metal Concat fallback",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONCAT_MULTI",
group: "metal",
summary: "Disable the multi-input Concat kernel; encode one input at a time",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONCURRENT",
group: "metal",
summary: "Open compute encoders in Concurrent dispatch so independent decode dispatches overlap",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONCURRENT_BARRIER_FRESH",
group: "metal",
summary: "Also barrier when no encoder is open yet, by opening it first (a Concurrent encoder is not implicitly ordered after the previous one)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONCURRENT_FENCE_ALL",
group: "metal",
summary: "Fence every concurrent dispatch instead of only the data-dependent ones",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONCURRENT_IDX_HI",
group: "metal",
summary: "Upper bound (exclusive) of the thunk index range allowed to use Concurrent dispatch",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONCURRENT_IDX_LO",
group: "metal",
summary: "Lower bound (inclusive) of the thunk index range allowed to use Concurrent dispatch",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONCURRENT_OPAQUE",
group: "metal",
summary: "Comma-separated thunk kinds to force-fence, to bisect which op is sensitive to the encoder dispatch type",
kind: EnvKind::String,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONCURRENT_SPLIT_ENC",
group: "metal",
summary: "Give every thunk its own encoder so the dispatch type can be chosen per thunk (bisection aid)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONCURRENT_STATS",
group: "metal",
summary: "Report encoders opened vs thunks dispatched, to tell whether Concurrent can overlap anything",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONV3D",
group: "metal",
summary: "Force a Conv3D kernel: naive (scalar) or gemm (skip the c-tiled variant)",
kind: EnvKind::Enum(&["naive", "gemm"]),
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONV3D_FOLDS",
group: "metal",
summary: "Enable the Conv3D concat/upsample/leaky thunk folds (off: they regress arena liveness)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONV3D_MIN_M",
group: "metal",
summary: "Minimum Conv3D output-position count before the tiled implicit-GEMM kernel is used (default 8)",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONV_BWD_IMPLICIT",
group: "metal",
summary: "Use implicit GEMM for conv backward-weight; set 0 for explicit",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_DEBUG",
group: "metal",
summary: "Extra Metal validity checks and diagnostics",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_DEQUANT_GPU_DISABLE",
group: "metal",
summary: "Disable Metal GPU GGUF/MLX dequant (host / legacy path)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &["RLX_DISABLE_METAL_DEQUANT_GPU"],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_EXPAND_CONV3D_BIAS",
group: "metal",
summary: "Let a rank-5 no-activation FusedConvBiasAct keep its expanded bias instead of declining the fusion",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_FOLD_AUDIT",
group: "metal",
summary: "Assert no two arena-sharing nodes have overlapping FOLD-ADJUSTED live ranges",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_FUSE_L2NORM",
group: "metal",
summary: "Fuse the ggml L2_NORM op chain into one L2NormLastDim dispatch (default on; =0 disables)",
kind: EnvKind::BoolOr { default: true },
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_G8_0_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused G8_0 decode GEMV and dequantize to a scratch buffer first",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_G8_0_SCALAR",
group: "metal",
summary: "Fall back to the scalar byte inner loop in the G8_0 GEMV (A/B only)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_NO_CONV3D_CONCAT_FOLD",
group: "metal",
summary: "Disable folding a channel-wise Concat into the Conv3D gather",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_NO_CONV3D_FOLDS",
group: "metal",
summary: "Disable every Conv3D epilogue thunk fold at once (bisect the fold path as a whole)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_NO_CONV3D_LEAKY_FOLD",
group: "metal",
summary: "Disable folding LeakyReLU into the Conv3D kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_NO_CONV3D_UPSAMPLE_FOLD",
group: "metal",
summary: "Disable folding a nearest-neighbour Upsample into the Conv3D gather",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_PARAMS",
group: "metal",
summary: "Apple kernel tuning spec, `k=v,...` (stages/sync/precision/tile/encode)",
kind: EnvKind::String,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q2_0_SCALAR",
group: "metal",
summary: "Fall back to the scalar byte inner loop in the Q2_0 GEMVs instead of the shared q2_0_dot16 (A/B only)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q3K_SG_DISABLE",
group: "metal",
summary: "Disable the simdgroup Q3_K decode GEMV, using the one-thread-per-row kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q41_SG_DISABLE",
group: "metal",
summary: "Disable the simdgroup Q4_1 decode GEMV, using the one-thread-per-row kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q5K_SG_DISABLE",
group: "metal",
summary: "Disable the simdgroup Q5_K decode GEMV, using the one-thread-per-row kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_MLX_DEQUANT_GPU_DISABLE",
group: "quant",
summary: "Force host-staged MLX DequantMatMul on all GPU backends",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_METAL_DEQUANT_MATMUL_LEGACY",
group: "metal",
summary: "Use pre-fused dequant+matmul path (materializes weights)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_DISABLE_NARROW_ROPE_FUSE",
group: "metal",
summary: "Disable fusing Narrow into RoPE",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_DUMP_NODES",
group: "metal",
summary: "Dump per-node output buffers on Metal",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_DUMP_NODES_LIMIT",
group: "metal",
summary: "Cap how many nodes the Metal dump writes (default 4000)",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_EXTERNALIZE_QUANT",
group: "metal",
summary: "Keep quantized weights external rather than baking them in",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_EXT_TRACE",
group: "metal",
summary: "Trace Metal op-extension lowering",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_FA",
group: "metal",
summary: "Opt in to the Metal flash-attention kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_FFT_HOST_FALLBACK",
group: "metal",
summary: "Run FFT on the host instead of the Metal kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_FORCE_INLINE_PARAMS",
group: "metal",
summary: "Inline parameters into the encoder rather than binding buffers",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_FORCE_PIN_OUTPUT_ANCESTORS",
group: "metal",
summary: "Force output-ancestor buffers to stay pinned",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_FORCE_UNPIN_OUTPUT_ANCESTORS",
group: "metal",
summary: "Force output-ancestor buffers to be unpinned",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_FUSE_DECODE",
group: "metal",
summary: "Fuse the decode MLP on Metal; on by default, set 0 to opt out",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_FUSE_DECODE_GELU",
group: "metal",
summary: "Combined GeGLU decode fusion; set 0 to opt out",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_FUSE_DECODE_LOG",
group: "metal",
summary: "Log which decode fusions fired",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_FUSE_DEPTHWISE",
group: "metal",
summary: "Fuse depthwise conv1d in BSC layout; set 0 to disable",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_FUSE_GDN_NORM",
group: "metal",
summary: "Fuse the GatedDeltaNet gated norm; set 0 to disable",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_FUSE_RESIDUAL_RMS",
group: "metal",
summary: "Fuse residual add into RmsNorm; on by default, set 0 to opt out",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_GDN_CPU",
group: "metal",
summary: "Run GatedDeltaNet on the CPU instead of Metal",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_GDN_HOST_FALLBACK",
group: "metal",
summary: "Force the host fallback for GatedDeltaNet on Metal",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_HOST_FALLBACK",
group: "metal",
summary: "Run unsupported Metal ops on the host instead of failing",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_HOST_SLICE",
group: "metal",
summary: "Run Slice on the host instead of a Metal kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_HYBRID_BIG_ARENA",
group: "metal",
summary: "Allow the large hybrid arena instead of all-thunks (can OOM)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_IQ1M_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused IQ1_M matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_IQ1S_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused IQ1_S matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_IQ2S_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused IQ2_S matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_IQ2XS_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused IQ2_XS matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_IQ2XXS_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused IQ2_XXS matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_IQ3S_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused IQ3_S matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_IQ3XXS_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused IQ3_XXS matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_IQ4NL_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused IQ4_NL matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_LSTM_CPU",
group: "metal",
summary: "Run LSTM on the CPU instead of Metal",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_LSTM_HOST_FALLBACK",
group: "metal",
summary: "Run LSTM on the host instead of Metal",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_MPSGRAPH_BIG_ARENA",
group: "metal",
summary: "Allow MPSGraph to use the large arena",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_MPS_PROFILE",
group: "metal",
summary: "Time MPSGraph and hybrid dispatches",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_MPS_SDPA",
group: "metal",
summary: "Route SDPA through MPSGraph",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_NARROW_BATCH",
group: "metal",
summary: "Disable batching of adjacent Narrow ops",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_NO_FUSION",
group: "metal",
summary: "Skip every Metal pattern fusion",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_NO_SHARE",
group: "metal",
summary: "Disable sharing parameter buffers between compiled graphs",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_ONNX_QMATMUL_GPU",
group: "metal",
summary: "Run ONNX quantized matmul on the GPU; set 0 to host it",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_ONNX_QMATMUL_MIN_FLOPS",
group: "metal",
summary: "FLOP threshold above which ONNX qmatmul goes to the GPU",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_OUTPUT_TRACE",
group: "metal",
summary: "Trace each graph output read back from Metal",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_PIPELINE_CACHE",
group: "metal",
summary: "On-disk Metal pipeline-state cache directory",
kind: EnvKind::Path,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_PRECISE",
group: "metal",
summary: "Force the scalar fp32 Metal path instead of reduced-precision tensor units",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q1_0_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused Q1_0 matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q1_0_SG_DISABLE",
group: "metal",
summary: "Disable the simdgroup Q1_0 kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q1_DUAL_DISABLE",
group: "metal",
summary: "Disable the dual-output Q1 decode kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q2_0_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused Q2_0 matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q2_0_SG_DISABLE",
group: "metal",
summary: "Disable the simdgroup Q2_0 kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q2_DUAL_DISABLE",
group: "metal",
summary: "Disable the dual-output Q2 decode kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q40_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused Q4_0 matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q41_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused Q4_1 matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q4K_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused Q4_K matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q4K_GEMM_DISABLE",
group: "metal",
summary: "Disable the Q4_K GEMM kernel; use GEMV",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q4K_SG_DISABLE",
group: "metal",
summary: "Disable the simdgroup Q4_K kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q6K_GEMM_DISABLE",
group: "metal",
summary: "Disable the Q6_K GEMM kernel; use GEMV",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q80_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused Q8_0 matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_RNN_HOST_FALLBACK",
group: "metal",
summary: "Run RNN ops on the host instead of Metal",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SAMPLE_HOST",
group: "metal",
summary: "Run token sampling on the host instead of on device",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SDPA_DECODE_M1",
group: "metal",
summary: "Dedicated m=1 SDPA decode kernel; set 0 to disable",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SGEMM_MPS",
group: "metal",
summary: "Force the MPS sgemm path at shapes the cost model would route elsewhere",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SGEMM_PRECISE",
group: "metal",
summary: "Force the scalar fp32 sgemm path for precision-critical work",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SGEMM_VARIANT",
group: "metal",
summary: "Pin one Metal sgemm variant for an A/B (refused where ineligible)",
kind: EnvKind::Enum(&[
"mps",
"simd4x4",
"simd_4x4",
"4x4",
"simd64",
"simd_64",
"64x64",
"simd64splitk",
"simd64_splitk",
"splitk",
"simd",
"simd8",
"simd_8",
"padded",
"simd_padded",
"simdpadded",
"tiled",
"naive",
]),
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_NO_SGEMM64",
group: "metal",
summary: "Disable the default 64x64-tile Simd64 sgemm (fall back to Simd4x4) for aligned tall/short-K matmuls.",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SGEMM_SPLITK",
group: "metal",
summary: "Enable the split-K 64x64 sgemm for fat-K/small-MN shapes (dW=xT.dq); atomic-accumulate into a pre-zeroed C.",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SOFTMAX_TRACE",
group: "metal",
summary: "Trace softmax kernel selection",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SSM_CPU",
group: "metal",
summary: "Run the SSM recurrence on the CPU instead of Metal",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SSM_HOST_FALLBACK",
group: "metal",
summary: "Run the SSM recurrence on the host instead of Metal",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SYNTH_MPS_DISABLE",
group: "metal",
summary: "SynthMatMul m>8 prefill uses the fused kernel instead of reconstruct→MPS (A/B)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SYNTH_RECON_F16",
group: "metal",
summary: "SynthMatMul m>8 prefill reconstructs the weight in f16 → MPS hgemm (~1.3×, 2× smaller scratch, relaxed precision)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SYNTH_TILED",
group: "metal",
summary: "SynthMatMul m>8 uses the threadgroup-tiled fused kernel (zero-scratch/capturable; slower than recon→MPS)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SYNTH_TILED_F16",
group: "metal",
summary: "With RLX_METAL_SYNTH_TILED, use the f16 (simdgroup_half8x8) tiled kernel variant",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_THUNK_PROFILE",
group: "metal",
summary: "Time each Metal schedule thunk individually",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_TRACE",
group: "metal",
summary: "Log every Metal dispatch as it is encoded",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_TRACE_FAB",
group: "metal",
summary: "Trace fused-attention-block construction on Metal",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_UNFUSE_REGIONS",
group: "metal",
summary: "Unfuse ElementwiseRegion nodes before Metal lowering",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_MLX_COMPILE_MAX_NODES",
group: "mlx",
summary: "Max graph nodes MLX will compile as one Compiled unit (default 1536)",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_CUDA_ARCH",
group: "mlx",
summary: "CUDA arch to build MLX for",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_DEBUG_EVAL",
group: "mlx",
summary: "Log every MLX eval boundary",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_DEQUANT_CACHE_BYTES",
group: "mlx",
summary: "Byte budget for the MLX dequantized-weight cache",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_DEQUANT_CACHE_DISABLE",
group: "mlx",
summary: "Disable the MLX dequantized-weight cache",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_DEVICE",
group: "mlx",
summary: "MLX device label for the device benchmark",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_FUSE_CAP",
group: "mlx",
summary: "Max ops MLX will fuse into one kernel (default 12)",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_GGUF_HOST_FALLBACK",
group: "mlx",
summary: "Force host GGUF dequant on MLX",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_JOBS",
group: "mlx",
summary: "Parallel job count for the MLX CMake build",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_MODE",
group: "mlx",
summary: "eager | lazy | compiled execution mode",
kind: EnvKind::Enum(&["eager", "lazy", "compiled"]),
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_PARAM_VIEW",
group: "mlx",
summary: "Hold MLX parameters as views rather than owned arrays",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_PROFILE",
group: "mlx",
summary: "Dump the per-op-kind wall-time breakdown once",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_Q1_HOST",
group: "mlx",
summary: "Dequantize Q1 weights on the host instead of through MLX",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_Q1_MV_DISABLE",
group: "mlx",
summary: "Disable the Q1 matrix-vector kernel; use one-shot dequant plus sgemm",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_RNN_F16",
group: "mlx",
summary: "Run the RNN recurrence in fp16 compute on the Metal GPU",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_SDPA_REFERENCE",
group: "mlx",
summary: "Use reference SDPA composition for bisects",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_WARN_LAZY",
group: "mlx",
summary: "Lazy-evaluation warning level; set all for per-executable warnings",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MNIST_DIR",
group: "misc",
summary: "Directory holding the raw MNIST files",
kind: EnvKind::Path,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_MPSGRAPH_FORCE",
group: "metal",
summary: "Force this op through MPSGraph regardless of the cost model",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_MPSGRAPH_MIN_FLOPS",
group: "metal",
summary: "FLOP threshold above which MPSGraph is used",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_MPSGRAPH_PARAM_CONST",
group: "metal",
summary: "Bake parameters into MPSGraph as constants (opt-in)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_MPSGRAPH_PARAM_CONST_CAP",
group: "metal",
summary: "Byte cap on parameters frozen into MPSGraph constants",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_MPSGRAPH_TRACE",
group: "metal",
summary: "Log when the MPSGraph path fires",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_MPSG_TRACE",
group: "misc",
summary: "Verbose MPSGraph construction trace",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_MPS_ALIGN_DEBUG",
group: "metal",
summary: "Report MPS buffer alignment decisions",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_MPS_FP16",
group: "metal",
summary: "Cast both inputs of a 2-input MPS op to fp16",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_MPS_THRESHOLD_FLOP",
group: "metal",
summary: "FLOP cutoff above which matmul routes to MPS",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_NATIVE_FK_REGIONS",
group: "misc",
summary: "Keep TransformRegion/BatchElementwiseRegion in MIR for native lowering",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_NEMO_TEST_FILE",
group: "misc",
summary: "Path to a .nemo file for the NeMo tests",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_NO_FK_FUSION",
group: "misc",
summary: "Disable the FKL fusion passes",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_NO_IO_PEAKS_OUTPUT",
group: "compile",
summary: "Disable compile-time IO-gated peaks-only fusion",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_NO_NATIVE_FK_REGIONS",
group: "misc",
summary: "Force FK regions to decompose rather than lower natively",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_NO_SHARED_INPUT_MATMUL",
group: "misc",
summary: "Opt out of sharing one matmul across consumers of the same input",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_NO_SHUFFLE",
group: "misc",
summary: "Do not shuffle the training set between epochs",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_OCLOC",
group: "oneapi",
summary: "Path to Intel `ocloc` for OpenCL C ahead-of-time compilation",
kind: EnvKind::Path,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("oneapi"),
},
EnvVarEntry {
name: "RLX_ONEAPI_BUILD_KERNELS",
group: "oneapi",
summary: "Build native oneAPI SPIR-V kernels at build time (needs ocloc)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("oneapi"),
},
EnvVarEntry {
name: "RLX_ONEAPI_LOADER",
group: "oneapi",
summary: "Path to the Level Zero loader library",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("oneapi"),
},
EnvVarEntry {
name: "RLX_ONEAPI_OCLOC",
group: "oneapi",
summary: "Path to the ocloc kernel compiler (default ocloc)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("oneapi"),
},
EnvVarEntry {
name: "RLX_ONEAPI_OCLOC_DEVICE",
group: "oneapi",
summary: "Target device token passed to ocloc (default pvc)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("oneapi"),
},
EnvVarEntry {
name: "RLX_ONNX_ALLOW_RESIZE_STUB",
group: "onnx",
summary: "Import an unlowerable ONNX Resize as zeros instead of failing (runs, output is silently wrong)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_ONNX_BUNDLE",
group: "onnx",
summary: "Path to an exported ONNX bundle directory",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_ONNX_SEQUENCE_LENGTH",
group: "onnx",
summary: "Sequence length to restore and propagate for ONNX import",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_ONNX_TAP",
group: "onnx",
summary: "Comma-separated ONNX tensor names to tap during import",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_ONNX_TEST_MODEL",
group: "onnx",
summary: "Path to the ONNX model used by the import tests",
kind: EnvKind::String,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_ORT_CUDA_GRAPH",
group: "misc",
summary: "Enable CUDA graph capture in the ONNX Runtime CUDA provider",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_PARITY_DEVICE",
group: "misc",
summary: "Device the parity harness compares against: cuda (default) | rocm",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_PAR_THRESHOLD",
group: "cpu",
summary: "Minimum element count before an op dispatches in parallel",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_PHASE_TIMING",
group: "misc",
summary: "Print wall time for each compiler phase",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_PLAN_VERIFY",
group: "compile",
summary: "Re-derive every memory-plan offset with the O(N^3) candidate scan and assert the sweep agrees",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_PROBE_DYNAMIC",
group: "misc",
summary: "Import the probe model with dynamic shapes",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_PROBE_FEATURE_DIM",
group: "misc",
summary: "Feature dimension for the ONNX import probe",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_PROFILE_COMPILE",
group: "misc",
summary: "Report time spent compiling each LIR kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_PROFILE_THUNKS",
group: "misc",
summary: "Time each thunk to see which ops dominate a step",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_QNN_BACKEND_LIB",
group: "qnn",
summary: "Explicit path to the QNN backend library (libQnn*.so)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("qnn"),
},
EnvVarEntry {
name: "RLX_REQUIRE_DEVICE",
group: "debug",
summary: "Make a missing GPU a test FAILURE instead of a silent skip (rig runs)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_RIG_RUNTIME",
group: "misc",
summary: "Where the device sweep runs: local (default) or a rig host",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_ROCM_HSACO_CACHE",
group: "rocm",
summary: "Directory for the compiled HSACO cache",
kind: EnvKind::Path,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_IM2COL_HOST",
group: "rocm",
summary: "Run im2col on the host instead of the ROCm kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_INDEXING_HOST",
group: "rocm",
summary: "Force ND indexing (GatherND/GatherElements/Scatter*) to the CPU host route",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_LOG_FALLBACK",
group: "rocm",
summary: "Log every op that falls back off its native ROCm path",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_MFMA",
group: "rocm",
summary: "Use MFMA matrix cores where the arch supports them",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_NO_PACKED_BSHD_ATTN",
group: "rocm",
summary: "Disable the packed BSHD attention kernel on ROCm",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_PINNED_IO",
group: "rocm",
summary: "Use pinned host I/O for ROCm graph exec (default on in graph mode)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_SCHEDULE_MATMUL",
group: "rocm",
summary: "Source for the HIP matmul kernel: schedule-IR emitted vs handwritten",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_SMI_INDEX",
group: "rocm",
summary: "rocm-smi card index for contention gating (follows HIP_VISIBLE_DEVICES)",
kind: EnvKind::U64,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROPE_DEBUG",
group: "misc",
summary: "Print RoPE table strides and rotation indices",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_SDPA_THRESHOLD",
group: "cpu",
summary: "Sequence length above which SDPA switches from NEON dots to BLAS sgemm",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_SOFT_MEMORY_BUDGET_BYTES",
group: "misc",
summary: "Soft RAM budget in bytes, used when RLX_MAX_RAM_BYTES is unset",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_SOFT_MEMORY_FRACTION",
group: "misc",
summary: "Fraction of physical RAM used as the soft memory budget",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_SPD_JACOBI_SWEEPS",
group: "misc",
summary: "Jacobi sweep count for the SPD eigensolver",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_SPD_UNROLL",
group: "misc",
summary: "Use the natively unrolled SPD eigensolver rounds",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_STATIC_WEIGHT_PACK",
group: "gpu",
summary: "Materialise step-invariant weight packs once and skip after; `=0` opts out",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_TPU_HLO_DUMP",
group: "tpu",
summary: "Dump generated HLO for inspection",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("tpu"),
},
EnvVarEntry {
name: "RLX_TRACE_PERFETTO",
group: "misc",
summary: "Write a Perfetto trace to this path",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_TRACE_THUNK",
group: "misc",
summary: "Print each CPU thunk as it executes",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_UMAP_CUDA_FUSED_KNN",
group: "misc",
summary: "Use the fused GPU kNN path in UMAP",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_USE_ICB",
group: "misc",
summary: "Opt in to Metal indirect command buffers",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_VMATH_ACCURATE",
group: "cpu",
summary: "CPU vmath exp/tanh/log/sqrt: use Accelerate/libm accurate path (default 0 = SIMD *_fast)",
kind: EnvKind::BoolOr { default: false },
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_VERBOSE",
group: "debug",
summary: "Extra runtime logging",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_VULKAN_ARENA_DEBUG",
group: "vulkan",
summary: "Report Vulkan arena allocation and slot reuse",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_CHECK_BCAST",
group: "vulkan",
summary: "Verify broadcast strides before each Vulkan dispatch",
kind: EnvKind::String,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_CHECK_CAST",
group: "vulkan",
summary: "Verify dtype casts during Vulkan pipeline build",
kind: EnvKind::String,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_DEBUG",
group: "vulkan",
summary: "Extra Vulkan diagnostics and object labelling",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_DUMP_OPS",
group: "vulkan",
summary: "Dump the Vulkan dispatch schedule",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_FULLBARRIER",
group: "vulkan",
summary: "Insert a full barrier between every pair of dispatches",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_HOST_CONV",
group: "vulkan",
summary: "Run convolution on the host instead of a Vulkan kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_HOST_OPS",
group: "vulkan",
summary: "Comma-separated op classes forced onto the host, or all",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_INDEXING_HOST",
group: "vulkan",
summary: "Force ND indexing (GatherND/GatherElements/Scatter*) to the CPU host route",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_MATMUL",
group: "vulkan",
summary: "Vulkan matmul kernel: auto (default) | scalar | tiled",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_NOBARRIER",
group: "vulkan",
summary: "Drop all pipeline barriers (unsafe; for isolating sync bugs)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_NOCACHE",
group: "vulkan",
summary: "Disable the command-buffer cache; re-record every run",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_SCAN_ALL",
group: "vulkan",
summary: "Scan every buffer for NaN/Inf after each run",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_SCHEDULE_MATMUL",
group: "vulkan",
summary: "Source for the GLSL matmul kernel: schedule-IR emitted vs handwritten",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_SHARD_LOG",
group: "vulkan",
summary: "Log Vulkan buffer sharding decisions",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_VALIDATION",
group: "vulkan",
summary: "Enable the Khronos validation layer",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_WGPU_ACTIVE_TRACE",
group: "wgpu",
summary: "Trace each wgpu dispatch as it is submitted",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_ALLOW_SHARD",
group: "wgpu",
summary: "Warn instead of refusing when the activation arena would be striped across buffers (wrong results)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_ARENA_REPORT",
group: "wgpu",
summary: "Report which tensors dominate the activation arena after planning",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_CONCAT_HOST",
group: "wgpu",
summary: "Run Concat on the host instead of a wgpu kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_CONV3D_SCALAR",
group: "wgpu",
summary: "Force the one-element-per-thread Conv3D kernel instead of the 4-channel tiled one",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_CONV_IM2COL",
group: "wgpu",
summary: "Use the im2col convolution path on wgpu",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_COOP_F16_VK_DISABLE",
group: "wgpu",
summary: "Disable f16 cooperative-matrix matmul on Vulkan",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_COOP_F16_VK_ENABLE",
group: "wgpu",
summary: "Re-enable f16 cooperative-matrix matmul on Vulkan (at your risk)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_COOP_F16_VK_FORCE_WIDE",
group: "wgpu",
summary: "Force the wide cooperative-matrix tile",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_COOP_F16_VK_LARGE_N",
group: "wgpu",
summary: "Keep cooperative tensor cores at large N instead of falling back",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_COOP_F16_VK_LOAD_T",
group: "wgpu",
summary: "Use coopLoadT on B at N > 768 instead of coopLoad",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_COOP_F16_VK_NO_AUTO_WIDE",
group: "wgpu",
summary: "Disable automatic selection of the wide cooperative tile",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_COOP_F16_VK_NO_F32ACC",
group: "wgpu",
summary: "Accumulate cooperative f16 matmul in f16 instead of f32",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_COOP_F16_VK_OSC_THRESH",
group: "wgpu",
summary: "Oscillation threshold for cooperative tile selection (default 0.35)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_DEBUG",
group: "wgpu",
summary: "Extra wgpu diagnostics and validation",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_DEBUG_ATTN_ALIAS",
group: "wgpu",
summary: "Report when an attention output aliases its Q/K/V input",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_DEBUG_ATTN_MASK",
group: "wgpu",
summary: "Report the attention mask shape and kind chosen",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_DEBUG_OP_HIST",
group: "wgpu",
summary: "Print a histogram of op kinds after compilation",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_DUMP_FLAT",
group: "wgpu",
summary: "Also print this flat element index when dumping node stats",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_DUMP_IDS",
group: "wgpu",
summary: "Comma-separated node ids to dump",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_DUMP_INPUTS",
group: "wgpu",
summary: "Dump graph inputs before running",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_DUMP_NODES",
group: "wgpu",
summary: "Dump per-node output buffers for cross-backend diffing",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_DUMP_NODES_LIMIT",
group: "wgpu",
summary: "Cap how many nodes the wgpu dump writes (default 40)",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_DUMP_TAIL",
group: "wgpu",
summary: "Dump the tail of each buffer rather than the head",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_F16_WEIGHTS",
group: "wgpu",
summary: "Store wgpu weights in f16 to halve weight bandwidth",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_FORCE_COOP_F32",
group: "wgpu",
summary: "Force f32 cooperative matrices (unreliable on RTX, which lacks 8x8 f32)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_FORCE_INPUT_UPLOAD",
group: "wgpu",
summary: "Re-upload graph inputs every run instead of reusing device copies",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_GDN_HOST",
group: "wgpu",
summary: "Force GatedDeltaNet host fallback on wgpu (skip WGSL)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_INDEXING_HOST",
group: "wgpu",
summary: "Force ND indexing (GatherND/GatherElements/Scatter*) to the CPU host route",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_HOST_BUFFER_COPY",
group: "wgpu",
summary: "Route buffer-to-buffer copies through the host",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_IM2COL_MIN_COUT",
group: "wgpu",
summary: "Minimum output channels before im2col is used (default 64)",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_IM2COL_MIN_K",
group: "wgpu",
summary: "Minimum K before im2col is used (default 256)",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_IM2COL_MIN_SPATIAL",
group: "wgpu",
summary: "Minimum spatial extent before im2col is used (default 2048)",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_LARGE_BUFFERS",
group: "wgpu",
summary: "Allow buffers past the conservative wgpu size limit",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_LSTM_WIDE",
group: "gpu",
summary: "Re-enable the wide-hidden native WGSL LSTM (debugging the Apple exp defect)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("gpu"),
},
EnvVarEntry {
name: "RLX_WGPU_MATMUL_F32_ONLY",
group: "wgpu",
summary: "Restrict wgpu matmul to f32, disabling f16 paths",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_NAN_TRACE",
group: "wgpu",
summary: "Legacy wgpu NaN trace (prefer RLX_DEBUG_NANS)",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_NO_COOP_F16_VK",
group: "wgpu",
summary: "Disable f16 cooperative matrices in the wgpu Vulkan backend",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_NO_COOP_F32",
group: "wgpu",
summary: "Disable f32 cooperative matrices",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_NO_F16_MIRROR",
group: "wgpu",
summary: "Do not keep an f16 mirror of f32 weights",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_NO_F16_SHADOW",
group: "wgpu",
summary: "Do not allocate the f16 shadow buffer",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_NO_PACKED_BSHD_ATTN",
group: "wgpu",
summary: "Disable the packed BSHD attention kernel on wgpu",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_NO_TILED_CONV",
group: "wgpu",
summary: "Disable the tiled convolution kernel",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_PRINT_LIMITS",
group: "wgpu",
summary: "Print the adapter's reported limits at startup",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_Q1_0_GEMM_DISABLE",
group: "wgpu",
summary: "Disable the Q1_0 GEMM kernel on wgpu",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_SCHEDULE",
group: "wgpu",
summary: "Print the wgpu execution schedule",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_SHARD_CAP_MIB",
group: "wgpu",
summary: "Lower the per-shard activation-arena cap below the adapter's max_buffer_size (clamps smaller only) so the striping path is testable without multi-GiB allocations",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_SHARD_LOG",
group: "wgpu",
summary: "Log wgpu buffer sharding decisions",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_SHARE_WEIGHTS",
group: "wgpu",
summary: "Reuse GPU weight buffers across compiles with matching named layouts (default on; set 0 to disable)",
kind: EnvKind::Bool,
stability: EnvStability::Public,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_STRIPE_VERBOSE",
group: "wgpu",
summary: "Print every op whose operands straddle an arena stripe boundary",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_TEE_VETO_LOG",
group: "wgpu",
summary: "Log residual-LN tee folds vetoed as arena-unsafe",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_TILED_MIN_SPATIAL",
group: "wgpu",
summary: "Minimum spatial extent before the tiled conv path is used",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WINOGRAD",
group: "misc",
summary: "Opt in to the Winograd convolution path",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_WORKERS",
group: "cpu",
summary: "Thread-pool size (0 = auto)",
kind: EnvKind::U64,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_DISABLE_METAL_DEQUANT_GPU",
group: "metal",
summary: "Deprecated alias of `RLX_METAL_DEQUANT_GPU_DISABLE`",
kind: EnvKind::Bool,
stability: EnvStability::Deprecated { replace_with: "RLX_METAL_DEQUANT_GPU_DISABLE" },
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_XDNA_SKIP_PROBE",
group: "xdna",
summary: "Trust the toolchain paths instead of compiling a probe design during device selection",
kind: EnvKind::Bool,
stability: EnvStability::Bisect,
aliases: &[],
layer: EnvLayer::Backend("xdna"),
},
EnvVarEntry {
name: "RLX_AD_FUSED_MMBA_VJP",
group: "misc",
summary: "Opt in to the fused matmul+bias+activation VJP in autodiff",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_ARENA_CHECK",
group: "cpu",
summary: "Verify CPU arena plan invariants when building from a plan",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_ARENA_PIN_HOST_STRUCTURE",
group: "compile",
summary: "Pin host-visible structure across planner entry points",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_ATTN_BWD_SERIAL",
group: "cpu",
summary: "Run CPU attention backward serially instead of with Rayon",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_BAKE_OUT",
group: "misc",
summary: "Output path for the baked .rlx artifact",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_BAKE_PASSWORD",
group: "misc",
summary: "Password for an encrypted bake (never passed on the CLI)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_BLAS_LINK",
group: "cpu",
summary: "Escape hatch naming the BLAS/LAPACK libraries to link",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_BLAS_SEARCH",
group: "cpu",
summary: "Extra directories searched for the RLX_BLAS_LINK libraries",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_CONVINT_FLOAT_ACT",
group: "misc",
summary: "Route ConvInteger through the pre-quant float activation",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_COREML_DEBUG_IO",
group: "coreml",
summary: "Dump CoreML model inputs and outputs at finalize",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("coreml"),
},
EnvVarEntry {
name: "RLX_CPU_ARENA_REPORT",
group: "cpu",
summary: "Attribute the CPU arena's resident bytes to the nodes holding them",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_CPU_ATTN_BWD_FUSE",
group: "cpu",
summary: "Opt in to the fused CPU attention-backward path (default off)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_CPU_BNNS_BF16",
group: "cpu",
summary: "Opt in to the BNNS bf16 GEMM path (lossy f32 to bf16 downcast)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_CPU_BNNS_F16",
group: "cpu",
summary: "Opt in to the BNNS f16 GEMM path (half bandwidth)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_CPU_SME",
group: "cpu",
summary: "Opt in to the direct ARM SME2 GEMM path (Apple M4+)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_CPU_SME_BF16",
group: "cpu",
summary: "Opt in to the native SME bf16 path (Apple M4+)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_CPU_SME_W8A8",
group: "cpu",
summary: "Opt in to the wired SME W8A8 integer path (Apple M4+)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_CUDA_ATTN_RANK4",
group: "cuda",
summary: "Force the rank-4 CUDA attention-backward promotion",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_CAPTURE_DEBUG",
group: "cuda",
summary: "Check CUDA graph capture state once per step after dispatch",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_CONV_FORCE_GATHER",
group: "cuda",
summary: "Force the gather-based conv path instead of cuDNN",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_CONV_TRACE",
group: "cuda",
summary: "Trace which convolution path each conv takes",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_CONV_T_KERNEL",
group: "cuda",
summary: "Opt out of cuDNN for transposed conv; use the rlx kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_DUMP_IO",
group: "cuda",
summary: "Also dump Input/Param/Constant slots when dumping nodes",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_DYN_LSTM_HOST",
group: "cuda",
summary: "Run the dynamically-quantized LSTM on the host",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_DYN_LSTM_TRACE",
group: "cuda",
summary: "Trace the dynamically-quantized LSTM path",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_INDEXING_TRACE",
group: "cuda",
summary: "Trace indexing-op dispatch",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_INPUT_DIAG",
group: "cuda",
summary: "Print diagnostics for each graph input before the run",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_KDA_CHUNK",
group: "cuda",
summary: "Enable the chunked Kimi delta-attention kernel (read at launch)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_LEGACY_ATTENTION_ROW",
group: "cuda",
summary: "Use the legacy thread-per-row attention kernel instead of warp-per-row",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_LSTM_CUBLAS_TRACE",
group: "cuda",
summary: "Trace the cuBLAS LSTM gemm path",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_LSTM_CUDNN",
group: "cuda",
summary: "Use cuDNN for LSTM inference; set 0 to opt out",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_LSTM_CUDNN_TRACE",
group: "cuda",
summary: "Trace cuDNN LSTM descriptor setup",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_LSTM_DEBUG",
group: "cuda",
summary: "Print LSTM shapes and gate values during execution",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_MATMUL_PRECISE_MIN_K",
group: "cuda",
summary: "K above which RLX_CUDA_MATMUL_PRECISE compensation engages",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_NO_BCAST_BINARY",
group: "cuda",
summary: "Disable the broadcasting binary kernel; materialize instead",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_Q4K_GEMV_COOP",
group: "cuda",
summary: "Opt in to the cooperative block-per-row Q4_K GEMV",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_RNN_HOST_FALLBACK",
group: "cuda",
summary: "Force RNN ops onto the host D2H-CPU-H2D path",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_SEGMENTED_CAPTURE",
group: "cuda",
summary: "Opt in to segmented CUDA graph capture",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_SEGMENTED_CAPTURE_ENGAGE",
group: "cuda",
summary: "Actually replay segmented captures (needs RLX_CUDA_SEGMENTED_CAPTURE)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_SHARED_PARAMS",
group: "cuda",
summary: "Back shared parameters with one CUDA VMM region to cut HtoD traffic",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_SSM_HOST_FALLBACK",
group: "cuda",
summary: "Force SSM ops onto the host D2H-CPU-H2D path",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_STEP_PROFILE",
group: "cuda",
summary: "Print per-op ms/call for each execution step",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_TRANSPOSE_SHAPES",
group: "cuda",
summary: "Log the shapes of every transpose dispatched",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_UNIFIED",
group: "cuda",
summary: "Use CUDA unified memory so models larger than VRAM page over PCIe",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_CUDA_WHOLE_GRAPH_CAPTURE",
group: "cuda",
summary: "Capture the whole graph as one CUDA graph (needs RLX_CUDA_EXEC_MODE=graph)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_DEQUANT_CACHE_MAX_BYTES",
group: "cpu",
summary: "Total byte budget for the dequant cache",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_DEQUANT_PARALLEL",
group: "cpu",
summary: "Dequantize GGUF blocks in parallel; set 0 to serialize",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_DIM_DBG",
group: "misc",
summary: "Print dimension resolution during ONNX runtime emission",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_DISABLE_CSE",
group: "misc",
summary: "Disable common-subexpression elimination for an A/B",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_DISABLE_MPSGRAPH_HYBRID",
group: "metal",
summary: "Disable hybrid MPSGraph/native segmentation",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_DISABLE_TRANSPOSE_ELISION",
group: "misc",
summary: "Keep redundant transposes instead of eliding them (A/B + safety valve)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_DUMP_KERNELS",
group: "rocm",
summary: "Snapshot generated GPU kernel sources into this directory",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_DUMP_SCHED_DECODE",
group: "metal",
summary: "Dump the encode schedule for decode steps only",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_EGPU_CONFIG",
group: "egpu",
summary: "Path to the eGPU config file (overrides XDG_CONFIG_HOME/rlx/egpu.conf)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("egpu"),
},
EnvVarEntry {
name: "RLX_EGPU_SERVICE",
group: "egpu",
summary: "eGPU service endpoint used instead of the assumed default",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("egpu"),
},
EnvVarEntry {
name: "RLX_FW_DIR",
group: "egpu",
summary: "Where firmware is written (overrides XDG_CACHE_HOME/rlx/firmware)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("egpu"),
},
EnvVarEntry {
name: "RLX_GPU_VALIDATE_PARAMS",
group: "rocm",
summary: "Check each launch's argument count against the kernel signature",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_HD_PROFILE",
group: "gpu-host",
summary: "Accumulate nanoseconds spent in each host-delegate phase",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("gpu-host"),
},
EnvVarEntry {
name: "RLX_HF_CACHE",
group: "misc",
summary: "HuggingFace cache directory (default $HOME/.cache/rlx/hf)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_HF_MLX",
group: "metal",
summary: "Allow downloading MLX weights from HuggingFace on a cold cache",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_HF_MLX_REPO",
group: "misc",
summary: "HuggingFace repo id to pull MLX weights from",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_HF_REVISION",
group: "misc",
summary: "HuggingFace revision to fetch (default main)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_HIPCC",
group: "egpu",
summary: "Path to the hipcc compiler (default hipcc)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("egpu"),
},
EnvVarEntry {
name: "RLX_INSPECT_BINS",
group: "compile",
summary: "Histogram bin count for the op tap (default 24)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_INSPECT_OPS",
group: "compile",
summary: "Tag every op output for the inspection tap",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_KITTEN_IF_STUB_META",
group: "misc",
summary: "Stub the KittenTTS If-branch metadata (debug aid)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_KITTEN_INORM_ACTIVE",
group: "misc",
summary: "Treat rank-3 channel-axis-1 norms as KittenTTS instance norm",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_KVSTORE_PREAD",
group: "misc",
summary: "Read the KV store with pread instead of mmap (opt-in)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_KVSTORE_READ_STATS",
group: "misc",
summary: "Report KV store batched-read statistics",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_MADV",
group: "misc",
summary: "madvise policy for arena pages, trading footprint against refault cost",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_MAX_RAM_BYTES",
group: "misc",
summary: "Hard RAM budget in bytes (else RLX_SOFT_MEMORY_BUDGET_BYTES or physical RAM)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_MAX_RESIDENT_EXPERTS",
group: "misc",
summary: "Max resident MoE experts per layer (else derived from the RAM budget)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_MEM_VERIFY",
group: "compile",
summary: "Run the memory-planner invariant self-check",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_METAL_ARENA_DIAG",
group: "metal",
summary: "Report Metal arena slot assignment and reuse",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_ATTN_BWD_FUSE",
group: "metal",
summary: "Fuse Metal attention backward; set 0 to opt out",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_ATTN_BWD_FUSED",
group: "metal",
summary: "Use the fused Metal attention-backward kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_ATTN_BWD_FUSED_ZEROONLY",
group: "metal",
summary: "Run only the zeroing half of the fused attention backward (isolation aid)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_CONCURRENT_NOBARRIER",
group: "metal",
summary: "Drop barriers between concurrent Metal dispatches",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_DEQUANT_FORCE_GEMV",
group: "metal",
summary: "Force the GEMV dequant kernel even when m > 1",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_DUMP_BYTES",
group: "metal",
summary: "Sum bytes moved per op type (cache excluded)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_DUMP_MSL",
group: "metal",
summary: "Write every generated MSL source to this path",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_DW_SUM",
group: "metal",
summary: "Correctly-rounded near-f64 dW reduction, at ~3-4x the flops",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_FUSE_RESIDUAL_DUAL",
group: "metal",
summary: "Opt in to the dual-output fused residual",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_GDN_SG_DISABLE",
group: "metal",
summary: "Disable the simdgroup GatedDeltaNet kernel (state size 128)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_GEMV_KPART",
group: "metal",
summary: "K-partition small-N GEMVs across threadgroups (ceiling probe)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_GEMV_SPLITK",
group: "metal",
summary: "Split-K GEMV for m=1, n>=64; set 0 to disable",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_GROUPED_GEMV_DISABLE",
group: "metal",
summary: "Disable the grouped GEMV kernel and fall back to MPS",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_KERNEL_STATS",
group: "metal",
summary: "Report per-kernel dispatch statistics",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_LN_GAMMA_SIMD",
group: "metal",
summary: "Use the simdgroup reduction for LayerNorm gamma backward",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_LSTM_NATIVE_WIDE",
group: "metal",
summary: "Allow the native LSTM kernel at wide hidden sizes",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_MATMUL_TRANSPOSE_FOLD",
group: "metal",
summary: "Opt in to folding transposes into Metal matmul (default off)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_MATMUL_TRANSPOSE_FOLD_WEIGHTS",
group: "metal",
summary: "Fold transposes into weight operands of Metal matmul",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_MLP_SG",
group: "metal",
summary: "Simdgroup MLP kernel; set 0 to disable",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_NO_AUTO_WIDE_GEMM",
group: "metal",
summary: "Disable automatic selection of the wide GEMM path",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_PARAM_DIAG",
group: "metal",
summary: "Diagnose parameter uploads larger than 1 MB",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_PREFILL_COUNTERS",
group: "metal",
summary: "Report Metal prefill dispatch counters",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_PREFILL_FA",
group: "metal",
summary: "Flash-attention prefill kernel; set 0 to disable",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_PREFILL_FA_MMA",
group: "metal",
summary: "Opt in to the MMA flash-attention prefill kernel",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_PREFILL_HD256",
group: "metal",
summary: "Head-dim-256 prefill kernel; set 0 to disable",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_PREFILL_TRACE",
group: "metal",
summary: "Trace Metal prefill matmul dispatches (m > 1)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q2K_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused Q2_K matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q3K_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused Q3_K matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q40_SG_DISABLE",
group: "metal",
summary: "Disable the simdgroup Q4_0 kernels; use one thread per row",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q4K_GEMM_MAX_M",
group: "metal",
summary: "m above which Q4_K routes to GEMM instead of GEMV (default 32)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q4K_GEMM_XS_DISABLE",
group: "metal",
summary: "Disable the extra-small-tile Q4_K GEMM kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q5K_FUSED_DISABLE",
group: "metal",
summary: "Disable the fused Q5_K matmul kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q5K_GEMM_DISABLE",
group: "metal",
summary: "Disable the Q5_K GEMM kernel; use GEMV",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q6K_FUSED_DISABLE",
group: "metal",
summary: "Disable fused Q6_K matmul, falling back to dequant-to-scratch",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q6K_GEMM_XS_DISABLE",
group: "metal",
summary: "Disable the extra-small-tile Q6_K GEMM kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q6K_SG_DISABLE",
group: "metal",
summary: "Disable the simdgroup Q6_K/Q8_0 decode kernels",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q8_0_GEMM_DISABLE",
group: "metal",
summary: "Disable the Q8_0 GEMM kernel; use GEMV",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_Q8_0_SG_DISABLE",
group: "metal",
summary: "Disable the simdgroup Q8_0 kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_RB_DEBUG",
group: "metal",
summary: "Log register-blocked GEMM selection",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_RB_FUSED_GEMM",
group: "metal",
summary: "Use the register-blocked fused GEMM kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_REDUCE_SIMD",
group: "metal",
summary: "Use the simdgroup reduction kernel for axis reductions",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SDPA_FA2",
group: "metal",
summary: "Use the FlashAttention-2 style SDPA kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SDPA_H16",
group: "metal",
summary: "Opt in to the f16-scores SDPA kernel (implies SIMD softmax)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SDPA_HDSPLIT",
group: "metal",
summary: "Opt in to the head-dim-split SDPA variant",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SDPA_MMA",
group: "metal",
summary: "Use the MMA (simdgroup matrix) SDPA kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SDPA_OCCPAD",
group: "metal",
summary: "Pad SDPA dispatches for occupancy at seq > 1",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SDPA_SIMD",
group: "metal",
summary: "Use the simdgroup softmax SDPA kernel (f32 only)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_SDPA_SPLITK",
group: "metal",
summary: "Split SDPA along K across threadgroups",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_UNPIN_ALL",
group: "metal",
summary: "Do not pin output-ancestor buffers",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_VALIDATE_BINDINGS",
group: "metal",
summary: "Check each Metal dispatch's buffer bindings against the kernel signature",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_W8A8_ATTN",
group: "metal",
summary: "W8A8 decode attention: int8 Q-K integer dot with int8 V",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_W8A8_BLOCK",
group: "metal",
summary: "Use block-wise W8A8 quantization scales",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_W8A8_INCR",
group: "metal",
summary: "Token count for the incremental-quantize timing probe",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_W8A8_KMODE",
group: "metal",
summary: "W8A8 K source: i8 (default) or f32 for diagnostic isolation",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_W8A8_QMODE",
group: "metal",
summary: "W8A8 Q source: i8 integer dot (default) or f32",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_W8A8_VMODE",
group: "metal",
summary: "W8A8 V source: i8 (default) or f32 to isolate K error",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_METAL_WIDE_GEMM",
group: "metal",
summary: "Force the wide GEMM path wherever M makes it eligible",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_MLX_KEEP_WARM",
group: "mlx",
summary: "Skip the per-tensor WILLNEED/DONTNEED madvise on MLX loads",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MODELS_DIR",
group: "misc",
summary: "Directory searched for test model weights",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_MPSGRAPH_NO_SYNC_COMPILE",
group: "metal",
summary: "Restore the old nil-descriptor asynchronous MPSGraph compile",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_MPS_LOWP_ACT",
group: "metal",
summary: "Low-precision MPS activations: less bandwidth, some accuracy cost",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_NO_WEIGHT_CONCAT_FUSION",
group: "compile",
summary: "Skip the fusions that concatenate weights into one buffer",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_NVRTC",
group: "egpu",
summary: "Path to the NVRTC library for ahead-of-time baking",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("egpu"),
},
EnvVarEntry {
name: "RLX_PRECISION",
group: "misc",
summary: "Mixed-precision policy for the graph (default F32)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_PROBE_DEVICE",
group: "misc",
summary: "Device for the probe examples: cpu (default) | cuda | rocm | metal | gpu",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_PTXAS",
group: "egpu",
summary: "Path to ptxas for ahead-of-time baking (default ptxas)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("egpu"),
},
EnvVarEntry {
name: "RLX_Q4K_FUSED_MIN_N",
group: "cpu",
summary: "n above which the fused Q4_K kernel is used instead of dequant",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_Q4K_NO_DOTPROD",
group: "cpu",
summary: "Force the baseline Q4_K path instead of the dotprod intrinsics",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_QWEN3_BAKE_WEIGHTS",
group: "metal",
summary: "Legacy Metal-only alias for RLX_STATIC_WEIGHT_PACK (default on; `=0` opts out)",
kind: EnvKind::BoolOr { default: true },
stability: EnvStability::Deprecated {
replace_with: "RLX_STATIC_WEIGHT_PACK",
},
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_QWEN3_F16_KV",
group: "misc",
summary: "Keep the Qwen3 KV cache resident in f16",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_QWEN3_F16_WEIGHTS",
group: "misc",
summary: "Store Qwen3 weights in f16 to halve decode weight bandwidth",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_QWEN3_FUSED_QKV",
group: "misc",
summary: "Fuse Q/K/V into one DequantMatMul (packed weights only)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_QWEN3_GQA_NATIVE",
group: "misc",
summary: "Index K/V natively in the SDPA kernels instead of expanding for GQA",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_QWEN3_NO_INPLACE_KV",
group: "misc",
summary: "Concat the new KV row instead of appending it in place (A/B lever)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_RANK",
group: "misc",
summary: "This process's rank in a distributed run",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_ROCM_ARCH",
group: "rocm",
summary: "Explicit ROCm target arch override, e.g. gfx1100",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_ARENA_DUMP",
group: "rocm",
summary: "Dump the ROCm arena layout",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_ARENA_NO_REUSE",
group: "rocm",
summary: "Give every ROCm tensor one permanent arena slot; disables reuse",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_DISABLE_MIOPEN",
group: "rocm",
summary: "Disable MIOpen whether or not its kernel database is present",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_EXEC",
group: "rocm",
summary: "ROCm execution mode: stream | graph | ms<N>",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_FAST_FP_ATOMICS",
group: "rocm",
summary: "Compile HIP kernels with unsafe fast fp atomics",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_GEMV",
group: "rocm",
summary: "Opt in to the skinny-m split-K ROCm GEMV kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_GGUF_DIAG",
group: "rocm",
summary: "Diagnose GGUF dequant-matmul dispatch on ROCm",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_GGUF_FUSED_DISABLE",
group: "rocm",
summary: "Disable the fused GGUF matmul kernel on ROCm",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_GGUF_HOST",
group: "rocm",
summary: "Dequantize GGUF weights on the CPU instead of on device",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_LEGACY_ATTENTION_ROW",
group: "rocm",
summary: "Use the legacy thread-per-row attention kernel instead of warp-per-row",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_MIOPEN_FORCE",
group: "rocm",
summary: "Use MIOpen even with no kernel database present (for measuring)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_NO_TF32",
group: "rocm",
summary: "Disable XF32/TF32 on ROCm and compute in true f32",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_NO_VENDOR_GEMM",
group: "rocm",
summary: "Use rlx's own GEMM instead of rocBLAS",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_PARAM_DIAG",
group: "rocm",
summary: "Diagnose parameter uploads on ROCm",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_PARITY",
group: "rocm",
summary: "Tighten ROCm numerics for cross-backend parity",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_PROFILE_STEPS",
group: "rocm",
summary: "Attribute wall time to each op kind by syncing between steps",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_RNN_HOST_FALLBACK",
group: "rocm",
summary: "Force RNN ops onto the host on ROCm",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_SSM_HOST_FALLBACK",
group: "rocm",
summary: "Force SSM ops onto the host on ROCm",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_ROCM_TRACE_STEPS",
group: "rocm",
summary: "Print each step synchronously so a crashing kernel is identifiable",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("rocm"),
},
EnvVarEntry {
name: "RLX_SDPA_CASE",
group: "metal",
summary: "Case selector for the SDPA canonicalization crash test",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("metal"),
},
EnvVarEntry {
name: "RLX_SELSCAN_LEGACY_UNROLL",
group: "misc",
summary: "Use the legacy unrolled selective-scan VJP for an A/B",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Tooling,
},
EnvVarEntry {
name: "RLX_TEST_DEVICE",
group: "misc",
summary: "Device the cross-backend tests run against",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_TPU_BENCH",
group: "tpu",
summary: "Run the TPU benchmarks (off during a normal cargo test)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("tpu"),
},
EnvVarEntry {
name: "RLX_TPU_BENCH_SWEEP",
group: "tpu",
summary: "Run the slower TPU benchmark sweep",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("tpu"),
},
EnvVarEntry {
name: "RLX_TPU_MATMUL_HIGHEST",
group: "tpu",
summary: "Pin TPU matmul to highest precision instead of bf16 passes",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("tpu"),
},
EnvVarEntry {
name: "RLX_TUNE_PREFILTER",
group: "cuda",
summary: "Narrow the tuner's candidate tiles to the N cheapest by cost model",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cuda"),
},
EnvVarEntry {
name: "RLX_VULKAN_MLX_NATIVE",
group: "vulkan",
summary: "Route dense MLX matmuls back onto the SPIR-V kernels",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_VULKAN_SHARD_STAGE_MIB",
group: "vulkan",
summary: "Staging-window size in MiB for sharded Vulkan uploads",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("vulkan"),
},
EnvVarEntry {
name: "RLX_WGPU_ALL_PARAMS_WEIGHT",
group: "wgpu",
summary: "Park every parameter in the weight arena rather than splitting",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_CONV_HOST",
group: "wgpu",
summary: "Run convolution on the host instead of a wgpu kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_DBG_HOST_OP",
group: "wgpu",
summary: "Log each op routed to the host fallback",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_DBG_INT8_HOST",
group: "wgpu",
summary: "Log the host int8 dequant-matmul path",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_DISCRETE_HOST",
group: "wgpu",
summary: "Force full host execution on discrete adapters",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_EXPAND_HOST",
group: "wgpu",
summary: "Run Expand on the host instead of a wgpu kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_FORCE_HOST",
group: "wgpu",
summary: "Host every op, reproducing discrete-adapter behaviour on any adapter",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_GATHER_SPLIT",
group: "wgpu",
summary: "Split Gather into per-slice dispatches",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_GPU_NORM",
group: "wgpu",
summary: "Force norms onto the GPU (opposite of RLX_WGPU_HOST_NORM)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_HOST_EAGER_H2D",
group: "wgpu",
summary: "Flush deferred host-to-device copies eagerly rather than batching",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_HOST_MATMUL",
group: "wgpu",
summary: "Run matmul on the host instead of a wgpu kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_HOST_NORM",
group: "wgpu",
summary: "Force norms onto the host",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_MAX_BIND_MB",
group: "wgpu",
summary: "Cap the bind-group size in MiB to exercise the sharding path",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_NO_FOLD_MMBA",
group: "wgpu",
summary: "Disable folding matmul+bias+activation into one node",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_NO_FOLD_RESLN",
group: "wgpu",
summary: "Disable folding residual add into LayerNorm",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_ONE_OP_PER_PASS",
group: "wgpu",
summary: "Submit one op per compute pass (isolates sync bugs)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_QUIET_SHARD",
group: "wgpu",
summary: "Suppress the buffer-sharding log",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_SHARD_GPU",
group: "wgpu",
summary: "Keep sharded ops on the GPU instead of hosting them",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_SHARD_STAGE_MIB",
group: "wgpu",
summary: "Staging-window size in MiB for sharded wgpu uploads",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_TRANSPOSE_HOST",
group: "wgpu",
summary: "Run Transpose on the host instead of a wgpu kernel",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WG_LAYERS",
group: "misc",
summary: "Layer count for the whole-graph capture probe",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_WORKER_ERR_DIR",
group: "misc",
summary: "Directory distributed workers write captured errors to",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Runtime,
},
EnvVarEntry {
name: "RLX_XDNA_AIE_INCLUDE",
group: "xdna",
summary: "Include directory for the MLIR-AIE toolchain",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("xdna"),
},
EnvVarEntry {
name: "RLX_XDNA_CHAIN_CAP",
group: "xdna",
summary: "Max ops chained into one XDNA submission",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("xdna"),
},
EnvVarEntry {
name: "RLX_XDNA_GEMM",
group: "xdna",
summary: "GEMM dimensions for the XDNA example, as m,k,n",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("xdna"),
},
EnvVarEntry {
name: "RLX_XDNA_NO_RESIDENT",
group: "xdna",
summary: "Re-upload weights each call instead of keeping them resident on the NPU",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("xdna"),
},
EnvVarEntry {
name: "RLX_XDNA_SHIM",
group: "xdna",
summary: "Path to the XDNA shim library",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("xdna"),
},
EnvVarEntry {
name: "RLX_XDNA_TURBO",
group: "xdna",
summary: "Request NPU turbo clocks (needs root/DRM-master)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("xdna"),
},
EnvVarEntry {
name: "RLX_XDNA_UMQ",
group: "xdna",
summary: "Request a user-mode queue with a ring buffer object",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("xdna"),
},
EnvVarEntry {
name: "RLX_XDNA_UMQ_RING",
group: "xdna",
summary: "Ring-buffer size in bytes for the XDNA user-mode queue",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("xdna"),
},
EnvVarEntry {
name: "RLX_XDNA_XRT_LIB",
group: "xdna",
summary: "Explicit path to the XRT library (which ships only versioned sonames)",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("xdna"),
},
EnvVarEntry {
name: "RLX_MLX_GROUPED_ONDEVICE",
group: "mlx",
summary: "Keep grouped-MoE gather on device instead of host-staging (experiment)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_MLX_GROUPED_ONDEVICE_MAX_BYTES",
group: "mlx",
summary: "Byte budget guarding RLX_MLX_GROUPED_ONDEVICE (default 2e9)",
kind: EnvKind::U64,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("mlx"),
},
EnvVarEntry {
name: "RLX_COREML_DBG_CONV",
group: "coreml",
summary: "Trace CoreML conv placement decisions",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("coreml"),
},
EnvVarEntry {
name: "RLX_COREML_HOST_RESHAPED_REDUCE",
group: "coreml",
summary: "Host-execute Reduce over a reshaped Concat (MIL miscompiles it)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("coreml"),
},
EnvVarEntry {
name: "RLX_COREML_MAX_CONV_KERNEL",
group: "coreml",
summary: "Widest conv kernel left on MIL; wider convs go host",
kind: EnvKind::U64,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("coreml"),
},
EnvVarEntry {
name: "RLX_COREML_MAX_SCAN_LEN",
group: "coreml",
summary: "Longest SelectiveScan left on MIL before host fallback",
kind: EnvKind::U64,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("coreml"),
},
EnvVarEntry {
name: "RLX_CPU_DUMP_DIR",
group: "cpu",
summary: "Directory for RLX_CPU_DUMP_NODE_DATA output",
kind: EnvKind::Path,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_CPU_DUMP_NODE_DATA",
group: "cpu",
summary: "Write listed nodes' raw f32 arena buffers to files",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_CPU_DUMP_THUNKS",
group: "cpu",
summary: "Print compiled thunks in a <lo>:<hi> index range",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_CPU_ROW_TABS",
group: "cpu",
summary: "Per-row sum of |x| for one node id",
kind: EnvKind::U64,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_CPU_WATCH",
group: "cpu",
summary: "Report every thunk that changes <byte_offset>:<n_floats>",
kind: EnvKind::String,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("cpu"),
},
EnvVarEntry {
name: "RLX_DISABLE_GATED_RESIDUAL_FUSION",
group: "compile",
summary: "Keep gated residuals decomposed (fusion ablation)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_DISABLE_MATMUL_BIAS_FUSION",
group: "compile",
summary: "Disable matmul+bias+activation fusion (ablation)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_DISABLE_RESIDUAL_LN_FUSION",
group: "compile",
summary: "Disable residual+LayerNorm fusion (ablation)",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_PIN_OUTPUT_ANCESTORS",
group: "compile",
summary: "Give every output ancestor its own arena slot",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Compile,
},
EnvVarEntry {
name: "RLX_WGPU_DBG_COUNTERS",
group: "wgpu",
summary: "Print step/host/uniform/bind-group counters",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_DBG_TRANSPOSE",
group: "wgpu",
summary: "Trace transpose lowering decisions",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_HAZARD_REPORT",
group: "wgpu",
summary: "Report arena slot read/write hazards after planning",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_ROW_TABS",
group: "wgpu",
summary: "Per-row sum of |x| for one snapshot step",
kind: EnvKind::U64,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_SNAPSHOT",
group: "wgpu",
summary: "Replay the compiled schedule truncated at each step, reading every node as produced",
kind: EnvKind::Bool,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
EnvVarEntry {
name: "RLX_WGPU_SNAPSHOT_WATCH",
group: "wgpu",
summary: "Also report one chosen node id after every snapshot step",
kind: EnvKind::U64,
stability: EnvStability::Internal,
aliases: &[],
layer: EnvLayer::Backend("wgpu"),
},
]