<!-- GENERATED by docs/sphinx/bin/20-reference (p5) from Peta/WebGPU/mod.rs. Do not edit. -->

```{index} single: Peta::WebGPU
```

# Peta::WebGPU

```{pperl-module-badges} Peta::WebGPU
```

`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.

```perl
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:

```perl
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:

```perl
$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

```text
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:

```perl
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:

```perl
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:

```perl
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`.

```perl
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:

```perl
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.

```perl
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:

```perl
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:

```perl
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:

```perl
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`](WebGPU/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`](WebGPU/adapters)

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

```{toctree}
:hidden:
:maxdepth: 1

WebGPU/have_gpu
WebGPU/adapters
```