workgroup_size / gl_WorkGroupSize support

[LegNeato]

I see this is commented out:

rust-gpu/crates/rustc_codegen_spirv/src/symbols.rs

Line 62 in 30024f8

// ("workgroup_size", WorkgroupSize), -- constant

Is there a way to get this info without hardcoding values? You can't include constants / variables in the attribute, so I see no way to have this in one place and instead you have to keep two places in sync manually.

[nazar-pc]

I guess one can combine num_subgroups and subgroup_size, but I find it confusing that workgroup_size isn't a thing, as a developer I would have expected it to be available in arguments, even if it is in fact a constant.

Especially since it is apparently possible to not define number of threads in the workgroup at all, in which case it would be very handy to have workgroup_size readily available.

UPD: Annoyinggly, #[spirv(compute(threads(256)))] doesn't allow replacing 256 with a constant, which would would be a decent workaround, and actually a nice solution because the same constant can then be shared with the host side too.

[LegNeato]

no, supporting workgroupsize should be fine, it's the using a rust constant in the attr which is annoying.

[LegNeato]

@nazar-pc #298

[charles-r-earp]

The WorkgroupSize builtin is used to specify the number of threads, not as an input like the others. This is overrides LocalSize, and is superseded by LocalSizeId, introduced in 1.2. When using WorkgroupSize, LocalSize execution mode is used, but its value is ignored.

This might be easier as a function:

#[spirv(compute)]
pub fn main_cs(
    #[spirv(global_invocation_id)] id: UVec3,
    #[spirv(storage_buffer, descriptor_set = 0, binding = 0)] prime_indices: &mut [u32],
) {
let workgroup_size = UVec3::new(8, 4, 2);
spirv_std:arch::workgroup_size(workgroup_size); 

It would be nice if this was mapped into LocalSizeId when supported, as WorkgroupSize is deprecated.

[LegNeato]

You are talking about vulkan's WorkGroupSize, we are talking about https://registry.khronos.org/OpenGL-Refpages/gl4/html/gl_WorkGroupSize.xhtml. Those are different with different semantics AFAICT. I didn't need to annotate the synthesized constant with WorkGroupSize just for querying but while I was there I figured why not set it.

For specifying, I'd much rather do what is in #299

[charles-r-earp]

This

#version 450

layout(local_size_x = 10) in;

void main() {}

Compiles to this

; SPIR-V
; Version: 1.0
; Generator: Google Shaderc over Glslang; 10
; Bound: 11
; Schema: 0
               OpCapability Shader
          %1 = OpExtInstImport "GLSL.std.450"
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %main "main"
               OpExecutionMode %main LocalSize 10 1 1
               OpSource GLSL 450
               OpSourceExtension "GL_GOOGLE_cpp_style_line_directive"
               OpSourceExtension "GL_GOOGLE_include_directive"
               OpName %main "main"
               OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
       %void = OpTypeVoid
          %3 = OpTypeFunction %void
       %uint = OpTypeInt 32 0
     %v3uint = OpTypeVector %uint 3
    %uint_10 = OpConstant %uint 10
     %uint_1 = OpConstant %uint 1
%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_10 %uint_1 %uint_1
       %main = OpFunction %void None %3
          %5 = OpLabel
               OpReturn
               OpFunctionEnd