Uniform Alignment

Structuring a Uniform Buffer with WGSL Alignment Rules

Arrays and Alignment gave the alignment table. In practice, write the struct with its offsets in comments, order members from largest to smallest, and let 4-byte scalars fill the spare slot after a vec3f. JavaScript then writes each member at its byte offset through typed-array views of one ArrayBuffer, and a compute shader proves the layout by reading the members back:

A per-frame uniform struct packed from JavaScript and read back by WGSLJavaScript
const device = await (await navigator.gpu.requestAdapter()).requestDevice();
const B = GPUBufferUsage;
const code = /* wgsl */ `
struct Frame {
  viewProj: mat4x4f,     // offset  0, size 64
  eye: vec3f,            // offset 64: a vec3f aligns to 16
  time: f32,             // offset 76: the vec3f's spare 4 bytes
  viewport: vec2f,       // offset 80: a vec2f aligns to 8
  selected: u32,         // offset 88
}                        // size 96: rounded up to the 16-byte struct alignment
@group(0) @binding(0) var<uniform> frame: Frame;
@group(0) @binding(1) var<storage, read_write> out: array<f32, 5>;
@compute @workgroup_size(1) fn main() {
  out = array(frame.viewProj[3].z, frame.eye.y, frame.time, frame.viewport.x,
              f32(frame.selected));
}`;
const bytes = new ArrayBuffer(96);                     // one Frame, two views of its bytes
const f32 = new Float32Array(bytes), u32 = new Uint32Array(bytes);
f32.set([1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, -5, 1]);   // viewProj, column by column
f32.set([0, 1.6, 4], 64 / 4);                          // eye
f32[76 / 4] = 12.5;                                    // time
f32.set([480, 330], 80 / 4);                           // viewport
u32[88 / 4] = 4;                                       // selected book id
const uniforms = device.createBuffer({ size: 96, usage: B.UNIFORM | B.COPY_DST });
device.queue.writeBuffer(uniforms, 0, bytes);
const [out, read] = [B.STORAGE | B.COPY_SRC, B.COPY_DST | B.MAP_READ].map((usage) =>
  device.createBuffer({ size: 20, usage }));
const pipeline = device.createComputePipeline({ layout: 'auto',
  compute: { module: device.createShaderModule({ code }) } });
const encoder = device.createCommandEncoder(), pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
pass.setBindGroup(0, device.createBindGroup({ layout: pipeline.getBindGroupLayout(0),
  entries: [{ binding: 0, resource: uniforms }, { binding: 1, resource: out }] }));
pass.dispatchWorkgroups(1), pass.end();
encoder.copyBufferToBuffer(out, read), device.queue.submit([encoder.finish()]);
await read.mapAsync(GPUMapMode.READ);
console.log([...new Float32Array(read.getMappedRange())].join(', '));
Output
-5, 1.600000023841858, 12.5, 480, 4

Every member arrived intact (1.6 is not exact in 32 bits). The usual mistakes are a vec3f followed by another vec3f (the second starts 16 bytes later, not 12), mat3x3f (three 16-byte columns, 48 bytes, not 36) and a u32 written through the Float32Array, which stores the float's bit pattern instead of the integer. Keep offsets next to the WGSL, or generate both from one description as TypeGPU 320,420 does (TypeGPU).

A 96-byte Frame uniform packed from JavaScript at WGSL offsets, drawn word by word and proven by a compute read-backHTMLLive
<!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);
}

const code = /* wgsl */ `
struct Frame {
  viewProj: mat4x4f,     // offset  0, size 64
  eye: vec3f,            // offset 64: a vec3f aligns to 16
  time: f32,             // offset 76: the vec3f's spare 4 bytes
  viewport: vec2f,       // offset 80: a vec2f aligns to 8
  selected: u32,         // offset 88
}                        // size 96: rounded up to the 16-byte struct alignment
@group(0) @binding(0) var<uniform> frame: Frame;
@group(0) @binding(1) var<storage, read_write> out: array<f32, 5>;
@compute @workgroup_size(1) fn main() {
  out = array(frame.viewProj[3].z, frame.eye.y, frame.time, frame.viewport.x, f32(frame.selected));
}`;
// The same 96 bytes bound again with a different WGSL type: 24 raw words, one cell each.
const view = /* wgsl */ `
@group(0) @binding(0) var<uniform> words: array<vec4u, 6>;
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 px = vec2f(14 + f32(i % 8) * 72, 40 + f32(i / 8) * 64) + q * vec2f(68, 44);
  return Out(vec4f(px.x / 300 - 1, 1 - px.y / 165, 0, 1), i);
}
@fragment fn fs(in: Out) -> @location(0) vec4f {
  let word = words[in.i / 4][in.i % 4];
  var color = vec3f(0.30, 0.50, 0.80);                              // viewProj
  if (in.i >= 16u && in.i < 19u) { color = vec3f(0.30, 0.65, 0.40); } // eye
  if (in.i == 19u) { color = vec3f(0.90, 0.60, 0.20); }               // time
  if (in.i == 20u || in.i == 21u) { color = vec3f(0.60, 0.40, 0.70); } // viewport
  if (in.i == 22u) { color = vec3f(0.80, 0.35, 0.30); }               // selected
  if (in.i == 23u) { color = vec3f(0.85); }                           // padding
  return vec4f(select(mix(color, vec3f(1), 0.7), color, word != 0u), 1);   // pale when zero
}`;

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 bytes = new ArrayBuffer(96);                     // one Frame, two views of its bytes
  const f32 = new Float32Array(bytes), u32 = new Uint32Array(bytes);
  f32.set([1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, -5, 1]);   // viewProj, column by column
  f32.set([0, 1.6, 4], 64 / 4);                          // eye
  f32[76 / 4] = 12.5;                                    // time
  f32.set([480, 330], 80 / 4);                           // viewport
  u32[88 / 4] = 4;                                       // selected book id: an integer, via Uint32Array
  const uniforms = device.createBuffer({ size: 96, usage: B.UNIFORM | B.COPY_DST });
  device.queue.writeBuffer(uniforms, 0, bytes);
  const [out, read] = [B.STORAGE | B.COPY_SRC, B.COPY_DST | B.MAP_READ].map((usage) => device.createBuffer({ size: 20, usage }));
  const compute = device.createComputePipeline({ layout: 'auto', compute: { module: device.createShaderModule({ code }) } });
  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 pass = encoder.beginComputePass();
  pass.setPipeline(compute);
  pass.setBindGroup(0, device.createBindGroup({ layout: compute.getBindGroupLayout(0),
    entries: [{ binding: 0, resource: { buffer: uniforms } }, { binding: 1, resource: { buffer: out } }] }));
  pass.dispatchWorkgroups(1);
  pass.end();
  encoder.copyBufferToBuffer(out, 0, read, 0, 20);
  const draw = encoder.beginRenderPass({ colorAttachments: [{ view: context.getCurrentTexture().createView(),
    clearValue: [0.97, 0.96, 0.93, 1], loadOp: 'clear', storeOp: 'store' }] });
  draw.setPipeline(render);
  draw.setBindGroup(0, device.createBindGroup({ layout: render.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: uniforms } }] }));
  draw.draw(4, 24);
  draw.end();
  device.queue.submit([encoder.finish()]);
  await read.mapAsync(GPUMapMode.READ);
  const back = [...new Float32Array(read.getMappedRange())];
  read.unmap();

  ink.font = 'bold 12.5px system-ui, sans-serif'; ink.fillStyle = '#222';
  ink.fillText('struct Frame: 96 bytes = 24 words (8 per row)', 14, 26);
  ink.font = '10.5px ui-monospace, monospace'; ink.textAlign = 'center';
  const names = [...Array(16).fill('viewProj'), 'eye.x', 'eye.y', 'eye.z', 'time', 'viewport.x', 'viewport.y', 'selected', 'pad'];
  for (let i = 0; i < 24; i++) {
    const x = 14 + (i % 8) * 72 + 34, y = 40 + Math.floor(i / 8) * 64;
    const value = i === 22 ? u32[i] : +f32[i].toFixed(2);
    ink.fillStyle = '#111'; ink.fillText(String(value), x, y + 20);
    ink.fillStyle = '#333'; ink.fillText(i < 16 ? `[${i >> 2}][${i & 3}]` : names[i], x, y + 36);
    ink.fillStyle = '#777'; ink.fillText(`@${i * 4}`, x, y + 56);
  }
  ink.textAlign = 'left'; ink.font = '12px system-ui, sans-serif'; ink.fillStyle = '#1e6b3a';
  ink.fillText(`Read back by WGSL: viewProj[3].z = ${back[0]}, eye.y = ${back[1].toFixed(7)}, time = ${back[2]},`, 14, 262);
  ink.fillText(`viewport.x = ${back[3]}, selected = ${back[4]}: every member arrived intact (1.6 is not exact in f32).`, 14, 280);
  ink.fillStyle = '#8a2b2b';
  ink.fillText('Classic mistakes: vec3f then vec3f (16 bytes apart, not 12), mat3x3f (48 bytes, not 36),', 14, 306);
  ink.fillText('and a u32 written through the Float32Array (the float\'s bit pattern, not the integer).', 14, 322);
}
main();
</script>