Fix textureGather on texture_2d<u32/i32> (gfx-rs#2138)

* Fix textureGather on texture_2d<u32/i32>

* Add textureGather u32/i32 tests

* Update src/valid/expression.rs

Co-authored-by: Teodor Tanasoaia <[email protected]>

* Fix formatting

* undo analyzer change

Co-authored-by: Teodor Tanasoaia <[email protected]>

Author: JMS55

src/valid/expression.rs

@@ -384,6 +384,10 @@ impl super::Validator {
                         kind: crate::ScalarKind::Float,
                         multi: false,
                     } => false,
+                    crate::ImageClass::Sampled {
+                        kind: crate::ScalarKind::Uint | crate::ScalarKind::Sint,
+                        multi: false,
+                    } if gather.is_some() => false,
                     crate::ImageClass::Depth { multi: false } => true,
                     _ => return Err(ExpressionError::InvalidImageClass(class)),
                 };

tests/in/image.wgsl

@@ -45,14 +45,18 @@ var image_1d: texture_1d<f32>;
 @group(0) @binding(1)
 var image_2d: texture_2d<f32>;
 @group(0) @binding(2)
-var image_2d_array: texture_2d_array<f32>;
+var image_2d_u32: texture_2d<u32>;
 @group(0) @binding(3)
-var image_cube: texture_cube<f32>;
+var image_2d_i32: texture_2d<i32>;
 @group(0) @binding(4)
-var image_cube_array: texture_cube_array<f32>;
+var image_2d_array: texture_2d_array<f32>;
 @group(0) @binding(5)
-var image_3d: texture_3d<f32>;
+var image_cube: texture_cube<f32>;
 @group(0) @binding(6)
+var image_cube_array: texture_cube_array<f32>;
+@group(0) @binding(7)
+var image_3d: texture_3d<f32>;
+@group(0) @binding(8)
 var image_aa: texture_multisampled_2d<f32>;
 
 @vertex
@@ -71,7 +75,7 @@ fn queries() -> @builtin(position) vec4<f32> {
     let dim_3d_lod = textureDimensions(image_3d, 1);
     let dim_2s_ms = textureDimensions(image_aa);
 
-    let sum = dim_1d + dim_2d.y + dim_2d_lod.y + dim_2d_array.y + dim_2d_array_lod.y +
+    let sum = dim_1d + dim_2d.y + dim_2d_lod.y + dim_2d_array.y + dim_2d_array_lod.y + 
         dim_cube.y + dim_cube_lod.y + dim_cube_array.y + dim_cube_array_lod.y +
         dim_3d.z + dim_3d_lod.z;
     return vec4<f32>(f32(sum));
@@ -134,7 +138,12 @@ fn gather() -> @location(0) vec4<f32> {
     let s2d_offset = textureGather(3, image_2d, sampler_reg, tc, vec2<i32>(3, 1));
     let s2d_depth = textureGatherCompare(image_2d_depth, sampler_cmp, tc, dref);
     let s2d_depth_offset = textureGatherCompare(image_2d_depth, sampler_cmp, tc, dref, vec2<i32>(3, 1));
-    return s2d + s2d_offset + s2d_depth + s2d_depth_offset;
+
+    let u = textureGather(0, image_2d_u32, sampler_reg, tc);
+    let i = textureGather(0, image_2d_i32, sampler_reg, tc);
+    let f = vec4<f32>(u) + vec4<f32>(i);
+
+    return s2d + s2d_offset + s2d_depth + s2d_depth_offset + f;
 }
 
 @fragment

tests/out/glsl/image.gather.Fragment.glsl

@@ -6,6 +6,10 @@ precision highp int;
 
 uniform highp sampler2D _group_0_binding_1_fs;
 
+uniform highp usampler2D _group_0_binding_2_fs;
+
+uniform highp isampler2D _group_0_binding_3_fs;
+
 uniform highp sampler2DShadow _group_1_binding_2_fs;
 
 layout(location = 0) out vec4 _fs2p_location0;
@@ -16,7 +20,10 @@ void main() {
     vec4 s2d_offset = textureGatherOffset(_group_0_binding_1_fs, vec2(tc), ivec2(3, 1), 3);
     vec4 s2d_depth = textureGather(_group_1_binding_2_fs, vec2(tc), 0.5);
     vec4 s2d_depth_offset = textureGatherOffset(_group_1_binding_2_fs, vec2(tc), 0.5, ivec2(3, 1));
-    _fs2p_location0 = (((s2d + s2d_offset) + s2d_depth) + s2d_depth_offset);
+    uvec4 u = textureGather(_group_0_binding_2_fs, vec2(tc), 0);
+    ivec4 i = textureGather(_group_0_binding_3_fs, vec2(tc), 0);
+    vec4 f = (vec4(u) + vec4(i));
+    _fs2p_location0 = ((((s2d + s2d_offset) + s2d_depth) + s2d_depth_offset) + f);
     return;
 }
 

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

@@ -8,31 +8,31 @@ uniform highp sampler2D _group_0_binding_0_vs;
 
 uniform highp sampler2D _group_0_binding_1_vs;
 
-uniform highp sampler2DArray _group_0_binding_2_vs;
+uniform highp sampler2DArray _group_0_binding_4_vs;
 
-uniform highp samplerCube _group_0_binding_3_vs;
+uniform highp samplerCube _group_0_binding_5_vs;
 
-uniform highp samplerCubeArray _group_0_binding_4_vs;
+uniform highp samplerCubeArray _group_0_binding_6_vs;
 
-uniform highp sampler3D _group_0_binding_5_vs;
+uniform highp sampler3D _group_0_binding_7_vs;
 
-uniform highp sampler2DMS _group_0_binding_6_vs;
+uniform highp sampler2DMS _group_0_binding_8_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;
-    ivec2 dim_2d_array_lod = textureSize(_group_0_binding_2_vs, 1).xy;
-    ivec2 dim_cube = textureSize(_group_0_binding_3_vs, 0).xy;
-    ivec2 dim_cube_lod = textureSize(_group_0_binding_3_vs, 1).xy;
-    ivec2 dim_cube_array = textureSize(_group_0_binding_4_vs, 0).xy;
-    ivec2 dim_cube_array_lod = textureSize(_group_0_binding_4_vs, 1).xy;
-    ivec3 dim_3d = textureSize(_group_0_binding_5_vs, 0).xyz;
-    ivec3 dim_3d_lod = textureSize(_group_0_binding_5_vs, 1).xyz;
-    ivec2 dim_2s_ms = textureSize(_group_0_binding_6_vs).xy;
+    ivec2 dim_2d_array = textureSize(_group_0_binding_4_vs, 0).xy;
+    ivec2 dim_2d_array_lod = textureSize(_group_0_binding_4_vs, 1).xy;
+    ivec2 dim_cube = textureSize(_group_0_binding_5_vs, 0).xy;
+    ivec2 dim_cube_lod = textureSize(_group_0_binding_5_vs, 1).xy;
+    ivec2 dim_cube_array = textureSize(_group_0_binding_6_vs, 0).xy;
+    ivec2 dim_cube_array_lod = textureSize(_group_0_binding_6_vs, 1).xy;
+    ivec3 dim_3d = textureSize(_group_0_binding_7_vs, 0).xyz;
+    ivec3 dim_3d_lod = textureSize(_group_0_binding_7_vs, 1).xyz;
+    ivec2 dim_2s_ms = textureSize(_group_0_binding_8_vs).xy;
     int sum = ((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z);
     gl_Position = vec4(float(sum));
     gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);

tests/out/hlsl/image.hlsl

@@ -9,11 +9,13 @@ 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);
-Texture2DArray<float4> image_2d_array : register(t2);
-TextureCube<float4> image_cube : register(t3);
-TextureCubeArray<float4> image_cube_array : register(t4);
-Texture3D<float4> image_3d : register(t5);
-Texture2DMS<float4> image_aa : register(t6);
+Texture2D<uint4> image_2d_u32_ : register(t2);
+Texture2D<int4> image_2d_i32_ : register(t3);
+Texture2DArray<float4> image_2d_array : register(t4);
+TextureCube<float4> image_cube : register(t5);
+TextureCubeArray<float4> image_cube_array : register(t6);
+Texture3D<float4> image_3d : register(t7);
+Texture2DMS<float4> image_aa : register(t8);
 SamplerState sampler_reg : register(s0, space1);
 SamplerComparisonState sampler_cmp : register(s1, space1);
 Texture2D<float> image_2d_depth : register(t2, space1);
@@ -258,7 +260,10 @@ float4 gather() : SV_Target0
     float4 s2d_offset_1 = image_2d.GatherAlpha(sampler_reg, tc_2, int2(3, 1));
     float4 s2d_depth_1 = image_2d_depth.GatherCmp(sampler_cmp, tc_2, 0.5);
     float4 s2d_depth_offset = image_2d_depth.GatherCmp(sampler_cmp, tc_2, 0.5, int2(3, 1));
-    return (((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset);
+    uint4 u = image_2d_u32_.Gather(sampler_reg, tc_2);
+    int4 i = image_2d_i32_.Gather(sampler_reg, tc_2);
+    float4 f = (float4(u) + float4(i));
+    return ((((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset) + f);
 }
 
 float4 depth_no_comparison() : SV_Target0

tests/out/msl/image.msl

@@ -139,6 +139,8 @@ struct gatherOutput {
 };
 fragment gatherOutput gather(
   metal::texture2d<float, metal::access::sample> image_2d [[user(fake0)]]
+, metal::texture2d<uint, metal::access::sample> image_2d_u32_ [[user(fake0)]]
+, metal::texture2d<int, metal::access::sample> image_2d_i32_ [[user(fake0)]]
 , metal::sampler sampler_reg [[user(fake0)]]
 , metal::sampler sampler_cmp [[user(fake0)]]
 , metal::depth2d<float, metal::access::sample> image_2d_depth [[user(fake0)]]
@@ -148,7 +150,10 @@ fragment gatherOutput gather(
     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_9_);
-    return gatherOutput { ((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset };
+    metal::uint4 u = image_2d_u32_.gather(sampler_reg, tc_2);
+    metal::int4 i = image_2d_i32_.gather(sampler_reg, tc_2);
+    metal::float4 f = static_cast<metal::float4>(u) + static_cast<metal::float4>(i);
+    return gatherOutput { (((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset) + f };
 }