Porting GPU-Accelerated Graph Algorithms from CUDA to Apple Metal: 5 Critical Findings for Apple Silicon GPGPU (M1/M2/M3/M4)

porting_cuda_graph_algorithms_to_apple_metal.md

Porting GPU-Accelerated Graph Algorithms from CUDA to Apple Metal: Technical Findings

Abstract

This document presents technical findings from a ground-up port of GPU-accelerated shortest-path graph algorithms (Bellman-Ford, Delta-Stepping, SPFA) from NVIDIA CUDA C++ to Apple Metal Shading Language (MSL) on the M4 architecture. The target application is PCB autorouting (OrthoRoute), where graphs contain 2,000 to 401,800 nodes and up to 2 million edges. The final Metal implementation achieves 3.7x higher throughput than an RTX 3060 on the largest tested graph, with 111.4 billion edges per second and 831.5 GB/s effective memory bandwidth.

Five critical discoveries emerged during the port. Each required solving a problem that has no direct CUDA equivalent and limited public documentation.

---

1. Persistent Thread Structural Deadlock on Apple Silicon

The Problem

CUDA persistent thread queues are a common pattern: launch a single large grid, have threads spin-wait on a global work queue index, and process items as they arrive. This pattern avoids repeated kernel launch overhead.

On the Apple M4, launching a persistent grid that exceeds the physical GPU core concurrency causes an irrecoverable OS watchdog timeout (IOAF code 5 / gputimeout).

Root Cause

Apple Silicon GPU schedulers do not preempt compute threadgroups that are trapped in a spin-loop. The M4 has a 10-core GPU. If you launch a grid of 65,536 threads (128 threadgroups), only approximately the first 16 threadgroups physically execute. The remaining threadgroups are queued by the hardware scheduler, waiting for a core to become available. If the executing threadgroups spin-wait for data that must be produced by a queued threadgroup, the hardware deadlocks permanently.

Resolution

The persistent grid must be strictly bounded to the maximum concurrent hardware occupancy. On the M4, the safe limit is 8,192 threads (16 threadgroups of 512 threads each). This fully saturates all 10 GPU cores while leaving no threadgroups waiting in the hardware queue.

This constraint does not exist in CUDA because NVIDIA GPUs support preemptive scheduling of thread blocks.

---

2. Zero-Dispatch Software Grid Barrier

The Problem

Apple MSL provides threadgroup_barrier(mem_flags::mem_threadgroup) for synchronization within a single threadgroup. However, unlike CUDA Cooperative Groups (cg::this_grid().sync()), MSL has no built-in global grid barrier for compute kernels.

The naive workaround is to split each algorithm phase into separate kernel dispatches and synchronize via commandBuffer.commit() / waitUntilCompleted(). This incurs a minimum latency of approximately 38.3 microseconds per dispatch on the M4, which dominates total execution time when the algorithm requires hundreds of iterations.

Resolution

A software grid barrier was constructed using a device atomic_uint array. Because Apple Silicon uses a weakly-ordered memory model, simple atomic increments are insufficient. L1 caches across GPU cores can fall out of coherency, causing threads to read stale barrier state.

The solution uses threadgroup_barrier(mem_flags::mem_device) before reading global state to force L1 cache invalidation:

// 1. Flush local device writes to L2
threadgroup_barrier(mem_flags::mem_device);

// 2. One thread per threadgroup increments the global counter
if (tid == 0) {
    uint old = atomic_fetch_add_explicit(&barrier[0], 1, memory_order_relaxed);
    if (old == num_threadgroups - 1) {
        // Last threadgroup: reset counter, advance generation
        atomic_store_explicit(&barrier[0], 0, memory_order_relaxed);
        atomic_fetch_add_explicit(&barrier[1], 1, memory_order_relaxed);
    }
}

// 3. All threads spin on the generation counter
threadgroup_barrier(mem_flags::mem_device);
while (atomic_load_explicit(&barrier[1], memory_order_relaxed) < target_gen) { }
threadgroup_barrier(mem_flags::mem_device);

This eliminates all CPU-GPU round-trips during the iterative algorithm. The entire SSSP computation runs inside a single commandBuffer.commit() call.

---

3. SIMD-Group Block Stealing for Atomic Congestion Avoidance

The Problem

In a persistent work-stealing queue, 8,192 threads simultaneously execute atomic_fetch_add_explicit on a single device atomic_uint queue index. Even with L2 cache residency, this creates severe atomic serialization latency on the M4's memory controller.

Resolution

The M4 executes in SIMD-groups of 32 threads (analogous to CUDA warps). Instead of having every thread steal 1 work item, the first thread in each SIMD-group steals 32 items at once using simd_is_first() and simd_broadcast_first():

uint base_idx = 0;
if (simd_is_first()) {
    base_idx = atomic_fetch_add_explicit(steal_index, 32, memory_order_relaxed);
}
base_idx = simd_broadcast_first(base_idx);
uint my_task = base_idx + lane_id;

This reduces atomic contention on the queue index by exactly 32x.

---

4. Zero-Copy UMA Memory Mapping from Python to Metal

The Problem

In traditional CUDA workflows (CuPy, PyTorch), graph data follows this path:

Python NumPy array -> CPU RAM -> PCIe bus -> GPU VRAM

This transfer accounts for approximately 15-17% of total routing time for typical PCB graphs.

Resolution

Apple Silicon's Unified Memory Architecture (UMA) allows the CPU and GPU to share the same physical DRAM. By using MTLResourceStorageModeShared with new_buffer_with_bytes_no_copy, raw NumPy array memory pointers (accessed via Rust PyO3 PyReadonlyArray1) are mapped directly into Metal buffer objects.

When the Metal kernel finishes writing distances, the Python NumPy array instantly reflects the changes without any data copy. The PCIe transfer bottleneck is structurally eliminated.

This mapping is performed through a Rust intermediate layer using pyo3 and metal-rs:

let indptr_buf = device.new_buffer_with_bytes_no_copy(
    indptr_slice.as_ptr() as *const _,
    (indptr_slice.len() * mem::size_of::<i32>()) as u64,
    MTLResourceOptions::StorageModeShared,
    None,
);

---

5. Delta-Stepping Cache Amplification

The Problem

Naive Bellman-Ford examines every edge on every iteration. For large graphs, this means the GPU reads the full edge array from DRAM on each pass, limited by the M4's 120 GB/s memory bandwidth.

Resolution

Delta-Stepping partitions the frontier into buckets of width delta. Each iteration only processes nodes whose tentative distance falls within the current bucket. This means the working set of edges examined per iteration drops dramatically as the algorithm converges.

On the 401,800-node ClassAB MOSFET graph, Delta-Stepping processes only a small fraction of total edges per iteration. The accessed edges and distances fit entirely within the M4's L1/L2 cache hierarchy. The result is an effective bandwidth of 831.5 GB/s, which is 6.9x the M4's theoretical DRAM bandwidth of 120 GB/s. The multiplier is entirely explained by cache hit rates.

---

Performance Results

All benchmarks use corner-to-corner shortest path on CSR-format graphs derived from real PCB routing lattices. CUDA results were collected on Vast.ai instances (RTX 2080 Ti at $0.083/hr, RTX 3060 at $0.055/hr). Metal results were collected on a local Apple M4.

Traversal Time (microseconds, lower is better)

Graph (Nodes)	RTX 2080 Ti	RTX 3060	Apple M4 Metal
2,000	856	541	1,200
8,000	1,786	1,054	2,100
30,000	3,241	2,179	4,800
60,000	5,051	3,781	6,100
180,000	10,294	9,406	8,500
401,800	--	15,486	4,130

The crossover point occurs at approximately 150,000 nodes. Below that threshold, the CUDA GPUs have lower absolute latency due to higher clock speeds and more compute units. Above it, the M4's cache-amplified Delta-Stepping and zero-copy UMA dominate.

Throughput (Billion edges/sec, higher is better)

Graph (Nodes)	RTX 3060	Apple M4 Metal
2,000	0.020	0.257
30,000	0.073	4.817
180,000	0.111	30.034
401,800	0.124	111.4

The M4 throughput figures include total edge count divided by time. Because Delta-Stepping actively skips most edges via bucket-based frontier culling, the effective per-edge processing cost is extremely low.

Parity Verification

36 out of 36 parity tests pass across 6 graph sizes (2,000 to 180,000 nodes):

• Distances: bitwise float32 match against CUDA reference
• Paths: identical node sequences
• Reachable node counts: exact match

---

Key Takeaways

1. Apple Silicon GPU schedulers do not preempt. Persistent thread grids must fit within physical core concurrency or the hardware deadlocks.

2. MSL has no grid-level barrier. A software barrier using threadgroup_barrier(mem_flags::mem_device) with atomic generation counters is required for multi-iteration persistent kernels.

3. SIMD-group intrinsics are essential. simd_broadcast_first, simd_is_first, and simd_any reduce atomic contention by 32x and enable efficient work distribution.

4. UMA zero-copy is a structural advantage. Eliminating PCIe transfers saves 15-17% of total time and removes an entire class of synchronization complexity.

5. Algorithm choice matters more than hardware. Delta-Stepping on the M4 outperforms dense Bellman-Ford on discrete CUDA GPUs because it converts the problem from bandwidth-bound to cache-bound.

---

References

• McMurchie, L. and Ebeling, C. "PathFinder: A Negotiation-Based Performance-Driven Router for FPGAs." ACM/SIGDA FPGA, 1995.
• Meyer, U. and Sanders, P. "Delta-Stepping: A Parallelizable Shortest Path Algorithm." Journal of Algorithms, 2003.
• Apple. "Metal Shading Language Specification." Version 3.2, 2024.
• NVIDIA. "CUDA C++ Programming Guide." Version 12.6, 2024.