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

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

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.

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>.

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.

@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

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

• 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.

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.

Related

• Memory Alignment for Spatial Data Buffers — the full alignment reference this padding case belongs to.
• Structuring Uniform Buffers for Coordinate Alignment — the same 16-byte rules applied to uniform blocks.
• VRAM Budget Management Across Tile Zoom Levels — how per-vertex stride sets the VRAM cost of every tile.
• Spatial Aggregation in GPU Memory — consumes packed position buffers as aggregation input.

Up: Memory Alignment for Spatial Data Buffers