Skip to content

Commit 1dc879e

Browse files
committed
lcg64
1 parent 054d272 commit 1dc879e

9 files changed

Lines changed: 315 additions & 34 deletions

File tree

apps/typegpu-docs/src/examples/tests/uniformity/constants.ts

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,6 @@ import type { PRNGKey } from './prngs.ts';
22

33
export const gridSizes = [8, 16, 32, 64, 128, 256, 512, 1024];
44
export const initialGridSize = gridSizes[4];
5-
export const initialPRNG: PRNGKey = 'bpeter';
65
export const samplesPerThread = [1, 8, 16, 64, 256, 1024, 131072, 262144];
76
export const initialSamplesPerThread = samplesPerThread[0];
87
export const initialTakeAverage = false;

apps/typegpu-docs/src/examples/tests/uniformity/index.ts

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -69,9 +69,9 @@ const computeFn = (x: number, y: number) => {
6969
const gridSize = configUniform.$.gridSize;
7070

7171
if (configUniform.$.useSeed2 === 1) {
72-
randf.seed2(d.vec2f(x, y));
72+
randf.seed2(d.vec2f(x, y) + 1);
7373
} else {
74-
randf.seed(d.f32(x) * gridSize + d.f32(y));
74+
randf.seed(d.f32(x + 1) * gridSize + d.f32(y + 1));
7575
}
7676

7777
let i = d.u32(0);
@@ -98,7 +98,7 @@ const getComputePipeline = (key: PRNGKey) => {
9898
.with(randomGeneratorSlot, prngs[key].generator)
9999
.createGuardedComputePipeline(computeFn)
100100
.withPerformanceCallback((start, end) => {
101-
console.log(`[${key}] - ${Number(end - start) / 1000} ms.`);
101+
console.log(`[${key}] - ${Number(end - start) / 1_000_000} ms.`);
102102
});
103103
computePipelineCache.set(key, pipeline);
104104
}

apps/typegpu-docs/src/examples/tests/uniformity/prngs.ts

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
import { BPETER, LCG32, XOROSHIRO64STARSTAR, type StatefulGenerator } from '@typegpu/noise';
1+
import { BPETER, LCG32, LCG64, XOROSHIRO64STARSTAR, type StatefulGenerator } from '@typegpu/noise';
22

33
interface PRNGOptions {
44
name: string;
@@ -9,11 +9,11 @@ interface PRNGOptions {
99
export const prngs = {
1010
bpeter: { name: 'bpeter (default)', useSeed2: true, generator: BPETER },
1111
lcg32: { name: 'lcg32', useSeed2: false, generator: LCG32 },
12+
lcg64: { name: 'lcg64', useSeed2: true, generator: LCG64 },
1213
xoroshiro64: { name: 'xoroshiro64', useSeed2: true, generator: XOROSHIRO64STARSTAR },
1314
} as const satisfies Record<string, PRNGOptions>;
1415

1516
export type PRNGKey = keyof typeof prngs;
1617

1718
export const prngKeys = Object.keys(prngs) as PRNGKey[];
18-
1919
export const initialPRNG: PRNGKey = 'bpeter';

apps/typegpu-docs/tests/individual-example-tests/uniformity.test.ts

Lines changed: 146 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -70,10 +70,10 @@ describe('uniformity test example', () => {
7070
fn computeFn(x: u32, y: u32, _arg_2: u32) {
7171
let gridSize2 = configUniform.gridSize;
7272
if ((configUniform.useSeed2 == 1u)) {
73-
randSeed2(vec2f(f32(x), f32(y)));
73+
randSeed2((vec2f(f32(x), f32(y)) + 1f));
7474
}
7575
else {
76-
randSeed(((f32(x) * gridSize2) + f32(y)));
76+
randSeed(((f32((x + 1u)) * gridSize2) + f32((y + 1u))));
7777
}
7878
var i = 0u;
7979
let samplesPerThread = configUniform.samplesPerThread;
@@ -137,16 +137,16 @@ describe('uniformity test example', () => {
137137
seed_1(seed);
138138
}
139139
140-
fn u32To01Float(value: u32) -> f32 {
140+
fn u32To01F32(value: u32) -> f32 {
141141
let mantissa = (value >> 9u);
142142
let bits = (1065353216u | mantissa);
143143
let f = bitcast<f32>(bits);
144144
return (f - 1f);
145145
}
146146
147147
fn sample() -> f32 {
148-
seed = ((seed * 1664525u) + 1013904223u);
149-
return u32To01Float(seed);
148+
seed = ((1664525u * seed) + 1013904223u);
149+
return u32To01F32(seed);
150150
}
151151
152152
fn randFloat01() -> f32 {
@@ -158,10 +158,10 @@ describe('uniformity test example', () => {
158158
fn computeFn(x: u32, y: u32, _arg_2: u32) {
159159
let gridSize2 = configUniform.gridSize;
160160
if ((configUniform.useSeed2 == 1u)) {
161-
randSeed2(vec2f(f32(x), f32(y)));
161+
randSeed2((vec2f(f32(x), f32(y)) + 1f));
162162
}
163163
else {
164-
randSeed(((f32(x) * gridSize2) + f32(y)));
164+
randSeed(((f32((x + 1u)) * gridSize2) + f32((y + 1u))));
165165
}
166166
var i = 0u;
167167
let samplesPerThread = configUniform.samplesPerThread;
@@ -211,10 +211,16 @@ describe('uniformity test example', () => {
211211
return x;
212212
}
213213
214+
fn rotl(x: u32, k: u32) -> u32 {
215+
return ((x << k) | (x >> (32u - k)));
216+
}
217+
214218
var<private> seed: vec2u;
215219
216220
fn seed2(value: vec2f) {
217-
seed = vec2u(hash(u32(value.x)), hash(u32(value.y)));
221+
let hx = hash((u32(value.x) ^ 2135587861u));
222+
let hy = hash((u32(value.y) ^ 2654435769u));
223+
seed = vec2u(hash((hx ^ hy)), hash((rotl(hx, 16u) ^ hy)));
218224
}
219225
220226
fn randSeed2(seed: vec2f) {
@@ -225,10 +231,138 @@ describe('uniformity test example', () => {
225231
226232
}
227233
234+
fn u64Mul(a: vec2u, b: vec2u) -> vec2u {
235+
let all_1 = (a.x & 65535u);
236+
let alh = (a.x >> 16u);
237+
let ahl = (a.y & 65535u);
238+
let ahh = (a.y >> 16u);
239+
let bll = (b.x & 65535u);
240+
let blh = (b.x >> 16u);
241+
let bhl = (b.y & 65535u);
242+
let bhh = (b.y >> 16u);
243+
let row0_0 = (bll * all_1);
244+
let row0_1 = (bll * alh);
245+
let row0_2 = (bll * ahl);
246+
let row0_3 = (bll * ahh);
247+
let row1_0 = (blh * all_1);
248+
let row1_1 = (blh * alh);
249+
let row1_2 = (blh * ahl);
250+
let row2_0 = (bhl * all_1);
251+
let row2_1 = (bhl * alh);
252+
let row3_0 = (bhh * all_1);
253+
let r1 = (row0_0 & 65535u);
254+
var r2 = (((row0_0 >> 16u) + (row0_1 & 65535u)) + (row1_0 & 65535u));
255+
var r3 = (((((row0_1 >> 16u) + (row0_2 & 65535u)) + (row1_0 >> 16u)) + (row1_1 & 65535u)) + (row2_0 & 65535u));
256+
var r4 = (((((((row0_2 >> 16u) + (row0_3 & 65535u)) + (row1_1 >> 16u)) + (row1_2 & 65535u)) + (row2_0 >> 16u)) + (row2_1 & 65535u)) + (row3_0 & 65535u));
257+
r3 += (r2 >> 16u);
258+
r2 &= 65535u;
259+
r4 += (r3 >> 16u);
260+
r3 &= 65535u;
261+
r4 &= 65535u;
262+
return vec2u((r1 | (r2 << 16u)), (r3 | (r4 << 16u)));
263+
}
264+
265+
fn u64Add(a: vec2u, b: vec2u) -> vec2u {
266+
let rl = (a.x + b.x);
267+
let carry = u32(((rl < a.x) && (rl < b.x)));
268+
let rh = ((a.y + b.y) + carry);
269+
return vec2u(rl, rh);
270+
}
271+
272+
fn u32To01F32(value: u32) -> f32 {
273+
let mantissa = (value >> 9u);
274+
let bits = (1065353216u | mantissa);
275+
let f = bitcast<f32>(bits);
276+
return (f - 1f);
277+
}
278+
279+
fn sample() -> f32 {
280+
seed = u64Add(u64Mul(seed, vec2u(1284865837, 1481765933)), vec2u(1, 0));
281+
return u32To01F32(seed.y);
282+
}
283+
284+
fn randFloat01() -> f32 {
285+
return sample();
286+
}
287+
288+
@group(1) @binding(0) var texture: texture_storage_2d<r32float, write>;
289+
290+
fn computeFn(x: u32, y: u32, _arg_2: u32) {
291+
let gridSize2 = configUniform.gridSize;
292+
if ((configUniform.useSeed2 == 1u)) {
293+
randSeed2((vec2f(f32(x), f32(y)) + 1f));
294+
}
295+
else {
296+
randSeed(((f32((x + 1u)) * gridSize2) + f32((y + 1u))));
297+
}
298+
var i = 0u;
299+
let samplesPerThread = configUniform.samplesPerThread;
300+
var samples = 0f;
301+
while ((i < (samplesPerThread - 1u))) {
302+
samples += randFloat01();
303+
i += 1u;
304+
}
305+
var result = randFloat01();
306+
if ((configUniform.takeAverage == 1u)) {
307+
result = ((result + samples) / f32(samplesPerThread));
308+
}
309+
textureStore(texture, vec2u(x, y), vec4f(result, 0f, 0f, 0f));
310+
}
311+
312+
struct mainCompute_Input {
313+
@builtin(global_invocation_id) id: vec3u,
314+
}
315+
316+
@compute @workgroup_size(16, 16, 1) fn mainCompute(in: mainCompute_Input) {
317+
if (any(in.id >= sizeUniform)) {
318+
return;
319+
}
320+
computeFn(in.id.x, in.id.y, in.id.z);
321+
}
322+
323+
@group(0) @binding(0) var<uniform> sizeUniform: vec3u;
324+
325+
struct Config {
326+
gridSize: f32,
327+
canvasRatio: f32,
328+
useSeed2: u32,
329+
samplesPerThread: u32,
330+
takeAverage: u32,
331+
}
332+
333+
@group(0) @binding(1) var<uniform> configUniform: Config;
334+
335+
fn hash(v: u32) -> u32 {
336+
var x = (v ^ (v >> 17u));
337+
x *= 3982152891u;
338+
x ^= (x >> 11u);
339+
x *= 2890668881u;
340+
x ^= (x >> 15u);
341+
x *= 830770091u;
342+
x ^= (x >> 14u);
343+
return x;
344+
}
345+
228346
fn rotl(x: u32, k: u32) -> u32 {
229347
return ((x << k) | (x >> (32u - k)));
230348
}
231349
350+
var<private> seed: vec2u;
351+
352+
fn seed2(value: vec2f) {
353+
let hx = hash((u32(value.x) ^ 2135587861u));
354+
let hy = hash((u32(value.y) ^ 2654435769u));
355+
seed = vec2u(hash((hx ^ hy)), hash((rotl(hx, 16u) ^ hy)));
356+
}
357+
358+
fn randSeed2(seed: vec2f) {
359+
seed2(seed);
360+
}
361+
362+
fn randSeed(seed_1: f32) {
363+
364+
}
365+
232366
fn next() -> u32 {
233367
let s0 = seed[0i];
234368
var s1 = seed[1i];
@@ -238,7 +372,7 @@ describe('uniformity test example', () => {
238372
return (rotl((seed[0i] * 2654435771u), 5u) * 5u);
239373
}
240374
241-
fn u32To01Float(value: u32) -> f32 {
375+
fn u32To01F32(value: u32) -> f32 {
242376
let mantissa = (value >> 9u);
243377
let bits = (1065353216u | mantissa);
244378
let f = bitcast<f32>(bits);
@@ -247,7 +381,7 @@ describe('uniformity test example', () => {
247381
248382
fn sample() -> f32 {
249383
let r = next();
250-
return u32To01Float(r);
384+
return u32To01F32(r);
251385
}
252386
253387
fn randFloat01() -> f32 {
@@ -259,10 +393,10 @@ describe('uniformity test example', () => {
259393
fn computeFn(x: u32, y: u32, _arg_2: u32) {
260394
let gridSize2 = configUniform.gridSize;
261395
if ((configUniform.useSeed2 == 1u)) {
262-
randSeed2(vec2f(f32(x), f32(y)));
396+
randSeed2((vec2f(f32(x), f32(y)) + 1f));
263397
}
264398
else {
265-
randSeed(((f32(x) * gridSize2) + f32(y)));
399+
randSeed(((f32((x + 1u)) * gridSize2) + f32((y + 1u))));
266400
}
267401
var i = 0u;
268402
let samplesPerThread = configUniform.samplesPerThread;

packages/typegpu-noise/src/generator.ts

Lines changed: 20 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
import tgpu, { d, type TgpuFnShell, type TgpuSlot } from 'typegpu';
22
import { cos, dot, fract } from 'typegpu/std';
3-
import { hash, u32To01F32 } from './utils.ts';
3+
import { hash, rotl, u32To01F32, u64Add, u64Mul } from './utils.ts';
44

55
export interface StatefulGenerator {
66
seed?: (seed: number) => void;
@@ -56,13 +56,6 @@ export const BPETER: StatefulGenerator = (() => {
5656
export const XOROSHIRO64STARSTAR: StatefulGenerator = (() => {
5757
const seed = tgpu.privateVar(d.vec2u);
5858

59-
const rotl = tgpu.fn(
60-
[d.u32, d.u32],
61-
d.u32,
62-
)((x, k) => {
63-
return (x << k) | (x >> (32 - k));
64-
});
65-
6659
const next = tgpu.fn(
6760
[],
6861
d.u32,
@@ -77,7 +70,9 @@ export const XOROSHIRO64STARSTAR: StatefulGenerator = (() => {
7770

7871
return {
7972
seed2: tgpu.fn([d.vec2f])((value) => {
80-
seed.$ = d.vec2u(hash(d.u32(value.x)), hash(d.u32(value.y)));
73+
const hx = hash(d.u32(value.x) ^ 2135587861);
74+
const hy = hash(d.u32(value.y) ^ 2654435769);
75+
seed.$ = d.vec2u(hash(hx ^ hy), hash(rotl(hx, 16) ^ hy));
8176
}),
8277

8378
sample: randomGeneratorShell(() => {
@@ -94,34 +89,44 @@ export const XOROSHIRO64STARSTAR: StatefulGenerator = (() => {
9489
export const LCG32: StatefulGenerator = (() => {
9590
const seed = tgpu.privateVar(d.u32);
9691

92+
const multiplier = d.u32(1664525);
93+
const increment = d.u32(1013904223);
94+
9795
return {
9896
seed: tgpu.fn([d.f32])((value) => {
9997
seed.$ = hash(d.u32(value));
10098
}),
10199

102100
sample: randomGeneratorShell(() => {
103101
'use gpu';
104-
seed.$ = seed.$ * 1664525 + 1013904223; // % 2 ^ 32
102+
seed.$ = multiplier * seed.$ + increment; // % 2 ^ 32
105103
return u32To01F32(seed.$);
106104
}).$name('sample'),
107105
};
108106
})();
109107

110108
/**
111109
* Naive Linear Congruential Generator (LCG) with 64 bits state
110+
* Incorporated from: https://en.wikipedia.org/wiki/Linear_congruential_generator (Musl)
112111
*/
113112
export const LCG64: StatefulGenerator = (() => {
114-
const seed = tgpu.privateVar(d.u32);
113+
const seed = tgpu.privateVar(d.vec2u);
114+
115+
// 1481765933 * 2 ** 32 + 1284865837 = 6364136223846793005
116+
const multiplier = d.vec2u(1284865837, 1481765933);
117+
const increment = d.vec2u(1, 0);
115118

116119
return {
117-
seed: tgpu.fn([d.f32])((value) => {
118-
seed.$ = d.u32(value) * 1048577;
120+
seed2: tgpu.fn([d.vec2f])((value) => {
121+
const hx = hash(d.u32(value.x) ^ 2135587861);
122+
const hy = hash(d.u32(value.y) ^ 2654435769);
123+
seed.$ = d.vec2u(hash(hx ^ hy), hash(rotl(hx, 16) ^ hy));
119124
}),
120125

121126
sample: randomGeneratorShell(() => {
122127
'use gpu';
123-
seed.$ = seed.$ * 1664525 + 1013904223; // % 2 ^ 32
124-
return u32To01F32(seed.$);
128+
seed.$ = u64Add(u64Mul(seed.$, multiplier), increment);
129+
return u32To01F32(seed.$.y);
125130
}).$name('sample'),
126131
};
127132
})();

0 commit comments

Comments
 (0)