# 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:

```loup,ignore
@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):

````loup,ignore
@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.