diff --git a/apps/typegpu-docs/src/content/docs/fundamentals/utils.mdx b/apps/typegpu-docs/src/content/docs/fundamentals/utils.mdx index 4c8fe0056c..e49508267b 100644 --- a/apps/typegpu-docs/src/content/docs/fundamentals/utils.mdx +++ b/apps/typegpu-docs/src/content/docs/fundamentals/utils.mdx @@ -5,39 +5,8 @@ description: A list of various utilities provided by TypeGPU. ## *prepareDispatch* -The `prepareDispatch` function simplifies running simple computations on the GPU. -Under the hood, it wraps the callback in a `TgpuFn`, creates a compute pipeline, and returns a dispatch function to execute it. - -This can help reduce serialization overhead when initializing buffers with data. - -```ts twoslash -import tgpu, { prepareDispatch } from 'typegpu'; -import * as d from 'typegpu/data'; - -const root = await tgpu.init(); - -const Boid = d.struct({ - index: d.u32, - pos: d.vec3f, -}); - -// buffer of 2048 Boids -const boidsMutable = root.createMutable(d.arrayOf(Boid, 2048)); - -const dispatch = prepareDispatch(root, (x) => { - 'kernel'; - const boidData = Boid({ index: x, pos: d.vec3f() }); - boidsMutable.$[x] = boidData; -}); -// run callback for each x in range 0..2047 -dispatch(2048); -``` - -:::note -Remember to mark the callback with `'kernel'` directive to let TypeGPU know that this function is TGSL. -::: - -The returned dispatch function can be called multiple times. +The `prepareDispatch` function streamlines running simple computations on the GPU. +Under the hood, it wraps the callback in a `TgpuFn`, creates a compute pipeline, and returns an object with dispatch method that executes the pipeline. Since the pipeline is reused, there’s no additional overhead for subsequent calls. ```ts twoslash @@ -52,16 +21,20 @@ const doubleUp = prepareDispatch(root, (x) => { data.$[x] *= 2; }); -doubleUp(8); -doubleUp(8); -doubleUp(4); +doubleUp.dispatch(8); +doubleUp.dispatch(8); +doubleUp.dispatch(4); -// no need to call `onSubmittedWorkDone()` because the command encoder -// will queue the read after `doubleUp` anyway +// the command encoder will queue the read after `doubleUp` console.log(await data.read()); // [0, 8, 16, 24, 16, 20, 24, 28] ``` +:::note +Remember to mark the callback with `'kernel'` directive to let TypeGPU know that this function is TGSL. +::: + The callback can have up to three arguments (dimensions). +`prepareDispatch` can simplify writing a pipeline helping reduce serialization overhead when initializing buffers with data. Buffer initialization commonly uses random number generators. For that, you can use the [`@typegpu/noise`](TypeGPU/ecosystem/typegpu-noise) library. @@ -82,22 +55,56 @@ prepareDispatch(root, (x, y) => { 'kernel'; randf.seed2(d.vec2f(x, y).div(1024)); waterLevelMutable.$[x][y] = 10 + randf.sample(); -})(1024, 512); +}).dispatch(1024, 512); // callback will be called for x in range 0..1023 and y in range 0..511 // (optional) read values in JS console.log(await waterLevelMutable.read()); ``` -It is highly recommended NOT to use `dispatch` for: +Analogously to `TgpuComputePipeline`, the result of `prepareDispatch` can have bind groups bound using the `with` method. + +```ts twoslash +import tgpu, { prepareDispatch } from 'typegpu'; +import * as d from 'typegpu/data'; +import * as std from 'typegpu/std'; +const root = await tgpu.init(); +// ---cut--- +const layout = tgpu.bindGroupLayout({ + buffer: { storage: d.arrayOf(d.u32), access: 'mutable' }, +}); +const buffer1 = root + .createBuffer(d.arrayOf(d.u32, 3), [1, 2, 3]).$usage('storage'); +const buffer2 = root + .createBuffer(d.arrayOf(d.u32, 4), [2, 4, 8, 16]).$usage('storage'); +const bindGroup1 = root.createBindGroup(layout, { + buffer: buffer1, +}); +const bindGroup2 = root.createBindGroup(layout, { + buffer: buffer2, +}); + +const test = prepareDispatch(root, (x) => { + 'kernel'; + layout.$.buffer[x] *= 2; +}); + +test.with(layout, bindGroup1).dispatch(3); +test.with(layout, bindGroup2).dispatch(4); + +console.log(await buffer1.read()); // [2, 4, 6]; +console.log(await buffer2.read()); // [4, 8, 16, 32]; +``` + +It is recommended NOT to use `prepareDispatch` for: - More complex compute shaders. -When using `dispatch`, it is impossible to switch bind groups or to change workgroup sizes. +When using `prepareDispatch`, it is impossible to change workgroup sizes or to use [slots](/TypeGPU/fundamentals/slots). For such cases, a manually created pipeline would be more suitable. - Small calls. Usually, for small data the shader creation and dispatch is more costly than serialization. -Small buffers can be more efficiently initialized with `buffer.write()` method. +Small buffers can be more efficiently initialized with the `buffer.write()` method. :::note The default workgroup sizes are: @@ -122,14 +129,14 @@ import * as d from 'typegpu/data'; const root = await tgpu.init(); // ---cut--- const callCountMutable = root.createMutable(d.u32, 0); -const dispatch = prepareDispatch(root, () => { +const compute = prepareDispatch(root, () => { 'kernel'; callCountMutable.$ += 1; console.log('Call number', callCountMutable.$); }); -dispatch(); -dispatch(); +compute.dispatch(); +compute.dispatch(); // Eventually... // "[GPU] Call number 1" diff --git a/apps/typegpu-docs/src/examples/rendering/3d-fish/index.ts b/apps/typegpu-docs/src/examples/rendering/3d-fish/index.ts index 740cf6b873..8d92328a5d 100644 --- a/apps/typegpu-docs/src/examples/rendering/3d-fish/index.ts +++ b/apps/typegpu-docs/src/examples/rendering/3d-fish/index.ts @@ -98,7 +98,7 @@ function enqueuePresetChanges() { const buffer0mutable = fishDataBuffers[0].as('mutable'); const buffer1mutable = fishDataBuffers[1].as('mutable'); const seedUniform = root.createUniform(d.f32); -const randomizeFishPositionsDispatch = prepareDispatch(root, (x) => { +const randomizeFishPositionsOnGPU = prepareDispatch(root, (x) => { 'kernel'; randf.seed2(d.vec2f(d.f32(x), seedUniform.$)); const data = ModelData({ @@ -124,7 +124,7 @@ const randomizeFishPositionsDispatch = prepareDispatch(root, (x) => { const randomizeFishPositions = () => { seedUniform.write((performance.now() % 10000) / 10000); - randomizeFishPositionsDispatch(p.fishAmount); + randomizeFishPositionsOnGPU.dispatch(p.fishAmount); enqueuePresetChanges(); }; diff --git a/apps/typegpu-docs/src/examples/simple/increment/index.ts b/apps/typegpu-docs/src/examples/simple/increment/index.ts index 91fbddb0af..70a7e61f2a 100644 --- a/apps/typegpu-docs/src/examples/simple/increment/index.ts +++ b/apps/typegpu-docs/src/examples/simple/increment/index.ts @@ -6,14 +6,14 @@ const root = await tgpu.init(); const counter = root.createMutable(d.u32); // A 0-dimentional shader function -const dispatch = prepareDispatch(root, () => { +const incrementKernel = prepareDispatch(root, () => { 'kernel'; counter.$ += 1; }); async function increment() { // Dispatch and read the result - dispatch(); + incrementKernel.dispatch(); return await counter.read(); } diff --git a/apps/typegpu-docs/src/examples/tests/dispatch/index.ts b/apps/typegpu-docs/src/examples/tests/dispatch/index.ts index c49da2c8a5..08d5767102 100644 --- a/apps/typegpu-docs/src/examples/tests/dispatch/index.ts +++ b/apps/typegpu-docs/src/examples/tests/dispatch/index.ts @@ -16,7 +16,7 @@ async function test0d(): Promise { prepareDispatch(root, () => { 'kernel'; mutable.$ = 126; - })(); + }).dispatch(); const filled = await mutable.read(); return isEqual(filled, 126); } @@ -27,7 +27,7 @@ async function test1d(): Promise { prepareDispatch(root, (x) => { 'kernel'; mutable.$[x] = x; - })(...size); + }).dispatch(...size); const filled = await mutable.read(); return isEqual(filled, [0, 1, 2, 3, 4, 5, 6]); } @@ -40,7 +40,7 @@ async function test2d(): Promise { prepareDispatch(root, (x, y) => { 'kernel'; mutable.$[x][y] = d.vec2u(x, y); - })(...size); + }).dispatch(...size); const filled = await mutable.read(); return isEqual(filled, [ [d.vec2u(0, 0), d.vec2u(0, 1), d.vec2u(0, 2)], @@ -59,7 +59,7 @@ async function test3d(): Promise { prepareDispatch(root, (x, y, z) => { 'kernel'; mutable.$[x][y][z] = d.vec3u(x, y, z); - })(...size); + }).dispatch(...size); const filled = await mutable.read(); return isEqual(filled, [ [[d.vec3u(0, 0, 0), d.vec3u(0, 0, 1)]], @@ -72,7 +72,7 @@ async function testWorkgroupSize(): Promise { prepareDispatch(root, (x, y, z) => { 'kernel'; std.atomicAdd(mutable.$, 1); - })(4, 3, 2); + }).dispatch(4, 3, 2); const filled = await mutable.read(); return isEqual(filled, 4 * 3 * 2); } @@ -81,17 +81,47 @@ async function testMultipleDispatches(): Promise { const size = [7] as const; const mutable = root .createMutable(d.arrayOf(d.u32, size[0]), [0, 1, 2, 3, 4, 5, 6]); - const dispatch = prepareDispatch(root, (x: number) => { + const test = prepareDispatch(root, (x: number) => { 'kernel'; mutable.$[x] *= 2; }); - dispatch(6); - dispatch(2); - dispatch(4); + test.dispatch(6); + test.dispatch(2); + test.dispatch(4); const filled = await mutable.read(); return isEqual(filled, [0 * 8, 1 * 8, 2 * 4, 3 * 4, 4 * 2, 5 * 2, 6 * 1]); } +async function testDifferentBindGroups(): Promise { + const layout = tgpu.bindGroupLayout({ + buffer: { storage: d.arrayOf(d.u32), access: 'mutable' }, + }); + const buffer1 = root + .createBuffer(d.arrayOf(d.u32, 3), [1, 2, 3]).$usage('storage'); + const buffer2 = root + .createBuffer(d.arrayOf(d.u32, 4), [2, 4, 8, 16]).$usage('storage'); + const bindGroup1 = root.createBindGroup(layout, { + buffer: buffer1, + }); + const bindGroup2 = root.createBindGroup(layout, { + buffer: buffer2, + }); + + const test = prepareDispatch(root, () => { + 'kernel'; + for (let i = d.u32(); i < std.arrayLength(layout.$.buffer); i++) { + layout.$.buffer[i] *= 2; + } + }); + + test.with(layout, bindGroup1).dispatch(); + test.with(layout, bindGroup2).dispatch(); + + const filled1 = await buffer1.read(); + const filled2 = await buffer2.read(); + return isEqual(filled1, [2, 4, 6]) && isEqual(filled2, [4, 8, 16, 32]); +} + async function runTests(): Promise { let result = true; result = await test0d() && result; @@ -100,6 +130,7 @@ async function runTests(): Promise { result = await test3d() && result; result = await testWorkgroupSize() && result; result = await testMultipleDispatches() && result; + result = await testDifferentBindGroups() && result; return result; } diff --git a/apps/typegpu-docs/src/examples/tests/log-test/index.ts b/apps/typegpu-docs/src/examples/tests/log-test/index.ts index aaf0299ddf..67af56b7c8 100644 --- a/apps/typegpu-docs/src/examples/tests/log-test/index.ts +++ b/apps/typegpu-docs/src/examples/tests/log-test/index.ts @@ -20,21 +20,21 @@ export const controls = { prepareDispatch(root, () => { 'kernel'; console.log(d.u32(321)); - })(), + }).dispatch(), }, 'Multiple arguments': { onButtonClick: () => prepareDispatch(root, () => { 'kernel'; console.log(d.u32(1), d.vec3u(2, 3, 4), d.u32(5), d.u32(6)); - })(), + }).dispatch(), }, 'String literals': { onButtonClick: () => prepareDispatch(root, () => { 'kernel'; console.log(d.u32(2), 'plus', d.u32(3), 'equals', d.u32(5)); - })(), + }).dispatch(), }, 'Two logs': { onButtonClick: () => @@ -42,7 +42,7 @@ export const controls = { 'kernel'; console.log('First log.'); console.log('Second log.'); - })(), + }).dispatch(), }, 'Different types': { onButtonClick: () => @@ -86,25 +86,25 @@ export const controls = { } else { console.log("The 'shader-f16' flag is not enabled."); } - })(), + }).dispatch(), }, 'Two threads': { onButtonClick: () => prepareDispatch(root, (x) => { 'kernel'; console.log('Log from thread', x); - })(2), + }).dispatch(2), }, '100 dispatches': { onButtonClick: async () => { const indexUniform = root.createUniform(d.u32); - const dispatch = prepareDispatch(root, () => { + const test = prepareDispatch(root, () => { 'kernel'; console.log('Log from dispatch', indexUniform.$); }); for (let i = 0; i < 100; i++) { indexUniform.write(i); - dispatch(); + test.dispatch(); console.log(`dispatched ${i}`); } }, @@ -112,16 +112,16 @@ export const controls = { 'Varying size logs': { onButtonClick: async () => { const logCountUniform = root.createUniform(d.u32); - const dispatch = prepareDispatch(root, () => { + const test = prepareDispatch(root, () => { 'kernel'; for (let i = d.u32(); i < logCountUniform.$; i++) { console.log('Log index', d.u32(i) + 1, 'out of', logCountUniform.$); } }); logCountUniform.write(3); - dispatch(); + test.dispatch(); logCountUniform.write(1); - dispatch(); + test.dispatch(); }, }, 'Render pipeline': { @@ -179,16 +179,15 @@ export const controls = { console.log('Log 1 from thread', x); console.log('Log 2 from thread', x); console.log('Log 3 from thread', x); - })(16), + }).dispatch(16), }, 'Too much data': { onButtonClick: () => { - const dispatch = prepareDispatch(root, () => { - 'kernel'; - console.log(d.mat4x4f(), d.mat4x4f(), 1); - }); try { - dispatch(); + prepareDispatch(root, () => { + 'kernel'; + console.log(d.vec3u(), d.vec3u(), d.vec3u()); + }).dispatch(); } catch (err) { console.log(err); } diff --git a/packages/typegpu/src/prepareDispatch.ts b/packages/typegpu/src/prepareDispatch.ts index a863c374a6..b7d3e22538 100644 --- a/packages/typegpu/src/prepareDispatch.ts +++ b/packages/typegpu/src/prepareDispatch.ts @@ -1,18 +1,16 @@ import { builtin } from './builtin.ts'; import { computeFn } from './core/function/tgpuComputeFn.ts'; import { fn } from './core/function/tgpuFn.ts'; +import type { TgpuComputePipeline } from './core/pipeline/computePipeline.ts'; import type { TgpuRoot } from './core/root/rootTypes.ts'; import { u32 } from './data/numeric.ts'; import { vec3f, vec3u } from './data/vector.ts'; import type { v3u } from './data/wgslTypes.ts'; import { ceil } from './std/numeric.ts'; - -const workgroupSizeConfigs = [ - vec3u(1, 1, 1), - vec3u(256, 1, 1), - vec3u(16, 16, 1), - vec3u(8, 8, 4), -] as const; +import type { + TgpuBindGroup, + TgpuBindGroupLayout, +} from './tgpuBindGroupLayout.ts'; /** * Changes the given array to a vec of 3 numbers, filling missing values with 1. @@ -32,6 +30,48 @@ type DispatchForArgs = TArgs extends { length: infer TLength } : never : never; +class PreparedDispatch { + #pipeline: TgpuComputePipeline; + #createDispatch: (pipeline: TgpuComputePipeline) => DispatchForArgs; + constructor( + createDispatch: (pipeline: TgpuComputePipeline) => DispatchForArgs, + pipeline: TgpuComputePipeline, + ) { + this.#createDispatch = createDispatch; + this.#pipeline = pipeline; + } + + /** + * Returns a new PreparedDispatch with the specified bind group bound. + * Analogous to `TgpuComputePipeline.with()`. + */ + with( + bindGroupLayout: TgpuBindGroupLayout, + bindGroup: TgpuBindGroup, + ): PreparedDispatch { + return new PreparedDispatch( + this.#createDispatch, + this.#pipeline.with(bindGroupLayout, bindGroup), + ); + } + + /** + * Run the prepared dispatch. + * Unlike `TgpuComputePipeline.dispatchWorkgroups()`, + * this method takes in the number of threads to run in each dimension. + */ + get dispatch(): DispatchForArgs { + return this.#createDispatch(this.#pipeline); + } +} + +const workgroupSizeConfigs = [ + vec3u(1, 1, 1), + vec3u(256, 1, 1), + vec3u(16, 16, 1), + vec3u(8, 8, 4), +] as const; + /** * Creates a dispatch function for a compute pipeline. * @@ -43,7 +83,7 @@ type DispatchForArgs = TArgs extends { length: infer TLength } export function prepareDispatch( root: TgpuRoot, callback: (...args: TArgs) => undefined, -): DispatchForArgs { +): PreparedDispatch { if (callback.length >= 4) { throw new Error('Dispatch only supports up to three dimensions.'); } @@ -60,25 +100,30 @@ export function prepareDispatch( workgroupSize: workgroupSize, in: { id: builtin.globalInvocationId }, })`{ - if (any(in.id >= sizeUniform)) { - return; - } - wrappedCallback(in.id.x, in.id.y, in.id.z); - }`.$uses({ sizeUniform, wrappedCallback }); + if (any(in.id >= sizeUniform)) { + return; + } + wrappedCallback(in.id.x, in.id.y, in.id.z); +}`.$uses({ sizeUniform, wrappedCallback }); const pipeline = root['~unstable'] .withCompute(mainCompute) .createPipeline(); - return ((...size: (number | undefined)[]) => { - const sanitizedSize = toVec3(size); - const workgroupCount = ceil(vec3f(sanitizedSize).div(vec3f(workgroupSize))); - sizeUniform.write(sanitizedSize); - pipeline.dispatchWorkgroups( - workgroupCount.x, - workgroupCount.y, - workgroupCount.z, - ); - root['~unstable'].flush(); - }) as DispatchForArgs; + const createDispatch = (pipeline: TgpuComputePipeline) => + ((...size: (number | undefined)[]) => { + const sanitizedSize = toVec3(size); + const workgroupCount = ceil( + vec3f(sanitizedSize).div(vec3f(workgroupSize)), + ); + sizeUniform.write(sanitizedSize); + pipeline.dispatchWorkgroups( + workgroupCount.x, + workgroupCount.y, + workgroupCount.z, + ); + root['~unstable'].flush(); + }) as DispatchForArgs; + + return new PreparedDispatch(createDispatch, pipeline); } diff --git a/packages/typegpu/tests/examples/individual/3d-fish.test.ts b/packages/typegpu/tests/examples/individual/3d-fish.test.ts index b614daf176..941c947e3a 100644 --- a/packages/typegpu/tests/examples/individual/3d-fish.test.ts +++ b/packages/typegpu/tests/examples/individual/3d-fish.test.ts @@ -77,11 +77,11 @@ describe('3d fish example', () => { } @compute @workgroup_size(256, 1, 1) fn mainCompute_0(in: mainCompute_Input_12) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); + if (any(in.id >= sizeUniform_1)) { + return; } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } struct ModelData_2 { position: vec3f, diff --git a/packages/typegpu/tests/examples/individual/dispatch.test.ts b/packages/typegpu/tests/examples/individual/dispatch.test.ts index 0565e05110..fb29619861 100644 --- a/packages/typegpu/tests/examples/individual/dispatch.test.ts +++ b/packages/typegpu/tests/examples/individual/dispatch.test.ts @@ -30,11 +30,11 @@ describe('tgsl parsing test example', () => { } @compute @workgroup_size(1, 1, 1) fn mainCompute_0(in: mainCompute_Input_4) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); + if (any(in.id >= sizeUniform_1)) { + return; } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } @group(0) @binding(0) var sizeUniform_1: vec3u; @@ -49,11 +49,11 @@ describe('tgsl parsing test example', () => { } @compute @workgroup_size(256, 1, 1) fn mainCompute_0(in: mainCompute_Input_4) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); + if (any(in.id >= sizeUniform_1)) { + return; } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } @group(0) @binding(0) var sizeUniform_1: vec3u; @@ -68,11 +68,11 @@ describe('tgsl parsing test example', () => { } @compute @workgroup_size(16, 16, 1) fn mainCompute_0(in: mainCompute_Input_4) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); + if (any(in.id >= sizeUniform_1)) { + return; } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } @group(0) @binding(0) var sizeUniform_1: vec3u; @@ -87,11 +87,11 @@ describe('tgsl parsing test example', () => { } @compute @workgroup_size(8, 8, 4) fn mainCompute_0(in: mainCompute_Input_4) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); + if (any(in.id >= sizeUniform_1)) { + return; } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } @group(0) @binding(0) var sizeUniform_1: vec3u; @@ -106,11 +106,11 @@ describe('tgsl parsing test example', () => { } @compute @workgroup_size(8, 8, 4) fn mainCompute_0(in: mainCompute_Input_4) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); + if (any(in.id >= sizeUniform_1)) { + return; } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } @group(0) @binding(0) var sizeUniform_1: vec3u; @@ -125,11 +125,32 @@ describe('tgsl parsing test example', () => { } @compute @workgroup_size(256, 1, 1) fn mainCompute_0(in: mainCompute_Input_4) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); - }" + if (any(in.id >= sizeUniform_1)) { + return; + } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } + + @group(0) @binding(0) var sizeUniform_1: vec3u; + + @group(1) @binding(0) var buffer_3: array; + + fn wrappedCallback_2(_arg_0: u32, _arg_1: u32, _arg_2: u32) { + for (var i = 0u; (i < arrayLength(&buffer_3)); i++) { + buffer_3[i] *= 2; + } + } + + struct mainCompute_Input_4 { + @builtin(global_invocation_id) id: vec3u, + } + + @compute @workgroup_size(1, 1, 1) fn mainCompute_0(in: mainCompute_Input_4) { + if (any(in.id >= sizeUniform_1)) { + return; + } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + }" `); }); }); diff --git a/packages/typegpu/tests/examples/individual/increment.test.ts b/packages/typegpu/tests/examples/individual/increment.test.ts index 933ec45666..875453560d 100644 --- a/packages/typegpu/tests/examples/individual/increment.test.ts +++ b/packages/typegpu/tests/examples/individual/increment.test.ts @@ -31,11 +31,11 @@ describe('increment example', () => { } @compute @workgroup_size(1, 1, 1) fn mainCompute_0(in: mainCompute_Input_4) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); - }" + if (any(in.id >= sizeUniform_1)) { + return; + } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + }" `); }); }); diff --git a/packages/typegpu/tests/examples/individual/log-test.test.ts b/packages/typegpu/tests/examples/individual/log-test.test.ts index a372c06a07..cfcf7f74be 100644 --- a/packages/typegpu/tests/examples/individual/log-test.test.ts +++ b/packages/typegpu/tests/examples/individual/log-test.test.ts @@ -66,11 +66,11 @@ describe('console log example', () => { } @compute @workgroup_size(1, 1, 1) fn mainCompute_0(in: mainCompute_Input_8) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); + if (any(in.id >= sizeUniform_1)) { + return; } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } @group(0) @binding(0) var sizeUniform_1: vec3u; @@ -119,11 +119,11 @@ describe('console log example', () => { } @compute @workgroup_size(1, 1, 1) fn mainCompute_0(in: mainCompute_Input_9) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); + if (any(in.id >= sizeUniform_1)) { + return; } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } @group(0) @binding(0) var sizeUniform_1: vec3u; @@ -164,11 +164,11 @@ describe('console log example', () => { } @compute @workgroup_size(1, 1, 1) fn mainCompute_0(in: mainCompute_Input_8) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); + if (any(in.id >= sizeUniform_1)) { + return; } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } @group(0) @binding(0) var sizeUniform_1: vec3u; @@ -670,11 +670,11 @@ describe('console log example', () => { } @compute @workgroup_size(1, 1, 1) fn mainCompute_0(in: mainCompute_Input_54) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); + if (any(in.id >= sizeUniform_1)) { + return; } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } @group(0) @binding(0) var sizeUniform_1: vec3u; @@ -717,11 +717,11 @@ describe('console log example', () => { } @compute @workgroup_size(1, 1, 1) fn mainCompute_0(in: mainCompute_Input_8) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); + if (any(in.id >= sizeUniform_1)) { + return; } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } @group(0) @binding(0) var sizeUniform_1: vec3u; @@ -758,11 +758,11 @@ describe('console log example', () => { } @compute @workgroup_size(256, 1, 1) fn mainCompute_0(in: mainCompute_Input_8) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); + if (any(in.id >= sizeUniform_1)) { + return; } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } @group(0) @binding(0) var sizeUniform_1: vec3u; @@ -801,11 +801,11 @@ describe('console log example', () => { } @compute @workgroup_size(1, 1, 1) fn mainCompute_0(in: mainCompute_Input_9) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); + if (any(in.id >= sizeUniform_1)) { + return; } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } @group(0) @binding(0) var sizeUniform_1: vec3u; @@ -848,11 +848,11 @@ describe('console log example', () => { } @compute @workgroup_size(1, 1, 1) fn mainCompute_0(in: mainCompute_Input_9) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); + if (any(in.id >= sizeUniform_1)) { + return; } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + } struct mainVertex_Output_1 { @builtin(position) pos: vec4f, @@ -961,11 +961,11 @@ describe('console log example', () => { } @compute @workgroup_size(256, 1, 1) fn mainCompute_0(in: mainCompute_Input_10) { - if (any(in.id >= sizeUniform_1)) { - return; - } - wrappedCallback_2(in.id.x, in.id.y, in.id.z); - }" + if (any(in.id >= sizeUniform_1)) { + return; + } + wrappedCallback_2(in.id.x, in.id.y, in.id.z); + }" `); }); });