Rake

What Rust does for safety with unsafe {}, Rake does for speed with slow {}.

Rust's type system and borrow checker enforce memory safety in safe code, and inside an unsafe block the programmer is responsible for it. Rake's compiler checks that vector calculations become vector instructions, reviewing the generated assembly for you, and compilation fails if it can't keep a calculation vectorised. A slow { ... } block makes an explicit place for scalar work, and vector code resumes after the closing brace.

Rake is a SIMD, or vector, programming language for calculations that apply the same operation to many numbers at once. Its current profiles target CPUs and WebAssembly. The planned GPU profiles will preserve parallel work across warp lanes, put matrix products on the tensor cores in their declared formats, and compile each fused region to one kernel, or fail to compile.

• Documentation
• Tutorial
• Compiler on GitHub

Checked vectorisation

Suppose we want the square roots of a column of numbers, with zero for negative inputs. An ordinary C loop is easy to read:

#include <math.h>
#include <stddef.h>

void safe_root_c(const float *values, float *roots, size_t count)
{
    for (size_t i = 0; i < count; ++i) {
        if (values[i] >= 0.0f)
            roots[i] = sqrtf(values[i]);
        else
            roots[i] = 0.0f;
    }
}

GCC 14.3 turns this into a vector loop for AVX2. A similar loop with a different maths function fares worse:

void sin_array_scalar(const float *input, float *output, size_t count)
{
    for (size_t i = 0; i < count; ++i)
        output[i] = sinf(input[i]);
}

C still compiles it, but in this build the loop calls the scalar sinf library function once per element. The build succeeds anyway, so only vectorisation diagnostics or the assembly would show that the work became scalar.

Rake makes the vector requirement part of compilation:

scratch reject_sin(values: f32s) -> f32s:
  | result <| sin(values)
  result

$ rakec --verify-native --target x86-avx2 demo/soa-proof/rake/reject_sin.rk
native lowering: call to 'sin' is not supported by native scratch lowering

The profile has no vector lowering for this operation, so Rake stops at the line responsible, before the scalar loop reaches the application.

Explicit SIMD

The usual next step is to write the vector operations explicitly. For the square-root loop, C's AVX2 intrinsics expose the register width, lane mask, loads and stores, and the final partial vector needs its own memory mask:

#include <immintrin.h>
#include <stddef.h>

static __m256 safe_root_rack(__m256 value)
{
    const __m256 zero = _mm256_setzero_ps();
    const __m256 valid = _mm256_cmp_ps(value, zero, _CMP_GE_OQ);
    const __m256 input = _mm256_and_ps(valid, value);
    return _mm256_and_ps(valid, _mm256_sqrt_ps(input));
}

void safe_root_avx2(const float *values, float *roots, size_t count)
{
    size_t i = 0;
    for (; count - i >= 8; i += 8) {
        const __m256 value = _mm256_loadu_ps(values + i);
        _mm256_storeu_ps(roots + i, safe_root_rack(value));
    }
    if (i < count) {
        const __m256i lanes = _mm256_setr_epi32(0, 1, 2, 3, 4, 5, 6, 7);
        const __m256i remaining = _mm256_set1_epi32((int)(count - i));
        const __m256i active = _mm256_cmpgt_epi32(remaining, lanes);
        const __m256 value = _mm256_maskload_ps(values + i, active);
        _mm256_maskstore_ps(roots + i, active, safe_root_rack(value));
    }
}

ISPC supplies a parallel loop and manages its lanes, while its C-family syntax keeps the array indices and statements:

export void safe_roots(uniform float values[],
                       uniform float roots[],
                       uniform int count)
{
    foreach (i = 0 ... count) {
        float value = values[i];
        roots[i] = value >= 0.0f ? sqrt(value) : 0.0f;
    }
}

In Rake, we describe the eligible lanes and their calculation. The profile chooses the width, and the compiler handles the loads, stores and tail:

pack Samples {
  f32: value;
}

tine #valid(values: f32s) means values >= <0.0>

rake safe_root(values: f32s) -> f32s:
  through #valid(values) into rooted:
    sqrt(values)
  sweep:
    | #valid(values) => rooted
    | #valid(values) gaps => <0.0>

run roots(input: stack Samples, output: stack Samples) -> stack Samples:
  output with { value: safe_root(input.value) }

A polynomial sine

Suppose our algorithm only needs the sine of angles between −0.5 and 0.5 radians. A ninth-degree polynomial over that interval can replace the general library call, and because it needs only multiplies and adds, it passes Rake's vector checks:

scratch bounded_sine(values: f32s) -> f32s:
  let square = values * values
  let fifth = <0.0000027557319223985893> * square - <0.0001984126984126984>
  let fourth = fifth * square + <0.008333333333333333>
  let third = fourth * square - <0.16666666666666666>
  let second = third * square + <1.0>
  values * second

Measured execution

These bars return to the square-root example, with all four versions processing the same million-element input. C's automatic vectoriser already handles this loop, so its two vector versions and Rake finish at similar times.

The benchmark source, flags and raw samples include bitwise result comparisons, guard-page tail checks and disassemblies.

The programming model

Rake has no branches inside a rack. Its execution model has three steps: define the tines, rake data through them, then sweep the results.

One instruction, several numbers

Picture a person at a desk with a long tape passing over it. In a scalar loop, they add 1 to the number in front of them, then advance the tape and repeat. A vector operation gives them a wider desk with several numbers side by side, and linked pens apply the one instruction to all of them at once. This is SIMD, single instruction, multiple data, and each position is a lane. A processor holds the values in a vector register, and its width sets how many fit: a 128-bit register holds four 32-bit numbers, and a 256-bit register holds eight. Rake calls the row of values held together a rack.

Tines

This rake takes the square root of each nonnegative number, and for a negative number it takes the root of the magnitude and keeps the sign, so 16 becomes 4 and −4 becomes −2:

tine #nonnegative(values: f32s) means values >= <0.0>
tine #negative(values: f32s) means values < <0.0>

rake signed_root(values: f32s) -> f32s:
  through #nonnegative(values) into positive_roots:
    sqrt(values)

  through #negative(values) into negative_roots:
    -sqrt(-values)

  sweep:
    | #nonnegative(values) => positive_roots
    | #negative(values) => negative_roots
    | (#nonnegative(values) or #negative(values)) gaps => <0.0>

A tine, named after the prongs of a garden rake, is a labelled mask: a true-or-false choice for each lane. Read the first line as “the nonnegative tine catches numbers that are greater than or equal to zero”. <0.0> is a uniform, one scalar value shared by every lane, and the angle brackets mark every point where a scalar becomes a rack, so the cost of broadcasting is always visible:

values:        [16,    -4,     9,    -1]
#nonnegative: [true, false,  true, false]
#negative:    [false, true, false,  true]

Tines are defined outside the rake, so another rake can reuse them with different data. They store no global mask: applying #nonnegative(values) computes the mask for that particular input.

Rakes and through blocks

The rake keyword tells the compiler that signed_root chooses work for individual lanes. Its through blocks supply that work, and its sweep chooses the result. Read through #nonnegative(values) into positive_roots: as “calculate the nonnegative lanes using the body below, and bind those results to positive_roots”:

values:         [16, -4, 9, -1]
positive_roots: [ 4,  ·, 3,  ·]
negative_roots: [ ·, -2, ·, -1]

The dots are undefined lanes, and the compiler checks every later use of positive_roots against the mask where it's defined. The binding doesn't need a temporary array, so the value can stay in a vector register.

Sweeps

A sweep chooses the values that leave the function, trying its arms in order. The first two take each lane's root from the through block that computed it. The last arm gives zero to the lanes in neither mask, found by gaps, an exact mask inversion, so a NaN that passed neither comparison still gets a value. The result is [4, -2, 3, -1], with each value still in its original lane, and because sweep: is the rake's result form, it needs no return keyword.

Reading positive_roots in a negative lane would be a compile error. If a later calculation needs a whole intermediate rack, a through fallback such as else <0.0> defines every lane. A final _ arm catches unmatched lanes, and it's optional when the masks provably cover every lane, as these tines and their gaps do. Lanes outside a mask can't fail or raise a floating-point exception, because physical CPU profiles give inactive lanes benign operands and WebAssembly SIMD has no floating-point exception flags.

Lesson 8 lets you try an undefined-lane read, then add a fallback and inspect the result. Tines, through and sweeps defines the scope and masking rules.

Characteristics

Fused stages

Fusion is the general compiler term for combining pieces of work so they execute together. In | name <| expression, the expression flows from right to left into its name. It's still a binding, but the leading | marks it as a fused stage, and the compiler keeps consecutive stages together as one contiguous sequence of vector instructions. Here a scratch, a function of racks that doesn't need tines, calculates a new position in two stages:

scratch advance(positions: f32s, velocities: f32s, <dt: f32>) -> f32s:
  | step  <| velocities * <dt>
  | moved <| positions + step
  moved

The names keep the calculation legible, while on AVX2 the code becomes one vbroadcastss to share dt across the lanes and one vfmadd231ps, a fused multiply-add that rounds the combined calculation once. Strict WebAssembly SIMD has no such instruction, so that target keeps the multiply and add separate. Fused bindings defines which rewrites Rake permits.

Packs, racks and stacks

A pack describes one record, here a particle with a position, velocity and age:

pack Particles {
  f32: position, velocity;
  u8: age;
}

A stack stores many packs as columns, so 600 particles become 600 positions together, 600 velocities together and 600 ages together, and each column splits into racks at the target's SIMD width. Remember the hierarchy as "define our pack, then rack 'em and stack 'em".

A run applies scratches and rakes to whole stacks. This one feeds a column of positions into signed_root and returns the stack with the results:

pack Positions {
  f32: value;
}

run roots(positions: stack Positions) -> stack Positions:
  positions with { value: signed_root(positions.value) }

The stack carries its count, which needn't be a multiple of the rack width, and the final rack loads and stores only the records that exist. Calling the run moves the stack into it, so the run updates the column in place. A run can also choose which records to touch: where #young after its result would update only the young particles, while compact keeps only them, moved to the front of each column, and shrinks the count to match:

pack Particles {
  f32: position, velocity;
  u8: age;
}

run young_only(particles: stack Particles) -> stack Particles:
  let age = widen(particles.age)
  tine #young means age < <2>
  compact particles where #young

Which is faster depends on how many particles survive, so the source chooses.

Slow code

Programs need setup, records, state and calls to C as well as kernels. That code is marked slow, and it can't hold a rack. Inside a run, slow { ... } is a scoped escape for scalar work, and its scalar result becomes a rack only through an explicit broadcast:

run shift(values: []i32, out: mut []i32, <steps: i32>):
  let rack = values[<0>]
  let <offset: i32> = slow {
    total: i32 := 0
    for index from 0 up to steps:
      total <- total + index
    total
  }
  out[<0>] <- rack + <offset>

The loop computes a scalar offset, and after } one vector add applies it to every lane. The slow tier defines the boundary.

Verified output

rakec --verify-native disassembles what it built and checks it against the profile's rules: no hidden calls or stack work, one physical register per rack on physical targets, and only the permitted WebAssembly virtual instructions on wasm-simd128.

Notation

Most of Rake reads like any expression language. These marks are its own:

~~ evokes furrows in raked sand. Rake borrows let bindings, expression-based conditionals and (* ... *) block comments from OCaml, and indentation from Python, with a line ending in : opening a body. The tutorial explains each mark for readers who come from C or Python.

Targets

The native profiles also compile stack runs over f32, i32 and u32 columns, while general native runs are still work in progress. Primitives, operations, and targets lists what each profile compiles, the backend explains which parts of a program come from Rake and which use a platform C compiler, and the roadmap covers what comes next.

rakec --emit-c writes C for your own toolchain, embedding the kernel assembly on native targets and using SIMD intrinsics on WebAssembly. Rake can call C functions and exchange records, unions and typed callbacks across the C ABI, as the compiler command explains.

GPU execution: the design

On an NVIDIA GPU, the planned rack will span a warp of 32 threads, each holding one lane's value, and Rake will check that no hidden loop processes the rack serially inside one thread.

The first proposed profile, nvidia-ptx-sm120a, will emit inspectable PTX and build a cubin ahead of time with a pinned NVIDIA toolchain. Rake will verify that device artifact and load those exact bytes, with no unverified PTX JIT fallback.

The vectorisation rule carries over to the GPU's tensor cores, which multiply whole tiles of two matrices at once and do nothing else. A matrix product compiles to tensor-core instructions or the build fails, and slow { ... } marks one kept on the ordinary CUDA cores. A declared number format, from 4-bit NVFP4 to TF32, runs as itself: a format the profile has no instruction for is rejected rather than quietly widened. A fused region compiles to one kernel whose intermediate values stay on chip, so norms, activations and softmax don't pay for a round trip through device memory between each operation.

Some properties can be checked in the artifact, while others only show up when it runs:

The GPU design defines the proposed contract, including later SPIR-V/Vulkan and direct physical backends, and the comparison with ISPC and Bend 2 separates lane-parallel kernels from parallel fork–join tasks.

Projects using Rake

The Loong Game

An open source bot and toolkit for the 2026 UNSW Battlecode tournament. The Loong Game uses Rake to compile room-counting kernels to WebAssembly through C, at 437 points a turn against 435 for hand-written intrinsics.

• The choice: Rake for vector kernels
• Counting room faster: keeping it vectorised

Getting Rake

The compiler, rakec, is written in OCaml, and its repository explains how to build it in a pinned Nix development shell. The rakec command lists its modes and options. Rake is released under the MIT-SLOP licence, and it's pre-1.0, so the language and its binary boundaries may still change between versions.

Syntax highlighting

These Tree-sitter grammar packages provide parsing and highlighting for editor integrations. They don't include the compiler:

• PyPI · tree-sitter-rake
• npm · tree-sitter-rake
• Cargo · tree-sitter-rake