Don't switch on a single constant in output consumed by FXC #4514
Labels
area: naga back-end
Outputs of naga shader conversion
external: driver-bug
A driver is causing the bug, though we may still want to work around it
FXC bug
lang: GLSL
OpenGL Shading Language
lang: HLSL
D3D Shading Language
naga
Shader Translator
type: bug
Something isn't working
Comments
|
Maybe this should be happening in spv-in instead. |
JCapucho
added
kind: feature
lang: SPIR-V
|
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 I'll allow myself to leave some keywords for people who might stumble upon this, too:
|
|
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:
|
|
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: 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 |
|
Oh! The issue seems to easily reproduce with an addition to the triangle example. patchdiff --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);
}On DX12 FXC
On Vulkan
|
|
Seems like people have run into this with SPIRV-Cross as well KhronosGroup/SPIRV-Cross#1402. I think we should work around these cases in the backends. |
teoxoy
added
type: bug
|
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 |
jimblandy
changed the title
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 |
|
Sounds good to me, we can then take care of the issue with |
6 tasks
Is it worth opening an issue to track this? |
|
@Imberflur: Yes please! |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
glslcemits code likeOpSwitch 0for control flow reasons, for example when one branch ends with aOpKill:this produces the rather dodgy glsl code:
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.
The text was updated successfully, but these errors were encountered: