Merge pull request #514 from KhronosGroup/fix-510

Fix name aliasing for temporary variables.
This commit is contained in:
Hans-Kristian Arntzen 2018-03-24 04:52:18 +01:00 committed by GitHub
commit 761b06bc73
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
31 changed files with 362 additions and 174 deletions

View File

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

View File

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

View File

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

View File

@ -1,3 +1,5 @@
static const uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u);
ByteAddressBuffer _22 : register(t0); ByteAddressBuffer _22 : register(t0);
RWByteAddressBuffer _44 : register(u1); RWByteAddressBuffer _44 : register(u1);

View File

@ -3,6 +3,12 @@
using namespace metal; 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 struct SSBO
{ {
float a; float a;

View File

@ -3,6 +3,8 @@
using namespace metal; using namespace metal;
constant uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u);
kernel void main0() kernel void main0()
{ {
threadgroup_barrier(mem_flags::mem_threadgroup); threadgroup_barrier(mem_flags::mem_threadgroup);

View File

@ -3,6 +3,8 @@
using namespace metal; 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]]) 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]])
{ {
} }

View File

@ -6,6 +6,8 @@
using namespace metal; using namespace metal;
constant uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u);
struct SSBO struct SSBO
{ {
float in_data[1]; float in_data[1];

View File

@ -3,6 +3,8 @@
using namespace metal; using namespace metal;
constant uint3 gl_WorkGroupSize = uint3(32u, 1u, 1u);
struct T1 struct T1
{ {
packed_float3 a; packed_float3 a;

View File

@ -3,6 +3,8 @@
using namespace metal; using namespace metal;
constant uint3 gl_WorkGroupSize = uint3(32u, 1u, 1u);
struct T1 struct T1
{ {
packed_float3 a; packed_float3 a;

View File

@ -3,6 +3,8 @@
using namespace metal; using namespace metal;
constant uint3 gl_WorkGroupSize = uint3(4u, 4u, 1u);
struct SSBO struct SSBO
{ {
float out_data[1]; float out_data[1];

View File

@ -3,6 +3,8 @@
using namespace metal; using namespace metal;
constant uint3 gl_WorkGroupSize = uint3(4u, 1u, 1u);
struct SSBO struct SSBO
{ {
float in_data[1]; float in_data[1];

View File

@ -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]]) 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; 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; continue;
} }
} }
_89.out_data[gl_GlobalInvocationID.x] = _95; _89.out_data[gl_GlobalInvocationID.x] = _98;
} }

View File

@ -3,6 +3,9 @@
using namespace metal; using namespace metal;
constant int2 _184 = {};
constant int _199 = {};
struct main0_out struct main0_out
{ {
int FragColor [[color(0)]]; int FragColor [[color(0)]];
@ -12,56 +15,70 @@ fragment main0_out main0()
{ {
main0_out out = {}; main0_out out = {};
out.FragColor = 16; out.FragColor = 16;
for (int _140 = 0; _140 < 25; ) for (int _168 = 0; _168 < 25; )
{ {
out.FragColor += 10; out.FragColor += 10;
_140++; _168++;
continue; continue;
} }
for (int _141 = 1; _141 < 30; ) for (int _169 = 1; _169 < 30; )
{ {
out.FragColor += 11; out.FragColor += 11;
_141++; _169++;
continue; continue;
} }
int _142; int _170;
_142 = 0; _170 = 0;
for (; _142 < 20; ) for (; _170 < 20; )
{ {
out.FragColor += 12; out.FragColor += 12;
_142++; _170++;
continue; continue;
} }
int _62 = _142 + 3; int _62 = _170 + 3;
out.FragColor += _62; 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; out.FragColor += 13;
_143++; _171++;
continue; 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; return out;
} }

View File

@ -1,5 +1,5 @@
#version 310 es #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 layout(binding = 0, std430) buffer SSBO
{ {

View File

@ -10,41 +10,73 @@ layout(binding = 0, std140) uniform Foo
layout(location = 0) in vec3 fragWorld; layout(location = 0) in vec3 fragWorld;
layout(location = 0) out int _entryPointOutput; 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() 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);
} }

View File

@ -10,10 +10,10 @@ void main()
_50 = 0; _50 = 0;
for (; _50 < 4; _50++) 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; continue;
} }
} }

View File

@ -14,8 +14,8 @@ void main()
{ {
FragColor += vec4(float(_53)); FragColor += vec4(float(_53));
FragColor += vec4(float(_54)); FragColor += vec4(float(_54));
_53 += counter;
_54 += uint(counter); _54 += uint(counter);
_53 += counter;
continue; continue;
} }
} }

View File

@ -0,0 +1,6 @@
#version 450
void main()
{
}

View File

@ -6,7 +6,7 @@ layout(binding = 0, std430) buffer SSBO
float data; float data;
} _11; } _11;
float _180; float _183;
void main() void main()
{ {
@ -31,15 +31,15 @@ void main()
break; break;
} }
} }
float _181; float _180;
_181 = _180; _180 = _183;
for (int _179 = 0; _179 < 20; ) for (int _179 = 0; _179 < 20; )
{ {
_180 += 10.0;
_179++; _179++;
_181 += 10.0;
continue; continue;
} }
_11.data = _181;
_11.data = _180; _11.data = _180;
_11.data = _183;
} }

View File

@ -14,27 +14,27 @@ layout(binding = 1, std430) writeonly buffer SSBO2
void main() 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; 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; continue;
} }
} }
_89.out_data[gl_GlobalInvocationID.x] = _95; _89.out_data[gl_GlobalInvocationID.x] = _98;
} }

View File

@ -13,14 +13,14 @@ void main()
vec4 values3[2 * 3 * 1]; vec4 values3[2 * 3 * 1];
for (; _92 < 2; _92++) 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; continue;
} }
} }

View File

@ -4,58 +4,72 @@ precision highp int;
layout(location = 0) out mediump int FragColor; layout(location = 0) out mediump int FragColor;
ivec2 _184;
int _199;
void main() void main()
{ {
FragColor = 16; FragColor = 16;
for (int _140 = 0; _140 < 25; ) for (int _168 = 0; _168 < 25; )
{ {
FragColor += 10; FragColor += 10;
_140++; _168++;
continue; continue;
} }
for (int _141 = 1; _141 < 30; ) for (int _169 = 1; _169 < 30; )
{ {
FragColor += 11; FragColor += 11;
_141++; _169++;
continue; continue;
} }
int _142; int _170;
_142 = 0; _170 = 0;
for (; _142 < 20; ) for (; _170 < 20; )
{ {
FragColor += 12; FragColor += 12;
_142++; _170++;
continue; continue;
} }
mediump int _62 = _142 + 3; mediump int _62 = _170 + 3;
FragColor += _62; 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; FragColor += 13;
_143++; _171++;
continue; 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;
} }

View File

@ -9,17 +9,17 @@ layout(location = 1) flat in mediump int vB;
void main() void main()
{ {
FragColor = vec4(0.0); 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; continue;
} }
else else
{ {
_60 = ((vB + _58) == 40) ? 60 : _57; _58 = ((vB + _57) == 40) ? 60 : _60;
continue; continue;
} }
continue; continue;

View File

@ -50,7 +50,7 @@ layout(location = 0) in vec4 Position;
layout(location = 0) out vec3 EyeVec; layout(location = 0) out vec3 EyeVec;
layout(location = 1) out vec4 TexCoord; layout(location = 1) out vec4 TexCoord;
uvec4 _483; uvec4 _484;
void main() void main()
{ {
@ -69,42 +69,42 @@ void main()
{ {
_482 = 0u; _482 = 0u;
} }
uvec4 _445 = _483; uvec4 _445 = _484;
_445.x = _482; _445.x = _482;
bool _379 = _359.y < 32u; bool _379 = _359.y < 32u;
uint _484; uint _485;
if (_379) 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 else
{ {
_485 = 0u; _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 else
{ {
_486 = 0u; _487 = 0u;
}
uvec4 _457 = _451;
_457.z = _487;
uint _489;
if (_379)
{
_489 = _366.y;
}
else
{
_489 = 0u;
} }
uvec4 _463 = _457; uvec4 _463 = _457;
_463.w = _486; _463.w = _489;
vec4 _415 = vec4((_359.xyxy + _463) & (~_366).xxyy); 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 _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; vec2 _204 = _197 * _180.NormalTexCoordScale.zw;

View File

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

View File

@ -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

View File

@ -636,6 +636,10 @@ struct SPIRBlock : IVariant
// Used for handling complex continue blocks which have side effects. // Used for handling complex continue blocks which have side effects.
std::vector<std::pair<uint32_t, uint32_t>> declare_temporary; 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 struct Case
{ {
uint32_t value; uint32_t value;

View File

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

View File

@ -3091,6 +3091,7 @@ string CompilerGLSL::declare_temporary(uint32_t result_type, uint32_t result_id)
else else
{ {
// The result_id has not been made into an expression yet, so use flags interface. // 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)), " = "); 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<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) void CompilerGLSL::emit_block_chain(SPIRBlock &block)
{ {
propagate_loop_dominators(block); propagate_loop_dominators(block);
@ -8882,20 +8905,7 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block)
bool emitted_for_loop_header = false; bool emitted_for_loop_header = false;
bool force_complex_continue_block = 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)
{
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; SPIRBlock::ContinueBlockType continue_type = SPIRBlock::ContinueNone;
if (block.continue_block) if (block.continue_block)
@ -8944,6 +8954,11 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block)
else if (continue_type == SPIRBlock::DoWhileLoop) else if (continue_type == SPIRBlock::DoWhileLoop)
{ {
flush_undeclared_variables(block); 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"); statement("do");
begin_scope(); begin_scope();
@ -8957,6 +8972,11 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block)
get<SPIRBlock>(block.continue_block).complex_continue = true; get<SPIRBlock>(block.continue_block).complex_continue = true;
continue_type = SPIRBlock::ComplexLoop; 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 (;;)"); statement("for (;;)");
begin_scope(); begin_scope();

View File

@ -366,6 +366,7 @@ protected:
void emit_interface_block(const SPIRVariable &type); void emit_interface_block(const SPIRVariable &type);
void emit_flattened_io_block(const SPIRVariable &var, const char *qual); void emit_flattened_io_block(const SPIRVariable &var, const char *qual);
void emit_block_chain(SPIRBlock &block); 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); void emit_specialization_constant(const SPIRConstant &constant);
std::string emit_continue_block(uint32_t continue_block); std::string emit_continue_block(uint32_t continue_block);
bool attempt_emit_loop_header(SPIRBlock &block, SPIRBlock::Method method); bool attempt_emit_loop_header(SPIRBlock &block, SPIRBlock::Method method);