diff --git a/apps/typegpu-docs/src/content/docs/ecosystem/typegpu-noise.mdx b/apps/typegpu-docs/src/content/docs/ecosystem/typegpu-noise.mdx index 58b14520fb..631524bad4 100644 --- a/apps/typegpu-docs/src/content/docs/ecosystem/typegpu-noise.mdx +++ b/apps/typegpu-docs/src/content/docs/ecosystem/typegpu-noise.mdx @@ -277,34 +277,33 @@ const f = tgpu.computeFn({ workgroupSize: [1] })(() => { }); // ---cut--- import { + hash, + randomGeneratorShell, randomGeneratorSlot, + u32To01F32, type StatefulGenerator, } from '@typegpu/noise'; -const LCG: StatefulGenerator = (() => { +const LCG32: StatefulGenerator = (() => { const seed = tgpu.privateVar(d.u32); - const u32To01Float = tgpu.fn([d.u32], d.f32)((value) => { - const mantissa = value >> 9; - const bits = 0x3F800000 | mantissa; - const f = std.bitcastU32toF32(bits); - return f - 1; - }); - - return { - seed2: (value: d.v2f) => { - 'use gpu'; - seed.$ = d.u32(value.x * std.pow(32, 3) + value.y * std.pow(32, 2)); - }, - sample: () => { - 'use gpu'; - seed.$ = seed.$ * 1664525 + 1013904223; // % 2 ^ 32 - return u32To01Float(seed.$); - }, - }; + const multiplier = 1664525; + const increment = 1013904223; + + return { + seed: tgpu.fn([d.f32])((value) => { + seed.$ = hash(d.u32(value)); + }), + + sample: randomGeneratorShell(() => { + 'use gpu'; + seed.$ = multiplier * seed.$ + increment; // % 2 ^ 32 + return u32To01F32(seed.$); + }).$name('sample'), + }; })(); const pipeline = root - .with(randomGeneratorSlot, LCG) + .with(randomGeneratorSlot, LCG32) .createComputePipeline({ compute: f }); ``` diff --git a/apps/typegpu-docs/src/examples/tests/uniformity/constants.ts b/apps/typegpu-docs/src/examples/tests/uniformity/constants.ts index 82bccef947..38157392ec 100644 --- a/apps/typegpu-docs/src/examples/tests/uniformity/constants.ts +++ b/apps/typegpu-docs/src/examples/tests/uniformity/constants.ts @@ -1,6 +1,5 @@ -import { PRNG } from './prngs.ts'; - export const gridSizes = [8, 16, 32, 64, 128, 256, 512, 1024]; export const initialGridSize = gridSizes[4]; -export const initialPRNG = PRNG.BPETER; -export const prngs: PRNG[] = Object.values(PRNG); +export const samplesPerThread = [1, 8, 16, 64, 256, 1024, 131072, 262144]; +export const initialSamplesPerThread = samplesPerThread[0]; +export const initialTakeAverage = false; diff --git a/apps/typegpu-docs/src/examples/tests/uniformity/index.ts b/apps/typegpu-docs/src/examples/tests/uniformity/index.ts index 62267e1c4a..4e8d75c6d7 100644 --- a/apps/typegpu-docs/src/examples/tests/uniformity/index.ts +++ b/apps/typegpu-docs/src/examples/tests/uniformity/index.ts @@ -1,90 +1,162 @@ import { randf, randomGeneratorSlot } from '@typegpu/noise'; -import tgpu, { common, d, std, type TgpuRenderPipeline } from 'typegpu'; +import tgpu, { common, d, std, type TgpuGuardedComputePipeline } from 'typegpu'; import * as c from './constants.ts'; -import { getPRNG, type PRNG } from './prngs.ts'; +import { initialPRNG, prngKeys, prngs, type PRNGKey } from './prngs.ts'; import { defineControls } from '../../common/defineControls.ts'; -const root = await tgpu.init(); +const root = await tgpu.init({ device: { requiredFeatures: ['timestamp-query'] } }); const canvas = document.querySelector('canvas') as HTMLCanvasElement; const context = root.configureContext({ canvas, alphaMode: 'premultiplied' }); const presentationFormat = navigator.gpu.getPreferredCanvasFormat(); -const gridSizeUniform = root.createUniform(d.f32, c.initialGridSize); -const canvasRatioUniform = root.createUniform(d.f32, canvas.width / canvas.height); +const Config = d.struct({ + gridSize: d.f32, + canvasRatio: d.f32, + samplesPerThread: d.u32, + takeAverage: d.u32, +}); -const fragmentShader = tgpu.fragmentFn({ - in: { uv: d.vec2f }, - out: d.vec4f, -})((input) => { - 'use gpu'; - const uv = ((input.uv + 1) / 2) * d.vec2f(canvasRatioUniform.$, 1); - const gridedUV = std.floor(uv * gridSizeUniform.$); +const configUniform = root.createUniform(Config, { + gridSize: c.initialGridSize, + canvasRatio: canvas.width / canvas.height, + samplesPerThread: c.initialSamplesPerThread, + takeAverage: d.u32(c.initialTakeAverage), +}); - randf.seed2(gridedUV); +const layouts = { + compute: tgpu.bindGroupLayout({ + texture: { storageTexture: d.textureStorage2d('r32float', 'write-only') }, + }), + display: tgpu.bindGroupLayout({ + texture: { storageTexture: d.textureStorage2d('r32float', 'read-only') }, + }), +}; - return d.vec4f(d.vec3f(randf.sample()), 1); +const bindGroups = Object.fromEntries( + c.gridSizes.map((size) => { + const texture = root + .createTexture({ size: [size, size], format: 'r32float' }) + .$usage('storage', 'sampled'); + return [ + size, + { + compute: root.createBindGroup(layouts.compute, { texture }), + display: root.createBindGroup(layouts.display, { texture }), + }, + ]; + }), +); + +const displayPipeline = root.createRenderPipeline({ + vertex: common.fullScreenTriangle, + fragment: ({ uv }) => { + 'use gpu'; + const adjustedUv = uv * d.vec2f(configUniform.$.canvasRatio, 1); + const gridSize = configUniform.$.gridSize; + const coords = d.vec2u(std.floor(adjustedUv * gridSize)); + const value = std.textureLoad(layouts.display.$.texture, coords).r; + return d.vec4f(d.vec3f(value), 1); + }, + targets: { format: presentationFormat }, }); -const pipelineCache = new Map>(); -let prng: PRNG = c.initialPRNG; +const computeFn = (x: number, y: number) => { + 'use gpu'; + const gridSize = configUniform.$.gridSize; -const redraw = () => { - let pipeline = pipelineCache.get(prng); + if (!randomGeneratorSlot.$.seed2) { + randf.seed(d.f32(x + 1) * gridSize + d.f32(y + 1)); + } else { + randf.seed2(d.vec2f(x, y) + 1); + } + + let i = d.u32(0); + const samplesPerThread = configUniform.$.samplesPerThread; + let samples = d.f32(0); + while (i < samplesPerThread - 1) { + samples += randf.sample(); + i += 1; + } + + let result = randf.sample(); + if (configUniform.$.takeAverage === 1) { + result = (result + samples) / samplesPerThread; + } + + std.textureStore(layouts.compute.$.texture, d.vec2u(x, y), d.vec4f(result, 0, 0, 0)); +}; + +const computePipelineCache = new Map>(); +const getComputePipeline = (key: PRNGKey) => { + let pipeline = computePipelineCache.get(key); if (!pipeline) { - pipeline = root.with(randomGeneratorSlot, getPRNG(prng)).createRenderPipeline({ - vertex: common.fullScreenTriangle, - fragment: fragmentShader, - targets: { format: presentationFormat }, - }); - pipelineCache.set(prng, pipeline); + pipeline = root + .with(randomGeneratorSlot, prngs[key].generator) + .createGuardedComputePipeline(computeFn) + .withPerformanceCallback((start, end) => { + console.log(`[${key}] - ${Number(end - start) / 1_000_000} ms.`); + }); + computePipelineCache.set(key, pipeline); } + return pipeline; +}; - pipeline.withColorAttachment({ view: context }).draw(3); +let prng = initialPRNG; +let gridSize = c.initialGridSize; + +const redraw = () => { + getComputePipeline(prng).with(bindGroups[gridSize].compute).dispatchThreads(gridSize, gridSize); + displayPipeline.withColorAttachment({ view: context }).with(bindGroups[gridSize].display).draw(3); }; // #region Example controls & Cleanup export const controls = defineControls({ PRNG: { - initial: c.initialPRNG, - options: c.prngs, + initial: initialPRNG, + options: prngKeys, onSelectChange: (value) => { prng = value; redraw(); }, }, + 'Samples per thread': { + initial: c.initialSamplesPerThread, + options: c.samplesPerThread, + onSelectChange: (value) => { + configUniform.writePartial({ samplesPerThread: value }); + redraw(); + }, + }, + 'Take Average': { + initial: c.initialTakeAverage, + onToggleChange: (value) => { + configUniform.writePartial({ takeAverage: d.u32(value) }); + redraw(); + }, + }, 'Grid Size': { initial: c.initialGridSize, options: c.gridSizes, onSelectChange: (value) => { - gridSizeUniform.write(value); + gridSize = value; + configUniform.writePartial({ gridSize }); redraw(); }, }, + // this is the only place where some niche prngs are tested 'Test Resolution': import.meta.env.DEV && { onButtonClick: () => { - const namespace = tgpu['~unstable'].namespace(); - c.prngs - .map((prng) => - tgpu.resolve( - [ - root.with(randomGeneratorSlot, getPRNG(prng)).createRenderPipeline({ - vertex: common.fullScreenTriangle, - fragment: fragmentShader, - targets: { format: presentationFormat }, - }), - ], - { names: namespace }, - ), - ) - .map((r) => root.device.createShaderModule({ code: r })); + prngKeys + .map((key) => tgpu.resolve([getComputePipeline(key).pipeline])) + .forEach((r) => root.device.createShaderModule({ code: r })); }, }, }); const resizeObserver = new ResizeObserver(() => { - canvasRatioUniform.write(canvas.width / canvas.height); + configUniform.writePartial({ canvasRatio: canvas.width / canvas.height }); redraw(); }); resizeObserver.observe(canvas); diff --git a/apps/typegpu-docs/src/examples/tests/uniformity/lcg.ts b/apps/typegpu-docs/src/examples/tests/uniformity/lcg.ts deleted file mode 100644 index f6bffe867c..0000000000 --- a/apps/typegpu-docs/src/examples/tests/uniformity/lcg.ts +++ /dev/null @@ -1,28 +0,0 @@ -import type { StatefulGenerator } from '@typegpu/noise'; -import tgpu, { d, std } from 'typegpu'; - -export const LCG: StatefulGenerator = (() => { - const seed = tgpu.privateVar(d.u32); - - const u32To01Float = tgpu.fn( - [d.u32], - d.f32, - )((value) => { - const mantissa = value >> 9; - const bits = 0x3f800000 | mantissa; - const f = std.bitcastU32toF32(bits); - return f - 1; - }); - - return { - seed2: (value: d.v2f) => { - 'use gpu'; - seed.$ = d.u32(value.x * std.pow(32, 3) + value.y * std.pow(32, 2)); - }, - sample: () => { - 'use gpu'; - seed.$ = seed.$ * 1664525 + 1013904223; // % 2 ^ 32 - return u32To01Float(seed.$); - }, - }; -})(); diff --git a/apps/typegpu-docs/src/examples/tests/uniformity/prngs.ts b/apps/typegpu-docs/src/examples/tests/uniformity/prngs.ts index f176c77285..d43b1cc09b 100644 --- a/apps/typegpu-docs/src/examples/tests/uniformity/prngs.ts +++ b/apps/typegpu-docs/src/examples/tests/uniformity/prngs.ts @@ -1,17 +1,17 @@ -import { BPETER, type StatefulGenerator } from '@typegpu/noise'; +import { BPETER, LCG32, XOROSHIRO64STARSTAR, type StatefulGenerator } from '@typegpu/noise'; -import { LCG } from './lcg.ts'; +interface PRNGOptions { + name: string; + generator: StatefulGenerator; +} -export const PRNG = { - BPETER: 'bpeter (default)', - LCG: 'lcg', -} as const; +export const prngs = { + bpeter: { name: 'bpeter (default)', generator: BPETER }, + lcg32: { name: 'lcg32', generator: LCG32 }, + xoroshiro64: { name: 'xoroshiro64', generator: XOROSHIRO64STARSTAR }, +} as const satisfies Record; -export type PRNG = (typeof PRNG)[keyof typeof PRNG]; +export type PRNGKey = keyof typeof prngs; -const PRNG_MAP = { - [PRNG.BPETER]: BPETER, - [PRNG.LCG]: LCG, -}; - -export const getPRNG = (prng: PRNG): StatefulGenerator => PRNG_MAP[prng]; +export const prngKeys = Object.keys(prngs) as PRNGKey[]; +export const initialPRNG: PRNGKey = prngKeys[0]; diff --git a/apps/typegpu-docs/tests/individual-example-tests/uniformity.test.ts b/apps/typegpu-docs/tests/individual-example-tests/uniformity.test.ts index 11a61fe114..f8fdae5879 100644 --- a/apps/typegpu-docs/tests/individual-example-tests/uniformity.test.ts +++ b/apps/typegpu-docs/tests/individual-example-tests/uniformity.test.ts @@ -17,27 +17,22 @@ describe('uniformity test example', () => { name: 'uniformity', setupMocks: mockResizeObserver, controlTriggers: ['Test Resolution'], - expectedCalls: 2, + expectedCalls: 3, }, device, ); expect(shaderCodes).toMatchInlineSnapshot(` - "struct fullScreenTriangle_Output { - @builtin(position) pos: vec4f, - @location(0) uv: vec2f, - } - - @vertex fn fullScreenTriangle(@builtin(vertex_index) vertexIndex: u32) -> fullScreenTriangle_Output { - const pos = array(vec2f(-1, -1), vec2f(3, -1), vec2f(-1, 3)); - const uv = array(vec2f(0, 1), vec2f(2, 1), vec2f(0, -1)); + "@group(0) @binding(0) var sizeUniform: vec3u; - return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]); + struct Config { + gridSize: f32, + canvasRatio: f32, + samplesPerThread: u32, + takeAverage: u32, } - @group(0) @binding(0) var canvasRatioUniform: f32; - - @group(0) @binding(1) var gridSizeUniform: f32; + @group(0) @binding(1) var configUniform: Config; var seed: vec2f; @@ -61,52 +56,199 @@ describe('uniformity test example', () => { return sample(); } - struct fragmentShader_Input { - @location(0) uv: vec2f, + @group(1) @binding(0) var texture: texture_storage_2d; + + fn computeFn(x: u32, y: u32, _arg_2: u32) { + let gridSize = configUniform.gridSize; + { + randSeed2((vec2f(f32(x), f32(y)) + 1f)); + } + var i = 0u; + let samplesPerThread = configUniform.samplesPerThread; + var samples = 0f; + while ((i < (samplesPerThread - 1u))) { + samples += randFloat01(); + i += 1u; + } + var result = randFloat01(); + if ((configUniform.takeAverage == 1u)) { + result = ((result + samples) / f32(samplesPerThread)); + } + textureStore(texture, vec2u(x, y), vec4f(result, 0f, 0f, 0f)); + } + + @compute @workgroup_size(16, 16, 1) fn mainCompute(@builtin(global_invocation_id) id: vec3u) { + if (any(id >= sizeUniform)) { + return; + } + computeFn(id.x, id.y, id.z); + } + + @group(0) @binding(0) var sizeUniform: vec3u; + + struct Config { + gridSize: f32, + canvasRatio: f32, + samplesPerThread: u32, + takeAverage: u32, } - @fragment fn fragmentShader(_arg_0: fragmentShader_Input) -> @location(0) vec4f { - var uv = (((_arg_0.uv + 1f) / 2f) * vec2f(canvasRatioUniform, 1f)); - var gridedUV = floor((uv * gridSizeUniform)); - randSeed2(gridedUV); - return vec4f(vec3f(randFloat01()), 1f); + @group(0) @binding(1) var configUniform: Config; + + fn hash(value: u32) -> u32 { + var x = (value ^ (value >> 17u)); + x *= 3982152891u; + x ^= (x >> 11u); + x *= 2890668881u; + x ^= (x >> 15u); + x *= 830770091u; + x ^= (x >> 14u); + return x; } - var seed_1: u32; + var seed: u32; - fn seed2_1(value: vec2f) { - seed_1 = u32(((value.x * 32768f) + (value.y * 1024f))); + fn seed_1(value: f32) { + seed = hash(u32(value)); } - fn randSeed2_1(seed_1: vec2f) { - seed2_1(seed_1); + fn randSeed(seed: f32) { + seed_1(seed); } - fn u32To01Float(value: u32) -> f32 { + fn u32To01F32(value: u32) -> f32 { let mantissa = (value >> 9u); let bits = (1065353216u | mantissa); let f = bitcast(bits); return (f - 1f); } - fn sample_1() -> f32 { - seed_1 = ((seed_1 * 1664525u) + 1013904223u); - return u32To01Float(seed_1); + fn sample() -> f32 { + seed = ((1664525u * seed) + 1013904223u); + return u32To01F32(seed); } - fn randFloat01_1() -> f32 { - return sample_1(); + fn randFloat01() -> f32 { + return sample(); + } + + @group(1) @binding(0) var texture: texture_storage_2d; + + fn computeFn(x: u32, y: u32, _arg_2: u32) { + let gridSize = configUniform.gridSize; + { + randSeed(((f32((x + 1u)) * gridSize) + f32((y + 1u)))); + } + var i = 0u; + let samplesPerThread = configUniform.samplesPerThread; + var samples = 0f; + while ((i < (samplesPerThread - 1u))) { + samples += randFloat01(); + i += 1u; + } + var result = randFloat01(); + if ((configUniform.takeAverage == 1u)) { + result = ((result + samples) / f32(samplesPerThread)); + } + textureStore(texture, vec2u(x, y), vec4f(result, 0f, 0f, 0f)); + } + + @compute @workgroup_size(16, 16, 1) fn mainCompute(@builtin(global_invocation_id) id: vec3u) { + if (any(id >= sizeUniform)) { + return; + } + computeFn(id.x, id.y, id.z); + } + + @group(0) @binding(0) var sizeUniform: vec3u; + + struct Config { + gridSize: f32, + canvasRatio: f32, + samplesPerThread: u32, + takeAverage: u32, + } + + @group(0) @binding(1) var configUniform: Config; + + fn hash(value: u32) -> u32 { + var x = (value ^ (value >> 17u)); + x *= 3982152891u; + x ^= (x >> 11u); + x *= 2890668881u; + x ^= (x >> 15u); + x *= 830770091u; + x ^= (x >> 14u); + return x; + } + + fn rotl(x: u32, k: u32) -> u32 { + return ((x << k) | (x >> (32u - k))); + } + + var seed: vec2u; + + fn seed2(value: vec2f) { + let hx = hash((u32(value.x) ^ 2135587861u)); + let hy = hash((u32(value.y) ^ 2654435769u)); + seed = vec2u(hash((hx ^ hy)), hash((rotl(hx, 16u) ^ hy))); + } + + fn randSeed2(seed: vec2f) { + seed2(seed); + } + + fn next() -> u32 { + let s0 = seed[0i]; + var s1 = seed[1i]; + s1 ^= s0; + seed[0i] = ((rotl(s0, 26u) ^ s1) ^ (s1 << 9u)); + seed[1i] = rotl(s1, 13u); + return (rotl((seed[0i] * 2654435771u), 5u) * 5u); + } + + fn u32To01F32(value: u32) -> f32 { + let mantissa = (value >> 9u); + let bits = (1065353216u | mantissa); + let f = bitcast(bits); + return (f - 1f); + } + + fn sample() -> f32 { + let r = next(); + return u32To01F32(r); + } + + fn randFloat01() -> f32 { + return sample(); } - struct fragmentShader_Input_1 { - @location(0) uv: vec2f, + @group(1) @binding(0) var texture: texture_storage_2d; + + fn computeFn(x: u32, y: u32, _arg_2: u32) { + let gridSize = configUniform.gridSize; + { + randSeed2((vec2f(f32(x), f32(y)) + 1f)); + } + var i = 0u; + let samplesPerThread = configUniform.samplesPerThread; + var samples = 0f; + while ((i < (samplesPerThread - 1u))) { + samples += randFloat01(); + i += 1u; + } + var result = randFloat01(); + if ((configUniform.takeAverage == 1u)) { + result = ((result + samples) / f32(samplesPerThread)); + } + textureStore(texture, vec2u(x, y), vec4f(result, 0f, 0f, 0f)); } - @fragment fn fragmentShader_1(_arg_0: fragmentShader_Input_1) -> @location(0) vec4f { - var uv = (((_arg_0.uv + 1f) / 2f) * vec2f(canvasRatioUniform, 1f)); - var gridedUV = floor((uv * gridSizeUniform)); - randSeed2_1(gridedUV); - return vec4f(vec3f(randFloat01_1()), 1f); + @compute @workgroup_size(16, 16, 1) fn mainCompute(@builtin(global_invocation_id) id: vec3u) { + if (any(id >= sizeUniform)) { + return; + } + computeFn(id.x, id.y, id.z); }" `); }); diff --git a/packages/typegpu-noise/src/generator.ts b/packages/typegpu-noise/src/generator.ts index fa9d1153ef..59039becec 100644 --- a/packages/typegpu-noise/src/generator.ts +++ b/packages/typegpu-noise/src/generator.ts @@ -1,5 +1,5 @@ -import tgpu, { d, type TgpuFnShell, type TgpuSlot } from 'typegpu'; -import { cos, dot, fract } from 'typegpu/std'; +import tgpu, { d, type TgpuFnShell, type TgpuSlot, std } from 'typegpu'; +import { hash, rotl, u32To01F32 } from './utils.ts'; export interface StatefulGenerator { seed?: (seed: number) => void; @@ -39,15 +39,95 @@ export const BPETER: StatefulGenerator = (() => { sample: randomGeneratorShell(() => { 'use gpu'; - const a = dot(seed.$, d.vec2f(23.14077926, 232.61690225)); - const b = dot(seed.$, d.vec2f(54.47856553, 345.84153136)); - seed.$.x = fract(cos(a) * 136.8168); - seed.$.y = fract(cos(b) * 534.7645); + const a = std.dot(seed.$, d.vec2f(23.14077926, 232.61690225)); + const b = std.dot(seed.$, d.vec2f(54.47856553, 345.84153136)); + seed.$.x = std.fract(std.cos(a) * 136.8168); + seed.$.y = std.fract(std.cos(b) * 534.7645); return seed.$.y; }).$name('sample'), }; })(); +/** + * Incorporated from https://github.com/chaos-matters/chaos-master + * by deluksic and Komediruzecki + */ +export const XOROSHIRO64STARSTAR: StatefulGenerator = (() => { + const seed = tgpu.privateVar(d.vec2u); + + const next = tgpu.fn( + [], + d.u32, + )(() => { + const s0 = seed.$[0]; + let s1 = seed.$[1]; + s1 ^= s0; + seed.$[0] = rotl(s0, 26) ^ s1 ^ (s1 << 9); + seed.$[1] = rotl(s1, 13); + return rotl(seed.$[0] * 0x9e3779bb, 5) * 5; + }); + + return { + seed2: tgpu.fn([d.vec2f])((value) => { + const u32Value = std.bitcastF32toU32(value); + const hx = hash(u32Value.x ^ 0x4ab57dfb); + const hy = hash(u32Value.y ^ 0xacdeda47); + seed.$ = d.vec2u(hash(hx ^ hy), hash(rotl(hx, 16) ^ hy)); + }), + + seed3: tgpu.fn([d.vec3f])((value) => { + const u32Value = std.bitcastF32toU32(value); + const hx = hash(u32Value.x ^ 0x4ab57dfb); + const hy = hash(u32Value.y ^ 0xacdeda47); + const hz = hash(u32Value.z ^ 0xbca0294b); + seed.$ = d.vec2u(hash(hx ^ rotl(hz, 16)), hash(rotl(hy, 16) ^ hz)); + }), + + sample: randomGeneratorShell(() => { + 'use gpu'; + const r = next(); + return u32To01F32(r); + }).$name('sample'), + }; +})(); + +/** + * Naive Linear Congruential Generator (LCG) with 32 bits state + */ +export const LCG32: StatefulGenerator = (() => { + const seed = tgpu.privateVar(d.u32); + + const multiplier = 0x19660d; + const increment = 0x3c6ef35f; + + return { + seed: tgpu.fn([d.f32])((value) => { + seed.$ = hash(std.bitcastF32toU32(value)) ^ 0x4ab57dfb; + }), + + seed2: tgpu.fn([d.vec2f])((value) => { + const u32Value = std.bitcastF32toU32(value); + const hx = hash(u32Value.x ^ 0x4ab57dfb); + const hy = hash(u32Value.y ^ 0xacdeda47); + seed.$ = hash(hx ^ rotl(hy, 16)); + }), + + seed3: tgpu.fn([d.vec3f])((value) => { + const u32Value = std.bitcastF32toU32(value); + const hx = hash(u32Value.x ^ 0x4ab57dfb); + const hy = hash(u32Value.y ^ 0xacdeda47); + const hz = hash(u32Value.z ^ 0xbca0294b); + seed.$ = hash(hash(hx ^ rotl(hy, 16)) ^ hz); + }), + + sample: randomGeneratorShell(() => { + 'use gpu'; + seed.$ = multiplier * seed.$ + increment; // % 2 ^ 32 + return u32To01F32(seed.$); + }).$name('sample'), + }; +})(); + // The default (Can change between releases to improve uniformity). export const DefaultGenerator: StatefulGenerator = BPETER; diff --git a/packages/typegpu-noise/src/index.ts b/packages/typegpu-noise/src/index.ts index 97dc584cf3..b96c103267 100644 --- a/packages/typegpu-noise/src/index.ts +++ b/packages/typegpu-noise/src/index.ts @@ -145,11 +145,15 @@ export { BPETER, // The default (Can change between releases to improve uniformity). DefaultGenerator, + LCG32, + XOROSHIRO64STARSTAR, // --- randomGeneratorShell, randomGeneratorSlot, type StatefulGenerator, } from './generator.ts'; +export { hash, u32To01F32 } from './utils.ts'; + export * as perlin2d from './perlin-2d/index.ts'; export * as perlin3d from './perlin-3d/index.ts'; diff --git a/packages/typegpu-noise/src/random.ts b/packages/typegpu-noise/src/random.ts index e1ddd472b2..dfbf0f4b4d 100644 --- a/packages/typegpu-noise/src/random.ts +++ b/packages/typegpu-noise/src/random.ts @@ -9,8 +9,6 @@ const warnIfNotProvided = tgpu.comptime((seedFnName: keyof typeof randomGenerato if (!randomGeneratorSlot.$[seedFnName]) { console.warn(`Called \`randf.${seedFnName}\`, but it wasn't provided`); } - - return undefined; }); export const randSeed = tgpu.fn([d.f32])((seed) => { diff --git a/packages/typegpu-noise/src/utils.ts b/packages/typegpu-noise/src/utils.ts index 9e90781916..c5223b12d3 100644 --- a/packages/typegpu-noise/src/utils.ts +++ b/packages/typegpu-noise/src/utils.ts @@ -1,4 +1,4 @@ -import type { d } from 'typegpu'; +import tgpu, { d, std } from 'typegpu'; export type Prettify = { [K in keyof T]: T[K]; @@ -28,3 +28,47 @@ export function quinticDerivative(t: d.vecBase): d.vecBase { 'use gpu'; return 30 * t * t * (t * (t - 2) + 1); } + +/** + * Left circular shif of x by k positions. + */ +export const rotl = tgpu.fn( + [d.u32, d.u32], + d.u32, +)((x, k) => { + return (x << k) | (x >> (32 - k)); +}); + +/** + * Converts `u32` to `f32` value in the range `[0.0, 1.0)`. + */ +export const u32To01F32 = tgpu.fn( + [d.u32], + d.f32, +)((value) => { + const mantissa = value & 0x007fffff; + const bits = 0x3f800000 | mantissa; + const f = std.bitcastU32toF32(bits); + return f - 1; +}); + +/** + * Simple hashing function to scramble the seed. + * Keep in mind that `hash(0) -> 0`. + * + * Incorporated from https://github.com/chaos-matters/chaos-master + * by deluksic and Komediruzecki + */ +export const hash = tgpu.fn( + [d.u32], + d.u32, +)((value) => { + let x = value ^ (value >> 17); + x *= d.u32(0xed5ad4bb); + x ^= x >> 11; + x *= d.u32(0xac4c1b51); + x ^= x >> 15; + x *= d.u32(0x31848bab); + x ^= x >> 14; + return x; +}); diff --git a/packages/typegpu/src/core/root/init.ts b/packages/typegpu/src/core/root/init.ts index bc84ad408b..0ef85a27c9 100644 --- a/packages/typegpu/src/core/root/init.ts +++ b/packages/typegpu/src/core/root/init.ts @@ -152,6 +152,30 @@ export class TgpuGuardedComputePipelineImpl< ); } + withPerformanceCallback( + callback: (start: bigint, end: bigint) => void | Promise, + ): TgpuGuardedComputePipeline { + return new TgpuGuardedComputePipelineImpl( + this.#root, + this.#pipeline.withPerformanceCallback(callback), + this.#sizeUniform, + this.#workgroupSize, + ); + } + + withTimestampWrites(options: { + querySet: TgpuQuerySet<'timestamp'> | GPUQuerySet; + beginningOfPassWriteIndex?: number; + endOfPassWriteIndex?: number; + }): TgpuGuardedComputePipeline { + return new TgpuGuardedComputePipelineImpl( + this.#root, + this.#pipeline.withTimestampWrites(options), + this.#sizeUniform, + this.#workgroupSize, + ); + } + dispatchThreads(...threads: TArgs): void { const sanitizedSize = toVec3(threads); const workgroupCount = ceil(vec3f(sanitizedSize).div(vec3f(this.#workgroupSize))); diff --git a/packages/typegpu/src/core/root/rootTypes.ts b/packages/typegpu/src/core/root/rootTypes.ts index e6c5c8bf66..afa6cb08df 100644 --- a/packages/typegpu/src/core/root/rootTypes.ts +++ b/packages/typegpu/src/core/root/rootTypes.ts @@ -1,6 +1,7 @@ import type { AnyComputeBuiltin, AnyFragmentInputBuiltin, OmitBuiltins } from '../../builtin.ts'; import type { TgpuQuerySet } from '../../core/querySet/querySet.ts'; import type { AnyData, Disarray, UndecorateRecord } from '../../data/dataTypes.ts'; +import type { InstanceToSchema } from '../../data/instanceToSchema.ts'; import type { WgslComparisonSamplerProps, WgslSamplerProps } from '../../data/sampler.ts'; import type { AnyWgslData, @@ -12,6 +13,7 @@ import type { Void, WgslArray, } from '../../data/wgslTypes.ts'; +import type { TgpuNamable } from '../../shared/meta.ts'; import type { ExtractInvalidSchemaError, InferGPURecord, @@ -33,7 +35,13 @@ import type { ShaderGenerator } from '../../tgsl/shaderGenerator.ts'; import type { Unwrapper } from '../../unwrapper.ts'; import type { TgpuBuffer, VertexFlag } from '../buffer/buffer.ts'; import type { TgpuMutable, TgpuReadonly, TgpuUniform } from '../buffer/bufferShorthand.ts'; -import type { TgpuFixedComparisonSampler, TgpuFixedSampler } from '../sampler/sampler.ts'; +import type { + AnyAutoCustoms, + AutoFragmentIn, + AutoFragmentOut, + AutoVertexIn, + AutoVertexOut, +} from '../function/autoIO.ts'; import type { IORecord } from '../function/fnTypes.ts'; import type { FragmentInConstrained, @@ -44,6 +52,7 @@ import type { import type { TgpuVertexFn } from '../function/tgpuVertexFn.ts'; import type { TgpuComputePipeline } from '../pipeline/computePipeline.ts'; import type { FragmentOutToTargets, TgpuRenderPipeline } from '../pipeline/renderPipeline.ts'; +import type { TgpuFixedComparisonSampler, TgpuFixedSampler } from '../sampler/sampler.ts'; import type { Eventual, TgpuAccessor, TgpuMutableAccessor, TgpuSlot } from '../slot/slotTypes.ts'; import type { TgpuTexture } from '../texture/texture.ts'; import type { @@ -52,15 +61,6 @@ import type { } from '../vertexLayout/vertexAttribute.ts'; import type { TgpuVertexLayout } from '../vertexLayout/vertexLayout.ts'; import type { TgpuComputeFn } from './../function/tgpuComputeFn.ts'; -import type { TgpuNamable } from '../../shared/meta.ts'; -import type { - AnyAutoCustoms, - AutoFragmentIn, - AutoFragmentOut, - AutoVertexIn, - AutoVertexOut, -} from '../function/autoIO.ts'; -import type { InstanceToSchema } from '../../data/instanceToSchema.ts'; // ---------- // Public API @@ -80,6 +80,24 @@ export interface TgpuGuardedComputePipeline e */ with(encoder: GPUCommandEncoder): TgpuGuardedComputePipeline; + /** + * Returns a pipeline wrapper with the given performance callback attached. + * Analogous to `TgpuComputePipeline.withPerformanceCallback(callback)`. + */ + withPerformanceCallback( + callback: (start: bigint, end: bigint) => void | Promise, + ): TgpuGuardedComputePipeline; + + /** + * Returns a pipeline wrapper with the given timestamp writes configuration. + * Analogous to `TgpuComputePipeline.withTimestampWrites(options)`. + */ + withTimestampWrites(options: { + querySet: TgpuQuerySet<'timestamp'> | GPUQuerySet; + beginningOfPassWriteIndex?: number; + endOfPassWriteIndex?: number; + }): TgpuGuardedComputePipeline; + /** * Dispatches the pipeline. * Unlike `TgpuComputePipeline.dispatchWorkgroups()`, this method takes in the @@ -378,7 +396,7 @@ export interface WithBinding extends Withable { /** * Creates a compute pipeline that executes the given callback in an exact number of threads. - * This is different from `withCompute(...).createPipeline()` in that it does a bounds check on the + * This is different from `createComputePipeline()` in that it does a bounds check on the * thread id, where as regular pipelines do not and work in units of workgroups. * * @param callback A function converted to WGSL and executed on the GPU. diff --git a/packages/typegpu/src/data/numberOps.ts b/packages/typegpu/src/data/numberOps.ts index 55ad1b4787..85cf34e82c 100644 --- a/packages/typegpu/src/data/numberOps.ts +++ b/packages/typegpu/src/data/numberOps.ts @@ -27,3 +27,9 @@ export function bitcastU32toI32Impl(n: number): number { dataView.setUint32(0, n, true); return dataView.getInt32(0, true); } + +export function bitcastF32toU32Impl(n: number): number { + const dataView = new DataView(new ArrayBuffer(4)); + dataView.setFloat32(0, n, true); + return dataView.getUint32(0, true); +} diff --git a/packages/typegpu/src/data/vectorOps.ts b/packages/typegpu/src/data/vectorOps.ts index a05d1d1af2..3bd64a6aa0 100644 --- a/packages/typegpu/src/data/vectorOps.ts +++ b/packages/typegpu/src/data/vectorOps.ts @@ -1,5 +1,6 @@ import { mat2x2f, mat3x3f, mat4x4f } from './matrix.ts'; import { + bitcastF32toU32Impl, bitcastU32toF32Impl, bitcastU32toI32Impl, clamp, @@ -1162,4 +1163,22 @@ export const VectorOps = { v: T, ) => T extends wgsl.v2u ? wgsl.v2i : T extends wgsl.v3u ? wgsl.v3i : wgsl.v4i >, + + bitcastF32toU32: { + vec2f: (n: wgsl.v2f) => vec2u(bitcastF32toU32Impl(n.x), bitcastF32toU32Impl(n.y)), + vec3f: (n: wgsl.v3f) => + vec3u(bitcastF32toU32Impl(n.x), bitcastF32toU32Impl(n.y), bitcastF32toU32Impl(n.z)), + vec4f: (n: wgsl.v4f) => + vec4u( + bitcastF32toU32Impl(n.x), + bitcastF32toU32Impl(n.y), + bitcastF32toU32Impl(n.z), + bitcastF32toU32Impl(n.w), + ), + } as Record< + VecKind, + ( + v: T, + ) => T extends wgsl.v2f ? wgsl.v2u : T extends wgsl.v3f ? wgsl.v3u : wgsl.v4u + >, }; diff --git a/packages/typegpu/src/std/bitcast.ts b/packages/typegpu/src/std/bitcast.ts index 4039858dd9..3b66a2ca61 100644 --- a/packages/typegpu/src/std/bitcast.ts +++ b/packages/typegpu/src/std/bitcast.ts @@ -1,9 +1,13 @@ import { dualImpl } from '../core/function/dualImpl.ts'; import { stitch } from '../core/resolve/stitch.ts'; -import { bitcastU32toF32Impl, bitcastU32toI32Impl } from '../data/numberOps.ts'; +import { + bitcastF32toU32Impl, + bitcastU32toF32Impl, + bitcastU32toI32Impl, +} from '../data/numberOps.ts'; import { f32, i32, u32 } from '../data/numeric.ts'; import { isVec } from '../data/wgslTypes.ts'; -import { vec2f, vec2i, vec3f, vec3i, vec4f, vec4i } from '../data/vector.ts'; +import { vec2f, vec2i, vec2u, vec3f, vec3i, vec3u, vec4f, vec4i, vec4u } from '../data/vector.ts'; import { VectorOps } from '../data/vectorOps.ts'; import type { v2f, v2i, v2u, v3f, v3i, v3u, v4f, v4i, v4u } from '../data/wgslTypes.ts'; import { unify } from '../tgsl/conversion.ts'; @@ -65,3 +69,36 @@ export const bitcastU32toI32 = dualImpl({ }; }, }); + +export type BitcastF32toU32Overload = ((value: number) => number) & + ((value: v2f) => v2u) & + ((value: v3f) => v3u) & + ((value: v4f) => v4u); + +export const bitcastF32toU32 = dualImpl({ + name: 'bitcastF32toU32', + normalImpl: ((value) => { + if (typeof value === 'number') { + return bitcastF32toU32Impl(value); + } + return VectorOps.bitcastF32toU32[value.kind](value); + }) as BitcastF32toU32Overload, + codegenImpl: (_ctx, [n]) => { + return isVec(n.dataType) + ? stitch`bitcast(${n})` + : stitch`bitcast(${n})`; + }, + signature: (...arg) => { + const uargs = unify(arg, [f32]) ?? arg; + return { + argTypes: uargs, + returnType: isVec(uargs[0]) + ? uargs[0].type === 'vec2f' + ? vec2u + : uargs[0].type === 'vec3f' + ? vec3u + : vec4u + : u32, + }; + }, +}); diff --git a/packages/typegpu/src/std/index.ts b/packages/typegpu/src/std/index.ts index 185e4b3c08..515d7d2bca 100644 --- a/packages/typegpu/src/std/index.ts +++ b/packages/typegpu/src/std/index.ts @@ -184,6 +184,6 @@ export { export { extensionEnabled } from './extensions.ts'; -export { bitcastU32toF32, bitcastU32toI32 } from './bitcast.ts'; +export { bitcastU32toF32, bitcastU32toI32, bitcastF32toU32 } from './bitcast.ts'; export { range } from './range.ts'; diff --git a/packages/typegpu/tests/guardedComputePipeline.test.ts b/packages/typegpu/tests/guardedComputePipeline.test.ts index 065101fc82..58484d5e2e 100644 --- a/packages/typegpu/tests/guardedComputePipeline.test.ts +++ b/packages/typegpu/tests/guardedComputePipeline.test.ts @@ -1,4 +1,4 @@ -import { describe, expect } from 'vitest'; +import { describe, expect, vi } from 'vitest'; import { it } from 'typegpu-testing-utility'; import { getName } from '../src/shared/meta.ts'; import { bindGroupLayout } from '../src/tgpuBindGroupLayout.ts'; @@ -31,4 +31,34 @@ describe('TgpuGuardedComputePipeline', () => { expect(getName(pipeline)).toBe('myPipeline'); expect(getName(pipeline.pipeline)).toBe('myPipeline'); }); + + it('delegates `withPerformanceCallback` to the underlying pipeline', ({ root }) => { + const callback = vi.fn(); + const guarded = root.createGuardedComputePipeline(() => { + 'use gpu'; + }); + + const spy = vi.spyOn(guarded.pipeline, 'withPerformanceCallback'); + guarded.withPerformanceCallback(callback); + + expect(spy).toHaveBeenCalledWith(callback); + }); + + it('delegates `withTimestampWrites` to the underlying pipeline', ({ root }) => { + const querySet = root.createQuerySet('timestamp', 2); + const guarded = root.createGuardedComputePipeline(() => { + 'use gpu'; + }); + + const options = { + querySet, + beginningOfPassWriteIndex: 0, + endOfPassWriteIndex: 1, + }; + + const spy = vi.spyOn(guarded.pipeline, 'withTimestampWrites'); + guarded.withTimestampWrites(options); + + expect(spy).toHaveBeenCalledWith(options); + }); }); diff --git a/packages/typegpu/tests/std/bitcast.test.ts b/packages/typegpu/tests/std/bitcast.test.ts index dabcd382a4..a2a4c10e23 100644 --- a/packages/typegpu/tests/std/bitcast.test.ts +++ b/packages/typegpu/tests/std/bitcast.test.ts @@ -12,6 +12,9 @@ import { } from '../../src/data/vector.ts'; import tgpu, { d, std } from '../../src/index.js'; +// remember to pad with zeros to 8 hex symbols +const floatFromHex = (hex: string) => Buffer.from(hex, 'hex').readFloatBE(0); + describe('bitcast', () => { it('bitcastU32toF32', () => { // 1.0 in f32 @@ -37,6 +40,14 @@ describe('bitcast', () => { expect(i2).toBe(-2147483648); }); + it('bitcastF32toU32', () => { + const i1 = std.bitcastF32toU32(floatFromHex('00000001')); + expect(i1).toBe(1); + + const i2 = std.bitcastF32toU32(floatFromHex('7f800000')); + expect(i2).toBe(2139095040); + }); + it('bitcastU32toF32 vectors', () => { const v2 = vec2u(1065353216, 3212836864); // 1.0f, -1.0f const cast2 = std.bitcastU32toF32(v2); @@ -65,6 +76,25 @@ describe('bitcast', () => { expect(cast4).toEqual(vec4i(0, 1, -1, -2147483648)); }); + it('bitcastF32toU32 vectors', () => { + const v2 = vec2f(floatFromHex('7f800000'), floatFromHex('7fc00000')); // +inf, quiet nan + const cast2 = std.bitcastF32toU32(v2); + expect(cast2).toStrictEqual(vec2u(2139095040, 2143289344)); + + const v3 = vec3f(floatFromHex('ff800000'), floatFromHex('00000001'), floatFromHex('80000001')); + const cast3 = std.bitcastF32toU32(v3); + expect(cast3).toStrictEqual(vec3u(4286578688, 1, 2147483649)); + + const v4 = vec4f( + floatFromHex('84220925'), + floatFromHex('68800000'), + floatFromHex('48980780'), + floatFromHex('0000075a'), + ); + const cast4 = std.bitcastF32toU32(v4); + expect(cast4).toStrictEqual(vec4u(2216823077, 1753219072, 1217922944, 1882)); + }); + it('bitcastU32toF32 specials (NaN, infinities etc)', () => { // +0 const pz = std.bitcastU32toF32(0x00000000); @@ -128,6 +158,7 @@ describe('bitcast in shaders', () => { it('works for primitives', () => { const fnf32 = tgpu.fn([], d.f32)(() => std.bitcastU32toF32(1234)); const fni32 = tgpu.fn([], d.i32)(() => std.bitcastU32toI32(d.u32(2 ** 31))); + const fnu32 = tgpu.fn([d.f32], d.u32)((v) => std.bitcastF32toU32(v)); expect(tgpu.resolve([fnf32])).toMatchInlineSnapshot(` "fn fnf32() -> f32 { @@ -139,11 +170,17 @@ describe('bitcast in shaders', () => { return -2147483648i; }" `); + expect(tgpu.resolve([fnu32])).toMatchInlineSnapshot(` + "fn fnu32(v: f32) -> u32 { + return bitcast(v); + }" + `); }); it('works for vectors', () => { const fnvec4i = tgpu.fn([], d.vec4i)(() => std.bitcastU32toI32(vec4u(1, 2, 3, 4))); const fnvec4f = tgpu.fn([], d.vec4f)(() => std.bitcastU32toF32(vec4u(1, 2, 3, 4))); + const fnvec4u = tgpu.fn([d.vec4f], d.vec4u)((v) => std.bitcastF32toU32(v)); expect(tgpu.resolve([fnvec4i])).toMatchInlineSnapshot(` "fn fnvec4i() -> vec4i { @@ -155,5 +192,10 @@ describe('bitcast in shaders', () => { return vec4f(1.401298464324817e-45, 2.802596928649634e-45, 4.203895392974451e-45, 5.605193857299268e-45); }" `); + expect(tgpu.resolve([fnvec4u])).toMatchInlineSnapshot(` + "fn fnvec4u(v: vec4f) -> vec4u { + return bitcast(v); + }" + `); }); });