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 4e28e39aa6..1bbaeab418 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 deleted file mode 100644 index 82bccef947..0000000000 --- a/apps/typegpu-docs/src/examples/tests/uniformity/constants.ts +++ /dev/null @@ -1,6 +0,0 @@ -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); diff --git a/apps/typegpu-docs/src/examples/tests/uniformity/index.ts b/apps/typegpu-docs/src/examples/tests/uniformity/index.ts index 62267e1c4a..e37866009f 100644 --- a/apps/typegpu-docs/src/examples/tests/uniformity/index.ts +++ b/apps/typegpu-docs/src/examples/tests/uniformity/index.ts @@ -1,96 +1,296 @@ 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 { Camera, setupOrbitCamera } from '../../common/setup-orbit-camera.ts'; -import * as c from './constants.ts'; -import { getPRNG, type PRNG } from './prngs.ts'; +import { prngKeys, prngs, type PRNGKey } from './prngs.ts'; import { defineControls } from '../../common/defineControls.ts'; -const root = await tgpu.init(); +type Mode = '2d' | '3d'; + +const modes: Mode[] = ['2d', '3d']; +const gridSizes = [8, 16, 32, 64, 128, 256]; +const samplesPerThread = [1, 8, 16, 64, 256, 1024, 131072]; +const initialSamplesPerThread = samplesPerThread[0]; +const initialTakeAverage = false; +const initialMultiplier = 1; + +let mode = modes[1]; +let prng = prngKeys[0]; +let gridSize = gridSizes[2]; + +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, + samplesPerThread: d.i32, + takeAverage: d.i32, + multiplier: d.f32, + canvasRatio: d.f32, +}); + +const configUniform = root.createUniform(Config, { + gridSize, + samplesPerThread: initialSamplesPerThread, + takeAverage: d.i32(initialTakeAverage), + multiplier: initialMultiplier, + canvasRatio: canvas.width / canvas.height, +}); + +const layouts = { + compute: tgpu.bindGroupLayout({ + texture: { storageTexture: d.textureStorage3d('r32float', 'write-only') }, + }), + display: tgpu.bindGroupLayout({ + texture: { storageTexture: d.textureStorage3d('r32float', 'read-only') }, + }), +}; + +const bindGroups = Object.fromEntries( + gridSizes.map((size) => { + const texture = root + .createTexture({ size: [size, size, size], format: 'r32float', dimension: '3d' }) + .$usage('storage', 'sampled'); + return [ + size, + { + compute: root.createBindGroup(layouts.compute, { texture }), + display: root.createBindGroup(layouts.display, { texture }), + }, + ]; + }), +); -const fragmentShader = tgpu.fragmentFn({ - in: { uv: d.vec2f }, - out: d.vec4f, -})((input) => { +const modeSlot = tgpu.slot(); +const computeFn = tgpu.fn((x: number, y: number, z: number) => { 'use gpu'; - const uv = ((input.uv + 1) / 2) * d.vec2f(canvasRatioUniform.$, 1); - const gridedUV = std.floor(uv * gridSizeUniform.$); + const multiplier = configUniform.$.multiplier; + modeSlot.$ === 1 + ? randf.seed3((d.vec3f(x, y, z) - configUniform.$.gridSize / 2) * multiplier) + : randf.seed2((d.vec2f(x, y) - configUniform.$.gridSize / 2) * multiplier); - randf.seed2(gridedUV); + const samplesPerThread = configUniform.$.samplesPerThread; + const takeAverage = configUniform.$.takeAverage; + + let sum = d.f32(0); + for (let i = d.i32(0); i < samplesPerThread - 1; i++) { + sum += randf.sample(); + } - return d.vec4f(d.vec3f(randf.sample()), 1); + let result = randf.sample(); + result += sum * d.f32(takeAverage); + const denominator = d.f32(1 + (samplesPerThread - 1) * takeAverage); + result /= denominator; + + std.textureStore(layouts.compute.$.texture, d.vec3u(x, y, z), d.vec4f(result, 0, 0, 0)); }); -const pipelineCache = new Map>(); -let prng: PRNG = c.initialPRNG; +const computeFns = { + '2d': computeFn.with(modeSlot, 0), + '3d': computeFn.with(modeSlot, 1), +}; -const redraw = () => { - let pipeline = pipelineCache.get(prng); - if (!pipeline) { - pipeline = root.with(randomGeneratorSlot, getPRNG(prng)).createRenderPipeline({ - vertex: common.fullScreenTriangle, - fragment: fragmentShader, - targets: { format: presentationFormat }, - }); - pipelineCache.set(prng, pipeline); +const computePipelineCache = { + '2d': new Map>(), + '3d': new Map>(), +}; + +const getComputePipeline = (mode: Mode, key: PRNGKey) => { + const cache = computePipelineCache[mode]; + + let p = cache.get(key); + if (!p) { + p = root + .with(randomGeneratorSlot, prngs[key].generator) + .createGuardedComputePipeline(computeFns[mode]) + .withPerformanceCallback((start, end) => { + console.log(`[${key} ${mode}] ${Number(end - start) / 1_000_000} ms`); + }); + cache.set(key, p); } + return p; +}; + +const displayPipeline2d = 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, d.vec3u(coords, 0)).r; + return d.vec4f(d.vec3f(value), 1); + }, + targets: { format: presentationFormat }, +}); + +const cameraUniform = root.createUniform(Camera); +const BoxIntersection = d.struct({ tNear: d.f32, tFar: d.f32, hit: d.bool }); + +// based on: https://www.scratchapixel.com/lessons/3d-basic-rendering/minimal-ray-tracer-rendering-simple-shapes/ray-box-intersection.html +const getBoxIntersection = (rayOrigin: d.v3f, rayDir: d.v3f, boxMin: d.v3f, boxMax: d.v3f) => { + 'use gpu'; + const invDir = 1 / rayDir; + const t0 = (boxMin - rayOrigin) * invDir; + const t1 = (boxMax - rayOrigin) * invDir; + const tmin = std.min(t0, t1); + const tmax = std.max(t0, t1); + const tNear = std.max(tmin.x, tmin.y, tmin.z); + const tFar = std.min(tmax.x, tmax.y, tmax.z); + return BoxIntersection({ tNear, tFar, hit: tFar >= tNear }); +}; + +const STEPS = 64; +const displayPipeline3d = root.createRenderPipeline({ + vertex: common.fullScreenTriangle, + fragment: ({ uv }) => { + 'use gpu'; + const ndc = d.vec2f(uv.x * 2 - 1, 1 - uv.y * 2); + const invViewProj = cameraUniform.$.viewInverse * cameraUniform.$.projectionInverse; + const worldNear = invViewProj * d.vec4f(ndc, 0, 1); + const worldFar = invViewProj * d.vec4f(ndc, 1, 1); + const rayOrigin = worldNear.xyz / worldNear.w; + const rayDir = std.normalize(worldFar.xyz / worldFar.w - rayOrigin); + + const gridSize = configUniform.$.gridSize; + const boxMax = d.vec3f(gridSize); + const isect = getBoxIntersection(rayOrigin, rayDir, d.vec3f(0), boxMax); + if (!isect.hit) { + return d.vec4f(); + } + + const stepSize = (isect.tFar - isect.tNear) / STEPS; + const opacity = (stepSize / gridSize) * 3; + let transmittance = d.f32(1); + let accum = d.f32(); + let i = 0; + while (i < STEPS && transmittance > 1e-3) { + const t = isect.tNear + (d.f32(i) + 0.5) * stepSize; + const pos = rayOrigin + rayDir * t; + const value = std.textureLoad( + layouts.display.$.texture, + d.vec3u(std.clamp(pos, d.vec3f(0), boxMax - 1)), + ).r; + accum += value * opacity * transmittance; + transmittance *= 1 - opacity; + i += 1; + } - pipeline.withColorAttachment({ view: context }).draw(3); + return d.vec4f(d.vec3f(accum), 1 - transmittance); + }, + targets: { format: presentationFormat }, +}); + +const displayPipelines = { + '2d': displayPipeline2d, + '3d': displayPipeline3d, +}; + +const resample = () => { + configUniform.patch({ gridSize }); + + getComputePipeline(mode, prng) + .with(bindGroups[gridSize].compute) + .dispatchThreads(gridSize, gridSize, mode === '3d' ? gridSize : 1); +}; + +const redraw = () => { + displayPipelines[mode] + .withColorAttachment({ view: context }) + .with(bindGroups[gridSize].display) + .draw(3); }; -// #region Example controls & Cleanup +const { cleanupCamera, targetCamera } = setupOrbitCamera( + canvas, + { + initPos: d.vec4f(d.vec3f(2 * gridSize), 1), + target: d.vec4f(d.vec3f(0.5 * gridSize), 1), + minZoom: 10, + maxZoom: 1000, + }, + (updates) => { + cameraUniform.patch(updates); + redraw(); + }, +); + export const controls = defineControls({ + Mode: { + initial: mode, + options: modes, + onSelectChange: (value) => { + mode = value; + resample(); + redraw(); + }, + }, PRNG: { - initial: c.initialPRNG, - options: c.prngs, + initial: prng, + options: prngKeys, onSelectChange: (value) => { prng = value; + resample(); redraw(); }, }, 'Grid Size': { - initial: c.initialGridSize, - options: c.gridSizes, + initial: gridSize, + options: gridSizes, onSelectChange: (value) => { - gridSizeUniform.write(value); + gridSize = value; + targetCamera(d.vec4f(d.vec3f(2 * gridSize), 1), d.vec4f(d.vec3f(0.5 * gridSize), 1)); + resample(); + redraw(); + }, + }, + 'Samples per thread': { + initial: initialSamplesPerThread, + options: samplesPerThread, + onSelectChange: (value) => { + configUniform.patch({ samplesPerThread: value }); + resample(); + redraw(); + }, + }, + 'Take Average': { + initial: initialTakeAverage, + onToggleChange: (value) => { + configUniform.patch({ takeAverage: d.i32(value) }); + resample(); + redraw(); + }, + }, + 'Seed Multiplier': { + initial: initialMultiplier, + min: 0.00001, + max: 2000, + step: 1, + onSliderChange: (value) => { + configUniform.patch({ multiplier: value }); + resample(); redraw(); }, }, '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 })); + modes.forEach((mode) => { + prngKeys.forEach((key) => { + // don't care about display pipelines + root.device.createShaderModule({ + code: tgpu.resolve([getComputePipeline(mode, key).pipeline]), + }); + }); + }); }, }, }); -const resizeObserver = new ResizeObserver(() => { - canvasRatioUniform.write(canvas.width / canvas.height); - redraw(); -}); -resizeObserver.observe(canvas); - export function onCleanup() { - resizeObserver.disconnect(); + cleanupCamera(); root.destroy(); } 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..750db15c38 100644 --- a/apps/typegpu-docs/src/examples/tests/uniformity/prngs.ts +++ b/apps/typegpu-docs/src/examples/tests/uniformity/prngs.ts @@ -1,17 +1,16 @@ -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[]; diff --git a/apps/typegpu-docs/src/examples/tests/uniformity/thumbnail.png b/apps/typegpu-docs/src/examples/tests/uniformity/thumbnail.png index 0f55b3f149..e9c167f4c1 100644 Binary files a/apps/typegpu-docs/src/examples/tests/uniformity/thumbnail.png and b/apps/typegpu-docs/src/examples/tests/uniformity/thumbnail.png differ 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 445bfbfd9b..86d4b91b9d 100644 --- a/apps/typegpu-docs/tests/individual-example-tests/uniformity.test.ts +++ b/apps/typegpu-docs/tests/individual-example-tests/uniformity.test.ts @@ -17,7 +17,7 @@ describe('uniformity test example', () => { name: 'uniformity', setupMocks: mockResizeObserver, controlTriggers: ['Test Resolution'], - expectedCalls: 2, + expectedCalls: 3, }, device, ); @@ -35,9 +35,90 @@ describe('uniformity test example', () => { return fullScreenTriangle_Output(vec4f(pos[vertexIndex], 0, 1), uv[vertexIndex]); } - @group(0) @binding(0) var canvasRatioUniform: f32; + struct Camera { + position: vec4f, + targetPos: vec4f, + view: mat4x4f, + projection: mat4x4f, + viewInverse: mat4x4f, + projectionInverse: mat4x4f, + } + + @group(0) @binding(0) var cameraUniform: Camera; + + struct Config { + gridSize: f32, + samplesPerThread: i32, + takeAverage: i32, + multiplier: f32, + canvasRatio: f32, + } + + @group(0) @binding(1) var configUniform: Config; + + struct BoxIntersection { + tNear: f32, + tFar: f32, + hit: bool, + } + + fn getBoxIntersection(rayOrigin: vec3f, rayDir: vec3f, boxMin: vec3f, boxMax: vec3f) -> BoxIntersection { + let invDir = (1f / rayDir); + let t0 = ((boxMin - rayOrigin) * invDir); + let t1 = ((boxMax - rayOrigin) * invDir); + let tmin = min(t0, t1); + let tmax = max(t0, t1); + let tNear = max(max(tmin.x, tmin.y), tmin.z); + let tFar = min(min(tmax.x, tmax.y), tmax.z); + return BoxIntersection(tNear, tFar, (tFar >= tNear)); + } + + @group(1) @binding(0) var texture: texture_storage_3d; - @group(0) @binding(1) var gridSizeUniform: f32; + struct FragmentIn { + @location(0) uv: vec2f, + } + + @fragment fn fragment(_arg_0: FragmentIn) -> @location(0) vec4f { + let ndc = vec2f(((_arg_0.uv.x * 2f) - 1f), (1f - (_arg_0.uv.y * 2f))); + let invViewProj = (cameraUniform.viewInverse * cameraUniform.projectionInverse); + let worldNear = (invViewProj * vec4f(ndc, 0f, 1f)); + let worldFar = (invViewProj * vec4f(ndc, 1f, 1f)); + let rayOrigin = (worldNear.xyz / worldNear.w); + let rayDir = normalize(((worldFar.xyz / worldFar.w) - rayOrigin)); + let gridSize = configUniform.gridSize; + let boxMax = vec3f(gridSize); + let isect = getBoxIntersection(rayOrigin, rayDir, vec3f(), boxMax); + if (!isect.hit) { + return vec4f(); + } + let stepSize = ((isect.tFar - isect.tNear) / 64f); + let opacity = ((stepSize / gridSize) * 3f); + var transmittance = 1f; + var accum = 0f; + var i = 0; + while (((i < 64i) && (transmittance > 1e-3f))) { + let t = (isect.tNear + ((f32(i) + 0.5f) * stepSize)); + let pos = (rayOrigin + (rayDir * t)); + let value = textureLoad(texture, vec3u(clamp(pos, vec3f(), (boxMax - 1f)))).r; + accum += ((value * opacity) * transmittance); + transmittance *= (1f - opacity); + i += 1i; + } + return vec4f(vec3f(accum), (1f - transmittance)); + } + + @group(0) @binding(0) var sizeUniform: vec3u; + + struct Config { + gridSize: f32, + samplesPerThread: i32, + takeAverage: i32, + multiplier: f32, + canvasRatio: f32, + } + + @group(0) @binding(1) var configUniform: Config; var seed: vec2f; @@ -61,52 +142,432 @@ describe('uniformity test example', () => { return sample(); } - struct fragmentShader_Input { - @location(0) uv: vec2f, + @group(1) @binding(0) var texture: texture_storage_3d; + + fn computeFn(x: u32, y: u32, z: u32) { + let multiplier = configUniform.multiplier; + randSeed2(((vec2f(f32(x), f32(y)) - (configUniform.gridSize / 2f)) * multiplier)); + let samplesPerThread = configUniform.samplesPerThread; + let takeAverage = configUniform.takeAverage; + var sum = 0f; + for (var i = 0i; (i < (samplesPerThread - 1i)); i++) { + sum += randFloat01(); + } + var result = randFloat01(); + result += (sum * f32(takeAverage)); + let denominator = f32((1i + ((samplesPerThread - 1i) * takeAverage))); + result /= denominator; + textureStore(texture, vec3u(x, y, z), vec4f(result, 0f, 0f, 0f)); } - @fragment fn fragmentShader(_arg_0: fragmentShader_Input) -> @location(0) vec4f { - let uv = (((_arg_0.uv + 1f) / 2f) * vec2f(canvasRatioUniform, 1f)); - let gridedUV = floor((uv * gridSizeUniform)); - randSeed2(gridedUV); - return vec4f(vec3f(randFloat01()), 1f); + @compute @workgroup_size(1, 1, 1) fn mainCompute(@builtin(global_invocation_id) id: vec3u) { + if (any(id >= sizeUniform)) { + return; + } + computeFn(id.x, id.y, id.z); } - var seed_1: u32; + @group(0) @binding(0) var sizeUniform: vec3u; - fn seed2_1(value: vec2f) { - seed_1 = u32(((value.x * 32768f) + (value.y * 1024f))); + struct Config { + gridSize: f32, + samplesPerThread: i32, + takeAverage: i32, + multiplier: f32, + canvasRatio: f32, } - fn randSeed2_1(seed_1: vec2f) { - seed2_1(seed_1); + @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: u32; + + fn seed2(value: vec2f) { + let u32Value = bitcast(value); + let hx = hash((u32Value.x ^ 1253408251u)); + let hy = hash((u32Value.y ^ 2900286023u)); + seed = hash((hx ^ rotl(hy, 16u))); + } + + fn randSeed2(seed: vec2f) { + seed2(seed); } - fn u32To01Float(value: u32) -> f32 { - let mantissa = (value >> 9u); + fn u32To01F32(value: u32) -> f32 { + let mantissa = (value & 8388607u); 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(); } - struct fragmentShader_Input_1 { - @location(0) uv: vec2f, + @group(1) @binding(0) var texture: texture_storage_3d; + + fn computeFn(x: u32, y: u32, z: u32) { + let multiplier = configUniform.multiplier; + randSeed2(((vec2f(f32(x), f32(y)) - (configUniform.gridSize / 2f)) * multiplier)); + let samplesPerThread = configUniform.samplesPerThread; + let takeAverage = configUniform.takeAverage; + var sum = 0f; + for (var i = 0i; (i < (samplesPerThread - 1i)); i++) { + sum += randFloat01(); + } + var result = randFloat01(); + result += (sum * f32(takeAverage)); + let denominator = f32((1i + ((samplesPerThread - 1i) * takeAverage))); + result /= denominator; + textureStore(texture, vec3u(x, y, z), vec4f(result, 0f, 0f, 0f)); + } + + @compute @workgroup_size(1, 1, 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, + samplesPerThread: i32, + takeAverage: i32, + multiplier: f32, + canvasRatio: f32, + } + + @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 u32Value = bitcast(value); + let hx = hash((u32Value.x ^ 1253408251u)); + let hy = hash((u32Value.y ^ 2900286023u)); + 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 & 8388607u); + 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(); + } + + @group(1) @binding(0) var texture: texture_storage_3d; + + fn computeFn(x: u32, y: u32, z: u32) { + let multiplier = configUniform.multiplier; + randSeed2(((vec2f(f32(x), f32(y)) - (configUniform.gridSize / 2f)) * multiplier)); + let samplesPerThread = configUniform.samplesPerThread; + let takeAverage = configUniform.takeAverage; + var sum = 0f; + for (var i = 0i; (i < (samplesPerThread - 1i)); i++) { + sum += randFloat01(); + } + var result = randFloat01(); + result += (sum * f32(takeAverage)); + let denominator = f32((1i + ((samplesPerThread - 1i) * takeAverage))); + result /= denominator; + textureStore(texture, vec3u(x, y, z), vec4f(result, 0f, 0f, 0f)); + } + + @compute @workgroup_size(1, 1, 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, + samplesPerThread: i32, + takeAverage: i32, + multiplier: f32, + canvasRatio: f32, + } + + @group(0) @binding(1) var configUniform: Config; + + var seed: vec2f; + + fn seed3(value: vec3f) { + seed = (value.xy + vec2f(value.z)); + } + + fn randSeed3(seed: vec3f) { + seed3(seed); + } + + fn sample() -> f32 { + let a = dot(seed, vec2f(23.140779495239258, 232.6168975830078)); + let b = dot(seed, vec2f(54.47856521606445, 345.8415222167969)); + seed.x = fract((cos(a) * 136.8168f)); + seed.y = fract((cos(b) * 534.7645f)); + return seed.y; + } + + fn randFloat01() -> f32 { + return sample(); + } + + @group(1) @binding(0) var texture: texture_storage_3d; + + fn computeFn(x: u32, y: u32, z: u32) { + let multiplier = configUniform.multiplier; + randSeed3(((vec3f(f32(x), f32(y), f32(z)) - (configUniform.gridSize / 2f)) * multiplier)); + let samplesPerThread = configUniform.samplesPerThread; + let takeAverage = configUniform.takeAverage; + var sum = 0f; + for (var i = 0i; (i < (samplesPerThread - 1i)); i++) { + sum += randFloat01(); + } + var result = randFloat01(); + result += (sum * f32(takeAverage)); + let denominator = f32((1i + ((samplesPerThread - 1i) * takeAverage))); + result /= denominator; + textureStore(texture, vec3u(x, y, z), vec4f(result, 0f, 0f, 0f)); + } + + @compute @workgroup_size(1, 1, 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, + samplesPerThread: i32, + takeAverage: i32, + multiplier: f32, + canvasRatio: f32, + } + + @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: u32; + + fn seed3(value: vec3f) { + let u32Value = bitcast(value); + let hx = hash((u32Value.x ^ 1253408251u)); + let hy = hash((u32Value.y ^ 2900286023u)); + let hz = hash((u32Value.z ^ 3164612939u)); + seed = hash((hash((hx ^ rotl(hy, 16u))) ^ hz)); + } + + fn randSeed3(seed: vec3f) { + seed3(seed); + } + + fn u32To01F32(value: u32) -> f32 { + let mantissa = (value & 8388607u); + let bits = (1065353216u | mantissa); + let f = bitcast(bits); + return (f - 1f); + } + + fn sample() -> f32 { + seed = ((1664525u * seed) + 1013904223u); + return u32To01F32(seed); + } + + fn randFloat01() -> f32 { + return sample(); + } + + @group(1) @binding(0) var texture: texture_storage_3d; + + fn computeFn(x: u32, y: u32, z: u32) { + let multiplier = configUniform.multiplier; + randSeed3(((vec3f(f32(x), f32(y), f32(z)) - (configUniform.gridSize / 2f)) * multiplier)); + let samplesPerThread = configUniform.samplesPerThread; + let takeAverage = configUniform.takeAverage; + var sum = 0f; + for (var i = 0i; (i < (samplesPerThread - 1i)); i++) { + sum += randFloat01(); + } + var result = randFloat01(); + result += (sum * f32(takeAverage)); + let denominator = f32((1i + ((samplesPerThread - 1i) * takeAverage))); + result /= denominator; + textureStore(texture, vec3u(x, y, z), vec4f(result, 0f, 0f, 0f)); + } + + @compute @workgroup_size(1, 1, 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, + samplesPerThread: i32, + takeAverage: i32, + multiplier: f32, + canvasRatio: f32, + } + + @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 seed3(value: vec3f) { + let u32Value = bitcast(value); + let hx = hash((u32Value.x ^ 1253408251u)); + let hy = hash((u32Value.y ^ 2900286023u)); + let hz = hash((u32Value.z ^ 3164612939u)); + seed = vec2u(hash((hx ^ rotl(hz, 16u))), hash((rotl(hy, 16u) ^ hz))); + } + + fn randSeed3(seed: vec3f) { + seed3(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 & 8388607u); + 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(); + } + + @group(1) @binding(0) var texture: texture_storage_3d; + + fn computeFn(x: u32, y: u32, z: u32) { + let multiplier = configUniform.multiplier; + randSeed3(((vec3f(f32(x), f32(y), f32(z)) - (configUniform.gridSize / 2f)) * multiplier)); + let samplesPerThread = configUniform.samplesPerThread; + let takeAverage = configUniform.takeAverage; + var sum = 0f; + for (var i = 0i; (i < (samplesPerThread - 1i)); i++) { + sum += randFloat01(); + } + var result = randFloat01(); + result += (sum * f32(takeAverage)); + let denominator = f32((1i + ((samplesPerThread - 1i) * takeAverage))); + result /= denominator; + textureStore(texture, vec3u(x, y, z), vec4f(result, 0f, 0f, 0f)); } - @fragment fn fragmentShader_1(_arg_0: fragmentShader_Input_1) -> @location(0) vec4f { - let uv = (((_arg_0.uv + 1f) / 2f) * vec2f(canvasRatioUniform, 1f)); - let gridedUV = floor((uv * gridSizeUniform)); - randSeed2_1(gridedUV); - return vec4f(vec3f(randFloat01_1()), 1f); + @compute @workgroup_size(1, 1, 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; +});