diff --git a/examples/models/gltf/kira.glb b/examples/models/gltf/kira.glb index 2ce9e76631a12f..5c9028225fc203 100644 Binary files a/examples/models/gltf/kira.glb and b/examples/models/gltf/kira.glb differ diff --git a/examples/screenshots/webgl_animation_skinning_ik.jpg b/examples/screenshots/webgl_animation_skinning_ik.jpg index 58b36f93e37e55..bc0eaf77b75f41 100644 Binary files a/examples/screenshots/webgl_animation_skinning_ik.jpg and b/examples/screenshots/webgl_animation_skinning_ik.jpg differ diff --git a/src/nodes/gpgpu/SubgroupFunctionNode.js b/src/nodes/gpgpu/SubgroupFunctionNode.js index b8684023f1a6f0..b4f3682320c9b8 100644 --- a/src/nodes/gpgpu/SubgroupFunctionNode.js +++ b/src/nodes/gpgpu/SubgroupFunctionNode.js @@ -195,7 +195,7 @@ class SubgroupFunctionNode extends TempNode { } - static get SUBGROUP_EXCLUSIVE_AND() { + static get SUBGROUP_EXCLUSIVE_ADD() { return 'subgroupExclusiveAdd'; @@ -375,7 +375,7 @@ export const subgroupInclusiveAdd = /*@__PURE__*/ nodeProxyIntent( SubgroupFunct * @param {number} e - The value provided to the exclusive scan by the current invocation. * @return {number} The accumulated result of the exclusive scan operation. */ -export const subgroupExclusiveAdd = /*@__PURE__*/ nodeProxyIntent( SubgroupFunctionNode, SubgroupFunctionNode.SUBGROUP_EXCLUSIVE_AND ).setParameterLength( 1 ); +export const subgroupExclusiveAdd = /*@__PURE__*/ nodeProxyIntent( SubgroupFunctionNode, SubgroupFunctionNode.SUBGROUP_EXCLUSIVE_ADD ).setParameterLength( 1 ); /** * A reduction that multiplies e among all active invocations and returns that result. diff --git a/src/nodes/gpgpu/WorkgroupInfoNode.js b/src/nodes/gpgpu/WorkgroupInfoNode.js index 6a671095014d82..a739aa9c40ada4 100644 --- a/src/nodes/gpgpu/WorkgroupInfoNode.js +++ b/src/nodes/gpgpu/WorkgroupInfoNode.js @@ -94,6 +94,14 @@ class WorkgroupInfoNode extends Node { */ this.bufferCount = bufferCount; + /** + * Whether the node is atomic or not. + * + * @type {boolean} + * @default false + */ + this.isAtomic = false; + /** * This flag can be used for type testing. * @@ -170,6 +178,31 @@ class WorkgroupInfoNode extends Node { } + /** + * Defines whether the node is atomic or not. + * + * @param {boolean} value - The atomic flag. + * @return {WorkgroupInfoNode} A reference to this node. + */ + setAtomic( value ) { + + this.isAtomic = value; + + return this; + + } + + /** + * Convenience method for making this node atomic. + * + * @return {WorkgroupInfoNode} A reference to this node. + */ + toAtomic() { + + return this.setAtomic( true ); + + } + /** * The data type of the array buffer. @@ -217,7 +250,7 @@ class WorkgroupInfoNode extends Node { const name = ( this.name !== '' ) ? this.name : `${this.scope}Array_${this.id}`; - return builder.getScopedArray( name, this.scope.toLowerCase(), this.bufferType, this.bufferCount ); + return builder.getScopedArray( name, this.scope.toLowerCase(), this.bufferType, this.bufferCount, this.isAtomic ); } @@ -237,4 +270,3 @@ export default WorkgroupInfoNode; */ export const workgroupArray = ( type, count ) => new WorkgroupInfoNode( 'Workgroup', type, count ); - diff --git a/src/renderers/webgpu/nodes/WGSLNodeBuilder.js b/src/renderers/webgpu/nodes/WGSLNodeBuilder.js index 2741fb913266c7..141480f7109eec 100644 --- a/src/renderers/webgpu/nodes/WGSLNodeBuilder.js +++ b/src/renderers/webgpu/nodes/WGSLNodeBuilder.js @@ -1802,9 +1802,10 @@ ${ flowData.code } * @param {string} scope - The scope. * @param {string} bufferType - The buffer type. * @param {string} bufferCount - The buffer count. + * @param {boolean} isAtomic - Whether the array elements are atomic or not. * @return {string} The array name. */ - getScopedArray( name, scope, bufferType, bufferCount ) { + getScopedArray( name, scope, bufferType, bufferCount, isAtomic ) { if ( this.scopedArrays.has( name ) === false ) { @@ -1812,7 +1813,8 @@ ${ flowData.code } name, scope, bufferType, - bufferCount + bufferCount, + isAtomic } ); } @@ -1838,9 +1840,11 @@ ${ flowData.code } const snippets = []; - for ( const { name, scope, bufferType, bufferCount } of this.scopedArrays.values() ) { + for ( const { name, scope, bufferType, bufferCount, isAtomic } of this.scopedArrays.values() ) { - const type = this.getType( bufferType ); + let type = this.getType( bufferType ); + + if ( isAtomic === true ) type = `atomic<${type}>`; snippets.push( `var<${scope}> ${name}: array< ${type}, ${bufferCount} >;` ); diff --git a/test/unit/addons/tsl/GPUAtomicsStorage.tests.js b/test/unit/addons/tsl/GPUAtomicsStorage.tests.js new file mode 100644 index 00000000000000..10e9d93d351a82 --- /dev/null +++ b/test/unit/addons/tsl/GPUAtomicsStorage.tests.js @@ -0,0 +1,228 @@ +import { + Fn, instanceIndex, instancedArray, + atomicAdd, atomicSub, atomicMax, atomicMin, atomicAnd, atomicOr, atomicXor, + atomicLoad, atomicStore, + uint, shiftLeft, bitNot +} from 'three/tsl'; +import { rawComputeTest, readUintBuffer } from './gpu-raw-test-utils.js'; + +// Coverage for every `atomicFunc()`-family op (AtomicFunctionNode.js) on a +// *storage* buffer (`instancedArray(...).toAtomic()`) -- as opposed to +// mrdoob/three.js#34428's `workgroupArray(...).toAtomic()`, which is scoped +// to a single workgroup. A storage buffer is visible to *every* invocation +// across *every* workgroup in the dispatch, so these tests deliberately +// spread invocations across several workgroups (`workgroupSize` far smaller +// than `dispatchCount`) to exercise cross-workgroup atomicity, not just +// within-workgroup atomicity (already covered separately for the workgroup +// case). +// +// Each op that needs a specific starting value (`Sub`/`Max`/`Min`/`And`/`Or`/ +// `Xor`) is seeded with its own small "init" compute dispatch first, awaited +// (`computeAsync`) before the "op" dispatch runs, so the two never race each +// other -- only the op dispatch's *own* invocations are racing, which is +// exactly what's under test. +// +// Reading the buffer back afterwards uses the raw storage bytes directly +// (`getArrayBufferAsync` -- see `readUintBuffer`), not a further +// `atomicLoad()` kernel: `atomic` has the same in-memory layout as +// plain `u32`, atomics are a WGSL type-checking construct, not a different +// storage format, so a host-side readback after all GPU work has completed +// is exactly the final value. + +const WORKGROUP_SIZE = 8; + +function makeCounter() { + + return instancedArray( 1, 'uint' ).toAtomic(); + +} + +async function seed( renderer, counter, value ) { + + const kernel = Fn( () => { + + atomicStore( counter.element( uint( 0 ) ), uint( value ) ); + + } )().compute( 1 ); + + await renderer.computeAsync( kernel ); + +} + +export default QUnit.module( 'TSL', () => { + + QUnit.module( 'storage buffer atomics', () => { + + rawComputeTest( 'atomicAdd: concurrent adds across multiple workgroups sum exactly once each', {}, async ( { assert, renderer } ) => { + + const dispatchCount = 64; // 8 workgroups of WORKGROUP_SIZE + const counter = makeCounter(); + + const kernel = Fn( () => { + + atomicAdd( counter.element( uint( 0 ) ), uint( 1 ) ); + + } )().compute( dispatchCount, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const data = await readUintBuffer( renderer, counter ); + assert.strictEqual( data[ 0 ], dispatchCount, `expected ${ dispatchCount } (one add per invocation, across ${ dispatchCount / WORKGROUP_SIZE } workgroups)` ); + + } ); + + rawComputeTest( 'atomicSub: concurrent subs across multiple workgroups drain exactly once each', {}, async ( { assert, renderer } ) => { + + const dispatchCount = 64; + const counter = makeCounter(); + + await seed( renderer, counter, dispatchCount ); + + const kernel = Fn( () => { + + atomicSub( counter.element( uint( 0 ) ), uint( 1 ) ); + + } )().compute( dispatchCount, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const data = await readUintBuffer( renderer, counter ); + assert.strictEqual( data[ 0 ], 0, 'expected 0 (one sub per invocation, draining the seeded count exactly)' ); + + } ); + + rawComputeTest( 'atomicMax: concurrent max across multiple workgroups converges to the true maximum', {}, async ( { assert, renderer } ) => { + + const dispatchCount = 37; // deliberately not a multiple of WORKGROUP_SIZE + const counter = makeCounter(); + + await seed( renderer, counter, 0 ); + + const kernel = Fn( () => { + + atomicMax( counter.element( uint( 0 ) ), instanceIndex ); + + } )().compute( dispatchCount, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const data = await readUintBuffer( renderer, counter ); + assert.strictEqual( data[ 0 ], dispatchCount - 1, `expected ${ dispatchCount - 1 } (the largest instanceIndex)` ); + + } ); + + rawComputeTest( 'atomicMin: concurrent min across multiple workgroups converges to the true minimum', {}, async ( { assert, renderer } ) => { + + const dispatchCount = 37; + const counter = makeCounter(); + + await seed( renderer, counter, 0xffffffff ); + + const kernel = Fn( () => { + + atomicMin( counter.element( uint( 0 ) ), instanceIndex ); + + } )().compute( dispatchCount, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const data = await readUintBuffer( renderer, counter ); + assert.strictEqual( data[ 0 ], 0, 'expected 0 (the smallest instanceIndex)' ); + + } ); + + rawComputeTest( 'atomicAnd: each invocation clears one distinct bit, all clears land', {}, async ( { assert, renderer } ) => { + + // 32 invocations, each clearing a different one of the 32 bits -- + // only passes if every single invocation's AND actually took + // effect (a lost update would leave a stray 1 bit set). + const dispatchCount = 32; + const counter = makeCounter(); + + await seed( renderer, counter, 0xffffffff ); + + const kernel = Fn( () => { + + const bit = shiftLeft( uint( 1 ), instanceIndex ); + atomicAnd( counter.element( uint( 0 ) ), bitNot( bit ) ); + + } )().compute( dispatchCount, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const data = await readUintBuffer( renderer, counter ); + assert.strictEqual( data[ 0 ], 0, 'expected 0x00000000 (every one of the 32 bits cleared exactly once)' ); + + } ); + + rawComputeTest( 'atomicOr: each invocation sets one distinct bit, all sets land', {}, async ( { assert, renderer } ) => { + + const dispatchCount = 32; + const counter = makeCounter(); + + await seed( renderer, counter, 0 ); + + const kernel = Fn( () => { + + const bit = shiftLeft( uint( 1 ), instanceIndex ); + atomicOr( counter.element( uint( 0 ) ), bit ); + + } )().compute( dispatchCount, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const data = await readUintBuffer( renderer, counter ); + assert.strictEqual( data[ 0 ] >>> 0, 0xffffffff, 'expected 0xffffffff (every one of the 32 bits set exactly once)' ); + + } ); + + rawComputeTest( 'atomicXor: each invocation flips one distinct bit, all flips land', {}, async ( { assert, renderer } ) => { + + const dispatchCount = 32; + const counter = makeCounter(); + + await seed( renderer, counter, 0 ); + + const kernel = Fn( () => { + + const bit = shiftLeft( uint( 1 ), instanceIndex ); + atomicXor( counter.element( uint( 0 ) ), bit ); + + } )().compute( dispatchCount, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const data = await readUintBuffer( renderer, counter ); + assert.strictEqual( data[ 0 ] >>> 0, 0xffffffff, 'expected 0xffffffff (every one of the 32 bits toggled from 0 to 1 exactly once)' ); + + } ); + + rawComputeTest( 'atomicStore + atomicLoad: a store from one dispatch is visible to a later dispatch\'s loads', {}, async ( { assert, renderer } ) => { + + const dispatchCount = 16; + const counter = makeCounter(); + const output = instancedArray( dispatchCount, 'uint' ); + + await seed( renderer, counter, 424242 ); + + const kernel = Fn( () => { + + output.element( instanceIndex ).assign( atomicLoad( counter.element( uint( 0 ) ) ) ); + + } )().compute( dispatchCount, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const data = await readUintBuffer( renderer, output ); + + for ( let i = 0; i < dispatchCount; i ++ ) { + + assert.strictEqual( data[ i ], 424242, `invocation ${ i }: atomicLoad should read back the earlier atomicStore's value` ); + + } + + } ); + + } ); + +} ); diff --git a/test/unit/addons/tsl/GPUBarriers.tests.js b/test/unit/addons/tsl/GPUBarriers.tests.js new file mode 100644 index 00000000000000..dacdf4be1b8731 --- /dev/null +++ b/test/unit/addons/tsl/GPUBarriers.tests.js @@ -0,0 +1,126 @@ +import { + Fn, If, + instanceIndex, invocationLocalIndex, + instancedArray, storageBarrier, textureBarrier, + textureStore, storageTexture, + uvec2, vec4, uint, float +} from 'three/tsl'; +import { StorageTexture, FloatType } from 'three/webgpu'; +import { rawComputeTest, readUintBuffer } from './gpu-raw-test-utils.js'; + +// Coverage for the two barrier variants BarrierNode.js declares beyond +// `workgroupBarrier()` (already covered in GPUWorkgroupAtomic.tests.js): +// `storageBarrier()` and `textureBarrier()`. Neither had any test, and +// `textureBarrier()` in particular has no usage anywhere in the codebase +// (src, examples, or tests) to model against -- see the second test's +// comments for what had to be worked out from scratch. + +export default QUnit.module( 'TSL', () => { + + QUnit.module( 'barriers', () => { + + rawComputeTest( 'storageBarrier: a workgroup-mate\'s storage write is visible after the barrier', {}, async ( { assert, renderer } ) => { + + const workgroupSize = 8; + const dispatchCount = 32; // 4 workgroups of 8 + + const shared = instancedArray( dispatchCount, 'uint' ); + const output = instancedArray( dispatchCount, 'uint' ); + + const kernel = Fn( () => { + + // Every invocation writes its own local index into its own + // slot of a *storage* (not workgroup) buffer, then reads the + // slot written by the invocation "to its left" within the + // same workgroup (wrapping at 0) -- only correct if + // storageBarrier() actually ordered that write before this + // read for every invocation in the workgroup, not just + // program order for the writer itself. + shared.element( instanceIndex ).assign( invocationLocalIndex ); + + storageBarrier(); + + const neighborLocalId = invocationLocalIndex.add( uint( workgroupSize - 1 ) ).mod( uint( workgroupSize ) ); + const neighborGlobalSlot = instanceIndex.sub( invocationLocalIndex ).add( neighborLocalId ); + + output.element( instanceIndex ).assign( shared.element( neighborGlobalSlot ) ); + + } )().compute( dispatchCount, [ workgroupSize ] ); + + await renderer.computeAsync( kernel ); + + const data = await readUintBuffer( renderer, output ); + + for ( let i = 0; i < dispatchCount; i ++ ) { + + const localId = i % workgroupSize; + const expected = ( localId + workgroupSize - 1 ) % workgroupSize; + + assert.strictEqual( data[ i ], expected, `invocation ${ i } (local ${ localId }): should read its left neighbor's local index (${ expected }) via the storage buffer` ); + + } + + } ); + + rawComputeTest( 'textureBarrier: a workgroup-mate\'s storage-texture write is visible after the barrier', {}, async ( { assert, renderer } ) => { + + // No existing code in this repo calls textureBarrier() (checked: + // zero hits across src/examples/test besides its own + // declaration) -- built from the WGSL/WebGPU spec instead: + // textureBarrier() only orders *storage-texture* accesses within + // a workgroup, the texture analogue of storageBarrier() for + // buffers. `FloatType` avoids the default rgba8unorm storage + // format's quantization, so the round-tripped values below can + // be compared exactly rather than with a tolerance. + const width = 4; + const storageTex = new StorageTexture( width, 1 ); + storageTex.type = FloatType; + + const output = instancedArray( 1, 'float' ); + + const kernel = Fn( () => { + + // Each of the 4 invocations (one workgroup) writes a + // distinct, identifiable value into its own texel. + const coord = uvec2( instanceIndex, uint( 0 ) ); + textureStore( storageTex, coord, vec4( float( instanceIndex ).add( 1 ), 0, 0, 1 ) ).toReadWrite(); + + textureBarrier(); + + // Only invocation 0 reads all 4 texels back and sums them -- + // only correct if every other invocation's write above is + // actually visible to invocation 0 by the time it runs this, + // i.e. if textureBarrier() did its job. Reading a storage + // texture back (as opposed to a plain sampled texture) needs + // the dedicated `storageTexture()` node -- `texture()`/ + // `textureLoad()` build a *sampled*-texture node (with an + // implicit mip-level argument), which doesn't type-check + // against a `texture_storage_2d` binding in WGSL. + If( instanceIndex.equal( uint( 0 ) ), () => { + + let sum = storageTexture( storageTex, uvec2( uint( 0 ), uint( 0 ) ) ).setSampler( false ).toReadWrite().r; + + for ( let i = 1; i < width; i ++ ) { + + sum = sum.add( storageTexture( storageTex, uvec2( uint( i ), uint( 0 ) ) ).setSampler( false ).toReadWrite().r ); + + } + + output.element( uint( 0 ) ).assign( sum ); + + } ); + + } )().compute( width, [ width ] ); + + await renderer.computeAsync( kernel ); + + const data = new Float32Array( await renderer.getArrayBufferAsync( output.value ) ); + + // Texel values are 1, 2, 3, 4 -- sum = 10. + assert.strictEqual( data[ 0 ], 10, 'sum of all 4 texels, read back after textureBarrier(), should be exactly 10' ); + + } ); + + } ); + +} ); diff --git a/test/unit/addons/tsl/GPUComputeBuiltins.tests.js b/test/unit/addons/tsl/GPUComputeBuiltins.tests.js new file mode 100644 index 00000000000000..86c7c364c40905 --- /dev/null +++ b/test/unit/addons/tsl/GPUComputeBuiltins.tests.js @@ -0,0 +1,127 @@ +import { + Fn, instanceIndex, instancedArray, + localId, workgroupId, globalId, numWorkgroups, + uint +} from 'three/tsl'; +import { rawComputeTest, readUintBuffer } from './gpu-raw-test-utils.js'; + +// Coverage for the compute-scope builtins in ComputeBuiltinNode.js +// (`localId`, `workgroupId`, `globalId`, `numWorkgroups`) -- previously +// untested. `subgroupSize` (also declared in ComputeBuiltinNode.js) is +// covered in GPUSubgroup.tests.js instead, since -- unlike these four -- +// it requires the WebGPU `'subgroups'` feature to even be declared as a +// shader builtin (see `WGSLNodeBuilder.getSubgroupSize()`'s +// `enableSubGroups()` call) and so needs the same feature-detection gate as +// the rest of the subgroup functions. +// +// Verified against a 1D dispatch with an exact multiple of workgroupSize +// (`dispatchCount = workgroupSize * workgroupCount`), so every relationship +// below is unambiguous -- no partially-filled last workgroup to reason +// about. + +const WORKGROUP_SIZE = 8; +const WORKGROUP_COUNT = 5; +const DISPATCH_COUNT = WORKGROUP_SIZE * WORKGROUP_COUNT; // 40 + +export default QUnit.module( 'TSL', () => { + + QUnit.module( 'compute builtins', () => { + + rawComputeTest( 'localId, workgroupId, globalId and numWorkgroups match their WGSL definitions', {}, async ( { assert, renderer } ) => { + + const localIdOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const workgroupIdOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const globalIdOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const numWorkgroupsOut = instancedArray( DISPATCH_COUNT, 'uint' ); + + const kernel = Fn( () => { + + localIdOut.element( instanceIndex ).assign( localId.x ); + workgroupIdOut.element( instanceIndex ).assign( workgroupId.x ); + globalIdOut.element( instanceIndex ).assign( globalId.x ); + numWorkgroupsOut.element( instanceIndex ).assign( numWorkgroups.x ); + + } )().compute( DISPATCH_COUNT, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const localIdData = await readUintBuffer( renderer, localIdOut ); + const workgroupIdData = await readUintBuffer( renderer, workgroupIdOut ); + const globalIdData = await readUintBuffer( renderer, globalIdOut ); + const numWorkgroupsData = await readUintBuffer( renderer, numWorkgroupsOut ); + + for ( let i = 0; i < DISPATCH_COUNT; i ++ ) { + + const expectedLocalId = i % WORKGROUP_SIZE; + const expectedWorkgroupId = Math.floor( i / WORKGROUP_SIZE ); + + assert.strictEqual( localIdData[ i ], expectedLocalId, `invocation ${ i }: localId.x` ); + assert.strictEqual( workgroupIdData[ i ], expectedWorkgroupId, `invocation ${ i }: workgroupId.x` ); + // By the WGSL spec, global_invocation_id = workgroup_id * workgroup_size + local_invocation_id + // -- which for a 1D dispatch collapses to exactly the flat invocation index. + assert.strictEqual( globalIdData[ i ], i, `invocation ${ i }: globalId.x should equal workgroupId.x * ${ WORKGROUP_SIZE } + localId.x` ); + assert.strictEqual( numWorkgroupsData[ i ], WORKGROUP_COUNT, `invocation ${ i }: numWorkgroups.x should equal the dispatched workgroup count (${ WORKGROUP_COUNT })` ); + + } + + } ); + + rawComputeTest( 'workgroupId spans all three dimensions for a multi-dimensional dispatch', {}, async ( { assert, renderer } ) => { + + // A small 3x3x1 grid of single-invocation workgroups -- covers + // workgroupId.y (and confirms .z stays 0 for an unused dimension) + // which the 1D test above can't reach. + const gridX = 3; + const gridY = 3; + const count = gridX * gridY; + + const xOut = instancedArray( count, 'uint' ); + const yOut = instancedArray( count, 'uint' ); + const zOut = instancedArray( count, 'uint' ); + + const kernel = Fn( () => { + + // One invocation per workgroup, so instanceIndex directly + // addresses each workgroup's own output slot in row-major + // (x fastest) order, matching how `.compute()` dispatches a + // [gridX, gridY, 1] workgroup grid. + const slot = workgroupId.y.mul( uint( gridX ) ).add( workgroupId.x ); + + xOut.element( slot ).assign( workgroupId.x ); + yOut.element( slot ).assign( workgroupId.y ); + zOut.element( slot ).assign( workgroupId.z ); + + // `.compute()`'s first argument is either a plain invocation + // *count* (number -- dispatches ceil(count / invocationsPerWorkgroup) + // workgroups along X only) or, as used here, an explicit + // per-axis *dispatch size* (array -- one workgroup per grid + // cell, workgroupSize elements per workgroup) -- see + // `ComputeNode`'s `count` vs `dispatchSize` fields. + + } )().compute( [ gridX, gridY, 1 ], [ 1, 1, 1 ] ); + + await renderer.computeAsync( kernel ); + + const xData = await readUintBuffer( renderer, xOut ); + const yData = await readUintBuffer( renderer, yOut ); + const zData = await readUintBuffer( renderer, zOut ); + + for ( let gy = 0; gy < gridY; gy ++ ) { + + for ( let gx = 0; gx < gridX; gx ++ ) { + + const slot = gy * gridX + gx; + + assert.strictEqual( xData[ slot ], gx, `workgroup (${ gx }, ${ gy }): workgroupId.x` ); + assert.strictEqual( yData[ slot ], gy, `workgroup (${ gx }, ${ gy }): workgroupId.y` ); + assert.strictEqual( zData[ slot ], 0, `workgroup (${ gx }, ${ gy }): workgroupId.z should stay 0 (unused dimension)` ); + + } + + } + + } ); + + } ); + +} ); diff --git a/test/unit/addons/tsl/GPUSubgroup.tests.js b/test/unit/addons/tsl/GPUSubgroup.tests.js new file mode 100644 index 00000000000000..29f879cc8e0a42 --- /dev/null +++ b/test/unit/addons/tsl/GPUSubgroup.tests.js @@ -0,0 +1,413 @@ +import { + Fn, + instanceIndex, invocationSubgroupIndex, + instancedArray, subgroupSize, + subgroupAdd, subgroupMul, subgroupMin, subgroupMax, + subgroupAnd, subgroupOr, subgroupXor, + subgroupInclusiveAdd, subgroupExclusiveAdd, + subgroupElect, subgroupBallot, + subgroupBroadcast, subgroupShuffle, + uint, shiftLeft, bool +} from 'three/tsl'; +import { rawComputeTest, readUintBuffer, readIntBuffer } from './gpu-raw-test-utils.js'; + +// Coverage for SubgroupFunctionNode.js -- previously entirely untested (zero +// hits for any `subgroup*`/`quad*` TSL function anywhere in test/). Gated on +// `renderer.hasFeature('subgroups')`: many GPUs/software backends don't +// implement the WebGPU `'subgroups'` feature, in which case every test here +// soft-skips rather than failing (see `rawComputeTest`'s `requiredFeature` +// option) -- confirmed *not* the case in this sandbox (subgroupSize reads +// back as a real, non-skipped 32), so these do exercise real subgroup +// hardware/driver behavior here. +// +// Deliberately NOT covered here: +// - `subgroupAll()`/`subgroupAny()`: calling either with the one boolean +// argument the WGSL spec requires is rejected as declared in this +// codebase (wrong `setParameterLength`) -- see the sibling bugfix branch. +// - `subgroupBroadcastFirst()`: same story, wrong `setParameterLength` the +// other way -- see its own sibling bugfix branch. +// - `quadSwapX/Y/Diagonal` and `quadBroadcast` (also declared in +// SubgroupFunctionNode.js) -- they operate on 2x2 "quad" groupings that +// are a natural fit for fragment-shader derivatives and have no +// well-defined compute-shader lane grouping to test against +// independently of the function under test itself, unlike every function +// covered below. Left as a known gap, not a bug. +// +// Approach: which physical invocations land in the same subgroup, and in +// what lane order, is implementation-defined -- so every test here first +// reads back each invocation's own `invocationSubgroupIndex` (lane id) and +// `subgroupSize` (its subgroup's size) as ground truth topology, *from the +// same dispatch*, and derives the expected reduction/scan/broadcast value +// as a closed-form function of (laneId, groupSize) computed independently +// in JS -- never by re-deriving it from another subgroup call, so these +// can't degrade into test theater (see TSLMath.tests.js's header for the +// same principle applied to plain math functions). + +const WORKGROUP_SIZE = 64; +const WORKGROUP_COUNT = 4; // several workgroups, several subgroups per workgroup +const DISPATCH_COUNT = WORKGROUP_SIZE * WORKGROUP_COUNT; // 256 + +// `.compute(count, ws)` with a plain numeric `count` makes ComputeNode +// auto-insert an `if (instanceIndex < count) { ... }` bounds-check branch +// around the whole kernel body (see ComputeNode.js's `count` vs +// `dispatchSize` doc comments) -- even when, as here, count is an exact +// multiple of the dispatch so no invocation is ever actually excluded. +// WGSL's subgroup functions require being called from *uniform* control +// flow, and that auto-inserted branch is enough to violate it ("must only +// be called from subgroup uniform control flow" -- confirmed by triggering +// it during this file's development). Dispatching via the explicit +// `[workgroupCount, 1, 1]` array form instead skips that guard entirely. +const DISPATCH_SIZE = [ WORKGROUP_COUNT, 1, 1 ]; + +export default QUnit.module( 'TSL', () => { + + QUnit.module( 'subgroup functions', () => { + + rawComputeTest( 'subgroupSize is a plausible, non-skipped value', { requiredFeature: 'subgroups' }, async ( { assert, renderer } ) => { + + const output = instancedArray( DISPATCH_COUNT, 'uint' ); + + const kernel = Fn( () => { + + output.element( instanceIndex ).assign( subgroupSize ); + + } )().compute( DISPATCH_SIZE, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const data = await readUintBuffer( renderer, output ); + + for ( let i = 0; i < DISPATCH_COUNT; i ++ ) { + + assert.ok( data[ i ] >= 1 && data[ i ] <= 128, `invocation ${ i }: subgroupSize (${ data[ i ] }) should be a plausible subgroup size` ); + + } + + } ); + + rawComputeTest( 'subgroupAdd, subgroupMin, subgroupMax reduce across the whole subgroup', { requiredFeature: 'subgroups' }, async ( { assert, renderer } ) => { + + const laneIdOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const groupSizeOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const addOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const minOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const maxOut = instancedArray( DISPATCH_COUNT, 'uint' ); + + const kernel = Fn( () => { + + // value(lane) = lane + 1, so the sum has a simple closed form + // (a triangular number) and min/max are unambiguous (0 is + // never contributed, so min > 0 genuinely exercises the + // reduction rather than trivially reading back a contributed 0). + const value = invocationSubgroupIndex.add( uint( 1 ) ); + + laneIdOut.element( instanceIndex ).assign( invocationSubgroupIndex ); + groupSizeOut.element( instanceIndex ).assign( subgroupSize ); + addOut.element( instanceIndex ).assign( subgroupAdd( value ) ); + minOut.element( instanceIndex ).assign( subgroupMin( value ) ); + maxOut.element( instanceIndex ).assign( subgroupMax( value ) ); + + } )().compute( DISPATCH_SIZE, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const laneIdData = await readUintBuffer( renderer, laneIdOut ); + const groupSizeData = await readUintBuffer( renderer, groupSizeOut ); + const addData = await readUintBuffer( renderer, addOut ); + const minData = await readUintBuffer( renderer, minOut ); + const maxData = await readUintBuffer( renderer, maxOut ); + + for ( let i = 0; i < DISPATCH_COUNT; i ++ ) { + + const n = groupSizeData[ i ]; + const expectedSum = ( n * ( n + 1 ) ) / 2; // sum_{lane=0}^{n-1} (lane+1) + + assert.strictEqual( addData[ i ], expectedSum, `invocation ${ i } (lane ${ laneIdData[ i ] }, group size ${ n }): subgroupAdd` ); + assert.strictEqual( minData[ i ], 1, `invocation ${ i }: subgroupMin should be 1 (lane 0's value)` ); + assert.strictEqual( maxData[ i ], n, `invocation ${ i }: subgroupMax should be ${ n } (last lane's value)` ); + + } + + } ); + + rawComputeTest( 'subgroupMul multiplies contributions from exactly two lanes', { requiredFeature: 'subgroups' }, async ( { assert, renderer } ) => { + + const output = instancedArray( DISPATCH_COUNT, 'uint' ); + + const kernel = Fn( () => { + + // Every lane contributes the multiplicative identity (1) + // except lanes 0 and 1, which contribute 2 and 3 -- keeps the + // product exactly 6 regardless of subgroup size (a genuinely + // unbounded per-lane value would overflow uint32 on wide + // subgroups), while still exercising a real multi-lane + // combination rather than a single-contributor trivial case. + const value = invocationSubgroupIndex + .equal( uint( 0 ) ).select( uint( 2 ), + invocationSubgroupIndex.equal( uint( 1 ) ).select( uint( 3 ), uint( 1 ) ) ); + + output.element( instanceIndex ).assign( subgroupMul( value ) ); + + } )().compute( DISPATCH_SIZE, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const data = await readUintBuffer( renderer, output ); + + for ( let i = 0; i < DISPATCH_COUNT; i ++ ) { + + assert.strictEqual( data[ i ], 6, `invocation ${ i }: subgroupMul should be 2 * 3 * 1^(n-2) = 6` ); + + } + + } ); + + rawComputeTest( 'subgroupAnd, subgroupOr, subgroupXor combine one distinguishing bit per lane', { requiredFeature: 'subgroups' }, async ( { assert, renderer } ) => { + + const groupSizeOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const andOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const orOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const xorOut = instancedArray( DISPATCH_COUNT, 'uint' ); + + const kernel = Fn( () => { + + // Same "one bit per lane" construction as the storage-buffer + // atomic bitwise tests (GPUAtomicsStorage.tests.js), scoped + // to a subgroup instead of the whole dispatch: bit position + // wraps at 32 (a uint's width), matching what the JS-side + // expected-value computation below also wraps at. + const bit = shiftLeft( uint( 1 ), invocationSubgroupIndex.mod( uint( 32 ) ) ); + + groupSizeOut.element( instanceIndex ).assign( subgroupSize ); + // AND starts from all-ones and each lane clears its own bit + // (contributes ~bit, identity 0xffffffff elsewhere). + andOut.element( instanceIndex ).assign( subgroupAnd( bit.bitNot() ) ); + orOut.element( instanceIndex ).assign( subgroupOr( bit ) ); + xorOut.element( instanceIndex ).assign( subgroupXor( bit ) ); + + } )().compute( DISPATCH_SIZE, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const groupSizeData = await readUintBuffer( renderer, groupSizeOut ); + const andData = await readUintBuffer( renderer, andOut ); + const orData = await readUintBuffer( renderer, orOut ); + const xorData = await readUintBuffer( renderer, xorOut ); + + for ( let i = 0; i < DISPATCH_COUNT; i ++ ) { + + const n = groupSizeData[ i ]; + let expectedOr = 0; + let expectedXor = 0; + + for ( let lane = 0; lane < n; lane ++ ) { + + const bit = 1 << ( lane % 32 ); + expectedOr |= bit; + expectedXor ^= bit; + + } + + // AND of (~bit) over every lane clears exactly the bits that + // were contributed by at least one lane -- i.e. the bitwise + // complement of the OR result. + const expectedAnd = ( ~ expectedOr ) >>> 0; + + assert.strictEqual( andData[ i ] >>> 0, expectedAnd, `invocation ${ i } (group size ${ n }): subgroupAnd` ); + assert.strictEqual( orData[ i ] >>> 0, expectedOr >>> 0, `invocation ${ i } (group size ${ n }): subgroupOr` ); + assert.strictEqual( xorData[ i ] >>> 0, expectedXor >>> 0, `invocation ${ i } (group size ${ n }): subgroupXor` ); + + } + + } ); + + rawComputeTest( 'subgroupInclusiveAdd and subgroupExclusiveAdd compute correct prefix sums', { requiredFeature: 'subgroups' }, async ( { assert, renderer } ) => { + + const laneIdOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const inclusiveOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const exclusiveOut = instancedArray( DISPATCH_COUNT, 'uint' ); + + const kernel = Fn( () => { + + const value = invocationSubgroupIndex.add( uint( 1 ) ); // 1, 2, 3, ... + + laneIdOut.element( instanceIndex ).assign( invocationSubgroupIndex ); + inclusiveOut.element( instanceIndex ).assign( subgroupInclusiveAdd( value ) ); + exclusiveOut.element( instanceIndex ).assign( subgroupExclusiveAdd( value ) ); + + } )().compute( DISPATCH_SIZE, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const laneIdData = await readUintBuffer( renderer, laneIdOut ); + const inclusiveData = await readUintBuffer( renderer, inclusiveOut ); + const exclusiveData = await readUintBuffer( renderer, exclusiveOut ); + + for ( let i = 0; i < DISPATCH_COUNT; i ++ ) { + + const l = laneIdData[ i ]; + // inclusive prefix sum of (1, 2, ..., l+1) = (l+1)(l+2)/2 + const expectedInclusive = ( ( l + 1 ) * ( l + 2 ) ) / 2; + // exclusive prefix sum of (1, 2, ..., l) = l(l+1)/2 (0 at lane 0) + const expectedExclusive = ( l * ( l + 1 ) ) / 2; + + assert.strictEqual( inclusiveData[ i ], expectedInclusive, `invocation ${ i } (lane ${ l }): subgroupInclusiveAdd` ); + assert.strictEqual( exclusiveData[ i ], expectedExclusive, `invocation ${ i } (lane ${ l }): subgroupExclusiveAdd` ); + + } + + } ); + + // subgroupAll()/subgroupAny() are NOT covered here: calling either + // with the one boolean predicate argument the WGSL spec requires + // (`subgroupAll(e: bool) -> bool`) is rejected by this codebase with + // "parameter length exceeds limit" -- both are declared with + // `setParameterLength(0)` in SubgroupFunctionNode.js, so as shipped + // neither is callable for its actual purpose. See the sibling branch + // with the fix + the (now-passing) test for this. + + rawComputeTest( 'subgroupElect is true for exactly lane 0', { requiredFeature: 'subgroups' }, async ( { assert, renderer } ) => { + + const laneIdOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const electOut = instancedArray( DISPATCH_COUNT, 'uint' ); + + const kernel = Fn( () => { + + laneIdOut.element( instanceIndex ).assign( invocationSubgroupIndex ); + electOut.element( instanceIndex ).assign( subgroupElect().select( uint( 1 ), uint( 0 ) ) ); + + } )().compute( DISPATCH_SIZE, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const laneIdData = await readUintBuffer( renderer, laneIdOut ); + const electData = await readUintBuffer( renderer, electOut ); + + for ( let i = 0; i < DISPATCH_COUNT; i ++ ) { + + const expected = laneIdData[ i ] === 0 ? 1 : 0; + assert.strictEqual( electData[ i ], expected, `invocation ${ i } (lane ${ laneIdData[ i ] }): subgroupElect should be true only for lane 0` ); + + } + + } ); + + rawComputeTest( 'subgroupBallot sets exactly bits [0, groupSize) for an always-true predicate', { requiredFeature: 'subgroups' }, async ( { assert, renderer } ) => { + + const groupSizeOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const ballotXOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const ballotYOut = instancedArray( DISPATCH_COUNT, 'uint' ); + + const kernel = Fn( () => { + + const ballot = subgroupBallot( bool( true ) ); + + groupSizeOut.element( instanceIndex ).assign( subgroupSize ); + // A subgroup wider than 32 would need .z/.w too -- this + // sandbox's subgroupSize (32) only ever needs .x, and .y + // should stay 0, which the check below asserts explicitly. + ballotXOut.element( instanceIndex ).assign( ballot.x ); + ballotYOut.element( instanceIndex ).assign( ballot.y ); + + } )().compute( DISPATCH_SIZE, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const groupSizeData = await readUintBuffer( renderer, groupSizeOut ); + const ballotXData = await readUintBuffer( renderer, ballotXOut ); + const ballotYData = await readUintBuffer( renderer, ballotYOut ); + + for ( let i = 0; i < DISPATCH_COUNT; i ++ ) { + + const n = groupSizeData[ i ]; + + assert.ok( n <= 32, `invocation ${ i }: this test only checks .x/.y -- group size ${ n } would need .z/.w too` ); + + // Per the WGSL spec, subgroupBallot's nth bit corresponds to + // the invocation with subgroup_invocation_id == n -- so for + // an always-true predicate, bits [0, n) are set. + const expectedX = n >= 32 ? 0xffffffff : ( ( 1 << n ) - 1 ) >>> 0; + + assert.strictEqual( ballotXData[ i ] >>> 0, expectedX, `invocation ${ i } (group size ${ n }): subgroupBallot(true).x` ); + assert.strictEqual( ballotYData[ i ], 0, `invocation ${ i }: subgroupBallot(true).y should be 0 for a <=32-lane subgroup` ); + + } + + } ); + + // subgroupBroadcastFirst() is NOT covered here: calling it with the + // one argument the WGSL spec requires (`subgroupBroadcastFirst(e: T) + // -> T` -- no lane id, unlike subgroupBroadcast) is rejected by this + // codebase with "parameter length is less than minimum required" -- + // it's declared with `setParameterLength(2)` in + // SubgroupFunctionNode.js, so three.js auto-pads a bogus second + // argument, producing invalid WGSL. See the sibling branch with the + // fix + the (now-passing) test for this. + + rawComputeTest( 'subgroupBroadcast reads a specific lane\'s value', { requiredFeature: 'subgroups' }, async ( { assert, renderer } ) => { + + const broadcastOut = instancedArray( DISPATCH_COUNT, 'uint' ); + + const kernel = Fn( () => { + + const value = invocationSubgroupIndex.add( uint( 100 ) ); // 100, 101, 102, ... + + // Every lane asks for lane 0's value explicitly. + broadcastOut.element( instanceIndex ).assign( subgroupBroadcast( value, uint( 0 ) ) ); + + } )().compute( DISPATCH_SIZE, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const broadcastData = await readUintBuffer( renderer, broadcastOut ); + + for ( let i = 0; i < DISPATCH_COUNT; i ++ ) { + + assert.strictEqual( broadcastData[ i ], 100, `invocation ${ i }: subgroupBroadcast(value, 0) should read back lane 0's value (100)` ); + + } + + } ); + + rawComputeTest( 'subgroupShuffle reverses lane order within the subgroup', { requiredFeature: 'subgroups' }, async ( { assert, renderer } ) => { + + const laneIdOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const groupSizeOut = instancedArray( DISPATCH_COUNT, 'uint' ); + const shuffledOut = instancedArray( DISPATCH_COUNT, 'int' ); + + const kernel = Fn( () => { + + const laneId = invocationSubgroupIndex; + const targetLane = subgroupSize.sub( uint( 1 ) ).sub( laneId ); + + laneIdOut.element( instanceIndex ).assign( laneId ); + groupSizeOut.element( instanceIndex ).assign( subgroupSize ); + // Each lane fetches the value from its mirror-image lane + // (targetLane) -- so the value it fetches is that lane's own + // id, and the expected result is fully determined by + // (laneId, groupSize) alone. + shuffledOut.element( instanceIndex ).assign( subgroupShuffle( laneId.toInt(), targetLane ) ); + + } )().compute( DISPATCH_SIZE, [ WORKGROUP_SIZE ] ); + + await renderer.computeAsync( kernel ); + + const laneIdData = await readUintBuffer( renderer, laneIdOut ); + const groupSizeData = await readUintBuffer( renderer, groupSizeOut ); + const shuffledData = await readIntBuffer( renderer, shuffledOut ); + + for ( let i = 0; i < DISPATCH_COUNT; i ++ ) { + + const laneId = laneIdData[ i ]; + const n = groupSizeData[ i ]; + const expected = n - 1 - laneId; + + assert.strictEqual( shuffledData[ i ], expected, `invocation ${ i } (lane ${ laneId }, group size ${ n }): subgroupShuffle should fetch the mirror lane's id (${ expected })` ); + + } + + } ); + + } ); + +} ); diff --git a/test/unit/addons/tsl/GPUWorkgroupAtomic.tests.js b/test/unit/addons/tsl/GPUWorkgroupAtomic.tests.js new file mode 100644 index 00000000000000..1dc934da348298 --- /dev/null +++ b/test/unit/addons/tsl/GPUWorkgroupAtomic.tests.js @@ -0,0 +1,146 @@ +import { + Fn, + instanceIndex, invocationLocalIndex, + instancedArray, workgroupArray, workgroupBarrier, + atomicAdd, atomicLoad, uint +} from 'three/tsl'; +import { getSharedRenderer } from './gpu-test-utils.js'; + +// `readBuffer()` in gpu-test-utils.js always reinterprets the raw bytes as a +// `Float32Array`, which is wrong here: `output` below is a `uint` storage +// buffer, so its bytes must be read back as `Uint32Array` -- reinterpreting +// small uint values (e.g. 8) as float32 bit patterns would produce garbage, +// not the integers under test. +async function readUintBuffer( renderer, buffer ) { + + return new Uint32Array( await renderer.getArrayBufferAsync( buffer.value ) ); + +} + +// Coverage for workgroup-scoped shared arrays and atomics +// (`workgroupArray()`, `.toAtomic()`, `atomicAdd()`, `workgroupBarrier()`) -- +// see mrdoob/three.js#34428, which added `.toAtomic()` support to +// `workgroupArray()`. None of this is exercised by `gpuTest`/`gpuFuzzTest` +// (gpu-test-utils.js): those dispatch one compute invocation per assertion +// and never rely on multiple invocations of the *same* workgroup cooperating +// through shared/atomic memory, which is exactly the behavior these nodes +// add. So these tests build their own small compute kernels directly, +// dispatching several invocations per workgroup and reading the shared +// result back once per workgroup finishes. +// +// WebGPU-only: `workgroupArray()`/atomics aren't implemented for the WebGL2 +// fallback backend (no `getScopedArray()` there), so these only register +// against the 'webgpu' backend. + +export default QUnit.module( 'TSL', () => { + + QUnit.module( 'workgroup arrays and atomics', () => { + + QUnit.test( 'workgroupArray: plain (non-atomic) shared read/write survives a barrier', async ( assert ) => { + + const renderer = await getSharedRenderer( 'webgpu' ); + + if ( renderer === null ) { + + assert.ok( true, 'SKIPPED: "webgpu" backend is not available in this environment.' ); + return; + + } + + // Regression guard: PR #34428 threads an `isAtomic` flag through + // `WorkgroupInfoNode`/`WGSLNodeBuilder.getScopedArray()` -- this + // confirms the default (non-atomic) path still declares and + // round-trips a plain `array` in the `workgroup` address + // space exactly as before. + const workgroupSize = 8; + const dispatchCount = 32; // 4 workgroups of 8 + const output = instancedArray( dispatchCount, 'uint' ); + + const kernel = Fn( () => { + + const shared = workgroupArray( 'uint', workgroupSize ); + + // Every invocation writes a distinct non-zero value into its own + // slot, then reads back the value written by its right-hand + // neighbor. Using non-zero values ensures this cannot pass from + // workgroup memory's default initialization alone. + shared.element( invocationLocalIndex ).assign( invocationLocalIndex.add( uint( 1 ) ) ); + + workgroupBarrier(); + + const neighborLocalIndex = invocationLocalIndex.add( uint( 1 ) ).mod( uint( workgroupSize ) ); + output.element( instanceIndex ).assign( shared.element( neighborLocalIndex ) ); + + } )().compute( dispatchCount, [ workgroupSize ] ); + + await renderer.computeAsync( kernel ); + + const data = await readUintBuffer( renderer, output ); + + for ( let i = 0; i < dispatchCount; i ++ ) { + + const localIndex = i % workgroupSize; + const expected = ( ( localIndex + 1 ) % workgroupSize ) + 1; + + assert.strictEqual( data[ i ], expected, `invocation ${ i }: should read its right neighbor's non-zero value (${ expected })` ); + + } + + } ); + + QUnit.test( 'workgroupArray.toAtomic(): concurrent atomicAdd sums exactly once per invocation, per workgroup', async ( assert ) => { + + const renderer = await getSharedRenderer( 'webgpu' ); + + if ( renderer === null ) { + + assert.ok( true, 'SKIPPED: "webgpu" backend is not available in this environment.' ); + return; + + } + + const workgroupSize = 8; + const workgroupCount = 4; + const dispatchCount = workgroupSize * workgroupCount; + const output = instancedArray( dispatchCount, 'uint' ); + + const kernel = Fn( () => { + + // A single atomic counter, shared by the whole workgroup. WGSL + // zero-initializes `workgroup`-address-space variables (atomic + // ones included) once per workgroup, so no explicit reset is + // needed here -- and a plain `.assign()` on an atomic element + // wouldn't be valid WGSL anyway (writes must go through + // `atomicStore`/`atomicAdd`/etc.). + const counter = workgroupArray( 'uint', 1 ).toAtomic(); + + // Every invocation in the workgroup increments the same + // atomic slot concurrently -- this only produces the exact + // expected sum if `getScopedArray()` genuinely declared the + // element as `atomic` (PR #34428); a plain (non-atomic) + // `array` here would race and typically undercount. + atomicAdd( counter.element( uint( 0 ) ), uint( 1 ) ); + + workgroupBarrier(); + + // WGSL forbids reading an `atomic` element with a plain + // load/assign -- it must go through `atomicLoad()`. + output.element( instanceIndex ).assign( atomicLoad( counter.element( uint( 0 ) ) ) ); + + } )().compute( dispatchCount, [ workgroupSize ] ); + + await renderer.computeAsync( kernel ); + + const data = await readUintBuffer( renderer, output ); + + for ( let i = 0; i < dispatchCount; i ++ ) { + + assert.strictEqual( data[ i ], workgroupSize, `invocation ${ i }: atomic counter should equal workgroupSize (${ workgroupSize }) -- every invocation in its workgroup added exactly once` ); + + } + + } ); + + } ); + +} ); diff --git a/test/unit/addons/tsl/gpu-raw-test-utils.js b/test/unit/addons/tsl/gpu-raw-test-utils.js new file mode 100644 index 00000000000000..d6553673d445c2 --- /dev/null +++ b/test/unit/addons/tsl/gpu-raw-test-utils.js @@ -0,0 +1,64 @@ +// +// Shared boilerplate for GPU-native compute tests that need full control over +// dispatch shape (workgroupSize, multiple workgroups, raw typed-array +// readback) -- i.e. tests that can't be expressed as `gpuTest`/`gpuFuzzTest` +// assertions (gpu-test-utils.js), because those dispatch exactly one compute +// invocation per assertion and never let multiple invocations of the same +// workgroup cooperate through shared/atomic memory or inspect built-in +// workgroup/subgroup identity values. +// +// `rawComputeTest(name, options, run)` registers one QUnit.test that: +// - resolves the shared renderer for `options.backend` (default 'webgpu'), +// soft-skipping (not failing) if that backend isn't available here -- +// same policy as gpu-test-utils.js's `declareTest`. +// - soft-skips if `options.requiredFeature` is set and the renderer doesn't +// report it (`renderer.hasFeature(...)`) -- used for subgroup tests, which +// many GPUs/software backends don't implement. +// - calls `run({ assert, renderer })`; the test body builds and dispatches +// its own TSL compute kernel and reads results back itself. +// +import { getSharedRenderer } from './gpu-test-utils.js'; + +export function rawComputeTest( name, options, run ) { + + const { backend = 'webgpu', requiredFeature } = options; + + QUnit.test( name, async ( assert ) => { + + const renderer = await getSharedRenderer( backend ); + + if ( renderer === null ) { + + assert.ok( true, `SKIPPED: "${ backend }" backend is not available in this environment.` ); + return; + + } + + if ( requiredFeature !== undefined && renderer.hasFeature( requiredFeature ) !== true ) { + + assert.ok( true, `SKIPPED: "${ backend }" backend does not support required feature "${ requiredFeature }" in this environment.` ); + return; + + } + + await run( { assert, renderer } ); + + } ); + +} + +// Every uint/int storage buffer read back in these tests needs the *typed +// integer* view of the raw bytes, not `gpu-test-utils.js`'s `readBuffer()` +// (which always reinterprets as `Float32Array` -- reinterpreting a small +// uint like 8 as float32 bits produces garbage, not the integer under test). +export async function readUintBuffer( renderer, buffer ) { + + return new Uint32Array( await renderer.getArrayBufferAsync( buffer.value ) ); + +} + +export async function readIntBuffer( renderer, buffer ) { + + return new Int32Array( await renderer.getArrayBufferAsync( buffer.value ) ); + +} diff --git a/test/unit/addons/tsl/gpu-test-utils.js b/test/unit/addons/tsl/gpu-test-utils.js index 04fecfa6ea019a..34b2120be331f5 100644 --- a/test/unit/addons/tsl/gpu-test-utils.js +++ b/test/unit/addons/tsl/gpu-test-utils.js @@ -242,7 +242,7 @@ const BACKEND_OPTIONS = { // so availability is detected empirically per backend rather than assumed. const sharedRenderers = {}; -async function getSharedRenderer( backend ) { +export async function getSharedRenderer( backend ) { if ( BACKEND_OPTIONS[ backend ] === undefined ) { diff --git a/test/unit/three.addons.unit.js b/test/unit/three.addons.unit.js index 03df4a20ed2db9..bd5f2526427394 100644 --- a/test/unit/three.addons.unit.js +++ b/test/unit/three.addons.unit.js @@ -16,6 +16,11 @@ import './addons/loaders/USDLoader.tests.js'; import './addons/exporters/USDZExporter.tests.js'; import './addons/tsl/WebGLNodesHandler.tests.js'; import './addons/tsl/GPUTest.tests.js'; +import './addons/tsl/GPUAtomicsStorage.tests.js'; +import './addons/tsl/GPUComputeBuiltins.tests.js'; +import './addons/tsl/GPUBarriers.tests.js'; +import './addons/tsl/GPUSubgroup.tests.js'; +import './addons/tsl/GPUWorkgroupAtomic.tests.js'; import './addons/tsl/TSLDeterminant.tests.js'; import './addons/tsl/TSLFaceForward.tests.js'; import './addons/tsl/TSLGainPcurve.tests.js';