wgpu
wgpu copied to clipboard
gfx-rs
Don't switch on a single constant in output consumed by FXC
glslc emits code like OpSwitch 0 for control flow reasons, for example when one branch ends with a OpKill: 
this produces the rather dodgy glsl code:
switch(int(0u)) {
default:
...
if (((_e46.w * _e49.w) < 0.5)) {
discard;
} else {
break;
}
break;
}
}
On certain GPUs I've had this not render the fragment at all (need to double check).
Ideally we should check if there's only one default case and just not produce the switch in that case.
I can confirm - code like this breaks fragment shaders at least when running on DX12 + Nvidia; GL + Nvidia seems to work, Intel's iGPU also seem to be, uhm, "resilient".
For a few days now I've been dealing with my fragment shader breaking out of blue due do this thingie - while I guess naga is technically correct here (as in: the output code looks fine, it's DirectX / Nvidia driver overoptimizing it), could this be marked as bug, not an enhancement, for practical reasons? 😇
I'll allow myself to leave some keywords for people who might stumble upon this, too:
- chrome shader white background,
- chrome complex shader miscompilation,
- empty shader webgl.
Self-contained attempt at a HLSL test:
struct PSInput
{
float4 color : COLOR;
float4 a : COLOR2;
float4 b : COLOR3;
};
float4 PSMain(PSInput input) : SV_TARGET
{
switch(int(0u)) {
default:
float4 _e46 = input.a, _e49 = input.b;
if (((_e46.w * _e49.w) < 0.5)) {
discard;
} else {
break;
}
break;
}
return input.color;
}
Testing that on the shader playground:
- DXC's DXIL output
- looks fine (honestly it would be shocking, given DXC is a Clang front-end)
- FXC's DXBC output
- broken (no conditional discard), maybe
wgpushould warn stronger against using FXC - (it might even make more sense to write a DXBC backend than allow FXC to be used?)
- broken (no conditional discard), maybe
I think I'm seeing this break a hlsl shader on integrated AMD with FXC. Should I try to pursue fixing this or am I likely to encounter other issues with FXC anyway?
While someone would probably not write this, the switch on a constant with just a default case can be produced from wgsl too:
@compute @workgroup_size(1)
fn main() {
var a : i32;
switch 0 {
default {
a = 1;
}
}
a = 3;
}
HLSL output:
[numthreads(1, 1, 1)]
void main()
{
int a = (int)0;
switch(0) {
default: {
a = 1;
break;
}
}
a = 3;
return;
}
@Imberflur what do you mean by "break a hlsl shader"?
The output color of a fragment shader ends up being zero instead of the expected color and it seems to be fixed when I edited out the switch statement using PIX (keeping the body of the default case) or when commenting out the original piece of code in glsl that becomes the body of the switch. When stepping through the shader in Renderdoc/PIX it just skipped the switch statement entirely. I might be able to find the DXBC output in PIX if that would help?
In my case there is this switch being produced:
276 float3 _expr758 = phi_4151_;
277 float4 _expr760 = phi_4149_;
278 float3 _expr761 = _expr760.xyz;
279 switch(asint(0u)) {
280 default: {
281 uint _expr765 = global.member_14[0u];
282 if ((_expr765 == 1u)) {
283 float4 _expr769 = global.member_4;
284 float4 _expr778 = global_3.SampleLevel(global_4, (((_expr769.xy + _expr184.xy) + float2(16.0, 16.0)) / (float2(int2(NagaMipDimensions2D(global_3, 0))) * 32.0)), 0.0);
285 float _expr785 = global.member_6[3u];
286 float _expr789 = global.member_6[2u];
287 float _expr793 = global.member_4[2u];
288 float _expr797 = global.member_3[2u];
289 float _expr802 = (max((_expr797 + 1.0), floor(((((((_expr778.x * 0.00390625) + _expr778.y) * _expr785) + _expr789) - _expr793) + 1.0))) - _expr797);
290 float _expr807 = clamp((_expr802 / pow(max(_expr758.z, 0.0), 2.0)), 0.0, _expr186);
291 float _expr811 = global.member_8[2u];
292 float _expr817 = global.member_8[0u];
293 float _expr830 = global.member_9[2u];
294 phi_4172_ = lerp(((float3(0.0, 0.2, 0.5) * ((lerp(lerp((((_expr817 > 0.0)).xxx ? float3(8.0, 1.5, 0.15) : float3(5.0, 2.0, 1.15)), float3(5.0, 0.75, 0.2), (max(_expr811, 0.0)).xxx), float3(3.8, 3.0, 1.8), (max(-(_expr811), 0.0)).xxx) * max((0.5 - _expr811), 0.0)) + (float3(0.5, 0.5, 1.6) * (max((0.6 - _expr830), 0.0) * 0.1)))) * pow(0.99, max(((_expr802 * 12.0) - (_expr758.z * 200.0)), 0.0))), (_expr761 * exp(((float3(-0.6, -0.04, -0.01) * _expr807) * 0.1))), (pow(0.95, _expr807)).xxx);
295 break;
296 } else {
297 phi_4172_ = _expr761;
298 break;
299 }
300 break;
301 }
302 }
And the generated DXBC seems to skip it:
#line 276
243 0x00001A30: mov r6.xyz, r6.xyzx // r6.x <- _expr758.x; r6.y <- _expr758.y; r6.z <- _expr758.z
#line 303
244 0x00001A44: mov r1.xyz, r1.xyzx // r1.x <- _expr850.x; r1.y <- _expr850.y; r1.z <- _expr850.z
The full HLSL
1 struct NagaConstants {
2 int base_vertex;
3 int base_instance;
4 uint other;
5 };
6 ConstantBuffer<NagaConstants> _NagaConstants: register(b6);
7
8 struct type_10 {
9 row_major float4x4 member;
10 row_major float4x4 member_1;
11 row_major float4x4 member_2;
12 float4 member_3;
13 float4 member_4;
14 float4 member_5;
15 float4 member_6;
16 float4 member_7;
17 float4 member_8;
18 float4 member_9;
19 float4 member_10;
20 float4 member_11;
21 uint4 member_12;
22 float4 member_13;
23 uint4 member_14;
24 int4 member_15;
25 float4 member_16;
26 float4 member_17;
27 float2 member_18;
28 float member_19;
29 uint member_20;
30 float member_21;
31 float member_22;
32 int _end_pad_0;
33 int _end_pad_1;
34 };
35
36 struct type_19 {
37 float4 member;
38 float4 member_1;
39 };
40
41 struct type_21 {
42 type_19 member[20];
43 };
44
45 struct type_25 {
46 row_major float4x4 member;
47 };
48
49 cbuffer global : register(b0) { type_10 global; }
50 Texture2D<float4> global_1 : register(t0);
51 SamplerState global_2 : register(s0);
52 Texture2D<float4> global_3 : register(t1);
53 SamplerState global_4 : register(s1);
54 cbuffer global_5 : register(b1) { type_21 global_5; }
55 Texture2D<float4> global_6 : register(t9);
56 SamplerState global_7 : register(s9);
57 cbuffer global_8 : register(b5) { type_25 global_8; }
58 Texture2D<float4> global_9 : register(t8);
59 SamplerState global_10 : register(s8);
60 static float2 global_11 = (float2)0;
61 Texture2D<uint4> global_12 : register(t10);
62 static float4 global_13 = (float4)0;
63
64 struct FragmentInput_main {
65 float2 member : LOC0;
66 };
67
68 uint2 NagaMipDimensions2D(Texture2D<uint4> tex, uint mip_level)
69 {
70 uint4 ret;
71 tex.GetDimensions(mip_level, ret.x, ret.y, ret.z);
72 return ret.xy;
73 }
74
75 uint2 NagaMipDimensions2D(Texture2D<float4> tex, uint mip_level)
76 {
77 uint4 ret;
78 tex.GetDimensions(mip_level, ret.x, ret.y, ret.z);
79 return ret.xy;
80 }
81
82 void function()
83 {
84 float3 phi_4078_ = (float3)0;
85 float3 phi_4084_ = (float3)0;
86 float3 phi_4097_ = (float3)0;
87 float3 phi_4114_ = (float3)0;
88 float3 phi_4120_ = (float3)0;
89 float3 phi_4128_ = (float3)0;
90 float3 phi_4148_ = (float3)0;
91 float4 phi_4150_ = (float4)0;
92 float3 phi_4151_ = (float3)0;
93 float4 phi_4149_ = (float4)0;
94 float3 phi_4172_ = (float3)0;
95 float3 phi_4180_ = (float3)0;
96 float3 phi_4200_ = (float3)0;
97 uint phi_4195_ = (uint)0;
98 float phi_4196_ = (float)0;
99 float phi_4197_ = (float)0;
100 float3 local = (float3)0;
101 float3 phi_4201_ = (float3)0;
102 float3 local_1 = (float3)0;
103
104 float2 _expr147 = global_11;
105 float4 _expr148 = global_9.Sample(global_10, _expr147);
106 uint2 _expr151 = asuint(int2(NagaMipDimensions2D(global_12, 0)));
107 float2 _expr152 = float2(_expr151);
108 int2 _expr156 = (asint(_expr151) - int2(1, 1));
109 uint4 _expr158 = global_12.Load(int3(clamp(int2((_expr147 * _expr152)), int2(0, 0), _expr156), 0));
110 uint2 _expr161 = asuint(int2(NagaMipDimensions2D(global_6, 0)));
111 float4 _expr168 = global_6.Load(int3(clamp(int2((_expr147 * float2(_expr161))), int2(0, 0), (asint(_expr161) - int2(1, 1))), 0));
112 float2 _expr172 = (((_expr147 * 2.0) - float2(1.0, 1.0)) * float2(1.0, -1.0));
113 float4x4 _expr177 = global_8.member;
114 float4 _expr178 = mul(float4(_expr172.x, _expr172.y, _expr168.x, 1.0), _expr177);
115 float4 _expr181 = (_expr178 / (_expr178.w).xxxx);
116 float3 _expr182 = _expr181.xyz;
117 float4 _expr184 = global.member_3;
118 float3 _expr185 = _expr184.xyz;
119 float _expr186 = distance(_expr182, _expr185);
120 float3 _expr189 = ((_expr182 - _expr185) / (_expr186).xxx);
121 phi_4151_ = _expr189;
122 phi_4149_ = _expr148;
123 if ((_expr148.w < 1.0)) {
124 uint2 _expr194 = asuint(int2(NagaMipDimensions2D(global_9, 0)));
125 float4 _expr197 = global.member_4;
126 float2 _expr198 = _expr197.xy;
127 float2 _expr200 = ((_expr181.xy + _expr198) * 0.1);
128 float _expr203 = global.member_10[0u];
129 float _expr204 = (_expr203 * 0.2);
130 float _expr207 = (_expr204 + (_expr181.x * 0.01));
131 float _expr211 = (_expr207 * 15.0);
132 uint _expr213 = uint(trunc(_expr211));
133 float2 _expr239 = float3(_expr200.x, _expr200.y, _expr207).xy;
134 float4 _expr241 = global_1.SampleLevel(global_2, (_expr239 + float2((float((((3961281721u * _expr213) ^ 3883706873u) * 2047667443u)) * 2.3283064e-10), (float((((3961281721u * (_expr213 + 73u)) ^ 3883706873u) * 2047667443u)) * 2.3283064e-10))), 0.0);
135 float4 _expr244 = global_1.SampleLevel(global_2, (_expr239 + float2((float((((3961281721u * (_expr213 + 1u)) ^ 3883706873u) * 2047667443u)) * 2.3283064e-10), (float((((3961281721u * (_expr213 + 74u)) ^ 3883706873u) * 2047667443u)) * 2.3283064e-10))), 0.0);
136 float2 _expr251 = ((_expr181.yx + _expr197.yx) * 0.1);
137 float _expr254 = (_expr204 + (_expr181.y * 0.01));
138 float _expr258 = (_expr254 * 15.0);
139 uint _expr260 = uint(trunc(_expr258));
140 float2 _expr286 = float3(_expr251.x, _expr251.y, _expr254).xy;
141 float4 _expr288 = global_1.SampleLevel(global_2, (_expr286 + float2((float((((3961281721u * _expr260) ^ 3883706873u) * 2047667443u)) * 2.3283064e-10), (float((((3961281721u * (_expr260 + 73u)) ^ 3883706873u) * 2047667443u)) * 2.3283064e-10))), 0.0);
142 float4 _expr291 = global_1.SampleLevel(global_2, (_expr286 + float2((float((((3961281721u * (_expr260 + 1u)) ^ 3883706873u) * 2047667443u)) * 2.3283064e-10), (float((((3961281721u * (_expr260 + 74u)) ^ 3883706873u) * 2047667443u)) * 2.3283064e-10))), 0.0);
143 if ((_expr158.w == 2u)) {
144 float2 _expr305 = ((((float2(lerp(_expr241.x, _expr244.x, frac(_expr211)), lerp(_expr288.x, _expr291.x, frac(_expr258))) - float2(0.5, 0.5)) * ((_expr189.z < 0.0) ? _expr148.w : 1.0)) * 1.5) / (_expr186).xx);
145 phi_4078_ = normalize((_expr189 + float3(_expr305.x, _expr305.y, 0.0)));
146 } else {
147 phi_4078_ = _expr189;
148 }
149 float3 _expr312 = phi_4078_;
150 float4x4 _expr314 = global.member_2;
151 float3 _expr315 = (_expr185 + _expr312);
152 float4 _expr320 = mul(float4(_expr315.x, _expr315.y, _expr315.z, 1.0), _expr314);
153 float2 _expr328 = ((((_expr320.xy / (max(_expr320.w, 0.0)).xx) * 0.5) * float2(1.0, -1.0)) + float2(0.5, 0.5));
154 float2 _expr338 = float2(_expr194);
155 int2 _expr342 = (asint(_expr194) - int2(1, 1));
156 float4 _expr344 = global_9.Load(int3(clamp(int2((lerp(_expr147, _expr328, (clamp(((1.0 - (abs((_expr328.y - 0.5)) * 2.0)) * 5.0), 0.0, 1.0)).xx) * _expr338)), int2(0, 0), _expr342), 0));
157 bool _expr346 = (_expr344.w < 1.0);
158 float3 _expr348 = ((_expr346).xxx ? _expr312 : _expr189);
159 float4 _expr350 = ((_expr346).xxxx ? _expr344 : _expr148);
160 if ((_expr158.w != 0u)) {
161 float3 _expr356 = reflect(_expr348, ((float3(_expr158.xyz) * float3(0.007874016, 0.007874016, 0.007874016)) - float3(1.0, 1.0, 1.0)));
162 phi_4084_ = _expr356;
163 if ((_expr356.z <= 0.0)) {
164 phi_4084_ = normalize((_expr356 * (float3(1.0, 1.0, 1.0) - abs(float3(0.0, 0.0, 1.0)))));
165 }
166 float3 _expr364 = phi_4084_;
167 float3 _expr365 = (_expr185 + _expr364);
168 float4 _expr370 = mul(float4(_expr365.x, _expr365.y, _expr365.z, 1.0), _expr314);
169 float2 _expr379 = clamp(((((_expr370.xy / (max(_expr370.w, 0.0)).xx) * 0.5) * float2(1.0, -1.0)) + float2(0.5, 0.5)), float2(0.0, 0.0), float2(1.0, 1.0));
170 uint2 _expr382 = asuint(int2(NagaMipDimensions2D(global_6, 0)));
171 float4 _expr389 = global_6.Load(int3(clamp(int2((_expr379 * float2(_expr382))), int2(0, 0), (asint(_expr382) - int2(1, 1))), 0));
172 float2 _expr393 = (((_expr379 * 2.0) - float2(1.0, 1.0)) * float2(1.0, -1.0));
173 float4 _expr397 = mul(float4(_expr393.x, _expr393.y, _expr389.x, 1.0), _expr177);
174 float3 _expr401 = (_expr397 / (_expr397.w).xxxx).xyz;
175 float _expr414 = (_expr186 * 0.5);
176 float _expr418 = min(clamp(((1.0 - (max(abs((_expr379.y - 0.5)), abs((_expr379.x - 0.5))) * 2.0)) * 6.0), 0.0, 1.0), clamp(((distance(_expr401, _expr185) - _expr414) / _expr414), 0.0, 1.0));
177 if ((_expr418 > 0.0)) {
178 uint4 _expr574 = global_12.Load(int3(clamp(int2((_expr379 * _expr152)), int2(0, 0), _expr156), 0));
179 float4 _expr576 = global.member_8;
180 float4 _expr579 = global.member_9;
181 float _expr605 = global.member_8[0u];
182 bool3 _expr607 = ((_expr605 > 0.0)).xxx;
183 float _expr612 = global.member_8[2u];
184 float _expr613 = max(_expr612, 0.0);
185 float3 _expr615 = (pow(_expr613, 0.2)).xxx;
186 float3 _expr619 = (max(-(_expr612), 0.0)).xxx;
187 float _expr632 = max(_expr364.z, 0.0);
188 float3 _expr634 = lerp(lerp(lerp(lerp((_expr607 ? float3(2.5, 0.3, 0.1) : float3(1.2, 0.3, 0.2)), float3(0.001, 0.005, 0.02), (pow(_expr613, 0.1)).xxx), float3(0.18, 0.28, 0.6), _expr619), lerp(lerp(float3(0.0, 0.1, 0.23), float3(0.002, 0.004, 0.004), _expr615), float3(0.1, 0.2, 0.3), _expr619), (max(-(_expr364.z), 0.0)).xxx), lerp(lerp((_expr607 ? float3(1.06, 0.1, 0.2) : float3(0.1, 0.1, 0.1)), float3(0.001, 0.001, 0.0025), _expr615), float3(0.1, 0.5, 0.9), _expr619), (_expr632).xxx);
189 float3 _expr638 = (min(((_expr634 + (((lerp((((_expr576.x > 0.0)).xxx ? float3(10.2, 3.0, 0.1) : float3(8.2, 3.0, 2.1)), float3(0.25, 0.25, 0.001), (pow(max(-(_expr576.z), 0.0), 0.5)).xxx) * 0.01) * 25.0) * pow(max(dot(_expr364, -(_expr576.xyz)), 0.0), 30.0))) + (float3(0.75, 0.75, 2.5) * pow(max(dot(_expr364, -(_expr579.xyz)), 0.0), 50.0))), float3(1.0, 1.0, 1.0)) * 1.0);
190 if ((_expr574.w == 0u)) {
191 phi_4114_ = _expr638;
192 } else {
193 float4 _expr644 = global_9.Load(int3(clamp(int2((_expr379 * _expr338)), int2(0, 0), _expr342), 0));
194 phi_4114_ = lerp(_expr638, _expr644.xyz, (_expr418).xxx);
195 }
196 float3 _expr649 = phi_4114_;
197 float _expr650 = distance(_expr401, _expr182);
198 switch(asint(0u)) {
199 default: {
200 uint _expr654 = global.member_14[0u];
201 if ((_expr654 == 1u)) {
202 float4 _expr664 = global_3.SampleLevel(global_4, (((_expr198 + _expr184.xy) + float2(16.0, 16.0)) / (float2(int2(NagaMipDimensions2D(global_3, 0))) * 32.0)), 0.0);
203 float _expr671 = global.member_6[3u];
204 float _expr675 = global.member_6[2u];
205 float _expr679 = global.member_4[2u];
206 float _expr683 = global.member_3[2u];
207 float _expr688 = (max((_expr683 + 1.0), floor(((((((_expr664.x * 0.00390625) + _expr664.y) * _expr671) + _expr675) - _expr679) + 1.0))) - _expr683);
208 float _expr691 = clamp((_expr688 / pow(_expr632, 2.0)), 0.0, _expr650);
209 float _expr702 = global.member_9[2u];
210 phi_4120_ = lerp(((float3(0.0, 0.2, 0.5) * ((lerp(lerp((_expr607 ? float3(8.0, 1.5, 0.15) : float3(5.0, 2.0, 1.15)), float3(5.0, 0.75, 0.2), (_expr613).xxx), float3(3.8, 3.0, 1.8), _expr619) * max((0.5 - _expr612), 0.0)) + (float3(0.5, 0.5, 1.6) * (max((0.6 - _expr702), 0.0) * 0.1)))) * pow(0.99, max(((_expr688 * 12.0) - (_expr364.z * 200.0)), 0.0))), (_expr649 * exp(((float3(-0.6, -0.04, -0.01) * _expr691) * 0.1))), (pow(0.95, _expr691)).xxx);
211 break;
212 } else {
213 phi_4120_ = _expr649;
214 break;
215 }
216 break;
217 }
218 }
219 float3 _expr722 = phi_4120_;
220 phi_4128_ = _expr722;
221 if ((_expr650 < 500000.0)) {
222 phi_4128_ = lerp(_expr634, _expr722, ((1.0 / exp((_expr650 * 0.0002)))).xxx);
223 }
224 float3 _expr730 = phi_4128_;
225 phi_4148_ = _expr730;
226 } else {
227 float4 _expr421 = global.member_8;
228 float4 _expr424 = global.member_9;
229 float _expr439 = dot(_expr364, -(_expr421.xyz));
230 float _expr452 = dot(_expr364, -(_expr424.xyz));
231 float _expr464 = global.member_8[0u];
232 bool3 _expr466 = ((_expr464 > 0.0)).xxx;
233 float _expr471 = global.member_8[2u];
234 float _expr472 = max(_expr471, 0.0);
235 float3 _expr474 = (pow(_expr472, 0.2)).xxx;
236 float3 _expr478 = (max(-(_expr471), 0.0)).xxx;
237 float _expr491 = max(_expr364.z, 0.0);
238 float3 _expr493 = lerp(lerp(lerp(lerp((_expr466 ? float3(2.5, 0.3, 0.1) : float3(1.2, 0.3, 0.2)), float3(0.001, 0.005, 0.02), (pow(_expr472, 0.1)).xxx), float3(0.18, 0.28, 0.6), _expr478), lerp(lerp(float3(0.0, 0.1, 0.23), float3(0.002, 0.004, 0.004), _expr474), float3(0.1, 0.2, 0.3), _expr478), (max(-(_expr364.z), 0.0)).xxx), lerp(lerp((_expr466 ? float3(1.06, 0.1, 0.2) : float3(0.1, 0.1, 0.1)), float3(0.001, 0.001, 0.0025), _expr474), float3(0.1, 0.5, 0.9), _expr478), (_expr491).xxx);
239 float3 _expr497 = min((((_expr493 + ((((lerp((((_expr421.x > 0.0)).xxx ? float3(10.2, 3.0, 0.1) : float3(8.2, 3.0, 2.1)), float3(0.25, 0.25, 0.001), (pow(max(-(_expr421.z), 0.0), 0.5)).xxx) * 0.01) * 25.0) * pow(max(_expr439, 0.0), 30.0)) + ((((float3(15.0, 9.0, 3.5) * clamp(((_expr439 - 0.99965) * 11428.572), 0.0, 1.0)) * 5.0) * 1.0) * 0.1))) + ((float3(0.75, 0.75, 2.5) * pow(max(_expr452, 0.0), 50.0)) + ((float3(175.0, 250.0, 375.0) * clamp(((_expr452 - 0.99965) * 11428.572), 0.0, 1.0)) * 0.05))) * 1.0), float3(1.0, 1.0, 1.0));
240 switch(asint(0u)) {
241 default: {
242 uint _expr501 = global.member_14[0u];
243 if ((_expr501 == 1u)) {
244 float4 _expr511 = global_3.SampleLevel(global_4, (((_expr198 + _expr184.xy) + float2(16.0, 16.0)) / (float2(int2(NagaMipDimensions2D(global_3, 0))) * 32.0)), 0.0);
245 float _expr518 = global.member_6[3u];
246 float _expr522 = global.member_6[2u];
247 float _expr526 = global.member_4[2u];
248 float _expr530 = global.member_3[2u];
249 float _expr535 = (max((_expr530 + 1.0), floor(((((((_expr511.x * 0.00390625) + _expr511.y) * _expr518) + _expr522) - _expr526) + 1.0))) - _expr530);
250 float _expr538 = clamp((_expr535 / pow(_expr491, 2.0)), 0.0, 100000.0);
251 float _expr549 = global.member_9[2u];
252 phi_4097_ = lerp(((float3(0.0, 0.2, 0.5) * ((lerp(lerp((_expr466 ? float3(8.0, 1.5, 0.15) : float3(5.0, 2.0, 1.15)), float3(5.0, 0.75, 0.2), (_expr472).xxx), float3(3.8, 3.0, 1.8), _expr478) * max((0.5 - _expr471), 0.0)) + (float3(0.5, 0.5, 1.6) * (max((0.6 - _expr549), 0.0) * 0.1)))) * pow(0.99, max(((_expr535 * 12.0) - (_expr364.z * 200.0)), 0.0))), (_expr497 * exp(((float3(-0.6, -0.04, -0.01) * _expr538) * 0.1))), (pow(0.95, _expr538)).xxx);
253 break;
254 } else {
255 phi_4097_ = _expr497;
256 break;
257 }
258 break;
259 }
260 }
261 float3 _expr569 = phi_4097_;
262 phi_4148_ = lerp(_expr493, _expr569, float3(2.0611537e-9, 2.0611537e-9, 2.0611537e-9));
263 }
264 float3 _expr732 = phi_4148_;
265 float3 _expr736 = lerp(_expr350.xyz, _expr732, (_expr350.w).xxx);
266 float4 _expr742 = float4(_expr736.x, ((float4)0).y, ((float4)0).z, ((float4)0).w);
267 float4 _expr748 = float4(_expr742.x, _expr736.y, _expr742.z, _expr742.w);
268 phi_4150_ = float4(_expr748.x, _expr748.y, _expr736.z, _expr748.w);
269 } else {
270 phi_4150_ = _expr350;
271 }
272 float4 _expr756 = phi_4150_;
273 phi_4151_ = _expr348;
274 phi_4149_ = _expr756;
275 }
276 float3 _expr758 = phi_4151_;
277 float4 _expr760 = phi_4149_;
278 float3 _expr761 = _expr760.xyz;
279 switch(asint(0u)) {
280 default: {
281 uint _expr765 = global.member_14[0u];
282 if ((_expr765 == 1u)) {
283 float4 _expr769 = global.member_4;
284 float4 _expr778 = global_3.SampleLevel(global_4, (((_expr769.xy + _expr184.xy) + float2(16.0, 16.0)) / (float2(int2(NagaMipDimensions2D(global_3, 0))) * 32.0)), 0.0);
285 float _expr785 = global.member_6[3u];
286 float _expr789 = global.member_6[2u];
287 float _expr793 = global.member_4[2u];
288 float _expr797 = global.member_3[2u];
289 float _expr802 = (max((_expr797 + 1.0), floor(((((((_expr778.x * 0.00390625) + _expr778.y) * _expr785) + _expr789) - _expr793) + 1.0))) - _expr797);
290 float _expr807 = clamp((_expr802 / pow(max(_expr758.z, 0.0), 2.0)), 0.0, _expr186);
291 float _expr811 = global.member_8[2u];
292 float _expr817 = global.member_8[0u];
293 float _expr830 = global.member_9[2u];
294 phi_4172_ = lerp(((float3(0.0, 0.2, 0.5) * ((lerp(lerp((((_expr817 > 0.0)).xxx ? float3(8.0, 1.5, 0.15) : float3(5.0, 2.0, 1.15)), float3(5.0, 0.75, 0.2), (max(_expr811, 0.0)).xxx), float3(3.8, 3.0, 1.8), (max(-(_expr811), 0.0)).xxx) * max((0.5 - _expr811), 0.0)) + (float3(0.5, 0.5, 1.6) * (max((0.6 - _expr830), 0.0) * 0.1)))) * pow(0.99, max(((_expr802 * 12.0) - (_expr758.z * 200.0)), 0.0))), (_expr761 * exp(((float3(-0.6, -0.04, -0.01) * _expr807) * 0.1))), (pow(0.95, _expr807)).xxx);
295 break;
296 } else {
297 phi_4172_ = _expr761;
298 break;
299 }
300 break;
301 }
302 }
303 float3 _expr850 = phi_4172_;
304 phi_4180_ = _expr850;
305 if ((_expr186 < 500000.0)) {
306 float _expr854 = global.member_8[0u];
307 bool3 _expr856 = ((_expr854 > 0.0)).xxx;
308 float _expr861 = global.member_8[2u];
309 float _expr862 = max(_expr861, 0.0);
310 float3 _expr864 = (pow(_expr862, 0.2)).xxx;
311 float3 _expr868 = (max(-(_expr861), 0.0)).xxx;
312 phi_4180_ = lerp(lerp(lerp(lerp(lerp((_expr856 ? float3(2.5, 0.3, 0.1) : float3(1.2, 0.3, 0.2)), float3(0.001, 0.005, 0.02), (pow(_expr862, 0.1)).xxx), float3(0.18, 0.28, 0.6), _expr868), lerp(lerp(float3(0.0, 0.1, 0.23), float3(0.002, 0.004, 0.004), _expr864), float3(0.1, 0.2, 0.3), _expr868), (max(-(_expr758.z), 0.0)).xxx), lerp(lerp((_expr856 ? float3(1.06, 0.1, 0.2) : float3(0.1, 0.1, 0.1)), float3(0.001, 0.001, 0.0025), _expr864), float3(0.1, 0.5, 0.9), _expr868), (max(_expr758.z, 0.0)).xxx), _expr850, ((1.0 / exp((_expr186 * 0.0002)))).xxx);
313 }
314 float3 _expr890 = phi_4180_;
315 float4 _expr896 = float4(_expr890.x, ((float4)0).y, ((float4)0).z, ((float4)0).w);
316 float4 _expr902 = float4(_expr896.x, _expr890.y, _expr896.z, _expr896.w);
317 float4 _expr910 = global.member_4;
318 float3 _expr912 = (_expr185 + _expr910.xyz);
319 phi_4200_ = float4(_expr902.x, _expr902.y, _expr890.z, _expr902.w).xyz;
320 phi_4195_ = 0u;
321 bool loop_init = true;
322 while(true) {
323 if (!loop_init) {
324 type_19 _expr924 = global_5.member[phi_4195_];
325 float3 _expr927 = _expr924.member.xyz;
326 float3 _expr934 = (_expr927 - (_expr912 + (_expr758 * min(max(dot((_expr927 - _expr912), _expr758), 0.0), _expr186))));
327 phi_4200_ = (phi_4200_ + (((_expr924.member_1.xyz * pow((1.0 / (0.025 + (((_expr934.x * _expr934.x) + (_expr934.y * _expr934.y)) + (_expr934.z * _expr934.z)))), 1.0)) * 0.002) * 0.2));
328 phi_4195_ = (phi_4195_ + asuint(1));
329 }
330 loop_init = false;
331 float3 _expr915 = phi_4200_;
332 uint _expr917 = phi_4195_;
333 uint _expr920 = global.member_12[0u];
334 local = _expr915;
335 local_1 = _expr915;
336 if ((_expr917 < _expr920)) {
337 continue;
338 } else {
339 break;
340 }
341 }
342 float _expr955 = global.member_17[3u];
343 float _expr958 = global.member_10[0u];
344 bool _expr959 = (_expr958 < _expr955);
345 if (_expr959) {
346 phi_4196_ = ((300000.0 - _expr955) + _expr958);
347 } else {
348 phi_4196_ = (_expr958 - _expr955);
349 }
350 float _expr964 = phi_4196_;
351 float3 _expr1041 = local_1;
352 phi_4201_ = _expr1041;
353 if ((_expr964 < 5.0)) {
354 float4 _expr967 = global.member_17;
355 if (_expr959) {
356 phi_4197_ = ((300000.0 - _expr955) + _expr958);
357 } else {
358 phi_4197_ = (_expr958 - _expr955);
359 }
360 float _expr973 = phi_4197_;
361 float _expr979 = global.member_7[0u];
362 float3 _expr985 = ((float3(0.2, 0.4, 1.0) * ((1000000.0 * max(0.0, (1.0 - _expr973))) * max(sin((_expr979 * 0.4)), 0.0))) * 0.003);
363 float3 _expr990 = (_expr967 + float4(0.0, 0.0, 25.0, 0.0)).xyz;
364 float3 _expr997 = (_expr990 - (_expr912 + (_expr758 * min(max(dot((_expr990 - _expr912), _expr758), 0.0), _expr186))));
365 float3 _expr1014 = local;
366 phi_4201_ = (_expr1014 + (((float4(_expr985.x, _expr985.y, _expr985.z, 1.0).xyz * pow((1.0 / (0.025 + (((_expr997.x * _expr997.x) + (_expr997.y * _expr997.y)) + (_expr997.z * _expr997.z)))), 1.0)) * 0.002) * 0.2));
367 }
368 float3 _expr1017 = phi_4201_;
369 global_13 = float4(_expr1017.x, _expr1017.y, _expr1017.z, 1.0);
370 return;
371 }
372
373 float4 main(FragmentInput_main fragmentinput_main) : SV_Target0
374 {
375 float2 param = fragmentinput_main.member;
376 global_11 = param;
377 function();
378 float4 _expr3 = global_13;
379 return _expr3;
380 }
Full DXBC
//
// Generated by Microsoft (R) HLSL Shader Compiler 10.1
//
//
// Buffer Definitions:
//
// cbuffer global
// {
//
// struct type_10
// {
//
// row_major float4x4 member; // Offset: 0
// row_major float4x4 member_1; // Offset: 64
// row_major float4x4 member_2; // Offset: 128
// float4 member_3; // Offset: 192
// float4 member_4; // Offset: 208
// float4 member_5; // Offset: 224
// float4 member_6; // Offset: 240
// float4 member_7; // Offset: 256
// float4 member_8; // Offset: 272
// float4 member_9; // Offset: 288
// float4 member_10; // Offset: 304
// float4 member_11; // Offset: 320
// uint4 member_12; // Offset: 336
// float4 member_13; // Offset: 352
// uint4 member_14; // Offset: 368
// int4 member_15; // Offset: 384
// float4 member_16; // Offset: 400
// float4 member_17; // Offset: 416
// float2 member_18; // Offset: 432
// float member_19; // Offset: 440
// uint member_20; // Offset: 444
// float member_21; // Offset: 448
// float member_22; // Offset: 452
// int _end_pad_0; // Offset: 456
// int _end_pad_1; // Offset: 460
//
// } global; // Offset: 0 Size: 464
//
// }
//
// cbuffer global_5
// {
//
// struct type_21
// {
//
// struct type_19
// {
//
// float4 member; // Offset: 0
// float4 member_1; // Offset: 16
//
// } member[20]; // Offset: 0
//
// } global_5; // Offset: 0 Size: 640
//
// }
//
// cbuffer global_8
// {
//
// struct type_25
// {
//
// row_major float4x4 member; // Offset: 0
//
// } global_8; // Offset: 0 Size: 64
//
// }
//
//
// Resource Bindings:
//
// Name Type Format Dim ID HLSL Bind Count
// ------------------------------ ---------- ------- ----------- ------- -------------- ------
// global_2 sampler NA NA S0 s0 1
// global_10 sampler NA NA S1 s8 1
// global_1 texture float4 2d T0 t0 1
// global_9 texture float4 2d T1 t8 1
// global_6 texture float4 2d T2 t9 1
// global_12 texture uint4 2d T3 t10 1
// global cbuffer NA NA CB0 cb0 1
// global_5 cbuffer NA NA CB1 cb1 1
// global_8 cbuffer NA NA CB2 cb5 1
//
//
//
// Input signature:
//
// Name Index Mask Register SysValue Format Used
// -------------------- ----- ------ -------- -------- ------- ------
// LOC 0 xy 0 NONE float xy
//
//
// Output signature:
//
// Name Index Mask Register SysValue Format Used
// -------------------- ----- ------ -------- -------- ------- ------
// SV_Target 0 xyzw 0 TARGET float xyzw
//
0x00000000: ps_5_1
0x00000008: dcl_globalFlags refactoringAllowed | skipOptimization
0x0000000C: dcl_constantbuffer CB0[0:0][27], immediateIndexed, space=0
0x00000028: dcl_constantbuffer CB1[1:1][40], dynamicIndexed, space=0
0x00000044: dcl_constantbuffer CB2[5:5][4], immediateIndexed, space=0
0x00000060: dcl_sampler S0[0:0], mode_default, space=0
0x00000078: dcl_sampler S1[8:8], mode_default, space=0
0x00000090: dcl_resource_texture2d (float,float,float,float) T0[0:0], space=0
0x000000AC: dcl_resource_texture2d (float,float,float,float) T1[8:8], space=0
0x000000C8: dcl_resource_texture2d (float,float,float,float) T2[9:9], space=0
0x000000E4: dcl_resource_texture2d (uint,uint,uint,uint) T3[10:10], space=0
0x00000100: dcl_input_ps linear v0.xy
0x0000010C: dcl_output o0.xyzw
0x00000118: dcl_temps 13
//
// Initial variable locations:
// v0.x <- fragmentinput_main.member.x; v0.y <- fragmentinput_main.member.y;
// o0.x <- <main return value>.x; o0.y <- <main return value>.y; o0.z <- <main return value>.z; o0.w <- <main return value>.w
//
#line 375 "D:\veloren\veloren\clouds-frag.glsl"
0 0x00000120: mov r0.xy, v0.xyxx // r0.x <- param.x; r0.y <- param.y
#line 376
1 0x00000134: mov r0.xy, r0.xyxx // r0.x <- global_11.x; r0.y <- global_11.y
#line 377
2 0x00000148: nop
#line 94
3 0x0000014C: itof r1.xyz, l(0, 0, 0, 0) // r1.x <- phi_4172_.x; r1.y <- phi_4172_.y; r1.z <- phi_4172_.z
#line 100
4 0x0000016C: itof r2.xyz, l(0, 0, 0, 0) // r2.x <- local.x; r2.y <- local.y; r2.z <- local.z
#line 102
5 0x0000018C: itof r3.xyz, l(0, 0, 0, 0) // r3.x <- local_1.x; r3.y <- local_1.y; r3.z <- local_1.z
#line 104
6 0x000001AC: mov r0.xy, r0.xyxx // r0.x <- _expr147.x; r0.y <- _expr147.y
#line 105
7 0x000001C0: sample r0.z, v0.xyxx, T1[8].xywz, S1[8] // r0.z <- _expr148.w
#line 106
8 0x000001EC: nop
9 0x000001F0: mov r0.w, l(0)
#line 71
10 0x00000204: resinfo_uint r4.xy, r0.w, T3[10].xyzw
11 0x00000224: mov r4.x, r4.x // r4.x <- ret.x
12 0x00000238: mov r4.y, r4.y // r4.y <- ret.y
#line 72
13 0x0000024C: mov r4.x, r4.x // r4.x <- <NagaMipDimensions2D return value>.x
14 0x00000260: mov r4.y, r4.y // r4.y <- <NagaMipDimensions2D return value>.y
#line 106
15 0x00000274: mov r4.xy, r4.xyxx // r4.x <- _expr151.x; r4.y <- _expr151.y
#line 107
16 0x00000288: utof r4.zw, r4.xxxy // r4.z <- _expr152.x; r4.w <- _expr152.y
#line 108
17 0x0000029C: ineg r5.xy, l(1, 1, 0, 0)
18 0x000002BC: iadd r4.xy, r4.xyxx, r5.xyxx // r4.x <- _expr156.x; r4.y <- _expr156.y
#line 109
19 0x000002D8: mul r4.zw, r0.xxxy, r4.zzzw
20 0x000002F4: ftoi r4.zw, r4.zzzw
21 0x00000308: imax r4.zw, r4.zzzw, l(0, 0, 0, 0)
22 0x00000330: imin r4.xy, r4.xyxx, r4.zwzz
23 0x0000034C: mov r4.zw, l(0,0,0,0)
24 0x0000036C: ld r0.w, r4.xyzw, T3[10].xyzw
25 0x0000038C: mov r0.w, r0.w // r0.w <- _expr158.w
#line 110
26 0x000003A0: nop
27 0x000003A4: mov r1.w, l(0)
#line 78
28 0x000003B8: resinfo_uint r4.xy, r1.w, T2[9].xyzw // r4.x <- ret.x; r4.y <- ret.y
#line 79
29 0x000003D8: mov r4.x, r4.x // r4.x <- <NagaMipDimensions2D return value>.x
30 0x000003EC: mov r4.y, r4.y // r4.y <- <NagaMipDimensions2D return value>.y
#line 110
31 0x00000400: mov r4.xy, r4.xyxx // r4.x <- _expr161.x; r4.y <- _expr161.y
#line 111
32 0x00000414: utof r4.zw, r4.xxxy
33 0x00000428: mul r4.zw, r0.xxxy, r4.zzzw
34 0x00000444: ftoi r4.zw, r4.zzzw
35 0x00000458: ineg r5.xy, l(1, 1, 0, 0)
36 0x00000478: iadd r4.xy, r4.xyxx, r5.xyxx
37 0x00000494: imax r4.zw, r4.zzzw, l(0, 0, 0, 0)
38 0x000004BC: imin r4.xy, r4.xyxx, r4.zwzz
39 0x000004D8: mov r4.zw, l(0,0,0,0)
40 0x000004F8: ld r4.z, r4.xyzw, T2[9].yzxw
41 0x00000518: mov r4.z, r4.z // r4.z <- _expr168.x
#line 112
42 0x0000052C: mul r5.xy, r0.xyxx, l(2.000000, 2.000000, 0.000000, 0.000000)
43 0x00000554: mov r5.zw, l(-0.000000,-0.000000,-1.000000,-1.000000)
44 0x00000574: add r5.xy, r5.zwzz, r5.xyxx
45 0x00000590: mul r4.xy, r5.xyxx, l(1.000000, -1.000000, 0.000000, 0.000000) // r4.x <- _expr172.x; r4.y <- _expr172.y
#line 113
46 0x000005B8: mov r5.x, CB2[5][0].x // r5.x <- _expr177._m00
47 0x000005D4: mov r6.x, CB2[5][0].y // r6.x <- _expr177._m01
48 0x000005F0: mov r7.x, CB2[5][0].z // r7.x <- _expr177._m02
49 0x0000060C: mov r8.x, CB2[5][0].w // r8.x <- _expr177._m03
50 0x00000628: mov r5.y, CB2[5][1].x // r5.y <- _expr177._m10
51 0x00000644: mov r6.y, CB2[5][1].y // r6.y <- _expr177._m11
52 0x00000660: mov r7.y, CB2[5][1].z // r7.y <- _expr177._m12
53 0x0000067C: mov r8.y, CB2[5][1].w // r8.y <- _expr177._m13
54 0x00000698: mov r5.z, CB2[5][2].x // r5.z <- _expr177._m20
55 0x000006B4: mov r6.z, CB2[5][2].y // r6.z <- _expr177._m21
56 0x000006D0: mov r7.z, CB2[5][2].z // r7.z <- _expr177._m22
57 0x000006EC: mov r8.z, CB2[5][2].w // r8.z <- _expr177._m23
58 0x00000708: mov r9.xyzw, CB2[5][3].xyzw // r9.x <- _expr177._m30; r9.y <- _expr177._m31; r9.z <- _expr177._m32; r9.w <- _expr177._m33
#line 114
59 0x00000724: mov r4.w, l(1.000000)
60 0x00000738: mov r5.w, r9.x
61 0x0000074C: dp4 r5.x, r4.xyzw, r5.xyzw // r5.x <- _expr178.x
62 0x00000768: mov r6.w, r9.y
63 0x0000077C: dp4 r5.y, r4.xyzw, r6.xyzw // r5.y <- _expr178.y
64 0x00000798: mov r7.w, r9.z
65 0x000007AC: dp4 r5.z, r4.xyzw, r7.xyzw // r5.z <- _expr178.z
66 0x000007C8: mov r8.w, r9.w
67 0x000007DC: dp4 r1.w, r4.xyzw, r8.xyzw // r1.w <- _expr178.w
#line 115
68 0x000007F8: div r4.xyz, r5.xyzx, r1.wwww // r4.x <- _expr181.x; r4.y <- _expr181.y; r4.z <- _expr181.z
#line 116
69 0x00000814: mov r4.xyz, r4.xyzx // r4.x <- _expr182.x; r4.y <- _expr182.y; r4.z <- _expr182.z
#line 117
70 0x00000828: mov r5.xyz, CB0[0][12].xyzx // r5.x <- _expr184.x; r5.y <- _expr184.y; r5.z <- _expr184.z
#line 118
71 0x00000844: mov r5.xyz, r5.xyzx // r5.x <- _expr185.x; r5.y <- _expr185.y; r5.z <- _expr185.z
#line 119
72 0x00000858: mov r6.xyz, -r5.xyzx
73 0x00000870: add r6.xyz, r4.xyzx, r6.xyzx
74 0x0000088C: dp3 r1.w, r6.xyzx, r6.xyzx
75 0x000008A8: sqrt r1.w, r1.w // r1.w <- _expr186
#line 120
76 0x000008BC: mov r6.xyz, -r5.xyzx
77 0x000008D4: add r6.xyz, r4.xyzx, r6.xyzx
78 0x000008F0: div r6.xyz, r6.xyzx, r1.wwww // r6.x <- _expr189.x; r6.y <- _expr189.y; r6.z <- _expr189.z
#line 121
79 0x0000090C: mov r6.xyz, r6.xyzx // r6.x <- phi_4151_.x; r6.y <- phi_4151_.y; r6.z <- phi_4151_.z
#line 123
80 0x00000920: lt r2.w, r0.z, l(1.000000)
81 0x0000093C: if_nz r2.w
#line 124
82 0x00000948: nop
83 0x0000094C: mov r2.w, l(0)
#line 78
84 0x00000960: resinfo_uint r4.zw, r2.w, T1[8].zwxy // r4.z <- ret.x; r4.w <- ret.y
#line 79
85 0x00000980: mov r4.z, r4.z // r4.z <- <NagaMipDimensions2D return value>.x
86 0x00000994: mov r4.w, r4.w // r4.w <- <NagaMipDimensions2D return value>.y
#line 124
87 0x000009A8: mov r4.zw, r4.zzzw // r4.z <- _expr194.x; r4.w <- _expr194.y
#line 125
88 0x000009BC: mov r7.xy, CB0[0][13].xyxx // r7.x <- _expr197.x; r7.y <- _expr197.y
#line 126
89 0x000009D8: mov r7.xy, r7.xyxx // r7.x <- _expr198.x; r7.y <- _expr198.y
#line 127
90 0x000009EC: add r7.zw, r4.xxxy, r7.xxxy
91 0x00000A08: mul r7.zw, r7.zzzw, l(0.000000, 0.000000, 0.100000, 0.100000) // r7.z <- _expr200.x; r7.w <- _expr200.y
#line 128
92 0x00000A30: mov r2.w, CB0[0][19].x // r2.w <- _expr203
#line 129
93 0x00000A4C: mul r2.w, r2.w, l(0.200000) // r2.w <- _expr204
#line 130
94 0x00000A68: mul r3.w, r4.x, l(0.010000)
95 0x00000A84: add r3.w, r2.w, r3.w // r3.w <- _expr207
#line 131
96 0x00000AA0: mul r3.w, r3.w, l(15.000000) // r3.w <- _expr211
#line 132
97 0x00000ABC: round_z r5.w, r3.w
98 0x00000AD0: ftou r5.w, r5.w // r5.w <- _expr213
#line 133
99 0x00000AE4: mov r7.zw, r7.zzzw // r7.z <- _expr239.x; r7.w <- _expr239.y
#line 134
100 0x00000AF8: imul null, r6.w, r5.w, l(0xec1c5cb9)
101 0x00000B18: xor r6.w, r6.w, l(0xe77ca9f9)
102 0x00000B34: imul null, r6.w, r6.w, l(0x7a0cecf3)
103 0x00000B54: utof r6.w, r6.w
104 0x00000B68: mul r8.x, r6.w, l(0.000000)
105 0x00000B84: iadd r6.w, r5.w, l(73)
106 0x00000BA0: imul null, r6.w, r6.w, l(0xec1c5cb9)
107 0x00000BC0: xor r6.w, r6.w, l(0xe77ca9f9)
108 0x00000BDC: imul null, r6.w, r6.w, l(0x7a0cecf3)
109 0x00000BFC: utof r6.w, r6.w
110 0x00000C10: mul r8.y, r6.w, l(0.000000)
111 0x00000C2C: add r8.xy, r7.zwzz, r8.xyxx
112 0x00000C48: sample_l r6.w, r8.xyxx, T0[0].yzwx, S0[0], l(0.000000) // r6.w <- _expr241.x
#line 135
113 0x00000C7C: iadd r8.x, r5.w, l(1)
114 0x00000C98: imul null, r8.x, r8.x, l(0xec1c5cb9)
115 0x00000CB8: xor r8.x, r8.x, l(0xe77ca9f9)
116 0x00000CD4: imul null, r8.x, r8.x, l(0x7a0cecf3)
117 0x00000CF4: utof r8.x, r8.x
118 0x00000D08: mul r8.x, r8.x, l(0.000000)
119 0x00000D24: iadd r5.w, r5.w, l(74)
120 0x00000D40: imul null, r5.w, r5.w, l(0xec1c5cb9)
121 0x00000D60: xor r5.w, r5.w, l(0xe77ca9f9)
122 0x00000D7C: imul null, r5.w, r5.w, l(0x7a0cecf3)
123 0x00000D9C: utof r5.w, r5.w
124 0x00000DB0: mul r8.y, r5.w, l(0.000000)
125 0x00000DCC: add r7.zw, r7.zzzw, r8.xxxy
126 0x00000DE8: sample_l r5.w, r7.zwzz, T0[0].yzwx, S0[0], l(0.000000) // r5.w <- _expr244.x
#line 136
127 0x00000E1C: add r7.xy, r4.yxyy, r7.yxyy
128 0x00000E38: mul r7.xy, r7.xyxx, l(0.100000, 0.100000, 0.000000, 0.000000) // r7.x <- _expr251.x; r7.y <- _expr251.y
#line 137
129 0x00000E60: mul r4.x, r4.y, l(0.010000)
130 0x00000E7C: add r2.w, r2.w, r4.x // r2.w <- _expr254
#line 138
131 0x00000E98: mul r2.w, r2.w, l(15.000000) // r2.w <- _expr258
#line 139
132 0x00000EB4: round_z r4.x, r2.w
133 0x00000EC8: ftou r4.x, r4.x // r4.x <- _expr260
#line 140
134 0x00000EDC: mov r7.xy, r7.xyxx // r7.x <- _expr286.x; r7.y <- _expr286.y
#line 141
135 0x00000EF0: imul null, r4.y, r4.x, l(0xec1c5cb9)
136 0x00000F10: xor r4.y, r4.y, l(0xe77ca9f9)
137 0x00000F2C: imul null, r4.y, r4.y, l(0x7a0cecf3)
138 0x00000F4C: utof r4.y, r4.y
139 0x00000F60: mul r8.x, r4.y, l(0.000000)
140 0x00000F7C: iadd r4.y, r4.x, l(73)
141 0x00000F98: imul null, r4.y, r4.y, l(0xec1c5cb9)
142 0x00000FB8: xor r4.y, r4.y, l(0xe77ca9f9)
143 0x00000FD4: imul null, r4.y, r4.y, l(0x7a0cecf3)
144 0x00000FF4: utof r4.y, r4.y
145 0x00001008: mul r8.y, r4.y, l(0.000000)
146 0x00001024: add r7.zw, r7.xxxy, r8.xxxy
147 0x00001040: sample_l r4.y, r7.zwzz, T0[0].yxzw, S0[0], l(0.000000) // r4.y <- _expr288.x
#line 142
148 0x00001074: iadd r7.z, r4.x, l(1)
149 0x00001090: imul null, r7.z, r7.z, l(0xec1c5cb9)
150 0x000010B0: xor r7.z, r7.z, l(0xe77ca9f9)
151 0x000010CC: imul null, r7.z, r7.z, l(0x7a0cecf3)
152 0x000010EC: utof r7.z, r7.z
153 0x00001100: mul r8.x, r7.z, l(0.000000)
154 0x0000111C: iadd r4.x, r4.x, l(74)
155 0x00001138: imul null, r4.x, r4.x, l(0xec1c5cb9)
156 0x00001158: xor r4.x, r4.x, l(0xe77ca9f9)
157 0x00001174: imul null, r4.x, r4.x, l(0x7a0cecf3)
158 0x00001194: utof r4.x, r4.x
159 0x000011A8: mul r8.y, r4.x, l(0.000000)
160 0x000011C4: add r7.xy, r7.xyxx, r8.xyxx
161 0x000011E0: sample_l r4.x, r7.xyxx, T0[0].xyzw, S0[0], l(0.000000) // r4.x <- _expr291.x
#line 143
162 0x00001214: ieq r7.x, r0.w, l(2)
163 0x00001230: if_nz r7.x
#line 144
164 0x0000123C: frc r3.w, r3.w
165 0x00001250: mov r7.x, -r6.w
166 0x00001268: add r5.w, r5.w, r7.x
167 0x00001284: mul r3.w, r3.w, r5.w
168 0x000012A0: add r7.x, r3.w, r6.w
169 0x000012BC: frc r2.w, r2.w
170 0x000012D0: mov r3.w, -r4.y
171 0x000012E8: add r3.w, r3.w, r4.x
172 0x00001304: mul r2.w, r2.w, r3.w
173 0x00001320: add r7.y, r2.w, r4.y
174 0x0000133C: mov r4.xy, l(-0.500000,-0.500000,-0.000000,-0.000000)
175 0x0000135C: add r4.xy, r4.xyxx, r7.xyxx
176 0x00001378: lt r2.w, r6.z, l(0.000000)
177 0x00001394: movc r0.z, r2.w, r0.z, l(1.000000)
178 0x000013B8: mul r4.xy, r0.zzzz, r4.xyxx
179 0x000013D4: mul r4.xy, r4.xyxx, l(1.500000, 1.500000, 0.000000, 0.000000)
180 0x000013FC: div r7.xy, r4.xyxx, r1.wwww // r7.x <- _expr305.x; r7.y <- _expr305.y
#line 145
181 0x00001418: mov r7.z, l(0)
182 0x0000142C: add r7.xyz, r6.xyzx, r7.xyzx
183 0x00001448: dp3 r0.z, r7.xyzx, r7.xyzx
184 0x00001464: rsq r0.z, r0.z
185 0x00001478: mul r7.xyz, r0.zzzz, r7.xyzx // r7.x <- phi_4078_.x; r7.y <- phi_4078_.y; r7.z <- phi_4078_.z
#line 146
186 0x00001494: else // Prior locations: r0.z <- _expr148.w; r3.w <- _expr211; r5.w <- _expr244.x; r2.w <- _expr258; r4.y <- _expr288.x; r4.x <- _expr291.x
#line 147
187 0x00001498: mov r7.xyz, r6.xyzx // r7.x <- phi_4078_.x; r7.y <- phi_4078_.y; r7.z <- phi_4078_.z
#line 148
188 0x000014AC: endif
#line 149
189 0x000014B0: mov r7.xyz, r7.xyzx // r7.x <- _expr312.x; r7.y <- _expr312.y; r7.z <- _expr312.z
#line 150
190 0x000014C4: mov r8.x, CB0[0][8].x // r8.x <- _expr314._m00
191 0x000014E0: mov r9.x, CB0[0][8].y // r9.x <- _expr314._m01
192 0x000014FC: mov r10.x, CB0[0][8].w // r10.x <- _expr314._m03
193 0x00001518: mov r8.y, CB0[0][9].x // r8.y <- _expr314._m10
194 0x00001534: mov r9.y, CB0[0][9].y // r9.y <- _expr314._m11
195 0x00001550: mov r10.y, CB0[0][9].w // r10.y <- _expr314._m13
196 0x0000156C: mov r8.z, CB0[0][10].x // r8.z <- _expr314._m20
197 0x00001588: mov r9.z, CB0[0][10].y // r9.z <- _expr314._m21
198 0x000015A4: mov r11.x, CB0[0][10].w // r11.x <- _expr314._m23
199 0x000015C0: mov r11.yzw, CB0[0][11].xxyw // r11.y <- _expr314._m30; r11.z <- _expr314._m31; r11.w <- _expr314._m33
#line 151
200 0x000015DC: add r12.xyz, r5.xyzx, r7.xyzx // r12.x <- _expr315.x; r12.y <- _expr315.y; r12.z <- _expr315.z
#line 152
201 0x000015F8: mov r12.w, l(1.000000)
202 0x0000160C: mov r8.w, r11.y
203 0x00001620: dp4 r4.x, r12.xyzw, r8.xyzw // r4.x <- _expr320.x
204 0x0000163C: mov r9.w, r11.z
205 0x00001650: dp4 r4.y, r12.xyzw, r9.xyzw // r4.y <- _expr320.y
206 0x0000166C: mov r10.zw, r11.xxxw
207 0x00001680: dp4 r0.z, r12.xyzw, r10.xyzw // r0.z <- _expr320.w
#line 153
208 0x0000169C: max r0.z, r0.z, l(0.000000)
209 0x000016B8: div r4.xy, r4.xyxx, r0.zzzz
210 0x000016D4: mul r4.xy, r4.xyxx, l(0.500000, 0.500000, 0.000000, 0.000000)
211 0x000016FC: mul r4.xy, r4.xyxx, l(1.000000, -1.000000, 0.000000, 0.000000)
212 0x00001724: add r4.xy, r4.xyxx, l(0.500000, 0.500000, 0.000000, 0.000000) // r4.x <- _expr328.x; r4.y <- _expr328.y
#line 154
213 0x0000174C: utof r8.xy, r4.zwzz // r8.x <- _expr338.x; r8.y <- _expr338.y
#line 155
214 0x00001760: ineg r8.zw, l(0, 0, 1, 1)
215 0x00001780: iadd r4.zw, r4.zzzw, r8.zzzw // r4.z <- _expr342.x; r4.w <- _expr342.y
#line 156
216 0x0000179C: mov r0.z, l(-0.500000)
217 0x000017B0: add r0.z, r0.z, r4.y
218 0x000017CC: mov r2.w, -r0.z
219 0x000017E4: max r0.z, r0.z, r2.w
220 0x00001800: mul r0.z, r0.z, l(2.000000)
221 0x0000181C: mov r0.z, -r0.z
222 0x00001834: add r0.z, r0.z, l(1.000000)
223 0x00001850: mul r0.z, r0.z, l(5.000000)
224 0x0000186C: max r0.z, r0.z, l(0.000000)
225 0x00001888: min r0.z, r0.z, l(1.000000)
226 0x000018A4: mov r8.zw, -r0.xxxy
227 0x000018BC: add r4.xy, r4.xyxx, r8.zwzz
228 0x000018D8: mul r4.xy, r0.zzzz, r4.xyxx
229 0x000018F4: add r0.xy, r0.xyxx, r4.xyxx
230 0x00001910: mul r0.xy, r8.xyxx, r0.xyxx
231 0x0000192C: ftoi r0.xy, r0.xyxx
232 0x00001940: imax r0.xy, r0.xyxx, l(0, 0, 0, 0)
233 0x00001968: imin r4.xy, r4.zwzz, r0.xyxx
234 0x00001984: mov r4.zw, l(0,0,0,0)
235 0x000019A4: ld r0.x, r4.xyzw, T1[8].wxyz // r0.x <- _expr344.w
#line 157
236 0x000019C4: lt r0.x, r0.x, l(1.000000) // r0.x <- _expr346
#line 158
237 0x000019E0: movc r6.xyz, r0.xxxx, r7.xyzx, r6.xyzx // r6.x <- _expr348.x; r6.y <- _expr348.y; r6.z <- _expr348.z
#line 160
238 0x00001A04: if_nz r0.w
#line 170
239 0x00001A10: nop
#line 271
240 0x00001A14: endif
#line 273
241 0x00001A18: mov r6.xyz, r6.xyzx // r6.x <- phi_4151_.x; r6.y <- phi_4151_.y; r6.z <- phi_4151_.z
#line 275
242 0x00001A2C: endif
#line 276
243 0x00001A30: mov r6.xyz, r6.xyzx // r6.x <- _expr758.x; r6.y <- _expr758.y; r6.z <- _expr758.z
#line 303
244 0x00001A44: mov r1.xyz, r1.xyzx // r1.x <- _expr850.x; r1.y <- _expr850.y; r1.z <- _expr850.z
#line 304
245 0x00001A58: mov r1.xyz, r1.xyzx // r1.x <- phi_4180_.x; r1.y <- phi_4180_.y; r1.z <- phi_4180_.z
#line 305
246 0x00001A6C: lt r0.x, r1.w, l(500000.000000)
247 0x00001A88: if_nz r0.x
#line 306
248 0x00001A94: mov r0.x, CB0[0][17].x // r0.x <- _expr854
#line 307
249 0x00001AB0: lt r0.xyz, l(0.000000, 0.000000, 0.000000, 0.000000), r0.xxxx // r0.z <- _expr856.x
250 0x00001AD8: mov r0.xyz, r0.xyzx // r0.x <- _expr856.x; r0.y <- _expr856.y; r0.z <- _expr856.z
#line 308
251 0x00001AEC: mov r0.w, CB0[0][17].z // r0.w <- _expr861
#line 309
252 0x00001B08: max r2.w, r0.w, l(0.000000) // r2.w <- _expr862
#line 310
253 0x00001B24: log r3.w, r2.w
254 0x00001B38: mul r3.w, r3.w, l(0.200000)
255 0x00001B54: exp r4.xyz, r3.wwww // r4.z <- _expr864.x
256 0x00001B68: mov r4.xyz, r4.xyzx // r4.x <- _expr864.x; r4.y <- _expr864.y; r4.z <- _expr864.z
#line 311
257 0x00001B7C: mov r0.w, -r0.w
258 0x00001B94: max r7.xyz, r0.wwww, l(0.000000, 0.000000, 0.000000, 0.000000) // r7.z <- _expr868.x
259 0x00001BBC: mov r7.xyz, r7.xyzx // r7.x <- _expr868.x; r7.y <- _expr868.y; r7.z <- _expr868.z
#line 312
260 0x00001BD0: movc r8.xyz, r0.xyzx, l(2.500000,0.300000,0.100000,0), l(1.200000,0.300000,0.200000,0)
261 0x00001C0C: log r0.w, r2.w
262 0x00001C20: mul r0.w, r0.w, l(0.100000)
263 0x00001C3C: exp r0.w, r0.w
264 0x00001C50: mov r9.xyz, -r8.xyzx
265 0x00001C68: add r9.xyz, r9.xyzx, l(0.001000, 0.005000, 0.020000, 0.000000)
266 0x00001C90: mul r9.xyz, r0.wwww, r9.xyzx
267 0x00001CAC: add r8.xyz, r8.xyzx, r9.xyzx
268 0x00001CC8: mov r9.xyz, -r8.xyzx
269 0x00001CE0: add r9.xyz, r9.xyzx, l(0.180000, 0.280000, 0.600000, 0.000000)
270 0x00001D08: mul r9.xyz, r7.xyzx, r9.xyzx
271 0x00001D24: add r8.xyz, r8.xyzx, r9.xyzx
272 0x00001D40: mov r9.xyz, l(-0.000000,-0.100000,-0.230000,-0.000000)
273 0x00001D60: add r9.xyz, r9.xyzx, l(0.002000, 0.004000, 0.004000, 0.000000)
274 0x00001D88: mul r9.xyz, r4.xyzx, r9.xyzx
275 0x00001DA4: add r9.xyz, r9.xyzx, l(0.000000, 0.100000, 0.230000, 0.000000)
276 0x00001DCC: mov r10.xyz, -r9.xyzx
277 0x00001DE4: add r10.xyz, r10.xyzx, l(0.100000, 0.200000, 0.300000, 0.000000)
278 0x00001E0C: mul r10.xyz, r7.xyzx, r10.xyzx
279 0x00001E28: add r9.xyz, r9.xyzx, r10.xyzx
280 0x00001E44: mov r0.w, -r6.z
281 0x00001E5C: max r0.w, r0.w, l(0.000000)
282 0x00001E78: mov r10.xyz, -r8.xyzx
283 0x00001E90: add r9.xyz, r9.xyzx, r10.xyzx
284 0x00001EAC: mul r9.xyz, r0.wwww, r9.xyzx
285 0x00001EC8: add r8.xyz, r8.xyzx, r9.xyzx
286 0x00001EE4: movc r0.xyz, r0.xyzx, l(1.060000,0.100000,0.200000,0), l(0.100000,0.100000,0.100000,0)
287 0x00001F20: mov r9.xyz, -r0.xyzx
288 0x00001F38: add r9.xyz, r9.xyzx, l(0.001000, 0.001000, 0.002500, 0.000000)
289 0x00001F60: mul r4.xyz, r4.xyzx, r9.xyzx
290 0x00001F7C: add r0.xyz, r0.xyzx, r4.xyzx
291 0x00001F98: mov r4.xyz, -r0.xyzx
292 0x00001FB0: add r4.xyz, r4.xyzx, l(0.100000, 0.500000, 0.900000, 0.000000)
293 0x00001FD8: mul r4.xyz, r4.xyzx, r7.xyzx
294 0x00001FF4: add r0.xyz, r0.xyzx, r4.xyzx
295 0x00002010: max r0.w, r6.z, l(0.000000)
296 0x0000202C: mov r4.xyz, -r8.xyzx
297 0x00002044: add r0.xyz, r0.xyzx, r4.xyzx
298 0x00002060: mul r0.xyz, r0.xyzx, r0.wwww
299 0x0000207C: add r0.xyz, r0.xyzx, r8.xyzx
300 0x00002098: mul r0.w, r1.w, l(0.000200)
301 0x000020B4: mul r0.w, r0.w, l(1.442695)
302 0x000020D0: exp r0.w, r0.w
303 0x000020E4: div r0.w, l(1.000000), r0.w
304 0x00002100: mov r4.xyz, -r0.xyzx
305 0x00002118: add r4.xyz, r1.xyzx, r4.xyzx
306 0x00002134: mul r4.xyz, r0.wwww, r4.xyzx
307 0x00002150: add r1.xyz, r0.xyzx, r4.xyzx
#line 313
308 0x0000216C: endif
#line 314
309 0x00002170: mov r1.xyz, r1.xyzx // r1.x <- _expr890.x; r1.y <- _expr890.y; r1.z <- _expr890.z
#line 315
310 0x00002184: mov r1.x, r1.x // r1.x <- _expr896.x
#line 316
311 0x00002198: mov r1.x, r1.x // r1.x <- _expr902.x
312 0x000021AC: mov r1.y, r1.y // r1.y <- _expr902.y
#line 317
313 0x000021C0: mov r0.xyz, CB0[0][13].xyzx // r0.x <- _expr910.x; r0.y <- _expr910.y; r0.z <- _expr910.z
#line 318
314 0x000021DC: add r0.xyz, r0.xyzx, r5.xyzx // r0.x <- _expr912.x; r0.y <- _expr912.y; r0.z <- _expr912.z
#line 319
315 0x000021F8: mov r1.xy, r1.xyxx // r1.x <- phi_4200_.x; r1.y <- phi_4200_.y
316 0x0000220C: mov r1.z, r1.z // r1.z <- phi_4200_.z
#line 320
317 0x00002220: mov r0.w, l(0) // r0.w <- phi_4195_
#line 321
318 0x00002234: mov r2.w, l(-1) // r2.w <- loop_init
#line 322
319 0x00002248: mov r4.xyz, r1.xyzx // r4.x <- phi_4200_.x; r4.y <- phi_4200_.y; r4.z <- phi_4200_.z
320 0x0000225C: mov r5.xyz, r2.xyzx // r5.x <- local.x; r5.y <- local.y; r5.z <- local.z
321 0x00002270: mov r7.xyz, r3.xyzx // r7.x <- local_1.x; r7.y <- local_1.y; r7.z <- local_1.z
322 0x00002284: mov r3.w, CB0[0][21].x // r3.w <- global.member_12.x; r3.w <- global.member_12.x
323 0x000022A0: mov r4.w, r0.w // r4.w <- phi_4195_
324 0x000022B4: mov r5.w, r2.w // r5.w <- loop_init
325 0x000022C8: loop
#line 323
326 0x000022CC: if_z r5.w
#line 324
327 0x000022D8: imul null, r6.w, r4.w, l(2)
328 0x000022F8: mov r8.xyz, CB1[1][r6.w + 0].xyzx // r8.x <- _expr924.member.x; r8.y <- _expr924.member.y; r8.z <- _expr924.member.z
329 0x00002318: mov r9.xyz, CB1[1][r6.w + 1].xyzx // r9.x <- _expr924.member_1.x; r9.y <- _expr924.member_1.y; r9.z <- _expr924.member_1.z
#line 325
330 0x0000233C: mov r8.xyz, r8.xyzx // r8.x <- _expr927.x; r8.y <- _expr927.y; r8.z <- _expr927.z
#line 326
331 0x00002350: mov r10.xyz, -r0.xyzx
332 0x00002368: add r10.xyz, r8.xyzx, r10.xyzx
333 0x00002384: dp3 r6.w, r10.xyzx, r6.xyzx
334 0x000023A0: max r6.w, r6.w, l(0.000000)
335 0x000023BC: min r6.w, r1.w, r6.w
336 0x000023D8: mul r10.xyz, r6.wwww, r6.xyzx
337 0x000023F4: add r10.xyz, r0.xyzx, r10.xyzx
338 0x00002410: mov r10.xyz, -r10.xyzx
339 0x00002428: add r8.xyz, r8.xyzx, r10.xyzx // r8.x <- _expr934.x; r8.y <- _expr934.y; r8.z <- _expr934.z
#line 327
340 0x00002444: mul r6.w, r8.x, r8.x
341 0x00002460: mul r8.x, r8.y, r8.y
342 0x0000247C: add r6.w, r6.w, r8.x
343 0x00002498: mul r8.x, r8.z, r8.z
344 0x000024B4: add r6.w, r6.w, r8.x
345 0x000024D0: add r6.w, r6.w, l(0.025000)
346 0x000024EC: div r6.w, l(1.000000), r6.w
347 0x00002508: mov r8.x, l(1.000000)
348 0x0000251C: mul r6.w, r6.w, r8.x
349 0x00002538: mul r8.xyz, r6.wwww, r9.xyzx
350 0x00002554: mul r8.xyz, r8.xyzx, l(0.002000, 0.002000, 0.002000, 0.000000)
351 0x0000257C: mul r8.xyz, r8.xyzx, l(0.200000, 0.200000, 0.200000, 0.000000)
352 0x000025A4: add r8.xyz, r4.xyzx, r8.xyzx // r8.x <- phi_4200_.x; r8.y <- phi_4200_.y; r8.z <- phi_4200_.z
#line 328
353 0x000025C0: mov r6.w, l(1)
354 0x000025D4: iadd r6.w, r4.w, r6.w // r6.w <- phi_4195_
#line 329
355 0x000025F0: else // Prior locations: r4.x <- phi_4200_.x; r4.y <- phi_4200_.y; r4.z <- phi_4200_.z; r4.w <- phi_4195_
356 0x000025F4: mov r8.xyz, r4.xyzx // r8.x <- phi_4200_.x; r8.y <- phi_4200_.y; r8.z <- phi_4200_.z
357 0x00002608: mov r6.w, r4.w // r6.w <- phi_4195_
358 0x0000261C: endif
#line 330
359 0x00002620: mov r8.w, l(0) // r8.w <- loop_init
#line 331
360 0x00002634: mov r8.xyz, r8.xyzx // r8.x <- _expr915.x; r8.y <- _expr915.y; r8.z <- _expr915.z
#line 332
361 0x00002648: mov r6.w, r6.w // r6.w <- _expr917
#line 333
362 0x0000265C: mov r3.w, r3.w // r3.w <- _expr920
#line 334
363 0x00002670: mov r9.xyz, r8.xyzx // r9.x <- local.x; r9.y <- local.y; r9.z <- local.z
#line 335
364 0x00002684: mov r9.xyz, r9.xyzx // r9.x <- local_1.x; r9.y <- local_1.y; r9.z <- local_1.z
#line 336
365 0x00002698: ult r9.w, r6.w, r3.w
366 0x000026B4: if_nz r9.w
367 0x000026C0: mov r4.w, r6.w // r4.w <- phi_4195_
368 0x000026D4: mov r9.w, r3.w // r9.w <- global.member_12.x; r9.w <- global.member_12.x
#line 337
369 0x000026E8: mov r4.xyz, r8.xyzx // r4.x <- phi_4200_.x; r4.y <- phi_4200_.y; r4.z <- phi_4200_.z
370 0x000026FC: mov r5.xyz, r9.xyzx // r5.x <- local.x; r5.y <- local.y; r5.z <- local.z
371 0x00002710: mov r7.xyz, r9.xyzx // r7.x <- local_1.x; r7.y <- local_1.y; r7.z <- local_1.z
372 0x00002724: mov r5.w, r8.w // r5.w <- loop_init
373 0x00002738: mov r3.w, r9.w // r3.w <- global.member_12.x; r3.w <- global.member_12.x
374 0x0000274C: continue
#line 338
375 0x00002750: else // Prior locations: r8.x <- phi_4200_.x; r8.y <- phi_4200_.y; r8.z <- phi_4200_.z; r6.w <- phi_4195_; r9.x <- local.x; r9.y <- local.y; r9.z <- local.z; r9.x <- local_1.x; r9.y <- local_1.y; r9.z <- local_1.z; r8.w <- loop_init
#line 339
376 0x00002754: mov r5.xyz, r9.xyzx // r5.x <- local.x; r5.y <- local.y; r5.z <- local.z
377 0x00002768: mov r7.xyz, r9.xyzx // r7.x <- local_1.x; r7.y <- local_1.y; r7.z <- local_1.z
378 0x0000277C: break
#line 340
379 0x00002780: endif
#line 341
380 0x00002784: mov r4.xyz, r8.xyzx // r4.x <- phi_4200_.x; r4.y <- phi_4200_.y; r4.z <- phi_4200_.z
381 0x00002798: mov r5.xyz, r9.xyzx // r5.x <- local.x; r5.y <- local.y; r5.z <- local.z
382 0x000027AC: mov r7.xyz, r9.xyzx // r7.x <- local_1.x; r7.y <- local_1.y; r7.z <- local_1.z
383 0x000027C0: mov r4.w, r6.w // r4.w <- phi_4195_
384 0x000027D4: mov r5.w, r8.w // r5.w <- loop_init
385 0x000027E8: endloop // r3.w <- global.member_12.x; r8.x <- phi_4200_.x; r8.y <- phi_4200_.y; r8.z <- phi_4200_.z; r6.w <- phi_4195_; r8.w <- loop_init; r8.x <- _expr915.x; r8.y <- _expr915.y; r8.z <- _expr915.z; r6.w <- _expr917; r3.w <- _expr920; r3.w <- global.member_12.x
#line 342
386 0x000027EC: mov r0.w, CB0[0][26].w // r0.w <- _expr955
#line 343
387 0x00002808: mov r1.x, CB0[0][19].x // r1.x <- _expr958
#line 344
388 0x00002824: lt r1.y, r1.x, r0.w // r1.y <- _expr959
#line 345
389 0x00002840: if_nz r1.y
#line 346
390 0x0000284C: mov r1.z, -r0.w
391 0x00002864: add r1.z, r1.z, l(300000.000000)
392 0x00002880: add r1.z, r1.x, r1.z // r1.z <- phi_4196_
#line 347
393 0x0000289C: else
#line 348
394 0x000028A0: mov r2.x, -r0.w
395 0x000028B8: add r1.z, r1.x, r2.x // r1.z <- phi_4196_
#line 349
396 0x000028D4: endif
#line 350
397 0x000028D8: mov r1.z, r1.z // r1.z <- _expr964
#line 351
398 0x000028EC: mov r7.xyz, r7.xyzx // r7.x <- _expr1041.x; r7.y <- _expr1041.y; r7.z <- _expr1041.z
#line 352
399 0x00002900: mov r7.xyz, r7.xyzx // r7.x <- phi_4201_.x; r7.y <- phi_4201_.y; r7.z <- phi_4201_.z
#line 353
400 0x00002914: lt r1.z, r1.z, l(5.000000)
401 0x00002930: if_nz r1.z
#line 354
402 0x0000293C: mov r2.xyz, CB0[0][26].xyzx // r2.x <- _expr967.x; r2.y <- _expr967.y; r2.z <- _expr967.z
#line 355
403 0x00002958: if_nz r1.y
#line 356
404 0x00002964: mov r1.y, -r0.w
405 0x0000297C: add r1.y, r1.y, l(300000.000000)
406 0x00002998: add r1.y, r1.x, r1.y // r1.y <- phi_4197_
#line 357
407 0x000029B4: else // Prior locations: r1.y <- _expr959
#line 358
408 0x000029B8: mov r0.w, -r0.w
409 0x000029D0: add r1.y, r0.w, r1.x // r1.y <- phi_4197_
#line 359
410 0x000029EC: endif
#line 360
411 0x000029F0: mov r1.y, r1.y // r1.y <- _expr973
#line 361
412 0x00002A04: mov r0.w, CB0[0][16].x // r0.w <- _expr979
#line 362
413 0x00002A20: mov r1.x, -r1.y
414 0x00002A38: add r1.x, r1.x, l(1.000000)
415 0x00002A54: max r1.x, r1.x, l(0.000000)
416 0x00002A70: mul r1.x, r1.x, l(1000000.000000)
417 0x00002A8C: mul r0.w, r0.w, l(0.400000)
418 0x00002AA8: sincos r0.w, null, r0.w
419 0x00002AC0: max r0.w, r0.w, l(0.000000)
420 0x00002ADC: mul r0.w, r0.w, r1.x
421 0x00002AF8: mul r1.xyz, r0.wwww, l(0.200000, 0.400000, 1.000000, 0.000000)
422 0x00002B20: mul r1.xyz, r1.xyzx, l(0.003000, 0.003000, 0.003000, 0.000000) // r1.x <- _expr985.x; r1.y <- _expr985.y; r1.z <- _expr985.z
#line 363
423 0x00002B48: add r2.xyz, r2.xyzx, l(0.000000, 0.000000, 25.000000, 0.000000) // r2.x <- _expr990.x; r2.y <- _expr990.y; r2.z <- _expr990.z
#line 364
424 0x00002B70: mov r3.xyz, -r0.xyzx
425 0x00002B88: add r3.xyz, r2.xyzx, r3.xyzx
426 0x00002BA4: dp3 r0.w, r3.xyzx, r6.xyzx
427 0x00002BC0: max r0.w, r0.w, l(0.000000)
428 0x00002BDC: min r0.w, r1.w, r0.w
429 0x00002BF8: mul r3.xyz, r0.wwww, r6.xyzx
430 0x00002C14: add r0.xyz, r0.xyzx, r3.xyzx
431 0x00002C30: mov r0.xyz, -r0.xyzx
432 0x00002C48: add r0.xyz, r0.xyzx, r2.xyzx // r0.x <- _expr997.x; r0.y <- _expr997.y; r0.z <- _expr997.z
#line 365
433 0x00002C64: mov r5.xyz, r5.xyzx // r5.x <- _expr1014.x; r5.y <- _expr1014.y; r5.z <- _expr1014.z
#line 366
434 0x00002C78: mul r0.x, r0.x, r0.x
435 0x00002C94: mul r0.y, r0.y, r0.y
436 0x00002CB0: add r0.x, r0.y, r0.x
437 0x00002CCC: mul r0.y, r0.z, r0.z
438 0x00002CE8: add r0.x, r0.y, r0.x
439 0x00002D04: add r0.x, r0.x, l(0.025000)
440 0x00002D20: div r0.x, l(1.000000), r0.x
441 0x00002D3C: mov r0.y, l(1.000000)
442 0x00002D50: mul r0.x, r0.x, r0.y
443 0x00002D6C: mul r0.xyz, r0.xxxx, r1.xyzx
444 0x00002D88: mul r0.xyz, r0.xyzx, l(0.002000, 0.002000, 0.002000, 0.000000)
445 0x00002DB0: mul r0.xyz, r0.xyzx, l(0.200000, 0.200000, 0.200000, 0.000000)
446 0x00002DD8: add r7.xyz, r0.xyzx, r5.xyzx
#line 367
447 0x00002DF4: endif
#line 368
448 0x00002DF8: mov r7.xyz, r7.xyzx // r7.x <- _expr1017.x; r7.y <- _expr1017.y; r7.z <- _expr1017.z
#line 369
449 0x00002E0C: mov r7.xyz, r7.xyzx // r7.x <- global_13.x; r7.y <- global_13.y; r7.z <- global_13.z; r7.x <- global_13.x; r7.y <- global_13.y; r7.z <- global_13.z
450 0x00002E20: mov r7.w, l(1.000000) // r7.w <- global_13.w; r7.w <- global_13.w
#line 378
451 0x00002E34: mov r7.xyzw, r7.xyzw // r7.x <- _expr3.x; r7.y <- _expr3.y; r7.z <- _expr3.z; r7.w <- _expr3.w
#line 379
452 0x00002E48: mov o0.xyzw, r7.xyzw
453 0x00002E5C: ret
// Approximately 454 instruction slots used
Oh! The issue seems to easily reproduce with an addition to the triangle example.
patch
diff --git a/examples/src/hello_triangle/mod.rs b/examples/src/hello_triangle/mod.rs
index 79162a695..a243d4d89 100644
--- a/examples/src/hello_triangle/mod.rs
+++ b/examples/src/hello_triangle/mod.rs
@@ -10,7 +10,10 @@ async fn run(event_loop: EventLoop<()>, window: Window) {
size.width = size.width.max(1);
size.height = size.height.max(1);
- let instance = wgpu::Instance::default();
+ let instance = wgpu::Instance::new(wgpu::InstanceDescriptor {
+ backends: wgpu::Backends::DX12,
+ ..Default::default()
+ });
let surface = instance.create_surface(&window).unwrap();
let adapter = instance
diff --git a/examples/src/hello_triangle/shader.wgsl b/examples/src/hello_triangle/shader.wgsl
index f84ccfe94..fa1c0e47f 100644
--- a/examples/src/hello_triangle/shader.wgsl
+++ b/examples/src/hello_triangle/shader.wgsl
@@ -7,5 +7,11 @@ fn vs_main(@builtin(vertex_index) in_vertex_index: u32) -> @builtin(position) ve
@fragment
fn fs_main() -> @location(0) vec4<f32> {
- return vec4<f32>(1.0, 0.0, 0.0, 1.0);
+ var x = 0.0;
+ switch 0 {
+ default {
+ x = 1.0;
+ }
+ }
+ return vec4<f32>(1.0, 0.0, x, 1.0);
}
cargo r --bin wgpu-examples hello_triangle
On DX12 FXC
On Vulkan
Seems like people have run into this with SPIRV-Cross as well https://github.com/KhronosGroup/SPIRV-Cross/issues/1402.
I think we should work around these cases in the backends.
oh no there are more switch related bugs in FXC https://bugs.chromium.org/p/tint/issues/detail?id=2126#c4
We absolutely should have execution tests for these cases that force fxc
var z = 0.0;
switch 0 {
case 0, 2, default {
z = 1.0;
}
}
multiple cases with a single body like this also aren't compiled correctly in FXC
So we can probably lower this to something like a do {} while(false) like the solution for the linked spirv-cross issue. This needs any contained continue statements to be forwarded to the outer loop somehow. Nevertheless, it looks like continues in switches already need special handling for FXC https://github.com/gfx-rs/wgpu/issues/4485
Sounds good to me, we can then take care of the issue with continue in #4485.
oh no there are more switch related bugs in FXC https://bugs.chromium.org/p/tint/issues/detail?id=2126#c4
Is it worth opening an issue to track this?
