diff --git a/apps/typegpu-docs/astro.config.mjs b/apps/typegpu-docs/astro.config.mjs index 209bd9df6..21e5b6d72 100644 --- a/apps/typegpu-docs/astro.config.mjs +++ b/apps/typegpu-docs/astro.config.mjs @@ -128,6 +128,11 @@ export default defineConfig({ slug: 'fundamentals/resolve', badge: { text: '0.3' }, }, + DEV && { + label: 'Vertex Layouts', + slug: 'fundamentals/vertex-layouts', + badge: { text: '0.3.3' }, + }, DEV && { label: 'Slots', slug: 'fundamentals/slots', diff --git a/apps/typegpu-docs/src/content/docs/fundamentals/roots.mdx b/apps/typegpu-docs/src/content/docs/fundamentals/roots.mdx index b14da3654..31b7d74b2 100644 --- a/apps/typegpu-docs/src/content/docs/fundamentals/roots.mdx +++ b/apps/typegpu-docs/src/content/docs/fundamentals/roots.mdx @@ -55,9 +55,13 @@ untyped value of a typed resource, use the `root.unwrap` function. | `root.unwrap(resource: TgpuBuffer)` | Returns a `GPUBuffer`. | | `root.unwrap(resource: TgpuBindGroupLayout)` | Returns a `GPUBindGroupLayout`. | | `root.unwrap(resource: TgpuBindGroup)` | Returns a `GPUBindGroup`. | +| `root.unwrap(resource: TgpuVertexLayout)` | Returns a `GPUVertexBufferLayout`. | {/* | `root.unwrap(resource: TgpuTexture)` | Returns a `GPUTexture`. | */} {/* | `root.unwrap(resource: TgpuReadonlyTexture \| TgpuWriteonlyTexture \| TgpuMutableTexture \| TgpuSampledTexture)` | Returns a `GPUTextureView`. | */} +:::note +To unwrap a `TgpuVertexLayout` make sure that all its attributes are marked with the appropriate location. +::: ## Destroying resources @@ -86,7 +90,7 @@ import React from 'react'; function SceneView() { const ref = useWebGPU(({ context, device, presentationFormat }) => { const root = tgpu.initFromDevice({ device }); - + // create all resources... }); @@ -135,4 +139,4 @@ class GameObject { // create all resources... } } -``` \ No newline at end of file +``` diff --git a/apps/typegpu-docs/src/content/docs/fundamentals/vertex-layouts.mdx b/apps/typegpu-docs/src/content/docs/fundamentals/vertex-layouts.mdx new file mode 100644 index 000000000..799d58511 --- /dev/null +++ b/apps/typegpu-docs/src/content/docs/fundamentals/vertex-layouts.mdx @@ -0,0 +1,189 @@ +--- +title: Vertex Layouts +description: A guide on how to create and use typed vertex layouts +draft: true +--- + +Typed vertex layouts are a way to describe the structure of vertex data in a typed manner. They are used to create a single source of truth for vertex data, which can be used in combination with [`tgpu.resolve`](/TypeGPU/fundamentals/resolve) and [`root.unwrap`](/TypeGPU/fundamentals/roots) to easily create and manage vertex buffers. + +## Creating a vertex layout + +To create a vertex layout, use the `tgpu.vertexLayout` function. It takes a function that takes a number and returns an array containing any data type, and a string that describes the type of the vertex layout. It can be either `vertex` or `instance` (default is `vertex`). + +```ts +import tgpu from 'typegpu'; +import * as d from 'typegpu/data'; + +const ParticleGeometry = d.struct({ + tilt: d.f32, + angle: d.f32, + color: d.vec4f, +}); + +const geometryLayout = tgpu + .vertexLayout((n: number) => d.arrayOf(ParticleGeometry, n), 'instance'); +``` + +## Utilizing loose schemas with vertex layouts + +The above example is great if the vertex buffer will also be used as a storage or uniform buffer. However, if you're only interested in the vertex data, you can use a loose schema instead. +It is not restricted by alignment rules and can utilize many useful [vertex formats](https://www.w3.org/TR/webgpu/#vertex-formats). To create loose schemas, use `d.unstruct` instead of `d.struct` and `d.disarrayOf` instead of `d.arrayOf`. +Inside of loose schemas you can use vertex formats as well as the usual data types. + +```ts +const LooseParticleGeometry = d.unstruct({ + tilt: d.f32, + angle: d.f32, + color: d.unorm8x4, // 4x8-bit unsigned normalized +}); +``` + +The size of `LooseParticleGeometry` will be 12 bytes, compared to 32 bytes of `ParticleGeometry`. This can be useful when you're working with large amounts of vertex data and want to save memory. + +:::tip[Aligning loose schemas] +Sometimes you might want to align the data in a loose schema due to external requirements or performance reasons. +You can do this by using the `d.align` function, just like in normal schemas. Even though loose schemas don't have alignment requirements, they will still respect any alignment you specify. + +```ts +const LooseParticleGeometry = d.unstruct({ + tilt: d.f32, + angle: d.f32, + color: d.align(16, d.unorm8x4), + // 4x8-bit unsigned normalized aligned to 16 bytes +}); +``` + +This will align the `color` field to 16 bytes, making the size of `LooseParticleGeometry` 20 bytes. +::: + +## Using vertex layouts + +You can utilize [`root.unwrap`](/TypeGPU/fundamentals/roots) to get the raw `GPUVertexBufferLayout` from a typed vertex layout. It will automatically calculate the stride and attributes for you, according to the vertex layout you provided. + +:::caution +Make sure that all attributes in the vertex layout are marked with the appropriate location. You can use the `d.location` function to specify the location of each attribute. +If you don't do this, the unwrapping will fail at runtime. +::: + +```ts +const ParticleGeometry = d.struct({ + tilt: d.location(0, d.f32), + angle: d.location(1, d.f32), + color: d.location(2, d.vec4f), +}); + +const geometryLayout = tgpu + .vertexLayout((n: number) => d.arrayOf(ParticleGeometry, n), 'instance'); + +const geometry = root.unwrap(geometryLayout); + +console.log(geometry); +//{ +// "arrayStride": 32, +// "stepMode": "instance", +// "attributes": [ +// { +// "format": "float32", +// "offset": 0, +// "shaderLocation": 0 +// }, +// { +// "format": "float32", +// "offset": 4, +// "shaderLocation": 1 +// }, +// { +// "format": "float32x4", +// "offset": 16, +// "shaderLocation": 2 +// } +// ] +//} +``` + +This will return a `GPUVertexBufferLayout` that can be used when creating a render pipeline. + +```diff lang=ts +const renderPipeline = device.createRenderPipeline({ + layout: device.createPipelineLayout({ + bindGroupLayouts: [root.unwrap(bindGroupLayout)], + }), + primitive: { + topology: 'triangle-strip', + }, + vertex: { + module: renderShader, +- buffers: [ +- { +- arrayStride: 32, +- stepMode: 'instance', +- attributes: [ +- { +- format: 'float32', +- offset: 0, +- shaderLocation: 0, +- }, +- { +- format: 'float32', +- offset: 4, +- shaderLocation: 1, +- }, +- { +- format: 'float32x4', +- offset: 16, +- shaderLocation: 2, +- }, +- ], +- }, +- ], ++ buffers: [root.unwrap(geometryLayout)], + }, + fragment: { + ... + }, +}); +``` + +If you are using a loose schema, you can now resolve it to get its WGSL representation. + +```ts +const LooseParticleGeometry = d.unstruct({ + tilt: d.location(0, d.f32), + angle: d.location(1, d.f32), + color: d.location(2, d.unorm8x4), +}); + +const sampleShader = ` + @vertex + fn main(particleGeometry: LooseParticleGeometry) -> @builtin(position) pos: vec4f { + return vec4f( + particleGeometry.tilt, + particleGeometry.angle, + particleGeometry.color.rgb, + 1.0 + ); + } +`; + +const wgslDefinition = tgpu.resolve({ + template: sampleShader, + externals: { LooseParticleGeometry } +}); + +console.log(wgslDefinition); +// struct LooseParticleGeometry_0 { +// @location(0) tilt: f32, +// @location(1) angle: f32, +// @location(2) color: vec4f, +// } +// +// @vertex +// fn main(particleGeometry: LooseParticleGeometry_0) -> @builtin(position) pos: vec4f { +// return vec4f( +// particleGeometry.tilt, +// particleGeometry.angle, +// particleGeometry.color.rgb, +// 1.0 +// ); +// } +``` diff --git a/apps/typegpu-docs/src/content/examples/simulation/boids/index.ts b/apps/typegpu-docs/src/content/examples/simulation/boids/index.ts index 3a94cddaf..1efda59f6 100644 --- a/apps/typegpu-docs/src/content/examples/simulation/boids/index.ts +++ b/apps/typegpu-docs/src/content/examples/simulation/boids/index.ts @@ -4,7 +4,7 @@ import * as d from 'typegpu/data'; const triangleAmount = 1000; const triangleSize = 0.03; -const renderCode = /* wgsl */ ` +const utilityFunctions = /* wgsl */ ` fn rotate(v: vec2f, angle: f32) -> vec2f { let pos = vec2( (v.x * cos(angle)) - (v.y * sin(angle)), @@ -16,20 +16,14 @@ const renderCode = /* wgsl */ ` fn getRotationFromVelocity(velocity: vec2f) -> f32 { return -atan2(velocity.x, velocity.y); }; +`; - struct TriangleData { - position : vec2f, - velocity : vec2f, - }; - +const renderCode = /* wgsl */ ` struct VertexOutput { @builtin(position) position : vec4f, @location(1) color : vec4f, }; - @binding(0) @group(0) var trianglePos : array; - @binding(1) @group(0) var colorPalette : vec3f; - @vertex fn mainVert(@builtin(instance_index) ii: u32, @location(0) v: vec2f) -> VertexOutput { let instanceInfo = trianglePos[ii]; @@ -56,24 +50,6 @@ const renderCode = /* wgsl */ ` `; const computeCode = /* wgsl */ ` - struct TriangleData { - position : vec2f, - velocity : vec2f, - }; - - struct Parameters { - separation_distance : f32, - separation_strength : f32, - alignment_distance : f32, - alignment_strength : f32, - cohesion_distance : f32, - cohesion_strength : f32, - }; - - @binding(0) @group(0) var currentTrianglePos : array; - @binding(1) @group(0) var nextTrianglePos : array; - @binding(2) @group(0) var params : Parameters; - @compute @workgroup_size(1) fn mainCompute(@builtin(global_invocation_id) gid: vec3u) { let index = gid.x; @@ -89,14 +65,14 @@ const computeCode = /* wgsl */ ` } var other = currentTrianglePos[i]; var dist = distance(instanceInfo.position, other.position); - if (dist < params.separation_distance) { + if (dist < params.separationDistance) { separation += instanceInfo.position - other.position; } - if (dist < params.alignment_distance) { + if (dist < params.alignmentDistance) { alignment += other.velocity; alignmentCount++; } - if (dist < params.cohesion_distance) { + if (dist < params.cohesionDistance) { cohesion += other.position; cohesionCount++; } @@ -108,9 +84,9 @@ const computeCode = /* wgsl */ ` cohesion = (cohesion / f32(cohesionCount)) - instanceInfo.position; } instanceInfo.velocity += - (separation * params.separation_strength) - + (alignment * params.alignment_strength) - + (cohesion * params.cohesion_strength); + (separation * params.separationStrength) + + (alignment * params.alignmentStrength) + + (cohesion * params.cohesionStrength); instanceInfo.velocity = normalize(instanceInfo.velocity) * clamp(length(instanceInfo.velocity), 0.0, 0.01); let triangleSize = ${triangleSize}; if (instanceInfo.position[0] > 1.0 + triangleSize) { @@ -223,13 +199,10 @@ const paramsBuffer = root .$usage('storage'); const triangleVertexBuffer = root - .createBuffer(d.arrayOf(d.f32, 6), [ - 0.0, - triangleSize, - -triangleSize / 2, - -triangleSize / 2, - triangleSize / 2, - -triangleSize / 2, + .createBuffer(d.arrayOf(d.vec2f, 3), [ + d.vec2f(0.0, triangleSize), + d.vec2f(-triangleSize / 2, -triangleSize / 2), + d.vec2f(triangleSize / 2, -triangleSize / 2), ]) .$usage('vertex'); @@ -266,18 +239,43 @@ const updateParams = (newOptions: BoidsOptions) => { paramsBuffer.write(newOptions); }; +const renderBindGroupLayout = tgpu.bindGroupLayout({ + trianglePos: { uniform: d.arrayOf(TriangleInfoStruct, triangleAmount) }, + colorPalette: { uniform: d.vec3f }, +}); + +const computeBindGroupLayout = tgpu.bindGroupLayout({ + currentTrianglePos: { + uniform: d.arrayOf(TriangleInfoStruct, triangleAmount), + }, + nextTrianglePos: { + storage: d.arrayOf(TriangleInfoStruct, triangleAmount), + access: 'mutable', + }, + params: { storage: Params }, +}); + const renderModule = root.device.createShaderModule({ - code: renderCode, + code: tgpu.resolve({ + template: renderCode.concat(utilityFunctions), + externals: { + ...renderBindGroupLayout.bound, + }, + }), }); const computeModule = root.device.createShaderModule({ - code: computeCode, + code: tgpu.resolve({ + template: computeCode, + externals: { + ...computeBindGroupLayout.bound, + }, + }), }); -const renderBindGroupLayout = tgpu.bindGroupLayout({ - trianglePos: { uniform: d.arrayOf(TriangleInfoStruct, triangleAmount) }, - colorPalette: { uniform: d.vec3f }, -}); +const vertexLayout = tgpu['~unstable'].vertexLayout((n) => + d.arrayOf(d.location(0, d.vec2f), n), +); const pipeline = root.device.createRenderPipeline({ layout: root.device.createPipelineLayout({ @@ -285,18 +283,7 @@ const pipeline = root.device.createRenderPipeline({ }), vertex: { module: renderModule, - buffers: [ - { - arrayStride: 2 * 4, - attributes: [ - { - shaderLocation: 0, - offset: 0, - format: 'float32x2' as const, - }, - ], - }, - ], + buffers: [root.unwrap(vertexLayout)], }, fragment: { module: renderModule, @@ -311,17 +298,6 @@ const pipeline = root.device.createRenderPipeline({ }, }); -const computeBindGroupLayout = tgpu.bindGroupLayout({ - currentTrianglePos: { - uniform: d.arrayOf(TriangleInfoStruct, triangleAmount), - }, - nextTrianglePos: { - storage: d.arrayOf(TriangleInfoStruct, triangleAmount), - access: 'mutable', - }, - params: { storage: Params }, -}); - const computePipeline = root.device.createComputePipeline({ layout: root.device.createPipelineLayout({ bindGroupLayouts: [root.unwrap(computeBindGroupLayout)], @@ -410,7 +386,7 @@ export const controls = { onButtonClick: () => updateParams(presets.blobs), }, - '⚛️ Particles': { + '⚛ Particles': { onButtonClick: () => updateParams(presets.particles), }, diff --git a/apps/typegpu-docs/src/content/examples/simulation/game-of-life/index.ts b/apps/typegpu-docs/src/content/examples/simulation/game-of-life/index.ts index 6e00e9dd2..fad943f91 100644 --- a/apps/typegpu-docs/src/content/examples/simulation/game-of-life/index.ts +++ b/apps/typegpu-docs/src/content/examples/simulation/game-of-life/index.ts @@ -23,12 +23,29 @@ let timestep = 4; let swap = false; let paused = false; -const computeShader = device.createShaderModule({ - code: ` -@binding(0) @group(0) var size: vec2u; -@binding(1) @group(0) var current: array; -@binding(2) @group(0) var next: array; +const bindGroupLayoutCompute = tgpu.bindGroupLayout({ + size: { + storage: d.vec2u, + access: 'readonly', + }, + current: { + storage: (arrayLength: number) => d.arrayOf(d.u32, arrayLength), + access: 'readonly', + }, + next: { + storage: (arrayLength: number) => d.arrayOf(d.u32, arrayLength), + access: 'mutable', + }, +}); +const bindGroupLayoutRender = tgpu.bindGroupLayout({ + size: { + uniform: d.vec2u, + }, +}); +const computeShader = device.createShaderModule({ + code: tgpu.resolve({ + template: ` override blockSize = 8; fn getIndex(x: u32, y: u32) -> u32 { @@ -43,8 +60,8 @@ fn getCell(x: u32, y: u32) -> u32 { } fn countNeighbors(x: u32, y: u32) -> u32 { - return getCell(x - 1, y - 1) + getCell(x, y - 1) + getCell(x + 1, y - 1) + - getCell(x - 1, y) + getCell(x + 1, y) + + return getCell(x - 1, y - 1) + getCell(x, y - 1) + getCell(x + 1, y - 1) + + getCell(x - 1, y) + getCell(x + 1, y) + getCell(x - 1, y + 1) + getCell(x, y + 1) + getCell(x + 1, y + 1); } @@ -53,38 +70,39 @@ fn main(@builtin(global_invocation_id) grid: vec3u) { let x = grid.x; let y = grid.y; let n = countNeighbors(x, y); - next[getIndex(x, y)] = select(u32(n == 3u), u32(n == 2u || n == 3u), getCell(x, y) == 1u); -} + next[getIndex(x, y)] = select(u32(n == 3u), u32(n == 2u || n == 3u), getCell(x, y) == 1u); +} `, + externals: { + ...bindGroupLayoutCompute.bound, + }, + }), }); const squareBuffer = root .createBuffer(d.arrayOf(d.u32, 8), [0, 0, 1, 0, 0, 1, 1, 1]) .$usage('vertex'); -const squareStride: GPUVertexBufferLayout = { - arrayStride: 2 * Uint32Array.BYTES_PER_ELEMENT, - stepMode: 'vertex', - attributes: [ - { - shaderLocation: 1, - offset: 0, - format: 'uint32x2', - }, - ], -}; +const squareVertexLayout = tgpu['~unstable'].vertexLayout( + (n: number) => d.arrayOf(d.location(1, d.vec2u), n), + 'vertex', +); -const vertexShader = device.createShaderModule({ - code: ` +const cellsVertexLayout = tgpu['~unstable'].vertexLayout( + (n: number) => d.arrayOf(d.location(0, d.u32), n), + 'instance', +); + +const renderShader = device.createShaderModule({ + code: tgpu.resolve({ + template: ` struct Out { @builtin(position) pos: vec4f, @location(0) cell: f32, } -@binding(0) @group(0) var size: vec2u; - @vertex -fn main(@builtin(instance_index) i: u32, @location(0) cell: u32, @location(1) pos: vec2u) -> Out { +fn vert(@builtin(instance_index) i: u32, @location(0) cell: u32, @location(1) pos: vec2u) -> Out { let w = size.x; let h = size.y; let x = (f32(i % w + pos.x) / f32(w) - 0.5) * 2. * f32(w) / f32(max(w, h)); @@ -94,61 +112,28 @@ fn main(@builtin(instance_index) i: u32, @location(0) cell: u32, @location(1) po vec4f(x, y, 0., 1.), f32(cell), ); -}`, -}); -const fragmentShader = device.createShaderModule({ - code: ` +} + @fragment -fn main(@location(0) cell: f32, @builtin(position) pos: vec4f) -> @location(0) vec4f { +fn frag(@location(0) cell: f32, @builtin(position) pos: vec4f) -> @location(0) vec4f { if (cell == 0.) { discard; } - + return vec4f( - max(f32(cell) * pos.x / 1024, 0), - max(f32(cell) * pos.y / 1024, 0), + max(f32(cell) * pos.x / 1024, 0), + max(f32(cell) * pos.y / 1024, 0), max(f32(cell) * (1 - pos.x / 1024), 0), 1. ); }`, + externals: { + ...bindGroupLayoutRender.bound, + }, + }), }); let commandEncoder: GPUCommandEncoder; -const cellsStride: GPUVertexBufferLayout = { - arrayStride: Uint32Array.BYTES_PER_ELEMENT, - stepMode: 'instance', - attributes: [ - { - shaderLocation: 0, - offset: 0, - format: 'uint32', - }, - ], -}; - -const layoutCompute = { - size: { - storage: d.vec2u, - access: 'readonly', - }, - current: { - storage: (arrayLength: number) => d.arrayOf(d.u32, arrayLength), - access: 'readonly', - }, - next: { - storage: (arrayLength: number) => d.arrayOf(d.u32, arrayLength), - access: 'mutable', - }, -} as const; -const groupLayout = { - size: { - uniform: d.vec2u, - }, -} as const; - -const bindGroupLayoutCompute = tgpu.bindGroupLayout(layoutCompute); -const bindGroupLayoutRender = tgpu.bindGroupLayout(groupLayout); - // compute pipeline const computePipeline = device.createComputePipeline({ layout: device.createPipelineLayout({ @@ -271,11 +256,11 @@ const renderPipeline = device.createRenderPipeline({ topology: 'triangle-strip', }, vertex: { - module: vertexShader, - buffers: [cellsStride, squareStride], + module: renderShader, + buffers: [root.unwrap(cellsVertexLayout), root.unwrap(squareVertexLayout)], }, fragment: { - module: fragmentShader, + module: renderShader, targets: [ { format: presentationFormat, diff --git a/packages/typegpu/src/core/resolve/externals.ts b/packages/typegpu/src/core/resolve/externals.ts index 2d403d899..2bfab415c 100644 --- a/packages/typegpu/src/core/resolve/externals.ts +++ b/packages/typegpu/src/core/resolve/externals.ts @@ -1,3 +1,4 @@ +import { isLooseData } from '../../data/dataTypes'; import { isWgslStruct } from '../../data/wgslTypes'; import { isNamable } from '../../namable'; import { type ResolutionCtx, isWgsl } from '../../types'; @@ -87,7 +88,7 @@ export function replaceExternalsInWgsl( wgsl: string, ): string { return Object.entries(externalMap).reduce((acc, [externalName, external]) => { - if (isWgsl(external)) { + if (isWgsl(external) || isLooseData(external)) { return acc.replaceAll( identifierRegex(externalName), ctx.resolve(external), diff --git a/packages/typegpu/src/core/resolve/resolveData.ts b/packages/typegpu/src/core/resolve/resolveData.ts index 86d942bd7..8e5d7bc2b 100644 --- a/packages/typegpu/src/core/resolve/resolveData.ts +++ b/packages/typegpu/src/core/resolve/resolveData.ts @@ -1,4 +1,11 @@ import { getAttributesString } from '../../data/attributes'; +import { + type AnyData, + type Disarray, + type Unstruct, + isLooseData, +} from '../../data/dataTypes'; +import { formatToWGSLType } from '../../data/vertexFormatData'; import type { AnyWgslData, BaseWgslData, @@ -27,6 +34,7 @@ import type { } from '../../data/wgslTypes'; import { assertExhaustive } from '../../shared/utilityTypes'; import type { ResolutionCtx } from '../../types'; +import { isAttribute } from '../vertexLayout/connectAttributesToShader'; /** * Schemas for which their `type` property directly @@ -116,6 +124,38 @@ ${Object.entries(struct.propTypes) return id; } +/** + * Resolves an unstruct (struct that does not align data by default) to its struct data counterpart. + * @param ctx - The resolution context. + * @param unstruct - The unstruct to resolve. + * + * @returns The resolved unstruct name. + * + * @example + * ```ts + * resolveUnstruct(ctx, { + * uv: d.float16x2, // -> d.vec2f after resolution + * color: d.snorm8x4, -> d.vec4f after resolution + * }); + * ``` + */ +function resolveUnstruct(ctx: ResolutionCtx, unstruct: Unstruct) { + const id = ctx.names.makeUnique(unstruct.label); + + ctx.addDeclaration(` +struct ${id} { +${Object.entries(unstruct.propTypes) + .map((prop) => + isAttribute(prop[1]) + ? resolveStructProperty(ctx, [prop[0], formatToWGSLType[prop[1].format]]) + : resolveStructProperty(ctx, prop), + ) + .join('')} +}\n`); + + return id; +} + /** * Resolves an array. * @param ctx - The resolution context. @@ -137,6 +177,18 @@ function resolveArray(ctx: ResolutionCtx, array: WgslArray) { : `array<${element}, ${array.elementCount}>`; } +function resolveDisarray(ctx: ResolutionCtx, disarray: Disarray) { + const element = ctx.resolve( + isAttribute(disarray.elementType) + ? formatToWGSLType[disarray.elementType.format] + : (disarray.elementType as AnyWgslData), + ); + + return disarray.elementCount === 0 + ? `array<${element}>` + : `array<${element}, ${disarray.elementCount}>`; +} + /** * Resolves a WGSL data-type schema to a string. * @param ctx - The resolution context. @@ -144,7 +196,27 @@ function resolveArray(ctx: ResolutionCtx, array: WgslArray) { * * @returns The resolved data-type string. */ -export function resolveData(ctx: ResolutionCtx, data: AnyWgslData): string { +export function resolveData(ctx: ResolutionCtx, data: AnyData): string { + if (isLooseData(data)) { + if (data.type === 'unstruct') { + return resolveUnstruct(ctx, data); + } + + if (data.type === 'disarray') { + return resolveDisarray(ctx, data); + } + + if (data.type === 'loose-decorated') { + return ctx.resolve( + isAttribute(data.inner) + ? formatToWGSLType[data.inner.format] + : data.inner, + ); + } + + return ctx.resolve(formatToWGSLType[data.type]); + } + if (isIdentityType(data)) { return data.type; } diff --git a/packages/typegpu/src/core/root/init.ts b/packages/typegpu/src/core/root/init.ts index 8901a3d86..b9eba81c7 100644 --- a/packages/typegpu/src/core/root/init.ts +++ b/packages/typegpu/src/core/root/init.ts @@ -65,6 +65,10 @@ import { isTexture, } from '../texture/texture'; import type { LayoutToAllowedAttribs } from '../vertexLayout/vertexAttribute'; +import { + type TgpuVertexLayout, + isVertexLayout, +} from '../vertexLayout/vertexLayout'; import type { CreateTextureOptions, CreateTextureResult, @@ -307,6 +311,7 @@ class TgpuRootImpl | TgpuMutableTexture | TgpuSampledTexture, ): GPUTextureView; + unwrap(resource: TgpuVertexLayout): GPUVertexBufferLayout; unwrap( resource: | TgpuComputePipeline @@ -317,14 +322,16 @@ class TgpuRootImpl | TgpuReadonlyTexture | TgpuWriteonlyTexture | TgpuMutableTexture - | TgpuSampledTexture, + | TgpuSampledTexture + | TgpuVertexLayout, ): | GPUComputePipeline | GPUBindGroupLayout | GPUBindGroup | GPUBuffer | GPUTexture - | GPUTextureView { + | GPUTextureView + | GPUVertexBufferLayout { if (isComputePipeline(resource)) { return (resource as unknown as INTERNAL_TgpuComputePipeline).rawPipeline; } @@ -355,6 +362,10 @@ class TgpuRootImpl return (resource as unknown as INTERNAL_TgpuFixedSampledTexture).unwrap(); } + if (isVertexLayout(resource)) { + return resource.vertexLayout; + } + throw new Error(`Unknown resource type: ${resource}`); } diff --git a/packages/typegpu/src/core/vertexLayout/connectAttributesToShader.ts b/packages/typegpu/src/core/vertexLayout/connectAttributesToShader.ts index f56fdecbd..609cddedc 100644 --- a/packages/typegpu/src/core/vertexLayout/connectAttributesToShader.ts +++ b/packages/typegpu/src/core/vertexLayout/connectAttributesToShader.ts @@ -16,9 +16,9 @@ export interface ConnectAttributesToShaderResult { bufferDefinitions: GPUVertexBufferLayout[]; } -function isAttribute( - value: unknown | T, -): value is T { +export function isAttribute< + T extends TgpuVertexAttrib & INTERNAL_TgpuVertexAttrib, +>(value: unknown | T): value is T { return typeof (value as T)?.format === 'string'; } diff --git a/packages/typegpu/src/core/vertexLayout/vertexLayout.ts b/packages/typegpu/src/core/vertexLayout/vertexLayout.ts index 197d79934..630a655a5 100644 --- a/packages/typegpu/src/core/vertexLayout/vertexLayout.ts +++ b/packages/typegpu/src/core/vertexLayout/vertexLayout.ts @@ -1,5 +1,9 @@ import { alignmentOf, customAlignmentOf } from '../../data/alignmentOf'; -import { isLooseDecorated, isUnstruct } from '../../data/dataTypes'; +import { + getCustomLocation, + isLooseDecorated, + isUnstruct, +} from '../../data/dataTypes'; import type { Disarray } from '../../data/dataTypes'; import { sizeOf } from '../../data/sizeOf'; import { isDecorated, isWgslStruct } from '../../data/wgslTypes'; @@ -31,6 +35,7 @@ export interface TgpuVertexLayout< readonly stride: number; readonly stepMode: 'vertex' | 'instance'; readonly attrib: ArrayToContainedAttribs; + readonly vertexLayout: GPUVertexBufferLayout; schemaForCount(n: number): TData; } @@ -58,6 +63,8 @@ export function isVertexLayout( // Implementation // -------------- +const defaultAttribEntry = Symbol('defaultAttribEntry'); + function dataToContainedAttribs< TLayoutData extends WgslArray | Disarray, TData extends BaseWgslData, @@ -65,12 +72,20 @@ function dataToContainedAttribs< layout: TgpuVertexLayout, data: TData, offset: number, + customLocationMap: Record, + key?: string, ): DataToContainedAttribs { if (isDecorated(data) || isLooseDecorated(data)) { + const customLocation = getCustomLocation(data); + if (customLocation !== undefined) { + customLocationMap[key ?? defaultAttribEntry] = customLocation; + } + return dataToContainedAttribs( layout, data.inner, roundUp(offset, customAlignmentOf(data)), + customLocationMap, ); } @@ -82,7 +97,13 @@ function dataToContainedAttribs< memberOffset = roundUp(memberOffset, alignmentOf(value)); const attrib = [ key, - dataToContainedAttribs(layout, value, memberOffset), + dataToContainedAttribs( + layout, + value, + memberOffset, + customLocationMap, + key, + ), ]; memberOffset += sizeOf(value); return attrib; @@ -98,7 +119,13 @@ function dataToContainedAttribs< memberOffset = roundUp(memberOffset, customAlignmentOf(value)); const attrib = [ key, - dataToContainedAttribs(layout, value, memberOffset), + dataToContainedAttribs( + layout, + value, + memberOffset, + customLocationMap, + key, + ), ]; memberOffset += sizeOf(value); return attrib; @@ -139,6 +166,7 @@ class TgpuVertexLayoutImpl public readonly resourceType = 'vertex-layout'; public readonly stride: number; public readonly attrib: ArrayToContainedAttribs; + private readonly _customLocationMap = {} as Record; private _label: string | undefined; @@ -153,13 +181,68 @@ class TgpuVertexLayoutImpl sizeOf(arraySchema.elementType), alignmentOf(arraySchema), ); - this.attrib = dataToContainedAttribs(this, arraySchema.elementType, 0); + this.attrib = dataToContainedAttribs( + this, + arraySchema.elementType, + 0, + this._customLocationMap, + ); } get label(): string | undefined { return this._label; } + get vertexLayout(): GPUVertexBufferLayout { + // If defaultAttribEntry is in the custom location map, + // it means that the vertex layout is based on a single attribute + if (this._customLocationMap[defaultAttribEntry] !== undefined) { + if ( + typeof this.attrib.format !== 'string' || + typeof this.attrib.offset !== 'number' + ) { + throw new Error( + 'Single attribute vertex layouts must have a format and offset.', + ); + } + + return { + arrayStride: this.stride, + stepMode: this.stepMode, + attributes: [ + { + format: this.attrib.format, + offset: this.attrib.offset, + shaderLocation: this._customLocationMap[defaultAttribEntry], + }, + ], + }; + } + + // check if all attributes have custom locations + const allAttributesHaveCustomLocations = Object.keys(this.attrib).every( + (key) => this._customLocationMap[key] !== undefined, + ); + + if (!allAttributesHaveCustomLocations) { + throw new Error( + 'All attributes must have custom locations in order to unwrap a vertex layout.', + ); + } + + return { + arrayStride: this.stride, + stepMode: this.stepMode, + attributes: [ + ...Object.entries(this.attrib).map(([key, attrib]) => ({ + format: attrib.format, + offset: attrib.offset, + shaderLocation: this._customLocationMap[key], + })), + ] as GPUVertexAttribute[], + }; + } + $name(label?: string | undefined): this { this._label = label; return this; diff --git a/packages/typegpu/src/data/dataTypes.ts b/packages/typegpu/src/data/dataTypes.ts index 03803fa14..a42e59c75 100644 --- a/packages/typegpu/src/data/dataTypes.ts +++ b/packages/typegpu/src/data/dataTypes.ts @@ -1,3 +1,4 @@ +import type { TgpuNamable } from '../namable'; import type { Infer, InferRecord } from '../shared/repr'; import { vertexFormats } from '../shared/vertexFormat'; import type { PackedData } from './vertexFormatData'; @@ -33,7 +34,8 @@ export interface Unstruct< string, wgsl.BaseWgslData >, -> { +> extends TgpuNamable { + readonly label?: string | undefined; readonly type: 'unstruct'; readonly propTypes: TProps; readonly '~repr': InferRecord; diff --git a/packages/typegpu/src/data/unstruct.ts b/packages/typegpu/src/data/unstruct.ts index 13020b20c..226ea0b7b 100644 --- a/packages/typegpu/src/data/unstruct.ts +++ b/packages/typegpu/src/data/unstruct.ts @@ -37,9 +37,20 @@ export const unstruct = >( class UnstructImpl> implements Unstruct { + private _label: string | undefined; + public readonly type = 'unstruct'; /** Type-token, not available at runtime */ public readonly '~repr'!: InferRecord; constructor(public readonly propTypes: TProps) {} + + get label() { + return this._label; + } + + $name(label: string) { + this._label = label; + return this; + } } diff --git a/packages/typegpu/src/data/vertexFormatData.ts b/packages/typegpu/src/data/vertexFormatData.ts index 9a3da1be4..682bec991 100644 --- a/packages/typegpu/src/data/vertexFormatData.ts +++ b/packages/typegpu/src/data/vertexFormatData.ts @@ -30,7 +30,7 @@ class TgpuVertexFormatDataImpl constructor(public readonly type: T) {} } -const formatToWGSLType = { +export const formatToWGSLType = { uint8: u32, uint8x2: vec2u, uint8x4: vec4u, diff --git a/packages/typegpu/src/resolutionCtx.ts b/packages/typegpu/src/resolutionCtx.ts index 13320679a..b2449577c 100644 --- a/packages/typegpu/src/resolutionCtx.ts +++ b/packages/typegpu/src/resolutionCtx.ts @@ -9,6 +9,7 @@ import { isProviding, isSlot, } from './core/slot/slotTypes'; +import { isLooseData } from './data'; import { type AnyWgslData, type BaseWgslData, @@ -474,7 +475,7 @@ class ResolutionCtxImpl implements ResolutionCtx { // If we got here, no item with the given slot-to-value combo exists in cache yet let result: string; - if (isWgslData(item)) { + if (isWgslData(item) || isLooseData(item)) { result = resolveData(this, item); } else if (isDerived(item) || isSlot(item)) { result = this.resolve(this.unwrap(item)); diff --git a/packages/typegpu/src/types.ts b/packages/typegpu/src/types.ts index 791db4c15..771f13431 100644 --- a/packages/typegpu/src/types.ts +++ b/packages/typegpu/src/types.ts @@ -20,6 +20,7 @@ import { import type { TgpuExternalTexture } from './core/texture/externalTexture'; import type { TgpuAnyTextureView, TgpuTexture } from './core/texture/texture'; import type { TgpuVar } from './core/variable/tgpuVariable'; +import type { AnyData } from './data'; import { type AnyMatInstance, type AnyVecInstance, @@ -53,7 +54,7 @@ export type ResolvableObject = | TgpuVar | AnyVecInstance | AnyMatInstance - | AnyWgslData + | AnyData // biome-ignore lint/suspicious/noExplicitAny: | TgpuFn; diff --git a/packages/typegpu/src/unwrapper.ts b/packages/typegpu/src/unwrapper.ts index 4b26a05b8..bfe696cbc 100644 --- a/packages/typegpu/src/unwrapper.ts +++ b/packages/typegpu/src/unwrapper.ts @@ -7,6 +7,7 @@ import type { TgpuTexture, TgpuWriteonlyTexture, } from './core/texture/texture'; +import type { TgpuVertexLayout } from './core/vertexLayout/vertexLayout'; import type { AnyData } from './data/dataTypes'; import type { TgpuBindGroup, TgpuBindGroupLayout } from './tgpuBindGroupLayout'; @@ -24,4 +25,5 @@ export interface Unwrapper { | TgpuMutableTexture | TgpuSampledTexture, ): GPUTextureView; + unwrap(resource: TgpuVertexLayout): GPUVertexBufferLayout; } diff --git a/packages/typegpu/tests/resolve.test.ts b/packages/typegpu/tests/resolve.test.ts index 0e4f970d9..f7ab2af94 100644 --- a/packages/typegpu/tests/resolve.test.ts +++ b/packages/typegpu/tests/resolve.test.ts @@ -176,6 +176,108 @@ describe('tgpu resolve', () => { ); }); + it('should resolve an unstruct to its corresponding struct', () => { + const vertexInfo = d.unstruct({ + color: d.snorm8x4, + colorHDR: d.unorm10_10_10_2, + position2d: d.float16x2, + }); + + const resolved = tgpu.resolve({ + template: 'fn foo() { var v: vertexInfo; }', + externals: { vertexInfo }, + names: 'strict', + }); + + expect(parse(resolved)).toEqual( + parse(` + struct vertexInfo { + color: vec4f, + colorHDR: vec4f, + position2d: vec2f, + } + fn foo() { var v: vertexInfo; } + `), + ); + }); + + it('should resolve an unstruct with a disarray to its corresponding struct', () => { + const vertexInfo = d.unstruct({ + color: d.snorm8x4, + colorHDR: d.unorm10_10_10_2, + position2d: d.float16x2, + extra: d.disarrayOf(d.snorm8x4, 16), + }); + + const resolved = tgpu.resolve({ + template: 'fn foo() { var v: vertexInfo; }', + externals: { vertexInfo }, + names: 'strict', + }); + + expect(parse(resolved)).toEqual( + parse(` + struct vertexInfo { + color: vec4f, + colorHDR: vec4f, + position2d: vec2f, + extra: array, + } + fn foo() { var v: vertexInfo; } + `), + ); + }); + + it('should resolve an unstruct with a complex nested structure', () => { + const vertexInfo = d.unstruct({ + color: d.snorm8x4, + colorHDR: d.unorm10_10_10_2, + position2d: d.float16x2, + extra: d + .unstruct({ + a: d.snorm8, + b: d.snorm8x4, + c: d.float16x2, + }) + .$name('extra'), + more: d.disarrayOf( + d.unstruct({ a: d.snorm8, b: d.snorm8x4 }).$name('more'), + 16, + ), + }); + + const resolved = tgpu.resolve({ + template: 'fn foo() { var v: vertexInfo; }', + externals: { vertexInfo }, + names: 'strict', + }); + + expect(parse(resolved)).toEqual( + parse(` + struct extra { + a: f32, + b: vec4f, + c: vec2f, + } + + struct more { + a: f32, + b: vec4f, + } + + struct vertexInfo { + color: vec4f, + colorHDR: vec4f, + position2d: vec2f, + extra: extra, + more: array, + } + + fn foo() { var v: vertexInfo; } + `), + ); + }); + it('should resolve object externals and replace their usages in template', () => { const getColor = tgpu['~unstable'] .fn([], d.vec3f) diff --git a/packages/typegpu/tests/root.test.ts b/packages/typegpu/tests/root.test.ts index e4b98cf8e..6e32003c0 100644 --- a/packages/typegpu/tests/root.test.ts +++ b/packages/typegpu/tests/root.test.ts @@ -1,4 +1,5 @@ import { describe, expect, vi } from 'vitest'; +import tgpu from '../src'; import * as d from '../src/data'; import { it } from './utils/extendedIt'; @@ -100,6 +101,64 @@ describe('TgpuRoot', () => { const buffer = root.createBuffer(d.u32, rawBuffer); expect(root.unwrap(buffer)).toBe(rawBuffer); }); + + it('should return the correct GPUVertexBufferLayout for a simple vertex layout', ({ + root, + }) => { + const vertexLayout = tgpu['~unstable'].vertexLayout( + (n) => d.arrayOf(d.location(0, d.vec2u), n), + 'vertex', + ); + + expect(root.unwrap(vertexLayout)).toEqual({ + arrayStride: 8, + stepMode: 'vertex', + attributes: [ + { + format: 'uint32x2', + offset: 0, + shaderLocation: 0, + }, + ], + }); + }); + + it('should return the correct GPUVertexBufferLayout for a complex vertex layout', ({ + root, + }) => { + const VertexData = d.unstruct({ + position: d.location(0, d.float32x3), + color: d.location(1, d.unorm10_10_10_2), + something: d.location(2, d.u32), + }); + + const vertexLayout = tgpu['~unstable'].vertexLayout( + (n) => d.disarrayOf(VertexData, n), + 'instance', + ); + + expect(root.unwrap(vertexLayout)).toEqual({ + arrayStride: 20, + stepMode: 'instance', + attributes: [ + { + format: 'float32x3', + offset: 0, + shaderLocation: 0, + }, + { + format: 'unorm10-10-10-2', + offset: 12, + shaderLocation: 1, + }, + { + format: 'uint32', + offset: 16, + shaderLocation: 2, + }, + ], + }); + }); }); // TODO: Adapt the tests to the new API