Bend 2 vs. CUDA & Metal: Manual GPU Shaders vs. the 128x128 Task Cube and the GPU Bimodality Law

When to write raw CUDA C++ or Metal Shading Language kernels vs. compiling Bend 2 (bend -o), explaining the GPU Bimodality Law across all 16 Apple M4 Max benchmarks.

Executive Summary: Bottom Line Up Front

Bend 2 is not a competitor to NVIDIA CUDA or Apple Metal; it is a high-level affine functional compiler that emits unified C, CUDA, and Metal shader code (bend prog.bend -o prog) (bendlang/bend).

1. Write raw CUDA C++ (.cu / PTX) or Metal Shading Language (.metal) when your workload is dense contiguous linear algebra (GEMM, FlashAttention, convolution stencils). Raw shaders saturate hardware Tensor Cores (mma.sync / simdgroup_matrix), explicit threadgroup shared memory (__shared__), and warp-shuffle intrinsics (__shfl_sync).
2. Write Bend 2 (f!(x) GPU dispatch) when your workload is a recursive divide-and-conquer tree, spatial grid, cellular automaton, particle simulation, or BVH raytracer. Instead of flattening pointer trees into index arrays, you write pure recursive functions. On an Apple M4 Max, Bend 2’s Metal backend achieves up to 123.9x GPU speedup (0.063 s on Conway’s Game of Life and 0.059 s on N-Body). It also verifies 3,200 formal proofs in 0.38 s (benchmark dataset).
3. Respect the GPU Bimodality Law: Our 16-workload Apple M4 Max benchmark suite proves that balanced parallel workloads run up to 123.9x faster on GPU. However, heavily branching irregular workloads (lexer, queens, symreg, hashmap) run 2.0x to 5.4x faster on 16 CPU cores (PAR-CPU) than on the GPU (PAR-GPU) (benchmark dataset).

1. How BendRT Compiles Functional Recursion to CUDA Warps and Metal Threadgroups

In traditional GPU programming, dynamic recursion and heap allocation inside a kernel either fail to compile or exhaust per-thread stack memory. Bend 2 (bend2/comp.ts and Lulzx/bendc) solves this by compiling .bend files into a flat state-machine shader backed by the 128x128 Task Cube (bendlang/bend):

• Zero GPU Call Stack (WL_CASE Worklist Loop): Recursive calls compile to register jumps inside a single shader loop rather than hardware call frames.
• Unified 8 TiB Address Space (MAP_NORESERVE): On Apple Silicon (Metal) and unified-memory architectures, CPU and GPU threads share the exact same 64-bit tagged pointer words ([rfc(1) | tag(7) | aux(16) | loc(40)]) with zero PCIe serialization (benchmark dataset).
• 16,384 Lock-Free Task Rings: Parallel forks (a b = f(x) g(y)) populate 128 × 128 = 16, 384 FIFO rings during the Grow phase and drain them without cross-lane locks during the Work phase (bendlang/bend).

2. The GPU Bimodality Law: All 16 Apple M4 Max Benchmarks (PAR-GPU vs. PAR-CPU)

Because Bend 2 runs the exact same source file on 1 CPU core (SEQ-CPU), 16 CPU cores (PAR-CPU), and Metal GPU warps (PAR-GPU), the M4 Max benchmarks reveal a clear hardware boundary (benchmark dataset):

Why 16 CPU Cores Beat the GPU on lexer, hashmap, queens, and symreg

1. SIMT Warp Divergence on Unbalanced Trees (queens, symreg): GPU hardware executes 32 threads in lockstep (a CUDA warp or Metal SIMD-group). In N-Queens backtracking (queens) and genetic AST evaluation (symreg), one lane prunes immediately while a neighboring lane in the same warp recurses 20 levels deeper, forcing the other 31 lanes to stall. Apple M4 Max’s 16 independent out-of-order CPU cores suffer zero warp divergence (benchmark dataset).
2. Pointer-Chasing Latency vs. Deep CPU Caches (hashmap, lexer): Traversing linked lists of characters (lexer) or trie nodes (hashmap) is memory-latency bound rather than compute bound, favoring the CPU’s low-latency L1/L2 caches (0.198 s PAR-CPU vs. 1.075 s PAR-GPU on lexer). Because Bend 2 lets you toggle CPU vs. GPU execution with a single ! annotation (or runtime flag), you can benchmark both targets in seconds without rewriting a single line of algorithm code (bendlang/bend).

3. Code Comparison: Flattened CUDA Array Indexing vs. Bend 2 f!(x) GPU Dispatch

Evaluating a recursive binary tree in raw CUDA C++ requires manually serializing the tree into flat device buffers (cudaMalloc). You must then manage parent/child integer index offsets and launch iterative level-by-level reductions from the host:

// CUDA C++: Manual flat-array reduction kernel (cannot express recursive algebraic trees directly)
__global__ void reduce_level_kernel(const uint32_t* __restrict__ in_nodes,
                                    uint32_t* __restrict__ out_nodes,
                                    size_t num_pairs) {
    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < num_pairs) {
        out_nodes[idx] = in_nodes[2 * idx] + in_nodes[2 * idx + 1];
    }
}

In Bend 2, you write standard algebraic recursion and append ! (sum!(lft)) to dispatch the parallel fork onto the GPU worklist shader via the 128 × 128 Task Cube (bendlang/bend):

type Tree:
  case @Leaf:
    val: U32
  case @Node:
    lft: Tree
    rgt: Tree

def sum(t: Tree) -> U32:
  match t:
    case @Leaf{val}:
      val
    case @Node{lft, rgt}:
      # '!' offloads recursive branches to Metal/CUDA GPU worklist lanes
      a b = sum!(lft) sum!(rgt)
      a + b

4. Architecture Decision Matrix: Raw CUDA / Metal vs. Bend 2

Related Guides

• Bend 2 Strengths, Weaknesses & Complete Decision Matrix
• Bend 2 vs. C++
• Bend 2 vs. Rust
• Bend 2 vs. Swift (Apple Silicon Metal)