← Back to Architecture & Comparison Guides
Published: 2026-10-23 Updated: 2026-10-23 Source: bend2.org

Bend 2 vs. C++ (CUDA, OpenMP, Kokkos): Manual GPU Kernels & Concepts vs. the 128x128 Task Cube

Comparing C++20/23 (std::execution, OpenMP, Kokkos, and CUDA C++) against Bend 2’s 128x128 Task Cube and BendTT dependent proofs across CPU and GPU workloads.


Executive Summary: Bottom Line Up Front

C++ (paired with CUDA, HIP, OpenMP, and Kokkos) is the incumbent workhorse of high-performance computing, game engines, and AI inference runtimes (llama.cpp, TensorRT, Unreal Engine). However, C++ forces developers to choose between expressive recursive data structures on the CPU and flat, manually synchronized pointer arrays on the GPU.

  1. Choose C++ / CUDA / Kokkos for dense numerical linear algebra kernels (GEMM, FlashAttention, stencil solvers), real-time 3D rendering pipelines (Vulkan/DirectX 12), and low-latency trading infrastructure. In those domains, explicit SIMD intrinsics (AVX-512, NEON), warp-level primitives (__shfl_sync), and custom arena allocators are non-negotiable.
  2. Choose Bend 2 (bend -o, LAWS.bend, PROOF.bend) when your parallel workload is a recursive divide-and-conquer tree, irregular graph traversal, BVH raytracer, symbolic AST evaluator, or verified state machine. Instead of maintaining separate C++ host code, OpenMP task pragmas, and __global__ CUDA kernels with explicit cudaMemcpy transfers, you write one pure functional .bend file. BendRT compiles that file to a 128x128 Task Cube (16,384 lock-free rings) that achieves up to 123.9x GPU speedup (0.063 s on Game of Life) (benchmark dataset). It also verifies 3,200 generic proof instantiations in 0.38 s (benchmark dataset).

1. Parallel Programming Ergonomics: CUDA / OpenMP / std::execution vs. BendRT

1.1 What It Takes to Parallelize a Recursive Tree in C++ vs. Bend 2

Suppose you want to build and reduce a recursive binary tree (such as a cryptographic Merkle tree or bitonic sorting network) across both 16 CPU cores and GPU warps:

Parallel Stack Required Engineering Effort for Recursive Trees Memory & Synchronization Hazards
CUDA C++ (__global__) Must flatten pointers into contiguous index arrays (int left_idx[]), write explicit host-to-device cudaMemcpy, and manage block/grid dimensions (<<<blocks, threads>>>). Warp divergence (32x serialization penalty), device heap exhaustion, race conditions on shared memory.
OpenMP (#pragma omp task) Requires #pragma omp parallel + #pragma omp single + manual cutoff thresholds (if (depth < 12)) so Chase-Lev work-stealing deques do not drown in task allocation overhead. Forgetting shared() or firstprivate() clauses introduces silent data races; #pragma omp target fails on recursive pointers.
C++20 std::execution::par_unseq Restricted to flat random-access iterators (std::vector). Cannot express recursive tree forks without manually flattening each tree level into a frontier vector. Undefined behavior if the lambda touches any unsynchronized state or throws an exception.
Kokkos (Kokkos::parallel_for) Provides portable Kokkos::View multi-dimensional arrays across CUDA/HIP/OpenMP, but requires template functor boilerplate and explicit Kokkos::deep_copy between HostSpace and CudaSpace. Designed for dense rectangular stencils and CSR sparse matrices, not algebraic sum types or ASTs.
Bend 2 (a b = f(x) g(y)) Write a one-line parallel let (a b = build(d-1, i*2) build(d-1, i*2+1)). Run on 16 CPU cores or GPU warps (sum!(tree)) from the exact same source file (source). Zero locks, zero data races, zero manual cudaMemcpy: confluence and affine lifetimes are proven at compile time (source).

1.2 Inside the BendRT 128x128 Task Cube (Valiant Bulk-Synchronous Rhythm)

How does Bend 2 execute recursive functional code on a GPU without dynamic work-stealing lock contention? In bend2/comp.ts, BendRT organizes parallel execution into a fixed 128 × 128 = 16, 384 Task Cube of lock-free FIFO rings (RING_LEN = 1024) operating in two bulk-synchronous phases (bendlang/bend):

  1. Grow Phase (Row-Major Breadth-First Dealing): While the active task frontier is below 214 (16,384 tasks), active threads pop ring heads, evaluate parallel forks (a b = ...), and deal child tasks round-robin across rows of the cube.
  2. Work Phase (Column-Major Straight-Line Draining): Once the 16,384 rings are populated, the cube transposes in O(1) time (moving zero bytes of memory) to column-major indexing. Every GPU lane or CPU worker drains its assigned column ring sequentially on a register-pinned value stack with zero cross-lane locking and zero work stealing.
  3. Atomic Fetch-Sub Join (HOLE = 2^{64}-1): When two child tasks finish, they write their 64-bit result words into the parent join frame and execute a single atomic release fetch-sub (atomic_sub_explicit(&j->rem, 1)). Whichever lane decrements rem to zero immediately runs the continuation (bendlang/bend).

2. Apple M4 Max Benchmarks: Native Sequential C/C++ Baseline vs. Bend 2 CPU & GPU

Across the 16 benchmarks in /data/bend-ecosystem.json on an Apple M4 Max, unified memory lets BendRT’s 8 TiB address space share pointers between CPU and Metal GPU warps with zero PCIe copy overhead (benchmark dataset):

Workload Category Native -O3 Baseline (c_sec) Bend SEQ-CPU (RSS) Bend PAR-CPU 16c (RSS) Bend PAR-GPU Metal (RSS) GPU Speedup vs. SEQ
Conway’s Game of Life Cellular Automata 6.776 s 7.803 s (2.1 MB) 0.647 s (2.7 MB) 0.063 s (16.7 MB) 123.9x (source)
N-Body Physics Particle Gravity 5.126 s 5.758 s (2.2 MB) 0.516 s (2.7 MB) 0.059 s (16.8 MB) 97.6x (source)
Mandelbrot Set Complex Plane 3.750 s 4.473 s (2.3 MB) 0.395 s (2.9 MB) 0.057 s (16.9 MB) 78.5x (source)
Merkle Tree Cryptographic Trees 3.423 s 4.222 s (130 MB) 0.422 s (131 MB) 0.071 s (17.5 MB) 59.5x (source)
3D Raytracer BVH Graphics 3.646 s 4.616 s (2.1 MB) 0.396 s (2.6 MB) 0.151 s (17.0 MB) 30.6x (source)
Procedural Terrain Simplex Noise 2.198 s 2.225 s (2.1 MB) 0.197 s (3.3 MB) 0.106 s (16.9 MB) 21.0x (source)
Breadth-First Search Graph Traversal 3.585 s 3.918 s (2.0 MB) 0.343 s (2.7 MB) 0.207 s (16.8 MB) 18.9x (source)
Levenshtein Edit Dist Wavefront DP 2.002 s 2.559 s (2.0 MB) 0.236 s (3.1 MB) 0.144 s (16.8 MB) 17.8x (source)
Tree Bitonic Sort Sorting Networks 5.719 s 6.043 s (146 MB) 0.799 s (156 MB) 0.450 s (17.0 MB) 13.4x (source)
Tree Matrix Multiply Quadtree Algebra 2.964 s 2.064 s (2.8 MB) 0.226 s (14.8 MB) 0.231 s (17.3 MB) 8.9x (source)

3. Type Systems: C++20 concepts & constexpr vs. BendTT Dependent Proofs

C++ programmers often use template metaprogramming, constexpr, and C++20 concepts to enforce compile-time constraints. However, C++ templates and BendTT solve fundamentally different verification problems:

Dimension C++20 concepts / constexpr / RAII Bend 2 (BendTT + LAWS.bend + PROOF.bend)
What It Verifies Syntactic shape (“does type T have .begin()?”) and finite compile-time constant evaluation (static_assert). Universal mathematical theorems over all possible runtime inputs (for xs: List(U32) {sort_len(xs) == len(xs)}) (source).
Memory Safety Manual discipline (std::unique_ptr, std::span); dangling references, iterator invalidation, and use-after-move remain possible. Compiler-enforced Affine Usage Wall: variables are used at most once; consumed constructors free immediately (source).
Data-Race Prevention None in the type system (std::thread and OpenMP allow unsynchronized shared mutation). Guaranteed by Church-Rosser confluence (mechanized in 20,981 lines of Lean 4) (source).
Verification Latency Deep template instantiations slow compilation to tens of seconds. Verifies 3,200 generic proof instantiations in 0.38 s on Apple M4 Max (source).
Interop Path extern "C" ABI bend kernel.bend -o kernel.c links directly into CMake/Clang C++ targets (source).