-
-
Notifications
You must be signed in to change notification settings - Fork 35.9k
Introduction of parameterizable workgroups for compute shaders and dispatch sizes for the renderer #31402
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
base: dev
Are you sure you want to change the base?
Introduction of parameterizable workgroups for compute shaders and dispatch sizes for the renderer #31402
Changes from 10 commits
d893334
334b61a
910ce62
58af151
227f25a
0d2d7fd
d44829a
2332ea5
14bd77f
b55b9e9
93850e9
9359fa8
b034372
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
Original file line number | Diff line number | Diff line change |
---|---|---|
|
@@ -1318,10 +1318,15 @@ class WebGPUBackend extends Backend { | |
* @param {Node} computeNode - The compute node. | ||
* @param {Array<BindGroup>} bindings - The bindings. | ||
* @param {ComputePipeline} pipeline - The compute pipeline. | ||
* @param {Array<number>} dispatchSize - Array with [x,y,z] values for dispatch. | ||
*/ | ||
compute( computeGroup, computeNode, bindings, pipeline ) { | ||
compute( computeGroup, computeNode, bindings, pipeline, dispatchSize ) { | ||
|
||
const computeNodeData = this.get( computeNode ); | ||
const { passEncoderGPU } = this.get( computeGroup ); | ||
const isValid = dispatchSize[ 0 ] > 0 && dispatchSize[ 1 ] > 0 && dispatchSize[ 2 ] > 0; | ||
|
||
dispatchSize = isValid ? dispatchSize : computeNodeData; | ||
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. It's valid to set @workgroup_size to 1 or 2 values when defining the parameters inline with the remaining dimensions implicitly being set to "1" (see the spec here) so it would be nice to allow for providing workgroup size of 1 or 2 to align with the native behavior and for the sake of ergonomics: const computeShader = wgslFn( `fn fragmentShader() -> void {}` );
// all are equivalant
const kernel = computeShader().compute( [ 16 ] );
const kernel = computeShader().compute( [ 16, 1 ] );
const kernel = computeShader().compute( [ 16, 1, 1 ] ); There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. I didn't know that from the W3C documentation, I'll take a look at it, thanks There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. This can be implemented in two ways. Since it makes no difference in performance, I prefer to check the length of the workgroupSize array before calling the wgsl code builder to fill it with ones if smaller than a three-valued array |
||
|
||
// pipeline | ||
|
||
|
@@ -1340,30 +1345,38 @@ class WebGPUBackend extends Backend { | |
|
||
} | ||
|
||
const maxComputeWorkgroupsPerDimension = this.device.limits.maxComputeWorkgroupsPerDimension; | ||
if ( isValid ) { | ||
|
||
const computeNodeData = this.get( computeNode ); | ||
passEncoderGPU.dispatchWorkgroups( | ||
dispatchSize[ 0 ], | ||
dispatchSize[ 1 ], | ||
dispatchSize[ 2 ] | ||
); | ||
|
||
if ( computeNodeData.dispatchSize === undefined ) computeNodeData.dispatchSize = { x: 0, y: 1, z: 1 }; | ||
} else { | ||
|
||
const { dispatchSize } = computeNodeData; | ||
const maxComputeWorkgroupsPerDimension = this.device.limits.maxComputeWorkgroupsPerDimension; | ||
|
||
if ( computeNode.dispatchCount > maxComputeWorkgroupsPerDimension ) { | ||
if ( computeNodeData.dispatchSize === undefined ) computeNodeData.dispatchSize = { x: 0, y: 1, z: 1 }; | ||
|
||
dispatchSize.x = Math.min( computeNode.dispatchCount, maxComputeWorkgroupsPerDimension ); | ||
dispatchSize.y = Math.ceil( computeNode.dispatchCount / maxComputeWorkgroupsPerDimension ); | ||
if ( computeNode.dispatchCount > maxComputeWorkgroupsPerDimension ) { | ||
|
||
} else { | ||
dispatchSize.x = Math.min( computeNode.dispatchCount, maxComputeWorkgroupsPerDimension ); | ||
dispatchSize.y = Math.ceil( computeNode.dispatchCount / maxComputeWorkgroupsPerDimension ); | ||
|
||
} else { | ||
|
||
dispatchSize.x = computeNode.dispatchCount; | ||
dispatchSize.x = computeNode.dispatchCount; | ||
|
||
} | ||
} | ||
|
||
passEncoderGPU.dispatchWorkgroups( | ||
dispatchSize.x, | ||
dispatchSize.y, | ||
dispatchSize.z | ||
); | ||
passEncoderGPU.dispatchWorkgroups( | ||
dispatchSize.x, | ||
dispatchSize.y, | ||
dispatchSize.z | ||
); | ||
|
||
} | ||
|
||
} | ||
|
||
|
Original file line number | Diff line number | Diff line change |
---|---|---|
|
@@ -1898,7 +1898,11 @@ ${ flowData.code } | |
|
||
} else { | ||
|
||
this.computeShader = this._getWGSLComputeCode( shadersData.compute, ( this.object.workgroupSize || [ 64 ] ).join( ', ' ) ); | ||
const workgroupSize = this.object.workgroupSize || [ 64, 1, 1 ]; | ||
|
||
if ( workgroupSize.length !== 3 ) throw new Error( 'workgroupSize must have 3 elements' ); | ||
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Likewise dispatchWorkgroups can also take 1 or 2 parameters with the remainder defaulting to 1: // all equivelant
renderer.computeAsync( kernel, [ 10 ] );
renderer.computeAsync( kernel, [ 10, 1 ] );
renderer.computeAsync( kernel, [ 10, 1, 1 ] ); I'm curious to hear other opinions There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Yes, that's simple. I can do it tonight ( CET ). If you have any other requests by then, I can include them as well |
||
|
||
this.computeShader = this._getWGSLComputeCode( shadersData.compute, workgroupSize ); | ||
|
||
} | ||
|
||
|
@@ -2103,6 +2107,8 @@ fn main( ${shaderData.varyings} ) -> ${shaderData.returnType} { | |
*/ | ||
_getWGSLComputeCode( shaderData, workgroupSize ) { | ||
|
||
const [ workgroupSizeX, workgroupSizeY, workgroupSizeZ ] = workgroupSize; | ||
|
||
return `${ this.getSignature() } | ||
// directives | ||
${shaderData.directives} | ||
|
@@ -2122,11 +2128,13 @@ ${shaderData.uniforms} | |
// codes | ||
${shaderData.codes} | ||
|
||
@compute @workgroup_size( ${workgroupSize} ) | ||
@compute @workgroup_size( ${workgroupSizeX}, ${workgroupSizeY}, ${workgroupSizeZ} ) | ||
fn main( ${shaderData.attributes} ) { | ||
|
||
// system | ||
instanceIndex = globalId.x + globalId.y * numWorkgroups.x * u32(${workgroupSize}) + globalId.z * numWorkgroups.x * numWorkgroups.y * u32(${workgroupSize}); | ||
instanceIndex = globalId.x | ||
+ globalId.y * (${workgroupSizeX} * numWorkgroups.x) | ||
+ globalId.z * (${workgroupSizeX} * numWorkgroups.x) * (${workgroupSizeY} * numWorkgroups.y); | ||
|
||
// vars | ||
${shaderData.vars} | ||
|
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Assuming the project agrees with the API change I'm wondering if we should deprecate the old code path where "count" is passed here?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I've set it up so that count is used if no workgroup is specified, so that existing user code and examples continue to run as before. But yes, I would also be in favor of removing the count from the code in the long term, with the usual transition time. This also significantly simplifies the compute call in the backend.