Packing vec3 Attributes Without Wasted Padding

The waste this page eliminates is the hidden fourth lane WGSL reserves for every vec3<f32>. A lon/lat/elev triplet is three f32 values — twelve bytes of real data — but WGSL’s memory layout rules give vec3<f32> a size of 12 bytes and an alignment of 16 bytes, so any array element or struct field that follows it is pushed to the next 16-byte boundary. Declare a storage array of vec3<f32> and every element occupies a 16-byte stride: four bytes per vertex are pure padding, a 33% VRAM tax and a 33% cut in memory bandwidth on a buffer that may hold tens of millions of terrain or point-cloud vertices. Worse, if you upload a tightly packed twelve-byte-stride buffer from the CPU and bind it as array<vec3<f32>>, the shader reads every element after the first at the wrong offset — the data is silently scrambled with no validation error. This page covers the two correct ways to carry three-component spatial coordinates without paying the padding: pack them as three separate scalar arrays that align at 4 bytes, or accept the 16-byte stride explicitly as a vec4<f32> and put the fourth lane to work.

Padding is the failure that bites first when the coordinate layout is wrong, so this sits under the memory alignment for spatial data buffers reference, alongside structuring uniform buffers for coordinate alignment. The same 16-byte stride that inflates a vertex buffer also determines how much VRAM a tile costs, so the packing choice here feeds directly into VRAM budget management across tile zoom levels.

The alignment rule and its three layouts

Three ways to store one 3D point Three byte strips share a ruler from zero to sixteen bytes. The first stores an array of vec3 f32: x, y and z fill twelve bytes and four bytes of padding round the element stride to sixteen. The second stores an array of vec4 f32: x, y and z fill twelve bytes and the fourth lane carries a real attribute such as intensity, so all sixteen bytes are used. The third is structure-of-arrays, three separate f32 arrays, where each point contributes four bytes to each of three buffers, twelve bytes in total with no padding anywhere. BYTE OFFSET 0 4 8 12 16 x y z pad vec3 f32 array 16 B · 4 B wasted the pad is unavoidable — vec3 aligns to 16 x y z w = attr vec4 f32 array 16 B · 0 wasted same 16 bytes, but the fourth lane carries intensity or a class id x y z 3 × f32 arrays 12 B · 0 wasted three buffers; each workgroup streams one axis contiguously Only split into structure-of-arrays when a kernel genuinely reads one axis at a time.
The middle layout is the one most teams should reach for. It costs exactly what the padded vec3 costs and hands back a usable fourth channel, which is why carrying intensity or a class id there is effectively free.

WGSL’s AlignOf and SizeOf rules fix the offsets. For the scalar and vector types involved:

Type Align (bytes) Size (bytes) Array element stride
f32 4 4 4
vec2<f32> 8 8 8
vec3<f32> 16 12 16 (padded)
vec4<f32> 16 16 16

The trap is the vec3<f32> row: size 12, align 16, so an array<vec3<f32>> has a 16-byte stride, not 12. Three layouts carry a lon/lat/elev triplet; only the first two avoid wasted memory.

Layout Bytes / vertex Padding When to use
array<vec3<f32>> 16 4 (25% of stride) Never for large buffers — the implicit pad is pure waste.
Three array<f32> (SoA) 12 0 Densest option; ideal when passes touch one component (e.g. elevation only).
array<vec4<f32>> 16 0 (4th lane used) When the fourth lane carries real data — weight, time, or a packed attribute.

Runnable reference implementation

The Structure-of-Arrays layout is the densest: pack lon, lat, and elev into three contiguous f32 arrays, each aligning at 4 bytes with zero stride padding. The packer below writes one interleaving-free buffer per component; the WGSL side binds them as three independent array<f32>.

typescript
interface Triplet { lon: number; lat: number; elev: number; }

// Structure-of-Arrays: three tightly packed f32 arrays, 12 bytes/vertex, zero padding.
function packSoA(points: Triplet[]): { lon: Float32Array; lat: Float32Array; elev: Float32Array } {
  const n = points.length;
  const lon = new Float32Array(n);
  const lat = new Float32Array(n);
  const elev = new Float32Array(n);
  for (let i = 0; i < n; i++) {
    lon[i] = points[i].lon;
    lat[i] = points[i].lat;
    elev[i] = points[i].elev;   // each array is 4-byte-strided, no gaps
  }
  return { lon, lat, elev };
}

// Alternative: padded vec4 with the 4th lane carrying real data (here, a weight).
// 16 bytes/vertex, but zero WASTED bytes because .w is used, not padding.
function packVec4(points: Triplet[], weight: (i: number) => number): Float32Array {
  const out = new Float32Array(points.length * 4);   // stride 4 floats = 16 bytes
  for (let i = 0; i < points.length; i++) {
    const b = i * 4;
    out[b + 0] = points[i].lon;
    out[b + 1] = points[i].lat;
    out[b + 2] = points[i].elev;
    out[b + 3] = weight(i);     // .w is a payload, not implicit padding
  }
  return out;
}

// Upload SoA components as three storage buffers.
function uploadSoA(device: GPUDevice, soa: ReturnType<typeof packSoA>) {
  const make = (data: Float32Array) => {
    const buf = device.createBuffer({
      size: data.byteLength,      // exact 12*n across the three, no rounding
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
    });
    device.queue.writeBuffer(buf, 0, data);
    return buf;
  };
  return { lon: make(soa.lon), lat: make(soa.lat), elev: make(soa.elev) };
}

The matching WGSL binds the three scalar arrays. Reassemble a vec3<f32> in registers when a pass needs all three; a pass that only reads elevation touches one third of the bandwidth.

wgsl
@group(0) @binding(0) var<storage, read> lon:  array<f32>;
@group(0) @binding(1) var<storage, read> lat:  array<f32>;
@group(0) @binding(2) var<storage, read> elev: array<f32>;

@compute @workgroup_size(256)
fn transform(@builtin(global_invocation_id) gid: vec3<u32>) {
    let i = gid.x;
    if (i >= arrayLength(&lon)) { return; }
    // Reassemble only in registers — the buffers stay 12-byte-strided in VRAM.
    let coord = vec3<f32>(lon[i], lat[i], elev[i]);
    // ... project / cull / write ...
}

// If you must keep an interleaved buffer, declare the stride explicitly as vec4
// so the CPU layout and the shader layout agree. NEVER bind a packed 12-byte
// buffer as array<vec3<f32>> — the 16-byte stride reads every element but the
// first from the wrong offset.
struct Vertex {
    pos: vec3<f32>,   // occupies lanes 0..2
    weight: f32,      // lane 3 — fills what would otherwise be padding
}
@group(0) @binding(3) var<storage, read> verts: array<Vertex>; // 16-byte stride, 0 waste

For the interleaved case, the struct Vertex layout is the safe form: vec3<f32> pos occupies the first three lanes and f32 weight fills the fourth, so the 16-byte stride carries a payload in every byte. The CPU-side packVec4 writes exactly this layout, so the offsets match and nothing is scrambled.

Parameter reference

Choosing between the three layouts A decision table with three columns judging each layout against four questions. For a kernel that reads whole positions, the padded vec3 array works, the vec4 array is the best fit, and structure-of-arrays is poor because it forces three separate reads. For a kernel that reads one axis, such as a height histogram, both vector layouts are wasteful and structure-of-arrays is the best fit. For upload from a Python packer, vec3 and vec4 are single-buffer writes while structure-of-arrays needs three coordinated writes. Bytes per point are sixteen, sixteen and twelve respectively. PICK BY ACCESS PATTERN array<vec3> array<vec4> 3 × array<f32> Reads whole positions works best fit three reads Reads one axis only wasteful wasteful best fit Upload from Python one write one write three writes Bytes per point 16 B 16 B 12 B Mixed workloads can keep both: a packed vec4 for rendering, plus one f32 array for the analytic pass that needs it.
There is no universally correct answer, only a correct answer per kernel. The trap is picking structure-of-arrays for elegance and then writing a shader that reads all three axes on every invocation — which turns one coalesced read into three.
Value Setting Guidance
vec3<f32> stride 16 bytes (fixed) WGSL rule, not tunable. Never assume 12; an array<vec3<f32>> element i starts at 16*i.
SoA buffer count 3 One array<f32> per component. Zero padding; best when passes read a subset of components.
SoA bytes/vertex 12 Sum of three f32. The densest correct layout for lon/lat/elev.
vec4 bytes/vertex 16 Only “free” if .w carries real data; otherwise it is the padded vec3 waste renamed.
writeBuffer size exact byteLength For SoA, 4 * n per component. For vec4, 16 * n. No 16-byte rounding needed at the buffer level.
Vertex-buffer arrayStride 12 (SoA attrs) or 16 (vec4) In a render pipeline’s vertex layout, match arrayStride to the packed stride, not to vec3 alignment.

Vertex-buffer attributes in a render pipeline are exempt from the 16-byte vec3 rule — a float32x3 attribute can sit at a 12-byte arrayStride because vertex fetch uses explicit offsets. The 16-byte trap applies to vec3<f32> inside storage and uniform structs. For the authoritative alignment and stride rules, see the WGSL specification and the W3C WebGPU specification.

Failure modes

Bytes per vertex, and what the fourth lane buys A bar chart in bytes per vertex for the three layouts. The vec3 array is sixteen bytes: twelve of coordinate payload and four of padding that carries nothing. The vec4 array is also sixteen bytes, but all sixteen are payload because the fourth lane holds an attribute. Structure-of-arrays is twelve bytes, all payload, spread across three buffers. BYTES PER VERTEX vec3 f32 array 16 B vec4 f32 array 16 B 3 × f32 arrays 12 B 0 4 8 12 16 20 B payload padding At forty million points that four-byte pad is 160 MiB of VRAM holding nothing.
Identical totals, different value. The first and second rows cost the same sixteen bytes; only the second one gets an attribute for the money.
  • Scrambled coordinates from a 12-byte upload bound as vec3. Uploading a tightly packed 12-byte-stride buffer and declaring it array<vec3<f32>> makes the shader read element i at offset 16*i, so every vertex after the first is misaligned — with no validation error. Detection: the first point is correct and the rest drift progressively off position. Fix: either pad the CPU buffer to a 16-byte stride (write a dummy .w) or switch to the SoA layout of three array<f32>.
  • Silent 33% VRAM inflation. Declaring array<vec3<f32>> for a large point cloud costs 16 bytes per vertex where 12 would do, quietly enlarging every tile buffer. Detection: measured buffer size is 16*n when you budgeted 12*n; VRAM headroom vanishes a zoom level early. Fix: use SoA scalar arrays, or fill the fourth lane with a real attribute so the stride is not waste.
  • Mismatched arrayStride in a vertex layout. Setting arrayStride: 16 for a packed 12-byte SoA vertex buffer (or vice versa) shifts the vertex fetch, reading each vertex from the wrong offset. Detection: geometry renders warped or exploded. Fix: set the render pipeline’s arrayStride to the exact packed stride — 12 for a float32x3 SoA attribute, 16 for a float32x4.
  • Struct field pushed past the buffer by trailing pad. A struct ending in vec3<f32> is rounded up to a 16-byte size, so an array of it strides at 16; sizing the buffer at 12*n leaves the last element out of bounds. Detection: a GPUValidationError that the bound range is too small, or a dropped final vertex. Fix: size the buffer at the struct’s true 16-byte SizeOf times n, or reorder fields so the vec3 is not last.

Backend / Python interop note

When the coordinates originate server-side, emit them in the layout the GPU binds so no client-side repacking is needed. For the SoA path, write three contiguous C-ordered float32 columns; for the vec4 path, interleave four float32 per vertex with the weight in lane three.

python
import numpy as np

# lon, lat, elev as separate 1-D float32 columns → three SoA storage buffers.
lon = coords[:, 0].astype(np.float32)
lat = coords[:, 1].astype(np.float32)
elev = coords[:, 2].astype(np.float32)
soa = [np.ascontiguousarray(c).tobytes() for c in (lon, lat, elev)]  # 12 bytes/vertex total

# Or a padded vec4 with a real weight in .w (16 bytes/vertex, zero wasted).
v4 = np.empty((len(coords), 4), dtype=np.float32)
v4[:, 0:3] = coords[:, 0:3]
v4[:, 3] = weights                      # .w carries data, not padding
payload = np.ascontiguousarray(v4).tobytes()

Never ship a raw (n, 3) float32 array expecting the shader to read it as array<vec3<f32>> — its 12-byte NumPy stride will not match the shader’s 16-byte stride. Split it into SoA columns or expand it to (n, 4) first.

Up: Memory Alignment for Spatial Data Buffers