diff --git a/checkout_glslang_spirv_tools.sh b/checkout_glslang_spirv_tools.sh index 9b816cbc..e71f5070 100755 --- a/checkout_glslang_spirv_tools.sh +++ b/checkout_glslang_spirv_tools.sh @@ -1,7 +1,7 @@ #!/bin/bash -GLSLANG_REV=845860d56513d95e15fe4820df7272f9687d076e -SPIRV_TOOLS_REV=340370eddbb9f0e7d26b6a4f7e22c1b98150e5e1 +GLSLANG_REV=7cec64fc42eba4587d2f3d73c03cb68ceb18dca4 +SPIRV_TOOLS_REV=fe9121f72144f1e1657cd21a55a1fec07c940e56 if [ -d external/glslang ]; then echo "Updating glslang to revision $GLSLANG_REV." diff --git a/reference/opt/shaders-hlsl/comp/barriers.comp b/reference/opt/shaders-hlsl/comp/barriers.comp index b24e7c48..7ac2a656 100644 --- a/reference/opt/shaders-hlsl/comp/barriers.comp +++ b/reference/opt/shaders-hlsl/comp/barriers.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u); + void comp_main() { GroupMemoryBarrier(); diff --git a/reference/opt/shaders-hlsl/comp/builtins.comp b/reference/opt/shaders-hlsl/comp/builtins.comp index 990fc853..7f88aa79 100644 --- a/reference/opt/shaders-hlsl/comp/builtins.comp +++ b/reference/opt/shaders-hlsl/comp/builtins.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(8u, 4u, 2u); + void comp_main() { } diff --git a/reference/opt/shaders-hlsl/comp/shared.comp b/reference/opt/shaders-hlsl/comp/shared.comp index a6b56f89..9831302a 100644 --- a/reference/opt/shaders-hlsl/comp/shared.comp +++ b/reference/opt/shaders-hlsl/comp/shared.comp @@ -1,3 +1,5 @@ +static const uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u); + ByteAddressBuffer _22 : register(t0); RWByteAddressBuffer _44 : register(u1); diff --git a/reference/opt/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp b/reference/opt/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp index 0c3b3059..5802ddac 100644 --- a/reference/opt/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp +++ b/reference/opt/shaders-msl/asm/comp/specialization-constant-workgroup.asm.comp @@ -3,6 +3,12 @@ using namespace metal; +constant uint _5_tmp [[function_constant(10)]]; +constant uint _5 = is_function_constant_defined(_5_tmp) ? _5_tmp : 9u; +constant uint _6_tmp [[function_constant(12)]]; +constant uint _6 = is_function_constant_defined(_6_tmp) ? _6_tmp : 4u; +constant uint3 gl_WorkGroupSize = uint3(_5, 20u, _6); + struct SSBO { float a; diff --git a/reference/opt/shaders-msl/comp/barriers.comp b/reference/opt/shaders-msl/comp/barriers.comp index 476a30a9..23a19914 100644 --- a/reference/opt/shaders-msl/comp/barriers.comp +++ b/reference/opt/shaders-msl/comp/barriers.comp @@ -3,6 +3,8 @@ using namespace metal; +constant uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u); + kernel void main0() { threadgroup_barrier(mem_flags::mem_threadgroup); diff --git a/reference/opt/shaders-msl/comp/builtins.comp b/reference/opt/shaders-msl/comp/builtins.comp index 82782202..189bb1b7 100644 --- a/reference/opt/shaders-msl/comp/builtins.comp +++ b/reference/opt/shaders-msl/comp/builtins.comp @@ -3,6 +3,8 @@ using namespace metal; +constant uint3 gl_WorkGroupSize = uint3(8u, 4u, 2u); + kernel void main0(uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_NumWorkGroups [[threadgroups_per_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) { } diff --git a/reference/opt/shaders-msl/comp/culling.comp b/reference/opt/shaders-msl/comp/culling.comp index f30e4853..b20480bb 100644 --- a/reference/opt/shaders-msl/comp/culling.comp +++ b/reference/opt/shaders-msl/comp/culling.comp @@ -6,6 +6,8 @@ using namespace metal; +constant uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u); + struct SSBO { float in_data[1]; diff --git a/reference/opt/shaders-msl/comp/packing-test-1.comp b/reference/opt/shaders-msl/comp/packing-test-1.comp index 073d8f26..0c52804e 100644 --- a/reference/opt/shaders-msl/comp/packing-test-1.comp +++ b/reference/opt/shaders-msl/comp/packing-test-1.comp @@ -3,6 +3,8 @@ using namespace metal; +constant uint3 gl_WorkGroupSize = uint3(32u, 1u, 1u); + struct T1 { packed_float3 a; diff --git a/reference/opt/shaders-msl/comp/packing-test-2.comp b/reference/opt/shaders-msl/comp/packing-test-2.comp index 8d5350b6..4f3ab465 100644 --- a/reference/opt/shaders-msl/comp/packing-test-2.comp +++ b/reference/opt/shaders-msl/comp/packing-test-2.comp @@ -3,6 +3,8 @@ using namespace metal; +constant uint3 gl_WorkGroupSize = uint3(32u, 1u, 1u); + struct T1 { packed_float3 a; diff --git a/reference/opt/shaders-msl/comp/shared-array-of-arrays.comp b/reference/opt/shaders-msl/comp/shared-array-of-arrays.comp index 98290cce..3884c22f 100644 --- a/reference/opt/shaders-msl/comp/shared-array-of-arrays.comp +++ b/reference/opt/shaders-msl/comp/shared-array-of-arrays.comp @@ -3,6 +3,8 @@ using namespace metal; +constant uint3 gl_WorkGroupSize = uint3(4u, 4u, 1u); + struct SSBO { float out_data[1]; diff --git a/reference/opt/shaders-msl/comp/shared.comp b/reference/opt/shaders-msl/comp/shared.comp index 425d7508..780287ba 100644 --- a/reference/opt/shaders-msl/comp/shared.comp +++ b/reference/opt/shaders-msl/comp/shared.comp @@ -3,6 +3,8 @@ using namespace metal; +constant uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u); + struct SSBO { float in_data[1]; diff --git a/reference/opt/shaders-msl/comp/torture-loop.comp b/reference/opt/shaders-msl/comp/torture-loop.comp index d12bd1d2..201aceca 100644 --- a/reference/opt/shaders-msl/comp/torture-loop.comp +++ b/reference/opt/shaders-msl/comp/torture-loop.comp @@ -16,27 +16,27 @@ struct SSBO2 kernel void main0(device SSBO& _24 [[buffer(0)]], device SSBO2& _89 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { - float4 _93; - _93 = _24.in_data[gl_GlobalInvocationID.x]; - for (int _94 = 0; (_94 + 1) < 10; ) + float4 _99; + _99 = _24.in_data[gl_GlobalInvocationID.x]; + for (int _93 = 0; (_93 + 1) < 10; ) { - _93 *= 2.0; - _94 += 2; + _99 *= 2.0; + _93 += 2; continue; } - float4 _95; - _95 = _93; - float4 _100; - for (uint _97 = 0u; _97 < 16u; _95 = _100, _97++) + float4 _98; + _98 = _99; + float4 _103; + for (uint _94 = 0u; _94 < 16u; _98 = _103, _94++) { - _100 = _95; - for (uint _101 = 0u; _101 < 30u; ) + _103 = _98; + for (uint _100 = 0u; _100 < 30u; ) { - _100 = _24.mvp * _100; - _101++; + _103 = _24.mvp * _103; + _100++; continue; } } - _89.out_data[gl_GlobalInvocationID.x] = _95; + _89.out_data[gl_GlobalInvocationID.x] = _98; } diff --git a/reference/opt/shaders-msl/frag/for-loop-init.frag b/reference/opt/shaders-msl/frag/for-loop-init.frag index 6068c2c0..71c156f8 100644 --- a/reference/opt/shaders-msl/frag/for-loop-init.frag +++ b/reference/opt/shaders-msl/frag/for-loop-init.frag @@ -3,6 +3,9 @@ using namespace metal; +constant int2 _184 = {}; +constant int _199 = {}; + struct main0_out { int FragColor [[color(0)]]; @@ -12,56 +15,70 @@ fragment main0_out main0() { main0_out out = {}; out.FragColor = 16; - for (int _140 = 0; _140 < 25; ) + for (int _168 = 0; _168 < 25; ) { out.FragColor += 10; - _140++; + _168++; continue; } - for (int _141 = 1; _141 < 30; ) + for (int _169 = 1; _169 < 30; ) { out.FragColor += 11; - _141++; + _169++; continue; } - int _142; - _142 = 0; - for (; _142 < 20; ) + int _170; + _170 = 0; + for (; _170 < 20; ) { out.FragColor += 12; - _142++; + _170++; continue; } - int _62 = _142 + 3; + int _62 = _170 + 3; out.FragColor += _62; - if (_62 == 40) + bool _68 = _62 == 40; + if (_68) { - for (int _143 = 0; _143 < 40; ) + for (int _171 = 0; _171 < 40; ) { out.FragColor += 13; - _143++; + _171++; continue; } - return out; } - out.FragColor += _62; - int2 _144; - _144 = int2(0); - for (; _144.x < 10; ) + else { - out.FragColor += _144.y; - int2 _139 = _144; - _139.x = _144.x + 4; - _144 = _139; - continue; + out.FragColor += _62; } - for (int _145 = _62; _145 < 40; ) + bool2 _211 = bool2(_68); + int2 _212 = int2(_211.x ? _184.x : _184.x, _211.y ? _184.y : _184.y); + bool _213 = _68 ? true : false; + bool2 _214 = bool2(_213); + if (!_213) { - out.FragColor += _145; - _145++; - continue; + int2 _177; + _177 = int2(_214.x ? _212.x : int2(0).x, _214.y ? _212.y : int2(0).y); + for (; _177.x < 10; ) + { + out.FragColor += _177.y; + int2 _167 = _177; + _167.x = _177.x + 4; + _177 = _167; + continue; + } + } + int _216 = _213 ? (_68 ? _199 : _199) : _62; + if (!_213) + { + for (int _191 = _216; _191 < 40; ) + { + out.FragColor += _191; + _191++; + continue; + } + out.FragColor += _216; } - out.FragColor += _62; return out; } diff --git a/reference/opt/shaders/asm/comp/specialization-constant-workgroup.asm.comp b/reference/opt/shaders/asm/comp/specialization-constant-workgroup.asm.comp index c1b84377..1b2285a8 100644 --- a/reference/opt/shaders/asm/comp/specialization-constant-workgroup.asm.comp +++ b/reference/opt/shaders/asm/comp/specialization-constant-workgroup.asm.comp @@ -1,5 +1,5 @@ #version 310 es -layout(local_size_x = 1, local_size_y = 20, local_size_z = 1) in; +layout(local_size_x = 9, local_size_y = 20, local_size_z = 4) in; layout(binding = 0, std430) buffer SSBO { diff --git a/reference/opt/shaders/asm/frag/loop-body-dominator-continue-access.asm.frag b/reference/opt/shaders/asm/frag/loop-body-dominator-continue-access.asm.frag index b64f16d4..9f7a1f87 100644 --- a/reference/opt/shaders/asm/frag/loop-body-dominator-continue-access.asm.frag +++ b/reference/opt/shaders/asm/frag/loop-body-dominator-continue-access.asm.frag @@ -10,41 +10,73 @@ layout(binding = 0, std140) uniform Foo layout(location = 0) in vec3 fragWorld; layout(location = 0) out int _entryPointOutput; -int GetCascade(vec3 fragWorldPosition) -{ - for (uint _151 = 0u; _151 < _11.shadowCascadesNum; _151++) - { - mat4 _157; - for (;;) - { - if (_11.test == 0) - { - _157 = mat4(vec4(0.5, 0.0, 0.0, 0.0), vec4(0.0, 0.5, 0.0, 0.0), vec4(0.0, 0.0, 0.5, 0.0), vec4(0.0, 0.0, 0.0, 1.0)); - break; - } - _157 = mat4(vec4(1.0, 0.0, 0.0, 0.0), vec4(0.0, 1.0, 0.0, 0.0), vec4(0.0, 0.0, 1.0, 0.0), vec4(0.0, 0.0, 0.0, 1.0)); - break; - } - vec4 _92 = (_157 * _11.lightVP[_151]) * vec4(fragWorldPosition, 1.0); - float _140 = _92.z; - float _144 = _92.x; - float _146 = _92.y; - if ((((_140 >= 0.0) && (_140 <= 1.0)) && (max(_144, _146) <= 1.0)) && (min(_144, _146) >= 0.0)) - { - return int(_151); - } - else - { - continue; - } - continue; - } - return -1; -} +mat4 _235; +int _245; void main() { - vec3 _123 = fragWorld; - _entryPointOutput = GetCascade(_123); + uint _229; + bool _231; + mat4 _234; + _234 = _235; + _231 = false; + _229 = 0u; + bool _251; + mat4 _232; + int _243; + bool _158; + for (;;) + { + _158 = _229 < _11.shadowCascadesNum; + if (_158) + { + bool _209 = _11.test == 0; + mat4 _233; + if (_209) + { + _233 = mat4(vec4(0.5, 0.0, 0.0, 0.0), vec4(0.0, 0.5, 0.0, 0.0), vec4(0.0, 0.0, 0.5, 0.0), vec4(0.0, 0.0, 0.0, 1.0)); + } + else + { + _233 = _234; + } + bool _250 = _209 ? true : _231; + if (!_250) + { + _232 = mat4(vec4(1.0, 0.0, 0.0, 0.0), vec4(0.0, 1.0, 0.0, 0.0), vec4(0.0, 0.0, 1.0, 0.0), vec4(0.0, 0.0, 0.0, 1.0)); + } + else + { + _232 = _233; + } + _251 = _250 ? _250 : true; + vec4 _171 = (_232 * _11.lightVP[_229]) * vec4(fragWorld, 1.0); + float _218 = _171.z; + float _222 = _171.x; + float _224 = _171.y; + if ((((_218 >= 0.0) && (_218 <= 1.0)) && (max(_222, _224) <= 1.0)) && (min(_222, _224) >= 0.0)) + { + _243 = int(_229); + break; + } + else + { + _234 = _232; + _231 = _251; + _229++; + continue; + } + _234 = _232; + _231 = _251; + _229++; + continue; + } + else + { + _243 = _245; + break; + } + } + _entryPointOutput = (_158 ? true : false) ? _243 : (-1); } diff --git a/reference/opt/shaders/asm/frag/loop-merge-to-continue.asm.frag b/reference/opt/shaders/asm/frag/loop-merge-to-continue.asm.frag index a1cbe932..faf32edc 100644 --- a/reference/opt/shaders/asm/frag/loop-merge-to-continue.asm.frag +++ b/reference/opt/shaders/asm/frag/loop-merge-to-continue.asm.frag @@ -10,10 +10,10 @@ void main() _50 = 0; for (; _50 < 4; _50++) { - for (int _53 = 0; _53 < 4; ) + for (int _51 = 0; _51 < 4; ) { - FragColor += vec4(v0[(_50 + _53) & 3]); - _53++; + FragColor += vec4(v0[(_50 + _51) & 3]); + _51++; continue; } } diff --git a/reference/opt/shaders/asm/frag/multi-for-loop-init.asm.frag b/reference/opt/shaders/asm/frag/multi-for-loop-init.asm.frag index b8e56995..6a177288 100644 --- a/reference/opt/shaders/asm/frag/multi-for-loop-init.asm.frag +++ b/reference/opt/shaders/asm/frag/multi-for-loop-init.asm.frag @@ -14,8 +14,8 @@ void main() { FragColor += vec4(float(_53)); FragColor += vec4(float(_54)); - _53 += counter; _54 += uint(counter); + _53 += counter; continue; } } diff --git a/reference/opt/shaders/asm/frag/temporary-name-alias.asm.frag b/reference/opt/shaders/asm/frag/temporary-name-alias.asm.frag new file mode 100644 index 00000000..05ce10ad --- /dev/null +++ b/reference/opt/shaders/asm/frag/temporary-name-alias.asm.frag @@ -0,0 +1,6 @@ +#version 450 + +void main() +{ +} + diff --git a/reference/opt/shaders/comp/cfg.comp b/reference/opt/shaders/comp/cfg.comp index 631f6fd3..45b219ec 100644 --- a/reference/opt/shaders/comp/cfg.comp +++ b/reference/opt/shaders/comp/cfg.comp @@ -6,7 +6,7 @@ layout(binding = 0, std430) buffer SSBO float data; } _11; -float _180; +float _183; void main() { @@ -31,15 +31,15 @@ void main() break; } } - float _181; - _181 = _180; + float _180; + _180 = _183; for (int _179 = 0; _179 < 20; ) { + _180 += 10.0; _179++; - _181 += 10.0; continue; } - _11.data = _181; _11.data = _180; + _11.data = _183; } diff --git a/reference/opt/shaders/comp/torture-loop.comp b/reference/opt/shaders/comp/torture-loop.comp index 245e84d3..5943966c 100644 --- a/reference/opt/shaders/comp/torture-loop.comp +++ b/reference/opt/shaders/comp/torture-loop.comp @@ -14,27 +14,27 @@ layout(binding = 1, std430) writeonly buffer SSBO2 void main() { - vec4 _93; - _93 = _24.in_data[gl_GlobalInvocationID.x]; - for (int _94 = 0; (_94 + 1) < 10; ) + vec4 _99; + _99 = _24.in_data[gl_GlobalInvocationID.x]; + for (int _93 = 0; (_93 + 1) < 10; ) { - _93 *= 2.0; - _94 += 2; + _99 *= 2.0; + _93 += 2; continue; } - vec4 _95; - _95 = _93; - vec4 _100; - for (uint _97 = 0u; _97 < 16u; _95 = _100, _97++) + vec4 _98; + _98 = _99; + vec4 _103; + for (uint _94 = 0u; _94 < 16u; _98 = _103, _94++) { - _100 = _95; - for (uint _101 = 0u; _101 < 30u; ) + _103 = _98; + for (uint _100 = 0u; _100 < 30u; ) { - _100 = _24.mvp * _100; - _101++; + _103 = _24.mvp * _103; + _100++; continue; } } - _89.out_data[gl_GlobalInvocationID.x] = _95; + _89.out_data[gl_GlobalInvocationID.x] = _98; } diff --git a/reference/opt/shaders/flatten/multi-dimensional.desktop.flatten_dim.frag b/reference/opt/shaders/flatten/multi-dimensional.desktop.flatten_dim.frag index ca8ab845..3657298f 100644 --- a/reference/opt/shaders/flatten/multi-dimensional.desktop.flatten_dim.frag +++ b/reference/opt/shaders/flatten/multi-dimensional.desktop.flatten_dim.frag @@ -13,14 +13,14 @@ void main() vec4 values3[2 * 3 * 1]; for (; _92 < 2; _92++) { - int _96; - _96 = 0; - for (; _96 < 3; _96++) + int _93; + _93 = 0; + for (; _93 < 3; _93++) { - for (int _98 = 0; _98 < 1; ) + for (int _95 = 0; _95 < 1; ) { - values3[_92 * 3 * 1 + _96 * 1 + _98] = texture(uTextures[_92 * 3 * 1 + _96 * 1 + _98], vUV); - _98++; + values3[_92 * 3 * 1 + _93 * 1 + _95] = texture(uTextures[_92 * 3 * 1 + _93 * 1 + _95], vUV); + _95++; continue; } } diff --git a/reference/opt/shaders/frag/for-loop-init.frag b/reference/opt/shaders/frag/for-loop-init.frag index f9701cc8..6d13815d 100644 --- a/reference/opt/shaders/frag/for-loop-init.frag +++ b/reference/opt/shaders/frag/for-loop-init.frag @@ -4,58 +4,72 @@ precision highp int; layout(location = 0) out mediump int FragColor; +ivec2 _184; +int _199; + void main() { FragColor = 16; - for (int _140 = 0; _140 < 25; ) + for (int _168 = 0; _168 < 25; ) { FragColor += 10; - _140++; + _168++; continue; } - for (int _141 = 1; _141 < 30; ) + for (int _169 = 1; _169 < 30; ) { FragColor += 11; - _141++; + _169++; continue; } - int _142; - _142 = 0; - for (; _142 < 20; ) + int _170; + _170 = 0; + for (; _170 < 20; ) { FragColor += 12; - _142++; + _170++; continue; } - mediump int _62 = _142 + 3; + mediump int _62 = _170 + 3; FragColor += _62; - if (_62 == 40) + bool _68 = _62 == 40; + if (_68) { - for (int _143 = 0; _143 < 40; ) + for (int _171 = 0; _171 < 40; ) { FragColor += 13; - _143++; + _171++; continue; } - return; } - FragColor += _62; - ivec2 _144; - _144 = ivec2(0); - for (; _144.x < 10; ) + else { - FragColor += _144.y; - ivec2 _139 = _144; - _139.x = _144.x + 4; - _144 = _139; - continue; + FragColor += _62; } - for (int _145 = _62; _145 < 40; ) + bool _213 = _68 ? true : false; + if (!_213) { - FragColor += _145; - _145++; - continue; + ivec2 _177; + _177 = mix(ivec2(0), mix(_184, _184, bvec2(_68)), bvec2(_213)); + for (; _177.x < 10; ) + { + FragColor += _177.y; + ivec2 _167 = _177; + _167.x = _177.x + 4; + _177 = _167; + continue; + } + } + int _216 = _213 ? (_68 ? _199 : _199) : _62; + if (!_213) + { + for (int _191 = _216; _191 < 40; ) + { + FragColor += _191; + _191++; + continue; + } + FragColor += _216; } - FragColor += _62; } diff --git a/reference/opt/shaders/frag/hoisted-temporary-use-continue-block-as-value.frag b/reference/opt/shaders/frag/hoisted-temporary-use-continue-block-as-value.frag index c6550dce..74db2bb8 100644 --- a/reference/opt/shaders/frag/hoisted-temporary-use-continue-block-as-value.frag +++ b/reference/opt/shaders/frag/hoisted-temporary-use-continue-block-as-value.frag @@ -9,17 +9,17 @@ layout(location = 1) flat in mediump int vB; void main() { FragColor = vec4(0.0); - int _60; - for (int _57 = 0, _58 = 0; _58 < vA; FragColor += vec4(1.0), _57 = _60, _58 += (_60 + 10)) + int _58; + for (int _57 = 0, _60 = 0; _57 < vA; FragColor += vec4(1.0), _60 = _58, _57 += (_58 + 10)) { - if ((vA + _58) == 20) + if ((vA + _57) == 20) { - _60 = 50; + _58 = 50; continue; } else { - _60 = ((vB + _58) == 40) ? 60 : _57; + _58 = ((vB + _57) == 40) ? 60 : _60; continue; } continue; diff --git a/reference/opt/shaders/vert/ocean.vert b/reference/opt/shaders/vert/ocean.vert index 53b69de2..f1b12fa0 100644 --- a/reference/opt/shaders/vert/ocean.vert +++ b/reference/opt/shaders/vert/ocean.vert @@ -50,7 +50,7 @@ layout(location = 0) in vec4 Position; layout(location = 0) out vec3 EyeVec; layout(location = 1) out vec4 TexCoord; -uvec4 _483; +uvec4 _484; void main() { @@ -69,42 +69,42 @@ void main() { _482 = 0u; } - uvec4 _445 = _483; + uvec4 _445 = _484; _445.x = _482; bool _379 = _359.y < 32u; - uint _484; + uint _485; if (_379) { - _484 = _366.x; - } - else - { - _484 = 0u; - } - uvec4 _451 = _445; - _451.y = _484; - uint _485; - if (_369) - { - _485 = _366.y; + _485 = _366.x; } else { _485 = 0u; } - uvec4 _457 = _451; - _457.z = _485; - uint _486; - if (_379) + uvec4 _451 = _445; + _451.y = _485; + uint _487; + if (_369) { - _486 = _366.y; + _487 = _366.y; } else { - _486 = 0u; + _487 = 0u; + } + uvec4 _457 = _451; + _457.z = _487; + uint _489; + if (_379) + { + _489 = _366.y; + } + else + { + _489 = 0u; } uvec4 _463 = _457; - _463.w = _486; + _463.w = _489; vec4 _415 = vec4((_359.xyxy + _463) & (~_366).xxyy); vec2 _197 = ((_53.Patches[(gl_InstanceID + SPIRV_Cross_BaseInstance)].Position.xz * _180.InvOceanSize_PatchScale.zw) + mix(_415.xy, _415.zw, vec2(_350 - _352))) * _180.InvOceanSize_PatchScale.xy; vec2 _204 = _197 * _180.NormalTexCoordScale.zw; diff --git a/reference/shaders/asm/frag/temporary-name-alias.asm.frag b/reference/shaders/asm/frag/temporary-name-alias.asm.frag new file mode 100644 index 00000000..927c0434 --- /dev/null +++ b/reference/shaders/asm/frag/temporary-name-alias.asm.frag @@ -0,0 +1,10 @@ +#version 450 + +void main() +{ + float constituent = float(0); + mat3 _mat3 = mat3(vec3(constituent), vec3(constituent), vec3(constituent)); + float constituent_1 = float(1); + _mat3 = mat3(vec3(constituent_1), vec3(constituent_1), vec3(constituent_1)); +} + diff --git a/shaders/asm/frag/temporary-name-alias.asm.frag b/shaders/asm/frag/temporary-name-alias.asm.frag new file mode 100644 index 00000000..fc3c0c36 --- /dev/null +++ b/shaders/asm/frag/temporary-name-alias.asm.frag @@ -0,0 +1,48 @@ +; SPIR-V +; Version: 1.2 +; Generator: Khronos; 0 +; Bound: 51 +; Schema: 0 + OpCapability Linkage + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint Fragment %EntryPoint_Main "main" + OpExecutionMode %EntryPoint_Main OriginUpperLeft + OpSource Unknown 100 + OpName %mat3 "mat3" + OpName %constituent "constituent" + OpName %constituent_0 "constituent" + OpName %constituent_1 "constituent" + OpName %constituent_2 "constituent" + OpName %constituent_3 "constituent" + OpName %constituent_4 "constituent" + OpName %constituent_5 "constituent" + OpName %constituent_6 "constituent" + OpName %EntryPoint_Main "EntryPoint_Main" + %void = OpTypeVoid +%_ptr_Function_void = OpTypePointer Function %void + %float = OpTypeFloat 32 + %int = OpTypeInt 32 1 + %v3float = OpTypeVector %float 3 +%mat3v3float = OpTypeMatrix %v3float 3 +%_ptr_Function_mat3v3float = OpTypePointer Function %mat3v3float + %14 = OpTypeFunction %void + %int_0 = OpConstant %int 0 + %int_1 = OpConstant %int 1 +%EntryPoint_Main = OpFunction %void None %14 + %45 = OpLabel + %mat3 = OpVariable %_ptr_Function_mat3v3float Function +%constituent = OpConvertSToF %float %int_0 +%constituent_0 = OpCompositeConstruct %v3float %constituent %constituent %constituent +%constituent_1 = OpCompositeConstruct %v3float %constituent %constituent %constituent +%constituent_2 = OpCompositeConstruct %v3float %constituent %constituent %constituent + %25 = OpCompositeConstruct %mat3v3float %constituent_0 %constituent_1 %constituent_2 + OpStore %mat3 %25 +%constituent_3 = OpConvertSToF %float %int_1 +%constituent_4 = OpCompositeConstruct %v3float %constituent_3 %constituent_3 %constituent_3 +%constituent_5 = OpCompositeConstruct %v3float %constituent_3 %constituent_3 %constituent_3 +%constituent_6 = OpCompositeConstruct %v3float %constituent_3 %constituent_3 %constituent_3 + %30 = OpCompositeConstruct %mat3v3float %constituent_4 %constituent_5 %constituent_6 + OpStore %mat3 %30 + OpReturn + OpFunctionEnd diff --git a/spirv_common.hpp b/spirv_common.hpp index acbf299a..1a102dc7 100644 --- a/spirv_common.hpp +++ b/spirv_common.hpp @@ -636,6 +636,10 @@ struct SPIRBlock : IVariant // Used for handling complex continue blocks which have side effects. std::vector> declare_temporary; + // Declare these temporaries, but only conditionally if this block turns out to be + // a complex loop header. + std::vector> potential_declare_temporary; + struct Case { uint32_t value; diff --git a/spirv_cross.cpp b/spirv_cross.cpp index e2fa3a1f..73126e32 100644 --- a/spirv_cross.cpp +++ b/spirv_cross.cpp @@ -3802,6 +3802,16 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry) auto &block_temporaries = this->get(dominating_block).declare_temporary; block_temporaries.emplace_back(handler.result_id_to_type[var.first], var.first); } + else if (blocks.size() > 1) + { + // Keep track of the temporary as we might have to declare this temporary. + // This can happen if the loop header dominates a temporary, but we have a complex fallback loop. + // In this case, the header is actually inside the for (;;) {} block, and we have problems. + // What we need to do is hoist the temporaries outside the for (;;) {} block in case the header block + // declares the temporary. + auto &block_temporaries = this->get(dominating_block).potential_declare_temporary; + block_temporaries.emplace_back(handler.result_id_to_type[var.first], var.first); + } } } diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index 9c574c28..2b4680f8 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -3091,6 +3091,7 @@ string CompilerGLSL::declare_temporary(uint32_t result_type, uint32_t result_id) else { // The result_id has not been made into an expression yet, so use flags interface. + add_local_variable_name(result_id); return join(flags_to_precision_qualifiers_glsl(type, flags), variable_decl(type, to_name(result_id)), " = "); } } @@ -8873,6 +8874,28 @@ void CompilerGLSL::flush_undeclared_variables(SPIRBlock &block) } } +void CompilerGLSL::emit_hoisted_temporaries(vector> &temporaries) +{ + // If we need to force temporaries for certain IDs due to continue blocks, do it before starting loop header. + // Need to sort these to ensure that reference output is stable. + sort(begin(temporaries), end(temporaries), + [](const pair &a, const pair &b) { return a.second < b.second; }); + + for (auto &tmp : temporaries) + { + add_local_variable_name(tmp.second); + auto flags = meta[tmp.second].decoration.decoration_flags; + auto &type = get(tmp.first); + statement(flags_to_precision_qualifiers_glsl(type, flags), variable_decl(type, to_name(tmp.second)), ";"); + + hoisted_temporaries.insert(tmp.second); + forced_temporaries.insert(tmp.second); + + // The temporary might be read from before it's assigned, set up the expression now. + set(tmp.second, to_name(tmp.second), tmp.first, true); + } +} + void CompilerGLSL::emit_block_chain(SPIRBlock &block) { propagate_loop_dominators(block); @@ -8882,20 +8905,7 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block) bool emitted_for_loop_header = false; bool force_complex_continue_block = false; - // If we need to force temporaries for certain IDs due to continue blocks, do it before starting loop header. - // Need to sort these to ensure that reference output is stable. - sort(begin(block.declare_temporary), end(block.declare_temporary), - [](const pair &a, const pair &b) { return a.second < b.second; }); - - for (auto &tmp : block.declare_temporary) - { - auto flags = meta[tmp.second].decoration.decoration_flags; - auto &type = get(tmp.first); - statement(flags_to_precision_qualifiers_glsl(type, flags), variable_decl(type, to_name(tmp.second)), ";"); - - // The temporary might be read from before it's assigned, set up the expression now. - set(tmp.second, to_name(tmp.second), tmp.first, true); - } + emit_hoisted_temporaries(block.declare_temporary); SPIRBlock::ContinueBlockType continue_type = SPIRBlock::ContinueNone; if (block.continue_block) @@ -8944,6 +8954,11 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block) else if (continue_type == SPIRBlock::DoWhileLoop) { flush_undeclared_variables(block); + // We have some temporaries where the loop header is the dominator. + // We risk a case where we have code like: + // for (;;) { create-temporary; break; } consume-temporary; + // so force-declare temporaries here. + emit_hoisted_temporaries(block.potential_declare_temporary); statement("do"); begin_scope(); @@ -8957,6 +8972,11 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block) get(block.continue_block).complex_continue = true; continue_type = SPIRBlock::ComplexLoop; + // We have some temporaries where the loop header is the dominator. + // We risk a case where we have code like: + // for (;;) { create-temporary; break; } consume-temporary; + // so force-declare temporaries here. + emit_hoisted_temporaries(block.potential_declare_temporary); statement("for (;;)"); begin_scope(); diff --git a/spirv_glsl.hpp b/spirv_glsl.hpp index 0243fb70..5547a3bf 100644 --- a/spirv_glsl.hpp +++ b/spirv_glsl.hpp @@ -366,6 +366,7 @@ protected: void emit_interface_block(const SPIRVariable &type); void emit_flattened_io_block(const SPIRVariable &var, const char *qual); void emit_block_chain(SPIRBlock &block); + void emit_hoisted_temporaries(std::vector> &temporaries); void emit_specialization_constant(const SPIRConstant &constant); std::string emit_continue_block(uint32_t continue_block); bool attempt_emit_loop_header(SPIRBlock &block, SPIRBlock::Method method);