array<T, N> has a fixed size; a runtime-sized array<T> lives only in storage, sized by the bound buffer (arrayLength() reads it). Structs get a constructor. Each member's offset must be a multiple of its alignment:
| Type | Align | Size |
|---|---|---|
| i32, u32, f32 | 4 | 4 |
| vec2f | 8 | 8 |
| vec3f | 16 | 12 |
| vec4f, mat4x4f column | 16 | 16 |
| array<T, N> | align of T | N x stride |
| struct | largest member align | end rounded up to its align |
An array's stride is the element size rounded up to its alignment: 16 bytes for array<vec3f>. The lab shows the rules by writing a struct and printing the raw words of the zeroed buffer:
// out: 16 x u32
struct Spine { id: u32, tint: vec3u, pages: u32 }
@group(0) @binding(0) var<storage, read_write> out: array<Spine>;
@compute @workgroup_size(1) fn main() {
out[0] = Spine(1, vec3u(31, 95, 139), 312);
out[1].id = arrayLength(&out); // 64 bytes / 32-byte stride
}Output
out: 1, 0, 0, 0, 31, 95, 139, 312, 2, 0, 0, 0, 0, 0, 0, 0
tint cannot start at byte 4 because a vec3u aligns to 16, so bytes 4-15 are padding, and pages fills the vec3's spare slot at byte 28: 32 bytes per struct. A tightly packed Float32Array knows none of this, the root of most "garbage in my uniform" bugs.
<!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', align = 'center', weight = '') {
ink.font = `${weight} ${size}px system-ui, sans-serif`; ink.fillStyle = color; ink.textAlign = align;
ink.fillText(text, x, y);
}
// The lab shader from the book: one struct, written by a compute shader.
const lab = /* wgsl */ `
struct Spine { id: u32, tint: vec3u, pages: u32 }
@group(0) @binding(0) var<storage, read_write> out: array<Spine>;
@compute @workgroup_size(1) fn main() {
out[0] = Spine(1, vec3u(31, 95, 139), 312);
out[1].id = arrayLength(&out); // 64 bytes / 32-byte stride
}`;
// A render pass reads the same 16 words and paints one cell per word.
const view = /* wgsl */ `
@group(0) @binding(0) var<storage> words: array<u32>;
const role = array(0u, 3u, 3u, 3u, 1u, 1u, 1u, 2u); // id, padding x3, tint.xyz, pages
struct Out { @builtin(position) pos: vec4f, @location(0) @interpolate(flat) word: 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 cell = vec2f(f32(i % 8), f32(i / 8));
let px = vec2f(62, 70) + cell * vec2f(66, 90) + q * vec2f(62, 60);
return Out(vec4f(px.x / 300 - 1, 1 - px.y / 170, 0, 1), i);
}
@fragment fn fs(in: Out) -> @location(0) vec4f {
let value = words[in.word];
let tint = vec3f(f32(words[4]), f32(words[5]), f32(words[6])) / 255; // the stored colour
switch role[in.word % 8] {
case 0u: { return vec4f(0.20, 0.45, 0.75, 1); } // id
case 1u: { return vec4f(select(vec3f(0.85), tint, value > 0), 1); } // tint.x, .y, .z
case 2u: { return vec4f(select(vec3f(0.85), vec3f(0.88, 0.55, 0.15), value > 0), 1); }
default: { // padding: hatched
let stripe = step(0.5, fract((in.pos.x + in.pos.y) / 10));
return vec4f(mix(vec3f(0.86), vec3f(0.76), 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 });
const B = GPUBufferUsage;
const out = device.createBuffer({ size: 64, usage: B.STORAGE | B.COPY_SRC }); // zero-filled
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 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, 16); // 16 words = 64 bytes
pass.end();
encoder.copyBufferToBuffer(out, 0, read, 0, 64);
device.queue.submit([encoder.finish()]);
await read.mapAsync(GPUMapMode.READ); // the numbers, for the labels
const words = new Uint32Array(read.getMappedRange().slice(0));
read.unmap();
label('struct Spine { id: u32, tint: vec3u, pages: u32 }: 32 bytes per element', 300, 24, 14, '#222', 'center', 'bold');
const names = ['id', 'pad', 'pad', 'pad', 'tint.x', 'tint.y', 'tint.z', 'pages'];
for (let i = 0; i < 16; i++) {
const x = 62 + (i % 8) * 66 + 31, y = 70 + Math.floor(i / 8) * 90;
label(`byte ${(i % 8) * 4}`, x, y - 6, 10.5, '#777');
label(String(words[i]), x, y + 30, 14, '#111', 'center', 'bold');
label(names[i % 8], x, y + 50, 11, '#333');
}
label('out[0]', 12, 104, 12, '#555', 'left'); label('out[1]', 12, 194, 12, '#555', 'left');
label('tint aligns to 16, so bytes 4-15 are padding; pages fills the vec3u\'s spare slot at byte 28.', 300, 270, 12, '#333');
label(`out[1].id = arrayLength(&out) = ${words[8]}. A tightly packed Float32Array [1, 31, 95, 139, 312]`, 300, 296, 12, '#8a2b2b');
label('would put tint at byte 4 and pages at byte 16: the root of "garbage in my uniform" bugs.', 300, 314, 12, '#8a2b2b');
}
main();
</script>