diff --git a/src/back/glsl/features.rs b/src/back/glsl/features.rs index d2e83e0040..74db74c70f 100644 --- a/src/back/glsl/features.rs +++ b/src/back/glsl/features.rs @@ -22,8 +22,6 @@ bitflags::bitflags! { /// Adds support for image load and early depth tests const IMAGE_LOAD_STORE = 1 << 8; const CONSERVATIVE_DEPTH = 1 << 9; - /// Isn't supported in ES - const TEXTURE_1D = 1 << 10; /// Interpolation and auxiliary qualifiers. Perspective, Flat, and /// Centroid are available in all GLSL versions we support. const NOPERSPECTIVE_QUALIFIER = 1 << 11; @@ -92,9 +90,6 @@ impl FeaturesManager { check_feature!(IMAGE_LOAD_STORE, 130, 310); check_feature!(CONSERVATIVE_DEPTH, 130, 300); check_feature!(CONSERVATIVE_DEPTH, 130, 300); - // 1D textures are supported by all core versions and aren't supported by an es versions - // so use 0 that way the check will always be false and can be optimized away - check_feature!(TEXTURE_1D, 0); check_feature!(NOPERSPECTIVE_QUALIFIER, 130); check_feature!(SAMPLE_QUALIFIER, 400, 320); // gl_ClipDistance is supported by core versions > 1.3 and aren't supported by an es versions without extensions @@ -298,8 +293,6 @@ impl<'a, W> Writer<'a, W> { } => { if arrayed && dim == ImageDimension::Cube { self.features.request(Features::CUBE_TEXTURES_ARRAY) - } else if dim == ImageDimension::D1 { - self.features.request(Features::TEXTURE_1D) } match class { diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 60e0d57b0d..f2406fa9cd 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -514,7 +514,7 @@ impl<'a, W: Write> Writer<'a, W> { // We treat images separately because they might require // writing the storage format TypeInner::Image { - dim, + mut dim, arrayed, class, } => { @@ -526,6 +526,11 @@ impl<'a, W: Write> Writer<'a, W> { } => Some((format, access)), _ => None, }; + + if dim == crate::ImageDimension::D1 && es { + dim = crate::ImageDimension::D2 + } + // Gether the location if needed let layout_binding = if self.options.version.supports_explicit_locations() { let br = global.binding.as_ref().unwrap(); @@ -1954,12 +1959,12 @@ impl<'a, W: Write> Writer<'a, W> { // We need to get the coordinates vector size to later build a vector that's `size + 1` // if `depth_ref` is some, if it isn't a vector we panic as that's not a valid expression - let size = match *ctx.info[coordinate].ty.inner_with(&self.module.types) { - TypeInner::Vector { size, .. } => size, + let mut coord_dim = match *ctx.info[coordinate].ty.inner_with(&self.module.types) { + TypeInner::Vector { size, .. } => size as u8, + TypeInner::Scalar { .. } => 1, _ => unreachable!(), }; - let mut coord_dim = size as u8; if array_index.is_some() { coord_dim += 1; } @@ -1968,9 +1973,16 @@ impl<'a, W: Write> Writer<'a, W> { coord_dim += 1; } + let tex_1d_hack = dim == crate::ImageDimension::D1 && self.options.version.is_es(); + let is_vec = tex_1d_hack || coord_dim != 1; // Compose a new texture coordinates vector - write!(self.out, "vec{}(", coord_dim)?; + if is_vec { + write!(self.out, "vec{}(", coord_dim + tex_1d_hack as u8)?; + } self.write_expr(coordinate, ctx)?; + if tex_1d_hack { + write!(self.out, ", 0.0")?; + } if let Some(expr) = array_index { write!(self.out, ", ")?; self.write_expr(expr, ctx)?; @@ -1981,7 +1993,9 @@ impl<'a, W: Write> Writer<'a, W> { self.write_expr(expr, ctx)?; } } - write!(self.out, ")")?; + if is_vec { + write!(self.out, ")")?; + } if cube_array_shadow { if let Some(expr) = depth_ref { @@ -2025,7 +2039,13 @@ impl<'a, W: Write> Writer<'a, W> { if let Some(constant) = offset { write!(self.out, ", ")?; + if tex_1d_hack { + write!(self.out, "ivec2(")?; + } self.write_constant(constant)?; + if tex_1d_hack { + write!(self.out, ", 0)")?; + } } // End the function @@ -2103,11 +2123,11 @@ impl<'a, W: Write> Writer<'a, W> { ImageClass::Sampled { .. } | ImageClass::Depth { .. } => { write!(self.out, "textureSize(")?; self.write_expr(image, ctx)?; - write!(self.out, ",")?; + write!(self.out, ", ")?; if let Some(expr) = level { self.write_expr(expr, ctx)?; } else { - write!(self.out, "0",)?; + write!(self.out, "0")?; } } ImageClass::Storage { .. } => { @@ -2115,7 +2135,10 @@ impl<'a, W: Write> Writer<'a, W> { self.write_expr(image, ctx)?; } } - write!(self.out, ").{}", &"xyz"[..components])?; + write!(self.out, ")")?; + if components != 1 || self.options.version.is_es() { + write!(self.out, ".{}", &"xyz"[..components])?; + } } crate::ImageQuery::NumLevels => { write!(self.out, "textureQueryLevels(",)?; @@ -2129,7 +2152,9 @@ impl<'a, W: Write> Writer<'a, W> { }; write!(self.out, "{}(", fun_name)?; self.write_expr(image, ctx)?; - write!(self.out, ",0).{}", back::COMPONENTS[components])?; + if components != 1 || self.options.version.is_es() { + write!(self.out, ", 0).{}", back::COMPONENTS[components])?; + } } crate::ImageQuery::NumSamples => { // assumes ARB_shader_texture_image_samples @@ -2473,7 +2498,14 @@ impl<'a, W: Write> Writer<'a, W> { write!(self.out, ")")?; } None => { + let tex_1d_hack = dim == IDim::D1 && self.options.version.is_es(); + if tex_1d_hack { + write!(self.out, "ivec2(")?; + } self.write_expr(coordinate, ctx)?; + if tex_1d_hack { + write!(self.out, ", 0.0)")?; + } } } Ok(()) diff --git a/tests/in/image.param.ron b/tests/in/image.param.ron index 91e6e85972..721b98fb2d 100644 --- a/tests/in/image.param.ron +++ b/tests/in/image.param.ron @@ -3,4 +3,5 @@ version: (1, 1), debug: true, ), + glsl_blacklist: ["depth_load", "levels_queries"] ) diff --git a/tests/in/image.wgsl b/tests/in/image.wgsl index c4621a407f..e9fd4f0fa9 100644 --- a/tests/in/image.wgsl +++ b/tests/in/image.wgsl @@ -23,10 +23,18 @@ fn main( let itc = dim * vec2(local_id.xy) % vec2(10, 20); let value1 = textureLoad(image_mipmapped_src, itc, i32(local_id.z)); let value2 = textureLoad(image_multisampled_src, itc, i32(local_id.z)); - let value3 = textureLoad(image_depth_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 + u32(value3) + value4 + value5); + textureStore(image_dst, itc.x, value1 + value2 + value4 + value5); +} + +[[stage(compute), workgroup_size(16, 1, 1)]] +fn depth_load([[builtin(local_invocation_id)]] local_id: vec3) { + let dim: vec2 = textureDimensions(image_storage_src); + let itc: vec2 = ((dim * vec2(local_id.xy)) % vec2(10, 20)); + let val: f32 = textureLoad(image_depth_multisampled_src, itc, i32(local_id.z)); + textureStore(image_dst, itc.x, vec4(u32(val))); + return; } [[group(0), binding(0)]] @@ -48,27 +56,34 @@ var image_aa: texture_multisampled_2d; fn queries() -> [[builtin(position)]] vec4 { let dim_1d = textureDimensions(image_1d); let dim_2d = textureDimensions(image_2d); - let num_levels_2d = textureNumLevels(image_2d); let dim_2d_lod = textureDimensions(image_2d, 1); let dim_2d_array = textureDimensions(image_2d_array); - let num_levels_2d_array = textureNumLevels(image_2d_array); let dim_2d_array_lod = textureDimensions(image_2d_array, 1); - let num_layers_2d = textureNumLayers(image_2d_array); let dim_cube = textureDimensions(image_cube); - let num_levels_cube = textureNumLevels(image_cube); let dim_cube_lod = textureDimensions(image_cube, 1); let dim_cube_array = textureDimensions(image_cube_array); - let num_levels_cube_array = textureNumLevels(image_cube_array); let dim_cube_array_lod = textureDimensions(image_cube_array, 1); - let num_layers_cube = textureNumLayers(image_cube_array); let dim_3d = textureDimensions(image_3d); - let num_levels_3d = textureNumLevels(image_3d); let dim_3d_lod = textureDimensions(image_3d, 1); - let num_samples_aa = textureNumSamples(image_aa); let sum = dim_1d + dim_2d.y + dim_2d_lod.y + dim_2d_array.y + dim_2d_array_lod.y + - num_layers_2d + dim_cube.y + dim_cube_lod.y + dim_cube_array.y + dim_cube_array_lod.y + - num_layers_cube + dim_3d.z + dim_3d_lod.z + num_samples_aa + + 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(sum)); +} + +[[stage(vertex)]] +fn levels_queries() -> [[builtin(position)]] vec4 { + let num_levels_2d = textureNumLevels(image_2d); + let num_levels_2d_array = textureNumLevels(image_2d_array); + let num_layers_2d = textureNumLayers(image_2d_array); + let num_levels_cube = textureNumLevels(image_cube); + let num_levels_cube_array = textureNumLevels(image_cube_array); + let num_layers_cube = textureNumLayers(image_cube_array); + let num_levels_3d = textureNumLevels(image_3d); + let num_samples_aa = textureNumSamples(image_aa); + + let sum = num_layers_2d + num_layers_cube + num_samples_aa + num_levels_2d + num_levels_2d_array + num_levels_3d + num_levels_cube + num_levels_cube_array; return vec4(f32(sum)); } @@ -80,11 +95,12 @@ var sampler_reg: sampler; fn sample() -> [[location(0)]] vec4 { let tc = vec2(0.5); let level = 2.3; + let s1d = textureSample(image_1d, sampler_reg, tc.x); let s2d = textureSample(image_2d, sampler_reg, tc); let s2d_offset = textureSample(image_2d, sampler_reg, tc, vec2(3, 1)); let s2d_level = textureSampleLevel(image_2d, sampler_reg, tc, level); let s2d_level_offset = textureSampleLevel(image_2d, sampler_reg, tc, level, vec2(3, 1)); - return s2d + s2d_offset + s2d_level + s2d_level_offset; + return s1d + s2d + s2d_offset + s2d_level + s2d_level_offset; } [[group(1), binding(1)]] diff --git a/tests/out/glsl/image.main.Compute.glsl b/tests/out/glsl/image.main.Compute.glsl new file mode 100644 index 0000000000..be6b7ed977 --- /dev/null +++ b/tests/out/glsl/image.main.Compute.glsl @@ -0,0 +1,31 @@ +#version 310 es +#extension GL_EXT_texture_cube_map_array : require + +precision highp float; +precision highp int; + +layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in; + +uniform highp usampler2D _group_0_binding_0; + +uniform highp usampler2DMS _group_0_binding_3; + +layout(rgba8ui) readonly uniform highp uimage2D _group_0_binding_1; + +uniform highp usampler2DArray _group_0_binding_5; + +layout(r32ui) writeonly uniform highp uimage2D _group_0_binding_2; + + +void main() { + uvec3 local_id = gl_LocalInvocationID; + ivec2 dim = imageSize(_group_0_binding_1).xy; + ivec2 itc = ((dim * ivec2(local_id.xy)) % ivec2(10, 20)); + uvec4 value1_ = texelFetch(_group_0_binding_0, itc, int(local_id.z)); + uvec4 value2_ = texelFetch(_group_0_binding_3, itc, int(local_id.z)); + uvec4 value4_ = imageLoad(_group_0_binding_1, itc); + uvec4 value5_ = texelFetch(_group_0_binding_5, ivec3(itc, int(local_id.z)), (int(local_id.z) + 1)); + imageStore(_group_0_binding_2, ivec2(itc.x, 0.0), (((value1_ + value2_) + value4_) + value5_)); + return; +} + diff --git a/tests/out/glsl/image.queries.Vertex.glsl b/tests/out/glsl/image.queries.Vertex.glsl new file mode 100644 index 0000000000..8f89d2dada --- /dev/null +++ b/tests/out/glsl/image.queries.Vertex.glsl @@ -0,0 +1,37 @@ +#version 310 es +#extension GL_EXT_texture_cube_map_array : require + +precision highp float; +precision highp int; + +uniform highp sampler2D _group_0_binding_0; + +uniform highp sampler2D _group_0_binding_1; + +uniform highp sampler2DArray _group_0_binding_2; + +uniform highp samplerCube _group_0_binding_3; + +uniform highp samplerCubeArray _group_0_binding_4; + +uniform highp sampler3D _group_0_binding_5; + + +void main() { + int dim_1d = textureSize(_group_0_binding_0, 0).x; + ivec2 dim_2d = textureSize(_group_0_binding_1, 0).xy; + ivec2 dim_2d_lod = textureSize(_group_0_binding_1, 1).xy; + ivec2 dim_2d_array = textureSize(_group_0_binding_2, 0).xy; + ivec2 dim_2d_array_lod = textureSize(_group_0_binding_2, 1).xy; + ivec2 dim_cube = textureSize(_group_0_binding_3, 0).xy; + ivec2 dim_cube_lod = textureSize(_group_0_binding_3, 1).xy; + ivec2 dim_cube_array = textureSize(_group_0_binding_4, 0).xy; + ivec2 dim_cube_array_lod = textureSize(_group_0_binding_4, 1).xy; + ivec3 dim_3d = textureSize(_group_0_binding_5, 0).xyz; + ivec3 dim_3d_lod = textureSize(_group_0_binding_5, 1).xyz; + 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); + return; +} + diff --git a/tests/out/glsl/image.sample.Fragment.glsl b/tests/out/glsl/image.sample.Fragment.glsl new file mode 100644 index 0000000000..a0d7946d0a --- /dev/null +++ b/tests/out/glsl/image.sample.Fragment.glsl @@ -0,0 +1,23 @@ +#version 310 es +#extension GL_EXT_texture_cube_map_array : require + +precision highp float; +precision highp int; + +uniform highp sampler2D _group_0_binding_0; + +uniform highp sampler2D _group_0_binding_1; + +layout(location = 0) out vec4 _fs2p_location0; + +void main() { + vec2 tc = vec2(0.5); + vec4 s1d = texture(_group_0_binding_0, vec2(tc.x, 0.0)); + vec4 s2d = texture(_group_0_binding_1, vec2(tc)); + vec4 s2d_offset = textureOffset(_group_0_binding_1, vec2(tc), ivec2(3, 1)); + vec4 s2d_level = textureLod(_group_0_binding_1, vec2(tc), 2.299999952316284); + vec4 s2d_level_offset = textureLodOffset(_group_0_binding_1, vec2(tc), 2.299999952316284, ivec2(3, 1)); + _fs2p_location0 = ((((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset); + return; +} + diff --git a/tests/out/glsl/image.sample_comparison.Fragment.glsl b/tests/out/glsl/image.sample_comparison.Fragment.glsl new file mode 100644 index 0000000000..4b52baac2c --- /dev/null +++ b/tests/out/glsl/image.sample_comparison.Fragment.glsl @@ -0,0 +1,18 @@ +#version 310 es +#extension GL_EXT_texture_cube_map_array : require + +precision highp float; +precision highp int; + +uniform highp sampler2DShadow _group_1_binding_2; + +layout(location = 0) out float _fs2p_location0; + +void main() { + vec2 tc = vec2(0.5); + float s2d_depth = texture(_group_1_binding_2, vec3(tc, 0.5)); + float s2d_depth_level = textureLod(_group_1_binding_2, vec3(tc, 0.5), 0.0); + _fs2p_location0 = (s2d_depth + s2d_depth_level); + return; +} + diff --git a/tests/out/hlsl/image.hlsl b/tests/out/hlsl/image.hlsl index 4eb297de5e..d6da396a5b 100644 --- a/tests/out/hlsl/image.hlsl +++ b/tests/out/hlsl/image.hlsl @@ -31,10 +31,19 @@ void main(uint3 local_id : SV_GroupThreadID) int2 itc = ((dim * int2(local_id.xy)) % int2(10, 20)); uint4 value1_ = image_mipmapped_src.Load(int3(itc, int(local_id.z))); uint4 value2_ = image_multisampled_src.Load(itc, int(local_id.z)); - float value3_ = image_depth_multisampled_src.Load(itc, int(local_id.z)).x; 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_) + uint4(uint(value3_).xxxx)) + value4_) + value5_); + image_dst[itc.x] = (((value1_ + value2_) + value4_) + value5_); + return; +} + +[numthreads(16, 1, 1)] +void depth_load(uint3 local_id1 : SV_GroupThreadID) +{ + int2 dim1 = NagaRWDimensions2D(image_storage_src); + int2 itc1 = ((dim1 * int2(local_id1.xy)) % int2(10, 20)); + float val = image_depth_multisampled_src.Load(itc1, int(local_id1.z)).x; + image_dst[itc1.x] = uint4(uint(val).xxxx); return; } @@ -52,13 +61,6 @@ int2 NagaDimensions2D(Texture2D tex) return ret.xy; } -int NagaNumLevels2D(Texture2D tex) -{ - uint4 ret; - tex.GetDimensions(0, ret.x, ret.y, ret.z); - return ret.z; -} - int2 NagaMipDimensions2D(Texture2D tex, uint mip_level) { uint4 ret; @@ -73,13 +75,6 @@ int2 NagaDimensions2DArray(Texture2DArray tex) return ret.xy; } -int NagaNumLevels2DArray(Texture2DArray tex) -{ - uint4 ret; - tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); - return ret.w; -} - int2 NagaMipDimensions2DArray(Texture2DArray tex, uint mip_level) { uint4 ret; @@ -87,13 +82,6 @@ int2 NagaMipDimensions2DArray(Texture2DArray tex, uint mip_level) return ret.xy; } -int NagaNumLayers2DArray(Texture2DArray tex) -{ - uint4 ret; - tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); - return ret.w; -} - int2 NagaDimensionsCube(TextureCube tex) { uint4 ret; @@ -101,13 +89,6 @@ int2 NagaDimensionsCube(TextureCube tex) return ret.xy; } -int NagaNumLevelsCube(TextureCube tex) -{ - uint4 ret; - tex.GetDimensions(0, ret.x, ret.y, ret.z); - return ret.z; -} - int2 NagaMipDimensionsCube(TextureCube tex, uint mip_level) { uint4 ret; @@ -122,13 +103,6 @@ int2 NagaDimensionsCubeArray(TextureCubeArray tex) return ret.xy; } -int NagaNumLevelsCubeArray(TextureCubeArray tex) -{ - uint4 ret; - tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); - return ret.w; -} - int2 NagaMipDimensionsCubeArray(TextureCubeArray tex, uint mip_level) { uint4 ret; @@ -136,13 +110,6 @@ int2 NagaMipDimensionsCubeArray(TextureCubeArray tex, uint mip_level) return ret.xy; } -int NagaNumLayersCubeArray(TextureCubeArray tex) -{ - uint4 ret; - tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); - return ret.w; -} - int3 NagaDimensions3D(Texture3D tex) { uint4 ret; @@ -150,18 +117,77 @@ int3 NagaDimensions3D(Texture3D tex) return ret.xyz; } -int NagaNumLevels3D(Texture3D tex) +int3 NagaMipDimensions3D(Texture3D tex, uint mip_level) +{ + uint4 ret; + tex.GetDimensions(mip_level, ret.x, ret.y, ret.z, ret.w); + return ret.xyz; +} + +float4 queries() : SV_Position +{ + int dim_1d = NagaDimensions1D(image_1d); + int2 dim_2d = NagaDimensions2D(image_2d); + int2 dim_2d_lod = NagaMipDimensions2D(image_2d, 1); + int2 dim_2d_array = NagaDimensions2DArray(image_2d_array); + int2 dim_2d_array_lod = NagaMipDimensions2DArray(image_2d_array, 1); + int2 dim_cube = NagaDimensionsCube(image_cube); + int2 dim_cube_lod = NagaMipDimensionsCube(image_cube, 1); + int2 dim_cube_array = NagaDimensionsCubeArray(image_cube_array); + int2 dim_cube_array_lod = NagaMipDimensionsCubeArray(image_cube_array, 1); + int3 dim_3d = NagaDimensions3D(image_3d); + int3 dim_3d_lod = NagaMipDimensions3D(image_3d, 1); + 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); + return float4(float(sum).xxxx); +} + +int NagaNumLevels2D(Texture2D tex) +{ + uint4 ret; + tex.GetDimensions(0, ret.x, ret.y, ret.z); + return ret.z; +} + +int NagaNumLevels2DArray(Texture2DArray tex) { uint4 ret; tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); return ret.w; } -int3 NagaMipDimensions3D(Texture3D tex, uint mip_level) +int NagaNumLayers2DArray(Texture2DArray tex) { uint4 ret; - tex.GetDimensions(mip_level, ret.x, ret.y, ret.z, ret.w); - return ret.xyz; + tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); + return ret.w; +} + +int NagaNumLevelsCube(TextureCube tex) +{ + uint4 ret; + tex.GetDimensions(0, ret.x, ret.y, ret.z); + return ret.z; +} + +int NagaNumLevelsCubeArray(TextureCubeArray tex) +{ + uint4 ret; + tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); + return ret.w; +} + +int NagaNumLayersCubeArray(TextureCubeArray tex) +{ + uint4 ret; + tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); + return ret.w; +} + +int NagaNumLevels3D(Texture3D tex) +{ + uint4 ret; + tex.GetDimensions(0, ret.x, ret.y, ret.z, ret.w); + return ret.w; } int NagaMSNumSamples2D(Texture2DMS tex) @@ -171,39 +197,29 @@ int NagaMSNumSamples2D(Texture2DMS tex) return ret.z; } -float4 queries() : SV_Position +float4 levels_queries() : SV_Position { - int dim_1d = NagaDimensions1D(image_1d); - int2 dim_2d = NagaDimensions2D(image_2d); int num_levels_2d = NagaNumLevels2D(image_2d); - int2 dim_2d_lod = NagaMipDimensions2D(image_2d, 1); - int2 dim_2d_array = NagaDimensions2DArray(image_2d_array); int num_levels_2d_array = NagaNumLevels2DArray(image_2d_array); - int2 dim_2d_array_lod = NagaMipDimensions2DArray(image_2d_array, 1); int num_layers_2d = NagaNumLayers2DArray(image_2d_array); - int2 dim_cube = NagaDimensionsCube(image_cube); int num_levels_cube = NagaNumLevelsCube(image_cube); - int2 dim_cube_lod = NagaMipDimensionsCube(image_cube, 1); - int2 dim_cube_array = NagaDimensionsCubeArray(image_cube_array); int num_levels_cube_array = NagaNumLevelsCubeArray(image_cube_array); - int2 dim_cube_array_lod = NagaMipDimensionsCubeArray(image_cube_array, 1); int num_layers_cube = NagaNumLayersCubeArray(image_cube_array); - int3 dim_3d = NagaDimensions3D(image_3d); int num_levels_3d = NagaNumLevels3D(image_3d); - int3 dim_3d_lod = NagaMipDimensions3D(image_3d, 1); int num_samples_aa = NagaMSNumSamples2D(image_aa); - int sum = ((((((((((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + num_layers_2d) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + num_layers_cube) + dim_3d.z) + dim_3d_lod.z) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array); - return float4(float(sum).xxxx); + int sum1 = (((((((num_layers_2d + num_layers_cube) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array); + return float4(float(sum1).xxxx); } float4 sample1() : SV_Target0 { float2 tc = float2(0.5.xx); + float4 s1d = image_1d.Sample(sampler_reg, tc.x); float4 s2d = image_2d.Sample(sampler_reg, tc); float4 s2d_offset = image_2d.Sample(sampler_reg, tc, int2(3, 1)); float4 s2d_level = image_2d.SampleLevel(sampler_reg, tc, 2.299999952316284); float4 s2d_level_offset = image_2d.SampleLevel(sampler_reg, tc, 2.299999952316284, int2(3, 1)); - return (((s2d + s2d_offset) + s2d_level) + s2d_level_offset); + return ((((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset); } float sample_comparison() : SV_Target0 diff --git a/tests/out/hlsl/image.hlsl.config b/tests/out/hlsl/image.hlsl.config index c5d0760d93..4f6c8ecdbc 100644 --- a/tests/out/hlsl/image.hlsl.config +++ b/tests/out/hlsl/image.hlsl.config @@ -1,3 +1,3 @@ -vertex=(queries:vs_5_1 ) +vertex=(queries:vs_5_1 levels_queries:vs_5_1 ) fragment=(sample1:ps_5_1 sample_comparison:ps_5_1 ) -compute=(main:cs_5_1 ) +compute=(main:cs_5_1 depth_load:cs_5_1 ) diff --git a/tests/out/msl/image.msl b/tests/out/msl/image.msl index 44fde4992e..1c8e5e4ea8 100644 --- a/tests/out/msl/image.msl +++ b/tests/out/msl/image.msl @@ -10,7 +10,6 @@ kernel void main1( metal::uint3 local_id [[thread_position_in_threadgroup]] , metal::texture2d image_mipmapped_src [[user(fake0)]] , metal::texture2d_ms image_multisampled_src [[user(fake0)]] -, metal::depth2d_ms image_depth_multisampled_src [[user(fake0)]] , metal::texture2d image_storage_src [[user(fake0)]] , metal::texture2d_array image_array_src [[user(fake0)]] , metal::texture1d image_dst [[user(fake0)]] @@ -19,16 +18,31 @@ kernel void main1( metal::int2 itc = (dim * static_cast(local_id.xy)) % metal::int2(10, 20); metal::uint4 value1_ = image_mipmapped_src.read(metal::uint2(itc), static_cast(local_id.z)); metal::uint4 value2_ = image_multisampled_src.read(metal::uint2(itc), static_cast(local_id.z)); - float value3_ = image_depth_multisampled_src.read(metal::uint2(itc), static_cast(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(local_id.z), static_cast(local_id.z) + 1); - image_dst.write((((value1_ + value2_) + metal::uint4(static_cast(value3_))) + value4_) + value5_, metal::uint(itc.x)); + image_dst.write(((value1_ + value2_) + value4_) + value5_, metal::uint(itc.x)); + return; +} + + +struct depth_loadInput { +}; +kernel void depth_load( + metal::uint3 local_id1 [[thread_position_in_threadgroup]] +, metal::depth2d_ms image_depth_multisampled_src [[user(fake0)]] +, metal::texture2d image_storage_src [[user(fake0)]] +, metal::texture1d image_dst [[user(fake0)]] +) { + metal::int2 dim1 = int2(image_storage_src.get_width(), image_storage_src.get_height()); + metal::int2 itc1 = (dim1 * static_cast(local_id1.xy)) % metal::int2(10, 20); + float val = image_depth_multisampled_src.read(metal::uint2(itc1), static_cast(local_id1.z)); + image_dst.write(metal::uint4(static_cast(val)), metal::uint(itc1.x)); return; } struct queriesOutput { - metal::float4 member1 [[position]]; + metal::float4 member2 [[position]]; }; vertex queriesOutput queries( metal::texture1d image_1d [[user(fake0)]] @@ -37,50 +51,67 @@ vertex queriesOutput queries( , metal::texturecube image_cube [[user(fake0)]] , metal::texturecube_array image_cube_array [[user(fake0)]] , metal::texture3d image_3d [[user(fake0)]] -, metal::texture2d_ms image_aa [[user(fake0)]] ) { int dim_1d = int(image_1d.get_width()); metal::int2 dim_2d = int2(image_2d.get_width(), image_2d.get_height()); - int num_levels_2d = int(image_2d.get_num_mip_levels()); 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()); - int num_levels_2d_array = int(image_2d_array.get_num_mip_levels()); metal::int2 dim_2d_array_lod = int2(image_2d_array.get_width(1), image_2d_array.get_height(1)); - int num_layers_2d = int(image_2d_array.get_array_size()); metal::int2 dim_cube = int2(image_cube.get_width()); - int num_levels_cube = int(image_cube.get_num_mip_levels()); metal::int2 dim_cube_lod = int2(image_cube.get_width(1)); metal::int2 dim_cube_array = int2(image_cube_array.get_width()); - int num_levels_cube_array = int(image_cube_array.get_num_mip_levels()); metal::int2 dim_cube_array_lod = int2(image_cube_array.get_width(1)); - int num_layers_cube = int(image_cube_array.get_array_size()); metal::int3 dim_3d = int3(image_3d.get_width(), image_3d.get_height(), image_3d.get_depth()); - int num_levels_3d = int(image_3d.get_num_mip_levels()); metal::int3 dim_3d_lod = int3(image_3d.get_width(1), image_3d.get_height(1), image_3d.get_depth(1)); - int num_samples_aa = int(image_aa.get_num_samples()); - int sum = (((((((((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + num_layers_2d) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + num_layers_cube) + dim_3d.z) + dim_3d_lod.z) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array; + 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; return queriesOutput { metal::float4(static_cast(sum)) }; } +struct levels_queriesOutput { + metal::float4 member3 [[position]]; +}; +vertex levels_queriesOutput levels_queries( + metal::texture2d image_2d [[user(fake0)]] +, metal::texture2d_array image_2d_array [[user(fake0)]] +, metal::texturecube image_cube [[user(fake0)]] +, metal::texturecube_array image_cube_array [[user(fake0)]] +, metal::texture3d image_3d [[user(fake0)]] +, metal::texture2d_ms image_aa [[user(fake0)]] +) { + int num_levels_2d = int(image_2d.get_num_mip_levels()); + int num_levels_2d_array = int(image_2d_array.get_num_mip_levels()); + int num_layers_2d = int(image_2d_array.get_array_size()); + int num_levels_cube = int(image_cube.get_num_mip_levels()); + int num_levels_cube_array = int(image_cube_array.get_num_mip_levels()); + int num_layers_cube = int(image_cube_array.get_array_size()); + int num_levels_3d = int(image_3d.get_num_mip_levels()); + int num_samples_aa = int(image_aa.get_num_samples()); + int sum1 = ((((((num_layers_2d + num_layers_cube) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array; + return levels_queriesOutput { metal::float4(static_cast(sum1)) }; +} + + struct sampleOutput { - metal::float4 member2 [[color(0)]]; + metal::float4 member4 [[color(0)]]; }; fragment sampleOutput sample( - metal::texture2d image_2d [[user(fake0)]] + metal::texture1d image_1d [[user(fake0)]] +, metal::texture2d image_2d [[user(fake0)]] , metal::sampler sampler_reg [[user(fake0)]] ) { 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_type8_); 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_type8_); - return sampleOutput { ((s2d + s2d_offset) + s2d_level) + s2d_level_offset }; + return sampleOutput { (((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset }; } struct sample_comparisonOutput { - float member3 [[color(0)]]; + float member5 [[color(0)]]; }; fragment sample_comparisonOutput sample_comparison( metal::sampler sampler_cmp [[user(fake0)]] diff --git a/tests/out/spv/image.spvasm b/tests/out/spv/image.spvasm index e0617cfe46..9fc7d518e7 100644 --- a/tests/out/spv/image.spvasm +++ b/tests/out/spv/image.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 215 +; Bound: 244 OpCapability SampledCubeArray OpCapability ImageQuery OpCapability Image1D @@ -10,12 +10,15 @@ OpCapability Sampled1D %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %69 "main" %66 -OpEntryPoint Vertex %115 "queries" %113 -OpEntryPoint Fragment %182 "sample" %181 -OpEntryPoint Fragment %203 "sample_comparison" %201 +OpEntryPoint GLCompute %107 "depth_load" %105 +OpEntryPoint Vertex %128 "queries" %126 +OpEntryPoint Vertex %176 "levels_queries" %175 +OpEntryPoint Fragment %205 "sample" %204 +OpEntryPoint Fragment %232 "sample_comparison" %230 OpExecutionMode %69 LocalSize 16 1 1 -OpExecutionMode %182 OriginUpperLeft -OpExecutionMode %203 OriginUpperLeft +OpExecutionMode %107 LocalSize 16 1 1 +OpExecutionMode %205 OriginUpperLeft +OpExecutionMode %232 OriginUpperLeft OpSource GLSL 450 OpName %31 "image_mipmapped_src" OpName %33 "image_multisampled_src" @@ -36,9 +39,12 @@ OpName %61 "sampler_cmp" OpName %63 "image_2d_depth" OpName %66 "local_id" OpName %69 "main" -OpName %115 "queries" -OpName %182 "sample" -OpName %203 "sample_comparison" +OpName %105 "local_id" +OpName %107 "depth_load" +OpName %128 "queries" +OpName %176 "levels_queries" +OpName %205 "sample" +OpName %232 "sample_comparison" OpDecorate %31 DescriptorSet 0 OpDecorate %31 Binding 0 OpDecorate %33 DescriptorSet 0 @@ -77,9 +83,11 @@ OpDecorate %61 Binding 1 OpDecorate %63 DescriptorSet 1 OpDecorate %63 Binding 2 OpDecorate %66 BuiltIn LocalInvocationId -OpDecorate %113 BuiltIn Position -OpDecorate %181 Location 0 -OpDecorate %201 Location 0 +OpDecorate %105 BuiltIn LocalInvocationId +OpDecorate %126 BuiltIn Position +OpDecorate %175 BuiltIn Position +OpDecorate %204 Location 0 +OpDecorate %230 Location 0 %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 10 @@ -146,166 +154,203 @@ OpDecorate %201 Location 0 %67 = OpTypePointer Input %18 %66 = OpVariable %67 Input %70 = OpTypeFunction %2 -%79 = OpTypeVector %12 2 -%87 = OpTypeVector %12 4 -%102 = OpTypeVector %4 3 -%114 = OpTypePointer Output %27 -%113 = OpVariable %114 Output -%124 = OpConstant %12 0 -%181 = OpVariable %114 Output -%186 = OpTypeVector %8 2 -%188 = OpTypeSampledImage %21 -%202 = OpTypePointer Output %8 -%201 = OpVariable %202 Output -%208 = OpTypeSampledImage %29 -%213 = OpConstant %8 0.0 +%78 = OpTypeVector %12 2 +%86 = OpTypeVector %12 4 +%97 = OpTypeVector %4 3 +%105 = OpVariable %67 Input +%127 = OpTypePointer Output %27 +%126 = OpVariable %127 Output +%136 = OpConstant %12 0 +%175 = OpVariable %127 Output +%204 = OpVariable %127 Output +%210 = OpTypeVector %8 2 +%213 = OpTypeSampledImage %20 +%216 = OpTypeSampledImage %21 +%231 = OpTypePointer Output %8 +%230 = OpVariable %231 Output +%237 = OpTypeSampledImage %29 +%242 = OpConstant %8 0.0 %69 = OpFunction %2 None %70 %65 = OpLabel %68 = OpLoad %18 %66 %71 = OpLoad %11 %31 %72 = OpLoad %13 %33 -%73 = OpLoad %14 %35 -%74 = OpLoad %15 %37 -%75 = OpLoad %16 %39 -%76 = OpLoad %17 %43 -OpBranch %77 -%77 = OpLabel -%78 = OpImageQuerySize %19 %74 -%80 = OpVectorShuffle %79 %68 %68 0 1 -%81 = OpBitcast %19 %80 -%82 = OpIMul %19 %78 %81 -%83 = OpCompositeConstruct %19 %3 %5 -%84 = OpSMod %19 %82 %83 -%85 = OpCompositeExtract %12 %68 2 -%86 = OpBitcast %4 %85 -%88 = OpImageFetch %87 %71 %84 Lod %86 -%89 = OpCompositeExtract %12 %68 2 -%90 = OpBitcast %4 %89 -%91 = OpImageFetch %87 %72 %84 Sample %90 +%73 = OpLoad %15 %37 +%74 = OpLoad %16 %39 +%75 = OpLoad %17 %43 +OpBranch %76 +%76 = OpLabel +%77 = OpImageQuerySize %19 %73 +%79 = OpVectorShuffle %78 %68 %68 0 1 +%80 = OpBitcast %19 %79 +%81 = OpIMul %19 %77 %80 +%82 = OpCompositeConstruct %19 %3 %5 +%83 = OpSMod %19 %81 %82 +%84 = OpCompositeExtract %12 %68 2 +%85 = OpBitcast %4 %84 +%87 = OpImageFetch %86 %71 %83 Lod %85 +%88 = OpCompositeExtract %12 %68 2 +%89 = OpBitcast %4 %88 +%90 = OpImageFetch %86 %72 %83 Sample %89 +%91 = OpImageRead %86 %73 %83 %92 = OpCompositeExtract %12 %68 2 %93 = OpBitcast %4 %92 -%94 = OpImageFetch %27 %73 %84 Sample %93 -%95 = OpCompositeExtract %8 %94 0 -%96 = OpImageRead %87 %74 %84 -%97 = OpCompositeExtract %12 %68 2 -%98 = OpBitcast %4 %97 -%99 = OpCompositeExtract %12 %68 2 -%100 = OpBitcast %4 %99 -%101 = OpIAdd %4 %100 %6 -%103 = OpCompositeConstruct %102 %84 %98 -%104 = OpImageFetch %87 %75 %103 Lod %101 -%105 = OpCompositeExtract %4 %84 0 -%106 = OpIAdd %87 %88 %91 -%107 = OpConvertFToU %12 %95 -%108 = OpCompositeConstruct %87 %107 %107 %107 %107 -%109 = OpIAdd %87 %106 %108 -%110 = OpIAdd %87 %109 %96 -%111 = OpIAdd %87 %110 %104 -OpImageWrite %76 %105 %111 +%94 = OpCompositeExtract %12 %68 2 +%95 = OpBitcast %4 %94 +%96 = OpIAdd %4 %95 %6 +%98 = OpCompositeConstruct %97 %83 %93 +%99 = OpImageFetch %86 %74 %98 Lod %96 +%100 = OpCompositeExtract %4 %83 0 +%101 = OpIAdd %86 %87 %90 +%102 = OpIAdd %86 %101 %91 +%103 = OpIAdd %86 %102 %99 +OpImageWrite %75 %100 %103 OpReturn OpFunctionEnd -%115 = OpFunction %2 None %70 -%112 = OpLabel -%116 = OpLoad %20 %45 -%117 = OpLoad %21 %47 -%118 = OpLoad %22 %49 -%119 = OpLoad %23 %51 -%120 = OpLoad %24 %53 -%121 = OpLoad %25 %55 -%122 = OpLoad %26 %57 -OpBranch %123 -%123 = OpLabel -%125 = OpImageQuerySizeLod %4 %116 %124 -%126 = OpImageQuerySizeLod %19 %117 %124 -%127 = OpImageQueryLevels %4 %117 -%128 = OpImageQuerySizeLod %19 %117 %6 -%129 = OpImageQuerySizeLod %102 %118 %124 -%130 = OpVectorShuffle %19 %129 %129 0 1 -%131 = OpImageQueryLevels %4 %118 -%132 = OpImageQuerySizeLod %102 %118 %6 -%133 = OpVectorShuffle %19 %132 %132 0 1 -%134 = OpImageQuerySizeLod %102 %118 %124 -%135 = OpCompositeExtract %4 %134 2 -%136 = OpImageQuerySizeLod %19 %119 %124 -%137 = OpImageQueryLevels %4 %119 -%138 = OpImageQuerySizeLod %19 %119 %6 -%139 = OpImageQuerySizeLod %102 %120 %124 -%140 = OpVectorShuffle %19 %139 %139 0 0 -%141 = OpImageQueryLevels %4 %120 -%142 = OpImageQuerySizeLod %102 %120 %6 -%143 = OpVectorShuffle %19 %142 %142 0 0 -%144 = OpImageQuerySizeLod %102 %120 %124 -%145 = OpCompositeExtract %4 %144 2 -%146 = OpImageQuerySizeLod %102 %121 %124 -%147 = OpImageQueryLevels %4 %121 -%148 = OpImageQuerySizeLod %102 %121 %6 -%149 = OpImageQuerySamples %4 %122 -%150 = OpCompositeExtract %4 %126 1 -%151 = OpIAdd %4 %125 %150 -%152 = OpCompositeExtract %4 %128 1 -%153 = OpIAdd %4 %151 %152 -%154 = OpCompositeExtract %4 %130 1 +%107 = OpFunction %2 None %70 +%104 = OpLabel +%106 = OpLoad %18 %105 +%108 = OpLoad %14 %35 +%109 = OpLoad %15 %37 +%110 = OpLoad %17 %43 +OpBranch %111 +%111 = OpLabel +%112 = OpImageQuerySize %19 %109 +%113 = OpVectorShuffle %78 %106 %106 0 1 +%114 = OpBitcast %19 %113 +%115 = OpIMul %19 %112 %114 +%116 = OpCompositeConstruct %19 %3 %5 +%117 = OpSMod %19 %115 %116 +%118 = OpCompositeExtract %12 %106 2 +%119 = OpBitcast %4 %118 +%120 = OpImageFetch %27 %108 %117 Sample %119 +%121 = OpCompositeExtract %8 %120 0 +%122 = OpCompositeExtract %4 %117 0 +%123 = OpConvertFToU %12 %121 +%124 = OpCompositeConstruct %86 %123 %123 %123 %123 +OpImageWrite %110 %122 %124 +OpReturn +OpFunctionEnd +%128 = OpFunction %2 None %70 +%125 = OpLabel +%129 = OpLoad %20 %45 +%130 = OpLoad %21 %47 +%131 = OpLoad %22 %49 +%132 = OpLoad %23 %51 +%133 = OpLoad %24 %53 +%134 = OpLoad %25 %55 +OpBranch %135 +%135 = OpLabel +%137 = OpImageQuerySizeLod %4 %129 %136 +%138 = OpImageQuerySizeLod %19 %130 %136 +%139 = OpImageQuerySizeLod %19 %130 %6 +%140 = OpImageQuerySizeLod %97 %131 %136 +%141 = OpVectorShuffle %19 %140 %140 0 1 +%142 = OpImageQuerySizeLod %97 %131 %6 +%143 = OpVectorShuffle %19 %142 %142 0 1 +%144 = OpImageQuerySizeLod %19 %132 %136 +%145 = OpImageQuerySizeLod %19 %132 %6 +%146 = OpImageQuerySizeLod %97 %133 %136 +%147 = OpVectorShuffle %19 %146 %146 0 0 +%148 = OpImageQuerySizeLod %97 %133 %6 +%149 = OpVectorShuffle %19 %148 %148 0 0 +%150 = OpImageQuerySizeLod %97 %134 %136 +%151 = OpImageQuerySizeLod %97 %134 %6 +%152 = OpCompositeExtract %4 %138 1 +%153 = OpIAdd %4 %137 %152 +%154 = OpCompositeExtract %4 %139 1 %155 = OpIAdd %4 %153 %154 -%156 = OpCompositeExtract %4 %133 1 +%156 = OpCompositeExtract %4 %141 1 %157 = OpIAdd %4 %155 %156 -%158 = OpIAdd %4 %157 %135 -%159 = OpCompositeExtract %4 %136 1 -%160 = OpIAdd %4 %158 %159 -%161 = OpCompositeExtract %4 %138 1 -%162 = OpIAdd %4 %160 %161 -%163 = OpCompositeExtract %4 %140 1 -%164 = OpIAdd %4 %162 %163 -%165 = OpCompositeExtract %4 %143 1 -%166 = OpIAdd %4 %164 %165 -%167 = OpIAdd %4 %166 %145 -%168 = OpCompositeExtract %4 %146 2 +%158 = OpCompositeExtract %4 %143 1 +%159 = OpIAdd %4 %157 %158 +%160 = OpCompositeExtract %4 %144 1 +%161 = OpIAdd %4 %159 %160 +%162 = OpCompositeExtract %4 %145 1 +%163 = OpIAdd %4 %161 %162 +%164 = OpCompositeExtract %4 %147 1 +%165 = OpIAdd %4 %163 %164 +%166 = OpCompositeExtract %4 %149 1 +%167 = OpIAdd %4 %165 %166 +%168 = OpCompositeExtract %4 %150 2 %169 = OpIAdd %4 %167 %168 -%170 = OpCompositeExtract %4 %148 2 +%170 = OpCompositeExtract %4 %151 2 %171 = OpIAdd %4 %169 %170 -%172 = OpIAdd %4 %171 %149 -%173 = OpIAdd %4 %172 %127 -%174 = OpIAdd %4 %173 %131 -%175 = OpIAdd %4 %174 %147 -%176 = OpIAdd %4 %175 %137 -%177 = OpIAdd %4 %176 %141 -%178 = OpConvertSToF %8 %177 -%179 = OpCompositeConstruct %27 %178 %178 %178 %178 -OpStore %113 %179 +%172 = OpConvertSToF %8 %171 +%173 = OpCompositeConstruct %27 %172 %172 %172 %172 +OpStore %126 %173 OpReturn OpFunctionEnd -%182 = OpFunction %2 None %70 -%180 = OpLabel -%183 = OpLoad %21 %47 -%184 = OpLoad %28 %59 -OpBranch %185 -%185 = OpLabel -%187 = OpCompositeConstruct %186 %7 %7 -%189 = OpSampledImage %188 %183 %184 -%190 = OpImageSampleImplicitLod %27 %189 %187 -%191 = OpSampledImage %188 %183 %184 -%192 = OpImageSampleImplicitLod %27 %191 %187 ConstOffset %30 -%193 = OpSampledImage %188 %183 %184 -%194 = OpImageSampleExplicitLod %27 %193 %187 Lod %9 -%195 = OpSampledImage %188 %183 %184 -%196 = OpImageSampleExplicitLod %27 %195 %187 Lod|ConstOffset %9 %30 -%197 = OpFAdd %27 %190 %192 -%198 = OpFAdd %27 %197 %194 -%199 = OpFAdd %27 %198 %196 -OpStore %181 %199 +%176 = OpFunction %2 None %70 +%174 = OpLabel +%177 = OpLoad %21 %47 +%178 = OpLoad %22 %49 +%179 = OpLoad %23 %51 +%180 = OpLoad %24 %53 +%181 = OpLoad %25 %55 +%182 = OpLoad %26 %57 +OpBranch %183 +%183 = OpLabel +%184 = OpImageQueryLevels %4 %177 +%185 = OpImageQueryLevels %4 %178 +%186 = OpImageQuerySizeLod %97 %178 %136 +%187 = OpCompositeExtract %4 %186 2 +%188 = OpImageQueryLevels %4 %179 +%189 = OpImageQueryLevels %4 %180 +%190 = OpImageQuerySizeLod %97 %180 %136 +%191 = OpCompositeExtract %4 %190 2 +%192 = OpImageQueryLevels %4 %181 +%193 = OpImageQuerySamples %4 %182 +%194 = OpIAdd %4 %187 %191 +%195 = OpIAdd %4 %194 %193 +%196 = OpIAdd %4 %195 %184 +%197 = OpIAdd %4 %196 %185 +%198 = OpIAdd %4 %197 %192 +%199 = OpIAdd %4 %198 %188 +%200 = OpIAdd %4 %199 %189 +%201 = OpConvertSToF %8 %200 +%202 = OpCompositeConstruct %27 %201 %201 %201 %201 +OpStore %175 %202 OpReturn OpFunctionEnd -%203 = OpFunction %2 None %70 -%200 = OpLabel -%204 = OpLoad %28 %61 -%205 = OpLoad %29 %63 -OpBranch %206 -%206 = OpLabel -%207 = OpCompositeConstruct %186 %7 %7 -%209 = OpSampledImage %208 %205 %204 -%210 = OpImageSampleDrefImplicitLod %8 %209 %207 %7 -%211 = OpSampledImage %208 %205 %204 -%212 = OpImageSampleDrefExplicitLod %8 %211 %207 %7 Lod %213 -%214 = OpFAdd %8 %210 %212 -OpStore %201 %214 +%205 = OpFunction %2 None %70 +%203 = OpLabel +%206 = OpLoad %20 %45 +%207 = OpLoad %21 %47 +%208 = OpLoad %28 %59 +OpBranch %209 +%209 = OpLabel +%211 = OpCompositeConstruct %210 %7 %7 +%212 = OpCompositeExtract %8 %211 0 +%214 = OpSampledImage %213 %206 %208 +%215 = OpImageSampleImplicitLod %27 %214 %212 +%217 = OpSampledImage %216 %207 %208 +%218 = OpImageSampleImplicitLod %27 %217 %211 +%219 = OpSampledImage %216 %207 %208 +%220 = OpImageSampleImplicitLod %27 %219 %211 ConstOffset %30 +%221 = OpSampledImage %216 %207 %208 +%222 = OpImageSampleExplicitLod %27 %221 %211 Lod %9 +%223 = OpSampledImage %216 %207 %208 +%224 = OpImageSampleExplicitLod %27 %223 %211 Lod|ConstOffset %9 %30 +%225 = OpFAdd %27 %215 %218 +%226 = OpFAdd %27 %225 %220 +%227 = OpFAdd %27 %226 %222 +%228 = OpFAdd %27 %227 %224 +OpStore %204 %228 +OpReturn +OpFunctionEnd +%232 = OpFunction %2 None %70 +%229 = OpLabel +%233 = OpLoad %28 %61 +%234 = OpLoad %29 %63 +OpBranch %235 +%235 = OpLabel +%236 = OpCompositeConstruct %210 %7 %7 +%238 = OpSampledImage %237 %234 %233 +%239 = OpImageSampleDrefImplicitLod %8 %238 %236 %7 +%240 = OpSampledImage %237 %234 %233 +%241 = OpImageSampleDrefExplicitLod %8 %240 %236 %7 Lod %242 +%243 = OpFAdd %8 %239 %241 +OpStore %230 %243 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/image.wgsl b/tests/out/wgsl/image.wgsl index 166faf6d6e..34d642a2f5 100644 --- a/tests/out/wgsl/image.wgsl +++ b/tests/out/wgsl/image.wgsl @@ -39,10 +39,18 @@ fn main([[builtin(local_invocation_id)]] local_id: vec3) { let itc: vec2 = ((dim * vec2(local_id.xy)) % vec2(10, 20)); let value1_: vec4 = textureLoad(image_mipmapped_src, itc, i32(local_id.z)); let value2_: vec4 = textureLoad(image_multisampled_src, itc, i32(local_id.z)); - let value3_: f32 = textureLoad(image_depth_multisampled_src, itc, i32(local_id.z)); let value4_: vec4 = textureLoad(image_storage_src, itc); let value5_: vec4 = textureLoad(image_array_src, itc, i32(local_id.z), (i32(local_id.z) + 1)); - textureStore(image_dst, itc.x, ((((value1_ + value2_) + vec4(u32(value3_))) + value4_) + value5_)); + textureStore(image_dst, itc.x, (((value1_ + value2_) + value4_) + value5_)); + return; +} + +[[stage(compute), workgroup_size(16, 1, 1)]] +fn depth_load([[builtin(local_invocation_id)]] local_id1: vec3) { + let dim1: vec2 = textureDimensions(image_storage_src); + let itc1: vec2 = ((dim1 * vec2(local_id1.xy)) % vec2(10, 20)); + let val: f32 = textureLoad(image_depth_multisampled_src, itc1, i32(local_id1.z)); + textureStore(image_dst, itc1.x, vec4(u32(val))); return; } @@ -50,35 +58,42 @@ fn main([[builtin(local_invocation_id)]] local_id: vec3) { fn queries() -> [[builtin(position)]] vec4 { let dim_1d: i32 = textureDimensions(image_1d); let dim_2d: vec2 = textureDimensions(image_2d); - let num_levels_2d: i32 = textureNumLevels(image_2d); let dim_2d_lod: vec2 = textureDimensions(image_2d, 1); let dim_2d_array: vec2 = textureDimensions(image_2d_array); - let num_levels_2d_array: i32 = textureNumLevels(image_2d_array); let dim_2d_array_lod: vec2 = textureDimensions(image_2d_array, 1); - let num_layers_2d: i32 = textureNumLayers(image_2d_array); let dim_cube: vec2 = textureDimensions(image_cube); - let num_levels_cube: i32 = textureNumLevels(image_cube); let dim_cube_lod: vec2 = textureDimensions(image_cube, 1); let dim_cube_array: vec2 = textureDimensions(image_cube_array); - let num_levels_cube_array: i32 = textureNumLevels(image_cube_array); let dim_cube_array_lod: vec2 = textureDimensions(image_cube_array, 1); - let num_layers_cube: i32 = textureNumLayers(image_cube_array); let dim_3d: vec3 = textureDimensions(image_3d); - let num_levels_3d: i32 = textureNumLevels(image_3d); let dim_3d_lod: vec3 = textureDimensions(image_3d, 1); - let num_samples_aa: i32 = textureNumSamples(image_aa); - let sum: i32 = ((((((((((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + num_layers_2d) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + num_layers_cube) + dim_3d.z) + dim_3d_lod.z) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array); + let sum: i32 = ((((((((((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(sum)); } +[[stage(vertex)]] +fn levels_queries() -> [[builtin(position)]] vec4 { + let num_levels_2d: i32 = textureNumLevels(image_2d); + let num_levels_2d_array: i32 = textureNumLevels(image_2d_array); + let num_layers_2d: i32 = textureNumLayers(image_2d_array); + let num_levels_cube: i32 = textureNumLevels(image_cube); + let num_levels_cube_array: i32 = textureNumLevels(image_cube_array); + let num_layers_cube: i32 = textureNumLayers(image_cube_array); + let num_levels_3d: i32 = textureNumLevels(image_3d); + let num_samples_aa: i32 = textureNumSamples(image_aa); + let sum1: i32 = (((((((num_layers_2d + num_layers_cube) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array); + return vec4(f32(sum1)); +} + [[stage(fragment)]] fn sample() -> [[location(0)]] vec4 { let tc: vec2 = vec2(0.5); + let s1d: vec4 = textureSample(image_1d, sampler_reg, tc.x); let s2d: vec4 = textureSample(image_2d, sampler_reg, tc); let s2d_offset: vec4 = textureSample(image_2d, sampler_reg, tc, vec2(3, 1)); let s2d_level: vec4 = textureSampleLevel(image_2d, sampler_reg, tc, 2.299999952316284); let s2d_level_offset: vec4 = textureSampleLevel(image_2d, sampler_reg, tc, 2.299999952316284, vec2(3, 1)); - return (((s2d + s2d_offset) + s2d_level) + s2d_level_offset); + return ((((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset); } [[stage(fragment)]] diff --git a/tests/snapshots.rs b/tests/snapshots.rs index 54e5f6d6b2..5725695e66 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -488,7 +488,7 @@ fn convert_wgsl() { ), ( "image", - Targets::SPIRV | Targets::METAL | Targets::HLSL | Targets::WGSL, + Targets::SPIRV | Targets::METAL | Targets::HLSL | Targets::WGSL | Targets::GLSL, ), ("extra", Targets::SPIRV | Targets::METAL | Targets::WGSL), (