summaryrefslogtreecommitdiff
path: root/src/shader_recompiler/frontend/ir
diff options
context:
space:
mode:
Diffstat (limited to 'src/shader_recompiler/frontend/ir')
-rw-r--r--src/shader_recompiler/frontend/ir/abstract_syntax_list.h58
-rw-r--r--src/shader_recompiler/frontend/ir/attribute.cpp454
-rw-r--r--src/shader_recompiler/frontend/ir/attribute.h250
-rw-r--r--src/shader_recompiler/frontend/ir/basic_block.cpp149
-rw-r--r--src/shader_recompiler/frontend/ir/basic_block.h185
-rw-r--r--src/shader_recompiler/frontend/ir/breadth_first_search.h56
-rw-r--r--src/shader_recompiler/frontend/ir/condition.cpp29
-rw-r--r--src/shader_recompiler/frontend/ir/condition.h60
-rw-r--r--src/shader_recompiler/frontend/ir/flow_test.cpp83
-rw-r--r--src/shader_recompiler/frontend/ir/flow_test.h62
-rw-r--r--src/shader_recompiler/frontend/ir/ir_emitter.cpp2017
-rw-r--r--src/shader_recompiler/frontend/ir/ir_emitter.h413
-rw-r--r--src/shader_recompiler/frontend/ir/microinstruction.cpp411
-rw-r--r--src/shader_recompiler/frontend/ir/modifiers.h49
-rw-r--r--src/shader_recompiler/frontend/ir/opcodes.cpp15
-rw-r--r--src/shader_recompiler/frontend/ir/opcodes.h110
-rw-r--r--src/shader_recompiler/frontend/ir/opcodes.inc550
-rw-r--r--src/shader_recompiler/frontend/ir/patch.cpp28
-rw-r--r--src/shader_recompiler/frontend/ir/patch.h149
-rw-r--r--src/shader_recompiler/frontend/ir/post_order.cpp46
-rw-r--r--src/shader_recompiler/frontend/ir/post_order.h14
-rw-r--r--src/shader_recompiler/frontend/ir/pred.h44
-rw-r--r--src/shader_recompiler/frontend/ir/program.cpp32
-rw-r--r--src/shader_recompiler/frontend/ir/program.h35
-rw-r--r--src/shader_recompiler/frontend/ir/reg.h332
-rw-r--r--src/shader_recompiler/frontend/ir/type.cpp38
-rw-r--r--src/shader_recompiler/frontend/ir/type.h61
-rw-r--r--src/shader_recompiler/frontend/ir/value.cpp99
-rw-r--r--src/shader_recompiler/frontend/ir/value.h398
29 files changed, 6227 insertions, 0 deletions
diff --git a/src/shader_recompiler/frontend/ir/abstract_syntax_list.h b/src/shader_recompiler/frontend/ir/abstract_syntax_list.h
new file mode 100644
index 000000000..b61773487
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/abstract_syntax_list.h
@@ -0,0 +1,58 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <vector>
8
9#include "shader_recompiler/frontend/ir/value.h"
10
11namespace Shader::IR {
12
13class Block;
14
15struct AbstractSyntaxNode {
16 enum class Type {
17 Block,
18 If,
19 EndIf,
20 Loop,
21 Repeat,
22 Break,
23 Return,
24 Unreachable,
25 };
26 union Data {
27 Block* block;
28 struct {
29 U1 cond;
30 Block* body;
31 Block* merge;
32 } if_node;
33 struct {
34 Block* merge;
35 } end_if;
36 struct {
37 Block* body;
38 Block* continue_block;
39 Block* merge;
40 } loop;
41 struct {
42 U1 cond;
43 Block* loop_header;
44 Block* merge;
45 } repeat;
46 struct {
47 U1 cond;
48 Block* merge;
49 Block* skip;
50 } break_node;
51 };
52
53 Data data{};
54 Type type{};
55};
56using AbstractSyntaxList = std::vector<AbstractSyntaxNode>;
57
58} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/attribute.cpp b/src/shader_recompiler/frontend/ir/attribute.cpp
new file mode 100644
index 000000000..4d0b8b8e5
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/attribute.cpp
@@ -0,0 +1,454 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <fmt/format.h>
6
7#include "shader_recompiler/exception.h"
8#include "shader_recompiler/frontend/ir/attribute.h"
9
10namespace Shader::IR {
11
12bool IsGeneric(Attribute attribute) noexcept {
13 return attribute >= Attribute::Generic0X && attribute <= Attribute::Generic31X;
14}
15
16u32 GenericAttributeIndex(Attribute attribute) {
17 if (!IsGeneric(attribute)) {
18 throw InvalidArgument("Attribute is not generic {}", attribute);
19 }
20 return (static_cast<u32>(attribute) - static_cast<u32>(Attribute::Generic0X)) / 4u;
21}
22
23u32 GenericAttributeElement(Attribute attribute) {
24 if (!IsGeneric(attribute)) {
25 throw InvalidArgument("Attribute is not generic {}", attribute);
26 }
27 return static_cast<u32>(attribute) % 4;
28}
29
30std::string NameOf(Attribute attribute) {
31 switch (attribute) {
32 case Attribute::PrimitiveId:
33 return "PrimitiveId";
34 case Attribute::Layer:
35 return "Layer";
36 case Attribute::ViewportIndex:
37 return "ViewportIndex";
38 case Attribute::PointSize:
39 return "PointSize";
40 case Attribute::PositionX:
41 return "Position.X";
42 case Attribute::PositionY:
43 return "Position.Y";
44 case Attribute::PositionZ:
45 return "Position.Z";
46 case Attribute::PositionW:
47 return "Position.W";
48 case Attribute::Generic0X:
49 return "Generic[0].X";
50 case Attribute::Generic0Y:
51 return "Generic[0].Y";
52 case Attribute::Generic0Z:
53 return "Generic[0].Z";
54 case Attribute::Generic0W:
55 return "Generic[0].W";
56 case Attribute::Generic1X:
57 return "Generic[1].X";
58 case Attribute::Generic1Y:
59 return "Generic[1].Y";
60 case Attribute::Generic1Z:
61 return "Generic[1].Z";
62 case Attribute::Generic1W:
63 return "Generic[1].W";
64 case Attribute::Generic2X:
65 return "Generic[2].X";
66 case Attribute::Generic2Y:
67 return "Generic[2].Y";
68 case Attribute::Generic2Z:
69 return "Generic[2].Z";
70 case Attribute::Generic2W:
71 return "Generic[2].W";
72 case Attribute::Generic3X:
73 return "Generic[3].X";
74 case Attribute::Generic3Y:
75 return "Generic[3].Y";
76 case Attribute::Generic3Z:
77 return "Generic[3].Z";
78 case Attribute::Generic3W:
79 return "Generic[3].W";
80 case Attribute::Generic4X:
81 return "Generic[4].X";
82 case Attribute::Generic4Y:
83 return "Generic[4].Y";
84 case Attribute::Generic4Z:
85 return "Generic[4].Z";
86 case Attribute::Generic4W:
87 return "Generic[4].W";
88 case Attribute::Generic5X:
89 return "Generic[5].X";
90 case Attribute::Generic5Y:
91 return "Generic[5].Y";
92 case Attribute::Generic5Z:
93 return "Generic[5].Z";
94 case Attribute::Generic5W:
95 return "Generic[5].W";
96 case Attribute::Generic6X:
97 return "Generic[6].X";
98 case Attribute::Generic6Y:
99 return "Generic[6].Y";
100 case Attribute::Generic6Z:
101 return "Generic[6].Z";
102 case Attribute::Generic6W:
103 return "Generic[6].W";
104 case Attribute::Generic7X:
105 return "Generic[7].X";
106 case Attribute::Generic7Y:
107 return "Generic[7].Y";
108 case Attribute::Generic7Z:
109 return "Generic[7].Z";
110 case Attribute::Generic7W:
111 return "Generic[7].W";
112 case Attribute::Generic8X:
113 return "Generic[8].X";
114 case Attribute::Generic8Y:
115 return "Generic[8].Y";
116 case Attribute::Generic8Z:
117 return "Generic[8].Z";
118 case Attribute::Generic8W:
119 return "Generic[8].W";
120 case Attribute::Generic9X:
121 return "Generic[9].X";
122 case Attribute::Generic9Y:
123 return "Generic[9].Y";
124 case Attribute::Generic9Z:
125 return "Generic[9].Z";
126 case Attribute::Generic9W:
127 return "Generic[9].W";
128 case Attribute::Generic10X:
129 return "Generic[10].X";
130 case Attribute::Generic10Y:
131 return "Generic[10].Y";
132 case Attribute::Generic10Z:
133 return "Generic[10].Z";
134 case Attribute::Generic10W:
135 return "Generic[10].W";
136 case Attribute::Generic11X:
137 return "Generic[11].X";
138 case Attribute::Generic11Y:
139 return "Generic[11].Y";
140 case Attribute::Generic11Z:
141 return "Generic[11].Z";
142 case Attribute::Generic11W:
143 return "Generic[11].W";
144 case Attribute::Generic12X:
145 return "Generic[12].X";
146 case Attribute::Generic12Y:
147 return "Generic[12].Y";
148 case Attribute::Generic12Z:
149 return "Generic[12].Z";
150 case Attribute::Generic12W:
151 return "Generic[12].W";
152 case Attribute::Generic13X:
153 return "Generic[13].X";
154 case Attribute::Generic13Y:
155 return "Generic[13].Y";
156 case Attribute::Generic13Z:
157 return "Generic[13].Z";
158 case Attribute::Generic13W:
159 return "Generic[13].W";
160 case Attribute::Generic14X:
161 return "Generic[14].X";
162 case Attribute::Generic14Y:
163 return "Generic[14].Y";
164 case Attribute::Generic14Z:
165 return "Generic[14].Z";
166 case Attribute::Generic14W:
167 return "Generic[14].W";
168 case Attribute::Generic15X:
169 return "Generic[15].X";
170 case Attribute::Generic15Y:
171 return "Generic[15].Y";
172 case Attribute::Generic15Z:
173 return "Generic[15].Z";
174 case Attribute::Generic15W:
175 return "Generic[15].W";
176 case Attribute::Generic16X:
177 return "Generic[16].X";
178 case Attribute::Generic16Y:
179 return "Generic[16].Y";
180 case Attribute::Generic16Z:
181 return "Generic[16].Z";
182 case Attribute::Generic16W:
183 return "Generic[16].W";
184 case Attribute::Generic17X:
185 return "Generic[17].X";
186 case Attribute::Generic17Y:
187 return "Generic[17].Y";
188 case Attribute::Generic17Z:
189 return "Generic[17].Z";
190 case Attribute::Generic17W:
191 return "Generic[17].W";
192 case Attribute::Generic18X:
193 return "Generic[18].X";
194 case Attribute::Generic18Y:
195 return "Generic[18].Y";
196 case Attribute::Generic18Z:
197 return "Generic[18].Z";
198 case Attribute::Generic18W:
199 return "Generic[18].W";
200 case Attribute::Generic19X:
201 return "Generic[19].X";
202 case Attribute::Generic19Y:
203 return "Generic[19].Y";
204 case Attribute::Generic19Z:
205 return "Generic[19].Z";
206 case Attribute::Generic19W:
207 return "Generic[19].W";
208 case Attribute::Generic20X:
209 return "Generic[20].X";
210 case Attribute::Generic20Y:
211 return "Generic[20].Y";
212 case Attribute::Generic20Z:
213 return "Generic[20].Z";
214 case Attribute::Generic20W:
215 return "Generic[20].W";
216 case Attribute::Generic21X:
217 return "Generic[21].X";
218 case Attribute::Generic21Y:
219 return "Generic[21].Y";
220 case Attribute::Generic21Z:
221 return "Generic[21].Z";
222 case Attribute::Generic21W:
223 return "Generic[21].W";
224 case Attribute::Generic22X:
225 return "Generic[22].X";
226 case Attribute::Generic22Y:
227 return "Generic[22].Y";
228 case Attribute::Generic22Z:
229 return "Generic[22].Z";
230 case Attribute::Generic22W:
231 return "Generic[22].W";
232 case Attribute::Generic23X:
233 return "Generic[23].X";
234 case Attribute::Generic23Y:
235 return "Generic[23].Y";
236 case Attribute::Generic23Z:
237 return "Generic[23].Z";
238 case Attribute::Generic23W:
239 return "Generic[23].W";
240 case Attribute::Generic24X:
241 return "Generic[24].X";
242 case Attribute::Generic24Y:
243 return "Generic[24].Y";
244 case Attribute::Generic24Z:
245 return "Generic[24].Z";
246 case Attribute::Generic24W:
247 return "Generic[24].W";
248 case Attribute::Generic25X:
249 return "Generic[25].X";
250 case Attribute::Generic25Y:
251 return "Generic[25].Y";
252 case Attribute::Generic25Z:
253 return "Generic[25].Z";
254 case Attribute::Generic25W:
255 return "Generic[25].W";
256 case Attribute::Generic26X:
257 return "Generic[26].X";
258 case Attribute::Generic26Y:
259 return "Generic[26].Y";
260 case Attribute::Generic26Z:
261 return "Generic[26].Z";
262 case Attribute::Generic26W:
263 return "Generic[26].W";
264 case Attribute::Generic27X:
265 return "Generic[27].X";
266 case Attribute::Generic27Y:
267 return "Generic[27].Y";
268 case Attribute::Generic27Z:
269 return "Generic[27].Z";
270 case Attribute::Generic27W:
271 return "Generic[27].W";
272 case Attribute::Generic28X:
273 return "Generic[28].X";
274 case Attribute::Generic28Y:
275 return "Generic[28].Y";
276 case Attribute::Generic28Z:
277 return "Generic[28].Z";
278 case Attribute::Generic28W:
279 return "Generic[28].W";
280 case Attribute::Generic29X:
281 return "Generic[29].X";
282 case Attribute::Generic29Y:
283 return "Generic[29].Y";
284 case Attribute::Generic29Z:
285 return "Generic[29].Z";
286 case Attribute::Generic29W:
287 return "Generic[29].W";
288 case Attribute::Generic30X:
289 return "Generic[30].X";
290 case Attribute::Generic30Y:
291 return "Generic[30].Y";
292 case Attribute::Generic30Z:
293 return "Generic[30].Z";
294 case Attribute::Generic30W:
295 return "Generic[30].W";
296 case Attribute::Generic31X:
297 return "Generic[31].X";
298 case Attribute::Generic31Y:
299 return "Generic[31].Y";
300 case Attribute::Generic31Z:
301 return "Generic[31].Z";
302 case Attribute::Generic31W:
303 return "Generic[31].W";
304 case Attribute::ColorFrontDiffuseR:
305 return "ColorFrontDiffuse.R";
306 case Attribute::ColorFrontDiffuseG:
307 return "ColorFrontDiffuse.G";
308 case Attribute::ColorFrontDiffuseB:
309 return "ColorFrontDiffuse.B";
310 case Attribute::ColorFrontDiffuseA:
311 return "ColorFrontDiffuse.A";
312 case Attribute::ColorFrontSpecularR:
313 return "ColorFrontSpecular.R";
314 case Attribute::ColorFrontSpecularG:
315 return "ColorFrontSpecular.G";
316 case Attribute::ColorFrontSpecularB:
317 return "ColorFrontSpecular.B";
318 case Attribute::ColorFrontSpecularA:
319 return "ColorFrontSpecular.A";
320 case Attribute::ColorBackDiffuseR:
321 return "ColorBackDiffuse.R";
322 case Attribute::ColorBackDiffuseG:
323 return "ColorBackDiffuse.G";
324 case Attribute::ColorBackDiffuseB:
325 return "ColorBackDiffuse.B";
326 case Attribute::ColorBackDiffuseA:
327 return "ColorBackDiffuse.A";
328 case Attribute::ColorBackSpecularR:
329 return "ColorBackSpecular.R";
330 case Attribute::ColorBackSpecularG:
331 return "ColorBackSpecular.G";
332 case Attribute::ColorBackSpecularB:
333 return "ColorBackSpecular.B";
334 case Attribute::ColorBackSpecularA:
335 return "ColorBackSpecular.A";
336 case Attribute::ClipDistance0:
337 return "ClipDistance[0]";
338 case Attribute::ClipDistance1:
339 return "ClipDistance[1]";
340 case Attribute::ClipDistance2:
341 return "ClipDistance[2]";
342 case Attribute::ClipDistance3:
343 return "ClipDistance[3]";
344 case Attribute::ClipDistance4:
345 return "ClipDistance[4]";
346 case Attribute::ClipDistance5:
347 return "ClipDistance[5]";
348 case Attribute::ClipDistance6:
349 return "ClipDistance[6]";
350 case Attribute::ClipDistance7:
351 return "ClipDistance[7]";
352 case Attribute::PointSpriteS:
353 return "PointSprite.S";
354 case Attribute::PointSpriteT:
355 return "PointSprite.T";
356 case Attribute::FogCoordinate:
357 return "FogCoordinate";
358 case Attribute::TessellationEvaluationPointU:
359 return "TessellationEvaluationPoint.U";
360 case Attribute::TessellationEvaluationPointV:
361 return "TessellationEvaluationPoint.V";
362 case Attribute::InstanceId:
363 return "InstanceId";
364 case Attribute::VertexId:
365 return "VertexId";
366 case Attribute::FixedFncTexture0S:
367 return "FixedFncTexture[0].S";
368 case Attribute::FixedFncTexture0T:
369 return "FixedFncTexture[0].T";
370 case Attribute::FixedFncTexture0R:
371 return "FixedFncTexture[0].R";
372 case Attribute::FixedFncTexture0Q:
373 return "FixedFncTexture[0].Q";
374 case Attribute::FixedFncTexture1S:
375 return "FixedFncTexture[1].S";
376 case Attribute::FixedFncTexture1T:
377 return "FixedFncTexture[1].T";
378 case Attribute::FixedFncTexture1R:
379 return "FixedFncTexture[1].R";
380 case Attribute::FixedFncTexture1Q:
381 return "FixedFncTexture[1].Q";
382 case Attribute::FixedFncTexture2S:
383 return "FixedFncTexture[2].S";
384 case Attribute::FixedFncTexture2T:
385 return "FixedFncTexture[2].T";
386 case Attribute::FixedFncTexture2R:
387 return "FixedFncTexture[2].R";
388 case Attribute::FixedFncTexture2Q:
389 return "FixedFncTexture[2].Q";
390 case Attribute::FixedFncTexture3S:
391 return "FixedFncTexture[3].S";
392 case Attribute::FixedFncTexture3T:
393 return "FixedFncTexture[3].T";
394 case Attribute::FixedFncTexture3R:
395 return "FixedFncTexture[3].R";
396 case Attribute::FixedFncTexture3Q:
397 return "FixedFncTexture[3].Q";
398 case Attribute::FixedFncTexture4S:
399 return "FixedFncTexture[4].S";
400 case Attribute::FixedFncTexture4T:
401 return "FixedFncTexture[4].T";
402 case Attribute::FixedFncTexture4R:
403 return "FixedFncTexture[4].R";
404 case Attribute::FixedFncTexture4Q:
405 return "FixedFncTexture[4].Q";
406 case Attribute::FixedFncTexture5S:
407 return "FixedFncTexture[5].S";
408 case Attribute::FixedFncTexture5T:
409 return "FixedFncTexture[5].T";
410 case Attribute::FixedFncTexture5R:
411 return "FixedFncTexture[5].R";
412 case Attribute::FixedFncTexture5Q:
413 return "FixedFncTexture[5].Q";
414 case Attribute::FixedFncTexture6S:
415 return "FixedFncTexture[6].S";
416 case Attribute::FixedFncTexture6T:
417 return "FixedFncTexture[6].T";
418 case Attribute::FixedFncTexture6R:
419 return "FixedFncTexture[6].R";
420 case Attribute::FixedFncTexture6Q:
421 return "FixedFncTexture[6].Q";
422 case Attribute::FixedFncTexture7S:
423 return "FixedFncTexture[7].S";
424 case Attribute::FixedFncTexture7T:
425 return "FixedFncTexture[7].T";
426 case Attribute::FixedFncTexture7R:
427 return "FixedFncTexture[7].R";
428 case Attribute::FixedFncTexture7Q:
429 return "FixedFncTexture[7].Q";
430 case Attribute::FixedFncTexture8S:
431 return "FixedFncTexture[8].S";
432 case Attribute::FixedFncTexture8T:
433 return "FixedFncTexture[8].T";
434 case Attribute::FixedFncTexture8R:
435 return "FixedFncTexture[8].R";
436 case Attribute::FixedFncTexture8Q:
437 return "FixedFncTexture[8].Q";
438 case Attribute::FixedFncTexture9S:
439 return "FixedFncTexture[9].S";
440 case Attribute::FixedFncTexture9T:
441 return "FixedFncTexture[9].T";
442 case Attribute::FixedFncTexture9R:
443 return "FixedFncTexture[9].R";
444 case Attribute::FixedFncTexture9Q:
445 return "FixedFncTexture[9].Q";
446 case Attribute::ViewportMask:
447 return "ViewportMask";
448 case Attribute::FrontFace:
449 return "FrontFace";
450 }
451 return fmt::format("<reserved attribute {}>", static_cast<int>(attribute));
452}
453
454} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/attribute.h b/src/shader_recompiler/frontend/ir/attribute.h
new file mode 100644
index 000000000..ca1199494
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/attribute.h
@@ -0,0 +1,250 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <fmt/format.h>
8
9#include "common/common_types.h"
10
11namespace Shader::IR {
12
13enum class Attribute : u64 {
14 PrimitiveId = 24,
15 Layer = 25,
16 ViewportIndex = 26,
17 PointSize = 27,
18 PositionX = 28,
19 PositionY = 29,
20 PositionZ = 30,
21 PositionW = 31,
22 Generic0X = 32,
23 Generic0Y = 33,
24 Generic0Z = 34,
25 Generic0W = 35,
26 Generic1X = 36,
27 Generic1Y = 37,
28 Generic1Z = 38,
29 Generic1W = 39,
30 Generic2X = 40,
31 Generic2Y = 41,
32 Generic2Z = 42,
33 Generic2W = 43,
34 Generic3X = 44,
35 Generic3Y = 45,
36 Generic3Z = 46,
37 Generic3W = 47,
38 Generic4X = 48,
39 Generic4Y = 49,
40 Generic4Z = 50,
41 Generic4W = 51,
42 Generic5X = 52,
43 Generic5Y = 53,
44 Generic5Z = 54,
45 Generic5W = 55,
46 Generic6X = 56,
47 Generic6Y = 57,
48 Generic6Z = 58,
49 Generic6W = 59,
50 Generic7X = 60,
51 Generic7Y = 61,
52 Generic7Z = 62,
53 Generic7W = 63,
54 Generic8X = 64,
55 Generic8Y = 65,
56 Generic8Z = 66,
57 Generic8W = 67,
58 Generic9X = 68,
59 Generic9Y = 69,
60 Generic9Z = 70,
61 Generic9W = 71,
62 Generic10X = 72,
63 Generic10Y = 73,
64 Generic10Z = 74,
65 Generic10W = 75,
66 Generic11X = 76,
67 Generic11Y = 77,
68 Generic11Z = 78,
69 Generic11W = 79,
70 Generic12X = 80,
71 Generic12Y = 81,
72 Generic12Z = 82,
73 Generic12W = 83,
74 Generic13X = 84,
75 Generic13Y = 85,
76 Generic13Z = 86,
77 Generic13W = 87,
78 Generic14X = 88,
79 Generic14Y = 89,
80 Generic14Z = 90,
81 Generic14W = 91,
82 Generic15X = 92,
83 Generic15Y = 93,
84 Generic15Z = 94,
85 Generic15W = 95,
86 Generic16X = 96,
87 Generic16Y = 97,
88 Generic16Z = 98,
89 Generic16W = 99,
90 Generic17X = 100,
91 Generic17Y = 101,
92 Generic17Z = 102,
93 Generic17W = 103,
94 Generic18X = 104,
95 Generic18Y = 105,
96 Generic18Z = 106,
97 Generic18W = 107,
98 Generic19X = 108,
99 Generic19Y = 109,
100 Generic19Z = 110,
101 Generic19W = 111,
102 Generic20X = 112,
103 Generic20Y = 113,
104 Generic20Z = 114,
105 Generic20W = 115,
106 Generic21X = 116,
107 Generic21Y = 117,
108 Generic21Z = 118,
109 Generic21W = 119,
110 Generic22X = 120,
111 Generic22Y = 121,
112 Generic22Z = 122,
113 Generic22W = 123,
114 Generic23X = 124,
115 Generic23Y = 125,
116 Generic23Z = 126,
117 Generic23W = 127,
118 Generic24X = 128,
119 Generic24Y = 129,
120 Generic24Z = 130,
121 Generic24W = 131,
122 Generic25X = 132,
123 Generic25Y = 133,
124 Generic25Z = 134,
125 Generic25W = 135,
126 Generic26X = 136,
127 Generic26Y = 137,
128 Generic26Z = 138,
129 Generic26W = 139,
130 Generic27X = 140,
131 Generic27Y = 141,
132 Generic27Z = 142,
133 Generic27W = 143,
134 Generic28X = 144,
135 Generic28Y = 145,
136 Generic28Z = 146,
137 Generic28W = 147,
138 Generic29X = 148,
139 Generic29Y = 149,
140 Generic29Z = 150,
141 Generic29W = 151,
142 Generic30X = 152,
143 Generic30Y = 153,
144 Generic30Z = 154,
145 Generic30W = 155,
146 Generic31X = 156,
147 Generic31Y = 157,
148 Generic31Z = 158,
149 Generic31W = 159,
150 ColorFrontDiffuseR = 160,
151 ColorFrontDiffuseG = 161,
152 ColorFrontDiffuseB = 162,
153 ColorFrontDiffuseA = 163,
154 ColorFrontSpecularR = 164,
155 ColorFrontSpecularG = 165,
156 ColorFrontSpecularB = 166,
157 ColorFrontSpecularA = 167,
158 ColorBackDiffuseR = 168,
159 ColorBackDiffuseG = 169,
160 ColorBackDiffuseB = 170,
161 ColorBackDiffuseA = 171,
162 ColorBackSpecularR = 172,
163 ColorBackSpecularG = 173,
164 ColorBackSpecularB = 174,
165 ColorBackSpecularA = 175,
166 ClipDistance0 = 176,
167 ClipDistance1 = 177,
168 ClipDistance2 = 178,
169 ClipDistance3 = 179,
170 ClipDistance4 = 180,
171 ClipDistance5 = 181,
172 ClipDistance6 = 182,
173 ClipDistance7 = 183,
174 PointSpriteS = 184,
175 PointSpriteT = 185,
176 FogCoordinate = 186,
177 TessellationEvaluationPointU = 188,
178 TessellationEvaluationPointV = 189,
179 InstanceId = 190,
180 VertexId = 191,
181 FixedFncTexture0S = 192,
182 FixedFncTexture0T = 193,
183 FixedFncTexture0R = 194,
184 FixedFncTexture0Q = 195,
185 FixedFncTexture1S = 196,
186 FixedFncTexture1T = 197,
187 FixedFncTexture1R = 198,
188 FixedFncTexture1Q = 199,
189 FixedFncTexture2S = 200,
190 FixedFncTexture2T = 201,
191 FixedFncTexture2R = 202,
192 FixedFncTexture2Q = 203,
193 FixedFncTexture3S = 204,
194 FixedFncTexture3T = 205,
195 FixedFncTexture3R = 206,
196 FixedFncTexture3Q = 207,
197 FixedFncTexture4S = 208,
198 FixedFncTexture4T = 209,
199 FixedFncTexture4R = 210,
200 FixedFncTexture4Q = 211,
201 FixedFncTexture5S = 212,
202 FixedFncTexture5T = 213,
203 FixedFncTexture5R = 214,
204 FixedFncTexture5Q = 215,
205 FixedFncTexture6S = 216,
206 FixedFncTexture6T = 217,
207 FixedFncTexture6R = 218,
208 FixedFncTexture6Q = 219,
209 FixedFncTexture7S = 220,
210 FixedFncTexture7T = 221,
211 FixedFncTexture7R = 222,
212 FixedFncTexture7Q = 223,
213 FixedFncTexture8S = 224,
214 FixedFncTexture8T = 225,
215 FixedFncTexture8R = 226,
216 FixedFncTexture8Q = 227,
217 FixedFncTexture9S = 228,
218 FixedFncTexture9T = 229,
219 FixedFncTexture9R = 230,
220 FixedFncTexture9Q = 231,
221 ViewportMask = 232,
222 FrontFace = 255,
223};
224
225constexpr size_t NUM_GENERICS = 32;
226
227[[nodiscard]] bool IsGeneric(Attribute attribute) noexcept;
228
229[[nodiscard]] u32 GenericAttributeIndex(Attribute attribute);
230
231[[nodiscard]] u32 GenericAttributeElement(Attribute attribute);
232
233[[nodiscard]] std::string NameOf(Attribute attribute);
234
235[[nodiscard]] constexpr IR::Attribute operator+(IR::Attribute attribute, size_t value) noexcept {
236 return static_cast<IR::Attribute>(static_cast<size_t>(attribute) + value);
237}
238
239} // namespace Shader::IR
240
241template <>
242struct fmt::formatter<Shader::IR::Attribute> {
243 constexpr auto parse(format_parse_context& ctx) {
244 return ctx.begin();
245 }
246 template <typename FormatContext>
247 auto format(const Shader::IR::Attribute& attribute, FormatContext& ctx) {
248 return fmt::format_to(ctx.out(), "{}", Shader::IR::NameOf(attribute));
249 }
250};
diff --git a/src/shader_recompiler/frontend/ir/basic_block.cpp b/src/shader_recompiler/frontend/ir/basic_block.cpp
new file mode 100644
index 000000000..7c08b25ce
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/basic_block.cpp
@@ -0,0 +1,149 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <algorithm>
6#include <initializer_list>
7#include <map>
8#include <memory>
9
10#include "common/bit_cast.h"
11#include "common/common_types.h"
12#include "shader_recompiler/frontend/ir/basic_block.h"
13#include "shader_recompiler/frontend/ir/value.h"
14
15namespace Shader::IR {
16
17Block::Block(ObjectPool<Inst>& inst_pool_) : inst_pool{&inst_pool_} {}
18
19Block::~Block() = default;
20
21void Block::AppendNewInst(Opcode op, std::initializer_list<Value> args) {
22 PrependNewInst(end(), op, args);
23}
24
25Block::iterator Block::PrependNewInst(iterator insertion_point, Opcode op,
26 std::initializer_list<Value> args, u32 flags) {
27 Inst* const inst{inst_pool->Create(op, flags)};
28 const auto result_it{instructions.insert(insertion_point, *inst)};
29
30 if (inst->NumArgs() != args.size()) {
31 throw InvalidArgument("Invalid number of arguments {} in {}", args.size(), op);
32 }
33 std::ranges::for_each(args, [inst, index = size_t{0}](const Value& arg) mutable {
34 inst->SetArg(index, arg);
35 ++index;
36 });
37 return result_it;
38}
39
40void Block::AddBranch(Block* block) {
41 if (std::ranges::find(imm_successors, block) != imm_successors.end()) {
42 throw LogicError("Successor already inserted");
43 }
44 if (std::ranges::find(block->imm_predecessors, this) != block->imm_predecessors.end()) {
45 throw LogicError("Predecessor already inserted");
46 }
47 imm_successors.push_back(block);
48 block->imm_predecessors.push_back(this);
49}
50
51static std::string BlockToIndex(const std::map<const Block*, size_t>& block_to_index,
52 Block* block) {
53 if (const auto it{block_to_index.find(block)}; it != block_to_index.end()) {
54 return fmt::format("{{Block ${}}}", it->second);
55 }
56 return fmt::format("$<unknown block {:016x}>", reinterpret_cast<u64>(block));
57}
58
59static size_t InstIndex(std::map<const Inst*, size_t>& inst_to_index, size_t& inst_index,
60 const Inst* inst) {
61 const auto [it, is_inserted]{inst_to_index.emplace(inst, inst_index + 1)};
62 if (is_inserted) {
63 ++inst_index;
64 }
65 return it->second;
66}
67
68static std::string ArgToIndex(std::map<const Inst*, size_t>& inst_to_index, size_t& inst_index,
69 const Value& arg) {
70 if (arg.IsEmpty()) {
71 return "<null>";
72 }
73 if (!arg.IsImmediate() || arg.IsIdentity()) {
74 return fmt::format("%{}", InstIndex(inst_to_index, inst_index, arg.Inst()));
75 }
76 switch (arg.Type()) {
77 case Type::U1:
78 return fmt::format("#{}", arg.U1() ? "true" : "false");
79 case Type::U8:
80 return fmt::format("#{}", arg.U8());
81 case Type::U16:
82 return fmt::format("#{}", arg.U16());
83 case Type::U32:
84 return fmt::format("#{}", arg.U32());
85 case Type::U64:
86 return fmt::format("#{}", arg.U64());
87 case Type::F32:
88 return fmt::format("#{}", arg.F32());
89 case Type::Reg:
90 return fmt::format("{}", arg.Reg());
91 case Type::Pred:
92 return fmt::format("{}", arg.Pred());
93 case Type::Attribute:
94 return fmt::format("{}", arg.Attribute());
95 default:
96 return "<unknown immediate type>";
97 }
98}
99
100std::string DumpBlock(const Block& block) {
101 size_t inst_index{0};
102 std::map<const Inst*, size_t> inst_to_index;
103 return DumpBlock(block, {}, inst_to_index, inst_index);
104}
105
106std::string DumpBlock(const Block& block, const std::map<const Block*, size_t>& block_to_index,
107 std::map<const Inst*, size_t>& inst_to_index, size_t& inst_index) {
108 std::string ret{"Block"};
109 if (const auto it{block_to_index.find(&block)}; it != block_to_index.end()) {
110 ret += fmt::format(" ${}", it->second);
111 }
112 ret += '\n';
113 for (const Inst& inst : block) {
114 const Opcode op{inst.GetOpcode()};
115 ret += fmt::format("[{:016x}] ", reinterpret_cast<u64>(&inst));
116 if (TypeOf(op) != Type::Void) {
117 ret += fmt::format("%{:<5} = {}", InstIndex(inst_to_index, inst_index, &inst), op);
118 } else {
119 ret += fmt::format(" {}", op); // '%00000 = ' -> 1 + 5 + 3 = 9 spaces
120 }
121 const size_t arg_count{inst.NumArgs()};
122 for (size_t arg_index = 0; arg_index < arg_count; ++arg_index) {
123 const Value arg{inst.Arg(arg_index)};
124 const std::string arg_str{ArgToIndex(inst_to_index, inst_index, arg)};
125 ret += arg_index != 0 ? ", " : " ";
126 if (op == Opcode::Phi) {
127 ret += fmt::format("[ {}, {} ]", arg_str,
128 BlockToIndex(block_to_index, inst.PhiBlock(arg_index)));
129 } else {
130 ret += arg_str;
131 }
132 if (op != Opcode::Phi) {
133 const Type actual_type{arg.Type()};
134 const Type expected_type{ArgTypeOf(op, arg_index)};
135 if (!AreTypesCompatible(actual_type, expected_type)) {
136 ret += fmt::format("<type error: {} != {}>", actual_type, expected_type);
137 }
138 }
139 }
140 if (TypeOf(op) != Type::Void) {
141 ret += fmt::format(" (uses: {})\n", inst.UseCount());
142 } else {
143 ret += '\n';
144 }
145 }
146 return ret;
147}
148
149} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/basic_block.h b/src/shader_recompiler/frontend/ir/basic_block.h
new file mode 100644
index 000000000..7e134b4c7
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/basic_block.h
@@ -0,0 +1,185 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <initializer_list>
8#include <map>
9#include <span>
10#include <vector>
11
12#include <boost/intrusive/list.hpp>
13
14#include "common/bit_cast.h"
15#include "common/common_types.h"
16#include "shader_recompiler/frontend/ir/condition.h"
17#include "shader_recompiler/frontend/ir/value.h"
18#include "shader_recompiler/object_pool.h"
19
20namespace Shader::IR {
21
22class Block {
23public:
24 using InstructionList = boost::intrusive::list<Inst>;
25 using size_type = InstructionList::size_type;
26 using iterator = InstructionList::iterator;
27 using const_iterator = InstructionList::const_iterator;
28 using reverse_iterator = InstructionList::reverse_iterator;
29 using const_reverse_iterator = InstructionList::const_reverse_iterator;
30
31 explicit Block(ObjectPool<Inst>& inst_pool_);
32 ~Block();
33
34 Block(const Block&) = delete;
35 Block& operator=(const Block&) = delete;
36
37 Block(Block&&) = default;
38 Block& operator=(Block&&) = default;
39
40 /// Appends a new instruction to the end of this basic block.
41 void AppendNewInst(Opcode op, std::initializer_list<Value> args);
42
43 /// Prepends a new instruction to this basic block before the insertion point.
44 iterator PrependNewInst(iterator insertion_point, Opcode op,
45 std::initializer_list<Value> args = {}, u32 flags = 0);
46
47 /// Adds a new branch to this basic block.
48 void AddBranch(Block* block);
49
50 /// Gets a mutable reference to the instruction list for this basic block.
51 [[nodiscard]] InstructionList& Instructions() noexcept {
52 return instructions;
53 }
54 /// Gets an immutable reference to the instruction list for this basic block.
55 [[nodiscard]] const InstructionList& Instructions() const noexcept {
56 return instructions;
57 }
58
59 /// Gets an immutable span to the immediate predecessors.
60 [[nodiscard]] std::span<Block* const> ImmPredecessors() const noexcept {
61 return imm_predecessors;
62 }
63 /// Gets an immutable span to the immediate successors.
64 [[nodiscard]] std::span<Block* const> ImmSuccessors() const noexcept {
65 return imm_successors;
66 }
67
68 /// Intrusively store the host definition of this instruction.
69 template <typename DefinitionType>
70 void SetDefinition(DefinitionType def) {
71 definition = Common::BitCast<u32>(def);
72 }
73
74 /// Return the intrusively stored host definition of this instruction.
75 template <typename DefinitionType>
76 [[nodiscard]] DefinitionType Definition() const noexcept {
77 return Common::BitCast<DefinitionType>(definition);
78 }
79
80 void SetSsaRegValue(IR::Reg reg, const Value& value) noexcept {
81 ssa_reg_values[RegIndex(reg)] = value;
82 }
83 const Value& SsaRegValue(IR::Reg reg) const noexcept {
84 return ssa_reg_values[RegIndex(reg)];
85 }
86
87 void SsaSeal() noexcept {
88 is_ssa_sealed = true;
89 }
90 [[nodiscard]] bool IsSsaSealed() const noexcept {
91 return is_ssa_sealed;
92 }
93
94 [[nodiscard]] bool empty() const {
95 return instructions.empty();
96 }
97 [[nodiscard]] size_type size() const {
98 return instructions.size();
99 }
100
101 [[nodiscard]] Inst& front() {
102 return instructions.front();
103 }
104 [[nodiscard]] const Inst& front() const {
105 return instructions.front();
106 }
107
108 [[nodiscard]] Inst& back() {
109 return instructions.back();
110 }
111 [[nodiscard]] const Inst& back() const {
112 return instructions.back();
113 }
114
115 [[nodiscard]] iterator begin() {
116 return instructions.begin();
117 }
118 [[nodiscard]] const_iterator begin() const {
119 return instructions.begin();
120 }
121 [[nodiscard]] iterator end() {
122 return instructions.end();
123 }
124 [[nodiscard]] const_iterator end() const {
125 return instructions.end();
126 }
127
128 [[nodiscard]] reverse_iterator rbegin() {
129 return instructions.rbegin();
130 }
131 [[nodiscard]] const_reverse_iterator rbegin() const {
132 return instructions.rbegin();
133 }
134 [[nodiscard]] reverse_iterator rend() {
135 return instructions.rend();
136 }
137 [[nodiscard]] const_reverse_iterator rend() const {
138 return instructions.rend();
139 }
140
141 [[nodiscard]] const_iterator cbegin() const {
142 return instructions.cbegin();
143 }
144 [[nodiscard]] const_iterator cend() const {
145 return instructions.cend();
146 }
147
148 [[nodiscard]] const_reverse_iterator crbegin() const {
149 return instructions.crbegin();
150 }
151 [[nodiscard]] const_reverse_iterator crend() const {
152 return instructions.crend();
153 }
154
155private:
156 /// Memory pool for instruction list
157 ObjectPool<Inst>* inst_pool;
158
159 /// List of instructions in this block
160 InstructionList instructions;
161
162 /// Block immediate predecessors
163 std::vector<Block*> imm_predecessors;
164 /// Block immediate successors
165 std::vector<Block*> imm_successors;
166
167 /// Intrusively store the value of a register in the block.
168 std::array<Value, NUM_REGS> ssa_reg_values;
169 /// Intrusively store if the block is sealed in the SSA pass.
170 bool is_ssa_sealed{false};
171
172 /// Intrusively stored host definition of this block.
173 u32 definition{};
174};
175
176using BlockList = std::vector<Block*>;
177
178[[nodiscard]] std::string DumpBlock(const Block& block);
179
180[[nodiscard]] std::string DumpBlock(const Block& block,
181 const std::map<const Block*, size_t>& block_to_index,
182 std::map<const Inst*, size_t>& inst_to_index,
183 size_t& inst_index);
184
185} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/breadth_first_search.h b/src/shader_recompiler/frontend/ir/breadth_first_search.h
new file mode 100644
index 000000000..a52ccbd58
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/breadth_first_search.h
@@ -0,0 +1,56 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <optional>
8#include <type_traits>
9#include <queue>
10
11#include <boost/container/small_vector.hpp>
12
13#include "shader_recompiler/frontend/ir/value.h"
14
15namespace Shader::IR {
16
17template <typename Pred>
18auto BreadthFirstSearch(const Value& value, Pred&& pred)
19 -> std::invoke_result_t<Pred, const Inst*> {
20 if (value.IsImmediate()) {
21 // Nothing to do with immediates
22 return std::nullopt;
23 }
24 // Breadth-first search visiting the right most arguments first
25 // Small vector has been determined from shaders in Super Smash Bros. Ultimate
26 boost::container::small_vector<const Inst*, 2> visited;
27 std::queue<const Inst*> queue;
28 queue.push(value.InstRecursive());
29
30 while (!queue.empty()) {
31 // Pop one instruction from the queue
32 const Inst* const inst{queue.front()};
33 queue.pop();
34 if (const std::optional result = pred(inst)) {
35 // This is the instruction we were looking for
36 return result;
37 }
38 // Visit the right most arguments first
39 for (size_t arg = inst->NumArgs(); arg--;) {
40 const Value arg_value{inst->Arg(arg)};
41 if (arg_value.IsImmediate()) {
42 continue;
43 }
44 // Queue instruction if it hasn't been visited
45 const Inst* const arg_inst{arg_value.InstRecursive()};
46 if (std::ranges::find(visited, arg_inst) == visited.end()) {
47 visited.push_back(arg_inst);
48 queue.push(arg_inst);
49 }
50 }
51 }
52 // SSA tree has been traversed and the result hasn't been found
53 return std::nullopt;
54}
55
56} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/condition.cpp b/src/shader_recompiler/frontend/ir/condition.cpp
new file mode 100644
index 000000000..fc18ea2a2
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/condition.cpp
@@ -0,0 +1,29 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <string>
6
7#include <fmt/format.h>
8
9#include "shader_recompiler/frontend/ir/condition.h"
10
11namespace Shader::IR {
12
13std::string NameOf(Condition condition) {
14 std::string ret;
15 if (condition.GetFlowTest() != FlowTest::T) {
16 ret = fmt::to_string(condition.GetFlowTest());
17 }
18 const auto [pred, negated]{condition.GetPred()};
19 if (!ret.empty()) {
20 ret += '&';
21 }
22 if (negated) {
23 ret += '!';
24 }
25 ret += fmt::to_string(pred);
26 return ret;
27}
28
29} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/condition.h b/src/shader_recompiler/frontend/ir/condition.h
new file mode 100644
index 000000000..aa8597c60
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/condition.h
@@ -0,0 +1,60 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <compare>
8#include <string>
9
10#include <fmt/format.h>
11
12#include "common/common_types.h"
13#include "shader_recompiler/frontend/ir/flow_test.h"
14#include "shader_recompiler/frontend/ir/pred.h"
15
16namespace Shader::IR {
17
18class Condition {
19public:
20 Condition() noexcept = default;
21
22 explicit Condition(FlowTest flow_test_, Pred pred_, bool pred_negated_ = false) noexcept
23 : flow_test{static_cast<u16>(flow_test_)}, pred{static_cast<u8>(pred_)},
24 pred_negated{pred_negated_ ? u8{1} : u8{0}} {}
25
26 explicit Condition(Pred pred_, bool pred_negated_ = false) noexcept
27 : Condition(FlowTest::T, pred_, pred_negated_) {}
28
29 explicit Condition(bool value) : Condition(Pred::PT, !value) {}
30
31 auto operator<=>(const Condition&) const noexcept = default;
32
33 [[nodiscard]] IR::FlowTest GetFlowTest() const noexcept {
34 return static_cast<IR::FlowTest>(flow_test);
35 }
36
37 [[nodiscard]] std::pair<IR::Pred, bool> GetPred() const noexcept {
38 return {static_cast<IR::Pred>(pred), pred_negated != 0};
39 }
40
41private:
42 u16 flow_test;
43 u8 pred;
44 u8 pred_negated;
45};
46
47std::string NameOf(Condition condition);
48
49} // namespace Shader::IR
50
51template <>
52struct fmt::formatter<Shader::IR::Condition> {
53 constexpr auto parse(format_parse_context& ctx) {
54 return ctx.begin();
55 }
56 template <typename FormatContext>
57 auto format(const Shader::IR::Condition& cond, FormatContext& ctx) {
58 return fmt::format_to(ctx.out(), "{}", Shader::IR::NameOf(cond));
59 }
60};
diff --git a/src/shader_recompiler/frontend/ir/flow_test.cpp b/src/shader_recompiler/frontend/ir/flow_test.cpp
new file mode 100644
index 000000000..6ebb4ad89
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/flow_test.cpp
@@ -0,0 +1,83 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <string>
6
7#include <fmt/format.h>
8
9#include "shader_recompiler/frontend/ir/flow_test.h"
10
11namespace Shader::IR {
12
13std::string NameOf(FlowTest flow_test) {
14 switch (flow_test) {
15 case FlowTest::F:
16 return "F";
17 case FlowTest::LT:
18 return "LT";
19 case FlowTest::EQ:
20 return "EQ";
21 case FlowTest::LE:
22 return "LE";
23 case FlowTest::GT:
24 return "GT";
25 case FlowTest::NE:
26 return "NE";
27 case FlowTest::GE:
28 return "GE";
29 case FlowTest::NUM:
30 return "NUM";
31 case FlowTest::NaN:
32 return "NAN";
33 case FlowTest::LTU:
34 return "LTU";
35 case FlowTest::EQU:
36 return "EQU";
37 case FlowTest::LEU:
38 return "LEU";
39 case FlowTest::GTU:
40 return "GTU";
41 case FlowTest::NEU:
42 return "NEU";
43 case FlowTest::GEU:
44 return "GEU";
45 case FlowTest::T:
46 return "T";
47 case FlowTest::OFF:
48 return "OFF";
49 case FlowTest::LO:
50 return "LO";
51 case FlowTest::SFF:
52 return "SFF";
53 case FlowTest::LS:
54 return "LS";
55 case FlowTest::HI:
56 return "HI";
57 case FlowTest::SFT:
58 return "SFT";
59 case FlowTest::HS:
60 return "HS";
61 case FlowTest::OFT:
62 return "OFT";
63 case FlowTest::CSM_TA:
64 return "CSM_TA";
65 case FlowTest::CSM_TR:
66 return "CSM_TR";
67 case FlowTest::CSM_MX:
68 return "CSM_MX";
69 case FlowTest::FCSM_TA:
70 return "FCSM_TA";
71 case FlowTest::FCSM_TR:
72 return "FCSM_TR";
73 case FlowTest::FCSM_MX:
74 return "FCSM_MX";
75 case FlowTest::RLE:
76 return "RLE";
77 case FlowTest::RGT:
78 return "RGT";
79 }
80 return fmt::format("<invalid flow test {}>", static_cast<int>(flow_test));
81}
82
83} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/flow_test.h b/src/shader_recompiler/frontend/ir/flow_test.h
new file mode 100644
index 000000000..09e113773
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/flow_test.h
@@ -0,0 +1,62 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <string>
8#include <fmt/format.h>
9
10#include "common/common_types.h"
11
12namespace Shader::IR {
13
14enum class FlowTest : u64 {
15 F,
16 LT,
17 EQ,
18 LE,
19 GT,
20 NE,
21 GE,
22 NUM,
23 NaN,
24 LTU,
25 EQU,
26 LEU,
27 GTU,
28 NEU,
29 GEU,
30 T,
31 OFF,
32 LO,
33 SFF,
34 LS,
35 HI,
36 SFT,
37 HS,
38 OFT,
39 CSM_TA,
40 CSM_TR,
41 CSM_MX,
42 FCSM_TA,
43 FCSM_TR,
44 FCSM_MX,
45 RLE,
46 RGT,
47};
48
49[[nodiscard]] std::string NameOf(FlowTest flow_test);
50
51} // namespace Shader::IR
52
53template <>
54struct fmt::formatter<Shader::IR::FlowTest> {
55 constexpr auto parse(format_parse_context& ctx) {
56 return ctx.begin();
57 }
58 template <typename FormatContext>
59 auto format(const Shader::IR::FlowTest& flow_test, FormatContext& ctx) {
60 return fmt::format_to(ctx.out(), "{}", Shader::IR::NameOf(flow_test));
61 }
62};
diff --git a/src/shader_recompiler/frontend/ir/ir_emitter.cpp b/src/shader_recompiler/frontend/ir/ir_emitter.cpp
new file mode 100644
index 000000000..13159a68d
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/ir_emitter.cpp
@@ -0,0 +1,2017 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include "common/bit_cast.h"
6#include "shader_recompiler/frontend/ir/ir_emitter.h"
7#include "shader_recompiler/frontend/ir/value.h"
8
9namespace Shader::IR {
10namespace {
11[[noreturn]] void ThrowInvalidType(Type type) {
12 throw InvalidArgument("Invalid type {}", type);
13}
14
15Value MakeLodClampPair(IREmitter& ir, const F32& bias_lod, const F32& lod_clamp) {
16 if (!bias_lod.IsEmpty() && !lod_clamp.IsEmpty()) {
17 return ir.CompositeConstruct(bias_lod, lod_clamp);
18 } else if (!bias_lod.IsEmpty()) {
19 return bias_lod;
20 } else if (!lod_clamp.IsEmpty()) {
21 return lod_clamp;
22 } else {
23 return Value{};
24 }
25}
26} // Anonymous namespace
27
28U1 IREmitter::Imm1(bool value) const {
29 return U1{Value{value}};
30}
31
32U8 IREmitter::Imm8(u8 value) const {
33 return U8{Value{value}};
34}
35
36U16 IREmitter::Imm16(u16 value) const {
37 return U16{Value{value}};
38}
39
40U32 IREmitter::Imm32(u32 value) const {
41 return U32{Value{value}};
42}
43
44U32 IREmitter::Imm32(s32 value) const {
45 return U32{Value{static_cast<u32>(value)}};
46}
47
48F32 IREmitter::Imm32(f32 value) const {
49 return F32{Value{value}};
50}
51
52U64 IREmitter::Imm64(u64 value) const {
53 return U64{Value{value}};
54}
55
56U64 IREmitter::Imm64(s64 value) const {
57 return U64{Value{static_cast<u64>(value)}};
58}
59
60F64 IREmitter::Imm64(f64 value) const {
61 return F64{Value{value}};
62}
63
64U1 IREmitter::ConditionRef(const U1& value) {
65 return Inst<U1>(Opcode::ConditionRef, value);
66}
67
68void IREmitter::Reference(const Value& value) {
69 Inst(Opcode::Reference, value);
70}
71
72void IREmitter::PhiMove(IR::Inst& phi, const Value& value) {
73 Inst(Opcode::PhiMove, Value{&phi}, value);
74}
75
76void IREmitter::Prologue() {
77 Inst(Opcode::Prologue);
78}
79
80void IREmitter::Epilogue() {
81 Inst(Opcode::Epilogue);
82}
83
84void IREmitter::DemoteToHelperInvocation() {
85 Inst(Opcode::DemoteToHelperInvocation);
86}
87
88void IREmitter::EmitVertex(const U32& stream) {
89 Inst(Opcode::EmitVertex, stream);
90}
91
92void IREmitter::EndPrimitive(const U32& stream) {
93 Inst(Opcode::EndPrimitive, stream);
94}
95
96void IREmitter::Barrier() {
97 Inst(Opcode::Barrier);
98}
99
100void IREmitter::WorkgroupMemoryBarrier() {
101 Inst(Opcode::WorkgroupMemoryBarrier);
102}
103
104void IREmitter::DeviceMemoryBarrier() {
105 Inst(Opcode::DeviceMemoryBarrier);
106}
107
108U32 IREmitter::GetReg(IR::Reg reg) {
109 return Inst<U32>(Opcode::GetRegister, reg);
110}
111
112void IREmitter::SetReg(IR::Reg reg, const U32& value) {
113 Inst(Opcode::SetRegister, reg, value);
114}
115
116U1 IREmitter::GetPred(IR::Pred pred, bool is_negated) {
117 if (pred == Pred::PT) {
118 return Imm1(!is_negated);
119 }
120 const U1 value{Inst<U1>(Opcode::GetPred, pred)};
121 if (is_negated) {
122 return Inst<U1>(Opcode::LogicalNot, value);
123 } else {
124 return value;
125 }
126}
127
128void IREmitter::SetPred(IR::Pred pred, const U1& value) {
129 if (pred != IR::Pred::PT) {
130 Inst(Opcode::SetPred, pred, value);
131 }
132}
133
134U1 IREmitter::GetGotoVariable(u32 id) {
135 return Inst<U1>(Opcode::GetGotoVariable, id);
136}
137
138void IREmitter::SetGotoVariable(u32 id, const U1& value) {
139 Inst(Opcode::SetGotoVariable, id, value);
140}
141
142U32 IREmitter::GetIndirectBranchVariable() {
143 return Inst<U32>(Opcode::GetIndirectBranchVariable);
144}
145
146void IREmitter::SetIndirectBranchVariable(const U32& value) {
147 Inst(Opcode::SetIndirectBranchVariable, value);
148}
149
150U32 IREmitter::GetCbuf(const U32& binding, const U32& byte_offset) {
151 return Inst<U32>(Opcode::GetCbufU32, binding, byte_offset);
152}
153
154Value IREmitter::GetCbuf(const U32& binding, const U32& byte_offset, size_t bitsize,
155 bool is_signed) {
156 switch (bitsize) {
157 case 8:
158 return Inst<U32>(is_signed ? Opcode::GetCbufS8 : Opcode::GetCbufU8, binding, byte_offset);
159 case 16:
160 return Inst<U32>(is_signed ? Opcode::GetCbufS16 : Opcode::GetCbufU16, binding, byte_offset);
161 case 32:
162 return Inst<U32>(Opcode::GetCbufU32, binding, byte_offset);
163 case 64:
164 return Inst(Opcode::GetCbufU32x2, binding, byte_offset);
165 default:
166 throw InvalidArgument("Invalid bit size {}", bitsize);
167 }
168}
169
170F32 IREmitter::GetFloatCbuf(const U32& binding, const U32& byte_offset) {
171 return Inst<F32>(Opcode::GetCbufF32, binding, byte_offset);
172}
173
174U1 IREmitter::GetZFlag() {
175 return Inst<U1>(Opcode::GetZFlag);
176}
177
178U1 IREmitter::GetSFlag() {
179 return Inst<U1>(Opcode::GetSFlag);
180}
181
182U1 IREmitter::GetCFlag() {
183 return Inst<U1>(Opcode::GetCFlag);
184}
185
186U1 IREmitter::GetOFlag() {
187 return Inst<U1>(Opcode::GetOFlag);
188}
189
190void IREmitter::SetZFlag(const U1& value) {
191 Inst(Opcode::SetZFlag, value);
192}
193
194void IREmitter::SetSFlag(const U1& value) {
195 Inst(Opcode::SetSFlag, value);
196}
197
198void IREmitter::SetCFlag(const U1& value) {
199 Inst(Opcode::SetCFlag, value);
200}
201
202void IREmitter::SetOFlag(const U1& value) {
203 Inst(Opcode::SetOFlag, value);
204}
205
206static U1 GetFlowTest(IREmitter& ir, FlowTest flow_test) {
207 switch (flow_test) {
208 case FlowTest::F:
209 return ir.Imm1(false);
210 case FlowTest::LT:
211 return ir.LogicalXor(ir.LogicalAnd(ir.GetSFlag(), ir.LogicalNot(ir.GetZFlag())),
212 ir.GetOFlag());
213 case FlowTest::EQ:
214 return ir.LogicalAnd(ir.LogicalNot(ir.GetSFlag()), ir.GetZFlag());
215 case FlowTest::LE:
216 return ir.LogicalXor(ir.GetSFlag(), ir.LogicalOr(ir.GetZFlag(), ir.GetOFlag()));
217 case FlowTest::GT:
218 return ir.LogicalAnd(ir.LogicalXor(ir.LogicalNot(ir.GetSFlag()), ir.GetOFlag()),
219 ir.LogicalNot(ir.GetZFlag()));
220 case FlowTest::NE:
221 return ir.LogicalNot(ir.GetZFlag());
222 case FlowTest::GE:
223 return ir.LogicalNot(ir.LogicalXor(ir.GetSFlag(), ir.GetOFlag()));
224 case FlowTest::NUM:
225 return ir.LogicalOr(ir.LogicalNot(ir.GetSFlag()), ir.LogicalNot(ir.GetZFlag()));
226 case FlowTest::NaN:
227 return ir.LogicalAnd(ir.GetSFlag(), ir.GetZFlag());
228 case FlowTest::LTU:
229 return ir.LogicalXor(ir.GetSFlag(), ir.GetOFlag());
230 case FlowTest::EQU:
231 return ir.GetZFlag();
232 case FlowTest::LEU:
233 return ir.LogicalOr(ir.LogicalXor(ir.GetSFlag(), ir.GetOFlag()), ir.GetZFlag());
234 case FlowTest::GTU:
235 return ir.LogicalXor(ir.LogicalNot(ir.GetSFlag()),
236 ir.LogicalOr(ir.GetZFlag(), ir.GetOFlag()));
237 case FlowTest::NEU:
238 return ir.LogicalOr(ir.GetSFlag(), ir.LogicalNot(ir.GetZFlag()));
239 case FlowTest::GEU:
240 return ir.LogicalXor(ir.LogicalOr(ir.LogicalNot(ir.GetSFlag()), ir.GetZFlag()),
241 ir.GetOFlag());
242 case FlowTest::T:
243 return ir.Imm1(true);
244 case FlowTest::OFF:
245 return ir.LogicalNot(ir.GetOFlag());
246 case FlowTest::LO:
247 return ir.LogicalNot(ir.GetCFlag());
248 case FlowTest::SFF:
249 return ir.LogicalNot(ir.GetSFlag());
250 case FlowTest::LS:
251 return ir.LogicalOr(ir.GetZFlag(), ir.LogicalNot(ir.GetCFlag()));
252 case FlowTest::HI:
253 return ir.LogicalAnd(ir.GetCFlag(), ir.LogicalNot(ir.GetZFlag()));
254 case FlowTest::SFT:
255 return ir.GetSFlag();
256 case FlowTest::HS:
257 return ir.GetCFlag();
258 case FlowTest::OFT:
259 return ir.GetOFlag();
260 case FlowTest::RLE:
261 return ir.LogicalOr(ir.GetSFlag(), ir.GetZFlag());
262 case FlowTest::RGT:
263 return ir.LogicalAnd(ir.LogicalNot(ir.GetSFlag()), ir.LogicalNot(ir.GetZFlag()));
264 case FlowTest::FCSM_TR:
265 LOG_WARNING(Shader, "(STUBBED) FCSM_TR");
266 return ir.Imm1(false);
267 case FlowTest::CSM_TA:
268 case FlowTest::CSM_TR:
269 case FlowTest::CSM_MX:
270 case FlowTest::FCSM_TA:
271 case FlowTest::FCSM_MX:
272 default:
273 throw NotImplementedException("Flow test {}", flow_test);
274 }
275}
276
277U1 IREmitter::Condition(IR::Condition cond) {
278 const FlowTest flow_test{cond.GetFlowTest()};
279 const auto [pred, is_negated]{cond.GetPred()};
280 if (flow_test == FlowTest::T) {
281 return GetPred(pred, is_negated);
282 }
283 return LogicalAnd(GetPred(pred, is_negated), GetFlowTest(*this, flow_test));
284}
285
286U1 IREmitter::GetFlowTestResult(FlowTest test) {
287 return GetFlowTest(*this, test);
288}
289
290F32 IREmitter::GetAttribute(IR::Attribute attribute) {
291 return GetAttribute(attribute, Imm32(0));
292}
293
294F32 IREmitter::GetAttribute(IR::Attribute attribute, const U32& vertex) {
295 return Inst<F32>(Opcode::GetAttribute, attribute, vertex);
296}
297
298void IREmitter::SetAttribute(IR::Attribute attribute, const F32& value, const U32& vertex) {
299 Inst(Opcode::SetAttribute, attribute, value, vertex);
300}
301
302F32 IREmitter::GetAttributeIndexed(const U32& phys_address) {
303 return GetAttributeIndexed(phys_address, Imm32(0));
304}
305
306F32 IREmitter::GetAttributeIndexed(const U32& phys_address, const U32& vertex) {
307 return Inst<F32>(Opcode::GetAttributeIndexed, phys_address, vertex);
308}
309
310void IREmitter::SetAttributeIndexed(const U32& phys_address, const F32& value, const U32& vertex) {
311 Inst(Opcode::SetAttributeIndexed, phys_address, value, vertex);
312}
313
314F32 IREmitter::GetPatch(Patch patch) {
315 return Inst<F32>(Opcode::GetPatch, patch);
316}
317
318void IREmitter::SetPatch(Patch patch, const F32& value) {
319 Inst(Opcode::SetPatch, patch, value);
320}
321
322void IREmitter::SetFragColor(u32 index, u32 component, const F32& value) {
323 Inst(Opcode::SetFragColor, Imm32(index), Imm32(component), value);
324}
325
326void IREmitter::SetSampleMask(const U32& value) {
327 Inst(Opcode::SetSampleMask, value);
328}
329
330void IREmitter::SetFragDepth(const F32& value) {
331 Inst(Opcode::SetFragDepth, value);
332}
333
334U32 IREmitter::WorkgroupIdX() {
335 return U32{CompositeExtract(Inst(Opcode::WorkgroupId), 0)};
336}
337
338U32 IREmitter::WorkgroupIdY() {
339 return U32{CompositeExtract(Inst(Opcode::WorkgroupId), 1)};
340}
341
342U32 IREmitter::WorkgroupIdZ() {
343 return U32{CompositeExtract(Inst(Opcode::WorkgroupId), 2)};
344}
345
346Value IREmitter::LocalInvocationId() {
347 return Inst(Opcode::LocalInvocationId);
348}
349
350U32 IREmitter::LocalInvocationIdX() {
351 return U32{CompositeExtract(Inst(Opcode::LocalInvocationId), 0)};
352}
353
354U32 IREmitter::LocalInvocationIdY() {
355 return U32{CompositeExtract(Inst(Opcode::LocalInvocationId), 1)};
356}
357
358U32 IREmitter::LocalInvocationIdZ() {
359 return U32{CompositeExtract(Inst(Opcode::LocalInvocationId), 2)};
360}
361
362U32 IREmitter::InvocationId() {
363 return Inst<U32>(Opcode::InvocationId);
364}
365
366U32 IREmitter::SampleId() {
367 return Inst<U32>(Opcode::SampleId);
368}
369
370U1 IREmitter::IsHelperInvocation() {
371 return Inst<U1>(Opcode::IsHelperInvocation);
372}
373
374F32 IREmitter::YDirection() {
375 return Inst<F32>(Opcode::YDirection);
376}
377
378U32 IREmitter::LaneId() {
379 return Inst<U32>(Opcode::LaneId);
380}
381
382U32 IREmitter::LoadGlobalU8(const U64& address) {
383 return Inst<U32>(Opcode::LoadGlobalU8, address);
384}
385
386U32 IREmitter::LoadGlobalS8(const U64& address) {
387 return Inst<U32>(Opcode::LoadGlobalS8, address);
388}
389
390U32 IREmitter::LoadGlobalU16(const U64& address) {
391 return Inst<U32>(Opcode::LoadGlobalU16, address);
392}
393
394U32 IREmitter::LoadGlobalS16(const U64& address) {
395 return Inst<U32>(Opcode::LoadGlobalS16, address);
396}
397
398U32 IREmitter::LoadGlobal32(const U64& address) {
399 return Inst<U32>(Opcode::LoadGlobal32, address);
400}
401
402Value IREmitter::LoadGlobal64(const U64& address) {
403 return Inst<Value>(Opcode::LoadGlobal64, address);
404}
405
406Value IREmitter::LoadGlobal128(const U64& address) {
407 return Inst<Value>(Opcode::LoadGlobal128, address);
408}
409
410void IREmitter::WriteGlobalU8(const U64& address, const U32& value) {
411 Inst(Opcode::WriteGlobalU8, address, value);
412}
413
414void IREmitter::WriteGlobalS8(const U64& address, const U32& value) {
415 Inst(Opcode::WriteGlobalS8, address, value);
416}
417
418void IREmitter::WriteGlobalU16(const U64& address, const U32& value) {
419 Inst(Opcode::WriteGlobalU16, address, value);
420}
421
422void IREmitter::WriteGlobalS16(const U64& address, const U32& value) {
423 Inst(Opcode::WriteGlobalS16, address, value);
424}
425
426void IREmitter::WriteGlobal32(const U64& address, const U32& value) {
427 Inst(Opcode::WriteGlobal32, address, value);
428}
429
430void IREmitter::WriteGlobal64(const U64& address, const IR::Value& vector) {
431 Inst(Opcode::WriteGlobal64, address, vector);
432}
433
434void IREmitter::WriteGlobal128(const U64& address, const IR::Value& vector) {
435 Inst(Opcode::WriteGlobal128, address, vector);
436}
437
438U32 IREmitter::LoadLocal(const IR::U32& word_offset) {
439 return Inst<U32>(Opcode::LoadLocal, word_offset);
440}
441
442void IREmitter::WriteLocal(const IR::U32& word_offset, const IR::U32& value) {
443 Inst(Opcode::WriteLocal, word_offset, value);
444}
445
446Value IREmitter::LoadShared(int bit_size, bool is_signed, const IR::U32& offset) {
447 switch (bit_size) {
448 case 8:
449 return Inst(is_signed ? Opcode::LoadSharedS8 : Opcode::LoadSharedU8, offset);
450 case 16:
451 return Inst(is_signed ? Opcode::LoadSharedS16 : Opcode::LoadSharedU16, offset);
452 case 32:
453 return Inst(Opcode::LoadSharedU32, offset);
454 case 64:
455 return Inst(Opcode::LoadSharedU64, offset);
456 case 128:
457 return Inst(Opcode::LoadSharedU128, offset);
458 }
459 throw InvalidArgument("Invalid bit size {}", bit_size);
460}
461
462void IREmitter::WriteShared(int bit_size, const IR::U32& offset, const IR::Value& value) {
463 switch (bit_size) {
464 case 8:
465 Inst(Opcode::WriteSharedU8, offset, value);
466 break;
467 case 16:
468 Inst(Opcode::WriteSharedU16, offset, value);
469 break;
470 case 32:
471 Inst(Opcode::WriteSharedU32, offset, value);
472 break;
473 case 64:
474 Inst(Opcode::WriteSharedU64, offset, value);
475 break;
476 case 128:
477 Inst(Opcode::WriteSharedU128, offset, value);
478 break;
479 default:
480 throw InvalidArgument("Invalid bit size {}", bit_size);
481 }
482}
483
484U1 IREmitter::GetZeroFromOp(const Value& op) {
485 return Inst<U1>(Opcode::GetZeroFromOp, op);
486}
487
488U1 IREmitter::GetSignFromOp(const Value& op) {
489 return Inst<U1>(Opcode::GetSignFromOp, op);
490}
491
492U1 IREmitter::GetCarryFromOp(const Value& op) {
493 return Inst<U1>(Opcode::GetCarryFromOp, op);
494}
495
496U1 IREmitter::GetOverflowFromOp(const Value& op) {
497 return Inst<U1>(Opcode::GetOverflowFromOp, op);
498}
499
500U1 IREmitter::GetSparseFromOp(const Value& op) {
501 return Inst<U1>(Opcode::GetSparseFromOp, op);
502}
503
504U1 IREmitter::GetInBoundsFromOp(const Value& op) {
505 return Inst<U1>(Opcode::GetInBoundsFromOp, op);
506}
507
508F16F32F64 IREmitter::FPAdd(const F16F32F64& a, const F16F32F64& b, FpControl control) {
509 if (a.Type() != b.Type()) {
510 throw InvalidArgument("Mismatching types {} and {}", a.Type(), b.Type());
511 }
512 switch (a.Type()) {
513 case Type::F16:
514 return Inst<F16>(Opcode::FPAdd16, Flags{control}, a, b);
515 case Type::F32:
516 return Inst<F32>(Opcode::FPAdd32, Flags{control}, a, b);
517 case Type::F64:
518 return Inst<F64>(Opcode::FPAdd64, Flags{control}, a, b);
519 default:
520 ThrowInvalidType(a.Type());
521 }
522}
523
524Value IREmitter::CompositeConstruct(const Value& e1, const Value& e2) {
525 if (e1.Type() != e2.Type()) {
526 throw InvalidArgument("Mismatching types {} and {}", e1.Type(), e2.Type());
527 }
528 switch (e1.Type()) {
529 case Type::U32:
530 return Inst(Opcode::CompositeConstructU32x2, e1, e2);
531 case Type::F16:
532 return Inst(Opcode::CompositeConstructF16x2, e1, e2);
533 case Type::F32:
534 return Inst(Opcode::CompositeConstructF32x2, e1, e2);
535 case Type::F64:
536 return Inst(Opcode::CompositeConstructF64x2, e1, e2);
537 default:
538 ThrowInvalidType(e1.Type());
539 }
540}
541
542Value IREmitter::CompositeConstruct(const Value& e1, const Value& e2, const Value& e3) {
543 if (e1.Type() != e2.Type() || e1.Type() != e3.Type()) {
544 throw InvalidArgument("Mismatching types {}, {}, and {}", e1.Type(), e2.Type(), e3.Type());
545 }
546 switch (e1.Type()) {
547 case Type::U32:
548 return Inst(Opcode::CompositeConstructU32x3, e1, e2, e3);
549 case Type::F16:
550 return Inst(Opcode::CompositeConstructF16x3, e1, e2, e3);
551 case Type::F32:
552 return Inst(Opcode::CompositeConstructF32x3, e1, e2, e3);
553 case Type::F64:
554 return Inst(Opcode::CompositeConstructF64x3, e1, e2, e3);
555 default:
556 ThrowInvalidType(e1.Type());
557 }
558}
559
560Value IREmitter::CompositeConstruct(const Value& e1, const Value& e2, const Value& e3,
561 const Value& e4) {
562 if (e1.Type() != e2.Type() || e1.Type() != e3.Type() || e1.Type() != e4.Type()) {
563 throw InvalidArgument("Mismatching types {}, {}, {}, and {}", e1.Type(), e2.Type(),
564 e3.Type(), e4.Type());
565 }
566 switch (e1.Type()) {
567 case Type::U32:
568 return Inst(Opcode::CompositeConstructU32x4, e1, e2, e3, e4);
569 case Type::F16:
570 return Inst(Opcode::CompositeConstructF16x4, e1, e2, e3, e4);
571 case Type::F32:
572 return Inst(Opcode::CompositeConstructF32x4, e1, e2, e3, e4);
573 case Type::F64:
574 return Inst(Opcode::CompositeConstructF64x4, e1, e2, e3, e4);
575 default:
576 ThrowInvalidType(e1.Type());
577 }
578}
579
580Value IREmitter::CompositeExtract(const Value& vector, size_t element) {
581 const auto read{[&](Opcode opcode, size_t limit) -> Value {
582 if (element >= limit) {
583 throw InvalidArgument("Out of bounds element {}", element);
584 }
585 return Inst(opcode, vector, Value{static_cast<u32>(element)});
586 }};
587 switch (vector.Type()) {
588 case Type::U32x2:
589 return read(Opcode::CompositeExtractU32x2, 2);
590 case Type::U32x3:
591 return read(Opcode::CompositeExtractU32x3, 3);
592 case Type::U32x4:
593 return read(Opcode::CompositeExtractU32x4, 4);
594 case Type::F16x2:
595 return read(Opcode::CompositeExtractF16x2, 2);
596 case Type::F16x3:
597 return read(Opcode::CompositeExtractF16x3, 3);
598 case Type::F16x4:
599 return read(Opcode::CompositeExtractF16x4, 4);
600 case Type::F32x2:
601 return read(Opcode::CompositeExtractF32x2, 2);
602 case Type::F32x3:
603 return read(Opcode::CompositeExtractF32x3, 3);
604 case Type::F32x4:
605 return read(Opcode::CompositeExtractF32x4, 4);
606 case Type::F64x2:
607 return read(Opcode::CompositeExtractF64x2, 2);
608 case Type::F64x3:
609 return read(Opcode::CompositeExtractF64x3, 3);
610 case Type::F64x4:
611 return read(Opcode::CompositeExtractF64x4, 4);
612 default:
613 ThrowInvalidType(vector.Type());
614 }
615}
616
617Value IREmitter::CompositeInsert(const Value& vector, const Value& object, size_t element) {
618 const auto insert{[&](Opcode opcode, size_t limit) {
619 if (element >= limit) {
620 throw InvalidArgument("Out of bounds element {}", element);
621 }
622 return Inst(opcode, vector, object, Value{static_cast<u32>(element)});
623 }};
624 switch (vector.Type()) {
625 case Type::U32x2:
626 return insert(Opcode::CompositeInsertU32x2, 2);
627 case Type::U32x3:
628 return insert(Opcode::CompositeInsertU32x3, 3);
629 case Type::U32x4:
630 return insert(Opcode::CompositeInsertU32x4, 4);
631 case Type::F16x2:
632 return insert(Opcode::CompositeInsertF16x2, 2);
633 case Type::F16x3:
634 return insert(Opcode::CompositeInsertF16x3, 3);
635 case Type::F16x4:
636 return insert(Opcode::CompositeInsertF16x4, 4);
637 case Type::F32x2:
638 return insert(Opcode::CompositeInsertF32x2, 2);
639 case Type::F32x3:
640 return insert(Opcode::CompositeInsertF32x3, 3);
641 case Type::F32x4:
642 return insert(Opcode::CompositeInsertF32x4, 4);
643 case Type::F64x2:
644 return insert(Opcode::CompositeInsertF64x2, 2);
645 case Type::F64x3:
646 return insert(Opcode::CompositeInsertF64x3, 3);
647 case Type::F64x4:
648 return insert(Opcode::CompositeInsertF64x4, 4);
649 default:
650 ThrowInvalidType(vector.Type());
651 }
652}
653
654Value IREmitter::Select(const U1& condition, const Value& true_value, const Value& false_value) {
655 if (true_value.Type() != false_value.Type()) {
656 throw InvalidArgument("Mismatching types {} and {}", true_value.Type(), false_value.Type());
657 }
658 switch (true_value.Type()) {
659 case Type::U1:
660 return Inst(Opcode::SelectU1, condition, true_value, false_value);
661 case Type::U8:
662 return Inst(Opcode::SelectU8, condition, true_value, false_value);
663 case Type::U16:
664 return Inst(Opcode::SelectU16, condition, true_value, false_value);
665 case Type::U32:
666 return Inst(Opcode::SelectU32, condition, true_value, false_value);
667 case Type::U64:
668 return Inst(Opcode::SelectU64, condition, true_value, false_value);
669 case Type::F32:
670 return Inst(Opcode::SelectF32, condition, true_value, false_value);
671 case Type::F64:
672 return Inst(Opcode::SelectF64, condition, true_value, false_value);
673 default:
674 throw InvalidArgument("Invalid type {}", true_value.Type());
675 }
676}
677
678template <>
679IR::U32 IREmitter::BitCast<IR::U32, IR::F32>(const IR::F32& value) {
680 return Inst<IR::U32>(Opcode::BitCastU32F32, value);
681}
682
683template <>
684IR::F32 IREmitter::BitCast<IR::F32, IR::U32>(const IR::U32& value) {
685 return Inst<IR::F32>(Opcode::BitCastF32U32, value);
686}
687
688template <>
689IR::U16 IREmitter::BitCast<IR::U16, IR::F16>(const IR::F16& value) {
690 return Inst<IR::U16>(Opcode::BitCastU16F16, value);
691}
692
693template <>
694IR::F16 IREmitter::BitCast<IR::F16, IR::U16>(const IR::U16& value) {
695 return Inst<IR::F16>(Opcode::BitCastF16U16, value);
696}
697
698template <>
699IR::U64 IREmitter::BitCast<IR::U64, IR::F64>(const IR::F64& value) {
700 return Inst<IR::U64>(Opcode::BitCastU64F64, value);
701}
702
703template <>
704IR::F64 IREmitter::BitCast<IR::F64, IR::U64>(const IR::U64& value) {
705 return Inst<IR::F64>(Opcode::BitCastF64U64, value);
706}
707
708U64 IREmitter::PackUint2x32(const Value& vector) {
709 return Inst<U64>(Opcode::PackUint2x32, vector);
710}
711
712Value IREmitter::UnpackUint2x32(const U64& value) {
713 return Inst<Value>(Opcode::UnpackUint2x32, value);
714}
715
716U32 IREmitter::PackFloat2x16(const Value& vector) {
717 return Inst<U32>(Opcode::PackFloat2x16, vector);
718}
719
720Value IREmitter::UnpackFloat2x16(const U32& value) {
721 return Inst(Opcode::UnpackFloat2x16, value);
722}
723
724U32 IREmitter::PackHalf2x16(const Value& vector) {
725 return Inst<U32>(Opcode::PackHalf2x16, vector);
726}
727
728Value IREmitter::UnpackHalf2x16(const U32& value) {
729 return Inst(Opcode::UnpackHalf2x16, value);
730}
731
732F64 IREmitter::PackDouble2x32(const Value& vector) {
733 return Inst<F64>(Opcode::PackDouble2x32, vector);
734}
735
736Value IREmitter::UnpackDouble2x32(const F64& value) {
737 return Inst<Value>(Opcode::UnpackDouble2x32, value);
738}
739
740F16F32F64 IREmitter::FPMul(const F16F32F64& a, const F16F32F64& b, FpControl control) {
741 if (a.Type() != b.Type()) {
742 throw InvalidArgument("Mismatching types {} and {}", a.Type(), b.Type());
743 }
744 switch (a.Type()) {
745 case Type::F16:
746 return Inst<F16>(Opcode::FPMul16, Flags{control}, a, b);
747 case Type::F32:
748 return Inst<F32>(Opcode::FPMul32, Flags{control}, a, b);
749 case Type::F64:
750 return Inst<F64>(Opcode::FPMul64, Flags{control}, a, b);
751 default:
752 ThrowInvalidType(a.Type());
753 }
754}
755
756F16F32F64 IREmitter::FPFma(const F16F32F64& a, const F16F32F64& b, const F16F32F64& c,
757 FpControl control) {
758 if (a.Type() != b.Type() || a.Type() != c.Type()) {
759 throw InvalidArgument("Mismatching types {}, {}, and {}", a.Type(), b.Type(), c.Type());
760 }
761 switch (a.Type()) {
762 case Type::F16:
763 return Inst<F16>(Opcode::FPFma16, Flags{control}, a, b, c);
764 case Type::F32:
765 return Inst<F32>(Opcode::FPFma32, Flags{control}, a, b, c);
766 case Type::F64:
767 return Inst<F64>(Opcode::FPFma64, Flags{control}, a, b, c);
768 default:
769 ThrowInvalidType(a.Type());
770 }
771}
772
773F16F32F64 IREmitter::FPAbs(const F16F32F64& value) {
774 switch (value.Type()) {
775 case Type::F16:
776 return Inst<F16>(Opcode::FPAbs16, value);
777 case Type::F32:
778 return Inst<F32>(Opcode::FPAbs32, value);
779 case Type::F64:
780 return Inst<F64>(Opcode::FPAbs64, value);
781 default:
782 ThrowInvalidType(value.Type());
783 }
784}
785
786F16F32F64 IREmitter::FPNeg(const F16F32F64& value) {
787 switch (value.Type()) {
788 case Type::F16:
789 return Inst<F16>(Opcode::FPNeg16, value);
790 case Type::F32:
791 return Inst<F32>(Opcode::FPNeg32, value);
792 case Type::F64:
793 return Inst<F64>(Opcode::FPNeg64, value);
794 default:
795 ThrowInvalidType(value.Type());
796 }
797}
798
799F16F32F64 IREmitter::FPAbsNeg(const F16F32F64& value, bool abs, bool neg) {
800 F16F32F64 result{value};
801 if (abs) {
802 result = FPAbs(result);
803 }
804 if (neg) {
805 result = FPNeg(result);
806 }
807 return result;
808}
809
810F32 IREmitter::FPCos(const F32& value) {
811 return Inst<F32>(Opcode::FPCos, value);
812}
813
814F32 IREmitter::FPSin(const F32& value) {
815 return Inst<F32>(Opcode::FPSin, value);
816}
817
818F32 IREmitter::FPExp2(const F32& value) {
819 return Inst<F32>(Opcode::FPExp2, value);
820}
821
822F32 IREmitter::FPLog2(const F32& value) {
823 return Inst<F32>(Opcode::FPLog2, value);
824}
825
826F32F64 IREmitter::FPRecip(const F32F64& value) {
827 switch (value.Type()) {
828 case Type::F32:
829 return Inst<F32>(Opcode::FPRecip32, value);
830 case Type::F64:
831 return Inst<F64>(Opcode::FPRecip64, value);
832 default:
833 ThrowInvalidType(value.Type());
834 }
835}
836
837F32F64 IREmitter::FPRecipSqrt(const F32F64& value) {
838 switch (value.Type()) {
839 case Type::F32:
840 return Inst<F32>(Opcode::FPRecipSqrt32, value);
841 case Type::F64:
842 return Inst<F64>(Opcode::FPRecipSqrt64, value);
843 default:
844 ThrowInvalidType(value.Type());
845 }
846}
847
848F32 IREmitter::FPSqrt(const F32& value) {
849 return Inst<F32>(Opcode::FPSqrt, value);
850}
851
852F16F32F64 IREmitter::FPSaturate(const F16F32F64& value) {
853 switch (value.Type()) {
854 case Type::F16:
855 return Inst<F16>(Opcode::FPSaturate16, value);
856 case Type::F32:
857 return Inst<F32>(Opcode::FPSaturate32, value);
858 case Type::F64:
859 return Inst<F64>(Opcode::FPSaturate64, value);
860 default:
861 ThrowInvalidType(value.Type());
862 }
863}
864
865F16F32F64 IREmitter::FPClamp(const F16F32F64& value, const F16F32F64& min_value,
866 const F16F32F64& max_value) {
867 if (value.Type() != min_value.Type() || value.Type() != max_value.Type()) {
868 throw InvalidArgument("Mismatching types {}, {}, and {}", value.Type(), min_value.Type(),
869 max_value.Type());
870 }
871 switch (value.Type()) {
872 case Type::F16:
873 return Inst<F16>(Opcode::FPClamp16, value, min_value, max_value);
874 case Type::F32:
875 return Inst<F32>(Opcode::FPClamp32, value, min_value, max_value);
876 case Type::F64:
877 return Inst<F64>(Opcode::FPClamp64, value, min_value, max_value);
878 default:
879 ThrowInvalidType(value.Type());
880 }
881}
882
883F16F32F64 IREmitter::FPRoundEven(const F16F32F64& value, FpControl control) {
884 switch (value.Type()) {
885 case Type::F16:
886 return Inst<F16>(Opcode::FPRoundEven16, Flags{control}, value);
887 case Type::F32:
888 return Inst<F32>(Opcode::FPRoundEven32, Flags{control}, value);
889 case Type::F64:
890 return Inst<F64>(Opcode::FPRoundEven64, Flags{control}, value);
891 default:
892 ThrowInvalidType(value.Type());
893 }
894}
895
896F16F32F64 IREmitter::FPFloor(const F16F32F64& value, FpControl control) {
897 switch (value.Type()) {
898 case Type::F16:
899 return Inst<F16>(Opcode::FPFloor16, Flags{control}, value);
900 case Type::F32:
901 return Inst<F32>(Opcode::FPFloor32, Flags{control}, value);
902 case Type::F64:
903 return Inst<F64>(Opcode::FPFloor64, Flags{control}, value);
904 default:
905 ThrowInvalidType(value.Type());
906 }
907}
908
909F16F32F64 IREmitter::FPCeil(const F16F32F64& value, FpControl control) {
910 switch (value.Type()) {
911 case Type::F16:
912 return Inst<F16>(Opcode::FPCeil16, Flags{control}, value);
913 case Type::F32:
914 return Inst<F32>(Opcode::FPCeil32, Flags{control}, value);
915 case Type::F64:
916 return Inst<F64>(Opcode::FPCeil64, Flags{control}, value);
917 default:
918 ThrowInvalidType(value.Type());
919 }
920}
921
922F16F32F64 IREmitter::FPTrunc(const F16F32F64& value, FpControl control) {
923 switch (value.Type()) {
924 case Type::F16:
925 return Inst<F16>(Opcode::FPTrunc16, Flags{control}, value);
926 case Type::F32:
927 return Inst<F32>(Opcode::FPTrunc32, Flags{control}, value);
928 case Type::F64:
929 return Inst<F64>(Opcode::FPTrunc64, Flags{control}, value);
930 default:
931 ThrowInvalidType(value.Type());
932 }
933}
934
935U1 IREmitter::FPEqual(const F16F32F64& lhs, const F16F32F64& rhs, FpControl control, bool ordered) {
936 if (lhs.Type() != rhs.Type()) {
937 throw InvalidArgument("Mismatching types {} and {}", lhs.Type(), rhs.Type());
938 }
939 switch (lhs.Type()) {
940 case Type::F16:
941 return Inst<U1>(ordered ? Opcode::FPOrdEqual16 : Opcode::FPUnordEqual16, Flags{control},
942 lhs, rhs);
943 case Type::F32:
944 return Inst<U1>(ordered ? Opcode::FPOrdEqual32 : Opcode::FPUnordEqual32, Flags{control},
945 lhs, rhs);
946 case Type::F64:
947 return Inst<U1>(ordered ? Opcode::FPOrdEqual64 : Opcode::FPUnordEqual64, Flags{control},
948 lhs, rhs);
949 default:
950 ThrowInvalidType(lhs.Type());
951 }
952}
953
954U1 IREmitter::FPNotEqual(const F16F32F64& lhs, const F16F32F64& rhs, FpControl control,
955 bool ordered) {
956 if (lhs.Type() != rhs.Type()) {
957 throw InvalidArgument("Mismatching types {} and {}", lhs.Type(), rhs.Type());
958 }
959 switch (lhs.Type()) {
960 case Type::F16:
961 return Inst<U1>(ordered ? Opcode::FPOrdNotEqual16 : Opcode::FPUnordNotEqual16,
962 Flags{control}, lhs, rhs);
963 case Type::F32:
964 return Inst<U1>(ordered ? Opcode::FPOrdNotEqual32 : Opcode::FPUnordNotEqual32,
965 Flags{control}, lhs, rhs);
966 case Type::F64:
967 return Inst<U1>(ordered ? Opcode::FPOrdNotEqual64 : Opcode::FPUnordNotEqual64,
968 Flags{control}, lhs, rhs);
969 default:
970 ThrowInvalidType(lhs.Type());
971 }
972}
973
974U1 IREmitter::FPLessThan(const F16F32F64& lhs, const F16F32F64& rhs, FpControl control,
975 bool ordered) {
976 if (lhs.Type() != rhs.Type()) {
977 throw InvalidArgument("Mismatching types {} and {}", lhs.Type(), rhs.Type());
978 }
979 switch (lhs.Type()) {
980 case Type::F16:
981 return Inst<U1>(ordered ? Opcode::FPOrdLessThan16 : Opcode::FPUnordLessThan16,
982 Flags{control}, lhs, rhs);
983 case Type::F32:
984 return Inst<U1>(ordered ? Opcode::FPOrdLessThan32 : Opcode::FPUnordLessThan32,
985 Flags{control}, lhs, rhs);
986 case Type::F64:
987 return Inst<U1>(ordered ? Opcode::FPOrdLessThan64 : Opcode::FPUnordLessThan64,
988 Flags{control}, lhs, rhs);
989 default:
990 ThrowInvalidType(lhs.Type());
991 }
992}
993
994U1 IREmitter::FPGreaterThan(const F16F32F64& lhs, const F16F32F64& rhs, FpControl control,
995 bool ordered) {
996 if (lhs.Type() != rhs.Type()) {
997 throw InvalidArgument("Mismatching types {} and {}", lhs.Type(), rhs.Type());
998 }
999 switch (lhs.Type()) {
1000 case Type::F16:
1001 return Inst<U1>(ordered ? Opcode::FPOrdGreaterThan16 : Opcode::FPUnordGreaterThan16,
1002 Flags{control}, lhs, rhs);
1003 case Type::F32:
1004 return Inst<U1>(ordered ? Opcode::FPOrdGreaterThan32 : Opcode::FPUnordGreaterThan32,
1005 Flags{control}, lhs, rhs);
1006 case Type::F64:
1007 return Inst<U1>(ordered ? Opcode::FPOrdGreaterThan64 : Opcode::FPUnordGreaterThan64,
1008 Flags{control}, lhs, rhs);
1009 default:
1010 ThrowInvalidType(lhs.Type());
1011 }
1012}
1013
1014U1 IREmitter::FPLessThanEqual(const F16F32F64& lhs, const F16F32F64& rhs, FpControl control,
1015 bool ordered) {
1016 if (lhs.Type() != rhs.Type()) {
1017 throw InvalidArgument("Mismatching types {} and {}", lhs.Type(), rhs.Type());
1018 }
1019 switch (lhs.Type()) {
1020 case Type::F16:
1021 return Inst<U1>(ordered ? Opcode::FPOrdLessThanEqual16 : Opcode::FPUnordLessThanEqual16,
1022 Flags{control}, lhs, rhs);
1023 case Type::F32:
1024 return Inst<U1>(ordered ? Opcode::FPOrdLessThanEqual32 : Opcode::FPUnordLessThanEqual32,
1025 Flags{control}, lhs, rhs);
1026 case Type::F64:
1027 return Inst<U1>(ordered ? Opcode::FPOrdLessThanEqual64 : Opcode::FPUnordLessThanEqual64,
1028 Flags{control}, lhs, rhs);
1029 default:
1030 ThrowInvalidType(lhs.Type());
1031 }
1032}
1033
1034U1 IREmitter::FPGreaterThanEqual(const F16F32F64& lhs, const F16F32F64& rhs, FpControl control,
1035 bool ordered) {
1036 if (lhs.Type() != rhs.Type()) {
1037 throw InvalidArgument("Mismatching types {} and {}", lhs.Type(), rhs.Type());
1038 }
1039 switch (lhs.Type()) {
1040 case Type::F16:
1041 return Inst<U1>(ordered ? Opcode::FPOrdGreaterThanEqual16
1042 : Opcode::FPUnordGreaterThanEqual16,
1043 Flags{control}, lhs, rhs);
1044 case Type::F32:
1045 return Inst<U1>(ordered ? Opcode::FPOrdGreaterThanEqual32
1046 : Opcode::FPUnordGreaterThanEqual32,
1047 Flags{control}, lhs, rhs);
1048 case Type::F64:
1049 return Inst<U1>(ordered ? Opcode::FPOrdGreaterThanEqual64
1050 : Opcode::FPUnordGreaterThanEqual64,
1051 Flags{control}, lhs, rhs);
1052 default:
1053 ThrowInvalidType(lhs.Type());
1054 }
1055}
1056
1057U1 IREmitter::FPIsNan(const F16F32F64& value) {
1058 switch (value.Type()) {
1059 case Type::F16:
1060 return Inst<U1>(Opcode::FPIsNan16, value);
1061 case Type::F32:
1062 return Inst<U1>(Opcode::FPIsNan32, value);
1063 case Type::F64:
1064 return Inst<U1>(Opcode::FPIsNan64, value);
1065 default:
1066 ThrowInvalidType(value.Type());
1067 }
1068}
1069
1070U1 IREmitter::FPOrdered(const F16F32F64& lhs, const F16F32F64& rhs) {
1071 if (lhs.Type() != rhs.Type()) {
1072 throw InvalidArgument("Mismatching types {} and {}", lhs.Type(), rhs.Type());
1073 }
1074 return LogicalAnd(LogicalNot(FPIsNan(lhs)), LogicalNot(FPIsNan(rhs)));
1075}
1076
1077U1 IREmitter::FPUnordered(const F16F32F64& lhs, const F16F32F64& rhs) {
1078 if (lhs.Type() != rhs.Type()) {
1079 throw InvalidArgument("Mismatching types {} and {}", lhs.Type(), rhs.Type());
1080 }
1081 return LogicalOr(FPIsNan(lhs), FPIsNan(rhs));
1082}
1083
1084F32F64 IREmitter::FPMax(const F32F64& lhs, const F32F64& rhs, FpControl control) {
1085 if (lhs.Type() != rhs.Type()) {
1086 throw InvalidArgument("Mismatching types {} and {}", lhs.Type(), rhs.Type());
1087 }
1088 switch (lhs.Type()) {
1089 case Type::F32:
1090 return Inst<F32>(Opcode::FPMax32, Flags{control}, lhs, rhs);
1091 case Type::F64:
1092 return Inst<F64>(Opcode::FPMax64, Flags{control}, lhs, rhs);
1093 default:
1094 ThrowInvalidType(lhs.Type());
1095 }
1096}
1097
1098F32F64 IREmitter::FPMin(const F32F64& lhs, const F32F64& rhs, FpControl control) {
1099 if (lhs.Type() != rhs.Type()) {
1100 throw InvalidArgument("Mismatching types {} and {}", lhs.Type(), rhs.Type());
1101 }
1102 switch (lhs.Type()) {
1103 case Type::F32:
1104 return Inst<F32>(Opcode::FPMin32, Flags{control}, lhs, rhs);
1105 case Type::F64:
1106 return Inst<F64>(Opcode::FPMin64, Flags{control}, lhs, rhs);
1107 default:
1108 ThrowInvalidType(lhs.Type());
1109 }
1110}
1111
1112U32U64 IREmitter::IAdd(const U32U64& a, const U32U64& b) {
1113 if (a.Type() != b.Type()) {
1114 throw InvalidArgument("Mismatching types {} and {}", a.Type(), b.Type());
1115 }
1116 switch (a.Type()) {
1117 case Type::U32:
1118 return Inst<U32>(Opcode::IAdd32, a, b);
1119 case Type::U64:
1120 return Inst<U64>(Opcode::IAdd64, a, b);
1121 default:
1122 ThrowInvalidType(a.Type());
1123 }
1124}
1125
1126U32U64 IREmitter::ISub(const U32U64& a, const U32U64& b) {
1127 if (a.Type() != b.Type()) {
1128 throw InvalidArgument("Mismatching types {} and {}", a.Type(), b.Type());
1129 }
1130 switch (a.Type()) {
1131 case Type::U32:
1132 return Inst<U32>(Opcode::ISub32, a, b);
1133 case Type::U64:
1134 return Inst<U64>(Opcode::ISub64, a, b);
1135 default:
1136 ThrowInvalidType(a.Type());
1137 }
1138}
1139
1140U32 IREmitter::IMul(const U32& a, const U32& b) {
1141 return Inst<U32>(Opcode::IMul32, a, b);
1142}
1143
1144U32U64 IREmitter::INeg(const U32U64& value) {
1145 switch (value.Type()) {
1146 case Type::U32:
1147 return Inst<U32>(Opcode::INeg32, value);
1148 case Type::U64:
1149 return Inst<U64>(Opcode::INeg64, value);
1150 default:
1151 ThrowInvalidType(value.Type());
1152 }
1153}
1154
1155U32 IREmitter::IAbs(const U32& value) {
1156 return Inst<U32>(Opcode::IAbs32, value);
1157}
1158
1159U32U64 IREmitter::ShiftLeftLogical(const U32U64& base, const U32& shift) {
1160 switch (base.Type()) {
1161 case Type::U32:
1162 return Inst<U32>(Opcode::ShiftLeftLogical32, base, shift);
1163 case Type::U64:
1164 return Inst<U64>(Opcode::ShiftLeftLogical64, base, shift);
1165 default:
1166 ThrowInvalidType(base.Type());
1167 }
1168}
1169
1170U32U64 IREmitter::ShiftRightLogical(const U32U64& base, const U32& shift) {
1171 switch (base.Type()) {
1172 case Type::U32:
1173 return Inst<U32>(Opcode::ShiftRightLogical32, base, shift);
1174 case Type::U64:
1175 return Inst<U64>(Opcode::ShiftRightLogical64, base, shift);
1176 default:
1177 ThrowInvalidType(base.Type());
1178 }
1179}
1180
1181U32U64 IREmitter::ShiftRightArithmetic(const U32U64& base, const U32& shift) {
1182 switch (base.Type()) {
1183 case Type::U32:
1184 return Inst<U32>(Opcode::ShiftRightArithmetic32, base, shift);
1185 case Type::U64:
1186 return Inst<U64>(Opcode::ShiftRightArithmetic64, base, shift);
1187 default:
1188 ThrowInvalidType(base.Type());
1189 }
1190}
1191
1192U32 IREmitter::BitwiseAnd(const U32& a, const U32& b) {
1193 return Inst<U32>(Opcode::BitwiseAnd32, a, b);
1194}
1195
1196U32 IREmitter::BitwiseOr(const U32& a, const U32& b) {
1197 return Inst<U32>(Opcode::BitwiseOr32, a, b);
1198}
1199
1200U32 IREmitter::BitwiseXor(const U32& a, const U32& b) {
1201 return Inst<U32>(Opcode::BitwiseXor32, a, b);
1202}
1203
1204U32 IREmitter::BitFieldInsert(const U32& base, const U32& insert, const U32& offset,
1205 const U32& count) {
1206 return Inst<U32>(Opcode::BitFieldInsert, base, insert, offset, count);
1207}
1208
1209U32 IREmitter::BitFieldExtract(const U32& base, const U32& offset, const U32& count,
1210 bool is_signed) {
1211 return Inst<U32>(is_signed ? Opcode::BitFieldSExtract : Opcode::BitFieldUExtract, base, offset,
1212 count);
1213}
1214
1215U32 IREmitter::BitReverse(const U32& value) {
1216 return Inst<U32>(Opcode::BitReverse32, value);
1217}
1218
1219U32 IREmitter::BitCount(const U32& value) {
1220 return Inst<U32>(Opcode::BitCount32, value);
1221}
1222
1223U32 IREmitter::BitwiseNot(const U32& value) {
1224 return Inst<U32>(Opcode::BitwiseNot32, value);
1225}
1226
1227U32 IREmitter::FindSMsb(const U32& value) {
1228 return Inst<U32>(Opcode::FindSMsb32, value);
1229}
1230
1231U32 IREmitter::FindUMsb(const U32& value) {
1232 return Inst<U32>(Opcode::FindUMsb32, value);
1233}
1234
1235U32 IREmitter::SMin(const U32& a, const U32& b) {
1236 return Inst<U32>(Opcode::SMin32, a, b);
1237}
1238
1239U32 IREmitter::UMin(const U32& a, const U32& b) {
1240 return Inst<U32>(Opcode::UMin32, a, b);
1241}
1242
1243U32 IREmitter::IMin(const U32& a, const U32& b, bool is_signed) {
1244 return is_signed ? SMin(a, b) : UMin(a, b);
1245}
1246
1247U32 IREmitter::SMax(const U32& a, const U32& b) {
1248 return Inst<U32>(Opcode::SMax32, a, b);
1249}
1250
1251U32 IREmitter::UMax(const U32& a, const U32& b) {
1252 return Inst<U32>(Opcode::UMax32, a, b);
1253}
1254
1255U32 IREmitter::IMax(const U32& a, const U32& b, bool is_signed) {
1256 return is_signed ? SMax(a, b) : UMax(a, b);
1257}
1258
1259U32 IREmitter::SClamp(const U32& value, const U32& min, const U32& max) {
1260 return Inst<U32>(Opcode::SClamp32, value, min, max);
1261}
1262
1263U32 IREmitter::UClamp(const U32& value, const U32& min, const U32& max) {
1264 return Inst<U32>(Opcode::UClamp32, value, min, max);
1265}
1266
1267U1 IREmitter::ILessThan(const U32& lhs, const U32& rhs, bool is_signed) {
1268 return Inst<U1>(is_signed ? Opcode::SLessThan : Opcode::ULessThan, lhs, rhs);
1269}
1270
1271U1 IREmitter::IEqual(const U32U64& lhs, const U32U64& rhs) {
1272 if (lhs.Type() != rhs.Type()) {
1273 throw InvalidArgument("Mismatching types {} and {}", lhs.Type(), rhs.Type());
1274 }
1275 switch (lhs.Type()) {
1276 case Type::U32:
1277 return Inst<U1>(Opcode::IEqual, lhs, rhs);
1278 case Type::U64: {
1279 // Manually compare the unpacked values
1280 const Value lhs_vector{UnpackUint2x32(lhs)};
1281 const Value rhs_vector{UnpackUint2x32(rhs)};
1282 return LogicalAnd(IEqual(IR::U32{CompositeExtract(lhs_vector, 0)},
1283 IR::U32{CompositeExtract(rhs_vector, 0)}),
1284 IEqual(IR::U32{CompositeExtract(lhs_vector, 1)},
1285 IR::U32{CompositeExtract(rhs_vector, 1)}));
1286 }
1287 default:
1288 ThrowInvalidType(lhs.Type());
1289 }
1290}
1291
1292U1 IREmitter::ILessThanEqual(const U32& lhs, const U32& rhs, bool is_signed) {
1293 return Inst<U1>(is_signed ? Opcode::SLessThanEqual : Opcode::ULessThanEqual, lhs, rhs);
1294}
1295
1296U1 IREmitter::IGreaterThan(const U32& lhs, const U32& rhs, bool is_signed) {
1297 return Inst<U1>(is_signed ? Opcode::SGreaterThan : Opcode::UGreaterThan, lhs, rhs);
1298}
1299
1300U1 IREmitter::INotEqual(const U32& lhs, const U32& rhs) {
1301 return Inst<U1>(Opcode::INotEqual, lhs, rhs);
1302}
1303
1304U1 IREmitter::IGreaterThanEqual(const U32& lhs, const U32& rhs, bool is_signed) {
1305 return Inst<U1>(is_signed ? Opcode::SGreaterThanEqual : Opcode::UGreaterThanEqual, lhs, rhs);
1306}
1307
1308U32 IREmitter::SharedAtomicIAdd(const U32& pointer_offset, const U32& value) {
1309 return Inst<U32>(Opcode::SharedAtomicIAdd32, pointer_offset, value);
1310}
1311
1312U32 IREmitter::SharedAtomicSMin(const U32& pointer_offset, const U32& value) {
1313 return Inst<U32>(Opcode::SharedAtomicSMin32, pointer_offset, value);
1314}
1315
1316U32 IREmitter::SharedAtomicUMin(const U32& pointer_offset, const U32& value) {
1317 return Inst<U32>(Opcode::SharedAtomicUMin32, pointer_offset, value);
1318}
1319
1320U32 IREmitter::SharedAtomicIMin(const U32& pointer_offset, const U32& value, bool is_signed) {
1321 return is_signed ? SharedAtomicSMin(pointer_offset, value)
1322 : SharedAtomicUMin(pointer_offset, value);
1323}
1324
1325U32 IREmitter::SharedAtomicSMax(const U32& pointer_offset, const U32& value) {
1326 return Inst<U32>(Opcode::SharedAtomicSMax32, pointer_offset, value);
1327}
1328
1329U32 IREmitter::SharedAtomicUMax(const U32& pointer_offset, const U32& value) {
1330 return Inst<U32>(Opcode::SharedAtomicUMax32, pointer_offset, value);
1331}
1332
1333U32 IREmitter::SharedAtomicIMax(const U32& pointer_offset, const U32& value, bool is_signed) {
1334 return is_signed ? SharedAtomicSMax(pointer_offset, value)
1335 : SharedAtomicUMax(pointer_offset, value);
1336}
1337
1338U32 IREmitter::SharedAtomicInc(const U32& pointer_offset, const U32& value) {
1339 return Inst<U32>(Opcode::SharedAtomicInc32, pointer_offset, value);
1340}
1341
1342U32 IREmitter::SharedAtomicDec(const U32& pointer_offset, const U32& value) {
1343 return Inst<U32>(Opcode::SharedAtomicDec32, pointer_offset, value);
1344}
1345
1346U32 IREmitter::SharedAtomicAnd(const U32& pointer_offset, const U32& value) {
1347 return Inst<U32>(Opcode::SharedAtomicAnd32, pointer_offset, value);
1348}
1349
1350U32 IREmitter::SharedAtomicOr(const U32& pointer_offset, const U32& value) {
1351 return Inst<U32>(Opcode::SharedAtomicOr32, pointer_offset, value);
1352}
1353
1354U32 IREmitter::SharedAtomicXor(const U32& pointer_offset, const U32& value) {
1355 return Inst<U32>(Opcode::SharedAtomicXor32, pointer_offset, value);
1356}
1357
1358U32U64 IREmitter::SharedAtomicExchange(const U32& pointer_offset, const U32U64& value) {
1359 switch (value.Type()) {
1360 case Type::U32:
1361 return Inst<U32>(Opcode::SharedAtomicExchange32, pointer_offset, value);
1362 case Type::U64:
1363 return Inst<U64>(Opcode::SharedAtomicExchange64, pointer_offset, value);
1364 default:
1365 ThrowInvalidType(pointer_offset.Type());
1366 }
1367}
1368
1369U32U64 IREmitter::GlobalAtomicIAdd(const U64& pointer_offset, const U32U64& value) {
1370 switch (value.Type()) {
1371 case Type::U32:
1372 return Inst<U32>(Opcode::GlobalAtomicIAdd32, pointer_offset, value);
1373 case Type::U64:
1374 return Inst<U64>(Opcode::GlobalAtomicIAdd64, pointer_offset, value);
1375 default:
1376 ThrowInvalidType(value.Type());
1377 }
1378}
1379
1380U32U64 IREmitter::GlobalAtomicSMin(const U64& pointer_offset, const U32U64& value) {
1381 switch (value.Type()) {
1382 case Type::U32:
1383 return Inst<U32>(Opcode::GlobalAtomicSMin32, pointer_offset, value);
1384 case Type::U64:
1385 return Inst<U64>(Opcode::GlobalAtomicSMin64, pointer_offset, value);
1386 default:
1387 ThrowInvalidType(value.Type());
1388 }
1389}
1390
1391U32U64 IREmitter::GlobalAtomicUMin(const U64& pointer_offset, const U32U64& value) {
1392 switch (value.Type()) {
1393 case Type::U32:
1394 return Inst<U32>(Opcode::GlobalAtomicUMin32, pointer_offset, value);
1395 case Type::U64:
1396 return Inst<U64>(Opcode::GlobalAtomicUMin64, pointer_offset, value);
1397 default:
1398 ThrowInvalidType(value.Type());
1399 }
1400}
1401
1402U32U64 IREmitter::GlobalAtomicIMin(const U64& pointer_offset, const U32U64& value, bool is_signed) {
1403 return is_signed ? GlobalAtomicSMin(pointer_offset, value)
1404 : GlobalAtomicUMin(pointer_offset, value);
1405}
1406
1407U32U64 IREmitter::GlobalAtomicSMax(const U64& pointer_offset, const U32U64& value) {
1408 switch (value.Type()) {
1409 case Type::U32:
1410 return Inst<U32>(Opcode::GlobalAtomicSMax32, pointer_offset, value);
1411 case Type::U64:
1412 return Inst<U64>(Opcode::GlobalAtomicSMax64, pointer_offset, value);
1413 default:
1414 ThrowInvalidType(value.Type());
1415 }
1416}
1417
1418U32U64 IREmitter::GlobalAtomicUMax(const U64& pointer_offset, const U32U64& value) {
1419 switch (value.Type()) {
1420 case Type::U32:
1421 return Inst<U32>(Opcode::GlobalAtomicUMax32, pointer_offset, value);
1422 case Type::U64:
1423 return Inst<U64>(Opcode::GlobalAtomicUMax64, pointer_offset, value);
1424 default:
1425 ThrowInvalidType(value.Type());
1426 }
1427}
1428
1429U32U64 IREmitter::GlobalAtomicIMax(const U64& pointer_offset, const U32U64& value, bool is_signed) {
1430 return is_signed ? GlobalAtomicSMax(pointer_offset, value)
1431 : GlobalAtomicUMax(pointer_offset, value);
1432}
1433
1434U32 IREmitter::GlobalAtomicInc(const U64& pointer_offset, const U32& value) {
1435 return Inst<U32>(Opcode::GlobalAtomicInc32, pointer_offset, value);
1436}
1437
1438U32 IREmitter::GlobalAtomicDec(const U64& pointer_offset, const U32& value) {
1439 return Inst<U32>(Opcode::GlobalAtomicDec32, pointer_offset, value);
1440}
1441
1442U32U64 IREmitter::GlobalAtomicAnd(const U64& pointer_offset, const U32U64& value) {
1443 switch (value.Type()) {
1444 case Type::U32:
1445 return Inst<U32>(Opcode::GlobalAtomicAnd32, pointer_offset, value);
1446 case Type::U64:
1447 return Inst<U64>(Opcode::GlobalAtomicAnd64, pointer_offset, value);
1448 default:
1449 ThrowInvalidType(value.Type());
1450 }
1451}
1452
1453U32U64 IREmitter::GlobalAtomicOr(const U64& pointer_offset, const U32U64& value) {
1454 switch (value.Type()) {
1455 case Type::U32:
1456 return Inst<U32>(Opcode::GlobalAtomicOr32, pointer_offset, value);
1457 case Type::U64:
1458 return Inst<U64>(Opcode::GlobalAtomicOr64, pointer_offset, value);
1459 default:
1460 ThrowInvalidType(value.Type());
1461 }
1462}
1463
1464U32U64 IREmitter::GlobalAtomicXor(const U64& pointer_offset, const U32U64& value) {
1465 switch (value.Type()) {
1466 case Type::U32:
1467 return Inst<U32>(Opcode::GlobalAtomicXor32, pointer_offset, value);
1468 case Type::U64:
1469 return Inst<U64>(Opcode::GlobalAtomicXor64, pointer_offset, value);
1470 default:
1471 ThrowInvalidType(value.Type());
1472 }
1473}
1474
1475U32U64 IREmitter::GlobalAtomicExchange(const U64& pointer_offset, const U32U64& value) {
1476 switch (value.Type()) {
1477 case Type::U32:
1478 return Inst<U32>(Opcode::GlobalAtomicExchange32, pointer_offset, value);
1479 case Type::U64:
1480 return Inst<U64>(Opcode::GlobalAtomicExchange64, pointer_offset, value);
1481 default:
1482 ThrowInvalidType(pointer_offset.Type());
1483 }
1484}
1485
1486F32 IREmitter::GlobalAtomicF32Add(const U64& pointer_offset, const Value& value,
1487 const FpControl control) {
1488 return Inst<F32>(Opcode::GlobalAtomicAddF32, Flags{control}, pointer_offset, value);
1489}
1490
1491Value IREmitter::GlobalAtomicF16x2Add(const U64& pointer_offset, const Value& value,
1492 const FpControl control) {
1493 return Inst(Opcode::GlobalAtomicAddF16x2, Flags{control}, pointer_offset, value);
1494}
1495
1496Value IREmitter::GlobalAtomicF16x2Min(const U64& pointer_offset, const Value& value,
1497 const FpControl control) {
1498 return Inst(Opcode::GlobalAtomicMinF16x2, Flags{control}, pointer_offset, value);
1499}
1500
1501Value IREmitter::GlobalAtomicF16x2Max(const U64& pointer_offset, const Value& value,
1502 const FpControl control) {
1503 return Inst(Opcode::GlobalAtomicMaxF16x2, Flags{control}, pointer_offset, value);
1504}
1505
1506U1 IREmitter::LogicalOr(const U1& a, const U1& b) {
1507 return Inst<U1>(Opcode::LogicalOr, a, b);
1508}
1509
1510U1 IREmitter::LogicalAnd(const U1& a, const U1& b) {
1511 return Inst<U1>(Opcode::LogicalAnd, a, b);
1512}
1513
1514U1 IREmitter::LogicalXor(const U1& a, const U1& b) {
1515 return Inst<U1>(Opcode::LogicalXor, a, b);
1516}
1517
1518U1 IREmitter::LogicalNot(const U1& value) {
1519 return Inst<U1>(Opcode::LogicalNot, value);
1520}
1521
1522U32U64 IREmitter::ConvertFToS(size_t bitsize, const F16F32F64& value) {
1523 switch (bitsize) {
1524 case 16:
1525 switch (value.Type()) {
1526 case Type::F16:
1527 return Inst<U32>(Opcode::ConvertS16F16, value);
1528 case Type::F32:
1529 return Inst<U32>(Opcode::ConvertS16F32, value);
1530 case Type::F64:
1531 return Inst<U32>(Opcode::ConvertS16F64, value);
1532 default:
1533 ThrowInvalidType(value.Type());
1534 }
1535 case 32:
1536 switch (value.Type()) {
1537 case Type::F16:
1538 return Inst<U32>(Opcode::ConvertS32F16, value);
1539 case Type::F32:
1540 return Inst<U32>(Opcode::ConvertS32F32, value);
1541 case Type::F64:
1542 return Inst<U32>(Opcode::ConvertS32F64, value);
1543 default:
1544 ThrowInvalidType(value.Type());
1545 }
1546 case 64:
1547 switch (value.Type()) {
1548 case Type::F16:
1549 return Inst<U64>(Opcode::ConvertS64F16, value);
1550 case Type::F32:
1551 return Inst<U64>(Opcode::ConvertS64F32, value);
1552 case Type::F64:
1553 return Inst<U64>(Opcode::ConvertS64F64, value);
1554 default:
1555 ThrowInvalidType(value.Type());
1556 }
1557 default:
1558 throw InvalidArgument("Invalid destination bitsize {}", bitsize);
1559 }
1560}
1561
1562U32U64 IREmitter::ConvertFToU(size_t bitsize, const F16F32F64& value) {
1563 switch (bitsize) {
1564 case 16:
1565 switch (value.Type()) {
1566 case Type::F16:
1567 return Inst<U32>(Opcode::ConvertU16F16, value);
1568 case Type::F32:
1569 return Inst<U32>(Opcode::ConvertU16F32, value);
1570 case Type::F64:
1571 return Inst<U32>(Opcode::ConvertU16F64, value);
1572 default:
1573 ThrowInvalidType(value.Type());
1574 }
1575 case 32:
1576 switch (value.Type()) {
1577 case Type::F16:
1578 return Inst<U32>(Opcode::ConvertU32F16, value);
1579 case Type::F32:
1580 return Inst<U32>(Opcode::ConvertU32F32, value);
1581 case Type::F64:
1582 return Inst<U32>(Opcode::ConvertU32F64, value);
1583 default:
1584 ThrowInvalidType(value.Type());
1585 }
1586 case 64:
1587 switch (value.Type()) {
1588 case Type::F16:
1589 return Inst<U64>(Opcode::ConvertU64F16, value);
1590 case Type::F32:
1591 return Inst<U64>(Opcode::ConvertU64F32, value);
1592 case Type::F64:
1593 return Inst<U64>(Opcode::ConvertU64F64, value);
1594 default:
1595 ThrowInvalidType(value.Type());
1596 }
1597 default:
1598 throw InvalidArgument("Invalid destination bitsize {}", bitsize);
1599 }
1600}
1601
1602U32U64 IREmitter::ConvertFToI(size_t bitsize, bool is_signed, const F16F32F64& value) {
1603 return is_signed ? ConvertFToS(bitsize, value) : ConvertFToU(bitsize, value);
1604}
1605
1606F16F32F64 IREmitter::ConvertSToF(size_t dest_bitsize, size_t src_bitsize, const Value& value,
1607 FpControl control) {
1608 switch (dest_bitsize) {
1609 case 16:
1610 switch (src_bitsize) {
1611 case 8:
1612 return Inst<F16>(Opcode::ConvertF16S8, Flags{control}, value);
1613 case 16:
1614 return Inst<F16>(Opcode::ConvertF16S16, Flags{control}, value);
1615 case 32:
1616 return Inst<F16>(Opcode::ConvertF16S32, Flags{control}, value);
1617 case 64:
1618 return Inst<F16>(Opcode::ConvertF16S64, Flags{control}, value);
1619 }
1620 break;
1621 case 32:
1622 switch (src_bitsize) {
1623 case 8:
1624 return Inst<F32>(Opcode::ConvertF32S8, Flags{control}, value);
1625 case 16:
1626 return Inst<F32>(Opcode::ConvertF32S16, Flags{control}, value);
1627 case 32:
1628 return Inst<F32>(Opcode::ConvertF32S32, Flags{control}, value);
1629 case 64:
1630 return Inst<F32>(Opcode::ConvertF32S64, Flags{control}, value);
1631 }
1632 break;
1633 case 64:
1634 switch (src_bitsize) {
1635 case 8:
1636 return Inst<F64>(Opcode::ConvertF64S8, Flags{control}, value);
1637 case 16:
1638 return Inst<F64>(Opcode::ConvertF64S16, Flags{control}, value);
1639 case 32:
1640 return Inst<F64>(Opcode::ConvertF64S32, Flags{control}, value);
1641 case 64:
1642 return Inst<F64>(Opcode::ConvertF64S64, Flags{control}, value);
1643 }
1644 break;
1645 }
1646 throw InvalidArgument("Invalid bit size combination dst={} src={}", dest_bitsize, src_bitsize);
1647}
1648
1649F16F32F64 IREmitter::ConvertUToF(size_t dest_bitsize, size_t src_bitsize, const Value& value,
1650 FpControl control) {
1651 switch (dest_bitsize) {
1652 case 16:
1653 switch (src_bitsize) {
1654 case 8:
1655 return Inst<F16>(Opcode::ConvertF16U8, Flags{control}, value);
1656 case 16:
1657 return Inst<F16>(Opcode::ConvertF16U16, Flags{control}, value);
1658 case 32:
1659 return Inst<F16>(Opcode::ConvertF16U32, Flags{control}, value);
1660 case 64:
1661 return Inst<F16>(Opcode::ConvertF16U64, Flags{control}, value);
1662 }
1663 break;
1664 case 32:
1665 switch (src_bitsize) {
1666 case 8:
1667 return Inst<F32>(Opcode::ConvertF32U8, Flags{control}, value);
1668 case 16:
1669 return Inst<F32>(Opcode::ConvertF32U16, Flags{control}, value);
1670 case 32:
1671 return Inst<F32>(Opcode::ConvertF32U32, Flags{control}, value);
1672 case 64:
1673 return Inst<F32>(Opcode::ConvertF32U64, Flags{control}, value);
1674 }
1675 break;
1676 case 64:
1677 switch (src_bitsize) {
1678 case 8:
1679 return Inst<F64>(Opcode::ConvertF64U8, Flags{control}, value);
1680 case 16:
1681 return Inst<F64>(Opcode::ConvertF64U16, Flags{control}, value);
1682 case 32:
1683 return Inst<F64>(Opcode::ConvertF64U32, Flags{control}, value);
1684 case 64:
1685 return Inst<F64>(Opcode::ConvertF64U64, Flags{control}, value);
1686 }
1687 break;
1688 }
1689 throw InvalidArgument("Invalid bit size combination dst={} src={}", dest_bitsize, src_bitsize);
1690}
1691
1692F16F32F64 IREmitter::ConvertIToF(size_t dest_bitsize, size_t src_bitsize, bool is_signed,
1693 const Value& value, FpControl control) {
1694 return is_signed ? ConvertSToF(dest_bitsize, src_bitsize, value, control)
1695 : ConvertUToF(dest_bitsize, src_bitsize, value, control);
1696}
1697
1698U32U64 IREmitter::UConvert(size_t result_bitsize, const U32U64& value) {
1699 switch (result_bitsize) {
1700 case 32:
1701 switch (value.Type()) {
1702 case Type::U32:
1703 // Nothing to do
1704 return value;
1705 case Type::U64:
1706 return Inst<U32>(Opcode::ConvertU32U64, value);
1707 default:
1708 break;
1709 }
1710 break;
1711 case 64:
1712 switch (value.Type()) {
1713 case Type::U32:
1714 return Inst<U64>(Opcode::ConvertU64U32, value);
1715 case Type::U64:
1716 // Nothing to do
1717 return value;
1718 default:
1719 break;
1720 }
1721 }
1722 throw NotImplementedException("Conversion from {} to {} bits", value.Type(), result_bitsize);
1723}
1724
1725F16F32F64 IREmitter::FPConvert(size_t result_bitsize, const F16F32F64& value, FpControl control) {
1726 switch (result_bitsize) {
1727 case 16:
1728 switch (value.Type()) {
1729 case Type::F16:
1730 // Nothing to do
1731 return value;
1732 case Type::F32:
1733 return Inst<F16>(Opcode::ConvertF16F32, Flags{control}, value);
1734 case Type::F64:
1735 throw LogicError("Illegal conversion from F64 to F16");
1736 default:
1737 break;
1738 }
1739 break;
1740 case 32:
1741 switch (value.Type()) {
1742 case Type::F16:
1743 return Inst<F32>(Opcode::ConvertF32F16, Flags{control}, value);
1744 case Type::F32:
1745 // Nothing to do
1746 return value;
1747 case Type::F64:
1748 return Inst<F32>(Opcode::ConvertF32F64, Flags{control}, value);
1749 default:
1750 break;
1751 }
1752 break;
1753 case 64:
1754 switch (value.Type()) {
1755 case Type::F16:
1756 throw LogicError("Illegal conversion from F16 to F64");
1757 case Type::F32:
1758 return Inst<F64>(Opcode::ConvertF64F32, Flags{control}, value);
1759 case Type::F64:
1760 // Nothing to do
1761 return value;
1762 default:
1763 break;
1764 }
1765 break;
1766 }
1767 throw NotImplementedException("Conversion from {} to {} bits", value.Type(), result_bitsize);
1768}
1769
1770Value IREmitter::ImageSampleImplicitLod(const Value& handle, const Value& coords, const F32& bias,
1771 const Value& offset, const F32& lod_clamp,
1772 TextureInstInfo info) {
1773 const Value bias_lc{MakeLodClampPair(*this, bias, lod_clamp)};
1774 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageSampleImplicitLod
1775 : Opcode::BindlessImageSampleImplicitLod};
1776 return Inst(op, Flags{info}, handle, coords, bias_lc, offset);
1777}
1778
1779Value IREmitter::ImageSampleExplicitLod(const Value& handle, const Value& coords, const F32& lod,
1780 const Value& offset, TextureInstInfo info) {
1781 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageSampleExplicitLod
1782 : Opcode::BindlessImageSampleExplicitLod};
1783 return Inst(op, Flags{info}, handle, coords, lod, offset);
1784}
1785
1786F32 IREmitter::ImageSampleDrefImplicitLod(const Value& handle, const Value& coords, const F32& dref,
1787 const F32& bias, const Value& offset,
1788 const F32& lod_clamp, TextureInstInfo info) {
1789 const Value bias_lc{MakeLodClampPair(*this, bias, lod_clamp)};
1790 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageSampleDrefImplicitLod
1791 : Opcode::BindlessImageSampleDrefImplicitLod};
1792 return Inst<F32>(op, Flags{info}, handle, coords, dref, bias_lc, offset);
1793}
1794
1795F32 IREmitter::ImageSampleDrefExplicitLod(const Value& handle, const Value& coords, const F32& dref,
1796 const F32& lod, const Value& offset,
1797 TextureInstInfo info) {
1798 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageSampleDrefExplicitLod
1799 : Opcode::BindlessImageSampleDrefExplicitLod};
1800 return Inst<F32>(op, Flags{info}, handle, coords, dref, lod, offset);
1801}
1802
1803Value IREmitter::ImageGather(const Value& handle, const Value& coords, const Value& offset,
1804 const Value& offset2, TextureInstInfo info) {
1805 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageGather : Opcode::BindlessImageGather};
1806 return Inst(op, Flags{info}, handle, coords, offset, offset2);
1807}
1808
1809Value IREmitter::ImageGatherDref(const Value& handle, const Value& coords, const Value& offset,
1810 const Value& offset2, const F32& dref, TextureInstInfo info) {
1811 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageGatherDref
1812 : Opcode::BindlessImageGatherDref};
1813 return Inst(op, Flags{info}, handle, coords, offset, offset2, dref);
1814}
1815
1816Value IREmitter::ImageFetch(const Value& handle, const Value& coords, const Value& offset,
1817 const U32& lod, const U32& multisampling, TextureInstInfo info) {
1818 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageFetch : Opcode::BindlessImageFetch};
1819 return Inst(op, Flags{info}, handle, coords, offset, lod, multisampling);
1820}
1821
1822Value IREmitter::ImageQueryDimension(const Value& handle, const IR::U32& lod) {
1823 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageQueryDimensions
1824 : Opcode::BindlessImageQueryDimensions};
1825 return Inst(op, handle, lod);
1826}
1827
1828Value IREmitter::ImageQueryLod(const Value& handle, const Value& coords, TextureInstInfo info) {
1829 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageQueryLod
1830 : Opcode::BindlessImageQueryLod};
1831 return Inst(op, Flags{info}, handle, coords);
1832}
1833
1834Value IREmitter::ImageGradient(const Value& handle, const Value& coords, const Value& derivates,
1835 const Value& offset, const F32& lod_clamp, TextureInstInfo info) {
1836 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageGradient
1837 : Opcode::BindlessImageGradient};
1838 return Inst(op, Flags{info}, handle, coords, derivates, offset, lod_clamp);
1839}
1840
1841Value IREmitter::ImageRead(const Value& handle, const Value& coords, TextureInstInfo info) {
1842 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageRead : Opcode::BindlessImageRead};
1843 return Inst(op, Flags{info}, handle, coords);
1844}
1845
1846void IREmitter::ImageWrite(const Value& handle, const Value& coords, const Value& color,
1847 TextureInstInfo info) {
1848 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageWrite : Opcode::BindlessImageWrite};
1849 Inst(op, Flags{info}, handle, coords, color);
1850}
1851
1852Value IREmitter::ImageAtomicIAdd(const Value& handle, const Value& coords, const Value& value,
1853 TextureInstInfo info) {
1854 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageAtomicIAdd32
1855 : Opcode::BindlessImageAtomicIAdd32};
1856 return Inst(op, Flags{info}, handle, coords, value);
1857}
1858
1859Value IREmitter::ImageAtomicSMin(const Value& handle, const Value& coords, const Value& value,
1860 TextureInstInfo info) {
1861 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageAtomicSMin32
1862 : Opcode::BindlessImageAtomicSMin32};
1863 return Inst(op, Flags{info}, handle, coords, value);
1864}
1865
1866Value IREmitter::ImageAtomicUMin(const Value& handle, const Value& coords, const Value& value,
1867 TextureInstInfo info) {
1868 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageAtomicUMin32
1869 : Opcode::BindlessImageAtomicUMin32};
1870 return Inst(op, Flags{info}, handle, coords, value);
1871}
1872
1873Value IREmitter::ImageAtomicIMin(const Value& handle, const Value& coords, const Value& value,
1874 bool is_signed, TextureInstInfo info) {
1875 return is_signed ? ImageAtomicSMin(handle, coords, value, info)
1876 : ImageAtomicUMin(handle, coords, value, info);
1877}
1878
1879Value IREmitter::ImageAtomicSMax(const Value& handle, const Value& coords, const Value& value,
1880 TextureInstInfo info) {
1881 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageAtomicSMax32
1882 : Opcode::BindlessImageAtomicSMax32};
1883 return Inst(op, Flags{info}, handle, coords, value);
1884}
1885
1886Value IREmitter::ImageAtomicUMax(const Value& handle, const Value& coords, const Value& value,
1887 TextureInstInfo info) {
1888 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageAtomicUMax32
1889 : Opcode::BindlessImageAtomicUMax32};
1890 return Inst(op, Flags{info}, handle, coords, value);
1891}
1892
1893Value IREmitter::ImageAtomicIMax(const Value& handle, const Value& coords, const Value& value,
1894 bool is_signed, TextureInstInfo info) {
1895 return is_signed ? ImageAtomicSMax(handle, coords, value, info)
1896 : ImageAtomicUMax(handle, coords, value, info);
1897}
1898
1899Value IREmitter::ImageAtomicInc(const Value& handle, const Value& coords, const Value& value,
1900 TextureInstInfo info) {
1901 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageAtomicInc32
1902 : Opcode::BindlessImageAtomicInc32};
1903 return Inst(op, Flags{info}, handle, coords, value);
1904}
1905
1906Value IREmitter::ImageAtomicDec(const Value& handle, const Value& coords, const Value& value,
1907 TextureInstInfo info) {
1908 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageAtomicDec32
1909 : Opcode::BindlessImageAtomicDec32};
1910 return Inst(op, Flags{info}, handle, coords, value);
1911}
1912
1913Value IREmitter::ImageAtomicAnd(const Value& handle, const Value& coords, const Value& value,
1914 TextureInstInfo info) {
1915 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageAtomicAnd32
1916 : Opcode::BindlessImageAtomicAnd32};
1917 return Inst(op, Flags{info}, handle, coords, value);
1918}
1919
1920Value IREmitter::ImageAtomicOr(const Value& handle, const Value& coords, const Value& value,
1921 TextureInstInfo info) {
1922 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageAtomicOr32
1923 : Opcode::BindlessImageAtomicOr32};
1924 return Inst(op, Flags{info}, handle, coords, value);
1925}
1926
1927Value IREmitter::ImageAtomicXor(const Value& handle, const Value& coords, const Value& value,
1928 TextureInstInfo info) {
1929 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageAtomicXor32
1930 : Opcode::BindlessImageAtomicXor32};
1931 return Inst(op, Flags{info}, handle, coords, value);
1932}
1933
1934Value IREmitter::ImageAtomicExchange(const Value& handle, const Value& coords, const Value& value,
1935 TextureInstInfo info) {
1936 const Opcode op{handle.IsImmediate() ? Opcode::BoundImageAtomicExchange32
1937 : Opcode::BindlessImageAtomicExchange32};
1938 return Inst(op, Flags{info}, handle, coords, value);
1939}
1940
1941U1 IREmitter::VoteAll(const U1& value) {
1942 return Inst<U1>(Opcode::VoteAll, value);
1943}
1944
1945U1 IREmitter::VoteAny(const U1& value) {
1946 return Inst<U1>(Opcode::VoteAny, value);
1947}
1948
1949U1 IREmitter::VoteEqual(const U1& value) {
1950 return Inst<U1>(Opcode::VoteEqual, value);
1951}
1952
1953U32 IREmitter::SubgroupBallot(const U1& value) {
1954 return Inst<U32>(Opcode::SubgroupBallot, value);
1955}
1956
1957U32 IREmitter::SubgroupEqMask() {
1958 return Inst<U32>(Opcode::SubgroupEqMask);
1959}
1960
1961U32 IREmitter::SubgroupLtMask() {
1962 return Inst<U32>(Opcode::SubgroupLtMask);
1963}
1964
1965U32 IREmitter::SubgroupLeMask() {
1966 return Inst<U32>(Opcode::SubgroupLeMask);
1967}
1968
1969U32 IREmitter::SubgroupGtMask() {
1970 return Inst<U32>(Opcode::SubgroupGtMask);
1971}
1972
1973U32 IREmitter::SubgroupGeMask() {
1974 return Inst<U32>(Opcode::SubgroupGeMask);
1975}
1976
1977U32 IREmitter::ShuffleIndex(const IR::U32& value, const IR::U32& index, const IR::U32& clamp,
1978 const IR::U32& seg_mask) {
1979 return Inst<U32>(Opcode::ShuffleIndex, value, index, clamp, seg_mask);
1980}
1981
1982U32 IREmitter::ShuffleUp(const IR::U32& value, const IR::U32& index, const IR::U32& clamp,
1983 const IR::U32& seg_mask) {
1984 return Inst<U32>(Opcode::ShuffleUp, value, index, clamp, seg_mask);
1985}
1986
1987U32 IREmitter::ShuffleDown(const IR::U32& value, const IR::U32& index, const IR::U32& clamp,
1988 const IR::U32& seg_mask) {
1989 return Inst<U32>(Opcode::ShuffleDown, value, index, clamp, seg_mask);
1990}
1991
1992U32 IREmitter::ShuffleButterfly(const IR::U32& value, const IR::U32& index, const IR::U32& clamp,
1993 const IR::U32& seg_mask) {
1994 return Inst<U32>(Opcode::ShuffleButterfly, value, index, clamp, seg_mask);
1995}
1996
1997F32 IREmitter::FSwizzleAdd(const F32& a, const F32& b, const U32& swizzle, FpControl control) {
1998 return Inst<F32>(Opcode::FSwizzleAdd, Flags{control}, a, b, swizzle);
1999}
2000
2001F32 IREmitter::DPdxFine(const F32& a) {
2002 return Inst<F32>(Opcode::DPdxFine, a);
2003}
2004
2005F32 IREmitter::DPdyFine(const F32& a) {
2006 return Inst<F32>(Opcode::DPdyFine, a);
2007}
2008
2009F32 IREmitter::DPdxCoarse(const F32& a) {
2010 return Inst<F32>(Opcode::DPdxCoarse, a);
2011}
2012
2013F32 IREmitter::DPdyCoarse(const F32& a) {
2014 return Inst<F32>(Opcode::DPdyCoarse, a);
2015}
2016
2017} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/ir_emitter.h b/src/shader_recompiler/frontend/ir/ir_emitter.h
new file mode 100644
index 000000000..53f7b3b06
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/ir_emitter.h
@@ -0,0 +1,413 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <cstring>
8#include <type_traits>
9
10#include "shader_recompiler/frontend/ir/attribute.h"
11#include "shader_recompiler/frontend/ir/basic_block.h"
12#include "shader_recompiler/frontend/ir/modifiers.h"
13#include "shader_recompiler/frontend/ir/value.h"
14
15namespace Shader::IR {
16
17class IREmitter {
18public:
19 explicit IREmitter(Block& block_) : block{&block_}, insertion_point{block->end()} {}
20 explicit IREmitter(Block& block_, Block::iterator insertion_point_)
21 : block{&block_}, insertion_point{insertion_point_} {}
22
23 Block* block;
24
25 [[nodiscard]] U1 Imm1(bool value) const;
26 [[nodiscard]] U8 Imm8(u8 value) const;
27 [[nodiscard]] U16 Imm16(u16 value) const;
28 [[nodiscard]] U32 Imm32(u32 value) const;
29 [[nodiscard]] U32 Imm32(s32 value) const;
30 [[nodiscard]] F32 Imm32(f32 value) const;
31 [[nodiscard]] U64 Imm64(u64 value) const;
32 [[nodiscard]] U64 Imm64(s64 value) const;
33 [[nodiscard]] F64 Imm64(f64 value) const;
34
35 U1 ConditionRef(const U1& value);
36 void Reference(const Value& value);
37
38 void PhiMove(IR::Inst& phi, const Value& value);
39
40 void Prologue();
41 void Epilogue();
42 void DemoteToHelperInvocation();
43 void EmitVertex(const U32& stream);
44 void EndPrimitive(const U32& stream);
45
46 [[nodiscard]] U32 GetReg(IR::Reg reg);
47 void SetReg(IR::Reg reg, const U32& value);
48
49 [[nodiscard]] U1 GetPred(IR::Pred pred, bool is_negated = false);
50 void SetPred(IR::Pred pred, const U1& value);
51
52 [[nodiscard]] U1 GetGotoVariable(u32 id);
53 void SetGotoVariable(u32 id, const U1& value);
54
55 [[nodiscard]] U32 GetIndirectBranchVariable();
56 void SetIndirectBranchVariable(const U32& value);
57
58 [[nodiscard]] U32 GetCbuf(const U32& binding, const U32& byte_offset);
59 [[nodiscard]] Value GetCbuf(const U32& binding, const U32& byte_offset, size_t bitsize,
60 bool is_signed);
61 [[nodiscard]] F32 GetFloatCbuf(const U32& binding, const U32& byte_offset);
62
63 [[nodiscard]] U1 GetZFlag();
64 [[nodiscard]] U1 GetSFlag();
65 [[nodiscard]] U1 GetCFlag();
66 [[nodiscard]] U1 GetOFlag();
67
68 void SetZFlag(const U1& value);
69 void SetSFlag(const U1& value);
70 void SetCFlag(const U1& value);
71 void SetOFlag(const U1& value);
72
73 [[nodiscard]] U1 Condition(IR::Condition cond);
74 [[nodiscard]] U1 GetFlowTestResult(FlowTest test);
75
76 [[nodiscard]] F32 GetAttribute(IR::Attribute attribute);
77 [[nodiscard]] F32 GetAttribute(IR::Attribute attribute, const U32& vertex);
78 void SetAttribute(IR::Attribute attribute, const F32& value, const U32& vertex);
79
80 [[nodiscard]] F32 GetAttributeIndexed(const U32& phys_address);
81 [[nodiscard]] F32 GetAttributeIndexed(const U32& phys_address, const U32& vertex);
82 void SetAttributeIndexed(const U32& phys_address, const F32& value, const U32& vertex);
83
84 [[nodiscard]] F32 GetPatch(Patch patch);
85 void SetPatch(Patch patch, const F32& value);
86
87 void SetFragColor(u32 index, u32 component, const F32& value);
88 void SetSampleMask(const U32& value);
89 void SetFragDepth(const F32& value);
90
91 [[nodiscard]] U32 WorkgroupIdX();
92 [[nodiscard]] U32 WorkgroupIdY();
93 [[nodiscard]] U32 WorkgroupIdZ();
94
95 [[nodiscard]] Value LocalInvocationId();
96 [[nodiscard]] U32 LocalInvocationIdX();
97 [[nodiscard]] U32 LocalInvocationIdY();
98 [[nodiscard]] U32 LocalInvocationIdZ();
99
100 [[nodiscard]] U32 InvocationId();
101 [[nodiscard]] U32 SampleId();
102 [[nodiscard]] U1 IsHelperInvocation();
103 [[nodiscard]] F32 YDirection();
104
105 [[nodiscard]] U32 LaneId();
106
107 [[nodiscard]] U32 LoadGlobalU8(const U64& address);
108 [[nodiscard]] U32 LoadGlobalS8(const U64& address);
109 [[nodiscard]] U32 LoadGlobalU16(const U64& address);
110 [[nodiscard]] U32 LoadGlobalS16(const U64& address);
111 [[nodiscard]] U32 LoadGlobal32(const U64& address);
112 [[nodiscard]] Value LoadGlobal64(const U64& address);
113 [[nodiscard]] Value LoadGlobal128(const U64& address);
114
115 void WriteGlobalU8(const U64& address, const U32& value);
116 void WriteGlobalS8(const U64& address, const U32& value);
117 void WriteGlobalU16(const U64& address, const U32& value);
118 void WriteGlobalS16(const U64& address, const U32& value);
119 void WriteGlobal32(const U64& address, const U32& value);
120 void WriteGlobal64(const U64& address, const IR::Value& vector);
121 void WriteGlobal128(const U64& address, const IR::Value& vector);
122
123 [[nodiscard]] U32 LoadLocal(const U32& word_offset);
124 void WriteLocal(const U32& word_offset, const U32& value);
125
126 [[nodiscard]] Value LoadShared(int bit_size, bool is_signed, const U32& offset);
127 void WriteShared(int bit_size, const U32& offset, const Value& value);
128
129 [[nodiscard]] U1 GetZeroFromOp(const Value& op);
130 [[nodiscard]] U1 GetSignFromOp(const Value& op);
131 [[nodiscard]] U1 GetCarryFromOp(const Value& op);
132 [[nodiscard]] U1 GetOverflowFromOp(const Value& op);
133 [[nodiscard]] U1 GetSparseFromOp(const Value& op);
134 [[nodiscard]] U1 GetInBoundsFromOp(const Value& op);
135
136 [[nodiscard]] Value CompositeConstruct(const Value& e1, const Value& e2);
137 [[nodiscard]] Value CompositeConstruct(const Value& e1, const Value& e2, const Value& e3);
138 [[nodiscard]] Value CompositeConstruct(const Value& e1, const Value& e2, const Value& e3,
139 const Value& e4);
140 [[nodiscard]] Value CompositeExtract(const Value& vector, size_t element);
141 [[nodiscard]] Value CompositeInsert(const Value& vector, const Value& object, size_t element);
142
143 [[nodiscard]] Value Select(const U1& condition, const Value& true_value,
144 const Value& false_value);
145
146 void Barrier();
147 void WorkgroupMemoryBarrier();
148 void DeviceMemoryBarrier();
149
150 template <typename Dest, typename Source>
151 [[nodiscard]] Dest BitCast(const Source& value);
152
153 [[nodiscard]] U64 PackUint2x32(const Value& vector);
154 [[nodiscard]] Value UnpackUint2x32(const U64& value);
155
156 [[nodiscard]] U32 PackFloat2x16(const Value& vector);
157 [[nodiscard]] Value UnpackFloat2x16(const U32& value);
158
159 [[nodiscard]] U32 PackHalf2x16(const Value& vector);
160 [[nodiscard]] Value UnpackHalf2x16(const U32& value);
161
162 [[nodiscard]] F64 PackDouble2x32(const Value& vector);
163 [[nodiscard]] Value UnpackDouble2x32(const F64& value);
164
165 [[nodiscard]] F16F32F64 FPAdd(const F16F32F64& a, const F16F32F64& b, FpControl control = {});
166 [[nodiscard]] F16F32F64 FPMul(const F16F32F64& a, const F16F32F64& b, FpControl control = {});
167 [[nodiscard]] F16F32F64 FPFma(const F16F32F64& a, const F16F32F64& b, const F16F32F64& c,
168 FpControl control = {});
169
170 [[nodiscard]] F16F32F64 FPAbs(const F16F32F64& value);
171 [[nodiscard]] F16F32F64 FPNeg(const F16F32F64& value);
172 [[nodiscard]] F16F32F64 FPAbsNeg(const F16F32F64& value, bool abs, bool neg);
173
174 [[nodiscard]] F32 FPCos(const F32& value);
175 [[nodiscard]] F32 FPSin(const F32& value);
176 [[nodiscard]] F32 FPExp2(const F32& value);
177 [[nodiscard]] F32 FPLog2(const F32& value);
178 [[nodiscard]] F32F64 FPRecip(const F32F64& value);
179 [[nodiscard]] F32F64 FPRecipSqrt(const F32F64& value);
180 [[nodiscard]] F32 FPSqrt(const F32& value);
181 [[nodiscard]] F16F32F64 FPSaturate(const F16F32F64& value);
182 [[nodiscard]] F16F32F64 FPClamp(const F16F32F64& value, const F16F32F64& min_value,
183 const F16F32F64& max_value);
184 [[nodiscard]] F16F32F64 FPRoundEven(const F16F32F64& value, FpControl control = {});
185 [[nodiscard]] F16F32F64 FPFloor(const F16F32F64& value, FpControl control = {});
186 [[nodiscard]] F16F32F64 FPCeil(const F16F32F64& value, FpControl control = {});
187 [[nodiscard]] F16F32F64 FPTrunc(const F16F32F64& value, FpControl control = {});
188
189 [[nodiscard]] U1 FPEqual(const F16F32F64& lhs, const F16F32F64& rhs, FpControl control = {},
190 bool ordered = true);
191 [[nodiscard]] U1 FPNotEqual(const F16F32F64& lhs, const F16F32F64& rhs, FpControl control = {},
192 bool ordered = true);
193 [[nodiscard]] U1 FPLessThan(const F16F32F64& lhs, const F16F32F64& rhs, FpControl control = {},
194 bool ordered = true);
195 [[nodiscard]] U1 FPGreaterThan(const F16F32F64& lhs, const F16F32F64& rhs,
196 FpControl control = {}, bool ordered = true);
197 [[nodiscard]] U1 FPLessThanEqual(const F16F32F64& lhs, const F16F32F64& rhs,
198 FpControl control = {}, bool ordered = true);
199 [[nodiscard]] U1 FPGreaterThanEqual(const F16F32F64& lhs, const F16F32F64& rhs,
200 FpControl control = {}, bool ordered = true);
201 [[nodiscard]] U1 FPIsNan(const F16F32F64& value);
202 [[nodiscard]] U1 FPOrdered(const F16F32F64& lhs, const F16F32F64& rhs);
203 [[nodiscard]] U1 FPUnordered(const F16F32F64& lhs, const F16F32F64& rhs);
204 [[nodiscard]] F32F64 FPMax(const F32F64& lhs, const F32F64& rhs, FpControl control = {});
205 [[nodiscard]] F32F64 FPMin(const F32F64& lhs, const F32F64& rhs, FpControl control = {});
206
207 [[nodiscard]] U32U64 IAdd(const U32U64& a, const U32U64& b);
208 [[nodiscard]] U32U64 ISub(const U32U64& a, const U32U64& b);
209 [[nodiscard]] U32 IMul(const U32& a, const U32& b);
210 [[nodiscard]] U32U64 INeg(const U32U64& value);
211 [[nodiscard]] U32 IAbs(const U32& value);
212 [[nodiscard]] U32U64 ShiftLeftLogical(const U32U64& base, const U32& shift);
213 [[nodiscard]] U32U64 ShiftRightLogical(const U32U64& base, const U32& shift);
214 [[nodiscard]] U32U64 ShiftRightArithmetic(const U32U64& base, const U32& shift);
215 [[nodiscard]] U32 BitwiseAnd(const U32& a, const U32& b);
216 [[nodiscard]] U32 BitwiseOr(const U32& a, const U32& b);
217 [[nodiscard]] U32 BitwiseXor(const U32& a, const U32& b);
218 [[nodiscard]] U32 BitFieldInsert(const U32& base, const U32& insert, const U32& offset,
219 const U32& count);
220 [[nodiscard]] U32 BitFieldExtract(const U32& base, const U32& offset, const U32& count,
221 bool is_signed = false);
222 [[nodiscard]] U32 BitReverse(const U32& value);
223 [[nodiscard]] U32 BitCount(const U32& value);
224 [[nodiscard]] U32 BitwiseNot(const U32& value);
225
226 [[nodiscard]] U32 FindSMsb(const U32& value);
227 [[nodiscard]] U32 FindUMsb(const U32& value);
228 [[nodiscard]] U32 SMin(const U32& a, const U32& b);
229 [[nodiscard]] U32 UMin(const U32& a, const U32& b);
230 [[nodiscard]] U32 IMin(const U32& a, const U32& b, bool is_signed);
231 [[nodiscard]] U32 SMax(const U32& a, const U32& b);
232 [[nodiscard]] U32 UMax(const U32& a, const U32& b);
233 [[nodiscard]] U32 IMax(const U32& a, const U32& b, bool is_signed);
234 [[nodiscard]] U32 SClamp(const U32& value, const U32& min, const U32& max);
235 [[nodiscard]] U32 UClamp(const U32& value, const U32& min, const U32& max);
236
237 [[nodiscard]] U1 ILessThan(const U32& lhs, const U32& rhs, bool is_signed);
238 [[nodiscard]] U1 IEqual(const U32U64& lhs, const U32U64& rhs);
239 [[nodiscard]] U1 ILessThanEqual(const U32& lhs, const U32& rhs, bool is_signed);
240 [[nodiscard]] U1 IGreaterThan(const U32& lhs, const U32& rhs, bool is_signed);
241 [[nodiscard]] U1 INotEqual(const U32& lhs, const U32& rhs);
242 [[nodiscard]] U1 IGreaterThanEqual(const U32& lhs, const U32& rhs, bool is_signed);
243
244 [[nodiscard]] U32 SharedAtomicIAdd(const U32& pointer_offset, const U32& value);
245 [[nodiscard]] U32 SharedAtomicSMin(const U32& pointer_offset, const U32& value);
246 [[nodiscard]] U32 SharedAtomicUMin(const U32& pointer_offset, const U32& value);
247 [[nodiscard]] U32 SharedAtomicIMin(const U32& pointer_offset, const U32& value, bool is_signed);
248 [[nodiscard]] U32 SharedAtomicSMax(const U32& pointer_offset, const U32& value);
249 [[nodiscard]] U32 SharedAtomicUMax(const U32& pointer_offset, const U32& value);
250 [[nodiscard]] U32 SharedAtomicIMax(const U32& pointer_offset, const U32& value, bool is_signed);
251 [[nodiscard]] U32 SharedAtomicInc(const U32& pointer_offset, const U32& value);
252 [[nodiscard]] U32 SharedAtomicDec(const U32& pointer_offset, const U32& value);
253 [[nodiscard]] U32 SharedAtomicAnd(const U32& pointer_offset, const U32& value);
254 [[nodiscard]] U32 SharedAtomicOr(const U32& pointer_offset, const U32& value);
255 [[nodiscard]] U32 SharedAtomicXor(const U32& pointer_offset, const U32& value);
256 [[nodiscard]] U32U64 SharedAtomicExchange(const U32& pointer_offset, const U32U64& value);
257
258 [[nodiscard]] U32U64 GlobalAtomicIAdd(const U64& pointer_offset, const U32U64& value);
259 [[nodiscard]] U32U64 GlobalAtomicSMin(const U64& pointer_offset, const U32U64& value);
260 [[nodiscard]] U32U64 GlobalAtomicUMin(const U64& pointer_offset, const U32U64& value);
261 [[nodiscard]] U32U64 GlobalAtomicIMin(const U64& pointer_offset, const U32U64& value,
262 bool is_signed);
263 [[nodiscard]] U32U64 GlobalAtomicSMax(const U64& pointer_offset, const U32U64& value);
264 [[nodiscard]] U32U64 GlobalAtomicUMax(const U64& pointer_offset, const U32U64& value);
265 [[nodiscard]] U32U64 GlobalAtomicIMax(const U64& pointer_offset, const U32U64& value,
266 bool is_signed);
267 [[nodiscard]] U32 GlobalAtomicInc(const U64& pointer_offset, const U32& value);
268 [[nodiscard]] U32 GlobalAtomicDec(const U64& pointer_offset, const U32& value);
269 [[nodiscard]] U32U64 GlobalAtomicAnd(const U64& pointer_offset, const U32U64& value);
270 [[nodiscard]] U32U64 GlobalAtomicOr(const U64& pointer_offset, const U32U64& value);
271 [[nodiscard]] U32U64 GlobalAtomicXor(const U64& pointer_offset, const U32U64& value);
272 [[nodiscard]] U32U64 GlobalAtomicExchange(const U64& pointer_offset, const U32U64& value);
273
274 [[nodiscard]] F32 GlobalAtomicF32Add(const U64& pointer_offset, const Value& value,
275 const FpControl control = {});
276 [[nodiscard]] Value GlobalAtomicF16x2Add(const U64& pointer_offset, const Value& value,
277 const FpControl control = {});
278 [[nodiscard]] Value GlobalAtomicF16x2Min(const U64& pointer_offset, const Value& value,
279 const FpControl control = {});
280 [[nodiscard]] Value GlobalAtomicF16x2Max(const U64& pointer_offset, const Value& value,
281 const FpControl control = {});
282
283 [[nodiscard]] U1 LogicalOr(const U1& a, const U1& b);
284 [[nodiscard]] U1 LogicalAnd(const U1& a, const U1& b);
285 [[nodiscard]] U1 LogicalXor(const U1& a, const U1& b);
286 [[nodiscard]] U1 LogicalNot(const U1& value);
287
288 [[nodiscard]] U32U64 ConvertFToS(size_t bitsize, const F16F32F64& value);
289 [[nodiscard]] U32U64 ConvertFToU(size_t bitsize, const F16F32F64& value);
290 [[nodiscard]] U32U64 ConvertFToI(size_t bitsize, bool is_signed, const F16F32F64& value);
291 [[nodiscard]] F16F32F64 ConvertSToF(size_t dest_bitsize, size_t src_bitsize, const Value& value,
292 FpControl control = {});
293 [[nodiscard]] F16F32F64 ConvertUToF(size_t dest_bitsize, size_t src_bitsize, const Value& value,
294 FpControl control = {});
295 [[nodiscard]] F16F32F64 ConvertIToF(size_t dest_bitsize, size_t src_bitsize, bool is_signed,
296 const Value& value, FpControl control = {});
297
298 [[nodiscard]] U32U64 UConvert(size_t result_bitsize, const U32U64& value);
299 [[nodiscard]] F16F32F64 FPConvert(size_t result_bitsize, const F16F32F64& value,
300 FpControl control = {});
301
302 [[nodiscard]] Value ImageSampleImplicitLod(const Value& handle, const Value& coords,
303 const F32& bias, const Value& offset,
304 const F32& lod_clamp, TextureInstInfo info);
305 [[nodiscard]] Value ImageSampleExplicitLod(const Value& handle, const Value& coords,
306 const F32& lod, const Value& offset,
307 TextureInstInfo info);
308 [[nodiscard]] F32 ImageSampleDrefImplicitLod(const Value& handle, const Value& coords,
309 const F32& dref, const F32& bias,
310 const Value& offset, const F32& lod_clamp,
311 TextureInstInfo info);
312 [[nodiscard]] F32 ImageSampleDrefExplicitLod(const Value& handle, const Value& coords,
313 const F32& dref, const F32& lod,
314 const Value& offset, TextureInstInfo info);
315 [[nodiscard]] Value ImageQueryDimension(const Value& handle, const IR::U32& lod);
316
317 [[nodiscard]] Value ImageQueryLod(const Value& handle, const Value& coords,
318 TextureInstInfo info);
319 [[nodiscard]] Value ImageGather(const Value& handle, const Value& coords, const Value& offset,
320 const Value& offset2, TextureInstInfo info);
321 [[nodiscard]] Value ImageGatherDref(const Value& handle, const Value& coords,
322 const Value& offset, const Value& offset2, const F32& dref,
323 TextureInstInfo info);
324 [[nodiscard]] Value ImageFetch(const Value& handle, const Value& coords, const Value& offset,
325 const U32& lod, const U32& multisampling, TextureInstInfo info);
326 [[nodiscard]] Value ImageGradient(const Value& handle, const Value& coords,
327 const Value& derivates, const Value& offset,
328 const F32& lod_clamp, TextureInstInfo info);
329 [[nodiscard]] Value ImageRead(const Value& handle, const Value& coords, TextureInstInfo info);
330 [[nodiscard]] void ImageWrite(const Value& handle, const Value& coords, const Value& color,
331 TextureInstInfo info);
332
333 [[nodiscard]] Value ImageAtomicIAdd(const Value& handle, const Value& coords,
334 const Value& value, TextureInstInfo info);
335 [[nodiscard]] Value ImageAtomicSMin(const Value& handle, const Value& coords,
336 const Value& value, TextureInstInfo info);
337 [[nodiscard]] Value ImageAtomicUMin(const Value& handle, const Value& coords,
338 const Value& value, TextureInstInfo info);
339 [[nodiscard]] Value ImageAtomicIMin(const Value& handle, const Value& coords,
340 const Value& value, bool is_signed, TextureInstInfo info);
341 [[nodiscard]] Value ImageAtomicSMax(const Value& handle, const Value& coords,
342 const Value& value, TextureInstInfo info);
343 [[nodiscard]] Value ImageAtomicUMax(const Value& handle, const Value& coords,
344 const Value& value, TextureInstInfo info);
345 [[nodiscard]] Value ImageAtomicIMax(const Value& handle, const Value& coords,
346 const Value& value, bool is_signed, TextureInstInfo info);
347 [[nodiscard]] Value ImageAtomicInc(const Value& handle, const Value& coords, const Value& value,
348 TextureInstInfo info);
349 [[nodiscard]] Value ImageAtomicDec(const Value& handle, const Value& coords, const Value& value,
350 TextureInstInfo info);
351 [[nodiscard]] Value ImageAtomicAnd(const Value& handle, const Value& coords, const Value& value,
352 TextureInstInfo info);
353 [[nodiscard]] Value ImageAtomicOr(const Value& handle, const Value& coords, const Value& value,
354 TextureInstInfo info);
355 [[nodiscard]] Value ImageAtomicXor(const Value& handle, const Value& coords, const Value& value,
356 TextureInstInfo info);
357 [[nodiscard]] Value ImageAtomicExchange(const Value& handle, const Value& coords,
358 const Value& value, TextureInstInfo info);
359 [[nodiscard]] U1 VoteAll(const U1& value);
360 [[nodiscard]] U1 VoteAny(const U1& value);
361 [[nodiscard]] U1 VoteEqual(const U1& value);
362 [[nodiscard]] U32 SubgroupBallot(const U1& value);
363 [[nodiscard]] U32 SubgroupEqMask();
364 [[nodiscard]] U32 SubgroupLtMask();
365 [[nodiscard]] U32 SubgroupLeMask();
366 [[nodiscard]] U32 SubgroupGtMask();
367 [[nodiscard]] U32 SubgroupGeMask();
368 [[nodiscard]] U32 ShuffleIndex(const IR::U32& value, const IR::U32& index, const IR::U32& clamp,
369 const IR::U32& seg_mask);
370 [[nodiscard]] U32 ShuffleUp(const IR::U32& value, const IR::U32& index, const IR::U32& clamp,
371 const IR::U32& seg_mask);
372 [[nodiscard]] U32 ShuffleDown(const IR::U32& value, const IR::U32& index, const IR::U32& clamp,
373 const IR::U32& seg_mask);
374 [[nodiscard]] U32 ShuffleButterfly(const IR::U32& value, const IR::U32& index,
375 const IR::U32& clamp, const IR::U32& seg_mask);
376 [[nodiscard]] F32 FSwizzleAdd(const F32& a, const F32& b, const U32& swizzle,
377 FpControl control = {});
378
379 [[nodiscard]] F32 DPdxFine(const F32& a);
380
381 [[nodiscard]] F32 DPdyFine(const F32& a);
382
383 [[nodiscard]] F32 DPdxCoarse(const F32& a);
384
385 [[nodiscard]] F32 DPdyCoarse(const F32& a);
386
387private:
388 IR::Block::iterator insertion_point;
389
390 template <typename T = Value, typename... Args>
391 T Inst(Opcode op, Args... args) {
392 auto it{block->PrependNewInst(insertion_point, op, {Value{args}...})};
393 return T{Value{&*it}};
394 }
395
396 template <typename T>
397 requires(sizeof(T) <= sizeof(u32) && std::is_trivially_copyable_v<T>) struct Flags {
398 Flags() = default;
399 Flags(T proxy_) : proxy{proxy_} {}
400
401 T proxy;
402 };
403
404 template <typename T = Value, typename FlagType, typename... Args>
405 T Inst(Opcode op, Flags<FlagType> flags, Args... args) {
406 u32 raw_flags{};
407 std::memcpy(&raw_flags, &flags.proxy, sizeof(flags.proxy));
408 auto it{block->PrependNewInst(insertion_point, op, {Value{args}...}, raw_flags)};
409 return T{Value{&*it}};
410 }
411};
412
413} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/microinstruction.cpp b/src/shader_recompiler/frontend/ir/microinstruction.cpp
new file mode 100644
index 000000000..3dfa5a880
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/microinstruction.cpp
@@ -0,0 +1,411 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <algorithm>
6#include <memory>
7
8#include "shader_recompiler/exception.h"
9#include "shader_recompiler/frontend/ir/type.h"
10#include "shader_recompiler/frontend/ir/value.h"
11
12namespace Shader::IR {
13namespace {
14void CheckPseudoInstruction(IR::Inst* inst, IR::Opcode opcode) {
15 if (inst && inst->GetOpcode() != opcode) {
16 throw LogicError("Invalid pseudo-instruction");
17 }
18}
19
20void SetPseudoInstruction(IR::Inst*& dest_inst, IR::Inst* pseudo_inst) {
21 if (dest_inst) {
22 throw LogicError("Only one of each type of pseudo-op allowed");
23 }
24 dest_inst = pseudo_inst;
25}
26
27void RemovePseudoInstruction(IR::Inst*& inst, IR::Opcode expected_opcode) {
28 if (inst->GetOpcode() != expected_opcode) {
29 throw LogicError("Undoing use of invalid pseudo-op");
30 }
31 inst = nullptr;
32}
33
34void AllocAssociatedInsts(std::unique_ptr<AssociatedInsts>& associated_insts) {
35 if (!associated_insts) {
36 associated_insts = std::make_unique<AssociatedInsts>();
37 }
38}
39} // Anonymous namespace
40
41Inst::Inst(IR::Opcode op_, u32 flags_) noexcept : op{op_}, flags{flags_} {
42 if (op == Opcode::Phi) {
43 std::construct_at(&phi_args);
44 } else {
45 std::construct_at(&args);
46 }
47}
48
49Inst::~Inst() {
50 if (op == Opcode::Phi) {
51 std::destroy_at(&phi_args);
52 } else {
53 std::destroy_at(&args);
54 }
55}
56
57bool Inst::MayHaveSideEffects() const noexcept {
58 switch (op) {
59 case Opcode::ConditionRef:
60 case Opcode::Reference:
61 case Opcode::PhiMove:
62 case Opcode::Prologue:
63 case Opcode::Epilogue:
64 case Opcode::Join:
65 case Opcode::DemoteToHelperInvocation:
66 case Opcode::Barrier:
67 case Opcode::WorkgroupMemoryBarrier:
68 case Opcode::DeviceMemoryBarrier:
69 case Opcode::EmitVertex:
70 case Opcode::EndPrimitive:
71 case Opcode::SetAttribute:
72 case Opcode::SetAttributeIndexed:
73 case Opcode::SetPatch:
74 case Opcode::SetFragColor:
75 case Opcode::SetSampleMask:
76 case Opcode::SetFragDepth:
77 case Opcode::WriteGlobalU8:
78 case Opcode::WriteGlobalS8:
79 case Opcode::WriteGlobalU16:
80 case Opcode::WriteGlobalS16:
81 case Opcode::WriteGlobal32:
82 case Opcode::WriteGlobal64:
83 case Opcode::WriteGlobal128:
84 case Opcode::WriteStorageU8:
85 case Opcode::WriteStorageS8:
86 case Opcode::WriteStorageU16:
87 case Opcode::WriteStorageS16:
88 case Opcode::WriteStorage32:
89 case Opcode::WriteStorage64:
90 case Opcode::WriteStorage128:
91 case Opcode::WriteLocal:
92 case Opcode::WriteSharedU8:
93 case Opcode::WriteSharedU16:
94 case Opcode::WriteSharedU32:
95 case Opcode::WriteSharedU64:
96 case Opcode::WriteSharedU128:
97 case Opcode::SharedAtomicIAdd32:
98 case Opcode::SharedAtomicSMin32:
99 case Opcode::SharedAtomicUMin32:
100 case Opcode::SharedAtomicSMax32:
101 case Opcode::SharedAtomicUMax32:
102 case Opcode::SharedAtomicInc32:
103 case Opcode::SharedAtomicDec32:
104 case Opcode::SharedAtomicAnd32:
105 case Opcode::SharedAtomicOr32:
106 case Opcode::SharedAtomicXor32:
107 case Opcode::SharedAtomicExchange32:
108 case Opcode::SharedAtomicExchange64:
109 case Opcode::GlobalAtomicIAdd32:
110 case Opcode::GlobalAtomicSMin32:
111 case Opcode::GlobalAtomicUMin32:
112 case Opcode::GlobalAtomicSMax32:
113 case Opcode::GlobalAtomicUMax32:
114 case Opcode::GlobalAtomicInc32:
115 case Opcode::GlobalAtomicDec32:
116 case Opcode::GlobalAtomicAnd32:
117 case Opcode::GlobalAtomicOr32:
118 case Opcode::GlobalAtomicXor32:
119 case Opcode::GlobalAtomicExchange32:
120 case Opcode::GlobalAtomicIAdd64:
121 case Opcode::GlobalAtomicSMin64:
122 case Opcode::GlobalAtomicUMin64:
123 case Opcode::GlobalAtomicSMax64:
124 case Opcode::GlobalAtomicUMax64:
125 case Opcode::GlobalAtomicAnd64:
126 case Opcode::GlobalAtomicOr64:
127 case Opcode::GlobalAtomicXor64:
128 case Opcode::GlobalAtomicExchange64:
129 case Opcode::GlobalAtomicAddF32:
130 case Opcode::GlobalAtomicAddF16x2:
131 case Opcode::GlobalAtomicAddF32x2:
132 case Opcode::GlobalAtomicMinF16x2:
133 case Opcode::GlobalAtomicMinF32x2:
134 case Opcode::GlobalAtomicMaxF16x2:
135 case Opcode::GlobalAtomicMaxF32x2:
136 case Opcode::StorageAtomicIAdd32:
137 case Opcode::StorageAtomicSMin32:
138 case Opcode::StorageAtomicUMin32:
139 case Opcode::StorageAtomicSMax32:
140 case Opcode::StorageAtomicUMax32:
141 case Opcode::StorageAtomicInc32:
142 case Opcode::StorageAtomicDec32:
143 case Opcode::StorageAtomicAnd32:
144 case Opcode::StorageAtomicOr32:
145 case Opcode::StorageAtomicXor32:
146 case Opcode::StorageAtomicExchange32:
147 case Opcode::StorageAtomicIAdd64:
148 case Opcode::StorageAtomicSMin64:
149 case Opcode::StorageAtomicUMin64:
150 case Opcode::StorageAtomicSMax64:
151 case Opcode::StorageAtomicUMax64:
152 case Opcode::StorageAtomicAnd64:
153 case Opcode::StorageAtomicOr64:
154 case Opcode::StorageAtomicXor64:
155 case Opcode::StorageAtomicExchange64:
156 case Opcode::StorageAtomicAddF32:
157 case Opcode::StorageAtomicAddF16x2:
158 case Opcode::StorageAtomicAddF32x2:
159 case Opcode::StorageAtomicMinF16x2:
160 case Opcode::StorageAtomicMinF32x2:
161 case Opcode::StorageAtomicMaxF16x2:
162 case Opcode::StorageAtomicMaxF32x2:
163 case Opcode::BindlessImageWrite:
164 case Opcode::BoundImageWrite:
165 case Opcode::ImageWrite:
166 case IR::Opcode::BindlessImageAtomicIAdd32:
167 case IR::Opcode::BindlessImageAtomicSMin32:
168 case IR::Opcode::BindlessImageAtomicUMin32:
169 case IR::Opcode::BindlessImageAtomicSMax32:
170 case IR::Opcode::BindlessImageAtomicUMax32:
171 case IR::Opcode::BindlessImageAtomicInc32:
172 case IR::Opcode::BindlessImageAtomicDec32:
173 case IR::Opcode::BindlessImageAtomicAnd32:
174 case IR::Opcode::BindlessImageAtomicOr32:
175 case IR::Opcode::BindlessImageAtomicXor32:
176 case IR::Opcode::BindlessImageAtomicExchange32:
177 case IR::Opcode::BoundImageAtomicIAdd32:
178 case IR::Opcode::BoundImageAtomicSMin32:
179 case IR::Opcode::BoundImageAtomicUMin32:
180 case IR::Opcode::BoundImageAtomicSMax32:
181 case IR::Opcode::BoundImageAtomicUMax32:
182 case IR::Opcode::BoundImageAtomicInc32:
183 case IR::Opcode::BoundImageAtomicDec32:
184 case IR::Opcode::BoundImageAtomicAnd32:
185 case IR::Opcode::BoundImageAtomicOr32:
186 case IR::Opcode::BoundImageAtomicXor32:
187 case IR::Opcode::BoundImageAtomicExchange32:
188 case IR::Opcode::ImageAtomicIAdd32:
189 case IR::Opcode::ImageAtomicSMin32:
190 case IR::Opcode::ImageAtomicUMin32:
191 case IR::Opcode::ImageAtomicSMax32:
192 case IR::Opcode::ImageAtomicUMax32:
193 case IR::Opcode::ImageAtomicInc32:
194 case IR::Opcode::ImageAtomicDec32:
195 case IR::Opcode::ImageAtomicAnd32:
196 case IR::Opcode::ImageAtomicOr32:
197 case IR::Opcode::ImageAtomicXor32:
198 case IR::Opcode::ImageAtomicExchange32:
199 return true;
200 default:
201 return false;
202 }
203}
204
205bool Inst::IsPseudoInstruction() const noexcept {
206 switch (op) {
207 case Opcode::GetZeroFromOp:
208 case Opcode::GetSignFromOp:
209 case Opcode::GetCarryFromOp:
210 case Opcode::GetOverflowFromOp:
211 case Opcode::GetSparseFromOp:
212 case Opcode::GetInBoundsFromOp:
213 return true;
214 default:
215 return false;
216 }
217}
218
219bool Inst::AreAllArgsImmediates() const {
220 if (op == Opcode::Phi) {
221 throw LogicError("Testing for all arguments are immediates on phi instruction");
222 }
223 return std::all_of(args.begin(), args.begin() + NumArgs(),
224 [](const IR::Value& value) { return value.IsImmediate(); });
225}
226
227Inst* Inst::GetAssociatedPseudoOperation(IR::Opcode opcode) {
228 if (!associated_insts) {
229 return nullptr;
230 }
231 switch (opcode) {
232 case Opcode::GetZeroFromOp:
233 CheckPseudoInstruction(associated_insts->zero_inst, Opcode::GetZeroFromOp);
234 return associated_insts->zero_inst;
235 case Opcode::GetSignFromOp:
236 CheckPseudoInstruction(associated_insts->sign_inst, Opcode::GetSignFromOp);
237 return associated_insts->sign_inst;
238 case Opcode::GetCarryFromOp:
239 CheckPseudoInstruction(associated_insts->carry_inst, Opcode::GetCarryFromOp);
240 return associated_insts->carry_inst;
241 case Opcode::GetOverflowFromOp:
242 CheckPseudoInstruction(associated_insts->overflow_inst, Opcode::GetOverflowFromOp);
243 return associated_insts->overflow_inst;
244 case Opcode::GetSparseFromOp:
245 CheckPseudoInstruction(associated_insts->sparse_inst, Opcode::GetSparseFromOp);
246 return associated_insts->sparse_inst;
247 case Opcode::GetInBoundsFromOp:
248 CheckPseudoInstruction(associated_insts->in_bounds_inst, Opcode::GetInBoundsFromOp);
249 return associated_insts->in_bounds_inst;
250 default:
251 throw InvalidArgument("{} is not a pseudo-instruction", opcode);
252 }
253}
254
255IR::Type Inst::Type() const {
256 return TypeOf(op);
257}
258
259void Inst::SetArg(size_t index, Value value) {
260 if (index >= NumArgs()) {
261 throw InvalidArgument("Out of bounds argument index {} in opcode {}", index, op);
262 }
263 const IR::Value arg{Arg(index)};
264 if (!arg.IsImmediate()) {
265 UndoUse(arg);
266 }
267 if (!value.IsImmediate()) {
268 Use(value);
269 }
270 if (op == Opcode::Phi) {
271 phi_args[index].second = value;
272 } else {
273 args[index] = value;
274 }
275}
276
277Block* Inst::PhiBlock(size_t index) const {
278 if (op != Opcode::Phi) {
279 throw LogicError("{} is not a Phi instruction", op);
280 }
281 if (index >= phi_args.size()) {
282 throw InvalidArgument("Out of bounds argument index {} in phi instruction");
283 }
284 return phi_args[index].first;
285}
286
287void Inst::AddPhiOperand(Block* predecessor, const Value& value) {
288 if (!value.IsImmediate()) {
289 Use(value);
290 }
291 phi_args.emplace_back(predecessor, value);
292}
293
294void Inst::Invalidate() {
295 ClearArgs();
296 ReplaceOpcode(Opcode::Void);
297}
298
299void Inst::ClearArgs() {
300 if (op == Opcode::Phi) {
301 for (auto& pair : phi_args) {
302 IR::Value& value{pair.second};
303 if (!value.IsImmediate()) {
304 UndoUse(value);
305 }
306 }
307 phi_args.clear();
308 } else {
309 for (auto& value : args) {
310 if (!value.IsImmediate()) {
311 UndoUse(value);
312 }
313 }
314 // Reset arguments to null
315 // std::memset was measured to be faster on MSVC than std::ranges:fill
316 std::memset(reinterpret_cast<char*>(&args), 0, sizeof(args));
317 }
318}
319
320void Inst::ReplaceUsesWith(Value replacement) {
321 Invalidate();
322 ReplaceOpcode(Opcode::Identity);
323 if (!replacement.IsImmediate()) {
324 Use(replacement);
325 }
326 args[0] = replacement;
327}
328
329void Inst::ReplaceOpcode(IR::Opcode opcode) {
330 if (opcode == IR::Opcode::Phi) {
331 throw LogicError("Cannot transition into Phi");
332 }
333 if (op == Opcode::Phi) {
334 // Transition out of phi arguments into non-phi
335 std::destroy_at(&phi_args);
336 std::construct_at(&args);
337 }
338 op = opcode;
339}
340
341void Inst::Use(const Value& value) {
342 Inst* const inst{value.Inst()};
343 ++inst->use_count;
344
345 std::unique_ptr<AssociatedInsts>& assoc_inst{inst->associated_insts};
346 switch (op) {
347 case Opcode::GetZeroFromOp:
348 AllocAssociatedInsts(assoc_inst);
349 SetPseudoInstruction(assoc_inst->zero_inst, this);
350 break;
351 case Opcode::GetSignFromOp:
352 AllocAssociatedInsts(assoc_inst);
353 SetPseudoInstruction(assoc_inst->sign_inst, this);
354 break;
355 case Opcode::GetCarryFromOp:
356 AllocAssociatedInsts(assoc_inst);
357 SetPseudoInstruction(assoc_inst->carry_inst, this);
358 break;
359 case Opcode::GetOverflowFromOp:
360 AllocAssociatedInsts(assoc_inst);
361 SetPseudoInstruction(assoc_inst->overflow_inst, this);
362 break;
363 case Opcode::GetSparseFromOp:
364 AllocAssociatedInsts(assoc_inst);
365 SetPseudoInstruction(assoc_inst->sparse_inst, this);
366 break;
367 case Opcode::GetInBoundsFromOp:
368 AllocAssociatedInsts(assoc_inst);
369 SetPseudoInstruction(assoc_inst->in_bounds_inst, this);
370 break;
371 default:
372 break;
373 }
374}
375
376void Inst::UndoUse(const Value& value) {
377 Inst* const inst{value.Inst()};
378 --inst->use_count;
379
380 std::unique_ptr<AssociatedInsts>& assoc_inst{inst->associated_insts};
381 switch (op) {
382 case Opcode::GetZeroFromOp:
383 AllocAssociatedInsts(assoc_inst);
384 RemovePseudoInstruction(assoc_inst->zero_inst, Opcode::GetZeroFromOp);
385 break;
386 case Opcode::GetSignFromOp:
387 AllocAssociatedInsts(assoc_inst);
388 RemovePseudoInstruction(assoc_inst->sign_inst, Opcode::GetSignFromOp);
389 break;
390 case Opcode::GetCarryFromOp:
391 AllocAssociatedInsts(assoc_inst);
392 RemovePseudoInstruction(assoc_inst->carry_inst, Opcode::GetCarryFromOp);
393 break;
394 case Opcode::GetOverflowFromOp:
395 AllocAssociatedInsts(assoc_inst);
396 RemovePseudoInstruction(assoc_inst->overflow_inst, Opcode::GetOverflowFromOp);
397 break;
398 case Opcode::GetSparseFromOp:
399 AllocAssociatedInsts(assoc_inst);
400 RemovePseudoInstruction(assoc_inst->sparse_inst, Opcode::GetSparseFromOp);
401 break;
402 case Opcode::GetInBoundsFromOp:
403 AllocAssociatedInsts(assoc_inst);
404 RemovePseudoInstruction(assoc_inst->in_bounds_inst, Opcode::GetInBoundsFromOp);
405 break;
406 default:
407 break;
408 }
409}
410
411} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/modifiers.h b/src/shader_recompiler/frontend/ir/modifiers.h
new file mode 100644
index 000000000..77cda1f8a
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/modifiers.h
@@ -0,0 +1,49 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include "common/bit_field.h"
8#include "common/common_types.h"
9#include "shader_recompiler/shader_info.h"
10
11namespace Shader::IR {
12
13enum class FmzMode : u8 {
14 DontCare, // Not specified for this instruction
15 FTZ, // Flush denorms to zero, NAN is propagated (D3D11, NVN, GL, VK)
16 FMZ, // Flush denorms to zero, x * 0 == 0 (D3D9)
17 None, // Denorms are not flushed, NAN is propagated (nouveau)
18};
19
20enum class FpRounding : u8 {
21 DontCare, // Not specified for this instruction
22 RN, // Round to nearest even,
23 RM, // Round towards negative infinity
24 RP, // Round towards positive infinity
25 RZ, // Round towards zero
26};
27
28struct FpControl {
29 bool no_contraction{false};
30 FpRounding rounding{FpRounding::DontCare};
31 FmzMode fmz_mode{FmzMode::DontCare};
32};
33static_assert(sizeof(FpControl) <= sizeof(u32));
34
35union TextureInstInfo {
36 u32 raw;
37 BitField<0, 16, u32> descriptor_index;
38 BitField<16, 3, TextureType> type;
39 BitField<19, 1, u32> is_depth;
40 BitField<20, 1, u32> has_bias;
41 BitField<21, 1, u32> has_lod_clamp;
42 BitField<22, 1, u32> relaxed_precision;
43 BitField<23, 2, u32> gather_component;
44 BitField<25, 2, u32> num_derivates;
45 BitField<27, 3, ImageFormat> image_format;
46};
47static_assert(sizeof(TextureInstInfo) <= sizeof(u32));
48
49} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/opcodes.cpp b/src/shader_recompiler/frontend/ir/opcodes.cpp
new file mode 100644
index 000000000..24d024ad7
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/opcodes.cpp
@@ -0,0 +1,15 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <string_view>
6
7#include "shader_recompiler/frontend/ir/opcodes.h"
8
9namespace Shader::IR {
10
11std::string_view NameOf(Opcode op) {
12 return Detail::META_TABLE[static_cast<size_t>(op)].name;
13}
14
15} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/opcodes.h b/src/shader_recompiler/frontend/ir/opcodes.h
new file mode 100644
index 000000000..9ab108292
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/opcodes.h
@@ -0,0 +1,110 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <algorithm>
8#include <array>
9#include <string_view>
10
11#include <fmt/format.h>
12
13#include "shader_recompiler/frontend/ir/type.h"
14
15namespace Shader::IR {
16
17enum class Opcode {
18#define OPCODE(name, ...) name,
19#include "opcodes.inc"
20#undef OPCODE
21};
22
23namespace Detail {
24struct OpcodeMeta {
25 std::string_view name;
26 Type type;
27 std::array<Type, 5> arg_types;
28};
29
30// using enum Type;
31constexpr Type Void{Type::Void};
32constexpr Type Opaque{Type::Opaque};
33constexpr Type Reg{Type::Reg};
34constexpr Type Pred{Type::Pred};
35constexpr Type Attribute{Type::Attribute};
36constexpr Type Patch{Type::Patch};
37constexpr Type U1{Type::U1};
38constexpr Type U8{Type::U8};
39constexpr Type U16{Type::U16};
40constexpr Type U32{Type::U32};
41constexpr Type U64{Type::U64};
42constexpr Type F16{Type::F16};
43constexpr Type F32{Type::F32};
44constexpr Type F64{Type::F64};
45constexpr Type U32x2{Type::U32x2};
46constexpr Type U32x3{Type::U32x3};
47constexpr Type U32x4{Type::U32x4};
48constexpr Type F16x2{Type::F16x2};
49constexpr Type F16x3{Type::F16x3};
50constexpr Type F16x4{Type::F16x4};
51constexpr Type F32x2{Type::F32x2};
52constexpr Type F32x3{Type::F32x3};
53constexpr Type F32x4{Type::F32x4};
54constexpr Type F64x2{Type::F64x2};
55constexpr Type F64x3{Type::F64x3};
56constexpr Type F64x4{Type::F64x4};
57
58constexpr OpcodeMeta META_TABLE[]{
59#define OPCODE(name_token, type_token, ...) \
60 { \
61 .name{#name_token}, \
62 .type = type_token, \
63 .arg_types{__VA_ARGS__}, \
64 },
65#include "opcodes.inc"
66#undef OPCODE
67};
68constexpr size_t CalculateNumArgsOf(Opcode op) {
69 const auto& arg_types{META_TABLE[static_cast<size_t>(op)].arg_types};
70 return static_cast<size_t>(
71 std::distance(arg_types.begin(), std::ranges::find(arg_types, Type::Void)));
72}
73
74constexpr u8 NUM_ARGS[]{
75#define OPCODE(name_token, type_token, ...) static_cast<u8>(CalculateNumArgsOf(Opcode::name_token)),
76#include "opcodes.inc"
77#undef OPCODE
78};
79} // namespace Detail
80
81/// Get return type of an opcode
82[[nodiscard]] inline Type TypeOf(Opcode op) noexcept {
83 return Detail::META_TABLE[static_cast<size_t>(op)].type;
84}
85
86/// Get the number of arguments an opcode accepts
87[[nodiscard]] inline size_t NumArgsOf(Opcode op) noexcept {
88 return static_cast<size_t>(Detail::NUM_ARGS[static_cast<size_t>(op)]);
89}
90
91/// Get the required type of an argument of an opcode
92[[nodiscard]] inline Type ArgTypeOf(Opcode op, size_t arg_index) noexcept {
93 return Detail::META_TABLE[static_cast<size_t>(op)].arg_types[arg_index];
94}
95
96/// Get the name of an opcode
97[[nodiscard]] std::string_view NameOf(Opcode op);
98
99} // namespace Shader::IR
100
101template <>
102struct fmt::formatter<Shader::IR::Opcode> {
103 constexpr auto parse(format_parse_context& ctx) {
104 return ctx.begin();
105 }
106 template <typename FormatContext>
107 auto format(const Shader::IR::Opcode& op, FormatContext& ctx) {
108 return format_to(ctx.out(), "{}", Shader::IR::NameOf(op));
109 }
110};
diff --git a/src/shader_recompiler/frontend/ir/opcodes.inc b/src/shader_recompiler/frontend/ir/opcodes.inc
new file mode 100644
index 000000000..d91098c80
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/opcodes.inc
@@ -0,0 +1,550 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5// opcode name, return type, arg1 type, arg2 type, arg3 type, arg4 type, arg4 type, ...
6OPCODE(Phi, Opaque, )
7OPCODE(Identity, Opaque, Opaque, )
8OPCODE(Void, Void, )
9OPCODE(ConditionRef, U1, U1, )
10OPCODE(Reference, Void, Opaque, )
11OPCODE(PhiMove, Void, Opaque, Opaque, )
12
13// Special operations
14OPCODE(Prologue, Void, )
15OPCODE(Epilogue, Void, )
16OPCODE(Join, Void, )
17OPCODE(DemoteToHelperInvocation, Void, )
18OPCODE(EmitVertex, Void, U32, )
19OPCODE(EndPrimitive, Void, U32, )
20
21// Barriers
22OPCODE(Barrier, Void, )
23OPCODE(WorkgroupMemoryBarrier, Void, )
24OPCODE(DeviceMemoryBarrier, Void, )
25
26// Context getters/setters
27OPCODE(GetRegister, U32, Reg, )
28OPCODE(SetRegister, Void, Reg, U32, )
29OPCODE(GetPred, U1, Pred, )
30OPCODE(SetPred, Void, Pred, U1, )
31OPCODE(GetGotoVariable, U1, U32, )
32OPCODE(SetGotoVariable, Void, U32, U1, )
33OPCODE(GetIndirectBranchVariable, U32, )
34OPCODE(SetIndirectBranchVariable, Void, U32, )
35OPCODE(GetCbufU8, U32, U32, U32, )
36OPCODE(GetCbufS8, U32, U32, U32, )
37OPCODE(GetCbufU16, U32, U32, U32, )
38OPCODE(GetCbufS16, U32, U32, U32, )
39OPCODE(GetCbufU32, U32, U32, U32, )
40OPCODE(GetCbufF32, F32, U32, U32, )
41OPCODE(GetCbufU32x2, U32x2, U32, U32, )
42OPCODE(GetAttribute, F32, Attribute, U32, )
43OPCODE(SetAttribute, Void, Attribute, F32, U32, )
44OPCODE(GetAttributeIndexed, F32, U32, U32, )
45OPCODE(SetAttributeIndexed, Void, U32, F32, U32, )
46OPCODE(GetPatch, F32, Patch, )
47OPCODE(SetPatch, Void, Patch, F32, )
48OPCODE(SetFragColor, Void, U32, U32, F32, )
49OPCODE(SetSampleMask, Void, U32, )
50OPCODE(SetFragDepth, Void, F32, )
51OPCODE(GetZFlag, U1, Void, )
52OPCODE(GetSFlag, U1, Void, )
53OPCODE(GetCFlag, U1, Void, )
54OPCODE(GetOFlag, U1, Void, )
55OPCODE(SetZFlag, Void, U1, )
56OPCODE(SetSFlag, Void, U1, )
57OPCODE(SetCFlag, Void, U1, )
58OPCODE(SetOFlag, Void, U1, )
59OPCODE(WorkgroupId, U32x3, )
60OPCODE(LocalInvocationId, U32x3, )
61OPCODE(InvocationId, U32, )
62OPCODE(SampleId, U32, )
63OPCODE(IsHelperInvocation, U1, )
64OPCODE(YDirection, F32, )
65
66// Undefined
67OPCODE(UndefU1, U1, )
68OPCODE(UndefU8, U8, )
69OPCODE(UndefU16, U16, )
70OPCODE(UndefU32, U32, )
71OPCODE(UndefU64, U64, )
72
73// Memory operations
74OPCODE(LoadGlobalU8, U32, Opaque, )
75OPCODE(LoadGlobalS8, U32, Opaque, )
76OPCODE(LoadGlobalU16, U32, Opaque, )
77OPCODE(LoadGlobalS16, U32, Opaque, )
78OPCODE(LoadGlobal32, U32, Opaque, )
79OPCODE(LoadGlobal64, U32x2, Opaque, )
80OPCODE(LoadGlobal128, U32x4, Opaque, )
81OPCODE(WriteGlobalU8, Void, Opaque, U32, )
82OPCODE(WriteGlobalS8, Void, Opaque, U32, )
83OPCODE(WriteGlobalU16, Void, Opaque, U32, )
84OPCODE(WriteGlobalS16, Void, Opaque, U32, )
85OPCODE(WriteGlobal32, Void, Opaque, U32, )
86OPCODE(WriteGlobal64, Void, Opaque, U32x2, )
87OPCODE(WriteGlobal128, Void, Opaque, U32x4, )
88
89// Storage buffer operations
90OPCODE(LoadStorageU8, U32, U32, U32, )
91OPCODE(LoadStorageS8, U32, U32, U32, )
92OPCODE(LoadStorageU16, U32, U32, U32, )
93OPCODE(LoadStorageS16, U32, U32, U32, )
94OPCODE(LoadStorage32, U32, U32, U32, )
95OPCODE(LoadStorage64, U32x2, U32, U32, )
96OPCODE(LoadStorage128, U32x4, U32, U32, )
97OPCODE(WriteStorageU8, Void, U32, U32, U32, )
98OPCODE(WriteStorageS8, Void, U32, U32, U32, )
99OPCODE(WriteStorageU16, Void, U32, U32, U32, )
100OPCODE(WriteStorageS16, Void, U32, U32, U32, )
101OPCODE(WriteStorage32, Void, U32, U32, U32, )
102OPCODE(WriteStorage64, Void, U32, U32, U32x2, )
103OPCODE(WriteStorage128, Void, U32, U32, U32x4, )
104
105// Local memory operations
106OPCODE(LoadLocal, U32, U32, )
107OPCODE(WriteLocal, Void, U32, U32, )
108
109// Shared memory operations
110OPCODE(LoadSharedU8, U32, U32, )
111OPCODE(LoadSharedS8, U32, U32, )
112OPCODE(LoadSharedU16, U32, U32, )
113OPCODE(LoadSharedS16, U32, U32, )
114OPCODE(LoadSharedU32, U32, U32, )
115OPCODE(LoadSharedU64, U32x2, U32, )
116OPCODE(LoadSharedU128, U32x4, U32, )
117OPCODE(WriteSharedU8, Void, U32, U32, )
118OPCODE(WriteSharedU16, Void, U32, U32, )
119OPCODE(WriteSharedU32, Void, U32, U32, )
120OPCODE(WriteSharedU64, Void, U32, U32x2, )
121OPCODE(WriteSharedU128, Void, U32, U32x4, )
122
123// Vector utility
124OPCODE(CompositeConstructU32x2, U32x2, U32, U32, )
125OPCODE(CompositeConstructU32x3, U32x3, U32, U32, U32, )
126OPCODE(CompositeConstructU32x4, U32x4, U32, U32, U32, U32, )
127OPCODE(CompositeExtractU32x2, U32, U32x2, U32, )
128OPCODE(CompositeExtractU32x3, U32, U32x3, U32, )
129OPCODE(CompositeExtractU32x4, U32, U32x4, U32, )
130OPCODE(CompositeInsertU32x2, U32x2, U32x2, U32, U32, )
131OPCODE(CompositeInsertU32x3, U32x3, U32x3, U32, U32, )
132OPCODE(CompositeInsertU32x4, U32x4, U32x4, U32, U32, )
133OPCODE(CompositeConstructF16x2, F16x2, F16, F16, )
134OPCODE(CompositeConstructF16x3, F16x3, F16, F16, F16, )
135OPCODE(CompositeConstructF16x4, F16x4, F16, F16, F16, F16, )
136OPCODE(CompositeExtractF16x2, F16, F16x2, U32, )
137OPCODE(CompositeExtractF16x3, F16, F16x3, U32, )
138OPCODE(CompositeExtractF16x4, F16, F16x4, U32, )
139OPCODE(CompositeInsertF16x2, F16x2, F16x2, F16, U32, )
140OPCODE(CompositeInsertF16x3, F16x3, F16x3, F16, U32, )
141OPCODE(CompositeInsertF16x4, F16x4, F16x4, F16, U32, )
142OPCODE(CompositeConstructF32x2, F32x2, F32, F32, )
143OPCODE(CompositeConstructF32x3, F32x3, F32, F32, F32, )
144OPCODE(CompositeConstructF32x4, F32x4, F32, F32, F32, F32, )
145OPCODE(CompositeExtractF32x2, F32, F32x2, U32, )
146OPCODE(CompositeExtractF32x3, F32, F32x3, U32, )
147OPCODE(CompositeExtractF32x4, F32, F32x4, U32, )
148OPCODE(CompositeInsertF32x2, F32x2, F32x2, F32, U32, )
149OPCODE(CompositeInsertF32x3, F32x3, F32x3, F32, U32, )
150OPCODE(CompositeInsertF32x4, F32x4, F32x4, F32, U32, )
151OPCODE(CompositeConstructF64x2, F64x2, F64, F64, )
152OPCODE(CompositeConstructF64x3, F64x3, F64, F64, F64, )
153OPCODE(CompositeConstructF64x4, F64x4, F64, F64, F64, F64, )
154OPCODE(CompositeExtractF64x2, F64, F64x2, U32, )
155OPCODE(CompositeExtractF64x3, F64, F64x3, U32, )
156OPCODE(CompositeExtractF64x4, F64, F64x4, U32, )
157OPCODE(CompositeInsertF64x2, F64x2, F64x2, F64, U32, )
158OPCODE(CompositeInsertF64x3, F64x3, F64x3, F64, U32, )
159OPCODE(CompositeInsertF64x4, F64x4, F64x4, F64, U32, )
160
161// Select operations
162OPCODE(SelectU1, U1, U1, U1, U1, )
163OPCODE(SelectU8, U8, U1, U8, U8, )
164OPCODE(SelectU16, U16, U1, U16, U16, )
165OPCODE(SelectU32, U32, U1, U32, U32, )
166OPCODE(SelectU64, U64, U1, U64, U64, )
167OPCODE(SelectF16, F16, U1, F16, F16, )
168OPCODE(SelectF32, F32, U1, F32, F32, )
169OPCODE(SelectF64, F64, U1, F64, F64, )
170
171// Bitwise conversions
172OPCODE(BitCastU16F16, U16, F16, )
173OPCODE(BitCastU32F32, U32, F32, )
174OPCODE(BitCastU64F64, U64, F64, )
175OPCODE(BitCastF16U16, F16, U16, )
176OPCODE(BitCastF32U32, F32, U32, )
177OPCODE(BitCastF64U64, F64, U64, )
178OPCODE(PackUint2x32, U64, U32x2, )
179OPCODE(UnpackUint2x32, U32x2, U64, )
180OPCODE(PackFloat2x16, U32, F16x2, )
181OPCODE(UnpackFloat2x16, F16x2, U32, )
182OPCODE(PackHalf2x16, U32, F32x2, )
183OPCODE(UnpackHalf2x16, F32x2, U32, )
184OPCODE(PackDouble2x32, F64, U32x2, )
185OPCODE(UnpackDouble2x32, U32x2, F64, )
186
187// Pseudo-operation, handled specially at final emit
188OPCODE(GetZeroFromOp, U1, Opaque, )
189OPCODE(GetSignFromOp, U1, Opaque, )
190OPCODE(GetCarryFromOp, U1, Opaque, )
191OPCODE(GetOverflowFromOp, U1, Opaque, )
192OPCODE(GetSparseFromOp, U1, Opaque, )
193OPCODE(GetInBoundsFromOp, U1, Opaque, )
194
195// Floating-point operations
196OPCODE(FPAbs16, F16, F16, )
197OPCODE(FPAbs32, F32, F32, )
198OPCODE(FPAbs64, F64, F64, )
199OPCODE(FPAdd16, F16, F16, F16, )
200OPCODE(FPAdd32, F32, F32, F32, )
201OPCODE(FPAdd64, F64, F64, F64, )
202OPCODE(FPFma16, F16, F16, F16, F16, )
203OPCODE(FPFma32, F32, F32, F32, F32, )
204OPCODE(FPFma64, F64, F64, F64, F64, )
205OPCODE(FPMax32, F32, F32, F32, )
206OPCODE(FPMax64, F64, F64, F64, )
207OPCODE(FPMin32, F32, F32, F32, )
208OPCODE(FPMin64, F64, F64, F64, )
209OPCODE(FPMul16, F16, F16, F16, )
210OPCODE(FPMul32, F32, F32, F32, )
211OPCODE(FPMul64, F64, F64, F64, )
212OPCODE(FPNeg16, F16, F16, )
213OPCODE(FPNeg32, F32, F32, )
214OPCODE(FPNeg64, F64, F64, )
215OPCODE(FPRecip32, F32, F32, )
216OPCODE(FPRecip64, F64, F64, )
217OPCODE(FPRecipSqrt32, F32, F32, )
218OPCODE(FPRecipSqrt64, F64, F64, )
219OPCODE(FPSqrt, F32, F32, )
220OPCODE(FPSin, F32, F32, )
221OPCODE(FPExp2, F32, F32, )
222OPCODE(FPCos, F32, F32, )
223OPCODE(FPLog2, F32, F32, )
224OPCODE(FPSaturate16, F16, F16, )
225OPCODE(FPSaturate32, F32, F32, )
226OPCODE(FPSaturate64, F64, F64, )
227OPCODE(FPClamp16, F16, F16, F16, F16, )
228OPCODE(FPClamp32, F32, F32, F32, F32, )
229OPCODE(FPClamp64, F64, F64, F64, F64, )
230OPCODE(FPRoundEven16, F16, F16, )
231OPCODE(FPRoundEven32, F32, F32, )
232OPCODE(FPRoundEven64, F64, F64, )
233OPCODE(FPFloor16, F16, F16, )
234OPCODE(FPFloor32, F32, F32, )
235OPCODE(FPFloor64, F64, F64, )
236OPCODE(FPCeil16, F16, F16, )
237OPCODE(FPCeil32, F32, F32, )
238OPCODE(FPCeil64, F64, F64, )
239OPCODE(FPTrunc16, F16, F16, )
240OPCODE(FPTrunc32, F32, F32, )
241OPCODE(FPTrunc64, F64, F64, )
242
243OPCODE(FPOrdEqual16, U1, F16, F16, )
244OPCODE(FPOrdEqual32, U1, F32, F32, )
245OPCODE(FPOrdEqual64, U1, F64, F64, )
246OPCODE(FPUnordEqual16, U1, F16, F16, )
247OPCODE(FPUnordEqual32, U1, F32, F32, )
248OPCODE(FPUnordEqual64, U1, F64, F64, )
249OPCODE(FPOrdNotEqual16, U1, F16, F16, )
250OPCODE(FPOrdNotEqual32, U1, F32, F32, )
251OPCODE(FPOrdNotEqual64, U1, F64, F64, )
252OPCODE(FPUnordNotEqual16, U1, F16, F16, )
253OPCODE(FPUnordNotEqual32, U1, F32, F32, )
254OPCODE(FPUnordNotEqual64, U1, F64, F64, )
255OPCODE(FPOrdLessThan16, U1, F16, F16, )
256OPCODE(FPOrdLessThan32, U1, F32, F32, )
257OPCODE(FPOrdLessThan64, U1, F64, F64, )
258OPCODE(FPUnordLessThan16, U1, F16, F16, )
259OPCODE(FPUnordLessThan32, U1, F32, F32, )
260OPCODE(FPUnordLessThan64, U1, F64, F64, )
261OPCODE(FPOrdGreaterThan16, U1, F16, F16, )
262OPCODE(FPOrdGreaterThan32, U1, F32, F32, )
263OPCODE(FPOrdGreaterThan64, U1, F64, F64, )
264OPCODE(FPUnordGreaterThan16, U1, F16, F16, )
265OPCODE(FPUnordGreaterThan32, U1, F32, F32, )
266OPCODE(FPUnordGreaterThan64, U1, F64, F64, )
267OPCODE(FPOrdLessThanEqual16, U1, F16, F16, )
268OPCODE(FPOrdLessThanEqual32, U1, F32, F32, )
269OPCODE(FPOrdLessThanEqual64, U1, F64, F64, )
270OPCODE(FPUnordLessThanEqual16, U1, F16, F16, )
271OPCODE(FPUnordLessThanEqual32, U1, F32, F32, )
272OPCODE(FPUnordLessThanEqual64, U1, F64, F64, )
273OPCODE(FPOrdGreaterThanEqual16, U1, F16, F16, )
274OPCODE(FPOrdGreaterThanEqual32, U1, F32, F32, )
275OPCODE(FPOrdGreaterThanEqual64, U1, F64, F64, )
276OPCODE(FPUnordGreaterThanEqual16, U1, F16, F16, )
277OPCODE(FPUnordGreaterThanEqual32, U1, F32, F32, )
278OPCODE(FPUnordGreaterThanEqual64, U1, F64, F64, )
279OPCODE(FPIsNan16, U1, F16, )
280OPCODE(FPIsNan32, U1, F32, )
281OPCODE(FPIsNan64, U1, F64, )
282
283// Integer operations
284OPCODE(IAdd32, U32, U32, U32, )
285OPCODE(IAdd64, U64, U64, U64, )
286OPCODE(ISub32, U32, U32, U32, )
287OPCODE(ISub64, U64, U64, U64, )
288OPCODE(IMul32, U32, U32, U32, )
289OPCODE(INeg32, U32, U32, )
290OPCODE(INeg64, U64, U64, )
291OPCODE(IAbs32, U32, U32, )
292OPCODE(ShiftLeftLogical32, U32, U32, U32, )
293OPCODE(ShiftLeftLogical64, U64, U64, U32, )
294OPCODE(ShiftRightLogical32, U32, U32, U32, )
295OPCODE(ShiftRightLogical64, U64, U64, U32, )
296OPCODE(ShiftRightArithmetic32, U32, U32, U32, )
297OPCODE(ShiftRightArithmetic64, U64, U64, U32, )
298OPCODE(BitwiseAnd32, U32, U32, U32, )
299OPCODE(BitwiseOr32, U32, U32, U32, )
300OPCODE(BitwiseXor32, U32, U32, U32, )
301OPCODE(BitFieldInsert, U32, U32, U32, U32, U32, )
302OPCODE(BitFieldSExtract, U32, U32, U32, U32, )
303OPCODE(BitFieldUExtract, U32, U32, U32, U32, )
304OPCODE(BitReverse32, U32, U32, )
305OPCODE(BitCount32, U32, U32, )
306OPCODE(BitwiseNot32, U32, U32, )
307
308OPCODE(FindSMsb32, U32, U32, )
309OPCODE(FindUMsb32, U32, U32, )
310OPCODE(SMin32, U32, U32, U32, )
311OPCODE(UMin32, U32, U32, U32, )
312OPCODE(SMax32, U32, U32, U32, )
313OPCODE(UMax32, U32, U32, U32, )
314OPCODE(SClamp32, U32, U32, U32, U32, )
315OPCODE(UClamp32, U32, U32, U32, U32, )
316OPCODE(SLessThan, U1, U32, U32, )
317OPCODE(ULessThan, U1, U32, U32, )
318OPCODE(IEqual, U1, U32, U32, )
319OPCODE(SLessThanEqual, U1, U32, U32, )
320OPCODE(ULessThanEqual, U1, U32, U32, )
321OPCODE(SGreaterThan, U1, U32, U32, )
322OPCODE(UGreaterThan, U1, U32, U32, )
323OPCODE(INotEqual, U1, U32, U32, )
324OPCODE(SGreaterThanEqual, U1, U32, U32, )
325OPCODE(UGreaterThanEqual, U1, U32, U32, )
326
327// Atomic operations
328OPCODE(SharedAtomicIAdd32, U32, U32, U32, )
329OPCODE(SharedAtomicSMin32, U32, U32, U32, )
330OPCODE(SharedAtomicUMin32, U32, U32, U32, )
331OPCODE(SharedAtomicSMax32, U32, U32, U32, )
332OPCODE(SharedAtomicUMax32, U32, U32, U32, )
333OPCODE(SharedAtomicInc32, U32, U32, U32, )
334OPCODE(SharedAtomicDec32, U32, U32, U32, )
335OPCODE(SharedAtomicAnd32, U32, U32, U32, )
336OPCODE(SharedAtomicOr32, U32, U32, U32, )
337OPCODE(SharedAtomicXor32, U32, U32, U32, )
338OPCODE(SharedAtomicExchange32, U32, U32, U32, )
339OPCODE(SharedAtomicExchange64, U64, U32, U64, )
340
341OPCODE(GlobalAtomicIAdd32, U32, U64, U32, )
342OPCODE(GlobalAtomicSMin32, U32, U64, U32, )
343OPCODE(GlobalAtomicUMin32, U32, U64, U32, )
344OPCODE(GlobalAtomicSMax32, U32, U64, U32, )
345OPCODE(GlobalAtomicUMax32, U32, U64, U32, )
346OPCODE(GlobalAtomicInc32, U32, U64, U32, )
347OPCODE(GlobalAtomicDec32, U32, U64, U32, )
348OPCODE(GlobalAtomicAnd32, U32, U64, U32, )
349OPCODE(GlobalAtomicOr32, U32, U64, U32, )
350OPCODE(GlobalAtomicXor32, U32, U64, U32, )
351OPCODE(GlobalAtomicExchange32, U32, U64, U32, )
352OPCODE(GlobalAtomicIAdd64, U64, U64, U64, )
353OPCODE(GlobalAtomicSMin64, U64, U64, U64, )
354OPCODE(GlobalAtomicUMin64, U64, U64, U64, )
355OPCODE(GlobalAtomicSMax64, U64, U64, U64, )
356OPCODE(GlobalAtomicUMax64, U64, U64, U64, )
357OPCODE(GlobalAtomicAnd64, U64, U64, U64, )
358OPCODE(GlobalAtomicOr64, U64, U64, U64, )
359OPCODE(GlobalAtomicXor64, U64, U64, U64, )
360OPCODE(GlobalAtomicExchange64, U64, U64, U64, )
361OPCODE(GlobalAtomicAddF32, F32, U64, F32, )
362OPCODE(GlobalAtomicAddF16x2, U32, U64, F16x2, )
363OPCODE(GlobalAtomicAddF32x2, U32, U64, F32x2, )
364OPCODE(GlobalAtomicMinF16x2, U32, U64, F16x2, )
365OPCODE(GlobalAtomicMinF32x2, U32, U64, F32x2, )
366OPCODE(GlobalAtomicMaxF16x2, U32, U64, F16x2, )
367OPCODE(GlobalAtomicMaxF32x2, U32, U64, F32x2, )
368
369OPCODE(StorageAtomicIAdd32, U32, U32, U32, U32, )
370OPCODE(StorageAtomicSMin32, U32, U32, U32, U32, )
371OPCODE(StorageAtomicUMin32, U32, U32, U32, U32, )
372OPCODE(StorageAtomicSMax32, U32, U32, U32, U32, )
373OPCODE(StorageAtomicUMax32, U32, U32, U32, U32, )
374OPCODE(StorageAtomicInc32, U32, U32, U32, U32, )
375OPCODE(StorageAtomicDec32, U32, U32, U32, U32, )
376OPCODE(StorageAtomicAnd32, U32, U32, U32, U32, )
377OPCODE(StorageAtomicOr32, U32, U32, U32, U32, )
378OPCODE(StorageAtomicXor32, U32, U32, U32, U32, )
379OPCODE(StorageAtomicExchange32, U32, U32, U32, U32, )
380OPCODE(StorageAtomicIAdd64, U64, U32, U32, U64, )
381OPCODE(StorageAtomicSMin64, U64, U32, U32, U64, )
382OPCODE(StorageAtomicUMin64, U64, U32, U32, U64, )
383OPCODE(StorageAtomicSMax64, U64, U32, U32, U64, )
384OPCODE(StorageAtomicUMax64, U64, U32, U32, U64, )
385OPCODE(StorageAtomicAnd64, U64, U32, U32, U64, )
386OPCODE(StorageAtomicOr64, U64, U32, U32, U64, )
387OPCODE(StorageAtomicXor64, U64, U32, U32, U64, )
388OPCODE(StorageAtomicExchange64, U64, U32, U32, U64, )
389OPCODE(StorageAtomicAddF32, F32, U32, U32, F32, )
390OPCODE(StorageAtomicAddF16x2, U32, U32, U32, F16x2, )
391OPCODE(StorageAtomicAddF32x2, U32, U32, U32, F32x2, )
392OPCODE(StorageAtomicMinF16x2, U32, U32, U32, F16x2, )
393OPCODE(StorageAtomicMinF32x2, U32, U32, U32, F32x2, )
394OPCODE(StorageAtomicMaxF16x2, U32, U32, U32, F16x2, )
395OPCODE(StorageAtomicMaxF32x2, U32, U32, U32, F32x2, )
396
397// Logical operations
398OPCODE(LogicalOr, U1, U1, U1, )
399OPCODE(LogicalAnd, U1, U1, U1, )
400OPCODE(LogicalXor, U1, U1, U1, )
401OPCODE(LogicalNot, U1, U1, )
402
403// Conversion operations
404OPCODE(ConvertS16F16, U32, F16, )
405OPCODE(ConvertS16F32, U32, F32, )
406OPCODE(ConvertS16F64, U32, F64, )
407OPCODE(ConvertS32F16, U32, F16, )
408OPCODE(ConvertS32F32, U32, F32, )
409OPCODE(ConvertS32F64, U32, F64, )
410OPCODE(ConvertS64F16, U64, F16, )
411OPCODE(ConvertS64F32, U64, F32, )
412OPCODE(ConvertS64F64, U64, F64, )
413OPCODE(ConvertU16F16, U32, F16, )
414OPCODE(ConvertU16F32, U32, F32, )
415OPCODE(ConvertU16F64, U32, F64, )
416OPCODE(ConvertU32F16, U32, F16, )
417OPCODE(ConvertU32F32, U32, F32, )
418OPCODE(ConvertU32F64, U32, F64, )
419OPCODE(ConvertU64F16, U64, F16, )
420OPCODE(ConvertU64F32, U64, F32, )
421OPCODE(ConvertU64F64, U64, F64, )
422OPCODE(ConvertU64U32, U64, U32, )
423OPCODE(ConvertU32U64, U32, U64, )
424OPCODE(ConvertF16F32, F16, F32, )
425OPCODE(ConvertF32F16, F32, F16, )
426OPCODE(ConvertF32F64, F32, F64, )
427OPCODE(ConvertF64F32, F64, F32, )
428OPCODE(ConvertF16S8, F16, U32, )
429OPCODE(ConvertF16S16, F16, U32, )
430OPCODE(ConvertF16S32, F16, U32, )
431OPCODE(ConvertF16S64, F16, U64, )
432OPCODE(ConvertF16U8, F16, U32, )
433OPCODE(ConvertF16U16, F16, U32, )
434OPCODE(ConvertF16U32, F16, U32, )
435OPCODE(ConvertF16U64, F16, U64, )
436OPCODE(ConvertF32S8, F32, U32, )
437OPCODE(ConvertF32S16, F32, U32, )
438OPCODE(ConvertF32S32, F32, U32, )
439OPCODE(ConvertF32S64, F32, U64, )
440OPCODE(ConvertF32U8, F32, U32, )
441OPCODE(ConvertF32U16, F32, U32, )
442OPCODE(ConvertF32U32, F32, U32, )
443OPCODE(ConvertF32U64, F32, U64, )
444OPCODE(ConvertF64S8, F64, U32, )
445OPCODE(ConvertF64S16, F64, U32, )
446OPCODE(ConvertF64S32, F64, U32, )
447OPCODE(ConvertF64S64, F64, U64, )
448OPCODE(ConvertF64U8, F64, U32, )
449OPCODE(ConvertF64U16, F64, U32, )
450OPCODE(ConvertF64U32, F64, U32, )
451OPCODE(ConvertF64U64, F64, U64, )
452
453// Image operations
454OPCODE(BindlessImageSampleImplicitLod, F32x4, U32, Opaque, Opaque, Opaque, )
455OPCODE(BindlessImageSampleExplicitLod, F32x4, U32, Opaque, Opaque, Opaque, )
456OPCODE(BindlessImageSampleDrefImplicitLod, F32, U32, Opaque, F32, Opaque, Opaque, )
457OPCODE(BindlessImageSampleDrefExplicitLod, F32, U32, Opaque, F32, Opaque, Opaque, )
458OPCODE(BindlessImageGather, F32x4, U32, Opaque, Opaque, Opaque, )
459OPCODE(BindlessImageGatherDref, F32x4, U32, Opaque, Opaque, Opaque, F32, )
460OPCODE(BindlessImageFetch, F32x4, U32, Opaque, Opaque, U32, Opaque, )
461OPCODE(BindlessImageQueryDimensions, U32x4, U32, U32, )
462OPCODE(BindlessImageQueryLod, F32x4, U32, Opaque, )
463OPCODE(BindlessImageGradient, F32x4, U32, Opaque, Opaque, Opaque, Opaque, )
464OPCODE(BindlessImageRead, U32x4, U32, Opaque, )
465OPCODE(BindlessImageWrite, Void, U32, Opaque, U32x4, )
466
467OPCODE(BoundImageSampleImplicitLod, F32x4, U32, Opaque, Opaque, Opaque, )
468OPCODE(BoundImageSampleExplicitLod, F32x4, U32, Opaque, Opaque, Opaque, )
469OPCODE(BoundImageSampleDrefImplicitLod, F32, U32, Opaque, F32, Opaque, Opaque, )
470OPCODE(BoundImageSampleDrefExplicitLod, F32, U32, Opaque, F32, Opaque, Opaque, )
471OPCODE(BoundImageGather, F32x4, U32, Opaque, Opaque, Opaque, )
472OPCODE(BoundImageGatherDref, F32x4, U32, Opaque, Opaque, Opaque, F32, )
473OPCODE(BoundImageFetch, F32x4, U32, Opaque, Opaque, U32, Opaque, )
474OPCODE(BoundImageQueryDimensions, U32x4, U32, U32, )
475OPCODE(BoundImageQueryLod, F32x4, U32, Opaque, )
476OPCODE(BoundImageGradient, F32x4, U32, Opaque, Opaque, Opaque, Opaque, )
477OPCODE(BoundImageRead, U32x4, U32, Opaque, )
478OPCODE(BoundImageWrite, Void, U32, Opaque, U32x4, )
479
480OPCODE(ImageSampleImplicitLod, F32x4, Opaque, Opaque, Opaque, Opaque, )
481OPCODE(ImageSampleExplicitLod, F32x4, Opaque, Opaque, Opaque, Opaque, )
482OPCODE(ImageSampleDrefImplicitLod, F32, Opaque, Opaque, F32, Opaque, Opaque, )
483OPCODE(ImageSampleDrefExplicitLod, F32, Opaque, Opaque, F32, Opaque, Opaque, )
484OPCODE(ImageGather, F32x4, Opaque, Opaque, Opaque, Opaque, )
485OPCODE(ImageGatherDref, F32x4, Opaque, Opaque, Opaque, Opaque, F32, )
486OPCODE(ImageFetch, F32x4, Opaque, Opaque, Opaque, U32, Opaque, )
487OPCODE(ImageQueryDimensions, U32x4, Opaque, U32, )
488OPCODE(ImageQueryLod, F32x4, Opaque, Opaque, )
489OPCODE(ImageGradient, F32x4, Opaque, Opaque, Opaque, Opaque, Opaque, )
490OPCODE(ImageRead, U32x4, Opaque, Opaque, )
491OPCODE(ImageWrite, Void, Opaque, Opaque, U32x4, )
492
493// Atomic Image operations
494
495OPCODE(BindlessImageAtomicIAdd32, U32, U32, Opaque, U32, )
496OPCODE(BindlessImageAtomicSMin32, U32, U32, Opaque, U32, )
497OPCODE(BindlessImageAtomicUMin32, U32, U32, Opaque, U32, )
498OPCODE(BindlessImageAtomicSMax32, U32, U32, Opaque, U32, )
499OPCODE(BindlessImageAtomicUMax32, U32, U32, Opaque, U32, )
500OPCODE(BindlessImageAtomicInc32, U32, U32, Opaque, U32, )
501OPCODE(BindlessImageAtomicDec32, U32, U32, Opaque, U32, )
502OPCODE(BindlessImageAtomicAnd32, U32, U32, Opaque, U32, )
503OPCODE(BindlessImageAtomicOr32, U32, U32, Opaque, U32, )
504OPCODE(BindlessImageAtomicXor32, U32, U32, Opaque, U32, )
505OPCODE(BindlessImageAtomicExchange32, U32, U32, Opaque, U32, )
506
507OPCODE(BoundImageAtomicIAdd32, U32, U32, Opaque, U32, )
508OPCODE(BoundImageAtomicSMin32, U32, U32, Opaque, U32, )
509OPCODE(BoundImageAtomicUMin32, U32, U32, Opaque, U32, )
510OPCODE(BoundImageAtomicSMax32, U32, U32, Opaque, U32, )
511OPCODE(BoundImageAtomicUMax32, U32, U32, Opaque, U32, )
512OPCODE(BoundImageAtomicInc32, U32, U32, Opaque, U32, )
513OPCODE(BoundImageAtomicDec32, U32, U32, Opaque, U32, )
514OPCODE(BoundImageAtomicAnd32, U32, U32, Opaque, U32, )
515OPCODE(BoundImageAtomicOr32, U32, U32, Opaque, U32, )
516OPCODE(BoundImageAtomicXor32, U32, U32, Opaque, U32, )
517OPCODE(BoundImageAtomicExchange32, U32, U32, Opaque, U32, )
518
519OPCODE(ImageAtomicIAdd32, U32, Opaque, Opaque, U32, )
520OPCODE(ImageAtomicSMin32, U32, Opaque, Opaque, U32, )
521OPCODE(ImageAtomicUMin32, U32, Opaque, Opaque, U32, )
522OPCODE(ImageAtomicSMax32, U32, Opaque, Opaque, U32, )
523OPCODE(ImageAtomicUMax32, U32, Opaque, Opaque, U32, )
524OPCODE(ImageAtomicInc32, U32, Opaque, Opaque, U32, )
525OPCODE(ImageAtomicDec32, U32, Opaque, Opaque, U32, )
526OPCODE(ImageAtomicAnd32, U32, Opaque, Opaque, U32, )
527OPCODE(ImageAtomicOr32, U32, Opaque, Opaque, U32, )
528OPCODE(ImageAtomicXor32, U32, Opaque, Opaque, U32, )
529OPCODE(ImageAtomicExchange32, U32, Opaque, Opaque, U32, )
530
531// Warp operations
532OPCODE(LaneId, U32, )
533OPCODE(VoteAll, U1, U1, )
534OPCODE(VoteAny, U1, U1, )
535OPCODE(VoteEqual, U1, U1, )
536OPCODE(SubgroupBallot, U32, U1, )
537OPCODE(SubgroupEqMask, U32, )
538OPCODE(SubgroupLtMask, U32, )
539OPCODE(SubgroupLeMask, U32, )
540OPCODE(SubgroupGtMask, U32, )
541OPCODE(SubgroupGeMask, U32, )
542OPCODE(ShuffleIndex, U32, U32, U32, U32, U32, )
543OPCODE(ShuffleUp, U32, U32, U32, U32, U32, )
544OPCODE(ShuffleDown, U32, U32, U32, U32, U32, )
545OPCODE(ShuffleButterfly, U32, U32, U32, U32, U32, )
546OPCODE(FSwizzleAdd, F32, F32, F32, U32, )
547OPCODE(DPdxFine, F32, F32, )
548OPCODE(DPdyFine, F32, F32, )
549OPCODE(DPdxCoarse, F32, F32, )
550OPCODE(DPdyCoarse, F32, F32, )
diff --git a/src/shader_recompiler/frontend/ir/patch.cpp b/src/shader_recompiler/frontend/ir/patch.cpp
new file mode 100644
index 000000000..4c956a970
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/patch.cpp
@@ -0,0 +1,28 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include "shader_recompiler/exception.h"
6#include "shader_recompiler/frontend/ir/patch.h"
7
8namespace Shader::IR {
9
10bool IsGeneric(Patch patch) noexcept {
11 return patch >= Patch::Component0 && patch <= Patch::Component119;
12}
13
14u32 GenericPatchIndex(Patch patch) {
15 if (!IsGeneric(patch)) {
16 throw InvalidArgument("Patch {} is not generic", patch);
17 }
18 return (static_cast<u32>(patch) - static_cast<u32>(Patch::Component0)) / 4;
19}
20
21u32 GenericPatchElement(Patch patch) {
22 if (!IsGeneric(patch)) {
23 throw InvalidArgument("Patch {} is not generic", patch);
24 }
25 return (static_cast<u32>(patch) - static_cast<u32>(Patch::Component0)) % 4;
26}
27
28} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/patch.h b/src/shader_recompiler/frontend/ir/patch.h
new file mode 100644
index 000000000..6d66ff0d6
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/patch.h
@@ -0,0 +1,149 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include "common/common_types.h"
8
9namespace Shader::IR {
10
11enum class Patch : u64 {
12 TessellationLodLeft,
13 TessellationLodTop,
14 TessellationLodRight,
15 TessellationLodBottom,
16 TessellationLodInteriorU,
17 TessellationLodInteriorV,
18 ComponentPadding0,
19 ComponentPadding1,
20 Component0,
21 Component1,
22 Component2,
23 Component3,
24 Component4,
25 Component5,
26 Component6,
27 Component7,
28 Component8,
29 Component9,
30 Component10,
31 Component11,
32 Component12,
33 Component13,
34 Component14,
35 Component15,
36 Component16,
37 Component17,
38 Component18,
39 Component19,
40 Component20,
41 Component21,
42 Component22,
43 Component23,
44 Component24,
45 Component25,
46 Component26,
47 Component27,
48 Component28,
49 Component29,
50 Component30,
51 Component31,
52 Component32,
53 Component33,
54 Component34,
55 Component35,
56 Component36,
57 Component37,
58 Component38,
59 Component39,
60 Component40,
61 Component41,
62 Component42,
63 Component43,
64 Component44,
65 Component45,
66 Component46,
67 Component47,
68 Component48,
69 Component49,
70 Component50,
71 Component51,
72 Component52,
73 Component53,
74 Component54,
75 Component55,
76 Component56,
77 Component57,
78 Component58,
79 Component59,
80 Component60,
81 Component61,
82 Component62,
83 Component63,
84 Component64,
85 Component65,
86 Component66,
87 Component67,
88 Component68,
89 Component69,
90 Component70,
91 Component71,
92 Component72,
93 Component73,
94 Component74,
95 Component75,
96 Component76,
97 Component77,
98 Component78,
99 Component79,
100 Component80,
101 Component81,
102 Component82,
103 Component83,
104 Component84,
105 Component85,
106 Component86,
107 Component87,
108 Component88,
109 Component89,
110 Component90,
111 Component91,
112 Component92,
113 Component93,
114 Component94,
115 Component95,
116 Component96,
117 Component97,
118 Component98,
119 Component99,
120 Component100,
121 Component101,
122 Component102,
123 Component103,
124 Component104,
125 Component105,
126 Component106,
127 Component107,
128 Component108,
129 Component109,
130 Component110,
131 Component111,
132 Component112,
133 Component113,
134 Component114,
135 Component115,
136 Component116,
137 Component117,
138 Component118,
139 Component119,
140};
141static_assert(static_cast<u64>(Patch::Component119) == 127);
142
143[[nodiscard]] bool IsGeneric(Patch patch) noexcept;
144
145[[nodiscard]] u32 GenericPatchIndex(Patch patch);
146
147[[nodiscard]] u32 GenericPatchElement(Patch patch);
148
149} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/post_order.cpp b/src/shader_recompiler/frontend/ir/post_order.cpp
new file mode 100644
index 000000000..16bc44101
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/post_order.cpp
@@ -0,0 +1,46 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <algorithm>
6
7#include <boost/container/flat_set.hpp>
8#include <boost/container/small_vector.hpp>
9
10#include "shader_recompiler/frontend/ir/basic_block.h"
11#include "shader_recompiler/frontend/ir/post_order.h"
12
13namespace Shader::IR {
14
15BlockList PostOrder(const AbstractSyntaxNode& root) {
16 boost::container::small_vector<Block*, 16> block_stack;
17 boost::container::flat_set<Block*> visited;
18 BlockList post_order_blocks;
19
20 if (root.type != AbstractSyntaxNode::Type::Block) {
21 throw LogicError("First node in abstract syntax list root is not a block");
22 }
23 Block* const first_block{root.data.block};
24 visited.insert(first_block);
25 block_stack.push_back(first_block);
26
27 while (!block_stack.empty()) {
28 Block* const block{block_stack.back()};
29 const auto visit{[&](Block* branch) {
30 if (!visited.insert(branch).second) {
31 return false;
32 }
33 // Calling push_back twice is faster than insert on MSVC
34 block_stack.push_back(block);
35 block_stack.push_back(branch);
36 return true;
37 }};
38 block_stack.pop_back();
39 if (std::ranges::none_of(block->ImmSuccessors(), visit)) {
40 post_order_blocks.push_back(block);
41 }
42 }
43 return post_order_blocks;
44}
45
46} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/post_order.h b/src/shader_recompiler/frontend/ir/post_order.h
new file mode 100644
index 000000000..07bfbadc3
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/post_order.h
@@ -0,0 +1,14 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include "shader_recompiler/frontend/ir/abstract_syntax_list.h"
8#include "shader_recompiler/frontend/ir/basic_block.h"
9
10namespace Shader::IR {
11
12BlockList PostOrder(const AbstractSyntaxNode& root);
13
14} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/pred.h b/src/shader_recompiler/frontend/ir/pred.h
new file mode 100644
index 000000000..4e7f32423
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/pred.h
@@ -0,0 +1,44 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <fmt/format.h>
8
9namespace Shader::IR {
10
11enum class Pred : u64 {
12 P0,
13 P1,
14 P2,
15 P3,
16 P4,
17 P5,
18 P6,
19 PT,
20};
21
22constexpr size_t NUM_USER_PREDS = 7;
23constexpr size_t NUM_PREDS = 8;
24
25[[nodiscard]] constexpr size_t PredIndex(Pred pred) noexcept {
26 return static_cast<size_t>(pred);
27}
28
29} // namespace Shader::IR
30
31template <>
32struct fmt::formatter<Shader::IR::Pred> {
33 constexpr auto parse(format_parse_context& ctx) {
34 return ctx.begin();
35 }
36 template <typename FormatContext>
37 auto format(const Shader::IR::Pred& pred, FormatContext& ctx) {
38 if (pred == Shader::IR::Pred::PT) {
39 return fmt::format_to(ctx.out(), "PT");
40 } else {
41 return fmt::format_to(ctx.out(), "P{}", static_cast<int>(pred));
42 }
43 }
44};
diff --git a/src/shader_recompiler/frontend/ir/program.cpp b/src/shader_recompiler/frontend/ir/program.cpp
new file mode 100644
index 000000000..3fc06f855
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/program.cpp
@@ -0,0 +1,32 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <map>
6#include <string>
7
8#include <fmt/format.h>
9
10#include "shader_recompiler/frontend/ir/basic_block.h"
11#include "shader_recompiler/frontend/ir/program.h"
12#include "shader_recompiler/frontend/ir/value.h"
13
14namespace Shader::IR {
15
16std::string DumpProgram(const Program& program) {
17 size_t index{0};
18 std::map<const IR::Inst*, size_t> inst_to_index;
19 std::map<const IR::Block*, size_t> block_to_index;
20
21 for (const IR::Block* const block : program.blocks) {
22 block_to_index.emplace(block, index);
23 ++index;
24 }
25 std::string ret;
26 for (const auto& block : program.blocks) {
27 ret += IR::DumpBlock(*block, block_to_index, inst_to_index, index) + '\n';
28 }
29 return ret;
30}
31
32} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/program.h b/src/shader_recompiler/frontend/ir/program.h
new file mode 100644
index 000000000..ebcaa8bc2
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/program.h
@@ -0,0 +1,35 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <array>
8#include <string>
9
10#include "shader_recompiler/frontend/ir/abstract_syntax_list.h"
11#include "shader_recompiler/frontend/ir/basic_block.h"
12#include "shader_recompiler/program_header.h"
13#include "shader_recompiler/shader_info.h"
14#include "shader_recompiler/stage.h"
15
16namespace Shader::IR {
17
18struct Program {
19 AbstractSyntaxList syntax_list;
20 BlockList blocks;
21 BlockList post_order_blocks;
22 Info info;
23 Stage stage{};
24 std::array<u32, 3> workgroup_size{};
25 OutputTopology output_topology{};
26 u32 output_vertices{};
27 u32 invocations{};
28 u32 local_memory_size{};
29 u32 shared_memory_size{};
30 bool is_geometry_passthrough{};
31};
32
33[[nodiscard]] std::string DumpProgram(const Program& program);
34
35} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/reg.h b/src/shader_recompiler/frontend/ir/reg.h
new file mode 100644
index 000000000..a4b635792
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/reg.h
@@ -0,0 +1,332 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <fmt/format.h>
8
9#include "common/common_types.h"
10#include "shader_recompiler/exception.h"
11
12namespace Shader::IR {
13
14enum class Reg : u64 {
15 R0,
16 R1,
17 R2,
18 R3,
19 R4,
20 R5,
21 R6,
22 R7,
23 R8,
24 R9,
25 R10,
26 R11,
27 R12,
28 R13,
29 R14,
30 R15,
31 R16,
32 R17,
33 R18,
34 R19,
35 R20,
36 R21,
37 R22,
38 R23,
39 R24,
40 R25,
41 R26,
42 R27,
43 R28,
44 R29,
45 R30,
46 R31,
47 R32,
48 R33,
49 R34,
50 R35,
51 R36,
52 R37,
53 R38,
54 R39,
55 R40,
56 R41,
57 R42,
58 R43,
59 R44,
60 R45,
61 R46,
62 R47,
63 R48,
64 R49,
65 R50,
66 R51,
67 R52,
68 R53,
69 R54,
70 R55,
71 R56,
72 R57,
73 R58,
74 R59,
75 R60,
76 R61,
77 R62,
78 R63,
79 R64,
80 R65,
81 R66,
82 R67,
83 R68,
84 R69,
85 R70,
86 R71,
87 R72,
88 R73,
89 R74,
90 R75,
91 R76,
92 R77,
93 R78,
94 R79,
95 R80,
96 R81,
97 R82,
98 R83,
99 R84,
100 R85,
101 R86,
102 R87,
103 R88,
104 R89,
105 R90,
106 R91,
107 R92,
108 R93,
109 R94,
110 R95,
111 R96,
112 R97,
113 R98,
114 R99,
115 R100,
116 R101,
117 R102,
118 R103,
119 R104,
120 R105,
121 R106,
122 R107,
123 R108,
124 R109,
125 R110,
126 R111,
127 R112,
128 R113,
129 R114,
130 R115,
131 R116,
132 R117,
133 R118,
134 R119,
135 R120,
136 R121,
137 R122,
138 R123,
139 R124,
140 R125,
141 R126,
142 R127,
143 R128,
144 R129,
145 R130,
146 R131,
147 R132,
148 R133,
149 R134,
150 R135,
151 R136,
152 R137,
153 R138,
154 R139,
155 R140,
156 R141,
157 R142,
158 R143,
159 R144,
160 R145,
161 R146,
162 R147,
163 R148,
164 R149,
165 R150,
166 R151,
167 R152,
168 R153,
169 R154,
170 R155,
171 R156,
172 R157,
173 R158,
174 R159,
175 R160,
176 R161,
177 R162,
178 R163,
179 R164,
180 R165,
181 R166,
182 R167,
183 R168,
184 R169,
185 R170,
186 R171,
187 R172,
188 R173,
189 R174,
190 R175,
191 R176,
192 R177,
193 R178,
194 R179,
195 R180,
196 R181,
197 R182,
198 R183,
199 R184,
200 R185,
201 R186,
202 R187,
203 R188,
204 R189,
205 R190,
206 R191,
207 R192,
208 R193,
209 R194,
210 R195,
211 R196,
212 R197,
213 R198,
214 R199,
215 R200,
216 R201,
217 R202,
218 R203,
219 R204,
220 R205,
221 R206,
222 R207,
223 R208,
224 R209,
225 R210,
226 R211,
227 R212,
228 R213,
229 R214,
230 R215,
231 R216,
232 R217,
233 R218,
234 R219,
235 R220,
236 R221,
237 R222,
238 R223,
239 R224,
240 R225,
241 R226,
242 R227,
243 R228,
244 R229,
245 R230,
246 R231,
247 R232,
248 R233,
249 R234,
250 R235,
251 R236,
252 R237,
253 R238,
254 R239,
255 R240,
256 R241,
257 R242,
258 R243,
259 R244,
260 R245,
261 R246,
262 R247,
263 R248,
264 R249,
265 R250,
266 R251,
267 R252,
268 R253,
269 R254,
270 RZ,
271};
272static_assert(static_cast<int>(Reg::RZ) == 255);
273
274constexpr size_t NUM_USER_REGS = 255;
275constexpr size_t NUM_REGS = 256;
276
277[[nodiscard]] constexpr Reg operator+(Reg reg, int num) {
278 if (reg == Reg::RZ) {
279 // Adding or subtracting registers from RZ yields RZ
280 return Reg::RZ;
281 }
282 const int result{static_cast<int>(reg) + num};
283 if (result >= static_cast<int>(Reg::RZ)) {
284 throw LogicError("Overflow on register arithmetic");
285 }
286 if (result < 0) {
287 throw LogicError("Underflow on register arithmetic");
288 }
289 return static_cast<Reg>(result);
290}
291
292[[nodiscard]] constexpr Reg operator-(Reg reg, int num) {
293 return reg + (-num);
294}
295
296constexpr Reg operator++(Reg& reg) {
297 reg = reg + 1;
298 return reg;
299}
300
301constexpr Reg operator++(Reg& reg, int) {
302 const Reg copy{reg};
303 reg = reg + 1;
304 return copy;
305}
306
307[[nodiscard]] constexpr size_t RegIndex(Reg reg) noexcept {
308 return static_cast<size_t>(reg);
309}
310
311[[nodiscard]] constexpr bool IsAligned(Reg reg, size_t align) {
312 return RegIndex(reg) % align == 0 || reg == Reg::RZ;
313}
314
315} // namespace Shader::IR
316
317template <>
318struct fmt::formatter<Shader::IR::Reg> {
319 constexpr auto parse(format_parse_context& ctx) {
320 return ctx.begin();
321 }
322 template <typename FormatContext>
323 auto format(const Shader::IR::Reg& reg, FormatContext& ctx) {
324 if (reg == Shader::IR::Reg::RZ) {
325 return fmt::format_to(ctx.out(), "RZ");
326 } else if (static_cast<int>(reg) >= 0 && static_cast<int>(reg) < 255) {
327 return fmt::format_to(ctx.out(), "R{}", static_cast<int>(reg));
328 } else {
329 throw Shader::LogicError("Invalid register with raw value {}", static_cast<int>(reg));
330 }
331 }
332};
diff --git a/src/shader_recompiler/frontend/ir/type.cpp b/src/shader_recompiler/frontend/ir/type.cpp
new file mode 100644
index 000000000..f28341bfe
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/type.cpp
@@ -0,0 +1,38 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include <array>
6#include <string>
7
8#include "shader_recompiler/frontend/ir/type.h"
9
10namespace Shader::IR {
11
12std::string NameOf(Type type) {
13 static constexpr std::array names{
14 "Opaque", "Label", "Reg", "Pred", "Attribute", "U1", "U8", "U16", "U32",
15 "U64", "F16", "F32", "F64", "U32x2", "U32x3", "U32x4", "F16x2", "F16x3",
16 "F16x4", "F32x2", "F32x3", "F32x4", "F64x2", "F64x3", "F64x4",
17 };
18 const size_t bits{static_cast<size_t>(type)};
19 if (bits == 0) {
20 return "Void";
21 }
22 std::string result;
23 for (size_t i = 0; i < names.size(); i++) {
24 if ((bits & (size_t{1} << i)) != 0) {
25 if (!result.empty()) {
26 result += '|';
27 }
28 result += names[i];
29 }
30 }
31 return result;
32}
33
34bool AreTypesCompatible(Type lhs, Type rhs) noexcept {
35 return lhs == rhs || lhs == Type::Opaque || rhs == Type::Opaque;
36}
37
38} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/type.h b/src/shader_recompiler/frontend/ir/type.h
new file mode 100644
index 000000000..294b230c4
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/type.h
@@ -0,0 +1,61 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <string>
8
9#include <fmt/format.h>
10
11#include "common/common_funcs.h"
12#include "shader_recompiler/exception.h"
13
14namespace Shader::IR {
15
16enum class Type {
17 Void = 0,
18 Opaque = 1 << 0,
19 Reg = 1 << 1,
20 Pred = 1 << 2,
21 Attribute = 1 << 3,
22 Patch = 1 << 4,
23 U1 = 1 << 5,
24 U8 = 1 << 6,
25 U16 = 1 << 7,
26 U32 = 1 << 8,
27 U64 = 1 << 9,
28 F16 = 1 << 10,
29 F32 = 1 << 11,
30 F64 = 1 << 12,
31 U32x2 = 1 << 13,
32 U32x3 = 1 << 14,
33 U32x4 = 1 << 15,
34 F16x2 = 1 << 16,
35 F16x3 = 1 << 17,
36 F16x4 = 1 << 18,
37 F32x2 = 1 << 19,
38 F32x3 = 1 << 20,
39 F32x4 = 1 << 21,
40 F64x2 = 1 << 22,
41 F64x3 = 1 << 23,
42 F64x4 = 1 << 24,
43};
44DECLARE_ENUM_FLAG_OPERATORS(Type)
45
46[[nodiscard]] std::string NameOf(Type type);
47
48[[nodiscard]] bool AreTypesCompatible(Type lhs, Type rhs) noexcept;
49
50} // namespace Shader::IR
51
52template <>
53struct fmt::formatter<Shader::IR::Type> {
54 constexpr auto parse(format_parse_context& ctx) {
55 return ctx.begin();
56 }
57 template <typename FormatContext>
58 auto format(const Shader::IR::Type& type, FormatContext& ctx) {
59 return fmt::format_to(ctx.out(), "{}", NameOf(type));
60 }
61};
diff --git a/src/shader_recompiler/frontend/ir/value.cpp b/src/shader_recompiler/frontend/ir/value.cpp
new file mode 100644
index 000000000..d365ea1bc
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/value.cpp
@@ -0,0 +1,99 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#include "shader_recompiler/frontend/ir/opcodes.h"
6#include "shader_recompiler/frontend/ir/value.h"
7
8namespace Shader::IR {
9
10Value::Value(IR::Inst* value) noexcept : type{Type::Opaque}, inst{value} {}
11
12Value::Value(IR::Reg value) noexcept : type{Type::Reg}, reg{value} {}
13
14Value::Value(IR::Pred value) noexcept : type{Type::Pred}, pred{value} {}
15
16Value::Value(IR::Attribute value) noexcept : type{Type::Attribute}, attribute{value} {}
17
18Value::Value(IR::Patch value) noexcept : type{Type::Patch}, patch{value} {}
19
20Value::Value(bool value) noexcept : type{Type::U1}, imm_u1{value} {}
21
22Value::Value(u8 value) noexcept : type{Type::U8}, imm_u8{value} {}
23
24Value::Value(u16 value) noexcept : type{Type::U16}, imm_u16{value} {}
25
26Value::Value(u32 value) noexcept : type{Type::U32}, imm_u32{value} {}
27
28Value::Value(f32 value) noexcept : type{Type::F32}, imm_f32{value} {}
29
30Value::Value(u64 value) noexcept : type{Type::U64}, imm_u64{value} {}
31
32Value::Value(f64 value) noexcept : type{Type::F64}, imm_f64{value} {}
33
34IR::Type Value::Type() const noexcept {
35 if (IsPhi()) {
36 // The type of a phi node is stored in its flags
37 return inst->Flags<IR::Type>();
38 }
39 if (IsIdentity()) {
40 return inst->Arg(0).Type();
41 }
42 if (type == Type::Opaque) {
43 return inst->Type();
44 }
45 return type;
46}
47
48bool Value::operator==(const Value& other) const {
49 if (type != other.type) {
50 return false;
51 }
52 switch (type) {
53 case Type::Void:
54 return true;
55 case Type::Opaque:
56 return inst == other.inst;
57 case Type::Reg:
58 return reg == other.reg;
59 case Type::Pred:
60 return pred == other.pred;
61 case Type::Attribute:
62 return attribute == other.attribute;
63 case Type::Patch:
64 return patch == other.patch;
65 case Type::U1:
66 return imm_u1 == other.imm_u1;
67 case Type::U8:
68 return imm_u8 == other.imm_u8;
69 case Type::U16:
70 case Type::F16:
71 return imm_u16 == other.imm_u16;
72 case Type::U32:
73 case Type::F32:
74 return imm_u32 == other.imm_u32;
75 case Type::U64:
76 case Type::F64:
77 return imm_u64 == other.imm_u64;
78 case Type::U32x2:
79 case Type::U32x3:
80 case Type::U32x4:
81 case Type::F16x2:
82 case Type::F16x3:
83 case Type::F16x4:
84 case Type::F32x2:
85 case Type::F32x3:
86 case Type::F32x4:
87 case Type::F64x2:
88 case Type::F64x3:
89 case Type::F64x4:
90 break;
91 }
92 throw LogicError("Invalid type {}", type);
93}
94
95bool Value::operator!=(const Value& other) const {
96 return !operator==(other);
97}
98
99} // namespace Shader::IR
diff --git a/src/shader_recompiler/frontend/ir/value.h b/src/shader_recompiler/frontend/ir/value.h
new file mode 100644
index 000000000..0c6bf684d
--- /dev/null
+++ b/src/shader_recompiler/frontend/ir/value.h
@@ -0,0 +1,398 @@
1// Copyright 2021 yuzu Emulator Project
2// Licensed under GPLv2 or any later version
3// Refer to the license.txt file included.
4
5#pragma once
6
7#include <array>
8#include <cstring>
9#include <memory>
10#include <type_traits>
11#include <utility>
12#include <vector>
13
14#include <boost/container/small_vector.hpp>
15#include <boost/intrusive/list.hpp>
16
17#include "common/assert.h"
18#include "common/bit_cast.h"
19#include "common/common_types.h"
20#include "shader_recompiler/exception.h"
21#include "shader_recompiler/frontend/ir/attribute.h"
22#include "shader_recompiler/frontend/ir/opcodes.h"
23#include "shader_recompiler/frontend/ir/patch.h"
24#include "shader_recompiler/frontend/ir/pred.h"
25#include "shader_recompiler/frontend/ir/reg.h"
26#include "shader_recompiler/frontend/ir/type.h"
27#include "shader_recompiler/frontend/ir/value.h"
28
29namespace Shader::IR {
30
31class Block;
32class Inst;
33
34struct AssociatedInsts;
35
36class Value {
37public:
38 Value() noexcept = default;
39 explicit Value(IR::Inst* value) noexcept;
40 explicit Value(IR::Reg value) noexcept;
41 explicit Value(IR::Pred value) noexcept;
42 explicit Value(IR::Attribute value) noexcept;
43 explicit Value(IR::Patch value) noexcept;
44 explicit Value(bool value) noexcept;
45 explicit Value(u8 value) noexcept;
46 explicit Value(u16 value) noexcept;
47 explicit Value(u32 value) noexcept;
48 explicit Value(f32 value) noexcept;
49 explicit Value(u64 value) noexcept;
50 explicit Value(f64 value) noexcept;
51
52 [[nodiscard]] bool IsIdentity() const noexcept;
53 [[nodiscard]] bool IsPhi() const noexcept;
54 [[nodiscard]] bool IsEmpty() const noexcept;
55 [[nodiscard]] bool IsImmediate() const noexcept;
56 [[nodiscard]] IR::Type Type() const noexcept;
57
58 [[nodiscard]] IR::Inst* Inst() const;
59 [[nodiscard]] IR::Inst* InstRecursive() const;
60 [[nodiscard]] IR::Value Resolve() const;
61 [[nodiscard]] IR::Reg Reg() const;
62 [[nodiscard]] IR::Pred Pred() const;
63 [[nodiscard]] IR::Attribute Attribute() const;
64 [[nodiscard]] IR::Patch Patch() const;
65 [[nodiscard]] bool U1() const;
66 [[nodiscard]] u8 U8() const;
67 [[nodiscard]] u16 U16() const;
68 [[nodiscard]] u32 U32() const;
69 [[nodiscard]] f32 F32() const;
70 [[nodiscard]] u64 U64() const;
71 [[nodiscard]] f64 F64() const;
72
73 [[nodiscard]] bool operator==(const Value& other) const;
74 [[nodiscard]] bool operator!=(const Value& other) const;
75
76private:
77 IR::Type type{};
78 union {
79 IR::Inst* inst{};
80 IR::Reg reg;
81 IR::Pred pred;
82 IR::Attribute attribute;
83 IR::Patch patch;
84 bool imm_u1;
85 u8 imm_u8;
86 u16 imm_u16;
87 u32 imm_u32;
88 f32 imm_f32;
89 u64 imm_u64;
90 f64 imm_f64;
91 };
92};
93static_assert(static_cast<u32>(IR::Type::Void) == 0, "memset relies on IR::Type being zero");
94static_assert(std::is_trivially_copyable_v<Value>);
95
96template <IR::Type type_>
97class TypedValue : public Value {
98public:
99 TypedValue() = default;
100
101 template <IR::Type other_type>
102 requires((other_type & type_) != IR::Type::Void) explicit(false)
103 TypedValue(const TypedValue<other_type>& value)
104 : Value(value) {}
105
106 explicit TypedValue(const Value& value) : Value(value) {
107 if ((value.Type() & type_) == IR::Type::Void) {
108 throw InvalidArgument("Incompatible types {} and {}", type_, value.Type());
109 }
110 }
111
112 explicit TypedValue(IR::Inst* inst_) : TypedValue(Value(inst_)) {}
113};
114
115class Inst : public boost::intrusive::list_base_hook<> {
116public:
117 explicit Inst(IR::Opcode op_, u32 flags_) noexcept;
118 ~Inst();
119
120 Inst& operator=(const Inst&) = delete;
121 Inst(const Inst&) = delete;
122
123 Inst& operator=(Inst&&) = delete;
124 Inst(Inst&&) = delete;
125
126 /// Get the number of uses this instruction has.
127 [[nodiscard]] int UseCount() const noexcept {
128 return use_count;
129 }
130
131 /// Determines whether this instruction has uses or not.
132 [[nodiscard]] bool HasUses() const noexcept {
133 return use_count > 0;
134 }
135
136 /// Get the opcode this microinstruction represents.
137 [[nodiscard]] IR::Opcode GetOpcode() const noexcept {
138 return op;
139 }
140
141 /// Determines if there is a pseudo-operation associated with this instruction.
142 [[nodiscard]] bool HasAssociatedPseudoOperation() const noexcept {
143 return associated_insts != nullptr;
144 }
145
146 /// Determines whether or not this instruction may have side effects.
147 [[nodiscard]] bool MayHaveSideEffects() const noexcept;
148
149 /// Determines whether or not this instruction is a pseudo-instruction.
150 /// Pseudo-instructions depend on their parent instructions for their semantics.
151 [[nodiscard]] bool IsPseudoInstruction() const noexcept;
152
153 /// Determines if all arguments of this instruction are immediates.
154 [[nodiscard]] bool AreAllArgsImmediates() const;
155
156 /// Gets a pseudo-operation associated with this instruction
157 [[nodiscard]] Inst* GetAssociatedPseudoOperation(IR::Opcode opcode);
158
159 /// Get the type this instruction returns.
160 [[nodiscard]] IR::Type Type() const;
161
162 /// Get the number of arguments this instruction has.
163 [[nodiscard]] size_t NumArgs() const {
164 return op == IR::Opcode::Phi ? phi_args.size() : NumArgsOf(op);
165 }
166
167 /// Get the value of a given argument index.
168 [[nodiscard]] Value Arg(size_t index) const noexcept {
169 if (op == IR::Opcode::Phi) {
170 return phi_args[index].second;
171 } else {
172 return args[index];
173 }
174 }
175
176 /// Set the value of a given argument index.
177 void SetArg(size_t index, Value value);
178
179 /// Get a pointer to the block of a phi argument.
180 [[nodiscard]] Block* PhiBlock(size_t index) const;
181 /// Add phi operand to a phi instruction.
182 void AddPhiOperand(Block* predecessor, const Value& value);
183
184 void Invalidate();
185 void ClearArgs();
186
187 void ReplaceUsesWith(Value replacement);
188
189 void ReplaceOpcode(IR::Opcode opcode);
190
191 template <typename FlagsType>
192 requires(sizeof(FlagsType) <= sizeof(u32) && std::is_trivially_copyable_v<FlagsType>)
193 [[nodiscard]] FlagsType Flags() const noexcept {
194 FlagsType ret;
195 std::memcpy(reinterpret_cast<char*>(&ret), &flags, sizeof(ret));
196 return ret;
197 }
198
199 template <typename FlagsType>
200 requires(sizeof(FlagsType) <= sizeof(u32) && std::is_trivially_copyable_v<FlagsType>)
201 [[nodiscard]] void SetFlags(FlagsType value) noexcept {
202 std::memcpy(&flags, &value, sizeof(value));
203 }
204
205 /// Intrusively store the host definition of this instruction.
206 template <typename DefinitionType>
207 void SetDefinition(DefinitionType def) {
208 definition = Common::BitCast<u32>(def);
209 }
210
211 /// Return the intrusively stored host definition of this instruction.
212 template <typename DefinitionType>
213 [[nodiscard]] DefinitionType Definition() const noexcept {
214 return Common::BitCast<DefinitionType>(definition);
215 }
216
217 /// Destructively remove one reference count from the instruction
218 /// Useful for register allocation
219 void DestructiveRemoveUsage() {
220 --use_count;
221 }
222
223 /// Destructively add usages to the instruction
224 /// Useful for register allocation
225 void DestructiveAddUsage(int count) {
226 use_count += count;
227 }
228
229private:
230 struct NonTriviallyDummy {
231 NonTriviallyDummy() noexcept {}
232 };
233
234 void Use(const Value& value);
235 void UndoUse(const Value& value);
236
237 IR::Opcode op{};
238 int use_count{};
239 u32 flags{};
240 u32 definition{};
241 union {
242 NonTriviallyDummy dummy{};
243 boost::container::small_vector<std::pair<Block*, Value>, 2> phi_args;
244 std::array<Value, 5> args;
245 };
246 std::unique_ptr<AssociatedInsts> associated_insts;
247};
248static_assert(sizeof(Inst) <= 128, "Inst size unintentionally increased");
249
250struct AssociatedInsts {
251 union {
252 Inst* in_bounds_inst;
253 Inst* sparse_inst;
254 Inst* zero_inst{};
255 };
256 Inst* sign_inst{};
257 Inst* carry_inst{};
258 Inst* overflow_inst{};
259};
260
261using U1 = TypedValue<Type::U1>;
262using U8 = TypedValue<Type::U8>;
263using U16 = TypedValue<Type::U16>;
264using U32 = TypedValue<Type::U32>;
265using U64 = TypedValue<Type::U64>;
266using F16 = TypedValue<Type::F16>;
267using F32 = TypedValue<Type::F32>;
268using F64 = TypedValue<Type::F64>;
269using U32U64 = TypedValue<Type::U32 | Type::U64>;
270using F32F64 = TypedValue<Type::F32 | Type::F64>;
271using U16U32U64 = TypedValue<Type::U16 | Type::U32 | Type::U64>;
272using F16F32F64 = TypedValue<Type::F16 | Type::F32 | Type::F64>;
273using UAny = TypedValue<Type::U8 | Type::U16 | Type::U32 | Type::U64>;
274
275inline bool Value::IsIdentity() const noexcept {
276 return type == Type::Opaque && inst->GetOpcode() == Opcode::Identity;
277}
278
279inline bool Value::IsPhi() const noexcept {
280 return type == Type::Opaque && inst->GetOpcode() == Opcode::Phi;
281}
282
283inline bool Value::IsEmpty() const noexcept {
284 return type == Type::Void;
285}
286
287inline bool Value::IsImmediate() const noexcept {
288 IR::Type current_type{type};
289 const IR::Inst* current_inst{inst};
290 while (current_type == Type::Opaque && current_inst->GetOpcode() == Opcode::Identity) {
291 const Value& arg{current_inst->Arg(0)};
292 current_type = arg.type;
293 current_inst = arg.inst;
294 }
295 return current_type != Type::Opaque;
296}
297
298inline IR::Inst* Value::Inst() const {
299 DEBUG_ASSERT(type == Type::Opaque);
300 return inst;
301}
302
303inline IR::Inst* Value::InstRecursive() const {
304 DEBUG_ASSERT(type == Type::Opaque);
305 if (IsIdentity()) {
306 return inst->Arg(0).InstRecursive();
307 }
308 return inst;
309}
310
311inline IR::Value Value::Resolve() const {
312 if (IsIdentity()) {
313 return inst->Arg(0).Resolve();
314 }
315 return *this;
316}
317
318inline IR::Reg Value::Reg() const {
319 DEBUG_ASSERT(type == Type::Reg);
320 return reg;
321}
322
323inline IR::Pred Value::Pred() const {
324 DEBUG_ASSERT(type == Type::Pred);
325 return pred;
326}
327
328inline IR::Attribute Value::Attribute() const {
329 DEBUG_ASSERT(type == Type::Attribute);
330 return attribute;
331}
332
333inline IR::Patch Value::Patch() const {
334 DEBUG_ASSERT(type == Type::Patch);
335 return patch;
336}
337
338inline bool Value::U1() const {
339 if (IsIdentity()) {
340 return inst->Arg(0).U1();
341 }
342 DEBUG_ASSERT(type == Type::U1);
343 return imm_u1;
344}
345
346inline u8 Value::U8() const {
347 if (IsIdentity()) {
348 return inst->Arg(0).U8();
349 }
350 DEBUG_ASSERT(type == Type::U8);
351 return imm_u8;
352}
353
354inline u16 Value::U16() const {
355 if (IsIdentity()) {
356 return inst->Arg(0).U16();
357 }
358 DEBUG_ASSERT(type == Type::U16);
359 return imm_u16;
360}
361
362inline u32 Value::U32() const {
363 if (IsIdentity()) {
364 return inst->Arg(0).U32();
365 }
366 DEBUG_ASSERT(type == Type::U32);
367 return imm_u32;
368}
369
370inline f32 Value::F32() const {
371 if (IsIdentity()) {
372 return inst->Arg(0).F32();
373 }
374 DEBUG_ASSERT(type == Type::F32);
375 return imm_f32;
376}
377
378inline u64 Value::U64() const {
379 if (IsIdentity()) {
380 return inst->Arg(0).U64();
381 }
382 DEBUG_ASSERT(type == Type::U64);
383 return imm_u64;
384}
385
386inline f64 Value::F64() const {
387 if (IsIdentity()) {
388 return inst->Arg(0).F64();
389 }
390 DEBUG_ASSERT(type == Type::F64);
391 return imm_f64;
392}
393
394[[nodiscard]] inline bool IsPhi(const Inst& inst) {
395 return inst.GetOpcode() == Opcode::Phi;
396}
397
398} // namespace Shader::IR
HIC SVNT DRACONES