Fix the compute limits, adds maxComputeWorkgroupSize

[Kangz]

Validation was incorrectly checking values were above limits instead of
below limits.

maxComputeWorkgroupSize is necessary because the 3rd component of
workgroup_size is almost always 64 in Vulkan devices (on mobile and
desktop).

---

Preview | Diff

[github-actions]

Previews, as seen when this build job started (7753ba0):
WebGPU | IDL
WGSL
Explainer

[Kangz]

Thanks for the suggestions, applied all of them.

[github-actions]

Previews, as seen when this build job started (e53244c):
WebGPU | IDL
WGSL
Explainer

[github-actions]

Previews, as seen when this build job started (51d1835):
WebGPU | IDL
WGSL
Explainer

[github-actions]

Previews, as seen when this build job started (52aa85e):
WebGPU | IDL
WGSL
Explainer

[github-actions]

Previews, as seen when this build job started (52aa85e):
WebGPU | IDL
WGSL
Explainer

[alan-baker]

Out of curiosity, why are the base limits for workgroup size and invocations higher than Vulkan? Vulkan only provides (128,128,64) for size and 128 for invocations by default.

[Kangz]

It is because looking at vulkan.gpuinfo.org the only devices that don't support (256, 256, 64) are Swiftshader (which I just fixed in this CL) and a few ARM devices that correspond to very old Mali.

We try to support a wide variety of Android Vulkan devices but are ok removing some percentage of them to raise WebGPU's core feature level. The others can be supported through lower-feature levels on GLES in the future.

[github-actions]

Previews, as seen when this build job started (5eb5a41):
WebGPU | IDL
WGSL
Explainer

[github-actions]

Previews, as seen when this build job started (8829596):
WebGPU | IDL
WGSL
Explainer

[kainino0x]

I'm curious whether there are any devices where maxComputeWorkgroupSize.x or maxComputeWorkgroupSize.y is less than maxComputeWorkgroupInvocations (or if there are, whether it matters). It seems unlikely, at least for x. Maybe we only need maxComputeWorkgroupInvocations and maxComputeWorkgroupSizeZ?

We should also add some kind of guarantee that the adapter limits make sense maxX <= maxXYZ && maxY <= maxXYZ && maxZ <= maxXYZ and probably maxXYZ <= maxX * maxY * maxZ though it matters less. (I probably wouldn't disallow setting weird limits in requestDevice, but we could.)

For example, on this random phone (first one I looked at), the max is [1024,1024,64] but total 512 (I guess you can use [1024,1,0], but [1024,1,1] is invalid??). If an app saw these limits it would be rather likely to try to use [1024,1,1] and fail unexpectedly.

Aside: now that I've been looking at it repeatedly, "workgroup invocations" confused me, it sounds like "number of workgroup invocations" not "number of invocations in a workgroup". Should we rename to, like, maxComputeWorkgroupVolume?

[Kangz]

+1 on using some computeWorkgroupVolume terminology and having some guarantees on the relations between some of the limits in the spec. However I'd rather not merge maxComputeWorkgroupSize.x and maxComputeWorkgroupSize.y together even if we're not able to find devices with different limits there: such devices might exist in the future, but also I think it is a net win for developer understanding to have the limit exposed as a "vector".

[kvark]

The observation that Kai makes is quite interesting. We shouldn't define limits "just in case" the future hardware comes and needs them. We should be ready to have Limits V2 at some point, anyway. Hopefully, not soon :)

[kvark]

Using the GPUExtent3D is a bit strange here. It's made for texture sizes specifically, hence depthOrArrayLayers field, but the compute workgroup grid has no direct relation to the texture sizes.

[Kangz]

It is readonly and described as [256, 256, 64] above. The intent is that it is always the array version. I agree that the parallels with texture size are weird because they should be non-existent, but we're really talking about a volume here so it makes sense to use Extent3D imho.

[kvark]

If the intent is to use the array, any reason to not put sequence<GPUIntegerCoordinate> in here?
I don't know how users are expected to deal with GPUExtent3D coming from us. It's one thing to pass in either a list or a structure. It's another thing to not know what to do with this type when it comes to you.

[Kangz]

The reason is because WebIDL doesn't have [GPUIntegerCoordinate; 3] to show the sequence is statically sized, while GPUExtent3D is clearly sized. But happy to change, this doesn't really matter.

[kvark]

Yeah, no strong feelings here, we can follow-up

[github-actions]

Previews, as seen when this build job started (2479ab4):
WebGPU | IDL
WGSL
Explainer

[kainino0x]

+1 on using some computeWorkgroupVolume terminology

Editors came up with "maxComputeInvocationsPerWorkgroup"

All together that would be:

• maxComputeWorkgroupStorageSize, maybe maxWorkgroupStorage
• maxComputeInvocationsPerWorkgroup, maybe maxInvocationsPerWorkgroup
• maxComputeWorkgroupSize, maybe maxComputeWorkgroupDimensions