An entry point is a function the pipeline calls, marked with its stage. A module may hold many, sharing helpers and structs; the pipeline picks one per stage with entryPoint, optional when there is only one:
| Stage | Called once per | Must return | Typical inputs |
|---|---|---|---|
| @vertex | vertex (per instance) | @builtin(position) | vertex and instance index, vertex buffers |
| @fragment | fragment (covered pixel) | @location(n) colors, or nothing | interpolated vertex outputs |
| @compute | invocation in a workgroup grid | nothing | invocation IDs |
Inputs and outputs carry @builtin or @location, usually grouped in a struct that the vertex stage returns and the fragment stage receives (First Shader Pair). A vertex shader returning only @location(0) vec4f failed with "a vertex shader must include the 'position' builtin in its return type"; a fragment shader returning nothing compiled.
<!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="340"></canvas>
<canvas id="labels" width="600" height="340"></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 = '', font = 'system-ui, sans-serif') {
ink.font = `${weight} ${size}px ${font}`; ink.fillStyle = color; ink.fillText(text, x, y);
}
// Shared structs and helpers, then one entry point per stage.
const code = /* wgsl */ `
struct Book { pos: vec2f, tint: vec3f } // 32-byte stride in storage
@group(0) @binding(0) var<uniform> time: f32;
@group(0) @binding(1) var<storage, read_write> books: array<Book>;
@compute @workgroup_size(6) // called once per invocation
fn place(@builtin(local_invocation_index) i: u32) {
let a = time + f32(i) * 1.0472;
books[i] = Book(vec2f(cos(a) * 0.55 - 0.35, sin(a) * 0.6), vec3f(0.3 + 0.12 * f32(i), 0.4, 0.8 - 0.1 * f32(i)));
}
@group(0) @binding(2) var<storage> shelf: array<Book>; // the same buffer, read-only
struct Varyings { @builtin(position) pos: vec4f, @location(0) tint: vec3f, @location(1) uv: vec2f }
@vertex // called once per vertex
fn vs(@builtin(vertex_index) v: u32, @builtin(instance_index) i: u32) -> Varyings {
let q = vec2f(f32(v & 1), f32(v >> 1));
let b = shelf[i];
return Varyings(vec4f(b.pos + (q - 0.5) * vec2f(0.12, 0.34), 0, 1), b.tint, q); // must return position
}
@fragment // called once per covered pixel
fn fs(in: Varyings) -> @location(0) vec4f {
let spine = smoothstep(0.0, 0.2, in.uv.x);
return vec4f(in.tint * (0.6 + 0.4 * spine), 1);
}`;
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 });
device.pushErrorScope('validation'); // a vertex shader without position: an expected error
const bad = await device.createShaderModule({ code: '@vertex fn vs() -> @location(0) vec4f { return vec4f(1); }' }).getCompilationInfo();
await device.popErrorScope();
const module = device.createShaderModule({ code }); // one module, three entry points
const compute = device.createComputePipeline({ layout: 'auto', compute: { module, entryPoint: 'place' } });
const render = device.createRenderPipeline({ layout: 'auto', primitive: { topology: 'triangle-strip' },
vertex: { module, entryPoint: 'vs' }, fragment: { module, entryPoint: 'fs', targets: [{ format }] } });
const B = GPUBufferUsage;
const time = device.createBuffer({ size: 4, usage: B.UNIFORM | B.COPY_DST });
const books = device.createBuffer({ size: 6 * 32, usage: B.STORAGE });
const computeGroup = device.createBindGroup({ layout: compute.getBindGroupLayout(0), entries: [
{ binding: 0, resource: { buffer: time } }, { binding: 1, resource: { buffer: books } }] });
const renderGroup = device.createBindGroup({ layout: render.getBindGroupLayout(0), entries: [{ binding: 2, resource: { buffer: books } }] });
const e = bad.messages.find((m) => m.type === 'error');
const notes = [['@compute place()', 'once per invocation in the grid', 'returns nothing; gets IDs'],
['@vertex vs()', 'once per vertex (per instance)', 'must return @builtin(position)'],
['@fragment fs()', 'once per covered pixel', 'returns @location(n) colours']];
function frame(now) {
device.queue.writeBuffer(time, 0, new Float32Array([now / 1500]));
const encoder = device.createCommandEncoder();
const cp = encoder.beginComputePass();
cp.setPipeline(compute);
cp.setBindGroup(0, computeGroup);
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, renderGroup);
pass.draw(4, 6);
pass.end();
device.queue.submit([encoder.finish()]);
requestAnimationFrame(frame);
}
notes.forEach(([name, when, rule], i) => {
label(name, 400, 60 + i * 70, 13, '#222', 'bold', 'ui-monospace, monospace');
label(when, 400, 80 + i * 70, 12, '#444'); label(rule, 400, 96 + i * 70, 12, '#1f4f8a');
});
label('A vertex shader returning only @location(0):', 12, 300, 12, '#222');
label(e ? e.message.slice(0, 84) : 'compiled?', 12, 318, 11.5, '#8a2b2b', '', 'ui-monospace, monospace');
requestAnimationFrame(frame);
}
main();
</script>