LFM2.5-24B-A2B on M1: cold prompts run at about 5 tokens a second, and a long one ignores SIGTERM

@Arminova

Reported by @Arminova on #1394 (receipt in this comment), on an M1 Pro (32 GB) with LiquidAI/LFM2.5-24B-A2B-MLX-4bit at 667408d4, no drafter. The same numbers show on main ee896951d, so the drafter stack in #1390 to #1394 didn't cause them.

Cold prompts are slow. Prompts from 16 to 930 tokens fit about 15 s before the first prompt token, then about 0.2 s a token (about 5 prompt tokens a second). Decode is normal (44 to 55 tok/s). The server's own startup line times LFM2's windows at about 110 rows a second at width 16, some 20 times what these prompts get, so the prompt path on M1 isn't using the batched windows it measures.

A long prompt can't be stopped. At --context 16384, a 1,864-token chat prompt took 404 s, and a following ~8k-token prompt didn't finish. The server ignored SIGINT and SIGTERM during it and needed SIGKILL (seen once). The whole-prompt pass has no cancel point.

What we know so far: on an M5 Ultra our 1.1.8 release cell runs this model's cold prompts at about 16,300 tokens a second at 2k and 14,500 at 16k, so this looks specific to M1 (and perhaps M2), where the window kernels take another path. We have no M1 to reproduce on, so a run with TF_* logging from an M1 owner will help once we have a candidate. Next steps on our side: find which prompt path M1 takes and why it falls back to rows one at a time, and give the whole-prompt pass a cancel point so a shutdown or a cancelled request ends at the next window.

Thanks @Arminova for the careful receipt.

/atc claim

/atc release

Found the cause on my M1 Pro (32 GB, macOS 27.0.1, Xcode 27.1, Zig 0.17.0, main ee896951d), with a one-function change that passes tf-lfm2-exact.

Where the time goes. tf-lfm2-bench on main: every core/tile_qmm call runs at 0.04 TFLOPS (dense and grouped, every shape). A 2,048-row chunk budgets 243.3 s, of which the tiles take all but 0.23 s; attention's own MMAs run at 1.98 TFLOPS. A 512-token prompt in 128-row chunks: 179.63 s wall, 179.60 s GPU.

What in the tile. Compiling the shipped tile_qmm.metal on its own and timing tf_tile_qmm_q4 at 2048x2048, 512 rows:

variant	ms	TFLOPS	output vs shipped
shipped	122.80	0.035	-
4-bit codes computed, no `tf_tq_pairs` table	122.30	0.035	bit-identical
MMAs on half	2.09	2.051	differs
A fragment as one `vec<bfloat, 2>` store	48.31	0.089	bit-identical
A and B fragments as `vec<bfloat, 2>` stores	2.13	2.016	bit-identical
A, B and the accumulator as vector stores (as `lfm2_mma`)	1.60	2.682	bit-identical

tf_tq_mmas writes each simdgroup_matrix<bfloat, 8, 8> fragment element by element (a[i].thread_elements()[0] = ap[0]); on this M1 that costs about 75x. lfm2_mma in lfm2.metal stores whole fragments through reinterpret_cast and is fast here.

Latest main 7c0688951 still has it: tile_qmm.metal is unchanged since ee896951d, and its source alone times the same on this M1 (120.52 ms shipped, 1.60 ms with the change below, bit-identical).

The change, in tf_tq_mmas only:

    for (int i = 0; i < TI; ++i) {
      const threadgroup bfloat* ap = xb + (sr + 8 * i + fm) * TP + kk + fn;
      reinterpret_cast<thread vec<bfloat, 2>&>(a[i].thread_elements()) = *(const threadgroup vec<bfloat, 2>*)ap;
    }
    for (int j = 0; j < TJ; ++j) {
      const threadgroup bfloat* bp = wb + (tf_tq_col<GLU>(sc, j) + fn) * TP + kk + fm;
      reinterpret_cast<thread vec<bfloat, 2>&>(b[j].thread_elements()) = vec<bfloat, 2>(bp[0], bp[TP]);
    }
    for (int i = 0; i < TI; ++i)
      for (int j = 0; j < TJ; ++j) {
        simdgroup_matrix<float, 8, 8> C;
        reinterpret_cast<thread float2&>(C.thread_elements()) = reinterpret_cast<thread float2&>(P[i][j].thread_elements());
        simdgroup_multiply_accumulate(C, a[i], b[j], C);
        reinterpret_cast<thread float2&>(P[i][j].thread_elements()) = reinterpret_cast<thread float2&>(C.thread_elements());
      }

With it, on the real weights (LiquidAI/LFM2.5-24B-A2B-MLX-4bit @ 667408d4):

tf-lfm2-exact: 72 forward/cache checks passed, 0 failed, 0 skipped, including "prompt 2047 in 2048-row chunks equals one row at a time".
tf-kernel-contract --only tile_qmm: 6 kernels pass gates 1-4, 0 fail, 0 limit problems.
zig build test -Dcpu=apple_m1: 7,547 of 7,593 passed, 45 skipped, 1 failed, the same counts as main without the change. The one failure is families.glm.eh_check "eh_proj on the tensor units: fp64 error no worse than the row kernel", on main too.
Cold prompts, --parallel 1 --context 4096, server ttft, the same probe as my receipt on lfm2: a cheaper DSpark hold (one concurrent encoder) #1394:

prompt tokens	main `ee896951d`	main + change
16	15.38 s	0.222 s
114	40.62 s	0.439 s
462	109.23 s	1.081 s
930	212.11 s	1.975 s
1,856	-	3.706 s
fit	15.0 s + 0.204 s a token	0.20 s + 0.0019 s a token (about 530 tok/s)

Also seen, not looked into: the server loaded in 11.9 s against 131 to 133 s, and one 128-token reply decoded at 77.6 tok/s against 44.1 on main.

The stop signals weren't ignored. The log shows shutting down (SIGINT): draining ... and then a second stop signal: 1 request(s) still in progress end "server restarting: retry shortly" at the next round boundary. prefillFn runs the whole prompt as one round with no stop check, so the next boundary was the end of the prompt, about 20 minutes away at the old speed. The change makes that seconds, but the cancel point you described is still missing.

Not tested: M2 or M3; other kernels that write bfloat fragments element by element (core/tile_qmv_rows.metal, qwen27/attention_prompt_sg.metal); the PR stack's test-lfm2 and --only lfm2 (main ee896951d has neither). Happy to run a branch with TF_* logging, or any other check, on this M1.

The change above is ready on a branch. I'm running the full checks once more to make sure it works, and I'll raise the PR tomorrow.

tile_qmm: store the MMA fragments whole, prompt tiles 62-78x faster o…

Navigation Menu

Repository navigation

LFM2.5-24B-A2B on M1: cold prompts run at about 5 tokens a second, and a long one ignores SIGTERM #1523

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

LFM2.5-24B-A2B on M1: cold prompts run at about 5 tokens a second, and a long one ignores SIGTERM #1523

Description

Activity

ashhart commented on Oct 28, 2026

ashhart commented on Oct 28, 2026

Arminova commented on Oct 29, 2026

Arminova commented on Oct 29, 2026

Metadata

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

Issue actions