blob: ea356a6ac4d8b4d80a0c834c4b58399148bd19be [file] [log] [blame]
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001/*
Hans-Kristian Arntzen318c17c2019-01-04 12:38:35 +01002 * Copyright 2018-2019 Arm Limited
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003 *
4 * Licensed under the Apache License, Version 2.0 (the "License");
5 * you may not use this file except in compliance with the License.
6 * You may obtain a copy of the License at
7 *
8 * http://www.apache.org/licenses/LICENSE-2.0
9 *
10 * Unless required by applicable law or agreed to in writing, software
11 * distributed under the License is distributed on an "AS IS" BASIS,
12 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 * See the License for the specific language governing permissions and
14 * limitations under the License.
15 */
16
17#include "spirv_parser.hpp"
18#include <assert.h>
19
20using namespace std;
21using namespace spv;
22
Hans-Kristian Arntzen9b92e682019-03-29 10:29:44 +010023namespace SPIRV_CROSS_NAMESPACE
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020024{
Hans-Kristian Arntzen3fe57d32019-04-09 12:46:23 +020025Parser::Parser(vector<uint32_t> spirv)
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020026{
27 ir.spirv = move(spirv);
28}
29
30Parser::Parser(const uint32_t *spirv_data, size_t word_count)
31{
Hans-Kristian Arntzen3fe57d32019-04-09 12:46:23 +020032 ir.spirv = vector<uint32_t>(spirv_data, spirv_data + word_count);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020033}
34
Hans-Kristian Arntzenfa42ed32018-11-15 10:51:01 +010035static bool decoration_is_string(Decoration decoration)
36{
37 switch (decoration)
38 {
39 case DecorationHlslSemanticGOOGLE:
40 return true;
41
42 default:
43 return false;
44 }
45}
46
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020047static inline uint32_t swap_endian(uint32_t v)
48{
49 return ((v >> 24) & 0x000000ffu) | ((v >> 8) & 0x0000ff00u) | ((v << 8) & 0x00ff0000u) | ((v << 24) & 0xff000000u);
50}
51
52static bool is_valid_spirv_version(uint32_t version)
53{
54 switch (version)
55 {
56 // Allow v99 since it tends to just work.
57 case 99:
58 case 0x10000: // SPIR-V 1.0
59 case 0x10100: // SPIR-V 1.1
60 case 0x10200: // SPIR-V 1.2
61 case 0x10300: // SPIR-V 1.3
62 return true;
63
64 default:
65 return false;
66 }
67}
68
69void Parser::parse()
70{
71 auto &spirv = ir.spirv;
72
73 auto len = spirv.size();
74 if (len < 5)
75 SPIRV_CROSS_THROW("SPIRV file too small.");
76
77 auto s = spirv.data();
78
79 // Endian-swap if we need to.
80 if (s[0] == swap_endian(MagicNumber))
81 transform(begin(spirv), end(spirv), begin(spirv), [](uint32_t c) { return swap_endian(c); });
82
83 if (s[0] != MagicNumber || !is_valid_spirv_version(s[1]))
84 SPIRV_CROSS_THROW("Invalid SPIRV format.");
85
86 uint32_t bound = s[3];
87 ir.set_id_bounds(bound);
88
89 uint32_t offset = 5;
90
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +020091 SmallVector<Instruction> instructions;
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020092 while (offset < len)
93 {
94 Instruction instr = {};
95 instr.op = spirv[offset] & 0xffff;
96 instr.count = (spirv[offset] >> 16) & 0xffff;
97
98 if (instr.count == 0)
99 SPIRV_CROSS_THROW("SPIR-V instructions cannot consume 0 words. Invalid SPIR-V file.");
100
101 instr.offset = offset + 1;
102 instr.length = instr.count - 1;
103
104 offset += instr.count;
105
106 if (offset > spirv.size())
107 SPIRV_CROSS_THROW("SPIR-V instruction goes out of bounds.");
108
109 instructions.push_back(instr);
110 }
111
112 for (auto &i : instructions)
113 parse(i);
114
115 if (current_function)
116 SPIRV_CROSS_THROW("Function was not terminated.");
117 if (current_block)
118 SPIRV_CROSS_THROW("Block was not terminated.");
119}
120
121const uint32_t *Parser::stream(const Instruction &instr) const
122{
123 // If we're not going to use any arguments, just return nullptr.
124 // We want to avoid case where we return an out of range pointer
125 // that trips debug assertions on some platforms.
126 if (!instr.length)
127 return nullptr;
128
129 if (instr.offset + instr.length > ir.spirv.size())
130 SPIRV_CROSS_THROW("Compiler::stream() out of range.");
131 return &ir.spirv[instr.offset];
132}
133
Hans-Kristian Arntzen3fe57d32019-04-09 12:46:23 +0200134static string extract_string(const vector<uint32_t> &spirv, uint32_t offset)
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200135{
136 string ret;
137 for (uint32_t i = offset; i < spirv.size(); i++)
138 {
139 uint32_t w = spirv[i];
140
141 for (uint32_t j = 0; j < 4; j++, w >>= 8)
142 {
143 char c = w & 0xff;
144 if (c == '\0')
145 return ret;
146 ret += c;
147 }
148 }
149
150 SPIRV_CROSS_THROW("String was not terminated before EOF");
151}
152
153void Parser::parse(const Instruction &instruction)
154{
155 auto *ops = stream(instruction);
156 auto op = static_cast<Op>(instruction.op);
157 uint32_t length = instruction.length;
158
159 switch (op)
160 {
lifpan00a765e2018-11-15 09:04:36 +0800161 case OpSourceContinued:
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200162 case OpSourceExtension:
163 case OpNop:
164 case OpLine:
165 case OpNoLine:
166 case OpString:
lifpan91610962018-11-13 14:28:38 +0800167 case OpModuleProcessed:
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200168 break;
169
Hans-Kristian Arntzen2cc374a2019-04-24 14:12:50 +0200170 case OpMemoryModel:
171 ir.addressing_model = static_cast<AddressingModel>(ops[0]);
172 ir.memory_model = static_cast<MemoryModel>(ops[1]);
173 break;
174
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200175 case OpSource:
176 {
177 auto lang = static_cast<SourceLanguage>(ops[0]);
178 switch (lang)
179 {
180 case SourceLanguageESSL:
181 ir.source.es = true;
182 ir.source.version = ops[1];
183 ir.source.known = true;
184 ir.source.hlsl = false;
185 break;
186
187 case SourceLanguageGLSL:
188 ir.source.es = false;
189 ir.source.version = ops[1];
190 ir.source.known = true;
191 ir.source.hlsl = false;
192 break;
193
194 case SourceLanguageHLSL:
195 // For purposes of cross-compiling, this is GLSL 450.
196 ir.source.es = false;
197 ir.source.version = 450;
198 ir.source.known = true;
199 ir.source.hlsl = true;
200 break;
201
202 default:
203 ir.source.known = false;
204 break;
205 }
206 break;
207 }
208
209 case OpUndef:
210 {
211 uint32_t result_type = ops[0];
212 uint32_t id = ops[1];
213 set<SPIRUndef>(id, result_type);
lifpan876627d2019-04-08 19:45:31 +0800214 if (current_block)
215 current_block->ops.push_back(instruction);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200216 break;
217 }
218
219 case OpCapability:
220 {
221 uint32_t cap = ops[0];
222 if (cap == CapabilityKernel)
223 SPIRV_CROSS_THROW("Kernel capability not supported.");
224
225 ir.declared_capabilities.push_back(static_cast<Capability>(ops[0]));
226 break;
227 }
228
229 case OpExtension:
230 {
231 auto ext = extract_string(ir.spirv, instruction.offset);
232 ir.declared_extensions.push_back(move(ext));
233 break;
234 }
235
236 case OpExtInstImport:
237 {
238 uint32_t id = ops[0];
239 auto ext = extract_string(ir.spirv, instruction.offset + 1);
240 if (ext == "GLSL.std.450")
241 set<SPIRExtension>(id, SPIRExtension::GLSL);
242 else if (ext == "SPV_AMD_shader_ballot")
243 set<SPIRExtension>(id, SPIRExtension::SPV_AMD_shader_ballot);
244 else if (ext == "SPV_AMD_shader_explicit_vertex_parameter")
245 set<SPIRExtension>(id, SPIRExtension::SPV_AMD_shader_explicit_vertex_parameter);
246 else if (ext == "SPV_AMD_shader_trinary_minmax")
247 set<SPIRExtension>(id, SPIRExtension::SPV_AMD_shader_trinary_minmax);
248 else if (ext == "SPV_AMD_gcn_shader")
249 set<SPIRExtension>(id, SPIRExtension::SPV_AMD_gcn_shader);
250 else
251 set<SPIRExtension>(id, SPIRExtension::Unsupported);
252
253 // Other SPIR-V extensions which have ExtInstrs are currently not supported.
254
255 break;
256 }
257
258 case OpEntryPoint:
259 {
260 auto itr =
261 ir.entry_points.insert(make_pair(ops[1], SPIREntryPoint(ops[1], static_cast<ExecutionModel>(ops[0]),
262 extract_string(ir.spirv, instruction.offset + 2))));
263 auto &e = itr.first->second;
264
265 // Strings need nul-terminator and consume the whole word.
266 uint32_t strlen_words = uint32_t((e.name.size() + 1 + 3) >> 2);
267 e.interface_variables.insert(end(e.interface_variables), ops + strlen_words + 2, ops + instruction.length);
268
269 // Set the name of the entry point in case OpName is not provided later.
270 ir.set_name(ops[1], e.name);
271
272 // If we don't have an entry, make the first one our "default".
273 if (!ir.default_entry_point)
274 ir.default_entry_point = ops[1];
275 break;
276 }
277
278 case OpExecutionMode:
279 {
280 auto &execution = ir.entry_points[ops[0]];
281 auto mode = static_cast<ExecutionMode>(ops[1]);
282 execution.flags.set(mode);
283
284 switch (mode)
285 {
286 case ExecutionModeInvocations:
287 execution.invocations = ops[2];
288 break;
289
290 case ExecutionModeLocalSize:
291 execution.workgroup_size.x = ops[2];
292 execution.workgroup_size.y = ops[3];
293 execution.workgroup_size.z = ops[4];
294 break;
295
296 case ExecutionModeOutputVertices:
297 execution.output_vertices = ops[2];
298 break;
299
300 default:
301 break;
302 }
303 break;
304 }
305
306 case OpName:
307 {
308 uint32_t id = ops[0];
309 ir.set_name(id, extract_string(ir.spirv, instruction.offset + 1));
310 break;
311 }
312
313 case OpMemberName:
314 {
315 uint32_t id = ops[0];
316 uint32_t member = ops[1];
317 ir.set_member_name(id, member, extract_string(ir.spirv, instruction.offset + 2));
318 break;
319 }
320
Hans-Kristian Arntzenfa42ed32018-11-15 10:51:01 +0100321 case OpDecorationGroup:
322 {
323 // Noop, this simply means an ID should be a collector of decorations.
324 // The meta array is already a flat array of decorations which will contain the relevant decorations.
325 break;
326 }
327
328 case OpGroupDecorate:
329 {
330 uint32_t group_id = ops[0];
331 auto &decorations = ir.meta[group_id].decoration;
332 auto &flags = decorations.decoration_flags;
333
334 // Copies decorations from one ID to another. Only copy decorations which are set in the group,
335 // i.e., we cannot just copy the meta structure directly.
336 for (uint32_t i = 1; i < length; i++)
337 {
338 uint32_t target = ops[i];
339 flags.for_each_bit([&](uint32_t bit) {
340 auto decoration = static_cast<Decoration>(bit);
341
342 if (decoration_is_string(decoration))
343 {
344 ir.set_decoration_string(target, decoration, ir.get_decoration_string(group_id, decoration));
345 }
346 else
347 {
348 ir.meta[target].decoration_word_offset[decoration] =
349 ir.meta[group_id].decoration_word_offset[decoration];
350 ir.set_decoration(target, decoration, ir.get_decoration(group_id, decoration));
351 }
352 });
353 }
354 break;
355 }
356
357 case OpGroupMemberDecorate:
358 {
359 uint32_t group_id = ops[0];
360 auto &flags = ir.meta[group_id].decoration.decoration_flags;
361
362 // Copies decorations from one ID to another. Only copy decorations which are set in the group,
363 // i.e., we cannot just copy the meta structure directly.
364 for (uint32_t i = 1; i + 1 < length; i += 2)
365 {
366 uint32_t target = ops[i + 0];
367 uint32_t index = ops[i + 1];
368 flags.for_each_bit([&](uint32_t bit) {
369 auto decoration = static_cast<Decoration>(bit);
370
371 if (decoration_is_string(decoration))
372 ir.set_member_decoration_string(target, index, decoration,
373 ir.get_decoration_string(group_id, decoration));
374 else
375 ir.set_member_decoration(target, index, decoration, ir.get_decoration(group_id, decoration));
376 });
377 }
378 break;
379 }
380
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200381 case OpDecorate:
382 case OpDecorateId:
383 {
Hans-Kristian Arntzenfa42ed32018-11-15 10:51:01 +0100384 // OpDecorateId technically supports an array of arguments, but our only supported decorations are single uint,
385 // so merge decorate and decorate-id here.
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200386 uint32_t id = ops[0];
387
388 auto decoration = static_cast<Decoration>(ops[1]);
389 if (length >= 3)
390 {
391 ir.meta[id].decoration_word_offset[decoration] = uint32_t(&ops[2] - ir.spirv.data());
392 ir.set_decoration(id, decoration, ops[2]);
393 }
394 else
395 ir.set_decoration(id, decoration);
396
397 break;
398 }
399
400 case OpDecorateStringGOOGLE:
401 {
402 uint32_t id = ops[0];
403 auto decoration = static_cast<Decoration>(ops[1]);
404 ir.set_decoration_string(id, decoration, extract_string(ir.spirv, instruction.offset + 2));
405 break;
406 }
407
408 case OpMemberDecorate:
409 {
410 uint32_t id = ops[0];
411 uint32_t member = ops[1];
412 auto decoration = static_cast<Decoration>(ops[2]);
413 if (length >= 4)
414 ir.set_member_decoration(id, member, decoration, ops[3]);
415 else
416 ir.set_member_decoration(id, member, decoration);
417 break;
418 }
419
420 case OpMemberDecorateStringGOOGLE:
421 {
422 uint32_t id = ops[0];
423 uint32_t member = ops[1];
424 auto decoration = static_cast<Decoration>(ops[2]);
425 ir.set_member_decoration_string(id, member, decoration, extract_string(ir.spirv, instruction.offset + 3));
426 break;
427 }
428
429 // Build up basic types.
430 case OpTypeVoid:
431 {
432 uint32_t id = ops[0];
433 auto &type = set<SPIRType>(id);
434 type.basetype = SPIRType::Void;
435 break;
436 }
437
438 case OpTypeBool:
439 {
440 uint32_t id = ops[0];
441 auto &type = set<SPIRType>(id);
442 type.basetype = SPIRType::Boolean;
443 type.width = 1;
444 break;
445 }
446
447 case OpTypeFloat:
448 {
449 uint32_t id = ops[0];
450 uint32_t width = ops[1];
451 auto &type = set<SPIRType>(id);
452 if (width == 64)
453 type.basetype = SPIRType::Double;
454 else if (width == 32)
455 type.basetype = SPIRType::Float;
456 else if (width == 16)
457 type.basetype = SPIRType::Half;
458 else
459 SPIRV_CROSS_THROW("Unrecognized bit-width of floating point type.");
460 type.width = width;
461 break;
462 }
463
464 case OpTypeInt:
465 {
466 uint32_t id = ops[0];
467 uint32_t width = ops[1];
lifpanb21525b2018-11-28 14:20:24 +0800468 bool signedness = ops[2] != 0;
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200469 auto &type = set<SPIRType>(id);
Hans-Kristian Arntzen2ed171e2019-01-30 14:49:55 +0100470 type.basetype = signedness ? to_signed_basetype(width) : to_unsigned_basetype(width);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200471 type.width = width;
472 break;
473 }
474
475 // Build composite types by "inheriting".
476 // NOTE: The self member is also copied! For pointers and array modifiers this is a good thing
477 // since we can refer to decorations on pointee classes which is needed for UBO/SSBO, I/O blocks in geometry/tess etc.
478 case OpTypeVector:
479 {
480 uint32_t id = ops[0];
481 uint32_t vecsize = ops[2];
482
483 auto &base = get<SPIRType>(ops[1]);
484 auto &vecbase = set<SPIRType>(id);
485
486 vecbase = base;
487 vecbase.vecsize = vecsize;
488 vecbase.self = id;
489 vecbase.parent_type = ops[1];
490 break;
491 }
492
493 case OpTypeMatrix:
494 {
495 uint32_t id = ops[0];
496 uint32_t colcount = ops[2];
497
498 auto &base = get<SPIRType>(ops[1]);
499 auto &matrixbase = set<SPIRType>(id);
500
501 matrixbase = base;
502 matrixbase.columns = colcount;
503 matrixbase.self = id;
504 matrixbase.parent_type = ops[1];
505 break;
506 }
507
508 case OpTypeArray:
509 {
510 uint32_t id = ops[0];
511 auto &arraybase = set<SPIRType>(id);
512
513 uint32_t tid = ops[1];
514 auto &base = get<SPIRType>(tid);
515
516 arraybase = base;
517 arraybase.parent_type = tid;
518
519 uint32_t cid = ops[2];
520 ir.mark_used_as_array_length(cid);
521 auto *c = maybe_get<SPIRConstant>(cid);
522 bool literal = c && !c->specialization;
523
524 arraybase.array_size_literal.push_back(literal);
525 arraybase.array.push_back(literal ? c->scalar() : cid);
526 // Do NOT set arraybase.self!
527 break;
528 }
529
530 case OpTypeRuntimeArray:
531 {
532 uint32_t id = ops[0];
533
534 auto &base = get<SPIRType>(ops[1]);
535 auto &arraybase = set<SPIRType>(id);
536
537 arraybase = base;
538 arraybase.array.push_back(0);
539 arraybase.array_size_literal.push_back(true);
540 arraybase.parent_type = ops[1];
541 // Do NOT set arraybase.self!
542 break;
543 }
544
545 case OpTypeImage:
546 {
547 uint32_t id = ops[0];
548 auto &type = set<SPIRType>(id);
549 type.basetype = SPIRType::Image;
550 type.image.type = ops[1];
551 type.image.dim = static_cast<Dim>(ops[2]);
552 type.image.depth = ops[3] == 1;
553 type.image.arrayed = ops[4] != 0;
554 type.image.ms = ops[5] != 0;
555 type.image.sampled = ops[6];
556 type.image.format = static_cast<ImageFormat>(ops[7]);
557 type.image.access = (length >= 9) ? static_cast<AccessQualifier>(ops[8]) : AccessQualifierMax;
558
559 if (type.image.sampled == 0)
560 SPIRV_CROSS_THROW("OpTypeImage Sampled parameter must not be zero.");
561
562 break;
563 }
564
565 case OpTypeSampledImage:
566 {
567 uint32_t id = ops[0];
568 uint32_t imagetype = ops[1];
569 auto &type = set<SPIRType>(id);
570 type = get<SPIRType>(imagetype);
571 type.basetype = SPIRType::SampledImage;
572 type.self = id;
573 break;
574 }
575
576 case OpTypeSampler:
577 {
578 uint32_t id = ops[0];
579 auto &type = set<SPIRType>(id);
580 type.basetype = SPIRType::Sampler;
581 break;
582 }
583
584 case OpTypePointer:
585 {
586 uint32_t id = ops[0];
587
588 auto &base = get<SPIRType>(ops[2]);
589 auto &ptrbase = set<SPIRType>(id);
590
591 ptrbase = base;
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200592 ptrbase.pointer = true;
Hans-Kristian Arntzend0b93722018-11-26 12:23:28 +0100593 ptrbase.pointer_depth++;
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200594 ptrbase.storage = static_cast<StorageClass>(ops[1]);
595
596 if (ptrbase.storage == StorageClassAtomicCounter)
597 ptrbase.basetype = SPIRType::AtomicCounter;
598
599 ptrbase.parent_type = ops[2];
600
601 // Do NOT set ptrbase.self!
602 break;
603 }
604
Hans-Kristian Arntzen2cc374a2019-04-24 14:12:50 +0200605 case OpTypeForwardPointer:
606 {
607 uint32_t id = ops[0];
608 auto &ptrbase = set<SPIRType>(id);
609 ptrbase.pointer = true;
610 ptrbase.pointer_depth++;
611 ptrbase.storage = static_cast<StorageClass>(ops[1]);
612
613 if (ptrbase.storage == StorageClassAtomicCounter)
614 ptrbase.basetype = SPIRType::AtomicCounter;
615
616 break;
617 }
618
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200619 case OpTypeStruct:
620 {
621 uint32_t id = ops[0];
622 auto &type = set<SPIRType>(id);
623 type.basetype = SPIRType::Struct;
624 for (uint32_t i = 1; i < length; i++)
625 type.member_types.push_back(ops[i]);
626
627 // Check if we have seen this struct type before, with just different
628 // decorations.
629 //
630 // Add workaround for issue #17 as well by looking at OpName for the struct
631 // types, which we shouldn't normally do.
632 // We should not normally have to consider type aliases like this to begin with
633 // however ... glslang issues #304, #307 cover this.
634
635 // For stripped names, never consider struct type aliasing.
636 // We risk declaring the same struct multiple times, but type-punning is not allowed
637 // so this is safe.
638 bool consider_aliasing = !ir.get_name(type.self).empty();
639 if (consider_aliasing)
640 {
641 for (auto &other : global_struct_cache)
642 {
643 if (ir.get_name(type.self) == ir.get_name(other) &&
644 types_are_logically_equivalent(type, get<SPIRType>(other)))
645 {
646 type.type_alias = other;
647 break;
648 }
649 }
650
651 if (type.type_alias == 0)
652 global_struct_cache.push_back(id);
653 }
654 break;
655 }
656
657 case OpTypeFunction:
658 {
659 uint32_t id = ops[0];
660 uint32_t ret = ops[1];
661
662 auto &func = set<SPIRFunctionPrototype>(id, ret);
663 for (uint32_t i = 2; i < length; i++)
664 func.parameter_types.push_back(ops[i]);
665 break;
666 }
667
Patrick Moursda39a7b2019-02-26 15:43:03 +0100668 case OpTypeAccelerationStructureNV:
669 {
670 uint32_t id = ops[0];
671 auto &type = set<SPIRType>(id);
672 type.basetype = SPIRType::AccelerationStructureNV;
673 break;
674 }
675
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200676 // Variable declaration
677 // All variables are essentially pointers with a storage qualifier.
678 case OpVariable:
679 {
680 uint32_t type = ops[0];
681 uint32_t id = ops[1];
682 auto storage = static_cast<StorageClass>(ops[2]);
683 uint32_t initializer = length == 4 ? ops[3] : 0;
684
685 if (storage == StorageClassFunction)
686 {
687 if (!current_function)
688 SPIRV_CROSS_THROW("No function currently in scope");
689 current_function->add_local_variable(id);
690 }
691
692 set<SPIRVariable>(id, type, storage, initializer);
693
694 // hlsl based shaders don't have those decorations. force them and then reset when reading/writing images
695 auto &ttype = get<SPIRType>(type);
696 if (ttype.basetype == SPIRType::BaseType::Image)
697 {
698 ir.set_decoration(id, DecorationNonWritable);
699 ir.set_decoration(id, DecorationNonReadable);
700 }
701
702 break;
703 }
704
705 // OpPhi
706 // OpPhi is a fairly magical opcode.
707 // It selects temporary variables based on which parent block we *came from*.
708 // In high-level languages we can "de-SSA" by creating a function local, and flush out temporaries to this function-local
709 // variable to emulate SSA Phi.
710 case OpPhi:
711 {
712 if (!current_function)
713 SPIRV_CROSS_THROW("No function currently in scope");
714 if (!current_block)
715 SPIRV_CROSS_THROW("No block currently in scope");
716
717 uint32_t result_type = ops[0];
718 uint32_t id = ops[1];
719
720 // Instead of a temporary, create a new function-wide temporary with this ID instead.
721 auto &var = set<SPIRVariable>(id, result_type, spv::StorageClassFunction);
722 var.phi_variable = true;
723
724 current_function->add_local_variable(id);
725
726 for (uint32_t i = 2; i + 2 <= length; i += 2)
727 current_block->phi_variables.push_back({ ops[i], ops[i + 1], id });
728 break;
729 }
730
731 // Constants
732 case OpSpecConstant:
733 case OpConstant:
734 {
735 uint32_t id = ops[1];
736 auto &type = get<SPIRType>(ops[0]);
737
738 if (type.width > 32)
739 set<SPIRConstant>(id, ops[0], ops[2] | (uint64_t(ops[3]) << 32), op == OpSpecConstant);
740 else
741 set<SPIRConstant>(id, ops[0], ops[2], op == OpSpecConstant);
742 break;
743 }
744
745 case OpSpecConstantFalse:
746 case OpConstantFalse:
747 {
748 uint32_t id = ops[1];
749 set<SPIRConstant>(id, ops[0], uint32_t(0), op == OpSpecConstantFalse);
750 break;
751 }
752
753 case OpSpecConstantTrue:
754 case OpConstantTrue:
755 {
756 uint32_t id = ops[1];
757 set<SPIRConstant>(id, ops[0], uint32_t(1), op == OpSpecConstantTrue);
758 break;
759 }
760
761 case OpConstantNull:
762 {
763 uint32_t id = ops[1];
764 uint32_t type = ops[0];
765 make_constant_null(id, type);
766 break;
767 }
768
769 case OpSpecConstantComposite:
770 case OpConstantComposite:
771 {
772 uint32_t id = ops[1];
773 uint32_t type = ops[0];
774
775 auto &ctype = get<SPIRType>(type);
776
777 // We can have constants which are structs and arrays.
778 // In this case, our SPIRConstant will be a list of other SPIRConstant ids which we
779 // can refer to.
780 if (ctype.basetype == SPIRType::Struct || !ctype.array.empty())
781 {
782 set<SPIRConstant>(id, type, ops + 2, length - 2, op == OpSpecConstantComposite);
783 }
784 else
785 {
786 uint32_t elements = length - 2;
787 if (elements > 4)
788 SPIRV_CROSS_THROW("OpConstantComposite only supports 1, 2, 3 and 4 elements.");
789
790 SPIRConstant remapped_constant_ops[4];
791 const SPIRConstant *c[4];
792 for (uint32_t i = 0; i < elements; i++)
793 {
794 // Specialization constants operations can also be part of this.
795 // We do not know their value, so any attempt to query SPIRConstant later
796 // will fail. We can only propagate the ID of the expression and use to_expression on it.
797 auto *constant_op = maybe_get<SPIRConstantOp>(ops[2 + i]);
Hans-Kristian Arntzendf3e21a2019-03-27 10:51:23 +0100798 auto *undef_op = maybe_get<SPIRUndef>(ops[2 + i]);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200799 if (constant_op)
800 {
801 if (op == OpConstantComposite)
802 SPIRV_CROSS_THROW("Specialization constant operation used in OpConstantComposite.");
803
804 remapped_constant_ops[i].make_null(get<SPIRType>(constant_op->basetype));
805 remapped_constant_ops[i].self = constant_op->self;
806 remapped_constant_ops[i].constant_type = constant_op->basetype;
807 remapped_constant_ops[i].specialization = true;
808 c[i] = &remapped_constant_ops[i];
809 }
Hans-Kristian Arntzendf3e21a2019-03-27 10:51:23 +0100810 else if (undef_op)
811 {
812 // Undefined, just pick 0.
813 remapped_constant_ops[i].make_null(get<SPIRType>(undef_op->basetype));
814 remapped_constant_ops[i].constant_type = undef_op->basetype;
815 c[i] = &remapped_constant_ops[i];
816 }
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200817 else
818 c[i] = &get<SPIRConstant>(ops[2 + i]);
819 }
820 set<SPIRConstant>(id, type, c, elements, op == OpSpecConstantComposite);
821 }
822 break;
823 }
824
825 // Functions
826 case OpFunction:
827 {
828 uint32_t res = ops[0];
829 uint32_t id = ops[1];
830 // Control
831 uint32_t type = ops[3];
832
833 if (current_function)
834 SPIRV_CROSS_THROW("Must end a function before starting a new one!");
835
836 current_function = &set<SPIRFunction>(id, res, type);
837 break;
838 }
839
840 case OpFunctionParameter:
841 {
842 uint32_t type = ops[0];
843 uint32_t id = ops[1];
844
845 if (!current_function)
846 SPIRV_CROSS_THROW("Must be in a function!");
847
848 current_function->add_parameter(type, id);
849 set<SPIRVariable>(id, type, StorageClassFunction);
850 break;
851 }
852
853 case OpFunctionEnd:
854 {
855 if (current_block)
856 {
857 // Very specific error message, but seems to come up quite often.
858 SPIRV_CROSS_THROW(
859 "Cannot end a function before ending the current block.\n"
860 "Likely cause: If this SPIR-V was created from glslang HLSL, make sure the entry point is valid.");
861 }
862 current_function = nullptr;
863 break;
864 }
865
866 // Blocks
867 case OpLabel:
868 {
869 // OpLabel always starts a block.
870 if (!current_function)
871 SPIRV_CROSS_THROW("Blocks cannot exist outside functions!");
872
873 uint32_t id = ops[0];
874
875 current_function->blocks.push_back(id);
876 if (!current_function->entry_block)
877 current_function->entry_block = id;
878
879 if (current_block)
880 SPIRV_CROSS_THROW("Cannot start a block before ending the current block.");
881
882 current_block = &set<SPIRBlock>(id);
883 break;
884 }
885
886 // Branch instructions end blocks.
887 case OpBranch:
888 {
889 if (!current_block)
890 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
891
892 uint32_t target = ops[0];
893 current_block->terminator = SPIRBlock::Direct;
894 current_block->next_block = target;
895 current_block = nullptr;
896 break;
897 }
898
899 case OpBranchConditional:
900 {
901 if (!current_block)
902 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
903
904 current_block->condition = ops[0];
905 current_block->true_block = ops[1];
906 current_block->false_block = ops[2];
907
908 current_block->terminator = SPIRBlock::Select;
909 current_block = nullptr;
910 break;
911 }
912
913 case OpSwitch:
914 {
915 if (!current_block)
916 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
917
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200918 current_block->terminator = SPIRBlock::MultiSelect;
919
920 current_block->condition = ops[0];
921 current_block->default_block = ops[1];
922
923 for (uint32_t i = 2; i + 2 <= length; i += 2)
924 current_block->cases.push_back({ ops[i], ops[i + 1] });
925
926 // If we jump to next block, make it break instead since we're inside a switch case block at that point.
927 ir.block_meta[current_block->next_block] |= ParsedIR::BLOCK_META_MULTISELECT_MERGE_BIT;
928
929 current_block = nullptr;
930 break;
931 }
932
933 case OpKill:
934 {
935 if (!current_block)
936 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
937 current_block->terminator = SPIRBlock::Kill;
938 current_block = nullptr;
939 break;
940 }
941
942 case OpReturn:
943 {
944 if (!current_block)
945 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
946 current_block->terminator = SPIRBlock::Return;
947 current_block = nullptr;
948 break;
949 }
950
951 case OpReturnValue:
952 {
953 if (!current_block)
954 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
955 current_block->terminator = SPIRBlock::Return;
956 current_block->return_value = ops[0];
957 current_block = nullptr;
958 break;
959 }
960
961 case OpUnreachable:
962 {
963 if (!current_block)
964 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
965 current_block->terminator = SPIRBlock::Unreachable;
966 current_block = nullptr;
967 break;
968 }
969
970 case OpSelectionMerge:
971 {
972 if (!current_block)
973 SPIRV_CROSS_THROW("Trying to modify a non-existing block.");
974
975 current_block->next_block = ops[0];
976 current_block->merge = SPIRBlock::MergeSelection;
977 ir.block_meta[current_block->next_block] |= ParsedIR::BLOCK_META_SELECTION_MERGE_BIT;
978
979 if (length >= 2)
980 {
981 if (ops[1] & SelectionControlFlattenMask)
982 current_block->hint = SPIRBlock::HintFlatten;
983 else if (ops[1] & SelectionControlDontFlattenMask)
984 current_block->hint = SPIRBlock::HintDontFlatten;
985 }
986 break;
987 }
988
989 case OpLoopMerge:
990 {
991 if (!current_block)
992 SPIRV_CROSS_THROW("Trying to modify a non-existing block.");
993
994 current_block->merge_block = ops[0];
995 current_block->continue_block = ops[1];
996 current_block->merge = SPIRBlock::MergeLoop;
997
998 ir.block_meta[current_block->self] |= ParsedIR::BLOCK_META_LOOP_HEADER_BIT;
999 ir.block_meta[current_block->merge_block] |= ParsedIR::BLOCK_META_LOOP_MERGE_BIT;
1000
1001 ir.continue_block_to_loop_header[current_block->continue_block] = current_block->self;
1002
1003 // Don't add loop headers to continue blocks,
1004 // which would make it impossible branch into the loop header since
1005 // they are treated as continues.
1006 if (current_block->continue_block != current_block->self)
1007 ir.block_meta[current_block->continue_block] |= ParsedIR::BLOCK_META_CONTINUE_BIT;
1008
1009 if (length >= 3)
1010 {
1011 if (ops[2] & LoopControlUnrollMask)
1012 current_block->hint = SPIRBlock::HintUnroll;
1013 else if (ops[2] & LoopControlDontUnrollMask)
1014 current_block->hint = SPIRBlock::HintDontUnroll;
1015 }
1016 break;
1017 }
1018
1019 case OpSpecConstantOp:
1020 {
1021 if (length < 3)
1022 SPIRV_CROSS_THROW("OpSpecConstantOp not enough arguments.");
1023
1024 uint32_t result_type = ops[0];
1025 uint32_t id = ops[1];
1026 auto spec_op = static_cast<Op>(ops[2]);
1027
1028 set<SPIRConstantOp>(id, result_type, spec_op, ops + 3, length - 3);
1029 break;
1030 }
1031
1032 // Actual opcodes.
1033 default:
1034 {
1035 if (!current_block)
1036 SPIRV_CROSS_THROW("Currently no block to insert opcode.");
1037
1038 current_block->ops.push_back(instruction);
1039 break;
1040 }
1041 }
1042}
1043
1044bool Parser::types_are_logically_equivalent(const SPIRType &a, const SPIRType &b) const
1045{
1046 if (a.basetype != b.basetype)
1047 return false;
1048 if (a.width != b.width)
1049 return false;
1050 if (a.vecsize != b.vecsize)
1051 return false;
1052 if (a.columns != b.columns)
1053 return false;
1054 if (a.array.size() != b.array.size())
1055 return false;
1056
1057 size_t array_count = a.array.size();
1058 if (array_count && memcmp(a.array.data(), b.array.data(), array_count * sizeof(uint32_t)) != 0)
1059 return false;
1060
1061 if (a.basetype == SPIRType::Image || a.basetype == SPIRType::SampledImage)
1062 {
1063 if (memcmp(&a.image, &b.image, sizeof(SPIRType::Image)) != 0)
1064 return false;
1065 }
1066
1067 if (a.member_types.size() != b.member_types.size())
1068 return false;
1069
1070 size_t member_types = a.member_types.size();
1071 for (size_t i = 0; i < member_types; i++)
1072 {
1073 if (!types_are_logically_equivalent(get<SPIRType>(a.member_types[i]), get<SPIRType>(b.member_types[i])))
1074 return false;
1075 }
1076
1077 return true;
1078}
1079
1080bool Parser::variable_storage_is_aliased(const SPIRVariable &v) const
1081{
1082 auto &type = get<SPIRType>(v.basetype);
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01001083
1084 auto *type_meta = ir.find_meta(type.self);
1085
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001086 bool ssbo = v.storage == StorageClassStorageBuffer ||
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01001087 (type_meta && type_meta->decoration.decoration_flags.get(DecorationBufferBlock));
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001088 bool image = type.basetype == SPIRType::Image;
1089 bool counter = type.basetype == SPIRType::AtomicCounter;
1090
1091 bool is_restrict;
1092 if (ssbo)
1093 is_restrict = ir.get_buffer_block_flags(v).get(DecorationRestrict);
1094 else
1095 is_restrict = ir.has_decoration(v.self, DecorationRestrict);
1096
1097 return !is_restrict && (ssbo || image || counter);
1098}
1099
1100void Parser::make_constant_null(uint32_t id, uint32_t type)
1101{
1102 auto &constant_type = get<SPIRType>(type);
1103
Chip Davis3bfb2f92018-12-03 02:06:33 -06001104 if (constant_type.pointer)
1105 {
1106 auto &constant = set<SPIRConstant>(id, type);
1107 constant.make_null(constant_type);
1108 }
1109 else if (!constant_type.array.empty())
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001110 {
1111 assert(constant_type.parent_type);
1112 uint32_t parent_id = ir.increase_bound_by(1);
1113 make_constant_null(parent_id, constant_type.parent_type);
1114
1115 if (!constant_type.array_size_literal.back())
1116 SPIRV_CROSS_THROW("Array size of OpConstantNull must be a literal.");
1117
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02001118 SmallVector<uint32_t> elements(constant_type.array.back());
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001119 for (uint32_t i = 0; i < constant_type.array.back(); i++)
1120 elements[i] = parent_id;
1121 set<SPIRConstant>(id, type, elements.data(), uint32_t(elements.size()), false);
1122 }
1123 else if (!constant_type.member_types.empty())
1124 {
1125 uint32_t member_ids = ir.increase_bound_by(uint32_t(constant_type.member_types.size()));
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02001126 SmallVector<uint32_t> elements(constant_type.member_types.size());
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001127 for (uint32_t i = 0; i < constant_type.member_types.size(); i++)
1128 {
1129 make_constant_null(member_ids + i, constant_type.member_types[i]);
1130 elements[i] = member_ids + i;
1131 }
1132 set<SPIRConstant>(id, type, elements.data(), uint32_t(elements.size()), false);
1133 }
1134 else
1135 {
1136 auto &constant = set<SPIRConstant>(id, type);
1137 constant.make_null(constant_type);
1138 }
1139}
1140
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02001141} // namespace SPIRV_CROSS_NAMESPACE