Attributes (@name or @name(arguments)) attach meaning to declarations. A compute entry point also needs @workgroup_size(x, y, z), with y and z defaulting to 1; without it: "a compute shader must include '@workgroup_size' in its attributes". The attributes of this chapter:
| Attribute | Applies to | Purpose |
|---|---|---|
| @vertex, @fragment, @compute | functions | Mark an entry point and its stage |
| @workgroup_size(x, y, z) | compute entry points | Invocations per workgroup; may use override values |
| @builtin(name) | parameters, members, returns | System values such as position, vertex_index |
| @location(n) | parameters, members, returns | User data between stages, or color target n |
| @align(n), @size(n) | struct members | Override layout to match a host format |
With struct Rec { a: u32, @align(8) b: u32 } the lab placed b at byte 8 instead of 4. Vertex Shader Built-ins to Interpolation cover the @builtin and @location values.
<!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="330"></canvas>
<canvas id="labels" width="600" height="330"></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', align = 'left', font = 'system-ui, sans-serif') {
ink.font = `${size}px ${font}`; ink.fillStyle = color; ink.textAlign = align; ink.fillText(text, x, y);
}
// Two structs that differ only by an attribute; each writes two elements into its own buffer.
const lab = /* wgsl */ `
struct Plain { a: u32, b: u32 }
struct Aligned { a: u32, @align(8) b: u32, @size(8) c: u32 }
@group(0) @binding(0) var<storage, read_write> plain: array<Plain>;
@group(0) @binding(1) var<storage, read_write> aligned: array<Aligned>;
@compute @workgroup_size(1) // required: y and z default to 1
fn main() {
plain[0] = Plain(1, 2); plain[1] = Plain(3, 4);
aligned[0] = Aligned(1, 2, 3); aligned[1] = Aligned(4, 5, 6);
}`;
// Paint every word of both buffers: blue where a value landed, hatched where nothing was written.
const view = /* wgsl */ `
@group(0) @binding(0) var<storage> plain: array<u32>; // 4 words
@group(0) @binding(1) var<storage> aligned: array<u32>; // 12 words
struct Out { @builtin(position) pos: vec4f, @location(0) @interpolate(flat) i: u32 }
@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 = select(0u, 1u, i >= 4);
let col = select(i, i - 4, i >= 4);
let px = vec2f(40 + f32(col) * 44, 80 + f32(row) * 120) + q * vec2f(40, 46);
return Out(vec4f(px.x / 300 - 1, 1 - px.y / 165, 0, 1), i);
}
@fragment fn fs(in: Out) -> @location(0) vec4f {
let value = select(plain[min(in.i, 3u)], aligned[max(in.i, 4u) - 4], in.i >= 4);
if (value != 0) { return vec4f(0.20, 0.45, 0.75, 1); }
let stripe = step(0.5, fract((in.pos.x + in.pos.y) / 10));
return vec4f(mix(vec3f(0.87), vec3f(0.77), stripe), 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'); // @compute without @workgroup_size: expected error
const noSize = await device.createShaderModule({ code: '@compute fn main() {}' }).getCompilationInfo();
await device.popErrorScope();
const B = GPUBufferUsage;
const plain = device.createBuffer({ size: 16, usage: B.STORAGE | B.COPY_SRC }); // 2 x 8 bytes
const aligned = device.createBuffer({ size: 48, usage: B.STORAGE | B.COPY_SRC }); // 2 x 24 bytes
const read = device.createBuffer({ size: 64, usage: B.COPY_DST | B.MAP_READ });
const compute = device.createComputePipeline({ layout: 'auto', compute: { module: device.createShaderModule({ code: lab }) } });
const module = device.createShaderModule({ code: view });
const render = device.createRenderPipeline({ layout: 'auto', primitive: { topology: 'triangle-strip' },
vertex: { module }, fragment: { module, targets: [{ format }] } });
const both = (pipeline) => device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [
{ binding: 0, resource: { buffer: plain } }, { binding: 1, resource: { buffer: aligned } }] });
const encoder = device.createCommandEncoder();
const cp = encoder.beginComputePass();
cp.setPipeline(compute);
cp.setBindGroup(0, both(compute));
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, both(render));
pass.draw(4, 16);
pass.end();
encoder.copyBufferToBuffer(plain, 0, read, 0, 16);
encoder.copyBufferToBuffer(aligned, 0, read, 16, 48);
device.queue.submit([encoder.finish()]);
await read.mapAsync(GPUMapMode.READ);
const w = new Uint32Array(read.getMappedRange().slice(0));
read.unmap();
label('struct Plain { a: u32, b: u32 }: size 8, two elements', 40, 64, 12.5, '#222', 'left', 'ui-monospace, monospace');
label('struct Aligned { a: u32, @align(8) b: u32, @size(8) c: u32 }: size 24', 40, 184, 12, '#222', 'left', 'ui-monospace, monospace');
for (let i = 0; i < 16; i++) {
const row = i < 4 ? 0 : 1, col = i < 4 ? i : i - 4;
const x = 60 + col * 44, y = 80 + row * 120;
label(String(col * 4), x, y + 62, 10, '#777', 'center');
label(w[i] ? String(w[i]) : '-', x, y + 29, 15, w[i] ? '#fff' : '#888', 'center');
}
label('byte offsets', 40, 158, 10, '#777'); label('byte offsets', 40, 278, 10, '#777');
label('@align(8) moves b from byte 4 to byte 8; @size(8) gives c 8 bytes; the size rounds up to 24.', 12, 300, 11.5, '#333');
const e = noSize.messages.find((m) => m.type === 'error');
label(e ? e.message : '', 12, 320, 11.5, '#8a2b2b');
}
main();
</script>