Chapter 23: SIMD and GPU

⚠ Planned — not in Loup 1.25. This chapter describes the design in the Language Reference (§1-2.6 and §1-16.8) and the Packages Reference (§21.29, §21.30). The toolchain does not implement it yet, so its listings are design sketches and are not verified.

The specification gives native programs two routes to data-parallel speed: explicit CPU vector types in native.simd, and GPU compute kernels written in a shading language inside a template fence and dispatched through native.gpu. Neither ships in 1.25.1. native.simd, native.gpu and native.webgpu are not standard-library modules ('native.simd' is not a standard-library module), Vec[T, N] is an undefined name, and a wgsl fence parses as an ordinary template fence that nothing can compile or dispatch.

What a native program can use today is the LLVM backend's own optimizer: a --release build compiles through clang -O2 with ThinLTO (chapter 24), which auto-vectorizes simple loops over lists of numbers without any annotation.

CPU SIMD: native.simd

The planned module centres on one generic value type, Vec[T, N]: N lanes of a fixed-width number type, mapped directly to a hardware vector register. sizeof(T) × N must be 16, 32 or 64 bytes (an SSE/NEON, AVX or AVX-512 register); other widths are compile errors. Vectors are immutable values. They are built with splat(x) (every lane x), of(...) or load(list, offset), written back with store(list, offset), combined lane-wise with + - * /, and reduced or transformed with functions such as dot, sqrt, min, max and fma. A dot product:

@mode(host.native)

from native.simd import (Vec, dot);

func dotProduct(a: list[float32], b: list[float32]): float32 {
    var sum: float32 = 0.0f32;
    var i = 0;
    while (i + 8 <= a.length()) {
        let va = Vec[float32, 8].load(a, i);
        let vb = Vec[float32, 8].load(b, i);
        sum = sum + dot(va, vb);
        i = i + 8;
    }
    while (i < a.length()) {
        sum = sum + a[i] * b[i];
        i = i + 1;
    }
    return sum;
}

Listing 23-1: A dot product with eight-lane vectors (design sketch)

The tail loop handles the elements that do not fill a whole vector. The @vectorize annotation is specified as a hint asking the compiler to vectorize a loop it would otherwise leave scalar.

SIMD is a host.native and freestanding feature. Using native.simd, Vec[T, N] or @vectorize in a host.dual file is specified as E_NATIVE_ONLY_IN_DUAL, because a host.dual program must also run as WebAssembly (§1-2.6).

GPU compute: shader fences

Three template-fence languages hold GPU shader source: wgsl (WebGPU), metal (Apple platforms only) and glsl (§1-16.8). A shader fence evaluates to the shader text; {$ … $} interpolations may only insert compile-time constants (int, float32, bool), and a runtime value is E_GPU_FENCE_RUNTIME. native.gpu compiles the text into a kernel and runs it over buffers (§21.30):

@mode(host.native)

from native.gpu import (GpuKernel, GpuBuffer, dispatch);

const WGSIZE: int = 256;

async func addOnGpu(a: list[float32], b: list[float32]): Task[list[float32]] {
    let kernel = GpuKernel.compile(```(wgsl):(none)
        @group(0) @binding(0) var<storage, read_write> out: array<f32>;
        @group(0) @binding(1) var<storage, read>       x:   array<f32>;
        @group(0) @binding(2) var<storage, read>       y:   array<f32>;

        @compute @workgroup_size({$ WGSIZE $}, 1, 1)
        fn main(@builtin(global_invocation_id) gid: vec3u) {
            let i = gid.x;
            if (i < arrayLength(&out)) {
                out[i] = x[i] + y[i];
            }
        }
    ```);

    let outBuf = GpuBuffer[float32].alloc(a.length());
    let xBuf = GpuBuffer[float32].alloc(a.length());
    let yBuf = GpuBuffer[float32].alloc(b.length());
    await xBuf.copyFromHost(a);
    await yBuf.copyFromHost(b);

    let groups = (a.length() + WGSIZE - 1) / WGSIZE;
    await dispatch(kernel, [groups], [outBuf, xBuf, yBuf]);
    return await outBuf.copyToHost();
}

Listing 23-2: Adding two vectors on the GPU (design sketch)

Buffers are bound in order (binding(0), binding(1), …). Copies and dispatch are async because the GPU runs independently of the CPU, so GPU code also waits on the native async executor (chapter 21). The backend is chosen at build time (--gpu-backend wgpu|cuda|metal|vulkan, default wgpu) or at run time. The same wgsl text can also feed a render pipeline through native.webgpu (§21.29), which is the layer meant to be shared with browser code.

Structs shared with a shader need an agreed layout. @repr(C) fixes a Loup struct's layout (chapter 19); the shader's struct must then follow WGSL's own alignment rules, which pad a vec3<f32> to 16 bytes, so the Loup side adds explicit padding fields.

When the design lands

The usual guidance applies. GPU compute pays off for thousands to millions of independent, arithmetic-heavy elements whose data stays on the GPU between steps; CPU SIMD wins for smaller inputs, branchy code, and work that has to return to the CPU often. Before either, write the straightforward loop and measure a --release build, since the optimizer may already vectorize it.

Summary

• native.simd (Vec[T, N], lane-wise operators, dot/sqrt/fma, @vectorize) is specified for host.native and freestanding and is ⚠ planned.
• GPU kernels are specified as wgsl, metal or glsl fences compiled by native.gpu.GpuKernel.compile and run with dispatch; native.gpu and native.webgpu are ⚠ planned, and depend on native async.
• Today, rely on the LLVM backend's auto-vectorization in --release builds.
• Chapter 24 covers building and distributing native programs.