[inureyes]

Summary

Affine 1-bit checkpoints (quantization.bits == 1, e.g. prism-ml/Bonsai-*-mlx-1bit) now load and generate. Before this change, main aborted on the first token of Bonsai-1.7B-mlx-1bit with libc++abi: terminating due to uncaught exception of type std::runtime_error: [metal::Device] Unable to load kernel affine_dequantize_float16_t_gs_128_b_1 (exit 134).

• New cpp/mlx_cxx_one_bit.cpp: Metal qmv kernel (M < 16, factored scale * masked_sum + bias * total_sum), tiled simdgroup qmm kernel (M >= 16, f32 accumulation), and a dequantize graph used as the reference, the non-Metal fallback, and the MLXCEL_ONE_BIT_KERNEL=0 kill switch.
• Routing lives in the bridge primitives (quantized_matmul in both layouts, quantized_linear_forward, _global_scale, dequantize, quantized_embedding), so every dense linear, embedding lookup and tied lm_head is covered, including callers outside layers.rs.
• UnifiedLinear::quantized_weight() / as_quantized_weight() answer None for 1-bit weights (same mechanism as runtime LoRA), so every fused helper that calls MLX's quantized_matmul internally takes its graph fallback. The su-scaled-rope and compiled SwiGLUMLP paths decline explicitly.
• Loader: 1 bit is accepted through a new LOADABLE_AFFINE_BITS; SUPPORTED_AFFINE_BITS stays the producer set (split-mtp, gemma4 overrides). 1-bit without .biases, 1-bit expert stacks (SwitchLinear) and 1-bit MLA (QuantizedMultiLinear) fail at load with the prefix named.
• Docs: supported-models.md quantization row, MLXCEL_ONE_BIT_KERNEL in environment-variables.md.

Differences from the issue text

• The Bonsai export ships f16 scales and biases, not f32. They are kept as shipped.
• forward_fused_rope_append needs no decline: it consumes project_concat, which goes through UnifiedLinear::forward.
• The explicit FFI entry points (one_bit_quantized_matmul, one_bit_dequantize) return Result and are used by tests; the load path reaches the same code through the bridge primitives, which keeps call sites unchanged.
• CUDA qmv port: pending. No CUDA host is reachable (GB10 runner down), and an unexecuted JIT kernel running by default would risk silent wrong output. CUDA and ROCm run the dequantize graph, which the unit tests cover.
• inspect memory estimate needed no change: it reads packed sizes from the safetensors header.

Test plan

• cargo test --release -p mlxcel-core --lib one_bit_tests: 17 pass. Kernels and fallback are each checked against a host f64 oracle (M 1/3/7/16/17/18/20/33/40, K 256 to 4096, N 64 to 1003 incl. unaligned, G 32/64/128, f16/bf16/f32), kill switch is bit-identical to the graph, rejections name the prefix.
• layers::tests (92), models::switch_layers::tests (26), cli_help_consistency, dead_doc_pointers, llama_compat_manifest pass. infer_bits_rejects_non_canonical_widths updated: it asserted that 1 bit is rejected, and now uses 7.
• cargo fmt --all -- --check, cargo clippy --release -p mlxcel-core --lib --tests -- -D warnings, same for -p mlxcel.
• Bonsai-1.7B-mlx-1bit, "The capital of France is", -n 64 -t 0: kernel and MLXCEL_ONE_BIT_KERNEL=0 outputs identical over all 64 tokens, answer "Paris".
• Bonsai-8B-mlx-1bit, chess prompt, -n 128 -t 0: kernel and fallback identical over 128 tokens, fluent English.
• mlxcel-server on the 1.7B, two concurrent greedy chat requests: "Paris is the capital of France." and a correct primary-colors answer.
• Decode throughput vs mlx-community/Qwen3-8B-4bit: deferred. The machine was shared with other GPU jobs during this run, so no timing is reported here.

Closes #1370