apple metal

In GPU programming, tension between fine-grained per-thread control and coarse-grained collective operations (e.g., Tensor Core instructions) undermines modularity and safety: collective primitives require coordinated execution across thread groups, yet encapsulated functions are invoked by individual threads. Method: We introduce Prism, a new language featuring *typed views*—a novel type-system mechanism that explicitly classifies thread behavior by control granularity (per-thread, group-wide, or global), enabling safe, modular abstraction of collective operations. Built upon the Bundl core calculus, Prism’s type-safe compiler supports precise, hardware-aware abstractions for accelerators like Tensor Cores. Contribution/Results: Evaluation shows Prism delivers strong type safety with zero runtime overhead, significantly improving GPU kernel correctness, maintainability, and developer productivity—without sacrificing performance.

Heterogeneous GPUs exhibit severe binary incompatibility due to divergent instruction sets, execution models, and driver stacks. To address this, we present the first cross-vendor binary-compatible system supporting NVIDIA, AMD, Intel, and Tenstorrent GPUs. Our approach comprises three key components: (1) an architecture-agnostic intermediate representation (IR) that unifies SIMT and MIMD execution semantics; (2) a dynamic binary translation runtime coupled with state-aware serialization for real-time migration; and (3) a unified abstraction layer ensuring consistent cross-platform semantics for threads, memory, and synchronization. Crucially, our system enables direct binary migration across all four GPU architectures without source or binary modification. Evaluation shows migration overhead under 8% and average performance degradation below 12%, effectively breaking down long-standing hardware-software co-design barriers across GPU vendors.

This work presents the first empirical investigation into the execution mechanism of fp8 matmul2d operations in Apple Metal 4.1 on the M4 Max GPU, revealing that they are implemented via software emulation on shader cores rather than dedicated hardware acceleration, utilizing an undocumented 8×8 tensor tiling layout. Through a meticulously designed microbenchmarking framework incorporating checksum-based validation, provenance tracking, throughput ceiling analysis, comparison against simdgroup_matrix primitives, and power attribution, we demonstrate that fp8 achieves only 0.94× the performance of fp16, confirming its primary benefit lies in memory footprint reduction. Leveraging these insights, we develop a hand-optimized fused kernel that attains performance gains of 6.5–12.9% in cache-resident scenarios.

Assessing performance portability across GPU programming models on heterogeneous NVIDIA and AMD hardware remains challenging due to fragmented benchmarks and irreproducible evaluation methodologies. Method: We conduct a systematic, cross-platform evaluation of seven programming models—CUDA, HIP, Kokkos, RAJA, OpenMP, OpenACC, and SYCL—using five interdisciplinary proxy applications. Leveraging a Spack-based automation framework, we ensure fully reproducible build, deployment, and benchmarking across real multi-vendor GPU systems. Contribution/Results: Our empirical study quantifies both performance consistency and migration overhead for each model. HIP and SYCL achieve superior performance on AMD GPUs; CUDA remains dominant on NVIDIA hardware; Kokkos and RAJA deliver balanced portability with moderate performance; OpenMP and OpenACC exhibit significant cross-platform performance degradation. This work provides the first unified, vendor-agnostic assessment of GPU programming models’ performance portability and establishes a rigorous, reproducible methodology to guide architecture selection for high-performance scientific software.

To address low GPU utilization and poor adaptability to diverse AI workloads caused by single-task execution in the large-model era, this paper proposes, for the first time, a systematic GPU multitasking paradigm. Inspired by CPU operating system resource management, we design an OS-like GPU resource management layer that enables dynamic resource partitioning, strong task isolation, and priority-aware scheduling. We formally define key requirements and core challenges—including fine-grained resource sharing, low-overhead context switching, and cross-task QoS guarantees—and outline concrete technical pathways to address them. This work establishes a theoretical foundation and a holistic architectural framework for evolving GPUs from single-task devices toward efficient, secure, and schedulable multitasking compute platforms. It advances the development of high-utilization, high-performance AI computing systems.

This work addresses the challenge of automatically generating high-performance GPU tiled kernels from high-level tensor algebra expressions, thereby alleviating the burden of manual optimization. The authors propose an end-to-end compilation framework that integrates layer-wise lowering, expression rewriting, automated schedule search, reduction fusion, and tiling optimizations. For the first time, this framework automatically discovers high-efficiency kernels—comparable to FlashAttention-3—from the mathematical specification of attention operators, while introducing a novel scheduler that preserves program structural regularity. Evaluated on GH200 and RTX 5090 GPUs, the generated kernels achieve up to 23% and 42% higher throughput, respectively, and match or surpass hand-optimized cuDNN kernels across multiple long-sequence configurations.

Developing GPU kernels in Rust presents a significant challenge in simultaneously achieving high performance and memory safety. This work proposes cuTile Rust, the first system to enforce Rust’s ownership semantics within GPU kernels. By introducing a tiled kernel design coupled with a novel ownership separation mechanism, cuTile Rust enables safe, idiomatic low-level control while unifying diverse execution models—including synchronous launches, asynchronous pipelining, and CUDA Graph replay. Evaluated on NVIDIA B200 hardware, the system attains 7 TB/s element-wise memory bandwidth and 2 PFlop/s GEMM performance, reaching 96% of cuBLAS throughput. Furthermore, the Grout inference engine built atop cuTile Rust achieves 171 and 82 tokens/s on Qwen3-4B and Qwen3-32B models, respectively, matching the performance of state-of-the-art frameworks such as vLLM and SGLang.

This study addresses the low utilization of modern GPU computing resources by systematically evaluating the performance, energy efficiency, and resource isolation characteristics of NVIDIA’s Multi-Process Service (MPS) and Multi-Instance GPU (MIG) technologies under concurrent application workloads. The experiments reveal a critical trade-off between MPS’s scheduling flexibility and MIG’s hardware-level isolation: MPS can improve performance by up to 30% and reduce energy consumption by approximately 20% in the absence of memory contention, yet suffers a 30% performance degradation under contention; MIG effectively mitigates resource contention but is constrained by its rigid configuration options and higher overhead. These findings provide empirical foundations for optimizing GPU co-execution strategies driven by application-specific workload characteristics.

This work addresses the challenge of type erasure in GPU binaries, where GPU compilers uniformly store all data types in the register file, thereby discarding type information and impeding disassembly and semantic recovery. To systematically resolve this issue, the paper introduces a constraint-propagation-based type inference method that integrates conflict detection, control-flow reconstruction, and multi-instruction pattern recognition to lift SASS binaries into typed LLVM IR. Evaluated on 24,437 GPU functions across eight benchmark suites, the approach achieves a 99.98% lifting success rate. Ablation studies further demonstrate that disabling type recovery reduces the x86 compatibility rate from 73.8% to 0%, underscoring the critical role of accurate type reconstruction in generating semantically correct intermediate representations.

This work addresses the inefficiencies of conventional GPU linear algebra libraries, which often fail to fully exploit hardware capabilities due to redundant memory accesses and runtime overhead. To overcome these limitations, the authors propose Bandicoot, a toolkit that leverages C++ template metaprogramming to perform expression fusion at compile time, thereby automatically generating highly optimized GPU kernels that saturate memory bandwidth—without relying on just-in-time compilation or runtime scheduling. Bandicoot provides an API compatible with Armadillo, facilitating straightforward migration of existing CPU codebases. Experimental results demonstrate that Bandicoot significantly outperforms PyTorch, TensorFlow, and JAX across multiple benchmarks, achieving substantial speedups in several scenarios.