Peta::WebGPU#

📦 std

Peta::WebGPU - compute and draw on the graphics card.

WebGPU is the modern successor to OpenGL: a portable API that reaches Vulkan, Metal and Direct3D 12 underneath, so the same program runs on whichever of those the machine has. Peta::WebGPU is pperl’s binding to it, built on the wgpu implementation.

use Peta::WebGPU;

my $gpu = Peta::WebGPU::Device->new;
printf "%s on %s\n", $gpu->info->{name}, $gpu->info->{backend};

my $shader = $gpu->shader(<<'WGSL');
@group(0) @binding(0) var<storage, read_write> data: array<u32>;
@compute @workgroup_size(1)
fn main(@builtin(global_invocation_id) id: vec3<u32>) {
    data[id.x] = data[id.x] * 2u;
}
WGSL

my $buf = $gpu->storage_buffer(pack 'L*', 1 .. 8);
$gpu->compute($shader, 'main', [$buf], 8);
print join(',', unpack 'L*', $buf->read), "\n";   # 2,4,6,8,10,12,14,16

This build may not have it#

The GPU backend is large, so it is a build option that is off by default. Peta::WebGPU::have_gpu() answers whether this pperl has it AND found a usable adapter, without dying, so a program can fall back to doing the work on the CPU:

if (Peta::WebGPU::have_gpu()) { ... } else { ... }

A pperl built without the option tells you so when you load the module, and names the flag to rebuild with.

Shaders are WGSL#

The programs that run on the card are written in WGSL, WebGPU’s own shading language - not GLSL. $gpu->shader($source) compiles one, and a syntax error in it raises a Perl exception with the compiler’s message rather than killing the program.

Errors#

Everything here reports failure as an exception, carrying the GPU driver’s own text, so eval is how you trap one. A call that succeeds and has nothing else to give back returns the object it was called on, so or die behaves and calls chain:

$gpu->compute($shader, 'main', [$buf], 8)->poll;

fork#

A device, and every buffer, shader, pipeline, surface and texture made from it, belongs to the process that created it. fork copies the objects but not the graphics driver’s threads, so in the child they cannot be used: any method on one dies with

this Peta::WebGPU::Device was created before fork; open a Peta::WebGPU::Device in the child process

The child may leave them alone and exit, or open a device of its own and work on that; the parent’s objects are unaffected either way. One limit: where the parent has opened the gl backend (the default all includes it), the child’s device comes from the other backends only, and asking for backend => 'gl' there fails - OpenGL’s display connection is shared with the parent and cannot be opened twice. A worker pool is best served by opening its device inside each worker:

my $pid = fork // die "fork: $!";
if (!$pid) {
    my $gpu = Peta::WebGPU::Device->new;    # this process's own
    ...
    exit 0;
}

Once a process has asked for an adapter or opened a device, a child forked from it ends, after Perl’s own END blocks and destructors, without running the exit handlers of the C libraries loaded before the graphics drivers: the drivers’ own handlers wait for threads the child does not have. Buffered output is flushed first.

Peta::WebGPU::Device#

Opening a device is the first thing any GPU program does. It picks an adapter, opens it, and gives you the one object everything else hangs off:

my $gpu = Peta::WebGPU::Device->new;                       # let it choose
my $gpu = Peta::WebGPU::Device->new(power => 'high-performance');
my $gpu = Peta::WebGPU::Device->new(backend => 'vulkan');

power is high-performance (prefer the discrete card) or low-power (prefer the integrated one, which saves battery and is often quicker to start). backend restricts the choice to one of vulkan, metal, dx12, gl or all. Opening fails with an exception if the machine has no usable adapter.

On gl, what the device can do depends on the OpenGL context the driver hands out. Drawing always works. Compute shaders need OpenGL 4.3 or OpenGL ES 3.1; on an older context (the NVIDIA driver gives 3.3) compute and pipeline die saying so, and vulkan is the backend to use.

  • $gpu->info - a hash reference describing the adapter actually opened: name, backend, device_type, driver, vendor, device.

  • $gpu->limits - what this device can do: max_buffer_size, max_compute_workgroups_per_dimension and the other ceilings you have to stay under. These are the real hardware limits, not a portable minimum, so a program that runs on one machine can still ask for more than another machine offers.

  • $gpu->shader($wgsl) - compile a shader.

  • $gpu->storage_buffer($data_or_size) - a buffer a shader can read and write, and that you can read back.

  • $gpu->uniform_buffer($data_or_size) - a small read-only buffer of parameters.

  • $gpu->pipeline($shader, $entry) - prepare a shader entry point to run. Worth keeping when you dispatch the same shader repeatedly.

  • $gpu->compute($shader, $entry, \@buffers, $x, $y, $z) - build the pipeline, bind and dispatch in one call.

  • $gpu->run($pipeline, \@buffers, $x, $y, $z) - dispatch a pipeline you already have.

  • $gpu->run($binding, $x, $y, $z) - dispatch a pipeline whose buffers are already bound ($pipeline->bind, below).

  • $gpu->begin, $gpu->commit and $gpu->batch(sub { ... }) - send several dispatches to the card at once; see “Many dispatches”.

  • $gpu->features - the optional features this device has, as a list of names: shader-f16 (a shader may say enable f16;) and timestamp-query. A device has each one its adapter offers. $gpu->has_feature($name) is the yes or no for one of them.

  • $gpu->poll - wait for everything submitted so far to finish.

  • $gpu->memory - { allocated => ..., reserved => ... }: the bytes this device’s buffers and textures occupy, and the bytes it holds from the driver to put them in. undef on a backend that keeps no such account (Vulkan and DirectX 12 do).

Device->new(memory_limit => 90) makes the device refuse a buffer that would take it past 90 percent of the memory the driver allows this process. The refusal is an Out of Memory exception from storage_buffer, which a program can catch; without the option the driver decides, and may take the device with it.

The buffers are bound in the order given, as @binding(0), @binding(1) and so on within @group(0).

Many dispatches#

Each run on its own is a round trip to the driver, and binding a list of buffers is work of its own. A program that dispatches the same few shaders over and over - a simulation step, a training step - saves both:

my $step = $pipeline->bind([$state, $params]);   # once

$gpu->batch(sub {                                # every step
    $gpu->run($step, $groups);
    $params->write(pack 'f', $rate);
    $gpu->run($step, $groups);
});

$pipeline->bind(\@buffers) returns a Peta::WebGPU::Binding: the pipeline with those buffers bound, which keeps the buffers alive and can be run any number of times.

$gpu->batch($code) records every run the sub makes and sends them to the card together when it returns. $gpu->begin and $gpu->commit are the same without a sub. Inside a batch:

  • a mistake in a run is reported when the batch is sent, by commit or batch, not by the run itself; for a wrong buffer list the message says which run of the batch it was;

  • $buf->write takes effect in order: dispatches before it see the old contents, dispatches after it the new;

  • $buf->read sends what is recorded so far, reads, and the batch goes on;

  • a sub that dies leaves nothing sent, and its exception reaches the caller of batch.

Batches do not nest.

$gpu->timed($code) is batch with a stopwatch on the card: it returns the nanoseconds the card spent on the batch’s dispatches, which is what a shader costs, without the time the program spent getting it there. It waits for the card to finish, and dies on a device without timestamp-query.

my $ns = $gpu->timed(sub { $gpu->run($step, $groups) });
printf "%.2f ms on the card\n", $ns / 1e6;

Peta::WebGPU::Buffer#

  • $buf->size - in bytes.

  • $buf->read - the contents, as a byte string. Unpack it exactly as the shader’s type says: L* for array<u32>, l* for array<i32>, f* for array<f32>.

  • $buf->read($offset, $length) - a part of the contents, both in bytes and both multiples of four; without $length, from $offset to the end. Reading the one number you need costs the card far less than reading the buffer it is in.

  • $buf->write($bytes, $offset) - replace part of it. A PDL ndarray of type float, long or ulong in place of $bytes is written as it is, so a large buffer can be filled piece by piece.

  • $buf->destroy - give the buffer’s memory back to the card now. A buffer is freed when its last reference goes, but a binding made with it is such a reference; destroy does not wait for those. The buffer cannot be used afterwards, and a binding that has it dies when run.

A buffer’s size is always rounded up to a multiple of four bytes, because that is the granularity the GPU copies at. Pack L, l or f and it never comes up; pack an odd number of bytes and size will be a little larger than what you passed, with the extra bytes zero. A write must be aligned the same way.

Reading is not free: it waits for the card to finish and copies the data back across the bus, so read once at the end rather than inside a loop.

Workgroups, not elements#

compute and run take the number of WORKGROUPS to launch, not the number of data elements. A shader declaring @workgroup_size(64) processes 64 elements per workgroup, so 1000 elements needs int(($n + 63) / 64) workgroups - and the shader must check its own index against the array length, because the last workgroup runs past the end:

my $wgsl = <<'WGSL';
@group(0) @binding(0) var<storage, read_write> v: array<f32>;
@compute @workgroup_size(64)
fn scale(@builtin(global_invocation_id) id: vec3<u32>) {
    if (id.x >= arrayLength(&v)) { return; }
    v[id.x] = v[id.x] * 2.0;
}
WGSL
my $buf = $gpu->storage_buffer(pack 'f*', @data);
$gpu->compute($gpu->shader($wgsl), 'scale', [$buf], int((@data + 63) / 64));

Drawing#

The other half of a graphics card. A program that draws does the same three things every frame: get somewhere to draw, record what to draw there, and hand it over.

use SDL2 qw(:all);
use Peta::WebGPU;

my $win  = SDL2::Window->new(title => 'triangle', w => 800, h => 600);
my $gpu  = Peta::WebGPU::Device->new(window => $win);
my $surf = $gpu->surface;

my $pipe = $gpu->render_pipeline($gpu->shader($WGSL),
    vertex => 'vs', fragment => 'fs', format => $surf->format,
    vertex_buffers => [ { attributes => ['float32x2', 'float32x3'] } ],
);
my $verts = $gpu->vertex_buffer(pack 'f*', @triangle);

while (running()) {
    my $frame = $surf->next_frame or next;
    $gpu->render_pass($frame, clear => [0.1, 0.2, 0.3, 1])
        ->pipeline($pipe)->vertices($verts)->draw(3)->end;
    $frame->present;
}

Every recording call returns the pass, so a frame reads as one sentence. end is what hands the work to the card; nothing is drawn before it.

Drawing without a window#

A render pass does not care whether its target is on screen. Give it a texture instead of a frame and the picture comes back as bytes:

my $gpu = Peta::WebGPU::Device->new;           # no window at all
my $img = $gpu->texture(width => 800, height => 600);
$gpu->render_pass($img, clear => [0, 0, 0, 1])
    ->pipeline($pipe)->vertices($verts)->draw(3)->end;
my $rgba = $img->read;                          # 800*600*4 bytes

This needs no display, no window manager and no compositor, so it is how a program renders on a server - and how this module’s own tests check what was drawn, pixel by pixel.

What a window has to be#

Peta::WebGPU opens no windows. It asks a window object two questions, and anything that answers them can be drawn on:

  • $win->wm_info - a hash reference with driver (x11 or wayland), display, and window or surface.

  • $win->size - the list ($width, $height).

SDL2::Window answers both, which is where a window normally comes from here. A window system this build cannot make a surface on is refused by name rather than failing later.

The window belongs at Device->new, not at surface, because the adapter has to be one that can actually present to that window - a machine with two cards can otherwise open the one that does not drive the screen.

Peta::WebGPU::Surface#

$gpu->surface - the window’s swap chain, configured and ready. $gpu->surface($other_window) gives a second window its own.

  • $surf->format - the pixel format the compositor wants, as a WebGPU name such as bgra8unorm-srgb. Pass it to render_pipeline so the two agree; they must.

  • $surf->size - the list ($width, $height).

  • $surf->configure($width, $height) - after the window was resized.

  • $surf->next_frame - the image to draw this frame into, or undef when the window is hidden or the compositor is busy. That is not an error: skip the frame and come back.

present_mode chooses how frames reach the screen: fifo waits for the display and never tears (the safe default), immediate does not wait, mailbox keeps only the newest. A mode the window does not offer is refused, with the ones it does offer in the message.

A swap chain lives on the window system’s connection. SDL_Quit, and SDL_QuitSubSystem for video, close that connection, so they take the swap chains of every surface first, with the frames not yet presented. The objects remain: size and format still answer, while configure, next_frame, present and a render pass on such a frame die saying that the window system was shut down. The device is not affected and goes on working off screen. A window object that is not SDL2’s has to outlive its surfaces by itself.

Peta::WebGPU::Frame#

One image out of the swap chain. It is a render target exactly as a texture is, and $frame->present puts it on the screen. A frame that goes out of scope unpresented is discarded, which is the right thing for a frame abandoned half-drawn. Presenting twice is refused.

End the pass before presenting: end is what submits the drawing.

Peta::WebGPU::Texture#

$gpu->texture(width => .., height => .., format => ..) - somewhere to draw that is not the screen. The format defaults to rgba8unorm, the same default render_pipeline has, so the two agree unless you say otherwise.

  • $tex->size - the list ($width, $height).

  • $tex->format - its format, by name.

  • $tex->read - the pixels, tightly packed, top row first. Four bytes per pixel for the 8-bit-per-channel formats, so unpack 'C*' gives R, G, B, A in that order.

Reading waits for the card and copies the image back across the bus, exactly as reading a buffer does. It belongs at the end, not in a loop.

$gpu->depth_buffer($width, $height) makes the other kind of target: a depth texture, which is what stops a far triangle painting over a near one. Pass it to render_pass as depth => $d and build the pipeline with depth => 1; both halves are needed, and wgpu says so if only one is there.

Render pipelines#

$gpu->render_pipeline($shader, %options) compiles one shader’s vertex and fragment entry points into everything the card needs to draw. It is worth keeping: build it once, use it every frame.

  • vertex / fragment - the entry point names in the WGSL. With one of each in the shader they can be left out.

  • format - the format of what is being drawn into. Default rgba8unorm; for a window it must be $surf->format.

  • vertex_buffers - the shape of the vertex data; see below.

  • topology - triangle-list (the default), triangle-strip, line-list, line-strip or point-list.

  • cull - none (the default), back or front, and front_face (ccw, the default, or cw), which decide which side of a triangle is its front and whether the other one is drawn at all.

  • depth => 1 - take part in depth testing.

  • blend => 1 - blend with what is already there, by alpha.

Describing the vertices#

Each entry of vertex_buffers describes one buffer, as a list of attributes. The short form is a list of format names, which are laid end to end at @location(0), @location(1) and so on:

vertex_buffers => [ { attributes => ['float32x3', 'float32x3'] } ],

That is the layout of a WGSL function taking consecutive locations, which is what almost every shader has. Anything left out is computed: offsets follow each other, locations count up, and the stride is where the last attribute ends. A hash gives the long form for the cases that need it:

vertex_buffers => [ {
    stride     => 32,
    attributes => [ { format => 'float32x3', location => 0, offset => 0 },
                    { format => 'float32x2', location => 5, offset => 20 } ],
} ],

Formats are WebGPU’s own names - float32, float32x2, float32x3, float32x4, and the uint32, sint32, unorm8x4 and float16x2 families. A name that is not one of them is refused with the whole list, because a typo here otherwise shows up as geometry that is nearly right.

Peta::WebGPU::Pass#

$gpu->render_pass($target, %options) starts recording. The target is a Frame or a Texture.

  • clear => [$r, $g, $b, $a] - the colour to start from, each component 0 to 1. clear => 0 keeps what is already there, which is how a second pass draws on top of the first. Left out altogether, it clears to black.

  • depth => $depth_texture - test against depth while drawing.

Then, in order:

  • $pass->pipeline($pipeline) - what to draw with. First, always: the bind groups take their layout from it.

  • $pass->bind(\@buffers, [$group]) - the uniform and storage buffers of @group(0), in binding order, exactly as compute takes them.

  • $pass->vertices($buffer, [$slot]) - the vertex data.

  • $pass->indices($buffer, ['uint16'|'uint32']) - the index data. 32-bit by default, because a mesh with more than 65536 vertices is ordinary.

  • $pass->draw($vertex_count, [$instances]) or $pass->draw_indexed($index_count, [$instances]).

  • $pass->end - stop recording and hand the work to the card.

Every one of them returns the pass, so they chain. A pass that is dropped without end throws its recording away rather than submitting half a frame.

$gpu->vertex_buffer($data) and $gpu->index_buffer($data) make the buffers, and take the same arguments storage_buffer does - a byte string, a size, or a PDL ndarray of a 32-bit type.

Functions#

Availability#

have_gpu#

True when this pperl can actually reach a graphics card: the build has the GPU option AND an adapter was found.

Adapters#

adapters#

Every graphics adapter this machine offers, as a list of hash references.