[msl-out] Don't give level in texture1d accesses.

Fixes gfx-rs#1642.

Since 1d textures cannot have mipmaps, MSL requires that the `level` argument to
texel accesses and dimension queries be a constexpr 0. For our purposes, just
omit the level argument altogether.

Makefile

@@ -36,7 +36,7 @@ validate-msl: $(SNAPSHOTS_BASE_OUT)/msl/*.msl
 	@set -e && for file in $^ ; do \
 		echo "Validating" $${file#"$(SNAPSHOTS_BASE_OUT)/"};	\
 		header=$$(head -n1 $${file});	\
-		cat $${file} | xcrun -sdk macosx metal -mmacosx-version-min=10.11 -std=macos-$${header:13:8} -x metal - -o /dev/null; \
+		cat $${file} | ``; \
 	done
 
 validate-glsl: $(SNAPSHOTS_BASE_OUT)/glsl/*.glsl

src/back/msl/writer.rs

@@ -582,7 +582,10 @@ impl<W: Write> Writer<W> {
         match dim {
             crate::ImageDimension::D1 => {
                 write!(self.out, "int(")?;
-                self.put_image_query(image, "width", level, context)?;
+                // Since 1D textures never have mipmaps, MSL requires that the
+                // `level` argument be a constexpr 0. It's simplest for us just
+                // to omit the level entirely.
+                self.put_image_query(image, "width", None, context)?;
                 write!(self.out, ")")?;
             }
             crate::ImageDimension::D2 => {
@@ -996,8 +999,18 @@ impl<W: Write> Writer<W> {
                     self.put_expression(expr, context, true)?;
                 }
                 if let Some(index) = index {
-                    write!(self.out, ", ")?;
-                    self.put_expression(index, context, true)?;
+                    // Metal requires that the `level` argument to
+                    // `texture1d::read` be a constexpr equal to zero.
+                    if let crate::TypeInner::Image {
+                        dim: crate::ImageDimension::D1,
+                        ..
+                    } = *context.resolve_type(image)
+                    {
+                        // The argument defaults to zero.
+                    } else {
+                        write!(self.out, ", ")?;
+                        self.put_expression(index, context, true)?
+                    }
                 }
                 write!(self.out, ")")?;
             }

tests/in/image.wgsl

@@ -10,6 +10,8 @@ var image_storage_src: texture_storage_2d<rgba8uint, read>;
 var image_array_src: texture_2d_array<u32>;
 [[group(0), binding(6)]]
 var image_dup_src: texture_storage_1d<r32uint,read>; // for #1307
+[[group(0), binding(7)]]
+var image_1d_src: texture_1d<u32>;
 [[group(0), binding(2)]]
 var image_dst: texture_storage_1d<r32uint,write>;
 
@@ -25,7 +27,8 @@ fn main(
     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);
-    textureStore(image_dst, itc.x, value1 + value2 + value4 + value5);
+    let value6 = textureLoad(image_1d_src, i32(local_id.x), i32(local_id.z));
+    textureStore(image_dst, itc.x, value1 + value2 + value4 + value5 + value6);
 }
 
 [[stage(compute), workgroup_size(16, 1, 1)]]
@@ -55,6 +58,7 @@ var image_aa: texture_multisampled_2d<f32>;
 [[stage(vertex)]]
 fn queries() -> [[builtin(position)]] vec4<f32> {
     let dim_1d = textureDimensions(image_1d);
+    let dim_1d_lod = textureDimensions(image_1d, i32(dim_1d));
     let dim_2d = textureDimensions(image_2d);
     let dim_2d_lod = textureDimensions(image_2d, 1);
     let dim_2d_array = textureDimensions(image_2d_array);

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

@@ -14,6 +14,8 @@ layout(rgba8ui) readonly uniform highp uimage2D _group_0_binding_1_cs;
 
 uniform highp usampler2DArray _group_0_binding_5_cs;
 
+uniform highp usampler2D _group_0_binding_7_cs;
+
 layout(r32ui) writeonly uniform highp uimage2D _group_0_binding_2_cs;
 
 
@@ -25,7 +27,8 @@ void main() {
     uvec4 value2_ = texelFetch(_group_0_binding_3_cs, itc, int(local_id.z));
     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));
-    imageStore(_group_0_binding_2_cs, ivec2(itc.x, 0.0), (((value1_ + value2_) + value4_) + value5_));
+    uvec4 value6_ = texelFetch(_group_0_binding_7_cs, ivec2(int(local_id.x), 0.0), int(local_id.z));
+    imageStore(_group_0_binding_2_cs, ivec2(itc.x, 0.0), ((((value1_ + value2_) + value4_) + value5_) + value6_));
     return;
 }
 

tests/out/glsl/image.queries.Vertex.glsl

@@ -19,6 +19,7 @@ uniform highp sampler3D _group_0_binding_5_vs;
 
 void main() {
     int dim_1d = textureSize(_group_0_binding_0_vs, 0).x;
+    int dim_1d_lod = textureSize(_group_0_binding_0_vs, int(dim_1d)).x;
     ivec2 dim_2d = textureSize(_group_0_binding_1_vs, 0).xy;
     ivec2 dim_2d_lod = textureSize(_group_0_binding_1_vs, 1).xy;
     ivec2 dim_2d_array = textureSize(_group_0_binding_2_vs, 0).xy;

tests/out/hlsl/image.hlsl

@@ -5,6 +5,7 @@ Texture2DMS<float> image_depth_multisampled_src : register(t4);
 RWTexture2D<uint4> image_storage_src : register(u1);
 Texture2DArray<uint4> image_array_src : register(t5);
 RWTexture1D<uint4> image_dup_src : register(u6);
+Texture1D<uint4> image_1d_src : register(t7);
 RWTexture1D<uint4> image_dst : register(u2);
 Texture1D<float4> image_1d : register(t0);
 Texture2D<float4> image_2d : register(t1);
@@ -33,7 +34,8 @@ void main(uint3 local_id : SV_GroupThreadID)
     uint4 value2_ = image_multisampled_src.Load(itc, int(local_id.z));
     uint4 value4_ = image_storage_src.Load(itc);
     uint4 value5_ = image_array_src.Load(int4(itc, int(local_id.z), (int(local_id.z) + 1)));
-    image_dst[itc.x] = (((value1_ + value2_) + value4_) + value5_);
+    uint4 value6_ = image_1d_src.Load(int2(int(local_id.x), int(local_id.z)));
+    image_dst[itc.x] = ((((value1_ + value2_) + value4_) + value5_) + value6_);
     return;
 }
 
@@ -54,6 +56,13 @@ int NagaDimensions1D(Texture1D<float4> tex)
     return ret.x;
 }
 
+int NagaMipDimensions1D(Texture1D<float4> tex, uint mip_level)
+{
+    uint4 ret;
+    tex.GetDimensions(mip_level, ret.x, ret.y);
+    return ret.x;
+}
+
 int2 NagaDimensions2D(Texture2D<float4> tex)
 {
     uint4 ret;
@@ -127,6 +136,7 @@ int3 NagaMipDimensions3D(Texture3D<float4> tex, uint mip_level)
 float4 queries() : SV_Position
 {
     int dim_1d = NagaDimensions1D(image_1d);
+    int dim_1d_lod = NagaMipDimensions1D(image_1d, int(dim_1d));
     int2 dim_2d = NagaDimensions2D(image_2d);
     int2 dim_2d_lod = NagaMipDimensions2D(image_2d, 1);
     int2 dim_2d_array = NagaDimensions2DArray(image_2d_array);

tests/out/msl/image.msl

@@ -2,7 +2,7 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
-constant metal::int2 const_type_8_ = {3, 1};
+constant metal::int2 const_type_9_ = {3, 1};
 
 struct main_Input {
 };
@@ -12,6 +12,7 @@ kernel void main_(
 , metal::texture2d_ms<uint, metal::access::read> image_multisampled_src [[user(fake0)]]
 , metal::texture2d<uint, metal::access::read> image_storage_src [[user(fake0)]]
 , metal::texture2d_array<uint, metal::access::sample> image_array_src [[user(fake0)]]
+, metal::texture1d<uint, metal::access::sample> image_1d_src [[user(fake0)]]
 , metal::texture1d<uint, metal::access::write> image_dst [[user(fake0)]]
 ) {
     metal::int2 dim = int2(image_storage_src.get_width(), image_storage_src.get_height());
@@ -20,7 +21,8 @@ kernel void main_(
     metal::uint4 value2_ = image_multisampled_src.read(metal::uint2(itc), static_cast<int>(local_id.z));
     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);
-    image_dst.write(((value1_ + value2_) + value4_) + value5_, metal::uint(itc.x));
+    metal::uint4 value6_ = image_1d_src.read(metal::uint(static_cast<int>(local_id.x)));
+    image_dst.write((((value1_ + value2_) + value4_) + value5_) + value6_, metal::uint(itc.x));
     return;
 }
 
@@ -53,6 +55,7 @@ vertex queriesOutput queries(
 , metal::texture3d<float, metal::access::sample> image_3d [[user(fake0)]]
 ) {
     int dim_1d = int(image_1d.get_width());
+    int dim_1d_lod = int(image_1d.get_width());
     metal::int2 dim_2d = int2(image_2d.get_width(), image_2d.get_height());
     metal::int2 dim_2d_lod = int2(image_2d.get_width(1), image_2d.get_height(1));
     metal::int2 dim_2d_array = int2(image_2d_array.get_width(), image_2d_array.get_height());
@@ -103,9 +106,9 @@ fragment sampleOutput sample(
     metal::float2 tc = metal::float2(0.5);
     metal::float4 s1d = image_1d.sample(sampler_reg, tc.x);
     metal::float4 s2d = image_2d.sample(sampler_reg, tc);
-    metal::float4 s2d_offset = image_2d.sample(sampler_reg, tc, const_type_8_);
+    metal::float4 s2d_offset = image_2d.sample(sampler_reg, tc, const_type_9_);
     metal::float4 s2d_level = image_2d.sample(sampler_reg, tc, metal::level(2.299999952316284));
-    metal::float4 s2d_level_offset = image_2d.sample(sampler_reg, tc, metal::level(2.299999952316284), const_type_8_);
+    metal::float4 s2d_level_offset = image_2d.sample(sampler_reg, tc, metal::level(2.299999952316284), const_type_9_);
     return sampleOutput { (((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset };
 }
 
@@ -135,9 +138,9 @@ fragment gatherOutput gather(
 ) {
     metal::float2 tc_2 = metal::float2(0.5);
     metal::float4 s2d_1 = image_2d.gather(sampler_reg, tc_2, int2(0), metal::component::y);
-    metal::float4 s2d_offset_1 = image_2d.gather(sampler_reg, tc_2, const_type_8_, metal::component::w);
+    metal::float4 s2d_offset_1 = image_2d.gather(sampler_reg, tc_2, const_type_9_, metal::component::w);
     metal::float4 s2d_depth_1 = image_2d_depth.gather_compare(sampler_cmp, tc_2, 0.5);
-    metal::float4 s2d_depth_offset = image_2d_depth.gather_compare(sampler_cmp, tc_2, 0.5, const_type_8_);
+    metal::float4 s2d_depth_offset = image_2d_depth.gather_compare(sampler_cmp, tc_2, 0.5, const_type_9_);
     return gatherOutput { ((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset };
 }