GPUWL

English · Português · Español

GPU compute for Luau, on any vendor. The library is Rust (wgpu + mlua); you write kernels in WGSL and drive them from Luau. One kernel runs on AMD, Intel, NVIDIA and Apple hardware without a line changing.

local dev = gpu.open()
local a = dev:buffer("f32", { 1, 2, 3 })
local b = dev:buffer("f32", { 10, 20, 30 })
local out = dev:buffer("f32", 3)

dev:kernel([[
@group(0) @binding(0) var<storage, read>       a: array<f32>;
@group(0) @binding(1) var<storage, read>       b: array<f32>;
@group(0) @binding(2) var<storage, read_write> out: array<f32>;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) id: vec3<u32>) {
	let i = id.x;
	if (i < arrayLength(&out)) { out[i] = a[i] + b[i]; }
}
]]):run({ a, b, out }, gpu.groups(3, 64))

print(table.concat(out:read(), ", ")) --> 11, 22, 33

---

Why not CUDA

CUDA is NVIDIA-only. There is no CUDA on an AMD or Intel card, so "CUDA for every GPU" cannot exist — the name belongs to one vendor's proprietary stack.

The portable equivalent is a compute shader, and that is what this uses. You write one WGSL kernel; wgpu translates it at runtime to whatever the machine actually has:

Backend	Where
Vulkan	AMD, Intel, NVIDIA — Windows and Linux
DirectX 12	Windows
Metal	macOS, iOS
OpenGL	fallback for older drivers

If you specifically need CUDA-only features — cuBLAS, cuDNN, tensor-core intrinsics, NVIDIA's profiler — use CUDA directly. This library is not that, and does not pretend to be.

---

Requirements

• Rust 1.85 or newer (2024 edition).
• A C++ compiler, because mlua builds Luau from source.
  • Windows: Visual Studio with the Desktop development with C++ workload.
  • Linux: build-essential (or your distro's gcc/g++).
  • macOS: Xcode command line tools.
• A GPU with a working driver. Any vendor. Software fallbacks (like Windows' Basic Render Driver) also work, just slowly.

Build

cargo build --release

This produces target/release/gpuwl, a runner that executes a Luau script with the gpu table already in scope — no require needed.

./target/release/gpuwl examples/02_vector_add.luau

On Windows:

.\target\release\gpuwl.exe examples\02_vector_add.luau

---

The mental model

Four ideas cover everything this library does.

1. A device is one GPU. gpu.open() picks the best one available (discrete over integrated over software). gpu.adapters() lists them all if you want to choose.

2. A buffer is a flat array of 32-bit numbers living in GPU memory. You declare its type — "f32", "u32" or "i32" — and either upload a Luau table or ask for n zeroed elements. Luau numbers are 64-bit floats, so values are converted on the way in and out.

3. A kernel is a WGSL function that runs once per invocation. Buffers are bound to @group(0) @binding(0), @binding(1), … in the order you pass them to run. Get that order wrong and you will read the wrong array — nothing else warns you.

4. Invocations come in workgroups. @workgroup_size(64) means 64 invocations per group. kernel:run(buffers, 10) dispatches 10 groups, so 640 invocations total. To cover n items you dispatch gpu.groups(n, 64) groups — which rounds up, so the last group usually runs past the end of your data. Every kernel needs a bounds guard:

let i = id.x;
if (i >= arrayLength(&out)) { return; }   // without this, the tail writes out of range

---

API

Module

Call	Returns
gpu.adapters()	array of adapter info tables (see below)
gpu.open(opts?)	a device. Without options: discrete > integrated > anything
gpu.groups(n, size)	ceil(n / size) — the workgroup count covering n items

gpu.open accepts any combination of:

gpu.open({ index = 2 })            -- 1-based, into the gpu.adapters() list
gpu.open({ name = "radeon" })      -- case-insensitive substring of the adapter name
gpu.open({ backend = "vulkan" })   -- "vulkan" | "dx12" | "metal" | "gl"

The environment variable WGPU_BACKEND overrides backend selection without touching code.

Device

Member	Meaning
dev.name	adapter name, e.g. "NVIDIA GeForce RTX 4070"
dev.backend	"vulkan", "dx12", "metal" or "gl"
dev.kind	"discrete", "integrated", "virtual", "cpu" or "other"
dev.info	table: name, backend, kind, vendor, device, driver, driver_info
dev.limits	table, see below
dev:buffer(kind, t)	kind is "f32", "u32" or "i32". t is a table to upload, or a positive number to allocate that many zeroed elements
dev:kernel(wgsl, entry?)	compiles a pipeline. entry defaults to "main"
dev:sync()	blocks until all queued work finished — the sync point for timing

dev.limits holds what the hardware will accept:

Field	Meaning
max_workgroups_per_dim	largest x, y or z you can pass to run
max_workgroup_size_x/y/z	largest @workgroup_size(...) per axis
max_invocations_per_workgroup	cap on x * y * z inside one workgroup
max_storage_buffers	how many buffers one kernel can bind
max_storage_buffer_bytes	largest single binding
max_buffer_bytes	largest allocation

Buffer

Member	Meaning
buf.len	element count
buf.kind	"f32", "u32" or "i32"
buf.bytes	len * 4
buf:read()	downloads and returns a Luau array. Blocks
buf:write(t)	uploads a table. May be shorter than the buffer, never longer

Kernel

Call	Meaning
kernel:run(buffers, x, y?, z?)	dispatches x * y * z workgroups. y and z default to 1

buffers is an array; element i binds to @binding(i-1). All buffers must come from the device that compiled the kernel.

---

Writing kernels

Bindings are always storage buffers in group 0. The bind group layout is inferred from your shader, so read and read_write both work exactly as you declare them:

@group(0) @binding(0) var<storage, read>       input:  array<f32>;
@group(0) @binding(1) var<storage, read_write> output: array<u32>;

Useful builtins:

@builtin(global_invocation_id) id: vec3<u32>	this invocation's index across the whole dispatch
@builtin(local_invocation_id)	index inside its workgroup
@builtin(workgroup_id)	which workgroup this is
arrayLength(&buf)	element count of a runtime-sized binding

There are no uniforms and no push constants. Pass scalars in a small buffer:

@group(0) @binding(3) var<storage, read> cfg: array<u32>;   // [width, height, iterations]

local cfg = dev:buffer("u32", { width, height, iterations })

Vector types work as long as the packing is tight — array<vec2<f32>> of n elements is just a buffer of 2n floats, which is how examples/05_nbody.luau stores positions.

self is a reserved word in WGSL. So are mut, ref, struct and a longer list; name a variable one of them and the shader fails to compile with a clear message.

---

Examples

Timings below were measured on an Intel Arc integrated GPU over Vulkan. Yours will differ.

File	What it teaches	Result here
01_info.luau	enumerating adapters and reading limits	found 4 adapters: Arc over vulkan/dx12/gl, plus a software one
02_vector_add.luau	the smallest complete kernel, verified element by element against the CPU	1,000,000 elements in 4.3 ms
03_matmul.luau	passing scalars in a config buffer, 2D dispatch, honest GPU-vs-CPU timing	256²: 4.1 ms vs 1080 ms in Luau (262×). 1024²: 64 GFLOP/s
04_mandelbrot.luau	one invocation per pixel, drawn as ASCII	120×40 render
05_nbody.luau	many dispatches over data that never leaves the GPU, ping-pong buffers	8192 bodies × 200 steps in 677 ms, ~19.8 billion interactions/s

Run any of them:

./target/release/gpuwl examples/05_nbody.luau

---

Embedding in your own Rust program

The runner is thin; the library is the point. Any mlua host can take the same gpu table:

let lua = mlua::Lua::new();
let gpu = gpuwl::module(&lua).map_err(|e| anyhow::anyhow!("{e}"))?;
lua.globals().set("gpu", gpu).map_err(|e| anyhow::anyhow!("{e}"))?;
lua.load(script).exec().map_err(|e| anyhow::anyhow!("{e}"))?;

mlua::Error is neither Send nor Sync, which is why it needs flattening before anyhow will take it.

The wgpu layer is public too, if you want it without Lua: gpuwl::gpu::{open, Dev, Kind}.

---

Tests

cargo test

Four tests: one pure unit test on number encoding, three end-to-end on a real adapter (compile a kernel, dispatch it, read the result back; reject a broken shader; reject a zero dispatch). The end-to-end ones skip themselves when no adapter is available, so they are safe on a headless CI box.

---

Troubleshooting

Symptom	Cause
no GPU adapter found	no driver wgpu can use. On Linux install the Vulkan ICD for your card; on Windows update the display driver
name 'x' is a reserved keyword	WGSL reserves more words than you would expect. Rename the variable
Results are the wrong array	your run({...}) order does not match your @binding(n) order
Last few elements are garbage	missing the if (i >= arrayLength(&out)) { return; } guard
workgroup count N on axis x is out of range	n / workgroup_size exceeded dev.limits.max_workgroups_per_dim. Have each invocation process several items
Build fails on mlua-sys	no C++ compiler. See Requirements
Everything runs but slowly	check dev.kind — you may have opened a "cpu" software adapter

---

Limits

This is a small library and stays one on purpose.

• Storage buffers only — no textures, no samplers, no uniforms, no push constants.
• One bind group. Everything lands in group 0.
• Element types are 32-bit f32 / u32 / i32. No f16, no f64.
• read() allocates a staging buffer per call. Fine for results; not for per-frame readback.
• Dispatches are submitted immediately; there is no command batching across run calls.
• Synchronous throughout. Nothing is async, nothing is threaded.

License

MIT. See LICENSE.