+
+
+
-
-
result matrix
-
+
+
Result Matrix
+
+
Press Compute to calculate the result
+
+
diff --git a/apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/index.ts b/apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/index.ts
index af6a8daf89..600db8e552 100644
--- a/apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/index.ts
+++ b/apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/index.ts
@@ -1,177 +1,278 @@
-import { tgpu, d } from 'typegpu';
+import { tgpu } from 'typegpu';
+import { INITIAL_MAX_MATRIX_SIZE, TILE_SIZE } from './params.ts';
+import { type CalculationStrategy, computeLayout, createMatrixData, MatrixInfo } from './types.ts';
+import { computeSharedMemory } from './computeShared.ts';
+import { computeSimple } from './computeSimple.ts';
+import { multiplyMatricesCPU } from './computeCpu.ts';
import { defineControls } from '../../common/defineControls.ts';
-const WORKGROUP_SIZE = [8, 8] as [number, number];
-const MAX_MATRIX_SIZE = 6;
-
-let firstRowCount = 3;
-let firstColumnCount = 4;
-let secondColumnCount = 2;
+const state = {
+ dimensions: { firstRowCount: 3, firstColumnCount: 4, secondColumnCount: 2 },
+ strategy: 'gpu-optimized' as CalculationStrategy,
+ bufferCapacity: INITIAL_MAX_MATRIX_SIZE ** 2,
+ matrices: {
+ first: [] as number[],
+ second: [] as number[],
+ result: [] as number[],
+ },
+ gpuTime: 0,
+};
-const MatrixStruct = d.struct({
- size: d.vec2f,
- numbers: d.arrayOf(d.f32, MAX_MATRIX_SIZE ** 2),
+const root = await tgpu.init({
+ device: {
+ optionalFeatures: ['timestamp-query'],
+ },
});
+const hasTimestampQuery = root.enabledFeatures.has('timestamp-query');
-const layout = tgpu.bindGroupLayout({
- firstMatrix: { storage: MatrixStruct, access: 'readonly' },
- secondMatrix: { storage: MatrixStruct, access: 'readonly' },
- resultMatrix: { storage: MatrixStruct, access: 'mutable' },
-});
+const buffers = {
+ first: root.createBuffer(createMatrixData(INITIAL_MAX_MATRIX_SIZE)).$usage('storage'),
+ second: root.createBuffer(createMatrixData(INITIAL_MAX_MATRIX_SIZE)).$usage('storage'),
+ result: root.createBuffer(createMatrixData(INITIAL_MAX_MATRIX_SIZE)).$usage('storage'),
+ info: root.createBuffer(MatrixInfo, state.dimensions).$usage('uniform'),
+};
+
+let bindGroup = createBindGroup();
-const shaderCode = /* wgsl */ `
+const pipelines = createPipelines();
-@group(0) @binding(0) var
firstMatrix: MatrixStruct;
-@group(0) @binding(1) var secondMatrix: MatrixStruct;
-@group(0) @binding(2) var resultMatrix: MatrixStruct;
+function createBindGroup() {
+ return root.createBindGroup(computeLayout, {
+ firstMatrix: buffers.first,
+ secondMatrix: buffers.second,
+ resultMatrix: buffers.result,
+ dimensions: buffers.info,
+ });
+}
+
+function createPipelines() {
+ const performanceCallback = (start: bigint, end: bigint) => {
+ state.gpuTime = Number(end - start) / 1_000_000;
+ };
-@compute @workgroup_size(${WORKGROUP_SIZE.join(', ')})
-fn main(@builtin(global_invocation_id) global_id: vec3u) {
- if (global_id.x >= u32(firstMatrix.size.x) || global_id.y >= u32(secondMatrix.size.y)) {
+ const optimized = root.createComputePipeline({
+ compute: computeSharedMemory,
+ });
+ const simple = root.createComputePipeline({
+ compute: computeSimple,
+ });
+
+ return {
+ 'gpu-optimized': hasTimestampQuery
+ ? optimized.withPerformanceCallback(performanceCallback)
+ : optimized,
+ 'gpu-simple': hasTimestampQuery ? simple.withPerformanceCallback(performanceCallback) : simple,
+ };
+}
+
+function resizeBuffersIfNeeded(requiredSize: number) {
+ if (requiredSize <= state.bufferCapacity) {
return;
}
- if (global_id.x + global_id.y == 0u) {
- resultMatrix.size = vec2(firstMatrix.size.x, secondMatrix.size.y);
+ while (state.bufferCapacity < requiredSize) {
+ state.bufferCapacity *= 2;
}
- let resultCell = vec2(global_id.x, global_id.y);
- var result = 0.0;
-
- for (var i = 0u; i < u32(firstMatrix.size.y); i = i + 1u) {
- let a = i + resultCell.x * u32(firstMatrix.size.y);
- let b = resultCell.y + i * u32(secondMatrix.size.y);
- result = result + firstMatrix.numbers[a] * secondMatrix.numbers[b];
+ for (const buffer of [buffers.first, buffers.second, buffers.result]) {
+ buffer.destroy();
}
- let index = resultCell.y + resultCell.x * u32(secondMatrix.size.y);
- resultMatrix.numbers[index] = result;
-}`;
-
-const root = await tgpu.init();
-const device = root.device;
-
-const firstMatrixBuffer = root.createBuffer(MatrixStruct).$usage('storage');
-const secondMatrixBuffer = root.createBuffer(MatrixStruct).$usage('storage');
-const resultMatrixBuffer = root.createBuffer(MatrixStruct).$usage('storage');
-
-const pipeline = device.createComputePipeline({
- layout: device.createPipelineLayout({
- bindGroupLayouts: [root.unwrap(layout)],
- }),
- compute: {
- module: device.createShaderModule({
- code: tgpu.resolve({
- template: shaderCode,
- externals: { MatrixStruct },
- }),
- }),
- },
-});
-
-const bindGroup = root.createBindGroup(layout, {
- firstMatrix: firstMatrixBuffer,
- secondMatrix: secondMatrixBuffer,
- resultMatrix: resultMatrixBuffer,
-});
+ buffers.first = root.createBuffer(createMatrixData(state.bufferCapacity)).$usage('storage');
+ buffers.second = root.createBuffer(createMatrixData(state.bufferCapacity)).$usage('storage');
+ buffers.result = root.createBuffer(createMatrixData(state.bufferCapacity)).$usage('storage');
-function createMatrix(size: d.v2f, initValue: (row: number, col: number) => number) {
- return {
- size: size,
- numbers: Array(size.x * size.y)
- .fill(0)
- .map((_, i) => initValue(Math.floor(i / size.y), i % size.y)),
- };
+ bindGroup = createBindGroup();
}
-async function run() {
- const firstMatrix = createMatrix(d.vec2f(firstRowCount, firstColumnCount), () =>
- Math.floor(Math.random() * 10),
- );
- const secondMatrix = createMatrix(d.vec2f(firstColumnCount, secondColumnCount), () =>
- Math.floor(Math.random() * 10),
- );
+function createMatrix(rows: number, cols: number) {
+ return Array.from({ length: rows * cols }, () => Math.floor(Math.random() * 10));
+}
- firstMatrixBuffer.write(firstMatrix);
- secondMatrixBuffer.write(secondMatrix);
+function generateMatrices() {
+ const { firstRowCount, firstColumnCount, secondColumnCount } = state.dimensions;
- const workgroupCountX = Math.ceil(firstMatrix.size.x / WORKGROUP_SIZE[0]);
- const workgroupCountY = Math.ceil(secondMatrix.size.y / WORKGROUP_SIZE[1]);
+ state.matrices.first = createMatrix(firstRowCount, firstColumnCount);
+ state.matrices.second = createMatrix(firstColumnCount, secondColumnCount);
- const encoder = device.createCommandEncoder();
+ updateInputDisplays();
+}
- const pass = encoder.beginComputePass();
- pass.setPipeline(pipeline);
- pass.setBindGroup(0, root.unwrap(bindGroup));
- pass.dispatchWorkgroups(workgroupCountX, workgroupCountY);
- pass.end();
+function updateInputDisplays() {
+ const { firstRowCount, firstColumnCount, secondColumnCount } = state.dimensions;
- device.queue.submit([encoder.finish()]);
+ printMatrixToHtml(firstTable, state.matrices.first, firstRowCount, firstColumnCount);
+ printMatrixToHtml(secondTable, state.matrices.second, firstColumnCount, secondColumnCount);
+}
- const multiplicationResult = await resultMatrixBuffer.read();
+let isComputing = false;
- printMatrixToHtml(firstTable, firstMatrix);
- printMatrixToHtml(secondTable, secondMatrix);
- printMatrixToHtml(resultTable, multiplicationResult);
+async function compute() {
+ if (isComputing) {
+ console.warn('Computation already in progress');
+ return;
+ }
+ isComputing = true;
+
+ try {
+ const { firstRowCount, firstColumnCount, secondColumnCount } = state.dimensions;
+ const maxSize = Math.max(
+ firstRowCount * firstColumnCount,
+ firstColumnCount * secondColumnCount,
+ firstRowCount * secondColumnCount,
+ );
+
+ const startTime = performance.now();
+
+ if (state.strategy === 'cpu') {
+ state.matrices.result = multiplyMatricesCPU(
+ state.matrices.first,
+ state.matrices.second,
+ firstRowCount,
+ firstColumnCount,
+ secondColumnCount,
+ );
+ } else {
+ resizeBuffersIfNeeded(maxSize);
+ buffers.info.write(state.dimensions);
+ buffers.first.write(state.matrices.first);
+ buffers.second.write(state.matrices.second);
+
+ const workgroupCount = {
+ x: Math.ceil(firstRowCount / TILE_SIZE),
+ y: Math.ceil(secondColumnCount / TILE_SIZE),
+ };
+
+ pipelines[state.strategy]
+ .with(bindGroup)
+ .dispatchWorkgroups(workgroupCount.x, workgroupCount.y);
+
+ await root.device.queue.onSubmittedWorkDone();
+
+ const gpuResult = await buffers.result.read();
+ state.matrices.result = gpuResult.slice(0, firstRowCount * secondColumnCount);
+ }
+
+ const totalTime = performance.now() - startTime;
+ const strategyName = {
+ 'gpu-optimized': 'GPU (Optimized)',
+ 'gpu-simple': 'GPU (Simple)',
+ cpu: 'CPU',
+ }[state.strategy];
+
+ const showGpu = state.strategy !== 'cpu' && hasTimestampQuery;
+ updateTimingDisplay(strategyName, totalTime, showGpu ? state.gpuTime : undefined);
+
+ printMatrixToHtml(resultTable, state.matrices.result, firstRowCount, secondColumnCount);
+ } finally {
+ isComputing = false;
+ }
}
-void run();
-
// #region UI
const firstTable = document.querySelector('.matrix-a') as HTMLDivElement;
const secondTable = document.querySelector('.matrix-b') as HTMLDivElement;
const resultTable = document.querySelector('.matrix-result') as HTMLDivElement;
+const timingDisplay = document.querySelector('.timing-content') as HTMLDivElement;
-function printMatrixToHtml(element: HTMLDivElement, matrix: d.Infer) {
- element.style.gridTemplateColumns = `repeat(${matrix.size.y}, 1fr)`;
- element.innerHTML = matrix.numbers
- .slice(0, matrix.size.x * matrix.size.y)
- .map((x) => `${x}
`)
- .join('');
+generateMatrices();
+
+function updateTimingDisplay(strategy: string, totalTime: number, gpuTime?: number) {
+ if (!timingDisplay) return;
+
+ timingDisplay.innerHTML =
+ gpuTime !== undefined
+ ? `${strategy} computation:
+
+ Total: ${totalTime.toFixed(2)}ms | GPU: ${gpuTime >= 0.01 ? gpuTime.toFixed(2) : '<0.01'}ms
+
`
+ : `${strategy} computation: ${totalTime.toFixed(2)}ms`;
+}
+
+function printMatrixToHtml(element: HTMLDivElement, matrix: number[], rows: number, cols: number) {
+ const isMobile = window.innerWidth <= 480;
+ const isTablet = window.innerWidth <= 768;
+ const maxDisplaySize = isMobile ? 3 : isTablet ? 5 : 8;
+ const shouldTruncate = rows > maxDisplaySize || cols > maxDisplaySize;
+
+ if (shouldTruncate) {
+ const displayRows = Math.min(rows, maxDisplaySize);
+ const displayCols = Math.min(cols, maxDisplaySize);
+ const needsColEllipsis = cols > maxDisplaySize;
+ const needsRowEllipsis = rows > maxDisplaySize;
+ const gridCols = displayCols + (needsColEllipsis ? 1 : 0);
+
+ element.style.gridTemplateColumns = `repeat(${gridCols}, 1fr)`;
+
+ let html = '';
+ for (let row = 0; row < displayRows; row++) {
+ for (let col = 0; col < displayCols; col++) {
+ html += `${matrix[col + row * cols]}
`;
+ }
+ if (needsColEllipsis) html += '...
';
+ }
+
+ if (needsRowEllipsis) {
+ for (let col = 0; col < gridCols; col++) {
+ html += '⋮
';
+ }
+ }
+
+ element.innerHTML = html;
+
+ const sizeInfo = document.createElement('div');
+ sizeInfo.textContent = `${rows}×${cols}`;
+ sizeInfo.style.cssText =
+ 'grid-column: 1 / -1; text-align: center; font-size: 0.8em; color: #666; margin-top: 0.5rem;';
+ element.appendChild(sizeInfo);
+ } else {
+ element.style.gridTemplateColumns = `repeat(${cols}, 1fr)`;
+ element.innerHTML = matrix
+ .slice(0, rows * cols)
+ .map((x) => `${x}
`)
+ .join('');
+ }
}
// #endregion
// #region Example controls & Cleanup
-const paramSettings = {
- min: 1,
- max: MAX_MATRIX_SIZE,
- step: 1,
-};
+const paramSettings = { min: 1, max: 512, step: 1 };
export const controls = defineControls({
- Reshuffle: {
- onButtonClick: () => {
- void run();
+ Reshuffle: { onButtonClick: () => generateMatrices() },
+ Compute: { onButtonClick: () => compute() },
+ strategy: {
+ initial: 'gpu-optimized',
+ options: ['gpu-optimized', 'gpu-simple', 'cpu'],
+ onSelectChange: (value) => {
+ state.strategy = value;
},
},
-
'#1 rows': {
- initial: firstRowCount,
+ initial: state.dimensions.firstRowCount,
...paramSettings,
onSliderChange: (value) => {
- firstRowCount = value;
- void run();
+ state.dimensions.firstRowCount = value;
+ generateMatrices();
},
},
-
'#1 columns': {
- initial: firstColumnCount,
+ initial: state.dimensions.firstColumnCount,
...paramSettings,
onSliderChange: (value) => {
- firstColumnCount = value;
- void run();
+ state.dimensions.firstColumnCount = value;
+ generateMatrices();
},
},
-
'#2 columns': {
- initial: secondColumnCount,
+ initial: state.dimensions.secondColumnCount,
...paramSettings,
onSliderChange: (value) => {
- secondColumnCount = value;
- void run();
+ state.dimensions.secondColumnCount = value;
+ generateMatrices();
},
},
});
diff --git a/apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/meta.json b/apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/meta.json
index 8198b5c710..704b1293a3 100644
--- a/apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/meta.json
+++ b/apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/meta.json
@@ -1,7 +1,7 @@
{
- "title": "Matrix Multiplication",
- "description": "Multiplies configurable matrices in a compute shader and displays the input and result data for comparison.",
+ "title": "Matrix Multiplication (next)",
+ "description": "Explores a newer tiled GPU matrix-multiplication implementation with configurable dimensions and workgroup behavior.",
"category": "algorithms",
"tags": ["linear algebra"],
- "coolFactor": 4
+ "coolFactor": 5
}
diff --git a/apps/typegpu-docs/src/examples/algorithms/matrix-next/params.ts b/apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/params.ts
similarity index 100%
rename from apps/typegpu-docs/src/examples/algorithms/matrix-next/params.ts
rename to apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/params.ts
diff --git a/apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/thumbnail.png b/apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/thumbnail.png
index b38ddb66ba..9e2a49b9a6 100644
Binary files a/apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/thumbnail.png and b/apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/thumbnail.png differ
diff --git a/apps/typegpu-docs/src/examples/algorithms/matrix-next/types.ts b/apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/types.ts
similarity index 100%
rename from apps/typegpu-docs/src/examples/algorithms/matrix-next/types.ts
rename to apps/typegpu-docs/src/examples/algorithms/matrix-multiplication/types.ts
diff --git a/apps/typegpu-docs/src/examples/algorithms/matrix-next/index.html b/apps/typegpu-docs/src/examples/algorithms/matrix-next/index.html
deleted file mode 100644
index 7c1b6f7ebd..0000000000
--- a/apps/typegpu-docs/src/examples/algorithms/matrix-next/index.html
+++ /dev/null
@@ -1,116 +0,0 @@
-Press Compute to time
-
-
-
-
-
-
-
-
Result Matrix
-
-
Press Compute to calculate the result
-
-
-
${strategy} computation:
-
- Total: ${totalTime.toFixed(2)}ms | GPU: ${gpuTime >= 0.01 ? gpuTime.toFixed(2) : '<0.01'}ms
-
`
- : `${strategy} computation: ${totalTime.toFixed(2)}ms`;
-}
-
-function printMatrixToHtml(element: HTMLDivElement, matrix: number[], rows: number, cols: number) {
- const isMobile = window.innerWidth <= 480;
- const isTablet = window.innerWidth <= 768;
- const maxDisplaySize = isMobile ? 3 : isTablet ? 5 : 8;
- const shouldTruncate = rows > maxDisplaySize || cols > maxDisplaySize;
-
- if (shouldTruncate) {
- const displayRows = Math.min(rows, maxDisplaySize);
- const displayCols = Math.min(cols, maxDisplaySize);
- const needsColEllipsis = cols > maxDisplaySize;
- const needsRowEllipsis = rows > maxDisplaySize;
- const gridCols = displayCols + (needsColEllipsis ? 1 : 0);
-
- element.style.gridTemplateColumns = `repeat(${gridCols}, 1fr)`;
-
- let html = '';
- for (let row = 0; row < displayRows; row++) {
- for (let col = 0; col < displayCols; col++) {
- html += `${matrix[col + row * cols]}
`;
- }
- if (needsColEllipsis) html += '...
';
- }
-
- if (needsRowEllipsis) {
- for (let col = 0; col < gridCols; col++) {
- html += '⋮
';
- }
- }
-
- element.innerHTML = html;
-
- const sizeInfo = document.createElement('div');
- sizeInfo.textContent = `${rows}×${cols}`;
- sizeInfo.style.cssText =
- 'grid-column: 1 / -1; text-align: center; font-size: 0.8em; color: #666; margin-top: 0.5rem;';
- element.appendChild(sizeInfo);
- } else {
- element.style.gridTemplateColumns = `repeat(${cols}, 1fr)`;
- element.innerHTML = matrix
- .slice(0, rows * cols)
- .map((x) => `${x}
`)
- .join('');
- }
-}
-
-// #endregion
-
-// #region Example controls & Cleanup
-
-const paramSettings = { min: 1, max: 512, step: 1 };
-
-export const controls = defineControls({
- Reshuffle: { onButtonClick: () => generateMatrices() },
- Compute: { onButtonClick: () => compute() },
- strategy: {
- initial: 'gpu-optimized',
- options: ['gpu-optimized', 'gpu-simple', 'cpu'],
- onSelectChange: (value) => {
- state.strategy = value;
- },
- },
- '#1 rows': {
- initial: state.dimensions.firstRowCount,
- ...paramSettings,
- onSliderChange: (value) => {
- state.dimensions.firstRowCount = value;
- generateMatrices();
- },
- },
- '#1 columns': {
- initial: state.dimensions.firstColumnCount,
- ...paramSettings,
- onSliderChange: (value) => {
- state.dimensions.firstColumnCount = value;
- generateMatrices();
- },
- },
- '#2 columns': {
- initial: state.dimensions.secondColumnCount,
- ...paramSettings,
- onSliderChange: (value) => {
- state.dimensions.secondColumnCount = value;
- generateMatrices();
- },
- },
-});
-
-export function onCleanup() {
- root.destroy();
-}
-
-// #endregion
diff --git a/apps/typegpu-docs/src/examples/algorithms/matrix-next/meta.json b/apps/typegpu-docs/src/examples/algorithms/matrix-next/meta.json
deleted file mode 100644
index 704b1293a3..0000000000
--- a/apps/typegpu-docs/src/examples/algorithms/matrix-next/meta.json
+++ /dev/null
@@ -1,7 +0,0 @@
-{
- "title": "Matrix Multiplication (next)",
- "description": "Explores a newer tiled GPU matrix-multiplication implementation with configurable dimensions and workgroup behavior.",
- "category": "algorithms",
- "tags": ["linear algebra"],
- "coolFactor": 5
-}
diff --git a/apps/typegpu-docs/src/examples/algorithms/matrix-next/thumbnail.png b/apps/typegpu-docs/src/examples/algorithms/matrix-next/thumbnail.png
deleted file mode 100644
index 9e2a49b9a6..0000000000
Binary files a/apps/typegpu-docs/src/examples/algorithms/matrix-next/thumbnail.png and /dev/null differ
diff --git a/apps/typegpu-docs/src/examples/rendering/render-bundles-with/geometry.ts b/apps/typegpu-docs/src/examples/rendering/render-bundles-with/geometry.ts
deleted file mode 100644
index 9feef5c193..0000000000
--- a/apps/typegpu-docs/src/examples/rendering/render-bundles-with/geometry.ts
+++ /dev/null
@@ -1,74 +0,0 @@
-import { d } from 'typegpu';
-import { Vertex } from './schemas.ts';
-
-export function createCubeGeometry(): d.Infer[] {
- const faces: [number[], number[]][] = [
- [
- [0, 0, 1],
- [0, 0, 1],
- ],
- [
- [0, 0, -1],
- [0, 0, -1],
- ],
- [
- [0, 1, 0],
- [0, 1, 0],
- ],
- [
- [0, -1, 0],
- [0, -1, 0],
- ],
- [
- [1, 0, 0],
- [1, 0, 0],
- ],
- [
- [-1, 0, 0],
- [-1, 0, 0],
- ],
- ];
-
- const verts: d.Infer[] = [];
-
- for (const [_offset, normal] of faces) {
- const n = d.vec3f(normal[0], normal[1], normal[2]);
- const absN = [Math.abs(normal[0]), Math.abs(normal[1]), Math.abs(normal[2])];
-
- let tangent: number[];
- let bitangent: number[];
- if (absN[2] === 1) {
- tangent = [1, 0, 0];
- bitangent = [0, 1, 0];
- } else if (absN[1] === 1) {
- tangent = [1, 0, 0];
- bitangent = [0, 0, 1];
- } else {
- tangent = [0, 0, 1];
- bitangent = [0, 1, 0];
- }
-
- const center = normal.map((v) => v * 0.5);
- const corners = [
- [-1, -1],
- [1, -1],
- [1, 1],
- [-1, -1],
- [1, 1],
- [-1, 1],
- ];
-
- for (const [u, v] of corners) {
- verts.push({
- position: d.vec3f(
- center[0] + (tangent[0] * u + bitangent[0] * v) * 0.5,
- center[1] + (tangent[1] * u + bitangent[1] * v) * 0.5,
- center[2] + (tangent[2] * u + bitangent[2] * v) * 0.5,
- ),
- normal: n,
- });
- }
- }
-
- return verts;
-}
diff --git a/apps/typegpu-docs/src/examples/rendering/render-bundles-with/index.html b/apps/typegpu-docs/src/examples/rendering/render-bundles-with/index.html
deleted file mode 100644
index 581d6789f8..0000000000
--- a/apps/typegpu-docs/src/examples/rendering/render-bundles-with/index.html
+++ /dev/null
@@ -1 +0,0 @@
-
diff --git a/apps/typegpu-docs/src/examples/rendering/render-bundles-with/index.ts b/apps/typegpu-docs/src/examples/rendering/render-bundles-with/index.ts
deleted file mode 100644
index baae4dfcc4..0000000000
--- a/apps/typegpu-docs/src/examples/rendering/render-bundles-with/index.ts
+++ /dev/null
@@ -1,219 +0,0 @@
-import { perlin2d } from '@typegpu/noise';
-import { tgpu, d } from 'typegpu';
-import * as m from 'wgpu-matrix';
-import { defineControls } from '../../common/defineControls.ts';
-import { setupOrbitCamera } from '../../common/setup-orbit-camera.ts';
-import { createCubeGeometry } from './geometry.ts';
-import {
- Camera,
- cameraLayout,
- Cube,
- cubeLayout,
- terrainLayout,
- TerrainParams,
- vertexLayout,
-} from './schemas.ts';
-import { fragmentFn, vertexFn } from './shaders.ts';
-
-const CUBE_COUNTS = [1024, 4096, 8192, 16384, 32768, 65536, 131072, 262144, 524288];
-const INITIAL_CUBE_COUNT = CUBE_COUNTS[0];
-const TERRAIN_SIZE = 50;
-const TERRAIN_HEIGHT = 6;
-const NOISE_SCALE = 0.3;
-
-const root = await tgpu.init();
-const canvas = document.querySelector('canvas') as HTMLCanvasElement;
-const context = root.configureContext({ canvas, alphaMode: 'premultiplied' });
-const presentationFormat = navigator.gpu.getPreferredCanvasFormat();
-
-const perlinCache = perlin2d.staticCache({ root, size: d.vec2u(16, 16) });
-
-const cubeVerts = createCubeGeometry();
-const VERTS_PER_CUBE = cubeVerts.length;
-
-const vertexBuffer = root
- .createBuffer(vertexLayout.schemaForCount(VERTS_PER_CUBE), cubeVerts)
- .$usage('vertex');
-
-const cameraBuffer = root.createBuffer(Camera).$usage('uniform');
-
-const terrainBuffer = root
- .createBuffer(TerrainParams, {
- terrainHeight: TERRAIN_HEIGHT,
- noiseScale: NOISE_SCALE,
- })
- .$usage('uniform');
-
-const { cleanupCamera } = setupOrbitCamera(
- canvas,
- {
- initPos: d.vec4f(0, 10, 30, 1),
- target: d.vec4f(0, 0, 0, 1),
- minZoom: 5,
- maxZoom: 100,
- },
- (state) => cameraBuffer.write(state),
-);
-
-const cameraBindGroup = root.createBindGroup(cameraLayout, {
- camera: cameraBuffer,
-});
-
-const terrainBindGroup = root.createBindGroup(terrainLayout, {
- terrain: terrainBuffer,
-});
-
-const pipeline = root.pipe(perlinCache.inject()).createRenderPipeline({
- attribs: vertexLayout.attrib,
- vertex: vertexFn,
- fragment: fragmentFn,
- depthStencil: {
- format: 'depth24plus',
- depthWriteEnabled: true,
- depthCompare: 'less',
- },
-});
-
-let depthTexture = root.device.createTexture({
- size: [canvas.width, canvas.height],
- format: 'depth24plus',
- usage: GPUTextureUsage.RENDER_ATTACHMENT,
-});
-
-let cubeCount = INITIAL_CUBE_COUNT;
-let cubeData: d.Infer[] = [];
-let cubeBuffer = root.createBuffer(d.arrayOf(Cube, INITIAL_CUBE_COUNT)).$usage('storage');
-let cubeBindGroup = root.createBindGroup(cubeLayout, { cubes: cubeBuffer });
-let renderBundle: GPURenderBundle;
-
-let prepared = pipeline
- .with(cameraBindGroup)
- .with(cubeBindGroup)
- .with(terrainBindGroup)
- .with(vertexLayout, vertexBuffer);
-
-function generateCubes(count: number) {
- cubeData = [];
-
- const gridRes = Math.ceil(Math.sqrt(count));
- const cubeScale = TERRAIN_SIZE / gridRes;
- const half = TERRAIN_SIZE / 2;
-
- for (let i = 0; i < count; i++) {
- const gx = i % gridRes;
- const gz = Math.floor(i / gridRes);
-
- const x = gx * cubeScale - half + cubeScale * 0.5;
- const z = gz * cubeScale - half + cubeScale * 0.5;
-
- const model = m.mat4.translation([x, 0, z], d.mat4x4f());
- m.mat4.scale(model, [cubeScale, cubeScale, cubeScale], model);
-
- cubeData.push({ model });
- }
-}
-
-function buildBundle(): GPURenderBundle {
- const bundleEncoder = root.device.createRenderBundleEncoder({
- colorFormats: [presentationFormat],
- depthStencilFormat: 'depth24plus',
- });
-
- const withEncoder = prepared.with(bundleEncoder);
- for (let i = 0; i < cubeCount; i++) {
- withEncoder.draw(VERTS_PER_CUBE, 1, 0, i);
- }
- return bundleEncoder.finish();
-}
-
-function setCubeCount(count: number) {
- cubeCount = count;
- generateCubes(count);
-
- cubeBuffer.destroy();
- cubeBuffer = root.createBuffer(d.arrayOf(Cube, count), cubeData).$usage('storage');
- cubeBindGroup = root.createBindGroup(cubeLayout, { cubes: cubeBuffer });
-
- // Rebuild the prepared pipeline since cubeBindGroup changed.
- prepared = pipeline
- .with(cameraBindGroup)
- .with(cubeBindGroup)
- .with(terrainBindGroup)
- .with(vertexLayout, vertexBuffer);
- renderBundle = buildBundle();
-}
-
-setCubeCount(INITIAL_CUBE_COUNT);
-
-let useBundles = true;
-
-let disposed = false;
-
-function frame() {
- if (disposed) return;
-
- if (depthTexture.width !== canvas.width || depthTexture.height !== canvas.height) {
- depthTexture.destroy();
- depthTexture = root.device.createTexture({
- size: [canvas.width, canvas.height],
- format: 'depth24plus',
- usage: GPUTextureUsage.RENDER_ATTACHMENT,
- });
- }
-
- const encoder = root['~unstable'].createCommandEncoder();
-
- const pass = encoder.beginRenderPass({
- colorAttachments: [
- {
- view: context,
- clearValue: [1, 0.85, 0.74, 1],
- },
- ],
- depthStencilAttachment: {
- view: depthTexture.createView(),
- },
- });
-
- if (useBundles) {
- pass.executeBundles([renderBundle]);
- } else {
- const withPass = prepared.with(pass);
- for (let i = 0; i < cubeCount; i++) {
- withPass.draw(VERTS_PER_CUBE, 1, 0, i);
- }
- }
-
- pass.end();
- encoder.submit();
-
- requestAnimationFrame(frame);
-}
-
-requestAnimationFrame(frame);
-
-// #region Example controls and cleanup
-
-export const controls = defineControls({
- 'cube count': {
- initial: INITIAL_CUBE_COUNT,
- options: CUBE_COUNTS,
- onSelectChange: (value: number) => {
- setCubeCount(value);
- },
- },
- 'use render bundles': {
- initial: true,
- onToggleChange: (value: boolean) => {
- useBundles = value;
- },
- },
-});
-
-export function onCleanup() {
- disposed = true;
- cleanupCamera();
- perlinCache.destroy();
- depthTexture.destroy();
- root.destroy();
-}
diff --git a/apps/typegpu-docs/src/examples/rendering/render-bundles-with/meta.json b/apps/typegpu-docs/src/examples/rendering/render-bundles-with/meta.json
deleted file mode 100644
index a087f3d98c..0000000000
--- a/apps/typegpu-docs/src/examples/rendering/render-bundles-with/meta.json
+++ /dev/null
@@ -1,7 +0,0 @@
-{
- "title": "Render Bundles (.with API)",
- "description": "Uses TypeGPU's fluent `.with` API to record reusable render bundles and efficiently redraw a large 3D scene.",
- "category": "rendering",
- "tags": ["3d", "rasterization", "performance"],
- "coolFactor": 5
-}
diff --git a/apps/typegpu-docs/src/examples/rendering/render-bundles-with/schemas.ts b/apps/typegpu-docs/src/examples/rendering/render-bundles-with/schemas.ts
deleted file mode 100644
index 38932294da..0000000000
--- a/apps/typegpu-docs/src/examples/rendering/render-bundles-with/schemas.ts
+++ /dev/null
@@ -1,31 +0,0 @@
-import { tgpu, d } from 'typegpu';
-
-export const Camera = d.struct({ position: d.vec4f, viewProjection: d.mat4x4f });
-
-export const Vertex = d.struct({
- position: d.vec3f,
- normal: d.vec3f,
-});
-
-export const Cube = d.struct({
- model: d.mat4x4f,
-});
-
-export const TerrainParams = d.struct({
- terrainHeight: d.f32,
- noiseScale: d.f32,
-});
-
-export const vertexLayout = tgpu.vertexLayout(d.arrayOf(Vertex));
-
-export const cameraLayout = tgpu.bindGroupLayout({
- camera: { uniform: Camera },
-});
-
-export const cubeLayout = tgpu.bindGroupLayout({
- cubes: { storage: d.arrayOf(Cube), access: 'readonly' },
-});
-
-export const terrainLayout = tgpu.bindGroupLayout({
- terrain: { uniform: TerrainParams },
-});
diff --git a/apps/typegpu-docs/src/examples/rendering/render-bundles-with/shaders.ts b/apps/typegpu-docs/src/examples/rendering/render-bundles-with/shaders.ts
deleted file mode 100644
index 8b9b22a466..0000000000
--- a/apps/typegpu-docs/src/examples/rendering/render-bundles-with/shaders.ts
+++ /dev/null
@@ -1,91 +0,0 @@
-import { perlin2d } from '@typegpu/noise';
-import { tgpu, d, std } from 'typegpu';
-import { cameraLayout, cubeLayout, terrainLayout, Vertex } from './schemas.ts';
-
-const LOW_COLOR = d.vec3f(0.28, 0.52, 0.3);
-const MID_COLOR = d.vec3f(0.76, 0.7, 0.5);
-const HIGH_COLOR = d.vec3f(0.95, 0.95, 0.97);
-
-const heightColor = tgpu.fn(
- [d.f32],
- d.vec3f,
-)((t) => {
- const clamped = std.clamp(t, 0, 1);
-
- const lowToMid = std.mix(LOW_COLOR, MID_COLOR, std.smoothstep(0, 0.4, clamped));
- const midToHigh = std.mix(lowToMid, HIGH_COLOR, std.smoothstep(0.4, 1, clamped));
-
- return midToHigh;
-});
-
-const fbm = tgpu.fn(
- [d.vec2f],
- d.f32,
-)((p) => {
- let value = perlin2d.sample(p) * 0.5;
- value += perlin2d.sample(p.mul(2)) * 0.25;
- value += perlin2d.sample(p.mul(4)) * 0.125;
- value += perlin2d.sample(p.mul(8)) * 0.0625;
- return value;
-});
-
-export const vertexFn = tgpu.vertexFn({
- in: {
- ...Vertex.propTypes,
- instanceIndex: d.builtin.instanceIndex,
- },
- out: {
- pos: d.builtin.position,
- worldNormal: d.vec3f,
- color: d.vec3f,
- worldPos: d.vec3f,
- },
-})((input) => {
- 'use gpu';
- const cube = cubeLayout.$.cubes[input.instanceIndex];
- const cam = cameraLayout.$.camera;
- const params = terrainLayout.$.terrain;
-
- const worldPos = cube.model * d.vec4f(input.position, 1);
-
- const cubeCenter = cube.model.columns[3];
- const noiseCoord = cubeCenter.xz * params.noiseScale;
- const height = fbm(noiseCoord) * params.terrainHeight;
-
- const displaced = d.vec4f(worldPos.x, worldPos.y + height, worldPos.z, 1);
-
- const worldNormal = std.normalize((cube.model * d.vec4f(input.normal, 0)).xyz);
-
- const t = height / params.terrainHeight + 0.5;
- const color = heightColor(t);
-
- return {
- pos: cam.viewProjection * displaced,
- worldNormal,
- color,
- worldPos: displaced.xyz,
- };
-});
-
-export const fragmentFn = tgpu.fragmentFn({
- in: {
- worldNormal: d.vec3f,
- color: d.vec3f,
- worldPos: d.vec3f,
- },
- out: d.vec4f,
-})((input) => {
- 'use gpu';
- const cam = cameraLayout.$.camera;
- const lightDir = std.normalize(d.vec3f(0.4, 1.0, 0.3));
-
- const diffuse = std.max(0, std.dot(input.worldNormal, lightDir));
- const ambient = 0.55;
-
- const viewDir = std.normalize(cam.position.xyz - input.worldPos);
- const halfDir = std.normalize(lightDir + viewDir);
- const specular = std.pow(std.max(0, std.dot(input.worldNormal, halfDir)), 32);
-
- const lit = input.color * (ambient + diffuse * 0.85) + specular * 0.3;
- return d.vec4f(lit, 1);
-});
diff --git a/apps/typegpu-docs/src/examples/rendering/render-bundles-with/thumbnail.png b/apps/typegpu-docs/src/examples/rendering/render-bundles-with/thumbnail.png
deleted file mode 100644
index b80128f505..0000000000
Binary files a/apps/typegpu-docs/src/examples/rendering/render-bundles-with/thumbnail.png and /dev/null differ
diff --git a/apps/typegpu-docs/tests/individual-example-tests/matrix-multiplication.test.ts b/apps/typegpu-docs/tests/individual-example-tests/matrix-multiplication.test.ts
index 93b2325816..fd5ba40cfe 100644
--- a/apps/typegpu-docs/tests/individual-example-tests/matrix-multiplication.test.ts
+++ b/apps/typegpu-docs/tests/individual-example-tests/matrix-multiplication.test.ts
@@ -6,50 +6,84 @@ import { describe, expect } from 'vitest';
import { it } from 'typegpu-testing-utility';
import { runExampleTest, setupCommonMocks } from './utils/baseTest.ts';
-describe('matrix multiplication example', () => {
+describe('matrix(next) example', () => {
setupCommonMocks();
it('should produce valid code', async ({ device }) => {
const shaderCodes = await runExampleTest(
{
category: 'algorithms',
- name: 'matrix-multiplication',
+ name: 'matrix-next',
+ controlTriggers: ['Compute'],
expectedCalls: 1,
},
device,
);
expect(shaderCodes).toMatchInlineSnapshot(`
- "struct MatrixStruct {
- size: vec2f,
- numbers: array,
+ "struct MatrixInfo {
+ firstRowCount: u32,
+ firstColumnCount: u32,
+ secondColumnCount: u32,
}
- @group(0) @binding(0) var firstMatrix: MatrixStruct;
- @group(0) @binding(1) var secondMatrix: MatrixStruct;
- @group(0) @binding(2) var resultMatrix: MatrixStruct;
+ @group(0) @binding(3) var dimensions_1: MatrixInfo;
- @compute @workgroup_size(8, 8)
- fn main(@builtin(global_invocation_id) global_id: vec3u) {
- if (global_id.x >= u32(firstMatrix.size.x) || global_id.y >= u32(secondMatrix.size.y)) {
- return;
- }
+ fn getTileIndex(row: u32, col: u32) -> u32 {
+ return (col + (row * 16u));
+ }
- if (global_id.x + global_id.y == 0u) {
- resultMatrix.size = vec2(firstMatrix.size.x, secondMatrix.size.y);
- }
+ fn getIndex(row: u32, col: u32, columns: u32) -> u32 {
+ return (col + (row * columns));
+ }
- let resultCell = vec2(global_id.x, global_id.y);
- var result = 0.0;
+ @group(0) @binding(0) var firstMatrix: array;
- for (var i = 0u; i < u32(firstMatrix.size.y); i = i + 1u) {
- let a = i + resultCell.x * u32(firstMatrix.size.y);
- let b = resultCell.y + i * u32(secondMatrix.size.y);
- result = result + firstMatrix.numbers[a] * secondMatrix.numbers[b];
- }
+ var tileA: array;
- let index = resultCell.y + resultCell.x * u32(secondMatrix.size.y);
- resultMatrix.numbers[index] = result;
+ @group(0) @binding(1) var secondMatrix: array;
+
+ var tileB: array;
+
+ @group(0) @binding(2) var resultMatrix: array;
+
+ @compute @workgroup_size(16, 16) fn computeSharedMemory(@builtin(local_invocation_id) lid: vec3u, @builtin(workgroup_id) wid: vec3u) {
+ let dimensions = (&dimensions_1);
+ let numTiles = u32((f32((((*dimensions).firstColumnCount + 16u) - 1u)) / 16f));
+ let globalRow = ((wid.x * 16u) + lid.x);
+ let globalCol = ((wid.y * 16u) + lid.y);
+ let localRow = lid.x;
+ let localCol = lid.y;
+ let tileIdx = getTileIndex(localRow, localCol);
+ var accumulatedResult = 0;
+ for (var tileIndex = 0u; (tileIndex < numTiles); tileIndex++) {
+ let matrixACol = ((tileIndex * 16u) + localCol);
+ var valueA = 0;
+ if (((globalRow < (*dimensions).firstRowCount) && (matrixACol < (*dimensions).firstColumnCount))) {
+ let indexA = getIndex(globalRow, matrixACol, (*dimensions).firstColumnCount);
+ valueA = firstMatrix[indexA];
+ }
+ tileA[tileIdx] = valueA;
+ let matrixBRow = ((tileIndex * 16u) + localRow);
+ var valueB = 0;
+ if (((matrixBRow < (*dimensions).firstColumnCount) && (globalCol < (*dimensions).secondColumnCount))) {
+ let indexB = getIndex(matrixBRow, globalCol, (*dimensions).secondColumnCount);
+ valueB = secondMatrix[indexB];
+ }
+ tileB[tileIdx] = valueB;
+ workgroupBarrier();
+ let effectiveTileSize = min(16u, ((*dimensions).firstColumnCount - (tileIndex * 16u)));
+ for (var k = 0u; (k < effectiveTileSize); k++) {
+ let tileA_element = tileA[getTileIndex(localRow, k)];
+ let tileB_element = tileB[getTileIndex(k, localCol)];
+ accumulatedResult += (tileA_element * tileB_element);
+ }
+ workgroupBarrier();
+ }
+ if (((globalRow < (*dimensions).firstRowCount) && (globalCol < (*dimensions).secondColumnCount))) {
+ let outputIndex = getIndex(globalRow, globalCol, (*dimensions).secondColumnCount);
+ resultMatrix[outputIndex] = accumulatedResult;
+ }
}"
`);
});
diff --git a/apps/typegpu-docs/tests/individual-example-tests/matrix-next.test.ts b/apps/typegpu-docs/tests/individual-example-tests/matrix-next.test.ts
deleted file mode 100644
index fd5ba40cfe..0000000000
--- a/apps/typegpu-docs/tests/individual-example-tests/matrix-next.test.ts
+++ /dev/null
@@ -1,90 +0,0 @@
-/**
- * @vitest-environment jsdom
- */
-
-import { describe, expect } from 'vitest';
-import { it } from 'typegpu-testing-utility';
-import { runExampleTest, setupCommonMocks } from './utils/baseTest.ts';
-
-describe('matrix(next) example', () => {
- setupCommonMocks();
-
- it('should produce valid code', async ({ device }) => {
- const shaderCodes = await runExampleTest(
- {
- category: 'algorithms',
- name: 'matrix-next',
- controlTriggers: ['Compute'],
- expectedCalls: 1,
- },
- device,
- );
-
- expect(shaderCodes).toMatchInlineSnapshot(`
- "struct MatrixInfo {
- firstRowCount: u32,
- firstColumnCount: u32,
- secondColumnCount: u32,
- }
-
- @group(0) @binding(3) var dimensions_1: MatrixInfo;
-
- fn getTileIndex(row: u32, col: u32) -> u32 {
- return (col + (row * 16u));
- }
-
- fn getIndex(row: u32, col: u32, columns: u32) -> u32 {
- return (col + (row * columns));
- }
-
- @group(0) @binding(0) var firstMatrix: array;
-
- var tileA: array;
-
- @group(0) @binding(1) var secondMatrix: array;
-
- var tileB: array;
-
- @group(0) @binding(2) var resultMatrix: array;
-
- @compute @workgroup_size(16, 16) fn computeSharedMemory(@builtin(local_invocation_id) lid: vec3u, @builtin(workgroup_id) wid: vec3u) {
- let dimensions = (&dimensions_1);
- let numTiles = u32((f32((((*dimensions).firstColumnCount + 16u) - 1u)) / 16f));
- let globalRow = ((wid.x * 16u) + lid.x);
- let globalCol = ((wid.y * 16u) + lid.y);
- let localRow = lid.x;
- let localCol = lid.y;
- let tileIdx = getTileIndex(localRow, localCol);
- var accumulatedResult = 0;
- for (var tileIndex = 0u; (tileIndex < numTiles); tileIndex++) {
- let matrixACol = ((tileIndex * 16u) + localCol);
- var valueA = 0;
- if (((globalRow < (*dimensions).firstRowCount) && (matrixACol < (*dimensions).firstColumnCount))) {
- let indexA = getIndex(globalRow, matrixACol, (*dimensions).firstColumnCount);
- valueA = firstMatrix[indexA];
- }
- tileA[tileIdx] = valueA;
- let matrixBRow = ((tileIndex * 16u) + localRow);
- var valueB = 0;
- if (((matrixBRow < (*dimensions).firstColumnCount) && (globalCol < (*dimensions).secondColumnCount))) {
- let indexB = getIndex(matrixBRow, globalCol, (*dimensions).secondColumnCount);
- valueB = secondMatrix[indexB];
- }
- tileB[tileIdx] = valueB;
- workgroupBarrier();
- let effectiveTileSize = min(16u, ((*dimensions).firstColumnCount - (tileIndex * 16u)));
- for (var k = 0u; (k < effectiveTileSize); k++) {
- let tileA_element = tileA[getTileIndex(localRow, k)];
- let tileB_element = tileB[getTileIndex(k, localCol)];
- accumulatedResult += (tileA_element * tileB_element);
- }
- workgroupBarrier();
- }
- if (((globalRow < (*dimensions).firstRowCount) && (globalCol < (*dimensions).secondColumnCount))) {
- let outputIndex = getIndex(globalRow, globalCol, (*dimensions).secondColumnCount);
- resultMatrix[outputIndex] = accumulatedResult;
- }
- }"
- `);
- });
-});