nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2024-11-22 23:04:07 +00:00
Code Issues Packages Projects Releases Wiki Activity

[wgsl-in] don't treat let declarations as const declarations

Browse Source
This commit is contained in:
teoxoy 2023-09-29 19:20:50 +02:00 committed by Teodor Tanasoaia
parent 3d346977e5
commit 55dd0e1fbc
24 changed files with 926 additions and 931 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

22
src/front/wgsl/lower/mod.rs
View File

@ -363,11 +363,16 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
fn const_access(&self, handle: Handle<crate::Expression>) -> Option<u32> {
match self.expr_type {
ExpressionContextType::Runtime(ref ctx) => self
.module
.to_ctx()
.eval_expr_to_u32_from(handle, ctx.naga_expressions)
.ok(),
ExpressionContextType::Runtime(ref ctx) => {
if !ctx.expression_constness.is_const(handle) {
return None;
}
self.module
.to_ctx()
.eval_expr_to_u32_from(handle, ctx.naga_expressions)
.ok()
}
ExpressionContextType::Constant => self.module.to_ctx().eval_expr_to_u32(handle).ok(),
}
}
@ -404,6 +409,12 @@ impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
) -> Result<crate::SwizzleComponent, Error<'source>> {
match self.expr_type {
ExpressionContextType::Runtime(ref rctx) => {
if !rctx.expression_constness.is_const(expr) {
return Err(Error::ExpectedConstExprConcreteIntegerScalar(
component_span,
));
}
let index = self
.module
.to_ctx()
@ -1076,6 +1087,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
emitter.start(ctx.naga_expressions);
let value = self.expression(l.init, ctx.as_expression(block, &mut emitter))?;
ctx.expression_constness.force_non_const(value);
let explicit_ty =
l.ty.map(|ty| self.resolve_ast_type(ty, ctx.as_global()))

5
src/proc/constant_evaluator.rs
View File

@ -37,6 +37,11 @@ impl ExpressionConstnessTracker {
}
}
/// Forces the the expression to not be const
pub fn force_non_const(&mut self, value: Handle<Expression>) {
self.inner.remove(value.index());
}
fn insert(&mut self, value: Handle<Expression>) {
self.inner.insert(value.index());
}

2
src/proc/mod.rs
View File

@ -612,7 +612,7 @@ impl GlobalCtx<'_> {
self.eval_expr_to_literal_from(handle, self.const_expressions)
}
pub(crate) fn eval_expr_to_literal_from(
fn eval_expr_to_literal_from(
&self,
handle: crate::Handle<crate::Expression>,
arena: &crate::Arena<crate::Expression>,

8
tests/in/const-exprs.wgsl
View File

@ -13,16 +13,12 @@ fn main() {
// Swizzle the value of nested Compose expressions.
fn swizzle_of_compose() {
let a = vec2(1, 2);
let b = vec2(3, 4);
out = vec4(a, b).wzyx; // should assign vec4(4, 3, 2, 1);
out = vec4(vec2(1, 2), vec2(3, 4)).wzyx; // should assign vec4(4, 3, 2, 1);
}
// Index the value of nested Compose expressions.
fn index_of_compose() {
let a = vec2(1, 2);
let b = vec2(3, 4);
out2 += vec4(a, b)[1]; // should assign 2
out2 += vec4(vec2(1, 2), vec2(3, 4))[1]; // should assign 2
}
// Index the value of Compose expressions nested three deep

2
tests/out/analysis/access.info.ron
View File

@ -2774,7 +2774,7 @@
non_uniform_result: None,
requirements: (""),
),
ref_count: 0,
ref_count: 1,
assignable_global: None,
ty: Value(Scalar(
kind: Uint,

2
tests/out/glsl/access.foo_vert.Vertex.glsl
View File

@ -132,7 +132,7 @@ void main() {
test_matrix_within_array_within_struct_accesses();
mat4x3 _matrix = _group_0_binding_0_vs._matrix;
uvec2 arr_1[2] = _group_0_binding_0_vs.arr;
float b = _group_0_binding_0_vs._matrix[3][0];
float b = _group_0_binding_0_vs._matrix[3u][0];
int a_1 = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value;
ivec2 c = _group_0_binding_2_vs;
float _e33 = read_from_private(foo);

8
tests/out/glsl/const-exprs.main.Compute.glsl
View File

@ -16,17 +16,13 @@ layout(std430) buffer type_1_block_1Compute { int _group_0_binding_1_cs; };
void swizzle_of_compose() {
ivec2 a = ivec2(1, 2);
ivec2 b = ivec2(3, 4);
_group_0_binding_0_cs = ivec4(4, 3, 2, 1);
return;
}
void index_of_compose() {
ivec2 a_1 = ivec2(1, 2);
ivec2 b_1 = ivec2(3, 4);
int _e7 = _group_0_binding_1_cs;
_group_0_binding_1_cs = (_e7 + 2);
int _e2 = _group_0_binding_1_cs;
_group_0_binding_1_cs = (_e2 + 2);
return;
}

136
tests/out/glsl/image.texture_sample.Fragment.glsl
View File

@ -18,74 +18,74 @@ void main() {
vec4 a = vec4(0.0);
vec2 tc = vec2(0.5);
vec3 tc3_ = vec3(0.5);
vec4 _e8 = texture(_group_0_binding_0_fs, vec2(0.5, 0.0));
vec4 _e9 = a;
a = (_e9 + _e8);
vec4 _e13 = texture(_group_0_binding_1_fs, vec2(tc));
vec4 _e14 = a;
a = (_e14 + _e13);
vec4 _e18 = textureOffset(_group_0_binding_1_fs, vec2(tc), ivec2(3, 1));
vec4 _e19 = a;
a = (_e19 + _e18);
vec4 _e23 = textureLod(_group_0_binding_1_fs, vec2(tc), 2.3);
vec4 _e24 = a;
a = (_e24 + _e23);
vec4 _e28 = textureLodOffset(_group_0_binding_1_fs, vec2(tc), 2.3, ivec2(3, 1));
vec4 _e29 = a;
a = (_e29 + _e28);
vec4 _e34 = textureOffset(_group_0_binding_1_fs, vec2(tc), ivec2(3, 1), 2.0);
vec4 _e35 = a;
a = (_e35 + _e34);
vec4 _e40 = texture(_group_0_binding_4_fs, vec3(tc, 0u));
vec4 _e41 = a;
a = (_e41 + _e40);
vec4 _e46 = textureOffset(_group_0_binding_4_fs, vec3(tc, 0u), ivec2(3, 1));
vec4 _e47 = a;
a = (_e47 + _e46);
vec4 _e52 = textureLod(_group_0_binding_4_fs, vec3(tc, 0u), 2.3);
vec4 _e53 = a;
a = (_e53 + _e52);
vec4 _e58 = textureLodOffset(_group_0_binding_4_fs, vec3(tc, 0u), 2.3, ivec2(3, 1));
vec4 _e59 = a;
a = (_e59 + _e58);
vec4 _e65 = textureOffset(_group_0_binding_4_fs, vec3(tc, 0u), ivec2(3, 1), 2.0);
vec4 _e66 = a;
a = (_e66 + _e65);
vec4 _e71 = texture(_group_0_binding_4_fs, vec3(tc, 0));
vec4 _e72 = a;
a = (_e72 + _e71);
vec4 _e77 = textureOffset(_group_0_binding_4_fs, vec3(tc, 0), ivec2(3, 1));
vec4 _e78 = a;
a = (_e78 + _e77);
vec4 _e83 = textureLod(_group_0_binding_4_fs, vec3(tc, 0), 2.3);
vec4 _e84 = a;
a = (_e84 + _e83);
vec4 _e89 = textureLodOffset(_group_0_binding_4_fs, vec3(tc, 0), 2.3, ivec2(3, 1));
vec4 _e90 = a;
a = (_e90 + _e89);
vec4 _e96 = textureOffset(_group_0_binding_4_fs, vec3(tc, 0), ivec2(3, 1), 2.0);
vec4 _e97 = a;
a = (_e97 + _e96);
vec4 _e102 = texture(_group_0_binding_6_fs, vec4(tc3_, 0u));
vec4 _e103 = a;
a = (_e103 + _e102);
vec4 _e108 = textureLod(_group_0_binding_6_fs, vec4(tc3_, 0u), 2.3);
vec4 _e109 = a;
a = (_e109 + _e108);
vec4 _e115 = texture(_group_0_binding_6_fs, vec4(tc3_, 0u), 2.0);
vec4 _e116 = a;
a = (_e116 + _e115);
vec4 _e121 = texture(_group_0_binding_6_fs, vec4(tc3_, 0));
vec4 _e122 = a;
a = (_e122 + _e121);
vec4 _e127 = textureLod(_group_0_binding_6_fs, vec4(tc3_, 0), 2.3);
vec4 _e128 = a;
a = (_e128 + _e127);
vec4 _e134 = texture(_group_0_binding_6_fs, vec4(tc3_, 0), 2.0);
vec4 _e135 = a;
a = (_e135 + _e134);
vec4 _e137 = a;
_fs2p_location0 = _e137;
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.3);
vec4 _e25 = a;
a = (_e25 + _e24);
vec4 _e29 = textureLodOffset(_group_0_binding_1_fs, vec2(tc), 2.3, 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.3);
vec4 _e54 = a;
a = (_e54 + _e53);
vec4 _e59 = textureLodOffset(_group_0_binding_4_fs, vec3(tc, 0u), 2.3, 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.3);
vec4 _e85 = a;
a = (_e85 + _e84);
vec4 _e90 = textureLodOffset(_group_0_binding_4_fs, vec3(tc, 0), 2.3, 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.3);
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.3);
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;
}

2
tests/out/hlsl/access.hlsl
View File

@ -256,7 +256,7 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position
test_matrix_within_array_within_struct_accesses();
float4x3 _matrix = float4x3(asfloat(bar.Load3(0+0)), asfloat(bar.Load3(0+16)), asfloat(bar.Load3(0+32)), asfloat(bar.Load3(0+48)));
uint2 arr_1[2] = Constructarray2_uint2_(asuint(bar.Load2(144+0)), asuint(bar.Load2(144+8)));
float b = asfloat(bar.Load(0+48+0));
float b = asfloat(bar.Load(0+3u*16+0));
int a_1 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160));
int2 c = asint(qux.Load2(0));
const float _e33 = read_from_private(foo);

8
tests/out/hlsl/const-exprs.hlsl
View File

@ -8,18 +8,14 @@ RWByteAddressBuffer out2_ : register(u1);
void swizzle_of_compose()
{
int2 a = int2(1, 2);
int2 b = int2(3, 4);
out_.Store4(0, asuint(int4(4, 3, 2, 1)));
return;
}
void index_of_compose()
{
int2 a_1 = int2(1, 2);
int2 b_1 = int2(3, 4);
int _expr7 = asint(out2_.Load(0));
out2_.Store(0, asuint((_expr7 + 2)));
int _expr2 = asint(out2_.Load(0));
out2_.Store(0, asuint((_expr2 + 2)));
return;
}

136
tests/out/hlsl/image.hlsl
View File

@ -246,74 +246,74 @@ float4 texture_sample() : SV_Target0
float2 tc = (0.5).xx;
float3 tc3_ = (0.5).xxx;
float4 _expr8 = image_1d.Sample(sampler_reg, 0.5);
float4 _expr9 = a;
a = (_expr9 + _expr8);
float4 _expr13 = image_2d.Sample(sampler_reg, tc);
float4 _expr14 = a;
a = (_expr14 + _expr13);
float4 _expr18 = image_2d.Sample(sampler_reg, tc, int2(int2(3, 1)));
float4 _expr19 = a;
a = (_expr19 + _expr18);
float4 _expr23 = image_2d.SampleLevel(sampler_reg, tc, 2.3);
float4 _expr24 = a;
a = (_expr24 + _expr23);
float4 _expr28 = image_2d.SampleLevel(sampler_reg, tc, 2.3, int2(int2(3, 1)));
float4 _expr29 = a;
a = (_expr29 + _expr28);
float4 _expr34 = image_2d.SampleBias(sampler_reg, tc, 2.0, int2(int2(3, 1)));
float4 _expr35 = a;
a = (_expr35 + _expr34);
float4 _expr40 = image_2d_array.Sample(sampler_reg, float3(tc, 0u));
float4 _expr41 = a;
a = (_expr41 + _expr40);
float4 _expr46 = image_2d_array.Sample(sampler_reg, float3(tc, 0u), int2(int2(3, 1)));
float4 _expr47 = a;
a = (_expr47 + _expr46);
float4 _expr52 = image_2d_array.SampleLevel(sampler_reg, float3(tc, 0u), 2.3);
float4 _expr53 = a;
a = (_expr53 + _expr52);
float4 _expr58 = image_2d_array.SampleLevel(sampler_reg, float3(tc, 0u), 2.3, int2(int2(3, 1)));
float4 _expr59 = a;
a = (_expr59 + _expr58);
float4 _expr65 = image_2d_array.SampleBias(sampler_reg, float3(tc, 0u), 2.0, int2(int2(3, 1)));
float4 _expr66 = a;
a = (_expr66 + _expr65);
float4 _expr71 = image_2d_array.Sample(sampler_reg, float3(tc, 0));
float4 _expr72 = a;
a = (_expr72 + _expr71);
float4 _expr77 = image_2d_array.Sample(sampler_reg, float3(tc, 0), int2(int2(3, 1)));
float4 _expr78 = a;
a = (_expr78 + _expr77);
float4 _expr83 = image_2d_array.SampleLevel(sampler_reg, float3(tc, 0), 2.3);
float4 _expr84 = a;
a = (_expr84 + _expr83);
float4 _expr89 = image_2d_array.SampleLevel(sampler_reg, float3(tc, 0), 2.3, int2(int2(3, 1)));
float4 _expr90 = a;
a = (_expr90 + _expr89);
float4 _expr96 = image_2d_array.SampleBias(sampler_reg, float3(tc, 0), 2.0, int2(int2(3, 1)));
float4 _expr97 = a;
a = (_expr97 + _expr96);
float4 _expr102 = image_cube_array.Sample(sampler_reg, float4(tc3_, 0u));
float4 _expr103 = a;
a = (_expr103 + _expr102);
float4 _expr108 = image_cube_array.SampleLevel(sampler_reg, float4(tc3_, 0u), 2.3);
float4 _expr109 = a;
a = (_expr109 + _expr108);
float4 _expr115 = image_cube_array.SampleBias(sampler_reg, float4(tc3_, 0u), 2.0);
float4 _expr116 = a;
a = (_expr116 + _expr115);
float4 _expr121 = image_cube_array.Sample(sampler_reg, float4(tc3_, 0));
float4 _expr122 = a;
a = (_expr122 + _expr121);
float4 _expr127 = image_cube_array.SampleLevel(sampler_reg, float4(tc3_, 0), 2.3);
float4 _expr128 = a;
a = (_expr128 + _expr127);
float4 _expr134 = image_cube_array.SampleBias(sampler_reg, float4(tc3_, 0), 2.0);
float4 _expr135 = a;
a = (_expr135 + _expr134);
float4 _expr137 = a;
return _expr137;
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(int2(3, 1)));
float4 _expr20 = a;
a = (_expr20 + _expr19);
float4 _expr24 = image_2d.SampleLevel(sampler_reg, tc, 2.3);
float4 _expr25 = a;
a = (_expr25 + _expr24);
float4 _expr29 = image_2d.SampleLevel(sampler_reg, tc, 2.3, int2(int2(3, 1)));
float4 _expr30 = a;
a = (_expr30 + _expr29);
float4 _expr35 = image_2d.SampleBias(sampler_reg, tc, 2.0, int2(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(int2(3, 1)));
float4 _expr48 = a;
a = (_expr48 + _expr47);
float4 _expr53 = image_2d_array.SampleLevel(sampler_reg, float3(tc, 0u), 2.3);
float4 _expr54 = a;
a = (_expr54 + _expr53);
float4 _expr59 = image_2d_array.SampleLevel(sampler_reg, float3(tc, 0u), 2.3, int2(int2(3, 1)));
float4 _expr60 = a;
a = (_expr60 + _expr59);
float4 _expr66 = image_2d_array.SampleBias(sampler_reg, float3(tc, 0u), 2.0, int2(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(int2(3, 1)));
float4 _expr79 = a;
a = (_expr79 + _expr78);
float4 _expr84 = image_2d_array.SampleLevel(sampler_reg, float3(tc, 0), 2.3);
float4 _expr85 = a;
a = (_expr85 + _expr84);
float4 _expr90 = image_2d_array.SampleLevel(sampler_reg, float3(tc, 0), 2.3, int2(int2(3, 1)));
float4 _expr91 = a;
a = (_expr91 + _expr90);
float4 _expr97 = image_2d_array.SampleBias(sampler_reg, float3(tc, 0), 2.0, int2(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.3);
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.3);
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

4
tests/out/ir/access.compact.ron
View File

@ -1697,9 +1697,9 @@
base: 13,
index: 0,
),
AccessIndex(
Access(
base: 14,
index: 3,
index: 12,
),
AccessIndex(
base: 15,

4
tests/out/ir/access.ron
View File

@ -1792,9 +1792,9 @@
base: 13,
index: 0,
),
AccessIndex(
Access(
base: 14,
index: 3,
index: 12,
),
AccessIndex(
base: 15,

2
tests/out/msl/access.msl
View File

@ -180,7 +180,7 @@ vertex foo_vertOutput foo_vert(
test_matrix_within_array_within_struct_accesses(nested_mat_cx2_);
metal::float4x3 _matrix = bar._matrix;
type_9 arr_1 = bar.arr;
float b = bar._matrix[3].x;
float b = bar._matrix[3u].x;
int a_1 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value;
metal::int2 c = qux;
float _e33 = read_from_private(foo);

8
tests/out/msl/const-exprs.msl
View File

@ -12,8 +12,6 @@ constant int TEST_CONSTANT_ALIAS_ADDITION = 8;
void swizzle_of_compose(
device metal::int4& out
) {
metal::int2 a = metal::int2(1, 2);
metal::int2 b = metal::int2(3, 4);
out = metal::int4(4, 3, 2, 1);
return;
}
@ -21,10 +19,8 @@ void swizzle_of_compose(
void index_of_compose(
device int& out2_
) {
metal::int2 a_1 = metal::int2(1, 2);
metal::int2 b_1 = metal::int2(3, 4);
int _e7 = out2_;
out2_ = _e7 + 2;
int _e2 = out2_;
out2_ = _e2 + 2;
return;
}

136
tests/out/msl/image.msl
View File

@ -119,74 +119,74 @@ fragment texture_sampleOutput texture_sample(
metal::float4 a = {};
metal::float2 tc = metal::float2(0.5);
metal::float3 tc3_ = metal::float3(0.5);
metal::float4 _e8 = image_1d.sample(sampler_reg, 0.5);
metal::float4 _e9 = a;
a = _e9 + _e8;
metal::float4 _e13 = image_2d.sample(sampler_reg, tc);
metal::float4 _e14 = a;
a = _e14 + _e13;
metal::float4 _e18 = image_2d.sample(sampler_reg, tc, metal::int2(3, 1));
metal::float4 _e19 = a;
a = _e19 + _e18;
metal::float4 _e23 = image_2d.sample(sampler_reg, tc, metal::level(2.3));
metal::float4 _e24 = a;
a = _e24 + _e23;
metal::float4 _e28 = image_2d.sample(sampler_reg, tc, metal::level(2.3), metal::int2(3, 1));
metal::float4 _e29 = a;
a = _e29 + _e28;
metal::float4 _e34 = image_2d.sample(sampler_reg, tc, metal::bias(2.0), metal::int2(3, 1));
metal::float4 _e35 = a;
a = _e35 + _e34;
metal::float4 _e40 = image_2d_array.sample(sampler_reg, tc, 0u);
metal::float4 _e41 = a;
a = _e41 + _e40;
metal::float4 _e46 = image_2d_array.sample(sampler_reg, tc, 0u, metal::int2(3, 1));
metal::float4 _e47 = a;
a = _e47 + _e46;
metal::float4 _e52 = image_2d_array.sample(sampler_reg, tc, 0u, metal::level(2.3));
metal::float4 _e53 = a;
a = _e53 + _e52;
metal::float4 _e58 = image_2d_array.sample(sampler_reg, tc, 0u, metal::level(2.3), metal::int2(3, 1));
metal::float4 _e59 = a;
a = _e59 + _e58;
metal::float4 _e65 = image_2d_array.sample(sampler_reg, tc, 0u, metal::bias(2.0), metal::int2(3, 1));
metal::float4 _e66 = a;
a = _e66 + _e65;
metal::float4 _e71 = image_2d_array.sample(sampler_reg, tc, 0);
metal::float4 _e72 = a;
a = _e72 + _e71;
metal::float4 _e77 = image_2d_array.sample(sampler_reg, tc, 0, metal::int2(3, 1));
metal::float4 _e78 = a;
a = _e78 + _e77;
metal::float4 _e83 = image_2d_array.sample(sampler_reg, tc, 0, metal::level(2.3));
metal::float4 _e84 = a;
a = _e84 + _e83;
metal::float4 _e89 = image_2d_array.sample(sampler_reg, tc, 0, metal::level(2.3), metal::int2(3, 1));
metal::float4 _e90 = a;
a = _e90 + _e89;
metal::float4 _e96 = image_2d_array.sample(sampler_reg, tc, 0, metal::bias(2.0), metal::int2(3, 1));
metal::float4 _e97 = a;
a = _e97 + _e96;
metal::float4 _e102 = image_cube_array.sample(sampler_reg, tc3_, 0u);
metal::float4 _e103 = a;
a = _e103 + _e102;
metal::float4 _e108 = image_cube_array.sample(sampler_reg, tc3_, 0u, metal::level(2.3));
metal::float4 _e109 = a;
a = _e109 + _e108;
metal::float4 _e115 = image_cube_array.sample(sampler_reg, tc3_, 0u, metal::bias(2.0));
metal::float4 _e116 = a;
a = _e116 + _e115;
metal::float4 _e121 = image_cube_array.sample(sampler_reg, tc3_, 0);
metal::float4 _e122 = a;
a = _e122 + _e121;
metal::float4 _e127 = image_cube_array.sample(sampler_reg, tc3_, 0, metal::level(2.3));
metal::float4 _e128 = a;
a = _e128 + _e127;
metal::float4 _e134 = image_cube_array.sample(sampler_reg, tc3_, 0, metal::bias(2.0));
metal::float4 _e135 = a;
a = _e135 + _e134;
metal::float4 _e137 = a;
return texture_sampleOutput { _e137 };
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, metal::int2(3, 1));
metal::float4 _e20 = a;
a = _e20 + _e19;
metal::float4 _e24 = image_2d.sample(sampler_reg, tc, metal::level(2.3));
metal::float4 _e25 = a;
a = _e25 + _e24;
metal::float4 _e29 = image_2d.sample(sampler_reg, tc, metal::level(2.3), metal::int2(3, 1));
metal::float4 _e30 = a;
a = _e30 + _e29;
metal::float4 _e35 = image_2d.sample(sampler_reg, tc, metal::bias(2.0), metal::int2(3, 1));
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, metal::int2(3, 1));
metal::float4 _e48 = a;
a = _e48 + _e47;
metal::float4 _e53 = image_2d_array.sample(sampler_reg, tc, 0u, metal::level(2.3));
metal::float4 _e54 = a;
a = _e54 + _e53;
metal::float4 _e59 = image_2d_array.sample(sampler_reg, tc, 0u, metal::level(2.3), metal::int2(3, 1));
metal::float4 _e60 = a;
a = _e60 + _e59;
metal::float4 _e66 = image_2d_array.sample(sampler_reg, tc, 0u, metal::bias(2.0), metal::int2(3, 1));
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, metal::int2(3, 1));
metal::float4 _e79 = a;
a = _e79 + _e78;
metal::float4 _e84 = image_2d_array.sample(sampler_reg, tc, 0, metal::level(2.3));
metal::float4 _e85 = a;
a = _e85 + _e84;
metal::float4 _e90 = image_2d_array.sample(sampler_reg, tc, 0, metal::level(2.3), metal::int2(3, 1));
metal::float4 _e91 = a;
a = _e91 + _e90;
metal::float4 _e97 = image_2d_array.sample(sampler_reg, tc, 0, metal::bias(2.0), metal::int2(3, 1));
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.3));
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.3));
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 };
}

207
tests/out/spv/const-exprs.spvasm
View File

@ -1,130 +1,127 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 85
; Bound: 82
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %75 "main"
OpExecutionMode %75 LocalSize 1 1 1
OpDecorate %8 DescriptorSet 0
OpDecorate %8 Binding 0
OpDecorate %9 Block
OpMemberDecorate %9 0 Offset 0
OpDecorate %11 DescriptorSet 0
OpDecorate %11 Binding 1
OpDecorate %12 Block
OpMemberDecorate %12 0 Offset 0
OpEntryPoint GLCompute %72 "main"
OpExecutionMode %72 LocalSize 1 1 1
OpDecorate %7 DescriptorSet 0
OpDecorate %7 Binding 0
OpDecorate %8 Block
OpMemberDecorate %8 0 Offset 0
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 1
OpDecorate %11 Block
OpMemberDecorate %11 0 Offset 0
%2 = OpTypeVoid
%4 = OpTypeInt 32 1
%3 = OpTypeVector %4 4
%5 = OpTypeVector %4 2
%6 = OpConstant %4 4
%7 = OpConstant %4 8
%9 = OpTypeStruct %3
%10 = OpTypePointer StorageBuffer %9
%8 = OpVariable %10 StorageBuffer
%12 = OpTypeStruct %4
%13 = OpTypePointer StorageBuffer %12
%11 = OpVariable %13 StorageBuffer
%16 = OpTypeFunction %2
%17 = OpTypePointer StorageBuffer %3
%19 = OpTypeInt 32 0
%18 = OpConstant %19 0
%21 = OpConstant %4 1
%22 = OpConstant %4 2
%23 = OpConstantComposite %5 %21 %22
%24 = OpConstant %4 3
%25 = OpConstantComposite %5 %24 %6
%26 = OpConstantComposite %3 %6 %24 %22 %21
%30 = OpTypePointer StorageBuffer %4
%38 = OpConstant %4 6
%45 = OpConstant %4 30
%46 = OpConstant %4 70
%48 = OpTypePointer Function %4
%50 = OpConstantNull %4
%52 = OpConstantNull %4
%67 = OpConstant %4 -4
%68 = OpConstantComposite %3 %67 %67 %67 %67
%15 = OpFunction %2 None %16
%14 = OpLabel
%20 = OpAccessChain %17 %8 %18
OpBranch %27
%27 = OpLabel
OpStore %20 %26
%5 = OpConstant %4 4
%6 = OpConstant %4 8
%8 = OpTypeStruct %3
%9 = OpTypePointer StorageBuffer %8
%7 = OpVariable %9 StorageBuffer
%11 = OpTypeStruct %4
%12 = OpTypePointer StorageBuffer %11
%10 = OpVariable %12 StorageBuffer
%15 = OpTypeFunction %2
%16 = OpTypePointer StorageBuffer %3
%18 = OpTypeInt 32 0
%17 = OpConstant %18 0
%20 = OpConstant %4 1
%21 = OpConstant %4 2
%22 = OpConstant %4 3
%23 = OpConstantComposite %3 %5 %22 %21 %20
%27 = OpTypePointer StorageBuffer %4
%35 = OpConstant %4 6
%42 = OpConstant %4 30
%43 = OpConstant %4 70
%45 = OpTypePointer Function %4
%47 = OpConstantNull %4
%49 = OpConstantNull %4
%64 = OpConstant %4 -4
%65 = OpConstantComposite %3 %64 %64 %64 %64
%14 = OpFunction %2 None %15
%13 = OpLabel
%19 = OpAccessChain %16 %7 %17
OpBranch %24
%24 = OpLabel
OpStore %19 %23
OpReturn
OpFunctionEnd
%29 = OpFunction %2 None %16
%28 = OpLabel
%31 = OpAccessChain %30 %11 %18
OpBranch %32
%26 = OpFunction %2 None %15
%25 = OpLabel
%28 = OpAccessChain %27 %10 %17
OpBranch %29
%29 = OpLabel
%30 = OpLoad %4 %28
%31 = OpIAdd %4 %30 %21
OpStore %28 %31
OpReturn
OpFunctionEnd
%33 = OpFunction %2 None %15
%32 = OpLabel
%33 = OpLoad %4 %31
%34 = OpIAdd %4 %33 %22
OpStore %31 %34
%34 = OpAccessChain %27 %10 %17
OpBranch %36
%36 = OpLabel
%37 = OpLoad %4 %34
%38 = OpIAdd %4 %37 %35
OpStore %34 %38
OpReturn
OpFunctionEnd
%36 = OpFunction %2 None %16
%35 = OpLabel
%37 = OpAccessChain %30 %11 %18
OpBranch %39
%40 = OpFunction %2 None %15
%39 = OpLabel
%40 = OpLoad %4 %37
%41 = OpIAdd %4 %40 %38
OpStore %37 %41
%46 = OpVariable %45 Function %47
%50 = OpVariable %45 Function %43
%44 = OpVariable %45 Function %42
%48 = OpVariable %45 Function %49
%41 = OpAccessChain %16 %7 %17
OpBranch %51
%51 = OpLabel
%52 = OpLoad %4 %44
OpStore %46 %52
%53 = OpLoad %4 %46
OpStore %48 %53
%54 = OpLoad %4 %44
%55 = OpLoad %4 %46
%56 = OpLoad %4 %48
%57 = OpLoad %4 %50
%58 = OpCompositeConstruct %3 %54 %55 %56 %57
%59 = OpLoad %3 %41
%60 = OpIAdd %3 %59 %58
OpStore %41 %60
OpReturn
OpFunctionEnd
%43 = OpFunction %2 None %16
%42 = OpLabel
%49 = OpVariable %48 Function %50
%53 = OpVariable %48 Function %46
%47 = OpVariable %48 Function %45
%51 = OpVariable %48 Function %52
%44 = OpAccessChain %17 %8 %18
OpBranch %54
%54 = OpLabel
%55 = OpLoad %4 %47
OpStore %49 %55
%56 = OpLoad %4 %49
OpStore %51 %56
%57 = OpLoad %4 %47
%58 = OpLoad %4 %49
%59 = OpLoad %4 %51
%60 = OpLoad %4 %53
%61 = OpCompositeConstruct %3 %57 %58 %59 %60
%62 = OpLoad %3 %44
%63 = OpIAdd %3 %62 %61
OpStore %44 %63
%62 = OpFunction %2 None %15
%61 = OpLabel
%63 = OpAccessChain %16 %7 %17
OpBranch %66
%66 = OpLabel
OpStore %63 %65
OpReturn
OpFunctionEnd
%65 = OpFunction %2 None %16
%64 = OpLabel
%66 = OpAccessChain %17 %8 %18
OpBranch %69
%69 = OpLabel
OpStore %66 %68
OpReturn
OpFunctionEnd
%71 = OpFunction %2 None %16
%68 = OpFunction %2 None %15
%67 = OpLabel
%69 = OpAccessChain %16 %7 %17
OpBranch %70
%70 = OpLabel
%72 = OpAccessChain %17 %8 %18
OpBranch %73
%73 = OpLabel
OpStore %72 %68
OpStore %69 %65
OpReturn
OpFunctionEnd
%75 = OpFunction %2 None %16
%74 = OpLabel
%76 = OpAccessChain %17 %8 %18
%77 = OpAccessChain %30 %11 %18
OpBranch %78
%78 = OpLabel
%79 = OpFunctionCall %2 %15
%80 = OpFunctionCall %2 %29
%81 = OpFunctionCall %2 %36
%82 = OpFunctionCall %2 %43
%83 = OpFunctionCall %2 %65
%84 = OpFunctionCall %2 %71
%72 = OpFunction %2 None %15
%71 = OpLabel
%73 = OpAccessChain %16 %7 %17
%74 = OpAccessChain %27 %10 %17
OpBranch %75
%75 = OpLabel
%76 = OpFunctionCall %2 %14
%77 = OpFunctionCall %2 %26
%78 = OpFunctionCall %2 %33
%79 = OpFunctionCall %2 %40
%80 = OpFunctionCall %2 %62
%81 = OpFunctionCall %2 %68
OpReturn
OpFunctionEnd

382
tests/out/spv/debug-symbol-terrain.spvasm
View File

@ -360,10 +360,10 @@ OpName %52 "x"
OpName %53 "permute3"
OpName %66 "v"
OpName %67 "snoise2"
OpName %90 "i"
OpName %93 "i1"
OpName %95 "x12"
OpName %98 "m"
OpName %86 "i"
OpName %89 "i1"
OpName %91 "x12"
OpName %94 "m"
OpName %203 "p"
OpName %204 "fbm"
OpName %209 "x"
@ -547,34 +547,30 @@ OpDecorate %582 Location 0
%71 = OpConstant %5 -0.57735026
%72 = OpConstant %5 0.024390243
%73 = OpConstantComposite %7 %69 %70 %71 %72
%74 = OpConstantComposite %6 %70 %70
%75 = OpConstantComposite %6 %69 %69
%76 = OpConstant %5 0.0
%77 = OpConstantComposite %6 %56 %76
%78 = OpConstantComposite %6 %76 %56
%79 = OpConstantComposite %7 %69 %69 %71 %71
%80 = OpConstantComposite %6 %58 %58
%81 = OpConstant %5 0.5
%82 = OpConstantComposite %4 %81 %81 %81
%83 = OpConstantComposite %4 %76 %76 %76
%84 = OpConstant %5 2.0
%85 = OpConstantComposite %4 %72 %72 %72
%86 = OpConstant %5 1.7928429
%87 = OpConstant %5 0.85373473
%88 = OpConstantComposite %4 %86 %86 %86
%89 = OpConstant %5 130.0
%91 = OpTypePointer Function %6
%92 = OpConstantNull %6
%94 = OpConstantNull %6
%96 = OpTypePointer Function %7
%97 = OpConstantNull %7
%99 = OpTypePointer Function %4
%100 = OpConstantNull %4
%114 = OpTypeBool
%117 = OpTypeVector %114 2
%126 = OpTypePointer Function %5
%127 = OpConstant %8 1
%136 = OpConstant %8 0
%74 = OpConstant %5 0.0
%75 = OpConstantComposite %6 %56 %74
%76 = OpConstantComposite %6 %74 %56
%77 = OpConstantComposite %6 %58 %58
%78 = OpConstant %5 0.5
%79 = OpConstantComposite %4 %78 %78 %78
%80 = OpConstantComposite %4 %74 %74 %74
%81 = OpConstant %5 2.0
%82 = OpConstant %5 1.7928429
%83 = OpConstant %5 0.85373473
%84 = OpConstantComposite %4 %82 %82 %82
%85 = OpConstant %5 130.0
%87 = OpTypePointer Function %6
%88 = OpConstantNull %6
%90 = OpConstantNull %6
%92 = OpTypePointer Function %7
%93 = OpConstantNull %7
%95 = OpTypePointer Function %4
%96 = OpConstantNull %4
%112 = OpTypeBool
%115 = OpTypeVector %112 2
%125 = OpTypePointer Function %5
%126 = OpConstant %8 1
%135 = OpConstant %8 0
%205 = OpConstant %8 5
%206 = OpConstant %5 0.01
%207 = OpConstant %5 100.0
@ -585,11 +581,11 @@ OpDecorate %582 Location 0
%258 = OpTypeFunction %4 %6 %6
%271 = OpTypeFunction %14 %6 %6
%272 = OpConstant %5 0.1
%273 = OpConstantComposite %6 %272 %76
%274 = OpConstantComposite %6 %76 %272
%273 = OpConstantComposite %6 %272 %74
%274 = OpConstantComposite %6 %74 %272
%275 = OpConstant %5 -0.1
%276 = OpConstantComposite %6 %275 %76
%277 = OpConstantComposite %6 %76 %275
%276 = OpConstantComposite %6 %275 %74
%277 = OpConstantComposite %6 %74 %275
%304 = OpTypeFunction %6 %8 %10 %11
%321 = OpTypeFunction %4 %6
%322 = OpConstant %5 23.0
@ -650,7 +646,7 @@ OpDecorate %582 Location 0
%585 = OpTypePointer Uniform %25
%587 = OpConstantComposite %4 %272 %272 %272
%588 = OpConstant %5 0.7
%589 = OpConstantComposite %4 %81 %272 %588
%589 = OpConstantComposite %4 %78 %272 %588
%590 = OpConstant %5 0.2
%591 = OpConstantComposite %4 %590 %590 %590
%593 = OpConstantNull %4
@ -673,163 +669,167 @@ OpFunctionEnd
%67 = OpFunction %5 None %68
%66 = OpFunctionParameter %6
%65 = OpLabel
%93 = OpVariable %91 Function %94
%98 = OpVariable %99 Function %100
%90 = OpVariable %91 Function %92
%95 = OpVariable %96 Function %97
OpBranch %101
%101 = OpLabel
%89 = OpVariable %87 Function %90
%94 = OpVariable %95 Function %96
%86 = OpVariable %87 Function %88
%91 = OpVariable %92 Function %93
OpBranch %97
%97 = OpLabel
OpLine %3 13 13
OpLine %3 14 24
%102 = OpDot %5 %66 %74
%103 = OpCompositeConstruct %6 %102 %102
%104 = OpFAdd %6 %66 %103
%105 = OpExtInst %6 %1 Floor %104
%98 = OpVectorShuffle %6 %73 %73 1 1
%99 = OpDot %5 %66 %98
%100 = OpCompositeConstruct %6 %99 %99
%101 = OpFAdd %6 %66 %100
%102 = OpExtInst %6 %1 Floor %101
OpLine %3 14 5
OpStore %90 %105
OpStore %86 %102
OpLine %3 15 14
%106 = OpLoad %6 %90
%107 = OpFSub %6 %66 %106
%108 = OpLoad %6 %90
%109 = OpDot %5 %108 %75
%110 = OpCompositeConstruct %6 %109 %109
%111 = OpFAdd %6 %107 %110
%103 = OpLoad %6 %86
%104 = OpFSub %6 %66 %103
%105 = OpLoad %6 %86
%106 = OpVectorShuffle %6 %73 %73 0 0
%107 = OpDot %5 %105 %106
%108 = OpCompositeConstruct %6 %107 %107
%109 = OpFAdd %6 %104 %108
OpLine %3 17 32
OpLine %3 17 25
%112 = OpCompositeExtract %5 %111 0
%113 = OpCompositeExtract %5 %111 1
%115 = OpFOrdLessThan %114 %112 %113
%118 = OpCompositeConstruct %117 %115 %115
%116 = OpSelect %6 %118 %78 %77
%110 = OpCompositeExtract %5 %109 0
%111 = OpCompositeExtract %5 %109 1
%113 = OpFOrdLessThan %112 %110 %111
%116 = OpCompositeConstruct %115 %113 %113
%114 = OpSelect %6 %116 %76 %75
OpLine %3 17 5
OpStore %93 %116
OpStore %89 %114
OpLine %3 18 26
%119 = OpVectorShuffle %7 %111 %111 0 1 0 1
%120 = OpFAdd %7 %119 %79
%121 = OpLoad %6 %93
%117 = OpVectorShuffle %7 %109 %109 0 1 0 1
%118 = OpVectorShuffle %7 %73 %73 0 0 2 2
%119 = OpFAdd %7 %117 %118
%120 = OpLoad %6 %89
OpLine %3 18 26
%122 = OpCompositeConstruct %7 %121 %76 %76
%123 = OpFSub %7 %120 %122
%121 = OpCompositeConstruct %7 %120 %74 %74
%122 = OpFSub %7 %119 %121
OpLine %3 18 5
OpStore %95 %123
OpStore %91 %122
OpLine %3 1 1
%124 = OpLoad %6 %90
%123 = OpLoad %6 %86
OpLine %3 19 9
%125 = OpFRem %6 %124 %80
%124 = OpFRem %6 %123 %77
OpLine %3 19 5
OpStore %90 %125
OpStore %86 %124
OpLine %3 20 31
%128 = OpAccessChain %126 %90 %127
%129 = OpLoad %5 %128
%127 = OpAccessChain %125 %86 %126
%128 = OpLoad %5 %127
OpLine %3 20 51
%130 = OpAccessChain %126 %93 %127
%131 = OpLoad %5 %130
%129 = OpAccessChain %125 %89 %126
%130 = OpLoad %5 %129
OpLine %3 20 31
%132 = OpCompositeConstruct %4 %76 %131 %56
%133 = OpCompositeConstruct %4 %129 %129 %129
%134 = OpFAdd %4 %133 %132
%131 = OpCompositeConstruct %4 %74 %130 %56
%132 = OpCompositeConstruct %4 %128 %128 %128
%133 = OpFAdd %4 %132 %131
OpLine %3 20 22
%135 = OpFunctionCall %4 %53 %134
%134 = OpFunctionCall %4 %53 %133
OpLine %3 20 22
%137 = OpAccessChain %126 %90 %136
%138 = OpLoad %5 %137
%139 = OpCompositeConstruct %4 %138 %138 %138
%140 = OpFAdd %4 %135 %139
%136 = OpAccessChain %125 %86 %135
%137 = OpLoad %5 %136
%138 = OpCompositeConstruct %4 %137 %137 %137
%139 = OpFAdd %4 %134 %138
OpLine %3 20 84
%141 = OpAccessChain %126 %93 %136
%142 = OpLoad %5 %141
%140 = OpAccessChain %125 %89 %135
%141 = OpLoad %5 %140
OpLine %3 20 22
%143 = OpCompositeConstruct %4 %76 %142 %56
%144 = OpFAdd %4 %140 %143
%142 = OpCompositeConstruct %4 %74 %141 %56
%143 = OpFAdd %4 %139 %142
OpLine %3 20 13
%145 = OpFunctionCall %4 %53 %144
%144 = OpFunctionCall %4 %53 %143
OpLine %3 21 28
%146 = OpDot %5 %111 %111
%147 = OpLoad %7 %95
%148 = OpVectorShuffle %6 %147 %147 0 1
%149 = OpLoad %7 %95
%150 = OpVectorShuffle %6 %149 %149 0 1
%151 = OpDot %5 %148 %150
%152 = OpLoad %7 %95
%153 = OpVectorShuffle %6 %152 %152 2 3
%154 = OpLoad %7 %95
%155 = OpVectorShuffle %6 %154 %154 2 3
%156 = OpDot %5 %153 %155
%157 = OpCompositeConstruct %4 %146 %151 %156
%158 = OpFSub %4 %82 %157
%145 = OpDot %5 %109 %109
%146 = OpLoad %7 %91
%147 = OpVectorShuffle %6 %146 %146 0 1
%148 = OpLoad %7 %91
%149 = OpVectorShuffle %6 %148 %148 0 1
%150 = OpDot %5 %147 %149
%151 = OpLoad %7 %91
%152 = OpVectorShuffle %6 %151 %151 2 3
%153 = OpLoad %7 %91
%154 = OpVectorShuffle %6 %153 %153 2 3
%155 = OpDot %5 %152 %154
%156 = OpCompositeConstruct %4 %145 %150 %155
%157 = OpFSub %4 %79 %156
OpLine %3 21 24
%159 = OpExtInst %4 %1 FMax %158 %83
%158 = OpExtInst %4 %1 FMax %157 %80
OpLine %3 21 5
OpStore %98 %159
OpStore %94 %158
OpLine %3 22 9
%160 = OpLoad %4 %98
%161 = OpLoad %4 %98
%162 = OpFMul %4 %160 %161
%159 = OpLoad %4 %94
%160 = OpLoad %4 %94
%161 = OpFMul %4 %159 %160
OpLine %3 22 5
OpStore %98 %162
OpStore %94 %161
OpLine %3 23 9
%163 = OpLoad %4 %98
%164 = OpLoad %4 %98
%165 = OpFMul %4 %163 %164
%162 = OpLoad %4 %94
%163 = OpLoad %4 %94
%164 = OpFMul %4 %162 %163
OpLine %3 23 5
OpStore %98 %165
OpStore %94 %164
OpLine %3 24 13
%166 = OpFMul %4 %145 %85
%165 = OpVectorShuffle %4 %73 %73 3 3 3
%166 = OpFMul %4 %144 %165
%167 = OpExtInst %4 %1 Fract %166
%168 = OpVectorTimesScalar %4 %167 %84
%168 = OpVectorTimesScalar %4 %167 %81
OpLine %3 24 13
%169 = OpFSub %4 %168 %57
OpLine %3 25 13
%170 = OpExtInst %4 %1 FAbs %169
OpLine %3 25 13
%171 = OpFSub %4 %170 %82
%171 = OpFSub %4 %170 %79
OpLine %3 26 14
%172 = OpFAdd %4 %169 %82
%172 = OpFAdd %4 %169 %79
%173 = OpExtInst %4 %1 Floor %172
OpLine %3 27 14
%174 = OpFSub %4 %169 %173
OpLine %3 1 1
%175 = OpLoad %4 %98
%175 = OpLoad %4 %94
OpLine %3 28 9
%176 = OpFMul %4 %174 %174
%177 = OpFMul %4 %171 %171
%178 = OpFAdd %4 %176 %177
%179 = OpVectorTimesScalar %4 %178 %87
%180 = OpFSub %4 %88 %179
%179 = OpVectorTimesScalar %4 %178 %83
%180 = OpFSub %4 %84 %179
%181 = OpFMul %4 %175 %180
OpLine %3 28 5
OpStore %98 %181
OpStore %94 %181
OpLine %3 29 13
%182 = OpCompositeExtract %5 %174 0
%183 = OpCompositeExtract %5 %111 0
%183 = OpCompositeExtract %5 %109 0
%184 = OpFMul %5 %182 %183
%185 = OpCompositeExtract %5 %171 0
%186 = OpCompositeExtract %5 %111 1
%186 = OpCompositeExtract %5 %109 1
%187 = OpFMul %5 %185 %186
%188 = OpFAdd %5 %184 %187
%189 = OpVectorShuffle %6 %174 %174 1 2
%190 = OpLoad %7 %95
%190 = OpLoad %7 %91
%191 = OpVectorShuffle %6 %190 %190 0 2
%192 = OpFMul %6 %189 %191
%193 = OpVectorShuffle %6 %171 %171 1 2
%194 = OpLoad %7 %95
%194 = OpLoad %7 %91
%195 = OpVectorShuffle %6 %194 %194 1 3
%196 = OpFMul %6 %193 %195
%197 = OpFAdd %6 %192 %196
%198 = OpCompositeConstruct %4 %188 %197
OpLine %3 30 12
%199 = OpLoad %4 %98
%199 = OpLoad %4 %94
%200 = OpDot %5 %199 %198
%201 = OpFMul %5 %89 %200
%201 = OpFMul %5 %85 %200
OpReturnValue %201
OpFunctionEnd
%204 = OpFunction %5 None %68
%203 = OpFunctionParameter %6
%202 = OpLabel
%211 = OpVariable %212 Function %76
%214 = OpVariable %215 Function %136
%209 = OpVariable %91 Function %210
%213 = OpVariable %212 Function %81
%211 = OpVariable %212 Function %74
%214 = OpVariable %215 Function %135
%209 = OpVariable %87 Function %210
%213 = OpVariable %212 Function %78
OpBranch %216
%216 = OpLabel
OpLine %3 36 13
@ -838,9 +838,9 @@ OpLine %3 36 5
OpStore %209 %217
OpLine %3 39 17
OpLine %3 40 24
%218 = OpExtInst %5 %1 Cos %81
%218 = OpExtInst %5 %1 Cos %78
OpLine %3 40 14
%219 = OpExtInst %5 %1 Sin %81
%219 = OpExtInst %5 %1 Sin %78
%220 = OpCompositeConstruct %6 %218 %219
OpLine %3 41 15
%221 = OpCompositeExtract %5 %220 0
@ -859,7 +859,7 @@ OpBranch %231
%231 = OpLabel
OpLine %3 43 22
%233 = OpLoad %8 %214
%234 = OpULessThan %114 %233 %205
%234 = OpULessThan %112 %233 %205
OpLine %3 43 21
OpSelectionMerge %235 None
OpBranchConditional %234 %235 %236
@ -883,14 +883,14 @@ OpLine %3 45 13
%245 = OpLoad %6 %209
%246 = OpMatrixTimesVector %6 %228 %245
OpLine %3 45 13
%247 = OpVectorTimesScalar %6 %246 %84
%247 = OpVectorTimesScalar %6 %246 %81
%248 = OpFAdd %6 %247 %208
OpLine %3 45 9
OpStore %209 %248
OpLine %3 1 1
%249 = OpLoad %5 %213
OpLine %3 46 13
%250 = OpFMul %5 %249 %81
%250 = OpFMul %5 %249 %78
OpLine %3 46 9
OpStore %213 %250
OpBranch %238
@ -900,7 +900,7 @@ OpBranch %232
OpLine %3 1 1
%251 = OpLoad %8 %214
OpLine %3 43 43
%252 = OpIAdd %8 %251 %127
%252 = OpIAdd %8 %251 %126
OpLine %3 43 39
OpStore %214 %252
OpBranch %229
@ -968,7 +968,7 @@ OpLine %3 92 14
OpLine %3 94 14
%296 = OpFAdd %4 %293 %295
OpLine %3 94 13
%297 = OpVectorTimesScalar %4 %296 %81
%297 = OpVectorTimesScalar %4 %296 %78
OpLine %3 96 12
%298 = OpCompositeConstruct %14 %279 %297
OpReturnValue %298
@ -984,12 +984,12 @@ OpLine %3 101 9
%306 = OpConvertUToF %5 %300
%307 = OpCompositeExtract %8 %301 0
OpLine %3 101 9
%308 = OpIAdd %8 %307 %127
%308 = OpIAdd %8 %307 %126
%309 = OpConvertUToF %5 %308
%310 = OpFRem %5 %306 %309
%311 = OpCompositeExtract %8 %301 0
OpLine %3 100 12
%312 = OpIAdd %8 %311 %127
%312 = OpIAdd %8 %311 %126
%313 = OpUDiv %8 %300 %312
%314 = OpConvertUToF %5 %313
%315 = OpCompositeConstruct %6 %310 %314
@ -1005,41 +1005,41 @@ OpBranch %328
OpLine %3 270 9
%329 = OpFunctionCall %5 %67 %319
OpLine %3 270 9
%330 = OpFMul %5 %329 %81
%330 = OpFMul %5 %329 %78
OpLine %3 270 9
%331 = OpFAdd %5 %330 %81
%331 = OpFAdd %5 %330 %78
OpLine %3 271 17
%332 = OpFAdd %6 %319 %324
OpLine %3 271 9
%333 = OpFunctionCall %5 %67 %332
OpLine %3 271 9
%334 = OpFMul %5 %333 %81
%334 = OpFMul %5 %333 %78
OpLine %3 271 9
%335 = OpFAdd %5 %334 %81
%335 = OpFAdd %5 %334 %78
OpLine %3 272 17
%336 = OpFAdd %6 %319 %327
OpLine %3 272 9
%337 = OpFunctionCall %5 %67 %336
OpLine %3 272 9
%338 = OpFMul %5 %337 %81
%338 = OpFMul %5 %337 %78
OpLine %3 269 12
%339 = OpFAdd %5 %338 %81
%339 = OpFAdd %5 %338 %78
%340 = OpCompositeConstruct %4 %331 %335 %339
OpReturnValue %340
OpFunctionEnd
%345 = OpFunction %2 None %346
%341 = OpLabel
%344 = OpLoad %19 %342
%348 = OpAccessChain %347 %29 %136
%348 = OpAccessChain %347 %29 %135
OpBranch %353
%353 = OpLabel
OpLine %3 111 22
%354 = OpCompositeExtract %8 %344 0
OpLine %3 113 36
%356 = OpAccessChain %355 %348 %136
%356 = OpAccessChain %355 %348 %135
%357 = OpLoad %10 %356
OpLine %3 113 59
%359 = OpAccessChain %358 %348 %127
%359 = OpAccessChain %358 %348 %126
%360 = OpLoad %11 %359
OpLine %3 113 13
%361 = OpFunctionCall %6 %303 %354 %357 %360
@ -1050,22 +1050,22 @@ OpLine %3 115 51
OpLine %3 115 33
%367 = OpFunctionCall %14 %270 %361 %366
OpLine %3 115 5
%368 = OpAccessChain %363 %32 %136 %354
%368 = OpAccessChain %363 %32 %135 %354
OpStore %368 %367
OpLine %3 118 23
%369 = OpCompositeExtract %8 %344 0
OpLine %3 118 23
%370 = OpIMul %8 %369 %349
OpLine %3 120 25
%372 = OpAccessChain %371 %348 %136 %136
%372 = OpAccessChain %371 %348 %135 %135
%373 = OpLoad %8 %372
OpLine %3 120 25
%374 = OpAccessChain %371 %348 %136 %127
%374 = OpAccessChain %371 %348 %135 %126
%375 = OpLoad %8 %374
%376 = OpIMul %8 %373 %375
OpLine %3 120 9
%377 = OpIMul %8 %376 %349
%378 = OpUGreaterThanEqual %114 %370 %377
%378 = OpUGreaterThanEqual %112 %370 %377
OpLine %3 120 5
OpSelectionMerge %379 None
OpBranchConditional %378 %380 %379
@ -1075,60 +1075,60 @@ OpReturn
OpLine %3 122 28
%381 = OpCompositeExtract %8 %344 0
OpLine %3 122 15
%382 = OpAccessChain %371 %348 %136 %136
%382 = OpAccessChain %371 %348 %135 %135
%383 = OpLoad %8 %382
%384 = OpUDiv %8 %381 %383
%385 = OpIAdd %8 %354 %384
OpLine %3 123 15
%386 = OpIAdd %8 %385 %127
%386 = OpIAdd %8 %385 %126
OpLine %3 124 15
%387 = OpAccessChain %371 %348 %136 %136
%387 = OpAccessChain %371 %348 %135 %135
%388 = OpLoad %8 %387
%389 = OpIAdd %8 %385 %388
OpLine %3 124 15
%390 = OpIAdd %8 %389 %127
%390 = OpIAdd %8 %389 %126
OpLine %3 125 15
%391 = OpIAdd %8 %390 %127
%391 = OpIAdd %8 %390 %126
OpLine %3 127 5
OpLine %3 127 5
%394 = OpAccessChain %393 %34 %136 %370
%394 = OpAccessChain %393 %34 %135 %370
OpStore %394 %385
OpLine %3 128 5
OpLine %3 128 5
%395 = OpIAdd %8 %370 %127
%395 = OpIAdd %8 %370 %126
OpLine %3 128 5
%396 = OpAccessChain %393 %34 %136 %395
%396 = OpAccessChain %393 %34 %135 %395
OpStore %396 %390
OpLine %3 129 5
OpLine %3 129 5
%397 = OpIAdd %8 %370 %350
OpLine %3 129 5
%398 = OpAccessChain %393 %34 %136 %397
%398 = OpAccessChain %393 %34 %135 %397
OpStore %398 %391
OpLine %3 130 5
OpLine %3 130 5
%399 = OpIAdd %8 %370 %351
OpLine %3 130 5
%400 = OpAccessChain %393 %34 %136 %399
%400 = OpAccessChain %393 %34 %135 %399
OpStore %400 %385
OpLine %3 131 5
OpLine %3 131 5
%401 = OpIAdd %8 %370 %352
OpLine %3 131 5
%402 = OpAccessChain %393 %34 %136 %401
%402 = OpAccessChain %393 %34 %135 %401
OpStore %402 %391
OpLine %3 132 5
OpLine %3 132 5
%403 = OpIAdd %8 %370 %205
OpLine %3 132 5
%404 = OpAccessChain %393 %34 %136 %403
%404 = OpAccessChain %393 %34 %135 %403
OpStore %404 %386
OpReturn
OpFunctionEnd
%415 = OpFunction %2 None %346
%405 = OpLabel
%408 = OpLoad %8 %406
%417 = OpAccessChain %416 %36 %136
%417 = OpAccessChain %416 %36 %135
OpBranch %420
%420 = OpLabel
OpLine %3 161 19
@ -1139,7 +1139,7 @@ OpLine %3 161 13
%423 = OpUMod %8 %422 %350
%424 = OpConvertUToF %5 %423
OpLine %3 162 19
%425 = OpIAdd %8 %408 %127
%425 = OpIAdd %8 %408 %126
OpLine %3 162 18
%426 = OpUDiv %8 %425 %351
OpLine %3 162 13
@ -1148,10 +1148,10 @@ OpLine %3 162 13
OpLine %3 163 14
%429 = OpCompositeConstruct %6 %424 %428
OpLine %3 165 30
%430 = OpVectorTimesScalar %6 %429 %84
%430 = OpVectorTimesScalar %6 %429 %81
%431 = OpFAdd %6 %419 %430
OpLine %3 165 20
%432 = OpCompositeConstruct %7 %431 %76 %56
%432 = OpCompositeConstruct %7 %431 %74 %56
OpLine %3 168 21
%433 = OpCompositeExtract %5 %429 0
OpLine %3 168 21
@ -1183,13 +1183,13 @@ OpReturn
OpFunctionEnd
%465 = OpFunction %2 None %346
%453 = OpLabel
%468 = OpVariable %212 Function %76
%469 = OpVariable %215 Function %136
%468 = OpVariable %212 Function %74
%469 = OpVariable %215 Function %135
%456 = OpLoad %8 %455
%459 = OpLoad %7 %457
%462 = OpLoad %6 %460
%454 = OpCompositeConstruct %21 %456 %459 %462
%466 = OpAccessChain %416 %36 %136
%466 = OpAccessChain %416 %36 %135
OpBranch %470
%470 = OpLabel
OpLine %3 181 17
@ -1226,10 +1226,10 @@ OpLine %3 182 22
OpLine %3 183 22
%495 = OpUMod %8 %490 %349
OpLine %3 185 36
%496 = OpAccessChain %355 %466 %136
%496 = OpAccessChain %355 %466 %135
%497 = OpLoad %10 %496
OpLine %3 185 57
%498 = OpAccessChain %358 %466 %127
%498 = OpAccessChain %358 %466 %126
%499 = OpLoad %11 %498
OpLine %3 185 13
%500 = OpFunctionCall %6 %303 %494 %497 %499
@ -1287,20 +1287,20 @@ OpBranch %504
OpBranch %504
%504 = OpLabel
OpLine %3 200 15
%524 = OpAccessChain %371 %466 %136 %136
%524 = OpAccessChain %371 %466 %135 %135
%525 = OpLoad %8 %524
%526 = OpUDiv %8 %494 %525
%527 = OpIAdd %8 %494 %526
OpLine %3 201 15
%528 = OpIAdd %8 %527 %127
%528 = OpIAdd %8 %527 %126
OpLine %3 202 15
%529 = OpAccessChain %371 %466 %136 %136
%529 = OpAccessChain %371 %466 %135 %135
%530 = OpLoad %8 %529
%531 = OpIAdd %8 %527 %530
OpLine %3 202 15
%532 = OpIAdd %8 %531 %127
%532 = OpIAdd %8 %531 %126
OpLine %3 203 15
%533 = OpIAdd %8 %532 %127
%533 = OpIAdd %8 %532 %126
OpLine %3 206 5
OpSelectionMerge %534 None
OpSwitch %495 %539 0 %535 3 %535 2 %536 4 %536 1 %537 5 %538
@ -1344,11 +1344,11 @@ OpFunctionEnd
%551 = OpLoad %4 %549
%553 = OpLoad %4 %552
%548 = OpCompositeConstruct %14 %551 %553
%560 = OpAccessChain %559 %39 %136
%560 = OpAccessChain %559 %39 %135
OpBranch %561
%561 = OpLabel
OpLine %3 254 25
%563 = OpAccessChain %562 %560 %127
%563 = OpAccessChain %562 %560 %126
%564 = OpLoad %23 %563
%565 = OpCompositeExtract %4 %548 0
OpLine %3 254 25
@ -1369,30 +1369,30 @@ OpReturn
OpFunctionEnd
%583 = OpFunction %2 None %346
%574 = OpLabel
%592 = OpVariable %99 Function %593
%592 = OpVariable %95 Function %593
%577 = OpLoad %7 %576
%579 = OpLoad %4 %578
%581 = OpLoad %4 %580
%575 = OpCompositeConstruct %26 %577 %579 %581
%584 = OpAccessChain %559 %39 %136
%586 = OpAccessChain %585 %42 %136
%584 = OpAccessChain %559 %39 %135
%586 = OpAccessChain %585 %42 %135
OpBranch %594
%594 = OpLabel
OpLine %3 278 28
OpLine %3 278 17
%595 = OpCompositeExtract %4 %575 2
%596 = OpExtInst %4 %1 Fract %595
%597 = OpExtInst %4 %1 SmoothStep %83 %587 %596
%597 = OpExtInst %4 %1 SmoothStep %80 %587 %596
OpLine %3 278 5
OpStore %592 %597
OpLine %3 279 17
OpLine %3 279 13
%598 = OpAccessChain %126 %592 %136
%598 = OpAccessChain %125 %592 %135
%599 = OpLoad %5 %598
%600 = OpAccessChain %126 %592 %127
%600 = OpAccessChain %125 %592 %126
%601 = OpLoad %5 %600
%602 = OpFMul %5 %599 %601
%603 = OpAccessChain %126 %592 %350
%603 = OpAccessChain %125 %592 %350
%604 = OpLoad %5 %603
%605 = OpFMul %5 %602 %604
%606 = OpCompositeConstruct %4 %605 %605 %605
@ -1400,17 +1400,17 @@ OpLine %3 279 13
OpLine %3 279 5
OpStore %592 %607
OpLine %3 282 25
%609 = OpAccessChain %608 %586 %127
%609 = OpAccessChain %608 %586 %126
%610 = OpLoad %4 %609
%611 = OpVectorTimesScalar %4 %610 %272
OpLine %3 284 21
%612 = OpAccessChain %608 %586 %136
%612 = OpAccessChain %608 %586 %135
%613 = OpLoad %4 %612
%614 = OpCompositeExtract %4 %575 2
%615 = OpFSub %4 %613 %614
%616 = OpExtInst %4 %1 Normalize %615
OpLine %3 285 20
%618 = OpAccessChain %617 %584 %136
%618 = OpAccessChain %617 %584 %135
%619 = OpLoad %7 %618
%620 = OpVectorShuffle %4 %619 %619 0 1 2
%621 = OpCompositeExtract %4 %575 2
@ -1423,20 +1423,20 @@ OpLine %3 288 32
%626 = OpCompositeExtract %4 %575 1
%627 = OpDot %5 %626 %616
OpLine %3 288 28
%628 = OpExtInst %5 %1 FMax %627 %76
%628 = OpExtInst %5 %1 FMax %627 %74
OpLine %3 289 25
%629 = OpAccessChain %608 %586 %127
%629 = OpAccessChain %608 %586 %126
%630 = OpLoad %4 %629
%631 = OpVectorTimesScalar %4 %630 %628
OpLine %3 291 37
%632 = OpCompositeExtract %4 %575 1
%633 = OpDot %5 %632 %625
OpLine %3 291 33
%634 = OpExtInst %5 %1 FMax %633 %76
%634 = OpExtInst %5 %1 FMax %633 %74
OpLine %3 291 29
%635 = OpExtInst %5 %1 Pow %634 %323
OpLine %3 292 26
%636 = OpAccessChain %608 %586 %127
%636 = OpAccessChain %608 %586 %126
%637 = OpLoad %4 %636
%638 = OpVectorTimesScalar %4 %637 %635
OpLine %3 294 18

561
tests/out/spv/image.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 517
; Bound: 518
OpCapability Shader
OpCapability Image1D
OpCapability Sampled1D
@ -14,15 +14,15 @@ OpEntryPoint GLCompute %169 "depth_load" %167
OpEntryPoint Vertex %189 "queries" %187
OpEntryPoint Vertex %241 "levels_queries" %240
OpEntryPoint Fragment %270 "texture_sample" %269
OpEntryPoint Fragment %416 "texture_sample_comparison" %414
OpEntryPoint Fragment %472 "gather" %471
OpEntryPoint Fragment %506 "depth_no_comparison" %505
OpEntryPoint Fragment %417 "texture_sample_comparison" %415
OpEntryPoint Fragment %473 "gather" %472
OpEntryPoint Fragment %507 "depth_no_comparison" %506
OpExecutionMode %78 LocalSize 16 1 1
OpExecutionMode %169 LocalSize 16 1 1
OpExecutionMode %270 OriginUpperLeft
OpExecutionMode %416 OriginUpperLeft
OpExecutionMode %472 OriginUpperLeft
OpExecutionMode %506 OriginUpperLeft
OpExecutionMode %417 OriginUpperLeft
OpExecutionMode %473 OriginUpperLeft
OpExecutionMode %507 OriginUpperLeft
OpName %31 "image_mipmapped_src"
OpName %33 "image_multisampled_src"
OpName %35 "image_depth_multisampled_src"
@ -53,10 +53,10 @@ OpName %189 "queries"
OpName %241 "levels_queries"
OpName %270 "texture_sample"
OpName %284 "a"
OpName %416 "texture_sample_comparison"
OpName %421 "a"
OpName %472 "gather"
OpName %506 "depth_no_comparison"
OpName %417 "texture_sample_comparison"
OpName %422 "a"
OpName %473 "gather"
OpName %507 "depth_no_comparison"
OpDecorate %31 DescriptorSet 0
OpDecorate %31 Binding 0
OpDecorate %33 DescriptorSet 0
@ -109,9 +109,9 @@ OpDecorate %167 BuiltIn LocalInvocationId
OpDecorate %187 BuiltIn Position
OpDecorate %240 BuiltIn Position
OpDecorate %269 Location 0
OpDecorate %414 Location 0
OpDecorate %471 Location 0
OpDecorate %505 Location 0
OpDecorate %415 Location 0
OpDecorate %472 Location 0
OpDecorate %506 Location 0
%2 = OpTypeVoid
%4 = OpTypeInt 32 0
%3 = OpTypeImage %4 2D 0 0 0 1 Unknown
@ -209,25 +209,25 @@ OpDecorate %505 Location 0
%283 = OpConstant %14 0
%285 = OpTypePointer Function %23
%286 = OpConstantNull %23
%288 = OpTypeSampledImage %15
%293 = OpTypeSampledImage %16
%314 = OpTypeSampledImage %18
%375 = OpTypeSampledImage %20
%415 = OpTypePointer Output %7
%414 = OpVariable %415 Output
%422 = OpTypePointer Function %7
%423 = OpConstantNull %7
%425 = OpTypeSampledImage %25
%430 = OpTypeSampledImage %26
%443 = OpTypeSampledImage %27
%450 = OpConstant %7 0.0
%471 = OpVariable %188 Output
%482 = OpConstant %4 1
%485 = OpConstant %4 3
%490 = OpTypeSampledImage %3
%493 = OpTypeVector %14 4
%494 = OpTypeSampledImage %17
%505 = OpVariable %188 Output
%289 = OpTypeSampledImage %15
%294 = OpTypeSampledImage %16
%315 = OpTypeSampledImage %18
%376 = OpTypeSampledImage %20
%416 = OpTypePointer Output %7
%415 = OpVariable %416 Output
%423 = OpTypePointer Function %7
%424 = OpConstantNull %7
%426 = OpTypeSampledImage %25
%431 = OpTypeSampledImage %26
%444 = OpTypeSampledImage %27
%451 = OpConstant %7 0.0
%472 = OpVariable %188 Output
%483 = OpConstant %4 1
%486 = OpConstant %4 3
%491 = OpTypeSampledImage %3
%494 = OpTypeVector %14 4
%495 = OpTypeSampledImage %17
%506 = OpVariable %188 Output
%78 = OpFunction %2 None %79
%74 = OpLabel
%77 = OpLoad %12 %75
@ -434,258 +434,259 @@ OpFunctionEnd
%275 = OpLoad %24 %64
OpBranch %287
%287 = OpLabel
%289 = OpSampledImage %288 %271 %275
%290 = OpImageSampleImplicitLod %23 %289 %276
%291 = OpLoad %23 %284
%292 = OpFAdd %23 %291 %290
OpStore %284 %292
%294 = OpSampledImage %293 %272 %275
%295 = OpImageSampleImplicitLod %23 %294 %278
%296 = OpLoad %23 %284
%297 = OpFAdd %23 %296 %295
OpStore %284 %297
%298 = OpSampledImage %293 %272 %275
%299 = OpImageSampleImplicitLod %23 %298 %278 ConstOffset %30
%300 = OpLoad %23 %284
%301 = OpFAdd %23 %300 %299
OpStore %284 %301
%302 = OpSampledImage %293 %272 %275
%303 = OpImageSampleExplicitLod %23 %302 %278 Lod %281
%304 = OpLoad %23 %284
%305 = OpFAdd %23 %304 %303
OpStore %284 %305
%306 = OpSampledImage %293 %272 %275
%307 = OpImageSampleExplicitLod %23 %306 %278 Lod|ConstOffset %281 %30
%308 = OpLoad %23 %284
%309 = OpFAdd %23 %308 %307
OpStore %284 %309
%310 = OpSampledImage %293 %272 %275
%311 = OpImageSampleImplicitLod %23 %310 %278 Bias|ConstOffset %282 %30
%312 = OpLoad %23 %284
%313 = OpFAdd %23 %312 %311
OpStore %284 %313
%315 = OpConvertUToF %7 %198
%316 = OpCompositeConstruct %279 %278 %315
%317 = OpSampledImage %314 %273 %275
%318 = OpImageSampleImplicitLod %23 %317 %316
%319 = OpLoad %23 %284
%320 = OpFAdd %23 %319 %318
OpStore %284 %320
%321 = OpConvertUToF %7 %198
%322 = OpCompositeConstruct %279 %278 %321
%323 = OpSampledImage %314 %273 %275
%324 = OpImageSampleImplicitLod %23 %323 %322 ConstOffset %30
%325 = OpLoad %23 %284
%326 = OpFAdd %23 %325 %324
OpStore %284 %326
%327 = OpConvertUToF %7 %198
%328 = OpCompositeConstruct %279 %278 %327
%329 = OpSampledImage %314 %273 %275
%330 = OpImageSampleExplicitLod %23 %329 %328 Lod %281
%331 = OpLoad %23 %284
%332 = OpFAdd %23 %331 %330
OpStore %284 %332
%333 = OpConvertUToF %7 %198
%334 = OpCompositeConstruct %279 %278 %333
%335 = OpSampledImage %314 %273 %275
%336 = OpImageSampleExplicitLod %23 %335 %334 Lod|ConstOffset %281 %30
%337 = OpLoad %23 %284
%338 = OpFAdd %23 %337 %336
OpStore %284 %338
%339 = OpConvertUToF %7 %198
%340 = OpCompositeConstruct %279 %278 %339
%341 = OpSampledImage %314 %273 %275
%342 = OpImageSampleImplicitLod %23 %341 %340 Bias|ConstOffset %282 %30
%343 = OpLoad %23 %284
%344 = OpFAdd %23 %343 %342
OpStore %284 %344
%345 = OpConvertSToF %7 %283
%346 = OpCompositeConstruct %279 %278 %345
%347 = OpSampledImage %314 %273 %275
%348 = OpImageSampleImplicitLod %23 %347 %346
%349 = OpLoad %23 %284
%350 = OpFAdd %23 %349 %348
OpStore %284 %350
%351 = OpConvertSToF %7 %283
%352 = OpCompositeConstruct %279 %278 %351
%353 = OpSampledImage %314 %273 %275
%354 = OpImageSampleImplicitLod %23 %353 %352 ConstOffset %30
%355 = OpLoad %23 %284
%356 = OpFAdd %23 %355 %354
OpStore %284 %356
%357 = OpConvertSToF %7 %283
%358 = OpCompositeConstruct %279 %278 %357
%359 = OpSampledImage %314 %273 %275
%360 = OpImageSampleExplicitLod %23 %359 %358 Lod %281
%361 = OpLoad %23 %284
%362 = OpFAdd %23 %361 %360
OpStore %284 %362
%363 = OpConvertSToF %7 %283
%364 = OpCompositeConstruct %279 %278 %363
%365 = OpSampledImage %314 %273 %275
%366 = OpImageSampleExplicitLod %23 %365 %364 Lod|ConstOffset %281 %30
%367 = OpLoad %23 %284
%368 = OpFAdd %23 %367 %366
OpStore %284 %368
%369 = OpConvertSToF %7 %283
%370 = OpCompositeConstruct %279 %278 %369
%371 = OpSampledImage %314 %273 %275
%372 = OpImageSampleImplicitLod %23 %371 %370 Bias|ConstOffset %282 %30
%373 = OpLoad %23 %284
%374 = OpFAdd %23 %373 %372
OpStore %284 %374
%376 = OpConvertUToF %7 %198
%377 = OpCompositeConstruct %23 %280 %376
%378 = OpSampledImage %375 %274 %275
%379 = OpImageSampleImplicitLod %23 %378 %377
%380 = OpLoad %23 %284
%381 = OpFAdd %23 %380 %379
OpStore %284 %381
%382 = OpConvertUToF %7 %198
%383 = OpCompositeConstruct %23 %280 %382
%384 = OpSampledImage %375 %274 %275
%385 = OpImageSampleExplicitLod %23 %384 %383 Lod %281
%386 = OpLoad %23 %284
%387 = OpFAdd %23 %386 %385
OpStore %284 %387
%388 = OpConvertUToF %7 %198
%389 = OpCompositeConstruct %23 %280 %388
%390 = OpSampledImage %375 %274 %275
%391 = OpImageSampleImplicitLod %23 %390 %389 Bias %282
%392 = OpLoad %23 %284
%393 = OpFAdd %23 %392 %391
OpStore %284 %393
%394 = OpConvertSToF %7 %283
%395 = OpCompositeConstruct %23 %280 %394
%396 = OpSampledImage %375 %274 %275
%397 = OpImageSampleImplicitLod %23 %396 %395
%398 = OpLoad %23 %284
%399 = OpFAdd %23 %398 %397
OpStore %284 %399
%400 = OpConvertSToF %7 %283
%401 = OpCompositeConstruct %23 %280 %400
%402 = OpSampledImage %375 %274 %275
%403 = OpImageSampleExplicitLod %23 %402 %401 Lod %281
%404 = OpLoad %23 %284
%405 = OpFAdd %23 %404 %403
OpStore %284 %405
%406 = OpConvertSToF %7 %283
%407 = OpCompositeConstruct %23 %280 %406
%408 = OpSampledImage %375 %274 %275
%409 = OpImageSampleImplicitLod %23 %408 %407 Bias %282
%410 = OpLoad %23 %284
%411 = OpFAdd %23 %410 %409
OpStore %284 %411
%412 = OpLoad %23 %284
OpStore %269 %412
%288 = OpCompositeExtract %7 %278 0
%290 = OpSampledImage %289 %271 %275
%291 = OpImageSampleImplicitLod %23 %290 %288
%292 = OpLoad %23 %284
%293 = OpFAdd %23 %292 %291
OpStore %284 %293
%295 = OpSampledImage %294 %272 %275
%296 = OpImageSampleImplicitLod %23 %295 %278
%297 = OpLoad %23 %284
%298 = OpFAdd %23 %297 %296
OpStore %284 %298
%299 = OpSampledImage %294 %272 %275
%300 = OpImageSampleImplicitLod %23 %299 %278 ConstOffset %30
%301 = OpLoad %23 %284
%302 = OpFAdd %23 %301 %300
OpStore %284 %302
%303 = OpSampledImage %294 %272 %275
%304 = OpImageSampleExplicitLod %23 %303 %278 Lod %281
%305 = OpLoad %23 %284
%306 = OpFAdd %23 %305 %304
OpStore %284 %306
%307 = OpSampledImage %294 %272 %275
%308 = OpImageSampleExplicitLod %23 %307 %278 Lod|ConstOffset %281 %30
%309 = OpLoad %23 %284
%310 = OpFAdd %23 %309 %308
OpStore %284 %310
%311 = OpSampledImage %294 %272 %275
%312 = OpImageSampleImplicitLod %23 %311 %278 Bias|ConstOffset %282 %30
%313 = OpLoad %23 %284
%314 = OpFAdd %23 %313 %312
OpStore %284 %314
%316 = OpConvertUToF %7 %198
%317 = OpCompositeConstruct %279 %278 %316
%318 = OpSampledImage %315 %273 %275
%319 = OpImageSampleImplicitLod %23 %318 %317
%320 = OpLoad %23 %284
%321 = OpFAdd %23 %320 %319
OpStore %284 %321
%322 = OpConvertUToF %7 %198
%323 = OpCompositeConstruct %279 %278 %322
%324 = OpSampledImage %315 %273 %275
%325 = OpImageSampleImplicitLod %23 %324 %323 ConstOffset %30
%326 = OpLoad %23 %284
%327 = OpFAdd %23 %326 %325
OpStore %284 %327
%328 = OpConvertUToF %7 %198
%329 = OpCompositeConstruct %279 %278 %328
%330 = OpSampledImage %315 %273 %275
%331 = OpImageSampleExplicitLod %23 %330 %329 Lod %281
%332 = OpLoad %23 %284
%333 = OpFAdd %23 %332 %331
OpStore %284 %333
%334 = OpConvertUToF %7 %198
%335 = OpCompositeConstruct %279 %278 %334
%336 = OpSampledImage %315 %273 %275
%337 = OpImageSampleExplicitLod %23 %336 %335 Lod|ConstOffset %281 %30
%338 = OpLoad %23 %284
%339 = OpFAdd %23 %338 %337
OpStore %284 %339
%340 = OpConvertUToF %7 %198
%341 = OpCompositeConstruct %279 %278 %340
%342 = OpSampledImage %315 %273 %275
%343 = OpImageSampleImplicitLod %23 %342 %341 Bias|ConstOffset %282 %30
%344 = OpLoad %23 %284
%345 = OpFAdd %23 %344 %343
OpStore %284 %345
%346 = OpConvertSToF %7 %283
%347 = OpCompositeConstruct %279 %278 %346
%348 = OpSampledImage %315 %273 %275
%349 = OpImageSampleImplicitLod %23 %348 %347
%350 = OpLoad %23 %284
%351 = OpFAdd %23 %350 %349
OpStore %284 %351
%352 = OpConvertSToF %7 %283
%353 = OpCompositeConstruct %279 %278 %352
%354 = OpSampledImage %315 %273 %275
%355 = OpImageSampleImplicitLod %23 %354 %353 ConstOffset %30
%356 = OpLoad %23 %284
%357 = OpFAdd %23 %356 %355
OpStore %284 %357
%358 = OpConvertSToF %7 %283
%359 = OpCompositeConstruct %279 %278 %358
%360 = OpSampledImage %315 %273 %275
%361 = OpImageSampleExplicitLod %23 %360 %359 Lod %281
%362 = OpLoad %23 %284
%363 = OpFAdd %23 %362 %361
OpStore %284 %363
%364 = OpConvertSToF %7 %283
%365 = OpCompositeConstruct %279 %278 %364
%366 = OpSampledImage %315 %273 %275
%367 = OpImageSampleExplicitLod %23 %366 %365 Lod|ConstOffset %281 %30
%368 = OpLoad %23 %284
%369 = OpFAdd %23 %368 %367
OpStore %284 %369
%370 = OpConvertSToF %7 %283
%371 = OpCompositeConstruct %279 %278 %370
%372 = OpSampledImage %315 %273 %275
%373 = OpImageSampleImplicitLod %23 %372 %371 Bias|ConstOffset %282 %30
%374 = OpLoad %23 %284
%375 = OpFAdd %23 %374 %373
OpStore %284 %375
%377 = OpConvertUToF %7 %198
%378 = OpCompositeConstruct %23 %280 %377
%379 = OpSampledImage %376 %274 %275
%380 = OpImageSampleImplicitLod %23 %379 %378
%381 = OpLoad %23 %284
%382 = OpFAdd %23 %381 %380
OpStore %284 %382
%383 = OpConvertUToF %7 %198
%384 = OpCompositeConstruct %23 %280 %383
%385 = OpSampledImage %376 %274 %275
%386 = OpImageSampleExplicitLod %23 %385 %384 Lod %281
%387 = OpLoad %23 %284
%388 = OpFAdd %23 %387 %386
OpStore %284 %388
%389 = OpConvertUToF %7 %198
%390 = OpCompositeConstruct %23 %280 %389
%391 = OpSampledImage %376 %274 %275
%392 = OpImageSampleImplicitLod %23 %391 %390 Bias %282
%393 = OpLoad %23 %284
%394 = OpFAdd %23 %393 %392
OpStore %284 %394
%395 = OpConvertSToF %7 %283
%396 = OpCompositeConstruct %23 %280 %395
%397 = OpSampledImage %376 %274 %275
%398 = OpImageSampleImplicitLod %23 %397 %396
%399 = OpLoad %23 %284
%400 = OpFAdd %23 %399 %398
OpStore %284 %400
%401 = OpConvertSToF %7 %283
%402 = OpCompositeConstruct %23 %280 %401
%403 = OpSampledImage %376 %274 %275
%404 = OpImageSampleExplicitLod %23 %403 %402 Lod %281
%405 = OpLoad %23 %284
%406 = OpFAdd %23 %405 %404
OpStore %284 %406
%407 = OpConvertSToF %7 %283
%408 = OpCompositeConstruct %23 %280 %407
%409 = OpSampledImage %376 %274 %275
%410 = OpImageSampleImplicitLod %23 %409 %408 Bias %282
%411 = OpLoad %23 %284
%412 = OpFAdd %23 %411 %410
OpStore %284 %412
%413 = OpLoad %23 %284
OpStore %269 %413
OpReturn
OpFunctionEnd
%416 = OpFunction %2 None %79
%413 = OpLabel
%421 = OpVariable %422 Function %423
%417 = OpLoad %24 %66
%418 = OpLoad %25 %68
%419 = OpLoad %26 %70
%420 = OpLoad %27 %72
OpBranch %424
%424 = OpLabel
%426 = OpSampledImage %425 %418 %417
%427 = OpImageSampleDrefImplicitLod %7 %426 %278 %276
%428 = OpLoad %7 %421
%429 = OpFAdd %7 %428 %427
OpStore %421 %429
%431 = OpConvertUToF %7 %198
%432 = OpCompositeConstruct %279 %278 %431
%433 = OpSampledImage %430 %419 %417
%434 = OpImageSampleDrefImplicitLod %7 %433 %432 %276
%435 = OpLoad %7 %421
%436 = OpFAdd %7 %435 %434
OpStore %421 %436
%437 = OpConvertSToF %7 %283
%438 = OpCompositeConstruct %279 %278 %437
%439 = OpSampledImage %430 %419 %417
%440 = OpImageSampleDrefImplicitLod %7 %439 %438 %276
%441 = OpLoad %7 %421
%442 = OpFAdd %7 %441 %440
OpStore %421 %442
%444 = OpSampledImage %443 %420 %417
%445 = OpImageSampleDrefImplicitLod %7 %444 %280 %276
%446 = OpLoad %7 %421
%447 = OpFAdd %7 %446 %445
OpStore %421 %447
%448 = OpSampledImage %425 %418 %417
%449 = OpImageSampleDrefExplicitLod %7 %448 %278 %276 Lod %450
%451 = OpLoad %7 %421
%452 = OpFAdd %7 %451 %449
OpStore %421 %452
%453 = OpConvertUToF %7 %198
%454 = OpCompositeConstruct %279 %278 %453
%455 = OpSampledImage %430 %419 %417
%456 = OpImageSampleDrefExplicitLod %7 %455 %454 %276 Lod %450
%457 = OpLoad %7 %421
%458 = OpFAdd %7 %457 %456
OpStore %421 %458
%459 = OpConvertSToF %7 %283
%460 = OpCompositeConstruct %279 %278 %459
%461 = OpSampledImage %430 %419 %417
%462 = OpImageSampleDrefExplicitLod %7 %461 %460 %276 Lod %450
%463 = OpLoad %7 %421
%464 = OpFAdd %7 %463 %462
OpStore %421 %464
%465 = OpSampledImage %443 %420 %417
%466 = OpImageSampleDrefExplicitLod %7 %465 %280 %276 Lod %450
%467 = OpLoad %7 %421
%468 = OpFAdd %7 %467 %466
OpStore %421 %468
%469 = OpLoad %7 %421
OpStore %414 %469
%417 = OpFunction %2 None %79
%414 = OpLabel
%422 = OpVariable %423 Function %424
%418 = OpLoad %24 %66
%419 = OpLoad %25 %68
%420 = OpLoad %26 %70
%421 = OpLoad %27 %72
OpBranch %425
%425 = OpLabel
%427 = OpSampledImage %426 %419 %418
%428 = OpImageSampleDrefImplicitLod %7 %427 %278 %276
%429 = OpLoad %7 %422
%430 = OpFAdd %7 %429 %428
OpStore %422 %430
%432 = OpConvertUToF %7 %198
%433 = OpCompositeConstruct %279 %278 %432
%434 = OpSampledImage %431 %420 %418
%435 = OpImageSampleDrefImplicitLod %7 %434 %433 %276
%436 = OpLoad %7 %422
%437 = OpFAdd %7 %436 %435
OpStore %422 %437
%438 = OpConvertSToF %7 %283
%439 = OpCompositeConstruct %279 %278 %438
%440 = OpSampledImage %431 %420 %418
%441 = OpImageSampleDrefImplicitLod %7 %440 %439 %276
%442 = OpLoad %7 %422
%443 = OpFAdd %7 %442 %441
OpStore %422 %443
%445 = OpSampledImage %444 %421 %418
%446 = OpImageSampleDrefImplicitLod %7 %445 %280 %276
%447 = OpLoad %7 %422
%448 = OpFAdd %7 %447 %446
OpStore %422 %448
%449 = OpSampledImage %426 %419 %418
%450 = OpImageSampleDrefExplicitLod %7 %449 %278 %276 Lod %451
%452 = OpLoad %7 %422
%453 = OpFAdd %7 %452 %450
OpStore %422 %453
%454 = OpConvertUToF %7 %198
%455 = OpCompositeConstruct %279 %278 %454
%456 = OpSampledImage %431 %420 %418
%457 = OpImageSampleDrefExplicitLod %7 %456 %455 %276 Lod %451
%458 = OpLoad %7 %422
%459 = OpFAdd %7 %458 %457
OpStore %422 %459
%460 = OpConvertSToF %7 %283
%461 = OpCompositeConstruct %279 %278 %460
%462 = OpSampledImage %431 %420 %418
%463 = OpImageSampleDrefExplicitLod %7 %462 %461 %276 Lod %451
%464 = OpLoad %7 %422
%465 = OpFAdd %7 %464 %463
OpStore %422 %465
%466 = OpSampledImage %444 %421 %418
%467 = OpImageSampleDrefExplicitLod %7 %466 %280 %276 Lod %451
%468 = OpLoad %7 %422
%469 = OpFAdd %7 %468 %467
OpStore %422 %469
%470 = OpLoad %7 %422
OpStore %415 %470
OpReturn
OpFunctionEnd
%472 = OpFunction %2 None %79
%470 = OpLabel
%473 = OpLoad %16 %49
%474 = OpLoad %3 %51
%475 = OpLoad %17 %52
%476 = OpLoad %24 %64
%477 = OpLoad %24 %66
%478 = OpLoad %25 %68
OpBranch %479
%479 = OpLabel
%480 = OpSampledImage %293 %473 %476
%481 = OpImageGather %23 %480 %278 %482
%483 = OpSampledImage %293 %473 %476
%484 = OpImageGather %23 %483 %278 %485 ConstOffset %30
%486 = OpSampledImage %425 %478 %477
%487 = OpImageDrefGather %23 %486 %278 %276
%488 = OpSampledImage %425 %478 %477
%489 = OpImageDrefGather %23 %488 %278 %276 ConstOffset %30
%491 = OpSampledImage %490 %474 %476
%492 = OpImageGather %98 %491 %278 %198
%495 = OpSampledImage %494 %475 %476
%496 = OpImageGather %493 %495 %278 %198
%497 = OpConvertUToF %23 %492
%498 = OpConvertSToF %23 %496
%499 = OpFAdd %23 %497 %498
%500 = OpFAdd %23 %481 %484
%501 = OpFAdd %23 %500 %487
%502 = OpFAdd %23 %501 %489
%503 = OpFAdd %23 %502 %499
OpStore %471 %503
%473 = OpFunction %2 None %79
%471 = OpLabel
%474 = OpLoad %16 %49
%475 = OpLoad %3 %51
%476 = OpLoad %17 %52
%477 = OpLoad %24 %64
%478 = OpLoad %24 %66
%479 = OpLoad %25 %68
OpBranch %480
%480 = OpLabel
%481 = OpSampledImage %294 %474 %477
%482 = OpImageGather %23 %481 %278 %483
%484 = OpSampledImage %294 %474 %477
%485 = OpImageGather %23 %484 %278 %486 ConstOffset %30
%487 = OpSampledImage %426 %479 %478
%488 = OpImageDrefGather %23 %487 %278 %276
%489 = OpSampledImage %426 %479 %478
%490 = OpImageDrefGather %23 %489 %278 %276 ConstOffset %30
%492 = OpSampledImage %491 %475 %477
%493 = OpImageGather %98 %492 %278 %198
%496 = OpSampledImage %495 %476 %477
%497 = OpImageGather %494 %496 %278 %198
%498 = OpConvertUToF %23 %493
%499 = OpConvertSToF %23 %497
%500 = OpFAdd %23 %498 %499
%501 = OpFAdd %23 %482 %485
%502 = OpFAdd %23 %501 %488
%503 = OpFAdd %23 %502 %490
%504 = OpFAdd %23 %503 %500
OpStore %472 %504
OpReturn
OpFunctionEnd
%506 = OpFunction %2 None %79
%504 = OpLabel
%507 = OpLoad %24 %64
%508 = OpLoad %25 %68
OpBranch %509
%509 = OpLabel
%510 = OpSampledImage %425 %508 %507
%511 = OpImageSampleImplicitLod %23 %510 %278
%512 = OpCompositeExtract %7 %511 0
%513 = OpSampledImage %425 %508 %507
%514 = OpImageGather %23 %513 %278 %198
%515 = OpCompositeConstruct %23 %512 %512 %512 %512
%516 = OpFAdd %23 %515 %514
OpStore %505 %516
%507 = OpFunction %2 None %79
%505 = OpLabel
%508 = OpLoad %24 %64
%509 = OpLoad %25 %68
OpBranch %510
%510 = OpLabel
%511 = OpSampledImage %426 %509 %508
%512 = OpImageSampleImplicitLod %23 %511 %278
%513 = OpCompositeExtract %7 %512 0
%514 = OpSampledImage %426 %509 %508
%515 = OpImageGather %23 %514 %278 %198
%516 = OpCompositeConstruct %23 %513 %513 %513 %513
%517 = OpFAdd %23 %516 %515
OpStore %506 %517
OpReturn
OpFunctionEnd

2
tests/out/wgsl/access.wgsl
View File

@ -138,7 +138,7 @@ fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
test_matrix_within_array_within_struct_accesses();
let _matrix = bar._matrix;
let arr_1 = bar.arr;
let b = bar._matrix[3][0];
let b = bar._matrix[3u][0];
let a_1 = bar.data[(arrayLength((&bar.data)) - 2u)].value;
let c = qux;
let data_pointer = (&bar.data[0].value);

8
tests/out/wgsl/const-exprs.wgsl
View File

@ -9,17 +9,13 @@ var<storage, read_write> out: vec4<i32>;
var<storage, read_write> out2_: i32;
fn swizzle_of_compose() {
let a = vec2<i32>(1, 2);
let b = vec2<i32>(3, 4);
out = vec4<i32>(4, 3, 2, 1);
return;
}
fn index_of_compose() {
let a_1 = vec2<i32>(1, 2);
let b_1 = vec2<i32>(3, 4);
let _e7 = out2_;
out2_ = (_e7 + 2);
let _e2 = out2_;
out2_ = (_e2 + 2);
return;
}

136
tests/out/wgsl/image.wgsl
View File

@ -112,74 +112,74 @@ fn texture_sample() -> @location(0) vec4<f32> {
let tc = vec2(0.5);
let tc3_ = vec3(0.5);
let _e8 = textureSample(image_1d, sampler_reg, 0.5);
let _e9 = a;
a = (_e9 + _e8);
let _e13 = textureSample(image_2d, sampler_reg, tc);
let _e14 = a;
a = (_e14 + _e13);
let _e18 = textureSample(image_2d, sampler_reg, tc, vec2<i32>(3, 1));
let _e19 = a;
a = (_e19 + _e18);
let _e23 = textureSampleLevel(image_2d, sampler_reg, tc, 2.3);
let _e24 = a;
a = (_e24 + _e23);
let _e28 = textureSampleLevel(image_2d, sampler_reg, tc, 2.3, vec2<i32>(3, 1));
let _e29 = a;
a = (_e29 + _e28);
let _e34 = textureSampleBias(image_2d, sampler_reg, tc, 2.0, vec2<i32>(3, 1));
let _e35 = a;
a = (_e35 + _e34);
let _e40 = textureSample(image_2d_array, sampler_reg, tc, 0u);
let _e41 = a;
a = (_e41 + _e40);
let _e46 = textureSample(image_2d_array, sampler_reg, tc, 0u, vec2<i32>(3, 1));
let _e47 = a;
a = (_e47 + _e46);
let _e52 = textureSampleLevel(image_2d_array, sampler_reg, tc, 0u, 2.3);
let _e53 = a;
a = (_e53 + _e52);
let _e58 = textureSampleLevel(image_2d_array, sampler_reg, tc, 0u, 2.3, vec2<i32>(3, 1));
let _e59 = a;
a = (_e59 + _e58);
let _e65 = textureSampleBias(image_2d_array, sampler_reg, tc, 0u, 2.0, vec2<i32>(3, 1));
let _e66 = a;
a = (_e66 + _e65);
let _e71 = textureSample(image_2d_array, sampler_reg, tc, 0);
let _e72 = a;
a = (_e72 + _e71);
let _e77 = textureSample(image_2d_array, sampler_reg, tc, 0, vec2<i32>(3, 1));
let _e78 = a;
a = (_e78 + _e77);
let _e83 = textureSampleLevel(image_2d_array, sampler_reg, tc, 0, 2.3);
let _e84 = a;
a = (_e84 + _e83);
let _e89 = textureSampleLevel(image_2d_array, sampler_reg, tc, 0, 2.3, vec2<i32>(3, 1));
let _e90 = a;
a = (_e90 + _e89);
let _e96 = textureSampleBias(image_2d_array, sampler_reg, tc, 0, 2.0, vec2<i32>(3, 1));
let _e97 = a;
a = (_e97 + _e96);
let _e102 = textureSample(image_cube_array, sampler_reg, tc3_, 0u);
let _e103 = a;
a = (_e103 + _e102);
let _e108 = textureSampleLevel(image_cube_array, sampler_reg, tc3_, 0u, 2.3);
let _e109 = a;
a = (_e109 + _e108);
let _e115 = textureSampleBias(image_cube_array, sampler_reg, tc3_, 0u, 2.0);
let _e116 = a;
a = (_e116 + _e115);
let _e121 = textureSample(image_cube_array, sampler_reg, tc3_, 0);
let _e122 = a;
a = (_e122 + _e121);
let _e127 = textureSampleLevel(image_cube_array, sampler_reg, tc3_, 0, 2.3);
let _e128 = a;
a = (_e128 + _e127);
let _e134 = textureSampleBias(image_cube_array, sampler_reg, tc3_, 0, 2.0);
let _e135 = a;
a = (_e135 + _e134);
let _e137 = a;
return _e137;
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<i32>(3, 1));
let _e20 = a;
a = (_e20 + _e19);
let _e24 = textureSampleLevel(image_2d, sampler_reg, tc, 2.3);
let _e25 = a;
a = (_e25 + _e24);
let _e29 = textureSampleLevel(image_2d, sampler_reg, tc, 2.3, vec2<i32>(3, 1));
let _e30 = a;
a = (_e30 + _e29);
let _e35 = textureSampleBias(image_2d, sampler_reg, tc, 2.0, vec2<i32>(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<i32>(3, 1));
let _e48 = a;
a = (_e48 + _e47);
let _e53 = textureSampleLevel(image_2d_array, sampler_reg, tc, 0u, 2.3);
let _e54 = a;
a = (_e54 + _e53);
let _e59 = textureSampleLevel(image_2d_array, sampler_reg, tc, 0u, 2.3, vec2<i32>(3, 1));
let _e60 = a;
a = (_e60 + _e59);
let _e66 = textureSampleBias(image_2d_array, sampler_reg, tc, 0u, 2.0, vec2<i32>(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<i32>(3, 1));
let _e79 = a;
a = (_e79 + _e78);
let _e84 = textureSampleLevel(image_2d_array, sampler_reg, tc, 0, 2.3);
let _e85 = a;
a = (_e85 + _e84);
let _e90 = textureSampleLevel(image_2d_array, sampler_reg, tc, 0, 2.3, vec2<i32>(3, 1));
let _e91 = a;
a = (_e91 + _e90);
let _e97 = textureSampleBias(image_2d_array, sampler_reg, tc, 0, 2.0, vec2<i32>(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.3);
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.3);
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

2
tests/out/wgsl/type-alias.wgsl
View File

@ -3,7 +3,7 @@ fn main() {
let c = vec3(0.0);
let b = vec3<f32>(vec2(0.0), 0.0);
let d = vec3<f32>(vec2(0.0), 0.0);
let e = vec3<i32>(vec2<i32>(0, 0), 0);
let e = vec3<i32>(d);
let f = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
let g = mat3x3<f32>(a, a, a);
}

74
tests/wgsl-errors.rs
View File

@ -1759,47 +1759,47 @@ fn assign_to_let() {
"###,
);
// check(
// "
// fn f() {
// let a = array(1, 2);
// a[0] = 1;
// }
// ",
// r###"error: invalid left-hand side of assignment
// ┌─ wgsl:3:17
// │
// 3 │ let a = array(1, 2);
// │ ^ this is an immutable binding
// 4 │ a[0] = 1;
// │ ^^^^ cannot assign to this expression
// │
// = note: consider declaring 'a' with `var` instead of `let`
check(
"
fn f() {
let a = array(1, 2);
a[0] = 1;
}
",
r###"error: invalid left-hand side of assignment
┌─ wgsl:3:17
│
3 │ let a = array(1, 2);
│ ^ this is an immutable binding
4 │ a[0] = 1;
│ ^^^^ cannot assign to this expression
│
= note: consider declaring 'a' with `var` instead of `let`
// "###,
// );
"###,
);
// check(
// "
// struct S { a: i32 }
check(
"
struct S { a: i32 }
// fn f() {
// let a = S(10);
// a.a = 20;
// }
// ",
// r###"error: invalid left-hand side of assignment
// ┌─ wgsl:5:17
// │
// 5 │ let a = S(10);
// │ ^ this is an immutable binding
// 6 │ a.a = 20;
// │ ^^^ cannot assign to this expression
// │
// = note: consider declaring 'a' with `var` instead of `let`
fn f() {
let a = S(10);
a.a = 20;
}
",
r###"error: invalid left-hand side of assignment
┌─ wgsl:5:17
│
5 │ let a = S(10);
│ ^ this is an immutable binding
6 │ a.a = 20;
│ ^^^ cannot assign to this expression
│
= note: consider declaring 'a' with `var` instead of `let`
// "###,
// );
"###,
);
}
#[test]