From f59668ccfaf7bdb3a7e43d84363a21c77357b2fe Mon Sep 17 00:00:00 2001 From: daxpedda Date: Fri, 7 Apr 2023 16:01:55 +0200 Subject: [PATCH] Allow `array_index` to be unsigned (#2298) --- src/back/spv/image.rs | 52 +- src/valid/expression.rs | 4 +- src/valid/function.rs | 2 +- tests/in/bounds-check-image-restrict.wgsl | 27 +- tests/in/bounds-check-image-rzsw.wgsl | 27 +- tests/in/image.wgsl | 59 +- ...age-restrict.fragment_shader.Fragment.glsl | 49 +- ...k-image-rzsw.fragment_shader.Fragment.glsl | 46 +- tests/out/glsl/image.main.Compute.glsl | 10 +- .../glsl/image.texture_sample.Fragment.glsl | 81 +- ...ge.texture_sample_comparison.Fragment.glsl | 36 +- tests/out/hlsl/image.hlsl | 136 +- tests/out/msl/bounds-check-image-restrict.msl | 108 +- tests/out/msl/bounds-check-image-rzsw.msl | 110 +- tests/out/msl/image.msl | 134 +- .../spv/bounds-check-image-restrict.spvasm | 723 ++++++----- tests/out/spv/bounds-check-image-rzsw.spvasm | 873 +++++++------ tests/out/spv/image.spvasm | 1151 ++++++++++------- tests/out/spv/shadow.spvasm | 2 +- tests/out/wgsl/image.wgsl | 135 +- 20 files changed, 2336 insertions(+), 1429 deletions(-) diff --git a/src/back/spv/image.rs b/src/back/spv/image.rs index dc4f249949..27f3520502 100644 --- a/src/back/spv/image.rs +++ b/src/back/spv/image.rs @@ -314,41 +314,47 @@ impl<'w> BlockContext<'w> { }; // Convert the index to the coordinate component type, if necessary. - 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, + let array_index_id = self.cached[array_index]; + let ty = &self.fun_info[array_index].ty; + let inner_ty = ty.inner_with(&self.ir_module.types); + let array_index_kind = if let Ti::Scalar { kind, width: 4 } = *inner_ty { + debug_assert!(matches!( + kind, + crate::ScalarKind::Sint | crate::ScalarKind::Uint )); - reconciled_id + kind } else { - let component_type_id = self.get_type_id(LookupType::Local(LocalType::Value { + unreachable!("we only allow i32 and u32"); + }; + let cast = match (component_kind, array_index_kind) { + (crate::ScalarKind::Sint, crate::ScalarKind::Sint) + | (crate::ScalarKind::Uint, crate::ScalarKind::Uint) => None, + (crate::ScalarKind::Sint, crate::ScalarKind::Uint) + | (crate::ScalarKind::Uint, crate::ScalarKind::Sint) => Some(spirv::Op::Bitcast), + (crate::ScalarKind::Float, crate::ScalarKind::Sint) => Some(spirv::Op::ConvertSToF), + (crate::ScalarKind::Float, crate::ScalarKind::Uint) => Some(spirv::Op::ConvertUToF), + (crate::ScalarKind::Bool, _) => unreachable!("we don't allow bool for component"), + (_, crate::ScalarKind::Bool | crate::ScalarKind::Float) => { + unreachable!("we don't allow bool or float for array index") + } + }; + let reconciled_array_index_id = if let Some(cast) = cast { + let component_ty_id = self.get_type_id(LookupType::Local(LocalType::Value { vector_size: None, kind: component_kind, width: 4, pointer_space: None, })); - let reconciled_id = self.gen_id(); block.body.push(Instruction::unary( - spirv::Op::ConvertUToF, - component_type_id, + cast, + component_ty_id, reconciled_id, - array_index_i32_id, + array_index_id, )); reconciled_id + } else { + array_index_id }; // Find the SPIR-V type for the combined coordinates/index vector. diff --git a/src/valid/expression.rs b/src/valid/expression.rs index 1a91fe4d0a..3a81707904 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -359,7 +359,7 @@ impl super::Validator { if let Some(expr) = array_index { match resolver[expr] { Ti::Scalar { - kind: Sk::Sint, + kind: Sk::Sint | Sk::Uint, width: _, } => {} _ => return Err(ExpressionError::InvalidImageArrayIndexType(expr)), @@ -548,7 +548,7 @@ impl super::Validator { if let Some(expr) = array_index { match resolver[expr] { Ti::Scalar { - kind: Sk::Sint, + kind: Sk::Sint | Sk::Uint, width: _, } => {} _ => return Err(ExpressionError::InvalidImageArrayIndexType(expr)), diff --git a/src/valid/function.rs b/src/valid/function.rs index 737f33dc28..8ebcddd9e3 100644 --- a/src/valid/function.rs +++ b/src/valid/function.rs @@ -758,7 +758,7 @@ impl super::Validator { if let Some(expr) = array_index { match *context.resolve_type(expr, &self.valid_expression_set)? { Ti::Scalar { - kind: crate::ScalarKind::Sint, + kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint, width: _, } => {} _ => { diff --git a/tests/in/bounds-check-image-restrict.wgsl b/tests/in/bounds-check-image-restrict.wgsl index 0cce20eafe..4bd8d117c8 100644 --- a/tests/in/bounds-check-image-restrict.wgsl +++ b/tests/in/bounds-check-image-restrict.wgsl @@ -15,7 +15,11 @@ fn test_textureLoad_2d(coords: vec2, level: i32) -> vec4 { @group(0) @binding(2) var image_2d_array: texture_2d_array; -fn test_textureLoad_2d_array(coords: vec2, index: i32, level: i32) -> vec4 { +fn test_textureLoad_2d_array_u(coords: vec2, index: u32, level: i32) -> vec4 { + return textureLoad(image_2d_array, coords, index, level); +} + +fn test_textureLoad_2d_array_s(coords: vec2, index: i32, level: i32) -> vec4 { return textureLoad(image_2d_array, coords, index, level); } @@ -43,7 +47,11 @@ fn test_textureLoad_depth_2d(coords: vec2, level: i32) -> f32 { @group(0) @binding(6) var image_depth_2d_array: texture_depth_2d_array; -fn test_textureLoad_depth_2d_array(coords: vec2, index: i32, level: i32) -> f32 { +fn test_textureLoad_depth_2d_array_u(coords: vec2, index: u32, level: i32) -> f32 { + return textureLoad(image_depth_2d_array, coords, index, level); +} + +fn test_textureLoad_depth_2d_array_s(coords: vec2, index: i32, level: i32) -> f32 { return textureLoad(image_depth_2d_array, coords, index, level); } @@ -71,7 +79,11 @@ fn test_textureStore_2d(coords: vec2, value: vec4) { @group(0) @binding(10) var image_storage_2d_array: texture_storage_2d_array; -fn test_textureStore_2d_array(coords: vec2, array_index: i32, value: vec4) { +fn test_textureStore_2d_array_u(coords: vec2, array_index: u32, value: vec4) { + textureStore(image_storage_2d_array, coords, array_index, value); +} + +fn test_textureStore_2d_array_s(coords: vec2, array_index: i32, value: vec4) { textureStore(image_storage_2d_array, coords, array_index, value); } @@ -88,16 +100,19 @@ fn test_textureStore_3d(coords: vec3, value: vec4) { fn fragment_shader() -> @location(0) vec4 { test_textureLoad_1d(0, 0); test_textureLoad_2d(vec2(), 0); - test_textureLoad_2d_array(vec2(), 0, 0); + test_textureLoad_2d_array_u(vec2(), 0u, 0); + test_textureLoad_2d_array_s(vec2(), 0, 0); test_textureLoad_3d(vec3(), 0); test_textureLoad_multisampled_2d(vec2(), 0); // Not yet implemented for GLSL: // test_textureLoad_depth_2d(vec2(), 0); - // test_textureLoad_depth_2d_array(vec2(), 0, 0); + // test_textureLoad_depth_2d_array_u(vec2(), 0u, 0); + // test_textureLoad_depth_2d_array_s(vec2(), 0, 0); // test_textureLoad_depth_multisampled_2d(vec2(), 0); test_textureStore_1d(0, vec4()); test_textureStore_2d(vec2(), vec4()); - test_textureStore_2d_array(vec2(), 0, vec4()); + test_textureStore_2d_array_u(vec2(), 0u, vec4()); + test_textureStore_2d_array_s(vec2(), 0, vec4()); test_textureStore_3d(vec3(), vec4()); return vec4(0.,0.,0.,0.); diff --git a/tests/in/bounds-check-image-rzsw.wgsl b/tests/in/bounds-check-image-rzsw.wgsl index 0cce20eafe..4bd8d117c8 100644 --- a/tests/in/bounds-check-image-rzsw.wgsl +++ b/tests/in/bounds-check-image-rzsw.wgsl @@ -15,7 +15,11 @@ fn test_textureLoad_2d(coords: vec2, level: i32) -> vec4 { @group(0) @binding(2) var image_2d_array: texture_2d_array; -fn test_textureLoad_2d_array(coords: vec2, index: i32, level: i32) -> vec4 { +fn test_textureLoad_2d_array_u(coords: vec2, index: u32, level: i32) -> vec4 { + return textureLoad(image_2d_array, coords, index, level); +} + +fn test_textureLoad_2d_array_s(coords: vec2, index: i32, level: i32) -> vec4 { return textureLoad(image_2d_array, coords, index, level); } @@ -43,7 +47,11 @@ fn test_textureLoad_depth_2d(coords: vec2, level: i32) -> f32 { @group(0) @binding(6) var image_depth_2d_array: texture_depth_2d_array; -fn test_textureLoad_depth_2d_array(coords: vec2, index: i32, level: i32) -> f32 { +fn test_textureLoad_depth_2d_array_u(coords: vec2, index: u32, level: i32) -> f32 { + return textureLoad(image_depth_2d_array, coords, index, level); +} + +fn test_textureLoad_depth_2d_array_s(coords: vec2, index: i32, level: i32) -> f32 { return textureLoad(image_depth_2d_array, coords, index, level); } @@ -71,7 +79,11 @@ fn test_textureStore_2d(coords: vec2, value: vec4) { @group(0) @binding(10) var image_storage_2d_array: texture_storage_2d_array; -fn test_textureStore_2d_array(coords: vec2, array_index: i32, value: vec4) { +fn test_textureStore_2d_array_u(coords: vec2, array_index: u32, value: vec4) { + textureStore(image_storage_2d_array, coords, array_index, value); +} + +fn test_textureStore_2d_array_s(coords: vec2, array_index: i32, value: vec4) { textureStore(image_storage_2d_array, coords, array_index, value); } @@ -88,16 +100,19 @@ fn test_textureStore_3d(coords: vec3, value: vec4) { fn fragment_shader() -> @location(0) vec4 { test_textureLoad_1d(0, 0); test_textureLoad_2d(vec2(), 0); - test_textureLoad_2d_array(vec2(), 0, 0); + test_textureLoad_2d_array_u(vec2(), 0u, 0); + test_textureLoad_2d_array_s(vec2(), 0, 0); test_textureLoad_3d(vec3(), 0); test_textureLoad_multisampled_2d(vec2(), 0); // Not yet implemented for GLSL: // test_textureLoad_depth_2d(vec2(), 0); - // test_textureLoad_depth_2d_array(vec2(), 0, 0); + // test_textureLoad_depth_2d_array_u(vec2(), 0u, 0); + // test_textureLoad_depth_2d_array_s(vec2(), 0, 0); // test_textureLoad_depth_multisampled_2d(vec2(), 0); test_textureStore_1d(0, vec4()); test_textureStore_2d(vec2(), vec4()); - test_textureStore_2d_array(vec2(), 0, vec4()); + test_textureStore_2d_array_u(vec2(), 0u, vec4()); + test_textureStore_2d_array_s(vec2(), 0, vec4()); test_textureStore_3d(vec3(), vec4()); return vec4(0.,0.,0.,0.); diff --git a/tests/in/image.wgsl b/tests/in/image.wgsl index b61b101646..2bae8f9d80 100644 --- a/tests/in/image.wgsl +++ b/tests/in/image.wgsl @@ -23,14 +23,16 @@ fn main(@builtin(local_invocation_id) local_id: vec3) { 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)); + let value5 = textureLoad(image_array_src, itc, local_id.z, i32(local_id.z) + 1); + let value6 = textureLoad(image_array_src, itc, i32(local_id.z), i32(local_id.z) + 1); + let value7 = textureLoad(image_1d_src, i32(local_id.x), i32(local_id.z)); // loads with uvec2 coords. let value1u = textureLoad(image_mipmapped_src, vec2(itc), i32(local_id.z)); let value2u = textureLoad(image_multisampled_src, vec2(itc), i32(local_id.z)); let value4u = textureLoad(image_storage_src, vec2(itc)); - let value5u = textureLoad(image_array_src, vec2(itc), i32(local_id.z), i32(local_id.z) + 1); - let value6u = textureLoad(image_1d_src, u32(local_id.x), i32(local_id.z)); + let value5u = textureLoad(image_array_src, vec2(itc), local_id.z, i32(local_id.z) + 1); + let value6u = textureLoad(image_array_src, vec2(itc), i32(local_id.z), i32(local_id.z) + 1); + let value7u = 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. @@ -109,14 +111,32 @@ var sampler_reg: sampler; @fragment fn texture_sample() -> @location(0) vec4 { let tc = vec2(0.5); + let tc3 = vec3(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)); - let s2d_bias_offset = textureSampleBias(image_2d, sampler_reg, tc, 2.0, vec2(3, 1)); - return s1d + s2d + s2d_offset + s2d_level + s2d_level_offset; + var a: vec4; + a += textureSample(image_1d, sampler_reg, tc.x); + a += textureSample(image_2d, sampler_reg, tc); + a += textureSample(image_2d, sampler_reg, tc, vec2(3, 1)); + a += textureSampleLevel(image_2d, sampler_reg, tc, level); + a += textureSampleLevel(image_2d, sampler_reg, tc, level, vec2(3, 1)); + a += textureSampleBias(image_2d, sampler_reg, tc, 2.0, vec2(3, 1)); + a += textureSample(image_2d_array, sampler_reg, tc, 0u); + a += textureSample(image_2d_array, sampler_reg, tc, 0u, vec2(3, 1)); + a += textureSampleLevel(image_2d_array, sampler_reg, tc, 0u, level); + a += textureSampleLevel(image_2d_array, sampler_reg, tc, 0u, level, vec2(3, 1)); + a += textureSampleBias(image_2d_array, sampler_reg, tc, 0u, 2.0, vec2(3, 1)); + a += textureSample(image_2d_array, sampler_reg, tc, 0); + a += textureSample(image_2d_array, sampler_reg, tc, 0, vec2(3, 1)); + a += textureSampleLevel(image_2d_array, sampler_reg, tc, 0, level); + a += textureSampleLevel(image_2d_array, sampler_reg, tc, 0, level, vec2(3, 1)); + a += textureSampleBias(image_2d_array, sampler_reg, tc, 0, 2.0, vec2(3, 1)); + a += textureSample(image_cube_array, sampler_reg, tc3, 0u); + a += textureSampleLevel(image_cube_array, sampler_reg, tc3, 0u, level); + a += textureSampleBias(image_cube_array, sampler_reg, tc3, 0u, 2.0); + a += textureSample(image_cube_array, sampler_reg, tc3, 0); + a += textureSampleLevel(image_cube_array, sampler_reg, tc3, 0, level); + a += textureSampleBias(image_cube_array, sampler_reg, tc3, 0, 2.0); + return a; } @group(1) @binding(1) @@ -124,16 +144,25 @@ var sampler_cmp: sampler_comparison; @group(1) @binding(2) var image_2d_depth: texture_depth_2d; @group(1) @binding(3) +var image_2d_array_depth: texture_depth_2d_array; +@group(1) @binding(4) var image_cube_depth: texture_depth_cube; @fragment fn texture_sample_comparison() -> @location(0) f32 { let tc = vec2(0.5); + let tc3 = vec3(0.5); let dref = 0.5; - let s2d_depth = textureSampleCompare(image_2d_depth, sampler_cmp, tc, dref); - let s2d_depth_level = textureSampleCompareLevel(image_2d_depth, sampler_cmp, tc, dref); - let scube_depth_level = textureSampleCompareLevel(image_cube_depth, sampler_cmp, vec3(0.5), dref); - return s2d_depth + s2d_depth_level; + var a: f32; + a += textureSampleCompare(image_2d_depth, sampler_cmp, tc, dref); + a += textureSampleCompare(image_2d_array_depth, sampler_cmp, tc, 0u, dref); + a += textureSampleCompare(image_2d_array_depth, sampler_cmp, tc, 0, dref); + a += textureSampleCompare(image_cube_depth, sampler_cmp, tc3, dref); + a += textureSampleCompareLevel(image_2d_depth, sampler_cmp, tc, dref); + a += textureSampleCompareLevel(image_2d_array_depth, sampler_cmp, tc, 0u, dref); + a += textureSampleCompareLevel(image_2d_array_depth, sampler_cmp, tc, 0, dref); + a += textureSampleCompareLevel(image_cube_depth, sampler_cmp, tc3, dref); + return a; } @fragment diff --git a/tests/out/glsl/bounds-check-image-restrict.fragment_shader.Fragment.glsl b/tests/out/glsl/bounds-check-image-restrict.fragment_shader.Fragment.glsl index 8196d197dd..ba7854322f 100644 --- a/tests/out/glsl/bounds-check-image-restrict.fragment_shader.Fragment.glsl +++ b/tests/out/glsl/bounds-check-image-restrict.fragment_shader.Fragment.glsl @@ -32,53 +32,66 @@ vec4 test_textureLoad_2d(ivec2 coords_1, int level_1) { return _e3; } -vec4 test_textureLoad_2d_array(ivec2 coords_2, int index, int level_2) { +vec4 test_textureLoad_2d_array_u(ivec2 coords_2, uint index, int level_2) { int _e4_clamped_lod = clamp(level_2, 0, textureQueryLevels(_group_0_binding_2_fs) - 1); vec4 _e4 = texelFetch(_group_0_binding_2_fs, clamp(ivec3(coords_2, index), ivec3(0), textureSize(_group_0_binding_2_fs, _e4_clamped_lod) - ivec3(1)), _e4_clamped_lod); return _e4; } -vec4 test_textureLoad_3d(ivec3 coords_3, int level_3) { - int _e3_clamped_lod = clamp(level_3, 0, textureQueryLevels(_group_0_binding_3_fs) - 1); - vec4 _e3 = texelFetch(_group_0_binding_3_fs, clamp(coords_3, ivec3(0), textureSize(_group_0_binding_3_fs, _e3_clamped_lod) - ivec3(1)), _e3_clamped_lod); +vec4 test_textureLoad_2d_array_s(ivec2 coords_3, int index_1, int level_3) { + int _e4_clamped_lod = clamp(level_3, 0, textureQueryLevels(_group_0_binding_2_fs) - 1); + vec4 _e4 = texelFetch(_group_0_binding_2_fs, clamp(ivec3(coords_3, index_1), ivec3(0), textureSize(_group_0_binding_2_fs, _e4_clamped_lod) - ivec3(1)), _e4_clamped_lod); + return _e4; +} + +vec4 test_textureLoad_3d(ivec3 coords_4, int level_4) { + int _e3_clamped_lod = clamp(level_4, 0, textureQueryLevels(_group_0_binding_3_fs) - 1); + vec4 _e3 = texelFetch(_group_0_binding_3_fs, clamp(coords_4, ivec3(0), textureSize(_group_0_binding_3_fs, _e3_clamped_lod) - ivec3(1)), _e3_clamped_lod); return _e3; } -vec4 test_textureLoad_multisampled_2d(ivec2 coords_4, int _sample) { - vec4 _e3 = texelFetch(_group_0_binding_4_fs, clamp(coords_4, ivec2(0), textureSize(_group_0_binding_4_fs) - ivec2(1)), clamp(_sample, 0, textureSamples(_group_0_binding_4_fs) - 1) +vec4 test_textureLoad_multisampled_2d(ivec2 coords_5, int _sample) { + vec4 _e3 = texelFetch(_group_0_binding_4_fs, clamp(coords_5, ivec2(0), textureSize(_group_0_binding_4_fs) - ivec2(1)), clamp(_sample, 0, textureSamples(_group_0_binding_4_fs) - 1) ); return _e3; } -void test_textureStore_1d(int coords_8, vec4 value) { - imageStore(_group_0_binding_8_fs, coords_8, value); +void test_textureStore_1d(int coords_10, vec4 value) { + imageStore(_group_0_binding_8_fs, coords_10, value); return; } -void test_textureStore_2d(ivec2 coords_9, vec4 value_1) { - imageStore(_group_0_binding_9_fs, coords_9, value_1); +void test_textureStore_2d(ivec2 coords_11, vec4 value_1) { + imageStore(_group_0_binding_9_fs, coords_11, value_1); return; } -void test_textureStore_2d_array(ivec2 coords_10, int array_index, vec4 value_2) { - imageStore(_group_0_binding_10_fs, ivec3(coords_10, array_index), value_2); +void test_textureStore_2d_array_u(ivec2 coords_12, uint array_index, vec4 value_2) { + imageStore(_group_0_binding_10_fs, ivec3(coords_12, array_index), value_2); return; } -void test_textureStore_3d(ivec3 coords_11, vec4 value_3) { - imageStore(_group_0_binding_11_fs, coords_11, value_3); +void test_textureStore_2d_array_s(ivec2 coords_13, int array_index_1, vec4 value_3) { + imageStore(_group_0_binding_10_fs, ivec3(coords_13, array_index_1), value_3); + return; +} + +void test_textureStore_3d(ivec3 coords_14, vec4 value_4) { + imageStore(_group_0_binding_11_fs, coords_14, value_4); return; } void main() { vec4 _e2 = test_textureLoad_1d(0, 0); vec4 _e5 = test_textureLoad_2d(ivec2(0, 0), 0); - vec4 _e9 = test_textureLoad_2d_array(ivec2(0, 0), 0, 0); - vec4 _e12 = test_textureLoad_3d(ivec3(0, 0, 0), 0); - vec4 _e15 = test_textureLoad_multisampled_2d(ivec2(0, 0), 0); + vec4 _e9 = test_textureLoad_2d_array_u(ivec2(0, 0), 0u, 0); + vec4 _e13 = test_textureLoad_2d_array_s(ivec2(0, 0), 0, 0); + vec4 _e16 = test_textureLoad_3d(ivec3(0, 0, 0), 0); + vec4 _e19 = test_textureLoad_multisampled_2d(ivec2(0, 0), 0); test_textureStore_1d(0, vec4(0.0, 0.0, 0.0, 0.0)); test_textureStore_2d(ivec2(0, 0), vec4(0.0, 0.0, 0.0, 0.0)); - test_textureStore_2d_array(ivec2(0, 0), 0, vec4(0.0, 0.0, 0.0, 0.0)); + test_textureStore_2d_array_u(ivec2(0, 0), 0u, vec4(0.0, 0.0, 0.0, 0.0)); + test_textureStore_2d_array_s(ivec2(0, 0), 0, vec4(0.0, 0.0, 0.0, 0.0)); test_textureStore_3d(ivec3(0, 0, 0), vec4(0.0, 0.0, 0.0, 0.0)); _fs2p_location0 = vec4(0.0, 0.0, 0.0, 0.0); return; diff --git a/tests/out/glsl/bounds-check-image-rzsw.fragment_shader.Fragment.glsl b/tests/out/glsl/bounds-check-image-rzsw.fragment_shader.Fragment.glsl index 50d1590d29..7d8aba65c6 100644 --- a/tests/out/glsl/bounds-check-image-rzsw.fragment_shader.Fragment.glsl +++ b/tests/out/glsl/bounds-check-image-rzsw.fragment_shader.Fragment.glsl @@ -30,50 +30,62 @@ vec4 test_textureLoad_2d(ivec2 coords_1, int level_1) { return _e3; } -vec4 test_textureLoad_2d_array(ivec2 coords_2, int index, int level_2) { +vec4 test_textureLoad_2d_array_u(ivec2 coords_2, uint index, int level_2) { vec4 _e4 = (level_2 < textureQueryLevels(_group_0_binding_2_fs) && all(lessThan(ivec3(coords_2, index), textureSize(_group_0_binding_2_fs, level_2))) ? texelFetch(_group_0_binding_2_fs, ivec3(coords_2, index), level_2) : vec4(0.0)); return _e4; } -vec4 test_textureLoad_3d(ivec3 coords_3, int level_3) { - vec4 _e3 = (level_3 < textureQueryLevels(_group_0_binding_3_fs) && all(lessThan(coords_3, textureSize(_group_0_binding_3_fs, level_3))) ? texelFetch(_group_0_binding_3_fs, coords_3, level_3) : vec4(0.0)); +vec4 test_textureLoad_2d_array_s(ivec2 coords_3, int index_1, int level_3) { + vec4 _e4 = (level_3 < textureQueryLevels(_group_0_binding_2_fs) && all(lessThan(ivec3(coords_3, index_1), textureSize(_group_0_binding_2_fs, level_3))) ? texelFetch(_group_0_binding_2_fs, ivec3(coords_3, index_1), level_3) : vec4(0.0)); + return _e4; +} + +vec4 test_textureLoad_3d(ivec3 coords_4, int level_4) { + vec4 _e3 = (level_4 < textureQueryLevels(_group_0_binding_3_fs) && all(lessThan(coords_4, textureSize(_group_0_binding_3_fs, level_4))) ? texelFetch(_group_0_binding_3_fs, coords_4, level_4) : vec4(0.0)); return _e3; } -vec4 test_textureLoad_multisampled_2d(ivec2 coords_4, int _sample) { - vec4 _e3 = (_sample < textureSamples(_group_0_binding_4_fs) && all(lessThan(coords_4, textureSize(_group_0_binding_4_fs))) ? texelFetch(_group_0_binding_4_fs, coords_4, _sample) : vec4(0.0)); +vec4 test_textureLoad_multisampled_2d(ivec2 coords_5, int _sample) { + vec4 _e3 = (_sample < textureSamples(_group_0_binding_4_fs) && all(lessThan(coords_5, textureSize(_group_0_binding_4_fs))) ? texelFetch(_group_0_binding_4_fs, coords_5, _sample) : vec4(0.0)); return _e3; } -void test_textureStore_1d(int coords_8, vec4 value) { - imageStore(_group_0_binding_8_fs, coords_8, value); +void test_textureStore_1d(int coords_10, vec4 value) { + imageStore(_group_0_binding_8_fs, coords_10, value); return; } -void test_textureStore_2d(ivec2 coords_9, vec4 value_1) { - imageStore(_group_0_binding_9_fs, coords_9, value_1); +void test_textureStore_2d(ivec2 coords_11, vec4 value_1) { + imageStore(_group_0_binding_9_fs, coords_11, value_1); return; } -void test_textureStore_2d_array(ivec2 coords_10, int array_index, vec4 value_2) { - imageStore(_group_0_binding_10_fs, ivec3(coords_10, array_index), value_2); +void test_textureStore_2d_array_u(ivec2 coords_12, uint array_index, vec4 value_2) { + imageStore(_group_0_binding_10_fs, ivec3(coords_12, array_index), value_2); return; } -void test_textureStore_3d(ivec3 coords_11, vec4 value_3) { - imageStore(_group_0_binding_11_fs, coords_11, value_3); +void test_textureStore_2d_array_s(ivec2 coords_13, int array_index_1, vec4 value_3) { + imageStore(_group_0_binding_10_fs, ivec3(coords_13, array_index_1), value_3); + return; +} + +void test_textureStore_3d(ivec3 coords_14, vec4 value_4) { + imageStore(_group_0_binding_11_fs, coords_14, value_4); return; } void main() { vec4 _e2 = test_textureLoad_1d(0, 0); vec4 _e5 = test_textureLoad_2d(ivec2(0, 0), 0); - vec4 _e9 = test_textureLoad_2d_array(ivec2(0, 0), 0, 0); - vec4 _e12 = test_textureLoad_3d(ivec3(0, 0, 0), 0); - vec4 _e15 = test_textureLoad_multisampled_2d(ivec2(0, 0), 0); + vec4 _e9 = test_textureLoad_2d_array_u(ivec2(0, 0), 0u, 0); + vec4 _e13 = test_textureLoad_2d_array_s(ivec2(0, 0), 0, 0); + vec4 _e16 = test_textureLoad_3d(ivec3(0, 0, 0), 0); + vec4 _e19 = test_textureLoad_multisampled_2d(ivec2(0, 0), 0); test_textureStore_1d(0, vec4(0.0, 0.0, 0.0, 0.0)); test_textureStore_2d(ivec2(0, 0), vec4(0.0, 0.0, 0.0, 0.0)); - test_textureStore_2d_array(ivec2(0, 0), 0, vec4(0.0, 0.0, 0.0, 0.0)); + test_textureStore_2d_array_u(ivec2(0, 0), 0u, vec4(0.0, 0.0, 0.0, 0.0)); + test_textureStore_2d_array_s(ivec2(0, 0), 0, vec4(0.0, 0.0, 0.0, 0.0)); test_textureStore_3d(ivec3(0, 0, 0), vec4(0.0, 0.0, 0.0, 0.0)); _fs2p_location0 = vec4(0.0, 0.0, 0.0, 0.0); return; diff --git a/tests/out/glsl/image.main.Compute.glsl b/tests/out/glsl/image.main.Compute.glsl index b03b7ee9cf..78324dd4cc 100644 --- a/tests/out/glsl/image.main.Compute.glsl +++ b/tests/out/glsl/image.main.Compute.glsl @@ -26,13 +26,15 @@ void main() { uvec4 value1_ = texelFetch(_group_0_binding_0_cs, itc, int(local_id.z)); 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)); - uvec4 value6_ = texelFetch(_group_0_binding_7_cs, ivec2(int(local_id.x), 0), int(local_id.z)); + uvec4 value5_ = texelFetch(_group_0_binding_5_cs, ivec3(itc, local_id.z), (int(local_id.z) + 1)); + uvec4 value6_ = texelFetch(_group_0_binding_5_cs, ivec3(itc, int(local_id.z)), (int(local_id.z) + 1)); + uvec4 value7_ = 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)); + uvec4 value5u = texelFetch(_group_0_binding_5_cs, ivec3(uvec2(itc), local_id.z), (int(local_id.z) + 1)); + uvec4 value6u = texelFetch(_group_0_binding_5_cs, ivec3(uvec2(itc), int(local_id.z)), (int(local_id.z) + 1)); + uvec4 value7u = 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; diff --git a/tests/out/glsl/image.texture_sample.Fragment.glsl b/tests/out/glsl/image.texture_sample.Fragment.glsl index 5ebec0fc45..72aa8fa166 100644 --- a/tests/out/glsl/image.texture_sample.Fragment.glsl +++ b/tests/out/glsl/image.texture_sample.Fragment.glsl @@ -8,17 +8,84 @@ uniform highp sampler2D _group_0_binding_0_fs; uniform highp sampler2D _group_0_binding_1_fs; +uniform highp sampler2DArray _group_0_binding_4_fs; + +uniform highp samplerCubeArray _group_0_binding_6_fs; + layout(location = 0) out vec4 _fs2p_location0; void main() { + vec4 a = vec4(0.0); vec2 tc = vec2(0.5); - vec4 s1d = texture(_group_0_binding_0_fs, vec2(tc.x, 0.0)); - vec4 s2d = texture(_group_0_binding_1_fs, vec2(tc)); - vec4 s2d_offset = textureOffset(_group_0_binding_1_fs, vec2(tc), ivec2(3, 1)); - vec4 s2d_level = textureLod(_group_0_binding_1_fs, vec2(tc), 2.299999952316284); - vec4 s2d_level_offset = textureLodOffset(_group_0_binding_1_fs, vec2(tc), 2.299999952316284, ivec2(3, 1)); - vec4 s2d_bias_offset = textureOffset(_group_0_binding_1_fs, vec2(tc), ivec2(3, 1), 2.0); - _fs2p_location0 = ((((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset); + vec3 tc3_ = vec3(0.5); + vec4 _e9 = texture(_group_0_binding_0_fs, vec2(tc.x, 0.0)); + vec4 _e10 = a; + a = (_e10 + _e9); + vec4 _e14 = texture(_group_0_binding_1_fs, vec2(tc)); + vec4 _e15 = a; + a = (_e15 + _e14); + vec4 _e19 = textureOffset(_group_0_binding_1_fs, vec2(tc), ivec2(3, 1)); + vec4 _e20 = a; + a = (_e20 + _e19); + vec4 _e24 = textureLod(_group_0_binding_1_fs, vec2(tc), 2.299999952316284); + vec4 _e25 = a; + a = (_e25 + _e24); + vec4 _e29 = textureLodOffset(_group_0_binding_1_fs, vec2(tc), 2.299999952316284, ivec2(3, 1)); + vec4 _e30 = a; + a = (_e30 + _e29); + vec4 _e35 = textureOffset(_group_0_binding_1_fs, vec2(tc), ivec2(3, 1), 2.0); + vec4 _e36 = a; + a = (_e36 + _e35); + vec4 _e41 = texture(_group_0_binding_4_fs, vec3(tc, 0u)); + vec4 _e42 = a; + a = (_e42 + _e41); + vec4 _e47 = textureOffset(_group_0_binding_4_fs, vec3(tc, 0u), ivec2(3, 1)); + vec4 _e48 = a; + a = (_e48 + _e47); + vec4 _e53 = textureLod(_group_0_binding_4_fs, vec3(tc, 0u), 2.299999952316284); + vec4 _e54 = a; + a = (_e54 + _e53); + vec4 _e59 = textureLodOffset(_group_0_binding_4_fs, vec3(tc, 0u), 2.299999952316284, ivec2(3, 1)); + vec4 _e60 = a; + a = (_e60 + _e59); + vec4 _e66 = textureOffset(_group_0_binding_4_fs, vec3(tc, 0u), ivec2(3, 1), 2.0); + vec4 _e67 = a; + a = (_e67 + _e66); + vec4 _e72 = texture(_group_0_binding_4_fs, vec3(tc, 0)); + vec4 _e73 = a; + a = (_e73 + _e72); + vec4 _e78 = textureOffset(_group_0_binding_4_fs, vec3(tc, 0), ivec2(3, 1)); + vec4 _e79 = a; + a = (_e79 + _e78); + vec4 _e84 = textureLod(_group_0_binding_4_fs, vec3(tc, 0), 2.299999952316284); + vec4 _e85 = a; + a = (_e85 + _e84); + vec4 _e90 = textureLodOffset(_group_0_binding_4_fs, vec3(tc, 0), 2.299999952316284, ivec2(3, 1)); + vec4 _e91 = a; + a = (_e91 + _e90); + vec4 _e97 = textureOffset(_group_0_binding_4_fs, vec3(tc, 0), ivec2(3, 1), 2.0); + vec4 _e98 = a; + a = (_e98 + _e97); + vec4 _e103 = texture(_group_0_binding_6_fs, vec4(tc3_, 0u)); + vec4 _e104 = a; + a = (_e104 + _e103); + vec4 _e109 = textureLod(_group_0_binding_6_fs, vec4(tc3_, 0u), 2.299999952316284); + vec4 _e110 = a; + a = (_e110 + _e109); + vec4 _e116 = texture(_group_0_binding_6_fs, vec4(tc3_, 0u), 2.0); + vec4 _e117 = a; + a = (_e117 + _e116); + vec4 _e122 = texture(_group_0_binding_6_fs, vec4(tc3_, 0)); + vec4 _e123 = a; + a = (_e123 + _e122); + vec4 _e128 = textureLod(_group_0_binding_6_fs, vec4(tc3_, 0), 2.299999952316284); + vec4 _e129 = a; + a = (_e129 + _e128); + vec4 _e135 = texture(_group_0_binding_6_fs, vec4(tc3_, 0), 2.0); + vec4 _e136 = a; + a = (_e136 + _e135); + vec4 _e138 = a; + _fs2p_location0 = _e138; return; } diff --git a/tests/out/glsl/image.texture_sample_comparison.Fragment.glsl b/tests/out/glsl/image.texture_sample_comparison.Fragment.glsl index ab7fda49f4..1dc303ed6f 100644 --- a/tests/out/glsl/image.texture_sample_comparison.Fragment.glsl +++ b/tests/out/glsl/image.texture_sample_comparison.Fragment.glsl @@ -6,16 +6,42 @@ precision highp int; uniform highp sampler2DShadow _group_1_binding_2_fs; -uniform highp samplerCubeShadow _group_1_binding_3_fs; +uniform highp sampler2DArrayShadow _group_1_binding_3_fs; + +uniform highp samplerCubeShadow _group_1_binding_4_fs; layout(location = 0) out float _fs2p_location0; void main() { + float a_1 = 0.0; vec2 tc = vec2(0.5); - float s2d_depth = texture(_group_1_binding_2_fs, vec3(tc, 0.5)); - float s2d_depth_level = textureLod(_group_1_binding_2_fs, vec3(tc, 0.5), 0.0); - float scube_depth_level = textureGrad(_group_1_binding_3_fs, vec4(vec3(0.5), 0.5), vec3(0.0), vec3(0.0)); - _fs2p_location0 = (s2d_depth + s2d_depth_level); + vec3 tc3_ = vec3(0.5); + float _e8 = texture(_group_1_binding_2_fs, vec3(tc, 0.5)); + float _e9 = a_1; + a_1 = (_e9 + _e8); + float _e14 = texture(_group_1_binding_3_fs, vec4(tc, 0u, 0.5)); + float _e15 = a_1; + a_1 = (_e15 + _e14); + float _e20 = texture(_group_1_binding_3_fs, vec4(tc, 0, 0.5)); + float _e21 = a_1; + a_1 = (_e21 + _e20); + float _e25 = texture(_group_1_binding_4_fs, vec4(tc3_, 0.5)); + float _e26 = a_1; + a_1 = (_e26 + _e25); + float _e30 = textureLod(_group_1_binding_2_fs, vec3(tc, 0.5), 0.0); + float _e31 = a_1; + a_1 = (_e31 + _e30); + float _e36 = textureGrad(_group_1_binding_3_fs, vec4(tc, 0u, 0.5), vec2(0.0), vec2(0.0)); + float _e37 = a_1; + a_1 = (_e37 + _e36); + float _e42 = textureGrad(_group_1_binding_3_fs, vec4(tc, 0, 0.5), vec2(0.0), vec2(0.0)); + float _e43 = a_1; + a_1 = (_e43 + _e42); + float _e47 = textureGrad(_group_1_binding_4_fs, vec4(tc3_, 0.5), vec3(0.0), vec3(0.0)); + float _e48 = a_1; + a_1 = (_e48 + _e47); + float _e50 = a_1; + _fs2p_location0 = _e50; return; } diff --git a/tests/out/hlsl/image.hlsl b/tests/out/hlsl/image.hlsl index 3d774fa73a..79e490e249 100644 --- a/tests/out/hlsl/image.hlsl +++ b/tests/out/hlsl/image.hlsl @@ -19,7 +19,8 @@ Texture2DMS image_aa : register(t8); SamplerState sampler_reg : register(s0, space1); SamplerComparisonState sampler_cmp : register(s1, space1); Texture2D image_2d_depth : register(t2, space1); -TextureCube image_cube_depth : register(t3, space1); +Texture2DArray image_2d_array_depth : register(t3, space1); +TextureCube image_cube_depth : register(t4, space1); uint2 NagaRWDimensions2D(RWTexture2D tex) { @@ -36,13 +37,15 @@ void main(uint3 local_id : SV_GroupThreadID) uint4 value1_ = image_mipmapped_src.Load(int3(itc, int(local_id.z))); 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))); - uint4 value6_ = image_1d_src.Load(int2(int(local_id.x), int(local_id.z))); + uint4 value5_ = image_array_src.Load(int4(itc, local_id.z, (int(local_id.z) + 1))); + uint4 value6_ = image_array_src.Load(int4(itc, int(local_id.z), (int(local_id.z) + 1))); + uint4 value7_ = 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))); + uint4 value5u = image_array_src.Load(int4(uint2(itc), local_id.z, (int(local_id.z) + 1))); + uint4 value6u = image_array_src.Load(int4(uint2(itc), int(local_id.z), (int(local_id.z) + 1))); + uint4 value7u = 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; @@ -240,42 +243,131 @@ float4 levels_queries() : SV_Position float4 texture_sample() : SV_Target0 { + float4 a = (float4)0; + float2 tc = (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)); - float4 s2d_bias_offset = image_2d.SampleBias(sampler_reg, tc, 2.0, int2(3, 1)); - return ((((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset); + float3 tc3_ = (0.5).xxx; + float4 _expr9 = image_1d.Sample(sampler_reg, tc.x); + float4 _expr10 = a; + a = (_expr10 + _expr9); + float4 _expr14 = image_2d.Sample(sampler_reg, tc); + float4 _expr15 = a; + a = (_expr15 + _expr14); + float4 _expr19 = image_2d.Sample(sampler_reg, tc, int2(3, 1)); + float4 _expr20 = a; + a = (_expr20 + _expr19); + float4 _expr24 = image_2d.SampleLevel(sampler_reg, tc, 2.299999952316284); + float4 _expr25 = a; + a = (_expr25 + _expr24); + float4 _expr29 = image_2d.SampleLevel(sampler_reg, tc, 2.299999952316284, int2(3, 1)); + float4 _expr30 = a; + a = (_expr30 + _expr29); + float4 _expr35 = image_2d.SampleBias(sampler_reg, tc, 2.0, int2(3, 1)); + float4 _expr36 = a; + a = (_expr36 + _expr35); + float4 _expr41 = image_2d_array.Sample(sampler_reg, float3(tc, 0u)); + float4 _expr42 = a; + a = (_expr42 + _expr41); + float4 _expr47 = image_2d_array.Sample(sampler_reg, float3(tc, 0u), int2(3, 1)); + float4 _expr48 = a; + a = (_expr48 + _expr47); + float4 _expr53 = image_2d_array.SampleLevel(sampler_reg, float3(tc, 0u), 2.299999952316284); + float4 _expr54 = a; + a = (_expr54 + _expr53); + float4 _expr59 = image_2d_array.SampleLevel(sampler_reg, float3(tc, 0u), 2.299999952316284, int2(3, 1)); + float4 _expr60 = a; + a = (_expr60 + _expr59); + float4 _expr66 = image_2d_array.SampleBias(sampler_reg, float3(tc, 0u), 2.0, int2(3, 1)); + float4 _expr67 = a; + a = (_expr67 + _expr66); + float4 _expr72 = image_2d_array.Sample(sampler_reg, float3(tc, 0)); + float4 _expr73 = a; + a = (_expr73 + _expr72); + float4 _expr78 = image_2d_array.Sample(sampler_reg, float3(tc, 0), int2(3, 1)); + float4 _expr79 = a; + a = (_expr79 + _expr78); + float4 _expr84 = image_2d_array.SampleLevel(sampler_reg, float3(tc, 0), 2.299999952316284); + float4 _expr85 = a; + a = (_expr85 + _expr84); + float4 _expr90 = image_2d_array.SampleLevel(sampler_reg, float3(tc, 0), 2.299999952316284, int2(3, 1)); + float4 _expr91 = a; + a = (_expr91 + _expr90); + float4 _expr97 = image_2d_array.SampleBias(sampler_reg, float3(tc, 0), 2.0, int2(3, 1)); + float4 _expr98 = a; + a = (_expr98 + _expr97); + float4 _expr103 = image_cube_array.Sample(sampler_reg, float4(tc3_, 0u)); + float4 _expr104 = a; + a = (_expr104 + _expr103); + float4 _expr109 = image_cube_array.SampleLevel(sampler_reg, float4(tc3_, 0u), 2.299999952316284); + float4 _expr110 = a; + a = (_expr110 + _expr109); + float4 _expr116 = image_cube_array.SampleBias(sampler_reg, float4(tc3_, 0u), 2.0); + float4 _expr117 = a; + a = (_expr117 + _expr116); + float4 _expr122 = image_cube_array.Sample(sampler_reg, float4(tc3_, 0)); + float4 _expr123 = a; + a = (_expr123 + _expr122); + float4 _expr128 = image_cube_array.SampleLevel(sampler_reg, float4(tc3_, 0), 2.299999952316284); + float4 _expr129 = a; + a = (_expr129 + _expr128); + float4 _expr135 = image_cube_array.SampleBias(sampler_reg, float4(tc3_, 0), 2.0); + float4 _expr136 = a; + a = (_expr136 + _expr135); + float4 _expr138 = a; + return _expr138; } float texture_sample_comparison() : SV_Target0 { + float a_1 = (float)0; + float2 tc_1 = (0.5).xx; - float s2d_depth = image_2d_depth.SampleCmp(sampler_cmp, tc_1, 0.5); - float s2d_depth_level = image_2d_depth.SampleCmpLevelZero(sampler_cmp, tc_1, 0.5); - float scube_depth_level = image_cube_depth.SampleCmpLevelZero(sampler_cmp, (0.5).xxx, 0.5); - return (s2d_depth + s2d_depth_level); + float3 tc3_1 = (0.5).xxx; + float _expr8 = image_2d_depth.SampleCmp(sampler_cmp, tc_1, 0.5); + float _expr9 = a_1; + a_1 = (_expr9 + _expr8); + float _expr14 = image_2d_array_depth.SampleCmp(sampler_cmp, float3(tc_1, 0u), 0.5); + float _expr15 = a_1; + a_1 = (_expr15 + _expr14); + float _expr20 = image_2d_array_depth.SampleCmp(sampler_cmp, float3(tc_1, 0), 0.5); + float _expr21 = a_1; + a_1 = (_expr21 + _expr20); + float _expr25 = image_cube_depth.SampleCmp(sampler_cmp, tc3_1, 0.5); + float _expr26 = a_1; + a_1 = (_expr26 + _expr25); + float _expr30 = image_2d_depth.SampleCmpLevelZero(sampler_cmp, tc_1, 0.5); + float _expr31 = a_1; + a_1 = (_expr31 + _expr30); + float _expr36 = image_2d_array_depth.SampleCmpLevelZero(sampler_cmp, float3(tc_1, 0u), 0.5); + float _expr37 = a_1; + a_1 = (_expr37 + _expr36); + float _expr42 = image_2d_array_depth.SampleCmpLevelZero(sampler_cmp, float3(tc_1, 0), 0.5); + float _expr43 = a_1; + a_1 = (_expr43 + _expr42); + float _expr47 = image_cube_depth.SampleCmpLevelZero(sampler_cmp, tc3_1, 0.5); + float _expr48 = a_1; + a_1 = (_expr48 + _expr47); + float _expr50 = a_1; + return _expr50; } float4 gather() : SV_Target0 { float2 tc_2 = (0.5).xx; - float4 s2d_1 = image_2d.GatherGreen(sampler_reg, tc_2); - 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 = image_2d.GatherGreen(sampler_reg, tc_2); + float4 s2d_offset = image_2d.GatherAlpha(sampler_reg, tc_2, int2(3, 1)); + float4 s2d_depth = 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)); 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); + return ((((s2d + s2d_offset) + s2d_depth) + s2d_depth_offset) + f); } float4 depth_no_comparison() : SV_Target0 { float2 tc_3 = (0.5).xx; - float s2d_2 = image_2d_depth.Sample(sampler_reg, tc_3); + float s2d_1 = image_2d_depth.Sample(sampler_reg, tc_3); float4 s2d_gather = image_2d_depth.Gather(sampler_reg, tc_3); - return ((s2d_2).xxxx + s2d_gather); + return ((s2d_1).xxxx + s2d_gather); } diff --git a/tests/out/msl/bounds-check-image-restrict.msl b/tests/out/msl/bounds-check-image-restrict.msl index c1e5b99df0..f52fc844e0 100644 --- a/tests/out/msl/bounds-check-image-restrict.msl +++ b/tests/out/msl/bounds-check-image-restrict.msl @@ -5,7 +5,7 @@ using metal::uint; constant metal::int2 const_type_4_ = {0, 0}; -constant metal::int3 const_type_7_ = {0, 0, 0}; +constant metal::int3 const_type_8_ = {0, 0, 0}; constant metal::float4 const_type_2_ = {0.0, 0.0, 0.0, 0.0}; metal::float4 test_textureLoad_1d( @@ -27,9 +27,9 @@ metal::float4 test_textureLoad_2d( return _e3; } -metal::float4 test_textureLoad_2d_array( +metal::float4 test_textureLoad_2d_array_u( metal::int2 coords_2, - int index, + uint index, int level_2, metal::texture2d_array image_2d_array ) { @@ -38,89 +38,121 @@ metal::float4 test_textureLoad_2d_array( return _e4; } -metal::float4 test_textureLoad_3d( - metal::int3 coords_3, +metal::float4 test_textureLoad_2d_array_s( + metal::int2 coords_3, + int index_1, int level_3, + metal::texture2d_array image_2d_array +) { + uint clamped_lod_e4 = metal::min(uint(level_3), image_2d_array.get_num_mip_levels() - 1); + metal::float4 _e4 = image_2d_array.read(metal::min(metal::uint2(coords_3), metal::uint2(image_2d_array.get_width(clamped_lod_e4), image_2d_array.get_height(clamped_lod_e4)) - 1), metal::min(uint(index_1), image_2d_array.get_array_size() - 1), clamped_lod_e4); + return _e4; +} + +metal::float4 test_textureLoad_3d( + metal::int3 coords_4, + int level_4, metal::texture3d image_3d ) { - uint clamped_lod_e3 = metal::min(uint(level_3), image_3d.get_num_mip_levels() - 1); - metal::float4 _e3 = image_3d.read(metal::min(metal::uint3(coords_3), metal::uint3(image_3d.get_width(clamped_lod_e3), image_3d.get_height(clamped_lod_e3), image_3d.get_depth(clamped_lod_e3)) - 1), clamped_lod_e3); + uint clamped_lod_e3 = metal::min(uint(level_4), image_3d.get_num_mip_levels() - 1); + metal::float4 _e3 = image_3d.read(metal::min(metal::uint3(coords_4), metal::uint3(image_3d.get_width(clamped_lod_e3), image_3d.get_height(clamped_lod_e3), image_3d.get_depth(clamped_lod_e3)) - 1), clamped_lod_e3); return _e3; } metal::float4 test_textureLoad_multisampled_2d( - metal::int2 coords_4, + metal::int2 coords_5, int _sample, metal::texture2d_ms image_multisampled_2d ) { - metal::float4 _e3 = image_multisampled_2d.read(metal::min(metal::uint2(coords_4), metal::uint2(image_multisampled_2d.get_width(), image_multisampled_2d.get_height()) - 1), metal::min(uint(_sample), image_multisampled_2d.get_num_samples() - 1)); + metal::float4 _e3 = image_multisampled_2d.read(metal::min(metal::uint2(coords_5), metal::uint2(image_multisampled_2d.get_width(), image_multisampled_2d.get_height()) - 1), metal::min(uint(_sample), image_multisampled_2d.get_num_samples() - 1)); return _e3; } float test_textureLoad_depth_2d( - metal::int2 coords_5, - int level_4, + metal::int2 coords_6, + int level_5, metal::depth2d image_depth_2d ) { - uint clamped_lod_e3 = metal::min(uint(level_4), image_depth_2d.get_num_mip_levels() - 1); - float _e3 = image_depth_2d.read(metal::min(metal::uint2(coords_5), metal::uint2(image_depth_2d.get_width(clamped_lod_e3), image_depth_2d.get_height(clamped_lod_e3)) - 1), clamped_lod_e3); + uint clamped_lod_e3 = metal::min(uint(level_5), image_depth_2d.get_num_mip_levels() - 1); + float _e3 = image_depth_2d.read(metal::min(metal::uint2(coords_6), metal::uint2(image_depth_2d.get_width(clamped_lod_e3), image_depth_2d.get_height(clamped_lod_e3)) - 1), clamped_lod_e3); return _e3; } -float test_textureLoad_depth_2d_array( - metal::int2 coords_6, - int index_1, - int level_5, +float test_textureLoad_depth_2d_array_u( + metal::int2 coords_7, + uint index_2, + int level_6, metal::depth2d_array image_depth_2d_array ) { - uint clamped_lod_e4 = metal::min(uint(level_5), image_depth_2d_array.get_num_mip_levels() - 1); - float _e4 = image_depth_2d_array.read(metal::min(metal::uint2(coords_6), metal::uint2(image_depth_2d_array.get_width(clamped_lod_e4), image_depth_2d_array.get_height(clamped_lod_e4)) - 1), metal::min(uint(index_1), image_depth_2d_array.get_array_size() - 1), clamped_lod_e4); + uint clamped_lod_e4 = metal::min(uint(level_6), image_depth_2d_array.get_num_mip_levels() - 1); + float _e4 = image_depth_2d_array.read(metal::min(metal::uint2(coords_7), metal::uint2(image_depth_2d_array.get_width(clamped_lod_e4), image_depth_2d_array.get_height(clamped_lod_e4)) - 1), metal::min(uint(index_2), image_depth_2d_array.get_array_size() - 1), clamped_lod_e4); + return _e4; +} + +float test_textureLoad_depth_2d_array_s( + metal::int2 coords_8, + int index_3, + int level_7, + metal::depth2d_array image_depth_2d_array +) { + uint clamped_lod_e4 = metal::min(uint(level_7), image_depth_2d_array.get_num_mip_levels() - 1); + float _e4 = image_depth_2d_array.read(metal::min(metal::uint2(coords_8), metal::uint2(image_depth_2d_array.get_width(clamped_lod_e4), image_depth_2d_array.get_height(clamped_lod_e4)) - 1), metal::min(uint(index_3), image_depth_2d_array.get_array_size() - 1), clamped_lod_e4); return _e4; } float test_textureLoad_depth_multisampled_2d( - metal::int2 coords_7, + metal::int2 coords_9, int _sample_1, metal::depth2d_ms image_depth_multisampled_2d ) { - float _e3 = image_depth_multisampled_2d.read(metal::min(metal::uint2(coords_7), metal::uint2(image_depth_multisampled_2d.get_width(), image_depth_multisampled_2d.get_height()) - 1), metal::min(uint(_sample_1), image_depth_multisampled_2d.get_num_samples() - 1)); + float _e3 = image_depth_multisampled_2d.read(metal::min(metal::uint2(coords_9), metal::uint2(image_depth_multisampled_2d.get_width(), image_depth_multisampled_2d.get_height()) - 1), metal::min(uint(_sample_1), image_depth_multisampled_2d.get_num_samples() - 1)); return _e3; } void test_textureStore_1d( - int coords_8, + int coords_10, metal::float4 value, metal::texture1d image_storage_1d ) { - image_storage_1d.write(value, metal::min(uint(coords_8), image_storage_1d.get_width() - 1)); + image_storage_1d.write(value, metal::min(uint(coords_10), image_storage_1d.get_width() - 1)); return; } void test_textureStore_2d( - metal::int2 coords_9, + metal::int2 coords_11, metal::float4 value_1, metal::texture2d image_storage_2d ) { - image_storage_2d.write(value_1, metal::min(metal::uint2(coords_9), metal::uint2(image_storage_2d.get_width(), image_storage_2d.get_height()) - 1)); + image_storage_2d.write(value_1, metal::min(metal::uint2(coords_11), metal::uint2(image_storage_2d.get_width(), image_storage_2d.get_height()) - 1)); return; } -void test_textureStore_2d_array( - metal::int2 coords_10, - int array_index, +void test_textureStore_2d_array_u( + metal::int2 coords_12, + uint array_index, metal::float4 value_2, metal::texture2d_array image_storage_2d_array ) { - image_storage_2d_array.write(value_2, metal::min(metal::uint2(coords_10), metal::uint2(image_storage_2d_array.get_width(), image_storage_2d_array.get_height()) - 1), metal::min(uint(array_index), image_storage_2d_array.get_array_size() - 1)); + image_storage_2d_array.write(value_2, metal::min(metal::uint2(coords_12), metal::uint2(image_storage_2d_array.get_width(), image_storage_2d_array.get_height()) - 1), metal::min(uint(array_index), image_storage_2d_array.get_array_size() - 1)); + return; +} + +void test_textureStore_2d_array_s( + metal::int2 coords_13, + int array_index_1, + metal::float4 value_3, + metal::texture2d_array image_storage_2d_array +) { + image_storage_2d_array.write(value_3, metal::min(metal::uint2(coords_13), metal::uint2(image_storage_2d_array.get_width(), image_storage_2d_array.get_height()) - 1), metal::min(uint(array_index_1), image_storage_2d_array.get_array_size() - 1)); return; } void test_textureStore_3d( - metal::int3 coords_11, - metal::float4 value_3, + metal::int3 coords_14, + metal::float4 value_4, metal::texture3d image_storage_3d ) { - image_storage_3d.write(value_3, metal::min(metal::uint3(coords_11), metal::uint3(image_storage_3d.get_width(), image_storage_3d.get_height(), image_storage_3d.get_depth()) - 1)); + image_storage_3d.write(value_4, metal::min(metal::uint3(coords_14), metal::uint3(image_storage_3d.get_width(), image_storage_3d.get_height(), image_storage_3d.get_depth()) - 1)); return; } @@ -140,12 +172,14 @@ fragment fragment_shaderOutput fragment_shader( ) { metal::float4 _e2 = test_textureLoad_1d(0, 0, image_1d); metal::float4 _e5 = test_textureLoad_2d(const_type_4_, 0, image_2d); - metal::float4 _e9 = test_textureLoad_2d_array(const_type_4_, 0, 0, image_2d_array); - metal::float4 _e12 = test_textureLoad_3d(const_type_7_, 0, image_3d); - metal::float4 _e15 = test_textureLoad_multisampled_2d(const_type_4_, 0, image_multisampled_2d); + metal::float4 _e9 = test_textureLoad_2d_array_u(const_type_4_, 0u, 0, image_2d_array); + metal::float4 _e13 = test_textureLoad_2d_array_s(const_type_4_, 0, 0, image_2d_array); + metal::float4 _e16 = test_textureLoad_3d(const_type_8_, 0, image_3d); + metal::float4 _e19 = test_textureLoad_multisampled_2d(const_type_4_, 0, image_multisampled_2d); test_textureStore_1d(0, const_type_2_, image_storage_1d); test_textureStore_2d(const_type_4_, const_type_2_, image_storage_2d); - test_textureStore_2d_array(const_type_4_, 0, const_type_2_, image_storage_2d_array); - test_textureStore_3d(const_type_7_, const_type_2_, image_storage_3d); + test_textureStore_2d_array_u(const_type_4_, 0u, const_type_2_, image_storage_2d_array); + test_textureStore_2d_array_s(const_type_4_, 0, const_type_2_, image_storage_2d_array); + test_textureStore_3d(const_type_8_, const_type_2_, image_storage_3d); return fragment_shaderOutput { metal::float4(0.0, 0.0, 0.0, 0.0) }; } diff --git a/tests/out/msl/bounds-check-image-rzsw.msl b/tests/out/msl/bounds-check-image-rzsw.msl index eeb03c9849..db631d7135 100644 --- a/tests/out/msl/bounds-check-image-rzsw.msl +++ b/tests/out/msl/bounds-check-image-rzsw.msl @@ -11,7 +11,7 @@ struct DefaultConstructible { }; constant metal::int2 const_type_4_ = {0, 0}; -constant metal::int3 const_type_7_ = {0, 0, 0}; +constant metal::int3 const_type_8_ = {0, 0, 0}; constant metal::float4 const_type_2_ = {0.0, 0.0, 0.0, 0.0}; metal::float4 test_textureLoad_1d( @@ -32,9 +32,9 @@ metal::float4 test_textureLoad_2d( return _e3; } -metal::float4 test_textureLoad_2d_array( +metal::float4 test_textureLoad_2d_array_u( metal::int2 coords_2, - int index, + uint index, int level_2, metal::texture2d_array image_2d_array ) { @@ -42,93 +42,125 @@ metal::float4 test_textureLoad_2d_array( return _e4; } -metal::float4 test_textureLoad_3d( - metal::int3 coords_3, +metal::float4 test_textureLoad_2d_array_s( + metal::int2 coords_3, + int index_1, int level_3, + metal::texture2d_array image_2d_array +) { + metal::float4 _e4 = (uint(level_3) < image_2d_array.get_num_mip_levels() && uint(index_1) < image_2d_array.get_array_size() && metal::all(metal::uint2(coords_3) < metal::uint2(image_2d_array.get_width(level_3), image_2d_array.get_height(level_3))) ? image_2d_array.read(metal::uint2(coords_3), index_1, level_3): DefaultConstructible()); + return _e4; +} + +metal::float4 test_textureLoad_3d( + metal::int3 coords_4, + int level_4, metal::texture3d image_3d ) { - metal::float4 _e3 = (uint(level_3) < image_3d.get_num_mip_levels() && metal::all(metal::uint3(coords_3) < metal::uint3(image_3d.get_width(level_3), image_3d.get_height(level_3), image_3d.get_depth(level_3))) ? image_3d.read(metal::uint3(coords_3), level_3): DefaultConstructible()); + metal::float4 _e3 = (uint(level_4) < image_3d.get_num_mip_levels() && metal::all(metal::uint3(coords_4) < metal::uint3(image_3d.get_width(level_4), image_3d.get_height(level_4), image_3d.get_depth(level_4))) ? image_3d.read(metal::uint3(coords_4), level_4): DefaultConstructible()); return _e3; } metal::float4 test_textureLoad_multisampled_2d( - metal::int2 coords_4, + metal::int2 coords_5, int _sample, metal::texture2d_ms image_multisampled_2d ) { - metal::float4 _e3 = (uint(_sample) < image_multisampled_2d.get_num_samples() && metal::all(metal::uint2(coords_4) < metal::uint2(image_multisampled_2d.get_width(), image_multisampled_2d.get_height())) ? image_multisampled_2d.read(metal::uint2(coords_4), _sample): DefaultConstructible()); + metal::float4 _e3 = (uint(_sample) < image_multisampled_2d.get_num_samples() && metal::all(metal::uint2(coords_5) < metal::uint2(image_multisampled_2d.get_width(), image_multisampled_2d.get_height())) ? image_multisampled_2d.read(metal::uint2(coords_5), _sample): DefaultConstructible()); return _e3; } float test_textureLoad_depth_2d( - metal::int2 coords_5, - int level_4, + metal::int2 coords_6, + int level_5, metal::depth2d image_depth_2d ) { - float _e3 = (uint(level_4) < image_depth_2d.get_num_mip_levels() && metal::all(metal::uint2(coords_5) < metal::uint2(image_depth_2d.get_width(level_4), image_depth_2d.get_height(level_4))) ? image_depth_2d.read(metal::uint2(coords_5), level_4): DefaultConstructible()); + float _e3 = (uint(level_5) < image_depth_2d.get_num_mip_levels() && metal::all(metal::uint2(coords_6) < metal::uint2(image_depth_2d.get_width(level_5), image_depth_2d.get_height(level_5))) ? image_depth_2d.read(metal::uint2(coords_6), level_5): DefaultConstructible()); return _e3; } -float test_textureLoad_depth_2d_array( - metal::int2 coords_6, - int index_1, - int level_5, +float test_textureLoad_depth_2d_array_u( + metal::int2 coords_7, + uint index_2, + int level_6, metal::depth2d_array image_depth_2d_array ) { - float _e4 = (uint(level_5) < image_depth_2d_array.get_num_mip_levels() && uint(index_1) < image_depth_2d_array.get_array_size() && metal::all(metal::uint2(coords_6) < metal::uint2(image_depth_2d_array.get_width(level_5), image_depth_2d_array.get_height(level_5))) ? image_depth_2d_array.read(metal::uint2(coords_6), index_1, level_5): DefaultConstructible()); + float _e4 = (uint(level_6) < image_depth_2d_array.get_num_mip_levels() && uint(index_2) < image_depth_2d_array.get_array_size() && metal::all(metal::uint2(coords_7) < metal::uint2(image_depth_2d_array.get_width(level_6), image_depth_2d_array.get_height(level_6))) ? image_depth_2d_array.read(metal::uint2(coords_7), index_2, level_6): DefaultConstructible()); + return _e4; +} + +float test_textureLoad_depth_2d_array_s( + metal::int2 coords_8, + int index_3, + int level_7, + metal::depth2d_array image_depth_2d_array +) { + float _e4 = (uint(level_7) < image_depth_2d_array.get_num_mip_levels() && uint(index_3) < image_depth_2d_array.get_array_size() && metal::all(metal::uint2(coords_8) < metal::uint2(image_depth_2d_array.get_width(level_7), image_depth_2d_array.get_height(level_7))) ? image_depth_2d_array.read(metal::uint2(coords_8), index_3, level_7): DefaultConstructible()); return _e4; } float test_textureLoad_depth_multisampled_2d( - metal::int2 coords_7, + metal::int2 coords_9, int _sample_1, metal::depth2d_ms image_depth_multisampled_2d ) { - float _e3 = (uint(_sample_1) < image_depth_multisampled_2d.get_num_samples() && metal::all(metal::uint2(coords_7) < metal::uint2(image_depth_multisampled_2d.get_width(), image_depth_multisampled_2d.get_height())) ? image_depth_multisampled_2d.read(metal::uint2(coords_7), _sample_1): DefaultConstructible()); + float _e3 = (uint(_sample_1) < image_depth_multisampled_2d.get_num_samples() && metal::all(metal::uint2(coords_9) < metal::uint2(image_depth_multisampled_2d.get_width(), image_depth_multisampled_2d.get_height())) ? image_depth_multisampled_2d.read(metal::uint2(coords_9), _sample_1): DefaultConstructible()); return _e3; } void test_textureStore_1d( - int coords_8, + int coords_10, metal::float4 value, metal::texture1d image_storage_1d ) { - if (uint(coords_8) < image_storage_1d.get_width()) { - image_storage_1d.write(value, uint(coords_8)); + if (uint(coords_10) < image_storage_1d.get_width()) { + image_storage_1d.write(value, uint(coords_10)); } return; } void test_textureStore_2d( - metal::int2 coords_9, + metal::int2 coords_11, metal::float4 value_1, metal::texture2d image_storage_2d ) { - if (metal::all(metal::uint2(coords_9) < metal::uint2(image_storage_2d.get_width(), image_storage_2d.get_height()))) { - image_storage_2d.write(value_1, metal::uint2(coords_9)); + if (metal::all(metal::uint2(coords_11) < metal::uint2(image_storage_2d.get_width(), image_storage_2d.get_height()))) { + image_storage_2d.write(value_1, metal::uint2(coords_11)); } return; } -void test_textureStore_2d_array( - metal::int2 coords_10, - int array_index, +void test_textureStore_2d_array_u( + metal::int2 coords_12, + uint array_index, metal::float4 value_2, metal::texture2d_array image_storage_2d_array ) { - if (uint(array_index) < image_storage_2d_array.get_array_size() && metal::all(metal::uint2(coords_10) < metal::uint2(image_storage_2d_array.get_width(), image_storage_2d_array.get_height()))) { - image_storage_2d_array.write(value_2, metal::uint2(coords_10), array_index); + if (uint(array_index) < image_storage_2d_array.get_array_size() && metal::all(metal::uint2(coords_12) < metal::uint2(image_storage_2d_array.get_width(), image_storage_2d_array.get_height()))) { + image_storage_2d_array.write(value_2, metal::uint2(coords_12), array_index); + } + return; +} + +void test_textureStore_2d_array_s( + metal::int2 coords_13, + int array_index_1, + metal::float4 value_3, + metal::texture2d_array image_storage_2d_array +) { + if (uint(array_index_1) < image_storage_2d_array.get_array_size() && metal::all(metal::uint2(coords_13) < metal::uint2(image_storage_2d_array.get_width(), image_storage_2d_array.get_height()))) { + image_storage_2d_array.write(value_3, metal::uint2(coords_13), array_index_1); } return; } void test_textureStore_3d( - metal::int3 coords_11, - metal::float4 value_3, + metal::int3 coords_14, + metal::float4 value_4, metal::texture3d image_storage_3d ) { - if (metal::all(metal::uint3(coords_11) < metal::uint3(image_storage_3d.get_width(), image_storage_3d.get_height(), image_storage_3d.get_depth()))) { - image_storage_3d.write(value_3, metal::uint3(coords_11)); + if (metal::all(metal::uint3(coords_14) < metal::uint3(image_storage_3d.get_width(), image_storage_3d.get_height(), image_storage_3d.get_depth()))) { + image_storage_3d.write(value_4, metal::uint3(coords_14)); } return; } @@ -149,12 +181,14 @@ fragment fragment_shaderOutput fragment_shader( ) { metal::float4 _e2 = test_textureLoad_1d(0, 0, image_1d); metal::float4 _e5 = test_textureLoad_2d(const_type_4_, 0, image_2d); - metal::float4 _e9 = test_textureLoad_2d_array(const_type_4_, 0, 0, image_2d_array); - metal::float4 _e12 = test_textureLoad_3d(const_type_7_, 0, image_3d); - metal::float4 _e15 = test_textureLoad_multisampled_2d(const_type_4_, 0, image_multisampled_2d); + metal::float4 _e9 = test_textureLoad_2d_array_u(const_type_4_, 0u, 0, image_2d_array); + metal::float4 _e13 = test_textureLoad_2d_array_s(const_type_4_, 0, 0, image_2d_array); + metal::float4 _e16 = test_textureLoad_3d(const_type_8_, 0, image_3d); + metal::float4 _e19 = test_textureLoad_multisampled_2d(const_type_4_, 0, image_multisampled_2d); test_textureStore_1d(0, const_type_2_, image_storage_1d); test_textureStore_2d(const_type_4_, const_type_2_, image_storage_2d); - test_textureStore_2d_array(const_type_4_, 0, const_type_2_, image_storage_2d_array); - test_textureStore_3d(const_type_7_, const_type_2_, image_storage_3d); + test_textureStore_2d_array_u(const_type_4_, 0u, const_type_2_, image_storage_2d_array); + test_textureStore_2d_array_s(const_type_4_, 0, const_type_2_, image_storage_2d_array); + test_textureStore_3d(const_type_8_, const_type_2_, image_storage_3d); return fragment_shaderOutput { metal::float4(0.0, 0.0, 0.0, 0.0) }; } diff --git a/tests/out/msl/image.msl b/tests/out/msl/image.msl index b0de96238f..685867175b 100644 --- a/tests/out/msl/image.msl +++ b/tests/out/msl/image.msl @@ -22,13 +22,15 @@ kernel void main_( 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)); 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); - metal::uint4 value6_ = image_1d_src.read(uint(static_cast(local_id.x))); + metal::uint4 value5_ = image_array_src.read(metal::uint2(itc), local_id.z, static_cast(local_id.z) + 1); + metal::uint4 value6_ = image_array_src.read(metal::uint2(itc), static_cast(local_id.z), static_cast(local_id.z) + 1); + metal::uint4 value7_ = image_1d_src.read(uint(static_cast(local_id.x))); metal::uint4 value1u = image_mipmapped_src.read(metal::uint2(static_cast(itc)), static_cast(local_id.z)); metal::uint4 value2u = image_multisampled_src.read(metal::uint2(static_cast(itc)), static_cast(local_id.z)); metal::uint4 value4u = image_storage_src.read(metal::uint2(static_cast(itc))); - metal::uint4 value5u = image_array_src.read(metal::uint2(static_cast(itc)), static_cast(local_id.z), static_cast(local_id.z) + 1); - metal::uint4 value6u = image_1d_src.read(uint(static_cast(local_id.x))); + metal::uint4 value5u = image_array_src.read(metal::uint2(static_cast(itc)), local_id.z, static_cast(local_id.z) + 1); + metal::uint4 value6u = image_array_src.read(metal::uint2(static_cast(itc)), static_cast(local_id.z), static_cast(local_id.z) + 1); + metal::uint4 value7u = image_1d_src.read(uint(static_cast(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(itc.x))); return; @@ -111,16 +113,81 @@ struct texture_sampleOutput { fragment texture_sampleOutput texture_sample( metal::texture1d image_1d [[user(fake0)]] , metal::texture2d image_2d [[user(fake0)]] +, metal::texture2d_array image_2d_array [[user(fake0)]] +, metal::texturecube_array image_cube_array [[user(fake0)]] , metal::sampler sampler_reg [[user(fake0)]] ) { + metal::float4 a = {}; 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_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_9_); - metal::float4 s2d_bias_offset = image_2d.sample(sampler_reg, tc, metal::bias(2.0), const_type_9_); - return texture_sampleOutput { (((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset }; + metal::float3 tc3_ = metal::float3(0.5); + metal::float4 _e9 = image_1d.sample(sampler_reg, tc.x); + metal::float4 _e10 = a; + a = _e10 + _e9; + metal::float4 _e14 = image_2d.sample(sampler_reg, tc); + metal::float4 _e15 = a; + a = _e15 + _e14; + metal::float4 _e19 = image_2d.sample(sampler_reg, tc, const_type_9_); + metal::float4 _e20 = a; + a = _e20 + _e19; + metal::float4 _e24 = image_2d.sample(sampler_reg, tc, metal::level(2.299999952316284)); + metal::float4 _e25 = a; + a = _e25 + _e24; + metal::float4 _e29 = image_2d.sample(sampler_reg, tc, metal::level(2.299999952316284), const_type_9_); + metal::float4 _e30 = a; + a = _e30 + _e29; + metal::float4 _e35 = image_2d.sample(sampler_reg, tc, metal::bias(2.0), const_type_9_); + metal::float4 _e36 = a; + a = _e36 + _e35; + metal::float4 _e41 = image_2d_array.sample(sampler_reg, tc, 0u); + metal::float4 _e42 = a; + a = _e42 + _e41; + metal::float4 _e47 = image_2d_array.sample(sampler_reg, tc, 0u, const_type_9_); + metal::float4 _e48 = a; + a = _e48 + _e47; + metal::float4 _e53 = image_2d_array.sample(sampler_reg, tc, 0u, metal::level(2.299999952316284)); + metal::float4 _e54 = a; + a = _e54 + _e53; + metal::float4 _e59 = image_2d_array.sample(sampler_reg, tc, 0u, metal::level(2.299999952316284), const_type_9_); + metal::float4 _e60 = a; + a = _e60 + _e59; + metal::float4 _e66 = image_2d_array.sample(sampler_reg, tc, 0u, metal::bias(2.0), const_type_9_); + metal::float4 _e67 = a; + a = _e67 + _e66; + metal::float4 _e72 = image_2d_array.sample(sampler_reg, tc, 0); + metal::float4 _e73 = a; + a = _e73 + _e72; + metal::float4 _e78 = image_2d_array.sample(sampler_reg, tc, 0, const_type_9_); + metal::float4 _e79 = a; + a = _e79 + _e78; + metal::float4 _e84 = image_2d_array.sample(sampler_reg, tc, 0, metal::level(2.299999952316284)); + metal::float4 _e85 = a; + a = _e85 + _e84; + metal::float4 _e90 = image_2d_array.sample(sampler_reg, tc, 0, metal::level(2.299999952316284), const_type_9_); + metal::float4 _e91 = a; + a = _e91 + _e90; + metal::float4 _e97 = image_2d_array.sample(sampler_reg, tc, 0, metal::bias(2.0), const_type_9_); + metal::float4 _e98 = a; + a = _e98 + _e97; + metal::float4 _e103 = image_cube_array.sample(sampler_reg, tc3_, 0u); + metal::float4 _e104 = a; + a = _e104 + _e103; + metal::float4 _e109 = image_cube_array.sample(sampler_reg, tc3_, 0u, metal::level(2.299999952316284)); + metal::float4 _e110 = a; + a = _e110 + _e109; + metal::float4 _e116 = image_cube_array.sample(sampler_reg, tc3_, 0u, metal::bias(2.0)); + metal::float4 _e117 = a; + a = _e117 + _e116; + metal::float4 _e122 = image_cube_array.sample(sampler_reg, tc3_, 0); + metal::float4 _e123 = a; + a = _e123 + _e122; + metal::float4 _e128 = image_cube_array.sample(sampler_reg, tc3_, 0, metal::level(2.299999952316284)); + metal::float4 _e129 = a; + a = _e129 + _e128; + metal::float4 _e135 = image_cube_array.sample(sampler_reg, tc3_, 0, metal::bias(2.0)); + metal::float4 _e136 = a; + a = _e136 + _e135; + metal::float4 _e138 = a; + return texture_sampleOutput { _e138 }; } @@ -130,13 +197,38 @@ struct texture_sample_comparisonOutput { fragment texture_sample_comparisonOutput texture_sample_comparison( metal::sampler sampler_cmp [[user(fake0)]] , metal::depth2d image_2d_depth [[user(fake0)]] +, metal::depth2d_array image_2d_array_depth [[user(fake0)]] , metal::depthcube image_cube_depth [[user(fake0)]] ) { + float a_1 = {}; metal::float2 tc_1 = metal::float2(0.5); - float s2d_depth = image_2d_depth.sample_compare(sampler_cmp, tc_1, 0.5); - float s2d_depth_level = image_2d_depth.sample_compare(sampler_cmp, tc_1, 0.5); - float scube_depth_level = image_cube_depth.sample_compare(sampler_cmp, metal::float3(0.5), 0.5); - return texture_sample_comparisonOutput { s2d_depth + s2d_depth_level }; + metal::float3 tc3_1 = metal::float3(0.5); + float _e8 = image_2d_depth.sample_compare(sampler_cmp, tc_1, 0.5); + float _e9 = a_1; + a_1 = _e9 + _e8; + float _e14 = image_2d_array_depth.sample_compare(sampler_cmp, tc_1, 0u, 0.5); + float _e15 = a_1; + a_1 = _e15 + _e14; + float _e20 = image_2d_array_depth.sample_compare(sampler_cmp, tc_1, 0, 0.5); + float _e21 = a_1; + a_1 = _e21 + _e20; + float _e25 = image_cube_depth.sample_compare(sampler_cmp, tc3_1, 0.5); + float _e26 = a_1; + a_1 = _e26 + _e25; + float _e30 = image_2d_depth.sample_compare(sampler_cmp, tc_1, 0.5); + float _e31 = a_1; + a_1 = _e31 + _e30; + float _e36 = image_2d_array_depth.sample_compare(sampler_cmp, tc_1, 0u, 0.5); + float _e37 = a_1; + a_1 = _e37 + _e36; + float _e42 = image_2d_array_depth.sample_compare(sampler_cmp, tc_1, 0, 0.5); + float _e43 = a_1; + a_1 = _e43 + _e42; + float _e47 = image_cube_depth.sample_compare(sampler_cmp, tc3_1, 0.5); + float _e48 = a_1; + a_1 = _e48 + _e47; + float _e50 = a_1; + return texture_sample_comparisonOutput { _e50 }; } @@ -152,14 +244,14 @@ fragment gatherOutput gather( , metal::depth2d image_2d_depth [[user(fake0)]] ) { metal::float2 tc_2 = metal::float2(0.5); - metal::float4 s2d_1 = image_2d.gather(sampler_reg, tc_2, metal::int2(0), metal::component::y); - 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 = image_2d.gather(sampler_reg, tc_2, metal::int2(0), metal::component::y); + metal::float4 s2d_offset = image_2d.gather(sampler_reg, tc_2, const_type_9_, metal::component::w); + metal::float4 s2d_depth = 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_); 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(u) + static_cast(i); - return gatherOutput { (((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset) + f }; + return gatherOutput { (((s2d + s2d_offset) + s2d_depth) + s2d_depth_offset) + f }; } @@ -171,7 +263,7 @@ fragment depth_no_comparisonOutput depth_no_comparison( , metal::depth2d image_2d_depth [[user(fake0)]] ) { metal::float2 tc_3 = metal::float2(0.5); - float s2d_2 = image_2d_depth.sample(sampler_reg, tc_3); + float s2d_1 = image_2d_depth.sample(sampler_reg, tc_3); metal::float4 s2d_gather = image_2d_depth.gather(sampler_reg, tc_3); - return depth_no_comparisonOutput { metal::float4(s2d_2) + s2d_gather }; + return depth_no_comparisonOutput { metal::float4(s2d_1) + s2d_gather }; } diff --git a/tests/out/spv/bounds-check-image-restrict.spvasm b/tests/out/spv/bounds-check-image-restrict.spvasm index 51d53f3d03..09d6ad2f40 100644 --- a/tests/out/spv/bounds-check-image-restrict.spvasm +++ b/tests/out/spv/bounds-check-image-restrict.spvasm @@ -1,380 +1,457 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 244 +; Bound: 299 OpCapability ImageQuery OpCapability Image1D OpCapability Shader OpCapability Sampled1D %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Fragment %222 "fragment_shader" %220 -OpExecutionMode %222 OriginUpperLeft +OpEntryPoint Fragment %275 "fragment_shader" %273 +OpExecutionMode %275 OriginUpperLeft OpSource GLSL 450 -OpName %25 "image_1d" -OpName %27 "image_2d" -OpName %29 "image_2d_array" -OpName %31 "image_3d" -OpName %33 "image_multisampled_2d" -OpName %35 "image_depth_2d" -OpName %37 "image_depth_2d_array" -OpName %39 "image_depth_multisampled_2d" -OpName %41 "image_storage_1d" -OpName %43 "image_storage_2d" -OpName %45 "image_storage_2d_array" -OpName %47 "image_storage_3d" -OpName %50 "coords" -OpName %51 "level" -OpName %52 "test_textureLoad_1d" -OpName %65 "coords" -OpName %66 "level" -OpName %67 "test_textureLoad_2d" -OpName %80 "coords" -OpName %81 "index" -OpName %82 "level" -OpName %83 "test_textureLoad_2d_array" -OpName %97 "coords" -OpName %98 "level" -OpName %99 "test_textureLoad_3d" -OpName %112 "coords" -OpName %113 "_sample" -OpName %114 "test_textureLoad_multisampled_2d" -OpName %126 "coords" -OpName %127 "level" -OpName %128 "test_textureLoad_depth_2d" -OpName %142 "coords" -OpName %143 "index" -OpName %144 "level" -OpName %145 "test_textureLoad_depth_2d_array" -OpName %160 "coords" -OpName %161 "_sample" -OpName %162 "test_textureLoad_depth_multisampled_2d" -OpName %175 "coords" -OpName %176 "value" -OpName %177 "test_textureStore_1d" -OpName %185 "coords" -OpName %186 "value" -OpName %187 "test_textureStore_2d" -OpName %196 "coords" -OpName %197 "array_index" -OpName %198 "value" -OpName %199 "test_textureStore_2d_array" -OpName %209 "coords" -OpName %210 "value" -OpName %211 "test_textureStore_3d" -OpName %222 "fragment_shader" -OpDecorate %25 DescriptorSet 0 -OpDecorate %25 Binding 0 +OpName %27 "image_1d" +OpName %29 "image_2d" +OpName %31 "image_2d_array" +OpName %33 "image_3d" +OpName %35 "image_multisampled_2d" +OpName %37 "image_depth_2d" +OpName %39 "image_depth_2d_array" +OpName %41 "image_depth_multisampled_2d" +OpName %43 "image_storage_1d" +OpName %45 "image_storage_2d" +OpName %47 "image_storage_2d_array" +OpName %49 "image_storage_3d" +OpName %52 "coords" +OpName %53 "level" +OpName %54 "test_textureLoad_1d" +OpName %67 "coords" +OpName %68 "level" +OpName %69 "test_textureLoad_2d" +OpName %82 "coords" +OpName %83 "index" +OpName %84 "level" +OpName %85 "test_textureLoad_2d_array_u" +OpName %100 "coords" +OpName %101 "index" +OpName %102 "level" +OpName %103 "test_textureLoad_2d_array_s" +OpName %117 "coords" +OpName %118 "level" +OpName %119 "test_textureLoad_3d" +OpName %132 "coords" +OpName %133 "_sample" +OpName %134 "test_textureLoad_multisampled_2d" +OpName %146 "coords" +OpName %147 "level" +OpName %148 "test_textureLoad_depth_2d" +OpName %162 "coords" +OpName %163 "index" +OpName %164 "level" +OpName %165 "test_textureLoad_depth_2d_array_u" +OpName %181 "coords" +OpName %182 "index" +OpName %183 "level" +OpName %184 "test_textureLoad_depth_2d_array_s" +OpName %199 "coords" +OpName %200 "_sample" +OpName %201 "test_textureLoad_depth_multisampled_2d" +OpName %214 "coords" +OpName %215 "value" +OpName %216 "test_textureStore_1d" +OpName %224 "coords" +OpName %225 "value" +OpName %226 "test_textureStore_2d" +OpName %235 "coords" +OpName %236 "array_index" +OpName %237 "value" +OpName %238 "test_textureStore_2d_array_u" +OpName %249 "coords" +OpName %250 "array_index" +OpName %251 "value" +OpName %252 "test_textureStore_2d_array_s" +OpName %262 "coords" +OpName %263 "value" +OpName %264 "test_textureStore_3d" +OpName %275 "fragment_shader" OpDecorate %27 DescriptorSet 0 -OpDecorate %27 Binding 1 +OpDecorate %27 Binding 0 OpDecorate %29 DescriptorSet 0 -OpDecorate %29 Binding 2 +OpDecorate %29 Binding 1 OpDecorate %31 DescriptorSet 0 -OpDecorate %31 Binding 3 +OpDecorate %31 Binding 2 OpDecorate %33 DescriptorSet 0 -OpDecorate %33 Binding 4 +OpDecorate %33 Binding 3 OpDecorate %35 DescriptorSet 0 -OpDecorate %35 Binding 5 +OpDecorate %35 Binding 4 OpDecorate %37 DescriptorSet 0 -OpDecorate %37 Binding 6 +OpDecorate %37 Binding 5 OpDecorate %39 DescriptorSet 0 -OpDecorate %39 Binding 7 -OpDecorate %41 NonReadable +OpDecorate %39 Binding 6 OpDecorate %41 DescriptorSet 0 -OpDecorate %41 Binding 8 +OpDecorate %41 Binding 7 OpDecorate %43 NonReadable OpDecorate %43 DescriptorSet 0 -OpDecorate %43 Binding 9 +OpDecorate %43 Binding 8 OpDecorate %45 NonReadable OpDecorate %45 DescriptorSet 0 -OpDecorate %45 Binding 10 +OpDecorate %45 Binding 9 OpDecorate %47 NonReadable OpDecorate %47 DescriptorSet 0 -OpDecorate %47 Binding 11 -OpDecorate %220 Location 0 +OpDecorate %47 Binding 10 +OpDecorate %49 NonReadable +OpDecorate %49 DescriptorSet 0 +OpDecorate %49 Binding 11 +OpDecorate %273 Location 0 %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 0 -%6 = OpTypeFloat 32 -%5 = OpConstant %6 0.0 -%7 = OpTypeImage %6 1D 0 0 0 1 Unknown -%8 = OpTypeVector %6 4 -%9 = OpTypeImage %6 2D 0 0 0 1 Unknown -%10 = OpTypeVector %4 2 -%11 = OpTypeImage %6 2D 0 1 0 1 Unknown -%12 = OpTypeImage %6 3D 0 0 0 1 Unknown -%13 = OpTypeVector %4 3 -%14 = OpTypeImage %6 2D 0 0 1 1 Unknown -%15 = OpTypeImage %6 2D 1 0 0 1 Unknown -%16 = OpTypeImage %6 2D 1 1 0 1 Unknown -%17 = OpTypeImage %6 2D 1 0 1 1 Unknown -%18 = OpTypeImage %6 1D 0 0 0 2 Rgba8 -%19 = OpTypeImage %6 2D 0 0 0 2 Rgba8 -%20 = OpTypeImage %6 2D 0 1 0 2 Rgba8 -%21 = OpTypeImage %6 3D 0 0 0 2 Rgba8 -%22 = OpConstantComposite %10 %3 %3 -%23 = OpConstantComposite %13 %3 %3 %3 -%24 = OpConstantComposite %8 %5 %5 %5 %5 -%26 = OpTypePointer UniformConstant %7 -%25 = OpVariable %26 UniformConstant +%6 = OpTypeInt 32 0 +%5 = OpConstant %6 0 +%8 = OpTypeFloat 32 +%7 = OpConstant %8 0.0 +%9 = OpTypeImage %8 1D 0 0 0 1 Unknown +%10 = OpTypeVector %8 4 +%11 = OpTypeImage %8 2D 0 0 0 1 Unknown +%12 = OpTypeVector %4 2 +%13 = OpTypeImage %8 2D 0 1 0 1 Unknown +%14 = OpTypeImage %8 3D 0 0 0 1 Unknown +%15 = OpTypeVector %4 3 +%16 = OpTypeImage %8 2D 0 0 1 1 Unknown +%17 = OpTypeImage %8 2D 1 0 0 1 Unknown +%18 = OpTypeImage %8 2D 1 1 0 1 Unknown +%19 = OpTypeImage %8 2D 1 0 1 1 Unknown +%20 = OpTypeImage %8 1D 0 0 0 2 Rgba8 +%21 = OpTypeImage %8 2D 0 0 0 2 Rgba8 +%22 = OpTypeImage %8 2D 0 1 0 2 Rgba8 +%23 = OpTypeImage %8 3D 0 0 0 2 Rgba8 +%24 = OpConstantComposite %12 %3 %3 +%25 = OpConstantComposite %15 %3 %3 %3 +%26 = OpConstantComposite %10 %7 %7 %7 %7 %28 = OpTypePointer UniformConstant %9 %27 = OpVariable %28 UniformConstant %30 = OpTypePointer UniformConstant %11 %29 = OpVariable %30 UniformConstant -%32 = OpTypePointer UniformConstant %12 +%32 = OpTypePointer UniformConstant %13 %31 = OpVariable %32 UniformConstant %34 = OpTypePointer UniformConstant %14 %33 = OpVariable %34 UniformConstant -%36 = OpTypePointer UniformConstant %15 +%36 = OpTypePointer UniformConstant %16 %35 = OpVariable %36 UniformConstant -%38 = OpTypePointer UniformConstant %16 +%38 = OpTypePointer UniformConstant %17 %37 = OpVariable %38 UniformConstant -%40 = OpTypePointer UniformConstant %17 +%40 = OpTypePointer UniformConstant %18 %39 = OpVariable %40 UniformConstant -%42 = OpTypePointer UniformConstant %18 +%42 = OpTypePointer UniformConstant %19 %41 = OpVariable %42 UniformConstant -%44 = OpTypePointer UniformConstant %19 +%44 = OpTypePointer UniformConstant %20 %43 = OpVariable %44 UniformConstant -%46 = OpTypePointer UniformConstant %20 +%46 = OpTypePointer UniformConstant %21 %45 = OpVariable %46 UniformConstant -%48 = OpTypePointer UniformConstant %21 +%48 = OpTypePointer UniformConstant %22 %47 = OpVariable %48 UniformConstant -%53 = OpTypeFunction %8 %4 %4 -%57 = OpConstant %4 1 -%68 = OpTypeFunction %8 %10 %4 -%75 = OpConstantComposite %10 %57 %57 -%84 = OpTypeFunction %8 %10 %4 %4 -%92 = OpConstantComposite %13 %57 %57 %57 -%100 = OpTypeFunction %8 %13 %4 -%107 = OpConstantComposite %13 %57 %57 %57 -%121 = OpConstantComposite %10 %57 %57 -%129 = OpTypeFunction %6 %10 %4 -%136 = OpConstantComposite %10 %57 %57 -%146 = OpTypeFunction %6 %10 %4 %4 -%154 = OpConstantComposite %13 %57 %57 %57 -%169 = OpConstantComposite %10 %57 %57 -%178 = OpTypeFunction %2 %4 %8 -%188 = OpTypeFunction %2 %10 %8 -%192 = OpConstantComposite %10 %57 %57 -%200 = OpTypeFunction %2 %10 %4 %8 -%205 = OpConstantComposite %13 %57 %57 %57 -%212 = OpTypeFunction %2 %13 %8 -%216 = OpConstantComposite %13 %57 %57 %57 -%221 = OpTypePointer Output %8 -%220 = OpVariable %221 Output -%223 = OpTypeFunction %2 -%52 = OpFunction %8 None %53 -%50 = OpFunctionParameter %4 -%51 = OpFunctionParameter %4 -%49 = OpLabel -%54 = OpLoad %7 %25 -OpBranch %55 -%55 = OpLabel -%56 = OpImageQueryLevels %4 %54 -%58 = OpISub %4 %56 %57 -%59 = OpExtInst %4 %1 UMin %51 %58 -%60 = OpImageQuerySizeLod %4 %54 %59 -%61 = OpISub %4 %60 %57 -%62 = OpExtInst %4 %1 UMin %50 %61 -%63 = OpImageFetch %8 %54 %62 Lod %59 -OpReturnValue %63 +%50 = OpTypePointer UniformConstant %23 +%49 = OpVariable %50 UniformConstant +%55 = OpTypeFunction %10 %4 %4 +%59 = OpConstant %4 1 +%70 = OpTypeFunction %10 %12 %4 +%77 = OpConstantComposite %12 %59 %59 +%86 = OpTypeFunction %10 %12 %6 %4 +%95 = OpConstantComposite %15 %59 %59 %59 +%104 = OpTypeFunction %10 %12 %4 %4 +%112 = OpConstantComposite %15 %59 %59 %59 +%120 = OpTypeFunction %10 %15 %4 +%127 = OpConstantComposite %15 %59 %59 %59 +%141 = OpConstantComposite %12 %59 %59 +%149 = OpTypeFunction %8 %12 %4 +%156 = OpConstantComposite %12 %59 %59 +%166 = OpTypeFunction %8 %12 %6 %4 +%175 = OpConstantComposite %15 %59 %59 %59 +%185 = OpTypeFunction %8 %12 %4 %4 +%193 = OpConstantComposite %15 %59 %59 %59 +%208 = OpConstantComposite %12 %59 %59 +%217 = OpTypeFunction %2 %4 %10 +%227 = OpTypeFunction %2 %12 %10 +%231 = OpConstantComposite %12 %59 %59 +%239 = OpTypeFunction %2 %12 %6 %10 +%245 = OpConstantComposite %15 %59 %59 %59 +%253 = OpTypeFunction %2 %12 %4 %10 +%258 = OpConstantComposite %15 %59 %59 %59 +%265 = OpTypeFunction %2 %15 %10 +%269 = OpConstantComposite %15 %59 %59 %59 +%274 = OpTypePointer Output %10 +%273 = OpVariable %274 Output +%276 = OpTypeFunction %2 +%54 = OpFunction %10 None %55 +%52 = OpFunctionParameter %4 +%53 = OpFunctionParameter %4 +%51 = OpLabel +%56 = OpLoad %9 %27 +OpBranch %57 +%57 = OpLabel +%58 = OpImageQueryLevels %4 %56 +%60 = OpISub %4 %58 %59 +%61 = OpExtInst %4 %1 UMin %53 %60 +%62 = OpImageQuerySizeLod %4 %56 %61 +%63 = OpISub %4 %62 %59 +%64 = OpExtInst %4 %1 UMin %52 %63 +%65 = OpImageFetch %10 %56 %64 Lod %61 +OpReturnValue %65 OpFunctionEnd -%67 = OpFunction %8 None %68 -%65 = OpFunctionParameter %10 -%66 = OpFunctionParameter %4 -%64 = OpLabel -%69 = OpLoad %9 %27 -OpBranch %70 -%70 = OpLabel -%71 = OpImageQueryLevels %4 %69 -%72 = OpISub %4 %71 %57 -%73 = OpExtInst %4 %1 UMin %66 %72 -%74 = OpImageQuerySizeLod %10 %69 %73 -%76 = OpISub %10 %74 %75 -%77 = OpExtInst %10 %1 UMin %65 %76 -%78 = OpImageFetch %8 %69 %77 Lod %73 -OpReturnValue %78 +%69 = OpFunction %10 None %70 +%67 = OpFunctionParameter %12 +%68 = OpFunctionParameter %4 +%66 = OpLabel +%71 = OpLoad %11 %29 +OpBranch %72 +%72 = OpLabel +%73 = OpImageQueryLevels %4 %71 +%74 = OpISub %4 %73 %59 +%75 = OpExtInst %4 %1 UMin %68 %74 +%76 = OpImageQuerySizeLod %12 %71 %75 +%78 = OpISub %12 %76 %77 +%79 = OpExtInst %12 %1 UMin %67 %78 +%80 = OpImageFetch %10 %71 %79 Lod %75 +OpReturnValue %80 OpFunctionEnd -%83 = OpFunction %8 None %84 -%80 = OpFunctionParameter %10 -%81 = OpFunctionParameter %4 -%82 = OpFunctionParameter %4 -%79 = OpLabel -%85 = OpLoad %11 %29 -OpBranch %86 -%86 = OpLabel -%87 = OpCompositeConstruct %13 %80 %81 -%88 = OpImageQueryLevels %4 %85 -%89 = OpISub %4 %88 %57 -%90 = OpExtInst %4 %1 UMin %82 %89 -%91 = OpImageQuerySizeLod %13 %85 %90 -%93 = OpISub %13 %91 %92 -%94 = OpExtInst %13 %1 UMin %87 %93 -%95 = OpImageFetch %8 %85 %94 Lod %90 -OpReturnValue %95 +%85 = OpFunction %10 None %86 +%82 = OpFunctionParameter %12 +%83 = OpFunctionParameter %6 +%84 = OpFunctionParameter %4 +%81 = OpLabel +%87 = OpLoad %13 %31 +OpBranch %88 +%88 = OpLabel +%89 = OpBitcast %4 %83 +%90 = OpCompositeConstruct %15 %82 %89 +%91 = OpImageQueryLevels %4 %87 +%92 = OpISub %4 %91 %59 +%93 = OpExtInst %4 %1 UMin %84 %92 +%94 = OpImageQuerySizeLod %15 %87 %93 +%96 = OpISub %15 %94 %95 +%97 = OpExtInst %15 %1 UMin %90 %96 +%98 = OpImageFetch %10 %87 %97 Lod %93 +OpReturnValue %98 OpFunctionEnd -%99 = OpFunction %8 None %100 -%97 = OpFunctionParameter %13 -%98 = OpFunctionParameter %4 -%96 = OpLabel -%101 = OpLoad %12 %31 -OpBranch %102 -%102 = OpLabel -%103 = OpImageQueryLevels %4 %101 -%104 = OpISub %4 %103 %57 -%105 = OpExtInst %4 %1 UMin %98 %104 -%106 = OpImageQuerySizeLod %13 %101 %105 -%108 = OpISub %13 %106 %107 -%109 = OpExtInst %13 %1 UMin %97 %108 -%110 = OpImageFetch %8 %101 %109 Lod %105 -OpReturnValue %110 +%103 = OpFunction %10 None %104 +%100 = OpFunctionParameter %12 +%101 = OpFunctionParameter %4 +%102 = OpFunctionParameter %4 +%99 = OpLabel +%105 = OpLoad %13 %31 +OpBranch %106 +%106 = OpLabel +%107 = OpCompositeConstruct %15 %100 %101 +%108 = OpImageQueryLevels %4 %105 +%109 = OpISub %4 %108 %59 +%110 = OpExtInst %4 %1 UMin %102 %109 +%111 = OpImageQuerySizeLod %15 %105 %110 +%113 = OpISub %15 %111 %112 +%114 = OpExtInst %15 %1 UMin %107 %113 +%115 = OpImageFetch %10 %105 %114 Lod %110 +OpReturnValue %115 OpFunctionEnd -%114 = OpFunction %8 None %68 -%112 = OpFunctionParameter %10 -%113 = OpFunctionParameter %4 -%111 = OpLabel -%115 = OpLoad %14 %33 -OpBranch %116 +%119 = OpFunction %10 None %120 +%117 = OpFunctionParameter %15 +%118 = OpFunctionParameter %4 %116 = OpLabel -%117 = OpImageQuerySamples %4 %115 -%118 = OpISub %4 %117 %57 -%119 = OpExtInst %4 %1 UMin %113 %118 -%120 = OpImageQuerySize %10 %115 -%122 = OpISub %10 %120 %121 -%123 = OpExtInst %10 %1 UMin %112 %122 -%124 = OpImageFetch %8 %115 %123 Sample %119 -OpReturnValue %124 +%121 = OpLoad %14 %33 +OpBranch %122 +%122 = OpLabel +%123 = OpImageQueryLevels %4 %121 +%124 = OpISub %4 %123 %59 +%125 = OpExtInst %4 %1 UMin %118 %124 +%126 = OpImageQuerySizeLod %15 %121 %125 +%128 = OpISub %15 %126 %127 +%129 = OpExtInst %15 %1 UMin %117 %128 +%130 = OpImageFetch %10 %121 %129 Lod %125 +OpReturnValue %130 OpFunctionEnd -%128 = OpFunction %6 None %129 -%126 = OpFunctionParameter %10 -%127 = OpFunctionParameter %4 -%125 = OpLabel -%130 = OpLoad %15 %35 -OpBranch %131 +%134 = OpFunction %10 None %70 +%132 = OpFunctionParameter %12 +%133 = OpFunctionParameter %4 %131 = OpLabel -%132 = OpImageQueryLevels %4 %130 -%133 = OpISub %4 %132 %57 -%134 = OpExtInst %4 %1 UMin %127 %133 -%135 = OpImageQuerySizeLod %10 %130 %134 -%137 = OpISub %10 %135 %136 -%138 = OpExtInst %10 %1 UMin %126 %137 -%139 = OpImageFetch %8 %130 %138 Lod %134 -%140 = OpCompositeExtract %6 %139 0 -OpReturnValue %140 +%135 = OpLoad %16 %35 +OpBranch %136 +%136 = OpLabel +%137 = OpImageQuerySamples %4 %135 +%138 = OpISub %4 %137 %59 +%139 = OpExtInst %4 %1 UMin %133 %138 +%140 = OpImageQuerySize %12 %135 +%142 = OpISub %12 %140 %141 +%143 = OpExtInst %12 %1 UMin %132 %142 +%144 = OpImageFetch %10 %135 %143 Sample %139 +OpReturnValue %144 OpFunctionEnd -%145 = OpFunction %6 None %146 -%142 = OpFunctionParameter %10 -%143 = OpFunctionParameter %4 -%144 = OpFunctionParameter %4 -%141 = OpLabel -%147 = OpLoad %16 %37 -OpBranch %148 -%148 = OpLabel -%149 = OpCompositeConstruct %13 %142 %143 -%150 = OpImageQueryLevels %4 %147 -%151 = OpISub %4 %150 %57 -%152 = OpExtInst %4 %1 UMin %144 %151 -%153 = OpImageQuerySizeLod %13 %147 %152 -%155 = OpISub %13 %153 %154 -%156 = OpExtInst %13 %1 UMin %149 %155 -%157 = OpImageFetch %8 %147 %156 Lod %152 -%158 = OpCompositeExtract %6 %157 0 -OpReturnValue %158 +%148 = OpFunction %8 None %149 +%146 = OpFunctionParameter %12 +%147 = OpFunctionParameter %4 +%145 = OpLabel +%150 = OpLoad %17 %37 +OpBranch %151 +%151 = OpLabel +%152 = OpImageQueryLevels %4 %150 +%153 = OpISub %4 %152 %59 +%154 = OpExtInst %4 %1 UMin %147 %153 +%155 = OpImageQuerySizeLod %12 %150 %154 +%157 = OpISub %12 %155 %156 +%158 = OpExtInst %12 %1 UMin %146 %157 +%159 = OpImageFetch %10 %150 %158 Lod %154 +%160 = OpCompositeExtract %8 %159 0 +OpReturnValue %160 OpFunctionEnd -%162 = OpFunction %6 None %129 -%160 = OpFunctionParameter %10 -%161 = OpFunctionParameter %4 -%159 = OpLabel -%163 = OpLoad %17 %39 -OpBranch %164 -%164 = OpLabel -%165 = OpImageQuerySamples %4 %163 -%166 = OpISub %4 %165 %57 -%167 = OpExtInst %4 %1 UMin %161 %166 -%168 = OpImageQuerySize %10 %163 -%170 = OpISub %10 %168 %169 -%171 = OpExtInst %10 %1 UMin %160 %170 -%172 = OpImageFetch %8 %163 %171 Sample %167 -%173 = OpCompositeExtract %6 %172 0 -OpReturnValue %173 +%165 = OpFunction %8 None %166 +%162 = OpFunctionParameter %12 +%163 = OpFunctionParameter %6 +%164 = OpFunctionParameter %4 +%161 = OpLabel +%167 = OpLoad %18 %39 +OpBranch %168 +%168 = OpLabel +%169 = OpBitcast %4 %163 +%170 = OpCompositeConstruct %15 %162 %169 +%171 = OpImageQueryLevels %4 %167 +%172 = OpISub %4 %171 %59 +%173 = OpExtInst %4 %1 UMin %164 %172 +%174 = OpImageQuerySizeLod %15 %167 %173 +%176 = OpISub %15 %174 %175 +%177 = OpExtInst %15 %1 UMin %170 %176 +%178 = OpImageFetch %10 %167 %177 Lod %173 +%179 = OpCompositeExtract %8 %178 0 +OpReturnValue %179 OpFunctionEnd -%177 = OpFunction %2 None %178 -%175 = OpFunctionParameter %4 -%176 = OpFunctionParameter %8 -%174 = OpLabel -%179 = OpLoad %18 %41 -OpBranch %180 +%184 = OpFunction %8 None %185 +%181 = OpFunctionParameter %12 +%182 = OpFunctionParameter %4 +%183 = OpFunctionParameter %4 %180 = OpLabel -%181 = OpImageQuerySize %4 %179 -%182 = OpISub %4 %181 %57 -%183 = OpExtInst %4 %1 UMin %175 %182 -OpImageWrite %179 %183 %176 -OpReturn +%186 = OpLoad %18 %39 +OpBranch %187 +%187 = OpLabel +%188 = OpCompositeConstruct %15 %181 %182 +%189 = OpImageQueryLevels %4 %186 +%190 = OpISub %4 %189 %59 +%191 = OpExtInst %4 %1 UMin %183 %190 +%192 = OpImageQuerySizeLod %15 %186 %191 +%194 = OpISub %15 %192 %193 +%195 = OpExtInst %15 %1 UMin %188 %194 +%196 = OpImageFetch %10 %186 %195 Lod %191 +%197 = OpCompositeExtract %8 %196 0 +OpReturnValue %197 OpFunctionEnd -%187 = OpFunction %2 None %188 -%185 = OpFunctionParameter %10 -%186 = OpFunctionParameter %8 -%184 = OpLabel -%189 = OpLoad %19 %43 -OpBranch %190 -%190 = OpLabel -%191 = OpImageQuerySize %10 %189 -%193 = OpISub %10 %191 %192 -%194 = OpExtInst %10 %1 UMin %185 %193 -OpImageWrite %189 %194 %186 -OpReturn +%201 = OpFunction %8 None %149 +%199 = OpFunctionParameter %12 +%200 = OpFunctionParameter %4 +%198 = OpLabel +%202 = OpLoad %19 %41 +OpBranch %203 +%203 = OpLabel +%204 = OpImageQuerySamples %4 %202 +%205 = OpISub %4 %204 %59 +%206 = OpExtInst %4 %1 UMin %200 %205 +%207 = OpImageQuerySize %12 %202 +%209 = OpISub %12 %207 %208 +%210 = OpExtInst %12 %1 UMin %199 %209 +%211 = OpImageFetch %10 %202 %210 Sample %206 +%212 = OpCompositeExtract %8 %211 0 +OpReturnValue %212 OpFunctionEnd -%199 = OpFunction %2 None %200 -%196 = OpFunctionParameter %10 -%197 = OpFunctionParameter %4 -%198 = OpFunctionParameter %8 -%195 = OpLabel -%201 = OpLoad %20 %45 -OpBranch %202 -%202 = OpLabel -%203 = OpCompositeConstruct %13 %196 %197 -%204 = OpImageQuerySize %13 %201 -%206 = OpISub %13 %204 %205 -%207 = OpExtInst %13 %1 UMin %203 %206 -OpImageWrite %201 %207 %198 -OpReturn -OpFunctionEnd -%211 = OpFunction %2 None %212 -%209 = OpFunctionParameter %13 -%210 = OpFunctionParameter %8 -%208 = OpLabel -%213 = OpLoad %21 %47 -OpBranch %214 -%214 = OpLabel -%215 = OpImageQuerySize %13 %213 -%217 = OpISub %13 %215 %216 -%218 = OpExtInst %13 %1 UMin %209 %217 -OpImageWrite %213 %218 %210 -OpReturn -OpFunctionEnd -%222 = OpFunction %2 None %223 +%216 = OpFunction %2 None %217 +%214 = OpFunctionParameter %4 +%215 = OpFunctionParameter %10 +%213 = OpLabel +%218 = OpLoad %20 %43 +OpBranch %219 %219 = OpLabel -%224 = OpLoad %7 %25 -%225 = OpLoad %9 %27 -%226 = OpLoad %11 %29 -%227 = OpLoad %12 %31 -%228 = OpLoad %14 %33 -%229 = OpLoad %18 %41 -%230 = OpLoad %19 %43 -%231 = OpLoad %20 %45 -%232 = OpLoad %21 %47 -OpBranch %233 -%233 = OpLabel -%234 = OpFunctionCall %8 %52 %3 %3 -%235 = OpFunctionCall %8 %67 %22 %3 -%236 = OpFunctionCall %8 %83 %22 %3 %3 -%237 = OpFunctionCall %8 %99 %23 %3 -%238 = OpFunctionCall %8 %114 %22 %3 -%239 = OpFunctionCall %2 %177 %3 %24 -%240 = OpFunctionCall %2 %187 %22 %24 -%241 = OpFunctionCall %2 %199 %22 %3 %24 -%242 = OpFunctionCall %2 %211 %23 %24 -%243 = OpCompositeConstruct %8 %5 %5 %5 %5 -OpStore %220 %243 +%220 = OpImageQuerySize %4 %218 +%221 = OpISub %4 %220 %59 +%222 = OpExtInst %4 %1 UMin %214 %221 +OpImageWrite %218 %222 %215 +OpReturn +OpFunctionEnd +%226 = OpFunction %2 None %227 +%224 = OpFunctionParameter %12 +%225 = OpFunctionParameter %10 +%223 = OpLabel +%228 = OpLoad %21 %45 +OpBranch %229 +%229 = OpLabel +%230 = OpImageQuerySize %12 %228 +%232 = OpISub %12 %230 %231 +%233 = OpExtInst %12 %1 UMin %224 %232 +OpImageWrite %228 %233 %225 +OpReturn +OpFunctionEnd +%238 = OpFunction %2 None %239 +%235 = OpFunctionParameter %12 +%236 = OpFunctionParameter %6 +%237 = OpFunctionParameter %10 +%234 = OpLabel +%240 = OpLoad %22 %47 +OpBranch %241 +%241 = OpLabel +%242 = OpBitcast %4 %236 +%243 = OpCompositeConstruct %15 %235 %242 +%244 = OpImageQuerySize %15 %240 +%246 = OpISub %15 %244 %245 +%247 = OpExtInst %15 %1 UMin %243 %246 +OpImageWrite %240 %247 %237 +OpReturn +OpFunctionEnd +%252 = OpFunction %2 None %253 +%249 = OpFunctionParameter %12 +%250 = OpFunctionParameter %4 +%251 = OpFunctionParameter %10 +%248 = OpLabel +%254 = OpLoad %22 %47 +OpBranch %255 +%255 = OpLabel +%256 = OpCompositeConstruct %15 %249 %250 +%257 = OpImageQuerySize %15 %254 +%259 = OpISub %15 %257 %258 +%260 = OpExtInst %15 %1 UMin %256 %259 +OpImageWrite %254 %260 %251 +OpReturn +OpFunctionEnd +%264 = OpFunction %2 None %265 +%262 = OpFunctionParameter %15 +%263 = OpFunctionParameter %10 +%261 = OpLabel +%266 = OpLoad %23 %49 +OpBranch %267 +%267 = OpLabel +%268 = OpImageQuerySize %15 %266 +%270 = OpISub %15 %268 %269 +%271 = OpExtInst %15 %1 UMin %262 %270 +OpImageWrite %266 %271 %263 +OpReturn +OpFunctionEnd +%275 = OpFunction %2 None %276 +%272 = OpLabel +%277 = OpLoad %9 %27 +%278 = OpLoad %11 %29 +%279 = OpLoad %13 %31 +%280 = OpLoad %14 %33 +%281 = OpLoad %16 %35 +%282 = OpLoad %20 %43 +%283 = OpLoad %21 %45 +%284 = OpLoad %22 %47 +%285 = OpLoad %23 %49 +OpBranch %286 +%286 = OpLabel +%287 = OpFunctionCall %10 %54 %3 %3 +%288 = OpFunctionCall %10 %69 %24 %3 +%289 = OpFunctionCall %10 %85 %24 %5 %3 +%290 = OpFunctionCall %10 %103 %24 %3 %3 +%291 = OpFunctionCall %10 %119 %25 %3 +%292 = OpFunctionCall %10 %134 %24 %3 +%293 = OpFunctionCall %2 %216 %3 %26 +%294 = OpFunctionCall %2 %226 %24 %26 +%295 = OpFunctionCall %2 %238 %24 %5 %26 +%296 = OpFunctionCall %2 %252 %24 %3 %26 +%297 = OpFunctionCall %2 %264 %25 %26 +%298 = OpCompositeConstruct %10 %7 %7 %7 %7 +OpStore %273 %298 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/spv/bounds-check-image-rzsw.spvasm b/tests/out/spv/bounds-check-image-rzsw.spvasm index 3c8e79009a..ce64a3dfc5 100644 --- a/tests/out/spv/bounds-check-image-rzsw.spvasm +++ b/tests/out/spv/bounds-check-image-rzsw.spvasm @@ -1,454 +1,549 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 274 +; Bound: 336 OpCapability ImageQuery OpCapability Image1D OpCapability Shader OpCapability Sampled1D %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Fragment %252 "fragment_shader" %250 -OpExecutionMode %252 OriginUpperLeft +OpEntryPoint Fragment %312 "fragment_shader" %310 +OpExecutionMode %312 OriginUpperLeft OpSource GLSL 450 -OpName %25 "image_1d" -OpName %27 "image_2d" -OpName %29 "image_2d_array" -OpName %31 "image_3d" -OpName %33 "image_multisampled_2d" -OpName %35 "image_depth_2d" -OpName %37 "image_depth_2d_array" -OpName %39 "image_depth_multisampled_2d" -OpName %41 "image_storage_1d" -OpName %43 "image_storage_2d" -OpName %45 "image_storage_2d_array" -OpName %47 "image_storage_3d" -OpName %50 "coords" -OpName %51 "level" -OpName %52 "test_textureLoad_1d" -OpName %68 "coords" -OpName %69 "level" -OpName %70 "test_textureLoad_2d" -OpName %87 "coords" -OpName %88 "index" -OpName %89 "level" -OpName %90 "test_textureLoad_2d_array" -OpName %108 "coords" -OpName %109 "level" -OpName %110 "test_textureLoad_3d" -OpName %126 "coords" -OpName %127 "_sample" -OpName %128 "test_textureLoad_multisampled_2d" -OpName %143 "coords" -OpName %144 "level" -OpName %145 "test_textureLoad_depth_2d" -OpName %162 "coords" -OpName %163 "index" -OpName %164 "level" -OpName %165 "test_textureLoad_depth_2d_array" -OpName %183 "coords" -OpName %184 "_sample" -OpName %185 "test_textureLoad_depth_multisampled_2d" -OpName %201 "coords" -OpName %202 "value" -OpName %203 "test_textureStore_1d" -OpName %212 "coords" -OpName %213 "value" -OpName %214 "test_textureStore_2d" -OpName %224 "coords" -OpName %225 "array_index" -OpName %226 "value" -OpName %227 "test_textureStore_2d_array" -OpName %238 "coords" -OpName %239 "value" -OpName %240 "test_textureStore_3d" -OpName %252 "fragment_shader" -OpDecorate %25 DescriptorSet 0 -OpDecorate %25 Binding 0 +OpName %27 "image_1d" +OpName %29 "image_2d" +OpName %31 "image_2d_array" +OpName %33 "image_3d" +OpName %35 "image_multisampled_2d" +OpName %37 "image_depth_2d" +OpName %39 "image_depth_2d_array" +OpName %41 "image_depth_multisampled_2d" +OpName %43 "image_storage_1d" +OpName %45 "image_storage_2d" +OpName %47 "image_storage_2d_array" +OpName %49 "image_storage_3d" +OpName %52 "coords" +OpName %53 "level" +OpName %54 "test_textureLoad_1d" +OpName %70 "coords" +OpName %71 "level" +OpName %72 "test_textureLoad_2d" +OpName %89 "coords" +OpName %90 "index" +OpName %91 "level" +OpName %92 "test_textureLoad_2d_array_u" +OpName %111 "coords" +OpName %112 "index" +OpName %113 "level" +OpName %114 "test_textureLoad_2d_array_s" +OpName %131 "coords" +OpName %132 "level" +OpName %133 "test_textureLoad_3d" +OpName %149 "coords" +OpName %150 "_sample" +OpName %151 "test_textureLoad_multisampled_2d" +OpName %166 "coords" +OpName %167 "level" +OpName %168 "test_textureLoad_depth_2d" +OpName %185 "coords" +OpName %186 "index" +OpName %187 "level" +OpName %188 "test_textureLoad_depth_2d_array_u" +OpName %207 "coords" +OpName %208 "index" +OpName %209 "level" +OpName %210 "test_textureLoad_depth_2d_array_s" +OpName %228 "coords" +OpName %229 "_sample" +OpName %230 "test_textureLoad_depth_multisampled_2d" +OpName %246 "coords" +OpName %247 "value" +OpName %248 "test_textureStore_1d" +OpName %257 "coords" +OpName %258 "value" +OpName %259 "test_textureStore_2d" +OpName %269 "coords" +OpName %270 "array_index" +OpName %271 "value" +OpName %272 "test_textureStore_2d_array_u" +OpName %284 "coords" +OpName %285 "array_index" +OpName %286 "value" +OpName %287 "test_textureStore_2d_array_s" +OpName %298 "coords" +OpName %299 "value" +OpName %300 "test_textureStore_3d" +OpName %312 "fragment_shader" OpDecorate %27 DescriptorSet 0 -OpDecorate %27 Binding 1 +OpDecorate %27 Binding 0 OpDecorate %29 DescriptorSet 0 -OpDecorate %29 Binding 2 +OpDecorate %29 Binding 1 OpDecorate %31 DescriptorSet 0 -OpDecorate %31 Binding 3 +OpDecorate %31 Binding 2 OpDecorate %33 DescriptorSet 0 -OpDecorate %33 Binding 4 +OpDecorate %33 Binding 3 OpDecorate %35 DescriptorSet 0 -OpDecorate %35 Binding 5 +OpDecorate %35 Binding 4 OpDecorate %37 DescriptorSet 0 -OpDecorate %37 Binding 6 +OpDecorate %37 Binding 5 OpDecorate %39 DescriptorSet 0 -OpDecorate %39 Binding 7 -OpDecorate %41 NonReadable +OpDecorate %39 Binding 6 OpDecorate %41 DescriptorSet 0 -OpDecorate %41 Binding 8 +OpDecorate %41 Binding 7 OpDecorate %43 NonReadable OpDecorate %43 DescriptorSet 0 -OpDecorate %43 Binding 9 +OpDecorate %43 Binding 8 OpDecorate %45 NonReadable OpDecorate %45 DescriptorSet 0 -OpDecorate %45 Binding 10 +OpDecorate %45 Binding 9 OpDecorate %47 NonReadable OpDecorate %47 DescriptorSet 0 -OpDecorate %47 Binding 11 -OpDecorate %250 Location 0 +OpDecorate %47 Binding 10 +OpDecorate %49 NonReadable +OpDecorate %49 DescriptorSet 0 +OpDecorate %49 Binding 11 +OpDecorate %310 Location 0 %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 0 -%6 = OpTypeFloat 32 -%5 = OpConstant %6 0.0 -%7 = OpTypeImage %6 1D 0 0 0 1 Unknown -%8 = OpTypeVector %6 4 -%9 = OpTypeImage %6 2D 0 0 0 1 Unknown -%10 = OpTypeVector %4 2 -%11 = OpTypeImage %6 2D 0 1 0 1 Unknown -%12 = OpTypeImage %6 3D 0 0 0 1 Unknown -%13 = OpTypeVector %4 3 -%14 = OpTypeImage %6 2D 0 0 1 1 Unknown -%15 = OpTypeImage %6 2D 1 0 0 1 Unknown -%16 = OpTypeImage %6 2D 1 1 0 1 Unknown -%17 = OpTypeImage %6 2D 1 0 1 1 Unknown -%18 = OpTypeImage %6 1D 0 0 0 2 Rgba8 -%19 = OpTypeImage %6 2D 0 0 0 2 Rgba8 -%20 = OpTypeImage %6 2D 0 1 0 2 Rgba8 -%21 = OpTypeImage %6 3D 0 0 0 2 Rgba8 -%22 = OpConstantComposite %10 %3 %3 -%23 = OpConstantComposite %13 %3 %3 %3 -%24 = OpConstantComposite %8 %5 %5 %5 %5 -%26 = OpTypePointer UniformConstant %7 -%25 = OpVariable %26 UniformConstant +%6 = OpTypeInt 32 0 +%5 = OpConstant %6 0 +%8 = OpTypeFloat 32 +%7 = OpConstant %8 0.0 +%9 = OpTypeImage %8 1D 0 0 0 1 Unknown +%10 = OpTypeVector %8 4 +%11 = OpTypeImage %8 2D 0 0 0 1 Unknown +%12 = OpTypeVector %4 2 +%13 = OpTypeImage %8 2D 0 1 0 1 Unknown +%14 = OpTypeImage %8 3D 0 0 0 1 Unknown +%15 = OpTypeVector %4 3 +%16 = OpTypeImage %8 2D 0 0 1 1 Unknown +%17 = OpTypeImage %8 2D 1 0 0 1 Unknown +%18 = OpTypeImage %8 2D 1 1 0 1 Unknown +%19 = OpTypeImage %8 2D 1 0 1 1 Unknown +%20 = OpTypeImage %8 1D 0 0 0 2 Rgba8 +%21 = OpTypeImage %8 2D 0 0 0 2 Rgba8 +%22 = OpTypeImage %8 2D 0 1 0 2 Rgba8 +%23 = OpTypeImage %8 3D 0 0 0 2 Rgba8 +%24 = OpConstantComposite %12 %3 %3 +%25 = OpConstantComposite %15 %3 %3 %3 +%26 = OpConstantComposite %10 %7 %7 %7 %7 %28 = OpTypePointer UniformConstant %9 %27 = OpVariable %28 UniformConstant %30 = OpTypePointer UniformConstant %11 %29 = OpVariable %30 UniformConstant -%32 = OpTypePointer UniformConstant %12 +%32 = OpTypePointer UniformConstant %13 %31 = OpVariable %32 UniformConstant %34 = OpTypePointer UniformConstant %14 %33 = OpVariable %34 UniformConstant -%36 = OpTypePointer UniformConstant %15 +%36 = OpTypePointer UniformConstant %16 %35 = OpVariable %36 UniformConstant -%38 = OpTypePointer UniformConstant %16 +%38 = OpTypePointer UniformConstant %17 %37 = OpVariable %38 UniformConstant -%40 = OpTypePointer UniformConstant %17 +%40 = OpTypePointer UniformConstant %18 %39 = OpVariable %40 UniformConstant -%42 = OpTypePointer UniformConstant %18 +%42 = OpTypePointer UniformConstant %19 %41 = OpVariable %42 UniformConstant -%44 = OpTypePointer UniformConstant %19 +%44 = OpTypePointer UniformConstant %20 %43 = OpVariable %44 UniformConstant -%46 = OpTypePointer UniformConstant %20 +%46 = OpTypePointer UniformConstant %21 %45 = OpVariable %46 UniformConstant -%48 = OpTypePointer UniformConstant %21 +%48 = OpTypePointer UniformConstant %22 %47 = OpVariable %48 UniformConstant -%53 = OpTypeFunction %8 %4 %4 -%56 = OpTypeBool -%57 = OpConstantNull %8 -%71 = OpTypeFunction %8 %10 %4 -%74 = OpConstantNull %8 -%80 = OpTypeVector %56 2 -%91 = OpTypeFunction %8 %10 %4 %4 -%95 = OpConstantNull %8 -%101 = OpTypeVector %56 3 -%111 = OpTypeFunction %8 %13 %4 -%114 = OpConstantNull %8 -%131 = OpConstantNull %8 -%146 = OpTypeFunction %6 %10 %4 -%149 = OpConstantNull %8 -%166 = OpTypeFunction %6 %10 %4 %4 -%170 = OpConstantNull %8 -%188 = OpConstantNull %8 -%204 = OpTypeFunction %2 %4 %8 -%215 = OpTypeFunction %2 %10 %8 -%228 = OpTypeFunction %2 %10 %4 %8 -%241 = OpTypeFunction %2 %13 %8 -%251 = OpTypePointer Output %8 -%250 = OpVariable %251 Output -%253 = OpTypeFunction %2 -%52 = OpFunction %8 None %53 -%50 = OpFunctionParameter %4 -%51 = OpFunctionParameter %4 -%49 = OpLabel -%54 = OpLoad %7 %25 -OpBranch %55 -%55 = OpLabel -%58 = OpImageQueryLevels %4 %54 -%59 = OpULessThan %56 %51 %58 -OpSelectionMerge %60 None -OpBranchConditional %59 %61 %60 -%61 = OpLabel -%62 = OpImageQuerySizeLod %4 %54 %51 -%63 = OpULessThan %56 %50 %62 -OpBranchConditional %63 %64 %60 -%64 = OpLabel -%65 = OpImageFetch %8 %54 %50 Lod %51 -OpBranch %60 -%60 = OpLabel -%66 = OpPhi %8 %57 %55 %57 %61 %65 %64 -OpReturnValue %66 +%50 = OpTypePointer UniformConstant %23 +%49 = OpVariable %50 UniformConstant +%55 = OpTypeFunction %10 %4 %4 +%58 = OpTypeBool +%59 = OpConstantNull %10 +%73 = OpTypeFunction %10 %12 %4 +%76 = OpConstantNull %10 +%82 = OpTypeVector %58 2 +%93 = OpTypeFunction %10 %12 %6 %4 +%98 = OpConstantNull %10 +%104 = OpTypeVector %58 3 +%115 = OpTypeFunction %10 %12 %4 %4 +%119 = OpConstantNull %10 +%134 = OpTypeFunction %10 %15 %4 +%137 = OpConstantNull %10 +%154 = OpConstantNull %10 +%169 = OpTypeFunction %8 %12 %4 +%172 = OpConstantNull %10 +%189 = OpTypeFunction %8 %12 %6 %4 +%194 = OpConstantNull %10 +%211 = OpTypeFunction %8 %12 %4 %4 +%215 = OpConstantNull %10 +%233 = OpConstantNull %10 +%249 = OpTypeFunction %2 %4 %10 +%260 = OpTypeFunction %2 %12 %10 +%273 = OpTypeFunction %2 %12 %6 %10 +%288 = OpTypeFunction %2 %12 %4 %10 +%301 = OpTypeFunction %2 %15 %10 +%311 = OpTypePointer Output %10 +%310 = OpVariable %311 Output +%313 = OpTypeFunction %2 +%54 = OpFunction %10 None %55 +%52 = OpFunctionParameter %4 +%53 = OpFunctionParameter %4 +%51 = OpLabel +%56 = OpLoad %9 %27 +OpBranch %57 +%57 = OpLabel +%60 = OpImageQueryLevels %4 %56 +%61 = OpULessThan %58 %53 %60 +OpSelectionMerge %62 None +OpBranchConditional %61 %63 %62 +%63 = OpLabel +%64 = OpImageQuerySizeLod %4 %56 %53 +%65 = OpULessThan %58 %52 %64 +OpBranchConditional %65 %66 %62 +%66 = OpLabel +%67 = OpImageFetch %10 %56 %52 Lod %53 +OpBranch %62 +%62 = OpLabel +%68 = OpPhi %10 %59 %57 %59 %63 %67 %66 +OpReturnValue %68 OpFunctionEnd -%70 = OpFunction %8 None %71 -%68 = OpFunctionParameter %10 -%69 = OpFunctionParameter %4 -%67 = OpLabel -%72 = OpLoad %9 %27 -OpBranch %73 -%73 = OpLabel -%75 = OpImageQueryLevels %4 %72 -%76 = OpULessThan %56 %69 %75 -OpSelectionMerge %77 None -OpBranchConditional %76 %78 %77 -%78 = OpLabel -%79 = OpImageQuerySizeLod %10 %72 %69 -%81 = OpULessThan %80 %68 %79 -%82 = OpAll %56 %81 -OpBranchConditional %82 %83 %77 -%83 = OpLabel -%84 = OpImageFetch %8 %72 %68 Lod %69 -OpBranch %77 -%77 = OpLabel -%85 = OpPhi %8 %74 %73 %74 %78 %84 %83 -OpReturnValue %85 +%72 = OpFunction %10 None %73 +%70 = OpFunctionParameter %12 +%71 = OpFunctionParameter %4 +%69 = OpLabel +%74 = OpLoad %11 %29 +OpBranch %75 +%75 = OpLabel +%77 = OpImageQueryLevels %4 %74 +%78 = OpULessThan %58 %71 %77 +OpSelectionMerge %79 None +OpBranchConditional %78 %80 %79 +%80 = OpLabel +%81 = OpImageQuerySizeLod %12 %74 %71 +%83 = OpULessThan %82 %70 %81 +%84 = OpAll %58 %83 +OpBranchConditional %84 %85 %79 +%85 = OpLabel +%86 = OpImageFetch %10 %74 %70 Lod %71 +OpBranch %79 +%79 = OpLabel +%87 = OpPhi %10 %76 %75 %76 %80 %86 %85 +OpReturnValue %87 OpFunctionEnd -%90 = OpFunction %8 None %91 -%87 = OpFunctionParameter %10 -%88 = OpFunctionParameter %4 -%89 = OpFunctionParameter %4 -%86 = OpLabel -%92 = OpLoad %11 %29 -OpBranch %93 -%93 = OpLabel -%94 = OpCompositeConstruct %13 %87 %88 -%96 = OpImageQueryLevels %4 %92 -%97 = OpULessThan %56 %89 %96 -OpSelectionMerge %98 None -OpBranchConditional %97 %99 %98 -%99 = OpLabel -%100 = OpImageQuerySizeLod %13 %92 %89 -%102 = OpULessThan %101 %94 %100 -%103 = OpAll %56 %102 -OpBranchConditional %103 %104 %98 -%104 = OpLabel -%105 = OpImageFetch %8 %92 %94 Lod %89 -OpBranch %98 -%98 = OpLabel -%106 = OpPhi %8 %95 %93 %95 %99 %105 %104 -OpReturnValue %106 -OpFunctionEnd -%110 = OpFunction %8 None %111 -%108 = OpFunctionParameter %13 -%109 = OpFunctionParameter %4 +%92 = OpFunction %10 None %93 +%89 = OpFunctionParameter %12 +%90 = OpFunctionParameter %6 +%91 = OpFunctionParameter %4 +%88 = OpLabel +%94 = OpLoad %13 %31 +OpBranch %95 +%95 = OpLabel +%96 = OpBitcast %4 %90 +%97 = OpCompositeConstruct %15 %89 %96 +%99 = OpImageQueryLevels %4 %94 +%100 = OpULessThan %58 %91 %99 +OpSelectionMerge %101 None +OpBranchConditional %100 %102 %101 +%102 = OpLabel +%103 = OpImageQuerySizeLod %15 %94 %91 +%105 = OpULessThan %104 %97 %103 +%106 = OpAll %58 %105 +OpBranchConditional %106 %107 %101 %107 = OpLabel -%112 = OpLoad %12 %31 -OpBranch %113 -%113 = OpLabel -%115 = OpImageQueryLevels %4 %112 -%116 = OpULessThan %56 %109 %115 -OpSelectionMerge %117 None -OpBranchConditional %116 %118 %117 -%118 = OpLabel -%119 = OpImageQuerySizeLod %13 %112 %109 -%120 = OpULessThan %101 %108 %119 -%121 = OpAll %56 %120 -OpBranchConditional %121 %122 %117 -%122 = OpLabel -%123 = OpImageFetch %8 %112 %108 Lod %109 +%108 = OpImageFetch %10 %94 %97 Lod %91 +OpBranch %101 +%101 = OpLabel +%109 = OpPhi %10 %98 %95 %98 %102 %108 %107 +OpReturnValue %109 +OpFunctionEnd +%114 = OpFunction %10 None %115 +%111 = OpFunctionParameter %12 +%112 = OpFunctionParameter %4 +%113 = OpFunctionParameter %4 +%110 = OpLabel +%116 = OpLoad %13 %31 OpBranch %117 %117 = OpLabel -%124 = OpPhi %8 %114 %113 %114 %118 %123 %122 -OpReturnValue %124 +%118 = OpCompositeConstruct %15 %111 %112 +%120 = OpImageQueryLevels %4 %116 +%121 = OpULessThan %58 %113 %120 +OpSelectionMerge %122 None +OpBranchConditional %121 %123 %122 +%123 = OpLabel +%124 = OpImageQuerySizeLod %15 %116 %113 +%125 = OpULessThan %104 %118 %124 +%126 = OpAll %58 %125 +OpBranchConditional %126 %127 %122 +%127 = OpLabel +%128 = OpImageFetch %10 %116 %118 Lod %113 +OpBranch %122 +%122 = OpLabel +%129 = OpPhi %10 %119 %117 %119 %123 %128 %127 +OpReturnValue %129 OpFunctionEnd -%128 = OpFunction %8 None %71 -%126 = OpFunctionParameter %10 -%127 = OpFunctionParameter %4 -%125 = OpLabel -%129 = OpLoad %14 %33 -OpBranch %130 +%133 = OpFunction %10 None %134 +%131 = OpFunctionParameter %15 +%132 = OpFunctionParameter %4 %130 = OpLabel -%132 = OpImageQuerySamples %4 %129 -%133 = OpULessThan %56 %127 %132 -OpSelectionMerge %134 None -OpBranchConditional %133 %135 %134 -%135 = OpLabel -%136 = OpImageQuerySize %10 %129 -%137 = OpULessThan %80 %126 %136 -%138 = OpAll %56 %137 -OpBranchConditional %138 %139 %134 -%139 = OpLabel -%140 = OpImageFetch %8 %129 %126 Sample %127 -OpBranch %134 -%134 = OpLabel -%141 = OpPhi %8 %131 %130 %131 %135 %140 %139 -OpReturnValue %141 +%135 = OpLoad %14 %33 +OpBranch %136 +%136 = OpLabel +%138 = OpImageQueryLevels %4 %135 +%139 = OpULessThan %58 %132 %138 +OpSelectionMerge %140 None +OpBranchConditional %139 %141 %140 +%141 = OpLabel +%142 = OpImageQuerySizeLod %15 %135 %132 +%143 = OpULessThan %104 %131 %142 +%144 = OpAll %58 %143 +OpBranchConditional %144 %145 %140 +%145 = OpLabel +%146 = OpImageFetch %10 %135 %131 Lod %132 +OpBranch %140 +%140 = OpLabel +%147 = OpPhi %10 %137 %136 %137 %141 %146 %145 +OpReturnValue %147 OpFunctionEnd -%145 = OpFunction %6 None %146 -%143 = OpFunctionParameter %10 -%144 = OpFunctionParameter %4 -%142 = OpLabel -%147 = OpLoad %15 %35 -OpBranch %148 +%151 = OpFunction %10 None %73 +%149 = OpFunctionParameter %12 +%150 = OpFunctionParameter %4 %148 = OpLabel -%150 = OpImageQueryLevels %4 %147 -%151 = OpULessThan %56 %144 %150 -OpSelectionMerge %152 None -OpBranchConditional %151 %153 %152 +%152 = OpLoad %16 %35 +OpBranch %153 %153 = OpLabel -%154 = OpImageQuerySizeLod %10 %147 %144 -%155 = OpULessThan %80 %143 %154 -%156 = OpAll %56 %155 -OpBranchConditional %156 %157 %152 +%155 = OpImageQuerySamples %4 %152 +%156 = OpULessThan %58 %150 %155 +OpSelectionMerge %157 None +OpBranchConditional %156 %158 %157 +%158 = OpLabel +%159 = OpImageQuerySize %12 %152 +%160 = OpULessThan %82 %149 %159 +%161 = OpAll %58 %160 +OpBranchConditional %161 %162 %157 +%162 = OpLabel +%163 = OpImageFetch %10 %152 %149 Sample %150 +OpBranch %157 %157 = OpLabel -%158 = OpImageFetch %8 %147 %143 Lod %144 -OpBranch %152 -%152 = OpLabel -%159 = OpPhi %8 %149 %148 %149 %153 %158 %157 -%160 = OpCompositeExtract %6 %159 0 -OpReturnValue %160 +%164 = OpPhi %10 %154 %153 %154 %158 %163 %162 +OpReturnValue %164 OpFunctionEnd -%165 = OpFunction %6 None %166 -%162 = OpFunctionParameter %10 -%163 = OpFunctionParameter %4 -%164 = OpFunctionParameter %4 -%161 = OpLabel -%167 = OpLoad %16 %37 -OpBranch %168 -%168 = OpLabel -%169 = OpCompositeConstruct %13 %162 %163 -%171 = OpImageQueryLevels %4 %167 -%172 = OpULessThan %56 %164 %171 -OpSelectionMerge %173 None -OpBranchConditional %172 %174 %173 -%174 = OpLabel -%175 = OpImageQuerySizeLod %13 %167 %164 -%176 = OpULessThan %101 %169 %175 -%177 = OpAll %56 %176 -OpBranchConditional %177 %178 %173 -%178 = OpLabel -%179 = OpImageFetch %8 %167 %169 Lod %164 -OpBranch %173 -%173 = OpLabel -%180 = OpPhi %8 %170 %168 %170 %174 %179 %178 -%181 = OpCompositeExtract %6 %180 0 -OpReturnValue %181 +%168 = OpFunction %8 None %169 +%166 = OpFunctionParameter %12 +%167 = OpFunctionParameter %4 +%165 = OpLabel +%170 = OpLoad %17 %37 +OpBranch %171 +%171 = OpLabel +%173 = OpImageQueryLevels %4 %170 +%174 = OpULessThan %58 %167 %173 +OpSelectionMerge %175 None +OpBranchConditional %174 %176 %175 +%176 = OpLabel +%177 = OpImageQuerySizeLod %12 %170 %167 +%178 = OpULessThan %82 %166 %177 +%179 = OpAll %58 %178 +OpBranchConditional %179 %180 %175 +%180 = OpLabel +%181 = OpImageFetch %10 %170 %166 Lod %167 +OpBranch %175 +%175 = OpLabel +%182 = OpPhi %10 %172 %171 %172 %176 %181 %180 +%183 = OpCompositeExtract %8 %182 0 +OpReturnValue %183 OpFunctionEnd -%185 = OpFunction %6 None %146 -%183 = OpFunctionParameter %10 -%184 = OpFunctionParameter %4 -%182 = OpLabel -%186 = OpLoad %17 %39 -OpBranch %187 -%187 = OpLabel -%189 = OpImageQuerySamples %4 %186 -%190 = OpULessThan %56 %184 %189 -OpSelectionMerge %191 None -OpBranchConditional %190 %192 %191 -%192 = OpLabel -%193 = OpImageQuerySize %10 %186 -%194 = OpULessThan %80 %183 %193 -%195 = OpAll %56 %194 -OpBranchConditional %195 %196 %191 -%196 = OpLabel -%197 = OpImageFetch %8 %186 %183 Sample %184 +%188 = OpFunction %8 None %189 +%185 = OpFunctionParameter %12 +%186 = OpFunctionParameter %6 +%187 = OpFunctionParameter %4 +%184 = OpLabel +%190 = OpLoad %18 %39 OpBranch %191 %191 = OpLabel -%198 = OpPhi %8 %188 %187 %188 %192 %197 %196 -%199 = OpCompositeExtract %6 %198 0 -OpReturnValue %199 +%192 = OpBitcast %4 %186 +%193 = OpCompositeConstruct %15 %185 %192 +%195 = OpImageQueryLevels %4 %190 +%196 = OpULessThan %58 %187 %195 +OpSelectionMerge %197 None +OpBranchConditional %196 %198 %197 +%198 = OpLabel +%199 = OpImageQuerySizeLod %15 %190 %187 +%200 = OpULessThan %104 %193 %199 +%201 = OpAll %58 %200 +OpBranchConditional %201 %202 %197 +%202 = OpLabel +%203 = OpImageFetch %10 %190 %193 Lod %187 +OpBranch %197 +%197 = OpLabel +%204 = OpPhi %10 %194 %191 %194 %198 %203 %202 +%205 = OpCompositeExtract %8 %204 0 +OpReturnValue %205 OpFunctionEnd -%203 = OpFunction %2 None %204 -%201 = OpFunctionParameter %4 -%202 = OpFunctionParameter %8 -%200 = OpLabel -%205 = OpLoad %18 %41 -OpBranch %206 +%210 = OpFunction %8 None %211 +%207 = OpFunctionParameter %12 +%208 = OpFunctionParameter %4 +%209 = OpFunctionParameter %4 %206 = OpLabel -%207 = OpImageQuerySize %4 %205 -%208 = OpULessThan %56 %201 %207 -OpSelectionMerge %209 None -OpBranchConditional %208 %210 %209 -%210 = OpLabel -OpImageWrite %205 %201 %202 -OpBranch %209 -%209 = OpLabel -OpReturn -OpFunctionEnd -%214 = OpFunction %2 None %215 -%212 = OpFunctionParameter %10 -%213 = OpFunctionParameter %8 -%211 = OpLabel -%216 = OpLoad %19 %43 -OpBranch %217 -%217 = OpLabel -%218 = OpImageQuerySize %10 %216 -%219 = OpULessThan %80 %212 %218 -%220 = OpAll %56 %219 -OpSelectionMerge %221 None -OpBranchConditional %220 %222 %221 -%222 = OpLabel -OpImageWrite %216 %212 %213 -OpBranch %221 -%221 = OpLabel -OpReturn -OpFunctionEnd -%227 = OpFunction %2 None %228 -%224 = OpFunctionParameter %10 -%225 = OpFunctionParameter %4 -%226 = OpFunctionParameter %8 +%212 = OpLoad %18 %39 +OpBranch %213 +%213 = OpLabel +%214 = OpCompositeConstruct %15 %207 %208 +%216 = OpImageQueryLevels %4 %212 +%217 = OpULessThan %58 %209 %216 +OpSelectionMerge %218 None +OpBranchConditional %217 %219 %218 +%219 = OpLabel +%220 = OpImageQuerySizeLod %15 %212 %209 +%221 = OpULessThan %104 %214 %220 +%222 = OpAll %58 %221 +OpBranchConditional %222 %223 %218 %223 = OpLabel -%229 = OpLoad %20 %45 -OpBranch %230 -%230 = OpLabel -%231 = OpCompositeConstruct %13 %224 %225 -%232 = OpImageQuerySize %13 %229 -%233 = OpULessThan %101 %231 %232 -%234 = OpAll %56 %233 -OpSelectionMerge %235 None -OpBranchConditional %234 %236 %235 -%236 = OpLabel -OpImageWrite %229 %231 %226 -OpBranch %235 -%235 = OpLabel -OpReturn +%224 = OpImageFetch %10 %212 %214 Lod %209 +OpBranch %218 +%218 = OpLabel +%225 = OpPhi %10 %215 %213 %215 %219 %224 %223 +%226 = OpCompositeExtract %8 %225 0 +OpReturnValue %226 OpFunctionEnd -%240 = OpFunction %2 None %241 -%238 = OpFunctionParameter %13 -%239 = OpFunctionParameter %8 +%230 = OpFunction %8 None %169 +%228 = OpFunctionParameter %12 +%229 = OpFunctionParameter %4 +%227 = OpLabel +%231 = OpLoad %19 %41 +OpBranch %232 +%232 = OpLabel +%234 = OpImageQuerySamples %4 %231 +%235 = OpULessThan %58 %229 %234 +OpSelectionMerge %236 None +OpBranchConditional %235 %237 %236 %237 = OpLabel -%242 = OpLoad %21 %47 -OpBranch %243 -%243 = OpLabel -%244 = OpImageQuerySize %13 %242 -%245 = OpULessThan %101 %238 %244 -%246 = OpAll %56 %245 -OpSelectionMerge %247 None -OpBranchConditional %246 %248 %247 -%248 = OpLabel -OpImageWrite %242 %238 %239 -OpBranch %247 -%247 = OpLabel +%238 = OpImageQuerySize %12 %231 +%239 = OpULessThan %82 %228 %238 +%240 = OpAll %58 %239 +OpBranchConditional %240 %241 %236 +%241 = OpLabel +%242 = OpImageFetch %10 %231 %228 Sample %229 +OpBranch %236 +%236 = OpLabel +%243 = OpPhi %10 %233 %232 %233 %237 %242 %241 +%244 = OpCompositeExtract %8 %243 0 +OpReturnValue %244 +OpFunctionEnd +%248 = OpFunction %2 None %249 +%246 = OpFunctionParameter %4 +%247 = OpFunctionParameter %10 +%245 = OpLabel +%250 = OpLoad %20 %43 +OpBranch %251 +%251 = OpLabel +%252 = OpImageQuerySize %4 %250 +%253 = OpULessThan %58 %246 %252 +OpSelectionMerge %254 None +OpBranchConditional %253 %255 %254 +%255 = OpLabel +OpImageWrite %250 %246 %247 +OpBranch %254 +%254 = OpLabel OpReturn OpFunctionEnd -%252 = OpFunction %2 None %253 -%249 = OpLabel -%254 = OpLoad %7 %25 -%255 = OpLoad %9 %27 -%256 = OpLoad %11 %29 -%257 = OpLoad %12 %31 -%258 = OpLoad %14 %33 -%259 = OpLoad %18 %41 -%260 = OpLoad %19 %43 -%261 = OpLoad %20 %45 -%262 = OpLoad %21 %47 -OpBranch %263 -%263 = OpLabel -%264 = OpFunctionCall %8 %52 %3 %3 -%265 = OpFunctionCall %8 %70 %22 %3 -%266 = OpFunctionCall %8 %90 %22 %3 %3 -%267 = OpFunctionCall %8 %110 %23 %3 -%268 = OpFunctionCall %8 %128 %22 %3 -%269 = OpFunctionCall %2 %203 %3 %24 -%270 = OpFunctionCall %2 %214 %22 %24 -%271 = OpFunctionCall %2 %227 %22 %3 %24 -%272 = OpFunctionCall %2 %240 %23 %24 -%273 = OpCompositeConstruct %8 %5 %5 %5 %5 -OpStore %250 %273 +%259 = OpFunction %2 None %260 +%257 = OpFunctionParameter %12 +%258 = OpFunctionParameter %10 +%256 = OpLabel +%261 = OpLoad %21 %45 +OpBranch %262 +%262 = OpLabel +%263 = OpImageQuerySize %12 %261 +%264 = OpULessThan %82 %257 %263 +%265 = OpAll %58 %264 +OpSelectionMerge %266 None +OpBranchConditional %265 %267 %266 +%267 = OpLabel +OpImageWrite %261 %257 %258 +OpBranch %266 +%266 = OpLabel +OpReturn +OpFunctionEnd +%272 = OpFunction %2 None %273 +%269 = OpFunctionParameter %12 +%270 = OpFunctionParameter %6 +%271 = OpFunctionParameter %10 +%268 = OpLabel +%274 = OpLoad %22 %47 +OpBranch %275 +%275 = OpLabel +%276 = OpBitcast %4 %270 +%277 = OpCompositeConstruct %15 %269 %276 +%278 = OpImageQuerySize %15 %274 +%279 = OpULessThan %104 %277 %278 +%280 = OpAll %58 %279 +OpSelectionMerge %281 None +OpBranchConditional %280 %282 %281 +%282 = OpLabel +OpImageWrite %274 %277 %271 +OpBranch %281 +%281 = OpLabel +OpReturn +OpFunctionEnd +%287 = OpFunction %2 None %288 +%284 = OpFunctionParameter %12 +%285 = OpFunctionParameter %4 +%286 = OpFunctionParameter %10 +%283 = OpLabel +%289 = OpLoad %22 %47 +OpBranch %290 +%290 = OpLabel +%291 = OpCompositeConstruct %15 %284 %285 +%292 = OpImageQuerySize %15 %289 +%293 = OpULessThan %104 %291 %292 +%294 = OpAll %58 %293 +OpSelectionMerge %295 None +OpBranchConditional %294 %296 %295 +%296 = OpLabel +OpImageWrite %289 %291 %286 +OpBranch %295 +%295 = OpLabel +OpReturn +OpFunctionEnd +%300 = OpFunction %2 None %301 +%298 = OpFunctionParameter %15 +%299 = OpFunctionParameter %10 +%297 = OpLabel +%302 = OpLoad %23 %49 +OpBranch %303 +%303 = OpLabel +%304 = OpImageQuerySize %15 %302 +%305 = OpULessThan %104 %298 %304 +%306 = OpAll %58 %305 +OpSelectionMerge %307 None +OpBranchConditional %306 %308 %307 +%308 = OpLabel +OpImageWrite %302 %298 %299 +OpBranch %307 +%307 = OpLabel +OpReturn +OpFunctionEnd +%312 = OpFunction %2 None %313 +%309 = OpLabel +%314 = OpLoad %9 %27 +%315 = OpLoad %11 %29 +%316 = OpLoad %13 %31 +%317 = OpLoad %14 %33 +%318 = OpLoad %16 %35 +%319 = OpLoad %20 %43 +%320 = OpLoad %21 %45 +%321 = OpLoad %22 %47 +%322 = OpLoad %23 %49 +OpBranch %323 +%323 = OpLabel +%324 = OpFunctionCall %10 %54 %3 %3 +%325 = OpFunctionCall %10 %72 %24 %3 +%326 = OpFunctionCall %10 %92 %24 %5 %3 +%327 = OpFunctionCall %10 %114 %24 %3 %3 +%328 = OpFunctionCall %10 %133 %25 %3 +%329 = OpFunctionCall %10 %151 %24 %3 +%330 = OpFunctionCall %2 %248 %3 %26 +%331 = OpFunctionCall %2 %259 %24 %26 +%332 = OpFunctionCall %2 %272 %24 %5 %26 +%333 = OpFunctionCall %2 %287 %24 %3 %26 +%334 = OpFunctionCall %2 %300 %25 %26 +%335 = OpCompositeConstruct %10 %7 %7 %7 %7 +OpStore %310 %335 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/spv/image.spvasm b/tests/out/spv/image.spvasm index 7efb4483e0..040ded2af0 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: 376 +; Bound: 546 OpCapability SampledCubeArray OpCapability ImageQuery OpCapability Image1D @@ -9,105 +9,110 @@ OpCapability Shader OpCapability Sampled1D %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %84 "main" %81 -OpEntryPoint GLCompute %158 "depth_load" %156 -OpEntryPoint Vertex %180 "queries" %178 -OpEntryPoint Vertex %245 "levels_queries" %244 -OpEntryPoint Fragment %282 "texture_sample" %281 -OpEntryPoint Fragment %310 "texture_sample_comparison" %308 -OpEntryPoint Fragment %329 "gather" %328 -OpEntryPoint Fragment %364 "depth_no_comparison" %363 -OpExecutionMode %84 LocalSize 16 1 1 -OpExecutionMode %158 LocalSize 16 1 1 -OpExecutionMode %282 OriginUpperLeft -OpExecutionMode %310 OriginUpperLeft -OpExecutionMode %329 OriginUpperLeft -OpExecutionMode %364 OriginUpperLeft +OpEntryPoint GLCompute %89 "main" %86 +OpEntryPoint GLCompute %177 "depth_load" %175 +OpEntryPoint Vertex %199 "queries" %197 +OpEntryPoint Vertex %263 "levels_queries" %262 +OpEntryPoint Fragment %303 "texture_sample" %302 +OpEntryPoint Fragment %444 "texture_sample_comparison" %442 +OpEntryPoint Fragment %499 "gather" %498 +OpEntryPoint Fragment %534 "depth_no_comparison" %533 +OpExecutionMode %89 LocalSize 16 1 1 +OpExecutionMode %177 LocalSize 16 1 1 +OpExecutionMode %303 OriginUpperLeft +OpExecutionMode %444 OriginUpperLeft +OpExecutionMode %499 OriginUpperLeft +OpExecutionMode %534 OriginUpperLeft OpSource GLSL 450 -OpName %39 "image_mipmapped_src" -OpName %41 "image_multisampled_src" -OpName %43 "image_depth_multisampled_src" -OpName %45 "image_storage_src" -OpName %47 "image_array_src" -OpName %49 "image_dup_src" -OpName %51 "image_1d_src" -OpName %53 "image_dst" -OpName %55 "image_1d" -OpName %57 "image_2d" -OpName %59 "image_2d_u32" -OpName %60 "image_2d_i32" -OpName %62 "image_2d_array" -OpName %64 "image_cube" -OpName %66 "image_cube_array" -OpName %68 "image_3d" -OpName %70 "image_aa" -OpName %72 "sampler_reg" -OpName %74 "sampler_cmp" -OpName %76 "image_2d_depth" -OpName %78 "image_cube_depth" -OpName %81 "local_id" -OpName %84 "main" -OpName %156 "local_id" -OpName %158 "depth_load" -OpName %180 "queries" -OpName %245 "levels_queries" -OpName %282 "texture_sample" -OpName %310 "texture_sample_comparison" -OpName %329 "gather" -OpName %364 "depth_no_comparison" -OpDecorate %39 DescriptorSet 0 -OpDecorate %39 Binding 0 -OpDecorate %41 DescriptorSet 0 -OpDecorate %41 Binding 3 -OpDecorate %43 DescriptorSet 0 -OpDecorate %43 Binding 4 -OpDecorate %45 NonWritable -OpDecorate %45 DescriptorSet 0 -OpDecorate %45 Binding 1 -OpDecorate %47 DescriptorSet 0 -OpDecorate %47 Binding 5 -OpDecorate %49 NonWritable -OpDecorate %49 DescriptorSet 0 -OpDecorate %49 Binding 6 -OpDecorate %51 DescriptorSet 0 -OpDecorate %51 Binding 7 -OpDecorate %53 NonReadable -OpDecorate %53 DescriptorSet 0 -OpDecorate %53 Binding 2 -OpDecorate %55 DescriptorSet 0 -OpDecorate %55 Binding 0 -OpDecorate %57 DescriptorSet 0 -OpDecorate %57 Binding 1 -OpDecorate %59 DescriptorSet 0 -OpDecorate %59 Binding 2 +OpName %42 "image_mipmapped_src" +OpName %44 "image_multisampled_src" +OpName %46 "image_depth_multisampled_src" +OpName %48 "image_storage_src" +OpName %50 "image_array_src" +OpName %52 "image_dup_src" +OpName %54 "image_1d_src" +OpName %56 "image_dst" +OpName %58 "image_1d" +OpName %60 "image_2d" +OpName %62 "image_2d_u32" +OpName %63 "image_2d_i32" +OpName %65 "image_2d_array" +OpName %67 "image_cube" +OpName %69 "image_cube_array" +OpName %71 "image_3d" +OpName %73 "image_aa" +OpName %75 "sampler_reg" +OpName %77 "sampler_cmp" +OpName %79 "image_2d_depth" +OpName %81 "image_2d_array_depth" +OpName %83 "image_cube_depth" +OpName %86 "local_id" +OpName %89 "main" +OpName %175 "local_id" +OpName %177 "depth_load" +OpName %199 "queries" +OpName %263 "levels_queries" +OpName %298 "a" +OpName %303 "texture_sample" +OpName %438 "a" +OpName %444 "texture_sample_comparison" +OpName %499 "gather" +OpName %534 "depth_no_comparison" +OpDecorate %42 DescriptorSet 0 +OpDecorate %42 Binding 0 +OpDecorate %44 DescriptorSet 0 +OpDecorate %44 Binding 3 +OpDecorate %46 DescriptorSet 0 +OpDecorate %46 Binding 4 +OpDecorate %48 NonWritable +OpDecorate %48 DescriptorSet 0 +OpDecorate %48 Binding 1 +OpDecorate %50 DescriptorSet 0 +OpDecorate %50 Binding 5 +OpDecorate %52 NonWritable +OpDecorate %52 DescriptorSet 0 +OpDecorate %52 Binding 6 +OpDecorate %54 DescriptorSet 0 +OpDecorate %54 Binding 7 +OpDecorate %56 NonReadable +OpDecorate %56 DescriptorSet 0 +OpDecorate %56 Binding 2 +OpDecorate %58 DescriptorSet 0 +OpDecorate %58 Binding 0 OpDecorate %60 DescriptorSet 0 -OpDecorate %60 Binding 3 +OpDecorate %60 Binding 1 OpDecorate %62 DescriptorSet 0 -OpDecorate %62 Binding 4 -OpDecorate %64 DescriptorSet 0 -OpDecorate %64 Binding 5 -OpDecorate %66 DescriptorSet 0 -OpDecorate %66 Binding 6 -OpDecorate %68 DescriptorSet 0 -OpDecorate %68 Binding 7 -OpDecorate %70 DescriptorSet 0 -OpDecorate %70 Binding 8 -OpDecorate %72 DescriptorSet 1 -OpDecorate %72 Binding 0 -OpDecorate %74 DescriptorSet 1 -OpDecorate %74 Binding 1 -OpDecorate %76 DescriptorSet 1 -OpDecorate %76 Binding 2 -OpDecorate %78 DescriptorSet 1 -OpDecorate %78 Binding 3 -OpDecorate %81 BuiltIn LocalInvocationId -OpDecorate %156 BuiltIn LocalInvocationId -OpDecorate %178 BuiltIn Position -OpDecorate %244 BuiltIn Position -OpDecorate %281 Location 0 -OpDecorate %308 Location 0 -OpDecorate %328 Location 0 -OpDecorate %363 Location 0 +OpDecorate %62 Binding 2 +OpDecorate %63 DescriptorSet 0 +OpDecorate %63 Binding 3 +OpDecorate %65 DescriptorSet 0 +OpDecorate %65 Binding 4 +OpDecorate %67 DescriptorSet 0 +OpDecorate %67 Binding 5 +OpDecorate %69 DescriptorSet 0 +OpDecorate %69 Binding 6 +OpDecorate %71 DescriptorSet 0 +OpDecorate %71 Binding 7 +OpDecorate %73 DescriptorSet 0 +OpDecorate %73 Binding 8 +OpDecorate %75 DescriptorSet 1 +OpDecorate %75 Binding 0 +OpDecorate %77 DescriptorSet 1 +OpDecorate %77 Binding 1 +OpDecorate %79 DescriptorSet 1 +OpDecorate %79 Binding 2 +OpDecorate %81 DescriptorSet 1 +OpDecorate %81 Binding 3 +OpDecorate %83 DescriptorSet 1 +OpDecorate %83 Binding 4 +OpDecorate %86 BuiltIn LocalInvocationId +OpDecorate %175 BuiltIn LocalInvocationId +OpDecorate %197 BuiltIn Position +OpDecorate %262 BuiltIn Position +OpDecorate %302 Location 0 +OpDecorate %442 Location 0 +OpDecorate %498 Location 0 +OpDecorate %533 Location 0 %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 10 @@ -119,398 +124,598 @@ OpDecorate %363 Location 0 %10 = OpConstant %4 3 %11 = OpConstant %8 2.0 %13 = OpTypeInt 32 0 -%12 = OpTypeImage %13 2D 0 0 0 1 Unknown -%14 = OpTypeImage %13 2D 0 0 1 1 Unknown -%15 = OpTypeImage %8 2D 1 0 1 1 Unknown -%16 = OpTypeImage %13 2D 0 0 0 2 Rgba8ui -%17 = OpTypeImage %13 2D 0 1 0 1 Unknown -%18 = OpTypeImage %13 1D 0 0 0 2 R32ui -%19 = OpTypeImage %13 1D 0 0 0 1 Unknown -%20 = OpTypeVector %13 3 -%21 = OpTypeVector %4 2 -%22 = OpTypeVector %13 2 -%23 = OpTypeVector %13 4 -%24 = OpTypeImage %8 1D 0 0 0 1 Unknown -%25 = OpTypeImage %8 2D 0 0 0 1 Unknown -%26 = OpTypeImage %4 2D 0 0 0 1 Unknown -%27 = OpTypeImage %8 2D 0 1 0 1 Unknown -%28 = OpTypeImage %8 Cube 0 0 0 1 Unknown -%29 = OpTypeImage %8 Cube 0 1 0 1 Unknown -%30 = OpTypeImage %8 3D 0 0 0 1 Unknown -%31 = OpTypeImage %8 2D 0 0 1 1 Unknown -%32 = OpTypeVector %8 4 -%33 = OpTypeSampler -%34 = OpTypeVector %8 2 -%35 = OpTypeImage %8 2D 1 0 0 1 Unknown -%36 = OpTypeImage %8 Cube 1 0 0 1 Unknown +%12 = OpConstant %13 0 +%14 = OpConstant %4 0 +%15 = OpTypeImage %13 2D 0 0 0 1 Unknown +%16 = OpTypeImage %13 2D 0 0 1 1 Unknown +%17 = OpTypeImage %8 2D 1 0 1 1 Unknown +%18 = OpTypeImage %13 2D 0 0 0 2 Rgba8ui +%19 = OpTypeImage %13 2D 0 1 0 1 Unknown +%20 = OpTypeImage %13 1D 0 0 0 2 R32ui +%21 = OpTypeImage %13 1D 0 0 0 1 Unknown +%22 = OpTypeVector %13 3 +%23 = OpTypeVector %4 2 +%24 = OpTypeVector %13 2 +%25 = OpTypeVector %13 4 +%26 = OpTypeImage %8 1D 0 0 0 1 Unknown +%27 = OpTypeImage %8 2D 0 0 0 1 Unknown +%28 = OpTypeImage %4 2D 0 0 0 1 Unknown +%29 = OpTypeImage %8 2D 0 1 0 1 Unknown +%30 = OpTypeImage %8 Cube 0 0 0 1 Unknown +%31 = OpTypeImage %8 Cube 0 1 0 1 Unknown +%32 = OpTypeImage %8 3D 0 0 0 1 Unknown +%33 = OpTypeImage %8 2D 0 0 1 1 Unknown +%34 = OpTypeVector %8 4 +%35 = OpTypeSampler +%36 = OpTypeVector %8 2 %37 = OpTypeVector %8 3 -%38 = OpConstantComposite %21 %10 %6 -%40 = OpTypePointer UniformConstant %12 -%39 = OpVariable %40 UniformConstant -%42 = OpTypePointer UniformConstant %14 -%41 = OpVariable %42 UniformConstant -%44 = OpTypePointer UniformConstant %15 -%43 = OpVariable %44 UniformConstant -%46 = OpTypePointer UniformConstant %16 -%45 = OpVariable %46 UniformConstant -%48 = OpTypePointer UniformConstant %17 -%47 = OpVariable %48 UniformConstant -%50 = OpTypePointer UniformConstant %18 -%49 = OpVariable %50 UniformConstant -%52 = OpTypePointer UniformConstant %19 -%51 = OpVariable %52 UniformConstant -%54 = OpTypePointer UniformConstant %18 -%53 = OpVariable %54 UniformConstant -%56 = OpTypePointer UniformConstant %24 -%55 = OpVariable %56 UniformConstant -%58 = OpTypePointer UniformConstant %25 -%57 = OpVariable %58 UniformConstant -%59 = OpVariable %40 UniformConstant -%61 = OpTypePointer UniformConstant %26 +%38 = OpTypeImage %8 2D 1 0 0 1 Unknown +%39 = OpTypeImage %8 2D 1 1 0 1 Unknown +%40 = OpTypeImage %8 Cube 1 0 0 1 Unknown +%41 = OpConstantComposite %23 %10 %6 +%43 = OpTypePointer UniformConstant %15 +%42 = OpVariable %43 UniformConstant +%45 = OpTypePointer UniformConstant %16 +%44 = OpVariable %45 UniformConstant +%47 = OpTypePointer UniformConstant %17 +%46 = OpVariable %47 UniformConstant +%49 = OpTypePointer UniformConstant %18 +%48 = OpVariable %49 UniformConstant +%51 = OpTypePointer UniformConstant %19 +%50 = OpVariable %51 UniformConstant +%53 = OpTypePointer UniformConstant %20 +%52 = OpVariable %53 UniformConstant +%55 = OpTypePointer UniformConstant %21 +%54 = OpVariable %55 UniformConstant +%57 = OpTypePointer UniformConstant %20 +%56 = OpVariable %57 UniformConstant +%59 = OpTypePointer UniformConstant %26 +%58 = OpVariable %59 UniformConstant +%61 = OpTypePointer UniformConstant %27 %60 = OpVariable %61 UniformConstant -%63 = OpTypePointer UniformConstant %27 -%62 = OpVariable %63 UniformConstant -%65 = OpTypePointer UniformConstant %28 -%64 = OpVariable %65 UniformConstant -%67 = OpTypePointer UniformConstant %29 -%66 = OpVariable %67 UniformConstant -%69 = OpTypePointer UniformConstant %30 -%68 = OpVariable %69 UniformConstant -%71 = OpTypePointer UniformConstant %31 -%70 = OpVariable %71 UniformConstant -%73 = OpTypePointer UniformConstant %33 -%72 = OpVariable %73 UniformConstant -%75 = OpTypePointer UniformConstant %33 -%74 = OpVariable %75 UniformConstant -%77 = OpTypePointer UniformConstant %35 -%76 = OpVariable %77 UniformConstant -%79 = OpTypePointer UniformConstant %36 -%78 = OpVariable %79 UniformConstant -%82 = OpTypePointer Input %20 -%81 = OpVariable %82 Input -%85 = OpTypeFunction %2 -%112 = OpTypeVector %4 3 -%156 = OpVariable %82 Input -%179 = OpTypePointer Output %32 -%178 = OpVariable %179 Output -%189 = OpConstant %13 0 -%244 = OpVariable %179 Output -%281 = OpVariable %179 Output -%289 = OpTypeSampledImage %24 -%292 = OpTypeSampledImage %25 -%309 = OpTypePointer Output %8 -%308 = OpVariable %309 Output -%316 = OpTypeSampledImage %35 -%321 = OpConstant %8 0.0 -%323 = OpTypeSampledImage %36 -%328 = OpVariable %179 Output -%340 = OpConstant %13 1 -%343 = OpConstant %13 3 -%348 = OpTypeSampledImage %12 -%351 = OpTypeVector %4 4 -%352 = OpTypeSampledImage %26 -%363 = OpVariable %179 Output -%84 = OpFunction %2 None %85 -%80 = OpLabel -%83 = OpLoad %20 %81 -%86 = OpLoad %12 %39 -%87 = OpLoad %14 %41 -%88 = OpLoad %16 %45 -%89 = OpLoad %17 %47 -%90 = OpLoad %19 %51 -%91 = OpLoad %18 %53 -OpBranch %92 -%92 = OpLabel -%93 = OpImageQuerySize %21 %88 -%94 = OpBitcast %22 %93 -%95 = OpVectorShuffle %22 %83 %83 0 1 -%96 = OpIMul %22 %94 %95 -%97 = OpBitcast %21 %96 -%98 = OpCompositeConstruct %21 %3 %5 -%99 = OpSRem %21 %97 %98 -%100 = OpCompositeExtract %13 %83 2 -%101 = OpBitcast %4 %100 -%102 = OpImageFetch %23 %86 %99 Lod %101 -%103 = OpCompositeExtract %13 %83 2 -%104 = OpBitcast %4 %103 -%105 = OpImageFetch %23 %87 %99 Sample %104 -%106 = OpImageRead %23 %88 %99 -%107 = OpCompositeExtract %13 %83 2 -%108 = OpBitcast %4 %107 -%109 = OpCompositeExtract %13 %83 2 -%110 = OpBitcast %4 %109 -%111 = OpIAdd %4 %110 %6 -%113 = OpCompositeConstruct %112 %99 %108 -%114 = OpImageFetch %23 %89 %113 Lod %111 -%115 = OpCompositeExtract %13 %83 0 -%116 = OpBitcast %4 %115 -%117 = OpCompositeExtract %13 %83 2 -%118 = OpBitcast %4 %117 -%119 = OpImageFetch %23 %90 %116 Lod %118 -%120 = OpBitcast %22 %99 -%121 = OpCompositeExtract %13 %83 2 -%122 = OpBitcast %4 %121 -%123 = OpImageFetch %23 %86 %120 Lod %122 -%124 = OpBitcast %22 %99 -%125 = OpCompositeExtract %13 %83 2 -%126 = OpBitcast %4 %125 -%127 = OpImageFetch %23 %87 %124 Sample %126 -%128 = OpBitcast %22 %99 -%129 = OpImageRead %23 %88 %128 -%130 = OpBitcast %22 %99 -%131 = OpCompositeExtract %13 %83 2 -%132 = OpBitcast %4 %131 -%133 = OpCompositeExtract %13 %83 2 +%62 = OpVariable %43 UniformConstant +%64 = OpTypePointer UniformConstant %28 +%63 = OpVariable %64 UniformConstant +%66 = OpTypePointer UniformConstant %29 +%65 = OpVariable %66 UniformConstant +%68 = OpTypePointer UniformConstant %30 +%67 = OpVariable %68 UniformConstant +%70 = OpTypePointer UniformConstant %31 +%69 = OpVariable %70 UniformConstant +%72 = OpTypePointer UniformConstant %32 +%71 = OpVariable %72 UniformConstant +%74 = OpTypePointer UniformConstant %33 +%73 = OpVariable %74 UniformConstant +%76 = OpTypePointer UniformConstant %35 +%75 = OpVariable %76 UniformConstant +%78 = OpTypePointer UniformConstant %35 +%77 = OpVariable %78 UniformConstant +%80 = OpTypePointer UniformConstant %38 +%79 = OpVariable %80 UniformConstant +%82 = OpTypePointer UniformConstant %39 +%81 = OpVariable %82 UniformConstant +%84 = OpTypePointer UniformConstant %40 +%83 = OpVariable %84 UniformConstant +%87 = OpTypePointer Input %22 +%86 = OpVariable %87 Input +%90 = OpTypeFunction %2 +%117 = OpTypeVector %4 3 +%175 = OpVariable %87 Input +%198 = OpTypePointer Output %34 +%197 = OpVariable %198 Output +%262 = OpVariable %198 Output +%299 = OpTypePointer Function %34 +%300 = OpConstantNull %34 +%302 = OpVariable %198 Output +%313 = OpTypeSampledImage %26 +%318 = OpTypeSampledImage %27 +%339 = OpTypeSampledImage %29 +%400 = OpTypeSampledImage %31 +%439 = OpTypePointer Function %8 +%440 = OpConstantNull %8 +%443 = OpTypePointer Output %8 +%442 = OpVariable %443 Output +%452 = OpTypeSampledImage %38 +%457 = OpTypeSampledImage %39 +%470 = OpTypeSampledImage %40 +%477 = OpConstant %8 0.0 +%498 = OpVariable %198 Output +%510 = OpConstant %13 1 +%513 = OpConstant %13 3 +%518 = OpTypeSampledImage %15 +%521 = OpTypeVector %4 4 +%522 = OpTypeSampledImage %28 +%533 = OpVariable %198 Output +%89 = OpFunction %2 None %90 +%85 = OpLabel +%88 = OpLoad %22 %86 +%91 = OpLoad %15 %42 +%92 = OpLoad %16 %44 +%93 = OpLoad %18 %48 +%94 = OpLoad %19 %50 +%95 = OpLoad %21 %54 +%96 = OpLoad %20 %56 +OpBranch %97 +%97 = OpLabel +%98 = OpImageQuerySize %23 %93 +%99 = OpBitcast %24 %98 +%100 = OpVectorShuffle %24 %88 %88 0 1 +%101 = OpIMul %24 %99 %100 +%102 = OpBitcast %23 %101 +%103 = OpCompositeConstruct %23 %3 %5 +%104 = OpSRem %23 %102 %103 +%105 = OpCompositeExtract %13 %88 2 +%106 = OpBitcast %4 %105 +%107 = OpImageFetch %25 %91 %104 Lod %106 +%108 = OpCompositeExtract %13 %88 2 +%109 = OpBitcast %4 %108 +%110 = OpImageFetch %25 %92 %104 Sample %109 +%111 = OpImageRead %25 %93 %104 +%112 = OpCompositeExtract %13 %88 2 +%113 = OpCompositeExtract %13 %88 2 +%114 = OpBitcast %4 %113 +%115 = OpIAdd %4 %114 %6 +%116 = OpBitcast %4 %112 +%118 = OpCompositeConstruct %117 %104 %116 +%119 = OpImageFetch %25 %94 %118 Lod %115 +%120 = OpCompositeExtract %13 %88 2 +%121 = OpBitcast %4 %120 +%122 = OpCompositeExtract %13 %88 2 +%123 = OpBitcast %4 %122 +%124 = OpIAdd %4 %123 %6 +%125 = OpCompositeConstruct %117 %104 %121 +%126 = OpImageFetch %25 %94 %125 Lod %124 +%127 = OpCompositeExtract %13 %88 0 +%128 = OpBitcast %4 %127 +%129 = OpCompositeExtract %13 %88 2 +%130 = OpBitcast %4 %129 +%131 = OpImageFetch %25 %95 %128 Lod %130 +%132 = OpBitcast %24 %104 +%133 = OpCompositeExtract %13 %88 2 %134 = OpBitcast %4 %133 -%135 = OpIAdd %4 %134 %6 -%136 = OpBitcast %13 %132 -%137 = OpCompositeConstruct %20 %130 %136 -%138 = OpImageFetch %23 %89 %137 Lod %135 -%139 = OpCompositeExtract %13 %83 0 -%141 = OpCompositeExtract %13 %83 2 -%142 = OpBitcast %4 %141 -%143 = OpImageFetch %23 %90 %139 Lod %142 -%144 = OpCompositeExtract %4 %99 0 -%145 = OpIAdd %23 %102 %105 -%146 = OpIAdd %23 %145 %106 -%147 = OpIAdd %23 %146 %114 -%148 = OpIAdd %23 %147 %119 -OpImageWrite %91 %144 %148 -%149 = OpCompositeExtract %4 %99 0 -%150 = OpBitcast %13 %149 -%151 = OpIAdd %23 %123 %127 -%152 = OpIAdd %23 %151 %129 -%153 = OpIAdd %23 %152 %138 -%154 = OpIAdd %23 %153 %143 -OpImageWrite %91 %150 %154 +%135 = OpImageFetch %25 %91 %132 Lod %134 +%136 = OpBitcast %24 %104 +%137 = OpCompositeExtract %13 %88 2 +%138 = OpBitcast %4 %137 +%139 = OpImageFetch %25 %92 %136 Sample %138 +%140 = OpBitcast %24 %104 +%141 = OpImageRead %25 %93 %140 +%142 = OpBitcast %24 %104 +%143 = OpCompositeExtract %13 %88 2 +%144 = OpCompositeExtract %13 %88 2 +%145 = OpBitcast %4 %144 +%146 = OpIAdd %4 %145 %6 +%147 = OpCompositeConstruct %22 %142 %143 +%148 = OpImageFetch %25 %94 %147 Lod %146 +%149 = OpBitcast %24 %104 +%150 = OpCompositeExtract %13 %88 2 +%151 = OpBitcast %4 %150 +%152 = OpCompositeExtract %13 %88 2 +%153 = OpBitcast %4 %152 +%154 = OpIAdd %4 %153 %6 +%155 = OpBitcast %13 %151 +%156 = OpCompositeConstruct %22 %149 %155 +%157 = OpImageFetch %25 %94 %156 Lod %154 +%158 = OpCompositeExtract %13 %88 0 +%160 = OpCompositeExtract %13 %88 2 +%161 = OpBitcast %4 %160 +%162 = OpImageFetch %25 %95 %158 Lod %161 +%163 = OpCompositeExtract %4 %104 0 +%164 = OpIAdd %25 %107 %110 +%165 = OpIAdd %25 %164 %111 +%166 = OpIAdd %25 %165 %119 +%167 = OpIAdd %25 %166 %126 +OpImageWrite %96 %163 %167 +%168 = OpCompositeExtract %4 %104 0 +%169 = OpBitcast %13 %168 +%170 = OpIAdd %25 %135 %139 +%171 = OpIAdd %25 %170 %141 +%172 = OpIAdd %25 %171 %148 +%173 = OpIAdd %25 %172 %157 +OpImageWrite %96 %169 %173 OpReturn OpFunctionEnd -%158 = OpFunction %2 None %85 -%155 = OpLabel -%157 = OpLoad %20 %156 -%159 = OpLoad %15 %43 -%160 = OpLoad %16 %45 -%161 = OpLoad %18 %53 -OpBranch %162 -%162 = OpLabel -%163 = OpImageQuerySize %21 %160 -%164 = OpBitcast %22 %163 -%165 = OpVectorShuffle %22 %157 %157 0 1 -%166 = OpIMul %22 %164 %165 -%167 = OpBitcast %21 %166 -%168 = OpCompositeConstruct %21 %3 %5 -%169 = OpSRem %21 %167 %168 -%170 = OpCompositeExtract %13 %157 2 -%171 = OpBitcast %4 %170 -%172 = OpImageFetch %32 %159 %169 Sample %171 -%173 = OpCompositeExtract %8 %172 0 -%174 = OpCompositeExtract %4 %169 0 -%175 = OpConvertFToU %13 %173 -%176 = OpCompositeConstruct %23 %175 %175 %175 %175 -OpImageWrite %161 %174 %176 +%177 = OpFunction %2 None %90 +%174 = OpLabel +%176 = OpLoad %22 %175 +%178 = OpLoad %17 %46 +%179 = OpLoad %18 %48 +%180 = OpLoad %20 %56 +OpBranch %181 +%181 = OpLabel +%182 = OpImageQuerySize %23 %179 +%183 = OpBitcast %24 %182 +%184 = OpVectorShuffle %24 %176 %176 0 1 +%185 = OpIMul %24 %183 %184 +%186 = OpBitcast %23 %185 +%187 = OpCompositeConstruct %23 %3 %5 +%188 = OpSRem %23 %186 %187 +%189 = OpCompositeExtract %13 %176 2 +%190 = OpBitcast %4 %189 +%191 = OpImageFetch %34 %178 %188 Sample %190 +%192 = OpCompositeExtract %8 %191 0 +%193 = OpCompositeExtract %4 %188 0 +%194 = OpConvertFToU %13 %192 +%195 = OpCompositeConstruct %25 %194 %194 %194 %194 +OpImageWrite %180 %193 %195 OpReturn OpFunctionEnd -%180 = OpFunction %2 None %85 -%177 = OpLabel -%181 = OpLoad %24 %55 -%182 = OpLoad %25 %57 -%183 = OpLoad %27 %62 -%184 = OpLoad %28 %64 -%185 = OpLoad %29 %66 -%186 = OpLoad %30 %68 -%187 = OpLoad %31 %70 -OpBranch %188 -%188 = OpLabel -%190 = OpImageQuerySizeLod %4 %181 %189 -%191 = OpBitcast %13 %190 -%192 = OpBitcast %4 %191 -%193 = OpImageQuerySizeLod %4 %181 %192 -%194 = OpBitcast %13 %193 -%195 = OpImageQuerySizeLod %21 %182 %189 -%196 = OpBitcast %22 %195 -%197 = OpImageQuerySizeLod %21 %182 %6 -%198 = OpBitcast %22 %197 -%199 = OpImageQuerySizeLod %112 %183 %189 -%200 = OpBitcast %20 %199 -%201 = OpVectorShuffle %22 %200 %200 0 1 -%202 = OpImageQuerySizeLod %112 %183 %6 -%203 = OpBitcast %20 %202 -%204 = OpVectorShuffle %22 %203 %203 0 1 -%205 = OpImageQuerySizeLod %21 %184 %189 -%206 = OpBitcast %22 %205 -%207 = OpImageQuerySizeLod %21 %184 %6 -%208 = OpBitcast %22 %207 -%209 = OpImageQuerySizeLod %112 %185 %189 -%210 = OpBitcast %20 %209 -%211 = OpVectorShuffle %22 %210 %210 0 0 -%212 = OpImageQuerySizeLod %112 %185 %6 -%213 = OpBitcast %20 %212 -%214 = OpVectorShuffle %22 %213 %213 0 0 -%215 = OpImageQuerySizeLod %112 %186 %189 -%216 = OpBitcast %20 %215 -%217 = OpImageQuerySizeLod %112 %186 %6 -%218 = OpBitcast %20 %217 -%219 = OpImageQuerySize %21 %187 -%220 = OpBitcast %22 %219 -%221 = OpCompositeExtract %13 %196 1 -%222 = OpIAdd %13 %191 %221 -%223 = OpCompositeExtract %13 %198 1 -%224 = OpIAdd %13 %222 %223 -%225 = OpCompositeExtract %13 %201 1 -%226 = OpIAdd %13 %224 %225 -%227 = OpCompositeExtract %13 %204 1 -%228 = OpIAdd %13 %226 %227 -%229 = OpCompositeExtract %13 %206 1 -%230 = OpIAdd %13 %228 %229 -%231 = OpCompositeExtract %13 %208 1 -%232 = OpIAdd %13 %230 %231 -%233 = OpCompositeExtract %13 %211 1 -%234 = OpIAdd %13 %232 %233 -%235 = OpCompositeExtract %13 %214 1 -%236 = OpIAdd %13 %234 %235 -%237 = OpCompositeExtract %13 %216 2 -%238 = OpIAdd %13 %236 %237 -%239 = OpCompositeExtract %13 %218 2 -%240 = OpIAdd %13 %238 %239 -%241 = OpConvertUToF %8 %240 -%242 = OpCompositeConstruct %32 %241 %241 %241 %241 -OpStore %178 %242 +%199 = OpFunction %2 None %90 +%196 = OpLabel +%200 = OpLoad %26 %58 +%201 = OpLoad %27 %60 +%202 = OpLoad %29 %65 +%203 = OpLoad %30 %67 +%204 = OpLoad %31 %69 +%205 = OpLoad %32 %71 +%206 = OpLoad %33 %73 +OpBranch %207 +%207 = OpLabel +%208 = OpImageQuerySizeLod %4 %200 %12 +%209 = OpBitcast %13 %208 +%210 = OpBitcast %4 %209 +%211 = OpImageQuerySizeLod %4 %200 %210 +%212 = OpBitcast %13 %211 +%213 = OpImageQuerySizeLod %23 %201 %12 +%214 = OpBitcast %24 %213 +%215 = OpImageQuerySizeLod %23 %201 %6 +%216 = OpBitcast %24 %215 +%217 = OpImageQuerySizeLod %117 %202 %12 +%218 = OpBitcast %22 %217 +%219 = OpVectorShuffle %24 %218 %218 0 1 +%220 = OpImageQuerySizeLod %117 %202 %6 +%221 = OpBitcast %22 %220 +%222 = OpVectorShuffle %24 %221 %221 0 1 +%223 = OpImageQuerySizeLod %23 %203 %12 +%224 = OpBitcast %24 %223 +%225 = OpImageQuerySizeLod %23 %203 %6 +%226 = OpBitcast %24 %225 +%227 = OpImageQuerySizeLod %117 %204 %12 +%228 = OpBitcast %22 %227 +%229 = OpVectorShuffle %24 %228 %228 0 0 +%230 = OpImageQuerySizeLod %117 %204 %6 +%231 = OpBitcast %22 %230 +%232 = OpVectorShuffle %24 %231 %231 0 0 +%233 = OpImageQuerySizeLod %117 %205 %12 +%234 = OpBitcast %22 %233 +%235 = OpImageQuerySizeLod %117 %205 %6 +%236 = OpBitcast %22 %235 +%237 = OpImageQuerySize %23 %206 +%238 = OpBitcast %24 %237 +%239 = OpCompositeExtract %13 %214 1 +%240 = OpIAdd %13 %209 %239 +%241 = OpCompositeExtract %13 %216 1 +%242 = OpIAdd %13 %240 %241 +%243 = OpCompositeExtract %13 %219 1 +%244 = OpIAdd %13 %242 %243 +%245 = OpCompositeExtract %13 %222 1 +%246 = OpIAdd %13 %244 %245 +%247 = OpCompositeExtract %13 %224 1 +%248 = OpIAdd %13 %246 %247 +%249 = OpCompositeExtract %13 %226 1 +%250 = OpIAdd %13 %248 %249 +%251 = OpCompositeExtract %13 %229 1 +%252 = OpIAdd %13 %250 %251 +%253 = OpCompositeExtract %13 %232 1 +%254 = OpIAdd %13 %252 %253 +%255 = OpCompositeExtract %13 %234 2 +%256 = OpIAdd %13 %254 %255 +%257 = OpCompositeExtract %13 %236 2 +%258 = OpIAdd %13 %256 %257 +%259 = OpConvertUToF %8 %258 +%260 = OpCompositeConstruct %34 %259 %259 %259 %259 +OpStore %197 %260 OpReturn OpFunctionEnd -%245 = OpFunction %2 None %85 -%243 = OpLabel -%246 = OpLoad %25 %57 -%247 = OpLoad %27 %62 -%248 = OpLoad %28 %64 -%249 = OpLoad %29 %66 -%250 = OpLoad %30 %68 -%251 = OpLoad %31 %70 -OpBranch %252 -%252 = OpLabel -%253 = OpImageQueryLevels %4 %246 -%254 = OpBitcast %13 %253 -%255 = OpImageQueryLevels %4 %247 -%256 = OpBitcast %13 %255 -%257 = OpImageQuerySizeLod %112 %247 %189 -%258 = OpCompositeExtract %4 %257 2 -%259 = OpBitcast %13 %258 -%260 = OpImageQueryLevels %4 %248 -%261 = OpBitcast %13 %260 -%262 = OpImageQueryLevels %4 %249 -%263 = OpBitcast %13 %262 -%264 = OpImageQuerySizeLod %112 %249 %189 -%265 = OpCompositeExtract %4 %264 2 -%266 = OpBitcast %13 %265 -%267 = OpImageQueryLevels %4 %250 -%268 = OpBitcast %13 %267 -%269 = OpImageQuerySamples %4 %251 -%270 = OpBitcast %13 %269 -%271 = OpIAdd %13 %259 %266 -%272 = OpIAdd %13 %271 %270 -%273 = OpIAdd %13 %272 %254 -%274 = OpIAdd %13 %273 %256 -%275 = OpIAdd %13 %274 %268 -%276 = OpIAdd %13 %275 %261 -%277 = OpIAdd %13 %276 %263 -%278 = OpConvertUToF %8 %277 -%279 = OpCompositeConstruct %32 %278 %278 %278 %278 -OpStore %244 %279 +%263 = OpFunction %2 None %90 +%261 = OpLabel +%264 = OpLoad %27 %60 +%265 = OpLoad %29 %65 +%266 = OpLoad %30 %67 +%267 = OpLoad %31 %69 +%268 = OpLoad %32 %71 +%269 = OpLoad %33 %73 +OpBranch %270 +%270 = OpLabel +%271 = OpImageQueryLevels %4 %264 +%272 = OpBitcast %13 %271 +%273 = OpImageQueryLevels %4 %265 +%274 = OpBitcast %13 %273 +%275 = OpImageQuerySizeLod %117 %265 %12 +%276 = OpCompositeExtract %4 %275 2 +%277 = OpBitcast %13 %276 +%278 = OpImageQueryLevels %4 %266 +%279 = OpBitcast %13 %278 +%280 = OpImageQueryLevels %4 %267 +%281 = OpBitcast %13 %280 +%282 = OpImageQuerySizeLod %117 %267 %12 +%283 = OpCompositeExtract %4 %282 2 +%284 = OpBitcast %13 %283 +%285 = OpImageQueryLevels %4 %268 +%286 = OpBitcast %13 %285 +%287 = OpImageQuerySamples %4 %269 +%288 = OpBitcast %13 %287 +%289 = OpIAdd %13 %277 %284 +%290 = OpIAdd %13 %289 %288 +%291 = OpIAdd %13 %290 %272 +%292 = OpIAdd %13 %291 %274 +%293 = OpIAdd %13 %292 %286 +%294 = OpIAdd %13 %293 %279 +%295 = OpIAdd %13 %294 %281 +%296 = OpConvertUToF %8 %295 +%297 = OpCompositeConstruct %34 %296 %296 %296 %296 +OpStore %262 %297 OpReturn OpFunctionEnd -%282 = OpFunction %2 None %85 -%280 = OpLabel -%283 = OpLoad %24 %55 -%284 = OpLoad %25 %57 -%285 = OpLoad %33 %72 -OpBranch %286 -%286 = OpLabel -%287 = OpCompositeConstruct %34 %7 %7 -%288 = OpCompositeExtract %8 %287 0 -%290 = OpSampledImage %289 %283 %285 -%291 = OpImageSampleImplicitLod %32 %290 %288 -%293 = OpSampledImage %292 %284 %285 -%294 = OpImageSampleImplicitLod %32 %293 %287 -%295 = OpSampledImage %292 %284 %285 -%296 = OpImageSampleImplicitLod %32 %295 %287 ConstOffset %38 -%297 = OpSampledImage %292 %284 %285 -%298 = OpImageSampleExplicitLod %32 %297 %287 Lod %9 -%299 = OpSampledImage %292 %284 %285 -%300 = OpImageSampleExplicitLod %32 %299 %287 Lod|ConstOffset %9 %38 -%301 = OpSampledImage %292 %284 %285 -%302 = OpImageSampleImplicitLod %32 %301 %287 Bias|ConstOffset %11 %38 -%303 = OpFAdd %32 %291 %294 -%304 = OpFAdd %32 %303 %296 -%305 = OpFAdd %32 %304 %298 -%306 = OpFAdd %32 %305 %300 -OpStore %281 %306 +%303 = OpFunction %2 None %90 +%301 = OpLabel +%298 = OpVariable %299 Function %300 +%304 = OpLoad %26 %58 +%305 = OpLoad %27 %60 +%306 = OpLoad %29 %65 +%307 = OpLoad %31 %69 +%308 = OpLoad %35 %75 +OpBranch %309 +%309 = OpLabel +%310 = OpCompositeConstruct %36 %7 %7 +%311 = OpCompositeConstruct %37 %7 %7 %7 +%312 = OpCompositeExtract %8 %310 0 +%314 = OpSampledImage %313 %304 %308 +%315 = OpImageSampleImplicitLod %34 %314 %312 +%316 = OpLoad %34 %298 +%317 = OpFAdd %34 %316 %315 +OpStore %298 %317 +%319 = OpSampledImage %318 %305 %308 +%320 = OpImageSampleImplicitLod %34 %319 %310 +%321 = OpLoad %34 %298 +%322 = OpFAdd %34 %321 %320 +OpStore %298 %322 +%323 = OpSampledImage %318 %305 %308 +%324 = OpImageSampleImplicitLod %34 %323 %310 ConstOffset %41 +%325 = OpLoad %34 %298 +%326 = OpFAdd %34 %325 %324 +OpStore %298 %326 +%327 = OpSampledImage %318 %305 %308 +%328 = OpImageSampleExplicitLod %34 %327 %310 Lod %9 +%329 = OpLoad %34 %298 +%330 = OpFAdd %34 %329 %328 +OpStore %298 %330 +%331 = OpSampledImage %318 %305 %308 +%332 = OpImageSampleExplicitLod %34 %331 %310 Lod|ConstOffset %9 %41 +%333 = OpLoad %34 %298 +%334 = OpFAdd %34 %333 %332 +OpStore %298 %334 +%335 = OpSampledImage %318 %305 %308 +%336 = OpImageSampleImplicitLod %34 %335 %310 Bias|ConstOffset %11 %41 +%337 = OpLoad %34 %298 +%338 = OpFAdd %34 %337 %336 +OpStore %298 %338 +%340 = OpConvertUToF %8 %12 +%341 = OpCompositeConstruct %37 %310 %340 +%342 = OpSampledImage %339 %306 %308 +%343 = OpImageSampleImplicitLod %34 %342 %341 +%344 = OpLoad %34 %298 +%345 = OpFAdd %34 %344 %343 +OpStore %298 %345 +%346 = OpConvertUToF %8 %12 +%347 = OpCompositeConstruct %37 %310 %346 +%348 = OpSampledImage %339 %306 %308 +%349 = OpImageSampleImplicitLod %34 %348 %347 ConstOffset %41 +%350 = OpLoad %34 %298 +%351 = OpFAdd %34 %350 %349 +OpStore %298 %351 +%352 = OpConvertUToF %8 %12 +%353 = OpCompositeConstruct %37 %310 %352 +%354 = OpSampledImage %339 %306 %308 +%355 = OpImageSampleExplicitLod %34 %354 %353 Lod %9 +%356 = OpLoad %34 %298 +%357 = OpFAdd %34 %356 %355 +OpStore %298 %357 +%358 = OpConvertUToF %8 %12 +%359 = OpCompositeConstruct %37 %310 %358 +%360 = OpSampledImage %339 %306 %308 +%361 = OpImageSampleExplicitLod %34 %360 %359 Lod|ConstOffset %9 %41 +%362 = OpLoad %34 %298 +%363 = OpFAdd %34 %362 %361 +OpStore %298 %363 +%364 = OpConvertUToF %8 %12 +%365 = OpCompositeConstruct %37 %310 %364 +%366 = OpSampledImage %339 %306 %308 +%367 = OpImageSampleImplicitLod %34 %366 %365 Bias|ConstOffset %11 %41 +%368 = OpLoad %34 %298 +%369 = OpFAdd %34 %368 %367 +OpStore %298 %369 +%370 = OpConvertSToF %8 %14 +%371 = OpCompositeConstruct %37 %310 %370 +%372 = OpSampledImage %339 %306 %308 +%373 = OpImageSampleImplicitLod %34 %372 %371 +%374 = OpLoad %34 %298 +%375 = OpFAdd %34 %374 %373 +OpStore %298 %375 +%376 = OpConvertSToF %8 %14 +%377 = OpCompositeConstruct %37 %310 %376 +%378 = OpSampledImage %339 %306 %308 +%379 = OpImageSampleImplicitLod %34 %378 %377 ConstOffset %41 +%380 = OpLoad %34 %298 +%381 = OpFAdd %34 %380 %379 +OpStore %298 %381 +%382 = OpConvertSToF %8 %14 +%383 = OpCompositeConstruct %37 %310 %382 +%384 = OpSampledImage %339 %306 %308 +%385 = OpImageSampleExplicitLod %34 %384 %383 Lod %9 +%386 = OpLoad %34 %298 +%387 = OpFAdd %34 %386 %385 +OpStore %298 %387 +%388 = OpConvertSToF %8 %14 +%389 = OpCompositeConstruct %37 %310 %388 +%390 = OpSampledImage %339 %306 %308 +%391 = OpImageSampleExplicitLod %34 %390 %389 Lod|ConstOffset %9 %41 +%392 = OpLoad %34 %298 +%393 = OpFAdd %34 %392 %391 +OpStore %298 %393 +%394 = OpConvertSToF %8 %14 +%395 = OpCompositeConstruct %37 %310 %394 +%396 = OpSampledImage %339 %306 %308 +%397 = OpImageSampleImplicitLod %34 %396 %395 Bias|ConstOffset %11 %41 +%398 = OpLoad %34 %298 +%399 = OpFAdd %34 %398 %397 +OpStore %298 %399 +%401 = OpConvertUToF %8 %12 +%402 = OpCompositeConstruct %34 %311 %401 +%403 = OpSampledImage %400 %307 %308 +%404 = OpImageSampleImplicitLod %34 %403 %402 +%405 = OpLoad %34 %298 +%406 = OpFAdd %34 %405 %404 +OpStore %298 %406 +%407 = OpConvertUToF %8 %12 +%408 = OpCompositeConstruct %34 %311 %407 +%409 = OpSampledImage %400 %307 %308 +%410 = OpImageSampleExplicitLod %34 %409 %408 Lod %9 +%411 = OpLoad %34 %298 +%412 = OpFAdd %34 %411 %410 +OpStore %298 %412 +%413 = OpConvertUToF %8 %12 +%414 = OpCompositeConstruct %34 %311 %413 +%415 = OpSampledImage %400 %307 %308 +%416 = OpImageSampleImplicitLod %34 %415 %414 Bias %11 +%417 = OpLoad %34 %298 +%418 = OpFAdd %34 %417 %416 +OpStore %298 %418 +%419 = OpConvertSToF %8 %14 +%420 = OpCompositeConstruct %34 %311 %419 +%421 = OpSampledImage %400 %307 %308 +%422 = OpImageSampleImplicitLod %34 %421 %420 +%423 = OpLoad %34 %298 +%424 = OpFAdd %34 %423 %422 +OpStore %298 %424 +%425 = OpConvertSToF %8 %14 +%426 = OpCompositeConstruct %34 %311 %425 +%427 = OpSampledImage %400 %307 %308 +%428 = OpImageSampleExplicitLod %34 %427 %426 Lod %9 +%429 = OpLoad %34 %298 +%430 = OpFAdd %34 %429 %428 +OpStore %298 %430 +%431 = OpConvertSToF %8 %14 +%432 = OpCompositeConstruct %34 %311 %431 +%433 = OpSampledImage %400 %307 %308 +%434 = OpImageSampleImplicitLod %34 %433 %432 Bias %11 +%435 = OpLoad %34 %298 +%436 = OpFAdd %34 %435 %434 +OpStore %298 %436 +%437 = OpLoad %34 %298 +OpStore %302 %437 OpReturn OpFunctionEnd -%310 = OpFunction %2 None %85 -%307 = OpLabel -%311 = OpLoad %33 %74 -%312 = OpLoad %35 %76 -%313 = OpLoad %36 %78 -OpBranch %314 -%314 = OpLabel -%315 = OpCompositeConstruct %34 %7 %7 -%317 = OpSampledImage %316 %312 %311 -%318 = OpImageSampleDrefImplicitLod %8 %317 %315 %7 -%319 = OpSampledImage %316 %312 %311 -%320 = OpImageSampleDrefExplicitLod %8 %319 %315 %7 Lod %321 -%322 = OpCompositeConstruct %37 %7 %7 %7 -%324 = OpSampledImage %323 %313 %311 -%325 = OpImageSampleDrefExplicitLod %8 %324 %322 %7 Lod %321 -%326 = OpFAdd %8 %318 %320 -OpStore %308 %326 +%444 = OpFunction %2 None %90 +%441 = OpLabel +%438 = OpVariable %439 Function %440 +%445 = OpLoad %35 %77 +%446 = OpLoad %38 %79 +%447 = OpLoad %39 %81 +%448 = OpLoad %40 %83 +OpBranch %449 +%449 = OpLabel +%450 = OpCompositeConstruct %36 %7 %7 +%451 = OpCompositeConstruct %37 %7 %7 %7 +%453 = OpSampledImage %452 %446 %445 +%454 = OpImageSampleDrefImplicitLod %8 %453 %450 %7 +%455 = OpLoad %8 %438 +%456 = OpFAdd %8 %455 %454 +OpStore %438 %456 +%458 = OpConvertUToF %8 %12 +%459 = OpCompositeConstruct %37 %450 %458 +%460 = OpSampledImage %457 %447 %445 +%461 = OpImageSampleDrefImplicitLod %8 %460 %459 %7 +%462 = OpLoad %8 %438 +%463 = OpFAdd %8 %462 %461 +OpStore %438 %463 +%464 = OpConvertSToF %8 %14 +%465 = OpCompositeConstruct %37 %450 %464 +%466 = OpSampledImage %457 %447 %445 +%467 = OpImageSampleDrefImplicitLod %8 %466 %465 %7 +%468 = OpLoad %8 %438 +%469 = OpFAdd %8 %468 %467 +OpStore %438 %469 +%471 = OpSampledImage %470 %448 %445 +%472 = OpImageSampleDrefImplicitLod %8 %471 %451 %7 +%473 = OpLoad %8 %438 +%474 = OpFAdd %8 %473 %472 +OpStore %438 %474 +%475 = OpSampledImage %452 %446 %445 +%476 = OpImageSampleDrefExplicitLod %8 %475 %450 %7 Lod %477 +%478 = OpLoad %8 %438 +%479 = OpFAdd %8 %478 %476 +OpStore %438 %479 +%480 = OpConvertUToF %8 %12 +%481 = OpCompositeConstruct %37 %450 %480 +%482 = OpSampledImage %457 %447 %445 +%483 = OpImageSampleDrefExplicitLod %8 %482 %481 %7 Lod %477 +%484 = OpLoad %8 %438 +%485 = OpFAdd %8 %484 %483 +OpStore %438 %485 +%486 = OpConvertSToF %8 %14 +%487 = OpCompositeConstruct %37 %450 %486 +%488 = OpSampledImage %457 %447 %445 +%489 = OpImageSampleDrefExplicitLod %8 %488 %487 %7 Lod %477 +%490 = OpLoad %8 %438 +%491 = OpFAdd %8 %490 %489 +OpStore %438 %491 +%492 = OpSampledImage %470 %448 %445 +%493 = OpImageSampleDrefExplicitLod %8 %492 %451 %7 Lod %477 +%494 = OpLoad %8 %438 +%495 = OpFAdd %8 %494 %493 +OpStore %438 %495 +%496 = OpLoad %8 %438 +OpStore %442 %496 OpReturn OpFunctionEnd -%329 = OpFunction %2 None %85 -%327 = OpLabel -%330 = OpLoad %25 %57 -%331 = OpLoad %12 %59 -%332 = OpLoad %26 %60 -%333 = OpLoad %33 %72 -%334 = OpLoad %33 %74 -%335 = OpLoad %35 %76 -OpBranch %336 -%336 = OpLabel -%337 = OpCompositeConstruct %34 %7 %7 -%338 = OpSampledImage %292 %330 %333 -%339 = OpImageGather %32 %338 %337 %340 -%341 = OpSampledImage %292 %330 %333 -%342 = OpImageGather %32 %341 %337 %343 ConstOffset %38 -%344 = OpSampledImage %316 %335 %334 -%345 = OpImageDrefGather %32 %344 %337 %7 -%346 = OpSampledImage %316 %335 %334 -%347 = OpImageDrefGather %32 %346 %337 %7 ConstOffset %38 -%349 = OpSampledImage %348 %331 %333 -%350 = OpImageGather %23 %349 %337 %189 -%353 = OpSampledImage %352 %332 %333 -%354 = OpImageGather %351 %353 %337 %189 -%355 = OpConvertUToF %32 %350 -%356 = OpConvertSToF %32 %354 -%357 = OpFAdd %32 %355 %356 -%358 = OpFAdd %32 %339 %342 -%359 = OpFAdd %32 %358 %345 -%360 = OpFAdd %32 %359 %347 -%361 = OpFAdd %32 %360 %357 -OpStore %328 %361 +%499 = OpFunction %2 None %90 +%497 = OpLabel +%500 = OpLoad %27 %60 +%501 = OpLoad %15 %62 +%502 = OpLoad %28 %63 +%503 = OpLoad %35 %75 +%504 = OpLoad %35 %77 +%505 = OpLoad %38 %79 +OpBranch %506 +%506 = OpLabel +%507 = OpCompositeConstruct %36 %7 %7 +%508 = OpSampledImage %318 %500 %503 +%509 = OpImageGather %34 %508 %507 %510 +%511 = OpSampledImage %318 %500 %503 +%512 = OpImageGather %34 %511 %507 %513 ConstOffset %41 +%514 = OpSampledImage %452 %505 %504 +%515 = OpImageDrefGather %34 %514 %507 %7 +%516 = OpSampledImage %452 %505 %504 +%517 = OpImageDrefGather %34 %516 %507 %7 ConstOffset %41 +%519 = OpSampledImage %518 %501 %503 +%520 = OpImageGather %25 %519 %507 %12 +%523 = OpSampledImage %522 %502 %503 +%524 = OpImageGather %521 %523 %507 %12 +%525 = OpConvertUToF %34 %520 +%526 = OpConvertSToF %34 %524 +%527 = OpFAdd %34 %525 %526 +%528 = OpFAdd %34 %509 %512 +%529 = OpFAdd %34 %528 %515 +%530 = OpFAdd %34 %529 %517 +%531 = OpFAdd %34 %530 %527 +OpStore %498 %531 OpReturn OpFunctionEnd -%364 = OpFunction %2 None %85 -%362 = OpLabel -%365 = OpLoad %33 %72 -%366 = OpLoad %35 %76 -OpBranch %367 -%367 = OpLabel -%368 = OpCompositeConstruct %34 %7 %7 -%369 = OpSampledImage %316 %366 %365 -%370 = OpImageSampleImplicitLod %32 %369 %368 -%371 = OpCompositeExtract %8 %370 0 -%372 = OpSampledImage %316 %366 %365 -%373 = OpImageGather %32 %372 %368 %189 -%374 = OpCompositeConstruct %32 %371 %371 %371 %371 -%375 = OpFAdd %32 %374 %373 -OpStore %363 %375 +%534 = OpFunction %2 None %90 +%532 = OpLabel +%535 = OpLoad %35 %75 +%536 = OpLoad %38 %79 +OpBranch %537 +%537 = OpLabel +%538 = OpCompositeConstruct %36 %7 %7 +%539 = OpSampledImage %452 %536 %535 +%540 = OpImageSampleImplicitLod %34 %539 %538 +%541 = OpCompositeExtract %8 %540 0 +%542 = OpSampledImage %452 %536 %535 +%543 = OpImageGather %34 %542 %538 %12 +%544 = OpCompositeConstruct %34 %541 %541 %541 %541 +%545 = OpFAdd %34 %544 %543 +OpStore %533 %545 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/spv/shadow.spvasm b/tests/out/spv/shadow.spvasm index 15f18c4f55..15cbf4b918 100644 --- a/tests/out/spv/shadow.spvasm +++ b/tests/out/spv/shadow.spvasm @@ -223,7 +223,7 @@ OpReturnValue %7 %68 = OpBitcast %4 %48 %69 = OpCompositeExtract %6 %49 2 %70 = OpFMul %6 %69 %62 -%72 = OpConvertUToF %6 %68 +%72 = OpConvertSToF %6 %68 %73 = OpCompositeConstruct %20 %67 %72 %74 = OpSampledImage %71 %52 %53 %75 = OpImageSampleDrefExplicitLod %6 %74 %73 %70 Lod %5 diff --git a/tests/out/wgsl/image.wgsl b/tests/out/wgsl/image.wgsl index 0dcaf3a4d4..3409712bc7 100644 --- a/tests/out/wgsl/image.wgsl +++ b/tests/out/wgsl/image.wgsl @@ -39,6 +39,8 @@ var sampler_cmp: sampler_comparison; @group(1) @binding(2) var image_2d_depth: texture_depth_2d; @group(1) @binding(3) +var image_2d_array_depth: texture_depth_2d_array; +@group(1) @binding(4) var image_cube_depth: texture_depth_cube; @compute @workgroup_size(16, 1, 1) @@ -48,13 +50,15 @@ fn main(@builtin(local_invocation_id) local_id: vec3) { 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)); + let value5_ = textureLoad(image_array_src, itc, local_id.z, (i32(local_id.z) + 1)); + let value6_ = textureLoad(image_array_src, itc, i32(local_id.z), (i32(local_id.z) + 1)); + let value7_ = textureLoad(image_1d_src, i32(local_id.x), i32(local_id.z)); let value1u = textureLoad(image_mipmapped_src, vec2(itc), i32(local_id.z)); let value2u = textureLoad(image_multisampled_src, vec2(itc), i32(local_id.z)); let value4u = textureLoad(image_storage_src, vec2(itc)); - let value5u = textureLoad(image_array_src, vec2(itc), i32(local_id.z), (i32(local_id.z) + 1)); - let value6u = textureLoad(image_1d_src, u32(local_id.x), i32(local_id.z)); + let value5u = textureLoad(image_array_src, vec2(itc), local_id.z, (i32(local_id.z) + 1)); + let value6u = textureLoad(image_array_src, vec2(itc), i32(local_id.z), (i32(local_id.z) + 1)); + let value7u = textureLoad(image_1d_src, u32(local_id.x), i32(local_id.z)); textureStore(image_dst, itc.x, ((((value1_ + value2_) + value4_) + value5_) + value6_)); textureStore(image_dst, u32(itc.x), ((((value1u + value2u) + value4u) + value5u) + value6u)); return; @@ -104,42 +108,131 @@ fn levels_queries() -> @builtin(position) vec4 { @fragment fn texture_sample() -> @location(0) vec4 { + var a: vec4; + let tc = vec2(0.5); - 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, 2.299999952316284); - let s2d_level_offset = textureSampleLevel(image_2d, sampler_reg, tc, 2.299999952316284, vec2(3, 1)); - let s2d_bias_offset = textureSampleBias(image_2d, sampler_reg, tc, 2.0, vec2(3, 1)); - return ((((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset); + let tc3_ = vec3(0.5); + let _e9 = textureSample(image_1d, sampler_reg, tc.x); + let _e10 = a; + a = (_e10 + _e9); + let _e14 = textureSample(image_2d, sampler_reg, tc); + let _e15 = a; + a = (_e15 + _e14); + let _e19 = textureSample(image_2d, sampler_reg, tc, vec2(3, 1)); + let _e20 = a; + a = (_e20 + _e19); + let _e24 = textureSampleLevel(image_2d, sampler_reg, tc, 2.299999952316284); + let _e25 = a; + a = (_e25 + _e24); + let _e29 = textureSampleLevel(image_2d, sampler_reg, tc, 2.299999952316284, vec2(3, 1)); + let _e30 = a; + a = (_e30 + _e29); + let _e35 = textureSampleBias(image_2d, sampler_reg, tc, 2.0, vec2(3, 1)); + let _e36 = a; + a = (_e36 + _e35); + let _e41 = textureSample(image_2d_array, sampler_reg, tc, 0u); + let _e42 = a; + a = (_e42 + _e41); + let _e47 = textureSample(image_2d_array, sampler_reg, tc, 0u, vec2(3, 1)); + let _e48 = a; + a = (_e48 + _e47); + let _e53 = textureSampleLevel(image_2d_array, sampler_reg, tc, 0u, 2.299999952316284); + let _e54 = a; + a = (_e54 + _e53); + let _e59 = textureSampleLevel(image_2d_array, sampler_reg, tc, 0u, 2.299999952316284, vec2(3, 1)); + let _e60 = a; + a = (_e60 + _e59); + let _e66 = textureSampleBias(image_2d_array, sampler_reg, tc, 0u, 2.0, vec2(3, 1)); + let _e67 = a; + a = (_e67 + _e66); + let _e72 = textureSample(image_2d_array, sampler_reg, tc, 0); + let _e73 = a; + a = (_e73 + _e72); + let _e78 = textureSample(image_2d_array, sampler_reg, tc, 0, vec2(3, 1)); + let _e79 = a; + a = (_e79 + _e78); + let _e84 = textureSampleLevel(image_2d_array, sampler_reg, tc, 0, 2.299999952316284); + let _e85 = a; + a = (_e85 + _e84); + let _e90 = textureSampleLevel(image_2d_array, sampler_reg, tc, 0, 2.299999952316284, vec2(3, 1)); + let _e91 = a; + a = (_e91 + _e90); + let _e97 = textureSampleBias(image_2d_array, sampler_reg, tc, 0, 2.0, vec2(3, 1)); + let _e98 = a; + a = (_e98 + _e97); + let _e103 = textureSample(image_cube_array, sampler_reg, tc3_, 0u); + let _e104 = a; + a = (_e104 + _e103); + let _e109 = textureSampleLevel(image_cube_array, sampler_reg, tc3_, 0u, 2.299999952316284); + let _e110 = a; + a = (_e110 + _e109); + let _e116 = textureSampleBias(image_cube_array, sampler_reg, tc3_, 0u, 2.0); + let _e117 = a; + a = (_e117 + _e116); + let _e122 = textureSample(image_cube_array, sampler_reg, tc3_, 0); + let _e123 = a; + a = (_e123 + _e122); + let _e128 = textureSampleLevel(image_cube_array, sampler_reg, tc3_, 0, 2.299999952316284); + let _e129 = a; + a = (_e129 + _e128); + let _e135 = textureSampleBias(image_cube_array, sampler_reg, tc3_, 0, 2.0); + let _e136 = a; + a = (_e136 + _e135); + let _e138 = a; + return _e138; } @fragment fn texture_sample_comparison() -> @location(0) f32 { + var a_1: f32; + let tc_1 = vec2(0.5); - let s2d_depth = textureSampleCompare(image_2d_depth, sampler_cmp, tc_1, 0.5); - let s2d_depth_level = textureSampleCompareLevel(image_2d_depth, sampler_cmp, tc_1, 0.5); - let scube_depth_level = textureSampleCompareLevel(image_cube_depth, sampler_cmp, vec3(0.5), 0.5); - return (s2d_depth + s2d_depth_level); + let tc3_1 = vec3(0.5); + let _e8 = textureSampleCompare(image_2d_depth, sampler_cmp, tc_1, 0.5); + let _e9 = a_1; + a_1 = (_e9 + _e8); + let _e14 = textureSampleCompare(image_2d_array_depth, sampler_cmp, tc_1, 0u, 0.5); + let _e15 = a_1; + a_1 = (_e15 + _e14); + let _e20 = textureSampleCompare(image_2d_array_depth, sampler_cmp, tc_1, 0, 0.5); + let _e21 = a_1; + a_1 = (_e21 + _e20); + let _e25 = textureSampleCompare(image_cube_depth, sampler_cmp, tc3_1, 0.5); + let _e26 = a_1; + a_1 = (_e26 + _e25); + let _e30 = textureSampleCompareLevel(image_2d_depth, sampler_cmp, tc_1, 0.5); + let _e31 = a_1; + a_1 = (_e31 + _e30); + let _e36 = textureSampleCompareLevel(image_2d_array_depth, sampler_cmp, tc_1, 0u, 0.5); + let _e37 = a_1; + a_1 = (_e37 + _e36); + let _e42 = textureSampleCompareLevel(image_2d_array_depth, sampler_cmp, tc_1, 0, 0.5); + let _e43 = a_1; + a_1 = (_e43 + _e42); + let _e47 = textureSampleCompareLevel(image_cube_depth, sampler_cmp, tc3_1, 0.5); + let _e48 = a_1; + a_1 = (_e48 + _e47); + let _e50 = a_1; + return _e50; } @fragment fn gather() -> @location(0) vec4 { let tc_2 = vec2(0.5); - let s2d_1 = textureGather(1, image_2d, sampler_reg, tc_2); - let s2d_offset_1 = textureGather(3, image_2d, sampler_reg, tc_2, vec2(3, 1)); - let s2d_depth_1 = textureGatherCompare(image_2d_depth, sampler_cmp, tc_2, 0.5); + let s2d = textureGather(1, image_2d, sampler_reg, tc_2); + let s2d_offset = textureGather(3, image_2d, sampler_reg, tc_2, vec2(3, 1)); + let s2d_depth = textureGatherCompare(image_2d_depth, sampler_cmp, tc_2, 0.5); let s2d_depth_offset = textureGatherCompare(image_2d_depth, sampler_cmp, tc_2, 0.5, vec2(3, 1)); let u = textureGather(0, image_2d_u32_, sampler_reg, tc_2); let i = textureGather(0, image_2d_i32_, sampler_reg, tc_2); let f = (vec4(u) + vec4(i)); - return ((((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset) + f); + return ((((s2d + s2d_offset) + s2d_depth) + s2d_depth_offset) + f); } @fragment fn depth_no_comparison() -> @location(0) vec4 { let tc_3 = vec2(0.5); - let s2d_2 = textureSample(image_2d_depth, sampler_reg, tc_3); + let s2d_1 = textureSample(image_2d_depth, sampler_reg, tc_3); let s2d_gather = textureGather(image_2d_depth, sampler_reg, tc_3); - return (vec4(s2d_2) + s2d_gather); + return (vec4(s2d_1) + s2d_gather); }