diff --git a/apps/typegpu-docs/src/examples/rendering/mesh-primitives/blob.ts b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/blob.ts
new file mode 100644
index 0000000000..438cc61f5b
--- /dev/null
+++ b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/blob.ts
@@ -0,0 +1,19 @@
+import { d, std } from 'typegpu';
+import { meshes } from '@typegpu/geometry';
+import { paint } from './geometry.ts';
+import { sceneLayout } from './shading.ts';
+
+const surface = meshes.parametric(
+  {
+    at: (u, v) => {
+      'use gpu';
+      const time = sceneLayout.$.uniforms.time;
+      const p = meshes.surfaces.sphere(u, v);
+      const ripple = std.sin(p.x * 4 + time * 2) * std.sin(p.y * 3 - time) * std.sin(p.z * 5);
+      return p * (0.55 + 0.15 * ripple) + d.vec3f(0, 0.55, 0);
+    },
+  },
+  { cols: 96, rows: 64 },
+);
+
+export const blob = paint(surface, d.vec3f(0.9));
diff --git a/apps/typegpu-docs/src/examples/rendering/mesh-primitives/gallery.ts b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/gallery.ts
new file mode 100644
index 0000000000..6d588ec97b
--- /dev/null
+++ b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/gallery.ts
@@ -0,0 +1,14 @@
+import { std } from 'typegpu';
+import { meshes } from '@typegpu/geometry';
+import { place } from './geometry.ts';
+
+const floor = meshes.plane({ width: 6, depth: 6 });
+const twoSidedFloor = meshes.concat(floor, meshes.transform(floor, std.rotationX4(Math.PI)));
+
+export const gallery = meshes.concat(
+  place(twoSidedFloor, [0, -0.51, 0], [0.75, 0.75, 0.8]),
+  place(meshes.box(), [-1.5, 0, 1], [0.9, 0.35, 0.3]),
+  place(meshes.cylinder(), [1.5, 0, 1], [0.3, 0.6, 0.9]),
+  place(meshes.torus(), [1.5, -0.35, -1], [0.95, 0.75, 0.2]),
+  place(meshes.sphere(), [-1.5, 0, -1], [0.4, 0.85, 0.5]),
+);
diff --git a/apps/typegpu-docs/src/examples/rendering/mesh-primitives/geometry.ts b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/geometry.ts
new file mode 100644
index 0000000000..1b9485ac70
--- /dev/null
+++ b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/geometry.ts
@@ -0,0 +1,23 @@
+import { d, std } from 'typegpu';
+import { meshes } from '@typegpu/geometry';
+
+const Paint = d.struct({ color: d.vec3f });
+
+export function paint<V extends d.WgslStruct>(
+  g: meshes.IndexedGeometry<V>,
+  color: d.InferInput<d.Vec3f>,
+) {
+  const rgb = d.vec3f(color[0], color[1], color[2]);
+  return meshes.attach(g, Paint, () => {
+    'use gpu';
+    return Paint({ color: rgb });
+  });
+}
+
+export function place<V extends meshes.Placed>(
+  g: meshes.IndexedGeometry<V>,
+  at: d.InferInput<d.Vec3f>,
+  color: d.InferInput<d.Vec3f>,
+) {
+  return paint(meshes.transform(g, std.translation4(d.vec3f(at[0], at[1], at[2]))), color);
+}
diff --git a/apps/typegpu-docs/src/examples/rendering/mesh-primitives/index.html b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/index.html
new file mode 100644
index 0000000000..581d6789f8
--- /dev/null
+++ b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/index.html
@@ -0,0 +1 @@
+<canvas data-fit-to-container></canvas>
diff --git a/apps/typegpu-docs/src/examples/rendering/mesh-primitives/index.ts b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/index.ts
new file mode 100644
index 0000000000..a8dd4f1c4b
--- /dev/null
+++ b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/index.ts
@@ -0,0 +1,146 @@
+import { d, std, tgpu } from 'typegpu';
+import { meshes } from '@typegpu/geometry';
+import { mat4 } from 'wgpu-matrix';
+import { setupOrbitCamera } from '../../common/setup-orbit-camera.ts';
+import { blob } from './blob.ts';
+import { gallery } from './gallery.ts';
+import { place } from './geometry.ts';
+import {
+  fragment,
+  SceneUniforms,
+  sceneLayout,
+  shadowLayout,
+  worldToClip,
+  worldToLight,
+} from './shading.ts';
+import { spring } from './spring.ts';
+
+const root = await tgpu.init();
+const canvas = document.querySelector('canvas') as HTMLCanvasElement;
+const context = root.configureContext({ canvas, alphaMode: 'premultiplied' });
+
+const uniforms = root.createUniform(SceneUniforms);
+const sceneBindGroup = root.createBindGroup(sceneLayout, { uniforms });
+
+const shadowMap = root
+  .createTexture({ size: [2048, 2048], format: 'depth32float' })
+  .$usage('render', 'sampled');
+
+const shadowBindGroup = root.createBindGroup(shadowLayout, {
+  map: shadowMap,
+  sampler: root.createComparisonSampler({
+    compare: 'less-equal',
+    magFilter: 'linear',
+    minFilter: 'linear',
+  }),
+});
+
+function createDepth() {
+  return root
+    .createTexture({ size: [canvas.width, canvas.height], format: 'depth24plus' })
+    .$usage('transient');
+}
+
+let depth = createDepth();
+
+const scene = meshes.bakeIndices(
+  root,
+  meshes.concat(gallery, blob, place(spring, [0, 0.07, 2.2], [0.7, 0.45, 0.95])),
+);
+
+const shadow = root
+  .createRenderPipeline({
+    vertex: ({ $vertexIndex }) => {
+      'use gpu';
+      return { $position: worldToLight(scene.vertexAt($vertexIndex).position) };
+    },
+    depthStencil: {
+      format: 'depth32float',
+      depthWriteEnabled: true,
+      depthCompare: 'less',
+      depthBias: 1,
+      depthBiasSlopeScale: 2,
+    },
+    primitive: { cullMode: 'back' },
+  })
+  .with(sceneBindGroup)
+  .pipe(scene.inject());
+
+const lit = root
+  .createRenderPipeline({
+    vertex: ({ $vertexIndex }) => {
+      'use gpu';
+      const vertex = scene.vertexAt($vertexIndex);
+      return {
+        $position: worldToClip(vertex.position),
+        worldPos: vertex.position,
+        normal: vertex.normal,
+        color: vertex.color,
+      };
+    },
+    fragment,
+    depthStencil: { format: 'depth24plus', depthWriteEnabled: true, depthCompare: 'less' },
+    primitive: { cullMode: 'back' },
+  })
+  .with(sceneBindGroup)
+  .with(shadowBindGroup)
+  .pipe(scene.inject());
+
+const { cleanupCamera } = setupOrbitCamera(
+  canvas,
+  { initPos: d.vec4f(3.2, 2.2, 3.7, 1), target: d.vec4f(0, 0.3, 0, 1), minZoom: 2, maxZoom: 12 },
+  (camera) => uniforms.patch({ camera }),
+);
+
+const lightProjection = mat4.ortho(-4.5, 4.5, -4.5, 4.5, 0.5, 20, d.mat4x4f());
+
+function lightAt(seconds: number) {
+  const angle = seconds * 0.2;
+  const direction = std.normalize(d.vec3f(Math.cos(angle) * 0.7, -1, Math.sin(angle) * 0.7));
+  const eye = std.mul(direction, -10);
+  const view = mat4.lookAt(eye, d.vec3f(), d.vec3f(0, 1, 0), d.mat4x4f());
+
+  return { viewProj: std.mul(lightProjection, view), direction };
+}
+
+let frame: number;
+
+function render(now: number) {
+  const seconds = now * 0.001;
+  uniforms.patch({ light: lightAt(seconds), time: seconds });
+
+  const encoder = root['~unstable'].createCommandEncoder();
+
+  const shadowPass = encoder.beginRenderPass({ depthStencilAttachment: { view: shadowMap } });
+  shadow.with(shadowPass).drawIndexed(scene.indexCount);
+  shadowPass.end();
+
+  const pass = encoder.beginRenderPass({
+    colorAttachments: [{ view: context, clearValue: [0.08, 0.09, 0.12, 1] }],
+    depthStencilAttachment: { view: depth, depthStoreOp: 'discard' },
+  });
+
+  lit.with(pass).drawIndexed(scene.indexCount);
+  pass.end();
+
+  encoder.submit();
+
+  frame = requestAnimationFrame(render);
+}
+
+frame = requestAnimationFrame(render);
+
+const resizeObserver = new ResizeObserver(() => {
+  depth.destroy();
+  depth = createDepth();
+});
+
+resizeObserver.observe(canvas);
+
+export function onCleanup() {
+  cancelAnimationFrame(frame);
+  resizeObserver.disconnect();
+  cleanupCamera();
+
+  root.destroy();
+}
diff --git a/apps/typegpu-docs/src/examples/rendering/mesh-primitives/meta.json b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/meta.json
new file mode 100644
index 0000000000..9dd070d0f3
--- /dev/null
+++ b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/meta.json
@@ -0,0 +1,7 @@
+{
+  "title": "Mesh Primitives",
+  "description": "Combines procedural primitives, a rippling parametric blob and a spring into one mesh, evaluated in the vertex shader and lit by a circling shadow-casting light.",
+  "category": "rendering",
+  "tags": ["3d", "geometry", "rasterization"],
+  "coolFactor": 6
+}
diff --git a/apps/typegpu-docs/src/examples/rendering/mesh-primitives/shading.ts b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/shading.ts
new file mode 100644
index 0000000000..455f67b96a
--- /dev/null
+++ b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/shading.ts
@@ -0,0 +1,43 @@
+import { d, std, tgpu } from 'typegpu';
+
+const Camera = d.struct({ viewProjection: d.mat4x4f });
+const Light = d.struct({ viewProj: d.mat4x4f, direction: d.vec3f });
+
+export const SceneUniforms = d.struct({ camera: Camera, light: Light, time: d.f32 });
+export const sceneLayout = tgpu.bindGroupLayout({ uniforms: { uniform: SceneUniforms } });
+
+export const shadowLayout = tgpu.bindGroupLayout({
+  map: { texture: d.textureDepth2d() },
+  sampler: { sampler: 'comparison' },
+});
+
+export function worldToClip(position: d.v3f) {
+  'use gpu';
+  return sceneLayout.$.uniforms.camera.viewProjection * d.vec4f(position, 1);
+}
+
+export function worldToLight(position: d.v3f) {
+  'use gpu';
+  return sceneLayout.$.uniforms.light.viewProj * d.vec4f(position, 1);
+}
+
+function shadowAt(position: d.v3f, normal: d.v3f) {
+  'use gpu';
+  const clip = worldToLight(position + normal * 0.03);
+  const uv = clip.xy * d.vec2f(0.5, -0.5) + 0.5;
+
+  return std.textureSampleCompare(shadowLayout.$.map, shadowLayout.$.sampler, uv, clip.z);
+}
+
+export const fragment = tgpu.fragmentFn({
+  in: { worldPos: d.vec3f, normal: d.vec3f, color: d.vec3f },
+  out: d.vec4f,
+})(({ worldPos, normal, color }) => {
+  'use gpu';
+  const n = std.normalize(normal);
+  const toLight = std.neg(sceneLayout.$.uniforms.light.direction);
+  const diffuse = std.max(std.dot(n, toLight), 0) * shadowAt(worldPos, n);
+  const sky = n.y * 0.5 + 0.5;
+
+  return d.vec4f(color * (0.2 + 0.25 * sky + 0.6 * diffuse), 1);
+});
diff --git a/apps/typegpu-docs/src/examples/rendering/mesh-primitives/spring.ts b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/spring.ts
new file mode 100644
index 0000000000..6379750aaa
--- /dev/null
+++ b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/spring.ts
@@ -0,0 +1,92 @@
+import { d, std } from 'typegpu';
+import { meshes } from '@typegpu/geometry';
+import { sceneLayout } from './shading.ts';
+
+const TAU = Math.PI * 2;
+const turns = 4;
+const radius = 0.3;
+const tube = 0.07;
+const height = 1;
+
+const centerAt = (u: number) => {
+  'use gpu';
+  const phi = u * TAU * turns;
+  return d.vec3f(radius * std.cos(phi), (u - 0.5) * height, radius * std.sin(phi));
+};
+
+const radialAt = (u: number) => {
+  'use gpu';
+  const phi = u * TAU * turns;
+  return d.vec3f(std.cos(phi), 0, std.sin(phi));
+};
+
+const tangentAt = (u: number) => {
+  'use gpu';
+  const phi = u * TAU * turns;
+  return std.normalize(
+    d.vec3f(-radius * std.sin(phi) * TAU * turns, height, radius * std.cos(phi) * TAU * turns),
+  );
+};
+
+const offset = (u: number, v: number) => {
+  'use gpu';
+  const t = v * TAU;
+  const radial = radialAt(u);
+  const binormal = std.cross(radial, tangentAt(u));
+  return radial * std.cos(t) + binormal * std.sin(t);
+};
+
+const body = meshes.parametric(
+  {
+    at: (u, v) => {
+      'use gpu';
+      return centerAt(u) + offset(u, v) * tube;
+    },
+    normalAt: offset,
+  },
+  { cols: 48 * turns, rows: 12 },
+);
+
+const cap = meshes.parametric(
+  {
+    at: (u, v) => {
+      'use gpu';
+      return meshes.surfaces.disc(u, v) * tube;
+    },
+    normalAt: () => {
+      'use gpu';
+      return d.vec3f(0, 1, 0);
+    },
+  },
+  { cols: 12, rows: 1 },
+);
+
+const capAt = (u: number, sign: number) => {
+  const radial = radialAt(u);
+  const tangent = tangentAt(u).mul(sign);
+  const binormal = std.cross(radial, tangent);
+
+  const frame = d.mat4x4f(
+    d.vec4f(radial, 0),
+    d.vec4f(tangent, 0),
+    d.vec4f(binormal, 0),
+    d.vec4f(centerAt(u), 1),
+  );
+
+  return meshes.transform(cap, frame);
+};
+
+export const spring = meshes.map(
+  meshes.concat(body, capAt(0, -1), capAt(1, 1)),
+  meshes.Surface,
+  (vertex) => {
+    'use gpu';
+    const stretch = 1 + 0.3 * std.sin(sceneLayout.$.uniforms.time * 2);
+
+    return meshes.Surface({
+      position: vertex.position * d.vec3f(1, stretch, 1),
+      normal: std.normalize(vertex.normal / d.vec3f(1, stretch, 1)),
+      uv: vertex.uv,
+    });
+  },
+);
diff --git a/apps/typegpu-docs/src/examples/rendering/mesh-primitives/thumbnail.png b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/thumbnail.png
new file mode 100644
index 0000000000..9988292c05
Binary files /dev/null and b/apps/typegpu-docs/src/examples/rendering/mesh-primitives/thumbnail.png differ
diff --git a/apps/typegpu-docs/tests/individual-example-tests/mesh-primitives.test.ts b/apps/typegpu-docs/tests/individual-example-tests/mesh-primitives.test.ts
new file mode 100644
index 0000000000..4f1c29ec41
--- /dev/null
+++ b/apps/typegpu-docs/tests/individual-example-tests/mesh-primitives.test.ts
@@ -0,0 +1,46 @@
+/**
+ * @vitest-environment jsdom
+ */
+
+import { describe, expect } from 'vitest';
+import { it } from 'typegpu-testing-utility';
+import { runExampleTest, setupCommonMocks } from './utils/baseTest.ts';
+import { mockResizeObserver } from './utils/commonMocks.ts';
+
+describe('mesh primitives example', () => {
+  setupCommonMocks();
+
+  it('resolves the procedural scene and its shadows', async ({ device }) => {
+    const shaderCodes = await runExampleTest(
+      {
+        category: 'rendering',
+        name: 'mesh-primitives',
+        setupMocks: mockResizeObserver,
+        expectedCalls: 3,
+      },
+      device,
+    );
+
+    const uniforms = shaderCodes.match(
+      /struct (?:Camera|Light|SceneUniforms) \{[^}]*\}|@group[^\n]*var<uniform>[^\n]*/g,
+    );
+    expect([...new Set(uniforms)].join('\n\n')).toMatchInlineSnapshot(`
+      "struct Camera {
+        viewProjection: mat4x4f,
+      }
+
+      struct Light {
+        viewProj: mat4x4f,
+        direction: vec3f,
+      }
+
+      struct SceneUniforms {
+        camera: Camera,
+        light: Light,
+        time: f32,
+      }
+
+      @group(0) @binding(0) var<uniform> uniforms: SceneUniforms;"
+    `);
+  });
+});
diff --git a/packages/typegpu-geometry/build.config.ts b/packages/typegpu-geometry/build.config.ts
deleted file mode 100644
index 7f9f024f1f..0000000000
--- a/packages/typegpu-geometry/build.config.ts
+++ /dev/null
@@ -1,12 +0,0 @@
-import { type BuildConfig, defineBuildConfig } from 'unbuild';
-import typegpu from 'unplugin-typegpu/rollup';
-
-const Config: BuildConfig[] = defineBuildConfig({
-  hooks: {
-    'rollup:options': (_options, config) => {
-      config.plugins.push(typegpu({ include: [/\.ts$/] }));
-    },
-  },
-});
-
-export default Config;
diff --git a/packages/typegpu-geometry/package.json b/packages/typegpu-geometry/package.json
index 46d8c955a5..4c81b95ca9 100644
--- a/packages/typegpu-geometry/package.json
+++ b/packages/typegpu-geometry/package.json
@@ -23,27 +23,30 @@
       "./package.json": "./dist/package.json",
       ".": {
         "types": "./dist/index.d.ts",
-        "module": "./dist/index.mjs",
-        "import": "./dist/index.mjs",
-        "default": "./dist/index.cjs"
+        "module": "./dist/index.js",
+        "import": "./dist/index.js"
       }
     },
     "linkDirectory": false,
-    "main": "./dist/index.mjs",
+    "main": "./dist/index.js",
     "types": "./dist/index.d.ts"
   },
   "scripts": {
-    "build": "unbuild",
-    "test:types": "pnpm tsc --p ./tsconfig.json --noEmit",
+    "build": "tsdown",
+    "test": "vitest run",
+    "test:types": "pnpm tsc --p ./tsconfig.json --noEmit && pnpm tsc --p ./tsconfig.test.json --noEmit",
     "prepublishOnly": "tgpu-dev-cli prepack"
   },
   "devDependencies": {
     "@typegpu/tgpu-dev-cli": "workspace:*",
     "@webgpu/types": "catalog:types",
+    "jiti": "catalog:build",
+    "tsdown": "catalog:build",
     "typegpu": "workspace:*",
+    "typegpu-testing-utility": "workspace:*",
     "typescript": "catalog:types",
-    "unbuild": "catalog:build",
-    "unplugin-typegpu": "workspace:*"
+    "unplugin-typegpu": "workspace:*",
+    "vitest": "catalog:test"
   },
   "peerDependencies": {
     "typegpu": "workspace:^"
diff --git a/packages/typegpu-geometry/src/index.ts b/packages/typegpu-geometry/src/index.ts
index 33f48fa5c3..c34c8b211b 100644
--- a/packages/typegpu-geometry/src/index.ts
+++ b/packages/typegpu-geometry/src/index.ts
@@ -1,2 +1,3 @@
 export * from './circle.ts';
 export * from './lines/index.ts';
+export * as meshes from './mesh/index.ts';
diff --git a/packages/typegpu-geometry/src/lines/joins/bevel.ts b/packages/typegpu-geometry/src/lines/joins/bevel.ts
new file mode 100644
index 0000000000..11ef78140e
--- /dev/null
+++ b/packages/typegpu-geometry/src/lines/joins/bevel.ts
@@ -0,0 +1,16 @@
+import { f32, vec2f } from 'typegpu/data';
+import { mix } from 'typegpu/std';
+import type { JoinInput } from '../types.ts';
+
+/**
+ * Cuts the outside of a turn with a straight edge between the sides of the two segments.
+ */
+export function bevel(join: JoinInput, joinVertexIndex: number, maxJoinCount: number) {
+  'use gpu';
+  if (joinVertexIndex === 0) {
+    return vec2f(join.v);
+  }
+  const t = f32(joinVertexIndex) / f32(maxJoinCount);
+  const dir = mix(join.d, (join.start + join.end) * 0.5, t);
+  return join.C.position + dir * join.C.radius;
+}
diff --git a/packages/typegpu-geometry/src/lines/joins/index.ts b/packages/typegpu-geometry/src/lines/joins/index.ts
index 0254331b8c..c1e6e78607 100644
--- a/packages/typegpu-geometry/src/lines/joins/index.ts
+++ b/packages/typegpu-geometry/src/lines/joins/index.ts
@@ -1,4 +1,7 @@
+import { bevel } from './bevel.ts';
 import { miter } from './miter.ts';
 import { round } from './round.ts';
 
-export const joins = { miter, round };
+export { miterJoinLimitSlot } from './miter.ts';
+
+export const joins = { bevel, miter, round };
diff --git a/packages/typegpu-geometry/src/mesh/bake.ts b/packages/typegpu-geometry/src/mesh/bake.ts
new file mode 100644
index 0000000000..51a78b6766
--- /dev/null
+++ b/packages/typegpu-geometry/src/mesh/bake.ts
@@ -0,0 +1,280 @@
+import {
+  d,
+  std,
+  type IndexFlag,
+  type StorageFlag,
+  type TgpuBuffer,
+  type TgpuRenderPipeline,
+  type TgpuRoot,
+  type TgpuVertexLayout,
+  type VertexFlag,
+  type WithBinding,
+} from 'typegpu';
+import { createFill, type Fill, recordFills, type UpdateOptions } from './fill.ts';
+import {
+  type ConstructibleSchema,
+  type Geometry,
+  type IndexedGeometry,
+  isIndexed,
+  layoutOf,
+} from './geometry.ts';
+
+export type { RecordOptions, UpdateOptions } from './fill.ts';
+
+type HasIndexBuffer = Pick<
+  ReturnType<TgpuRenderPipeline['withIndexBuffer']>,
+  'hasIndexBuffer' | 'drawIndexed'
+>;
+
+export type VertexBuffer<V extends ConstructibleSchema> = TgpuBuffer<d.WgslArray<V>> &
+  VertexFlag &
+  StorageFlag;
+export type IndexBuffer = TgpuBuffer<d.WgslArray<d.U32>> & IndexFlag & StorageFlag;
+
+export interface Baked<V extends ConstructibleSchema> extends Geometry<V> {
+  readonly vertices: VertexBuffer<V>;
+  /**
+   * Index of the mesh's first vertex in `vertices`. Pass it as `baseVertex` to `drawIndexed` or as
+   * `firstVertex` to `draw`. `vertexAt` already adds it.
+   */
+  readonly baseVertex: number;
+  readonly layout: TgpuVertexLayout<d.WgslArray<V>>;
+  inject(): <P extends TgpuRenderPipeline>(pipeline: P) => P;
+  updateVertices(this: void, options?: UpdateOptions): void;
+  destroy(): void;
+}
+
+export interface BakedIndexed<V extends ConstructibleSchema> extends Baked<V>, IndexedGeometry<V> {
+  readonly indices: IndexBuffer;
+  /** Index of the mesh's first index in `indices`. Pass it as `firstIndex` when drawing */
+  readonly firstIndex: number;
+  inject(): <P extends TgpuRenderPipeline>(pipeline: P) => P & HasIndexBuffer;
+}
+
+export type BakeOptions<V extends ConstructibleSchema> = UpdateOptions & {
+  /** Writes into this buffer instead of creating one. `destroy()` leaves it alive */
+  vertices?: VertexBuffer<V>;
+  /** Writes into this buffer instead of creating one. `destroy()` leaves it alive */
+  indices?: IndexBuffer;
+  /** Writes the first vertex at this index of `vertices`. Indices stay relative to it. Default 0 */
+  baseVertex?: number;
+  /** Writes the first index at this index of `indices`. Default 0 */
+  firstIndex?: number;
+  /** Creates the fill pipelines from this configured root, for slots and accessors */
+  with?: WithBinding;
+};
+
+interface Prepared<T> {
+  mesh: T;
+  fills: Fill[];
+}
+
+function fits(kind: string, capacity: number, first: number, count: number) {
+  if (first + count > capacity) {
+    throw new Error(
+      `The ${kind} range ${first}..${first + count} exceeds the buffer capacity ${capacity}`,
+    );
+  }
+}
+
+function finish<T>(root: TgpuRoot, { mesh, fills }: Prepared<T>, options: UpdateOptions): T {
+  recordFills(root, fills, options);
+  return mesh;
+}
+
+async function finishAsync<T>(
+  root: TgpuRoot,
+  prepared: Prepared<T>,
+  options: UpdateOptions,
+): Promise<T> {
+  await Promise.all(prepared.fills.map((fill) => fill.initAsync()));
+  return finish(root, prepared, options);
+}
+
+export function bake<V extends ConstructibleSchema>(
+  root: TgpuRoot,
+  g: IndexedGeometry<V>,
+  options?: BakeOptions<V>,
+): BakedIndexed<V>;
+export function bake<V extends ConstructibleSchema>(
+  root: TgpuRoot,
+  g: Geometry<V>,
+  options?: BakeOptions<V>,
+): Baked<V>;
+export function bake<V extends ConstructibleSchema>(
+  root: TgpuRoot,
+  g: Geometry<V>,
+  options: BakeOptions<V> = {},
+): Baked<V> {
+  return finish(root, prepareBake(root, g, options), options);
+}
+
+export function bakeAsync<V extends ConstructibleSchema>(
+  root: TgpuRoot,
+  g: IndexedGeometry<V>,
+  options?: BakeOptions<V>,
+): Promise<BakedIndexed<V>>;
+export function bakeAsync<V extends ConstructibleSchema>(
+  root: TgpuRoot,
+  g: Geometry<V>,
+  options?: BakeOptions<V>,
+): Promise<Baked<V>>;
+export async function bakeAsync<V extends ConstructibleSchema>(
+  root: TgpuRoot,
+  g: Geometry<V>,
+  options: BakeOptions<V> = {},
+): Promise<Baked<V>> {
+  return finishAsync(root, prepareBake(root, g, options), options);
+}
+
+function prepareBake<V extends ConstructibleSchema>(
+  root: TgpuRoot,
+  g: Geometry<V>,
+  options: BakeOptions<V>,
+): Prepared<Baked<V>> {
+  const { schema, topology, vertexCount } = g;
+  const { baseVertex = 0 } = options;
+  const indexed = isIndexed(g) ? g : undefined;
+
+  const vertices =
+    options.vertices ??
+    (root
+      .createBuffer(d.arrayOf(schema, Math.max(1, vertexCount)) as d.WgslArray<d.AnyWgslData>)
+      .$usage('vertex', 'storage') as unknown as VertexBuffer<V>);
+  fits('vertex', vertices.dataType.elementCount, baseVertex, vertexCount);
+
+  const output = vertices.as('mutable');
+  const fillVertices = createFill(
+    options.with ?? root,
+    vertexCount,
+    (i) => {
+      'use gpu';
+      output.$[baseVertex === 0 ? i : baseVertex + i] = g.vertexAt(i);
+    },
+    options.bindGroups,
+  );
+
+  const readVertices = vertices.as('readonly');
+  const layout = layoutOf(schema);
+  const mesh = {
+    schema,
+    topology,
+    vertexCount,
+    vertexAt: (i: number) => {
+      'use gpu';
+      return std.copy(readVertices.$[baseVertex === 0 ? i : baseVertex + i]) as d.InferGPU<V>;
+    },
+    vertices,
+    baseVertex,
+    layout,
+    updateVertices: (update: UpdateOptions = {}) => recordFills(root, [fillVertices], update),
+  };
+
+  if (!indexed) {
+    return {
+      mesh: {
+        ...mesh,
+        inject: () => (pipeline) => pipeline.with(layout, vertices),
+        destroy: () => {
+          if (!options.vertices) {
+            vertices.destroy();
+          }
+        },
+      },
+      fills: [fillVertices],
+    };
+  }
+
+  const { mesh: indexOnly, fills } = prepareIndices(root, indexed, options);
+  const indexedMesh: BakedIndexed<V> = {
+    ...mesh,
+    indices: indexOnly.indices,
+    firstIndex: indexOnly.firstIndex,
+    indexCount: indexOnly.indexCount,
+    indexAt: indexOnly.indexAt,
+    inject: () => (pipeline) => pipeline.with(layout, vertices).withIndexBuffer(indexOnly.indices),
+    destroy: () => {
+      if (!options.vertices) {
+        vertices.destroy();
+      }
+      indexOnly.destroy();
+    },
+  };
+  return { mesh: indexedMesh, fills: [fillVertices, ...fills] };
+}
+
+export type BakeIndicesOptions = UpdateOptions & {
+  /** Writes into this buffer instead of creating one. `destroy()` leaves it alive */
+  indices?: IndexBuffer;
+  /** Writes the first index at this index of `indices`. Default 0 */
+  firstIndex?: number;
+  /** Creates the fill pipeline from this configured root, for slots and accessors */
+  with?: WithBinding;
+};
+
+export interface BakedIndices<V extends ConstructibleSchema> extends IndexedGeometry<V> {
+  readonly indices: IndexBuffer;
+  /** Index of the mesh's first index in `indices`. Pass it as `firstIndex` when drawing */
+  readonly firstIndex: number;
+  inject(): <P extends TgpuRenderPipeline>(pipeline: P) => P & HasIndexBuffer;
+  destroy(): void;
+}
+
+export function bakeIndices<V extends ConstructibleSchema>(
+  root: TgpuRoot,
+  g: IndexedGeometry<V>,
+  options: BakeIndicesOptions = {},
+): BakedIndices<V> {
+  return finish(root, prepareIndices(root, g, options), options);
+}
+
+export async function bakeIndicesAsync<V extends ConstructibleSchema>(
+  root: TgpuRoot,
+  g: IndexedGeometry<V>,
+  options: BakeIndicesOptions = {},
+): Promise<BakedIndices<V>> {
+  return finishAsync(root, prepareIndices(root, g, options), options);
+}
+
+function prepareIndices<V extends ConstructibleSchema>(
+  root: TgpuRoot,
+  g: IndexedGeometry<V>,
+  options: BakeIndicesOptions,
+): Prepared<BakedIndices<V>> {
+  const { firstIndex = 0 } = options;
+  const indices =
+    options.indices ??
+    root.createBuffer(d.arrayOf(d.u32, Math.max(1, g.indexCount))).$usage('index', 'storage');
+  fits('index', indices.dataType.elementCount, firstIndex, g.indexCount);
+
+  const output = indices.as('mutable');
+  const fill = createFill(
+    options.with ?? root,
+    g.indexCount,
+    (i) => {
+      'use gpu';
+      output.$[firstIndex === 0 ? i : firstIndex + i] = g.indexAt(i);
+    },
+    options.bindGroups,
+  );
+
+  const readIndices = indices.as('readonly');
+  return {
+    mesh: {
+      ...g,
+      indices,
+      firstIndex,
+      indexAt: (i: number) => {
+        'use gpu';
+        return readIndices.$[firstIndex === 0 ? i : firstIndex + i] as number;
+      },
+      inject: () => (pipeline) => pipeline.withIndexBuffer(indices),
+      destroy: () => {
+        if (!options.indices) {
+          indices.destroy();
+        }
+      },
+    },
+    fills: [fill],
+  };
+}
diff --git a/packages/typegpu-geometry/src/mesh/combinators.ts b/packages/typegpu-geometry/src/mesh/combinators.ts
new file mode 100644
index 0000000000..7a916ee8c8
--- /dev/null
+++ b/packages/typegpu-geometry/src/mesh/combinators.ts
@@ -0,0 +1,263 @@
+import { d, std, tgpu } from 'typegpu';
+import {
+  type ConstructibleSchema,
+  type Geometry,
+  type IndexedGeometry,
+  isIndexed,
+} from './geometry.ts';
+
+function reversedCorner(i: number) {
+  'use gpu';
+  const corner = i % 3;
+
+  if (corner === 1) {
+    return i + 1;
+  }
+  if (corner === 2) {
+    return i - 1;
+  }
+
+  return i;
+}
+
+export type Placed = d.WgslStruct<{ position: d.Vec3f; normal: d.Vec3f }>;
+
+export function transform<V extends Placed>(g: IndexedGeometry<V>, m: d.m4x4f): IndexedGeometry<V>;
+export function transform<V extends Placed>(g: Geometry<V>, m: d.m4x4f): Geometry<V>;
+export function transform<V extends Placed>(g: Geometry<V>, m: d.m4x4f): Geometry<V> {
+  const matrix = std.copy(m);
+  const c0 = matrix.columns[0].xyz;
+  const c1 = matrix.columns[1].xyz;
+  const c2 = matrix.columns[2].xyz;
+  const determinant = std.dot(c0, std.cross(c1, c2));
+  // The cofactor matrix is the inverse transpose scaled by the determinant. Unlike the inverse, it
+  // exists for transforms that flatten a shape, and its sign keeps reflected normals outwards.
+  const sign = determinant < 0 ? -1 : 1;
+  const normalM = d.mat3x3f(
+    std.cross(c1, c2).mul(sign),
+    std.cross(c2, c0).mul(sign),
+    std.cross(c0, c1).mul(sign),
+  );
+  const reflected = determinant < 0 && g.topology === 'triangle-list';
+  const reverseVertices = reflected && !isIndexed(g);
+
+  const result = {
+    ...g,
+    vertexAt: (i: number) => {
+      'use gpu';
+      let s = std.copy(g.vertexAt(reverseVertices ? reversedCorner(i) : i));
+      s.position = (matrix * d.vec4f(s.position, 1)).xyz;
+
+      const normal = normalM * s.normal;
+      if (std.dot(normal, normal) > 0) {
+        s.normal = std.normalize(normal);
+      } else {
+        s.normal = d.vec3f();
+      }
+
+      return s;
+    },
+  };
+
+  if (reflected && isIndexed(g)) {
+    const indexed: IndexedGeometry<V> = {
+      ...result,
+      indexCount: g.indexCount,
+      indexAt: (i: number) => {
+        'use gpu';
+        return g.indexAt(reversedCorner(i));
+      },
+    };
+    return indexed;
+  }
+
+  return result;
+}
+
+export function map<VIn extends ConstructibleSchema, V extends ConstructibleSchema>(
+  g: IndexedGeometry<VIn>,
+  schema: V,
+  fn: (vertex: d.InferGPU<VIn>) => d.InferGPU<V>,
+): IndexedGeometry<V>;
+export function map<VIn extends ConstructibleSchema, V extends ConstructibleSchema>(
+  g: Geometry<VIn>,
+  schema: V,
+  fn: (vertex: d.InferGPU<VIn>) => d.InferGPU<V>,
+): Geometry<V>;
+export function map<VIn extends ConstructibleSchema, V extends ConstructibleSchema>(
+  g: Geometry<VIn>,
+  schema: V,
+  fn: (vertex: d.InferGPU<VIn>) => d.InferGPU<V>,
+): Geometry<V> {
+  return {
+    ...g,
+    schema,
+    vertexAt: (i: number) => {
+      'use gpu';
+      return fn(g.vertexAt(i));
+    },
+  };
+}
+
+export type Attached<VIn extends d.WgslStruct, E extends d.WgslStruct> = d.WgslStruct<
+  VIn['propTypes'] & E['propTypes']
+>;
+
+const attachedSchemas = new WeakMap<d.WgslStruct, WeakMap<d.WgslStruct, d.WgslStruct>>();
+
+function attachedSchema<VIn extends d.WgslStruct, E extends d.WgslStruct>(
+  base: VIn,
+  extra: E,
+): Attached<VIn, E> {
+  for (const key of Object.keys(extra.propTypes)) {
+    if (Object.hasOwn(base.propTypes, key)) {
+      throw new Error(`attach cannot replace the existing field '${key}'; use map instead`);
+    }
+  }
+
+  let byExtra = attachedSchemas.get(base);
+  if (!byExtra) {
+    byExtra = new WeakMap();
+    attachedSchemas.set(base, byExtra);
+  }
+
+  let schema = byExtra.get(extra);
+  if (!schema) {
+    const Vertex = d.struct({ ...base.propTypes, ...extra.propTypes } as Record<
+      string,
+      d.AnyWgslData
+    >);
+    byExtra.set(extra, Vertex);
+    schema = Vertex;
+  }
+
+  return schema as Attached<VIn, E>;
+}
+
+export function attach<VIn extends d.WgslStruct, E extends d.WgslStruct>(
+  g: IndexedGeometry<VIn>,
+  extra: E,
+  fn: (vertex: d.InferGPU<VIn>) => d.InferGPU<E>,
+): IndexedGeometry<Attached<VIn, E>>;
+export function attach<VIn extends d.WgslStruct, E extends d.WgslStruct>(
+  g: Geometry<VIn>,
+  extra: E,
+  fn: (vertex: d.InferGPU<VIn>) => d.InferGPU<E>,
+): Geometry<Attached<VIn, E>>;
+export function attach<VIn extends d.WgslStruct, E extends d.WgslStruct>(
+  g: Geometry<VIn>,
+  extra: E,
+  fn: (vertex: d.InferGPU<VIn>) => d.InferGPU<E>,
+): Geometry<Attached<VIn, E>> {
+  const schema = attachedSchema(g.schema, extra);
+  const baseKeys = Object.keys(g.schema.propTypes);
+  const extraKeys = Object.keys(extra.propTypes);
+
+  return {
+    ...g,
+    schema,
+    vertexAt: (i: number) => {
+      'use gpu';
+      const s = g.vertexAt(i);
+      const e = fn(s);
+
+      let out: Record<string, unknown> = schema();
+      for (const key of tgpu.unroll(baseKeys)) {
+        out[key] = std.copy(s[key]);
+      }
+      for (const key of tgpu.unroll(extraKeys)) {
+        out[key] = std.copy(e[key]);
+      }
+
+      return out as d.InferGPU<Attached<VIn, E>>;
+    },
+  };
+}
+
+export function unindexed<V extends ConstructibleSchema>(g: Geometry<V>): Geometry<V> {
+  if (!isIndexed(g)) {
+    return g;
+  }
+
+  return {
+    schema: g.schema,
+    topology: g.topology,
+    vertexCount: g.indexCount,
+    vertexAt: (i: number) => {
+      'use gpu';
+      return g.vertexAt(g.indexAt(i));
+    },
+  };
+}
+
+export function concat<V extends ConstructibleSchema>(...parts: Geometry<V>[]): IndexedGeometry<V> {
+  const first = parts[0];
+  if (!first) {
+    throw new Error('concat needs at least one geometry');
+  }
+
+  const schema = first.schema;
+  const topology = first.topology;
+
+  for (const g of parts) {
+    if (g.schema !== schema) {
+      throw new Error('concat needs one vertex schema across all of its parts');
+    }
+    if (g.topology !== topology) {
+      throw new Error(
+        `concat needs one topology across all of its parts, got '${topology}' and '${g.topology}'`,
+      );
+    }
+  }
+
+  const vertexCount = parts.reduce((n, g) => n + g.vertexCount, 0);
+
+  const vertexAt = (i: number) => {
+    'use gpu';
+    let local = i;
+    for (const g of tgpu.unroll(parts)) {
+      if (local < g.vertexCount) {
+        return g.vertexAt(local);
+      }
+      local -= g.vertexCount;
+    }
+
+    return schema() as d.InferGPU<V>;
+  };
+
+  let offset = 0;
+  const entries = parts.map((g) => {
+    const indexCount = isIndexed(g) ? g.indexCount : g.vertexCount;
+    const indexAt = isIndexed(g)
+      ? g.indexAt
+      : (i: number) => {
+          'use gpu';
+          return i;
+        };
+    const entry = { indexAt, indexCount, vertexOffset: offset };
+    offset += g.vertexCount;
+    return entry;
+  });
+
+  const result: IndexedGeometry<V> = {
+    schema,
+    topology,
+    vertexCount,
+    vertexAt,
+    indexCount: entries.reduce((n, e) => n + e.indexCount, 0),
+    indexAt: (i: number) => {
+      'use gpu';
+      let local = i;
+      for (const e of tgpu.unroll(entries)) {
+        if (local < e.indexCount) {
+          return e.indexAt(local) + e.vertexOffset;
+        }
+        local -= e.indexCount;
+      }
+
+      return d.u32(0);
+    },
+  };
+
+  return result;
+}
diff --git a/packages/typegpu-geometry/src/mesh/fill.ts b/packages/typegpu-geometry/src/mesh/fill.ts
new file mode 100644
index 0000000000..0580e2e948
--- /dev/null
+++ b/packages/typegpu-geometry/src/mesh/fill.ts
@@ -0,0 +1,106 @@
+import {
+  d,
+  tgpu,
+  type TgpuBindGroup,
+  type TgpuCommandEncoder,
+  type TgpuComputePass,
+  type TgpuRoot,
+  type WithBinding,
+} from 'typegpu';
+
+interface EncoderOptions {
+  /** Records the work as one compute pass on this encoder. Nothing is submitted */
+  encoder: GPUCommandEncoder | TgpuCommandEncoder;
+  pass?: never;
+}
+
+interface PassOptions {
+  encoder?: never;
+  /** Records the work into this pass. Nothing is submitted and the pass is not ended */
+  pass: GPUComputePassEncoder | TgpuComputePass;
+}
+
+interface SubmitOptions {
+  encoder?: never;
+  pass?: never;
+}
+
+/** Controls where the work is recorded. Defaults to a standalone submit */
+export type RecordOptions = EncoderOptions | PassOptions | SubmitOptions;
+
+export type UpdateOptions = RecordOptions & {
+  /** Bind groups for this call, applied over the ones given when baking */
+  bindGroups?: readonly TgpuBindGroup[];
+};
+
+export interface Fill {
+  initAsync(): Promise<void>;
+  record(pass: TgpuComputePass, bindGroups?: readonly TgpuBindGroup[]): void;
+}
+
+export function createFill(
+  root: WithBinding,
+  count: number,
+  write: (index: number) => void,
+  bindGroups: readonly TgpuBindGroup[] = [],
+): Fill {
+  const x = Math.min(Math.ceil(count / 64), 65535);
+  const y = x === 0 ? 0 : Math.ceil(count / (x * 64));
+
+  const compute = tgpu.computeFn({
+    workgroupSize: [64],
+    in: { gid: d.builtin.globalInvocationId },
+  })(({ gid }) => {
+    'use gpu';
+    const i = gid.x + gid.y * (x * 64);
+
+    if (i < count) {
+      write(i);
+    }
+  });
+
+  let pipeline = root.createComputePipeline({ compute });
+  for (const group of bindGroups) {
+    pipeline = pipeline.with(group);
+  }
+
+  return {
+    initAsync: async () => {
+      if (count > 0) {
+        await pipeline.initAsync();
+      }
+    },
+    record: (pass, groups = []) => {
+      if (count === 0) {
+        return;
+      }
+
+      let target = pipeline;
+      for (const group of groups) {
+        target = target.with(group);
+      }
+      target.with(pass).dispatchWorkgroups(x, y);
+    },
+  };
+}
+
+export function recordFills(root: TgpuRoot, fills: readonly Fill[], options: UpdateOptions) {
+  const record = (pass: TgpuComputePass) => {
+    for (const fill of fills) {
+      fill.record(pass, options.bindGroups);
+    }
+  };
+
+  if (options.pass) {
+    record(options.pass as TgpuComputePass);
+    return;
+  }
+
+  const encoder = options.encoder ?? root['~unstable'].createCommandEncoder();
+  const pass = encoder.beginComputePass() as TgpuComputePass;
+  record(pass);
+  pass.end();
+  if (!options.encoder) {
+    (encoder as TgpuCommandEncoder).submit();
+  }
+}
diff --git a/packages/typegpu-geometry/src/mesh/geometry.ts b/packages/typegpu-geometry/src/mesh/geometry.ts
new file mode 100644
index 0000000000..52c9ffb458
--- /dev/null
+++ b/packages/typegpu-geometry/src/mesh/geometry.ts
@@ -0,0 +1,43 @@
+import { d, tgpu, type TgpuVertexLayout } from 'typegpu';
+
+/** Default vertex schema for primitives */
+export const Surface = d.struct({ position: d.vec3f, normal: d.vec3f, uv: d.vec2f });
+
+/** Vertex schemas a 'use gpu' function can return */
+export type ConstructibleSchema = Extract<d.AnyWgslData, () => unknown>;
+
+export type Topology = 'triangle-list' | 'line-list' | 'point-list';
+
+export interface Geometry<V extends ConstructibleSchema = typeof Surface> {
+  readonly schema: V;
+  readonly topology: Topology;
+  readonly vertexCount: number;
+
+  vertexAt(this: void, i: number): d.InferGPU<V>;
+}
+
+export interface IndexedGeometry<
+  V extends ConstructibleSchema = typeof Surface,
+> extends Geometry<V> {
+  readonly indexCount: number;
+
+  indexAt(this: void, i: number): number;
+}
+
+const layouts = new WeakMap<d.AnyWgslData, TgpuVertexLayout>();
+
+export function isIndexed<V extends ConstructibleSchema>(g: Geometry<V>): g is IndexedGeometry<V> {
+  return typeof (g as IndexedGeometry<V>).indexAt === 'function';
+}
+
+export function layoutOf<V extends ConstructibleSchema>(
+  schema: V,
+): TgpuVertexLayout<d.WgslArray<V>> {
+  let layout = layouts.get(schema);
+  if (!layout) {
+    layout = tgpu.vertexLayout(d.arrayOf(schema));
+    layouts.set(schema, layout);
+  }
+
+  return layout as TgpuVertexLayout<d.WgslArray<V>>;
+}
diff --git a/packages/typegpu-geometry/src/mesh/index.ts b/packages/typegpu-geometry/src/mesh/index.ts
new file mode 100644
index 0000000000..a43e3bb078
--- /dev/null
+++ b/packages/typegpu-geometry/src/mesh/index.ts
@@ -0,0 +1,8 @@
+export * from './bake.ts';
+export { attach, concat, map, transform, unindexed } from './combinators.ts';
+export type { Attached, Placed } from './combinators.ts';
+export { isIndexed, Surface } from './geometry.ts';
+export type { ConstructibleSchema, Geometry, IndexedGeometry, Topology } from './geometry.ts';
+export { parametric } from './parametric.ts';
+export type { Parametric, ParametricOptions } from './parametric.ts';
+export * from './primitives.ts';
diff --git a/packages/typegpu-geometry/src/mesh/math.ts b/packages/typegpu-geometry/src/mesh/math.ts
new file mode 100644
index 0000000000..a23e01c3cc
--- /dev/null
+++ b/packages/typegpu-geometry/src/mesh/math.ts
@@ -0,0 +1,19 @@
+import { d } from 'typegpu';
+
+export const TAU = Math.PI * 2;
+
+export const X = d.vec3f(1, 0, 0);
+export const Y = d.vec3f(0, 1, 0);
+export const Z = d.vec3f(0, 0, 1);
+export const NX = d.vec3f(-1, 0, 0);
+export const NY = d.vec3f(0, -1, 0);
+export const NZ = d.vec3f(0, 0, -1);
+
+export function frame(tangent: d.v3f, normal: d.v3f, bitangent: d.v3f, center: d.v3f) {
+  return d.mat4x4f(
+    d.vec4f(tangent, 0),
+    d.vec4f(normal, 0),
+    d.vec4f(bitangent, 0),
+    d.vec4f(center, 1),
+  );
+}
diff --git a/packages/typegpu-geometry/src/mesh/parametric.ts b/packages/typegpu-geometry/src/mesh/parametric.ts
new file mode 100644
index 0000000000..ae84912268
--- /dev/null
+++ b/packages/typegpu-geometry/src/mesh/parametric.ts
@@ -0,0 +1,84 @@
+import { d, std, tgpu } from 'typegpu';
+import { type IndexedGeometry, Surface } from './geometry.ts';
+
+const QUAD = tgpu.const(d.arrayOf(d.vec2u, 6), [
+  d.vec2u(0, 0),
+  d.vec2u(1, 0),
+  d.vec2u(1, 1),
+  d.vec2u(0, 0),
+  d.vec2u(1, 1),
+  d.vec2u(0, 1),
+]);
+
+const EPS = 1e-3;
+
+export interface Parametric {
+  at(this: void, u: number, v: number): d.v3f;
+  normalAt?(this: void, u: number, v: number): d.v3f;
+}
+
+export interface ParametricOptions {
+  cols: number;
+  rows: number;
+}
+
+export function parametric(
+  surface: Parametric,
+  { cols, rows }: ParametricOptions,
+): IndexedGeometry {
+  const at = surface.at;
+  const normalAt =
+    surface.normalAt ??
+    ((u: number, v: number) => {
+      'use gpu';
+      const vv = std.clamp(v, 2 * EPS, 1 - 2 * EPS);
+      const du = at(std.min(u + EPS, 1), vv) - at(std.max(u - EPS, 0), vv);
+      const dv = at(u, vv + EPS) - at(u, vv - EPS);
+      const normal = std.cross(du, dv);
+      if (std.dot(normal, normal) > 0) {
+        return std.normalize(normal);
+      }
+
+      return d.vec3f();
+    });
+
+  return sampleGrid(
+    (u, v) => {
+      'use gpu';
+      return Surface({ position: at(u, v), normal: normalAt(u, v), uv: d.vec2f(u, v) });
+    },
+    { cols, rows },
+  );
+}
+
+export function sampleGrid(
+  sample: (u: number, v: number) => d.Infer<typeof Surface>,
+  { cols, rows }: ParametricOptions,
+): IndexedGeometry {
+  return {
+    schema: Surface,
+    topology: 'triangle-list',
+    vertexCount: (cols + 1) * (rows + 1),
+    indexCount: cols * rows * 6,
+
+    vertexAt: (i: number) => {
+      'use gpu';
+      const col = i % (cols + 1);
+      const row = std.intdiv(i, cols + 1);
+      const u = col / cols;
+      const v = row / rows;
+
+      return sample(u, v);
+    },
+
+    indexAt: (i: number) => {
+      'use gpu';
+      const cell = std.intdiv(i, 6);
+      const corner = QUAD.$[i % 6] as d.v2u;
+      const col = (cell % cols) + corner.x;
+      const row = std.intdiv(cell, cols) + corner.y;
+
+      return row * (cols + 1) + col;
+    },
+  };
+}
diff --git a/packages/typegpu-geometry/src/mesh/primitives.ts b/packages/typegpu-geometry/src/mesh/primitives.ts
new file mode 100644
index 0000000000..05a70554a8
--- /dev/null
+++ b/packages/typegpu-geometry/src/mesh/primitives.ts
@@ -0,0 +1,260 @@
+import { d, std, tgpu } from 'typegpu';
+import { concat, transform } from './combinators.ts';
+import { type IndexedGeometry, Surface } from './geometry.ts';
+import { frame, NX, NY, NZ, TAU, X, Y, Z } from './math.ts';
+import { parametric, sampleGrid } from './parametric.ts';
+
+export const surfaces = {
+  sphere: tgpu.fn(
+    [d.f32, d.f32],
+    d.vec3f,
+  )((u, v) => {
+    'use gpu';
+    const theta = v * Math.PI;
+    const phi = u * TAU;
+    const sinTheta = std.sin(theta);
+
+    return d.vec3f(sinTheta * std.cos(phi), std.cos(theta), sinTheta * std.sin(phi));
+  }),
+  plane: tgpu.fn(
+    [d.f32, d.f32],
+    d.vec3f,
+  )((u, v) => {
+    'use gpu';
+    return d.vec3f(u - 0.5, 0, 0.5 - v);
+  }),
+  disc: tgpu.fn(
+    [d.f32, d.f32],
+    d.vec3f,
+  )((u, v) => {
+    'use gpu';
+    const phi = u * TAU;
+    return d.vec3f(v * std.cos(phi), 0, v * std.sin(phi));
+  }),
+  cylinder: tgpu.fn(
+    [d.f32, d.f32],
+    d.vec3f,
+  )((u, v) => {
+    'use gpu';
+    const phi = u * TAU;
+    return d.vec3f(std.cos(phi), 0.5 - v, std.sin(phi));
+  }),
+  torus: tgpu.fn(
+    [d.f32, d.f32, d.f32, d.f32],
+    d.vec3f,
+  )((u, v, radius, tube) => {
+    'use gpu';
+    const phi = u * TAU;
+    const t = v * TAU;
+    const ring = radius + tube * std.cos(t);
+
+    return d.vec3f(ring * std.cos(phi), -tube * std.sin(t), ring * std.sin(phi));
+  }),
+};
+
+export interface SphereOptions {
+  radius?: number;
+  /** Segments around the vertical axis, as in `cylinder` */
+  radialSegments?: number;
+  /** Segments from pole to pole, at least 2 */
+  heightSegments?: number;
+}
+
+export function sphere({
+  radius = 0.5,
+  radialSegments = 32,
+  heightSegments = 16,
+}: SphereOptions = {}): IndexedGeometry {
+  const grid = sampleGrid(
+    (u, v) => {
+      'use gpu';
+      const normal = surfaces.sphere(u, v);
+      return Surface({ position: normal * radius, normal, uv: d.vec2f(u, v) });
+    },
+    { cols: radialSegments, rows: heightSegments },
+  );
+  const fan = radialSegments * 3;
+
+  // The grid's indices are valid, but in the pole rows one triangle of every quad collapses to a
+  // point. Skip those so the mesh has no zero-area triangles
+  return {
+    ...grid,
+    indexCount: 6 * radialSegments * (heightSegments - 1),
+    indexAt: (i: number) => {
+      'use gpu';
+      if (i < fan) {
+        return grid.indexAt(std.intdiv(i, 3) * 6 + (i % 3) + 3);
+      }
+      if (i >= grid.indexCount - 3 * fan) {
+        const last = i - (grid.indexCount - 3 * fan);
+        return grid.indexAt(
+          (heightSegments - 1) * radialSegments * 6 + std.intdiv(last, 3) * 6 + (last % 3),
+        );
+      }
+      return grid.indexAt(i + fan);
+    },
+  };
+}
+
+export interface PlaneOptions {
+  width?: number;
+  depth?: number;
+  widthSegments?: number;
+  depthSegments?: number;
+}
+
+export function plane({
+  width = 1,
+  depth = 1,
+  widthSegments = 1,
+  depthSegments = 1,
+}: PlaneOptions = {}): IndexedGeometry {
+  return parametric(
+    {
+      at: (u, v) => {
+        'use gpu';
+        return surfaces.plane(u, v) * d.vec3f(width, 1, depth);
+      },
+      normalAt: () => {
+        'use gpu';
+        return d.vec3f(0, 1, 0);
+      },
+    },
+    { cols: widthSegments, rows: depthSegments },
+  );
+}
+
+export interface BoxOptions {
+  width?: number;
+  height?: number;
+  depth?: number;
+  widthSegments?: number;
+  heightSegments?: number;
+  depthSegments?: number;
+}
+
+export function box({
+  width = 1,
+  height = 1,
+  depth = 1,
+  widthSegments = 1,
+  heightSegments = 1,
+  depthSegments = 1,
+}: BoxOptions = {}): IndexedGeometry {
+  const top = plane({ width, depth, widthSegments, depthSegments });
+  const front = plane({ width, depth: height, widthSegments, depthSegments: heightSegments });
+  const side = plane({
+    width: depth,
+    depth: height,
+    widthSegments: depthSegments,
+    depthSegments: heightSegments,
+  });
+
+  const x = width / 2;
+  const y = height / 2;
+  const z = depth / 2;
+
+  return concat(
+    transform(top, frame(X, Y, Z, d.vec3f(0, y, 0))),
+    transform(top, frame(X, NY, NZ, d.vec3f(0, -y, 0))),
+    transform(front, frame(X, Z, NY, d.vec3f(0, 0, z))),
+    transform(front, frame(NX, NZ, NY, d.vec3f(0, 0, -z))),
+    transform(side, frame(NZ, X, NY, d.vec3f(x, 0, 0))),
+    transform(side, frame(Z, NX, NY, d.vec3f(-x, 0, 0))),
+  );
+}
+
+export interface CylinderOptions {
+  radius?: number;
+  height?: number;
+  radialSegments?: number;
+  heightSegments?: number;
+  caps?: boolean;
+}
+
+export function cylinder({
+  radius = 0.5,
+  height = 1,
+  radialSegments = 32,
+  heightSegments = 1,
+  caps = true,
+}: CylinderOptions = {}): IndexedGeometry {
+  const side = sampleGrid(
+    (u, v) => {
+      'use gpu';
+      const p = surfaces.cylinder(u, v);
+      return Surface({
+        position: p * d.vec3f(radius, height, radius),
+        normal: d.vec3f(p.x, 0, p.z),
+        uv: d.vec2f(u, v),
+      });
+    },
+    { cols: radialSegments, rows: heightSegments },
+  );
+
+  if (!caps) {
+    return side;
+  }
+
+  const cap: IndexedGeometry = {
+    schema: Surface,
+    topology: 'triangle-list',
+    vertexCount: radialSegments + 1,
+    indexCount: radialSegments * 3,
+    vertexAt: (i: number) => {
+      'use gpu';
+      let direction = d.vec3f();
+      if (i > 0) {
+        direction = surfaces.disc((i - 1) / radialSegments, 1);
+      }
+      return Surface({
+        position: direction * radius,
+        normal: d.vec3f(0, 1, 0),
+        uv: direction.xz * 0.5 + 0.5,
+      });
+    },
+    indexAt: (i: number) => {
+      'use gpu';
+      const triangle = std.intdiv(i, 3);
+      if (i % 3 === 0) {
+        return 0;
+      }
+      if (i % 3 === 1) {
+        return ((triangle + 1) % radialSegments) + 1;
+      }
+      return triangle + 1;
+    },
+  };
+
+  return concat(
+    side,
+    transform(cap, frame(X, Y, Z, d.vec3f(0, height / 2, 0))),
+    transform(cap, frame(X, NY, NZ, d.vec3f(0, -height / 2, 0))),
+  );
+}
+
+export interface TorusOptions {
+  radius?: number;
+  tube?: number;
+  ringSegments?: number;
+  tubeSegments?: number;
+}
+
+export function torus({
+  radius = 0.35,
+  tube = 0.15,
+  ringSegments = 48,
+  tubeSegments = 24,
+}: TorusOptions = {}): IndexedGeometry {
+  return sampleGrid(
+    (u, v) => {
+      'use gpu';
+      return Surface({
+        position: surfaces.torus(u, v, radius, tube),
+        normal: surfaces.torus(u, v, 0, 1),
+        uv: d.vec2f(u, v),
+      });
+    },
+    { cols: ringSegments, rows: tubeSegments },
+  );
+}
diff --git a/packages/typegpu-geometry/tests/bake.test.ts b/packages/typegpu-geometry/tests/bake.test.ts
new file mode 100644
index 0000000000..1b7dc637cd
--- /dev/null
+++ b/packages/typegpu-geometry/tests/bake.test.ts
@@ -0,0 +1,313 @@
+import { expect } from 'vitest';
+import { d, tgpu } from 'typegpu';
+import { it } from 'typegpu-testing-utility';
+import { meshes } from '@typegpu/geometry';
+import { scalars, shaderCodes } from './fixtures.ts';
+
+it('maps and unindexes in a shader', () => {
+  const doubled = meshes.map(scalars, d.f32, (value) => {
+    'use gpu';
+    return value * 2;
+  });
+  const flat = meshes.unindexed(doubled);
+  const read = tgpu.fn([d.u32], d.f32)(flat.vertexAt);
+
+  expect(tgpu.resolve([read])).toMatchInlineSnapshot(`
+    "fn indexAt(i: u32) -> u32 {
+      return (i % 4u);
+    }
+
+    fn vertexAt_2(i: u32) -> f32 {
+      return f32(i);
+    }
+
+    fn item(value: f32) -> f32 {
+      return (value * 2f);
+    }
+
+    fn vertexAt_1(i: u32) -> f32 {
+      return item(vertexAt_2(i));
+    }
+
+    fn vertexAt(i: u32) -> f32 {
+      return vertexAt_1(indexAt(i));
+    }"
+  `);
+});
+
+it('prepares vertex and index fills before submitting', ({ root, device }) => {
+  meshes.bake(root, scalars);
+
+  expect(shaderCodes(device).join('\n\n')).toMatchInlineSnapshot(`
+    "@group(0) @binding(0) var<storage, read_write> output: array<f32, 4>;
+
+    fn vertexAt(i: u32) -> f32 {
+      return f32(i);
+    }
+
+    fn item(i: u32) {
+      output[i] = vertexAt(i);
+    }
+
+    @compute @workgroup_size(64) fn compute(@builtin(global_invocation_id) gid: vec3u) {
+      let i = (gid.x + (gid.y * 64u));
+      if ((i < 4u)) {
+        item(i);
+      }
+    }
+
+    @group(0) @binding(0) var<storage, read_write> output: array<u32, 6>;
+
+    fn indexAt(i: u32) -> u32 {
+      return (i % 4u);
+    }
+
+    fn item(i: u32) {
+      output[i] = indexAt(i);
+    }
+
+    @compute @workgroup_size(64) fn compute(@builtin(global_invocation_id) gid: vec3u) {
+      let i = (gid.x + (gid.y * 64u));
+      if ((i < 6u)) {
+        item(i);
+      }
+    }"
+  `);
+});
+
+it('injects the vertex and index buffers', ({ root, renderPassEncoder }) => {
+  const mesh = meshes.bake(root, meshes.plane());
+  const pipeline = root
+    .createRenderPipeline({
+      attribs: mesh.layout.attrib,
+      vertex: ({ position }) => {
+        'use gpu';
+        return { $position: d.vec4f(position, 1) };
+      },
+    })
+    .pipe(mesh.inject());
+
+  pipeline.with(renderPassEncoder).drawIndexed(mesh.indexCount);
+
+  expect(mesh.layout).toBe(meshes.bake(root, meshes.plane()).layout);
+  expect(renderPassEncoder.mock.setVertexBuffer).toHaveBeenCalledWith(
+    0,
+    root.unwrap(mesh.vertices),
+    undefined,
+    undefined,
+  );
+  expect(renderPassEncoder.mock.setIndexBuffer).toHaveBeenCalledWith(
+    root.unwrap(mesh.indices),
+    'uint32',
+    undefined,
+    undefined,
+  );
+});
+
+it('retains caller-owned buffers and uses geometry counts instead of capacity', ({ root }) => {
+  const vertices = root.createBuffer(d.arrayOf(d.f32, 64)).$usage('vertex', 'storage');
+  const indices = root.createBuffer(d.arrayOf(d.u32, 64)).$usage('index', 'storage');
+  const mesh = meshes.bake(root, scalars, { vertices, indices });
+  mesh.destroy();
+
+  expect([mesh.vertexCount, mesh.indexCount]).toEqual([4, 6]);
+  expect(mesh.vertices).toBe(vertices);
+  expect(mesh.indices).toBe(indices);
+  expect(vertices.destroyed).toBe(false);
+  expect(indices.destroyed).toBe(false);
+});
+
+it('exposes writable vertices and allows mapped reads in another shader', ({ root }) => {
+  const mesh = meshes.bake(root, scalars);
+  const copied = meshes.map(mesh, d.f32, (value) => {
+    'use gpu';
+    return value + 1;
+  });
+  const vertices = mesh.vertices.as('mutable');
+  const edit = tgpu.computeFn({ workgroupSize: [1] })(() => {
+    'use gpu';
+    vertices.$[0] = 1;
+  });
+  const read = tgpu.fn([d.u32], d.f32)(copied.vertexAt);
+
+  expect(tgpu.resolve([edit, read])).toMatchInlineSnapshot(`
+    "@group(0) @binding(0) var<storage, read_write> output: array<f32, 4>;
+
+    @compute @workgroup_size(1) fn edit() {
+      output[0i] = 1f;
+    }
+
+    @group(0) @binding(1) var<storage, read> output_1: array<f32, 4>;
+
+    fn vertexAt_1(i: u32) -> f32 {
+      return output_1[i];
+    }
+
+    fn item(value: f32) -> f32 {
+      return (value + 1f);
+    }
+
+    fn vertexAt(i: u32) -> f32 {
+      return item(vertexAt_1(i));
+    }"
+  `);
+});
+
+it('records vertex and index fills into one pass of a given encoder', ({
+  root,
+  device,
+  commandEncoder,
+  computePassEncoder,
+}) => {
+  const mesh = meshes.bake(root, scalars, { encoder: commandEncoder });
+
+  expect(commandEncoder.mock.beginComputePass).toHaveBeenCalledOnce();
+  expect(computePassEncoder.mock.dispatchWorkgroups).toHaveBeenCalledTimes(2);
+  expect(computePassEncoder.mock.end).toHaveBeenCalledOnce();
+
+  mesh.updateVertices({ pass: computePassEncoder });
+
+  expect(computePassEncoder.mock.dispatchWorkgroups).toHaveBeenCalledTimes(3);
+  expect(computePassEncoder.mock.end).toHaveBeenCalledOnce();
+  expect(device.mock.queue.submit).not.toHaveBeenCalled();
+
+  // @ts-expect-error: work goes either into an encoder or into an open pass
+  mesh.updateVertices({ encoder: commandEncoder, pass: computePassEncoder });
+});
+
+it('bakes several meshes into shared buffers at offsets', ({ root, device }) => {
+  const vertices = root.createBuffer(d.arrayOf(d.f32, 8)).$usage('vertex', 'storage');
+  const indices = root.createBuffer(d.arrayOf(d.u32, 12)).$usage('index', 'storage');
+  const first = meshes.bake(root, scalars, { vertices, indices });
+  const second = meshes.bake(root, scalars, { vertices, indices, baseVertex: 4, firstIndex: 6 });
+
+  expect([first.baseVertex, first.firstIndex]).toEqual([0, 0]);
+  expect([second.baseVertex, second.firstIndex]).toEqual([4, 6]);
+  expect(() => meshes.bake(root, scalars, { vertices, indices, baseVertex: 5 })).toThrow(
+    'The vertex range 5..9 exceeds the buffer capacity 8',
+  );
+
+  const fills = shaderCodes(device).slice(2).join('\n\n');
+  expect(fills).toContain('vertices[(4u + i)] = vertexAt(i);');
+  expect(fills).toContain('indices[(6u + i)] = indexAt(i);');
+
+  const read = tgpu.fn(
+    [d.u32],
+    d.f32,
+  )((i) => {
+    'use gpu';
+    return second.vertexAt(second.indexAt(i));
+  });
+  expect(tgpu.resolve([read])).toMatchInlineSnapshot(`
+    "@group(0) @binding(0) var<storage, read> indices: array<u32, 12>;
+
+    fn indexAt(i: u32) -> u32 {
+      return indices[(6u + i)];
+    }
+
+    @group(0) @binding(1) var<storage, read> vertices: array<f32, 8>;
+
+    fn vertexAt(i: u32) -> f32 {
+      return vertices[(4u + i)];
+    }
+
+    fn read_1(i: u32) -> f32 {
+      return vertexAt(indexAt(i));
+    }"
+  `);
+});
+
+it('bakes only the indices and keeps reading vertices from the source', ({ root, device }) => {
+  const baked = meshes.bakeIndices(root, scalars);
+  const read = tgpu.fn(
+    [d.u32],
+    d.f32,
+  )((i) => {
+    'use gpu';
+    return baked.vertexAt(baked.indexAt(i));
+  });
+
+  expect(shaderCodes(device).join('\n\n')).toMatchInlineSnapshot(`
+    "@group(0) @binding(0) var<storage, read_write> output: array<u32, 6>;
+
+    fn indexAt(i: u32) -> u32 {
+      return (i % 4u);
+    }
+
+    fn item(i: u32) {
+      output[i] = indexAt(i);
+    }
+
+    @compute @workgroup_size(64) fn compute(@builtin(global_invocation_id) gid: vec3u) {
+      let i = (gid.x + (gid.y * 64u));
+      if ((i < 6u)) {
+        item(i);
+      }
+    }"
+  `);
+  expect(tgpu.resolve([read])).toMatchInlineSnapshot(`
+    "@group(0) @binding(0) var<storage, read> output: array<u32, 6>;
+
+    fn indexAt(i: u32) -> u32 {
+      return output[i];
+    }
+
+    fn vertexAt(i: u32) -> f32 {
+      return f32(i);
+    }
+
+    fn read_1(i: u32) -> f32 {
+      return vertexAt(indexAt(i));
+    }"
+  `);
+});
+
+it('injects only the index buffer', ({ root, renderPassEncoder }) => {
+  const baked = meshes.bakeIndices(root, scalars);
+  const pipeline = root
+    .createRenderPipeline({
+      vertex: () => {
+        'use gpu';
+        return { $position: d.vec4f() };
+      },
+    })
+    .pipe(baked.inject());
+
+  pipeline.with(renderPassEncoder).drawIndexed(baked.indexCount);
+
+  expect(renderPassEncoder.mock.setVertexBuffer).not.toHaveBeenCalled();
+  expect(renderPassEncoder.mock.setIndexBuffer).toHaveBeenCalledWith(
+    root.unwrap(baked.indices),
+    'uint32',
+    undefined,
+    undefined,
+  );
+});
+
+it('bakes indices into a caller-owned buffer at an offset', ({ root, device }) => {
+  const indices = root.createBuffer(d.arrayOf(d.u32, 12)).$usage('index', 'storage');
+  const baked = meshes.bakeIndices(root, scalars, { indices, firstIndex: 6 });
+  baked.destroy();
+
+  expect([baked.firstIndex, baked.indexCount]).toEqual([6, 6]);
+  expect(baked.indices).toBe(indices);
+  expect(indices.destroyed).toBe(false);
+  expect(shaderCodes(device).join('\n\n')).toContain('indices[(6u + i)] = indexAt(i);');
+  expect(() => meshes.bakeIndices(root, scalars, { indices, firstIndex: 7 })).toThrow(
+    'The index range 7..13 exceeds the buffer capacity 12',
+  );
+});
+
+it('records the index fill into a given encoder', ({
+  root,
+  device,
+  commandEncoder,
+  computePassEncoder,
+}) => {
+  meshes.bakeIndices(root, scalars, { encoder: commandEncoder });
+
+  expect(commandEncoder.mock.beginComputePass).toHaveBeenCalledOnce();
+  expect(computePassEncoder.mock.dispatchWorkgroups).toHaveBeenCalledOnce();
+  expect(computePassEncoder.mock.end).toHaveBeenCalledOnce();
+  expect(device.mock.queue.submit).not.toHaveBeenCalled();
+});
diff --git a/packages/typegpu-geometry/tests/fixtures.ts b/packages/typegpu-geometry/tests/fixtures.ts
new file mode 100644
index 0000000000..b32ff1c745
--- /dev/null
+++ b/packages/typegpu-geometry/tests/fixtures.ts
@@ -0,0 +1,22 @@
+import { vi } from 'vitest';
+import { d } from 'typegpu';
+import type { meshes } from '@typegpu/geometry';
+
+export const scalars: meshes.IndexedGeometry<d.F32> = {
+  schema: d.f32,
+  topology: 'triangle-list',
+  vertexCount: 4,
+  indexCount: 6,
+  vertexAt: (i) => {
+    'use gpu';
+    return d.f32(i);
+  },
+  indexAt: (i) => {
+    'use gpu';
+    return i % 4;
+  },
+};
+
+export function shaderCodes(device: GPUDevice) {
+  return vi.mocked(device.createShaderModule).mock.calls.map(([{ code }]) => code);
+}
diff --git a/packages/typegpu-geometry/tests/lines.test.ts b/packages/typegpu-geometry/tests/lines.test.ts
new file mode 100644
index 0000000000..684a048e97
--- /dev/null
+++ b/packages/typegpu-geometry/tests/lines.test.ts
@@ -0,0 +1,47 @@
+import { expect } from 'vitest';
+import { d, std, tgpu } from 'typegpu';
+import { it } from 'typegpu-testing-utility';
+import {
+  JoinInput,
+  joins,
+  joinSlot,
+  LineControlPoint,
+  miterJoinLimitSlot,
+  polylineVariableWidth,
+} from '@typegpu/geometry';
+
+const turn = JoinInput({
+  C: LineControlPoint({ position: d.vec2f(1, 2), radius: 0.5 }),
+  v: d.vec2f(1, 2),
+  d: d.vec2f(1, 0),
+  fw: d.vec2f(0, 1),
+  cross: d.vec2f(1, 0),
+  start: d.vec2f(1, 0),
+  end: d.vec2f(0, 1),
+});
+
+it('cuts a bevel join straight between the sides of the two segments', () => {
+  const points = [0, 1, 2, 3, 4].map((i) => joins.bevel(turn, i, 4));
+  const start = d.vec2f(1.5, 2);
+  const end = d.vec2f(1, 2.5);
+
+  expect(points[0]).toEqual(turn.v);
+  expect(points[4]).toEqual(std.mix(start, end, 0.5));
+  for (const point of points.slice(1)) {
+    expect(std.dot(point.sub(start), d.vec2f(1, 1))).toBeCloseTo(0);
+  }
+});
+
+it('configures joins through exported slots', () => {
+  const vertex = tgpu.fn(
+    [LineControlPoint, d.u32],
+    d.vec2f,
+  )((point, i) => {
+    'use gpu';
+    return polylineVariableWidth(point, point, point, point, i, 4).vertexPosition;
+  });
+
+  const code = tgpu.resolve([vertex.with(joinSlot, joins.miter).with(miterJoinLimitSlot, 3)]);
+
+  expect(code).toContain('miterLimit(miterPointBisectorWhenClockwise(join.start, join.end), 3i)');
+});
diff --git a/packages/typegpu-geometry/tests/mesh.test.ts b/packages/typegpu-geometry/tests/mesh.test.ts
new file mode 100644
index 0000000000..8ad3c6b5d2
--- /dev/null
+++ b/packages/typegpu-geometry/tests/mesh.test.ts
@@ -0,0 +1,211 @@
+import { describe, expect } from 'vitest';
+import { d, std } from 'typegpu';
+import { it } from 'typegpu-testing-utility';
+import { meshes } from '@typegpu/geometry';
+
+function windingAgreesWithNormals(g: meshes.IndexedGeometry) {
+  if (g.topology !== 'triangle-list') {
+    throw new Error('Invalid topology.');
+  }
+  for (let t = 0; t < g.indexCount; t += 3) {
+    const a = g.vertexAt(g.indexAt(t));
+    const b = g.vertexAt(g.indexAt(t + 1));
+    const c = g.vertexAt(g.indexAt(t + 2));
+    const face = std.cross(b.position.sub(a.position), c.position.sub(a.position));
+    if (std.length(face) < 1e-9) {
+      continue;
+    }
+    const n = a.normal.add(b.normal).add(c.normal);
+    if (std.dot(face, n) <= 0) {
+      return false;
+    }
+  }
+  return true;
+}
+
+describe('meshes on the CPU', () => {
+  it.each([
+    ['plane', meshes.plane()],
+    ['sphere', meshes.sphere({ radialSegments: 8, heightSegments: 4 })],
+    ['box', meshes.box()],
+    ['cylinder', meshes.cylinder({ radialSegments: 6 })],
+    ['torus', meshes.torus({ ringSegments: 6, tubeSegments: 4 })],
+  ])('%s fits a unit cube with outward unit normals', (_, shape) => {
+    expect(windingAgreesWithNormals(shape)).toBe(true);
+
+    for (let i = 0; i < shape.vertexCount; i++) {
+      const { position, normal } = shape.vertexAt(i);
+      expect(
+        Math.max(Math.abs(position.x), Math.abs(position.y), Math.abs(position.z)),
+      ).toBeLessThanOrEqual(0.5 + 1e-6);
+      expect(std.length(normal)).toBeCloseTo(1);
+    }
+  });
+
+  it('scales each box dimension independently', () => {
+    const slab = meshes.box({ width: 2, height: 0.5, depth: 4 });
+    let extent = d.vec3f();
+    for (let i = 0; i < slab.vertexCount; i++) {
+      extent = std.max(extent, std.abs(slab.vertexAt(i).position));
+    }
+
+    expect([...extent]).toEqual([1, 0.25, 2]);
+  });
+
+  it('keeps collapsed surface normals zero through transforms', () => {
+    const collapsed = meshes.parametric(
+      {
+        at: () => {
+          'use gpu';
+          return d.vec3f();
+        },
+      },
+      { cols: 1, rows: 1 },
+    );
+    const moved = meshes.transform(collapsed, std.translation4(d.vec3f(1, 2, 3)));
+
+    expect(collapsed.vertexAt(0).normal).toEqual(d.vec3f());
+    expect(moved.vertexAt(0).normal).toEqual(d.vec3f());
+  });
+
+  it('keeps concat indexed and offsets the indices of later parts', () => {
+    const a = meshes.plane();
+    const b = meshes.transform(meshes.plane(), std.translation4(d.vec3f(0, 1, 0)));
+    const g = meshes.concat(a, b);
+
+    expect(g.vertexCount).toBe(8);
+    expect(g.indexCount).toBe(12);
+    expect(g.indexAt(6)).toBe(4);
+    expect(g.vertexAt(g.indexAt(6)).position.y).toBeCloseTo(1);
+
+    const flat = meshes.unindexed(g);
+    expect(flat.vertexCount).toBe(12);
+    expect(flat.vertexAt(11).position.y).toBeCloseTo(1);
+
+    const mixed = meshes.concat(a, { ...meshes.unindexed(b), vertexCount: 3 });
+    expect([mixed.vertexCount, mixed.indexCount]).toEqual([7, 9]);
+    expect([6, 7, 8].map(mixed.indexAt)).toEqual([4, 5, 6]);
+  });
+
+  it('returns zero vertices past the end of concatenated geometry', () => {
+    const plane = meshes.plane();
+    const scalar = meshes.map(plane, d.f32, (vertex) => {
+      'use gpu';
+      return vertex.position.x;
+    });
+    const scalars: meshes.IndexedGeometry<d.F32> = meshes.concat(scalar, scalar);
+    const surfaces = meshes.concat(plane, plane);
+
+    expect(scalars.vertexAt(scalars.vertexCount)).toBe(0);
+    expect(surfaces.vertexAt(surfaces.vertexCount)).toEqual(meshes.Surface());
+  });
+
+  it('attaches fields and shares the schema between calls', () => {
+    const Paint = d.struct({ color: d.vec3f, id: d.u32 });
+    const paint = (g: meshes.IndexedGeometry, id: number) =>
+      meshes.attach(g, Paint, (s) => {
+        'use gpu';
+        return Paint({ color: d.vec3f(s.uv, 0), id });
+      });
+    const a = paint(meshes.plane(), 1);
+    const b = paint(meshes.sphere({ radialSegments: 4, heightSegments: 2 }), 2);
+    const v = a.vertexAt(3);
+    const scene = meshes.concat(a, b);
+
+    expect(a.schema).toBe(b.schema);
+    expect(Object.keys(a.schema.propTypes)).toEqual(['position', 'normal', 'uv', 'color', 'id']);
+    expect(v.position.x).toBeCloseTo(0.5);
+    expect(v.color.x).toBeCloseTo(1);
+    expect(v.id).toBe(1);
+    expect(scene.vertexAt(4).id).toBe(2);
+  });
+
+  it('throws when attaching a field that is already present', () => {
+    const Paint = d.struct({ position: d.vec3f });
+    expect(() =>
+      meshes.attach(meshes.plane(), Paint, (s) => {
+        'use gpu';
+        return Paint({ position: d.vec3f(s.uv, 0) });
+      }),
+    ).toThrowErrorMatchingInlineSnapshot(
+      `[Error: attach cannot replace the existing field 'position'; use map instead]`,
+    );
+  });
+
+  it('preserves outward winding under reflections', () => {
+    for (const source of [meshes.plane(), meshes.unindexed(meshes.plane())]) {
+      const flat = meshes.unindexed(meshes.transform(source, std.scaling4(d.vec3f(-2, 3, 0.5))));
+      const a = flat.vertexAt(0),
+        b = flat.vertexAt(1),
+        c = flat.vertexAt(2);
+      const face = std.cross(b.position.sub(a.position), c.position.sub(a.position));
+
+      expect(std.dot(face, a.normal)).toBeGreaterThan(0);
+      expect(std.length(a.normal)).toBeCloseTo(1);
+    }
+  });
+
+  it('keeps normals of transforms that flatten a shape', () => {
+    const decal = meshes.transform(meshes.plane(), std.scaling4(d.vec3f(2, 0, 1)));
+    expect(decal.vertexAt(0).normal).toEqual(d.vec3f(0, 1, 0));
+
+    const coin = meshes.transform(meshes.sphere(), std.scaling4(d.vec3f(1, 0, -1)));
+    for (let i = 0; i < coin.vertexCount; i++) {
+      const { position, normal } = coin.vertexAt(i);
+      expect(position.y).toBe(0);
+      if (std.length(normal) > 0) {
+        expect(Math.abs(normal.y)).toBeCloseTo(1);
+      }
+    }
+  });
+
+  it('transforms a custom schema and keeps its other fields', () => {
+    const Tagged = d.struct({ position: d.vec3f, normal: d.vec3f, tag: d.f32 });
+    const tagged = meshes.map(meshes.plane(), Tagged, (s) => {
+      'use gpu';
+      return Tagged({ position: s.position, normal: s.normal, tag: s.uv.x });
+    });
+    const moved = meshes.transform(tagged, std.translation4(d.vec3f(0, 2, 0)));
+
+    expect(moved.schema).toBe(Tagged);
+    const v = moved.vertexAt(1);
+    expect(v.position.y).toBe(2);
+    expect(v.tag).toBe(1);
+  });
+
+  it('samples normals inside the unit square', () => {
+    const shape = meshes.parametric(
+      {
+        at: (u, v) => {
+          'use gpu';
+          return d.vec3f(std.sqrt(u), 0, 1 - v);
+        },
+      },
+      { cols: 1, rows: 1 },
+    );
+    for (let i = 0; i < shape.vertexCount; i++) {
+      expect(std.distance(shape.vertexAt(i).normal, d.vec3f(0, 1, 0))).toBeCloseTo(0);
+    }
+  });
+
+  it('does not mutate cached source vertices when transforming', () => {
+    const vertex = meshes.plane().vertexAt(0);
+    const shape = { ...meshes.plane(), vertexAt: () => vertex };
+    const moved = meshes.transform(shape, std.translation4(d.vec3f(0, 2, 0)));
+    expect(moved.vertexAt(0).position.y).toBe(2);
+    expect(moved.vertexAt(0).position.y).toBe(2);
+    expect(vertex.position.y).toBe(0);
+  });
+
+  it.each([meshes.sphere(), meshes.cylinder()])(
+    'omits collapsed pole and cap triangles',
+    (shape) => {
+      for (let i = 0; i < shape.indexCount; i += 3) {
+        const a = shape.vertexAt(shape.indexAt(i)).position;
+        const b = shape.vertexAt(shape.indexAt(i + 1)).position;
+        const c = shape.vertexAt(shape.indexAt(i + 2)).position;
+        expect(std.length(std.cross(b.sub(a), c.sub(a)))).toBeGreaterThan(1e-9);
+      }
+    },
+  );
+});
diff --git a/packages/typegpu-geometry/tests/updates.test.ts b/packages/typegpu-geometry/tests/updates.test.ts
new file mode 100644
index 0000000000..b471987db6
--- /dev/null
+++ b/packages/typegpu-geometry/tests/updates.test.ts
@@ -0,0 +1,114 @@
+import { expect } from 'vitest';
+import { it } from 'typegpu-testing-utility';
+import { d, tgpu } from 'typegpu';
+import { meshes } from '@typegpu/geometry';
+import { scalars, shaderCodes } from './fixtures.ts';
+
+it('bakes indices from baked geometry without reading GPU values on the CPU', ({ root }) => {
+  const source = meshes.bake(root, scalars);
+  const mesh = meshes.bakeIndices(root, source);
+  mesh.destroy();
+
+  expect([mesh.vertexCount, mesh.indexCount]).toEqual([4, 6]);
+  expect(mesh.indices.destroyed).toBe(true);
+  expect(source.vertices.destroyed).toBe(false);
+  expect(source.indices.destroyed).toBe(false);
+});
+
+it('specializes the vertex evaluator through a configured pipeline builder', ({ root }) => {
+  const height = tgpu.accessor(d.f32);
+  const source = {
+    ...scalars,
+    vertexAt: () => {
+      'use gpu';
+      return height.$;
+    },
+  };
+  const mesh = meshes.bake(root, source, { with: root.with(height, 2) });
+  expect(mesh.vertexCount).toBe(4);
+});
+
+it('overrides source bindings for one update', ({ root, device, commandEncoder }) => {
+  const config = tgpu.bindGroupLayout({ height: { uniform: d.f32 } });
+  const initial = root.createBindGroup(config, { height: root.createUniform(d.f32, 2) });
+  const override = root.createBindGroup(config, { height: root.createUniform(d.f32, 3) });
+  const mesh = meshes.bake(
+    root,
+    {
+      schema: d.f32,
+      topology: 'point-list',
+      vertexCount: 1,
+      vertexAt: () => {
+        'use gpu';
+        return config.$.height;
+      },
+    },
+    { bindGroups: [initial] },
+  );
+  const pass = commandEncoder.mock.beginComputePass.mock.results[0]?.value;
+  pass.setBindGroup.mockClear();
+
+  mesh.updateVertices({ bindGroups: [override] });
+  mesh.updateVertices();
+
+  const inputs = pass.setBindGroup.mock.calls
+    .map((call: unknown[]) => call[1])
+    .filter((group: unknown) => group === root.unwrap(initial) || group === root.unwrap(override));
+  expect(inputs).toEqual([root.unwrap(override), root.unwrap(initial)]);
+  expect(device.mock.createComputePipeline).toHaveBeenCalledOnce();
+});
+
+it('spans multiple dispatch rows without evaluating vertices', ({
+  root,
+  device,
+  commandEncoder,
+}) => {
+  const position = tgpu.accessor(d.f32);
+  meshes.bakeIndices(root, {
+    ...scalars,
+    indexCount: 64 * 65535 + 1,
+    vertexAt: () => {
+      'use gpu';
+      return position.$;
+    },
+  });
+  const pass = commandEncoder.mock.beginComputePass.mock.results[0]?.value;
+
+  expect(device.mock.createBuffer).toHaveBeenCalledOnce();
+  expect(pass.dispatchWorkgroups).toHaveBeenCalledWith(65535, 2, undefined);
+  expect(shaderCodes(device)[0]).toMatchInlineSnapshot(`
+    "@group(0) @binding(0) var<storage, read_write> output: array<u32, 4194241>;
+
+    fn indexAt(i: u32) -> u32 {
+      return (i % 4u);
+    }
+
+    fn item(i: u32) {
+      output[i] = indexAt(i);
+    }
+
+    @compute @workgroup_size(64) fn compute(@builtin(global_invocation_id) gid: vec3u) {
+      let i = (gid.x + (gid.y * 4194240u));
+      if ((i < 4194241u)) {
+        item(i);
+      }
+    }"
+  `);
+});
+
+it('prepares index-only baking asynchronously without evaluating vertices', async ({
+  root,
+  device,
+}) => {
+  const position = tgpu.accessor(d.f32);
+  const mesh = await meshes.bakeIndicesAsync(root, {
+    ...scalars,
+    vertexAt: () => {
+      'use gpu';
+      return position.$;
+    },
+  });
+  expect(mesh.indexCount).toBe(6);
+  expect(device.mock.createComputePipelineAsync).toHaveBeenCalledOnce();
+  expect(device.mock.createComputePipeline).not.toHaveBeenCalled();
+});
diff --git a/packages/typegpu-geometry/tsconfig.test.json b/packages/typegpu-geometry/tsconfig.test.json
new file mode 100644
index 0000000000..f3b3ce014a
--- /dev/null
+++ b/packages/typegpu-geometry/tsconfig.test.json
@@ -0,0 +1,4 @@
+{
+  "extends": "./tsconfig.json",
+  "include": ["src/**/*", "tests/**/*", "vitest.config.mts"]
+}
diff --git a/packages/typegpu-geometry/tsdown.config.ts b/packages/typegpu-geometry/tsdown.config.ts
new file mode 100644
index 0000000000..07ad9608d2
--- /dev/null
+++ b/packages/typegpu-geometry/tsdown.config.ts
@@ -0,0 +1,14 @@
+import { defineConfig } from 'tsdown';
+import typegpu from 'unplugin-typegpu/rolldown';
+
+export default defineConfig({
+  entry: ['src/index.ts'],
+  outDir: 'dist',
+  format: 'esm',
+  dts: true,
+  platform: 'neutral',
+  unbundle: true,
+  sourcemap: false,
+  target: false,
+  plugins: [typegpu({ include: [/\.ts$/], exclude: [/\.d\.ts$/] })],
+});
diff --git a/packages/typegpu-geometry/vitest.config.mts b/packages/typegpu-geometry/vitest.config.mts
new file mode 100644
index 0000000000..3367210cff
--- /dev/null
+++ b/packages/typegpu-geometry/vitest.config.mts
@@ -0,0 +1,16 @@
+import { createJiti } from 'jiti';
+import type TypeGPUPlugin from 'unplugin-typegpu/vite';
+import { defineConfig } from 'vitest/config';
+import { typegpuBuiltAliases } from 'typegpu-testing-utility/config';
+
+const jiti = createJiti(import.meta.url);
+const typegpu = await jiti.import<typeof TypeGPUPlugin>('unplugin-typegpu/vite', {
+  default: true,
+});
+
+export default defineConfig({
+  plugins: [typegpu({ forceTgpuAlias: 'tgpu', earlyPruning: false })],
+  resolve: {
+    alias: typegpuBuiltAliases(),
+  },
+});
diff --git a/pnpm-lock.yaml b/pnpm-lock.yaml
index 2d47e15f24..013756a312 100644
--- a/pnpm-lock.yaml
+++ b/pnpm-lock.yaml
@@ -781,18 +781,27 @@ importers:
       '@webgpu/types':
         specifier: catalog:types
         version: 0.1.71
+      jiti:
+        specifier: catalog:build
+        version: 2.6.1
+      tsdown:
+        specifier: catalog:build
+        version: 0.15.12(@emnapi/core@1.9.2)(@emnapi/runtime@1.9.2)(tsover@6.0.2)(unrun@0.2.31(@emnapi/core@1.9.2)(@emnapi/runtime@1.9.2))
       typegpu:
         specifier: workspace:*
         version: link:../typegpu
+      typegpu-testing-utility:
+        specifier: workspace:*
+        version: link:../typegpu-testing-utility
       typescript:
         specifier: npm:tsover@^6.0.2
         version: tsover@6.0.2
-      unbuild:
-        specifier: catalog:build
-        version: 3.5.0(tsover@6.0.2)
       unplugin-typegpu:
         specifier: workspace:*
         version: link:../unplugin-typegpu
+      vitest:
+        specifier: catalog:test
+        version: 4.1.11(@types/node@24.10.0)(@vitest/browser-preview@4.1.11)(@vitest/coverage-v8@4.1.11)(jsdom@27.0.0(postcss@8.5.8))(vite@8.0.5(@emnapi/core@1.9.2)(@emnapi/runtime@1.9.2)(@types/node@24.10.0)(esbuild@0.27.5)(jiti@2.6.1)(terser@5.51.2)(tsx@4.20.6)(yaml@2.9.1))
     publishDirectory: dist
 
   packages/typegpu-gl: