[naga] Allow textureLoad's sample index arg to be unsigned.

Permit the `sample_index` argument of `textureLoad` overloads to be
either `i32` or `u32`.

Author: jimblandy

CHANGELOG.md

@@ -52,6 +52,7 @@ Naga now infers the correct binding layout when a resource appears only in an as
 
 - Add polyfills for `dot4U8Packed` and `dot4I8Packed` for all backends. By @robamler in [#7494](https://github.com/gfx-rs/wgpu/pull/7494).
 - Add polyfilled `pack4x{I,U}8Clamped` built-ins to all backends and WGSL frontend. By @ErichDonGubler in [#7546](https://github.com/gfx-rs/wgpu/pull/7546).
+- Allow textureLoad's sample index arg to be unsigned. By @jimblandy in [#7625](https://github.com/gfx-rs/wgpu/pull/7625).
 
 #### DX12
 

naga/src/valid/expression.rs

@@ -660,7 +660,7 @@ impl super::Validator {
                 match (sample, class.is_multisampled()) {
                     (None, false) => {}
                     (Some(sample), true) => {
-                        if resolver[sample].scalar_kind() != Some(Sk::Sint) {
+                        if !matches!(resolver[sample], Ti::Scalar(Sc::I32 | Sc::U32)) {
                             return Err(ExpressionError::InvalidImageOtherIndexType(sample));
                         }
                     }

naga/tests/in/wgsl/image.wgsl

@@ -24,13 +24,15 @@ fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
     // doing the same thing as the line above, but with u32, as textureLoad must also support unsigned integers.
     let value1_2 = textureLoad(image_mipmapped_src, itc, u32(local_id.z));
     let value2 = textureLoad(image_multisampled_src, itc, i32(local_id.z));
+    let value3 = textureLoad(image_multisampled_src, itc, u32(local_id.z));
     let value4 = textureLoad(image_storage_src, itc);
     let value5 = textureLoad(image_array_src, itc, local_id.z, i32(local_id.z) + 1);
     let value6 = textureLoad(image_array_src, itc, i32(local_id.z), i32(local_id.z) + 1);
     let value7 = textureLoad(image_1d_src, i32(local_id.x), i32(local_id.z));
     // loads with uvec2 coords.
     let value1u = textureLoad(image_mipmapped_src, vec2<u32>(itc), i32(local_id.z));
     let value2u = textureLoad(image_multisampled_src, vec2<u32>(itc), i32(local_id.z));
+    let value3u = textureLoad(image_multisampled_src, vec2<u32>(itc), u32(local_id.z));
     let value4u = textureLoad(image_storage_src, vec2<u32>(itc));
     let value5u = textureLoad(image_array_src, vec2<u32>(itc), local_id.z, i32(local_id.z) + 1);
     let value6u = textureLoad(image_array_src, vec2<u32>(itc), i32(local_id.z), i32(local_id.z) + 1);

naga/tests/out/glsl/wgsl-image.main.Compute.glsl

@@ -22,12 +22,14 @@ void main() {
     uvec4 value1_ = texelFetch(_group_0_binding_0_cs, itc, int(local_id.z));
     uvec4 value1_2_ = texelFetch(_group_0_binding_0_cs, itc, int(uint(local_id.z)));
     uvec4 value2_ = texelFetch(_group_0_binding_3_cs, itc, int(local_id.z));
+    uvec4 value3_ = texelFetch(_group_0_binding_3_cs, itc, int(uint(local_id.z)));
     uvec4 value4_ = imageLoad(_group_0_binding_1_cs, itc);
     uvec4 value5_ = texelFetch(_group_0_binding_5_cs, ivec3(itc, local_id.z), (int(local_id.z) + 1));
     uvec4 value6_ = texelFetch(_group_0_binding_5_cs, ivec3(itc, int(local_id.z)), (int(local_id.z) + 1));
     uvec4 value7_ = texelFetch(_group_0_binding_7_cs, int(local_id.x), int(local_id.z));
     uvec4 value1u = texelFetch(_group_0_binding_0_cs, ivec2(uvec2(itc)), int(local_id.z));
     uvec4 value2u = texelFetch(_group_0_binding_3_cs, ivec2(uvec2(itc)), int(local_id.z));
+    uvec4 value3u = texelFetch(_group_0_binding_3_cs, ivec2(uvec2(itc)), int(uint(local_id.z)));
     uvec4 value4u = imageLoad(_group_0_binding_1_cs, ivec2(uvec2(itc)));
     uvec4 value5u = texelFetch(_group_0_binding_5_cs, ivec3(uvec2(itc), local_id.z), (int(local_id.z) + 1));
     uvec4 value6u = texelFetch(_group_0_binding_5_cs, ivec3(uvec2(itc), int(local_id.z)), (int(local_id.z) + 1));

naga/tests/out/hlsl/wgsl-image.hlsl

@@ -44,12 +44,14 @@ void main(uint3 local_id : SV_GroupThreadID)
     uint4 value1_ = image_mipmapped_src.Load(int3(itc, int(local_id.z)));
     uint4 value1_2_ = image_mipmapped_src.Load(int3(itc, int(uint(local_id.z))));
     uint4 value2_ = image_multisampled_src.Load(itc, int(local_id.z));
+    uint4 value3_ = image_multisampled_src.Load(itc, uint(local_id.z));
     uint4 value4_ = image_storage_src.Load(itc);
     uint4 value5_ = image_array_src.Load(int4(itc, local_id.z, asint(asuint(int(local_id.z)) + asuint(int(1)))));
     uint4 value6_ = image_array_src.Load(int4(itc, int(local_id.z), asint(asuint(int(local_id.z)) + asuint(int(1)))));
     uint4 value7_ = image_1d_src.Load(int2(int(local_id.x), int(local_id.z)));
     uint4 value1u = image_mipmapped_src.Load(int3(uint2(itc), int(local_id.z)));
     uint4 value2u = image_multisampled_src.Load(uint2(itc), int(local_id.z));
+    uint4 value3u = image_multisampled_src.Load(uint2(itc), uint(local_id.z));
     uint4 value4u = image_storage_src.Load(uint2(itc));
     uint4 value5u = image_array_src.Load(int4(uint2(itc), local_id.z, asint(asuint(int(local_id.z)) + asuint(int(1)))));
     uint4 value6u = image_array_src.Load(int4(uint2(itc), int(local_id.z), asint(asuint(int(local_id.z)) + asuint(int(1)))));

naga/tests/out/msl/wgsl-image.msl

@@ -26,12 +26,14 @@ kernel void main_(
     metal::uint4 value1_ = image_mipmapped_src.read(metal::uint2(itc), static_cast<int>(local_id.z));
     metal::uint4 value1_2_ = image_mipmapped_src.read(metal::uint2(itc), static_cast<uint>(local_id.z));
     metal::uint4 value2_ = image_multisampled_src.read(metal::uint2(itc), static_cast<int>(local_id.z));
+    metal::uint4 value3_ = image_multisampled_src.read(metal::uint2(itc), static_cast<uint>(local_id.z));
     metal::uint4 value4_ = image_storage_src.read(metal::uint2(itc));
     metal::uint4 value5_ = image_array_src.read(metal::uint2(itc), local_id.z, as_type<int>(as_type<uint>(static_cast<int>(local_id.z)) + as_type<uint>(1)));
     metal::uint4 value6_ = image_array_src.read(metal::uint2(itc), static_cast<int>(local_id.z), as_type<int>(as_type<uint>(static_cast<int>(local_id.z)) + as_type<uint>(1)));
     metal::uint4 value7_ = image_1d_src.read(uint(static_cast<int>(local_id.x)));
     metal::uint4 value1u = image_mipmapped_src.read(metal::uint2(static_cast<metal::uint2>(itc)), static_cast<int>(local_id.z));
     metal::uint4 value2u = image_multisampled_src.read(metal::uint2(static_cast<metal::uint2>(itc)), static_cast<int>(local_id.z));
+    metal::uint4 value3u = image_multisampled_src.read(metal::uint2(static_cast<metal::uint2>(itc)), static_cast<uint>(local_id.z));
     metal::uint4 value4u = image_storage_src.read(metal::uint2(static_cast<metal::uint2>(itc)));
     metal::uint4 value5u = image_array_src.read(metal::uint2(static_cast<metal::uint2>(itc)), local_id.z, as_type<int>(as_type<uint>(static_cast<int>(local_id.z)) + as_type<uint>(1)));
     metal::uint4 value6u = image_array_src.read(metal::uint2(static_cast<metal::uint2>(itc)), static_cast<int>(local_id.z), as_type<int>(as_type<uint>(static_cast<int>(local_id.z)) + as_type<uint>(1)));