Introduction

Trueno (Spanish: "thunder") is a high-performance Rust library providing unified compute primitives across three execution targets: CPU SIMD, GPU, and WebAssembly. The name reflects the library's mission: to deliver thunderous performance through intelligent hardware acceleration.

The Problem: Performance vs Portability

Modern applications face a critical tradeoff:

• Hand-optimized assembly: Maximum performance (2-50x speedup), but unmaintainable and platform-specific
• Portable high-level code: Easy to write and maintain, but leaves performance on the table
• Unsafe SIMD intrinsics: Good performance, but riddled with unsafe code and platform-specific complexity

Traditional approaches force you to choose between performance, safety, and portability. Trueno chooses all three.

The Solution: Write Once, Optimize Everywhere

Trueno's core philosophy is write once, optimize everywhere:

use trueno::Vector;

// Single API call, multiple backend implementations
let a = Vector::from_slice(&[1.0, 2.0, 3.0, 4.0]);
let b = Vector::from_slice(&[5.0, 6.0, 7.0, 8.0]);
let result = a.add(&b)?;

// Automatically selects best backend:
// - AVX2 on modern Intel/AMD (4-8x speedup)
// - NEON on ARM64 (2-4x speedup)
// - GPU for large workloads (10-50x speedup)
// - WASM SIMD128 in browsers (2x speedup)

Key Features

1. Multi-Target Execution

Trueno runs on three execution targets with a unified API:

2. Runtime Backend Selection

Trueno automatically selects the best available backend at runtime:

┌─────────────────────────────────────────────────┐
│           Trueno Public API (Safe)              │
│  compute(), map(), reduce(), transform()        │
└─────────────────────────────────────────────────┘
                      │
        ┌─────────────┼─────────────┐
        ▼             ▼             ▼
   ┌────────┐   ┌─────────┐   ┌──────────┐
   │  SIMD  │   │   GPU   │   │   WASM   │
   │ Backend│   │ Backend │   │  Backend │
   └────────┘   └─────────┘   └──────────┘
        │             │             │
   ┌────┴────┐   ┌────┴────┐   ┌───┴─────┐
   │ Runtime │   │CUDA/wgpu│   │ SIMD128 │
   │ Detect  │   │ Compute │   │ Portable│
   └─────────┘   └─────────┘   └─────────┘
   │  │  │  │       │
   SSE2 AVX NEON   PTX
      AVX512     (trueno-gpu)

Backend Selection Priority:

1. GPU (if available + workload > 100K elements)
2. AVX-512 (if CPU supports)
3. AVX2 (if CPU supports)
4. AVX (if CPU supports)
5. SSE2 (baseline x86_64)
6. NEON (ARM64)
7. SIMD128 (WASM)
8. Scalar fallback

3. Zero Unsafe in Public API

All unsafe code is isolated to backend implementations:

// ✅ SAFE public API
pub fn add(&self, other: &Self) -> Result<Self> {
    // Safe bounds checking, validation
    if self.len() != other.len() {
        return Err(TruenoError::SizeMismatch { ... });
    }

    // ❌ UNSAFE internal implementation (isolated)
    #[cfg(target_arch = "x86_64")]
    if is_x86_feature_detected!("avx2") {
        unsafe { self.add_avx2(other) }
    } else {
        self.add_scalar(other) // Safe fallback
    }
}

Safety guarantees:

• Public API is 100% safe Rust
• All bounds checked before dispatching to backends
• Miri validation for undefined behavior
• 286 documented SAFETY invariants in backend code

4. Proven Performance

Trueno delivers 2-50x speedups over scalar code:

All benchmarks validated with:

• Coefficient of variation < 5%
• 100+ iterations for statistical significance
• No regressions > 5% vs baseline

5. Extreme TDD Quality

Trueno is built with EXTREME TDD methodology:

• >90% test coverage (verified with cargo llvm-cov)
• Property-based testing (commutativity, associativity, distributivity)
• Backend equivalence tests (scalar vs SIMD vs GPU produce identical results)
• Mutation testing (>80% mutation kill rate with cargo mutants)
• Zero tolerance for defects (all quality gates must pass)

Real-World Impact: The FFmpeg Case Study

FFmpeg (the world's most-used video codec library) contains:

• 390 assembly files (~180,000 lines, 11% of codebase)
• Platform-specific implementations for x86, ARM, MIPS, PowerPC
• Speedups: SSE2 (2-4x), AVX2 (4-8x), AVX-512 (8-16x)

Problems with hand-written assembly:

• ❌ Unsafe (raw pointers, no bounds checking)
• ❌ Unmaintainable (390 files, must update all platforms)
• ❌ Non-portable (separate implementations per CPU)
• ❌ Expertise barrier (requires assembly knowledge)

Trueno's value proposition:

• ✅ Safety: Zero unsafe in public API
• ✅ Portability: Single source → x86/ARM/WASM/GPU
• ✅ Performance: 85-95% of hand-tuned assembly
• ✅ Maintainability: Rust type system catches errors at compile time

Who Should Use Trueno?

Trueno is designed for:

1. ML/AI Engineers - NumPy-like compute primitives for Rust (use with aprender for training)
2. Systems Programmers - Eliminate unsafe SIMD intrinsics
3. Game Developers - Fast vector math for physics/graphics
4. Scientific Computing - High-performance numerical operations
5. WebAssembly Developers - Portable SIMD for browsers/edge
6. Transpiler Authors - Safe SIMD target for Depyler/Decy/Ruchy

Design Principles

Trueno follows five core principles:

1. Write once, optimize everywhere - Single algorithm, multiple backends
2. Safety via type system - Zero unsafe in public API
3. Performance must be proven - Every optimization validated with benchmarks (≥10% speedup)
4. Extreme TDD - >90% coverage, mutation testing, property-based tests
5. Toyota Way - Kaizen (continuous improvement), Jidoka (built-in quality)

What's Next?

• Getting Started - Install Trueno and run your first program
• Architecture - Understand the multi-backend design
• API Reference - Explore available operations
• Performance - See benchmark results and optimization techniques
• Examples - Learn from real-world use cases

Project Status

Trueno is under active development at Pragmatic AI Labs:

• Current Version: 0.1.0 (Phase 1: Vector operations)
• License: MIT/Apache-2.0 dual-licensed
• Repository: github.com/paiml/trueno
• Issues: github.com/paiml/trueno/issues

Scope:

• Trueno: Compute primitives (vectors, matrices, SIMD, GPU) - NumPy equivalent
• Aprender: ML framework with autograd and training - PyTorch equivalent

Trueno is the compute backend for higher-level ML libraries. For neural networks and training, see aprender.

Join us in building the future of safe, high-performance compute!

Installation

This guide covers installing Trueno and its dependencies.

Prerequisites

Rust Toolchain

Trueno requires Rust 1.70 or later. Install via rustup:

curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh
rustup update stable

Verify installation:

rustc --version  # Should be >= 1.70.0
cargo --version

Platform-Specific Requirements

Linux

# Ubuntu/Debian
sudo apt-get update
sudo apt-get install build-essential pkg-config

# Fedora/RHEL
sudo dnf install gcc pkg-config

macOS

# Install Xcode Command Line Tools
xcode-select --install

Windows

Install Visual Studio 2022 with:

• Desktop development with C++
• Windows 10/11 SDK

Optional: GPU Support

For GPU acceleration, install graphics drivers:

NVIDIA (CUDA/Vulkan):

# Ubuntu/Debian
sudo apt-get install nvidia-driver-535 vulkan-tools

# Verify
vulkaninfo

AMD (Vulkan):

# Ubuntu/Debian
sudo apt-get install mesa-vulkan-drivers vulkan-tools

# Verify
vulkaninfo

Intel (Vulkan):

# Ubuntu/Debian
sudo apt-get install intel-media-va-driver vulkan-tools

macOS (Metal): Metal support is built-in on macOS 10.13+. No additional installation required.

Installing Trueno

From crates.io (Recommended)

Add Trueno to your Cargo.toml:

[dependencies]
trueno = "0.1"

Or use cargo add:

cargo add trueno

From GitHub (Development)

For the latest development version:

[dependencies]
trueno = { git = "https://github.com/paiml/trueno", branch = "main" }

With Specific Features

Trueno supports feature flags for selective compilation:

[dependencies]
# Default: SIMD backends only (no GPU)
trueno = "0.1"

# Enable GPU support
trueno = { version = "0.1", features = ["gpu"] }

# Enable all features
trueno = { version = "0.1", features = ["gpu", "wasm"] }

# Minimal (scalar only, for testing)
trueno = { version = "0.1", default-features = false }

Available features:

• gpu - Enable GPU backend via wgpu (adds ~5MB to binary)
• wasm - Enable WebAssembly SIMD128 support
• f16 - Enable half-precision (f16) support (requires nightly)

Verifying Installation

Create a test project:

cargo new trueno-test
cd trueno-test

Add Trueno to Cargo.toml:

[dependencies]
trueno = "0.1"

Replace src/main.rs with:

use trueno::Vector;

fn main() -> Result<(), Box<dyn std::error::Error>> {
    // Create two vectors
    let a = Vector::from_slice(&[1.0, 2.0, 3.0, 4.0]);
    let b = Vector::from_slice(&[5.0, 6.0, 7.0, 8.0]);

    // Add them (uses best available SIMD backend)
    let result = a.add(&b)?;

    println!("Result: {:?}", result.as_slice());
    // Output: [6.0, 8.0, 10.0, 12.0]

    // Check which backend was used
    println!("Backend: {:?}", a.backend());

    Ok(())
}

Run the test:

cargo run --release

Expected output:

Result: [6.0, 8.0, 10.0, 12.0]
Backend: Avx2  # (or Sse2, Neon, etc. depending on your CPU)

Development Installation

For contributing to Trueno or running tests:

# Clone repository
git clone https://github.com/paiml/trueno.git
cd trueno

# Build with all features
cargo build --all-features --release

# Run tests
cargo test --all-features

# Run benchmarks
cargo bench

# Generate coverage report
cargo llvm-cov --all-features --workspace

Development Dependencies

Install additional tools for development:

# Code coverage
cargo install cargo-llvm-cov

# Mutation testing
cargo install cargo-mutants

# Benchmarking (included in Cargo.toml dev-dependencies)
# criterion is automatically available

# Formatting and linting (included with rustup)
rustup component add rustfmt clippy

Platform-Specific Notes

x86_64 (Intel/AMD)

Trueno automatically detects and uses the best available SIMD instruction set:

• SSE2: Baseline (guaranteed on all x86_64)
• AVX: Sandy Bridge+ (2011+)
• AVX2: Haswell+ (2013+)
• AVX-512: Zen4, Sapphire Rapids+ (2022+)

Check your CPU features:

# Linux
cat /proc/cpuinfo | grep flags

# macOS
sysctl -a | grep cpu.features

# Windows (PowerShell)
Get-WmiObject -Class Win32_Processor | Select-Object -Property Name, Features

ARM64 (Apple Silicon, AWS Graviton)

Trueno uses NEON SIMD on ARM64:

• Apple M1/M2/M3: Full NEON support (128-bit)
• AWS Graviton2/3: Full NEON support
• Raspberry Pi 4: Limited NEON support

WebAssembly

For WASM targets:

# Install wasm32 target
rustup target add wasm32-unknown-unknown

# Build for WASM
cargo build --target wasm32-unknown-unknown --release

# Enable SIMD128 (requires nightly for now)
rustup toolchain install nightly
cargo +nightly build --target wasm32-unknown-unknown \
    -Z build-std=std,panic_abort \
    --release

Troubleshooting

"No suitable backend found" error

If you see this error, Trueno couldn't detect any SIMD support. Possible causes:

1. Running on ancient CPU (pre-2011 x86_64):

   • Solution: Use Backend::Scalar explicitly

2. Cross-compiling without proper target configuration:

   • Solution: Set RUSTFLAGS for target CPU:

     RUSTFLAGS="-C target-cpu=native" cargo build --release
     

3. WASM without SIMD128:

   • Solution: Enable SIMD in browser flags or use scalar fallback

GPU not detected

If GPU is available but not being used:

1. Check Vulkan/Metal installation:

   # Linux/Windows
   vulkaninfo
   
   # macOS - Metal is built-in, check system version
   sw_vers  # Should be >= 10.13
   

2. Verify GPU feature flag:

   trueno = { version = "0.1", features = ["gpu"] }
   

3. Check workload size (GPU only used for 100K+ elements):

   let large = Vector::from_slice(&vec![1.0; 200_000]);
   println!("Backend: {:?}", large.backend());
   // Should show: Gpu

Compilation errors

Error: feature 'avx512' requires nightly

• Trueno uses stable Rust. This error indicates you're on an old rustc version.
• Solution: rustup update stable

Error: wgpu fails to compile

• This is usually a missing system dependency.
• Solution (Ubuntu): sudo apt-get install libvulkan-dev

Error: Link errors on Windows

• Solution: Install Visual Studio 2022 with C++ build tools

Next Steps

Now that Trueno is installed:

• Quick Start - Run your first program
• Core Concepts - Understand key abstractions
• API Reference - Explore available operations

Quick Start

Get up and running with Trueno in 5 minutes.

Your First Trueno Program

Let's build a simple vector addition program that automatically uses the best available SIMD backend.

Create a New Project

cargo new trueno-quickstart
cd trueno-quickstart

Add Trueno Dependency

Edit Cargo.toml:

[dependencies]
trueno = "0.1"

Write the Code

Replace src/main.rs:

use trueno::Vector;

fn main() -> Result<(), Box<dyn std::error::Error>> {
    // Create vectors from slices
    let a = Vector::from_slice(&[1.0, 2.0, 3.0, 4.0, 5.0]);
    let b = Vector::from_slice(&[10.0, 20.0, 30.0, 40.0, 50.0]);

    // Element-wise addition (automatically uses AVX2/SSE2/NEON)
    let sum = a.add(&b)?;
    println!("a + b = {:?}", sum.as_slice());
    // Output: [11.0, 22.0, 33.0, 44.0, 55.0]

    // Element-wise multiplication
    let product = a.mul(&b)?;
    println!("a * b = {:?}", product.as_slice());
    // Output: [10.0, 40.0, 90.0, 160.0, 250.0]

    // Dot product (reduction operation)
    let dot = a.dot(&b)?;
    println!("a · b = {}", dot);
    // Output: 550.0

    // Check which backend was selected
    println!("Using backend: {:?}", a.backend());

    Ok(())
}

Run It

cargo run --release

Expected output:

a + b = [11.0, 22.0, 33.0, 44.0, 55.0]
a * b = [10.0, 40.0, 90.0, 160.0, 250.0]
a · b = 550.0
Using backend: Avx2  # (varies by CPU)

Understanding What Just Happened

Let's break down the magic:

1. Automatic Backend Selection

let a = Vector::from_slice(&[1.0, 2.0, 3.0, 4.0, 5.0]);

When you create a Vector, Trueno:

1. Detects your CPU features (AVX2, SSE2, NEON, etc.)
2. Selects the best available backend
3. Stores this choice with the vector (no repeated detection)

Backend priority:

• ✅ AVX2 (4-8x faster) if available
• ✅ SSE2 (2-4x faster) as x86_64 baseline
• ✅ NEON (2-4x faster) on ARM64
• ✅ Scalar fallback (always works)

2. Safe, High-Level API

let sum = a.add(&b)?;  // Returns Result<Vector>

Trueno's API is:

• 100% safe Rust - No unsafe in user code
• Bounds-checked - Size mismatches caught at runtime
• Ergonomic - Uses ? operator for error handling

3. Zero-Copy Performance

println!("{:?}", sum.as_slice());

as_slice() returns a reference to internal data - no allocation or copying.

Common Operations

Element-Wise Operations

use trueno::Vector;

let a = Vector::from_slice(&[1.0, 2.0, 3.0, 4.0]);
let b = Vector::from_slice(&[5.0, 6.0, 7.0, 8.0]);

// Arithmetic
let sum = a.add(&b)?;      // [6.0, 8.0, 10.0, 12.0]
let diff = a.sub(&b)?;     // [-4.0, -4.0, -4.0, -4.0]
let prod = a.mul(&b)?;     // [5.0, 12.0, 21.0, 32.0]
let quot = a.div(&b)?;     // [0.2, 0.33, 0.43, 0.5]

// Scalar operations
let scaled = a.mul_scalar(2.0)?;  // [2.0, 4.0, 6.0, 8.0]
let offset = a.add_scalar(10.0)?; // [11.0, 12.0, 13.0, 14.0]

Reduction Operations

let v = Vector::from_slice(&[1.0, 2.0, 3.0, 4.0]);

let sum = v.sum();      // 10.0
let mean = v.mean();    // 2.5
let min = v.min();      // 1.0
let max = v.max();      // 4.0

Transformation Operations

let v = Vector::from_slice(&[1.0, 2.0, 3.0, 4.0]);

// Map function over elements
let squared = v.map(|x| x * x)?;  // [1.0, 4.0, 9.0, 16.0]

// Filter elements
let filtered = v.filter(|x| x > 2.0)?;  // [3.0, 4.0]

// Apply activation functions (coming in Phase 3)
// let activated = v.relu()?;
// let normalized = v.softmax()?;

Error Handling

Trueno uses Rust's Result type for robust error handling:

use trueno::{Vector, TruenoError};

fn safe_divide() -> Result<Vector, TruenoError> {
    let a = Vector::from_slice(&[10.0, 20.0, 30.0]);
    let b = Vector::from_slice(&[2.0, 4.0]);  // Wrong size!

    // This returns Err(TruenoError::SizeMismatch)
    a.div(&b)
}

fn main() {
    match safe_divide() {
        Ok(result) => println!("Result: {:?}", result),
        Err(TruenoError::SizeMismatch { expected, actual }) => {
            eprintln!("Size mismatch: expected {}, got {}", expected, actual);
        }
        Err(e) => eprintln!("Error: {}", e),
    }
}

Performance Tips

1. Use Release Mode

Always benchmark in release mode:

# ❌ Debug mode (10-100x slower!)
cargo run

# ✅ Release mode (full optimizations)
cargo run --release

2. Large Workloads for GPU

GPU backend only activates for large vectors (100K+ elements):

// ❌ Too small for GPU (uses SIMD)
let small = Vector::from_slice(&vec![1.0; 1000]);

// ✅ Large enough for GPU
let large = Vector::from_slice(&vec![1.0; 200_000]);

3. Batch Operations

Chain operations to minimize allocations:

// ❌ Multiple allocations
let temp1 = a.add(&b)?;
let temp2 = temp1.mul(&c)?;
let result = temp2.sub(&d)?;

// ✅ Better: use `map` for complex expressions
let result = a.zip(&b, &c, |a_i, b_i, c_i| {
    (a_i + b_i) * c_i - d_i
})?;

4. Reuse Buffers

For hot loops, reuse output buffers:

let mut output = Vector::zeros(1000);

for i in 0..iterations {
    // Writes into existing buffer (no allocation)
    a.add_into(&b, &mut output)?;
}

What's Next?

Now that you've run your first Trueno program:

• Core Concepts - Understand backends, safety, and performance
• First Program - Build a more complex example
• API Reference - Explore all available operations
• Examples - Real-world use cases

First Program

Let's build a complete image processing program using Trueno to demonstrate real-world usage.

Project: Brightness Adjustment Tool

We'll create a CLI tool that adjusts image brightness using SIMD-accelerated vector operations.

[Content to be added: Complete example with image loading, vector processing, benchmarking]

Next Steps

• Architecture Overview
• API Reference
• Examples

Core Concepts

Understanding Trueno's fundamental concepts will help you write efficient, safe code.

The Vector Type

Vector<T> is Trueno's core abstraction:

use trueno::Vector;

let v = Vector::from_slice(&[1.0, 2.0, 3.0, 4.0]);

Key properties:

• Generic over numeric types: f32, f64, i32, i64
• Immutable by default (functional style)
• Backend selected at creation time (no repeated detection)
• Zero-copy views with as_slice()

Backend Selection

Trueno automatically selects the best backend when you create a Vector:

// Automatic backend selection
let v = Vector::from_slice(&[1.0; 1000]);
println!("{:?}", v.backend());  // Avx2, Sse2, Neon, etc.

// Manual backend override (for testing/profiling)
let v = Vector::with_backend(&[1.0; 1000], Backend::Scalar);

Selection priority:

1. GPU (if workload >100K elements and GPU available)
2. AVX-512 (if CPU supports)
3. AVX2 (if CPU supports)
4. AVX (if CPU supports)
5. SSE2 (x86_64 baseline)
6. NEON (ARM64)
7. Scalar fallback

Safety Model

Trueno maintains safety through three layers:

Layer 1: Type System

// Compile-time type safety
let a = Vector::from_slice(&[1.0f32, 2.0, 3.0]);
let b = Vector::from_slice(&[4.0f64, 5.0, 6.0]);

// ❌ Compile error: type mismatch
// let result = a.add(&b);

Layer 2: Runtime Validation

// Runtime size checking
let a = Vector::from_slice(&[1.0, 2.0, 3.0]);
let b = Vector::from_slice(&[4.0, 5.0]);

// Returns Err(SizeMismatch)
let result = a.add(&b);

Layer 3: Unsafe Isolation

All unsafe code is isolated to backend implementations:

// ✅ 100% safe public API
pub fn add(&self, other: &Self) -> Result<Self> {
    validate_sizes(self, other)?;  // Safe
    
    match self.backend {
        Backend::Avx2 => unsafe { self.add_avx2(other) },  // ❌ Unsafe (internal only)
        Backend::Scalar => self.add_scalar(other),  // ✅ Safe
    }
}

Error Handling

Trueno uses Rust's Result type for robust error handling:

use trueno::{Vector, TruenoError};

fn process_vectors() -> Result<Vector, TruenoError> {
    let a = Vector::from_slice(&[1.0, 2.0, 3.0]);
    let b = Vector::from_slice(&[4.0, 5.0, 6.0]);
    
    let sum = a.add(&b)?;  // Propagate errors with ?
    let product = sum.mul_scalar(2.0)?;
    
    Ok(product)
}

Error types:

• SizeMismatch - Vectors have incompatible sizes
• BackendError - Backend initialization failed
• GpuError - GPU operation failed
• InvalidInput - Invalid parameters (NaN, infinity)

Performance Model

Understanding Trueno's performance characteristics helps you write efficient code.

Operation Complexity

Operations fall into three categories:

Low complexity (add, sub, mul, div):

• Prefer SIMD for >1K elements
• Memory-bandwidth limited
• Expect 1.1-2x speedup

Medium complexity (dot, sum, max):

• SIMD shines here (3-5x speedup)
• Compute-bound, not memory-bound
• Use SIMD even for 100 elements

High complexity (tanh, exp, log):

• Excellent SIMD performance (6-9x speedup)
• Compute-intensive operations
• Consider GPU for >100K elements

Backend Overhead

Each backend has different overhead characteristics:

Next Steps

• First Program - Build a complete example
• Architecture Overview - Deep dive into backends
• API Reference - Explore all operations

Overview

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Backend Selection

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Multi Backend Design

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Simd Backends

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Sse2 Backend

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Avx Backend

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Avx512 Backend

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Neon Backend

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Wasm Backend

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

GPU Backend

Trueno provides two GPU acceleration options:

1. wgpu (Cross-platform) - Vulkan, Metal, DX12, WebGPU via wgpu
2. CUDA (NVIDIA) - Native PTX code generation via trueno-gpu

CUDA Support (trueno-gpu)

For NVIDIA GPUs, trueno-gpu provides pure Rust PTX code generation without requiring LLVM, nvcc, or external toolchains.

Quick Start with CUDA

use trueno_gpu::ptx::{PtxModule, PtxKernel, PtxType};
use trueno_gpu::kernels::{GemmKernel, Kernel};

// Generate optimized GEMM kernel
let kernel = GemmKernel::tiled(1024, 1024, 1024, 32);
let ptx = kernel.emit_ptx();

// PTX can be loaded via CUDA driver API
println!("{}", ptx);

Running CUDA Examples

# PTX code generation (no GPU required)
cargo run -p trueno-gpu --example ptx_quickstart
cargo run -p trueno-gpu --example gemm_kernel

# CUDA runtime examples (requires NVIDIA GPU)
cargo run -p trueno-gpu --example cuda_monitor
cargo run -p trueno-gpu --example flash_attention_cuda

Pre-built CUDA Kernels

See PTX Code Generation for detailed documentation.

wgpu Support (Cross-Platform)

For cross-platform GPU compute, Trueno uses wgpu, supporting Vulkan, Metal, DX12, and WebGPU.

Overview

The wgpu backend enables massive parallelism for compute-heavy operations like matrix multiplication. It supports both native platforms (Linux, macOS, Windows) and WebAssembly (via WebGPU in browsers).

Key Features

• Cross-platform: Single codebase for native and WASM
• Async-first: All operations have async variants for non-blocking execution
• Sync wrappers: Native platforms get convenient sync APIs
• Automatic fallback: Falls back to SIMD when GPU unavailable

Platform Support

Note: WASM cannot use sync APIs because JavaScript's single-threaded model prohibits blocking the main thread.

Feature Flags

[dependencies]
trueno = { version = "0.7.3", features = ["gpu"] }      # Native GPU
trueno = { version = "0.7.3", features = ["gpu-wasm"] } # WASM GPU (WebGPU)

Feature Differences

API Design

Sync API (Native Only)

use trueno::backends::gpu::GpuDevice;

// Initialize device
let device = GpuDevice::new()?;

// Check availability
if GpuDevice::is_available() {
    // Execute operations
    device.matmul(&a, &b, &mut result, m, k, n)?;
    device.relu(&input, &mut output)?;
    let dot = device.dot(&a, &b)?;
}

Async API (All Platforms)

use trueno::backends::gpu::GpuDevice;

// Initialize device
let device = GpuDevice::new_async().await?;

// Check availability
if GpuDevice::is_available_async().await {
    // Execute operations
    device.matmul_async(&a, &b, &mut result, m, k, n).await?;
    device.relu_async(&input, &mut output).await?;
    let dot = device.dot_async(&a, &b).await?;
}

Runtime Detection

use trueno::backends::gpu::runtime;

if runtime::sync_available() {
    // Can use sync APIs (native only)
    let device = GpuDevice::new()?;
} else {
    // Must use async APIs (WASM)
    let device = GpuDevice::new_async().await?;
}

Available Operations

Element-wise Operations

Vector Operations

Matrix Operations

WebGPU for WASM

The gpu-wasm feature enables GPU compute in browsers via WebGPU. This is particularly useful for:

• Browser-based ML inference: Run models client-side
• Interactive visualizations: GPU-accelerated data processing
• Scientific computing in browsers: Heavy computations without server round-trips

Example: trueno-viz

trueno-viz demonstrates Trueno's WebGPU capabilities for browser-based visualization:

// In WASM context, use async API
#[wasm_bindgen]
pub async fn process_data(input: &[f32]) -> Result<Vec<f32>, JsValue> {
    let device = GpuDevice::new_async().await
        .map_err(|e| JsValue::from_str(&e))?;

    let mut output = vec![0.0; input.len()];
    device.relu_async(input, &mut output).await
        .map_err(|e| JsValue::from_str(&e))?;

    Ok(output)
}

WASM Build Configuration

# Cargo.toml
[target.'cfg(target_arch = "wasm32")'.dependencies]
trueno = { version = "0.7.3", features = ["gpu-wasm"] }
wasm-bindgen = "0.2"
wasm-bindgen-futures = "0.4"

Build with:

wasm-pack build --target web --features gpu-wasm

Batch API

For chaining multiple GPU operations, use the batch API to minimize transfer overhead:

use trueno::backends::gpu::{GpuDevice, GpuCommandBatch};

let device = GpuDevice::new()?;
let mut batch = GpuCommandBatch::new(device);

// Queue operations (no GPU execution yet)
let input = batch.upload(&[1.0, 2.0, -3.0, 4.0]);
let a = batch.relu(input);
let b = batch.scale(a, 2.0);

// Execute batch in single GPU round-trip
batch.execute().await?;

// Read result
let result = batch.read(b).await?;

See GPU Performance for detailed batch API documentation.

Performance Considerations

When to Use GPU

✅ Use GPU for:

• Matrix multiplication >500×500
• 2D convolutions with large kernels
• Batched operations (multiple ops chained)

❌ Use SIMD instead for:

• Vector operations (add, mul, dot)
• Small matrices (<500×500)
• Single operations (transfer overhead dominates)

Transfer Overhead

GPU operations incur ~3.5ms fixed overhead per operation:

This overhead makes GPU slower than SIMD for simple operations. See GPU Performance for benchmarks.

Implementation Details

Runtime Module

The runtime module (src/backends/gpu/runtime.rs) provides platform-specific async runtime helpers:

// Native: Uses pollster for blocking
#[cfg(all(feature = "gpu", not(target_arch = "wasm32")))]
pub fn block_on<F: Future>(f: F) -> F::Output {
    pollster::block_on(f)
}

// Check if sync operations are available
pub const fn sync_available() -> bool {
    #[cfg(not(target_arch = "wasm32"))]
    { true }
    #[cfg(target_arch = "wasm32")]
    { false }
}

// WASM: Spawn async tasks
#[cfg(all(feature = "gpu-wasm", target_arch = "wasm32"))]
pub fn spawn_local<F: Future<Output = ()> + 'static>(f: F) {
    wasm_bindgen_futures::spawn_local(f);
}

Conditional Compilation

Sync methods are only available on native platforms:

#[cfg(all(feature = "gpu", not(target_arch = "wasm32")))]
pub fn relu(&self, input: &[f32], result: &mut [f32]) -> Result<(), String> {
    runtime::block_on(self.relu_async(input, result))
}

// Async always available
pub async fn relu_async(&self, input: &[f32], result: &mut [f32]) -> Result<(), String> {
    // Implementation
}

Next Steps

• GPU Performance - Detailed benchmarks and thresholds
• WASM Backend - SIMD128 for non-GPU WASM
• Backend Selection - How Trueno chooses backends

PTX Code Generation (trueno-gpu)

trueno-gpu provides pure Rust PTX (Parallel Thread Execution) code generation for NVIDIA GPUs. This enables GPU kernel development without requiring LLVM, nvcc, or any external dependencies.

Philosophy

Own the Stack - Build everything from first principles for complete control, auditability, and reproducibility.

Quick Start

use trueno_gpu::ptx::{PtxModule, PtxKernel, PtxType};

// Create a PTX module
let module = PtxModule::new()
    .version(8, 0)      // PTX ISA 8.0
    .target("sm_70")    // Volta+
    .address_size(64);  // 64-bit addressing

// Build a kernel with the fluent builder API
let kernel = PtxKernel::new("my_kernel")
    .param(PtxType::U64, "data_ptr")
    .param(PtxType::U32, "n")
    .build(|ctx| {
        // Generate PTX instructions
        let tid = ctx.special_reg(trueno_gpu::ptx::PtxReg::TidX);
        // ... more instructions
        ctx.ret();
    });

// Emit PTX source
let ptx_source = module.add_kernel(kernel).emit();

Module Structure

A PTX module consists of:

• Header: Version, target architecture, address size
• Declarations: Register declarations, shared memory
• Kernels: One or more entry points

Version and Target

// PTX ISA 8.0 for Ampere and newer
.version(8, 0)

// Target compute capability
.target("sm_70")  // Volta
.target("sm_75")  // Turing
.target("sm_80")  // Ampere
.target("sm_89")  // Ada Lovelace
.target("sm_90")  // Hopper

Kernel Builder API

The KernelBuilder provides a fluent API for generating PTX instructions:

Special Registers

// Thread and block IDs
ctx.special_reg(PtxReg::TidX);    // %tid.x
ctx.special_reg(PtxReg::TidY);    // %tid.y
ctx.special_reg(PtxReg::CtaIdX);  // %ctaid.x (block ID)
ctx.special_reg(PtxReg::NtidX);   // %ntid.x (block size)

Arithmetic Operations

// Integer arithmetic
ctx.add_u32(a, b);
ctx.mul_wide_u32(a, b);     // 32x32 -> 64 bit
ctx.mad_lo_u32(a, b, c);    // a*b + c (low 32 bits)

// Floating point
ctx.add_f32(a, b);
ctx.mul_f32(a, b);
ctx.fma_f32(a, b, c);       // Fused multiply-add

Memory Operations

// Load from global memory
let value = ctx.ld_global_f32(addr);

// Store to global memory
ctx.st_global_f32(addr, value);

// Load kernel parameters
let param = ctx.load_param_u32("param_name");
let ptr = ctx.load_param_u64("ptr_param");

Control Flow

// Predicated branch
let pred = ctx.setp_ge_u32(idx, n);  // idx >= n
ctx.branch_if(pred, "exit");

// Unconditional branch
ctx.branch("loop_start");

// Labels
ctx.label("loop_start");
ctx.label("exit");

// Return
ctx.ret();

Pre-built Kernels

trueno-gpu includes optimized kernel generators:

GEMM (Matrix Multiplication)

use trueno_gpu::kernels::{GemmKernel, Kernel};

// Naive GEMM (for correctness testing)
let kernel = GemmKernel::naive(1024, 1024, 1024);

// Tiled GEMM (shared memory optimization)
let kernel = GemmKernel::tiled(1024, 1024, 1024, 32);

// Tensor Core GEMM (SM 7.0+)
let kernel = GemmKernel::tensor_core(1024, 1024, 1024);

// Generate PTX
let ptx = kernel.emit_ptx();

Softmax

use trueno_gpu::kernels::{SoftmaxKernel, Kernel};

let kernel = SoftmaxKernel::new(1024);  // Vector length
let ptx = kernel.emit_ptx();

Bias + Activation (Epilogue Kernel)

Fused bias addition with optional activation function, commonly used as an epilogue after GEMM:

use trueno_gpu::kernels::{BiasActivationKernel, Activation, Kernel};

// Bias only (no activation)
let kernel = BiasActivationKernel::new(4096, 256);  // n=4096, bias_size=256

// Bias + ReLU
let kernel = BiasActivationKernel::new(4096, 256).with_relu();

// Bias + GELU (Transformer default)
let kernel = BiasActivationKernel::new(4096, 256).with_gelu();

// Custom activation via builder
let kernel = BiasActivationKernel::new(4096, 256)
    .with_activation(Activation::GELU);

let ptx = kernel.emit_ptx();

Note: The bias_size is baked into the kernel at generation time for efficiency. The kernel computes output[i] += bias[i % bias_size].

# Run the example
cargo run -p trueno-gpu --example bias_activation

# Run property tests and falsification tests
cargo test -p trueno-gpu bias_activation

# Run deep bug hunt (includes BiasActivation)
cargo run -p trueno-explain --example deep_bug_hunt

Testing: BiasActivationKernel includes 22 tests covering:

• Unit tests for configuration and PTX structure
• Property-based tests (proptest) for randomized validation
• Falsification tests verifying bounds checks, bias modulo, and activation correctness
• Mutation testing: 100% coverage (2 caught by tests, 4 caught by type system)

Quantized GEMM (Q4_K, Q5_K, Q6_K)

Optimized kernels for quantized inference with GGML-compatible formats:

use trueno_gpu::kernels::{QuantizeKernel, Q5KKernel, Q6KKernel, Kernel};

// Q4_K: 4-bit quantization (144 bytes per 256 values)
let q4k = QuantizeKernel::ggml(1024, 1024, 4096);

// Q5_K: 5-bit quantization (176 bytes per 256 values) - PARITY-116
let q5k = Q5KKernel::new(1024, 1024, 4096);

// Q6_K: 6-bit quantization (210 bytes per 256 values) - PARITY-117
let q6k = Q6KKernel::new(1024, 1024, 4096);

let ptx = q5k.emit_ptx();

Memory Management

use trueno_gpu::memory::{MemoryPool, PoolConfig, GpuBuffer};

// Create memory pool
let config = PoolConfig::new(1024 * 1024 * 1024);  // 1GB
let pool = MemoryPool::new(config);

// Allocate buffer
let buffer: GpuBuffer<f32> = GpuBuffer::new(1024);

Backend Detection

use trueno_gpu::backend::{detect_backend, Backend};

let backend = detect_backend();
println!("Using backend: {}", backend.name());
println!("Available: {}", backend.is_available());

Running Examples

# PTX quickstart - vector addition kernel
cargo run -p trueno-gpu --example ptx_quickstart

# GEMM kernel generation
cargo run -p trueno-gpu --example gemm_kernel

# Bias + Activation epilogue kernel
cargo run -p trueno-gpu --example bias_activation

# Quantized GEMM (Q5_K/Q6_K)
cargo run -p trueno-gpu --example q5k_q6k_gemm

PTX Type System

State Spaces

Best Practices

1. Minimize global memory access - Use shared memory for data reuse
2. Coalesce memory accesses - Adjacent threads access adjacent memory
3. Use FMA instructions - fma_f32 is faster than separate mul+add
4. Avoid branch divergence - Keep warps executing the same path
5. Maximize occupancy - Balance register usage vs parallelism

Feature Flags

[dependencies]
trueno-gpu = { version = "0.1", features = ["cuda"] }

• default - PTX generation only (no CUDA runtime required)
• cuda - Enable CUDA driver FFI for actual execution

Resources

• PTX ISA Documentation
• CUDA C Programming Guide
• trueno-gpu Examples

PTX Register Allocation Architecture

This chapter explains trueno-gpu's approach to register allocation, which delegates physical register assignment to NVIDIA's ptxas compiler. This is a pragmatic design that leverages 30+ years of GPU compiler optimization.

The Traditional Compiler Problem

In traditional compilers (like LLVM for x86), you must map an infinite number of variables to a finite set of physical registers (e.g., RAX, RDI, RSI on x86-64). This requires complex algorithms:

• Graph Coloring: Model register interference as a graph, color with K colors (K = number of physical registers)
• Linear Scan: Faster but less optimal allocation for JIT compilers

These algorithms are complex to implement correctly and require significant engineering effort.

Trueno's Strategy: Virtual Registers + ptxas

Trueno takes a different approach that leverages PTX's design as a virtual ISA:

┌─────────────────────────────────────────────────────────────┐
│  Trueno PTX Builder (Rust)                                  │
│  - Allocates unlimited virtual registers (%f0, %f1, ...)    │
│  - Tracks liveness for pressure REPORTING                   │
│  - Emits SSA-style PTX                                      │
└─────────────────────────────────────────────────────────────┘
                             │
                        PTX Source
                             │
                             ▼
┌─────────────────────────────────────────────────────────────┐
│  NVIDIA ptxas (JIT Compiler)                                │
│  - Graph coloring for physical register allocation          │
│  - Register spilling to local memory if needed              │
│  - Dead code elimination, constant folding, etc.            │
└─────────────────────────────────────────────────────────────┘
                             │
                        SASS Binary
                             │
                             ▼
┌─────────────────────────────────────────────────────────────┐
│  GPU Execution                                              │
└─────────────────────────────────────────────────────────────┘

How It Works

1. Virtual Register Allocation: Each operation allocates a new virtual register with a monotonically increasing ID:

// In trueno-gpu's KernelBuilder
pub fn add_f32(&mut self, a: VirtualReg, b: VirtualReg) -> VirtualReg {
    // Allocate NEW virtual register (SSA style)
    let dst = self.registers.allocate_virtual(PtxType::F32);
    self.instructions.push(
        PtxInstruction::new(PtxOp::Add, PtxType::F32)
            .dst(Operand::Reg(dst))
            .src(Operand::Reg(a))
            .src(Operand::Reg(b))
    );
    dst  // Return %f2, %f3, %f4, etc.
}

2. Per-Type Namespaces: PTX requires separate register namespaces per type:

3. Emitted PTX: The builder emits register declarations and instructions:

.visible .entry vector_add(
    .param .u64 a_ptr,
    .param .u64 b_ptr,
    .param .u64 c_ptr,
    .param .u32 n
) {
    .reg .f32  %f<3>;    // Virtual registers %f0, %f1, %f2
    .reg .u32  %r<5>;    // Virtual registers %r0-4
    .reg .u64  %rd<7>;   // Virtual registers %rd0-6
    .reg .pred %p<1>;    // Predicate register %p0

    // Instructions use virtual registers
    mov.u32 %r0, %tid.x;
    mov.u32 %r1, %ctaid.x;
    // ...
    add.rn.f32 %f2, %f0, %f1;
    // ...
}

4. ptxas Does the Rest: NVIDIA's ptxas compiler:
   • Builds an interference graph from virtual register liveness
   • Performs graph coloring to assign physical registers
   • Generates spill code if necessary (to .local memory)
   • Applies optimization passes

Why This Design?

1. Pragmatism (Avoid Muda)

NVIDIA has invested 30+ years into GPU compiler optimization. Reimplementing graph coloring would be:

• Redundant (ptxas already does it)
• Inferior (we can't match NVIDIA's GPU-specific knowledge)
• Wasteful engineering effort (Muda in Toyota terms)

2. PTX is Designed for This

PTX (Parallel Thread Execution) is explicitly designed as a virtual ISA:

• Unlimited virtual registers
• SSA (Static Single Assignment) form
• Meant to be lowered by a backend compiler

From the PTX ISA documentation:

"PTX defines a virtual machine and ISA for general purpose parallel thread execution."

3. Focus on What Matters

Trueno focuses on:

• Algorithm correctness: Ensuring SIMD/GPU operations produce correct results
• High-level optimization: Tiling, kernel fusion, memory access patterns
• Developer experience: Safe, ergonomic Rust API

Low-level optimization (register allocation, instruction scheduling) is delegated to specialized tools.

Register Pressure Monitoring

While we don't perform graph coloring, we DO track liveness for diagnostics:

pub struct RegisterAllocator {
    type_counters: HashMap<PtxType, u32>,
    live_ranges: HashMap<(PtxType, u32), LiveRange>,
    spill_count: usize,  // Muda tracking
}

impl RegisterAllocator {
    pub fn pressure_report(&self) -> RegisterPressure {
        RegisterPressure {
            max_live: self.allocated.len(),
            spill_count: self.spill_count,
            utilization: max_live as f64 / 256.0,
        }
    }
}

Why Track Pressure?

1. Developer Warnings: Alert when kernels exceed 256 registers/thread
2. Occupancy Estimation: High register usage reduces concurrent threads
3. Performance Debugging: Identify kernels that may suffer from register spills

GPU Register Limits

Occupancy Impact: If a kernel uses 64 registers/thread, an SM with 65,536 registers can run 1024 threads. If it uses 128 registers/thread, only 512 threads can run concurrently.

In-Place Operations for Register Reuse

For loops and accumulators, SSA-style allocation wastes registers:

// SSA style - allocates new register each iteration
for _ in 0..1000 {
    let new_sum = ctx.add_f32(sum, val);  // New register each time!
    sum = new_sum;
}

We provide in-place operations that reuse registers:

// In-place style - reuses existing register
let acc = ctx.mov_f32_imm(0.0);  // Allocate once
for _ in 0..1000 {
    ctx.add_f32_inplace(acc, val);  // Reuses %f0
}

Available In-Place Operations

Potential Future Enhancements

The current design delegates all register allocation to ptxas. Potential future enhancements (tracked in GitHub Issue #66):

1. Greedy Register Reuse

For kernels exceeding 256 registers, we could implement simple liveness-based reuse:

// Hypothetical future API
let allocator = RegisterAllocator::new()
    .with_reuse_strategy(ReuseStrategy::Greedy);

This would reuse %r2 after its last use, reducing virtual register count.

2. ptxas Output Parsing

Parse cuobjdump --dump-resource-usage output to validate:

• Expected vs actual register usage
• Spill detection
• Occupancy calculation

3. Occupancy Calculator

Integrate NVIDIA's occupancy calculator to predict SM utilization before runtime.

Best Practices

1. Use In-Place Operations for Loops

// Good - register reuse
let i = ctx.mov_u32_imm(0);
ctx.label("loop");
// ... loop body ...
ctx.add_u32_inplace(i, 1);  // Reuses %r0
ctx.branch("loop");

// Bad - register explosion
let mut i = ctx.mov_u32_imm(0);
ctx.label("loop");
// ... loop body ...
i = ctx.add_u32(i, 1);  // New register each iteration!
ctx.branch("loop");

2. Limit Unroll Factors

Each unrolled iteration adds registers. Balance throughput vs pressure:

// High pressure - 8x unroll
for i in 0..8 {
    let val = ctx.ld_global_f32(addr[i]);
    ctx.fma_f32_inplace(acc, val, weights[i]);
}

// Lower pressure - 4x unroll (often sufficient)
for i in 0..4 {
    let val = ctx.ld_global_f32(addr[i]);
    ctx.fma_f32_inplace(acc, val, weights[i]);
}

3. Use Shared Memory for Large Temporaries

Instead of keeping many values in registers, stage through shared memory:

// Use shared memory tile instead of many registers
let tile = ctx.alloc_shared::<f32>(TILE_SIZE * TILE_SIZE);

4. Monitor Kernel Complexity

For complex kernels, check register pressure:

let pressure = kernel.registers.pressure_report();
if pressure.utilization > 0.5 {
    eprintln!("Warning: High register pressure ({:.0}%)",
              pressure.utilization * 100.0);
}

Running the Example

cargo run -p trueno-gpu --example register_allocation

This demonstrates:

1. Simple kernel with low register pressure
2. Complex kernel with higher pressure (unrolled dot product)
3. In-place operations for register reuse
4. Architectural trade-offs

References

• PTX ISA Documentation
• CUDA Occupancy Calculator
• GitHub Issue #66: Liveness-Based Register Reuse
• Example: trueno-gpu/examples/register_allocation.rs

PTX Optimization Passes

This chapter documents the PTX optimization passes in trueno-gpu, aligned with NVIDIA's official CUDA Tile IR (CUDA Toolkit 13.1).

Overview

The trueno_gpu::ptx::optimize module provides four optimization passes:

FMA Fusion Pass

The FMA (Fused Multiply-Add) fusion pass detects mul + add instruction patterns and fuses them into a single fma instruction.

Benefits

• Latency: Single instruction instead of two
• Precision: Single rounding operation (IEEE 754 compliant)
• Throughput: Utilizes GPU FMA units efficiently

Example

use trueno_gpu::ptx::optimize::fma_fusion;
use trueno_gpu::ptx::{Operand, PtxInstruction, PtxOp, PtxType, VirtualReg};

// Create mul + add pattern
let r0 = VirtualReg::new(0, PtxType::F32);
let r1 = VirtualReg::new(1, PtxType::F32);
let r2 = VirtualReg::new(2, PtxType::F32);
let r3 = VirtualReg::new(3, PtxType::F32);

let mul = PtxInstruction::new(PtxOp::Mul, PtxType::F32)
    .dst(Operand::Reg(r2.clone()))
    .src(Operand::Reg(r0.clone()))
    .src(Operand::Reg(r1.clone()));

let add = PtxInstruction::new(PtxOp::Add, PtxType::F32)
    .dst(Operand::Reg(r3))
    .src(Operand::Reg(r2))
    .src(Operand::ImmF32(1.0));

// Fuse to single FMA instruction
let fused = fma_fusion::pass(vec![mul, add]);
assert_eq!(fused.len(), 1); // mul + add → fma

Academic Reference

Based on Click & Paleczny (1995) "A Simple Graph-Based Intermediate Representation" for SSA pattern matching.

Loop Splitting Pass

The loop splitting pass analyzes conditional loops and identifies opportunities to split them at condition boundaries, eliminating branch divergence in GPU warps.

Heavy Operations

The following operations trigger split profitability:

• Ld - Memory loads
• St - Memory stores
• WmmaMma - Tensor Core MMA
• WmmaLoadA, WmmaLoadB, WmmaLoadC - WMMA fragment loads
• WmmaStoreD - WMMA fragment stores

Example

use trueno_gpu::ptx::optimize::loop_split;
use trueno_gpu::ptx::{PtxInstruction, PtxOp, PtxType, CmpOp};

// Check profitability
let heavy_op = PtxInstruction::new(PtxOp::Ld, PtxType::F32);
assert!(loop_split::is_split_profitable(&[heavy_op], 10));

let light_op = PtxInstruction::new(PtxOp::Add, PtxType::F32);
assert!(!loop_split::is_split_profitable(&[light_op], 10));

// Split point alignment for non-unit steps
assert_eq!(loop_split::align_split_point(5, 0, 4), 8);
assert_eq!(loop_split::align_split_point(8, 0, 4), 8);

// Loop predicate conversion
assert_eq!(
    loop_split::LoopPredicate::from_cmp_op(CmpOp::Lt),
    Some(loop_split::LoopPredicate::LessThan)
);

NVIDIA Reference

Aligned with LoopSplit.cpp from NVIDIA CUDA Tile IR (CUDA Toolkit 13.1).

Token-Based Ordering (TKO)

Token-Based Ordering provides explicit memory dependency tracking, enabling compiler-driven barrier elimination.

Memory Ordering Semantics

Memory Scopes

Example

use trueno_gpu::ptx::optimize::tko;

// Create tokens for memory operations
let t1 = tko::Token::new();
let t2 = tko::Token::new();
let t3 = tko::Token::new();

// Join tokens at synchronization point
let joined = tko::join_tokens(&[t1, t2, t3]);

// Memory ordering
let ordering = tko::MemoryOrdering::Acquire;
assert_eq!(ordering.to_ptx_modifier(), ".acquire");

// Memory scope
let scope = tko::MemoryScope::Device;
assert_eq!(scope.to_ptx_scope(), ".gpu");

// Token graph with cycle detection
let mut graph = tko::TokenGraph::new();
let ta = tko::Token::new();
let tb = tko::Token::new();
let tc = tko::Token::new();

graph.create_token(ta);
graph.create_token(tb);
graph.create_token(tc);
graph.add_dependency(tb, ta);
graph.add_dependency(tc, tb);

assert!(!graph.has_cycle()); // No deadlock

graph.add_dependency(ta, tc);
assert!(graph.has_cycle()); // DEADLOCK!

NVIDIA Reference

Aligned with memory_consistency_ops.mlir from NVIDIA CUDA Tile IR.

Tile Validation

Tile validation enforces constraints to prevent register pressure issues and compilation hangs.

Constraints

1. Power-of-two dimensions: Required for efficient GPU scheduling
2. Maximum tile elements: 16M elements to prevent register spills
3. Maximum single dimension: 4096 to prevent degenerate shapes

WMMA Valid Shapes

Example

use trueno_gpu::ptx::optimize::tile_validation;
use trueno_gpu::ptx::WmmaShape;

// Valid shapes
assert!(tile_validation::validate_shape(&[16, 16]).is_ok());
assert!(tile_validation::validate_shape(&[32, 32]).is_ok());
assert!(tile_validation::validate_shape(&[64, 64]).is_ok());

// Invalid shapes
assert!(tile_validation::validate_shape(&[17, 16]).is_err()); // Not power of two
assert!(tile_validation::validate_shape(&[100, 100]).is_err());

// WMMA shapes
let valid_wmma = WmmaShape::M16N16K16;
assert!(tile_validation::validate_wmma_shape(&valid_wmma).is_ok());

let invalid_wmma = WmmaShape { m: 24, n: 24, k: 16 };
assert!(tile_validation::validate_wmma_shape(&invalid_wmma).is_err());

Academic Reference

Based on Volkov & Demmel (2008) "Benchmarking GPUs to Tune Dense Linear Algebra".

Running the Example

cargo run --example ptx_optimize

Output:

╔══════════════════════════════════════════════════════════════╗
║     PTX Optimization Passes (NVIDIA CUDA Tile IR Aligned)    ║
╚══════════════════════════════════════════════════════════════╝

1️⃣  FMA FUSION PASS
   Input:  2 instructions (mul + add)
   Output: 1 instruction (fma)

2️⃣  LOOP SPLITTING PASS
   Heavy ops trigger split: true
   Light ops trigger split: false

3️⃣  TOKEN-BASED ORDERING (TKO)
   Tokens created with unique IDs
   Cycle detection: working

4️⃣  TILE VALIDATION
   Power-of-two shapes: OK
   Invalid shapes: rejected

✅ All optimization demos completed successfully!

Specification

Full specification: cuda-tile-behavior.md (v1.4.0)

Coverage

Runtime Detection

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Vector Operations

The Vector<T> type is the core data structure in Trueno, providing SIMD-accelerated operations on contiguous arrays of floating-point numbers.

Creating Vectors

use trueno::{Vector, Backend};

// From a slice (uses best available backend)
let v = Vector::<f32>::from_slice(&[1.0, 2.0, 3.0, 4.0]);

// With explicit backend
let v_scalar = Vector::<f32>::from_slice_with_backend(
    &[1.0, 2.0, 3.0],
    Backend::Scalar
);

// From Vec
let v = Vector::<f32>::from_vec(vec![1.0, 2.0, 3.0, 4.0]);

Element-wise Operations

All element-wise operations return a new Vector with the same length.

let a = Vector::<f32>::from_slice(&[1.0, 2.0, 3.0]);
let b = Vector::<f32>::from_slice(&[4.0, 5.0, 6.0]);

// Arithmetic
let sum = a.add(&b)?;      // [5.0, 7.0, 9.0]
let diff = a.sub(&b)?;     // [-3.0, -3.0, -3.0]
let prod = a.mul(&b)?;     // [4.0, 10.0, 18.0]
let quot = a.div(&b)?;     // [0.25, 0.4, 0.5]

// Scalar operations
let scaled = a.scale(2.0)?; // [2.0, 4.0, 6.0]

// Math functions
let sqrts = a.sqrt()?;
let exps = a.exp()?;
let logs = a.ln()?;

Reduction Operations

Reductions collapse a vector to a single value.

let v = Vector::<f32>::from_slice(&[1.0, 2.0, 3.0, 4.0]);

let total = v.sum()?;        // 10.0
let maximum = v.max()?;      // 4.0
let minimum = v.min()?;      // 1.0
let dot = a.dot(&b)?;        // Dot product

// Norms
let l1 = v.norm_l1()?;       // Manhattan norm
let l2 = v.norm_l2()?;       // Euclidean norm
let linf = v.norm_linf()?;   // Max absolute value

// Argmax/Argmin
let idx_max = v.argmax()?;   // Index of max element
let idx_min = v.argmin()?;   // Index of min element

Activation Functions

Common neural network activations, optimized for ML inference.

let x = Vector::<f32>::from_slice(&[-2.0, -1.0, 0.0, 1.0, 2.0]);

// Classic activations
let relu = x.relu()?;
let sigmoid = x.sigmoid()?;
let tanh_v = x.tanh_activation()?;

// Modern activations (Transformer era)
let gelu = x.gelu()?;       // BERT, GPT
let swish = x.swish()?;     // EfficientNet
let mish = x.mish()?;       // YOLOv4

// Variants
let leaky = x.leaky_relu(0.01)?;
let elu = x.elu(1.0)?;
let selu = x.selu()?;

Layer Normalization

For transformer architectures.

let hidden = Vector::<f32>::from_slice(&[1.0, 2.0, 3.0, 4.0]);
let gamma = Vector::<f32>::from_slice(&[1.0, 1.0, 1.0, 1.0]); // scale
let beta = Vector::<f32>::from_slice(&[0.0, 0.0, 0.0, 0.0]);  // shift

let normalized = hidden.layer_norm(&gamma, &beta, 1e-5)?;
// Output has mean ≈ 0, variance ≈ 1

Similarity Metrics

For ML applications like recommendation systems.

let a = Vector::<f32>::from_slice(&[1.0, 2.0, 3.0]);
let b = Vector::<f32>::from_slice(&[4.0, 5.0, 6.0]);

let cosine = a.cosine_similarity(&b)?;  // [-1, 1]
let euclidean = a.euclidean_distance(&b)?;
let manhattan = a.manhattan_distance(&b)?;

Backend Selection

Vectors automatically use the best available SIMD backend.

use trueno::{select_best_available_backend, OperationType};

// Check what's available
let backend = select_best_available_backend();
println!("Using: {:?}", backend); // e.g., AVX2

// Operation-aware selection (memory-bound vs compute-bound)
let mem_backend = select_backend_for_operation(OperationType::MemoryBound);
let compute_backend = select_backend_for_operation(OperationType::ComputeBound);

Performance Characteristics

See Performance Guide for detailed analysis.

Matrix Operations

The Matrix<T> type provides 2D matrix operations with SIMD acceleration.

Creating Matrices

use trueno::Matrix;

// From dimensions (uninitialized)
let m = Matrix::<f32>::new(3, 4);

// From Vec with dimensions
let m = Matrix::<f32>::from_vec(2, 3, vec![
    1.0, 2.0, 3.0,
    4.0, 5.0, 6.0,
])?;

// Special matrices
let zeros = Matrix::<f32>::zeros(3, 3);
let identity = Matrix::<f32>::identity(4);

Basic Properties

let m = Matrix::<f32>::from_vec(2, 3, vec![1.0, 2.0, 3.0, 4.0, 5.0, 6.0])?;

m.rows();        // 2
m.cols();        // 3
m.len();         // 6 (total elements)
m.as_slice();    // &[f32] view of data
m.get(0, 1);     // Some(2.0)
m.get_mut(1, 2); // Mutable access

Matrix Multiplication

let a = Matrix::<f32>::from_vec(2, 3, vec![1.0, 2.0, 3.0, 4.0, 5.0, 6.0])?;
let b = Matrix::<f32>::from_vec(3, 2, vec![7.0, 8.0, 9.0, 10.0, 11.0, 12.0])?;

// Matrix-matrix multiplication: [2×3] × [3×2] = [2×2]
let c = a.matmul(&b)?;

Matrix-Vector Multiplication

use trueno::Vector;

let m = Matrix::<f32>::from_vec(3, 4, vec![/* 12 elements */])?;
let v = Vector::<f32>::from_slice(&[1.0, 2.0, 3.0, 4.0]);

// Matrix × Vector: [3×4] × [4×1] = [3×1]
let result = m.matvec(&v)?;

// Vector × Matrix: [1×3] × [3×4] = [1×4]
let v2 = Vector::<f32>::from_slice(&[1.0, 2.0, 3.0]);
let result = m.vecmat(&v2)?;

Transpose

let m = Matrix::<f32>::from_vec(2, 3, vec![1.0, 2.0, 3.0, 4.0, 5.0, 6.0])?;

// [2×3] → [3×2]
let mt = m.transpose();

Convolution (2D)

For image processing and CNNs.

let image = Matrix::<f32>::from_vec(5, 5, /* 25 elements */)?;
let kernel = Matrix::<f32>::from_vec(3, 3, vec![
    1.0, 0.0, -1.0,
    2.0, 0.0, -2.0,
    1.0, 0.0, -1.0,
])?; // Sobel edge detection

let edges = image.convolve2d(&kernel)?;

Embedding Lookup

For NLP models (word embeddings).

// Embedding table: vocab_size × embedding_dim
let embeddings = Matrix::<f32>::from_vec(1000, 128, /* ... */)?;

// Token indices
let tokens: Vec<usize> = vec![42, 7, 256, 13];

// Lookup: returns [4×128] matrix
let token_embeddings = embeddings.embedding_lookup(&tokens)?;

Batched Matrix Multiplication (3D Tensors)

For batch processing of independent matrix multiplications:

// Shape: [batch, m, k] @ [batch, k, n] -> [batch, m, n]
let batch = 4;
let m = 32;
let k = 64;
let n = 32;

// Flattened input tensors
let a_data: Vec<f32> = vec![0.0; batch * m * k];
let b_data: Vec<f32> = vec![0.0; batch * k * n];

let result = Matrix::batched_matmul(&a_data, &b_data, batch, m, k, n)?;
// Result: Vec<f32> with shape [batch, m, n]

Batched 4D Matrix Multiplication (Attention Pattern)

For multi-head attention in transformers:

// Shape: [batch, heads, m, k] @ [batch, heads, k, n] -> [batch, heads, m, n]
// This is the exact pattern for Q @ K^T and attn @ V in attention

let batch = 1;
let heads = 12;  // Number of attention heads
let seq_len = 512;
let head_dim = 64;

// Q: [batch, heads, seq_len, head_dim]
let q_data: Vec<f32> = vec![0.0; batch * heads * seq_len * head_dim];
// K^T: [batch, heads, head_dim, seq_len] (already transposed)
let kt_data: Vec<f32> = vec![0.0; batch * heads * head_dim * seq_len];

// Compute attention scores: Q @ K^T
let attn_scores = Matrix::batched_matmul_4d(
    &q_data,
    &kt_data,
    batch,
    heads,
    seq_len,   // m
    head_dim,  // k
    seq_len,   // n
)?;
// Result: [batch, heads, seq_len, seq_len] attention scores

This is critical for transformer performance - each (batch, head) pair is processed independently using SIMD matmul.

GPU Acceleration

For large matrices, use the GPU backend.

use trueno::GpuBackend;

let mut gpu = GpuBackend::new();
let a = Matrix::<f32>::from_vec(1024, 1024, /* ... */)?;
let b = Matrix::<f32>::from_vec(1024, 1024, /* ... */)?;

// GPU-accelerated matmul
let c = gpu.matmul(&a, &b)?;

Performance Tips

1. Matrix multiplication: O(n³) - GPU beneficial for n > 500
2. Convolution: Use separable kernels when possible
3. Memory layout: Row-major storage for cache efficiency
4. Batch operations: Group small matrices for GPU efficiency

See the GPU Performance Guide for details.

Eigendecomposition

The SymmetricEigen type provides eigendecomposition for symmetric matrices, essential for PCA, spectral clustering, and scientific computing.

Basic Usage

use trueno::{Matrix, SymmetricEigen};

// Create a symmetric matrix
let m = Matrix::<f32>::from_vec(3, 3, vec![
    4.0, 2.0, 0.0,
    2.0, 5.0, 3.0,
    0.0, 3.0, 6.0,
])?;

// Compute eigendecomposition
let eigen = SymmetricEigen::new(&m)?;

// Access results
let eigenvalues = eigen.eigenvalues();     // Sorted descending
let eigenvectors = eigen.eigenvectors();   // As matrix (columns = eigenvectors)

Eigenvalues

Eigenvalues are returned in descending order (PCA convention).

let eigen = SymmetricEigen::new(&covariance_matrix)?;

// Largest eigenvalue first
let principal = eigen.eigenvalues()[0];

// Variance explained by first PC
let total_variance: f32 = eigen.eigenvalues().iter().sum();
let explained = eigen.eigenvalues()[0] / total_variance;
println!("First PC explains {:.1}% of variance", explained * 100.0);

Eigenvectors

Eigenvectors form an orthonormal basis.

let eigen = SymmetricEigen::new(&m)?;

// Get i-th eigenvector as a Vector
let v0 = eigen.eigenvector(0)?;

// Eigenvectors are orthonormal
let dot = v0.dot(&eigen.eigenvector(1)?)?;
assert!(dot.abs() < 1e-5); // ≈ 0

// Unit length
let norm = v0.norm_l2()?;
assert!((norm - 1.0).abs() < 1e-5); // ≈ 1

Verification

Verify A × v = λ × v for each eigenpair.

let eigen = SymmetricEigen::new(&m)?;

for i in 0..eigen.len() {
    let lambda = eigen.eigenvalues()[i];
    let v = eigen.eigenvector(i)?;

    let av = m.matvec(&v)?;
    let lambda_v = v.scale(lambda)?;

    let error: f32 = av.sub(&lambda_v)?
        .as_slice()
        .iter()
        .map(|x| x.abs())
        .sum();

    assert!(error < 1e-5, "Eigenpair {} invalid", i);
}

Reconstruction

Reconstruct the original matrix: A = V × D × Vᵀ

let eigen = SymmetricEigen::new(&m)?;

// V * diag(eigenvalues) * V^T should equal original matrix
let reconstructed = eigen.reconstruct();
let error = m.frobenius_distance(&reconstructed);
assert!(error < 1e-5);

GPU Acceleration

For large matrices, use GPU backend.

use trueno::GpuBackend;

let mut gpu = GpuBackend::new();
let large = Matrix::<f32>::from_vec(256, 256, /* ... */)?;

let (eigenvalues, eigenvectors) = gpu.symmetric_eigen(
    large.as_slice(),
    256
)?;

Algorithm Details

Trueno uses the Jacobi eigenvalue algorithm:

• Numerically stable: Based on Golub & Van Loan formulation
• Convergence: Quadratic convergence for well-conditioned matrices
• SIMD-optimized: Jacobi rotations use SIMD where beneficial
• Accuracy: Results match nalgebra to 1e-5 tolerance

Performance

Use Cases

1. PCA (Principal Component Analysis)

   let cov = compute_covariance(&data);
   let eigen = SymmetricEigen::new(&cov)?;
   let top_k_components = &eigen.eigenvalues()[0..k];

2. Spectral Clustering

   let laplacian = compute_graph_laplacian(&adjacency);
   let eigen = SymmetricEigen::new(&laplacian)?;
   let fiedler_vector = eigen.eigenvector(1)?; // 2nd smallest

3. Vibration Analysis

   let stiffness = compute_stiffness_matrix(&structure);
   let eigen = SymmetricEigen::new(&stiffness)?;
   let natural_frequencies: Vec<f32> = eigen.eigenvalues()
       .iter()
       .map(|&λ| λ.sqrt())
       .collect();

Element Wise

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Reductions

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Transformations

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Error Handling

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Backend Api

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

GPU Monitoring

This chapter covers trueno's GPU monitoring capabilities as defined in TRUENO-SPEC-010.

Overview

Trueno provides comprehensive GPU monitoring through two complementary approaches:

1. Cross-platform wgpu backend - Works on any system with Vulkan, Metal, or DX12
2. Native CUDA backend - Direct access to NVIDIA GPU information via CUDA Driver API

Quick Start

use trueno::monitor::{GpuMonitor, GpuDeviceInfo, MonitorConfig};

// Enumerate all available GPUs
let devices = GpuDeviceInfo::enumerate()?;
for dev in &devices {
    println!("[{}] {} ({:.2} GB)", dev.index, dev.name, dev.vram_gb());
}

// Create a monitor with history buffer
let monitor = GpuMonitor::new(0, MonitorConfig::default())?;

// Collect metrics over time
for _ in 0..10 {
    let metrics = monitor.collect()?;
    println!("Memory: {:.1}% used", metrics.memory.usage_percent());
}

Feature Flags

Enable features in your Cargo.toml:

[dependencies]
trueno = { version = "0.8", features = ["gpu", "cuda-monitor"] }

Device Discovery

GpuDeviceInfo

Represents a discovered GPU device:

pub struct GpuDeviceInfo {
    pub index: usize,
    pub name: String,
    pub vendor: GpuVendor,
    pub backend: GpuBackend,
    pub vram_total: u64,
    pub compute_capability: Option<(u32, u32)>,
    pub driver_version: Option<String>,
}

Methods:

• enumerate() -> Result<Vec<GpuDeviceInfo>, MonitorError> - List all GPUs
• vram_gb() -> f64 - Get VRAM in gigabytes
• supports_cuda() -> bool - Check CUDA support

GpuVendor

GPU manufacturer identification:

pub enum GpuVendor {
    Nvidia,
    Amd,
    Intel,
    Apple,
    Unknown(u32),
}

PCI Vendor ID Mapping:

GpuBackend

Graphics/compute backend:

pub enum GpuBackend {
    Vulkan,
    Metal,
    Dx12,
    Cuda,
    WebGpu,
    OpenGl,
    Cpu,
}

Memory Monitoring

GpuMemoryMetrics

Real-time memory statistics:

pub struct GpuMemoryMetrics {
    pub total: u64,      // Total VRAM in bytes
    pub used: u64,       // Used VRAM in bytes
    pub free: u64,       // Free VRAM in bytes
}

Methods:

• usage_percent() -> f64 - Memory utilization (0.0-100.0)
• available_gb() -> f64 - Free memory in GB

GpuMonitor

The GpuMonitor provides continuous monitoring with a ring buffer for history:

// Configure monitoring
let config = MonitorConfig {
    poll_interval: Duration::from_millis(100),
    history_size: 1000,
};

// Create monitor for device 0
let monitor = GpuMonitor::new(0, config)?;

// Collect a sample
let metrics = monitor.collect()?;

// Get sample age
println!("Sample age: {:?}", metrics.age());

// Check history
println!("History size: {}", monitor.sample_count());

MonitorConfig

pub struct MonitorConfig {
    pub poll_interval: Duration,  // Default: 100ms
    pub history_size: usize,      // Default: 1000
}

GpuMetrics

Complete metrics snapshot:

pub struct GpuMetrics {
    pub memory: GpuMemoryMetrics,
    pub utilization: GpuUtilization,
    pub thermal: GpuThermalMetrics,
    pub power: GpuPowerMetrics,
    pub clock: GpuClockMetrics,
    pub pcie: GpuPcieMetrics,
    pub timestamp: Instant,
}

CUDA Native Monitoring

For NVIDIA GPUs, enable cuda-monitor for accurate device information via the CUDA Driver API:

use trueno::monitor::{
    cuda_monitor_available,
    enumerate_cuda_devices,
    query_cuda_memory,
};

// Check availability
if cuda_monitor_available() {
    // Enumerate CUDA devices
    let devices = enumerate_cuda_devices()?;

    // Query real-time memory
    let mem = query_cuda_memory(0)?;
    println!("CUDA Memory: {:.1}% used", mem.usage_percent());
}

Why CUDA Native?

trueno-gpu Module

For direct CUDA access without the trueno facade:

use trueno_gpu::monitor::{CudaDeviceInfo, CudaMemoryInfo};
use trueno_gpu::driver::CudaContext;

// Query device info
let info = CudaDeviceInfo::query(0)?;
println!("GPU: {} ({:.2} GB)", info.name, info.total_memory_gb());

// Create context and query memory
let ctx = CudaContext::new(0)?;
let mem = CudaMemoryInfo::query(&ctx)?;
println!("Memory: {}", mem);  // "8192 / 24576 MB (33.3% used)"

Examples

Run the GPU Monitor Demo

# Cross-platform (wgpu)
cargo run --example gpu_monitor_demo --features gpu

# With CUDA (NVIDIA)
cargo run --example gpu_monitor_demo --features "gpu,cuda-monitor"

Run the CUDA Monitor Example

cargo run -p trueno-gpu --example cuda_monitor --features cuda

Error Handling

pub enum MonitorError {
    NoDevice,           // No GPU found
    DeviceNotFound(u32), // Specific device not found
    BackendError(String), // Backend-specific error
    ContextError(String), // Context creation failed
}

Performance Considerations

• Poll Interval: Set poll_interval based on your monitoring needs. 100ms is good for visualization; 1s is sufficient for logging.
• History Size: The ring buffer is fixed-size. Larger sizes consume more memory but allow longer history analysis.
• CUDA Context: Creating a CUDA context has overhead. Reuse GpuMonitor instances when possible.

References

• TRUENO-SPEC-010: GPU Monitoring, Tracing, and Visualization
• Nickolls et al. (2008): GPU parallel computing model
• CUDA Driver API: cuDeviceGetName, cuDeviceTotalMem, cuMemGetInfo

Hash Functions

Trueno provides SIMD-optimized hash functions designed for high-performance key-value store operations. The hash module uses the FxHash algorithm with automatic backend selection for optimal performance.

Overview

The hash module is designed for:

• Fast key hashing in KV stores
• Consistent hashing for distributed systems
• Shard/partition key assignment
• Cache key generation

API Reference

hash_key

Hash a string key to a 64-bit value.

use trueno::hash_key;

let hash = hash_key("user:1001");
println!("Hash: 0x{:016x}", hash);

Signature:

pub fn hash_key(key: &str) -> u64

Properties:

• Deterministic: Same input always produces same output
• Fast: Optimized for short keys typical in KV stores
• Non-cryptographic: Not suitable for security purposes

hash_bytes

Hash raw bytes to a 64-bit value.

use trueno::hash_bytes;

let data = b"binary data";
let hash = hash_bytes(data);

Signature:

pub fn hash_bytes(bytes: &[u8]) -> u64

hash_keys_batch

Hash multiple keys using SIMD acceleration. Automatically selects the best backend for the current CPU.

use trueno::hash_keys_batch;

let keys = ["user:1", "user:2", "user:3", "user:4"];
let hashes = hash_keys_batch(&keys);

for (key, hash) in keys.iter().zip(hashes.iter()) {
    println!("{} -> 0x{:016x}", key, hash);
}

Signature:

pub fn hash_keys_batch(keys: &[&str]) -> Vec<u64>

Performance: Batch hashing is significantly faster than individual calls when processing multiple keys. The speedup depends on the SIMD backend:

• AVX-512: Up to 8x speedup
• AVX2: Up to 4x speedup
• SSE2: Up to 2x speedup
• Scalar: Baseline (no vectorization)

hash_keys_batch_with_backend

Hash multiple keys with explicit backend selection.

use trueno::{hash_keys_batch_with_backend, Backend};

let keys = ["a", "b", "c", "d"];

// Force scalar backend (useful for testing)
let scalar_hashes = hash_keys_batch_with_backend(&keys, Backend::Scalar);

// Use automatic selection (recommended)
let auto_hashes = hash_keys_batch_with_backend(&keys, Backend::Auto);

// Results are identical regardless of backend
assert_eq!(scalar_hashes, auto_hashes);

Signature:

pub fn hash_keys_batch_with_backend(keys: &[&str], backend: Backend) -> Vec<u64>

Use Cases

Partition/Shard Assignment

use trueno::hash_keys_batch;

let keys = ["order:1001", "order:1002", "order:1003", "order:1004"];
let hashes = hash_keys_batch(&keys);

let num_partitions = 4;
for (key, hash) in keys.iter().zip(hashes.iter()) {
    let partition = hash % num_partitions;
    println!("{} -> partition {}", key, partition);
}

Consistent Key Distribution

The FxHash algorithm provides good distribution for typical key patterns:

use trueno::hash_key;

// Sequential keys still distribute well
for i in 0..10 {
    let key = format!("item:{}", i);
    let hash = hash_key(&key);
    println!("{}: 0x{:016x}", key, hash);
}

Integration with trueno-db

The hash functions are re-exported by trueno-db for use with its KV store:

use trueno_db::kv::{hash_key, hash_keys_batch, KvStore, MemoryKvStore};

// Hash-based key lookup
let store = MemoryKvStore::new();
let key = "session:abc123";
let hash = hash_key(key);
println!("Key '{}' has hash 0x{:016x}", key, hash);

Algorithm Details

Trueno uses the FxHash algorithm, which is:

• Extremely fast for small inputs (typical KV keys)
• Non-cryptographic (not suitable for security)
• Deterministic across platforms
• Well-suited for hash tables and bloom filters

Constants:

const FX_HASH_K: u64 = 0x517cc1b727220a95;

The algorithm processes input in 8-byte chunks using multiply-rotate operations, with special handling for the tail bytes.

Backend Selection

The Backend enum controls SIMD acceleration:

Runtime detection ensures the correct backend is used even when Auto is specified.

Performance Benchmarks

Typical performance on modern x86_64 hardware (10,000 keys):

The exact speedup depends on:

• Key length (shorter keys benefit more from batching)
• CPU SIMD capabilities
• Memory access patterns

Example: Complete Demo

use trueno::{hash_key, hash_keys_batch, hash_keys_batch_with_backend, Backend};

fn main() {
    // Single key hashing
    let key = "hello";
    let hash = hash_key(key);
    println!("hash_key({:?}) = 0x{:016x}", key, hash);

    // Batch hashing
    let keys = ["user:1", "user:2", "user:3", "user:4"];
    let hashes = hash_keys_batch(&keys);
    for (k, h) in keys.iter().zip(hashes.iter()) {
        println!("{} -> 0x{:016x}", k, h);
    }

    // Backend comparison
    let scalar = hash_keys_batch_with_backend(&keys, Backend::Scalar);
    let auto = hash_keys_batch_with_backend(&keys, Backend::Auto);
    assert_eq!(scalar, auto, "All backends produce identical results");
}

Run the example:

cargo run --example hash_demo

Benchmarks Overview

This chapter presents comprehensive benchmark results for Trueno across different backends and workload sizes.

Latest Benchmark Results

Date: 2025-11-18 Platform: x86_64 Linux (AVX2-capable) Compiler: rustc 1.83 (release mode, opt-level=3, LTO=true) Tool: Criterion.rs (statistical benchmarking)

Executive Summary

Trueno's SIMD and GPU backends deliver 2-8x speedups for most operations, with exceptional performance on reduction and compute-intensive operations.

Key Findings

• Average speedup: 178.5% across all operations
• Best speedup: 8.8x (tanh activation, AVX2, 100 elements)
• Operations meeting ≥10% target: 66.7%
• Reduction operations: 200-400% speedup (dot, sum, max)
• Activation functions: 120-880% speedup (relu, tanh)
• Element-wise ops: 3-115% speedup (varies by operation and size)

Benchmark Results by Operation

Reduction Operations (Excellent Performance)

Reduction operations show exceptional SIMD performance due to parallel accumulation:

Why reduction operations excel:

• Combines multiple operations in SIMD lanes (4-8 parallel accumulations)
• No memory write bottleneck (single scalar result)
• Horizontal reduction is highly optimized
• Minimal overhead from setup/cleanup

Activation Functions (Good to Excellent Performance)

Activation functions benefit from SIMD, especially for compute-intensive operations:

Why activation functions perform well:

• Compute-intensive (tanh requires exp calculations)
• SIMD processes 4-8 elements in parallel
• No data dependencies between elements
• AVX2 benefits from wider registers (8 f32 vs 4 for SSE2)

Element-Wise Operations (Mixed Performance)

Element-wise operations show variable performance, often limited by memory bandwidth:

Why element-wise ops show limited speedups:

• Memory bandwidth bottleneck: Simple operations (add, mul) are memory-bound, not compute-bound
• Cache effects: Small workloads fit in L1 cache, scalar loop is efficient
• Large workloads: Both scalar and SIMD become memory-bound
• Overhead: SIMD setup/cleanup costs hurt small workloads (<1000 elements)

Performance by Backend

SSE2 (128-bit SIMD)

Availability: Guaranteed on all x86_64 CPUs Register width: 128 bits (4 × f32 or 2 × f64) Typical speedup: 2-4x for reduction ops, 1.05-1.15x for element-wise

Best operations:

• ✅ Reduction (dot, sum, max): 3-4.5x
• ✅ Activation functions (tanh, relu): 1.2-7.4x
• ⚠️ Element-wise (add, mul): 1.03-1.10x

Limitations:

• Limited to 4-way parallelism
• Some operations (div, sigmoid) show regressions
• Memory bandwidth limited for large workloads

AVX2 (256-bit SIMD)

Availability: Intel Haswell+ (2013+), AMD Zen+ (2018+) Register width: 256 bits (8 × f32 or 4 × f64) Typical speedup: 4-8x for reduction ops, 1.10-1.15x for element-wise

Best operations:

• ✅ Activation functions (tanh): 8.8x
• ✅ Scalar operations (scale): 1.15x
• ✅ Reduction (expected 2x over SSE2, not yet benchmarked)

Advantages over SSE2:

• 2x wider registers (8 vs 4 elements)
• FMA (fused multiply-add) instructions
• Better memory bandwidth utilization

GPU (WebGPU via wgpu)

Availability: Systems with Vulkan/Metal/DX12 support Typical speedup: 16-81x for large matrix operations (>500×500)

IMPORTANT: Empirical RTX 4090 benchmarking revealed that GPU has 3.5ms fixed transfer overhead, making it slower than SIMD for vector operations at ALL sizes.

GPU Performance Summary (2025-11-23, RTX 4090):

• ✅ Matrix multiplication: 81x speedup on 1000×1000
• ❌ Vector operations: 2000x+ slower than SIMD due to transfer overhead
• 🎯 Recommendation: GPU only for matrix ops >500×500, otherwise use SIMD

Current Thresholds:

GPU Transfer Overhead: ~3.5ms per operation for CPU↔GPU↔CPU transfer

See GPU Performance for detailed RTX 4090 benchmark results and analysis.

Performance by Workload Size

Small (100 elements)

Recommended backend: SSE2 or Scalar SIMD benefit: 5-10% for most ops, 120-650% for activation/reduction

At small sizes, SIMD overhead (setup, remainder handling) can exceed benefits for simple operations.

Medium (1K-10K elements)

Recommended backend: SSE2/AVX2 SIMD benefit: 3-440% depending on operation

Sweet spot for SIMD: large enough to amortize overhead, small enough to avoid memory bottlenecks.

Large (100K+ elements)

Recommended backend: GPU (if available), otherwise AVX2 SIMD benefit: 0-400% (memory-bound for simple ops, good for reductions)

At large sizes:

• Element-wise ops become memory-bound
• Reduction ops still benefit from SIMD
• GPU provides best performance if transfer overhead is justified

Benchmark Methodology

Tool: Criterion.rs

All benchmarks use Criterion.rs for statistical rigor:

• Samples: 100 per benchmark
• Warmup: 3 seconds
• Measurement: 5 seconds
• Outlier detection: Automated
• Statistical analysis: Mean, median, standard deviation

Test Data

• Sequential floats: (i as f32) * 0.5
• Workload sizes: 100, 1000, 10000, 100000 elements
• Backend comparison: Scalar vs SSE2 vs AVX2 vs GPU

Environment

• CPU: x86_64 with AVX2 support
• RAM: 16GB+ (prevents swapping)
• Compiler flags: -C opt-level=3 -C lto=true -C codegen-units=1
• CPU affinity: Pinned to single core (reduces variance)
• Background processes: Minimized

Quality Standards

Every benchmark must meet these criteria:

1. Coefficient of Variation (CV) < 5% - Consistent results across runs
2. No regressions >5% - SIMD should not be slower than scalar
3. Statistical significance - 100+ samples for reliable mean/median
4. Baseline comparison - Always compare against scalar implementation

Interpreting Results

Speedup calculation: (scalar_time / simd_time)

Reproducing Benchmarks

Run all benchmarks:

cargo bench --bench vector_ops

Run specific operation:

cargo bench --bench vector_ops -- dot

Generate HTML report:

cargo bench --bench vector_ops
open target/criterion/report/index.html

Compare against baseline:

# Save current results as baseline
cargo bench -- --save-baseline main

# Make changes, then compare
cargo bench -- --baseline main

Next Steps

• SIMD Performance - Deep dive into SIMD optimizations
• GPU Performance - GPU benchmarks and transfer overhead
• Optimization Guide - How to improve performance
• Profiling - Using perf, flamegraphs, and vtune

SIMD Performance Analysis

Date: 2025-11-18 System: x86_64 Linux (AVX2-capable) Benchmark Tool: Criterion.rs

This chapter provides a deep dive into Trueno's SIMD performance characteristics, analyzing when SIMD provides speedups and when it doesn't.

Executive Summary

Comprehensive benchmarking reveals mixed results across operations. While some operations show excellent SIMD speedups (tanh: 6.5-8.8x), many element-wise operations show minimal or negative speedups, especially for SSE2.

Key Findings

1. Activation functions (relu, tanh): Good to excellent SIMD speedups (1.2-8.8x)
2. Reduction operations (dot, sum, max): Excellent SIMD speedups (3-4.5x)
3. Element-wise operations (add, sub, div, fma): Minimal or negative SIMD benefit
4. SSE2 backend: Frequently slower than scalar for simple operations
5. Small workloads (<1000 elements): SIMD overhead often exceeds benefit

Performance by Operation Category

Excellent SIMD Performance (>5x speedup)

Why tanh excels:

• Compute-intensive operation (requires exp calculations)
• SIMD processes 4-8 exponentials in parallel
• No memory bottleneck (compute dominates)
• AVX2's wider registers (8 vs 4 elements) provide 2x improvement over SSE2

Good SIMD Performance (1.1-2x speedup)

Poor SIMD Performance (<1.1x or negative)

Root Cause Analysis

1. Memory Bandwidth Bottleneck

For simple operations, memory access dominates compute time. SIMD can't help with RAM speed.

2. SIMD Overhead for Small Workloads

Fixed ~20-50ns overhead per operation from setup, alignment checks, and remainder handling.

3. Suboptimal Implementations

Some operations (div, sigmoid) show regressions requiring investigation.

Next Steps

• Fix SSE2 div, sigmoid, fma, sub implementations
• Implement adaptive backend selection
• Benchmark against NumPy/PyTorch

Related Chapters

• Benchmarks Overview
• Optimization Guide
• Profiling

GPU Performance

This chapter presents empirical GPU performance findings from benchmarking on NVIDIA RTX 4090, documenting when GPU acceleration provides value versus SIMD.

Executive Summary

Date: 2025-11-23 Hardware: NVIDIA GeForce RTX 4090 (24GB VRAM) Driver: 570.195.03 Platform: Linux 6.8.0-87-generic Software: Trueno v0.7.0, wgpu v27.0.1

Key Findings

• ✅ GPU wins for matrix operations: 81x speedup on 1000×1000 matrix multiplication
• ❌ GPU fails for vector operations: 2000x+ slower than SIMD due to 3.5ms fixed overhead
• 🚀 SIMD vastly superior for vector ops: Zero transfer overhead, 200-400% speedup
• 💡 Hybrid approach recommended: Use SIMD by default, GPU only for matmul >500×500

GPU Transfer Overhead

Fixed Overhead Breakdown

Empirically measured per-operation costs:

Implications for Different Workload Sizes

Rule of thumb: GPU only becomes competitive when compute time >> 3.5ms.

Matrix Multiplication (GPU Excels)

Matrix multiplication has O(n³) complexity, which overwhelms the fixed 3.5ms overhead at large scales.

Benchmark Results

Why GPU Wins for Matrix Multiplication

Compute complexity dominates transfer cost:

• 100×100: 1M operations → 531µs scalar → GPU overhead too high
• 500×500: 125M operations → 77ms scalar → GPU wins at 4.6ms
• 1000×1000: 1B operations → 639ms scalar → GPU wins at 7.8ms

Threshold: GPU becomes competitive at >500×500 (250,000 elements).

Vector Operations (GPU Fails)

Simple vector operations are dominated by the 3.5ms fixed transfer overhead.

Vector Addition Results

Dot Product Results

Key finding: Even at 1M elements, GPU is still 62x slower than scalar due to transfer overhead. Reduction overhead compounds the problem.

Activation Functions

Activation functions are more compute-intensive than simple vector operations, but still suffer from transfer overhead.

ReLU (Simple Operation)

Sigmoid (Transcendental)

GELU (Very Compute-Heavy)

Key finding: Even compute-heavy operations like GELU and sigmoid are slower on GPU due to transfer overhead. At 1M elements, GPU barely reaches parity with scalar.

Softmax (Multi-Pass Algorithm)

Why softmax is even worse: Multi-pass algorithms require 3 GPU dispatches (max, exp, sum), compounding transfer overhead to ~10ms base cost.

SIMD vs GPU Comparison

Golden traces from Renacer v0.6.2 show SIMD baseline performance:

SIMD Performance (SSE2)

From golden_traces/performance_demo_summary.txt:

Head-to-Head Comparison

Key Insights

• ✅ SIMD dominates for vector operations at ALL sizes due to zero overhead
• ✅ GPU wins for matrix operations (O(n³) complexity) at large scales
• 💡 Hybrid approach: Use SIMD by default, GPU only for matmul >500×500

Current GPU Thresholds in Trueno

Based on empirical findings, Trueno uses these thresholds:

// src/vector.rs:1316
const GPU_THRESHOLD: usize = usize::MAX; // GPU DISABLED - 2-800x slower

// src/matrix.rs:268
const GPU_THRESHOLD: usize = 500; // Empirical: 2x at 500×500, 9.6x at 1000×1000

Rationale:

• Vector operations: Transfer overhead will always dominate → GPU disabled
• Matrix operations: O(n³) complexity amortizes overhead → GPU at 500×500

When to Use GPU

Use GPU when all of these conditions are met:

1. Operation complexity: O(n²) or higher (matrix multiplication, convolution)
2. Data size: >500×500 elements for matrix ops
3. Compute time: Operation takes >10ms on CPU
4. Batch processing: Multiple operations can be batched (future v2.0 API)

GPU is NOT recommended for:

• ❌ Vector operations (add, mul, dot, reduce) - use SIMD
• ❌ Activation functions (relu, sigmoid, tanh) - use SIMD
• ❌ Small matrices (<500×500) - overhead dominates
• ❌ Single operations - transfer overhead too high

GPU Tiled Reduction ✅ (v0.10.1)

Status: Validated on Metal (AMD Radeon Pro W5700X, Mac Pro 7,1)

The tiled reduction shader provides efficient GPU-based sum, max, and min operations using 16x16 workgroup tiles with two-phase reduction.

Metal Benchmark Results (2026-01-03)

Key Findings

• Consistent ~150 Melem/s throughput across all sizes on GPU
• ~8ms baseline overhead from CPU→GPU transfer
• CPU is 7-37x faster for standalone reductions (expected for O(n) ops)
• GPU wins for O(n³) operations like matmul, but loses for O(n) reductions

When GPU Tiled Reduction is Optimal

✅ Use GPU reduction when:

• Data is already resident on GPU (no transfer cost)
• Reduction is part of larger GPU compute pipeline
• Latency hiding in async GPU workloads

❌ Prefer SIMD when:

• Data starts on CPU (transfer overhead dominates)
• Standalone reduction operation
• Low-latency required

Metal Buffer Limits

CUDA PTX Validation ✅ (v0.10.1)

Status: Validated on NVIDIA GeForce RTX 4090 (Ada Lovelace, sm_89)

The trueno-gpu PTX code generation has been validated on real CUDA hardware, confirming JIT compilation and execution correctness.

RTX 4090 Validation Results (2026-01-03)

Kernel Generation Throughput

68,015 kernels/sec measured via bench_kernel_gen example.

Execution Verification

Simple Attention CUDA kernel verified with numerical accuracy:

• GPU execution: 134µs (16x16 sequence)
• Max difference: 2.98e-8 (vs CPU reference)
• Status: PASS

PTX Features Validated

• ✅ FMA fusion (mul+add → fma.rn.f32)
• ✅ F16 conversion (cvt.rn.f16.f32)
• ✅ Shared memory (smem with .align)
• ✅ WMMA Tensor Core ops
• ✅ Q4K quantization (4-bit dequantize)
• ✅ Tree reduction patterns
• ✅ Predicated execution (@%p bra)

Running CUDA Examples

# CUDA monitoring (device info, memory stats)
cargo run --example cuda_monitor --features cuda --release

# PTX generation benchmarks
cargo run --example bench_kernel_gen --features cuda --release

# Simple attention execution
cargo run --example simple_attention_cuda --features cuda --release

# Quantized GEMM PTX
cargo run --example q4k_gemm --features cuda --release

Example Usage

use trueno::backends::gpu::GpuBackend;

fn main() -> Result<(), String> {
    let mut gpu = GpuBackend::new();

    // Create 1000x1000 matrix
    let data: Vec<f32> = vec![1.0; 1_000_000];

    // GPU tiled sum reduction
    let sum = gpu.tiled_sum_2d_gpu(&data, 1000, 1000)?;
    println!("Sum: {}", sum);  // 1000000.0

    // GPU tiled max/min
    let max = gpu.tiled_max_2d_gpu(&data, 1000, 1000)?;
    let min = gpu.tiled_min_2d_gpu(&data, 1000, 1000)?;

    Ok(())
}

# Run the demonstration
cargo run --example gpu_tiled_reduction --features gpu --release

Benchmark Execution

# Run tiled reduction benchmarks
cargo bench --features gpu --bench gpu_reduction

Async Batch API ✅ (v0.3.0 - AVAILABLE NOW)

Status: Fully implemented and tested (previously documented as "Future v2.0")

The async batch API solves the transfer overhead problem by queuing multiple operations and executing them in a single batch, amortizing the 3.5ms overhead across all operations.

Transfer Overhead Reduction

Traditional Synchronous API (current default):

// ❌ 3 operations = 3 × 3.5ms = 10.5ms overhead
let a = gpu.vec_add(&input1, &input2)?;  // Upload → Compute → Download
let b = gpu.scale(&a, 2.0)?;             // Upload → Compute → Download
let c = gpu.relu(&b)?;                   // Upload → Compute → Download
// Total: 6 GPU transfers (3 uploads + 3 downloads)

Async Batch API (recommended for chained operations):

use trueno::backends::gpu::{GpuDevice, GpuCommandBatch};

// ✅ 3 operations = 1 × 3.5ms = 3.5ms overhead
let device = GpuDevice::new()?;
let mut batch = GpuCommandBatch::new(device);

// Queue operations (no GPU execution yet!)
let input = batch.upload(&[1.0, 2.0, -3.0, 4.0]);
let a = batch.add(input, other);
let b = batch.scale(a, 2.0);
let c = batch.relu(b);

// Execute entire batch in one GPU round-trip
batch.execute().await?;

// Read final result
let result = batch.read(c).await?;
// Total: 2 GPU transfers (1 upload + 1 download)

Performance Benefits

When to Use Batch API

✅ Use batch API when:

• Chaining multiple GPU operations (>2 ops)
• Processing large workloads where GPU is beneficial (matmul >500×500)
• Amortizing transfer overhead is critical

❌ Stick with traditional API when:

• Single operation only
• Interactive/real-time workloads requiring immediate results
• Workloads small enough that SIMD is faster anyway

Complete Example

See examples/gpu_batch_demo.rs for three comprehensive demonstrations:

1. Single Operation - Baseline batch API usage
2. Batched Operations - ReLU → Scale → Add pipeline
3. ML Pipeline - y = ReLU(x * W + b) simulation

# Run the demonstration
cargo run --example gpu_batch_demo --features gpu --release

Implementation Details

• Location: src/backends/gpu/batch.rs (1,008 lines)
• Tests: 8 comprehensive tests (all passing)
• Operations: relu, scale, add, mul, dot
• API: Fully async with tokio integration
• Safety: Type-safe buffer IDs prevent invalid operations

Future Enhancements (v0.4.0+)

While the batch API is complete, future improvements may include:

• Automatic optimization: Detect operation chains and auto-batch
• More operations: Expand beyond current 5 operations (relu, scale, add, mul, dot)
• Graph optimization: Reorder operations for maximum efficiency
• Multi-GPU: Distribute batches across multiple GPUs
• Persistent buffers: Reuse buffers across multiple batch executions

Hardware Details

GPU: NVIDIA GeForce RTX 4090
├─ Architecture: Ada Lovelace
├─ CUDA Cores: 16,384
├─ Memory: 24GB GDDR6X
├─ Memory Bandwidth: 1,008 GB/s
├─ Boost Clock: 2.52 GHz
└─ TDP: 450W

Driver: 570.195.03
Platform: Linux 6.8.0-87-generic (x86_64)

Validation and Testing

Quality Gates

• ✅ All 13 GPU operations benchmarked
• ✅ 4 size ranges tested per operation
• ✅ Statistical significance (10 samples, CV <5%)
• ✅ Comparison against scalar baseline
• ✅ Clippy: Zero warnings
• ✅ Coverage: 90.40% (≥90% threshold)
• ✅ GPU initialization verified
• ✅ Correctness tests pass

Golden Trace Integration

Performance budgets established via renacer.toml:

[performance.budgets]
# SIMD operations should complete in <2ms with <200 syscalls
backend_detection = { max_time_ms = 2.0, max_syscalls = 200 }
matrix_operations = { max_time_ms = 2.0, max_syscalls = 200 }
activation_functions = { max_time_ms = 2.0, max_syscalls = 200 }

Validation tests in tests/golden_trace_validation.rs ensure SIMD performance doesn't regress.

Recommendations

Immediate Actions

1. Use SIMD by default for all vector operations
2. Reserve GPU for matrix operations >500×500
3. Document transfer overhead prominently in API docs
4. Educate users that GPU is not always faster

Future Enhancements (v2.0)

1. Async batch API to amortize transfer overhead
2. Persistent GPU buffers for frequently-used data
3. Hybrid CPU/GPU scheduling with overlap
4. Profile-guided optimization for dynamic thresholds

References

• Full benchmark report: docs/gpu-benchmark-report-2025-11-23.md
• Golden traces: golden_traces/ directory
• Golden trace analysis: golden_traces/ANALYSIS.md
• SIMD performance: golden_traces/performance_demo_summary.txt
• Renacer configuration: renacer.toml
• GPU bug fix: Commit b5ca0af (missing device.poll() in wgpu v27)

WebGPU for WASM (v0.7.3)

Trueno v0.7.3 introduces the gpu-wasm feature enabling GPU compute in browsers via WebGPU.

Feature Flag

[target.'cfg(target_arch = "wasm32")'.dependencies]
trueno = { version = "0.7.3", features = ["gpu-wasm"] }

Platform Differences

Async-First Design

All GPU operations now have async variants (*_async) that work on both native and WASM:

// Works on all platforms
let device = GpuDevice::new_async().await?;
device.matmul_async(&a, &b, &mut result, m, k, n).await?;
device.relu_async(&input, &mut output).await?;

Runtime Detection

use trueno::backends::gpu::runtime;

if runtime::sync_available() {
    // Native: can use sync APIs
    let device = GpuDevice::new()?;
} else {
    // WASM: must use async
    let device = GpuDevice::new_async().await?;
}

Real-World Example: trueno-viz

trueno-viz demonstrates browser-based GPU compute with Trueno:

• WebGPU-accelerated matrix operations
• WASM-compiled Rust for client-side processing
• Interactive visualizations with GPU compute

See GPU Backend Architecture for complete WebGPU documentation.

Next Steps

• Backend Comparison - Detailed SIMD vs GPU trade-offs
• Benchmarks Overview - Complete benchmark methodology
• Optimization Guide - How to choose the right backend
• Profiling - Using Renacer for performance analysis

Optimization Guide

This chapter covers performance optimization techniques used in Trueno, with a focus on PTX code generation and kernel emission.

PTX Emission Optimization

The PTX code generator has been optimized to minimize memory allocations during kernel generation, achieving a 20.9% improvement in emission performance.

Key Optimizations

1. Pre-allocated String Capacity

Instead of growing the output string dynamically, we estimate the final size:

// Pre-allocate with estimated size: ~100 bytes per instruction + header overhead
let estimated_size = 512 + self.instructions.len() * 100;
let mut ptx = String::with_capacity(estimated_size);

This eliminates repeated reallocations as the PTX output grows.

2. Zero-Allocation Instruction Emission

The write_instruction() function writes directly to the output buffer instead of returning intermediate Strings:

// Before (allocates per instruction):
for instr in &self.instructions {
    ptx.push_str(&emit_instruction(instr));  // allocates String
}

// After (zero allocation):
for instr in &self.instructions {
    write_instruction(instr, &mut ptx);  // writes directly
}

3. Display Implementation for VirtualReg

Added Display trait implementation for zero-allocation register formatting:

impl fmt::Display for VirtualReg {
    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
        write!(f, "{}{}", self.ty.register_prefix(), self.id)
    }
}

// Now can use write! macro directly:
write!(out, "{}", vreg);  // No intermediate allocation

Performance Results

Kernel Generation Performance:

Throughput: 68,316 kernels/sec

Benchmarking

Run the kernel generation benchmark:

cargo run -p trueno-gpu --release --example bench_kernel_gen

General Optimization Principles

1. Minimize Allocations in Hot Paths

• Pre-allocate collections with known sizes
• Use &str instead of String where possible
• Use write! to write directly to buffers

2. Use Static Strings

Many PTX components are static and can use &'static str:

pub const fn to_ptx_string(self) -> &'static str {
    match self {
        Self::F32 => ".f32",
        Self::U32 => ".u32",
        // ...
    }
}

3. Avoid Intermediate Allocations

Instead of:

fn emit() -> String {
    format!("{}{}", prefix, suffix)  // allocates
}
out.push_str(&emit());  // pushes

Use:

fn write_to(out: &mut String) {
    out.push_str(prefix);
    out.push_str(suffix);  // no intermediate allocation
}

SIMD Backend Optimization

For SIMD backend optimizations, see:

• SIMD Performance
• Backend Comparison

GPU Performance

For GPU-specific optimizations, see:

• GPU Performance
• PTX Optimization Passes

Profiling

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Golden Trace Validation

Status: Integrated (v0.7.0) Tool: Renacer 0.6.2 Purpose: Performance regression detection via syscall tracing

Overview

Golden trace validation uses Renacer (pure Rust syscall tracer) to capture canonical execution traces for Trueno compute examples. These traces serve as performance baselines, enabling:

1. Regression Detection: Detect performance degradation via syscall count/latency budgets
2. PCIe Bottleneck Analysis: Identify inefficient GPU memory transfers
3. Build-Time Assertions: Enforce performance contracts in CI/CD
4. Root Cause Analysis: Correlate syscalls to Rust source code

Quick Start

1. Install Renacer

cargo install renacer --version 0.6.2

2. Capture Golden Traces

cd /path/to/trueno
./scripts/capture_golden_traces.sh

Output:

✅ Captured: golden_traces/backend_detection.json (0.73ms, 87 syscalls)
✅ Captured: golden_traces/matrix_operations.json (1.56ms, 168 syscalls)
✅ Captured: golden_traces/activation_functions.json (1.30ms, 159 syscalls)
✅ Captured: golden_traces/performance_demo.json (1.51ms, 138 syscalls)
✅ Captured: golden_traces/ml_similarity.json (0.82ms, 109 syscalls)

3. View Trace Summary

cat golden_traces/backend_detection_summary.txt

Example Output:

Syscall Summary:
write:     23 calls  (0.15ms total)
mmap:      13 calls  (0.21ms total)
mprotect:   6 calls  (0.08ms total)
munmap:     5 calls  (0.04ms total)
...
TOTAL:     87 calls  (0.73ms total)

Traced Operations

1. Backend Detection (backend_detection)

Purpose: Validate SIMD backend auto-selection (AVX-512 → AVX2 → SSE2 → Scalar)

Performance Budget:

• Runtime: <10ms
• Syscalls: <100
• Memory: <10MB

Actual Performance: ✅

• Runtime: 0.73ms (13× faster than budget)
• Syscalls: 87
• Top syscalls: write (23), mmap (13), mprotect (6)

Trace Capture:

renacer --format json -- ./target/release/examples/backend_detection > backend_detection.json

2. Matrix Operations (matrix_operations)

Purpose: Measure SIMD-accelerated matrix multiply and transpose overhead

Performance Budget:

• Runtime: <20ms
• Syscalls: <200

Actual Performance: ✅

• Runtime: 1.56ms (15× faster)
• Syscalls: 168

Key Insight: SIMD operations are compute-bound (minimal syscalls)

3. ML Activation Functions (activation_functions)

Purpose: Measure SIMD-accelerated activations (ReLU, sigmoid, tanh, GELU, swish)

Performance Budget:

• Runtime: <20ms
• Syscalls: <200

Actual Performance: ✅

• Runtime: 1.30ms
• Syscalls: 159

4. Performance Demo (performance_demo)

Purpose: Comprehensive benchmark across vector ops, matrix ops, and backend comparisons

Performance Budget:

• Runtime: <50ms
• Syscalls: <300

Actual Performance: ✅

• Runtime: 1.51ms (33× faster)
• Syscalls: 138

5. ML Similarity (ml_similarity)

Purpose: Measure vector similarity operations (cosine, Euclidean, Manhattan)

Performance Budget:

• Runtime: <20ms
• Syscalls: <200

Actual Performance: ✅ FASTEST

• Runtime: 0.82ms
• Syscalls: 109

Why Fast: Heavily optimized SIMD dot product, minimal allocations

Performance Assertions (renacer.toml)

Critical Path Latency

[[assertion]]
name = "example_startup_latency"
type = "critical_path"
max_duration_ms = 100
fail_on_violation = true
enabled = true

Rationale: Compute examples should complete quickly. 100ms allows for SIMD initialization and small-scale computations.

Violation Symptoms:

• SIMD overhead issues
• Unexpected I/O operations
• Debug build instead of release

Syscall Budget

[[assertion]]
name = "max_syscall_budget"
type = "span_count"
max_spans = 500
fail_on_violation = true
enabled = true

Rationale: SIMD operations are CPU-bound with minimal syscalls (mostly mmap for allocation). Budget prevents I/O regressions.

Typical Syscalls:

• write: stdout output (20-50 calls)
• mmap: vector/matrix allocation (10-30 calls)
• mprotect: memory permissions (5-10 calls)

Memory Allocation Budget

[[assertion]]
name = "memory_allocation_budget"
type = "memory_usage"
max_bytes = 104857600  # 100MB
tracking_mode = "allocations"
fail_on_violation = true
enabled = true

Rationale: Small examples should have minimal memory footprint. 100MB accommodates matrix allocations and SIMD buffers.

PCIe Bottleneck Detection

[[assertion]]
name = "detect_pcie_bottleneck"
type = "anti_pattern"
pattern = "PCIeBottleneck"
threshold = 0.7
fail_on_violation = false  # Warning only
enabled = true

Pattern Detected: GPU transfer time >> compute time

Symptoms:

• Many write/read syscalls to /dev/nvidia*
• High ioctl frequency for GPU operations
• Transfer overhead dominates (>70% of total time)

Example Warning:

⚠️  PCIe Bottleneck detected (confidence: 85%)
   GPU transfers: 45ms (90% of total time)
   Compute time:   5ms (10% of total time)
   Recommendation: Batch operations, keep data on GPU

Solution:

• Batch multiple operations
• Keep intermediate results on GPU
• Use larger workloads (amortize transfer costs)
• Trueno automatically disables GPU for small ops (v0.2.1+)

CI/CD Integration

GitHub Actions Workflow

Add to .github/workflows/ci.yml:

name: Golden Trace Validation

on: [push, pull_request]

jobs:
  validate-traces:
    runs-on: ubuntu-latest
    steps:
      - uses: actions/checkout@v3

      - name: Install Rust
        uses: dtolnay/rust-toolchain@stable

      - name: Install Renacer
        run: cargo install renacer --version 0.6.2

      - name: Build Examples (Release)
        run: cargo build --release --examples

      - name: Capture Golden Traces
        run: ./scripts/capture_golden_traces.sh

      - name: Run Performance Assertions
        run: |
          renacer --assert renacer.toml -- ./target/release/examples/backend_detection
          renacer --assert renacer.toml -- ./target/release/examples/matrix_operations
          renacer --assert renacer.toml -- ./target/release/examples/activation_functions

      - name: Upload Traces
        uses: actions/upload-artifact@v3
        with:
          name: golden-traces
          path: golden_traces/

CI Failure Example:

❌ Assertion 'example_startup_latency' FAILED
   Actual: 125ms
   Budget: 100ms
   Regression: +25%

⚠️  Build BLOCKED. SIMD overhead regression detected.

Advanced Usage

1. Source Code Correlation

Map syscalls to Rust source code:

renacer -s -- ./target/release/examples/backend_detection

Output:

write(1, "Backend: AVX2\n", 14) = 14  [src/lib.rs:245]
mmap(...) = 0x7f... [src/vector.rs:89]

Use Case: Identify which code paths trigger GPU initialization or excessive allocations.

2. OpenTelemetry Export

Visualize traces in Jaeger:

# Start Jaeger
docker run -d --name jaeger \
  -e COLLECTOR_OTLP_ENABLED=true \
  -p 4317:4317 \
  -p 16686:16686 \
  jaegertracing/all-in-one:latest

# Export trace
renacer --otlp http://localhost:4317 -- ./target/release/examples/performance_demo

# View in Jaeger UI
open http://localhost:16686

Use Case: Visualize syscall timelines for multi-operation pipelines.

3. Regression Analysis

Compare current execution against baseline:

# Capture current trace
renacer --format json -- ./target/release/examples/backend_detection > current.json

# Compare with golden
diff <(jq '.syscalls | length' golden_traces/backend_detection.json) \
     <(jq '.syscalls | length' current.json)

Expected: No difference in syscall count (±5% tolerance)

4. GPU Workload Analysis

For GPU-enabled builds:

# Build with GPU feature
cargo build --release --examples --features gpu

# Trace GPU example
renacer --format json -- ./target/release/examples/gpu_test > gpu_trace.json

# Filter GPU device operations
jq '.syscalls[] | select(.name == "ioctl" or .name == "write")' gpu_trace.json

Expected: GPU operations show as ioctl calls to /dev/nvidia0

Red Flag: If transfer syscalls dominate, GPU is inefficient for this workload size.

Toyota Way Principles

Andon (Stop the Line)

Implementation: Build-time assertions fail CI on regression

[[assertion]]
fail_on_violation = true  # ← Andon: Stop the CI pipeline

Poka-Yoke (Error-Proofing)

Implementation: Golden traces make expected patterns explicit

# Automated comparison prevents silent regressions
diff golden_traces/backend_detection.json new_trace.json

Jidoka (Autonomation)

Implementation: Automated quality enforcement without manual intervention

# GitHub Actions runs golden trace validation automatically
- name: Validate Performance
  run: ./scripts/capture_golden_traces.sh

Troubleshooting

Issue: Capture script fails with "Binary not found"

Solution:

cargo build --release --examples
./scripts/capture_golden_traces.sh

Issue: Performance regression detected

Diagnosis:

renacer --summary --timing -- ./target/release/examples/backend_detection
cat golden_traces/backend_detection_summary.txt

Common Causes:

• Debug build instead of release (cargo build --release)
• SIMD features disabled (check RUSTFLAGS)
• New dependencies (increase initialization overhead)

Issue: Syscall count regression

Diagnosis:

renacer -- ./target/release/examples/backend_detection > current_trace.txt
diff current_trace.txt golden_traces/backend_detection_summary.txt

Common Causes:

• New logging initialization (env_logger, tracing)
• Allocator changes (jemalloc → system allocator)
• Library updates (different I/O patterns)

Performance Baselines (v0.7.0)

Platform: x86_64 Linux, AVX2 backend, Release build

References

• Renacer GitHub
• Integration Report
• SIMD Performance Analysis
• Golden Trace Analysis

Last Updated: 2025-11-23 Renacer Version: 0.6.2 Trueno Version: 0.7.0

Targets

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Backend Comparison

This chapter compares Trueno's three execution backends (Scalar, SIMD, GPU) across different operation types and workload sizes, providing guidance on when to use each.

Backend Overview

Decision Matrix

Use this table to choose the optimal backend for your workload:

Scalar Backend

Characteristics

• Pros:

  • Zero overhead
  • Simple, maintainable code
  • Predictable performance
  • Works everywhere

• Cons:

  • No parallelism
  • Slowest for compute-heavy operations

When to Use Scalar

• Reference implementation for correctness testing
• Platforms without SIMD support (rare)
• Debugging (simpler code paths)
• Very small workloads (<100 elements) where SIMD overhead dominates

Performance

SIMD Backend

Characteristics

• Pros:

  • Zero transfer overhead
  • 2-4x speedup for most operations
  • Low latency (<10µs for typical ops)
  • Works on all modern CPUs

• Cons:

  • Limited parallelism (4-8 elements)
  • Complex implementation
  • Platform-specific code

SIMD Instruction Sets

SIMD Performance (SSE2)

From golden traces (golden_traces/performance_demo_summary.txt):

Why SIMD Excels

Zero overhead architecture:

• No data transfer (operates directly on CPU cache)
• No synchronization (single-threaded execution)
• Immediate execution (no queuing or dispatch)

Optimal for:

• ✅ Reduction operations (dot, sum, max): Parallel accumulation
• ✅ Compute-intensive ops (tanh, sigmoid): Amortizes instruction overhead
• ⚠️ Memory-bound ops (add, mul): Limited by RAM bandwidth, not compute

GPU Backend

Characteristics

• Pros:

  • Massive parallelism (thousands of cores)
  • 80x+ speedup for large matrix operations
  • Excellent for O(n³) algorithms

• Cons:

  • 3.5ms fixed overhead per operation
  • Requires PCIe transfer (CPU↔GPU)
  • Only beneficial for large workloads
  • Not always available

GPU Transfer Overhead

Critical limitation: Every GPU operation incurs ~3.5ms fixed cost:

GPU Performance (RTX 4090)

Vector operations (❌ GPU fails):

Matrix operations (✅ GPU wins):

Why GPU Fails for Vector Operations

Transfer overhead dominates:

• 10K vector add: 1.1µs compute vs 3500µs transfer → 3182x overhead
• 1M vector add: 96.5µs compute vs 3500µs transfer → 36x overhead

Even compute-heavy ops suffer:

• 1M sigmoid: 3.18ms compute vs 3.5ms transfer → Barely competitive

Why GPU Wins for Matrix Operations

O(n³) complexity overwhelms transfer cost:

• 500×500 matmul: 125M ops → 77ms scalar → GPU wins at 4.6ms (13x amortization)
• 1000×1000 matmul: 1B ops → 639ms scalar → GPU wins at 7.8ms (81x amortization)

GPU becomes competitive when: compute_time_scalar > 10 × transfer_overhead

For matrix multiplication:

• 500×500: 77ms compute >> 3.5ms transfer → GPU wins
• 100×100: 531µs compute << 3.5ms transfer → GPU loses

Backend Comparison by Operation Type

Element-Wise Operations (add, mul, scale)

Recommendation: Always use SIMD. Memory-bound, but SIMD has zero overhead.

Reduction Operations (dot, sum, max)

Recommendation: Always use SIMD. Excellent parallel accumulation, zero overhead.

Activation Functions (relu, sigmoid, tanh)

Recommendation: Always use SIMD, even for compute-heavy activations.

Matrix Multiplication

Recommendation: Use GPU for matrices >500×500, otherwise SIMD.

Threshold Guidelines

Current Trueno Thresholds

// Vector operations (src/vector.rs:1316)
const GPU_THRESHOLD: usize = usize::MAX; // GPU DISABLED

// Matrix operations (src/matrix.rs:268)
const GPU_THRESHOLD: usize = 500; // 500×500 minimum

Size-Based Recommendations

Operation Complexity Classes

Trueno categorizes operations by complexity:

pub enum OpComplexity {
    Low,    // Simple ops: add, mul (GPU disabled)
    Medium, // Moderate: dot, reduce (GPU at 100K+)
    High,   // Complex: matmul, conv2d (GPU at 500×500+)
}

Performance Validation

Golden Trace Baselines

Performance budgets in renacer.toml ensure SIMD doesn't regress:

[performance.budgets]
backend_detection = { max_time_ms = 2.0, max_syscalls = 200 }
matrix_operations = { max_time_ms = 2.0, max_syscalls = 200 }
activation_functions = { max_time_ms = 2.0, max_syscalls = 200 }

All SIMD operations must complete in <2ms with <200 syscalls.

Validation Tests

tests/golden_trace_validation.rs ensures:

• SIMD performance matches golden traces (±10%)
• No unexpected syscall patterns
• Runtime stays under budget

Future: Hybrid Scheduling (v2.0)

Current API forces a single backend per operation. Future hybrid scheduling will:

1. Profile operation characteristics at runtime
2. Dynamically select backend based on actual compute time
3. Batch GPU operations to amortize transfer overhead
4. Overlap CPU and GPU work for pipeline parallelism

Example future API:

let scheduler = HybridScheduler::new()
    .prefer_simd_threshold_ms(5.0)  // Use SIMD if op <5ms
    .gpu_batch_window_ms(10.0);     // Batch GPU ops within 10ms

scheduler.execute_pipeline(|pipe| {
    let a = pipe.add(&x, &y);       // SIMD (fast)
    let b = pipe.dot(&a, &z);       // SIMD (fast)
    let c = pipe.matmul(&b, &w);    // GPU (queued)
    let d = pipe.matmul(&c, &v);    // GPU (batched!)
    d
});

Recommendations Summary

For Vector Operations

1. Always use SIMD - Zero overhead, 2-4x speedup
2. Never use GPU - 2000x+ slower due to transfer overhead
3. Use scalar only for <100 elements or debugging

For Matrix Operations

1. Use SIMD for matrices <500×500
2. Use GPU for matrices ≥500×500 (16-81x speedup)
3. Consider batching multiple GPU operations in future

General Guidelines

• Latency-critical: Always SIMD (microsecond-scale)
• Throughput-critical: GPU for large batches, SIMD otherwise
• Portable: SIMD works everywhere (x86, ARM, WASM)
• Maximum performance: Profile and choose dynamically

References

• GPU Performance - Detailed GPU benchmarks (RTX 4090)
• SIMD Performance - SIMD optimization techniques
• Benchmarks Overview - Complete benchmark methodology
• Full report: docs/gpu-benchmark-report-2025-11-23.md
• Golden traces: golden_traces/ANALYSIS.md
• Configuration: renacer.toml

Philosophy

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Unsafe Code

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Safety Invariants

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Miri Validation

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Testing Correctness

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Backend Equivalence

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Vector Math

This chapter demonstrates Trueno's vector math capabilities using the quickstart and performance_demo examples.

Quick Start

Run the quickstart example to see all core vector operations:

cargo run --example quickstart

Basic Operations

use trueno::Vector;

// Create vectors
let a = Vector::from_slice(&[1.0, 2.0, 3.0, 4.0]);
let b = Vector::from_slice(&[5.0, 6.0, 7.0, 8.0]);

// Element-wise operations
let sum = a.add(&b)?;      // [6.0, 8.0, 10.0, 12.0]
let prod = a.mul(&b)?;     // [5.0, 12.0, 21.0, 32.0]

// Reductions
let dot = a.dot(&b)?;      // 70.0
let norm = a.norm_l2()?;   // 5.477...

// Statistical operations
let mean = a.mean()?;      // 2.5
let variance = a.variance()?;

Backend Selection

Trueno automatically selects the best available backend:

use trueno::{Vector, Backend};

// Auto backend (recommended)
let v = Vector::from_slice(&data);

// Force specific backend
let scalar = Vector::from_slice_with_backend(&data, Backend::Scalar);

Performance Comparison

Run the performance demo to see SIMD speedups:

cargo run --release --example performance_demo

Expected Results

Understanding the Results

Compute-intensive operations (dot product, sum, max) show significant speedups because SIMD can process 8 f32 values simultaneously.

Memory-bound operations (add, mul) show modest speedups because performance is limited by memory bandwidth, not computation.

ML Similarity Operations

Run the similarity example:

cargo run --example ml_similarity

Cosine Similarity

use trueno::Vector;

let query = Vector::from_slice(&[0.5, 0.8, 0.2]);
let document = Vector::from_slice(&[0.6, 0.7, 0.3]);

// Compute cosine similarity
let norm_q = query.norm_l2()?;
let norm_d = document.norm_l2()?;
let dot = query.dot(&document)?;
let similarity = dot / (norm_q * norm_d);

k-NN Classification

// Compute Euclidean distances
let diff = query.sub(&sample)?;
let dist_sq = diff.dot(&diff)?;
let distance = dist_sq.sqrt();

Layer Normalization

use trueno::Vector;

let input = Vector::from_slice(&[1.0, 2.0, 3.0, 4.0, 5.0]);

// Compute mean and variance
let mean = input.mean()?;
let centered = input.sub_scalar(mean)?;
let var = centered.dot(&centered)? / input.len() as f32;
let std = (var + 1e-5).sqrt();

// Normalize
let normalized = centered.mul_scalar(1.0 / std)?;

See Also

• Performance Demo - Detailed benchmarks
• ML Similarity - ML-specific operations
• Backend Selection - How backends are chosen

Matrix Operations

This chapter demonstrates Trueno's matrix operations using the matrix_operations example.

Running the Example

cargo run --example matrix_operations

Basic Matrix Operations

Creating Matrices

use trueno::Matrix;

// Create from row-major data
let a = Matrix::from_vec(2, 3, vec![
    1.0, 2.0, 3.0,  // Row 0
    4.0, 5.0, 6.0,  // Row 1
])?;

// Identity matrix
let identity = Matrix::identity(3);

// Zero matrix
let zeros = Matrix::zeros(2, 3);

Matrix Multiplication

// C = A × B
let a = Matrix::from_vec(2, 3, vec![1.0, 2.0, 3.0, 4.0, 5.0, 6.0])?;
let b = Matrix::from_vec(3, 2, vec![7.0, 8.0, 9.0, 10.0, 11.0, 12.0])?;

let c = a.matmul(&b)?;  // Result: 2×2 matrix

Matrix-Vector Multiplication

use trueno::{Matrix, Vector};

let weights = Matrix::from_vec(3, 4, weight_data)?;
let input = Vector::from_slice(&[1.0, 2.0, 3.0, 4.0]);

let output = weights.matvec(&input)?;  // Result: Vector of length 3

Transpose

let a = Matrix::from_vec(2, 3, vec![1.0, 2.0, 3.0, 4.0, 5.0, 6.0])?;
let at = a.transpose();  // Result: 3×2 matrix

Neural Network Layers

Linear Layer (Dense)

fn linear_layer(
    input: &Vector,
    weights: &Matrix,
    bias: &Vector,
) -> Result<Vector, TruenoError> {
    let output = weights.matvec(input)?;
    output.add(bias)
}

Batch Processing

// Process multiple samples through the same layer
let samples = vec![
    Vector::from_slice(&[0.2, -0.3, 0.5]),
    Vector::from_slice(&[0.3, 0.0, 0.1]),
    Vector::from_slice(&[0.0, 0.3, 0.4]),
];

for sample in &samples {
    let output = weights.matvec(sample)?;
    println!("Output: {:?}", output.as_slice());
}

Mathematical Properties

The example verifies key mathematical properties:

Identity Property

// I × v = v
let identity = Matrix::identity(3);
let v = Vector::from_slice(&[1.0, 2.0, 3.0]);
let result = identity.matvec(&v)?;
assert_eq!(result.as_slice(), v.as_slice());

Transpose Property

// (A × v)^T = v^T × A^T
// This is verified in the example

Zero Property

// A × 0 = 0
let zeros = Vector::from_slice(&[0.0, 0.0, 0.0, 0.0]);
let result = weights.matvec(&zeros)?;
// All elements should be 0

Batched Matrix Multiplication

3D Tensors (Batch Processing)

Process multiple matrix multiplications in a single call:

use trueno::Matrix;

// Shape: [batch, m, k] @ [batch, k, n] -> [batch, m, n]
let batch = 4;
let m = 32;
let k = 64;
let n = 32;

let a_data: Vec<f32> = vec![0.1; batch * m * k];
let b_data: Vec<f32> = vec![0.2; batch * k * n];

let result = Matrix::batched_matmul(&a_data, &b_data, batch, m, k, n)?;

4D Tensors (Multi-Head Attention)

The exact pattern used in transformer attention (Q @ K^T and attn @ V):

// Simulate multi-head attention: Q @ K^T
// Shape: [batch, heads, seq, head_dim] @ [batch, heads, head_dim, seq]
let batch = 1;
let heads = 12;
let seq_len = 512;
let head_dim = 64;

let q_data: Vec<f32> = vec![0.0; batch * heads * seq_len * head_dim];
let kt_data: Vec<f32> = vec![0.0; batch * heads * head_dim * seq_len];

let attn_scores = Matrix::batched_matmul_4d(
    &q_data,
    &kt_data,
    batch,
    heads,
    seq_len,   // m
    head_dim,  // k
    seq_len,   // n
)?;
// Output: [batch, heads, seq_len, seq_len] attention scores

This is critical for transformer inference performance - each (batch, head) pair is processed independently using SIMD matmul, achieving ~50 GFLOPS vs ~0.1 GFLOPS for naive implementation.

Performance Considerations

Blocking for Cache Efficiency

Trueno uses blocked matrix multiplication for better cache utilization:

// Automatic blocking for large matrices
let large_a = Matrix::from_vec(1024, 1024, data_a)?;
let large_b = Matrix::from_vec(1024, 1024, data_b)?;
let c = large_a.matmul(&large_b)?;  // Uses tiled algorithm internally

SIMD Acceleration

Matrix operations automatically use SIMD when beneficial:

• AVX2: Process 8 f32 values per instruction
• AVX-512: Process 16 f32 values per instruction
• Automatic fallback to scalar for small matrices

GPU Acceleration

For large matrices, enable GPU acceleration:

cargo run --release --features gpu --example matrix_operations

The GPU backend is automatically selected for matrices above the threshold (typically 256×256).

Benchmark Suite

Run the matrix benchmark suite:

cargo run --release --example benchmark_matrix_suite

This compares:

• Naive O(n³) multiplication
• SIMD-optimized blocked multiplication
• Parallel (rayon) multiplication

See Also

• Eigendecomposition - Symmetric eigenvalue decomposition
• GPU Backend - GPU acceleration details
• Performance Targets - Expected speedups

Neural Networks

This chapter demonstrates Trueno's neural network primitives using the activation_functions example.

Running the Example

cargo run --example activation_functions

Activation Functions

Trueno provides 11 activation functions commonly used in neural networks:

Basic Activations

use trueno::Vector;

let x = Vector::from_slice(&[0.5, -0.2, 1.2, -0.8, 2.1]);

// ReLU - Rectified Linear Unit
let relu = x.relu()?;  // max(0, x)

// Sigmoid - Logistic function
let sigmoid = x.sigmoid()?;  // 1 / (1 + exp(-x))

// Tanh - Hyperbolic tangent
let tanh_result = x.tanh_activation()?;  // (exp(x) - exp(-x)) / (exp(x) + exp(-x))

Advanced Activations

// GELU - Gaussian Error Linear Unit (Transformer default)
let gelu = x.gelu()?;

// Swish/SiLU - x * sigmoid(x) (EfficientNet)
let swish = x.swish()?;

// Mish - x * tanh(softplus(x)) (YOLOv4)
let mish = x.mish()?;

// SELU - Self-Normalizing ELU
let selu = x.selu()?;

// Hardswish - Efficient approximation (MobileNetV3)
let hardswish = x.hardswish()?;

// Softplus - Smooth ReLU approximation
let softplus = x.softplus()?;

// ELU - Exponential Linear Unit
let elu = x.elu(1.0)?;  // alpha = 1.0

// Leaky ReLU - ReLU with negative slope
let leaky = x.leaky_relu(0.01)?;  // alpha = 0.01

Softmax (Probability Distribution)

let logits = Vector::from_slice(&[1.0, 2.0, 3.0, 4.0, 5.0]);
let probs = logits.softmax()?;

// Properties:
// - All values in [0, 1]
// - Sum = 1.0

When to Use Each Activation

Building a Simple MLP

use trueno::{Vector, Matrix};

fn mlp_forward(
    input: &Vector,
    weights1: &Matrix,
    bias1: &Vector,
    weights2: &Matrix,
    bias2: &Vector,
) -> Result<Vector, TruenoError> {
    // Layer 1: Linear + ReLU
    let h1 = weights1.matvec(input)?;
    let h1 = h1.add(bias1)?;
    let h1 = h1.relu()?;

    // Layer 2: Linear + Softmax
    let h2 = weights2.matvec(&h1)?;
    let h2 = h2.add(bias2)?;
    h2.softmax()
}

Transformer Building Blocks

Layer Normalization

fn layer_norm(x: &Vector, gamma: &Vector, beta: &Vector) -> Result<Vector, TruenoError> {
    let mean = x.mean()?;
    let centered = x.sub_scalar(mean)?;
    let var = centered.dot(&centered)? / x.len() as f32;
    let std = (var + 1e-5).sqrt();

    let normalized = centered.mul_scalar(1.0 / std)?;
    let scaled = normalized.mul(gamma)?;
    scaled.add(beta)
}

Attention Scores

fn attention_scores(query: &Vector, key: &Vector) -> Result<f32, TruenoError> {
    let d_k = query.len() as f32;
    let score = query.dot(key)?;
    Ok(score / d_k.sqrt())
}

Performance Tips

Batching for Efficiency

// Process multiple samples together
let batch: Vec<Vector> = inputs
    .iter()
    .map(|x| x.relu().unwrap())
    .collect();

Fused Operations

// Fusing reduces memory bandwidth
// Instead of:
let h = x.relu()?.mul_scalar(scale)?;

// Use pre-scaled weights when possible

GPU Acceleration

For large batch sizes, use GPU:

cargo run --release --features gpu --example activation_functions

Fused Bias + Activation (GPU PTX)

For GPU inference, trueno-gpu provides a fused bias+activation kernel that combines bias addition with activation in a single kernel pass:

use trueno_gpu::kernels::{BiasActivationKernel, Kernel};

// Bias + GELU (common in Transformers)
let kernel = BiasActivationKernel::new(4096, 256).with_gelu();

// Bias + ReLU (common in CNNs)
let kernel = BiasActivationKernel::new(4096, 256).with_relu();

let ptx = kernel.emit_ptx();

This is typically used as an epilogue after GEMM operations, reducing memory bandwidth by avoiding intermediate writes.

cargo run -p trueno-gpu --example bias_activation

See Also

• Activation Functions Example - Full source
• ML Similarity - k-NN example
• Performance Demo - SIMD speedups

Image Processing

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Signal Processing

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Scientific Computing

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Contributing

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Extreme Tdd

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Testing

This chapter covers Trueno's comprehensive testing strategy and quality gates.

Overview

Trueno follows Extreme TDD principles with multiple layers of testing:

• Unit Tests: Correctness for all operations
• Property-Based Tests: Mathematical invariants (proptest)
• Backend Equivalence Tests: All backends produce identical results
• Mutation Testing: >80% mutation kill rate
• Coverage: 90%+ line coverage required

Running Tests

Quick Tests (Development)

# Fast tests with nextest (parallel execution)
make test-fast

# Run all tests with output
make test

# Verbose output (single-threaded)
make test-verbose

Coverage Commands

Trueno provides multiple coverage targets for different use cases:

# Standard coverage (fast, ~85%)
make coverage

# GPU-specific coverage (WGPU + CUDA tests)
make coverage-gpu

# Full coverage (fast tests + GPU tests sequentially)
make coverage-all

# View coverage summary
make coverage-summary

# Open HTML report in browser
make coverage-open

Coverage Targets

Test Categories

1. Unit Tests

Basic correctness tests for all operations:

#[test]
fn test_add_correctness() {
    let a = vec![1.0, 2.0, 3.0, 4.0];
    let b = vec![5.0, 6.0, 7.0, 8.0];
    let result = add_f32(&a, &b).unwrap();
    assert_eq!(result, vec![6.0, 8.0, 10.0, 12.0]);
}

#[test]
fn test_add_empty() {
    let result = add_f32(&[], &[]).unwrap();
    assert!(result.is_empty());
}

2. Property-Based Tests

Using proptest to verify mathematical properties:

use proptest::prelude::*;

proptest! {
    #[test]
    fn test_add_commutative(
        a in prop::collection::vec(-1000.0f32..1000.0, 1..1000),
        b in prop::collection::vec(-1000.0f32..1000.0, 1..1000)
    ) {
        let len = a.len().min(b.len());
        let a = &a[..len];
        let b = &b[..len];

        let result1 = add_f32(a, b).unwrap();
        let result2 = add_f32(b, a).unwrap();

        assert_eq!(result1, result2);
    }
}

3. Backend Equivalence Tests

Verify all backends produce identical results:

#[test]
fn test_backend_equivalence_add() {
    let a = vec![1.0f32; 10000];
    let b = vec![2.0f32; 10000];

    let scalar = add_vectors_scalar(&a, &b);
    let sse2 = unsafe { add_vectors_sse2(&a, &b) };
    let avx2 = unsafe { add_vectors_avx2(&a, &b) };

    // Allow small floating-point tolerance
    for i in 0..scalar.len() {
        assert!((scalar[i] - sse2[i]).abs() < 1e-5);
        assert!((scalar[i] - avx2[i]).abs() < 1e-5);
    }
}

Quality Gates

Pre-Commit Checklist

Every commit must pass:

# Full quality gate check
make quality-gates

# Individual checks
make lint          # Zero clippy warnings
make fmt-check     # Proper formatting
make test-fast     # All tests pass
make coverage      # >90% coverage

Tiered Testing (CI/CD)

# Tier 1: On-save (sub-second)
make tier1

# Tier 2: On-commit (1-5 minutes)
make tier2

# Tier 3: On-merge/nightly (hours)
make tier3

GPU Testing

GPU tests require special handling due to hardware dependencies:

# Check if GPU is available
cargo test --all-features test_gpu_backend_available_check

# Run GPU-specific tests
make coverage-gpu

# GPU tests use shared device pattern for faster execution
# See: src/backends/gpu/batch.rs

GPU Test Patterns

GPU tests use a shared device to reduce initialization overhead:

use std::sync::OnceLock;

static SHARED_DEVICE: OnceLock<Option<GpuDevice>> = OnceLock::new();

fn get_shared_device() -> Option<GpuDevice> {
    SHARED_DEVICE
        .get_or_init(|| {
            if GpuDevice::is_available() {
                GpuDevice::new().ok()
            } else {
                None
            }
        })
        .clone()
}

#[test]
fn test_gpu_operation() {
    let Some(device) = get_shared_device() else {
        eprintln!("GPU not available, skipping");
        return;
    };
    // Test with device...
}

Mutation Testing

Verify test quality with mutation testing:

# Run mutation testing (target: >80% kill rate)
make mutate

# Or directly with cargo-mutants
cargo mutants --timeout 60 --minimum-pass-rate 80

Nextest Configuration

Trueno uses cargo-nextest for parallel test execution. Configuration is in .config/nextest.toml:

[profile.default]
slow-timeout = { period = "30s", terminate-after = 2 }
test-threads = "num-cpus"

[profile.coverage]
slow-timeout = { period = "20s", terminate-after = 2 }
# Exclude slow async GPU batch tests from fast coverage
default-filter = "not test(/test_matmul_parallel_1024/) and not test(/batch::tests::test_all_batch_operations/)"

Troubleshooting

Coverage Too Low

1. Check which files have low coverage:

   make coverage
   # Look at the detailed breakdown
   

2. For GPU code, run GPU-specific coverage:

   make coverage-gpu
   

Tests Timing Out

1. Increase timeout in .config/nextest.toml
2. Use --test-threads=1 for GPU tests
3. Check for resource contention

GPU Tests Failing

1. Verify GPU availability:

   cargo test --all-features test_gpu_backend_available_check
   

2. Check CUDA installation (for CUDA tests):

   nvidia-smi
   

3. Run GPU tests in isolation:

   cargo test --all-features -- --test-threads=1 gpu
   

Unit Tests

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Property Based Tests

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Backend Equivalence Tests

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Mutation Testing

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Benchmarking

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Quality Gates

Trueno enforces rigorous quality gates following Toyota Production System principles (Jidoka, Genchi Genbutsu). This chapter documents the quality enforcement mechanisms implemented in TRUENO-SPEC-013.

Overview

Quality gates are automated checks that must pass before code can be committed or merged:

Coverage Requirements

Minimum Thresholds

Daily Commits:  ≥90% line coverage
Releases:       ≥95% line coverage (TRUENO-SPEC-013)

Running Coverage

# Generate coverage report (<5 minutes)
make coverage

# View HTML report
open target/coverage/html/index.html

Coverage Breakdown

The coverage report shows per-crate metrics:

trueno:     92.44%  (core library)
trueno-gpu: 93.12%  (GPU/CUDA backend)

Technical Notes

Coverage instrumentation requires disabling the mold linker:

# The Makefile handles this automatically:
# 1. Backs up ~/.cargo/config.toml
# 2. Runs tests with llvm-cov
# 3. Restores config

Smoke Tests (TRUENO-SPEC-013)

Smoke tests verify backend equivalence across SIMD, WGPU, and CUDA:

# Run all smoke tests
make smoke

# Individual backend tests
cargo test --test smoke_e2e smoke_simd -- --nocapture
cargo test --test smoke_e2e smoke_wgpu --features gpu -- --nocapture

Smoke Test Categories

1. SIMD Backend Tests

   • Vector add, mul, dot product
   • ReLU, Softmax activations
   • L2 norm computation

2. WGPU Backend Tests (requires GPU)

   • Vector operations (100K+ elements)
   • Matrix multiplication (256x256+)

3. Backend Equivalence Tests

   • Scalar vs Auto backend comparison
   • Floating-point tolerance: 1e-5

4. Edge Case Tests (Poka-Yoke)

   • Empty inputs
   • Single element
   • Non-aligned sizes (17 elements)
   • NaN/Infinity propagation

Pixel FKR Tests (Falsification Kernel Regression)

Pixel FKR tests catch GPU kernel bugs using Popperian falsification methodology:

# Run all pixel FKR tests
make pixel-fkr-all

# Individual suites
make pixel-scalar-fkr   # Baseline truth
make pixel-simd-fkr     # SIMD vs scalar
make pixel-wgpu-fkr     # WGPU vs scalar
make pixel-ptx-fkr      # PTX validation (CUDA)

FKR Test Suites

Realizer Operations Tested

• RMS Normalization
• SiLU Activation
• Softmax
• RoPE (Rotary Position Embedding)
• Causal Mask
• Q4_K Dequantization

Pre-Commit Hook

The pre-commit hook (.git/hooks/pre-commit) enforces all quality gates:

# Gates checked on every commit:
1. PMAT TDG regression check
2. PMAT TDG quality check (B+ minimum)
3. Bashrs linting (Makefile, shell scripts)
4. Coverage threshold (≥90%)

Bypassing (Not Recommended)

# Only for emergencies - document why in commit message
git commit --no-verify

Tiered Quality Workflow

Trueno uses a tiered approach inspired by certeza (97.7% mutation score):

Tier 1: On-Save (Sub-second)

make tier1
# Checks: cargo check, clippy (lib), unit tests, property tests (10 cases)

Tier 2: On-Commit (1-5 minutes)

make tier2
# Checks: fmt, full clippy, all tests, property tests (256 cases), coverage, TDG

Tier 3: On-Merge/Nightly (Hours)

make tier3
# Checks: tier2 + mutation testing (80%), security audit, full benchmarks

PMAT Integration

PMAT (Pragmatic AI Labs Multi-Agent Toolkit) provides Technical Debt Grading:

# Check TDG grade
pmat analyze tdg --min-grade B+

# Repository health score
pmat repo-score . --min-score 90

TDG Grading Scale

Toyota Way Principles

Jidoka (Built-in Quality)

Quality is built in through:

• Pre-commit hooks that stop defects immediately
• Automated testing at every tier
• No bypass without explicit override

Genchi Genbutsu (Go and See)

• Smoke tests run actual code on real hardware
• Pixel FKR tests verify visual output
• No simulation - real GPU execution

Poka-Yoke (Error Prevention)

• Edge case tests prevent common bugs
• Type system enforces API contracts
• Clippy warnings are errors

Quick Reference

# Full quality check
make quality-spec-013

# Coverage only
make coverage

# Smoke tests
make smoke

# Pixel FKR
make pixel-fkr-all

# Tier 2 (pre-commit)
make tier2

See Also

• Testing - Test infrastructure details
• Extreme TDD - TDD methodology
• TRUENO-SPEC-013 - Full specification

SATD Remediation Guide

This guide documents the process and patterns for identifying and fixing Self-Admitted Technical Debt (SATD) in trueno-gpu kernels.

What is SATD?

Self-Admitted Technical Debt (SATD) refers to code where developers have explicitly acknowledged shortcuts or incomplete implementations through comments. Common SATD markers include:

• // TODO
• // FIXME
• // HACK
• // Simplified
• // Placeholder
• // Exit after one iteration for simplicity

The Stubbed Loop Pattern

The most critical SATD pattern in GPU kernels is the stubbed loop:

// ANTI-PATTERN: Stubbed Loop (SATD)
let counter = ctx.mov_u32_imm(0);
ctx.label("loop_start");
let done = ctx.setp_ge_u32(counter, max);
ctx.branch_if(done, "loop_end");

// ... loop body ...

let _next = ctx.add_u32(counter, 1);  // INCREMENT DISCARDED!
ctx.branch("loop_end");                // WRONG: exits immediately!

ctx.label("loop_end");

Why it's dangerous:

• Loop executes only once regardless of input size
• Produces mathematically incorrect results
• Silently fails on real data (works on trivial test cases)

The Fix Pattern

Correct loop implementation uses in-place updates:

// CORRECT: Proper Loop
let counter = ctx.mov_u32_imm(0);
ctx.label("loop_start");
let done = ctx.setp_ge_u32(counter, max);
ctx.branch_if(done, "loop_end");

// ... loop body ...

ctx.add_u32_inplace(counter, 1);  // IN-PLACE UPDATE
ctx.branch("loop_start");          // BRANCH BACK TO LOOP

ctx.label("loop_end");

TRUENO-SATD-001 Fixes

The following SATD issues were identified and fixed:

1. quantize.rs: K-loop (Lines 232-233)

Before:

let _k_next = ctx.add_u32(k_block, 1);
ctx.branch("k_block_done");  // Simplified - single iteration

After:

ctx.add_u32_inplace(k_block, 1);
ctx.branch("k_block_loop");

2. quantize.rs: Shuffle Broadcast (Line 226)

Before:

let broadcast_sum = ctx.shfl_down_f32(block_sum, 0, mask);  // No-op!

After:

let broadcast_sum = ctx.shfl_idx_f32(block_sum, 0, mask);  // Proper broadcast

Why: shfl_down(x, 0) is a no-op (shifts by 0). Use shfl_idx(x, 0) to broadcast lane 0's value.

3. softmax.rs: Max-Reduce (Lines 214-215)

Before:

let _next_stride = ctx.add_u32(stride_reg, 0);  // placeholder
ctx.branch("max_reduce_done");  // Exit after one iteration

After:

ctx.shr_u32_inplace(stride_reg, 1);  // Halve stride
ctx.branch("max_reduce_loop");        // Loop back

4. softmax.rs: Sum-Reduce

Similar fix applied to sum reduction loop.

Testing SATD Fixes (EXTREME TDD)

Every SATD fix requires falsifiable tests:

#[test]
fn test_kloop_branches_back_to_loop_start() {
    let kernel = QuantizeKernel::new(64, 64, 128);
    let ptx = kernel.emit_ptx();

    let has_loop_back = ptx.contains("bra k_block_loop");

    assert!(
        has_loop_back,
        "FALSIFIED: K-loop does not branch back to loop start"
    );
}

Running the Example

Verify SATD fixes with:

cargo run --example satd_kernels

Expected output:

╔══════════════════════════════════════════════════════════════╗
║     SATD Remediation: Fixed Kernel Examples                  ║
╚══════════════════════════════════════════════════════════════╝

K-loop fix verified: ✓ PASS
Shuffle fix verified: ✓ PASS
Max-reduce fix verified: ✓ PASS
Sum-reduce fix verified: ✓ PASS
Stride halving verified: ✓ PASS

In-Place Update Methods

The PTX builder provides in-place update methods for loops:

Prevention Checklist

Before committing GPU kernel code:

1. Search for SATD comments: grep -r "Simplified\|TODO\|FIXME" src/kernels/
2. Verify loop structure: Branch targets should be loop_start, not loop_done
3. Check in-place updates: Loop counters use _inplace methods
4. Run SATD tests: cargo test test_kloop test_shuffle test_reduce
5. Run example: cargo run --example satd_kernels

References

• Specification: docs/specifications/fix-stubbed-kernel-loops-enhanced-monitoring-pixel-level-gpu-stress-testing-probar.md
• Academic: Potdar & Shihab (2014), "An exploratory study on self-admitted technical debt"
• Related: PARITY-040 (WMMA Infrastructure), PARITY-041 (Q4_K GGML Format)

PTX Best Practices

This document covers PTX assembly generation best practices learned from development and debugging of trueno-gpu CUDA kernels.

Register Types

U8 Registers Are Not Supported

Issue: PTX does not support 8-bit register types (.u8, .s8).

Incorrect:

.reg .u8 %rs<1>;  // ERROR: Invalid register type
ld.global.u8 %rs0, [%rd0];

Correct:

.reg .u16 %rh<1>;  // Minimum register size is 16-bit
ld.global.u8 %rh0, [%rd0];  // Load zero-extends to 16-bit

The ld.global.u8 instruction is valid, but it must store into a 16-bit or larger register. The loaded byte is zero-extended.

Half-Precision (F16) Operations

Loading F16 Values

Issue: PTX uses .b16 (binary 16-bit) for half-precision loads, not .f16.

Incorrect:

ld.global.f16 %h0, [%rd0];  // ERROR: Invalid type for load

Correct:

ld.global.b16 %h0, [%rd0];  // Load 16-bit binary value

F16 to F32 Conversion

Issue: Converting from f16 to f32 is exact and does NOT require a rounding modifier.

Incorrect:

cvt.rn.f32.f16 %f0, %h0;  // ERROR: Illegal rounding modifier

Correct:

cvt.f32.f16 %f0, %h0;  // No rounding needed (exact conversion)

Note: The reverse conversion (f32 → f16) DOES require a rounding modifier:

cvt.rn.f16.f32 %h0, %f0;  // Correct: rounding needed for narrowing

Bitwise Operations

AND, OR, XOR Types

Issue: PTX requires .b32 (binary) type for bitwise operations, not .u32.

Incorrect:

and.u32 %r2, %r0, %r1;  // ERROR: Invalid type
or.u32 %r2, %r0, %r1;   // ERROR: Invalid type

Correct:

and.b32 %r2, %r0, %r1;  // Use .b32 for bitwise ops
or.b32 %r2, %r0, %r1;
xor.b32 %r2, %r0, %r1;

Warp Shuffle Operations

Shuffle Width Parameter

Issue: The width parameter in shfl.sync.idx must be a power of 2 (1, 2, 4, 8, 16, or 32).

Incorrect:

shfl.sync.idx.b32 %f0, %f1, 0, 31, 0xFFFFFFFF;  // ERROR: 31 is not power of 2

Correct:

shfl.sync.idx.b32 %f0, %f1, 0, 32, 0xFFFFFFFF;  // 32 is valid

Warp Participation

Issue: shfl.sync with mask 0xFFFFFFFF requires ALL 32 threads in the warp to execute the instruction simultaneously.

If some threads exit early (e.g., via @%p bra exit), the remaining threads cannot perform shuffles.

Solution: Use address clamping to ensure all threads access valid memory, then skip only the final store for out-of-bounds threads:

// Clamp addresses for all threads
min.u32 %r_clamped_row, %r_global_row, %r_m_minus_1;
min.u32 %r_clamped_col, %r_global_col, %r_n_minus_1;

// All threads participate in computation and shuffles
// ...shuffle reduction code...

// Only in-bounds threads store
@%p_row_oob bra exit;
@%p_col_oob bra exit;
st.global.f32 [%rd_out], %f_result;
exit:
    ret;

Memory Alignment

4-Byte Alignment for U32 Loads

Issue: ld.global.u32 requires the address to be 4-byte aligned.

Incorrect:

// If header has 2-byte f16 scale at offset 0, and we try to read
// another u32 at offset 2, it will be misaligned
add.u64 %rd1, %rd0, 2;
ld.global.u32 %r0, [%rd1];  // ERROR: Misaligned access

Correct: Use smaller loads for misaligned data:

ld.global.b16 %rh0, [%rd0];  // Load 2-byte aligned data

Testing PTX

Always validate generated PTX with ptxas:

ptxas -arch=sm_89 -v kernel.ptx -o kernel.cubin

Use compute-sanitizer for runtime memory access checking:

compute-sanitizer --tool memcheck ./your_program

References

• PTX ISA Reference
• GitHub Issue #67 - U8 register bug
• GitHub Issue #68 - F16 load/convert bug

PTX Bug Detection

The trueno-explain crate provides static analysis for PTX (NVIDIA GPU assembly) to detect common bugs and performance issues before runtime.

Overview

Hand-written PTX is error-prone. The PTX bug analyzer catches:

Quick Start

use trueno_explain::{PtxBugAnalyzer, BugSeverity};

// Analyze PTX string
let ptx = include_str!("kernel.ptx");
let result = PtxBugAnalyzer::new().analyze(ptx);

// Check for bugs
if result.has_bugs() {
    println!("{}", result.format_report());
}

// Check specific severity
let critical = result.count_by_severity(BugSeverity::Critical);
assert_eq!(critical, 0, "No P0 bugs allowed!");

Analyzer Modes

Default Mode

Standard analysis - catches obvious bugs:

let analyzer = PtxBugAnalyzer::new();
let result = analyzer.analyze(ptx);

Strict Mode

Catches more potential issues (may have false positives):

let analyzer = PtxBugAnalyzer::strict();
let result = analyzer.analyze(ptx);

With Whitelist

Suppress known acceptable warnings:

use trueno_explain::PtxBugClass;

let analyzer = PtxBugAnalyzer::new()
    .with_whitelist("tensor_core*", PtxBugClass::HighRegisterPressure,
        "Tensor core kernels need high registers");

Quantized Kernel Whitelist

Pre-configured for quantized kernels (q4k, q5k, q6k, q8k):

// Suppresses HighRegisterPressure for quantized kernels
let analyzer = PtxBugAnalyzer::with_quantized_whitelist();

Examples

Run Deep Bug Hunt

Analyze all trueno-gpu kernels:

cargo run -p trueno-explain --example deep_bug_hunt

Output:

SUMMARY: 30 kernels analyzed
  Total bugs: 16
  P0 Critical: 0
  P1 High: 16
  P2 Medium: 0

BUGS BY CLASS:
  HIGH_REG_PRESSURE         : 16

Analyze External PTX

Analyze hand-rolled PTX from another project:

cargo run -p trueno-explain --example analyze_realizar

Output:

REALIZAR PTX SUMMARY
  Files analyzed: 4
  Total bugs: 18
  P0 Critical: 0
  P1 High: 15
  P2 Medium: 3

Inspect PTX Details

Deep dive into specific kernel PTX:

cargo run -p trueno-explain --example ptx_inspector

Bug Classes in Detail

P0 Critical - Correctness Bugs

SharedMemU64Addressing

Problem: Using 64-bit registers for shared memory addressing.

// BAD: %rd0 is 64-bit
st.shared.f32 [%rd0], %f0;

// GOOD: %r0 is 32-bit
st.shared.f32 [%r0], %f0;

Impact: Undefined behavior, potential silent corruption.

MissingBarrierSync

Problem: No bar.sync between shared memory write and read.

// BAD: Race condition!
st.shared.f32 [%r0], %f0;
ld.shared.f32 %f1, [%r1];  // May read stale data

// GOOD: Barrier ensures visibility
st.shared.f32 [%r0], %f0;
bar.sync 0;
ld.shared.f32 %f1, [%r1];

Impact: Race condition, non-deterministic results.

P1 High - Performance Bugs

HighRegisterPressure

Problem: >64 registers per thread reduces occupancy.

Register count: 120
Max occupancy: 65536 / (120 * 32) = 17 warps/SM (53%)

Impact: Reduced parallelism, lower throughput.

Fix: Reduce live variables, split kernel, or accept lower occupancy for compute-bound kernels.

PlaceholderCode

Problem: Comments indicate incomplete implementation.

// Detected patterns:
// "simplified"
// "omitted"
// "placeholder"
// "for now"
// "TODO"

Impact: Kernel may produce incorrect results or have missing functionality.

P2 Medium - Optimization Opportunities

DeadCode

Problem: Unreachable code after unconditional branch/return.

// BAD: add.f32 is unreachable
ret;
add.f32 %f0, %f1, %f2;

// BAD: mul.f32 is unreachable
bra skip;
mul.f32 %f0, %f1, %f2;
skip:

Impact: Code bloat, wasted compilation time.

UnoptimizedMemoryPattern

Problem: Multiple single-element loads that could be vectorized.

// BAD: 4 separate loads
ld.global.f32 %f0, [%rd0];
ld.global.f32 %f1, [%rd0+4];
ld.global.f32 %f2, [%rd0+8];
ld.global.f32 %f3, [%rd0+12];

// GOOD: Single vectorized load
ld.global.v4.f32 {%f0, %f1, %f2, %f3}, [%rd0];

Impact: 4x memory bandwidth reduction.

Integration with CI

Add PTX bug detection to your CI pipeline:

# .github/workflows/ptx-analysis.yml
- name: PTX Bug Analysis
  run: |
    cargo run -p trueno-explain --example deep_bug_hunt
    # Fail if any P0 bugs found
    cargo test -p trueno-explain --test ptx_bug_hunting

Writing Bug-Free PTX

Use trueno-gpu kernel generators instead of hand-writing PTX:

use trueno_gpu::kernels::{GemmKernel, Kernel};

// Generated PTX is verified bug-free
let kernel = GemmKernel::tiled(1024, 1024, 1024, 32);
let ptx = kernel.emit_ptx();

// Verify with analyzer
let result = PtxBugAnalyzer::new().analyze(&ptx);
assert!(result.is_valid());

API Reference

PtxBugAnalyzer

impl PtxBugAnalyzer {
    /// Create default analyzer
    pub fn new() -> Self;

    /// Create strict mode analyzer
    pub fn strict() -> Self;

    /// Pre-configured whitelist for quantized kernels
    pub fn with_quantized_whitelist() -> Self;

    /// Add whitelist entry
    pub fn with_whitelist(
        self,
        kernel_pattern: &str,  // e.g., "q4k*"
        bug_class: PtxBugClass,
        reason: &str
    ) -> Self;

    /// Analyze PTX and return report
    pub fn analyze(&self, ptx: &str) -> PtxBugReport;
}

PtxBugReport

impl PtxBugReport {
    /// Check if any bugs found
    pub fn has_bugs(&self) -> bool;

    /// Check for specific bug class
    pub fn has_bug(&self, class: &PtxBugClass) -> bool;

    /// Check if kernel is valid (no P0/P1 bugs)
    pub fn is_valid(&self) -> bool;

    /// Count bugs by severity
    pub fn count_by_severity(&self, severity: BugSeverity) -> usize;

    /// Get formatted report string
    pub fn format_report(&self) -> String;
}

See Also

• PTX Best Practices
• PTX Register Allocation
• PTX Code Generation

Code Review

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Simd Intrinsics

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Phase 2 Micro-Kernel: Achieving NumPy Performance Parity

Overview

The Phase 2 micro-kernel implementation represents a major performance milestone for Trueno: achieving parity with highly optimized BLAS libraries (NumPy/OpenBLAS) while maintaining a pure Rust codebase with zero external dependencies.

Achievement Summary:

• 256×256 matmul: 538 μs (vs NumPy 574 μs = 6% faster)
• 128×128 matmul: 72 μs (vs NumPy 463 μs = 6.4× faster)
• Improvement: 2.3-2.6× faster than Trueno v0.5.0
• Implementation: Pure Rust with AVX2/FMA intrinsics
• Safety: 100% safe public API, unsafe isolated to backends

Motivation

The Performance Gap

Prior to Phase 2, Trueno's matrix multiplication performance was:

• 128×128: 166 μs (2.79× faster than NumPy) ✅
• 256×256: 1391 μs (2.4× slower than NumPy) ❌

The performance cliff at 256×256 was caused by:

1. Sub-optimal memory access patterns
2. Cache inefficiency for larger matrices
3. Missed opportunities for register blocking
4. Sequential row processing (no parallelism within blocks)

Design Goals

1. Match BLAS Performance: Achieve ≤600 μs at 256×256 (NumPy baseline: 574 μs)
2. Pure Rust: No external C/BLAS dependencies
3. Zero Regressions: Maintain or improve performance at all matrix sizes
4. Safe API: Keep public API 100% safe
5. Maintainability: Clear, documented code with comprehensive tests

Implementation Strategy

Micro-Kernel Architecture

The micro-kernel is the computational core that processes a small block of the output matrix. Our design uses a 4×1 micro-kernel:

Input:  4 rows of matrix A (each length K)
        1 column of matrix B (length K)
Output: 4 scalar dot products

Processing: Simultaneously compute 4 dot products using AVX2 SIMD

Key Advantages:

• Register Blocking: Keep 4 accumulators in YMM registers (no memory traffic)
• Memory Efficiency: Load B column once, reuse for 4 A rows (4× bandwidth reduction)
• FMA Instructions: Fused multiply-add for 3× throughput vs separate ops
• Parallelism: 4 independent dot products computed in parallel

Algorithm Overview

fn matmul_simd(A: &Matrix, B: &Matrix) -> Matrix {
    // 1. Transpose B for cache-friendly access
    let B_T = B.transpose();

    // 2. L2 cache blocking (64×64 blocks)
    for (i_block, j_block, k_block) in blocks {

        // 3. Micro-kernel: Process 4 rows at a time
        for i in (i_block..i_end).step_by(4) {
            let a_rows = [A[i], A[i+1], A[i+2], A[i+3]];

            for j in j_block..j_end {
                let b_col = B_T[j];

                // 4×1 micro-kernel computes 4 dot products
                let dots = microkernel_4x1_avx2(a_rows, b_col);

                // Accumulate results
                result[i][j]   += dots[0];
                result[i+1][j] += dots[1];
                result[i+2][j] += dots[2];
                result[i+3][j] += dots[3];
            }
        }
    }
}

AVX2 Micro-Kernel Implementation

Core Function

#[target_feature(enable = "avx2,fma")]
#[inline]
unsafe fn matmul_microkernel_4x1_avx2(
    a_rows: [&[f32]; 4],  // 4 rows of A
    b_col: &[f32],        // 1 column of B (transposed)
    results: &mut [f32; 4],
) {
    use std::arch::x86_64::*;

    let len = b_col.len();
    let chunks = len / 8;  // AVX2 processes 8 f32 elements

    // Step 1: Initialize accumulators (stay in registers)
    let mut acc0 = _mm256_setzero_ps();
    let mut acc1 = _mm256_setzero_ps();
    let mut acc2 = _mm256_setzero_ps();
    let mut acc3 = _mm256_setzero_ps();

    // Step 2: Main SIMD loop (processes 8 elements per iteration)
    for i in 0..chunks {
        let offset = i * 8;

        // Load B column ONCE (critical optimization)
        let b_vec = _mm256_loadu_ps(b_col.as_ptr().add(offset));

        // Load A rows and FMA (Fused Multiply-Add)
        let a0_vec = _mm256_loadu_ps(a_rows[0].as_ptr().add(offset));
        acc0 = _mm256_fmadd_ps(a0_vec, b_vec, acc0);  // acc0 += a0 * b

        let a1_vec = _mm256_loadu_ps(a_rows[1].as_ptr().add(offset));
        acc1 = _mm256_fmadd_ps(a1_vec, b_vec, acc1);

        let a2_vec = _mm256_loadu_ps(a_rows[2].as_ptr().add(offset));
        acc2 = _mm256_fmadd_ps(a2_vec, b_vec, acc2);

        let a3_vec = _mm256_loadu_ps(a_rows[3].as_ptr().add(offset));
        acc3 = _mm256_fmadd_ps(a3_vec, b_vec, acc3);
    }

    // Step 3: Horizontal sum (reduce 8 elements to 1 scalar)
    results[0] = horizontal_sum_avx2(acc0);
    results[1] = horizontal_sum_avx2(acc1);
    results[2] = horizontal_sum_avx2(acc2);
    results[3] = horizontal_sum_avx2(acc3);

    // Step 4: Handle remainder (non-multiple of 8)
    let remainder_start = chunks * 8;
    if remainder_start < len {
        for i in remainder_start..len {
            results[0] += a_rows[0][i] * b_col[i];
            results[1] += a_rows[1][i] * b_col[i];
            results[2] += a_rows[2][i] * b_col[i];
            results[3] += a_rows[3][i] * b_col[i];
        }
    }
}

Horizontal Sum Helper

The horizontal sum reduces 8 f32 values in a YMM register to a single scalar:

#[target_feature(enable = "avx2")]
#[inline]
unsafe fn horizontal_sum_avx2(v: __m256) -> f32 {
    use std::arch::x86_64::*;

    // Step 1: Sum upper and lower 128-bit lanes
    //   [a7, a6, a5, a4 | a3, a2, a1, a0]
    //   → [a7+a3, a6+a2, a5+a1, a4+a0]
    let sum128 = _mm_add_ps(
        _mm256_castps256_ps128(v),        // Lower 128 bits
        _mm256_extractf128_ps(v, 1),      // Upper 128 bits
    );

    // Step 2: Horizontal add within 128-bit lane
    //   [a7+a3, a6+a2, a5+a1, a4+a0]
    //   → [a7+a3+a6+a2, a5+a1+a4+a0, ...]
    let sum64 = _mm_hadd_ps(sum128, sum128);

    // Step 3: Horizontal add again
    //   → [a7+a6+a5+a4+a3+a2+a1+a0, ...]
    let sum32 = _mm_hadd_ps(sum64, sum64);

    // Step 4: Extract final scalar
    _mm_cvtss_f32(sum32)
}

Performance Analysis

Benchmark Results

Why the Micro-Kernel Works

1. Register Blocking

• 4 YMM accumulators stay in CPU registers
• Zero memory traffic during accumulation
• Theoretical peak: 16 FLOPs/cycle (AVX2 FMA)

2. Memory Bandwidth Optimization

• B column loaded once per 4 A rows
• Bandwidth reduction: 4×
• Effective throughput: ~50 GB/s on modern CPUs

3. FMA (Fused Multiply-Add)

Traditional: acc = acc + (a * b)   // 2 ops, 2 cycles
FMA:        acc = fmadd(a, b, acc) // 1 op, 1 cycle
Speedup:    3× throughput

4. Cache-Aware Blocking

• L2 blocks: 64×64 (fit in 256 KB L2 cache)
• Transposed B ensures sequential access
• Cache miss rate: <2%

Performance Model

Theoretical Peak (AVX2 + FMA):

• FLOP rate: 16 FLOP/cycle (2 FMAs × 8 wide)
• CPU @ 3.0 GHz: 48 GFLOPS
• 256×256 matmul: 2×256³ = 33.5 MFLOPs
• Expected time: 33.5M / 48G = 0.7 ms

Actual Performance:

• Measured: 538 μs
• Efficiency: 0.538 / 0.7 = 77% of theoretical peak

Efficiency Breakdown:

• Memory bandwidth: 15%
• Cache misses: 5%
• Remainder handling: 2%
• Instruction scheduling: 1%

Testing Strategy

Unit Tests

Comprehensive micro-kernel testing with 11 test cases:

#[test]
fn test_matmul_microkernel_4x1_avx2() {
    // Test 1: Simple dot products
    // Test 2: Identity-like pattern
    // Test 3: Non-aligned sizes (remainder handling)
    // Test 4: Mixed positive/negative values
    // Test 5: Zero accumulation
    // Test 6: FMA correctness verification
}

#[test]
fn test_horizontal_sum_avx2() {
    // Test 1: All ones
    // Test 2: Sequence 1..8
    // Test 3: Alternating signs
    // Test 4: Large values
    // Test 5: Mixed positive/negative
}

Backend Equivalence

Verify micro-kernel produces identical results to naive implementation:

#[test]
fn test_matmul_simd_equivalence_large() {
    let a = Matrix::from_vec(256, 256, test_data_a);
    let b = Matrix::from_vec(256, 256, test_data_b);

    let naive = a.matmul_naive(&b);
    let simd = a.matmul_simd(&b);

    // Floating-point tolerance: <1e-3 for accumulated values
    assert_matrices_equal(naive, simd, 1e-3);
}

Coverage

• Overall: 90.63% line coverage (Trueno library)
• Micro-kernel: 100% coverage
• Tests added: 240+ lines (2 comprehensive test functions)

Integration

Dispatch Logic

The micro-kernel is automatically selected for AVX2/AVX512 backends:

impl Matrix<f32> {
    pub fn matmul(&self, other: &Matrix<f32>) -> Result<Matrix<f32>> {
        match self.backend {
            Backend::AVX2 | Backend::AVX512 => {
                // Use micro-kernel for optimal performance
                self.matmul_simd(other)
            }
            Backend::SSE2 | Backend::NEON => {
                // Use standard SIMD path
                self.matmul_simd(other)
            }
            _ => {
                // Scalar fallback
                self.matmul_naive(other)
            }
        }
    }
}

Automatic Fallback

For matrices with non-multiple-of-4 rows, the implementation automatically falls back to standard SIMD processing for the remainder:

// Process 4 rows at a time
let mut i = ii;
while i + 4 <= i_end {
    // Use micro-kernel
    matmul_microkernel_4x1_avx2(...);
    i += 4;
}

// Handle remainder rows (<4)
for i in i..i_end {
    // Standard SIMD path
    avx2_dot_product(...);
}

Lessons Learned

What Worked

1. Register Blocking: Keeping accumulators in registers eliminated memory bottleneck
2. FMA Instructions: 3× throughput improvement was critical
3. 4×1 Micro-Kernel: Sweet spot between complexity and performance
4. B Transposition: Sequential memory access patterns crucial for cache efficiency

What Didn't Work

1. 3-Level Blocking: Extra loop nesting caused 7% regression

   • Root cause: Instruction cache pollution
   • Solution: Stick with 2-level blocking (L2 only)

2. 8×8 Micro-Kernel: Ran out of YMM registers

   • AVX2 has 16 YMM registers (8 for accumulators, 8 for inputs)
   • 8×8 needs 64 accumulators → register spilling
   • Solution: 4×1 is optimal for AVX2

3. Vertical Micro-Kernel (1 row × 4 cols): Poor cache behavior

   • Requires 4 B columns (scattered memory access)
   • Solution: Horizontal micro-kernel with transposed B

Trade-offs

Future Optimizations

Phase 3: Larger Matrices (512×512+)

Target: Within 1.5× of NumPy for 512×512 matrices

Strategies:

1. 8×1 micro-kernel for AVX-512 (32 f32 wide)
2. 3-level cache blocking (L3: 256×256, L2: 64×64)
3. Multi-threading with rayon for very large matrices

ARM NEON Micro-Kernel

Target: Match AVX2 performance on ARM64

Strategy:

• 4×1 micro-kernel using NEON intrinsics (128-bit, 4 f32 wide)
• FMA using vfmaq_f32 instruction
• Expected speedup: 2-3× vs current NEON path

GPU Integration

Target: 10-50× for matrices >1024×1024

Strategy:

• Automatic GPU dispatch for large matrices
• Tile-based GPU kernel (16×16 or 32×32 tiles)
• Overlap CPU computation with PCIe transfer

Conclusion

The Phase 2 micro-kernel demonstrates that pure Rust can match highly optimized BLAS libraries while maintaining:

• ✅ Zero external dependencies
• ✅ Safe public API
• ✅ Portable code (x86/ARM/WASM)
• ✅ Maintainable implementation

Key Takeaway: With careful algorithm design and SIMD optimization, Rust can achieve performance parity with hand-tuned C/assembly code.

References

• BLIS: BLIS micro-kernel design
• Rust SIMD: std::arch x86_64 intrinsics
• Trueno Benchmarks: v0.6.0 benchmark summary
• CHANGELOG: v0.6.0 release notes

Implemented in Trueno v0.6.0 (2025-11-21) Zero excuses. Zero defects. EXTREME TDD.

GPU Compute Shaders

Trueno uses WGSL (WebGPU Shading Language) compute shaders for cross-platform GPU acceleration via wgpu. This chapter covers the shader architecture, memory hierarchy abstractions, and tiled reduction algorithms.

Memory Hierarchy Abstractions

Based on the cuda-tile-behavior.md specification (Section 3.2), Trueno provides two key abstractions for efficient GPU memory access:

TensorView

TensorView<T> provides a structured view into GPU buffer memory with shape, stride, and layout metadata. It enables zero-copy operations like slicing and transposition.

use trueno::backends::gpu::{TensorView, MemoryLayout};

// Create a 4D tensor view (batch=2, channels=3, height=32, width=32)
let view: TensorView<f32> = TensorView::new([2, 3, 32, 32]);

println!("Shape: {:?}", view.shape());     // [2, 3, 32, 32]
println!("Strides: {:?}", view.strides()); // [3072, 1024, 32, 1]
println!("Elements: {}", view.numel());    // 6144

// Create with explicit strides for non-contiguous views
let transposed = TensorView::<f32>::with_strides(
    [32, 32, 3, 2],
    [1, 32, 1024, 3072]
);
assert!(!transposed.is_contiguous());

// Change memory layout
let col_major = TensorView::new([4, 4, 1, 1])
    .with_layout(MemoryLayout::ColumnMajor);

PartitionView

PartitionView<T> divides a tensor into tiles for efficient GPU workgroup distribution:

use trueno::backends::gpu::{TensorView, PartitionView};

// Partition a 64x64 tensor into 16x16 tiles
let tensor: TensorView<f32> = TensorView::new([64, 64, 1, 1]);
let partition: PartitionView<f32> = PartitionView::new(tensor, [16, 16, 1, 1]);

println!("Tile count: {:?}", partition.tile_count());  // [4, 4, 1, 1]
println!("Total tiles: {}", partition.total_tiles());  // 16

// Handle non-aligned dimensions (100x100 with 16x16 tiles)
let non_aligned: TensorView<f32> = TensorView::new([100, 100, 1, 1]);
let partition2: PartitionView<f32> = PartitionView::new(non_aligned, [16, 16, 1, 1]);

// Edge tiles are automatically detected
if let Some(tile_info) = partition2.get_tile([6, 6, 0, 0]) {
    println!("Edge tile size: {:?}", tile_info.size);  // [4, 4, 1, 1]
    println!("Is edge tile: {}", tile_info.is_edge);   // true
}

Tiled Reduction Algorithms

Trueno implements 16x16 tile-based reduction algorithms inspired by GPU workgroup patterns:

TILE_SIZE Constant

use trueno::backends::gpu::TILE_SIZE;

// TILE_SIZE = 16 matches standard GPU workgroup size
// This enables efficient shared memory usage and warp/wavefront alignment

Tiled Sum, Max, Min

use trueno::backends::gpu::{tiled_sum_2d, tiled_max_2d, tiled_min_2d};

// 32x32 matrix data (row-major)
let data: Vec<f32> = (1..=1024).map(|x| x as f32).collect();

// Tiled sum reduction
let sum = tiled_sum_2d(&data, 32, 32);
assert!((sum - 524800.0).abs() < 1e-3);

// Tiled max reduction
let max_data = vec![1.0, 5.0, 3.0, 9.0, 2.0, 7.0, 8.0, 4.0, 6.0];
let max = tiled_max_2d(&max_data, 3, 3);
assert!((max - 9.0).abs() < 1e-5);

// Tiled min reduction
let min_data = vec![5.0, 3.0, 7.0, -1.0, 9.0, 2.0];
let min = tiled_min_2d(&min_data, 2, 3);
assert!((min - (-1.0)).abs() < 1e-5);

Reduction Algorithm

The tiled reduction uses a tree-based pattern:

1. Load Phase: Each workgroup loads a 16x16 tile into shared memory
2. Row Reduction: 16 -> 8 -> 4 -> 2 -> 1 (horizontal)
3. Column Reduction: 16 -> 8 -> 4 -> 2 -> 1 (vertical)
4. Combine Phase: Partial results from all tiles are combined

Tile (16x16 elements)
┌────────────────────────────────────────┐
│ Row reduction: 16 -> 8 -> 4 -> 2 -> 1  │
│                                        │
│  [x x x x x x x x x x x x x x x x]     │
│        │                               │
│  [x x x x x x x x]  (step 1: +8)       │
│        │                               │
│  [x x x x]          (step 2: +4)       │
│        │                               │
│  [x x]              (step 3: +2)       │
│        │                               │
│  [x]                (step 4: +1)       │
│                                        │
│ Then column reduction on first column  │
└────────────────────────────────────────┘

Custom Reduction Operations

You can implement custom reductions using the ReduceOp trait:

use trueno::backends::gpu::{tiled_reduce_2d, ReduceOp, SumOp, MaxOp, MinOp};

// Built-in operations
let sum = tiled_reduce_2d::<SumOp>(&data, width, height);
let max = tiled_reduce_2d::<MaxOp>(&data, width, height);
let min = tiled_reduce_2d::<MinOp>(&data, width, height);

// ReduceOp trait for custom operations:
// - identity(): Starting value (0 for sum, -inf for max, inf for min)
// - combine(a, b): Binary operation to combine two values

WGSL Shader Architecture

Element-wise Operations

Element-wise shaders process one element per thread:

@compute @workgroup_size(256)
fn relu_kernel(
    @builtin(global_invocation_id) global_id: vec3<u32>
) {
    let idx = global_id.x;
    if (idx >= arrayLength(&input)) {
        return;
    }
    output[idx] = max(0.0, input[idx]);
}

Reduction Shaders

Reduction shaders use shared memory and tree reduction:

var<workgroup> tile: array<array<f32, 16>, 16>;

@compute @workgroup_size(16, 16)
fn tiled_sum_kernel(
    @builtin(global_invocation_id) global_id: vec3<u32>,
    @builtin(local_invocation_id) local_id: vec3<u32>,
    @builtin(workgroup_id) wg_id: vec3<u32>
) {
    // Load to shared memory
    let gx = global_id.x;
    let gy = global_id.y;
    let lx = local_id.x;
    let ly = local_id.y;

    if (gx < width && gy < height) {
        tile[ly][lx] = input[gy * width + gx];
    } else {
        tile[ly][lx] = 0.0;  // Identity for sum
    }
    workgroupBarrier();

    // Row reduction: 16 -> 8 -> 4 -> 2 -> 1
    if (lx < 8u) { tile[ly][lx] += tile[ly][lx + 8u]; }
    workgroupBarrier();
    if (lx < 4u) { tile[ly][lx] += tile[ly][lx + 4u]; }
    workgroupBarrier();
    if (lx < 2u) { tile[ly][lx] += tile[ly][lx + 2u]; }
    workgroupBarrier();
    if (lx < 1u) { tile[ly][lx] += tile[ly][lx + 1u]; }
    workgroupBarrier();

    // Column reduction on first column
    if (lx == 0u && ly < 8u) { tile[ly][0] += tile[ly + 8u][0]; }
    workgroupBarrier();
    // ... continue tree reduction

    // Write partial result
    if (lx == 0u && ly == 0u) {
        let tile_idx = wg_id.y * tiles_x + wg_id.x;
        partials[tile_idx] = tile[0][0];
    }
}

Performance Characteristics

Metal Validation Results (2026-01-03)

Validated on AMD Radeon Pro W5700X (Mac Pro 7,1):

Key finding: Consistent ~150 Melem/s throughput demonstrates efficient tiled reduction algorithm regardless of input size.

Best Practices

1. Use power-of-two tile sizes - Enables efficient memory coalescing and avoids bank conflicts
2. Prefer 16x16 workgroups - Matches warp/wavefront size on most GPUs
3. Minimize global memory access - Load once to shared memory, compute locally
4. Handle edge tiles - Use identity elements for out-of-bounds values
5. Use CPU fallback for validation - The tiled reduction CPU implementation mirrors GPU algorithm

Running Examples

# Run the GPU tiled reduction demo
cargo run --example gpu_tiled_reduction --features gpu --release

# Run GPU batch operations demo
cargo run --example gpu_batch_demo --features gpu --release

# Run tiled reduction benchmarks
cargo bench --features gpu --bench gpu_reduction

Related Documentation

• cuda-tile-behavior.md - Full specification
• Performance Targets - Expected speedups
• Backend Selection - When GPU is selected

Memory Alignment

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Vectorization Patterns

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Portability

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Ffmpeg Case Study

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Ruchy

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Depyler

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Decy

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Ruchy Lambda

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Ruchy Docker

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Pmat

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Design Philosophy

[Content to be added]

This chapter will cover:

• Overview and key concepts
• Implementation details
• Best practices
• Examples and use cases

Placeholder

This section is currently under development. Please check back later or refer to the source code and inline documentation for now.

Trueno: Multi-Target High-Performance Compute Library

Specification v1.0.0

Status: Draft Created: 2025-11-15 Author: Pragmatic AI Labs Quality Standard: EXTREME TDD (>90% coverage), Toyota Way, PMAT A+ grade

1. Executive Summary

Trueno (Spanish: "thunder") is a Rust library providing unified, high-performance compute primitives across three execution targets:

1. CPU SIMD - x86 (SSE2/AVX/AVX2/AVX-512), ARM (NEON), WASM (SIMD128)
2. GPU - Vulkan/Metal/DX12/WebGPU via wgpu
3. WebAssembly - Portable SIMD128 for browser/edge deployment

Design Principles:

• Write once, optimize everywhere: Single algorithm, multiple backends
• Runtime dispatch: Auto-select best implementation based on CPU features
• Zero unsafe in public API: Safety via type system, unsafe isolated in backend
• Benchmarked performance: Every optimization must prove ≥10% speedup
• Extreme TDD: >90% test coverage, mutation testing, property-based tests

1.1 Ecosystem Integration

Trueno is designed to integrate seamlessly with the Pragmatic AI Labs transpiler ecosystem:

Primary Integration Targets:

1. Ruchy - Language-level vector operations

   • Native Vector type in Ruchy syntax transpiles to trueno calls
   • Enables NumPy-like performance without Python overhead
   • Example: let v = Vector([1.0, 2.0]) + Vector([3.0, 4.0]) → trueno::Vector::add()

2. Depyler (Python → Rust transpiler)

   • Transpile NumPy array operations to trueno
   • Replace numpy.add() → trueno::Vector::add()
   • Achieve native performance for scientific Python code
   • Example: np.dot(a, b) → trueno::Vector::dot(&a, &b)

3. Decy (C → Rust transpiler)

   • Transpile C SIMD intrinsics to trueno safe API
   • Replace _mm256_add_ps() → trueno::Vector::add()
   • Eliminate unsafe blocks from transpiled C code
   • Example: FFmpeg SIMD code → safe trueno equivalents

Deployment Targets:

4. ruchy-lambda - AWS Lambda compute optimization

   • Drop-in performance boost for data processing functions
   • Auto-select AVX2 on Lambda (x86_64 baseline)
   • Improve cold start benchmarks via faster compute

5. ruchy-docker - Cross-language benchmarking

   • Add trueno benchmarks alongside C/Rust/Python baselines
   • Prove transpiler-generated code matches hand-written performance
   • Demonstrate SIMD/GPU speedups across platforms

Quality Enforcement:

6. paiml-mcp-agent-toolkit (PMAT) - Quality gates
   • Pre-commit hooks enforce >90% coverage
   • TDG grading (target: A- / 92+)
   • Repository health scoring (target: 90/110)
   • Mutation testing (target: 80% kill rate)
   • SATD detection and management

Unified Performance Story:

Python/C Code
     ↓
Depyler/Decy (transpile)
     ↓
Safe Rust + trueno (optimize)
     ↓
Deploy: Lambda/Docker/WASM (benchmark)
     ↓
PMAT (quality gate)

2. Architecture Overview

2.1 Target Execution Model

┌─────────────────────────────────────────────────┐
│           Trueno Public API (Safe)              │
│  compute(), map(), reduce(), transform()        │
└─────────────────────────────────────────────────┘
                      │
        ┌─────────────┼─────────────┐
        ▼             ▼             ▼
   ┌────────┐   ┌─────────┐   ┌──────────┐
   │  SIMD  │   │   GPU   │   │   WASM   │
   │ Backend│   │ Backend │   │  Backend │
   └────────┘   └─────────┘   └──────────┘
        │             │             │
   ┌────┴────┐   ┌────┴────┐   ┌───┴─────┐
   │ Runtime │   │  wgpu   │   │ SIMD128 │
   │ Detect  │   │ Compute │   │ Portable│
   └─────────┘   └─────────┘   └─────────┘
   │  │  │  │
   SSE2 AVX  NEON AVX512

2.2 Runtime Target Selection

Priority Order (best → fallback):

1. GPU (if available + workload size > threshold)
2. AVX-512 (if CPU supports)
3. AVX2 (if CPU supports)
4. AVX (if CPU supports)
5. SSE2 (baseline x86_64)
6. NEON (ARM64)
7. SIMD128 (WASM)
8. Scalar fallback

Selection Algorithm:

if gpu_available() && workload_size > GPU_THRESHOLD {
    gpu_backend()
} else if is_x86_feature_detected!("avx512f") {
    avx512_backend()
} else if is_x86_feature_detected!("avx2") {
    avx2_backend()
} else {
    sse2_backend()  // x86_64 baseline
}

3. Core Operations (MVP)

3.1 Phase 1: Vector Operations

Target: Demonstrate SIMD/GPU/WASM parity

API Example:

use trueno::compute::Vector;

let a = Vector::from_slice(&[1.0f32; 1024]);
let b = Vector::from_slice(&[2.0f32; 1024]);

// Auto-selects best backend (AVX2/GPU/WASM)
let result = a.add(&b)?;
assert_eq!(result[0], 3.0);

// Force specific backend (testing/benchmarking)
let result_avx2 = a.add_with_backend(&b, Backend::AVX2)?;
let result_gpu = a.add_with_backend(&b, Backend::GPU)?;

3.2 Phase 2: Matrix Operations

3.3 Phase 3: Image Processing