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]

I've been working on making a constant work. It's hard because rustc doesn't support constants in attrs, so the underlying compiler infra isn't there.

[nazar-pc]

What about creating a constant or variable at the beginning of the function and removing the argument?

Basically translating this:

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

Into this:

#[spirv(compute(threads(8, 4, 2)))]
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);

Could make it constant too, but compiler should be able to inline it anyway and using a constant will require linter suppression due to variable being not uppercase.

[nazar-pc]

The only challenge with my example above is when number of threads is not specified, then I think it'll have to spit out something in SPIR-V, though it is certainly outside of my extremely limited knowledge, but it should at least cover the simplest cases.

If that works, then I'd love to also see things like global_size and global_linear_id emulated too (or maybe adding additional better named options in addition to those that are in the spec since names in the spec are horrible and not even consistent).

[LegNeato]

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

[LegNeato]

@nazar-pc #298