Rigel: Reverse-Engineering the Metal 4.1 Tensor Compute Path on the Apple M4 Max GPU

• URL: http://arxiv.org/abs/2606.12765v1
• Date: Thu, 11 Jun 2026 00:10:23 GMT
• Title: Rigel: Reverse-Engineering the Metal 4.1 Tensor Compute Path on the Apple M4 Max GPU
• Abstract summary: We present Rigel, an empirical characterization of this path on a single Apple M4 Max.<n>We reconstruct the opaque 8x8 cooperative_tensor fragment layout Apple documents nowhere.<n>Acting on the characterization, a hand-fused GEMM + bias + GELU kernel beats the decomposed path by +6.5-12.9% in the cache-resident regime.
• Score: 3.151184728006369
• License: http://creativecommons.org/licenses/by/4.0/
• Abstract: Apple's Metal 4.1 exposes a tensor compute path: the Metal Performance Primitives (MPP) matmul2d operation over cooperative_tensor fragments, whose interface is documented but whose hardware behavior is deliberately hidden. The specification states which data-type rows are supported, never whether they are hardware-accelerated, where the operation physically executes, what its accumulator width is, or how it partitions matrix fragments across threads. We present Rigel, an empirical characterization of this path on a single Apple M4 Max (a pre-neural-accelerator generation). Using a checksum-gated, provenance-tracked microbenchmark harness, Rigel recovers eleven facts the v4.1 specification hides or contradicts. The headline finding: the Metal 4.1 fp8 (E4M3) matmul2d is emulated, not accelerated: it sustains 0.94x the throughput of fp16 despite reading half the operand bytes, so on M4 it is a memory-footprint feature, not a performance feature. We further show, via a three-signal triangulation (throughput ceiling, comparison against simdgroup_matrix, and per-rail power attribution), that matmul2d executes entirely on the GPU shader cores with no dedicated matrix datapath and no evidence of Apple Neural Engine routing; that it accumulates in >=fp32; and we reconstruct the opaque 8x8 cooperative_tensor fragment layout Apple documents nowhere. Acting on the characterization, a hand-fused GEMM + bias + GELU kernel beats the decomposed path by +6.5-12.9% in the cache-resident regime. All findings are reproducible from committed MIT-licensed code and per-cell CSVs.

Related papers

• FusionML: Prefill, Not Decode - Mechanism and Boundaries of CPU+GPU Co-Execution on Unified-Memory Apple Silicon [0.0]
  Previous attempts, including our own, failed or produced precision-confounded wins.<n>We identify the cause: MLX's lazy-graph scheduler emphserializes cross-stream workWhenever a CPU-stream operation consumes an unmaterialized GPU result inside one evaluation graph, a row-split matmul runs x1.38 faster with materialized inputs.<n>We characterize the boundaries equally carefully: decode cannot benefit, bound by shared bandwidth co-execution does not add.
  arXiv  Detail & Related papers  (2026-07-24T11:42:17Z)
• BaseRT: Advancing Best-in-Class LLM Inference with Apple M5 Neural Accelerators [0.0]
  We show that BaseRT exploits the Metal inference for large language models on Apple Silicon.<n>On an Apple M5 Pro, BaseRT delivers up to $6.4times$ higher prompt-processing throughput than llama.<n>We show that the M5's tensor cores are the decisive lever for prompt processing on Apple Silicon.
  arXiv  Detail & Related papers  (2026-07-21T06:42:18Z)
• A Modern Multimodal Assistant on a 6 GB 2011 GPU: Stage-Validated, All-GPU CUDA Inference for Fermi [0.0]
  This report keeps the hardware and asks what a model that fits can do.<n>We deploy Mini-V-4.6, a modern multimodal assistant pairing a SigLIP2 encoder and window-attention merger.<n>The system answers an image question end-to-end in 1.7s.
  arXiv  Detail & Related papers  (2026-07-16T04:48:44Z)
• MonaVec: A Training-Free Embedded Vector Search Kernel for Edge and Offline AI Systems [0.0]
  MonaVec is a deterministic, embedded vector-search kernel for edge and offline AI.<n>It targets on-device RAG, offline embedded retrieval -- the niche occupies for relational data: one file, one call, runs anywhere.
  arXiv  Detail & Related papers  (2026-06-17T18:00:54Z)
• FLASH-MAXSIM: IO-Aware Fused Kernels for Late-Interaction Scoring [2.159285655678094]
  Flash-MaxSim is an IO-aware kernel that computes exactly the same scores without ever materializing the tensor.<n>It is up to 3.9x faster on an A100 than naive PyTorch at matched precision, uses up to 16x less inference memory and 28x less training memory.
  arXiv  Detail & Related papers  (2026-05-28T07:38:27Z)
• PySIFT: GPU-Resident Deterministic SIFT for Deep Learning Vision Pipelines [0.22940141855172033]
  We show that classical SIFT with DSP multi-scale pooling outperforms neural descriptor and orientation replacements.<n>We present PySIFT, the first fully GPU-resident SIFT, implemented in CuPy/Numba kernels with DLPack zero-copy handoff to downstream DL frameworks.
  arXiv  Detail & Related papers  (2026-05-18T05:24:55Z)
• A Controlled Study of Memory Hierarchy Transitions in Quantum Circuit Simulation on Apple M4 Pro Unified Memory Architecture [0.0]
  State-vector quantum circuit simulation is memory-bandwidth bound.<n>We address this using the Apple M4 Pro Unified Memory Architecture.<n>We show that peak streaming bandwidth does not predict simulation speedup for non-contiguous memory access patterns.
  arXiv  Detail & Related papers  (2026-05-09T08:22:55Z)
• ZenBrain: A Neuroscience-Inspired 7-Layer Memory Architecture for Autonomous AI Systems [51.56484100374058]
  On LongMemEval-500, ZenBrain matches a long-context oracle's binary-judge accuracy to within 4.5 pp.<n> ZenBrain is a 7-layer neuroscience-inspired memory architecture.
  arXiv  Detail & Related papers  (2026-04-26T20:39:19Z)
• FlashSampling: Fast and Memory-Efficient Exact Sampling [62.5203057469482]
  FlashSampling is an exact sampling primitive that fuses sampling into the LM-head matmul and never materializes the logits tensor in.<n>Across H100, H200, B200, and B300 GPUs, FlashSampling speeds up kernel-level decode workloads.<n>In end-to-end vLLM experiments, it reduces time per output token by up to $19%$ on the models we test.
  arXiv  Detail & Related papers  (2026-03-16T19:37:08Z)
• RFX: High-Performance Random Forests with GPU Acceleration and QLORA Compression [0.0]
  RFX (Random Forests X), where X stands for compression or quantization, presents a production-ready implementation of Breiman and Cutler's Random Forest classification methodology in Python.<n>RFX v1.0 provides complete classification: out-of-bag error estimation, overall and local importance measures, proximity matrices with QLORA compression, case-wise analysis, and interactive visualization (rfviz)<n> Regression, unsupervised learning, CLIQUE importance, and RF-GAP proximity are planned for v2.0.
  arXiv  Detail & Related papers  (2025-11-23T12:00:33Z)
• FuseSampleAgg: Fused Neighbor Sampling and Aggregation for Mini-batch GNNs [51.56484100374058]
  FuseSampleAgg fuses neighbor and sampling mean aggregation into a single pass for GraphSAGE.<n>Operator is deterministic, integrates with standard PyTorchs, and ships with scripts that reproduce all tables and figures from CSV logs.
  arXiv  Detail & Related papers  (2025-11-17T17:57:18Z)
• A Simple Linear Patch Revives Layer-Pruned Large Language Models [58.056251480151104]
  Layer pruning has emerged as a widely used technique for compressing large language models (LLMs)<n>textscLinearPatch fuses two operations into one matrix multiply at the pruning interface.<n>The patch can be further refined with 5K unlabeled samples via memory-efficient offline distillation, pushing the retention to 95.16% within only 30 minutes on a single GPU.
  arXiv  Detail & Related papers  (2025-05-30T15:06:08Z)
• MIB: A Mechanistic Interpretability Benchmark [77.35046700898326]
  We propose MIB, a Mechanistic Interpretability Benchmark, with two tracks spanning four tasks and five models.<n>Using MIB, we find that attribution and mask optimization methods perform best on circuit localization.<n>For causal variable localization, we find that the supervised DAS method performs best, while SAE features are not better than neurons.
  arXiv  Detail & Related papers  (2025-04-17T17:55:45Z)
• INR-Arch: A Dataflow Architecture and Compiler for Arbitrary-Order Gradient Computations in Implicit Neural Representation Processing [66.00729477511219]
  Given a function represented as a computation graph, traditional architectures face challenges in efficiently computing its nth-order gradient. We introduce INR-Arch, a framework that transforms the computation graph of an nth-order gradient into a hardware-optimized dataflow architecture. We present results that demonstrate 1.8-4.8x and 1.5-3.6x speedup compared to CPU and GPU baselines respectively.
  arXiv  Detail & Related papers  (2023-08-11T04:24:39Z)
• HDCC: A Hyperdimensional Computing compiler for classification on embedded systems and high-performance computing [58.720142291102135]
  This work introduces the name compiler, the first open-source compiler that translates high-level descriptions of HDC classification methods into optimized C code. name is designed like a modern compiler, featuring an intuitive and descriptive input language, an intermediate representation (IR), and a retargetable backend. To substantiate these claims, we conducted experiments with HDCC on several of the most popular datasets in the HDC literature.
  arXiv  Detail & Related papers  (2023-04-24T19:16:03Z)

This list is automatically generated from the titles and abstracts of the papers in this site.

This site does not guarantee the quality of this site (including all information) and is not responsible for any consequences.