GPU compute

In one line#

BASIC can run WebGPU compute shaders: you put numbers in a buffer on the GPU, run a small WGSL program over all of them at once, and read the result back into a memblock. Fourteen commands, in the Compute category, do the whole job.

⚠️ Engine only. GPU compute works when you press Play in the editor, on a GPU and browser that support WebGPU compute. The HTML export and the Bevy builds do not have it. Write the CPU fallback before you rely on it.

When it is available#

Call ComputeSupported() first. It answers on every target, and every other Compute command assumes it said 1.

There is no silent emulation: a Compute command either runs on the GPU or stops with an error that names the target. The pattern is always the same: test ComputeSupported(), and take a CPU path when it is 0 (see the fallback pattern below).

The pieces#

Steps 6 and 7 return jobs, the same jobs as the rest of the asynchronous commands. Poll them with PlatformJobDone, PlatformJobError and PlatformJobResultNumber.

Storage buffers#

buffer = CreateComputeBuffer(<elementCount>, <type$>)
CreateComputeBuffer(<bufferId>, <elementCount>, <type$>)
ComputeBufferFromMemblock(<buffer>, <memblock>)

• Three element types: "f32", "i32", "u32". All three are 4 bytes. A buffer of n elements is n * 4 bytes.
• The memblock must have exactly the buffer's size. A 16-byte buffer takes a 16-byte memblock, no more, no less. Anything else stops the program.
• The bytes are copied as they are. Write f32 values with MemblockFloat, i32 and u32 values with MemblockInt, at offsets 0, 4, 8….
• One buffer holds at most 128 MiB. That is the common minimum of mobile GPUs.
• The second form takes the id you choose; it must not be in use.

WGSL pipelines#

shader = CreateComputeShader(<wgsl$>, <bindingsJson$> [, <entryPoint$>])
CreateComputeShader(<shaderId>, <wgsl$>, <bindingsJson$> [, <entryPoint$>])
ComputeShaderBuffer(<shader>, <name$>, <buffer>)

• The kernel is raw WGSL. It is not GLSL and it is not translated. It must contain an @compute entry point, named main unless you pass another name.

• The bindings JSON maps each WGSL variable to its slot, exactly as the WGSL declares it:

  {"data": {"group": 0, "binding": 0}}
  

  One entry per storage variable, at least one, no slot used twice.

• ComputeShaderBuffer binds a buffer to a variable by its WGSL name.

• Storage buffers are the only bindings. There are no uniforms and no textures. To pass a parameter, such as a time step, put it in a small storage buffer of its own.

• CreateComputeShader only checks the outline: the source contains @compute, the JSON is valid, the slots are distinct. The WGSL itself compiles in the background, and a WGSL error comes back on the dispatch job, through PlatformJobError.

In BASIC strings, write the JSON quotes as \":

bindings$ = "{\"data\":{\"group\":0,\"binding\":0}}"
Console(bindings$)          ' {"data":{"group":0,"binding":0}}

Dispatch, and its place in the frame#

job = ComputeDispatchAsync(<shader>, <groupsX> [, <groupsY> [, <groupsZ>]])

The arguments are workgroup counts, not element counts. With @workgroup_size(64) and n elements, ask for (n + 63) DIV 64 groups, and let the kernel skip the indices past the end.

A dispatch does not run when you call it. It is queued, and the queue is sent to the GPU just before the next frame is drawn.

• The queue keeps its order. Two dispatches issued one after the other run in that order, even when the second kernel's pipeline is ready first. A second kernel can read what the first one wrote.
• The job is done when the work has been sent. Wait for it with PlatformJobDone, then check PlatformJobError (empty on success).
• A pipeline that is still not ready after 10 seconds fails its job.
• Waiting needs a program that can wait. The usual loop is WHILE PlatformJobDone(job) = 0 … Sync() … ENDWHILE. Its Sync() must really yield: the program needs a DO … LOOP or an END, or the loop must be in a label. Otherwise the tab freezes, see The game loop.

Readback into a memblock#

job = ComputeBufferReadAsync(<buffer>)
job = ComputeBufferReadAsync(<buffer>, <newMemblockId>)
memblock = PlatformJobResultNumber(job)

• The read makes a new memblock with the buffer's bytes. The job's number is its id.
• The second form chooses that id, and the id must be free. Reading into a memblock that exists stops the program: delete it first.
• ⚠️ The read is not queued behind the dispatches. Start it only once the dispatch job is done, or you read the data as it was before the kernel ran.
• A memblock holds at most 100 000 000 bytes, so a buffer you read back must stay under 25 000 000 elements, below the 128 MiB a buffer can hold.
• Delete the memblocks you no longer need, and release finished jobs with PlatformJobRelease.

A complete example#

Two kernels on one buffer: the first doubles each value, the second adds one. Both are queued at once, and the queue order guarantees the result.

' 1, 2, 3, 4 become 3, 5, 7, 9 on the GPU.
IF ComputeSupported() = 0
  Console("No WebGPU compute on this target")
  END
ENDIF

' The data: four f32 values, 4 bytes each.
input = CreateMemblock(16)
FOR i = 0 TO 3
  MemblockFloat(input, i * 4, i + 1)
NEXT

' A storage buffer of the same byte size, filled from the memblock.
buffer = CreateComputeBuffer(4, "f32")
ComputeBufferFromMemblock(buffer, input)

' Two kernels that share their first lines. One invocation per element.
head$ = "@group(0) @binding(0) var<storage, read_write> data: array<f32>; "
head$ = head$ + "@compute @workgroup_size(64) "
head$ = head$ + "fn main(@builtin(global_invocation_id) id: vec3<u32>) { "
head$ = head$ + "if (id.x >= arrayLength(&data)) { return; } "
bindings$ = "{\"data\":{\"group\":0,\"binding\":0}}"
doubler = CreateComputeShader(head$ + "data[id.x] = data[id.x] * 2.0; }", bindings$)
adder = CreateComputeShader(head$ + "data[id.x] = data[id.x] + 1.0; }", bindings$)
ComputeShaderBuffer(doubler, "data", buffer)
ComputeShaderBuffer(adder, "data", buffer)

' Queue both. They run in this order, before the next frame is drawn.
groups = (4 + 63) DIV 64
job1 = ComputeDispatchAsync(doubler, groups)
job2 = ComputeDispatchAsync(adder, groups)
WHILE PlatformJobDone(job2) = 0
  Sync()
ENDWHILE
IF PlatformJobError(job1) <> "" OR PlatformJobError(job2) <> ""
  Console("dispatch failed: " + PlatformJobError(job1) + PlatformJobError(job2))
  END
ENDIF

' Read back. The job's number is the id of a new memblock.
readJob = ComputeBufferReadAsync(buffer)
WHILE PlatformJobDone(readJob) = 0
  Sync()
ENDWHILE
result = PlatformJobResultNumber(readJob)
FOR i = 0 TO 3
  Console(Str(MemblockFloat(result, i * 4)))   ' 3, then 5, 7, 9
NEXT

' Free everything.
DeleteMemblock(result)
DeleteMemblock(input)
DeleteComputeShader(doubler)
DeleteComputeShader(adder)
DeleteComputeBuffer(buffer)
END

What makes it work:

• The END lines. They let Sync() yield inside the WHILE loops. Without them the program has no DO loop, no END, and the loops would freeze the tab.
• job2 is enough to wait for. The queue runs in order, so when the second dispatch has been sent, the first one has been sent too.
• The read starts after the dispatches are done.
• arrayLength(&data) guards the end. A workgroup of 64 runs 64 invocations; only the first 4 have an element.

The fallback pattern#

Keep the data in a memblock on both paths. The GPU path uploads it and reads it back; the CPU path works on it in place. Everything after the branch is the same code.

' Double 1024 values: on the GPU when possible, on the CPU otherwise.
count = 1024
data = CreateMemblock(count * 4)
FOR i = 0 TO count - 1
  MemblockFloat(data, i * 4, i)
NEXT

IF ComputeSupported()
  GOSUB DoubleOnGpu
ELSE
  DoubleOnCpu(data, count)
ENDIF
Console(Str(MemblockFloat(data, 10 * 4)))   ' 20, on either path
END

FUNCTION DoubleOnCpu(mem, n)
  FOR k = 0 TO n - 1
    MemblockFloat(mem, k * 4, MemblockFloat(mem, k * 4) * 2)
  NEXT
ENDFUNCTION

DoubleOnGpu:
  buf = CreateComputeBuffer(count, "f32")
  ComputeBufferFromMemblock(buf, data)
  wgsl$ = "@group(0) @binding(0) var<storage, read_write> data: array<f32>; "
  wgsl$ = wgsl$ + "@compute @workgroup_size(64) "
  wgsl$ = wgsl$ + "fn main(@builtin(global_invocation_id) id: vec3<u32>) { "
  wgsl$ = wgsl$ + "if (id.x >= arrayLength(&data)) { return; } "
  wgsl$ = wgsl$ + "data[id.x] = data[id.x] * 2.0; }"
  kernel = CreateComputeShader(wgsl$, "{\"data\":{\"group\":0,\"binding\":0}}")
  ComputeShaderBuffer(kernel, "data", buf)
  job = ComputeDispatchAsync(kernel, (count + 63) DIV 64)
  WHILE PlatformJobDone(job) = 0
    Sync()
  ENDWHILE
  ' Read back under the same id: free it first, then let the read recreate it.
  DeleteMemblock(data)
  readJob = ComputeBufferReadAsync(buf, data)
  WHILE PlatformJobDone(readJob) = 0
    Sync()
  ENDWHILE
  DeleteComputeShader(kernel)
  DeleteComputeBuffer(buf)
RETURN

• The GPU path is a label, not a FUNCTION: a FUNCTION cannot wait. GOSUB to a label that waits needs a DO loop or an END in the program.
• The CPU path is an ordinary FUNCTION: it does not wait.
• In a real program, also check PlatformJobError after each wait, as in the complete example.

To keep the GPU code out of the exports altogether, put it between #IF ENGINE and #ENDIF: the editor defines ENGINE, the HTML and Bevy builds do not. See the preprocessor.

Lifetime#

• Stop destroys everything. Buffers and shaders are freed when the program stops, and their ids start again at 1 on the next run.
• Dispatches still queued at Stop fail with Compute session ended before queued work was dispatched.
• Delete what you no longer need during a run: GPU memory is not collected for you while the program runs.

How it differs from classic dialects#

AGK, DarkBASIC and PureBasic have no GPU compute. This is an aukimi addition, built on WebGPU. Two consequences for a reader coming from them:

• The kernels are WGSL, WebGPU's shading language. GLSL compute shaders do not load.
• Everything is asynchronous. A dispatch and a readback are jobs you wait for across frames, never calls that block until the GPU is done.

The Compute commands#

Limits and common snags#

• Test ComputeSupported() first, and write the CPU path.
• The upload memblock has exactly the buffer's size.
• Dispatch counts are workgroups. Guard the end of the array in WGSL.
• Wait for the dispatch job before you start the read.
• A readback id must be free. Delete the old memblock first.
• Sync() in the wait loops must yield: DO … LOOP, END, or a label.
• Storage buffers only. Parameters go in a small buffer of their own.