SIGSEGV in mlx::core::metal::Device::end_encoding() during LoRA training on M5 Pro

[barrettj]

Environment

• pMetal: 0.4.0 (cargo install)
• macOS: 26.4 (25E243)
• Hardware: Apple M5 Pro, 48GB
• MLX (Python, for reference): 0.31.0

Reproduction

pmetal train \
  --model Qwen/Qwen3-0.6B \
  --dataset forms-train.jsonl \
  --eval-dataset forms-val.jsonl \
  --lora-r 16 --lora-alpha 16 \
  --learning-rate 1e-4 --epochs 2 --batch-size 8 \
  --lr-schedule cosine --warmup-steps 100 \
  --no-sequence-packing \
  --output /tmp/pmetal-test

Dataset is instruction/output format JSONL, ~4,600 records, short sequences (~50-400 tokens).

Behavior

• Training starts fine, throughput ~1,300 tokens/s at step 10-50
• Throughput degrades to ~800 tokens/s by step 60-100
• SIGSEGV (null pointer dereference) at step 100, immediately after eval + checkpoint write
• Exit code 139

Crash log (from macOS DiagnosticReports)

Exception: EXC_BAD_ACCESS (SIGSEGV)
Subtype: KERN_INVALID_ADDRESS at 0x0000000000000000

Faulting thread stack:
  pmetal :: mlx::core::metal::Device::end_encoding(int)
  pmetal :: mlx::core::gpu::eval(mlx::core::array&)
  pmetal :: mlx::core::eval_impl(std::__1::vector<mlx::core::array, ...>, bool)
  pmetal :: mlx::core::eval(std::__1::vector<mlx::core::array, ...>)
  pmetal :: mlx_eval
  pmetal :: pmetal_mlx_rs::transforms::eval
  pmetal :: pmetal_trainer::training_loop::TrainingLoop::accumulate_gradients
  pmetal :: pmetal_trainer::training_loop::run_metal_fused::...::run_metal_fused
  pmetal :: pmetal_trainer::orchestrator::run_lora_path

Root cause analysis

pmetal-mlx-sys 0.2.4 bundles MLX v0.30.6 via FetchContent in mlx-c/CMakeLists.txt:

GIT_TAG v0.30.6

MLX 0.31.0 includes two relevant fixes:

• Fix residency set with user provided buffer ml-explore/mlx#3108 — "Fix residency set with user provided buffer" — fixes Metal buffer tracking that can cause null encoder in end_encoding()
• Fix precision in Metal fused attention ml-explore/mlx#3119 — "Fix precision in Metal fused attention"

The M5 Pro (Apple10 Pro gen 17) is new hardware that likely exposes these bugs more readily than older chips.

Suggested fix

Bump the MLX dependency in pmetal-mlx-sys from v0.30.6 to ≥v0.31.0. The workspace currently pins pmetal-mlx-sys = "=0.2.4" in the root Cargo.toml.

Note: the upstream mlx-rs repo has already upgraded to MLX-C 0.4.0 (commit 1deb45a).

Additional notes

• The same training config completes successfully with mlx-lm (which uses MLX 0.31.0 via Python) in 545s on the same machine
• Sequence packing mode (--no-sequence-packing omitted) also showed very low throughput (125 tok/s) before we disabled it, though that may be a separate issue with batch sizing

[barrettj]

Update: MLX bump to 0.31.1 does NOT fix the crash

We patched pmetal-mlx-sys locally to use MLX-C 0.6.0 (which bundles MLX v0.31.1) and rebuilt pMetal from source. The two required code changes were adding global_scale (nullable) params to mlx_quantize and mlx_dequantize in pmetal-mlx-rs/src/ops/quantization.rs.

Build succeeded, but the SIGSEGV still occurs at exactly step 100 (the eval + checkpoint point). Same null deref at address 0x0, same stack shape. The crash is not in upstream MLX — it's in pMetal's own Metal fused optimizer or checkpoint path.

Additional findings

1. Throughput degradation (unpacked mode)

Even without packing, throughput degrades ~35% over 100 steps on M5 Pro:

Step	tokens/s
10	1,081
30	946
50	835
70	811
90	796
100	874 (then crash)

This degradation does NOT happen on M4 Max (per testing on another machine). Suggests a Metal memory or buffer management issue specific to Apple10 Pro (M5) GPU architecture.

2. Sequence packing is broken on M5 Pro

With default packing enabled (--no-sequence-packing NOT set):

• 4,611 sequences pack into 48 batches (~96 seqs/batch) at 8192 max_seq_len
• Throughput: ~125 tokens/s (vs ~1,000+ unpacked, vs ~2,000+ on M4 Max with packing)
• Steps take so long that training is 10x slower than without packing

Root cause appears to be that max_seq_len auto-detects to 8192 from the model config, but our sequences are only 50-400 tokens. The packer creates 8192-token batches with O(n²) attention cost. A --max-seq-len 512 flag might help, but the 125 tok/s still seems too low even for that batch size.

This may warrant a separate issue since it affects packing generally on M5, not just the crash.

3. Summary of M5 Pro issues

Three distinct problems, all M5-specific:

1. SIGSEGV at eval/checkpoint (step 100) — occurs with both MLX 0.30.6 and 0.31.1
2. Throughput degradation — 1,081 → 796 tok/s over 100 steps (unpacked)
3. Sequence packing 10x slower — 125 tok/s packed vs 1,000+ unpacked

All three work correctly on M4 Max. The M5 Pro (Apple10 Pro gen 17, 20 GPU cores) appears to expose Metal buffer lifecycle or fused kernel issues not present on M4.

[nicholasjpaterno]

@barrettj thanks for opening this issue, we're gearing up to release 0.5.0 which should address this issue. I'll see if i can ship a 0.4.1 patch in the meantime. The 0.5.0 has a total revamp for Metal 4 (M5 support).

[nicholasjpaterno]

@barrettj Are you able to test by building the latest main? I don't have an M5 device to test with but I've overhauled the Metal 4 (M5 chips) based on the available documentation. If you have the bandwidth to give it a run, please let me know if you bump into any issues!

[barrettj]

I should have time to check later tonight - is there anything specific you’d like me to run/test or just trying using it in my workflow like I did before?Sent from my iPhoneOn Apr 9, 2026, at 5:48 AM, Nick Paterno ***@***.***> wrote:nicholasjpaterno left a comment (Epistates/pmetal#14) @barrettj Are you able to test by building the latest main? I don't have an M5 device to test with but I've overhauled the Metal 4 (M5 chips) based on the available documentation. If you have the bandwidth to give it a run, please let me know if you bump into any issues! —Reply to this email directly, view it on GitHub, or unsubscribe.You are receiving this because you were mentioned.Message ID: ***@***.***>

[nicholasjpaterno]

@barrettj no rush and I appreciate any feedback would be great! It was a giant overhaul for metal 4 and I dont have an M5 chip to test with just yet! Just a simple training run would be great - the same run that caused the failure #14 (comment) but also inference as well.