AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Merge pull request #496 from KhronosGroup/fix-494

Browse Source 
Handle Phi expressions being invalidated at end of a block.
This commit is contained in:
Hans-Kristian Arntzen 2018-03-09 15:00:22 +01:00 committed by GitHub
commit 236b2fa47c
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
35 changed files with 839 additions and 163 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

57
reference/opt/shaders-hlsl/asm/frag/implicit-read-dep-phi.asm.frag Normal file
View File

@ -0,0 +1,57 @@
Texture2D<float4> uImage : register(t0);
SamplerState _uImage_sampler : register(s0);
static float4 v0;
static float4 FragColor;
struct SPIRV_Cross_Input
{
float4 v0 : TEXCOORD0;
};
struct SPIRV_Cross_Output
{
float4 FragColor : SV_Target0;
};
void frag_main()
{
float phi;
float4 _36;
int _51;
_51 = 0;
phi = 1.0f;
_36 = float4(1.0f, 2.0f, 1.0f, 2.0f);
for (;;)
{
FragColor = _36;
if (_51 < 4)
{
if (v0[_51] > 0.0f)
{
float2 _48 = phi.xx;
_51++;
phi += 2.0f;
_36 = uImage.SampleLevel(_uImage_sampler, _48, 0.0f);
continue;
}
else
{
break;
}
}
else
{
break;
}
}
}
SPIRV_Cross_Output main(SPIRV_Cross_Input stage_input)
{
v0 = stage_input.v0;
frag_main();
SPIRV_Cross_Output stage_output;
stage_output.FragColor = FragColor;
return stage_output;
}

18
reference/opt/shaders-msl/asm/comp/bitcast_sar.asm.comp
View File

@ -17,13 +17,15 @@ struct _4
kernel void main0(device _3& _5 [[buffer(0)]], device _4& _6 [[buffer(1)]])
{
_6._m0 = uint4(int4(_5._m1) >> _5._m0);
_6._m0 = uint4(_5._m0 >> int4(_5._m1));
_6._m0 = uint4(int4(_5._m1) >> int4(_5._m1));
_6._m0 = uint4(_5._m0 >> _5._m0);
_6._m1 = int4(_5._m1) >> int4(_5._m1);
_6._m1 = _5._m0 >> _5._m0;
_6._m1 = int4(_5._m1) >> _5._m0;
_6._m1 = _5._m0 >> int4(_5._m1);
int4 _22 = _5._m0;
uint4 _23 = _5._m1;
_6._m0 = uint4(int4(_23) >> _22);
_6._m0 = uint4(_22 >> int4(_23));
_6._m0 = uint4(int4(_23) >> int4(_23));
_6._m0 = uint4(_22 >> _22);
_6._m1 = int4(_23) >> int4(_23);
_6._m1 = _22 >> _22;
_6._m1 = int4(_23) >> _22;
_6._m1 = _22 >> int4(_23);
}

18
reference/opt/shaders-msl/asm/comp/bitcast_sdiv.asm.comp
View File

@ -17,13 +17,15 @@ struct _4
kernel void main0(device _3& _5 [[buffer(0)]], device _4& _6 [[buffer(1)]])
{
_6._m0 = uint4(int4(_5._m1) / _5._m0);
_6._m0 = uint4(_5._m0 / int4(_5._m1));
_6._m0 = uint4(int4(_5._m1) / int4(_5._m1));
_6._m0 = uint4(_5._m0 / _5._m0);
_6._m1 = int4(_5._m1) / int4(_5._m1);
_6._m1 = _5._m0 / _5._m0;
_6._m1 = int4(_5._m1) / _5._m0;
_6._m1 = _5._m0 / int4(_5._m1);
int4 _22 = _5._m0;
uint4 _23 = _5._m1;
_6._m0 = uint4(int4(_23) / _22);
_6._m0 = uint4(_22 / int4(_23));
_6._m0 = uint4(int4(_23) / int4(_23));
_6._m0 = uint4(_22 / _22);
_6._m1 = int4(_23) / int4(_23);
_6._m1 = _22 / _22;
_6._m1 = int4(_23) / _22;
_6._m1 = _22 / int4(_23);
}

18
reference/opt/shaders-msl/asm/comp/bitcast_slr.asm.comp
View File

@ -17,13 +17,15 @@ struct _4
kernel void main0(device _3& _5 [[buffer(0)]], device _4& _6 [[buffer(1)]])
{
_6._m0 = _5._m1 >> uint4(_5._m0);
_6._m0 = uint4(_5._m0) >> _5._m1;
_6._m0 = _5._m1 >> _5._m1;
_6._m0 = uint4(_5._m0) >> uint4(_5._m0);
_6._m1 = int4(_5._m1 >> _5._m1);
_6._m1 = int4(uint4(_5._m0) >> uint4(_5._m0));
_6._m1 = int4(_5._m1 >> uint4(_5._m0));
_6._m1 = int4(uint4(_5._m0) >> _5._m1);
int4 _22 = _5._m0;
uint4 _23 = _5._m1;
_6._m0 = _23 >> uint4(_22);
_6._m0 = uint4(_22) >> _23;
_6._m0 = _23 >> _23;
_6._m0 = uint4(_22) >> uint4(_22);
_6._m1 = int4(_23 >> _23);
_6._m1 = int4(uint4(_22) >> uint4(_22));
_6._m1 = int4(_23 >> uint4(_22));
_6._m1 = int4(uint4(_22) >> _23);
}

50
reference/opt/shaders-msl/asm/frag/implicit-read-dep-phi.asm.frag Normal file
View File

@ -0,0 +1,50 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct main0_in
{
float4 v0 [[user(locn0)]];
};
struct main0_out
{
float4 FragColor [[color(0)]];
};
fragment main0_out main0(main0_in in [[stage_in]], texture2d<float> uImage [[texture(0)]], sampler uImageSmplr [[sampler(0)]])
{
main0_out out = {};
float phi;
float4 _36;
int _51;
_51 = 0;
phi = 1.0;
_36 = float4(1.0, 2.0, 1.0, 2.0);
for (;;)
{
out.FragColor = _36;
if (_51 < 4)
{
if (in.v0[_51] > 0.0)
{
float2 _48 = float2(phi);
_51++;
phi += 2.0;
_36 = uImage.sample(uImageSmplr, _48, level(0.0));
continue;
}
else
{
break;
}
}
else
{
break;
}
}
return out;
}

3
reference/opt/shaders-msl/asm/frag/inliner-dominator-inside-loop.asm.frag
View File

@ -154,6 +154,7 @@ fragment main0_out main0(main0_in in [[stage_in]], constant CB0& _19 [[buffer(0)
_193 = _192;
break;
} while (false);
float4 _194 = _193 * 1.0;
float4 _220;
do
{
@ -179,7 +180,7 @@ fragment main0_out main0(main0_in in [[stage_in]], constant CB0& _19 [[buffer(0)
float3 _253 = float3(_252.x, _252.y, _232.z);
float2 _255 = _253.xy * _165;
float3 _256 = float3(_255.x, _255.y, _253.z);
float3 _271 = ((in.IN_Color.xyz * (_193 * 1.0).xyz) * (1.0 + (_256.x * 0.300000011920928955078125))) * (StudsMapTexture.sample(StudsMapSampler, _156.UvStuds).x * 2.0);
float3 _271 = ((in.IN_Color.xyz * _194.xyz) * (1.0 + (_256.x * 0.300000011920928955078125))) * (StudsMapTexture.sample(StudsMapSampler, _156.UvStuds).x * 2.0);
float4 _298;
do
{

36
reference/opt/shaders-msl/asm/frag/vector-shuffle-oom.asm.frag
View File

@ -114,7 +114,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_129 = _109;
}
float3 _133 = float4(0.0).xyz + (_129 * 0.5);
float3 _130 = _129 * 0.5;
float3 _133 = float4(0.0).xyz + _130;
float4 _134 = float4(_133.x, _133.y, _133.z, float4(0.0).w);
_28 _135 = _77;
_135._m0 = _134;
@ -130,7 +131,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_176 = _156;
}
float3 _180 = _134.xyz + (_176 * 0.5);
float3 _177 = _176 * 0.5;
float3 _180 = _134.xyz + _177;
float4 _181 = float4(_180.x, _180.y, _180.z, _134.w);
_28 _182 = _135;
_182._m0 = _181;
@ -146,7 +148,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_223 = _203;
}
float3 _227 = _181.xyz + (_223 * 0.75);
float3 _224 = _223 * 0.75;
float3 _227 = _181.xyz + _224;
float4 _228 = float4(_227.x, _227.y, _227.z, _181.w);
_28 _229 = _182;
_229._m0 = _228;
@ -162,7 +165,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_270 = _250;
}
float3 _274 = _228.xyz + (_270 * 0.5);
float3 _271 = _270 * 0.5;
float3 _274 = _228.xyz + _271;
float4 _275 = float4(_274.x, _274.y, _274.z, _228.w);
_28 _276 = _229;
_276._m0 = _275;
@ -178,7 +182,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_317 = _297;
}
float3 _321 = _275.xyz + (_317 * 0.5);
float3 _318 = _317 * 0.5;
float3 _321 = _275.xyz + _318;
float4 _322 = float4(_321.x, _321.y, _321.z, _275.w);
_28 _323 = _276;
_323._m0 = _322;
@ -194,7 +199,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_364 = _344;
}
float3 _368 = _322.xyz + (_364 * 0.75);
float3 _365 = _364 * 0.75;
float3 _368 = _322.xyz + _365;
float4 _369 = float4(_368.x, _368.y, _368.z, _322.w);
_28 _370 = _323;
_370._m0 = _369;
@ -210,7 +216,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_411 = _391;
}
float3 _415 = _369.xyz + (_411 * 1.0);
float3 _412 = _411 * 1.0;
float3 _415 = _369.xyz + _412;
float4 _416 = float4(_415.x, _415.y, _415.z, _369.w);
_28 _417 = _370;
_417._m0 = _416;
@ -226,7 +233,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_458 = _438;
}
float3 _462 = _416.xyz + (_458 * 0.75);
float3 _459 = _458 * 0.75;
float3 _462 = _416.xyz + _459;
float4 _463 = float4(_462.x, _462.y, _462.z, _416.w);
_28 _464 = _417;
_464._m0 = _463;
@ -242,7 +250,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_505 = _485;
}
float3 _509 = _463.xyz + (_505 * 0.5);
float3 _506 = _505 * 0.5;
float3 _509 = _463.xyz + _506;
float4 _510 = float4(_509.x, _509.y, _509.z, _463.w);
_28 _511 = _464;
_511._m0 = _510;
@ -258,7 +267,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_552 = _532;
}
float3 _556 = _510.xyz + (_552 * 0.5);
float3 _553 = _552 * 0.5;
float3 _556 = _510.xyz + _553;
float4 _557 = float4(_556.x, _556.y, _556.z, _510.w);
_28 _558 = _511;
_558._m0 = _557;
@ -274,7 +284,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_599 = _579;
}
float3 _603 = _557.xyz + (_599 * 0.75);
float3 _600 = _599 * 0.75;
float3 _603 = _557.xyz + _600;
float4 _604 = float4(_603.x, _603.y, _603.z, _557.w);
_28 _605 = _558;
_605._m0 = _604;
@ -290,7 +301,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_646 = _626;
}
float3 _650 = _604.xyz + (_646 * 0.5);
float3 _647 = _646 * 0.5;
float3 _650 = _604.xyz + _647;
float4 _651 = float4(_650.x, _650.y, _650.z, _604.w);
_28 _652 = _605;
_652._m0 = _651;

10
reference/opt/shaders/asm/comp/bitcast_iequal.asm.comp
View File

@ -15,10 +15,12 @@ layout(binding = 1, std430) buffer _4_6
void main()
{
bvec4 _34 = equal(ivec4(_5._m1), _5._m0);
bvec4 _35 = equal(_5._m0, ivec4(_5._m1));
bvec4 _36 = equal(_5._m1, _5._m1);
bvec4 _37 = equal(_5._m0, _5._m0);
ivec4 _30 = _5._m0;
uvec4 _31 = _5._m1;
bvec4 _34 = equal(ivec4(_31), _30);
bvec4 _35 = equal(_30, ivec4(_31));
bvec4 _36 = equal(_31, _31);
bvec4 _37 = equal(_30, _30);
_6._m0 = mix(uvec4(0u), uvec4(1u), _34);
_6._m0 = mix(uvec4(0u), uvec4(1u), _35);
_6._m0 = mix(uvec4(0u), uvec4(1u), _36);

18
reference/opt/shaders/asm/comp/bitcast_sar.asm.comp
View File

@ -15,13 +15,15 @@ layout(binding = 1, std430) buffer _4_6
void main()
{
_6._m0 = uvec4(ivec4(_5._m1) >> _5._m0);
_6._m0 = uvec4(_5._m0 >> ivec4(_5._m1));
_6._m0 = uvec4(ivec4(_5._m1) >> ivec4(_5._m1));
_6._m0 = uvec4(_5._m0 >> _5._m0);
_6._m1 = ivec4(_5._m1) >> ivec4(_5._m1);
_6._m1 = _5._m0 >> _5._m0;
_6._m1 = ivec4(_5._m1) >> _5._m0;
_6._m1 = _5._m0 >> ivec4(_5._m1);
ivec4 _22 = _5._m0;
uvec4 _23 = _5._m1;
_6._m0 = uvec4(ivec4(_23) >> _22);
_6._m0 = uvec4(_22 >> ivec4(_23));
_6._m0 = uvec4(ivec4(_23) >> ivec4(_23));
_6._m0 = uvec4(_22 >> _22);
_6._m1 = ivec4(_23) >> ivec4(_23);
_6._m1 = _22 >> _22;
_6._m1 = ivec4(_23) >> _22;
_6._m1 = _22 >> ivec4(_23);
}

18
reference/opt/shaders/asm/comp/bitcast_sdiv.asm.comp
View File

@ -15,13 +15,15 @@ layout(binding = 1, std430) buffer _4_6
void main()
{
_6._m0 = uvec4(ivec4(_5._m1) / _5._m0);
_6._m0 = uvec4(_5._m0 / ivec4(_5._m1));
_6._m0 = uvec4(ivec4(_5._m1) / ivec4(_5._m1));
_6._m0 = uvec4(_5._m0 / _5._m0);
_6._m1 = ivec4(_5._m1) / ivec4(_5._m1);
_6._m1 = _5._m0 / _5._m0;
_6._m1 = ivec4(_5._m1) / _5._m0;
_6._m1 = _5._m0 / ivec4(_5._m1);
ivec4 _22 = _5._m0;
uvec4 _23 = _5._m1;
_6._m0 = uvec4(ivec4(_23) / _22);
_6._m0 = uvec4(_22 / ivec4(_23));
_6._m0 = uvec4(ivec4(_23) / ivec4(_23));
_6._m0 = uvec4(_22 / _22);
_6._m1 = ivec4(_23) / ivec4(_23);
_6._m1 = _22 / _22;
_6._m1 = ivec4(_23) / _22;
_6._m1 = _22 / ivec4(_23);
}

18
reference/opt/shaders/asm/comp/bitcast_slr.asm.comp
View File

@ -15,13 +15,15 @@ layout(binding = 1, std430) buffer _4_6
void main()
{
_6._m0 = _5._m1 >> uvec4(_5._m0);
_6._m0 = uvec4(_5._m0) >> _5._m1;
_6._m0 = _5._m1 >> _5._m1;
_6._m0 = uvec4(_5._m0) >> uvec4(_5._m0);
_6._m1 = ivec4(_5._m1 >> _5._m1);
_6._m1 = ivec4(uvec4(_5._m0) >> uvec4(_5._m0));
_6._m1 = ivec4(_5._m1 >> uvec4(_5._m0));
_6._m1 = ivec4(uvec4(_5._m0) >> _5._m1);
ivec4 _22 = _5._m0;
uvec4 _23 = _5._m1;
_6._m0 = _23 >> uvec4(_22);
_6._m0 = uvec4(_22) >> _23;
_6._m0 = _23 >> _23;
_6._m0 = uvec4(_22) >> uvec4(_22);
_6._m1 = ivec4(_23 >> _23);
_6._m1 = ivec4(uvec4(_22) >> uvec4(_22));
_6._m1 = ivec4(_23 >> uvec4(_22));
_6._m1 = ivec4(uvec4(_22) >> _23);
}

40
reference/opt/shaders/asm/frag/implicit-read-dep-phi.asm.frag Normal file
View File

@ -0,0 +1,40 @@
#version 450
layout(binding = 0) uniform sampler2D uImage;
layout(location = 0) in vec4 v0;
layout(location = 0) out vec4 FragColor;
void main()
{
float phi;
vec4 _36;
int _51;
_51 = 0;
phi = 1.0;
_36 = vec4(1.0, 2.0, 1.0, 2.0);
for (;;)
{
FragColor = _36;
if (_51 < 4)
{
if (v0[_51] > 0.0)
{
vec2 _48 = vec2(phi);
_51++;
phi += 2.0;
_36 = textureLod(uImage, _48, 0.0);
continue;
}
else
{
break;
}
}
else
{
break;
}
}
}

3
reference/opt/shaders/asm/frag/inliner-dominator-inside-loop.asm.frag
View File

@ -147,6 +147,7 @@ void main()
_193 = _192;
break;
} while (false);
vec4 _194 = _193 * 1.0;
vec4 _220;
do
{
@ -172,7 +173,7 @@ void main()
vec3 _253 = vec3(_252.x, _252.y, _232.z);
vec2 _255 = _253.xy * _165;
vec3 _256 = vec3(_255.x, _255.y, _253.z);
vec3 _271 = ((IN_Color.xyz * (_193 * 1.0).xyz) * (1.0 + (_256.x * 0.300000011920928955078125))) * (texture(SPIRV_Cross_CombinedStudsMapTextureStudsMapSampler, _156.UvStuds).x * 2.0);
vec3 _271 = ((IN_Color.xyz * _194.xyz) * (1.0 + (_256.x * 0.300000011920928955078125))) * (texture(SPIRV_Cross_CombinedStudsMapTextureStudsMapSampler, _156.UvStuds).x * 2.0);
vec4 _298;
do
{

36
reference/opt/shaders/asm/frag/vector-shuffle-oom.asm.frag
View File

@ -110,7 +110,8 @@ void main()
{
_129 = _109;
}
vec3 _133 = vec4(0.0).xyz + (_129 * 0.5);
vec3 _130 = _129 * 0.5;
vec3 _133 = vec4(0.0).xyz + _130;
vec4 _134 = vec4(_133.x, _133.y, _133.z, vec4(0.0).w);
_28 _135 = _77;
_135._m0 = _134;
@ -126,7 +127,8 @@ void main()
{
_176 = _156;
}
vec3 _180 = _134.xyz + (_176 * 0.5);
vec3 _177 = _176 * 0.5;
vec3 _180 = _134.xyz + _177;
vec4 _181 = vec4(_180.x, _180.y, _180.z, _134.w);
_28 _182 = _135;
_182._m0 = _181;
@ -142,7 +144,8 @@ void main()
{
_223 = _203;
}
vec3 _227 = _181.xyz + (_223 * 0.75);
vec3 _224 = _223 * 0.75;
vec3 _227 = _181.xyz + _224;
vec4 _228 = vec4(_227.x, _227.y, _227.z, _181.w);
_28 _229 = _182;
_229._m0 = _228;
@ -158,7 +161,8 @@ void main()
{
_270 = _250;
}
vec3 _274 = _228.xyz + (_270 * 0.5);
vec3 _271 = _270 * 0.5;
vec3 _274 = _228.xyz + _271;
vec4 _275 = vec4(_274.x, _274.y, _274.z, _228.w);
_28 _276 = _229;
_276._m0 = _275;
@ -174,7 +178,8 @@ void main()
{
_317 = _297;
}
vec3 _321 = _275.xyz + (_317 * 0.5);
vec3 _318 = _317 * 0.5;
vec3 _321 = _275.xyz + _318;
vec4 _322 = vec4(_321.x, _321.y, _321.z, _275.w);
_28 _323 = _276;
_323._m0 = _322;
@ -190,7 +195,8 @@ void main()
{
_364 = _344;
}
vec3 _368 = _322.xyz + (_364 * 0.75);
vec3 _365 = _364 * 0.75;
vec3 _368 = _322.xyz + _365;
vec4 _369 = vec4(_368.x, _368.y, _368.z, _322.w);
_28 _370 = _323;
_370._m0 = _369;
@ -206,7 +212,8 @@ void main()
{
_411 = _391;
}
vec3 _415 = _369.xyz + (_411 * 1.0);
vec3 _412 = _411 * 1.0;
vec3 _415 = _369.xyz + _412;
vec4 _416 = vec4(_415.x, _415.y, _415.z, _369.w);
_28 _417 = _370;
_417._m0 = _416;
@ -222,7 +229,8 @@ void main()
{
_458 = _438;
}
vec3 _462 = _416.xyz + (_458 * 0.75);
vec3 _459 = _458 * 0.75;
vec3 _462 = _416.xyz + _459;
vec4 _463 = vec4(_462.x, _462.y, _462.z, _416.w);
_28 _464 = _417;
_464._m0 = _463;
@ -238,7 +246,8 @@ void main()
{
_505 = _485;
}
vec3 _509 = _463.xyz + (_505 * 0.5);
vec3 _506 = _505 * 0.5;
vec3 _509 = _463.xyz + _506;
vec4 _510 = vec4(_509.x, _509.y, _509.z, _463.w);
_28 _511 = _464;
_511._m0 = _510;
@ -254,7 +263,8 @@ void main()
{
_552 = _532;
}
vec3 _556 = _510.xyz + (_552 * 0.5);
vec3 _553 = _552 * 0.5;
vec3 _556 = _510.xyz + _553;
vec4 _557 = vec4(_556.x, _556.y, _556.z, _510.w);
_28 _558 = _511;
_558._m0 = _557;
@ -270,7 +280,8 @@ void main()
{
_599 = _579;
}
vec3 _603 = _557.xyz + (_599 * 0.75);
vec3 _600 = _599 * 0.75;
vec3 _603 = _557.xyz + _600;
vec4 _604 = vec4(_603.x, _603.y, _603.z, _557.w);
_28 _605 = _558;
_605._m0 = _604;
@ -286,7 +297,8 @@ void main()
{
_646 = _626;
}
vec3 _650 = _604.xyz + (_646 * 0.5);
vec3 _647 = _646 * 0.5;
vec3 _650 = _604.xyz + _647;
vec4 _651 = vec4(_650.x, _650.y, _650.z, _604.w);
_28 _652 = _605;
_652._m0 = _651;

56
reference/shaders-hlsl/asm/frag/implicit-read-dep-phi.asm.frag Normal file
View File

@ -0,0 +1,56 @@
Texture2D<float4> uImage : register(t0);
SamplerState _uImage_sampler : register(s0);
static float4 v0;
static float4 FragColor;
struct SPIRV_Cross_Input
{
float4 v0 : TEXCOORD0;
};
struct SPIRV_Cross_Output
{
float4 FragColor : SV_Target0;
};
void frag_main()
{
int i = 0;
float phi;
float4 _36;
phi = 1.0f;
_36 = float4(1.0f, 2.0f, 1.0f, 2.0f);
for (;;)
{
FragColor = _36;
if (i < 4)
{
if (v0[i] > 0.0f)
{
float2 _48 = phi.xx;
i++;
phi += 2.0f;
_36 = uImage.SampleLevel(_uImage_sampler, _48, 0.0f);
continue;
}
else
{
break;
}
}
else
{
break;
}
}
}
SPIRV_Cross_Output main(SPIRV_Cross_Input stage_input)
{
v0 = stage_input.v0;
frag_main();
SPIRV_Cross_Output stage_output;
stage_output.FragColor = FragColor;
return stage_output;
}

18
reference/shaders-msl/asm/comp/bitcast_sar.asm.comp
View File

@ -17,13 +17,15 @@ struct _4
kernel void main0(device _3& _5 [[buffer(0)]], device _4& _6 [[buffer(1)]])
{
_6._m0 = uint4(int4(_5._m1) >> _5._m0);
_6._m0 = uint4(_5._m0 >> int4(_5._m1));
_6._m0 = uint4(int4(_5._m1) >> int4(_5._m1));
_6._m0 = uint4(_5._m0 >> _5._m0);
_6._m1 = int4(_5._m1) >> int4(_5._m1);
_6._m1 = _5._m0 >> _5._m0;
_6._m1 = int4(_5._m1) >> _5._m0;
_6._m1 = _5._m0 >> int4(_5._m1);
int4 _22 = _5._m0;
uint4 _23 = _5._m1;
_6._m0 = uint4(int4(_23) >> _22);
_6._m0 = uint4(_22 >> int4(_23));
_6._m0 = uint4(int4(_23) >> int4(_23));
_6._m0 = uint4(_22 >> _22);
_6._m1 = int4(_23) >> int4(_23);
_6._m1 = _22 >> _22;
_6._m1 = int4(_23) >> _22;
_6._m1 = _22 >> int4(_23);
}

18
reference/shaders-msl/asm/comp/bitcast_sdiv.asm.comp
View File

@ -17,13 +17,15 @@ struct _4
kernel void main0(device _3& _5 [[buffer(0)]], device _4& _6 [[buffer(1)]])
{
_6._m0 = uint4(int4(_5._m1) / _5._m0);
_6._m0 = uint4(_5._m0 / int4(_5._m1));
_6._m0 = uint4(int4(_5._m1) / int4(_5._m1));
_6._m0 = uint4(_5._m0 / _5._m0);
_6._m1 = int4(_5._m1) / int4(_5._m1);
_6._m1 = _5._m0 / _5._m0;
_6._m1 = int4(_5._m1) / _5._m0;
_6._m1 = _5._m0 / int4(_5._m1);
int4 _22 = _5._m0;
uint4 _23 = _5._m1;
_6._m0 = uint4(int4(_23) / _22);
_6._m0 = uint4(_22 / int4(_23));
_6._m0 = uint4(int4(_23) / int4(_23));
_6._m0 = uint4(_22 / _22);
_6._m1 = int4(_23) / int4(_23);
_6._m1 = _22 / _22;
_6._m1 = int4(_23) / _22;
_6._m1 = _22 / int4(_23);
}

18
reference/shaders-msl/asm/comp/bitcast_slr.asm.comp
View File

@ -17,13 +17,15 @@ struct _4
kernel void main0(device _3& _5 [[buffer(0)]], device _4& _6 [[buffer(1)]])
{
_6._m0 = _5._m1 >> uint4(_5._m0);
_6._m0 = uint4(_5._m0) >> _5._m1;
_6._m0 = _5._m1 >> _5._m1;
_6._m0 = uint4(_5._m0) >> uint4(_5._m0);
_6._m1 = int4(_5._m1 >> _5._m1);
_6._m1 = int4(uint4(_5._m0) >> uint4(_5._m0));
_6._m1 = int4(_5._m1 >> uint4(_5._m0));
_6._m1 = int4(uint4(_5._m0) >> _5._m1);
int4 _22 = _5._m0;
uint4 _23 = _5._m1;
_6._m0 = _23 >> uint4(_22);
_6._m0 = uint4(_22) >> _23;
_6._m0 = _23 >> _23;
_6._m0 = uint4(_22) >> uint4(_22);
_6._m1 = int4(_23 >> _23);
_6._m1 = int4(uint4(_22) >> uint4(_22));
_6._m1 = int4(_23 >> uint4(_22));
_6._m1 = int4(uint4(_22) >> _23);
}

49
reference/shaders-msl/asm/frag/implicit-read-dep-phi.asm.frag Normal file
View File

@ -0,0 +1,49 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct main0_in
{
float4 v0 [[user(locn0)]];
};
struct main0_out
{
float4 FragColor [[color(0)]];
};
fragment main0_out main0(main0_in in [[stage_in]], texture2d<float> uImage [[texture(0)]], sampler uImageSmplr [[sampler(0)]])
{
main0_out out = {};
int i = 0;
float phi;
float4 _36;
phi = 1.0;
_36 = float4(1.0, 2.0, 1.0, 2.0);
for (;;)
{
out.FragColor = _36;
if (i < 4)
{
if (in.v0[i] > 0.0)
{
float2 _48 = float2(phi);
i++;
phi += 2.0;
_36 = uImage.sample(uImageSmplr, _48, level(0.0));
continue;
}
else
{
break;
}
}
else
{
break;
}
}
return out;
}

3
reference/shaders-msl/asm/frag/inliner-dominator-inside-loop.asm.frag
View File

@ -154,6 +154,7 @@ fragment main0_out main0(main0_in in [[stage_in]], constant CB0& _19 [[buffer(0)
_193 = _192;
break;
} while (false);
float4 _194 = _193 * 1.0;
float4 _220;
do
{
@ -179,7 +180,7 @@ fragment main0_out main0(main0_in in [[stage_in]], constant CB0& _19 [[buffer(0)
float3 _253 = float3(_252.x, _252.y, _232.z);
float2 _255 = _253.xy * _165;
float3 _256 = float3(_255.x, _255.y, _253.z);
float3 _271 = ((in.IN_Color.xyz * (_193 * 1.0).xyz) * (1.0 + (_256.x * 0.300000011920928955078125))) * (StudsMapTexture.sample(StudsMapSampler, _156.UvStuds).x * 2.0);
float3 _271 = ((in.IN_Color.xyz * _194.xyz) * (1.0 + (_256.x * 0.300000011920928955078125))) * (StudsMapTexture.sample(StudsMapSampler, _156.UvStuds).x * 2.0);
float4 _298;
do
{

36
reference/shaders-msl/asm/frag/vector-shuffle-oom.asm.frag
View File

@ -114,7 +114,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_129 = _109;
}
float3 _133 = float4(0.0).xyz + (_129 * 0.5);
float3 _130 = _129 * 0.5;
float3 _133 = float4(0.0).xyz + _130;
float4 _134 = float4(_133.x, _133.y, _133.z, float4(0.0).w);
_28 _135 = _77;
_135._m0 = _134;
@ -130,7 +131,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_176 = _156;
}
float3 _180 = _134.xyz + (_176 * 0.5);
float3 _177 = _176 * 0.5;
float3 _180 = _134.xyz + _177;
float4 _181 = float4(_180.x, _180.y, _180.z, _134.w);
_28 _182 = _135;
_182._m0 = _181;
@ -146,7 +148,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_223 = _203;
}
float3 _227 = _181.xyz + (_223 * 0.75);
float3 _224 = _223 * 0.75;
float3 _227 = _181.xyz + _224;
float4 _228 = float4(_227.x, _227.y, _227.z, _181.w);
_28 _229 = _182;
_229._m0 = _228;
@ -162,7 +165,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_270 = _250;
}
float3 _274 = _228.xyz + (_270 * 0.5);
float3 _271 = _270 * 0.5;
float3 _274 = _228.xyz + _271;
float4 _275 = float4(_274.x, _274.y, _274.z, _228.w);
_28 _276 = _229;
_276._m0 = _275;
@ -178,7 +182,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_317 = _297;
}
float3 _321 = _275.xyz + (_317 * 0.5);
float3 _318 = _317 * 0.5;
float3 _321 = _275.xyz + _318;
float4 _322 = float4(_321.x, _321.y, _321.z, _275.w);
_28 _323 = _276;
_323._m0 = _322;
@ -194,7 +199,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_364 = _344;
}
float3 _368 = _322.xyz + (_364 * 0.75);
float3 _365 = _364 * 0.75;
float3 _368 = _322.xyz + _365;
float4 _369 = float4(_368.x, _368.y, _368.z, _322.w);
_28 _370 = _323;
_370._m0 = _369;
@ -210,7 +216,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_411 = _391;
}
float3 _415 = _369.xyz + (_411 * 1.0);
float3 _412 = _411 * 1.0;
float3 _415 = _369.xyz + _412;
float4 _416 = float4(_415.x, _415.y, _415.z, _369.w);
_28 _417 = _370;
_417._m0 = _416;
@ -226,7 +233,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_458 = _438;
}
float3 _462 = _416.xyz + (_458 * 0.75);
float3 _459 = _458 * 0.75;
float3 _462 = _416.xyz + _459;
float4 _463 = float4(_462.x, _462.y, _462.z, _416.w);
_28 _464 = _417;
_464._m0 = _463;
@ -242,7 +250,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_505 = _485;
}
float3 _509 = _463.xyz + (_505 * 0.5);
float3 _506 = _505 * 0.5;
float3 _509 = _463.xyz + _506;
float4 _510 = float4(_509.x, _509.y, _509.z, _463.w);
_28 _511 = _464;
_511._m0 = _510;
@ -258,7 +267,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_552 = _532;
}
float3 _556 = _510.xyz + (_552 * 0.5);
float3 _553 = _552 * 0.5;
float3 _556 = _510.xyz + _553;
float4 _557 = float4(_556.x, _556.y, _556.z, _510.w);
_28 _558 = _511;
_558._m0 = _557;
@ -274,7 +284,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_599 = _579;
}
float3 _603 = _557.xyz + (_599 * 0.75);
float3 _600 = _599 * 0.75;
float3 _603 = _557.xyz + _600;
float4 _604 = float4(_603.x, _603.y, _603.z, _557.w);
_28 _605 = _558;
_605._m0 = _604;
@ -290,7 +301,8 @@ fragment main0_out main0(constant _6& _7 [[buffer(0)]], constant _10& _11 [[buff
{
_646 = _626;
}
float3 _650 = _604.xyz + (_646 * 0.5);
float3 _647 = _646 * 0.5;
float3 _650 = _604.xyz + _647;
float4 _651 = float4(_650.x, _650.y, _650.z, _604.w);
_28 _652 = _605;
_652._m0 = _651;

10
reference/shaders/asm/comp/bitcast_iequal.asm.comp
View File

@ -15,10 +15,12 @@ layout(binding = 1, std430) buffer _4_6
void main()
{
bvec4 _34 = equal(ivec4(_5._m1), _5._m0);
bvec4 _35 = equal(_5._m0, ivec4(_5._m1));
bvec4 _36 = equal(_5._m1, _5._m1);
bvec4 _37 = equal(_5._m0, _5._m0);
ivec4 _30 = _5._m0;
uvec4 _31 = _5._m1;
bvec4 _34 = equal(ivec4(_31), _30);
bvec4 _35 = equal(_30, ivec4(_31));
bvec4 _36 = equal(_31, _31);
bvec4 _37 = equal(_30, _30);
_6._m0 = mix(uvec4(0u), uvec4(1u), _34);
_6._m0 = mix(uvec4(0u), uvec4(1u), _35);
_6._m0 = mix(uvec4(0u), uvec4(1u), _36);

18
reference/shaders/asm/comp/bitcast_sar.asm.comp
View File

@ -15,13 +15,15 @@ layout(binding = 1, std430) buffer _4_6
void main()
{
_6._m0 = uvec4(ivec4(_5._m1) >> _5._m0);
_6._m0 = uvec4(_5._m0 >> ivec4(_5._m1));
_6._m0 = uvec4(ivec4(_5._m1) >> ivec4(_5._m1));
_6._m0 = uvec4(_5._m0 >> _5._m0);
_6._m1 = ivec4(_5._m1) >> ivec4(_5._m1);
_6._m1 = _5._m0 >> _5._m0;
_6._m1 = ivec4(_5._m1) >> _5._m0;
_6._m1 = _5._m0 >> ivec4(_5._m1);
ivec4 _22 = _5._m0;
uvec4 _23 = _5._m1;
_6._m0 = uvec4(ivec4(_23) >> _22);
_6._m0 = uvec4(_22 >> ivec4(_23));
_6._m0 = uvec4(ivec4(_23) >> ivec4(_23));
_6._m0 = uvec4(_22 >> _22);
_6._m1 = ivec4(_23) >> ivec4(_23);
_6._m1 = _22 >> _22;
_6._m1 = ivec4(_23) >> _22;
_6._m1 = _22 >> ivec4(_23);
}

18
reference/shaders/asm/comp/bitcast_sdiv.asm.comp
View File

@ -15,13 +15,15 @@ layout(binding = 1, std430) buffer _4_6
void main()
{
_6._m0 = uvec4(ivec4(_5._m1) / _5._m0);
_6._m0 = uvec4(_5._m0 / ivec4(_5._m1));
_6._m0 = uvec4(ivec4(_5._m1) / ivec4(_5._m1));
_6._m0 = uvec4(_5._m0 / _5._m0);
_6._m1 = ivec4(_5._m1) / ivec4(_5._m1);
_6._m1 = _5._m0 / _5._m0;
_6._m1 = ivec4(_5._m1) / _5._m0;
_6._m1 = _5._m0 / ivec4(_5._m1);
ivec4 _22 = _5._m0;
uvec4 _23 = _5._m1;
_6._m0 = uvec4(ivec4(_23) / _22);
_6._m0 = uvec4(_22 / ivec4(_23));
_6._m0 = uvec4(ivec4(_23) / ivec4(_23));
_6._m0 = uvec4(_22 / _22);
_6._m1 = ivec4(_23) / ivec4(_23);
_6._m1 = _22 / _22;
_6._m1 = ivec4(_23) / _22;
_6._m1 = _22 / ivec4(_23);
}

18
reference/shaders/asm/comp/bitcast_slr.asm.comp
View File

@ -15,13 +15,15 @@ layout(binding = 1, std430) buffer _4_6
void main()
{
_6._m0 = _5._m1 >> uvec4(_5._m0);
_6._m0 = uvec4(_5._m0) >> _5._m1;
_6._m0 = _5._m1 >> _5._m1;
_6._m0 = uvec4(_5._m0) >> uvec4(_5._m0);
_6._m1 = ivec4(_5._m1 >> _5._m1);
_6._m1 = ivec4(uvec4(_5._m0) >> uvec4(_5._m0));
_6._m1 = ivec4(_5._m1 >> uvec4(_5._m0));
_6._m1 = ivec4(uvec4(_5._m0) >> _5._m1);
ivec4 _22 = _5._m0;
uvec4 _23 = _5._m1;
_6._m0 = _23 >> uvec4(_22);
_6._m0 = uvec4(_22) >> _23;
_6._m0 = _23 >> _23;
_6._m0 = uvec4(_22) >> uvec4(_22);
_6._m1 = ivec4(_23 >> _23);
_6._m1 = ivec4(uvec4(_22) >> uvec4(_22));
_6._m1 = ivec4(_23 >> uvec4(_22));
_6._m1 = ivec4(uvec4(_22) >> _23);
}

39
reference/shaders/asm/frag/implicit-read-dep-phi.asm.frag Normal file
View File

@ -0,0 +1,39 @@
#version 450
layout(binding = 0) uniform sampler2D uImage;
layout(location = 0) in vec4 v0;
layout(location = 0) out vec4 FragColor;
void main()
{
int i = 0;
float phi;
vec4 _36;
phi = 1.0;
_36 = vec4(1.0, 2.0, 1.0, 2.0);
for (;;)
{
FragColor = _36;
if (i < 4)
{
if (v0[i] > 0.0)
{
vec2 _48 = vec2(phi);
i++;
phi += 2.0;
_36 = textureLod(uImage, _48, 0.0);
continue;
}
else
{
break;
}
}
else
{
break;
}
}
}

3
reference/shaders/asm/frag/inliner-dominator-inside-loop.asm.frag
View File

@ -147,6 +147,7 @@ void main()
_193 = _192;
break;
} while (false);
vec4 _194 = _193 * 1.0;
vec4 _220;
do
{
@ -172,7 +173,7 @@ void main()
vec3 _253 = vec3(_252.x, _252.y, _232.z);
vec2 _255 = _253.xy * _165;
vec3 _256 = vec3(_255.x, _255.y, _253.z);
vec3 _271 = ((IN_Color.xyz * (_193 * 1.0).xyz) * (1.0 + (_256.x * 0.300000011920928955078125))) * (texture(SPIRV_Cross_CombinedStudsMapTextureStudsMapSampler, _156.UvStuds).x * 2.0);
vec3 _271 = ((IN_Color.xyz * _194.xyz) * (1.0 + (_256.x * 0.300000011920928955078125))) * (texture(SPIRV_Cross_CombinedStudsMapTextureStudsMapSampler, _156.UvStuds).x * 2.0);
vec4 _298;
do
{

36
reference/shaders/asm/frag/vector-shuffle-oom.asm.frag
View File

@ -110,7 +110,8 @@ void main()
{
_129 = _109;
}
vec3 _133 = vec4(0.0).xyz + (_129 * 0.5);
vec3 _130 = _129 * 0.5;
vec3 _133 = vec4(0.0).xyz + _130;
vec4 _134 = vec4(_133.x, _133.y, _133.z, vec4(0.0).w);
_28 _135 = _77;
_135._m0 = _134;
@ -126,7 +127,8 @@ void main()
{
_176 = _156;
}
vec3 _180 = _134.xyz + (_176 * 0.5);
vec3 _177 = _176 * 0.5;
vec3 _180 = _134.xyz + _177;
vec4 _181 = vec4(_180.x, _180.y, _180.z, _134.w);
_28 _182 = _135;
_182._m0 = _181;
@ -142,7 +144,8 @@ void main()
{
_223 = _203;
}
vec3 _227 = _181.xyz + (_223 * 0.75);
vec3 _224 = _223 * 0.75;
vec3 _227 = _181.xyz + _224;
vec4 _228 = vec4(_227.x, _227.y, _227.z, _181.w);
_28 _229 = _182;
_229._m0 = _228;
@ -158,7 +161,8 @@ void main()
{
_270 = _250;
}
vec3 _274 = _228.xyz + (_270 * 0.5);
vec3 _271 = _270 * 0.5;
vec3 _274 = _228.xyz + _271;
vec4 _275 = vec4(_274.x, _274.y, _274.z, _228.w);
_28 _276 = _229;
_276._m0 = _275;
@ -174,7 +178,8 @@ void main()
{
_317 = _297;
}
vec3 _321 = _275.xyz + (_317 * 0.5);
vec3 _318 = _317 * 0.5;
vec3 _321 = _275.xyz + _318;
vec4 _322 = vec4(_321.x, _321.y, _321.z, _275.w);
_28 _323 = _276;
_323._m0 = _322;
@ -190,7 +195,8 @@ void main()
{
_364 = _344;
}
vec3 _368 = _322.xyz + (_364 * 0.75);
vec3 _365 = _364 * 0.75;
vec3 _368 = _322.xyz + _365;
vec4 _369 = vec4(_368.x, _368.y, _368.z, _322.w);
_28 _370 = _323;
_370._m0 = _369;
@ -206,7 +212,8 @@ void main()
{
_411 = _391;
}
vec3 _415 = _369.xyz + (_411 * 1.0);
vec3 _412 = _411 * 1.0;
vec3 _415 = _369.xyz + _412;
vec4 _416 = vec4(_415.x, _415.y, _415.z, _369.w);
_28 _417 = _370;
_417._m0 = _416;
@ -222,7 +229,8 @@ void main()
{
_458 = _438;
}
vec3 _462 = _416.xyz + (_458 * 0.75);
vec3 _459 = _458 * 0.75;
vec3 _462 = _416.xyz + _459;
vec4 _463 = vec4(_462.x, _462.y, _462.z, _416.w);
_28 _464 = _417;
_464._m0 = _463;
@ -238,7 +246,8 @@ void main()
{
_505 = _485;
}
vec3 _509 = _463.xyz + (_505 * 0.5);
vec3 _506 = _505 * 0.5;
vec3 _509 = _463.xyz + _506;
vec4 _510 = vec4(_509.x, _509.y, _509.z, _463.w);
_28 _511 = _464;
_511._m0 = _510;
@ -254,7 +263,8 @@ void main()
{
_552 = _532;
}
vec3 _556 = _510.xyz + (_552 * 0.5);
vec3 _553 = _552 * 0.5;
vec3 _556 = _510.xyz + _553;
vec4 _557 = vec4(_556.x, _556.y, _556.z, _510.w);
_28 _558 = _511;
_558._m0 = _557;
@ -270,7 +280,8 @@ void main()
{
_599 = _579;
}
vec3 _603 = _557.xyz + (_599 * 0.75);
vec3 _600 = _599 * 0.75;
vec3 _603 = _557.xyz + _600;
vec4 _604 = vec4(_603.x, _603.y, _603.z, _557.w);
_28 _605 = _558;
_605._m0 = _604;
@ -286,7 +297,8 @@ void main()
{
_646 = _626;
}
vec3 _650 = _604.xyz + (_646 * 0.5);
vec3 _647 = _646 * 0.5;
vec3 _650 = _604.xyz + _647;
vec4 _651 = vec4(_650.x, _650.y, _650.z, _604.w);
_28 _652 = _605;
_652._m0 = _651;

3
reference/shaders/tesc/water_tess.tesc
View File

@ -27,7 +27,8 @@ bool frustum_cull(vec2 p0)
vec3 f0 = vec3(dot(_41.uFrustum[0], vec4(center, 1.0)), dot(_41.uFrustum[1], vec4(center, 1.0)), dot(_41.uFrustum[2], vec4(center, 1.0)));
vec3 f1 = vec3(dot(_41.uFrustum[3], vec4(center, 1.0)), dot(_41.uFrustum[4], vec4(center, 1.0)), dot(_41.uFrustum[5], vec4(center, 1.0)));
vec3 _199 = f0;
bool _205 = any(lessThanEqual(_199, vec3(-radius)));
float _200 = radius;
bool _205 = any(lessThanEqual(_199, vec3(-_200)));
bool _215;
if (!_205)
{

81
shaders-hlsl/asm/frag/implicit-read-dep-phi.asm.frag Normal file
View File

@ -0,0 +1,81 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 3
; Bound: 60
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %v0 %FragColor
OpExecutionMode %main OriginUpperLeft
OpSource GLSL 450
OpName %main "main"
OpName %phi "phi"
OpName %i "i"
OpName %v0 "v0"
OpName %FragColor "FragColor"
OpName %uImage "uImage"
OpDecorate %v0 Location 0
OpDecorate %FragColor Location 0
OpDecorate %uImage DescriptorSet 0
OpDecorate %uImage Binding 0
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%_ptr_Function_float = OpTypePointer Function %float
%float_1 = OpConstant %float 1
%int = OpTypeInt 32 1
%_ptr_Function_int = OpTypePointer Function %int
%int_0 = OpConstant %int 0
%int_4 = OpConstant %int 4
%bool = OpTypeBool
%v4float = OpTypeVector %float 4
%_ptr_Input_v4float = OpTypePointer Input %v4float
%v0 = OpVariable %_ptr_Input_v4float Input
%_ptr_Input_float = OpTypePointer Input %float
%float_0 = OpConstant %float 0
%_ptr_Output_v4float = OpTypePointer Output %v4float
%FragColor = OpVariable %_ptr_Output_v4float Output
%36 = OpTypeImage %float 2D 0 0 0 1 Unknown
%37 = OpTypeSampledImage %36
%_ptr_UniformConstant_37 = OpTypePointer UniformConstant %37
%uImage = OpVariable %_ptr_UniformConstant_37 UniformConstant
%v2float = OpTypeVector %float 2
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%float_2 = OpConstant %float 2
%int_1 = OpConstant %int 1
%float_1_vec = OpConstantComposite %v4float %float_1 %float_2 %float_1 %float_2
%main = OpFunction %void None %3
%5 = OpLabel
%i = OpVariable %_ptr_Function_int Function
OpStore %i %int_0
OpBranch %loop_header
%loop_header = OpLabel
%phi = OpPhi %float %float_1 %5 %phi_plus_2 %continue_block
%tex_phi = OpPhi %v4float %float_1_vec %5 %texture_load_result %continue_block
OpLoopMerge %merge_block %continue_block None
OpBranch %loop_body
%loop_body = OpLabel
OpStore %FragColor %tex_phi
%19 = OpLoad %int %i
%22 = OpSLessThan %bool %19 %int_4
OpBranchConditional %22 %15 %merge_block
%15 = OpLabel
%26 = OpLoad %int %i
%28 = OpAccessChain %_ptr_Input_float %v0 %26
%29 = OpLoad %float %28
%31 = OpFOrdGreaterThan %bool %29 %float_0
OpBranchConditional %31 %continue_block %merge_block
%continue_block = OpLabel
%40 = OpLoad %37 %uImage
%43 = OpCompositeConstruct %v2float %phi %phi
%texture_load_result = OpImageSampleExplicitLod %v4float %40 %43 Lod %float_0
%phi_plus_2 = OpFAdd %float %phi %float_2
%54 = OpLoad %int %i
%56 = OpIAdd %int %54 %int_1
OpStore %i %56
OpBranch %loop_header
%merge_block = OpLabel
OpReturn
OpFunctionEnd

81
shaders-msl/asm/frag/implicit-read-dep-phi.asm.frag Normal file
View File

@ -0,0 +1,81 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 3
; Bound: 60
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %v0 %FragColor
OpExecutionMode %main OriginUpperLeft
OpSource GLSL 450
OpName %main "main"
OpName %phi "phi"
OpName %i "i"
OpName %v0 "v0"
OpName %FragColor "FragColor"
OpName %uImage "uImage"
OpDecorate %v0 Location 0
OpDecorate %FragColor Location 0
OpDecorate %uImage DescriptorSet 0
OpDecorate %uImage Binding 0
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%_ptr_Function_float = OpTypePointer Function %float
%float_1 = OpConstant %float 1
%int = OpTypeInt 32 1
%_ptr_Function_int = OpTypePointer Function %int
%int_0 = OpConstant %int 0
%int_4 = OpConstant %int 4
%bool = OpTypeBool
%v4float = OpTypeVector %float 4
%_ptr_Input_v4float = OpTypePointer Input %v4float
%v0 = OpVariable %_ptr_Input_v4float Input
%_ptr_Input_float = OpTypePointer Input %float
%float_0 = OpConstant %float 0
%_ptr_Output_v4float = OpTypePointer Output %v4float
%FragColor = OpVariable %_ptr_Output_v4float Output
%36 = OpTypeImage %float 2D 0 0 0 1 Unknown
%37 = OpTypeSampledImage %36
%_ptr_UniformConstant_37 = OpTypePointer UniformConstant %37
%uImage = OpVariable %_ptr_UniformConstant_37 UniformConstant
%v2float = OpTypeVector %float 2
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%float_2 = OpConstant %float 2
%int_1 = OpConstant %int 1
%float_1_vec = OpConstantComposite %v4float %float_1 %float_2 %float_1 %float_2
%main = OpFunction %void None %3
%5 = OpLabel
%i = OpVariable %_ptr_Function_int Function
OpStore %i %int_0
OpBranch %loop_header
%loop_header = OpLabel
%phi = OpPhi %float %float_1 %5 %phi_plus_2 %continue_block
%tex_phi = OpPhi %v4float %float_1_vec %5 %texture_load_result %continue_block
OpLoopMerge %merge_block %continue_block None
OpBranch %loop_body
%loop_body = OpLabel
OpStore %FragColor %tex_phi
%19 = OpLoad %int %i
%22 = OpSLessThan %bool %19 %int_4
OpBranchConditional %22 %15 %merge_block
%15 = OpLabel
%26 = OpLoad %int %i
%28 = OpAccessChain %_ptr_Input_float %v0 %26
%29 = OpLoad %float %28
%31 = OpFOrdGreaterThan %bool %29 %float_0
OpBranchConditional %31 %continue_block %merge_block
%continue_block = OpLabel
%40 = OpLoad %37 %uImage
%43 = OpCompositeConstruct %v2float %phi %phi
%texture_load_result = OpImageSampleExplicitLod %v4float %40 %43 Lod %float_0
%phi_plus_2 = OpFAdd %float %phi %float_2
%54 = OpLoad %int %i
%56 = OpIAdd %int %54 %int_1
OpStore %i %56
OpBranch %loop_header
%merge_block = OpLabel
OpReturn
OpFunctionEnd

81
shaders/asm/frag/implicit-read-dep-phi.asm.frag Normal file
View File

@ -0,0 +1,81 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 3
; Bound: 60
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %v0 %FragColor
OpExecutionMode %main OriginUpperLeft
OpSource GLSL 450
OpName %main "main"
OpName %phi "phi"
OpName %i "i"
OpName %v0 "v0"
OpName %FragColor "FragColor"
OpName %uImage "uImage"
OpDecorate %v0 Location 0
OpDecorate %FragColor Location 0
OpDecorate %uImage DescriptorSet 0
OpDecorate %uImage Binding 0
%void = OpTypeVoid
%3 = OpTypeFunction %void
%float = OpTypeFloat 32
%_ptr_Function_float = OpTypePointer Function %float
%float_1 = OpConstant %float 1
%int = OpTypeInt 32 1
%_ptr_Function_int = OpTypePointer Function %int
%int_0 = OpConstant %int 0
%int_4 = OpConstant %int 4
%bool = OpTypeBool
%v4float = OpTypeVector %float 4
%_ptr_Input_v4float = OpTypePointer Input %v4float
%v0 = OpVariable %_ptr_Input_v4float Input
%_ptr_Input_float = OpTypePointer Input %float
%float_0 = OpConstant %float 0
%_ptr_Output_v4float = OpTypePointer Output %v4float
%FragColor = OpVariable %_ptr_Output_v4float Output
%36 = OpTypeImage %float 2D 0 0 0 1 Unknown
%37 = OpTypeSampledImage %36
%_ptr_UniformConstant_37 = OpTypePointer UniformConstant %37
%uImage = OpVariable %_ptr_UniformConstant_37 UniformConstant
%v2float = OpTypeVector %float 2
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%float_2 = OpConstant %float 2
%int_1 = OpConstant %int 1
%float_1_vec = OpConstantComposite %v4float %float_1 %float_2 %float_1 %float_2
%main = OpFunction %void None %3
%5 = OpLabel
%i = OpVariable %_ptr_Function_int Function
OpStore %i %int_0
OpBranch %loop_header
%loop_header = OpLabel
%phi = OpPhi %float %float_1 %5 %phi_plus_2 %continue_block
%tex_phi = OpPhi %v4float %float_1_vec %5 %texture_load_result %continue_block
OpLoopMerge %merge_block %continue_block None
OpBranch %loop_body
%loop_body = OpLabel
OpStore %FragColor %tex_phi
%19 = OpLoad %int %i
%22 = OpSLessThan %bool %19 %int_4
OpBranchConditional %22 %15 %merge_block
%15 = OpLabel
%26 = OpLoad %int %i
%28 = OpAccessChain %_ptr_Input_float %v0 %26
%29 = OpLoad %float %28
%31 = OpFOrdGreaterThan %bool %29 %float_0
OpBranchConditional %31 %continue_block %merge_block
%continue_block = OpLabel
%40 = OpLoad %37 %uImage
%43 = OpCompositeConstruct %v2float %phi %phi
%texture_load_result = OpImageSampleExplicitLod %v4float %40 %43 Lod %float_0
%phi_plus_2 = OpFAdd %float %phi %float_2
%54 = OpLoad %int %i
%56 = OpIAdd %int %54 %int_1
OpStore %i %56
OpBranch %loop_header
%merge_block = OpLabel
OpReturn
OpFunctionEnd

9
spirv_cross.cpp
View File

@ -2647,6 +2647,14 @@ void Compiler::inherit_expression_dependencies(uint32_t dst, uint32_t source_exp
}
auto &e = get<SPIRExpression>(dst);
auto *phi = maybe_get<SPIRVariable>(source_expression);
if (phi && phi->phi_variable)
{
// We have used a phi variable, which can change at the end of the block,
// so make sure we take a dependency on this phi variable.
phi->dependees.push_back(dst);
}
auto *s = maybe_get<SPIRExpression>(source_expression);
if (!s)
return;
@ -2659,6 +2667,7 @@ void Compiler::inherit_expression_dependencies(uint32_t dst, uint32_t source_exp
e_deps.insert(end(e_deps), begin(s_deps), end(s_deps));
// Eliminate duplicated dependencies.
sort(begin(e_deps), end(e_deps));
e_deps.erase(unique(begin(e_deps), end(e_deps)), end(e_deps));
}

49
spirv_glsl.cpp
View File

@ -3239,6 +3239,8 @@ void CompilerGLSL::emit_binary_op_cast(uint32_t result_type, uint32_t result_id,
expr += join(cast_op0, " ", op, " ", cast_op1);
emit_op(result_type, result_id, expr, should_forward(op0) && should_forward(op1));
inherit_expression_dependencies(result_id, op0);
inherit_expression_dependencies(result_id, op1);
}
void CompilerGLSL::emit_unary_func_op(uint32_t result_type, uint32_t result_id, uint32_t op0, const char *op)
@ -3280,6 +3282,8 @@ void CompilerGLSL::emit_binary_func_op_cast(uint32_t result_type, uint32_t resul
}
emit_op(result_type, result_id, expr, should_forward(op0) && should_forward(op1));
inherit_expression_dependencies(result_id, op0);
inherit_expression_dependencies(result_id, op1);
}
void CompilerGLSL::emit_trinary_func_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1,
@ -3503,6 +3507,9 @@ void CompilerGLSL::emit_mix_op(uint32_t result_type, uint32_t id, uint32_t left,
}
emit_op(result_type, id, expr, should_forward(left) && should_forward(right) && should_forward(lerp));
inherit_expression_dependencies(id, left);
inherit_expression_dependencies(id, right);
inherit_expression_dependencies(id, lerp);
}
else
emit_trinary_func_op(result_type, id, left, right, lerp, "mix");
@ -3588,6 +3595,8 @@ void CompilerGLSL::emit_texture_op(const Instruction &i)
if (i.offset + length > spirv.size())
SPIRV_CROSS_THROW("Compiler::parse() opcode out of range.");
vector<uint32_t> inherited_expressions;
uint32_t result_type = ops[0];
uint32_t id = ops[1];
uint32_t img = ops[2];
@ -3599,6 +3608,8 @@ void CompilerGLSL::emit_texture_op(const Instruction &i)
bool fetch = false;
const uint32_t *opt = nullptr;
inherited_expressions.push_back(coord);
switch (op)
{
case OpImageSampleDrefImplicitLod:
@ -3677,6 +3688,9 @@ void CompilerGLSL::emit_texture_op(const Instruction &i)
break;
}
if (dref)
inherited_expressions.push_back(dref);
if (proj)
coord_components++;
if (imgtype.image.arrayed)
@ -3702,6 +3716,7 @@ void CompilerGLSL::emit_texture_op(const Instruction &i)
if (length && (flags & flag))
{
v = *opt++;
inherited_expressions.push_back(v);
length--;
}
};
@ -3725,6 +3740,8 @@ void CompilerGLSL::emit_texture_op(const Instruction &i)
expr += ")";
emit_op(result_type, id, expr, forward);
for (auto &inherit : inherited_expressions)
inherit_expression_dependencies(id, inherit);
}
// Returns the function name for a texture sampling function for the specified image and sampling characteristics.
@ -5803,7 +5820,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
{
uint32_t result_type = ops[0];
uint32_t id = ops[1];
const auto *elems = &ops[2];
const auto * const elems = &ops[2];
length -= 2;
bool forward = true;
@ -5887,6 +5904,8 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
}
emit_op(result_type, id, constructor_op, forward);
for (uint32_t i = 0; i < length; i++)
inherit_expression_dependencies(id, elems[i]);
break;
}
@ -5915,6 +5934,8 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
auto expr = access_chain_internal(ops[2], &ops[3], 1, false);
emit_op(result_type, id, expr, should_forward(ops[2]));
inherit_expression_dependencies(id, ops[2]);
inherit_expression_dependencies(id, ops[3]);
break;
}
@ -5951,12 +5972,14 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
// from expression causing it to be forced to an actual temporary in GLSL.
auto expr = access_chain_internal(ops[2], &ops[3], length, true, true);
auto &e = emit_op(result_type, id, expr, true, !expression_is_forwarded(ops[2]));
inherit_expression_dependencies(id, ops[2]);
e.base_expression = ops[2];
}
else
{
auto expr = access_chain_internal(ops[2], &ops[3], length, true);
emit_op(result_type, id, expr, should_forward(ops[2]), !expression_is_forwarded(ops[2]));
inherit_expression_dependencies(id, ops[2]);
}
break;
}
@ -6008,6 +6031,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
// For pointer types, we copy the pointer itself.
statement(declare_temporary(result_type, id), to_expression(rhs), ";");
set<SPIRExpression>(id, to_name(id), result_type, true);
inherit_expression_dependencies(id, rhs);
}
else
{
@ -6082,6 +6106,8 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
// We inherit the forwardedness from our arguments to avoid flushing out to temporaries when it's not really needed.
emit_op(result_type, id, expr, should_fwd, trivial_forward);
inherit_expression_dependencies(id, vec0);
inherit_expression_dependencies(id, vec1);
break;
}
@ -6483,6 +6509,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
}
emit_op(result_type, id, op, should_forward(arg));
inherit_expression_dependencies(id, arg);
break;
}
@ -6960,6 +6987,10 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
}
else
emit_op(result_type, id, imgexpr, false);
inherit_expression_dependencies(id, ops[2]);
if (type.image.ms)
inherit_expression_dependencies(id, ops[5]);
break;
}
@ -8312,6 +8343,8 @@ void CompilerGLSL::flush_phi(uint32_t from, uint32_t to)
if (!optimize_read_modify_write(lhs, rhs))
statement(lhs, " = ", rhs, ";");
}
register_write(phi.function_variable);
}
}
}
@ -8405,9 +8438,6 @@ void CompilerGLSL::branch(uint32_t from, uint32_t cond, uint32_t true_block, uin
bool true_sub = !is_conditional(true_block);
bool false_sub = !is_conditional(false_block);
// It is possible that a selection merge target also serves as a break/continue block.
// We will not emit break or continue here, but defer that to the outer scope.
if (true_sub)
{
statement("if (", to_expression(cond), ")");
@ -8415,7 +8445,7 @@ void CompilerGLSL::branch(uint32_t from, uint32_t cond, uint32_t true_block, uin
branch(from, true_block);
end_scope();
if (false_sub)
if (false_sub || is_continue(false_block) || is_break(false_block))
{
statement("else");
begin_scope();
@ -8438,7 +8468,14 @@ void CompilerGLSL::branch(uint32_t from, uint32_t cond, uint32_t true_block, uin
branch(from, false_block);
end_scope();
if (flush_phi_required(from, true_block))
if (is_continue(true_block) || is_break(true_block))
{
statement("else");
begin_scope();
branch(from, true_block);
end_scope();
}
else if (flush_phi_required(from, true_block))
{
statement("else");
begin_scope();

15
spirv_hlsl.cpp
View File

@ -2299,6 +2299,8 @@ void CompilerHLSL::emit_texture_op(const Instruction &i)
if (i.offset + length > spirv.size())
SPIRV_CROSS_THROW("Compiler::parse() opcode out of range.");
vector<uint32_t> inherited_expressions;
uint32_t result_type = ops[0];
uint32_t id = ops[1];
uint32_t img = ops[2];
@ -2311,6 +2313,8 @@ void CompilerHLSL::emit_texture_op(const Instruction &i)
auto *combined_image = maybe_get<SPIRCombinedImageSampler>(img);
auto img_expr = to_expression(combined_image ? combined_image->image : img);
inherited_expressions.push_back(coord);
switch (op)
{
case OpImageSampleDrefImplicitLod:
@ -2384,6 +2388,9 @@ void CompilerHLSL::emit_texture_op(const Instruction &i)
break;
}
if (dref)
inherited_expressions.push_back(dref);
if (proj)
coord_components++;
if (imgtype.image.arrayed)
@ -2410,6 +2417,7 @@ void CompilerHLSL::emit_texture_op(const Instruction &i)
if (length && (flags & flag))
{
v = *opt++;
inherited_expressions.push_back(v);
length--;
}
};
@ -2694,6 +2702,9 @@ void CompilerHLSL::emit_texture_op(const Instruction &i)
{
emit_op(result_type, id, expr, forward, false);
}
for (auto &inherit : inherited_expressions)
inherit_expression_dependencies(id, inherit);
}
string CompilerHLSL::to_resource_binding(const SPIRVariable &var)
@ -3911,6 +3922,10 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
}
else
emit_op(result_type, id, imgexpr, false);
inherit_expression_dependencies(id, ops[2]);
if (type.image.ms)
inherit_expression_dependencies(id, ops[5]);
break;
}