Issue 13406 in skia: Add a directive for workgroup size

[ethan… via monorail]

Status: Accepted
Owner: ethann...@google.com
Priority: Medium
Type: Defect

New issue 13406 by ethann...@google.com: Add a directive for workgroup size
https://bugs.chromium.org/p/skia/issues/detail?id=13406

We are currently relying on the compute shader workgroup size being set via APIs on the Skia side of things, but in shading languages other than Metal, the workgroup size is specified in the shader itself.

We should probably handle this by adding a language directive allowing us to specify this, and then feeding it back to Skia somehow.

--
You received this message because:
1. The project was configured to send all issue notifications to this address

You may adjust your notification preferences at:
https://bugs.chromium.org/hosting/settings

[arman… via monorail]

Updates:
Cc: johns...@google.com brian...@google.com arma...@google.com
Labels: GPU-Compute Area-SkSL
Owner: ----

Comment #1 on issue 13406 by arma...@google.com: Add a directive for workgroup size
https://bugs.chromium.org/p/skia/issues/detail?id=13406#c1

(No comment was entered for this change.)

[arman… via monorail]

Updates:
Cc: -arma...@google.com
Owner: arma...@google.com

Comment #2 on issue 13406 by arma...@google.com: Add a directive for workgroup size
https://bugs.chromium.org/p/skia/issues/detail?id=13406#c2

Assigning to myself for tracking purposes.

[arman… via monorail]

Issue 13406: Add a directive for workgroup size
https://bugs.chromium.org/p/skia/issues/detail?id=13406

This issue is now blocking issue 13578.
See https://bugs.chromium.org/p/skia/issues/detail?id=13578

[arman… via monorail]

Updates:
Cc: jvan...@google.com

Comment #4 on issue 13406 by arma...@google.com: Add a directive for workgroup size
https://bugs.chromium.org/p/skia/issues/detail?id=13406#c4

In Metal, Vulkan, WebGPU, and OpenGL compute all impose restrictions on the workgroup size (and the total number of invocations per workgroup). At least in Vulkan and OpenGL, a shader that specifies dimensions that exceed the restrictions can result in a shader compilation error. In all backends the client has to query the platform for allowable dimensions and make sure the local size used in the shaders (and in the dispatch call with Metal) are within the limits.

WGSL, SPIR-V 1.2, and GLSL (using GL_KHR_vulkan_glsl) provide ways to specify the workgroup size dimensions via specialization constants as well as statically. We should consider supporting the specialization constant approach in SkSL if we want to support executing a compute pipeline state using different dimensions without having to recompile it (though I imagine we'll often end up selecting the same set of values for a particular GPU).

graphite::ComputePipeline or its backend subclasses may inject the workgroup size directive during pipeline creation. I think the SkSL backends should generate an error if they don't support a particular workgroup size syntax (static vs specialization constant) instead of simply ignoring it, which means each ComputePipeline subclass may generate different SkSL (e.g. MtlComputePipeline may not inject any workgroup dimensions, a hypothetical OpenGL compute pipeline may only allow static declarations, etc). Either way, the pipeline object should calculate/restrict the appropriate dimensions based on GPU capabilities to make the effective number available for command encoding on Metal.

----
Syntax notes:

GLSL, WGSL, and SPIR-V allow the width/height/depth dimensions to be specified individually. An omitted dimension is assigned a default value of 1.

GLSL uses program-scope layout specifiers:
1. With static values (from https://www.khronos.org/opengl/wiki/Compute_Shader#Local_size):
layout(local_size_x = X​, local_size_y = Y​, local_size_z = Z​) in;

2. With specialization constants (from https://github.com/KhronosGroup/GLSL/blob/master/extensions/khr/GL_KHR_vulkan_glsl.txt). Y/height is not a specialization constant in this example:
layout(local_size_x_id = 18, local_size_z_id = 19) in;

WGSL uses entry point attributes (from https://www.w3.org/TR/WGSL/#entry-point-attributes):
1. With static values:
@compute @workgroup_size(x, y, z)
fn main() {...}

2. With "pipeline-overridable constants" (specifying only width; height and depth are 1):
@id(42) override block_width = 12u; // Default value is 12, slot of the constant is 42
@compute @workgroup_size(block_width)

SPIR-V uses per-entry-point execution modes using the OpExecutionMode:
1. For static values, LocalSize using literals for x/y/z
2. For specialization constants, LocalSizeId using IDs referencing x/y/z.

[johns… via monorail]

Updates:
Owner: johns...@google.com

Comment #5 on issue 13406 by johns...@google.com: Add a directive for workgroup size
https://bugs.chromium.org/p/skia/issues/detail?id=13406#c5

(No comment was entered for this change.)

[johns… via monorail]

Updates:
Labels: Type-Feature

Comment #6 on issue 13406 by johns...@google.com: Add a directive for workgroup size
https://bugs.chromium.org/p/skia/issues/detail?id=13406#c6

[johns… via monorail]

Updates:
Owner: arma...@google.com

Comment #7 on issue 13406 by johns...@google.com: Add a directive for workgroup size
https://bugs.chromium.org/p/skia/issues/detail?id=13406#c7

[arman… via monorail]

Updates:
Status: Started

Comment #8 on issue 13406 by arma...@google.com: Add a directive for workgroup size
https://bugs.chromium.org/p/skia/issues/detail?id=13406#c8