Allow u32 coordinates for textureStore/textureLoad

src/back/glsl/mod.rs

@@ -3144,12 +3144,32 @@ impl<'a, W: Write> Writer<'a, W> {
             }
             // Otherwise write just the expression (and the 1D hack if needed)
             None => {
+                let uvec_size = match *ctx.info[coordinate].ty.inner_with(&self.module.types) {
+                    TypeInner::Scalar {
+                        kind: crate::ScalarKind::Uint,
+                        ..
+                    } => Some(None),
+                    TypeInner::Vector {
+                        size,
+                        kind: crate::ScalarKind::Uint,
+                        ..
+                    } => Some(Some(size as u32)),
+                    _ => None,
+                };
                 if tex_1d_hack {
                     write!(self.out, "ivec2(")?;
+                } else if uvec_size.is_some() {
+                    match uvec_size {
+                        Some(None) => write!(self.out, "int(")?,
+                        Some(Some(size)) => write!(self.out, "ivec{}(", size)?,
+                        _ => {}
+                    }
                 }
                 self.write_expr(coordinate, ctx)?;
                 if tex_1d_hack {
                     write!(self.out, ", 0)")?;
+                } else if uvec_size.is_some() {
+                    write!(self.out, ")")?;
                 }
             }
         }

src/back/spv/image.rs

@@ -317,6 +317,22 @@ impl<'w> BlockContext<'w> {
         let array_index_i32_id = self.cached[array_index];
         let reconciled_array_index_id = if component_kind == crate::ScalarKind::Sint {
             array_index_i32_id
+        } else if component_kind == crate::ScalarKind::Uint {
+            let u32_id = self.get_type_id(LookupType::Local(LocalType::Value {
+                vector_size: None,
+                kind: crate::ScalarKind::Uint,
+                width: 4,
+                pointer_space: None,
+            }));
+
+            let reconciled_id = self.gen_id();
+            block.body.push(Instruction::unary(
+                spirv::Op::Bitcast,
+                u32_id,
+                reconciled_id,
+                array_index_i32_id,
+            ));
+            reconciled_id
         } else {
             let component_type_id = self.get_type_id(LookupType::Local(LocalType::Value {
                 vector_size: None,

src/valid/mod.rs

@@ -220,17 +220,17 @@ impl crate::TypeInner {
     const fn image_storage_coordinates(&self) -> Option<crate::ImageDimension> {
         match *self {
             Self::Scalar {
-                kind: crate::ScalarKind::Sint,
+                kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
                 ..
             } => Some(crate::ImageDimension::D1),
             Self::Vector {
                 size: crate::VectorSize::Bi,
-                kind: crate::ScalarKind::Sint,
+                kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
                 ..
             } => Some(crate::ImageDimension::D2),
             Self::Vector {
                 size: crate::VectorSize::Tri,
-                kind: crate::ScalarKind::Sint,
+                kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
                 ..
             } => Some(crate::ImageDimension::D3),
             _ => None,

tests/in/image.wgsl

@@ -19,12 +19,22 @@ var image_dst: texture_storage_1d<r32uint,write>;
 fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
     let dim = textureDimensions(image_storage_src);
     let itc = dim * vec2<i32>(local_id.xy) % vec2<i32>(10, 20);
+    // loads with ivec2 coords.
     let value1 = textureLoad(image_mipmapped_src, itc, i32(local_id.z));
     let value2 = textureLoad(image_multisampled_src, itc, i32(local_id.z));
     let value4 = textureLoad(image_storage_src, itc);
     let value5 = textureLoad(image_array_src, itc, i32(local_id.z), i32(local_id.z) + 1);
     let value6 = 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 value4u = textureLoad(image_storage_src, vec2<u32>(itc));
+    let value5u = textureLoad(image_array_src, vec2<u32>(itc), i32(local_id.z), i32(local_id.z) + 1);
+    let value6u = textureLoad(image_1d_src, u32(local_id.x), i32(local_id.z));
+    // store with ivec2 coords.
     textureStore(image_dst, itc.x, value1 + value2 + value4 + value5 + value6);
+    // store with uvec2 coords.
+    textureStore(image_dst, u32(itc.x), value1u + value2u + value4u + value5u + value6u);
 }
 
 @compute @workgroup_size(16, 1, 1)

tests/out/glsl/image.main.Compute.glsl

@@ -28,7 +28,13 @@ void main() {
     uvec4 value4_ = imageLoad(_group_0_binding_1_cs, itc);
     uvec4 value5_ = texelFetch(_group_0_binding_5_cs, ivec3(itc, int(local_id.z)), (int(local_id.z) + 1));
     uvec4 value6_ = texelFetch(_group_0_binding_7_cs, ivec2(int(local_id.x), 0), 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 value4u = imageLoad(_group_0_binding_1_cs, ivec2(uvec2(itc)));
+    uvec4 value5u = texelFetch(_group_0_binding_5_cs, ivec3(uvec2(itc), int(local_id.z)), (int(local_id.z) + 1));
+    uvec4 value6u = texelFetch(_group_0_binding_7_cs, ivec2(uint(local_id.x), 0), int(local_id.z));
     imageStore(_group_0_binding_2_cs, ivec2(itc.x, 0), ((((value1_ + value2_) + value4_) + value5_) + value6_));
+    imageStore(_group_0_binding_2_cs, ivec2(uint(itc.x), 0), ((((value1u + value2u) + value4u) + value5u) + value6u));
     return;
 }
 

tests/out/hlsl/image.hlsl

@@ -38,7 +38,13 @@ void main(uint3 local_id : SV_GroupThreadID)
     uint4 value4_ = image_storage_src.Load(itc);
     uint4 value5_ = image_array_src.Load(int4(itc, int(local_id.z), (int(local_id.z) + 1)));
     uint4 value6_ = 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 value4u = image_storage_src.Load(uint2(itc));
+    uint4 value5u = image_array_src.Load(int4(uint2(itc), int(local_id.z), (int(local_id.z) + 1)));
+    uint4 value6u = image_1d_src.Load(int2(uint(local_id.x), int(local_id.z)));
     image_dst[itc.x] = ((((value1_ + value2_) + value4_) + value5_) + value6_);
+    image_dst[uint(itc.x)] = ((((value1u + value2u) + value4u) + value5u) + value6u);
     return;
 }
 

tests/out/msl/image.msl

@@ -24,7 +24,13 @@ kernel void main_(
     metal::uint4 value4_ = image_storage_src.read(metal::uint2(itc));
     metal::uint4 value5_ = image_array_src.read(metal::uint2(itc), static_cast<int>(local_id.z), static_cast<int>(local_id.z) + 1);
     metal::uint4 value6_ = 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 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)), static_cast<int>(local_id.z), static_cast<int>(local_id.z) + 1);
+    metal::uint4 value6u = image_1d_src.read(uint(static_cast<uint>(local_id.x)));
     image_dst.write((((value1_ + value2_) + value4_) + value5_) + value6_, uint(itc.x));
+    image_dst.write((((value1u + value2u) + value4u) + value5u) + value6u, uint(static_cast<uint>(itc.x)));
     return;
 }