Compute Kernels

GPU compute lets scripts run a function across thousands of threads in parallel – typical uses are heightmap textures, particle simulations, vertex displacement on a dynamic mesh, or "how many triangles did marching cubes emit this frame" atomic counters.

Apple GPUs only, for now

Compute kernels are currently available on Apple GPUs only – web / WebXR support is planned but not yet available. Today you author the kernel directly in Metal Shading Language (MSL), which makes this the one corner of the scripting API that isn't fully engine-agnostic. Always gate compute behind environment.features.has('compute') (Feature Detection) so a script degrades cleanly on devices – and platforms – that don't support it yet.

// Attach to: onStart

if (!environment.features.has('compute')) {
    console.warn('compute not supported – bailing');
    return;
}

const heightMap = await Texture.compute({
    pixelFormat: 'rgba8Unorm',
    width: 256, height: 256
});

const kernel = await Kernel.fromSource({
    source: `
        #include <metal_stdlib>
        using namespace metal;
        struct Uniforms { float time; };
        kernel void wave(constant Uniforms& U                       [[buffer(0)]],
                         texture2d<float, access::write> outTex     [[texture(0)]],
                         uint2 gid                                  [[thread_position_in_grid]])
        {
            float2 uv = float2(gid) / float2(outTex.get_width(), outTex.get_height());
            float r = 0.5 + 0.5 * sin(U.time + uv.x * 6.28);
            outTex.write(float4(r, uv.y, 0.5, 1.0), gid);
        }`,
    functionName: 'wave'
});

const uniforms = new Float32Array(1);
scene.on('render', function () {
    uniforms[0] = scene.time;
    kernel.run({
        uniforms: uniforms,
        outputTextures: { outTex: heightMap.write },
        threadGroups: [16, 16, 1]
    });
});

Dispatch entry points

Two ways to dispatch a kernel:

• mesh.runCompute(kernel, options) – when the kernel writes mesh attribute buffers (storage: "compute" attributes) or the mesh's index buffer. The mesh sits in buffer(1...) slots.
• kernel.run(options) – standalone compute. No mesh involvement. Use for texture generation, particle simulation, atomic counters, anything that doesn't displace mesh geometry.

Both paths share the same options shape; kernel.run simply skips the mesh-related slots.

MSL slot ordering – strict and positional

Author-side bindings map to MSL [[buffer(N)]] / [[texture(N)]] indices in a fixed order:

Slot	Source	Notes
buffer(0)	uniforms: Float32Array	Always slot 0, packed as raw bytes
buffer(1..A)	outputs[] (mesh attributes)	Only for mesh.runCompute. Order matches the array – outputs: ['position', 'normal'] binds position at 1, normal at 2
buffer(A+1)	mesh index buffer	Only when outputIndices: true and the mesh has indexCapacity > 0
buffer(A+2..)	inputBuffers then outputBuffers	Both maps merged into one continuous block, inputs first
texture(0..)	inputTextures then outputTextures	Separate slot space from buffers, inputs first

Match your MSL signature to this order. Slot indices the kernel declares ([[buffer(N)]]) must match what the runtime binds.

Common pitfalls

"No outputs declared" warning

kernel.run: no outputs declared (outputs / outputIndices / outputBuffers / outputTextures all empty).
Kernel will run but nothing it writes is observable – likely a bug.

A dispatch without any write target runs the kernel but has no observable effect – the GPU reads inputs, computes values, then throws them away. Almost always a missed wiring. Add the destination map you meant to (outputBuffers: { triCount: counter }, outputTextures: { outTex: heightMap.write }, etc.).

Slot ordering off by one

MSL kernels declare slot indices by hand (constant Uniforms& U [[buffer(0)]]). If you change the JS dispatch options – add an attribute to outputs, toggle outputIndices, add an inputBuffer – every later slot index shifts. Always re-check the MSL signature when you change the dispatch options.

Texture token swap (.read vs .write)

A texture handle has two binding tokens: .read (sample inside the kernel) and .write (write to inside the kernel). They route to different Metal APIs under the hood. Passing .read where you meant .write (or vice versa) means the kernel binds the texture in the wrong access mode – usually a silent no-op for writes, or undefined sampling for reads. Token name = what the kernel does with it.

Uniforms must be a Float32Array

uniforms is packed as raw bytes. The runtime only accepts a Float32Array. Plain JS arrays and other typed arrays trigger a warn and skip. To pack mixed types (a float and an integer index, say), reinterpret memory inside MSL or use an inputBuffers entry instead.

outputs only valid with a mesh

outputs: [...] and outputIndices: true are mesh-attribute writes. They mean nothing for kernel.run(...) (no mesh) and are silently ignored there. If you want raw output, use outputBuffers: { ... } with a Buffer handle.

Forgetting await on factory promises

Texture.compute(...), Kernel.fromSource(...), and Kernel.fromAsset(...) all return Promises. Without await, downstream code touches a Promise object – kernel.run(...) fails, texture.write is undefined. The _promisify wrapper logs nothing helpful here; just remember the await.

Buffer size mismatches on writeSync

Buffer.writeSync: undersize write – source has 2 elements, buffer holds 4.
Copying 2 elements; tail of buffer UNTOUCHED.

The warn is loud on purpose. Seeding partial data is usually a bug – match src.length to the buffer's declared element count.

Capability gate

Always gate compute calls with environment.features.has('compute') – devices below the Apple 7 GPU family or below iOS 18 / visionOS 2 / macOS 15 return false. The Swift bridge throws cleanly if you slip past, but the JS-side gate keeps your script structure honest.