Update glslang/SPIRV-Tools on Travis.

checkout_glslang_spirv_tools.sh

@@ -1,7 +1,7 @@
 #!/bin/bash
 
-GLSLANG_REV=845860d56513d95e15fe4820df7272f9687d076e
+GLSLANG_REV=7cec64fc42eba4587d2f3d73c03cb68ceb18dca4
-SPIRV_TOOLS_REV=340370eddbb9f0e7d26b6a4f7e22c1b98150e5e1
+SPIRV_TOOLS_REV=fe9121f72144f1e1657cd21a55a1fec07c940e56
 
 if [ -d external/glslang ]; then
 	echo "Updating glslang to revision $GLSLANG_REV."

reference/opt/shaders-hlsl/comp/barriers.comp

@@ -1,3 +1,5 @@
+static const uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u);
+
 void comp_main()
 {
     GroupMemoryBarrier();

reference/opt/shaders-hlsl/comp/builtins.comp

@@ -1,3 +1,5 @@
+static const uint3 gl_WorkGroupSize = uint3(8u, 4u, 2u);
+
 void comp_main()
 {
 }

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);
 

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;

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);

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]])
 {
 }

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];

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;

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;

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];

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];

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;
+    float4 _99;
-    _93 = _24.in_data[gl_GlobalInvocationID.x];
+    _99 = _24.in_data[gl_GlobalInvocationID.x];
-    for (int _94 = 0; (_94 + 1) < 10; )
+    for (int _93 = 0; (_93 + 1) < 10; )
     {
-        _93 *= 2.0;
+        _99 *= 2.0;
-        _94 += 2;
+        _93 += 2;
         continue;
     }
-    float4 _95;
+    float4 _98;
-    _95 = _93;
+    _98 = _99;
-    float4 _100;
+    float4 _103;
-    for (uint _97 = 0u; _97 < 16u; _95 = _100, _97++)
+    for (uint _94 = 0u; _94 < 16u; _98 = _103, _94++)
     {
-        _100 = _95;
+        _103 = _98;
-        for (uint _101 = 0u; _101 < 30u; )
+        for (uint _100 = 0u; _100 < 30u; )
         {
-            _100 = _24.mvp * _100;
+            _103 = _24.mvp * _103;
-            _101++;
+            _100++;
             continue;
         }
     }
-    _89.out_data[gl_GlobalInvocationID.x] = _95;
+    _89.out_data[gl_GlobalInvocationID.x] = _98;
 }
 

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;
+    int _170;
-    _142 = 0;
+    _170 = 0;
-    for (; _142 < 20; )
+    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;
+    else
-    int2 _144;
-    _144 = int2(0);
-    for (; _144.x < 10; )
     {
-        out.FragColor += _144.y;
+        out.FragColor += _62;
-        int2 _139 = _144;
-        _139.x = _144.x + 4;
-        _144 = _139;
-        continue;
     }
-    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;
+        int2 _177;
-        _145++;
+        _177 = int2(_214.x ? _212.x : int2(0).x, _214.y ? _212.y : int2(0).y);
-        continue;
+        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;
 }
 

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
 {

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)
+mat4 _235;
-{
+int _245;
-    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;
-}
 
 void main()
 {
-    vec3 _123 = fragWorld;
+    uint _229;
-    _entryPointOutput = GetCascade(_123);
+    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);
 }
 

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]);
+            FragColor += vec4(v0[(_50 + _51) & 3]);
-            _53++;
+            _51++;
             continue;
         }
     }

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;
     }
 }

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;
+    float _180;
-    _181 = _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;
 }
 

reference/opt/shaders/comp/torture-loop.comp

@@ -14,27 +14,27 @@ layout(binding = 1, std430) writeonly buffer SSBO2
 
 void main()
 {
-    vec4 _93;
+    vec4 _99;
-    _93 = _24.in_data[gl_GlobalInvocationID.x];
+    _99 = _24.in_data[gl_GlobalInvocationID.x];
-    for (int _94 = 0; (_94 + 1) < 10; )
+    for (int _93 = 0; (_93 + 1) < 10; )
     {
-        _93 *= 2.0;
+        _99 *= 2.0;
-        _94 += 2;
+        _93 += 2;
         continue;
     }
-    vec4 _95;
+    vec4 _98;
-    _95 = _93;
+    _98 = _99;
-    vec4 _100;
+    vec4 _103;
-    for (uint _97 = 0u; _97 < 16u; _95 = _100, _97++)
+    for (uint _94 = 0u; _94 < 16u; _98 = _103, _94++)
     {
-        _100 = _95;
+        _103 = _98;
-        for (uint _101 = 0u; _101 < 30u; )
+        for (uint _100 = 0u; _100 < 30u; )
         {
-            _100 = _24.mvp * _100;
+            _103 = _24.mvp * _103;
-            _101++;
+            _100++;
             continue;
         }
     }
-    _89.out_data[gl_GlobalInvocationID.x] = _95;
+    _89.out_data[gl_GlobalInvocationID.x] = _98;
 }
 

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;
+        int _93;
-        _96 = 0;
+        _93 = 0;
-        for (; _96 < 3; _96++)
+        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);
+                values3[_92 * 3 * 1 + _93 * 1 + _95] = texture(uTextures[_92 * 3 * 1 + _93 * 1 + _95], vUV);
-                _98++;
+                _95++;
                 continue;
             }
         }

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;
+    int _170;
-    _142 = 0;
+    _170 = 0;
-    for (; _142 < 20; )
+    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;
+    else
-    ivec2 _144;
-    _144 = ivec2(0);
-    for (; _144.x < 10; )
     {
-        FragColor += _144.y;
+        FragColor += _62;
-        ivec2 _139 = _144;
-        _139.x = _144.x + 4;
-        _144 = _139;
-        continue;
     }
-    for (int _145 = _62; _145 < 40; )
+    bool _213 = _68 ? true : false;
+    if (!_213)
     {
-        FragColor += _145;
+        ivec2 _177;
-        _145++;
+        _177 = mix(ivec2(0), mix(_184, _184, bvec2(_68)), bvec2(_213));
-        continue;
+        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;
 }
 

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;
+    int _58;
-    for (int _57 = 0, _58 = 0; _58 < vA; FragColor += vec4(1.0), _57 = _60, _58 += (_60 + 10))
+    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;

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;
+        _485 = _366.x;
-    }
-    else
-    {
-        _484 = 0u;
-    }
-    uvec4 _451 = _445;
-    _451.y = _484;
-    uint _485;
-    if (_369)
-    {
-        _485 = _366.y;
     }
     else
     {
         _485 = 0u;
     }
-    uvec4 _457 = _451;
+    uvec4 _451 = _445;
-    _457.z = _485;
+    _451.y = _485;
-    uint _486;
+    uint _487;
-    if (_379)
+    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;

spirv_common.hpp

@@ -636,6 +636,10 @@ struct SPIRBlock : IVariant
 	// Used for handling complex continue blocks which have side effects.
 	std::vector<std::pair<uint32_t, uint32_t>> declare_temporary;
 
+	// Declare these temporaries, but only conditionally if this block turns out to be
+	// a complex loop header.
+	std::vector<std::pair<uint32_t, uint32_t>> potential_declare_temporary;
+
 	struct Case
 	{
 		uint32_t value;

spirv_cross.cpp

@@ -3802,6 +3802,16 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry)
 				auto &block_temporaries = this->get<SPIRBlock>(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<SPIRBlock>(dominating_block).potential_declare_temporary;
+				block_temporaries.emplace_back(handler.result_id_to_type[var.first], var.first);
+			}
 		}
 	}
 

spirv_glsl.cpp

@@ -8874,6 +8874,28 @@ void CompilerGLSL::flush_undeclared_variables(SPIRBlock &block)
 	}
 }
 
+void CompilerGLSL::emit_hoisted_temporaries(vector<pair<uint32_t, uint32_t>> &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<uint32_t, uint32_t> &a, const pair<uint32_t, uint32_t> &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<SPIRType>(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<SPIRExpression>(tmp.second, to_name(tmp.second), tmp.first, true);
+	}
+}
+
 void CompilerGLSL::emit_block_chain(SPIRBlock &block)
 {
 	propagate_loop_dominators(block);
@@ -8883,21 +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.
+	emit_hoisted_temporaries(block.declare_temporary);
-	// Need to sort these to ensure that reference output is stable.
-	sort(begin(block.declare_temporary), end(block.declare_temporary),
-	     [](const pair<uint32_t, uint32_t> &a, const pair<uint32_t, uint32_t> &b) { return a.second < b.second; });
-
-	for (auto &tmp : block.declare_temporary)
-	{
-		add_local_variable_name(tmp.second);
-		auto flags = meta[tmp.second].decoration.decoration_flags;
-		auto &type = get<SPIRType>(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<SPIRExpression>(tmp.second, to_name(tmp.second), tmp.first, true);
-	}
 
 	SPIRBlock::ContinueBlockType continue_type = SPIRBlock::ContinueNone;
 	if (block.continue_block)
@@ -8946,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();
 
@@ -8959,6 +8972,14 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block)
 		get<SPIRBlock>(block.continue_block).complex_continue = true;
 		continue_type = SPIRBlock::ComplexLoop;
 
+		sort(begin(block.potential_declare_temporary), end(block.potential_declare_temporary),
+		     [](const pair<uint32_t, uint32_t> &a, const pair<uint32_t, uint32_t> &b) { return a.second < b.second; });
+
+		// 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();
 

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<std::pair<uint32_t, uint32_t>> &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);