Introduction

Enki is a heterogeneous compute platform for Rust. It enables writing algorithms in standard Rust syntax and executing them natively across CPU cores or compiling them just-in-time (JIT) onto GPU hardware.

#![allow(unused)]
fn main() {
use enki::*;

#[nam]
fn transform(space: &Space, in_val: &f32, out_val: &mut f32, scale: f32) {
    *out_val = *in_val * scale;
}
}

The goal of the project is to explore how far a modern systems language can bridge the gap between host logic and accelerator compute without introducing foreign languages, complex binding tables, or abandoning Rust’s type-safety invariants.

Motivation and Context

General-purpose GPU computing (GPGPU) has historically developed around dedicated ecosystems:

• Specialized Platforms (CUDA, OpenCL): While powerful, these frameworks require writing kernels in vendor-specific dialects (CUDA C++) or older C99 subsets (OpenCL), introducing language barriers and separate compilation pipelines. Newer initiatives like Mojo approach this by designing an entirely new language for accelerator programming.
• Graphics Shading Languages (GLSL, HLSL, WGSL): These languages were architected primarily around the graphics rendering pipeline (rasterizers, texture sampling, fragment processing). While compute shaders exist in these languages, they operate outside the host language’s type system and rely on manual descriptor bindings.

Enki does not seek to replace or compete with hardware-specialized frameworks like CUDA on their own terms. Instead, it investigates a specific systems-programming question:

Can standard Rust serve as both the host language and the accelerator language, without introducing a separate shading language, without descriptor set management, and with runtime invariants that preserve Rust’s reference semantics?

Architectural Foundations

The platform is structured around three main principles:

• Unified Language Surface: Functions marked with #[nam] are standard Rust functions. They can be invoked sequentially or in parallel on CPU threads (e.g., via rayon) for testing and debugging, or dispatched across an execution grid on GPU silicon.
• Physical Memory Addressing: Enki targets Vulkan 1.3 and uses 64-bit Buffer Device Addresses (BDA). Buffers are referenced directly through GPU virtual addresses, completely bypassing traditional descriptor set management.
• Dynamic Invariant Checking: The runtime includes a verification component (BorrowEngine) that checks dispatch arguments before command buffer submission to detect common concurrency bugs, such as overlapping mutable slice access and container size mismatches.

Project Status and Engineering Scope

Enki is currently an alpha-stage project (v0.1). It serves as an active exploration into heterogeneous systems design. When evaluating the platform, the following engineering boundaries should be noted:

Invariant Checking vs. Formal Verification

Enki implements dynamic invariant checking at dispatch time. For example, it verifies that sub-slices passed to a kernel do not overlap in memory, and that buffer allocations are sufficiently large for the requested dispatch domain.

These checks act as a pragmatic runtime guardrail. They do not constitute a formal, mathematically verified soundness proof for arbitrary parallel access patterns. When dispatches are executed in unchecked mode (run_unchecked), synchronization correctness remains the responsibility of the developer.

API Stability

The runtime interfaces, internal compiler passes, and data structures are subject to refinement as the system matures. The platform is designed for experimentation, evaluation, and community feedback.

System Requirements

Running Enki requires hardware and driver support for modern Vulkan features:

• Rust: Stable toolchain (1.80 or newer).
• Vulkan Driver: Support for Vulkan 1.3 or higher, with the following capabilities:
  • 64-bit Buffer Device Addresses (VK_KHR_buffer_device_address)
  • Timeline Semaphores (VK_KHR_timeline_semaphore)
  • Synchronization2 (VK_KHR_synchronization2)
  • 64-bit Shader Integers (shaderInt64)
• Supported Platforms: Linux and Windows.

Organization of This Book

• Getting Started: Environment setup, compilation prerequisites, and writing your first compute dispatch.
• The Mental Model: How Rust types and references map to physical GPU memory structures.
• Memory & Resource System: Details on GpuVec, buffer slicing, uniforms (GpuParam), and hardware atomics.
• Compiler Invariants: The constraints enforced by the JIT compiler backend when lowering Rust to GPU bytecode.
• The GPU BorrowEngine: The mechanics of dynamic contract matching and spatial safety checks.
• Graphics & Under the Hood: Interactive window presentation and the graph-based barrier generation engine.
• Diagnostics Reference: Catalog of diagnostic error codes and messages.

Getting Started with Enki

This section guides you through the practical fundamentals of installing, configuring, and executing compute workloads with Enki.

• Hardware & Toolchain Setup: Verifying driver support for Vulkan 1.3, understanding the required compiler flags (--emit=llvm-bc, opt-level = 2), and configuring the workflow.
• Your First Nam & Dual Debugging: Writing a vector scaling compute kernel, allocating VRAM buffers, dispatching to GPU hardware, and validating output using native CPU testing.

Hardware & Toolchain Setup

This chapter covers setting up your development environment, verifying hardware capabilities, and configuring the compilation toolchain for Enki.

1. Graphics Driver Verification

Enki targets Vulkan 1.3 directly through physical device addresses.

Note on Vulkan SDK: You do not need to install the full LunarG Vulkan SDK to build or run Enki applications. A standard, up-to-date graphics driver provided by your GPU vendor (NVIDIA, AMD, or Intel) with Vulkan 1.3 support is completely sufficient.

To verify that your active GPU driver supports Vulkan 1.3 and the required features, you can run the standard diagnostic tool:

vulkaninfo --summary

Ensure that:

• The reported vulkanVersion is at least 1.3.xxx.
• 64-bit Buffer Device Addresses (bufferDeviceAddress) are supported.

2. Adding Enki to Your Project

Create a new binary project using Cargo:

cargo new my_enki_app
cd my_enki_app

Add the enki-gpu runtime to your dependencies:

cargo add enki-gpu

If you plan to perform vector or matrix mathematics, we also recommend adding an external algebraic crate such as glam:

cargo add glam

3. The Execution Workflow

GPU JIT compilation requires two critical compiler flags to lower Rust code onto graphics silicon:

1. --emit=llvm-bc: Generates the LLVM bitcode representation required by the JIT backend (parsu).
2. opt-level = 2 & codegen-units = 1: Optimization is mandatory to flatten high-level Rust abstractions into native silicon primitives, while single code generation units ensure monolithic bitcode generation.

Enki provides two distinct workflows to manage these flags:

Path A: Automatic Project Configuration (cargo run)

This is the standard and recommended workflow. On the very first execution of your project:

cargo run

Enki will detect that .cargo/config.toml is missing the required flags and will prompt you:

warning: missing required compilation profile and LLVM bitcode for GPU JIT synthesis
  --> append configuration to `.cargo/config.toml`? [Y/n]

Pressing Enter (or typing Y) will automatically append the required profile blocks to .cargo/config.toml and re-run your application:

# --- Added by Enki for GPU JIT compilation ---
[profile.dev]
opt-level = 2
codegen-units = 1

[profile.dev.package."*"]
opt-level = 2

[profile.release]
opt-level = 2
codegen-units = 1

[profile.release.package."*"]
opt-level = 2

[target.'cfg(all())']
rustflags = ["--emit=llvm-bc"]
# ---------------------------------------------

From this point forward, you can build and run using standard cargo run and cargo run --release.

Path B: Non-Intrusive Execution (enki run)

If you prefer not to let the runtime modify your project’s .cargo/config.toml file, you can use the official CLI runner, cargo-enki.

Install the CLI tool:

cargo install cargo-enki

Now, execute your project using the enki subcommand:

cargo enki run
# or directly:
enki run

The CLI tool automatically injects the necessary compiler flags and optimization profiles via environment variables during invocation, leaving your project directory completely untouched.

4. The Compiler Backend (parsu)

Enki compiles Rust functions into optimized GPU compute pipelines through its JIT backend, parsu.

On your first dispatch, Enki will automatically download the pre-compiled parsu binary matching your operating system from official releases and cache it locally. No manual compiler toolchain installation is required.

Next Steps

With your environment configured, proceed to Your First Nam & Dual Debugging to write and dispatch your first compute kernel.

Your First Nam & Dual Debugging

In Enki, GPU compute kernels are called Nams (derived from the Sumerian concept of nam—physical decrees or governing principles).

A #[nam] is declared as a standard Rust function. It can be compiled and dispatched across parallel GPU execution grids, or invoked on host CPU threads for unit testing and verification.

1. Writing the Complete Example

Open src/main.rs in your project and replace its contents with the following:

use enki::*;
use glam::Vec2;

const COUNT: usize = 5;

// 1. Declare the compute kernel with #[nam]
#[nam]
fn scale_vectors(_space: &Space, input: &Vec2, output: &mut Vec2, factor: f32) {
    *output = *input * factor;
}

fn main() {
    // 2. Initialize the headless GPU runtime
    let enki = Enki::init();

    // 3. Allocate physical data directly in GPU VRAM
    let in_gpu = gpu_vec![
        Vec2::new(1.0, 2.0),
        Vec2::new(3.0, 4.0),
        Vec2::new(5.0, 6.0),
        Vec2::new(7.0, 8.0),
        Vec2::new(9.0, 10.0),
    ];
    let mut out_gpu = gpu_vec![Vec2::ZERO; COUNT];

    let factor = 2.5f32;

    // 4. Record and submit the GPU execution flow
    enki.flow(|_| {
        scale_vectors.run(
            &Space::gpu_x(COUNT),
            &in_gpu,
            &mut out_gpu,
            GpuParam::new(factor),
        );
    });

    // 5. Dual CPU Execution: Run the exact same function on CPU threads
    let in_cpu = in_gpu.to_vec();
    let mut out_cpu = vec![Vec2::ZERO; COUNT];

    for i in 0..COUNT {
        scale_vectors(&Space::cpu_x(i, COUNT), &in_cpu[i], &mut out_cpu[i], factor);
    }

    // 6. Verify bit-for-bit equivalence
    assert_eq!(&out_cpu[..], &out_gpu.to_vec()[..]);

    println!("GPU Results: {:?}", out_gpu.to_vec());
    println!("Execution verified: CPU and GPU outputs match identically!");
}

Run the application:

cargo run

2. Anatomy of the Dispatch

Let us examine each component of this program:

The #[nam] Signature

#![allow(unused)]
fn main() {
#[nam]
fn scale_vectors(_space: &Space, input: &Vec2, output: &mut Vec2, factor: f32)
}

• The Execution Context (&Space): The first argument to every nam must be &Space. It provides execution metadata such as invocation coordinates (space.x), spatial grid bounds, and tile indices.
• The SPMD Per-Cell Principle: Notice that input is declared as &Vec2 and output as &mut Vec2. During execution, Enki maps each thread in the grid exclusively to a single element in the vector.
• By-Value Parameters: factor: f32 is a scalar uniform. Values passed by value are packed into the engine’s internal uniform arena (ParamArena).

Allocating GPU Memory (GpuVec)

#![allow(unused)]
fn main() {
let in_gpu = gpu_vec![Vec2::new(1.0, 2.0), ...];
let mut out_gpu = gpu_vec![Vec2::ZERO; COUNT];
}

GpuVec<T> allocates an owned contiguous buffer in GPU VRAM referenced via a 64-bit Buffer Device Address (BDA). The gpu_vec! macro mirrors the ergonomics of std::vec!.

The Active Flow (enki.flow)

#![allow(unused)]
fn main() {
enki.flow(|flow| {
    scale_vectors.run(
        &Space::gpu_x(COUNT),
        &in_gpu,
        &mut out_gpu,
        GpuParam::new(factor),
    );
});
}

• A flow records a sequence of compute and transfer operations into a command buffer synchronized via timeline semaphores.
• Space::gpu_x(COUNT) defines a 1D execution grid of COUNT threads.
• GpuParam::new(factor): To satisfy the host-to-device contract, any uniform passed by value must be wrapped in GpuParam::new(...) on the host side.

3. The Dual CPU/GPU Debugging Workflow

Because a #[nam] function is a standard Rust function, it can be executed directly on CPU threads:

#![allow(unused)]
fn main() {
for i in 0..COUNT {
    scale_vectors(&Space::cpu_x(i, COUNT), &in_cpu[i], &mut out_cpu[i], factor);
}
}

Or executed in parallel using multi-threaded iterators like rayon:

#![allow(unused)]
fn main() {
use rayon::prelude::*;

out_cpu.par_iter_mut().enumerate().for_each(|(i, out)| {
    scale_vectors(&Space::cpu_x(i, COUNT), &in_cpu[i], out, factor);
});
}

This dual-execution model allows developers to:

1. Write standard Rust #[test] unit tests that execute in headless CI environments without requiring physical GPU hardware.
2. Step through kernel logic using conventional CPU debuggers (gdb, lldb).

4. The Boundaries of CPU Equivalence (An Active Research Area)

While bit-for-bit equivalence between CPU and GPU execution is achieved for element-wise operations, it is critical to understand where CPU simulation currently diverges from silicon execution.

Where Equivalence Holds (The Deterministic Safe Zone)

Equivalence between host CPU iteration and GPU dispatches is guaranteed for independent, element-wise algorithms:

• Per-cell data transformations (&T to &mut T).
• Pure mathematical evaluations (matrix math, color conversions, ray generation).
• Algorithms with strictly isolated, unshared memory access per invocation.

Where Divergence Occurs (Hardware-Coupled Features)

Divergence arises when algorithms utilize specialized graphics silicon hardware features:

1. On-Chip Shared Memory (GpuTileMem) & Barriers (Space::sync()): On GPU silicon, GpuTileMem allocates high-speed on-chip SRAM shared across cooperative threads in a workgroup. Threads execute in lockstep and synchronize at physical memory barriers (Space::sync()).

   A standard sequential CPU loop (for i in 0..N) does not emulate cooperative thread pausing. When iteration 0 encounters Space::sync(), it cannot suspend its state to wait for iteration 100 to reach the same barrier.

2. Concurrent Atomics (GpuAtomic): When hundreds of GPU threads perform concurrent atomic operations (such as .fetch_add()), the resolution order depends on non-deterministic hardware warp scheduling. A sequential CPU loop processes operations in a rigid linear order.

3. Floating-Point Rounding Modes: GPUs commonly execute fused multiply-add (FMA) instructions and fast-math reciprocal approximations that may yield slight least-significant-bit discrepancies compared to x86 CPU FP32 pipelines.

Engineering Note: An Active Research Track
Simulating cooperative workgroup memory and barrier synchronization accurately on the CPU requires complex compiler transformations (such as continuation-passing loop splitting or coroutine fiber scheduling).

Developing an automated, zero-overhead CPU workgroup simulation harness that accurately mirrors GpuTileMem and Space::sync() across host threads is an active area of ongoing research within the Enki project. Currently, bit-for-bit testing is intended primarily for SPMD safe-mode workloads.

The Mental Model

To write effective GPU algorithms with Enki, developers do not need to learn a foreign shader language. However, it is essential to understand how standard Rust types map onto physical GPU hardware.

This section deconstructs the conceptual and architectural bridges between host Rust and graphics silicon:

• Rust References to 64-bit BDA: How Enki completely eliminates Vulkan descriptor sets by routing 64-bit Buffer Device Addresses (BDA) through push constants.
• The SPMD Per-Cell Principle: The conceptual distinction between 1:1 scalar cell access (&mut T) and unconstrained global slices (&mut [T]).
• Spatial Execution & Topology: Navigating multi-dimensional problem spaces (1D, 2D, 3D), workgroup tiles, and coordinate helper functions using Space.

Rust References to 64-bit BDA

In traditional GPU programming models (Vulkan, DirectX 12, WebGPU), passing memory to a compute shader requires constructing Descriptor Tables:

[Host Code] ──> Descriptor Pool ──> Descriptor Set ──> Pipeline Layout ──> Shader Binding

This indirection introduces fragile binding indices (@binding(0)), descriptor allocation overhead, and a disconnect between host and device type systems.

Enki takes an alternative approach enabled by Vulkan 1.3: physical Buffer Device Addresses (BDA).

1. What is Buffer Device Addressing (BDA)?

On modern GPU hardware, VRAM allocations reside in a unified 64-bit virtual address space. With the VK_KHR_buffer_device_address extension:

• Every buffer allocated on the GPU possesses a native 64-bit hardware address (u64).
• Compute shaders can dereference these addresses directly through physical machine registers, exactly like pointers in standard CPU architectures.

By leveraging BDA, Enki bypasses descriptor sets entirely. Buffers are treated as direct memory pointers rather than bound resources.

2. The Ingress Architecture: 32-Byte Push Constants

When you dispatch a #[nam] inside an enki.flow, how do these 64-bit addresses reach the GPU threads?

Instead of binding descriptor sets, Enki configures a single, high-speed 32-byte Push Constant block that is passed directly to graphics hardware registers:

                          32-Byte Push Constant Register
┌───────────────────┬──────────────┬──────────────┬──────────────┬─────────┬───────────────────┐
│ Bytes 0..8        │ Bytes 8..12  │ Bytes 12..16 │ Bytes 16..20 │ 20..24  │ Bytes 24..32      │
│ ParamArena BDA    │ Grid Size X  │ Grid Size Y  │ Grid Size Z  │ Padding │ Stack Buffer BDA  │
└─────────┬─────────┴──────────────┴──────────────┴──────────────┴─────────┴───────────────────┘
          │
          ▼
┌──────────────────────────────────────────────────────────────────────────────────────────────┐
│ GPU ParamArena (Contiguous Uniform Memory)                                                   │
│                                                                                              │
│  [0..8]   Arg 0: Base BDA of GpuVec A (0x7F00_1200)                                          │
│  [8..24]  Arg 1: Slice BDA + Element Count (0x7F00_3400, 1024)                               │
│  [24..28] Arg 2: By-Value Uniform Float (e.g. dt: 0.016f32)                                  │
└──────────────────────────────────────────────────────────────────────────────────────────────┘

1. ParamArena BDA (Bytes 0..8): Points to a linear uniform buffer allocated on the device. All function arguments—whether they are scalar uniforms or buffer device addresses—are packed sequentially into this arena.
2. Grid Dimensions (Bytes 8..20): Conveys the physical $(X, Y, Z)$ bounds of the current dispatch domain as three 32-bit integers.
3. Stack BDA (Bytes 24..32): Points to dedicated VRAM stack space if the kernel requires dynamic stack frames (aligned after 4-byte padding).

This means zero descriptor set allocations and zero pipeline layout rebinding occur across dispatches.

3. How Rust Types Map to Silicon Memory

When you declare arguments in a #[nam] function, Enki maps them into physical silicon structures according to standard Rust semantics:

In the next chapter, we explore how this direct addressing model forms the foundation of The SPMD Per-Cell Principle.

The SPMD Per-Cell Principle

Most GPU architectures operate on the Single Program, Multiple Data (SPMD) execution model. In this model, thousands of physical threads execute the same compiled instruction stream simultaneously, each operating on distinct data elements.

Enki leverages Rust’s type system to express SPMD execution directly, eliminating manual thread-index calculations.

1. Traditional Indexing vs. Type-Level SPMD

In traditional GPU compute environments (CUDA, OpenCL, WGSL), every kernel receives raw array pointers and requires developers to manually compute linear thread coordinates:

The Traditional Approach (CUDA / C99)

__global__ void scale_kernel(float* in, float* out, float factor, int total_elements) {
    // Manual 1D grid stride calculation
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    
    // Manual boundary guarding
    if (i < total_elements) {
        out[i] = in[i] * factor; // Unchecked raw array access
    }
}

This model places several burdens on the developer:

1. Arithmetic errors in the index calculation cause silent memory corruption or out-of-bounds faults.
2. The language has no concept of whether a thread has exclusive ownership of index i or if other threads are writing to the same location concurrently.

The Enki Approach (Type-Level SPMD)

In Enki, you write the kernel as if it operates on a single scalar element:

#![allow(unused)]
fn main() {
#[nam]
fn scale_nam(_space: &Space, in_val: &f32, out_val: &mut f32, factor: f32) {
    *out_val = *in_val * factor;
}
}

When you dispatch this nam with scale_nam.run(&Space::gpu_x(COUNT), &in_vec, &mut out_vec, GpuParam::new(factor)):

• Thread Isolation: The JIT compiler and runtime lower in_val: &f32 to address base_in + (thread_id * 4) and out_val: &mut f32 to base_out + (thread_id * 4).
• Zero Manual Indexing: Each thread receives an isolated reference to its own element.
• Physical Safety by Construction: Because thread A accesses element A and thread B accesses element B, concurrent write operations are spatially disjoint. Rust’s rule—that a mutable reference &mut T represents exclusive access—is preserved across parallel silicon cores.

2. Per-Cell References vs. Indexed Slices

Not all GPU algorithms are strictly element-wise. Algorithms such as stencils, convolutions, physical simulations, and raymarchers require threads to inspect neighboring elements or traverse arbitrary memory locations.

Enki differentiates these access modes through Rust’s container types:

3. Dispatch Modes: .run() vs. .run_unchecked()

Enki provides two dispatch methods to balance safety and performance:

Safe Mode (.run())

Used for standard dispatches where data flow is element-wise (&mut T) or reads are shared (&[T]). In this mode, the runtime ensures that thread writes are provably isolated and mutually disjoint.

#![allow(unused)]
fn main() {
// Safe, automated 1:1 per-cell dispatch
scale_nam.run(&Space::gpu_x(COUNT), &in_vec, &mut out_vec, GpuParam::new(factor));
}

Unchecked Mode (.run_unchecked())

When your algorithm requires arbitrary index writes across a mutable slice (&mut [T])—such as scattering particles into a screen buffer or custom tiled sorting—you dispatch using unsafe { nam.run_unchecked(...) }:

#![allow(unused)]
fn main() {
unsafe {
    render_nam.run_unchecked(
        &Space::gpu_x(COUNT),
        &particles,
        &mut screen_buffer.as_mut_slice(),
    );
}
}

The unsafe block explicitly documents that writes across the slice are coordinated manually by your algorithm, mirroring standard Rust systems programming semantics.

Spatial Execution & Topology

Every #[nam] function declared in Enki takes an immutable reference &Space as its first parameter:

#![allow(unused)]
fn main() {
#[nam]
fn my_nam(space: &Space, ...) { ... }
}

The Space struct represents the spatial execution topology of your workload. It encapsulates global invocation coordinates, local tile (workgroup) geometry, problem boundaries, linear indexing helpers, and graphics projections.

1. Defining Execution Domains

Enki provides ergonomic constructors to define 1D, 2D, or 3D problem spaces:

#![allow(unused)]
fn main() {
// 1D Domain (e.g. 10 million elements)
let space_1d = Space::gpu_x(10_000_000);

// 2D Domain (e.g. 1920x1080 framebuffer)
let space_2d = Space::gpu_xy(1920, 1080);

// 3D Domain (e.g. 256x256x128 volumetric simulation)
let space_3d = Space::gpu_xyz(256, 256, 128);
}

Workgroup Tiling Configuration

By default, Enki automatically derives optimal workgroup tile dimensions based on your physical GPU’s hardware profile (such as subgroup size and compute unit limits).

If your algorithm requires specific tile dimensions (for example, when using shared tile memory GpuTileMem), you can chain the .tile(...) method:

#![allow(unused)]
fn main() {
// 1D: Fixed tile of 256 threads
let space = Space::gpu_x(1_000_000).tile(256);

// 2D: 16x16 tile rectangles
let space = Space::gpu_xy(1920, 1080).tile([16, 16]);

// 3D: 8x8x4 volumetric blocks
let space = Space::gpu_xyz(256, 256, 128).tile([8, 8, 4]);
}

2. Navigating Coordinates

Inside a #[nam], Space provides structured coordinates describing where the current thread is executing

Global Coordinates

• space.x, space.y, space.z: Global invocation coordinates along each axis.
• space.pos_xy() / space.pos_xyz(): Returns tuple pairs or triplets of global coordinates.
• space.index(): Computes the standard 1D linear flat-memory index:

  index = x + y * size_x + z * size_x * size_y
  

Local Tile Coordinates

• space.cell_x, space.cell_y, space.cell_z: The relative cell coordinate of the thread inside its current tile.
• space.tile_x, space.tile_y, space.tile_z: The index of the active tile within the global tile grid.
• space.cell_index(): Computes the 1D linear index of the thread inside its tile scratchpad.

3. Boundary & Coordination Helpers

When grid dimensions are not exact multiples of workgroup tile sizes, trailing threads may be dispatched beyond the logical problem size. Space provides boundary-guarding helpers:

#![allow(unused)]
fn main() {
#[nam]
fn guarded_kernel(space: &Space, pixel: &mut u32) {
    // Exit if the thread falls outside the logical image boundaries
    if !space.in_bounds_xy() {
        return;
    }

    *pixel = 0xFF00FF00;
}
}

• space.in_bounds_x() / space.in_bounds_xy() / space.in_bounds(): Returns true if the thread is strictly within (size_x, size_y, size_z).
• space.is_tile_leader(): Returns true only for the first cell in a tile (cell == (0, 0, 0)). Useful for intra-tile reductions when writing a single aggregated result back to global memory after space.sync().
• space.is_global_leader(): Returns true strictly for the global origin (pos == (0, 0, 0)).

4. Graphics & Color Projections

For raytracing, procedural textures, and presentation pipelines, Space provides pre-calculated normalized coordinates:

• space.uv(): Computes normalized UV coordinates in [0.0, 1.0] with standard half-pixel centering (+0.5):

  u = (x + 0.5) / size_x
  v = (y + 0.5) / size_y
  

• space.ndc(): Computes Normalized Device Coordinates in [-1.0, 1.0] for ray direction generation.
• space.aspect_ratio(): Computes the aspect ratio (size_x / size_y).

Direct Color Packing

To write pixels directly to presentation buffers without third-party color conversion routines, Space includes direct color packing methods:

#![allow(unused)]
fn main() {
// Encodes RGB floats (0.0..1.0) into packed 32-bit integer (0xAARRGGBB)
*pixel = space.set_rgb_color(r, g, b);

// Encodes RGBA with custom alpha
*pixel = space.set_rgba_color(r, g, b, a);
}

5. The CPU Counterpart

When testing or debugging kernels on the CPU, you construct a single-point Space instance representing the active thread:

#![allow(unused)]
fn main() {
for y in 0..HEIGHT {
    for x in 0..WIDTH {
        // Construct a 2D CPU execution point
        let space_cpu = Space::cpu_xy(x, y, WIDTH, HEIGHT);
        
        my_nam(&space_cpu, &mut host_buffer[x + y * WIDTH]);
    }
}
}

This ensures that all coordinate methods (space.x, space.uv(), space.index()) return identical values regardless of whether execution occurs on physical GPU hardware or on host CPU threads.

Memory & Resource System

GPU hardware operates on distinct physical memory tiers with differing latency, bandwidth, and scope characteristics:

• Global VRAM (High Bandwidth Memory): Dedicated device memory accessible by all GPU compute units.
• ParamArena (Uniform Memory): High-speed linear payload space used for passing configurations and addresses into shader registers.
• On-Chip Shared Memory (SRAM / LDS): Ultra-low-latency scratchpad memory physically local to each compute workgroup.

Enki mirrors these hardware tiers through five distinct Rust abstractions:

Chapter Overview

• Owned VRAM with GpuVec: Allocation, capacity management, physical cloning, and timeline synchronization.
• Zero-Cost Slicing (Slice & SliceMut): View creation, BDA address arithmetic, and disjoint mutable splitting.
• By-Value Uniforms (GpuParam): Packing arbitrary structs and configuration types into the uniform arena.
• Hardware Atomics (GpuAtomic): Managing concurrent scalar and vector atomics across GPU threads.
• On-Chip Scratchpad (GpuTileMem): Leveraging local workgroup SRAM and coordinating access with Space::sync().

Owned VRAM with GpuVec

GpuVec<T> represents a contiguous, growable array allocated directly in GPU VRAM. It serves as the primary owned container for large datasets in Enki.

#![allow(unused)]
fn main() {
use enki::*;

// Allocate directly in GPU device memory
let mut buffer: GpuVec<f32> = gpu_vec![1.0, 2.0, 3.0, 4.0];
}

1. Allocation and Construction

GpuVec<T> mirrors standard Vec<T> constructors while allocating memory through the physical device allocator (apsu):

• GpuVec::new() / with_capacity(cap): Allocates a storage buffer with Buffer Device Address (BDA) capability.
• GpuVec::from_slice(&[T]): Allocates VRAM and performs a synchronous host-to-device upload.
• GpuVec::from_elem(value, count): Allocates and initializes count elements with value.
• GpuVec::zeroed(count): Allocates and zeroes out device memory via transfer commands.

Zero-Sized Types (ZST): GPU silicon requires every memory access to have a concrete physical stride. Attempting to instantiate a GpuVec<()> or any struct with a byte size of zero will assert at initialization.

2. Ownership and Deep Cloning

GpuVec<T> follows standard Rust move semantics. When you call .clone() on a GpuVec:

#![allow(unused)]
fn main() {
let v1 = gpu_vec![42.0f32; 1000];
let v2 = v1.clone(); // Physical VRAM-to-VRAM copy
}

Enki allocates an entirely distinct physical buffer with a unique 64-bit BDA and unique engine slot ID, then executes a hardware memory transfer (vkCmdCopyBuffer) to copy the contents. Mutating v2 on the GPU has zero impact on v1.

3. Host Synchronization & Timeline Semaphores

Because GPU execution is asynchronous, GpuVec<T> tracks its execution phase relative to the engine’s Vulkan timeline semaphores:

              GPU Execution Timeline (Time Progresses ──>)
──[Submitted Batch (Timeline: 42)]──────[GPU Reaches 42 (Idle)]──>
             ▲                                      ▲
             │                                      │
   host.get(i) called here                GPU finishes work
   (CPU automatically blocks)             (Readback completes)

When you read data back to the host CPU:

#![allow(unused)]
fn main() {
// Reads a single element
let val: Option<f32> = buffer.get(0);

// Reads all elements back into a standard heap Vec<T>
let host_data: Vec<f32> = buffer.to_vec();
}

Enki automatically queries the underlying timeline semaphore. If the buffer is currently involved in executing GPU dispatches, the CPU thread automatically waits until the GPU finishes writing before transferring bytes across the PCIe bus.

4. Phase Invariants: Recording vs. Execution (error[E2001])

While reading back from an idle or in-flight buffer on the CPU is safe, doing so inside an active flow recording block is strictly forbidden:

#![allow(unused)]
fn main() {
enki.flow(|flow| {
    my_nam.run(&space, &mut buffer);

    // ERROR: Illegal host readback during command recording
    let val = buffer.get(0); 
});
}

Why is this rejected?

Inside enki.flow, compute commands are actively being recorded into a Vulkan command buffer—they have not yet been submitted to the GPU queue.

Attempting to read back buffer.get(0) at this moment would require the CPU to wait for work that has not even been submitted, causing a permanent CPU-GPU deadlock.

To prevent this, the runtime immediately aborts with a diagnostic error:

error[E2001]: illegal host readback during GPU command recording phase
   --> src/main.rs:302:49
    |
302 |                     let cpu_pixels = gpu_pixels.to_vec();
    |                                                 ^^^^^^^^ host readback attempted here
    |
    = note: dispatches inside `enki.flow` are recorded into a command buffer
            and have not been submitted to the GPU; waiting for results here
            causes a permanent CPU-GPU deadlock.
    = note: `to_vec()` was invoked on the CPU while recording GPU commands
            inside `enki.flow`
    = help: move readback operations (`.to_vec()`, `.get()`, or `println!`)
            outside the `enki.flow` closure.

Zero-Cost Slicing (Slice & SliceMut)

While GpuVec<T> owns physical allocations in VRAM, algorithms frequently need to operate on sub-ranges of data (such as image tiles, particle partitions, or sub-matrices) without copying memory.

Enki provides Slice<'a, T> (read-only) and SliceMut<'a, T> (exclusive mutable) to enable zero-allocation, borrowed sub-views over existing GPU buffers.

#![allow(unused)]
fn main() {
use enki::*;

let mut buffer = gpu_vec![0u32; 1024];

// Zero-allocation borrowed sub-views
let first_half: Slice<u32> = buffer.slice(..512);
let mut second_half: SliceMut<u32> = buffer.slice_mut(512..);
}

1. Zero-Cost Physical Address Arithmetic

Unlike traditional graphics APIs where creating a sub-buffer requires creating new Vulkan buffer views (VkBufferView) or allocating fresh descriptor sets, slicing in Enki is a purely mathematical zero-cost abstraction:

Parent GpuVec (Base BDA: 0x7F00_0000, Stride: 4 bytes)
┌──────────────────────────────────────┬──────────────────────────────────────┐
│ Elements [0..512]                    │ Elements [512..1024]                 │
└──────────────────────────────────────┴──────────────────────────────────────┘
 ▲                                      ▲
 │                                      │
 Slice A: 0x7F00_0000                   Slice B: 0x7F00_0000 + (512 * 4)
 (Device Address: 0x7F00_0000)          (Device Address: 0x7F00_0800)

When you slice a container:

device_address = parent_bda + (element_offset * stride)

Creating, cloning, or passing a slice performs zero GPU allocations, zero syscalls, and zero driver submissions. It simply passes a 16-byte fat pointer (bda: u64, count: u64) into the dispatch ingress.

2. Immutable Slices (Slice<'a, T>)

Slice<'a, T> represents a read-only view into a sub-range of GPU memory:

• Implements Clone: Multiple read-only slices derived from the same parent buffer can freely overlap and exist simultaneously.
• Kernel Binding: In a #[nam] function, passing &Slice<T> binds as an indexed global slice:

  #![allow(unused)]
  fn main() {
  #[nam]
  fn read_lookup(_space: &Space, table: &[f32], target: &mut f32) {
      *target = table[42]; // Free random read access
  }
  }

Slicing Sub-ranges

Slices can be recursively partitioned into smaller sub-views:

#![allow(unused)]
fn main() {
let view = buffer.slice(100..500);
let sub_view = view.slice(0..50); // Covers elements 100..150 of parent
let (left, right) = view.split_at(200);
}

3. Mutable Slices (SliceMut<'a, T>)

SliceMut<'a, T> enforces Rust’s exclusive write semantics over a VRAM sub-range:

• Does NOT implement Clone: Enforces single-writer exclusivity.
• Reborrowing: A SliceMut can be reborrowed as an immutable Slice via .as_slice() or consumed via .into_slice().
• Kernel Binding: In a #[nam], passing &mut SliceMut<T> binds as a global read-write slice &mut [T].

Disjoint Partitioning with split_at_mut

To safely divide a mutable buffer into two independent mutable slices, use split_at_mut:

#![allow(unused)]
fn main() {
let mut buffer = gpu_vec![0.0f32; 1000];
let mut slice = buffer.as_mut_slice();

// Guarantees mathematically disjoint ranges [0..500) and [500..1000)
let (mut left, mut right) = slice.split_at_mut(500);
}

4. Host Manipulation Methods

Both slice types support direct data inspection and modification from the host CPU. Like GpuVec, host readbacks automatically wait on the GPU timeline if the underlying buffer is in-flight:

#![allow(unused)]
fn main() {
let mut slice = buffer.slice_mut(0..100);

// Set single element
slice.set(0, 42.0);

// Bulk upload from CPU slice
let host_data = vec![1.0; 100];
slice.copy_from_slice(&host_data);

// Read back to host Vec
let cpu_copy: Vec<f32> = slice.to_vec();

// Clone into independent physical VRAM allocation
let separate_gpu_vec: GpuVec<f32> = slice.to_gpu_vec();
}

Phase Rule: Slices track the execution phase of their underlying physical allocation. Calling .to_vec() or .copy_from_slice() on a slice inside an active enki.flow block halts execution with error[E2001].

5. Runtime Bounds Verification (error[E2005])

Attempting to slice beyond the element capacity of a container halts execution before command buffers are recorded:

error[E2005]: GPU slice index out of bounds
  --> src/main.rs:18:28
   |
18 |     let sub = buffer.slice(500..2000);
   |                            ^^^^^^^^^^ invalid slice range specified here
   |
   = note: GPU slice bounds must reside strictly within the allocated VRAM buffer limits.
   = help: verify that range indices satisfy `start <= end` and `end <= slice.len()`.

By-Value Uniforms (GpuParam)

Compute kernels frequently require configuration parameters, transform matrices, time deltas, and camera settings.

In traditional GPU APIs, passing these variables requires writing uniform buffer objects (UBOs) or push constant structures with strict hardware alignment padding rules (std140 / std430).

Enki simplifies this through GpuParam<T> and the engine’s linear ParamArena.

1. What is GpuParam<T>?

GpuParam<T> is a transparent wrapper type that marks an argument as a by-value uniform:

#![allow(unused)]
fn main() {
use enki::*;

let dt = GpuParam::new(0.016f32);
}

Arbitrary Struct Support

GpuParam<T> is not restricted to primitive numbers (f32, u32). It supports any arbitrary Rust composite struct:

#![allow(unused)]
fn main() {
use glam::{Mat4, Vec3};

#[derive(Clone, Copy, Debug)]
pub struct CameraSettings {
    pub view_projection: Mat4,
    pub eye_position: Vec3,
    pub exposure: f32,
    pub max_bounces: u32,
}
}

2. The Host-to-Device Contract

Notice the intentional symmetry between how the argument is passed on the host versus how it is declared inside the kernel:

Inside the #[nam] Function

Inside the kernel signature, you declare the parameter directly by value as T:

#![allow(unused)]
fn main() {
#[nam]
fn raymarch_nam(
    space: &Space,
    pixel: &mut u32,
    camera: CameraSettings, // Received directly by value!
    delta_time: f32,        // Received directly by value!
) {
    let ray_dir = camera.eye_position;
    // ...
}
}

On the Host Dispatch Side

When calling .run() or .run_unchecked(), wrap the host instances in GpuParam::new(...):

#![allow(unused)]
fn main() {
let cam = CameraSettings { /* ... */ };
let dt = 0.016f32;

raymarch_nam.run(
    &space,
    &mut screen_buffer,
    GpuParam::new(cam),
    GpuParam::new(dt),
);
}

On Host CPU Invocations

When calling the kernel directly on host CPU threads for unit testing, pass the instances directly without GpuParam:

#![allow(unused)]
fn main() {
// In CPU tests, pass by value directly
raymarch_nam(&space_cpu, &mut cpu_pixel, cam, dt);
}

3. Under the Hood: The ParamArena

GpuParam<T> does not allocate a dedicated VkBuffer for each uniform. Dedicated GPU allocations introduce heavy PCIe allocation overhead.

Instead, the runtime maintains a high-speed pre-allocated ParamArena:

Host Memory                         GPU ParamArena (Linear Device Memory)
┌────────────────────────┐          ┌──────────────────────────────────────────────┐
│ cam: CameraSettings    │ ──Upload─► [Offset 0]   CameraSettings (72 bytes + pad) │
│ dt:  0.016f32          │          │ [Offset 80]  delta_time (4 bytes + pad)     │
└────────────────────────┘          └──────────────────────────────────────────────┘
                                                           ▲
                                                           │
                                   PushConstant[0..8] (BDA Base Address)

1. Sequential Packing: Prior to dispatch, the runtime packs the raw byte payloads of all by-value arguments sequentially into the ParamArena.

2. 8-Byte Alignment: Each uniform entry is automatically aligned to 8-byte boundaries: arena_size = (byte_size + 7) & !7

3. Register Pointer: The 64-bit BDA pointer to the arena base is passed into the first 8 bytes of the GPU push constant register. Invocations load uniforms directly from registers.

4. Hardware Limits & Configuration (error[E3010])

The default capacity of the ParamArena is typically 2 MB. If your dispatch passes large configurations that exceed the available arena capacity, execution halts before recording:

error[E3010]: nam parameters payload exceeds ParamArena capacity
  = note: the total byte footprint of by-value parameters (`GpuParam`) passed to
          this nam exceeds the pre-allocated capacity of the ParamArena.
  = note: configured arena size: x.xx MB
  = note: device maximum allocation capacity: x.xx MB (allocation clamped to
          prevent VRAM allocation fault)
  = help: increase the arena size using `Enki::builder().param_arena_size(...)`
          or pass large arrays by reference (`&GpuVec` or `&[T]`) instead of by
          value.

Customizing Arena Capacity

To customize the capacity allocated for uniform payloads, configure the engine builder during initialization:

#![allow(unused)]
fn main() {
let enki = Enki::builder()
    .param_arena_size(8 * 1024 * 1024) // Allocate 8 MB for uniforms
    .init();
}

Hardware Atomics (GpuAtomic & GpuAtomicVec)

Parallel GPU algorithms frequently require multi-threaded coordination, such as global compaction counters, work-queue distribution, and category binning.

In traditional shading languages, developers must rely on specialized shader intrinsics (e.g. atomicAdd(), atomicCompSwap()).

Enki maps GPU hardware atomic instructions directly to Rust’s standard library atomic types (core::sync::atomic).

1. The GpuAtomicTarget Trait

Enki defines the GpuAtomicTarget trait for integer primitives that have native hardware atomic support on graphics silicon:

Inside a #[nam], you interact with atomics using standard Rust atomic methods: .fetch_add(), .fetch_sub(), .fetch_min(), .fetch_max(), .load(), and .store().

2. Scalar Atomics (GpuAtomic<T>)

GpuAtomic<T> represents a single, isolated hardware atomic variable allocated in GPU VRAM:

use enki::*;
use core::sync::atomic::{AtomicU32, Ordering};

#[nam]
fn count_active_particles(space: &Space, speed: &f32, counter: &AtomicU32) {
    if !space.in_bounds_x() {
        return;
    }

    if *speed > 10.0 {
        // Direct hardware atomic addition on silicon
        counter.fetch_add(1, Ordering::Relaxed);
    }
}

fn main() {
    let enki = Enki::init();

    // Allocate an atomic counter in VRAM initialized to 0
    let counter = GpuAtomic::new(0u32);
    let speeds = gpu_vec![5.0f32, 12.0, 3.0, 15.0, 8.0];

    enki.flow(|_| {
        count_active_particles.run(&Space::gpu_x(5), &speeds, &counter);
    });

    // Read back to CPU (automatically waits on timeline)
    let total_active = counter.get();
    println!("Active particles (>10.0): {}", total_active);
    assert_eq!(total_active, 2);
}

3. Atomic Arrays (GpuAtomicVec<T>)

When algorithms require a shared array of atomic variables (such as category classification or spatial hash grids), use GpuAtomicVec<T>.

In this example, 1,000 parallel threads classify sensor temperature readings into 4 distinct alert buckets:

use enki::*;
use core::sync::atomic::{AtomicU32, Ordering};

const NUM_CATEGORIES: usize = 4;

#[nam]
fn classify_sensors(space: &Space, temp: &f32, buckets: &[AtomicU32]) {
    if !space.in_bounds_x() {
        return;
    }

    // Determine category: 0 = Normal, 1 = Warning, 2 = High, 3 = Critical
    let bucket_idx = if *temp < 25.0 {
        0
    } else if *temp < 50.0 {
        1
    } else if *temp < 75.0 {
        2
    } else {
        3
    };

    // Concurrent multi-threaded atomic increment into the target bucket
    buckets[bucket_idx].fetch_add(1, Ordering::Relaxed);
}

fn main() {
    let enki = Enki::init();

    // 1. Allocate an array of 4 atomic counters in VRAM
    let buckets = GpuAtomicVec::<u32>::new(&[0u32; NUM_CATEGORIES]);

    // 2. Prepare 1,000 sensor readings
    let mut host_readings = Vec::with_capacity(1000);
    for i in 0..1000 {
        host_readings.push((i as f32 * 0.1) % 100.0);
    }
    let sensor_data = GpuVec::from_slice(&host_readings);

    // 3. Dispatch across 1,000 threads in Safe Mode
    enki.flow(|_| {
        classify_sensors.run(&Space::gpu_x(1000), &sensor_data, &buckets);
    });

    // 4. Read back the 4 buckets to the host
    let results = buckets.to_vec().unwrap();
    println!("Sensor Classification Results:");
    println!("  Normal   (<25.0C):  {}", results[0]);
    println!("  Warning  (25..50C): {}", results[1]);
    println!("  High     (50..75C): {}", results[2]);
    println!("  Critical (>=75.0C): {}", results[3]);

    // Total classified items must equal 1,000
    assert_eq!(results.iter().sum::<u32>(), 1000);
}

4. Parallel Safety in Safe Mode

In standard dispatches, passing a mutable slice (SliceMut) across parallel threads triggers restrictions because unconstrained concurrent writes to arbitrary indices introduce data races.

With GpuAtomic<T> and GpuAtomicVec<T>, the situation is fundamentally different:

• Hardware atomic operations are atomic by definition at the silicon transistor level.
• Even if hundreds of threads attempt to execute .fetch_add(1) on the exact same bucket simultaneously, the GPU memory controller serializes the memory requests cleanly.

For this reason, Enki’s BorrowEngine explicitly permits concurrent shared atomic references in Safe Mode (.run()).

On-Chip Scratchpad Memory (GpuTileMem)

Modern graphics hardware features a hierarchy of memory tiers. While global VRAM provides high capacity, latency to access it remains relatively high.

To hide memory latency in communication-heavy algorithms (such as stencils, convolutions, Fourier transforms, and pairwise interactions), GPU architectures include dedicated on-chip Local Data Share (LDS) / Shared Memory (SRAM) physically situated adjacent to the compute execution cores.

Enki exposes this hardware capability through GpuTileMem<T, N> and intra-tile synchronization barriers.

1. Architectural Characteristics

GpuTileMem<T, N> differs fundamentally from GpuVec<T> and GpuParam<T>:

• Zero Allocation Footprint: GpuTileMem does not allocate physical VRAM and incurs a 0-byte footprint in the engine’s ParamArena. It simply informs the JIT compiler to allocate space in the workgroup’s hardware shared storage registers.
• Workgroup-Local Scope: Memory allocated through GpuTileMem is shared exclusively among threads executing within the same spatial tile. Threads in different tiles cannot observe each other’s tile memory.
• Compile-Time Sizing: The element capacity N must be a compile-time constant.

Inside a #[nam], passing &GpuTileMem<T, N> binds as a mutable fixed-size array reference:

#![allow(unused)]
fn main() {
&mut [T; N]
}

2. Cooperative Tiling Pattern

The standard architectural pattern for shared memory involves cooperative loading, synchronization, and consumption:

#![allow(unused)]
fn main() {
use enki::*;

const TILE_SIZE: usize = 256;

#[nam]
fn tiled_convolution(
    space: &Space,
    global_input: &[f32],
    output: &mut f32,
    tile_cache: &mut [f32; TILE_SIZE], // On-chip scratchpad
) {
    // 1. Cooperative Load: Each thread in the tile loads one element into SRAM
    tile_cache[space.cell_x] = global_input[space.x];

    // 2. Intra-Tile Barrier: Wait for all 256 cells to complete writing
    space.sync();

    // 3. Compute: Read from shared memory without accessing global VRAM
    let left = if space.cell_x > 0 { tile_cache[space.cell_x - 1] } else { 0.0 };
    let center = tile_cache[space.cell_x];
    let right = if space.cell_x < TILE_SIZE - 1 { tile_cache[space.cell_x + 1] } else { 0.0 };

    *output = (left + center + right) / 3.0;
}
}

On the host, the dispatch must explicitly define a tile size matching the scratchpad capacity:

#![allow(unused)]
fn main() {
let tile_mem = GpuTileMem::<f32, TILE_SIZE>::new();

tiled_convolution.run(
    &Space::gpu_x(100_000).tile(TILE_SIZE),
    &input_slice,
    &mut output_vec,
    &tile_mem,
);
}

3. Control Flow Uniformity & Divergence Hazards (error[E0009])

On GPU silicon, physical thread groups execute instructions in lockstep (SIMD / SIMT). A memory barrier (space.sync()) instructs the hardware execution unit to pause all threads in the workgroup until every thread reaches that barrier instruction.

The Divergence Invariant

Because the barrier operates on the physical workgroup as a collective unit, every cell in the tile must reach space.sync() unconditionally.

If a barrier is placed inside a branch that evaluates differently across threads within the same tile (a non-uniform branch), the hardware enters a permanent deadlock: threads taking the branch pause waiting for non-branching threads, which will never arrive.

#![allow(unused)]
fn main() {
// COMPILE ERROR: Divergent tile synchronization
if space.cell_x < 128 {
    // Only half the tile reaches this barrier! Permanent GPU hang.
    space.sync();
}
}

This rule applies equally to boundary checks:

#![allow(unused)]
fn main() {
// COMPILE ERROR: Early exit before barrier
if !space.in_bounds_x() {
    return; // Trailing threads exit, leaving active threads hung at space.sync()
}

tile_cache[space.cell_x] = global_input[space.x];
space.sync();
}

Enki’s JIT compiler analyzes the control flow graph (CFG) for barrier reachability. If a barrier is found within non-uniform conditional blocks, compilation is rejected immediately:

error[E0009]: divergent tile synchronization detected inside #[nam]
  --> src/main.rs:18:9
   |
15 |     if space.cell_x < 128 {
   |        ------------------ branch condition is non-uniform across tile cells
...
18 |         space.sync();
   |         ^^^^^^^^^^^^ tile synchronization called here
   |
   = note: all cells in a tile must reach `space.sync()` concurrently; barriers inside non-uniform control flow cause permanent GPU hardware deadlocks.
   = help: move `space.sync()` outside the conditional block or ensure the branch condition is uniform across the entire tile.

4. Hardware Capacity Verification (error[E1006])

Physical GPU compute units have strict hardware limits on total shared memory capacity (typically 32 KB, 48 KB, or 64 KB per workgroup depending on the device architecture).

The BorrowEngine verifies requested capacity (N * size_of::<T>()) against the queried hardware profile prior to execution:

error[E1006]: tile shared memory capacity exceeded
  --> src/main.rs:12:9
   |
12 |         kernel.run(&space, &large_tile_mem);
   |         ^^^^^^ requested tile scratchpad memory is too large
   |
   = note: on-chip intra-tile shared scratchpad memory (`GpuTileMem`) is physically constrained per compute unit.
   = help: reduce the element count `N` in `GpuTileMem<T, N>` or use a more compact element type.

5. CPU Simulation Boundaries

As outlined in the introduction, GpuTileMem and Space::sync() represent an area where standard CPU iteration diverges from GPU execution.

When executing on silicon, threads run cooperatively and communicate through hardware memory fences. On the host, a standard sequential loop (for i in 0..N) executes iterations strictly one after another. Iteration 0 completes entirely before iteration 1 begins; therefore, calling space.sync() on a CPU thread simply performs a host compiler fence (compiler_fence) without suspending execution to wait for other iterations.

Accurately simulating cooperative workgroup barriers and shared memory on host CPU threads requires complex compiler transformations (such as loop splitting or fiber coroutine scheduling), which remain an active research track within Enki.

Compiler Invariants & Silicon Rules

Enki lowers standard Rust functions into GPU compute pipelines by processing the LLVM bitcode generated by rustc.

However, physical graphics silicon lacks an operating system, dynamic memory allocators, virtual memory paging, and stack unwinding machinery. Consequently, certain language features that are standard on host CPU threads cannot physically execute on a GPU compute core.

This section covers the constraints enforced by the JIT compilation backend:

• Rejected Patterns (E0001 - E0010): Language constructs rejected during compilation (dynamic heap allocations, recursion, panic unwinding, dynamic dispatch).
• Divergent Control Flow (E0009): The physical mechanics of warp deadlocks and why non-uniform barriers are strictly forbidden.
• The Crash Flight Recorder (E9999): How the compiler isolates Internal Compiler Errors (ICE) into encrypted diagnostic tokens.

Rejected Patterns on Silicon (E0001 - E0010)

When compiling a #[nam], Enki’s backend analyzes the function’s intermediate representation (IR) to ensure it can execute deterministically on parallel hardware.

The following language patterns cannot be lowered to graphics silicon and will trigger compiler diagnostics:

1. Dynamic Heap Allocation (error[E0002])

GPU compute units execute across thousands of concurrent execution cells without a dynamic memory heap manager (no malloc or free).

Using types that allocate heap memory dynamically (such as Vec<T>, String, or Box<T>) inside a #[nam] is rejected:

error[E0002]: dynamic heap allocation is not supported inside #[nam]
  --> src/main.rs:13:17
   |
13 |     let mut v = vec![0f32; N];
   |                 ^^^^^^^^^ heap allocation occurs here
   |
   = note: GPU nams execute in parallel across silicon cells without a dynamic
           memory heap manager.
   = help: use fixed-size stack arrays `[T; N]` or pre-allocate a `GpuVec` on
           the CPU before dispatch.

2. Panics and Stack Unwinding (error[E0003])

Graphics silicon lacks stack unwinding tables and landing pad machinery. Code running on the GPU must be non-panicking.

Calling panic!(), .unwrap(), .expect(), or panicking index operations inside a #[nam] will fail compilation:

error[E0003]: explicit panic or unwinding assertion inside #[nam]
  --> src/main.rs:22:9
   |
22 |         panic!("invalid condition");
   |         ^^^^^^^^^^^^^^^^^^^^^^^^^^^ panic or assertion occurs here
   |
   = note: GPU silicon lacks stack unwinding and landing pad machinery; execution must be non-panicking.
   = help: avoid panicking APIs (e.g. `.unwrap()`, `assert!`) and handle bounds via conditional checks.

3. Dynamic Dispatch & Trait Objects (error[E0010])

GPU pipelines require static control flow and register allocation known at compile time. Functions cannot be dispatched dynamically through runtime virtual method tables (&dyn Trait):

error[E0010]: dynamic dispatch and indirect calls are not supported inside #[nam]
  --> src/main.rs:30:5
   |
30 |     trait_obj.execute();
   |     ^^^^^^^^^^^^^^^^^^^ indirect function call occurs here
   |
   = note: GPU pipelines require static control flow and cannot invoke functions through runtime pointers or vtables (`&dyn Trait`).
   = help: use concrete types, static generics (`impl Trait`), or an `enum` with pattern matching (`match`).

4. Unbounded Recursion (error[E0005])

GPU hardware execution threads operate on fixed register allocations rather than dynamically growing call stacks:

error[E0005]: recursion is not supported inside #[nam]
  --> src/main.rs:15:5
   |
 7 | fn recursive_call(
   | ^^^^^^^^^^^^^^^^^^ recursive call occurs here
 ...
15 |     recursive_call(space, global_input, output, tile_cache);
   |     ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ recursive call occurs here
   |
   = note: GPU nams execute on fixed registers without a dynamic call stack to
           allocate new data at runtime.
   = help: rewrite recursive algorithms using iterative loops (`for` / `while`).

5. Host Operating System Calls (error[E0001])

GPU execution units are hardware accelerators isolated from the CPU’s operating system kernel. Attempting to invoke file I/O, network sockets, thread spawning, or standard output inside a #[nam] is rejected:

error[E0001]: unsupported CPU system operation inside #[nam]
  --> src/main.rs:15:5
   |
15 |     println!("{:?}", tile_cache);
   |     ^^^^^^^^^^^^^^^^^^^^^^^^^^^^ unsupported CPU call occurs here
   |
   = note: GPU nams execute isolated on graphics silicon without access to CPU
           OS services (I/O, filesystem, threads).
   = help: perform CPU operations in the main thread and transfer data using
           `GpuVec` or `GpuParam`.

Summary of Compiler Diagnostic Invariants

Divergent Control Flow & Barrier Deadlocks (error[E0009])

One of the most catastrophic failure modes on graphics silicon is the barrier deadlock. Unlike CPU architectures where waiting on a synchronization primitive can timeout or throw exceptions, an invalid GPU memory barrier can hang the physical compute unit entirely, triggering an operating system Timeout Detection and Recovery (TDR) event or device reset.

This chapter details the physical execution model behind thread divergence and how Enki statically prevents barrier deadlocks at compile time.

1. Physical Warp Execution & Branch Divergence

GPU hardware schedules threads in lockstep groups called Warps (NVIDIA: 32 threads) or Wavefronts (AMD / Intel: 32 or 64 threads). All threads within a warp advance under a single shared program counter.

When code encounters a conditional branch:

#![allow(unused)]
fn main() {
if space.cell_x < 16 {
    do_something();
} else {
    do_other();
}
}

The hardware cannot physically branch in two different directions at the same clock cycle. Instead, it executes both sides sequentially through thread masking:

1. Active mask is set for cells 0..16; other cells are disabled while do_something() executes.
2. Active mask is inverted for cells 16..32; cells 0..16 are disabled while do_other() executes.

This serialization is known as Branch Divergence. While performance degrades, execution remains functionally correct for independent computation.

2. The Anatomy of a Barrier Deadlock

A synchronization barrier (space.sync()) lowers to a hardware instruction (OpControlBarrier in SPIR-V) with workgroup execution scope.

The physical hardware contract of a barrier is absolute: every active thread in the workgroup must execute the barrier instruction before any thread can proceed past it.

The Divergent Barrier Scenario

Consider what happens when a barrier is placed within a non-uniform branch:

#![allow(unused)]
fn main() {
// CRITICAL FAULT: Divergent Barrier
if space.cell_x < 128 {
    tile_cache[space.cell_x] = input[space.x];
    space.sync(); // Threads 0..127 wait here
} else {
    // Threads 128..255 bypass the barrier entirely!
}
}

In a workgroup of 256 threads:

1. Threads 0..127 enter the if block, reach space.sync(), and halt their execution units waiting for the remaining 128 threads to arrive.
2. Threads 128..255 take the else branch, proceed forward, and never issue the barrier instruction.
3. The Deadlock: Threads 0..127 wait indefinitely for signals that can never be sent. The physical hardware compute unit locks permanently.

3. The Compiler Invariant (error[E0009])

To prevent device hangs, Enki’s JIT compiler analyzes the control flow graph (CFG) of the #[nam] during bitcode canonicalization.

If an invocation of space.sync() is reachable from a non-uniform branch condition (any condition that evaluates differently across cells in the tile), compilation halts immediately:

error[E0009]: divergent tile synchronization detected inside #[nam]
  --> src/main.rs:18:9
   |
15 |     if space.cell_x < 128 {
   |        ------------------ branch condition is non-uniform across tile cells
...
18 |         space.sync();
   |         ^^^^^^^^^^^^ tile synchronization called here
   |
   = note: all cells in a tile must reach `space.sync()` concurrently; barriers inside non-uniform control flow cause permanent GPU hardware deadlocks.
   = help: move `space.sync()` outside the conditional block or ensure the branch condition is uniform across the entire tile.

4. The Trailing Thread Boundary Trap

A common mistake in GPU programming occurs when guarding problem boundaries:

#![allow(unused)]
fn main() {
// WRONG PATTERN: Early exit causes deadlock in boundary tiles
if !space.in_bounds_x() {
    return; // Trailing threads exit the kernel early!
}

tile_cache[space.cell_x] = input[space.x];
space.sync(); // DEADLOCK: Trailing threads never reach this barrier!
}

If the global problem size is not an exact multiple of the tile size (for instance, 1000 elements partitioned into tiles of 256):

• The final tile contains 232 valid elements and 24 inactive trailing threads.
• The 24 trailing threads execute return and terminate.
• The 232 active threads hit space.sync() and deadlock waiting for the terminated threads.

The Correct Uniform Pattern

All threads within the tile must participate in the barrier. Clamp or zero out out-of-bounds memory loads instead of early-exiting:

#![allow(unused)]
fn main() {
// CORRECT: All threads reach the barrier unconditionally
let value = if space.in_bounds_x() {
    input[space.x]
} else {
    0.0 // Trailing threads load neutral dummy data
};

tile_cache[space.cell_x] = value;

// Uniform execution: all 256 cells execute the barrier together
space.sync();

if space.in_bounds_x() {
    output[space.x] = tile_cache[space.cell_x];
}
}

By ensuring that space.sync() is executed uniformly by all cells in the workgroup, the dispatch remains provably deadlock-free.

Internal Compiler Errors & Reporting (error[E9999])

During the active alpha phase (v0.1), the JIT compilation toolchain (parsu) undergoes continuous testing across diverse hardware architectures and complex Rust language constructs.

If the compiler encounters an unexpected invariant violation while lowering LLVM bitcode into GPU SPIR-V, it aborts execution with an Internal Compiler Error (ICE).

1. What is an Internal Compiler Error (ICE)?

An ICE does not indicate a syntax error or a type-safety violation in your code. Rather, it signifies that the JIT compiler backend encountered an edge case it could not resolve internally (such as an unhandled LLVM intrinsic, an unexpected register footprint, or an unsupported optimization graph state).

When this happens, the compiler catches the internal state, packages the compilation failure metadata, and outputs diagnostic error[E9999]:

error[E9999]: internal compiler error (ICE) in Parsu GPU backend (v0.1.0-alpha)
  --> src/main.rs:10:1
   |
10 | #[nam]
   | ^^^^^^ compiler hit an unexpected invariant condition
   |
   = note: this is a known limitation of the early alpha release on some architectures.
   = note: the compiler caught an internal failure and sealed the state into an encrypted token

  -----BEGIN ENKI CRASH FLIGHT RECORDER TOKEN-----
  eJy1V9tu2zYQfb9f8eAF... [SEALED DIAGNOSTIC PAYLOAD] ...bX2F1
  -----END ENKI CRASH FLIGHT RECORDER TOKEN-----

   = help: we would appreciate a bug report! Please open an issue and include your #[nam] code along with the diagnostic token below:
           https://github.com/enkiruntime/enki/issues

2. Reporting Compiler Issues

Because Enki is in active alpha, reporting these edge cases directly contributes to the stability and maturity of the platform.

If you encounter an E9999 diagnostic:

1. Isolate a Minimal Example: Try to isolate the minimal #[nam] function that triggers the condition.
2. Open a GitHub Issue: Visit github.com/enkiruntime/enki/issues.
3. Include the Full Diagnostic: Copy and paste the entire terminal output, including the sealed token block.

The diagnostic token allows me to reconstruct the compiler phase, target nam, and invariant failure without requiring access to your proprietary code.

The GPU BorrowEngine: Two-Tier Safety & Research Scope

Rust’s compile-time borrow checker statically guarantees memory safety on CPU threads by enforcing the aliasing XOR mutability invariant: memory may have multiple shared readers (&T), or exactly one exclusive writer (&mut T), but never both simultaneously.

Enki does not replace or bypass this mechanism. Instead, it is architected around a Two-Tier Cooperative Defense Model:

1. Tier 1 (Host Compile-Time): Leverages rustc’s native borrow checker to intercept aliasing bugs on the host CPU.
2. Tier 2 (Device Runtime): The BorrowEngine evaluates GPU-specific invariants where the host compiler has zero visibility.

1. Tier 1: Leveraging the Native Rust Borrow Checker

Enki’s resource APIs are designed with standard Rust lifetime signatures. For example, obtaining a mutable slice from a GpuVec requires an exclusive borrow of the parent container:

#![allow(unused)]
fn main() {
impl<T> GpuVec<T> {
    pub fn slice_mut<R: RangeBounds<usize>>(&mut self, range: R) -> SliceMut<'_, T>;
    pub fn slice<R: RangeBounds<usize>>(&self, range: R) -> Slice<'_, T>;
}
}

Because slice_mut borrows &mut self:

#![allow(unused)]
fn main() {
let mut vec = gpu_vec![0.0f32; 1000];

let mut slice_a = vec.slice_mut(0..600);
//                                   ^^^ mutable borrow occurs here
let slice_b = vec.slice(600..1000);
//                            ^^^ cannot borrow `vec` as immutable because it is also borrowed as mutable
my_nam.run(&Space::gpu_x(N), &mut slice_a, &slice_b);
//                                       ^^^^^^^^^^^^ mutable borrow later used here
}

rustc intercepts this attempt at compile time inside the developer’s editor. Enki deliberately relies on the host compiler as its first line of defense to eliminate trivial aliasing hazards before compilation even completes.

2. Tier 2: The Three Blind Spots of the Host Compiler

While rustc manages CPU references, it is physically unaware of graphics hardware architecture, execution grids, and deferred command queues.

The BorrowEngine intervenes at the dispatch boundary (.run()) to enforce invariants across three specific blind spots:

Blind Spot A: Spatial Execution Domains (error[E1008])

rustc verifies that a GpuVec is allocated and valid. However, the host compiler cannot inspect the dynamic dimensions of a GPU execution grid (Space::gpu_x(1024)).

If a buffer containing 512 elements is dispatched over 1024 threads in a 1:1 SPMD kernel:

• rustc considers the types valid and allows compilation.
• The GPU would attempt to calculate out-of-bounds Buffer Device Addresses for threads 512 through 1023.
• The BorrowEngine intercepts this at dispatch time and halts execution with error[E1008].

Blind Spot B: Temporal Presentation Latency (error[E1010])

In Vulkan, presentation is deferred: calling flow.present(&buffer) records a transfer command that executes asynchronously at the conclusion of the frame submission pipeline.

To rustc:

• flow.present(&buffer) takes a shared read-only reference &buffer whose lifetime ends immediately after the statement.
• Subsequent calls like clear.run(&space, &mut buffer) inside the same recording block are permitted by rustc.

To the GPU:

• Modifying the buffer after queuing it for display would overwrite the image before the swapchain copy command can read it.
• The BorrowEngine’s FrameBorrowLedger tracks this chronological lifecycle and rejects the subsequent mutation with error[E1010].

Blind Spot C: Unchecked Slice Intersections (error[E1007])

Certain advanced algorithms require passing multiple sub-slices derived from the same parent buffer. To permit this on the host, Enki provides explicit unchecked constructors that bypass host exclusivity:

#![allow(unused)]
fn main() {
// Bypasses host exclusivity by taking &self
let mut slice_a = unsafe { buffer.slice_mut_unchecked(0..600) };
let slice_b = buffer.slice(400..1000);
}

While rustc permits this because of the unsafe block, the BorrowEngine does not trust the developer blindly.

Before command buffers are submitted to the GPU, the engine computes the mathematical interval intersection:

overlap = [0..600) ∩ [400..1000) = [400..600)

If an overlap is detected on a mutable range, the BorrowEngine halts execution with error[E1007].

3. Engineering Status: Pragmatic Safety vs. Formal Proofs

It is essential to understand the scientific boundaries of this architecture:

An Active Systems Research Module

The BorrowEngine is an evolving component of the Enki runtime. Its algorithms and verification tables are actively being refined, expanded, and stress-tested.

Deterministic Invariant Checking, Not Formal Soundness

As outlined before, Enki does not claim to provide a formally verified, mathematically sound type system for arbitrary parallel execution.

Formal verification (such as theorem proving in Coq or formal memory models like CompCert) mathematically proves that all possible program paths are sound. The BorrowEngine, by contrast, is a compiler-grade pragmatic runtime guardrail. It intercepts known, concrete classes of spatial and temporal data race hazards at the dispatch boundary, providing a practical safety net without claiming mathematical infallibility.

Chapter Navigation

• Spatial Bounds & Range Collisions (E1007, E1008): Detailed analysis of interval intersection math and space domain capacity checks.
• Temporal Presentation & Dispatch Modes (E1010): Chronological frame tracking and the operational boundaries between Safe Mode and Unchecked Mode.

Spatial Safety & Range Collisions (E1007, E1008)

The spatial borrow checker (SpatialBorrowChecker) validates that memory containers passed to an execution grid do not violate physical memory boundaries or overlap concurrently during execution.

1. Space Domain Bounds (error[E1008])

In the 1:1 SPMD model, each thread in Space is mapped directly to its corresponding scalar element in a GpuVec (thread i accesses element i).

If the global execution domain requires more threads than the container contains elements, trailing threads would calculate out-of-bounds Buffer Device Addresses, causing a GPU page fault or invalid memory read.

Dynamic Verification

Before submitting the dispatch, the runtime verifies:

container.len() >= space.size_x * space.size_y * space.size_z

If the capacity is insufficient, execution halts with diagnostic error[E1008]:

error[E1008]: GpuVec capacity is smaller than the requested space domain
  --> src/main.rs:51:23
   |
51 |         scale_vectors.run(&Space::gpu_x(N), &mut vec);
   |                       ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ container contains fewer elements than required space cells
   |
   = note: each cell in `space` expects exclusive 1:1 access to its
           corresponding element; extra threads would access out-of-bounds
           memory.
   = note: argument 1 (`GpuVec<f32>`) contains only 1000 elements
   = note: space execution (dimensions: 10000x1x1) requires at least 10000
           elements (deficit of 9000 elements)
   = help: resize the `GpuVec` using `.resize(...)` to cover all space
           cells, or adjust the space domain dimensions.

2. Spatial Slice Disjointness (error[E1007])

When multiple slices derived from the same physical root buffer are passed to a single nam dispatch, the runtime must verify that concurrent writes do not collide.

The Interval Intersection Algorithm

For any pair of slices A and B sharing the same underlying allocation root ID:

• If both are immutable (&[T]), concurrent access is permitted.
• If either slice is mutable (&mut [T]), the engine computes their interval intersection:

overlap_start = max(A_start, B_start)
overlap_end   = min(A_end, B_end)

If overlap_start < overlap_end, an overlapping memory range exists. The runtime halts execution with diagnostic error[E1007]:

error[E1007]: conflicting access to overlapping slices in nam dispatch
  --> src/main.rs:20:24
   |
20 |         process_slices.run_unchecked(&Space::gpu_x(N), &mut slice_a, &slice_b);
   |                        ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ mutable slice overlaps with an existing slice of the same GpuVec
   |
   = note: parallel threads executing in `space` cannot safely write to
           memory that is concurrently being accessed by other threads.
   = note: argument 1 covers elements [400..1000]
   = note: argument 2 covers elements [0..500]
   = note: spatial collision occurs on 100 overlapping elements [400..500]
           of the parent GpuVec<f32>
   = help: ensure sub-slices derived from the same parent GpuVec are
           disjoint using `.split_at_mut()` or non-intersecting element
           ranges.

The Idiomatic Solution: split_at_mut

To guarantee spatial disjointness without runtime interval checking, use standard Rust partitioning methods such as split_at_mut():

#![allow(unused)]
fn main() {
let mut vec = gpu_vec![2.0f32; 1000]; // parent vector
let mut output = gpu_vec![0.0f32; 1000]; // output vector

vec.copy_from_slice_at(500, &[10.0f32; 500]); // filling the parent vector with 10.0f32 after index 500.

let mut slice = vec.as_mut_slice(); // getting a mutable slice of the parent vector.

// `left` covers elements [0..500] and it holds the value of `2.0f32` and `right` covers elements [500..1000] and it holds the value of `10.0f32`
let (mut left, mut right) = slice.split_at_mut(500); // splitting the slice into two at index 500

process_slices.run_unchecked(&Space::gpu_x(N), &mut left, &mut right, &mut output);
}

Temporal Safety & Dispatch Modes (E1010)

In addition to validating spatial boundaries within a single dispatch, the BorrowEngine tracks resource states chronologically across the entire active frame via the FrameBorrowLedger.

1. Temporal Presentation Hazards (error[E1010])

In real-time graphics and display pipelines, presenting a buffer to the screen is an asynchronous operation. When you call:

#![allow(unused)]
fn main() {
flow.present(&pixel_buffer);
}

Enki does not immediately copy the pixels to the swapchain. Instead, it marks the buffer as queued for display in the active command buffer. The physical transfer occurs at the conclusion of the frame submission pipeline.

The Mutation Hazard

If subsequent dispatches within the same flow attempt to mutate the buffer after queuing it for presentation:

Frame Execution Timeline (Recording Phase)
──[nam: processing `pixels`]──► [flow.present(pixels)] ──► [nam: Clear Image (&mut pixels)] ──► [Submit]
                                                            HAZARD: E1010
                                                            The intended presentation output would be corrupted!

Because the GPU executes the recorded commands in order, mutating the buffer after calling flow.present() would overwrite the image before the swapchain copy command can read it.

For example, if you write this pattern of code:

#![allow(unused)]
fn main() {
enki.flow(|flow| {
    process_image.run(&Space::gpu_xy(W, H), &left, &right, &mut output);
    flow.present(&output);
    clear_image.run(&Space::gpu_xy(W, H), &mut output);
});
}

The FrameBorrowLedger records presentation queues and halts execution with diagnostic error[E1010]:

error[E1010]: cannot mutably borrow GpuVec after queuing it for presentation
  --> src/main.rs:31:21
   |
31 |         clear_image.run(&Space::gpu_xy(W, H), &mut output);
   |                     ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ mutable borrow occurs after presentation in the same flow
   |
   = note: screen transfer copies (`.present(...)`) are deferred and executed at the very end of the flow; modifying the
           container afterwards will overwrite the intended screen output.
   = note: argument 1 (parent GpuVec #5) was already queued for screen presentation via `.present()` earlier in this
           flow
   = help: ensure `.present()` is the final operation performed on this container within the active flow.

2. Dispatch Modes: Safe vs. Unchecked

Enki provides two execution modes to accommodate different algorithmic requirements:

Safe Mode (.run())

Safe mode is the default dispatch path. It enforces the following invariants:

• 1:1 SPMD Isolation: Guarantees that thread i writes exclusively to cell i.
• Slice Disjointness: Rejects dispatches with overlapping mutable intervals (E1007).
• Domain Coverage: Rejects dispatches where container capacities do not satisfy the problem space (E1008).
• Presentation Immutability: Rejects mutable borrows on presented buffers (E1010).

Unchecked Mode (.run_unchecked())

Certain algorithms require arbitrary or indirect write operations across an entire mutable slice (&mut [T])—such as software rasterization, particle splatting, or custom parallel radix sorting.

In such cases, the compiler cannot statically or dynamically prove that concurrent writes will not collide. You explicitly opt out of safe mode:

#![allow(unused)]
fn main() {
unsafe {
    scatter_kernel.run_unchecked(&space, &mut global_slice);
}
}

What run_unchecked Disables (and What it Retains)

• Disabled: It bypasses restrictions on passing unrestricted mutable slices across parallel threads.
• Retained: It does not disable all runtime validation. Spatial domain checks, interval collision detection, and temporal presentation tracking remain actively enforced by the BorrowEngine.

The unsafe block explicitly signifies that spatial race-freedom inside the slice is guaranteed by the developer’s indexing algorithm.

Graphics & The Barrier Engine

While Enki is a general-purpose compute platform, it natively integrates compute execution with real-time display presentation and automated hardware synchronization.

This section covers two high-level systems:

• Real-Time Presentation (flow.present): Integrating compute pipelines directly with display windows (GLFW) and managing dynamic swapchain resizing.
• The Automatic Barrier Solver (TTRD): How Enki leverages Rust reference mutability to completely eliminate manual Vulkan pipeline barriers and memory hazards behind the scenes.

Real-Time Presentation (flow.present)

Traditional graphics programming enforces a strict separation between compute pipelines and rendering pipelines. Drawing pixels on screen typically requires setting up render passes, vertex buffers, rasterizers, and fragment shaders.

Enki unifies this through direct buffer presentation: you write pixels into a standard contiguous VRAM buffer (GpuVec<u32>) using a compute nam, and queue it directly to the display window’s swapchain in a single statement.

1. Adding Dependencies

To build a real-time windowed application, add enki-gpu and glfw to your project:

cargo add enki-gpu glfw

Your Cargo.toml dependencies should include:

[dependencies]
enki-gpu = "0.1"
glfw = "0.59"

2. Window Setup & The NoApi Invariant

Vulkan manages swapchain images directly through native platform display surfaces. Therefore, you must explicitly instruct GLFW not to create an OpenGL context:

#![allow(unused)]
fn main() {
let mut glfw = glfw::init(glfw::fail_on_errors).expect("Failed to initialize GLFW");

// MANDATORY: Disable OpenGL context creation
glfw.window_hint(glfw::WindowHint::ClientApi(glfw::ClientApiHint::NoApi));
glfw.window_hint(glfw::WindowHint::Resizable(true));

let (mut window, events) = glfw
    .create_window(WIDTH, HEIGHT, "Enki Display", glfw::WindowMode::Windowed)
    .expect("Failed to create GLFW window");

// Enable polling for resize and keyboard events
window.set_framebuffer_size_polling(true);
window.set_key_polling(true);
}

Passing the Window to Enki

Enki’s windowed initializer expects an atomic reference counted (Arc<W>) window instance implementing raw_window_handle::HasWindowHandle and HasDisplayHandle.

Because GLFW’s Window implements these traits natively, pass it directly:

#![allow(unused)]
fn main() {
let window_handle = Arc::new(window);

// Initialize Enki in windowed mode
let enki = Enki::init_windowed(window_handle.clone(), WIDTH, HEIGHT);
}

3. The Complete, Runnable Application

Below is a complete, self-contained src/main.rs that renders an animated real-time procedural color plasma at 60 FPS and handles dynamic window resizing cleanly without mutable borrowing conflicts:

use enki::*;
use glfw::{Action, Context, Key, WindowEvent, WindowHint, WindowMode};
use std::sync::Arc;
use std::time::Instant;

const WIDTH: u32 = 1280;
const HEIGHT: u32 = 720;

// 1. Declare the compute nam that generates pixel colors
#[nam]
fn render_plasma(space: &Space, pixel: &mut u32, time: f32) {
    if !space.in_bounds_xy() {
        return;
    }

    let (u, v) = space.uv();

    // Compute animated color waves
    let r = ((u * 10.0 + time).sin() * 0.5 + 0.5);
    let g = ((v * 10.0 - time * 1.5).cos() * 0.5 + 0.5);
    let b = (((u + v) * 8.0 + time * 2.0).sin() * 0.5 + 0.5);

    // Pack RGB floats into 32-bit presentation integer (0xAARRGGBB)
    *pixel = space.set_rgb_color(r, g, b);
}

fn main() {
    // 2. Initialize GLFW with NoApi
    let mut glfw = glfw::init(glfw::fail_on_errors).expect("Failed to initialize GLFW");
    glfw.window_hint(WindowHint::ClientApi(glfw::ClientApiHint::NoApi));
    glfw.window_hint(WindowHint::Resizable(true));

    let mut current_width = WIDTH;
    let mut current_height = HEIGHT;

    let (mut window, events) = glfw
        .create_window(current_width, current_height, "Enki Plasma", WindowMode::Windowed)
        .expect("Failed to create window");

    window.set_framebuffer_size_polling(true);
    window.set_key_polling(true);

    let window_handle = Arc::new(window);

    // 3. Initialize Enki windowed runtime
    let enki = Enki::init_windowed(window_handle.clone(), current_width, current_height);

    // 4. Allocate the physical VRAM framebuffer
    let mut screen = gpu_vec![0u32; (current_width * current_height) as usize];

    let start_time = Instant::now();
    let mut is_running = true;

    // 5. Main presentation loop
    while is_running && !window_handle.should_close() {
        glfw.poll_events();

        for (_, event) in glfw::flush_messages(&events) {
            match event {
                WindowEvent::Close => is_running = false,
                WindowEvent::Key(Key::Escape, _, Action::Press, _) => is_running = false,
                _ => {}
            }
        }

        // Handle dynamic swapchain recreation if window dimensions changed
        let (fb_w, fb_h) = window_handle.get_framebuffer_size();
        if fb_w > 0 && fb_h > 0 && (fb_w as u32 != current_width || fb_h as u32 != current_height) {
            current_width = fb_w as u32;
            current_height = fb_h as u32;

            // Recreate Vulkan swapchain images
            let _ = enki.resize(current_width, current_height);

            // Reallocate VRAM screen buffer to match new extent
            screen = GpuVec::zeroed((current_width * current_height) as usize);
        }

        let time = start_time.elapsed().as_secs_f32();

        // 6. Record and submit the compute-to-display flow
        enki.flow(|flow| {
            render_plasma.run(
                &Space::gpu_xy(current_width as usize, current_height as usize),
                &mut screen,
                GpuParam::new(time),
            );

            // Queue buffer for direct presentation
            flow.present(&screen);
        });
    }
}

Run the application:

cargo run

A window will appear rendering an animated procedural plasma directly from the GPU compute cores at your monitor’s native refresh rate.

4. The Presentation Lifecycle

When flow.present(&screen) is recorded:

Host CPU Flow Recording
┌──────────────────────────────────────┐
│ render_plasma.run(&space, &mut screen) │ ──► Compute writes pixels into VRAM
├──────────────────────────────────────┤
│ flow.present(&screen)                 │ ──► Schedules hardware copy: Buffer ──► Swapchain Image
└──────────────────────────────────────┘
                   │
                   ▼ (Queue Submission)
Vulkan Hardware Timeline
──[Compute Shader Finishes]──► [vkCmdCopyBufferToImage] ──► [vkQueuePresentKHR] ──► Display

1. Copy Command: The engine injects an asynchronous transfer command (vkCmdCopyBufferToImage) from the GpuVec into the current swapchain image.
2. Synchronization: The engine automatically coordinates the timeline semaphore with the swapchain’s image-acquired and render-finished binary semaphores.
3. Atomic Display: The frame is presented via vkQueuePresentKHR once memory transfers complete.

5. Dimension Verification (error[E1008])

Swapchain presentation requires an exact 1:1 pixel mapping. Attempting to present a buffer that contains fewer elements than the active window extent (width * height) halts execution immediately:

error[E1008]: GpuVec capacity is smaller than the requested space domain
  --> src/main.rs:88:27
   |
88 |             render_plasma.run(
   |                           ^^^^ container contains fewer elements than required space cells
   |
   = note: each cell in `space` expects exclusive 1:1 access to its
           corresponding element; extra threads would access
           out-of-bounds memory.
   = note: argument 1 (`GpuVec<u32>`) contains only 1439999 elements
   = note: space execution (dimensions: 1600x900x1) requires at least
           1440000 elements (deficit of 1 elements)
   = help: resize the `GpuVec` using `.resize(...)` to cover all space
           cells, or adjust the space domain dimensions.

The Automatic Barrier Solver (TTRD)

In raw Vulkan programming, managing synchronization is notoriously difficult. Developers must manually insert execution and memory barriers (vkCmdPipelineBarrier2), explicitly specifying pipeline stages and access masks to prevent data hazards.

Manual synchronization presents two major risks:

1. Under-synchronization: Missing a barrier causes silent data corruption, screen tearing, and non-deterministic race conditions.
2. Over-synchronization: Inserting redundant or overly broad barriers serializes the GPU execution units, creating pipeline bubbles that destroy performance.

Enki eliminates manual synchronization entirely. Through its Transitive Reduction Dependency Solver (TTRD), the runtime derives the mathematically optimal set of pipeline barriers automatically.

Note: If you are not interested in the underlying runtime implementation, you can safely ignore this section. This is an optional architectural overview.

1. Leveraging Rust Mutability Semantics

How does the runtime know what memory barriers are required without developer annotations?

Enki leverages the semantic guarantees already encoded in Rust’s type system:

• When a #[nam] declares &T or &[T], the argument is tagged internally as AccessIntent::Read.
• When a #[nam] declares &mut T or &mut [T], the argument is tagged as AccessIntent::Write (or ReadWrite).

Because each physical buffer allocation in VRAM possesses a unique internal engine slot ID, the runtime tracks the chronological memory access history for every resource across all tasks queued within an enki.flow.

2. Classical Data Hazards

As tasks are pushed to the active queue, the engine builds a directed dependency graph by detecting classical hazard conditions on each memory slot:

• RAW (Read-After-Write): Task A writes to a buffer; Task B subsequently reads from it. Task B must wait for Task A’s writes to be made visible.
• WAR (Write-After-Read): Task A reads from a buffer; Task B subsequently writes to it. Task B must not overwrite memory before Task A completes reading.
• WAW (Write-After-Write): Task A writes to a buffer; Task B subsequently overwrites it. Write orders must be strictly preserved.

3. The Transitive Reduction Algorithm

In a multi-pass compute workflow, naive hazard tracking creates numerous redundant dependencies.

For example, consider three sequential tasks operating on the same buffer:

• Task A writes data (Write).
• Task B reads and updates data (Read / Write).
• Task C reads the final data (Read).

A naive dependency system inserts three barriers: A B, B C, and A C.

However, because Task B already depends on Task A, and Task C depends on Task B, the direct dependency from A C is mathematically redundant. Inserting a barrier for A C forces the GPU to stall unnecessarily.

Graph Reduction via Reachability

Enki applies graph-theoretic Transitive Reduction using depth-first search (DFS) reachability analysis:

If a directed path already exists between task u and task v through an intermediate task w (u to w to v), the direct edge u to v is pruned from the dependency graph.

After reduction, the remaining edges represent the minimum necessary and sufficient set of hardware barriers.