&x gives a pointer, ptr<AS, T> (plus an access mode for storage), and *p dereferences it. Pointers let functions modify their caller's variables; they cannot be stored in a var, compared or turned into integers, which keeps WGSL memory-safe:
// out: 3 x u32
@group(0) @binding(0) var<storage, read_write> out: array<u32>;
fn restock(p: ptr<function, u32>) { *p += 10; }
fn fill(dst: ptr<storage, array<u32>, read_write>, i: u32) { dst[i] = i * 100; }
@compute @workgroup_size(1) fn main() {
var stock = 5u;
let p = &stock; // a pointer held in a let
restock(p);
restock(&stock);
out[0] = stock;
fill(&out, 1); // a pointer to a storage buffer
let q = &out[2]; // a pointer to one element
*q = *p - 18;
}Output
out: 25, 100, 7
Two conveniences come from extensions Chrome 154 1 reports: unrestricted_pointer_parameters lets fill() take a storage pointer (the base language allows only function and private ones), and pointer_composite_access allows dst[i] instead of (*dst)[i]. Portable shaders pass storage data by index instead.
<!doctype html>
<style>
body { margin: 0; background: #f7f4ee; font: 14px system-ui, sans-serif; }
.stage { position: relative; width: 100%; max-width: 600px; }
.stage canvas { display: block; width: 100%; }
.stage canvas + canvas { position: absolute; inset: 0; pointer-events: none; }
</style>
<div class="stage">
<canvas id="view" width="600" height="320"></canvas>
<canvas id="labels" width="600" height="320"></canvas>
</div>
<script>
const canvas = document.getElementById('view');
const ink = document.getElementById('labels').getContext('2d');
function showMessage(text) { // 2D fallback when WebGPU is missing
const ctx = canvas.getContext('2d');
ctx.fillStyle = '#fbeaea'; ctx.fillRect(0, 0, canvas.width, canvas.height);
ctx.fillStyle = '#8a2b2b'; ctx.font = '18px system-ui, sans-serif'; ctx.textAlign = 'center';
ctx.fillText(text, canvas.width / 2, canvas.height / 2);
}
function label(text, x, y, size = 12, color = '#2b2b2b', weight = '') {
ink.font = `${weight} ${size}px system-ui, sans-serif`; ink.fillStyle = color; ink.fillText(text, x, y);
}
// fill() takes a storage pointer: that needs the unrestricted_pointer_parameters extension.
const withStoragePointer = `
fn fill(dst: ptr<storage, array<u32>, read_write>, i: u32) { dst[i] = i * 100; }`;
const portable = `
fn fill(i: u32) { out[i] = i * 100; } // portable: pass an index, not a pointer`;
const lab = (extended) => /* wgsl */ `
@group(0) @binding(0) var<storage, read_write> out: array<u32>;
fn restock(p: ptr<function, u32>) { *p += 10; }
${extended ? withStoragePointer : portable}
@compute @workgroup_size(1) fn main() {
var stock = 5u;
let p = &stock; // a pointer held in a let
restock(p);
restock(&stock);
out[0] = stock;
${extended ? 'fill(&out, 1); // a pointer to a storage buffer' : 'fill(1);'}
let q = &out[2]; // a pointer to one element
*q = *p - 18;
}`;
// The render pass reads out[] directly: one spine per unit of stock in rows 0 and 2.
const view = /* wgsl */ `
@group(0) @binding(0) var<storage> out: array<u32>;
struct Out { @builtin(position) pos: vec4f, @location(0) color: vec4f }
@vertex fn vs(@builtin(vertex_index) v: u32, @builtin(instance_index) i: u32) -> Out {
let q = vec2f(f32(v & 1), f32(v >> 1));
let row = i / 100u; let k = i % 100u;
var px: vec2f; var color = vec4f(0);
if (row == 1u) { // out[1]: one long bar, 4 px per unit
px = vec2f(150, 150) + q * vec2f(f32(out[1]) * 4.2 * f32(k == 0u), 36);
color = vec4f(0.88, 0.55, 0.15, 1);
} else { // rows 0 and 2: a spine per unit
let shown = f32(k < out[row]);
let tints = array(vec3f(0.08, 0.40, 0.75), vec3f(0.55, 0.30, 0.65), vec3f(0.16, 0.56, 0.30));
px = vec2f(150 + f32(k) * 17, 60 + f32(row) * 90 + (1 - q.y) * 50 * shown) + vec2f(q.x * 14 * shown, f32(k % 3) * 4 * shown);
color = vec4f(tints[k % 3], 1);
}
return Out(vec4f(px.x / 300 - 1, 1 - px.y / 160, 0, 1), color);
}
@fragment fn fs(in: Out) -> @location(0) vec4f { return in.color; }`;
async function main() {
const adapter = await navigator.gpu?.requestAdapter();
if (!adapter) return showMessage('WebGPU is not available in this browser');
const device = await adapter.requestDevice();
const context = canvas.getContext('webgpu');
const format = navigator.gpu.getPreferredCanvasFormat();
context.configure({ device, format });
const features = navigator.gpu.wgslLanguageFeatures ?? new Set();
const extended = features.has('unrestricted_pointer_parameters') && features.has('pointer_composite_access');
const B = GPUBufferUsage;
const out = device.createBuffer({ size: 12, usage: B.STORAGE | B.COPY_SRC });
const read = device.createBuffer({ size: 12, usage: B.COPY_DST | B.MAP_READ });
const compute = device.createComputePipeline({ layout: 'auto', compute: { module: device.createShaderModule({ code: lab(extended) }) } });
const module = device.createShaderModule({ code: view });
const render = device.createRenderPipeline({ layout: 'auto', primitive: { topology: 'triangle-strip' },
vertex: { module }, fragment: { module, targets: [{ format }] } });
const encoder = device.createCommandEncoder();
const cp = encoder.beginComputePass();
cp.setPipeline(compute);
cp.setBindGroup(0, device.createBindGroup({ layout: compute.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: out } }] }));
cp.dispatchWorkgroups(1);
cp.end();
const pass = encoder.beginRenderPass({ colorAttachments: [{ view: context.getCurrentTexture().createView(),
clearValue: [0.97, 0.96, 0.93, 1], loadOp: 'clear', storeOp: 'store' }] });
pass.setPipeline(render);
pass.setBindGroup(0, device.createBindGroup({ layout: render.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: out } }] }));
pass.draw(4, 300); // 3 rows of up to 100 instances
pass.end();
encoder.copyBufferToBuffer(out, 0, read, 0, 12);
device.queue.submit([encoder.finish()]);
await read.mapAsync(GPUMapMode.READ);
const [a, b, c] = new Uint32Array(read.getMappedRange());
read.unmap();
label(`out[0] = ${a}`, 14, 84, 15, '#222', 'bold'); label('5, then restock(p),', 14, 102, 11, '#555'); label('then restock(&stock)', 14, 116, 11, '#555');
label(`out[1] = ${b}`, 14, 170, 15, '#222', 'bold'); label(extended ? 'fill(&out, 1): ptr<storage>' : 'fill(1): by index', 14, 188, 11, '#555');
label(`out[2] = ${c}`, 14, 264, 15, '#222', 'bold'); label('*q = *p - 18 via &out[2]', 14, 282, 11, '#555');
label(extended ? 'This browser has unrestricted_pointer_parameters and pointer_composite_access.'
: 'No unrestricted_pointer_parameters here: the portable index version ran.', 14, 312, 11.5, '#1f4f8a');
}
main();
</script>