WebGPU for Mojo — pure bindings to wgpu-native, with RAII GPU objects and a 5-line compute facade.
Compute and graphics on every backend Vulkan / Metal / DX12 supports — from a single Mojo program.
![]() |
![]() |
![]() |
example-plasmafullscreen fragment shader |
example-metaballs2D signed distance fields |
example-raymarchraymarched 3D SDF |
- Pure Mojo — no Python in the hot path.
from wgpu import …resolves at compile time. - RAII by default — every GPU object frees its native handle when it goes out of scope.
- Two altitudes — a low-level API that maps 1:1 to
webgpu.h, and aGPUfacade that does compile + dispatch + read-back in a handful of lines. - Windows too —
RenderCanvaswraps GLFW so the same device drives an interactive surface. - Verified downstream — CI installs the package into a fresh project on every push, so the quickstart below is exactly what runs in CI.
Use wgpu-mojo as a package in your own pixi project — three steps.
# pixi.toml
[workspace]
channels = ["https://conda.modular.com/max-nightly", "conda-forge"]
preview = ["pixi-build"]pixi add --git https://github.com/Hundo1018/wgpu-mojo wgpu-mojoThis builds wgpu.mojopkg and installs it into your environment.
The package needs libwgpu_native + its callback bridge at runtime. Run this inside your activated pixi env:
curl -fsSL https://raw.githubusercontent.com/Hundo1018/wgpu-mojo/main/scripts/setup-native.sh | bashWhat the script does
- Downloads
libwgpu_nativev29 from wgpu-native releases - Compiles the Mojo callback bridge
libwgpu_mojo_cbfrom source - Installs both into
$CONDA_PREFIX/lib/ - Compiles the GLFW window bridge
libglfw_input_cbif GLFW is present - Requires
curl,unzip,gcc— all standard in a conda environment
pixi run mojo -c 'from wgpu import Instance; _ = Instance(); print("wgpu-mojo OK")'Vector addition on the GPU, read back to the CPU — using the high-level GPU facade
(no manual bind groups, pipeline layouts, or lifetime pins):
from wgpu.gpu import GPU
from wgpu import WGPUBufferUsage
comptime N = 1024
comptime ADD_WGSL = """
@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> c : array<f32>;
@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
let i = gid.x;
if i < arrayLength(&a) { c[i] = a[i] + b[i]; }
}
"""
def main() raises:
var gpu = GPU.wgpu()
var a = gpu.buffer[Float32](N, WGPUBufferUsage.STORAGE | WGPUBufferUsage.COPY_DST)
var b = gpu.buffer[Float32](N, WGPUBufferUsage.STORAGE | WGPUBufferUsage.COPY_DST)
var c = gpu.buffer[Float32](N, WGPUBufferUsage.STORAGE | WGPUBufferUsage.COPY_SRC)
var xs = List[Float32](capacity=N)
var ys = List[Float32](capacity=N)
for i in range(N):
xs.append(Float32(i))
ys.append(Float32(i) * 2.0)
gpu.write(a, xs)
gpu.write(b, ys)
var prog = gpu.compile_compute(ADD_WGSL, entry_point="main", n_storage_buffers=3)
gpu.dispatch(prog^, [a.handle(), b.handle(), c.handle()], N // 64)
var result = gpu.read[Float32](c)
print(result[0], result[1], result[2]) # 0.0 3.0 6.0Full source: examples/compute_add_v2.mojo · runs with pixi run example-compute-v2.
| Run | Window? | Shows |
|---|---|---|
pixi run example-compute-v2 |
– | Vector add via the GPU facade (start here) |
pixi run example-compute |
– | Same, spelled out with the low-level API |
pixi run example-enumerate |
– | List every GPU adapter / backend |
pixi run example-clear |
✓ | Cornflower-blue window — minimal GPU smoke test |
pixi run hello · example-triangle |
✓ | The classic RGB "hello triangle" |
pixi run example-plasma |
✓ | Fragment-shader host — edit one WGSL shade function to make art |
pixi run example-metaballs |
✓ | 2D signed-distance fields (metaballs) |
pixi run example-raymarch |
✓ | Raymarched 3D SDF scene (camera orbit + lighting) |
pixi run example-texture-sample |
✓ | Texture + sampler on a fullscreen quad |
pixi run example-fire-sim |
✓ | Doom-style fire: compute + render ping-pong |
pixi run example-input |
✓ | GLFW keyboard / mouse polling |
examples/plasma.mojo is a fullscreen fragment-shader host that
hands your WGSL a small set of uniforms — resolution, time, mouse, frame. Copy it,
edit only the shade function, and climb the ladder:
plasma → 2D SDF → raymarching.
The three GIFs above are rendered headlessly on the GPU and read back frame by frame (
pixi run render-gifs) — not screen-recorded, because the wgpu surface presents via a GPU flip that screen grabbers can't capture.
fn shade(frag_coord: vec2<f32>) -> vec3<f32> {
let uv = frag_coord / U.resolution.xy;
return vec3<f32>(uv, 0.5 + 0.5 * sin(U.time));
}from wgpu.instance import Instance
from wgpu._ffi.structs import WGPUColor
from wgpu.rendercanvas import RenderCanvas
def main() raises:
var instance = Instance()
var adapter = instance.request_adapter()
var device = adapter.request_device()
var canvas = RenderCanvas(adapter, device, 800, 600, "hello triangle")
var shader = device.create_shader_module_wgsl(TRIANGLE_WGSL, "tri")
var pl = device.create_pipeline_layout(List[OpaquePointer[MutUntrackedOrigin]](), "layout")
var pipeline = device.create_render_pipeline(
shader, "vs_main", "fs_main", canvas.surface_format(), pl,
primitive_topology=UInt32(4), # TriangleStrip
)
while canvas.is_open():
canvas.poll()
var frame = canvas.next_frame()
if not frame.is_renderable():
continue
var enc = device.create_command_encoder("frame")
var rpass = enc.begin_surface_clear_pass(
frame.texture,
WGPUColor(Float64(0), Float64(0), Float64(0), Float64(1)),
"pass",
)
rpass.set_pipeline(pipeline)
rpass.draw(UInt32(3), UInt32(1), UInt32(0), UInt32(0))
rpass^.end()
device.queue_submit(enc^.finish())
canvas.present()Full source (with the WGSL): examples/triangle_window.mojo.
Every type is re-exported from wgpu, so from wgpu import Instance always works.
| Module | Provides |
|---|---|
wgpu.gpu |
GPU — high-level facade: buffer, write, compile_compute, dispatch, read |
wgpu.instance |
Instance — entry point, adapter selection |
wgpu.adapter |
Adapter — device creation |
wgpu.device |
Device — factory for every GPU object, submits work |
wgpu.buffer |
Buffer — GPU memory, mapping, typed read-back |
wgpu.texture |
Texture, TextureView |
wgpu.shader |
ShaderModule — WGSL compilation |
wgpu.pipeline |
ComputePipeline, RenderPipeline |
wgpu.command |
CommandEncoder, CommandBuffer |
wgpu.compute_pass · wgpu.render_pass |
ComputePassEncoder, RenderPassEncoder |
wgpu.bind_group |
BindGroup, BindGroupLayout (+ BGL entry helpers) |
wgpu.rendercanvas |
RenderCanvas — GLFW window + surface |
wgpu.diagnostics |
preflight() — library load status + adapter list |
Wrappers are RAII — GPU objects free themselves at end of scope. Two rules to remember
(both handled for you by the GPU facade):
# 1 — encoders must be explicitly finished:
var enc = device.create_command_encoder("enc")
var cpass = enc.begin_compute_pass("pass")
cpass^.end() # consume the pass
device.queue_submit(enc^.finish())
# 2 — pin resources that must outlive an async GPU call:
device.queue_submit(cmd)
_ = pipeline^ # prevent ASAP-drop before the GPU finishes
_ = bind_group^
device.poll(True)git clone https://github.com/Hundo1018/wgpu-mojo
cd wgpu-mojo
pixi run build-callbacks # download wgpu-native + compile the C bridges
pixi run test # non-GPU unit tests (no hardware needed)
pixi run check-compile # compile-check every test & example
pixi run hello # GPU smoke test — RGB triangle windowbuild-callbacks does the same job as setup-native.sh, but reads the native version from
ffi/wgpu-native-meta/wgpu-native-git-tag and
uses the bridge sources already in the repo. See docs/CI_CD.md for the
full CI/release pipeline.
| Platform | GPU driver |
|---|---|
Linux (linux-64) |
Vulkan: mesa-vulkan-drivers + libvulkan1, or NVIDIA proprietary |
macOS (osx-arm64) |
Metal — built in, nothing to install |
| Windows | D3D12 or Vulkan — usually present with vendor drivers |
linux-64 and osx-arm64 are first-class via pixi. osx-x86_64 / win-x64 can be built
manually — see conda.recipe/recipe.yaml.
from wgpu.diagnostics import preflight
print(preflight()) # search paths, load status, wgpu-native version, adapters

