@perryts/webgpu

Native WebGPU bindings — backed by wgpu — for the Perry TypeScript-to-native compiler.

What this is

A Perry "native library" package: a Rust crate that exports extern "C" symbols the Perry compiler links into your TypeScript program. From your TS code you import { ... } from "@perryts/webgpu"; under the hood every call resolves to a direct call into the bundled staticlib — no Node addon, no IPC, no JSON round-trip on the hot path.

The point of WebGPU as a Perry binding is portability: shaders + pipelines authored against the browser's WebGPU spec run unmodified under Perry. Same WGSL, same descriptor objects, same async lifecycle.

This package contains:

• src/lib.rs — the Rust crate wrapping wgpu and exporting js_webgpu_*extern "C" symbols
• src/index.ts — the TypeScript surface (functions + spec descriptor types + runtime constants like GPUBufferUsage)
• Cargo.toml — staticlib build config consumed by the Perry linker
• package.json — the perry.nativeLibrary manifest (60 functions)

Install

bun add @perryts/webgpu
# or
npm install @perryts/webgpu

The package's package.json declares a perry.nativeLibrary block (see the manifest spec) which Perry's compiler reads at link time to discover the staticlib + extern "C" symbols. No post-install build step — Perry compiles the Rust crate as part of your project's build.

Quick start — compute shader hello world

The canonical compute-shader smoke test: dispatch a workgroup that doubles every element of an input buffer, then map the output buffer back and read the result.

import{requestAdapter,adapterRequestDevice,deviceCreateBuffer,deviceCreateShaderModule,deviceCreateBindGroupLayout,deviceCreatePipelineLayout,deviceCreateBindGroup,deviceCreateComputePipeline,deviceCreateCommandEncoder,commandEncoderBeginComputePass,commandEncoderCopyBufferToBuffer,commandEncoderFinish,computePassSetPipeline,computePassSetBindGroup,computePassDispatchWorkgroups,computePassEnd,queueSubmit,queueWriteBuffer,bufferMapAsync,bufferGetMappedRange,bufferUnmap,devicePoll,GPUBufferUsage,GPUShaderStage,GPUMapMode,}from"@perryts/webgpu";constadapter=awaitrequestAdapter();const{ device, queue }=awaitadapterRequestDevice(adapter);// Input data: [1, 2, 3, 4] as u32 little-endian.constinput=newUint8Array([1,0,0,0,2,0,0,0,3,0,0,0,4,0,0,0]);constinBuf=deviceCreateBuffer(device,{size: 16,usage: GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_DST});constoutBuf=deviceCreateBuffer(device,{size: 16,usage: GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC});conststagingBuf=deviceCreateBuffer(device,{size: 16,usage: GPUBufferUsage.COPY_DST|GPUBufferUsage.MAP_READ});queueWriteBuffer(queue,inBuf,0,input);constshader=deviceCreateShaderModule(device,` @group(0) @binding(0) var<storage, read> input : array<u32>; @group(0) @binding(1) var<storage, read_write> output : array<u32>; @compute @workgroup_size(4) fn main(@builtin(global_invocation_id) id: vec3u) { output[id.x] = input[id.x] * 2u; }`);constbgl=deviceCreateBindGroupLayout(device,{entries: [{binding: 0,visibility: GPUShaderStage.COMPUTE,buffer: {type: "read-only-storage"}},{binding: 1,visibility: GPUShaderStage.COMPUTE,buffer: {type: "storage"}},],});constbg=deviceCreateBindGroup(device,{layout: bgl,entries: [{binding: 0,resource: {buffer: inBuf}},{binding: 1,resource: {buffer: outBuf}},],});constpipeline=deviceCreateComputePipeline(device,{layout: deviceCreatePipelineLayout(device,{bindGroupLayouts: [bgl]}),compute: {module: shader,entryPoint: "main"},});constenc=deviceCreateCommandEncoder(device);constpass=commandEncoderBeginComputePass(enc);computePassSetPipeline(pass,pipeline);computePassSetBindGroup(pass,0,bg);computePassDispatchWorkgroups(pass,1);computePassEnd(pass);commandEncoderCopyBufferToBuffer(enc,outBuf,0,stagingBuf,0,16);queueSubmit(queue,JSON.stringify([commandEncoderFinish(enc)]));awaitbufferMapAsync(stagingBuf,GPUMapMode.READ);devicePoll(device);// ← spec-extra: native runtime needs the explicit pollconstbytes=bufferGetMappedRange(stagingBuf);console.log(newUint32Array(bytes.buffer,bytes.byteOffset,4));// → Uint32Array(4) [ 2, 4, 6, 8 ]bufferUnmap(stagingBuf);

Render-pipeline sketch

The render path mirrors the spec one-for-one — same descriptor shapes, same WGSL, same setVertexBuffer / draw rhythm. Sampler + textureView bind-group entries use a {sampler:n} / {textureView:n} wrapping so the FFI parser can disambiguate them from buffer bindings:

consttex=deviceCreateTexture(device,{size: {width: 256,height: 256},format: "rgba8unorm",usage: GPUTextureUsage.TEXTURE_BINDING|GPUTextureUsage.COPY_DST,});constview=textureCreateView(tex);// default viewconstsampler=deviceCreateSampler(device,{magFilter: "linear",minFilter: "linear"});constbg=deviceCreateBindGroup(device,{layout: bgl,entries: [{binding: 0,resource: { sampler }},// ← sampler-flavoured resource{binding: 1,resource: {textureView: view}},// ← textureView-flavoured{binding: 2,resource: {buffer: uniforms}},],});constpipeline=deviceCreateRenderPipeline(device,{layout: pipelineLayout,vertex: {module: shader,entryPoint: "vs_main",buffers: [{arrayStride: 32,stepMode: "vertex",attributes: [{format: "float32x3",offset: 0,shaderLocation: 0},{format: "float32x2",offset: 12,shaderLocation: 1},]},],},fragment: {module: shader,entryPoint: "fs_main",targets: [{format: "bgra8unorm"}]},primitive: {topology: "triangle-list",cullMode: "back"},depthStencil: {format: "depth32float",depthWriteEnabled: true,depthCompare: "less"},multisample: {count: 1},});constenc=deviceCreateCommandEncoder(device);constpass=commandEncoderBeginRenderPass(enc,{colorAttachments: [{view: targetView,loadOp: "clear",storeOp: "store",clearValue: {r: 0,g: 0,b: 0,a: 1}}],depthStencilAttachment: {view: depthView,depthLoadOp: "clear",depthStoreOp: "store",depthClearValue: 1.0},});renderPassSetPipeline(pass,pipeline);renderPassSetBindGroup(pass,0,bg);renderPassSetVertexBuffer(pass,0,vbuf);renderPassSetIndexBuffer(pass,ibuf,"uint16");renderPassDrawIndexed(pass,indexCount);renderPassEnd(pass);queueSubmit(queue,JSON.stringify([commandEncoderFinish(enc)]));

What's a "spec-faithful" binding

WebGPU is a W3C spec. The wgpu crate is the canonical native implementation of that spec. This package is a direct binding to wgpu — every js_webgpu_* FFI function is a thin wrapper around the corresponding wgpu method. Two consequences:

1. Bring your shader from a browser: WGSL is the same. Descriptor shapes (GPUBufferDescriptor, GPURenderPipelineDescriptor, …) are the same JSON-serialisable objects with identical field names. GPUBufferUsage.STORAGE is 0x80 here just like in the browser.
2. Wgpu's gotchas are the spec's gotchas. If wgpu rejects an unaligned offset, this binding rejects too — we don't paper over the difference.

The one un-spec-y thing is the flat function shape: deviceCreateBuffer(device, descriptor) instead of device.createBuffer(descriptor). Same arguments, just receiver-as-first-arg. A class wrapper that restores literal spec parity will land once Perry's compiler grows codegen support for class-method dispatch on registered handles (the "property/method tower used for Web Fetch handles" called out in #571).

Surface coverage

• Adapter / Device / Queue: requestAdapter, adapterRequestDevice, adapterDrop, deviceDestroy, devicePoll
• Buffer: deviceCreateBuffer, bufferDestroy, bufferMapAsync, bufferGetMappedRange, bufferUnmap
• Shader: deviceCreateShaderModule (WGSL only — same as browser)
• Bindings: deviceCreateBindGroupLayout / deviceCreatePipelineLayout / deviceCreateBindGroup (with buffer / sampler / textureView resources, plus dynamic offsets)
• Pipelines: compute + render, sync + async, plus getBindGroupLayout accessors
• Textures: deviceCreateTexture / textureCreateView / textureDestroy / deviceCreateSampler — full TextureFormat / TextureViewDimension / TextureAspect / FilterMode / AddressMode / CompareFunction surface
• Render pass: full method tower — setPipeline / setBindGroup (with dynamic offsets) / setVertexBuffer / setIndexBuffer / draw / drawIndexed / setViewport / setScissorRect / setBlendConstant / setStencilReference / beginOcclusionQuery / endOcclusionQuery / end
• Compute pass: setPipeline / setBindGroup (with dynamic offsets) / dispatchWorkgroups / end
• Command encoder: beginComputePass / beginRenderPass / copyBufferToBuffer / copyBufferToTexture / copyTextureToBuffer / copyTextureToTexture / resolveQuerySet / finish
• Queue: submit / writeBuffer / writeTexture / onSubmittedWorkDone (real device-poll, not a stub)
• QuerySet: deviceCreateQuerySet (occlusion + timestamp) / querySetDestroy
• Error scopes: devicePushErrorScope / devicePopErrorScope (real wgpu errorScope wired through, returning {type, message} JSON)
• Surface (on-screen): requestSurface / surfaceFromNativeView / surfaceGetViewPtr / surfaceGetPreferredFormat / surfaceConfigure / surfaceGetCurrentTexture / surfacePresent / surfaceUnconfigure / surfaceDrop — render a swapchain into a perry-ui window

On-screen surface

WebGPU presents to a window through a swapchain. The browser hides this behind canvas.getContext("webgpu"); natively, wgpu builds a Surface from a platform window/view handle. Perry programs are headless by default, so the on-screen path rides on perry-ui: its BloomView widget ("render-surface host for an external GPU renderer") reserves a GPU-capable native view in the window and hands back the platform handle. @perryts/webgpuwraps that handle into a swapchain — it never creates windows itself, so compute / render-to-texture programs still link without dragging in the UI toolkit.

The render loop is the spec's, plus one explicit surfacePresent (browsers present implicitly at task end; native wgpu needs the hand-back). A complete, runnable program is in examples/triangle-macos — the shape is:

import{App,BloomView,bloomViewGetNativeHandle,onFrame}from"perry/ui";import{requestAdapter,adapterRequestDevice,surfaceFromNativeView,surfaceGetPreferredFormat,surfaceConfigure,surfaceGetCurrentTexture,surfacePresent,textureCreateView,deviceCreateCommandEncoder,commandEncoderBeginRenderPass,renderPassSetPipeline,renderPassDraw,renderPassEnd,commandEncoderFinish,queueSubmit,}from"@perryts/webgpu";constview=BloomView(800,600);// perry-ui reserves the native viewconstsurface=surfaceFromNativeView(bloomViewGetNativeHandle(view));// wrap it// Surface-first: pick the adapter *against* the surface. Required on// d3d12/vulkan (the adapter/device must be presentation-compatible or// `surfaceConfigure` fails); harmless on Metal.constadapter=awaitrequestAdapter(surface);const{ device, queue }=awaitadapterRequestDevice(adapter);constformat=surfaceGetPreferredFormat(surface,adapter);surfaceConfigure(surface,{ device, format,width: 800,height: 600});// …build `pipeline` against `format` exactly as in the browser…onFrame(functionframe(){consttex=surfaceGetCurrentTexture(surface);// this frame's swapchain imageconsttarget=textureCreateView(tex);// → render-pass color attachmentconstenc=deviceCreateCommandEncoder(device);constpass=commandEncoderBeginRenderPass(enc,{colorAttachments: [{view: target,loadOp: "clear",storeOp: "store",clearValue: {r: 0,g: 0,b: 0,a: 1}}],});renderPassSetPipeline(pass,pipeline);renderPassDraw(pass,3);renderPassEnd(pass);queueSubmit(queue,JSON.stringify([commandEncoderFinish(enc)]));surfacePresent(surface);// hand the frame to the compositor});App({title: "WebGPU Triangle",width: 800,height: 600,body: view});

Main-thread note: onFrame fires on perry-ui's main-thread frame pump, so each getCurrentTexture + surfacePresent blocks it on GPU/vsync sync. Render only as often as the image actually changes — a static scene can render a short burst and stop re-registering; the CAMetalLayer keeps showing the last presented frame and the UI stays responsive. (See the example.)

surfaceFromNativeView(ptr) is the canonical seam, and the one a future "wrap an existing wgpu device" bloom entry point builds on (see below). requestSurface({ width, height }) is the inverse — the binding allocates its ownNSView and returns its pointer via surfaceGetViewPtr for the host to embed; it's macOS-only today, since other toolkits own view creation.

Requirements

The ergonomic camelCase API (requestAdapter, not js_webgpu_request_adapter) routes to native through Perry's native-library support, which needs a Perry build with ergonomic-export routing (PerryTS/perry#5621) and the json descriptor param type (#5626). The package manifest declares each target's crate + lib and marks descriptor params json accordingly.

Platform status

Platform	View creation (requestSurface)	Adopt existing (surfaceFromNativeView)	Backend
macOS	✅ implemented & verified (NSView → CAMetalLayer)	✅ (NSView*)	Metal
iOS	— (UIKit owns the view)	✅ (UIView*, must be CAMetalLayer-backed)	Metal
Windows	— (create the HWND via perry-ui)	✅ (HWND)	DX12 / Vulkan
Android	— (the framework owns the Surface)	✅ (ANativeWindow*)	Vulkan
Linux	—	⏳ needs a display handle alongside the window (X11/Wayland) — follow-up	Vulkan

macOS is verified end-to-end — examples/triangle-macos renders a live triangle into a perry-ui window. The other arms compile against their platform window handles but await on-device validation. Mobile and Windows invert view ownership (the toolkit/OS creates the view), so they go through surfaceFromNativeView with the toolkit's pointer rather than requestSurface. Linux still needs a richer entry point because a single pointer can't carry the X11/Wayland display connection wgpu also requires.

Windows notes (PerryTS/perry#5812):

• Adapter must be surface-compatible. d3d12/vulkan reject surfaceConfigure unless the adapter was picked against the surface, so pass it: create the surface first, then requestAdapter(surface) (see the render example above). Metal doesn't care.
• Link libraries. wgpu's d3d12/gl backends pull D3DCompile and wgl*; the manifest declares d3dcompiler + opengl32 under targets.windows.libs so the link resolves (previously LNK2019). Perry ≥ the #5812 build also auto-appends these when a d3d12/vulkan backend is declared.
• Mount-aware swapchain (the #5812 "gray window", fixed). perry-ui creates widget HWNDs under a hidden message-only parking window and only reparents them into the real app window when App() mounts the body — usually aftersurfaceConfigure built the swapchain. Two failure modes followed: presents against the parked (invisible) HWND are never consumed by DWM, so a pre-mount render loop wedges the main thread once the swapchain queue fills (~10 frames — the reported "N frames present fine, then nothing"); and on some driver/DWM combinations a swapchain created under the parking window never composites even after the mount. The binding now handles both: surfaceGetCurrentTexture returns 0 ("not presentable yet" — same retry contract as macOS's not-yet-on-screen state) while the view is not visible, and on the first acquire after a reparent it recreates the wgpu::Surface so the swapchain is rebuilt against the visible window tree. Verified end-to-end on Windows 11 x64 (d3d12). Debug knobs: PERRY_WEBGPU_DIAG=1 logs per-call window state; PERRY_WEBGPU_NO_REBIND=1 disables the reparent rebind (bisection only).

Out of scope

• External textures: depend on the canvas integration crate exposing ImageBitmap (#570 acceptance criterion); the descriptor parser treats the v0.2 wire format as a forwards-compatible no-op.

Bloom interop

Out of scope for this version — every adapter/device is freshly allocated. The bloom-side hook is a future addition: bloom will expose its wgpu::Device + wgpu::Queue via a stable internal API, and @perryts/webgpu will gain a "wrap an existing device" entry point so a user can drop into raw WebGPU for a custom pass inside a bloom app. The surface side already has its half of that seam — surfaceFromNativeView wraps a view bloom hands over — so the remaining work is the device/queue adoption path.

Versioning

Pre-1.0. The perry.nativeLibrary.abiVersion (currently 0.5) is a hard pin against Perry's perry-ffi ABI — bump in lockstep with the Perry release that the bindings target.

License

MIT — see LICENSE.