M1 Ultra profile + two negative results (TensorOps unavailable, M5 gates are no-ops)

[intelac]

First, thanks for open-sourcing this — it builds and runs cleanly on M1 Ultra with zero warnings, and it is roughly 2.15× faster than the Swift/MLX path on the same machine while using 44% less memory, at full BF16 rather than 4-bit. That is a big deal for people on older Apple Silicon.

I did not find M1/M2-series numbers anywhere in the issues (#5 and #8 cover M4 Max, #15 mentions M3 Ultra), so here is a complete profile from an M1 Ultra, plus two negative results that may save others some time.

Environment

• Apple M1 Ultra, 128 GB, macOS 26.2
• h3.c at 8974cc0
• Official MiniMaxAI/MiniMax-H3 snapshot, FL2VA/ only

Device: Apple M1 Ultra (applegpu_g13d)
  physical memory       128.0 GiB
  recommended GPU set   107.5 GiB
  max Metal buffer      80.6 GiB
  Apple GPU family      7
  Metal 4               yes
  unified memory        yes

Profile — 512×512, 22 frames, 20 steps, --layers 45 --reuse 2

Same preset as the README's balanced example (~16.7 s denoise on M5 Max):

Stage	wall	encode	wait	peak
Qwen text encoder	10.5 s	0.003 s	5.4 s	2.7 GiB
H3 DiT load	18.3 s	1.6 s	1.7 s	32.8 GiB
H3 DiT Euler denoise	86.6 s	56.7 s	29.8 s	32.8 GiB
audio VAE decode	2.6 s	2.3 s	0.2 s	0.3 GiB
video VAE decode	7.3 s	0.9 s	4.3 s	9.5 GiB
total	~125 s			32.8 GiB

So denoise is ≈5.2× slower than M5 Max on the identical preset. encode dominates over wait, i.e. compute-bound rather than bandwidth-bound — which matches expectations, since M1 Ultra actually has higher memory bandwidth (800 GB/s) than M5 Max but no per-core matrix engine.

Profile — 768×416, --steps 15 --layers 45 --reuse 2

Two frame counts, showing denoise scales close to linearly:

frames	duration	denoise	total wall	peak
141	5.875 s	822.9 s	15 min 14 s	34.1 GiB
243	10.125 s	1930.1 s	34 min 28 s	35.2 GiB

One observation: at 243 frames the video VAE decode was 105.3 s wall with only 0.9 s encode — 101.7 s of wait. Decode appears almost entirely wait-bound at higher frame counts on this hardware.

Negative result 1 — TensorOps is a hard hardware limit, not a gate

I wanted to check whether the is_m5() name match in h3_gpu.m was leaving performance on the table, so I compiled a minimal probe directly:

NSString *src = @"#include <metal_stdlib>\n#include <metal_tensor>\n"
                 "using namespace metal;\nusing namespace mpp::tensor_ops;\n"
                 "kernel void probe(device float* o [[buffer(0)]]) { o[0]=1.0f; }\n";
id<MTLLibrary> lib = [dev newLibraryWithSource:src options:o error:&err];

Result on M1 Ultra:

Device: Apple M1 Ultra
Metal4 family: YES
TensorOps compile: FAILED
Reason: program_source:4:17: error: use of undeclared identifier 'mpp'

So supportsFamily:MTLGPUFamilyMetal4 returns YES while mpp::tensor_ops does not exist at all — the framework is available but the namespace is gated on the GPU neural accelerators introduced in M5/A19. Forcing is_m5() true would just hit the existing fallback at h3_gpu.m:381-393. Your name-based check is semantically correct; there is no free performance hiding behind it.

Negative result 2 — the three M5 software gates are no-ops here

Three gates select M5 behaviour by device name and are overridable by env var:

• H3_GPU_SAMPLER (h3_dit.c:2584) — GPU sampler, M5-only by default
• H3_QWEN_PREFETCH_DEPTH (h3_text_encoder.c:137) — 3 on M5, 2 otherwise
• H3_DIT_COMMAND_BLOCKS (h3_dit.c:216) — blocks*3/5 on M5, 30 otherwise

A/B on 512×512 / 22 frames / 20 steps, DiT total wall:

run	config	DiT total
A	defaults (cold)	106.88 s
B	H3_GPU_SAMPLER=1	103.49 s
C	H3_QWEN_PREFETCH_DEPTH=3	103.46 s
D	B + C	103.55 s
E	D + H3_DIT_COMMAND_BLOCKS=27	103.39 s

B–E span 0.16 s — noise. The 3.4 s gap to A is cold-start (first-run shader compile), not the flags: run C did not enable the GPU sampler yet matched B exactly. H3_DIT_COMMAND_BLOCKS=27 did change submissions (66 → 77) with no timing effect. Prefetch depth 3 only raised text-encoder peak from 2.7 to 3.6 GiB with no speedup.

The actual question

Given that the compute path is closed on pre-M5 hardware, is there anything on the roadmap targeting the parts that aren't matrix-engine dependent? Specifically:

1. The wait fraction. Denoise showed ~34% wait (29.8 s of 86.6 s), and 243-frame video decode was ~97% wait. If any of that is schedulable rather than inherent, it would be the largest available win on this class of machine.
2. Whether --layers/--reuse defaults should differ pre-M5. The fast preset is clearly the practical operating point here; the exact preset (--layers 50 --reuse 1 --steps 20) came out at ~4.5 min/step at 243 frames, which is not usable.

Entirely understood if pre-M5 is out of scope — the README is explicit that the current work is M3 Max and M5 Max. In that case please treat this as a data point for anyone else landing here with an Ultra, and feel free to close.

Happy to run further profiles or test patches on this hardware if useful.