[naga wgsl-in] Convert textureStore values correctly.

Apply necessary automatic conversions to the `value` argument of
`textureStore`.

naga/src/front/wgsl/error.rs

@@ -176,6 +176,7 @@ pub(crate) enum Error<'a> {
         from_type: String,
         to_type: String,
     },
+    NotStorageTexture(Span),
     BadTextureSampleType {
         span: Span,
         scalar: Scalar,
@@ -536,6 +537,11 @@ impl<'a> Error<'a> {
                 labels: vec![(bad_span, "unknown scalar type".into())],
                 notes: vec!["Valid scalar types are f32, f64, i32, u32, bool".into()],
             },
+            Error::NotStorageTexture(bad_span) => ParseError {
+                message: "textureStore can only be applied to storage textures".to_string(),
+                labels: vec![(bad_span, "not a storage texture".into())],
+                notes: vec![],
+            },
             Error::BadTextureSampleType { span, scalar } => ParseError {
                 message: format!(
                     "texture sample type must be one of f32, i32 or u32, but found {}",

naga/src/front/wgsl/lower/mod.rs

@@ -2681,15 +2681,21 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
 
                             let coordinate = self.expression(args.next()?, ctx)?;
 
-                            let (_, arrayed) = ctx.image_data(image, image_span)?;
+                            let (class, arrayed) = ctx.image_data(image, image_span)?;
                             let array_index = arrayed
                                 .then(|| {
                                     args.min_args += 1;
                                     self.expression(args.next()?, ctx)
                                 })
                                 .transpose()?;
+                            let scalar = if let ir::ImageClass::Storage { format, .. } = class {
+                                format.into()
+                            } else {
+                                return Err(Box::new(Error::NotStorageTexture(image_span)));
+                            };
 
-                            let value = self.expression(args.next()?, ctx)?;
+                            let value =
+                                self.expression_with_leaf_scalar(args.next()?, scalar, ctx)?;
 
                             args.finish()?;
 

naga/tests/in/wgsl/abstract-types-texture.toml

@@ -1 +1,4 @@
 targets = "METAL"
+
+[msl]
+lang_version = [1, 2]

naga/tests/in/wgsl/abstract-types-texture.wgsl

@@ -18,3 +18,9 @@ fn depth() {
   _ = textureSampleCompare(d, c, vec2(1,2), 0);
   _ = textureGatherCompare(d, c, vec2(1,2), 0);
 }
+
+@group(0) @binding(4) var st: texture_storage_2d<rgba8unorm, read_write>;
+
+fn storage() {
+  textureStore(st, vec2(0,1), vec4(2,3,4,5));
+}

naga/tests/out/msl/wgsl-abstract-types-texture.msl

@@ -1,4 +1,4 @@
-// language: metal1.0
+// language: metal1.2
 #include <metal_stdlib>
 #include <simd/simd.h>
 
@@ -28,3 +28,10 @@ void depth(
     metal::float4 phony_8 = d.gather_compare(c, metal::float2(1.0, 2.0), 0.0);
     return;
 }
+
+void storage(
+    metal::texture2d<float, metal::access::read_write> st
+) {
+    st.write(metal::float4(2.0, 3.0, 4.0, 5.0), metal::uint2(metal::int2(0, 1)));
+    return;
+}