Massively Parallel Execution Without Threads: Deconstructing Bend and Interaction Combinators on GPUs

The Parallel Concurrency Bottleneck: Beyond pthreads and CUDA Kernels

Writing high-performance parallel software has traditionally felt like making a pact with the devil. If you write high-level, expressive code in Python, TypeScript, or Go, your execution model is bound to single-threaded runtimes, garbage collection pauses, or asynchronous event loops running atop OS threads. If you need raw, hardware-saturating throughput, you descend into C++, Rust, or CUDA, manually orchestrating thread blocks, shared memory banks, warp shuffles, and memory barriers (__syncthreads()).

At the hardware level, modern GPUs like the NVIDIA H100 or RTX 4090 feature tens of thousands of arithmetic logic units (ALUs) capable of extraordinary floating-point operations per second (FLOPS). Yet, orchestrating these cores requires developers to translate algorithms into rigid, flat, Single Instruction Multiple Threads (SIMT) grids. Recursive tree traversals, dynamic graph algorithms, and functional closures typically fall apart on SIMT architectures due to heavy warp divergence and non-linear memory coalescing issues.

Enter Bend, an open-source programming language developed by the team at Higher Order Company (HOC). Bend allows developers to write pure, high-level functional code—complete with dynamic types, pattern matching, closures, and arbitrary recursion—while running natively on thousands of GPU cores without manual thread allocation, synchronization primitives, or CUDA kernel definitions.

To understand how Bend pulls off this seemingly impossible feat, we must peel back decades of imperative computing assumptions and examine its foundational runtime: Interaction Combinators.

Traditional GPU Programming:
Algorithm -> Flatten to Grids/Warps -> Manual Thread Synchs -> Hardware Execution

Bend GPU Programming:
Algorithm -> Interaction Net Graph -> Local Parallel Rewrites -> Native CUDA Cores

Interaction Combinators: The Theoretical Engine of Threadless Parallelism

In classical computer science, computation is viewed through the lens of either the Turing Machine (state manipulation on a tape) or the Lambda Calculus (variable substitution and beta-reduction). Both paradigms have severe shortcomings when ported to massively parallel hardware:

Turing Machines assume a centralized, global state, which leads to locking, cache invalidation, and data races in concurrent setups.
Lambda Calculus suffers from non-local substitutions during beta-reduction, requiring complex environments or substitution machinery that serialize execution.

In 1997, French logician Yves Lafont introduced Interaction Combinators, a minimalist model of computation based on graph rewriting. An Interaction Net is a graph where nodes (agents) have a distinct set of ports: one principal port and zero or more auxiliary ports.

The Core Axiom: Local Rewriting Rules

Computation in an interaction net happens exclusively when two agents connect through their principal ports. This arrangement is termed an active pair. When an active pair interacts, it performs an atomic graph rewrite step. Crucially, this rewrite depends only on the two interacting agents and affects only the immediate neighbors connected to their auxiliary ports.

Because every rewrite rule is strictly local:

No global state exists: Active pairs across disparate regions of the graph mutate memory with total spatial isolation.
Confluence (The Church-Rosser Property): If a graph can reduce in multiple places, the order of reduction does not change the final normal form. There are no race conditions by mathematical definition.
Implicit Parallelism: If a graph contains 100,000 active pairs, all 100,000 pairs can be reduced simultaneously in a single clock cycle if hardware allows.

   Agent A (Principal) <====> Agent B (Principal)
        /    \                     /    \
    Port 1  Port 2             Port 1  Port 2
             |                        |
             +-------[REWRITE]--------+

Bend compiles down to HVM2 (Higher-Order Virtual Machine 2), a runtime engine implemented in low-level C and native CUDA that maps these interaction graph rewrites directly onto GPU hardware.

Inside Bend: Syntax, Semantics, and Affine Logic

Bend provides a syntax resembling Python or Rust, but its evaluation strategy is entirely functional. The runtime avoids traditional stack frames; every function call expands an interaction net graph, which HVM2 reduces across available GPU execution units.

Consider calculating the sum of elements in a binary tree. In Python, deep recursion will blow the call stack (RecursionError), and multi-threading requires manual task splitting across multiprocessing.Pool:

# Standard Python recursive tree sum (Single-threaded, stack-bound)
def tree_sum(node):
    if node is None:
        return 0
    return node.value + tree_sum(node.left) + tree_sum(node.right)

In Bend, parallel tree construction and reduction are written as declarative functional transformations. Because any branch split represents independent subgraphs, parallelism is automatically extracted without user intervention:

# Define the Tree ADT
type Tree:
  Node { ~left, ~right }
  Leaf { val }

# Construct a balanced tree in parallel
def gen_tree(depth):
  if depth == 0:
    return Tree/Leaf { val: 1 }
  else:
    # The ~ operator denotes an affine reference capable of parallel expansion
    return Tree/Node { 
      left: gen_tree(depth - 1), 
      right: gen_tree(depth - 1) 
    }

# Sum all nodes in parallel
def sum_tree(tree):
  match tree:
    case Tree/Leaf:
      return tree.val
    case Tree/Node:
      return sum_tree(tree.left) + sum_tree(tree.right)

def main():
  # Generate a tree of depth 24 and sum it
  tree = gen_tree(24)
  return sum_tree(tree)

When executed with bend run-cu main.bend, the compiler transforms gen_tree and sum_tree into interaction nodes. While the tree branches expand, active pairs form exponentially: 1, 2, 4, 8, ... up to 16,777,216 leaf nodes. Every branch point is an autonomous rewrite step that can be processed by any available CUDA core.

Mapping Interaction Nets to the GPU Memory Hierarchy

Executing millions of dynamic graph operations on a GPU is notoriously difficult because GPUs prefer static, continuous arrays over dynamic, pointer-chasing topologies. HVM2 solves this impedance mismatch through a low-overhead 64-bit node memory layout.

The Node Architecture

In HVM2, nodes are represented as compact 64-bit words stored inside a pre-allocated global device buffer:

 63      56 55      48 47                                0
+----------+----------+-----------------------------------+
| Tag (8b) | Port(8b) |       Target Pointer (48b)        |
+----------+----------+-----------------------------------+

Tag (8 bits): Specifies the agent type (e.g., Constructor CON, Duplication DUP, Number NUM, Eraser ERA).
Port (8 bits): Distinguishes between the principal port (0) and auxiliary ports (1, 2).
Target (48 bits): A flat memory address locating the paired node inside the device's unified memory space.

Eliminating Thread Contention with Work-Stealing Bags

Rather than maintaining a central locking queue of active pairs—which would bottleneck the GPU's memory controllers—HVM2 partitions the global graph space into node bags mapped per CUDA Streaming Multiprocessor (SM):

Local Bags in Shared Memory: Each warp processes its local active pairs inside GPU shared memory ($L1$).
Atomic Redirection: When a rewrite produces new nodes that exceed local shared memory thresholds, nodes spill over into global DRAM using vectorized, lockless 64-bit atomic operations (atomicCAS / atomicExch).
Lock-Free Execution: Because interaction rules are structurally symmetric, two threads rewriting adjacent nodes never invalidate shared cache lines simultaneously.

[ GPU Streaming Multiprocessor (SM) ]
  ├── Warp 0 ── Local Active Bag (L1 Shared Mem) ── Local Rewrites
  ├── Warp 1 ── Local Active Bag (L1 Shared Mem) ── Local Rewrites
  └── Overflow Stack ──> Global Device Memory (DRAM via atomicCAS)

Tackling Warp Divergence

SIMT hardware requires that threads within a 32-lane warp execute the same instruction. If Thread 0 reduces a CON-CON pair while Thread 1 reduces an ERA-NUM pair, warp divergence occurs, serializing execution.

HVM2 minimizes divergence by sorting active pairs by interaction type. The scheduler bins pairs by opcode before dispatching blocks, ensuring that every thread lane inside a warp performs identical pointer adjustments concurrently.

Practical Tutorial: Compiling and Benchmarking a Bend Pipeline

To see Bend's runtime behavior in practice, let's walk through installing the environment, compiling a compute-heavy benchmark, and profiling the GPU acceleration.

Step 1: Environment Setup

Ensure you have Rust (>= 1.75) and the CUDA Toolkit (>= 12.0) installed on your system with a functional nvcc compiler.

# Install the High-Order Virtual Machine (HVM2)
cargo install hvm

# Install the Bend compiler
cargo install bend-lang

# Verify your CUDA device setup
nvcc --version

Step 2: Implementing Parallel Bitonic Merge Sort

Let's write a parallel bitonic sorter in Bend. This algorithm maps cleanly onto interaction nets, as divide-and-conquer sorting naturally generates non-interfering sub-graphs.

Create a file named sort.bend:

type List:
  Nil
  Cons { head, ~tail }

def merge(a, b):
  match a:
    case List/Nil:
      return b
    case List/Cons:
      match b:
        case List/Nil:
          return a
        case List/Cons:
          if a.head <= b.head:
            return List/Cons { head: a.head, tail: merge(a.tail, b) }
          else:
            return List/Cons { head: b.head, tail: merge(a, b.tail) }

def split(list):
  match list:
    case List/Nil:
      return (List/Nil, List/Nil)
    case List/Cons:
      match list.tail:
        case List/Nil:
          return (list, List/Nil)
        case List/Cons:
          let (left, right) = split(list.tail.tail)
          return (List/Cons { head: list.head, tail: left }, 
                  List/Cons { head: list.tail.head, tail: right })

def sort(list):
  match list:
    case List/Nil:
      return List/Nil
    case List/Cons:
      match list.tail:
        case List/Nil:
          return list
        case List/Cons:
          let (left, right) = split(list)
          return merge(sort(left), sort(right))

def make_list(n):
  if n == 0:
    return List/Nil
  else:
    return List/Cons { head: n, tail: make_list(n - 1) }

def main():
  # Generate a reversed list of 16,384 elements and sort it
  data = make_list(16384)
  return sort(data)

Step 3: Running and Benchmarking Across Devices

Bend allows executing the exact same code across three backends:

Interpreted mode (development/debugging):
```
bend run sort.bend
```
C Native backend (multi-threaded CPU execution):
```
bend run-c sort.bend
```
CUDA backend (massively parallel GPU execution):
```
bend run-cu sort.bend
```

On a test workstation featuring an AMD Ryzen 9 7950X (16 cores, 32 threads) and an NVIDIA RTX 4090 (16,384 CUDA cores), the execution profiles highlight how Interaction Combinators scale without changing a line of code:

Notice that the code contains no pthread_create, no channel orchestration, and no raw CUDA C <<<blocks, threads>>> syntax. The concurrency is mined automatically from the graph reduction tree.

The Engineering Trade-offs: When NOT to Use Bend

While Interaction Combinators provide incredible theoretical advantages, current implementations face real-world hardware limits that systems architects must consider before deploying Bend to production pipelines.

1. Pointer Chasing vs. Tensor Contiguity

GPUs excel at dense linear algebra because memory layouts are continuous; a matrix multiplication loads unified blocks into registers via optimized tensor cores. Bend models all computation as graph nodes connected via 64-bit pointers. Even with HVM2's node-sorting optimizations, random memory reads across dynamic graphs incur higher memory latency than hand-tuned CUDA kernels utilizing FP16/BF16 matrix instructions.

Dense Tensors (Traditional CUDA):
[ Float32 | Float32 | Float32 | Float32 ] -> Unified Cache Line Fetch

Interaction Nets (Bend/HVM2):
[ Node A ] ----(64-bit ptr)----> [ Node B ] ----(64-bit ptr)----> [ Node C ]
(Requires memory coalescing mitigations to avoid bus thrashing)

2. Node Graph Memory Expansion

During intermediate reductions, duplicate agents (DUP) can cause the node graph to temporarily explode in size before collapsing into its normal form. If a graph exceeds physical GPU VRAM during expansion, the runtime will trigger an unrecoverable out-of-memory error. Managing allocation boundaries without a garbage collection pause remains one of the primary engineering challenges being addressed in the HVM2 codebase.

3. Lack of Native Hardware Intrinsics

Currently, Bend compiles to high-level interaction nodes rather than lower-level Tensor Core instructions (e.g., mma.sync). As a result, using Bend for training large dense transformer models will not beat heavily optimized libraries like CUTLASS, FlashAttention-2, or TensorRT. Bend shines brightest on irregular, branching, dynamic workloads—such as theorem proving, AST processing, physics simulations, and symbolic AI—where standard CUDA kernels fail due to branch divergence.

The Road Ahead: Compilers That Eradicate Concurrency Bugs

For decades, high-performance computing required choosing between developer velocity and raw hardware efficiency. Imperative languages delegate synchronization to the engineer, resulting in memory safety vulnerabilities, thread deadlocks, and race conditions that take weeks to track down with tools like Valgrind or NVIDIA Compute Sanitizer.

Bend demonstrates that Yves Lafont's Interaction Combinators are not merely a theoretical curiosity of mathematical logic, but a viable blueprint for next-generation system compilers. By proving that pure, local rewrite rules guarantee confluence, Bend unlocks threadless parallelism across hardware architectures containing tens of thousands of cores.

As hardware developers move toward multi-chiplet GPUs and wafer-scale processors containing millions of compute nodes, managing execution threads manually will become increasingly impractical. Declarative, graph-based execution architectures like Bend offer a glimpse into the future of software engineering: a world where developers focus solely on solving algorithmic problems, leaving synchronization, parallel scheduling, and hardware allocation to the underlying logic of mathematics.