1 | /* |
---|---|
2 | * Copyright 2015-2021 Arm Limited |
3 | * SPDX-License-Identifier: Apache-2.0 OR MIT |
4 | * |
5 | * Licensed under the Apache License, Version 2.0 (the "License"); |
6 | * you may not use this file except in compliance with the License. |
7 | * You may obtain a copy of the License at |
8 | * |
9 | * http://www.apache.org/licenses/LICENSE-2.0 |
10 | * |
11 | * Unless required by applicable law or agreed to in writing, software |
12 | * distributed under the License is distributed on an "AS IS" BASIS, |
13 | * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. |
14 | * See the License for the specific language governing permissions and |
15 | * limitations under the License. |
16 | */ |
17 | |
18 | /* |
19 | * At your option, you may choose to accept this material under either: |
20 | * 1. The Apache License, Version 2.0, found at <http://www.apache.org/licenses/LICENSE-2.0>, or |
21 | * 2. The MIT License, found at <http://opensource.org/licenses/MIT>. |
22 | */ |
23 | |
24 | #include "spirv_cross.hpp" |
25 | #include "GLSL.std.450.h" |
26 | #include "spirv_cfg.hpp" |
27 | #include "spirv_common.hpp" |
28 | #include "spirv_parser.hpp" |
29 | #include <algorithm> |
30 | #include <cstring> |
31 | #include <utility> |
32 | |
33 | using namespace std; |
34 | using namespace spv; |
35 | using namespace SPIRV_CROSS_NAMESPACE; |
36 | |
37 | Compiler::Compiler(vector<uint32_t> ir_) |
38 | { |
39 | Parser parser(std::move(ir_)); |
40 | parser.parse(); |
41 | set_ir(std::move(parser.get_parsed_ir())); |
42 | } |
43 | |
44 | Compiler::Compiler(const uint32_t *ir_, size_t word_count) |
45 | { |
46 | Parser parser(ir_, word_count); |
47 | parser.parse(); |
48 | set_ir(std::move(parser.get_parsed_ir())); |
49 | } |
50 | |
51 | Compiler::Compiler(const ParsedIR &ir_) |
52 | { |
53 | set_ir(ir_); |
54 | } |
55 | |
56 | Compiler::Compiler(ParsedIR &&ir_) |
57 | { |
58 | set_ir(std::move(ir_)); |
59 | } |
60 | |
61 | void Compiler::set_ir(ParsedIR &&ir_) |
62 | { |
63 | ir = std::move(ir_); |
64 | parse_fixup(); |
65 | } |
66 | |
67 | void Compiler::set_ir(const ParsedIR &ir_) |
68 | { |
69 | ir = ir_; |
70 | parse_fixup(); |
71 | } |
72 | |
73 | string Compiler::compile() |
74 | { |
75 | return ""; |
76 | } |
77 | |
78 | bool Compiler::variable_storage_is_aliased(const SPIRVariable &v) |
79 | { |
80 | auto &type = get<SPIRType>(id: v.basetype); |
81 | bool ssbo = v.storage == StorageClassStorageBuffer || |
82 | ir.meta[type.self].decoration.decoration_flags.get(bit: DecorationBufferBlock); |
83 | bool image = type.basetype == SPIRType::Image; |
84 | bool counter = type.basetype == SPIRType::AtomicCounter; |
85 | bool buffer_reference = type.storage == StorageClassPhysicalStorageBufferEXT; |
86 | |
87 | bool is_restrict; |
88 | if (ssbo) |
89 | is_restrict = ir.get_buffer_block_flags(var: v).get(bit: DecorationRestrict); |
90 | else |
91 | is_restrict = has_decoration(id: v.self, decoration: DecorationRestrict); |
92 | |
93 | return !is_restrict && (ssbo || image || counter || buffer_reference); |
94 | } |
95 | |
96 | bool Compiler::block_is_control_dependent(const SPIRBlock &block) |
97 | { |
98 | for (auto &i : block.ops) |
99 | { |
100 | auto ops = stream(instr: i); |
101 | auto op = static_cast<Op>(i.op); |
102 | |
103 | switch (op) |
104 | { |
105 | case OpFunctionCall: |
106 | { |
107 | uint32_t func = ops[2]; |
108 | if (function_is_control_dependent(func: get<SPIRFunction>(id: func))) |
109 | return true; |
110 | break; |
111 | } |
112 | |
113 | // Derivatives |
114 | case OpDPdx: |
115 | case OpDPdxCoarse: |
116 | case OpDPdxFine: |
117 | case OpDPdy: |
118 | case OpDPdyCoarse: |
119 | case OpDPdyFine: |
120 | case OpFwidth: |
121 | case OpFwidthCoarse: |
122 | case OpFwidthFine: |
123 | |
124 | // Anything implicit LOD |
125 | case OpImageSampleImplicitLod: |
126 | case OpImageSampleDrefImplicitLod: |
127 | case OpImageSampleProjImplicitLod: |
128 | case OpImageSampleProjDrefImplicitLod: |
129 | case OpImageSparseSampleImplicitLod: |
130 | case OpImageSparseSampleDrefImplicitLod: |
131 | case OpImageSparseSampleProjImplicitLod: |
132 | case OpImageSparseSampleProjDrefImplicitLod: |
133 | case OpImageQueryLod: |
134 | case OpImageDrefGather: |
135 | case OpImageGather: |
136 | case OpImageSparseDrefGather: |
137 | case OpImageSparseGather: |
138 | |
139 | // Anything subgroups |
140 | case OpGroupNonUniformElect: |
141 | case OpGroupNonUniformAll: |
142 | case OpGroupNonUniformAny: |
143 | case OpGroupNonUniformAllEqual: |
144 | case OpGroupNonUniformBroadcast: |
145 | case OpGroupNonUniformBroadcastFirst: |
146 | case OpGroupNonUniformBallot: |
147 | case OpGroupNonUniformInverseBallot: |
148 | case OpGroupNonUniformBallotBitExtract: |
149 | case OpGroupNonUniformBallotBitCount: |
150 | case OpGroupNonUniformBallotFindLSB: |
151 | case OpGroupNonUniformBallotFindMSB: |
152 | case OpGroupNonUniformShuffle: |
153 | case OpGroupNonUniformShuffleXor: |
154 | case OpGroupNonUniformShuffleUp: |
155 | case OpGroupNonUniformShuffleDown: |
156 | case OpGroupNonUniformIAdd: |
157 | case OpGroupNonUniformFAdd: |
158 | case OpGroupNonUniformIMul: |
159 | case OpGroupNonUniformFMul: |
160 | case OpGroupNonUniformSMin: |
161 | case OpGroupNonUniformUMin: |
162 | case OpGroupNonUniformFMin: |
163 | case OpGroupNonUniformSMax: |
164 | case OpGroupNonUniformUMax: |
165 | case OpGroupNonUniformFMax: |
166 | case OpGroupNonUniformBitwiseAnd: |
167 | case OpGroupNonUniformBitwiseOr: |
168 | case OpGroupNonUniformBitwiseXor: |
169 | case OpGroupNonUniformLogicalAnd: |
170 | case OpGroupNonUniformLogicalOr: |
171 | case OpGroupNonUniformLogicalXor: |
172 | case OpGroupNonUniformQuadBroadcast: |
173 | case OpGroupNonUniformQuadSwap: |
174 | |
175 | // Control barriers |
176 | case OpControlBarrier: |
177 | return true; |
178 | |
179 | default: |
180 | break; |
181 | } |
182 | } |
183 | |
184 | return false; |
185 | } |
186 | |
187 | bool Compiler::block_is_pure(const SPIRBlock &block) |
188 | { |
189 | // This is a global side effect of the function. |
190 | if (block.terminator == SPIRBlock::Kill || |
191 | block.terminator == SPIRBlock::TerminateRay || |
192 | block.terminator == SPIRBlock::IgnoreIntersection || |
193 | block.terminator == SPIRBlock::EmitMeshTasks) |
194 | return false; |
195 | |
196 | for (auto &i : block.ops) |
197 | { |
198 | auto ops = stream(instr: i); |
199 | auto op = static_cast<Op>(i.op); |
200 | |
201 | switch (op) |
202 | { |
203 | case OpFunctionCall: |
204 | { |
205 | uint32_t func = ops[2]; |
206 | if (!function_is_pure(func: get<SPIRFunction>(id: func))) |
207 | return false; |
208 | break; |
209 | } |
210 | |
211 | case OpCopyMemory: |
212 | case OpStore: |
213 | { |
214 | auto &type = expression_type(id: ops[0]); |
215 | if (type.storage != StorageClassFunction) |
216 | return false; |
217 | break; |
218 | } |
219 | |
220 | case OpImageWrite: |
221 | return false; |
222 | |
223 | // Atomics are impure. |
224 | case OpAtomicLoad: |
225 | case OpAtomicStore: |
226 | case OpAtomicExchange: |
227 | case OpAtomicCompareExchange: |
228 | case OpAtomicCompareExchangeWeak: |
229 | case OpAtomicIIncrement: |
230 | case OpAtomicIDecrement: |
231 | case OpAtomicIAdd: |
232 | case OpAtomicISub: |
233 | case OpAtomicSMin: |
234 | case OpAtomicUMin: |
235 | case OpAtomicSMax: |
236 | case OpAtomicUMax: |
237 | case OpAtomicAnd: |
238 | case OpAtomicOr: |
239 | case OpAtomicXor: |
240 | return false; |
241 | |
242 | // Geometry shader builtins modify global state. |
243 | case OpEndPrimitive: |
244 | case OpEmitStreamVertex: |
245 | case OpEndStreamPrimitive: |
246 | case OpEmitVertex: |
247 | return false; |
248 | |
249 | // Mesh shader functions modify global state. |
250 | // (EmitMeshTasks is a terminator). |
251 | case OpSetMeshOutputsEXT: |
252 | return false; |
253 | |
254 | // Barriers disallow any reordering, so we should treat blocks with barrier as writing. |
255 | case OpControlBarrier: |
256 | case OpMemoryBarrier: |
257 | return false; |
258 | |
259 | // Ray tracing builtins are impure. |
260 | case OpReportIntersectionKHR: |
261 | case OpIgnoreIntersectionNV: |
262 | case OpTerminateRayNV: |
263 | case OpTraceNV: |
264 | case OpTraceRayKHR: |
265 | case OpExecuteCallableNV: |
266 | case OpExecuteCallableKHR: |
267 | case OpRayQueryInitializeKHR: |
268 | case OpRayQueryTerminateKHR: |
269 | case OpRayQueryGenerateIntersectionKHR: |
270 | case OpRayQueryConfirmIntersectionKHR: |
271 | case OpRayQueryProceedKHR: |
272 | // There are various getters in ray query, but they are considered pure. |
273 | return false; |
274 | |
275 | // OpExtInst is potentially impure depending on extension, but GLSL builtins are at least pure. |
276 | |
277 | case OpDemoteToHelperInvocationEXT: |
278 | // This is a global side effect of the function. |
279 | return false; |
280 | |
281 | case OpExtInst: |
282 | { |
283 | uint32_t extension_set = ops[2]; |
284 | if (get<SPIRExtension>(id: extension_set).ext == SPIRExtension::GLSL) |
285 | { |
286 | auto op_450 = static_cast<GLSLstd450>(ops[3]); |
287 | switch (op_450) |
288 | { |
289 | case GLSLstd450Modf: |
290 | case GLSLstd450Frexp: |
291 | { |
292 | auto &type = expression_type(id: ops[5]); |
293 | if (type.storage != StorageClassFunction) |
294 | return false; |
295 | break; |
296 | } |
297 | |
298 | default: |
299 | break; |
300 | } |
301 | } |
302 | break; |
303 | } |
304 | |
305 | default: |
306 | break; |
307 | } |
308 | } |
309 | |
310 | return true; |
311 | } |
312 | |
313 | string Compiler::to_name(uint32_t id, bool allow_alias) const |
314 | { |
315 | if (allow_alias && ir.ids[id].get_type() == TypeType) |
316 | { |
317 | // If this type is a simple alias, emit the |
318 | // name of the original type instead. |
319 | // We don't want to override the meta alias |
320 | // as that can be overridden by the reflection APIs after parse. |
321 | auto &type = get<SPIRType>(id); |
322 | if (type.type_alias) |
323 | { |
324 | // If the alias master has been specially packed, we will have emitted a clean variant as well, |
325 | // so skip the name aliasing here. |
326 | if (!has_extended_decoration(id: type.type_alias, decoration: SPIRVCrossDecorationBufferBlockRepacked)) |
327 | return to_name(id: type.type_alias); |
328 | } |
329 | } |
330 | |
331 | auto &alias = ir.get_name(id); |
332 | if (alias.empty()) |
333 | return join(ts: "_", ts&: id); |
334 | else |
335 | return alias; |
336 | } |
337 | |
338 | bool Compiler::function_is_pure(const SPIRFunction &func) |
339 | { |
340 | for (auto block : func.blocks) |
341 | if (!block_is_pure(block: get<SPIRBlock>(id: block))) |
342 | return false; |
343 | |
344 | return true; |
345 | } |
346 | |
347 | bool Compiler::function_is_control_dependent(const SPIRFunction &func) |
348 | { |
349 | for (auto block : func.blocks) |
350 | if (block_is_control_dependent(block: get<SPIRBlock>(id: block))) |
351 | return true; |
352 | |
353 | return false; |
354 | } |
355 | |
356 | void Compiler::register_global_read_dependencies(const SPIRBlock &block, uint32_t id) |
357 | { |
358 | for (auto &i : block.ops) |
359 | { |
360 | auto ops = stream(instr: i); |
361 | auto op = static_cast<Op>(i.op); |
362 | |
363 | switch (op) |
364 | { |
365 | case OpFunctionCall: |
366 | { |
367 | uint32_t func = ops[2]; |
368 | register_global_read_dependencies(func: get<SPIRFunction>(id: func), id); |
369 | break; |
370 | } |
371 | |
372 | case OpLoad: |
373 | case OpImageRead: |
374 | { |
375 | // If we're in a storage class which does not get invalidated, adding dependencies here is no big deal. |
376 | auto *var = maybe_get_backing_variable(chain: ops[2]); |
377 | if (var && var->storage != StorageClassFunction) |
378 | { |
379 | auto &type = get<SPIRType>(id: var->basetype); |
380 | |
381 | // InputTargets are immutable. |
382 | if (type.basetype != SPIRType::Image && type.image.dim != DimSubpassData) |
383 | var->dependees.push_back(t: id); |
384 | } |
385 | break; |
386 | } |
387 | |
388 | default: |
389 | break; |
390 | } |
391 | } |
392 | } |
393 | |
394 | void Compiler::register_global_read_dependencies(const SPIRFunction &func, uint32_t id) |
395 | { |
396 | for (auto block : func.blocks) |
397 | register_global_read_dependencies(block: get<SPIRBlock>(id: block), id); |
398 | } |
399 | |
400 | SPIRVariable *Compiler::maybe_get_backing_variable(uint32_t chain) |
401 | { |
402 | auto *var = maybe_get<SPIRVariable>(id: chain); |
403 | if (!var) |
404 | { |
405 | auto *cexpr = maybe_get<SPIRExpression>(id: chain); |
406 | if (cexpr) |
407 | var = maybe_get<SPIRVariable>(id: cexpr->loaded_from); |
408 | |
409 | auto *access_chain = maybe_get<SPIRAccessChain>(id: chain); |
410 | if (access_chain) |
411 | var = maybe_get<SPIRVariable>(id: access_chain->loaded_from); |
412 | } |
413 | |
414 | return var; |
415 | } |
416 | |
417 | void Compiler::register_read(uint32_t expr, uint32_t chain, bool forwarded) |
418 | { |
419 | auto &e = get<SPIRExpression>(id: expr); |
420 | auto *var = maybe_get_backing_variable(chain); |
421 | |
422 | if (var) |
423 | { |
424 | e.loaded_from = var->self; |
425 | |
426 | // If the backing variable is immutable, we do not need to depend on the variable. |
427 | if (forwarded && !is_immutable(id: var->self)) |
428 | var->dependees.push_back(t: e.self); |
429 | |
430 | // If we load from a parameter, make sure we create "inout" if we also write to the parameter. |
431 | // The default is "in" however, so we never invalidate our compilation by reading. |
432 | if (var && var->parameter) |
433 | var->parameter->read_count++; |
434 | } |
435 | } |
436 | |
437 | void Compiler::register_write(uint32_t chain) |
438 | { |
439 | auto *var = maybe_get<SPIRVariable>(id: chain); |
440 | if (!var) |
441 | { |
442 | // If we're storing through an access chain, invalidate the backing variable instead. |
443 | auto *expr = maybe_get<SPIRExpression>(id: chain); |
444 | if (expr && expr->loaded_from) |
445 | var = maybe_get<SPIRVariable>(id: expr->loaded_from); |
446 | |
447 | auto *access_chain = maybe_get<SPIRAccessChain>(id: chain); |
448 | if (access_chain && access_chain->loaded_from) |
449 | var = maybe_get<SPIRVariable>(id: access_chain->loaded_from); |
450 | } |
451 | |
452 | auto &chain_type = expression_type(id: chain); |
453 | |
454 | if (var) |
455 | { |
456 | bool check_argument_storage_qualifier = true; |
457 | auto &type = expression_type(id: chain); |
458 | |
459 | // If our variable is in a storage class which can alias with other buffers, |
460 | // invalidate all variables which depend on aliased variables. And if this is a |
461 | // variable pointer, then invalidate all variables regardless. |
462 | if (get_variable_data_type(var: *var).pointer) |
463 | { |
464 | flush_all_active_variables(); |
465 | |
466 | if (type.pointer_depth == 1) |
467 | { |
468 | // We have a backing variable which is a pointer-to-pointer type. |
469 | // We are storing some data through a pointer acquired through that variable, |
470 | // but we are not writing to the value of the variable itself, |
471 | // i.e., we are not modifying the pointer directly. |
472 | // If we are storing a non-pointer type (pointer_depth == 1), |
473 | // we know that we are storing some unrelated data. |
474 | // A case here would be |
475 | // void foo(Foo * const *arg) { |
476 | // Foo *bar = *arg; |
477 | // bar->unrelated = 42; |
478 | // } |
479 | // arg, the argument is constant. |
480 | check_argument_storage_qualifier = false; |
481 | } |
482 | } |
483 | |
484 | if (type.storage == StorageClassPhysicalStorageBufferEXT || variable_storage_is_aliased(v: *var)) |
485 | flush_all_aliased_variables(); |
486 | else if (var) |
487 | flush_dependees(var&: *var); |
488 | |
489 | // We tried to write to a parameter which is not marked with out qualifier, force a recompile. |
490 | if (check_argument_storage_qualifier && var->parameter && var->parameter->write_count == 0) |
491 | { |
492 | var->parameter->write_count++; |
493 | force_recompile(); |
494 | } |
495 | } |
496 | else if (chain_type.pointer) |
497 | { |
498 | // If we stored through a variable pointer, then we don't know which |
499 | // variable we stored to. So *all* expressions after this point need to |
500 | // be invalidated. |
501 | // FIXME: If we can prove that the variable pointer will point to |
502 | // only certain variables, we can invalidate only those. |
503 | flush_all_active_variables(); |
504 | } |
505 | |
506 | // If chain_type.pointer is false, we're not writing to memory backed variables, but temporaries instead. |
507 | // This can happen in copy_logical_type where we unroll complex reads and writes to temporaries. |
508 | } |
509 | |
510 | void Compiler::flush_dependees(SPIRVariable &var) |
511 | { |
512 | for (auto expr : var.dependees) |
513 | invalid_expressions.insert(x: expr); |
514 | var.dependees.clear(); |
515 | } |
516 | |
517 | void Compiler::flush_all_aliased_variables() |
518 | { |
519 | for (auto aliased : aliased_variables) |
520 | flush_dependees(var&: get<SPIRVariable>(id: aliased)); |
521 | } |
522 | |
523 | void Compiler::flush_all_atomic_capable_variables() |
524 | { |
525 | for (auto global : global_variables) |
526 | flush_dependees(var&: get<SPIRVariable>(id: global)); |
527 | flush_all_aliased_variables(); |
528 | } |
529 | |
530 | void Compiler::flush_control_dependent_expressions(uint32_t block_id) |
531 | { |
532 | auto &block = get<SPIRBlock>(id: block_id); |
533 | for (auto &expr : block.invalidate_expressions) |
534 | invalid_expressions.insert(x: expr); |
535 | block.invalidate_expressions.clear(); |
536 | } |
537 | |
538 | void Compiler::flush_all_active_variables() |
539 | { |
540 | // Invalidate all temporaries we read from variables in this block since they were forwarded. |
541 | // Invalidate all temporaries we read from globals. |
542 | for (auto &v : current_function->local_variables) |
543 | flush_dependees(var&: get<SPIRVariable>(id: v)); |
544 | for (auto &arg : current_function->arguments) |
545 | flush_dependees(var&: get<SPIRVariable>(id: arg.id)); |
546 | for (auto global : global_variables) |
547 | flush_dependees(var&: get<SPIRVariable>(id: global)); |
548 | |
549 | flush_all_aliased_variables(); |
550 | } |
551 | |
552 | uint32_t Compiler::expression_type_id(uint32_t id) const |
553 | { |
554 | switch (ir.ids[id].get_type()) |
555 | { |
556 | case TypeVariable: |
557 | return get<SPIRVariable>(id).basetype; |
558 | |
559 | case TypeExpression: |
560 | return get<SPIRExpression>(id).expression_type; |
561 | |
562 | case TypeConstant: |
563 | return get<SPIRConstant>(id).constant_type; |
564 | |
565 | case TypeConstantOp: |
566 | return get<SPIRConstantOp>(id).basetype; |
567 | |
568 | case TypeUndef: |
569 | return get<SPIRUndef>(id).basetype; |
570 | |
571 | case TypeCombinedImageSampler: |
572 | return get<SPIRCombinedImageSampler>(id).combined_type; |
573 | |
574 | case TypeAccessChain: |
575 | return get<SPIRAccessChain>(id).basetype; |
576 | |
577 | default: |
578 | SPIRV_CROSS_THROW("Cannot resolve expression type."); |
579 | } |
580 | } |
581 | |
582 | const SPIRType &Compiler::expression_type(uint32_t id) const |
583 | { |
584 | return get<SPIRType>(id: expression_type_id(id)); |
585 | } |
586 | |
587 | bool Compiler::expression_is_lvalue(uint32_t id) const |
588 | { |
589 | auto &type = expression_type(id); |
590 | switch (type.basetype) |
591 | { |
592 | case SPIRType::SampledImage: |
593 | case SPIRType::Image: |
594 | case SPIRType::Sampler: |
595 | return false; |
596 | |
597 | default: |
598 | return true; |
599 | } |
600 | } |
601 | |
602 | bool Compiler::is_immutable(uint32_t id) const |
603 | { |
604 | if (ir.ids[id].get_type() == TypeVariable) |
605 | { |
606 | auto &var = get<SPIRVariable>(id); |
607 | |
608 | // Anything we load from the UniformConstant address space is guaranteed to be immutable. |
609 | bool pointer_to_const = var.storage == StorageClassUniformConstant; |
610 | return pointer_to_const || var.phi_variable || !expression_is_lvalue(id); |
611 | } |
612 | else if (ir.ids[id].get_type() == TypeAccessChain) |
613 | return get<SPIRAccessChain>(id).immutable; |
614 | else if (ir.ids[id].get_type() == TypeExpression) |
615 | return get<SPIRExpression>(id).immutable; |
616 | else if (ir.ids[id].get_type() == TypeConstant || ir.ids[id].get_type() == TypeConstantOp || |
617 | ir.ids[id].get_type() == TypeUndef) |
618 | return true; |
619 | else |
620 | return false; |
621 | } |
622 | |
623 | static inline bool storage_class_is_interface(spv::StorageClass storage) |
624 | { |
625 | switch (storage) |
626 | { |
627 | case StorageClassInput: |
628 | case StorageClassOutput: |
629 | case StorageClassUniform: |
630 | case StorageClassUniformConstant: |
631 | case StorageClassAtomicCounter: |
632 | case StorageClassPushConstant: |
633 | case StorageClassStorageBuffer: |
634 | return true; |
635 | |
636 | default: |
637 | return false; |
638 | } |
639 | } |
640 | |
641 | bool Compiler::is_hidden_variable(const SPIRVariable &var, bool include_builtins) const |
642 | { |
643 | if ((is_builtin_variable(var) && !include_builtins) || var.remapped_variable) |
644 | return true; |
645 | |
646 | // Combined image samplers are always considered active as they are "magic" variables. |
647 | if (find_if(first: begin(cont: combined_image_samplers), last: end(cont: combined_image_samplers), pred: [&var](const CombinedImageSampler &samp) { |
648 | return samp.combined_id == var.self; |
649 | }) != end(cont: combined_image_samplers)) |
650 | { |
651 | return false; |
652 | } |
653 | |
654 | // In SPIR-V 1.4 and up we must also use the active variable interface to disable global variables |
655 | // which are not part of the entry point. |
656 | if (ir.get_spirv_version() >= 0x10400 && var.storage != spv::StorageClassGeneric && |
657 | var.storage != spv::StorageClassFunction && !interface_variable_exists_in_entry_point(id: var.self)) |
658 | { |
659 | return true; |
660 | } |
661 | |
662 | return check_active_interface_variables && storage_class_is_interface(storage: var.storage) && |
663 | active_interface_variables.find(x: var.self) == end(cont: active_interface_variables); |
664 | } |
665 | |
666 | bool Compiler::is_builtin_type(const SPIRType &type) const |
667 | { |
668 | auto *type_meta = ir.find_meta(id: type.self); |
669 | |
670 | // We can have builtin structs as well. If one member of a struct is builtin, the struct must also be builtin. |
671 | if (type_meta) |
672 | for (auto &m : type_meta->members) |
673 | if (m.builtin) |
674 | return true; |
675 | |
676 | return false; |
677 | } |
678 | |
679 | bool Compiler::is_builtin_variable(const SPIRVariable &var) const |
680 | { |
681 | auto *m = ir.find_meta(id: var.self); |
682 | |
683 | if (var.compat_builtin || (m && m->decoration.builtin)) |
684 | return true; |
685 | else |
686 | return is_builtin_type(type: get<SPIRType>(id: var.basetype)); |
687 | } |
688 | |
689 | bool Compiler::is_member_builtin(const SPIRType &type, uint32_t index, BuiltIn *builtin) const |
690 | { |
691 | auto *type_meta = ir.find_meta(id: type.self); |
692 | |
693 | if (type_meta) |
694 | { |
695 | auto &memb = type_meta->members; |
696 | if (index < memb.size() && memb[index].builtin) |
697 | { |
698 | if (builtin) |
699 | *builtin = memb[index].builtin_type; |
700 | return true; |
701 | } |
702 | } |
703 | |
704 | return false; |
705 | } |
706 | |
707 | bool Compiler::is_scalar(const SPIRType &type) const |
708 | { |
709 | return type.basetype != SPIRType::Struct && type.vecsize == 1 && type.columns == 1; |
710 | } |
711 | |
712 | bool Compiler::is_vector(const SPIRType &type) const |
713 | { |
714 | return type.vecsize > 1 && type.columns == 1; |
715 | } |
716 | |
717 | bool Compiler::is_matrix(const SPIRType &type) const |
718 | { |
719 | return type.vecsize > 1 && type.columns > 1; |
720 | } |
721 | |
722 | bool Compiler::is_array(const SPIRType &type) const |
723 | { |
724 | return type.op == OpTypeArray || type.op == OpTypeRuntimeArray; |
725 | } |
726 | |
727 | bool Compiler::is_pointer(const SPIRType &type) const |
728 | { |
729 | return type.op == OpTypePointer && type.basetype != SPIRType::Unknown; // Ignore function pointers. |
730 | } |
731 | |
732 | bool Compiler::is_physical_pointer(const SPIRType &type) const |
733 | { |
734 | return type.op == OpTypePointer && type.storage == StorageClassPhysicalStorageBuffer; |
735 | } |
736 | |
737 | bool Compiler::is_physical_pointer_to_buffer_block(const SPIRType &type) const |
738 | { |
739 | return is_physical_pointer(type) && get_pointee_type(type).self == type.parent_type && |
740 | (has_decoration(id: type.self, decoration: DecorationBlock) || |
741 | has_decoration(id: type.self, decoration: DecorationBufferBlock)); |
742 | } |
743 | |
744 | bool Compiler::is_runtime_size_array(const SPIRType &type) |
745 | { |
746 | return type.op == OpTypeRuntimeArray; |
747 | } |
748 | |
749 | ShaderResources Compiler::get_shader_resources() const |
750 | { |
751 | return get_shader_resources(active_variables: nullptr); |
752 | } |
753 | |
754 | ShaderResources Compiler::get_shader_resources(const unordered_set<VariableID> &active_variables) const |
755 | { |
756 | return get_shader_resources(active_variables: &active_variables); |
757 | } |
758 | |
759 | bool Compiler::InterfaceVariableAccessHandler::handle(Op opcode, const uint32_t *args, uint32_t length) |
760 | { |
761 | uint32_t variable = 0; |
762 | switch (opcode) |
763 | { |
764 | // Need this first, otherwise, GCC complains about unhandled switch statements. |
765 | default: |
766 | break; |
767 | |
768 | case OpFunctionCall: |
769 | { |
770 | // Invalid SPIR-V. |
771 | if (length < 3) |
772 | return false; |
773 | |
774 | uint32_t count = length - 3; |
775 | args += 3; |
776 | for (uint32_t i = 0; i < count; i++) |
777 | { |
778 | auto *var = compiler.maybe_get<SPIRVariable>(id: args[i]); |
779 | if (var && storage_class_is_interface(storage: var->storage)) |
780 | variables.insert(x: args[i]); |
781 | } |
782 | break; |
783 | } |
784 | |
785 | case OpSelect: |
786 | { |
787 | // Invalid SPIR-V. |
788 | if (length < 5) |
789 | return false; |
790 | |
791 | uint32_t count = length - 3; |
792 | args += 3; |
793 | for (uint32_t i = 0; i < count; i++) |
794 | { |
795 | auto *var = compiler.maybe_get<SPIRVariable>(id: args[i]); |
796 | if (var && storage_class_is_interface(storage: var->storage)) |
797 | variables.insert(x: args[i]); |
798 | } |
799 | break; |
800 | } |
801 | |
802 | case OpPhi: |
803 | { |
804 | // Invalid SPIR-V. |
805 | if (length < 2) |
806 | return false; |
807 | |
808 | uint32_t count = length - 2; |
809 | args += 2; |
810 | for (uint32_t i = 0; i < count; i += 2) |
811 | { |
812 | auto *var = compiler.maybe_get<SPIRVariable>(id: args[i]); |
813 | if (var && storage_class_is_interface(storage: var->storage)) |
814 | variables.insert(x: args[i]); |
815 | } |
816 | break; |
817 | } |
818 | |
819 | case OpAtomicStore: |
820 | case OpStore: |
821 | // Invalid SPIR-V. |
822 | if (length < 1) |
823 | return false; |
824 | variable = args[0]; |
825 | break; |
826 | |
827 | case OpCopyMemory: |
828 | { |
829 | if (length < 2) |
830 | return false; |
831 | |
832 | auto *var = compiler.maybe_get<SPIRVariable>(id: args[0]); |
833 | if (var && storage_class_is_interface(storage: var->storage)) |
834 | variables.insert(x: args[0]); |
835 | |
836 | var = compiler.maybe_get<SPIRVariable>(id: args[1]); |
837 | if (var && storage_class_is_interface(storage: var->storage)) |
838 | variables.insert(x: args[1]); |
839 | break; |
840 | } |
841 | |
842 | case OpExtInst: |
843 | { |
844 | if (length < 3) |
845 | return false; |
846 | auto &extension_set = compiler.get<SPIRExtension>(id: args[2]); |
847 | switch (extension_set.ext) |
848 | { |
849 | case SPIRExtension::GLSL: |
850 | { |
851 | auto op = static_cast<GLSLstd450>(args[3]); |
852 | |
853 | switch (op) |
854 | { |
855 | case GLSLstd450InterpolateAtCentroid: |
856 | case GLSLstd450InterpolateAtSample: |
857 | case GLSLstd450InterpolateAtOffset: |
858 | { |
859 | auto *var = compiler.maybe_get<SPIRVariable>(id: args[4]); |
860 | if (var && storage_class_is_interface(storage: var->storage)) |
861 | variables.insert(x: args[4]); |
862 | break; |
863 | } |
864 | |
865 | case GLSLstd450Modf: |
866 | case GLSLstd450Fract: |
867 | { |
868 | auto *var = compiler.maybe_get<SPIRVariable>(id: args[5]); |
869 | if (var && storage_class_is_interface(storage: var->storage)) |
870 | variables.insert(x: args[5]); |
871 | break; |
872 | } |
873 | |
874 | default: |
875 | break; |
876 | } |
877 | break; |
878 | } |
879 | case SPIRExtension::SPV_AMD_shader_explicit_vertex_parameter: |
880 | { |
881 | enum AMDShaderExplicitVertexParameter |
882 | { |
883 | InterpolateAtVertexAMD = 1 |
884 | }; |
885 | |
886 | auto op = static_cast<AMDShaderExplicitVertexParameter>(args[3]); |
887 | |
888 | switch (op) |
889 | { |
890 | case InterpolateAtVertexAMD: |
891 | { |
892 | auto *var = compiler.maybe_get<SPIRVariable>(id: args[4]); |
893 | if (var && storage_class_is_interface(storage: var->storage)) |
894 | variables.insert(x: args[4]); |
895 | break; |
896 | } |
897 | |
898 | default: |
899 | break; |
900 | } |
901 | break; |
902 | } |
903 | default: |
904 | break; |
905 | } |
906 | break; |
907 | } |
908 | |
909 | case OpAccessChain: |
910 | case OpInBoundsAccessChain: |
911 | case OpPtrAccessChain: |
912 | case OpLoad: |
913 | case OpCopyObject: |
914 | case OpImageTexelPointer: |
915 | case OpAtomicLoad: |
916 | case OpAtomicExchange: |
917 | case OpAtomicCompareExchange: |
918 | case OpAtomicCompareExchangeWeak: |
919 | case OpAtomicIIncrement: |
920 | case OpAtomicIDecrement: |
921 | case OpAtomicIAdd: |
922 | case OpAtomicISub: |
923 | case OpAtomicSMin: |
924 | case OpAtomicUMin: |
925 | case OpAtomicSMax: |
926 | case OpAtomicUMax: |
927 | case OpAtomicAnd: |
928 | case OpAtomicOr: |
929 | case OpAtomicXor: |
930 | case OpArrayLength: |
931 | // Invalid SPIR-V. |
932 | if (length < 3) |
933 | return false; |
934 | variable = args[2]; |
935 | break; |
936 | } |
937 | |
938 | if (variable) |
939 | { |
940 | auto *var = compiler.maybe_get<SPIRVariable>(id: variable); |
941 | if (var && storage_class_is_interface(storage: var->storage)) |
942 | variables.insert(x: variable); |
943 | } |
944 | return true; |
945 | } |
946 | |
947 | unordered_set<VariableID> Compiler::get_active_interface_variables() const |
948 | { |
949 | // Traverse the call graph and find all interface variables which are in use. |
950 | unordered_set<VariableID> variables; |
951 | InterfaceVariableAccessHandler handler(*this, variables); |
952 | traverse_all_reachable_opcodes(block: get<SPIRFunction>(id: ir.default_entry_point), handler); |
953 | |
954 | ir.for_each_typed_id<SPIRVariable>(op: [&](uint32_t, const SPIRVariable &var) { |
955 | if (var.storage != StorageClassOutput) |
956 | return; |
957 | if (!interface_variable_exists_in_entry_point(id: var.self)) |
958 | return; |
959 | |
960 | // An output variable which is just declared (but uninitialized) might be read by subsequent stages |
961 | // so we should force-enable these outputs, |
962 | // since compilation will fail if a subsequent stage attempts to read from the variable in question. |
963 | // Also, make sure we preserve output variables which are only initialized, but never accessed by any code. |
964 | if (var.initializer != ID(0) || get_execution_model() != ExecutionModelFragment) |
965 | variables.insert(x: var.self); |
966 | }); |
967 | |
968 | // If we needed to create one, we'll need it. |
969 | if (dummy_sampler_id) |
970 | variables.insert(x: dummy_sampler_id); |
971 | |
972 | return variables; |
973 | } |
974 | |
975 | void Compiler::set_enabled_interface_variables(std::unordered_set<VariableID> active_variables) |
976 | { |
977 | active_interface_variables = std::move(active_variables); |
978 | check_active_interface_variables = true; |
979 | } |
980 | |
981 | ShaderResources Compiler::get_shader_resources(const unordered_set<VariableID> *active_variables) const |
982 | { |
983 | ShaderResources res; |
984 | |
985 | bool ssbo_instance_name = reflection_ssbo_instance_name_is_significant(); |
986 | |
987 | ir.for_each_typed_id<SPIRVariable>(op: [&](uint32_t, const SPIRVariable &var) { |
988 | auto &type = this->get<SPIRType>(id: var.basetype); |
989 | |
990 | // It is possible for uniform storage classes to be passed as function parameters, so detect |
991 | // that. To detect function parameters, check of StorageClass of variable is function scope. |
992 | if (var.storage == StorageClassFunction || !type.pointer) |
993 | return; |
994 | |
995 | if (active_variables && active_variables->find(x: var.self) == end(cont: *active_variables)) |
996 | return; |
997 | |
998 | // In SPIR-V 1.4 and up, every global must be present in the entry point interface list, |
999 | // not just IO variables. |
1000 | bool active_in_entry_point = true; |
1001 | if (ir.get_spirv_version() < 0x10400) |
1002 | { |
1003 | if (var.storage == StorageClassInput || var.storage == StorageClassOutput) |
1004 | active_in_entry_point = interface_variable_exists_in_entry_point(id: var.self); |
1005 | } |
1006 | else |
1007 | active_in_entry_point = interface_variable_exists_in_entry_point(id: var.self); |
1008 | |
1009 | if (!active_in_entry_point) |
1010 | return; |
1011 | |
1012 | bool is_builtin = is_builtin_variable(var); |
1013 | |
1014 | if (is_builtin) |
1015 | { |
1016 | if (var.storage != StorageClassInput && var.storage != StorageClassOutput) |
1017 | return; |
1018 | |
1019 | auto &list = var.storage == StorageClassInput ? res.builtin_inputs : res.builtin_outputs; |
1020 | BuiltInResource resource; |
1021 | |
1022 | if (has_decoration(id: type.self, decoration: DecorationBlock)) |
1023 | { |
1024 | resource.resource = { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, |
1025 | .name: get_remapped_declared_block_name(id: var.self, fallback_prefer_instance_name: false) }; |
1026 | |
1027 | for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++) |
1028 | { |
1029 | resource.value_type_id = type.member_types[i]; |
1030 | resource.builtin = BuiltIn(get_member_decoration(id: type.self, index: i, decoration: DecorationBuiltIn)); |
1031 | list.push_back(t: resource); |
1032 | } |
1033 | } |
1034 | else |
1035 | { |
1036 | bool strip_array = |
1037 | !has_decoration(id: var.self, decoration: DecorationPatch) && ( |
1038 | get_execution_model() == ExecutionModelTessellationControl || |
1039 | (get_execution_model() == ExecutionModelTessellationEvaluation && |
1040 | var.storage == StorageClassInput)); |
1041 | |
1042 | resource.resource = { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) }; |
1043 | |
1044 | if (strip_array && !type.array.empty()) |
1045 | resource.value_type_id = get_variable_data_type(var).parent_type; |
1046 | else |
1047 | resource.value_type_id = get_variable_data_type_id(var); |
1048 | |
1049 | assert(resource.value_type_id); |
1050 | |
1051 | resource.builtin = BuiltIn(get_decoration(id: var.self, decoration: DecorationBuiltIn)); |
1052 | list.push_back(t: std::move(resource)); |
1053 | } |
1054 | return; |
1055 | } |
1056 | |
1057 | // Input |
1058 | if (var.storage == StorageClassInput) |
1059 | { |
1060 | if (has_decoration(id: type.self, decoration: DecorationBlock)) |
1061 | { |
1062 | res.stage_inputs.push_back( |
1063 | t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, |
1064 | .name: get_remapped_declared_block_name(id: var.self, fallback_prefer_instance_name: false) }); |
1065 | } |
1066 | else |
1067 | res.stage_inputs.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) }); |
1068 | } |
1069 | // Subpass inputs |
1070 | else if (var.storage == StorageClassUniformConstant && type.image.dim == DimSubpassData) |
1071 | { |
1072 | res.subpass_inputs.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) }); |
1073 | } |
1074 | // Outputs |
1075 | else if (var.storage == StorageClassOutput) |
1076 | { |
1077 | if (has_decoration(id: type.self, decoration: DecorationBlock)) |
1078 | { |
1079 | res.stage_outputs.push_back( |
1080 | t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_remapped_declared_block_name(id: var.self, fallback_prefer_instance_name: false) }); |
1081 | } |
1082 | else |
1083 | res.stage_outputs.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) }); |
1084 | } |
1085 | // UBOs |
1086 | else if (type.storage == StorageClassUniform && has_decoration(id: type.self, decoration: DecorationBlock)) |
1087 | { |
1088 | res.uniform_buffers.push_back( |
1089 | t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_remapped_declared_block_name(id: var.self, fallback_prefer_instance_name: false) }); |
1090 | } |
1091 | // Old way to declare SSBOs. |
1092 | else if (type.storage == StorageClassUniform && has_decoration(id: type.self, decoration: DecorationBufferBlock)) |
1093 | { |
1094 | res.storage_buffers.push_back( |
1095 | t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_remapped_declared_block_name(id: var.self, fallback_prefer_instance_name: ssbo_instance_name) }); |
1096 | } |
1097 | // Modern way to declare SSBOs. |
1098 | else if (type.storage == StorageClassStorageBuffer) |
1099 | { |
1100 | res.storage_buffers.push_back( |
1101 | t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_remapped_declared_block_name(id: var.self, fallback_prefer_instance_name: ssbo_instance_name) }); |
1102 | } |
1103 | // Push constant blocks |
1104 | else if (type.storage == StorageClassPushConstant) |
1105 | { |
1106 | // There can only be one push constant block, but keep the vector in case this restriction is lifted |
1107 | // in the future. |
1108 | res.push_constant_buffers.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) }); |
1109 | } |
1110 | else if (type.storage == StorageClassShaderRecordBufferKHR) |
1111 | { |
1112 | res.shader_record_buffers.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_remapped_declared_block_name(id: var.self, fallback_prefer_instance_name: ssbo_instance_name) }); |
1113 | } |
1114 | // Atomic counters |
1115 | else if (type.storage == StorageClassAtomicCounter) |
1116 | { |
1117 | res.atomic_counters.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) }); |
1118 | } |
1119 | else if (type.storage == StorageClassUniformConstant) |
1120 | { |
1121 | if (type.basetype == SPIRType::Image) |
1122 | { |
1123 | // Images |
1124 | if (type.image.sampled == 2) |
1125 | { |
1126 | res.storage_images.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) }); |
1127 | } |
1128 | // Separate images |
1129 | else if (type.image.sampled == 1) |
1130 | { |
1131 | res.separate_images.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) }); |
1132 | } |
1133 | } |
1134 | // Separate samplers |
1135 | else if (type.basetype == SPIRType::Sampler) |
1136 | { |
1137 | res.separate_samplers.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) }); |
1138 | } |
1139 | // Textures |
1140 | else if (type.basetype == SPIRType::SampledImage) |
1141 | { |
1142 | res.sampled_images.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) }); |
1143 | } |
1144 | // Acceleration structures |
1145 | else if (type.basetype == SPIRType::AccelerationStructure) |
1146 | { |
1147 | res.acceleration_structures.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) }); |
1148 | } |
1149 | else |
1150 | { |
1151 | res.gl_plain_uniforms.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) }); |
1152 | } |
1153 | } |
1154 | }); |
1155 | |
1156 | return res; |
1157 | } |
1158 | |
1159 | bool Compiler::type_is_top_level_block(const SPIRType &type) const |
1160 | { |
1161 | if (type.basetype != SPIRType::Struct) |
1162 | return false; |
1163 | return has_decoration(id: type.self, decoration: DecorationBlock) || has_decoration(id: type.self, decoration: DecorationBufferBlock); |
1164 | } |
1165 | |
1166 | bool Compiler::type_is_block_like(const SPIRType &type) const |
1167 | { |
1168 | if (type_is_top_level_block(type)) |
1169 | return true; |
1170 | |
1171 | if (type.basetype == SPIRType::Struct) |
1172 | { |
1173 | // Block-like types may have Offset decorations. |
1174 | for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++) |
1175 | if (has_member_decoration(id: type.self, index: i, decoration: DecorationOffset)) |
1176 | return true; |
1177 | } |
1178 | |
1179 | return false; |
1180 | } |
1181 | |
1182 | void Compiler::parse_fixup() |
1183 | { |
1184 | // Figure out specialization constants for work group sizes. |
1185 | for (auto id_ : ir.ids_for_constant_or_variable) |
1186 | { |
1187 | auto &id = ir.ids[id_]; |
1188 | |
1189 | if (id.get_type() == TypeConstant) |
1190 | { |
1191 | auto &c = id.get<SPIRConstant>(); |
1192 | if (has_decoration(id: c.self, decoration: DecorationBuiltIn) && |
1193 | BuiltIn(get_decoration(id: c.self, decoration: DecorationBuiltIn)) == BuiltInWorkgroupSize) |
1194 | { |
1195 | // In current SPIR-V, there can be just one constant like this. |
1196 | // All entry points will receive the constant value. |
1197 | // WorkgroupSize take precedence over LocalSizeId. |
1198 | for (auto &entry : ir.entry_points) |
1199 | { |
1200 | entry.second.workgroup_size.constant = c.self; |
1201 | entry.second.workgroup_size.x = c.scalar(col: 0, row: 0); |
1202 | entry.second.workgroup_size.y = c.scalar(col: 0, row: 1); |
1203 | entry.second.workgroup_size.z = c.scalar(col: 0, row: 2); |
1204 | } |
1205 | } |
1206 | } |
1207 | else if (id.get_type() == TypeVariable) |
1208 | { |
1209 | auto &var = id.get<SPIRVariable>(); |
1210 | if (var.storage == StorageClassPrivate || var.storage == StorageClassWorkgroup || |
1211 | var.storage == StorageClassTaskPayloadWorkgroupEXT || |
1212 | var.storage == StorageClassOutput) |
1213 | { |
1214 | global_variables.push_back(t: var.self); |
1215 | } |
1216 | if (variable_storage_is_aliased(v: var)) |
1217 | aliased_variables.push_back(t: var.self); |
1218 | } |
1219 | } |
1220 | } |
1221 | |
1222 | void Compiler::update_name_cache(unordered_set<string> &cache_primary, const unordered_set<string> &cache_secondary, |
1223 | string &name) |
1224 | { |
1225 | if (name.empty()) |
1226 | return; |
1227 | |
1228 | const auto find_name = [&](const string &n) -> bool { |
1229 | if (cache_primary.find(x: n) != end(cont&: cache_primary)) |
1230 | return true; |
1231 | |
1232 | if (&cache_primary != &cache_secondary) |
1233 | if (cache_secondary.find(x: n) != end(cont: cache_secondary)) |
1234 | return true; |
1235 | |
1236 | return false; |
1237 | }; |
1238 | |
1239 | const auto insert_name = [&](const string &n) { cache_primary.insert(x: n); }; |
1240 | |
1241 | if (!find_name(name)) |
1242 | { |
1243 | insert_name(name); |
1244 | return; |
1245 | } |
1246 | |
1247 | uint32_t counter = 0; |
1248 | auto tmpname = name; |
1249 | |
1250 | bool use_linked_underscore = true; |
1251 | |
1252 | if (tmpname == "_") |
1253 | { |
1254 | // We cannot just append numbers, as we will end up creating internally reserved names. |
1255 | // Make it like _0_<counter> instead. |
1256 | tmpname += "0"; |
1257 | } |
1258 | else if (tmpname.back() == '_') |
1259 | { |
1260 | // The last_character is an underscore, so we don't need to link in underscore. |
1261 | // This would violate double underscore rules. |
1262 | use_linked_underscore = false; |
1263 | } |
1264 | |
1265 | // If there is a collision (very rare), |
1266 | // keep tacking on extra identifier until it's unique. |
1267 | do |
1268 | { |
1269 | counter++; |
1270 | name = tmpname + (use_linked_underscore ? "_": "") + convert_to_string(t: counter); |
1271 | } while (find_name(name)); |
1272 | insert_name(name); |
1273 | } |
1274 | |
1275 | void Compiler::update_name_cache(unordered_set<string> &cache, string &name) |
1276 | { |
1277 | update_name_cache(cache_primary&: cache, cache_secondary: cache, name); |
1278 | } |
1279 | |
1280 | void Compiler::set_name(ID id, const std::string &name) |
1281 | { |
1282 | ir.set_name(id, name); |
1283 | } |
1284 | |
1285 | const SPIRType &Compiler::get_type(TypeID id) const |
1286 | { |
1287 | return get<SPIRType>(id); |
1288 | } |
1289 | |
1290 | const SPIRType &Compiler::get_type_from_variable(VariableID id) const |
1291 | { |
1292 | return get<SPIRType>(id: get<SPIRVariable>(id).basetype); |
1293 | } |
1294 | |
1295 | uint32_t Compiler::get_pointee_type_id(uint32_t type_id) const |
1296 | { |
1297 | auto *p_type = &get<SPIRType>(id: type_id); |
1298 | if (p_type->pointer) |
1299 | { |
1300 | assert(p_type->parent_type); |
1301 | type_id = p_type->parent_type; |
1302 | } |
1303 | return type_id; |
1304 | } |
1305 | |
1306 | const SPIRType &Compiler::get_pointee_type(const SPIRType &type) const |
1307 | { |
1308 | auto *p_type = &type; |
1309 | if (p_type->pointer) |
1310 | { |
1311 | assert(p_type->parent_type); |
1312 | p_type = &get<SPIRType>(id: p_type->parent_type); |
1313 | } |
1314 | return *p_type; |
1315 | } |
1316 | |
1317 | const SPIRType &Compiler::get_pointee_type(uint32_t type_id) const |
1318 | { |
1319 | return get_pointee_type(type: get<SPIRType>(id: type_id)); |
1320 | } |
1321 | |
1322 | uint32_t Compiler::get_variable_data_type_id(const SPIRVariable &var) const |
1323 | { |
1324 | if (var.phi_variable || var.storage == spv::StorageClass::StorageClassAtomicCounter) |
1325 | return var.basetype; |
1326 | return get_pointee_type_id(type_id: var.basetype); |
1327 | } |
1328 | |
1329 | SPIRType &Compiler::get_variable_data_type(const SPIRVariable &var) |
1330 | { |
1331 | return get<SPIRType>(id: get_variable_data_type_id(var)); |
1332 | } |
1333 | |
1334 | const SPIRType &Compiler::get_variable_data_type(const SPIRVariable &var) const |
1335 | { |
1336 | return get<SPIRType>(id: get_variable_data_type_id(var)); |
1337 | } |
1338 | |
1339 | SPIRType &Compiler::get_variable_element_type(const SPIRVariable &var) |
1340 | { |
1341 | SPIRType *type = &get_variable_data_type(var); |
1342 | if (is_array(type: *type)) |
1343 | type = &get<SPIRType>(id: type->parent_type); |
1344 | return *type; |
1345 | } |
1346 | |
1347 | const SPIRType &Compiler::get_variable_element_type(const SPIRVariable &var) const |
1348 | { |
1349 | const SPIRType *type = &get_variable_data_type(var); |
1350 | if (is_array(type: *type)) |
1351 | type = &get<SPIRType>(id: type->parent_type); |
1352 | return *type; |
1353 | } |
1354 | |
1355 | bool Compiler::is_sampled_image_type(const SPIRType &type) |
1356 | { |
1357 | return (type.basetype == SPIRType::Image || type.basetype == SPIRType::SampledImage) && type.image.sampled == 1 && |
1358 | type.image.dim != DimBuffer; |
1359 | } |
1360 | |
1361 | void Compiler::set_member_decoration_string(TypeID id, uint32_t index, spv::Decoration decoration, |
1362 | const std::string &argument) |
1363 | { |
1364 | ir.set_member_decoration_string(id, index, decoration, argument); |
1365 | } |
1366 | |
1367 | void Compiler::set_member_decoration(TypeID id, uint32_t index, Decoration decoration, uint32_t argument) |
1368 | { |
1369 | ir.set_member_decoration(id, index, decoration, argument); |
1370 | } |
1371 | |
1372 | void Compiler::set_member_name(TypeID id, uint32_t index, const std::string &name) |
1373 | { |
1374 | ir.set_member_name(id, index, name); |
1375 | } |
1376 | |
1377 | const std::string &Compiler::get_member_name(TypeID id, uint32_t index) const |
1378 | { |
1379 | return ir.get_member_name(id, index); |
1380 | } |
1381 | |
1382 | void Compiler::set_qualified_name(uint32_t id, const string &name) |
1383 | { |
1384 | ir.meta[id].decoration.qualified_alias = name; |
1385 | } |
1386 | |
1387 | void Compiler::set_member_qualified_name(uint32_t type_id, uint32_t index, const std::string &name) |
1388 | { |
1389 | ir.meta[type_id].members.resize(new_size: max(a: ir.meta[type_id].members.size(), b: size_t(index) + 1)); |
1390 | ir.meta[type_id].members[index].qualified_alias = name; |
1391 | } |
1392 | |
1393 | const string &Compiler::get_member_qualified_name(TypeID type_id, uint32_t index) const |
1394 | { |
1395 | auto *m = ir.find_meta(id: type_id); |
1396 | if (m && index < m->members.size()) |
1397 | return m->members[index].qualified_alias; |
1398 | else |
1399 | return ir.get_empty_string(); |
1400 | } |
1401 | |
1402 | uint32_t Compiler::get_member_decoration(TypeID id, uint32_t index, Decoration decoration) const |
1403 | { |
1404 | return ir.get_member_decoration(id, index, decoration); |
1405 | } |
1406 | |
1407 | const Bitset &Compiler::get_member_decoration_bitset(TypeID id, uint32_t index) const |
1408 | { |
1409 | return ir.get_member_decoration_bitset(id, index); |
1410 | } |
1411 | |
1412 | bool Compiler::has_member_decoration(TypeID id, uint32_t index, Decoration decoration) const |
1413 | { |
1414 | return ir.has_member_decoration(id, index, decoration); |
1415 | } |
1416 | |
1417 | void Compiler::unset_member_decoration(TypeID id, uint32_t index, Decoration decoration) |
1418 | { |
1419 | ir.unset_member_decoration(id, index, decoration); |
1420 | } |
1421 | |
1422 | void Compiler::set_decoration_string(ID id, spv::Decoration decoration, const std::string &argument) |
1423 | { |
1424 | ir.set_decoration_string(id, decoration, argument); |
1425 | } |
1426 | |
1427 | void Compiler::set_decoration(ID id, Decoration decoration, uint32_t argument) |
1428 | { |
1429 | ir.set_decoration(id, decoration, argument); |
1430 | } |
1431 | |
1432 | void Compiler::set_extended_decoration(uint32_t id, ExtendedDecorations decoration, uint32_t value) |
1433 | { |
1434 | auto &dec = ir.meta[id].decoration; |
1435 | dec.extended.flags.set(decoration); |
1436 | dec.extended.values[decoration] = value; |
1437 | } |
1438 | |
1439 | void Compiler::set_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration, |
1440 | uint32_t value) |
1441 | { |
1442 | ir.meta[type].members.resize(new_size: max(a: ir.meta[type].members.size(), b: size_t(index) + 1)); |
1443 | auto &dec = ir.meta[type].members[index]; |
1444 | dec.extended.flags.set(decoration); |
1445 | dec.extended.values[decoration] = value; |
1446 | } |
1447 | |
1448 | static uint32_t get_default_extended_decoration(ExtendedDecorations decoration) |
1449 | { |
1450 | switch (decoration) |
1451 | { |
1452 | case SPIRVCrossDecorationResourceIndexPrimary: |
1453 | case SPIRVCrossDecorationResourceIndexSecondary: |
1454 | case SPIRVCrossDecorationResourceIndexTertiary: |
1455 | case SPIRVCrossDecorationResourceIndexQuaternary: |
1456 | case SPIRVCrossDecorationInterfaceMemberIndex: |
1457 | return ~(0u); |
1458 | |
1459 | default: |
1460 | return 0; |
1461 | } |
1462 | } |
1463 | |
1464 | uint32_t Compiler::get_extended_decoration(uint32_t id, ExtendedDecorations decoration) const |
1465 | { |
1466 | auto *m = ir.find_meta(id); |
1467 | if (!m) |
1468 | return 0; |
1469 | |
1470 | auto &dec = m->decoration; |
1471 | |
1472 | if (!dec.extended.flags.get(bit: decoration)) |
1473 | return get_default_extended_decoration(decoration); |
1474 | |
1475 | return dec.extended.values[decoration]; |
1476 | } |
1477 | |
1478 | uint32_t Compiler::get_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const |
1479 | { |
1480 | auto *m = ir.find_meta(id: type); |
1481 | if (!m) |
1482 | return 0; |
1483 | |
1484 | if (index >= m->members.size()) |
1485 | return 0; |
1486 | |
1487 | auto &dec = m->members[index]; |
1488 | if (!dec.extended.flags.get(bit: decoration)) |
1489 | return get_default_extended_decoration(decoration); |
1490 | return dec.extended.values[decoration]; |
1491 | } |
1492 | |
1493 | bool Compiler::has_extended_decoration(uint32_t id, ExtendedDecorations decoration) const |
1494 | { |
1495 | auto *m = ir.find_meta(id); |
1496 | if (!m) |
1497 | return false; |
1498 | |
1499 | auto &dec = m->decoration; |
1500 | return dec.extended.flags.get(bit: decoration); |
1501 | } |
1502 | |
1503 | bool Compiler::has_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const |
1504 | { |
1505 | auto *m = ir.find_meta(id: type); |
1506 | if (!m) |
1507 | return false; |
1508 | |
1509 | if (index >= m->members.size()) |
1510 | return false; |
1511 | |
1512 | auto &dec = m->members[index]; |
1513 | return dec.extended.flags.get(bit: decoration); |
1514 | } |
1515 | |
1516 | void Compiler::unset_extended_decoration(uint32_t id, ExtendedDecorations decoration) |
1517 | { |
1518 | auto &dec = ir.meta[id].decoration; |
1519 | dec.extended.flags.clear(bit: decoration); |
1520 | dec.extended.values[decoration] = 0; |
1521 | } |
1522 | |
1523 | void Compiler::unset_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) |
1524 | { |
1525 | ir.meta[type].members.resize(new_size: max(a: ir.meta[type].members.size(), b: size_t(index) + 1)); |
1526 | auto &dec = ir.meta[type].members[index]; |
1527 | dec.extended.flags.clear(bit: decoration); |
1528 | dec.extended.values[decoration] = 0; |
1529 | } |
1530 | |
1531 | StorageClass Compiler::get_storage_class(VariableID id) const |
1532 | { |
1533 | return get<SPIRVariable>(id).storage; |
1534 | } |
1535 | |
1536 | const std::string &Compiler::get_name(ID id) const |
1537 | { |
1538 | return ir.get_name(id); |
1539 | } |
1540 | |
1541 | const std::string Compiler::get_fallback_name(ID id) const |
1542 | { |
1543 | return join(ts: "_", ts&: id); |
1544 | } |
1545 | |
1546 | const std::string Compiler::get_block_fallback_name(VariableID id) const |
1547 | { |
1548 | auto &var = get<SPIRVariable>(id); |
1549 | if (get_name(id).empty()) |
1550 | return join(ts: "_", ts: get<SPIRType>(id: var.basetype).self, ts: "_", ts&: id); |
1551 | else |
1552 | return get_name(id); |
1553 | } |
1554 | |
1555 | const Bitset &Compiler::get_decoration_bitset(ID id) const |
1556 | { |
1557 | return ir.get_decoration_bitset(id); |
1558 | } |
1559 | |
1560 | bool Compiler::has_decoration(ID id, Decoration decoration) const |
1561 | { |
1562 | return ir.has_decoration(id, decoration); |
1563 | } |
1564 | |
1565 | const string &Compiler::get_decoration_string(ID id, Decoration decoration) const |
1566 | { |
1567 | return ir.get_decoration_string(id, decoration); |
1568 | } |
1569 | |
1570 | const string &Compiler::get_member_decoration_string(TypeID id, uint32_t index, Decoration decoration) const |
1571 | { |
1572 | return ir.get_member_decoration_string(id, index, decoration); |
1573 | } |
1574 | |
1575 | uint32_t Compiler::get_decoration(ID id, Decoration decoration) const |
1576 | { |
1577 | return ir.get_decoration(id, decoration); |
1578 | } |
1579 | |
1580 | void Compiler::unset_decoration(ID id, Decoration decoration) |
1581 | { |
1582 | ir.unset_decoration(id, decoration); |
1583 | } |
1584 | |
1585 | bool Compiler::get_binary_offset_for_decoration(VariableID id, spv::Decoration decoration, uint32_t &word_offset) const |
1586 | { |
1587 | auto *m = ir.find_meta(id); |
1588 | if (!m) |
1589 | return false; |
1590 | |
1591 | auto &word_offsets = m->decoration_word_offset; |
1592 | auto itr = word_offsets.find(x: decoration); |
1593 | if (itr == end(cont: word_offsets)) |
1594 | return false; |
1595 | |
1596 | word_offset = itr->second; |
1597 | return true; |
1598 | } |
1599 | |
1600 | bool Compiler::block_is_noop(const SPIRBlock &block) const |
1601 | { |
1602 | if (block.terminator != SPIRBlock::Direct) |
1603 | return false; |
1604 | |
1605 | auto &child = get<SPIRBlock>(id: block.next_block); |
1606 | |
1607 | // If this block participates in PHI, the block isn't really noop. |
1608 | for (auto &phi : block.phi_variables) |
1609 | if (phi.parent == block.self || phi.parent == child.self) |
1610 | return false; |
1611 | |
1612 | for (auto &phi : child.phi_variables) |
1613 | if (phi.parent == block.self) |
1614 | return false; |
1615 | |
1616 | // Verify all instructions have no semantic impact. |
1617 | for (auto &i : block.ops) |
1618 | { |
1619 | auto op = static_cast<Op>(i.op); |
1620 | |
1621 | switch (op) |
1622 | { |
1623 | // Non-Semantic instructions. |
1624 | case OpLine: |
1625 | case OpNoLine: |
1626 | break; |
1627 | |
1628 | case OpExtInst: |
1629 | { |
1630 | auto *ops = stream(instr: i); |
1631 | auto ext = get<SPIRExtension>(id: ops[2]).ext; |
1632 | |
1633 | bool ext_is_nonsemantic_only = |
1634 | ext == SPIRExtension::NonSemanticShaderDebugInfo || |
1635 | ext == SPIRExtension::SPV_debug_info || |
1636 | ext == SPIRExtension::NonSemanticGeneric; |
1637 | |
1638 | if (!ext_is_nonsemantic_only) |
1639 | return false; |
1640 | |
1641 | break; |
1642 | } |
1643 | |
1644 | default: |
1645 | return false; |
1646 | } |
1647 | } |
1648 | |
1649 | return true; |
1650 | } |
1651 | |
1652 | bool Compiler::block_is_loop_candidate(const SPIRBlock &block, SPIRBlock::Method method) const |
1653 | { |
1654 | // Tried and failed. |
1655 | if (block.disable_block_optimization || block.complex_continue) |
1656 | return false; |
1657 | |
1658 | if (method == SPIRBlock::MergeToSelectForLoop || method == SPIRBlock::MergeToSelectContinueForLoop) |
1659 | { |
1660 | // Try to detect common for loop pattern |
1661 | // which the code backend can use to create cleaner code. |
1662 | // for(;;) { if (cond) { some_body; } else { break; } } |
1663 | // is the pattern we're looking for. |
1664 | const auto *false_block = maybe_get<SPIRBlock>(id: block.false_block); |
1665 | const auto *true_block = maybe_get<SPIRBlock>(id: block.true_block); |
1666 | const auto *merge_block = maybe_get<SPIRBlock>(id: block.merge_block); |
1667 | |
1668 | bool false_block_is_merge = block.false_block == block.merge_block || |
1669 | (false_block && merge_block && execution_is_noop(from: *false_block, to: *merge_block)); |
1670 | |
1671 | bool true_block_is_merge = block.true_block == block.merge_block || |
1672 | (true_block && merge_block && execution_is_noop(from: *true_block, to: *merge_block)); |
1673 | |
1674 | bool positive_candidate = |
1675 | block.true_block != block.merge_block && block.true_block != block.self && false_block_is_merge; |
1676 | |
1677 | bool negative_candidate = |
1678 | block.false_block != block.merge_block && block.false_block != block.self && true_block_is_merge; |
1679 | |
1680 | bool ret = block.terminator == SPIRBlock::Select && block.merge == SPIRBlock::MergeLoop && |
1681 | (positive_candidate || negative_candidate); |
1682 | |
1683 | if (ret && positive_candidate && method == SPIRBlock::MergeToSelectContinueForLoop) |
1684 | ret = block.true_block == block.continue_block; |
1685 | else if (ret && negative_candidate && method == SPIRBlock::MergeToSelectContinueForLoop) |
1686 | ret = block.false_block == block.continue_block; |
1687 | |
1688 | // If we have OpPhi which depends on branches which came from our own block, |
1689 | // we need to flush phi variables in else block instead of a trivial break, |
1690 | // so we cannot assume this is a for loop candidate. |
1691 | if (ret) |
1692 | { |
1693 | for (auto &phi : block.phi_variables) |
1694 | if (phi.parent == block.self) |
1695 | return false; |
1696 | |
1697 | auto *merge = maybe_get<SPIRBlock>(id: block.merge_block); |
1698 | if (merge) |
1699 | for (auto &phi : merge->phi_variables) |
1700 | if (phi.parent == block.self) |
1701 | return false; |
1702 | } |
1703 | return ret; |
1704 | } |
1705 | else if (method == SPIRBlock::MergeToDirectForLoop) |
1706 | { |
1707 | // Empty loop header that just sets up merge target |
1708 | // and branches to loop body. |
1709 | bool ret = block.terminator == SPIRBlock::Direct && block.merge == SPIRBlock::MergeLoop && block_is_noop(block); |
1710 | |
1711 | if (!ret) |
1712 | return false; |
1713 | |
1714 | auto &child = get<SPIRBlock>(id: block.next_block); |
1715 | |
1716 | const auto *false_block = maybe_get<SPIRBlock>(id: child.false_block); |
1717 | const auto *true_block = maybe_get<SPIRBlock>(id: child.true_block); |
1718 | const auto *merge_block = maybe_get<SPIRBlock>(id: block.merge_block); |
1719 | |
1720 | bool false_block_is_merge = child.false_block == block.merge_block || |
1721 | (false_block && merge_block && execution_is_noop(from: *false_block, to: *merge_block)); |
1722 | |
1723 | bool true_block_is_merge = child.true_block == block.merge_block || |
1724 | (true_block && merge_block && execution_is_noop(from: *true_block, to: *merge_block)); |
1725 | |
1726 | bool positive_candidate = |
1727 | child.true_block != block.merge_block && child.true_block != block.self && false_block_is_merge; |
1728 | |
1729 | bool negative_candidate = |
1730 | child.false_block != block.merge_block && child.false_block != block.self && true_block_is_merge; |
1731 | |
1732 | ret = child.terminator == SPIRBlock::Select && child.merge == SPIRBlock::MergeNone && |
1733 | (positive_candidate || negative_candidate); |
1734 | |
1735 | if (ret) |
1736 | { |
1737 | auto *merge = maybe_get<SPIRBlock>(id: block.merge_block); |
1738 | if (merge) |
1739 | for (auto &phi : merge->phi_variables) |
1740 | if (phi.parent == block.self || phi.parent == child.false_block) |
1741 | return false; |
1742 | } |
1743 | |
1744 | return ret; |
1745 | } |
1746 | else |
1747 | return false; |
1748 | } |
1749 | |
1750 | bool Compiler::execution_is_noop(const SPIRBlock &from, const SPIRBlock &to) const |
1751 | { |
1752 | if (!execution_is_branchless(from, to)) |
1753 | return false; |
1754 | |
1755 | auto *start = &from; |
1756 | for (;;) |
1757 | { |
1758 | if (start->self == to.self) |
1759 | return true; |
1760 | |
1761 | if (!block_is_noop(block: *start)) |
1762 | return false; |
1763 | |
1764 | auto &next = get<SPIRBlock>(id: start->next_block); |
1765 | start = &next; |
1766 | } |
1767 | } |
1768 | |
1769 | bool Compiler::execution_is_branchless(const SPIRBlock &from, const SPIRBlock &to) const |
1770 | { |
1771 | auto *start = &from; |
1772 | for (;;) |
1773 | { |
1774 | if (start->self == to.self) |
1775 | return true; |
1776 | |
1777 | if (start->terminator == SPIRBlock::Direct && start->merge == SPIRBlock::MergeNone) |
1778 | start = &get<SPIRBlock>(id: start->next_block); |
1779 | else |
1780 | return false; |
1781 | } |
1782 | } |
1783 | |
1784 | bool Compiler::execution_is_direct_branch(const SPIRBlock &from, const SPIRBlock &to) const |
1785 | { |
1786 | return from.terminator == SPIRBlock::Direct && from.merge == SPIRBlock::MergeNone && from.next_block == to.self; |
1787 | } |
1788 | |
1789 | SPIRBlock::ContinueBlockType Compiler::continue_block_type(const SPIRBlock &block) const |
1790 | { |
1791 | // The block was deemed too complex during code emit, pick conservative fallback paths. |
1792 | if (block.complex_continue) |
1793 | return SPIRBlock::ComplexLoop; |
1794 | |
1795 | // In older glslang output continue block can be equal to the loop header. |
1796 | // In this case, execution is clearly branchless, so just assume a while loop header here. |
1797 | if (block.merge == SPIRBlock::MergeLoop) |
1798 | return SPIRBlock::WhileLoop; |
1799 | |
1800 | if (block.loop_dominator == BlockID(SPIRBlock::NoDominator)) |
1801 | { |
1802 | // Continue block is never reached from CFG. |
1803 | return SPIRBlock::ComplexLoop; |
1804 | } |
1805 | |
1806 | auto &dominator = get<SPIRBlock>(id: block.loop_dominator); |
1807 | |
1808 | if (execution_is_noop(from: block, to: dominator)) |
1809 | return SPIRBlock::WhileLoop; |
1810 | else if (execution_is_branchless(from: block, to: dominator)) |
1811 | return SPIRBlock::ForLoop; |
1812 | else |
1813 | { |
1814 | const auto *false_block = maybe_get<SPIRBlock>(id: block.false_block); |
1815 | const auto *true_block = maybe_get<SPIRBlock>(id: block.true_block); |
1816 | const auto *merge_block = maybe_get<SPIRBlock>(id: dominator.merge_block); |
1817 | |
1818 | // If we need to flush Phi in this block, we cannot have a DoWhile loop. |
1819 | bool flush_phi_to_false = false_block && flush_phi_required(from: block.self, to: block.false_block); |
1820 | bool flush_phi_to_true = true_block && flush_phi_required(from: block.self, to: block.true_block); |
1821 | if (flush_phi_to_false || flush_phi_to_true) |
1822 | return SPIRBlock::ComplexLoop; |
1823 | |
1824 | bool positive_do_while = block.true_block == dominator.self && |
1825 | (block.false_block == dominator.merge_block || |
1826 | (false_block && merge_block && execution_is_noop(from: *false_block, to: *merge_block))); |
1827 | |
1828 | bool negative_do_while = block.false_block == dominator.self && |
1829 | (block.true_block == dominator.merge_block || |
1830 | (true_block && merge_block && execution_is_noop(from: *true_block, to: *merge_block))); |
1831 | |
1832 | if (block.merge == SPIRBlock::MergeNone && block.terminator == SPIRBlock::Select && |
1833 | (positive_do_while || negative_do_while)) |
1834 | { |
1835 | return SPIRBlock::DoWhileLoop; |
1836 | } |
1837 | else |
1838 | return SPIRBlock::ComplexLoop; |
1839 | } |
1840 | } |
1841 | |
1842 | const SmallVector<SPIRBlock::Case> &Compiler::get_case_list(const SPIRBlock &block) const |
1843 | { |
1844 | uint32_t width = 0; |
1845 | |
1846 | // First we check if we can get the type directly from the block.condition |
1847 | // since it can be a SPIRConstant or a SPIRVariable. |
1848 | if (const auto *constant = maybe_get<SPIRConstant>(id: block.condition)) |
1849 | { |
1850 | const auto &type = get<SPIRType>(id: constant->constant_type); |
1851 | width = type.width; |
1852 | } |
1853 | else if (const auto *var = maybe_get<SPIRVariable>(id: block.condition)) |
1854 | { |
1855 | const auto &type = get<SPIRType>(id: var->basetype); |
1856 | width = type.width; |
1857 | } |
1858 | else if (const auto *undef = maybe_get<SPIRUndef>(id: block.condition)) |
1859 | { |
1860 | const auto &type = get<SPIRType>(id: undef->basetype); |
1861 | width = type.width; |
1862 | } |
1863 | else |
1864 | { |
1865 | auto search = ir.load_type_width.find(x: block.condition); |
1866 | if (search == ir.load_type_width.end()) |
1867 | { |
1868 | SPIRV_CROSS_THROW("Use of undeclared variable on a switch statement."); |
1869 | } |
1870 | |
1871 | width = search->second; |
1872 | } |
1873 | |
1874 | if (width > 32) |
1875 | return block.cases_64bit; |
1876 | |
1877 | return block.cases_32bit; |
1878 | } |
1879 | |
1880 | bool Compiler::traverse_all_reachable_opcodes(const SPIRBlock &block, OpcodeHandler &handler) const |
1881 | { |
1882 | handler.set_current_block(block); |
1883 | handler.rearm_current_block(block); |
1884 | |
1885 | // Ideally, perhaps traverse the CFG instead of all blocks in order to eliminate dead blocks, |
1886 | // but this shouldn't be a problem in practice unless the SPIR-V is doing insane things like recursing |
1887 | // inside dead blocks ... |
1888 | for (auto &i : block.ops) |
1889 | { |
1890 | auto ops = stream(instr: i); |
1891 | auto op = static_cast<Op>(i.op); |
1892 | |
1893 | if (!handler.handle(opcode: op, args: ops, length: i.length)) |
1894 | return false; |
1895 | |
1896 | if (op == OpFunctionCall) |
1897 | { |
1898 | auto &func = get<SPIRFunction>(id: ops[2]); |
1899 | if (handler.follow_function_call(func)) |
1900 | { |
1901 | if (!handler.begin_function_scope(ops, i.length)) |
1902 | return false; |
1903 | if (!traverse_all_reachable_opcodes(block: get<SPIRFunction>(id: ops[2]), handler)) |
1904 | return false; |
1905 | if (!handler.end_function_scope(ops, i.length)) |
1906 | return false; |
1907 | |
1908 | handler.rearm_current_block(block); |
1909 | } |
1910 | } |
1911 | } |
1912 | |
1913 | if (!handler.handle_terminator(block)) |
1914 | return false; |
1915 | |
1916 | return true; |
1917 | } |
1918 | |
1919 | bool Compiler::traverse_all_reachable_opcodes(const SPIRFunction &func, OpcodeHandler &handler) const |
1920 | { |
1921 | for (auto block : func.blocks) |
1922 | if (!traverse_all_reachable_opcodes(block: get<SPIRBlock>(id: block), handler)) |
1923 | return false; |
1924 | |
1925 | return true; |
1926 | } |
1927 | |
1928 | uint32_t Compiler::type_struct_member_offset(const SPIRType &type, uint32_t index) const |
1929 | { |
1930 | auto *type_meta = ir.find_meta(id: type.self); |
1931 | if (type_meta) |
1932 | { |
1933 | // Decoration must be set in valid SPIR-V, otherwise throw. |
1934 | auto &dec = type_meta->members[index]; |
1935 | if (dec.decoration_flags.get(bit: DecorationOffset)) |
1936 | return dec.offset; |
1937 | else |
1938 | SPIRV_CROSS_THROW("Struct member does not have Offset set."); |
1939 | } |
1940 | else |
1941 | SPIRV_CROSS_THROW("Struct member does not have Offset set."); |
1942 | } |
1943 | |
1944 | uint32_t Compiler::type_struct_member_array_stride(const SPIRType &type, uint32_t index) const |
1945 | { |
1946 | auto *type_meta = ir.find_meta(id: type.member_types[index]); |
1947 | if (type_meta) |
1948 | { |
1949 | // Decoration must be set in valid SPIR-V, otherwise throw. |
1950 | // ArrayStride is part of the array type not OpMemberDecorate. |
1951 | auto &dec = type_meta->decoration; |
1952 | if (dec.decoration_flags.get(bit: DecorationArrayStride)) |
1953 | return dec.array_stride; |
1954 | else |
1955 | SPIRV_CROSS_THROW("Struct member does not have ArrayStride set."); |
1956 | } |
1957 | else |
1958 | SPIRV_CROSS_THROW("Struct member does not have ArrayStride set."); |
1959 | } |
1960 | |
1961 | uint32_t Compiler::type_struct_member_matrix_stride(const SPIRType &type, uint32_t index) const |
1962 | { |
1963 | auto *type_meta = ir.find_meta(id: type.self); |
1964 | if (type_meta) |
1965 | { |
1966 | // Decoration must be set in valid SPIR-V, otherwise throw. |
1967 | // MatrixStride is part of OpMemberDecorate. |
1968 | auto &dec = type_meta->members[index]; |
1969 | if (dec.decoration_flags.get(bit: DecorationMatrixStride)) |
1970 | return dec.matrix_stride; |
1971 | else |
1972 | SPIRV_CROSS_THROW("Struct member does not have MatrixStride set."); |
1973 | } |
1974 | else |
1975 | SPIRV_CROSS_THROW("Struct member does not have MatrixStride set."); |
1976 | } |
1977 | |
1978 | size_t Compiler::get_declared_struct_size(const SPIRType &type) const |
1979 | { |
1980 | if (type.member_types.empty()) |
1981 | SPIRV_CROSS_THROW("Declared struct in block cannot be empty."); |
1982 | |
1983 | // Offsets can be declared out of order, so we need to deduce the actual size |
1984 | // based on last member instead. |
1985 | uint32_t member_index = 0; |
1986 | size_t highest_offset = 0; |
1987 | for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++) |
1988 | { |
1989 | size_t offset = type_struct_member_offset(type, index: i); |
1990 | if (offset > highest_offset) |
1991 | { |
1992 | highest_offset = offset; |
1993 | member_index = i; |
1994 | } |
1995 | } |
1996 | |
1997 | size_t size = get_declared_struct_member_size(struct_type: type, index: member_index); |
1998 | return highest_offset + size; |
1999 | } |
2000 | |
2001 | size_t Compiler::get_declared_struct_size_runtime_array(const SPIRType &type, size_t array_size) const |
2002 | { |
2003 | if (type.member_types.empty()) |
2004 | SPIRV_CROSS_THROW("Declared struct in block cannot be empty."); |
2005 | |
2006 | size_t size = get_declared_struct_size(type); |
2007 | auto &last_type = get<SPIRType>(id: type.member_types.back()); |
2008 | if (!last_type.array.empty() && last_type.array_size_literal[0] && last_type.array[0] == 0) // Runtime array |
2009 | size += array_size * type_struct_member_array_stride(type, index: uint32_t(type.member_types.size() - 1)); |
2010 | |
2011 | return size; |
2012 | } |
2013 | |
2014 | uint32_t Compiler::evaluate_spec_constant_u32(const SPIRConstantOp &spec) const |
2015 | { |
2016 | auto &result_type = get<SPIRType>(id: spec.basetype); |
2017 | if (result_type.basetype != SPIRType::UInt && result_type.basetype != SPIRType::Int && |
2018 | result_type.basetype != SPIRType::Boolean) |
2019 | { |
2020 | SPIRV_CROSS_THROW( |
2021 | "Only 32-bit integers and booleans are currently supported when evaluating specialization constants.\n"); |
2022 | } |
2023 | |
2024 | if (!is_scalar(type: result_type)) |
2025 | SPIRV_CROSS_THROW("Spec constant evaluation must be a scalar.\n"); |
2026 | |
2027 | uint32_t value = 0; |
2028 | |
2029 | const auto eval_u32 = [&](uint32_t id) -> uint32_t { |
2030 | auto &type = expression_type(id); |
2031 | if (type.basetype != SPIRType::UInt && type.basetype != SPIRType::Int && type.basetype != SPIRType::Boolean) |
2032 | { |
2033 | SPIRV_CROSS_THROW("Only 32-bit integers and booleans are currently supported when evaluating " |
2034 | "specialization constants.\n"); |
2035 | } |
2036 | |
2037 | if (!is_scalar(type)) |
2038 | SPIRV_CROSS_THROW("Spec constant evaluation must be a scalar.\n"); |
2039 | if (const auto *c = this->maybe_get<SPIRConstant>(id)) |
2040 | return c->scalar(); |
2041 | else |
2042 | return evaluate_spec_constant_u32(spec: this->get<SPIRConstantOp>(id)); |
2043 | }; |
2044 | |
2045 | #define binary_spec_op(op, binary_op) \ |
2046 | case Op##op: \ |
2047 | value = eval_u32(spec.arguments[0]) binary_op eval_u32(spec.arguments[1]); \ |
2048 | break |
2049 | #define binary_spec_op_cast(op, binary_op, type) \ |
2050 | case Op##op: \ |
2051 | value = uint32_t(type(eval_u32(spec.arguments[0])) binary_op type(eval_u32(spec.arguments[1]))); \ |
2052 | break |
2053 | |
2054 | // Support the basic opcodes which are typically used when computing array sizes. |
2055 | switch (spec.opcode) |
2056 | { |
2057 | binary_spec_op(IAdd, +); |
2058 | binary_spec_op(ISub, -); |
2059 | binary_spec_op(IMul, *); |
2060 | binary_spec_op(BitwiseAnd, &); |
2061 | binary_spec_op(BitwiseOr, |); |
2062 | binary_spec_op(BitwiseXor, ^); |
2063 | binary_spec_op(LogicalAnd, &); |
2064 | binary_spec_op(LogicalOr, |); |
2065 | binary_spec_op(ShiftLeftLogical, <<); |
2066 | binary_spec_op(ShiftRightLogical, >>); |
2067 | binary_spec_op_cast(ShiftRightArithmetic, >>, int32_t); |
2068 | binary_spec_op(LogicalEqual, ==); |
2069 | binary_spec_op(LogicalNotEqual, !=); |
2070 | binary_spec_op(IEqual, ==); |
2071 | binary_spec_op(INotEqual, !=); |
2072 | binary_spec_op(ULessThan, <); |
2073 | binary_spec_op(ULessThanEqual, <=); |
2074 | binary_spec_op(UGreaterThan, >); |
2075 | binary_spec_op(UGreaterThanEqual, >=); |
2076 | binary_spec_op_cast(SLessThan, <, int32_t); |
2077 | binary_spec_op_cast(SLessThanEqual, <=, int32_t); |
2078 | binary_spec_op_cast(SGreaterThan, >, int32_t); |
2079 | binary_spec_op_cast(SGreaterThanEqual, >=, int32_t); |
2080 | #undef binary_spec_op |
2081 | #undef binary_spec_op_cast |
2082 | |
2083 | case OpLogicalNot: |
2084 | value = uint32_t(!eval_u32(spec.arguments[0])); |
2085 | break; |
2086 | |
2087 | case OpNot: |
2088 | value = ~eval_u32(spec.arguments[0]); |
2089 | break; |
2090 | |
2091 | case OpSNegate: |
2092 | value = uint32_t(-int32_t(eval_u32(spec.arguments[0]))); |
2093 | break; |
2094 | |
2095 | case OpSelect: |
2096 | value = eval_u32(spec.arguments[0]) ? eval_u32(spec.arguments[1]) : eval_u32(spec.arguments[2]); |
2097 | break; |
2098 | |
2099 | case OpUMod: |
2100 | { |
2101 | uint32_t a = eval_u32(spec.arguments[0]); |
2102 | uint32_t b = eval_u32(spec.arguments[1]); |
2103 | if (b == 0) |
2104 | SPIRV_CROSS_THROW("Undefined behavior in UMod, b == 0.\n"); |
2105 | value = a % b; |
2106 | break; |
2107 | } |
2108 | |
2109 | case OpSRem: |
2110 | { |
2111 | auto a = int32_t(eval_u32(spec.arguments[0])); |
2112 | auto b = int32_t(eval_u32(spec.arguments[1])); |
2113 | if (b == 0) |
2114 | SPIRV_CROSS_THROW("Undefined behavior in SRem, b == 0.\n"); |
2115 | value = a % b; |
2116 | break; |
2117 | } |
2118 | |
2119 | case OpSMod: |
2120 | { |
2121 | auto a = int32_t(eval_u32(spec.arguments[0])); |
2122 | auto b = int32_t(eval_u32(spec.arguments[1])); |
2123 | if (b == 0) |
2124 | SPIRV_CROSS_THROW("Undefined behavior in SMod, b == 0.\n"); |
2125 | auto v = a % b; |
2126 | |
2127 | // Makes sure we match the sign of b, not a. |
2128 | if ((b < 0 && v > 0) || (b > 0 && v < 0)) |
2129 | v += b; |
2130 | value = v; |
2131 | break; |
2132 | } |
2133 | |
2134 | case OpUDiv: |
2135 | { |
2136 | uint32_t a = eval_u32(spec.arguments[0]); |
2137 | uint32_t b = eval_u32(spec.arguments[1]); |
2138 | if (b == 0) |
2139 | SPIRV_CROSS_THROW("Undefined behavior in UDiv, b == 0.\n"); |
2140 | value = a / b; |
2141 | break; |
2142 | } |
2143 | |
2144 | case OpSDiv: |
2145 | { |
2146 | auto a = int32_t(eval_u32(spec.arguments[0])); |
2147 | auto b = int32_t(eval_u32(spec.arguments[1])); |
2148 | if (b == 0) |
2149 | SPIRV_CROSS_THROW("Undefined behavior in SDiv, b == 0.\n"); |
2150 | value = a / b; |
2151 | break; |
2152 | } |
2153 | |
2154 | default: |
2155 | SPIRV_CROSS_THROW("Unsupported spec constant opcode for evaluation.\n"); |
2156 | } |
2157 | |
2158 | return value; |
2159 | } |
2160 | |
2161 | uint32_t Compiler::evaluate_constant_u32(uint32_t id) const |
2162 | { |
2163 | if (const auto *c = maybe_get<SPIRConstant>(id)) |
2164 | return c->scalar(); |
2165 | else |
2166 | return evaluate_spec_constant_u32(spec: get<SPIRConstantOp>(id)); |
2167 | } |
2168 | |
2169 | size_t Compiler::get_declared_struct_member_size(const SPIRType &struct_type, uint32_t index) const |
2170 | { |
2171 | if (struct_type.member_types.empty()) |
2172 | SPIRV_CROSS_THROW("Declared struct in block cannot be empty."); |
2173 | |
2174 | auto &flags = get_member_decoration_bitset(id: struct_type.self, index); |
2175 | auto &type = get<SPIRType>(id: struct_type.member_types[index]); |
2176 | |
2177 | switch (type.basetype) |
2178 | { |
2179 | case SPIRType::Unknown: |
2180 | case SPIRType::Void: |
2181 | case SPIRType::Boolean: // Bools are purely logical, and cannot be used for externally visible types. |
2182 | case SPIRType::AtomicCounter: |
2183 | case SPIRType::Image: |
2184 | case SPIRType::SampledImage: |
2185 | case SPIRType::Sampler: |
2186 | SPIRV_CROSS_THROW("Querying size for object with opaque size."); |
2187 | |
2188 | default: |
2189 | break; |
2190 | } |
2191 | |
2192 | if (type.pointer && type.storage == StorageClassPhysicalStorageBuffer) |
2193 | { |
2194 | // Check if this is a top-level pointer type, and not an array of pointers. |
2195 | if (type.pointer_depth > get<SPIRType>(id: type.parent_type).pointer_depth) |
2196 | return 8; |
2197 | } |
2198 | |
2199 | if (!type.array.empty()) |
2200 | { |
2201 | // For arrays, we can use ArrayStride to get an easy check. |
2202 | bool array_size_literal = type.array_size_literal.back(); |
2203 | uint32_t array_size = array_size_literal ? type.array.back() : evaluate_constant_u32(id: type.array.back()); |
2204 | return type_struct_member_array_stride(type: struct_type, index) * array_size; |
2205 | } |
2206 | else if (type.basetype == SPIRType::Struct) |
2207 | { |
2208 | return get_declared_struct_size(type); |
2209 | } |
2210 | else |
2211 | { |
2212 | unsigned vecsize = type.vecsize; |
2213 | unsigned columns = type.columns; |
2214 | |
2215 | // Vectors. |
2216 | if (columns == 1) |
2217 | { |
2218 | size_t component_size = type.width / 8; |
2219 | return vecsize * component_size; |
2220 | } |
2221 | else |
2222 | { |
2223 | uint32_t matrix_stride = type_struct_member_matrix_stride(type: struct_type, index); |
2224 | |
2225 | // Per SPIR-V spec, matrices must be tightly packed and aligned up for vec3 accesses. |
2226 | if (flags.get(bit: DecorationRowMajor)) |
2227 | return matrix_stride * vecsize; |
2228 | else if (flags.get(bit: DecorationColMajor)) |
2229 | return matrix_stride * columns; |
2230 | else |
2231 | SPIRV_CROSS_THROW("Either row-major or column-major must be declared for matrices."); |
2232 | } |
2233 | } |
2234 | } |
2235 | |
2236 | bool Compiler::BufferAccessHandler::handle(Op opcode, const uint32_t *args, uint32_t length) |
2237 | { |
2238 | if (opcode != OpAccessChain && opcode != OpInBoundsAccessChain && opcode != OpPtrAccessChain) |
2239 | return true; |
2240 | |
2241 | bool ptr_chain = (opcode == OpPtrAccessChain); |
2242 | |
2243 | // Invalid SPIR-V. |
2244 | if (length < (ptr_chain ? 5u : 4u)) |
2245 | return false; |
2246 | |
2247 | if (args[2] != id) |
2248 | return true; |
2249 | |
2250 | // Don't bother traversing the entire access chain tree yet. |
2251 | // If we access a struct member, assume we access the entire member. |
2252 | uint32_t index = compiler.get<SPIRConstant>(id: args[ptr_chain ? 4 : 3]).scalar(); |
2253 | |
2254 | // Seen this index already. |
2255 | if (seen.find(x: index) != end(cont&: seen)) |
2256 | return true; |
2257 | seen.insert(x: index); |
2258 | |
2259 | auto &type = compiler.expression_type(id); |
2260 | uint32_t offset = compiler.type_struct_member_offset(type, index); |
2261 | |
2262 | size_t range; |
2263 | // If we have another member in the struct, deduce the range by looking at the next member. |
2264 | // This is okay since structs in SPIR-V can have padding, but Offset decoration must be |
2265 | // monotonically increasing. |
2266 | // Of course, this doesn't take into account if the SPIR-V for some reason decided to add |
2267 | // very large amounts of padding, but that's not really a big deal. |
2268 | if (index + 1 < type.member_types.size()) |
2269 | { |
2270 | range = compiler.type_struct_member_offset(type, index: index + 1) - offset; |
2271 | } |
2272 | else |
2273 | { |
2274 | // No padding, so just deduce it from the size of the member directly. |
2275 | range = compiler.get_declared_struct_member_size(struct_type: type, index); |
2276 | } |
2277 | |
2278 | ranges.push_back(t: { .index: index, .offset: offset, .range: range }); |
2279 | return true; |
2280 | } |
2281 | |
2282 | SmallVector<BufferRange> Compiler::get_active_buffer_ranges(VariableID id) const |
2283 | { |
2284 | SmallVector<BufferRange> ranges; |
2285 | BufferAccessHandler handler(*this, ranges, id); |
2286 | traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler); |
2287 | return ranges; |
2288 | } |
2289 | |
2290 | bool Compiler::types_are_logically_equivalent(const SPIRType &a, const SPIRType &b) const |
2291 | { |
2292 | if (a.basetype != b.basetype) |
2293 | return false; |
2294 | if (a.width != b.width) |
2295 | return false; |
2296 | if (a.vecsize != b.vecsize) |
2297 | return false; |
2298 | if (a.columns != b.columns) |
2299 | return false; |
2300 | if (a.array.size() != b.array.size()) |
2301 | return false; |
2302 | |
2303 | size_t array_count = a.array.size(); |
2304 | if (array_count && memcmp(s1: a.array.data(), s2: b.array.data(), n: array_count * sizeof(uint32_t)) != 0) |
2305 | return false; |
2306 | |
2307 | if (a.basetype == SPIRType::Image || a.basetype == SPIRType::SampledImage) |
2308 | { |
2309 | if (memcmp(s1: &a.image, s2: &b.image, n: sizeof(SPIRType::Image)) != 0) |
2310 | return false; |
2311 | } |
2312 | |
2313 | if (a.member_types.size() != b.member_types.size()) |
2314 | return false; |
2315 | |
2316 | size_t member_types = a.member_types.size(); |
2317 | for (size_t i = 0; i < member_types; i++) |
2318 | { |
2319 | if (!types_are_logically_equivalent(a: get<SPIRType>(id: a.member_types[i]), b: get<SPIRType>(id: b.member_types[i]))) |
2320 | return false; |
2321 | } |
2322 | |
2323 | return true; |
2324 | } |
2325 | |
2326 | const Bitset &Compiler::get_execution_mode_bitset() const |
2327 | { |
2328 | return get_entry_point().flags; |
2329 | } |
2330 | |
2331 | void Compiler::set_execution_mode(ExecutionMode mode, uint32_t arg0, uint32_t arg1, uint32_t arg2) |
2332 | { |
2333 | auto &execution = get_entry_point(); |
2334 | |
2335 | execution.flags.set(mode); |
2336 | switch (mode) |
2337 | { |
2338 | case ExecutionModeLocalSize: |
2339 | execution.workgroup_size.x = arg0; |
2340 | execution.workgroup_size.y = arg1; |
2341 | execution.workgroup_size.z = arg2; |
2342 | break; |
2343 | |
2344 | case ExecutionModeLocalSizeId: |
2345 | execution.workgroup_size.id_x = arg0; |
2346 | execution.workgroup_size.id_y = arg1; |
2347 | execution.workgroup_size.id_z = arg2; |
2348 | break; |
2349 | |
2350 | case ExecutionModeInvocations: |
2351 | execution.invocations = arg0; |
2352 | break; |
2353 | |
2354 | case ExecutionModeOutputVertices: |
2355 | execution.output_vertices = arg0; |
2356 | break; |
2357 | |
2358 | case ExecutionModeOutputPrimitivesEXT: |
2359 | execution.output_primitives = arg0; |
2360 | break; |
2361 | |
2362 | default: |
2363 | break; |
2364 | } |
2365 | } |
2366 | |
2367 | void Compiler::unset_execution_mode(ExecutionMode mode) |
2368 | { |
2369 | auto &execution = get_entry_point(); |
2370 | execution.flags.clear(bit: mode); |
2371 | } |
2372 | |
2373 | uint32_t Compiler::get_work_group_size_specialization_constants(SpecializationConstant &x, SpecializationConstant &y, |
2374 | SpecializationConstant &z) const |
2375 | { |
2376 | auto &execution = get_entry_point(); |
2377 | x = { .id: 0, .constant_id: 0 }; |
2378 | y = { .id: 0, .constant_id: 0 }; |
2379 | z = { .id: 0, .constant_id: 0 }; |
2380 | |
2381 | // WorkgroupSize builtin takes precedence over LocalSize / LocalSizeId. |
2382 | if (execution.workgroup_size.constant != 0) |
2383 | { |
2384 | auto &c = get<SPIRConstant>(id: execution.workgroup_size.constant); |
2385 | |
2386 | if (c.m.c[0].id[0] != ID(0)) |
2387 | { |
2388 | x.id = c.m.c[0].id[0]; |
2389 | x.constant_id = get_decoration(id: c.m.c[0].id[0], decoration: DecorationSpecId); |
2390 | } |
2391 | |
2392 | if (c.m.c[0].id[1] != ID(0)) |
2393 | { |
2394 | y.id = c.m.c[0].id[1]; |
2395 | y.constant_id = get_decoration(id: c.m.c[0].id[1], decoration: DecorationSpecId); |
2396 | } |
2397 | |
2398 | if (c.m.c[0].id[2] != ID(0)) |
2399 | { |
2400 | z.id = c.m.c[0].id[2]; |
2401 | z.constant_id = get_decoration(id: c.m.c[0].id[2], decoration: DecorationSpecId); |
2402 | } |
2403 | } |
2404 | else if (execution.flags.get(bit: ExecutionModeLocalSizeId)) |
2405 | { |
2406 | auto &cx = get<SPIRConstant>(id: execution.workgroup_size.id_x); |
2407 | if (cx.specialization) |
2408 | { |
2409 | x.id = execution.workgroup_size.id_x; |
2410 | x.constant_id = get_decoration(id: execution.workgroup_size.id_x, decoration: DecorationSpecId); |
2411 | } |
2412 | |
2413 | auto &cy = get<SPIRConstant>(id: execution.workgroup_size.id_y); |
2414 | if (cy.specialization) |
2415 | { |
2416 | y.id = execution.workgroup_size.id_y; |
2417 | y.constant_id = get_decoration(id: execution.workgroup_size.id_y, decoration: DecorationSpecId); |
2418 | } |
2419 | |
2420 | auto &cz = get<SPIRConstant>(id: execution.workgroup_size.id_z); |
2421 | if (cz.specialization) |
2422 | { |
2423 | z.id = execution.workgroup_size.id_z; |
2424 | z.constant_id = get_decoration(id: execution.workgroup_size.id_z, decoration: DecorationSpecId); |
2425 | } |
2426 | } |
2427 | |
2428 | return execution.workgroup_size.constant; |
2429 | } |
2430 | |
2431 | uint32_t Compiler::get_execution_mode_argument(spv::ExecutionMode mode, uint32_t index) const |
2432 | { |
2433 | auto &execution = get_entry_point(); |
2434 | switch (mode) |
2435 | { |
2436 | case ExecutionModeLocalSizeId: |
2437 | if (execution.flags.get(bit: ExecutionModeLocalSizeId)) |
2438 | { |
2439 | switch (index) |
2440 | { |
2441 | case 0: |
2442 | return execution.workgroup_size.id_x; |
2443 | case 1: |
2444 | return execution.workgroup_size.id_y; |
2445 | case 2: |
2446 | return execution.workgroup_size.id_z; |
2447 | default: |
2448 | return 0; |
2449 | } |
2450 | } |
2451 | else |
2452 | return 0; |
2453 | |
2454 | case ExecutionModeLocalSize: |
2455 | switch (index) |
2456 | { |
2457 | case 0: |
2458 | if (execution.flags.get(bit: ExecutionModeLocalSizeId) && execution.workgroup_size.id_x != 0) |
2459 | return get<SPIRConstant>(id: execution.workgroup_size.id_x).scalar(); |
2460 | else |
2461 | return execution.workgroup_size.x; |
2462 | case 1: |
2463 | if (execution.flags.get(bit: ExecutionModeLocalSizeId) && execution.workgroup_size.id_y != 0) |
2464 | return get<SPIRConstant>(id: execution.workgroup_size.id_y).scalar(); |
2465 | else |
2466 | return execution.workgroup_size.y; |
2467 | case 2: |
2468 | if (execution.flags.get(bit: ExecutionModeLocalSizeId) && execution.workgroup_size.id_z != 0) |
2469 | return get<SPIRConstant>(id: execution.workgroup_size.id_z).scalar(); |
2470 | else |
2471 | return execution.workgroup_size.z; |
2472 | default: |
2473 | return 0; |
2474 | } |
2475 | |
2476 | case ExecutionModeInvocations: |
2477 | return execution.invocations; |
2478 | |
2479 | case ExecutionModeOutputVertices: |
2480 | return execution.output_vertices; |
2481 | |
2482 | case ExecutionModeOutputPrimitivesEXT: |
2483 | return execution.output_primitives; |
2484 | |
2485 | default: |
2486 | return 0; |
2487 | } |
2488 | } |
2489 | |
2490 | ExecutionModel Compiler::get_execution_model() const |
2491 | { |
2492 | auto &execution = get_entry_point(); |
2493 | return execution.model; |
2494 | } |
2495 | |
2496 | bool Compiler::is_tessellation_shader(ExecutionModel model) |
2497 | { |
2498 | return model == ExecutionModelTessellationControl || model == ExecutionModelTessellationEvaluation; |
2499 | } |
2500 | |
2501 | bool Compiler::is_vertex_like_shader() const |
2502 | { |
2503 | auto model = get_execution_model(); |
2504 | return model == ExecutionModelVertex || model == ExecutionModelGeometry || |
2505 | model == ExecutionModelTessellationControl || model == ExecutionModelTessellationEvaluation; |
2506 | } |
2507 | |
2508 | bool Compiler::is_tessellation_shader() const |
2509 | { |
2510 | return is_tessellation_shader(model: get_execution_model()); |
2511 | } |
2512 | |
2513 | bool Compiler::is_tessellating_triangles() const |
2514 | { |
2515 | return get_execution_mode_bitset().get(bit: ExecutionModeTriangles); |
2516 | } |
2517 | |
2518 | void Compiler::set_remapped_variable_state(VariableID id, bool remap_enable) |
2519 | { |
2520 | get<SPIRVariable>(id).remapped_variable = remap_enable; |
2521 | } |
2522 | |
2523 | bool Compiler::get_remapped_variable_state(VariableID id) const |
2524 | { |
2525 | return get<SPIRVariable>(id).remapped_variable; |
2526 | } |
2527 | |
2528 | void Compiler::set_subpass_input_remapped_components(VariableID id, uint32_t components) |
2529 | { |
2530 | get<SPIRVariable>(id).remapped_components = components; |
2531 | } |
2532 | |
2533 | uint32_t Compiler::get_subpass_input_remapped_components(VariableID id) const |
2534 | { |
2535 | return get<SPIRVariable>(id).remapped_components; |
2536 | } |
2537 | |
2538 | void Compiler::add_implied_read_expression(SPIRExpression &e, uint32_t source) |
2539 | { |
2540 | auto itr = find(first: begin(cont&: e.implied_read_expressions), last: end(cont&: e.implied_read_expressions), val: ID(source)); |
2541 | if (itr == end(cont&: e.implied_read_expressions)) |
2542 | e.implied_read_expressions.push_back(t: source); |
2543 | } |
2544 | |
2545 | void Compiler::add_implied_read_expression(SPIRAccessChain &e, uint32_t source) |
2546 | { |
2547 | auto itr = find(first: begin(cont&: e.implied_read_expressions), last: end(cont&: e.implied_read_expressions), val: ID(source)); |
2548 | if (itr == end(cont&: e.implied_read_expressions)) |
2549 | e.implied_read_expressions.push_back(t: source); |
2550 | } |
2551 | |
2552 | void Compiler::add_active_interface_variable(uint32_t var_id) |
2553 | { |
2554 | active_interface_variables.insert(x: var_id); |
2555 | |
2556 | // In SPIR-V 1.4 and up we must also track the interface variable in the entry point. |
2557 | if (ir.get_spirv_version() >= 0x10400) |
2558 | { |
2559 | auto &vars = get_entry_point().interface_variables; |
2560 | if (find(first: begin(cont&: vars), last: end(cont&: vars), val: VariableID(var_id)) == end(cont&: vars)) |
2561 | vars.push_back(t: var_id); |
2562 | } |
2563 | } |
2564 | |
2565 | void Compiler::inherit_expression_dependencies(uint32_t dst, uint32_t source_expression) |
2566 | { |
2567 | // Don't inherit any expression dependencies if the expression in dst |
2568 | // is not a forwarded temporary. |
2569 | if (forwarded_temporaries.find(x: dst) == end(cont&: forwarded_temporaries) || |
2570 | forced_temporaries.find(x: dst) != end(cont&: forced_temporaries)) |
2571 | { |
2572 | return; |
2573 | } |
2574 | |
2575 | auto &e = get<SPIRExpression>(id: dst); |
2576 | auto *phi = maybe_get<SPIRVariable>(id: source_expression); |
2577 | if (phi && phi->phi_variable) |
2578 | { |
2579 | // We have used a phi variable, which can change at the end of the block, |
2580 | // so make sure we take a dependency on this phi variable. |
2581 | phi->dependees.push_back(t: dst); |
2582 | } |
2583 | |
2584 | auto *s = maybe_get<SPIRExpression>(id: source_expression); |
2585 | if (!s) |
2586 | return; |
2587 | |
2588 | auto &e_deps = e.expression_dependencies; |
2589 | auto &s_deps = s->expression_dependencies; |
2590 | |
2591 | // If we depend on a expression, we also depend on all sub-dependencies from source. |
2592 | e_deps.push_back(t: source_expression); |
2593 | e_deps.insert(itr: end(cont&: e_deps), insert_begin: begin(cont&: s_deps), insert_end: end(cont&: s_deps)); |
2594 | |
2595 | // Eliminate duplicated dependencies. |
2596 | sort(first: begin(cont&: e_deps), last: end(cont&: e_deps)); |
2597 | e_deps.erase(start_erase: unique(first: begin(cont&: e_deps), last: end(cont&: e_deps)), end_erase: end(cont&: e_deps)); |
2598 | } |
2599 | |
2600 | SmallVector<EntryPoint> Compiler::get_entry_points_and_stages() const |
2601 | { |
2602 | SmallVector<EntryPoint> entries; |
2603 | for (auto &entry : ir.entry_points) |
2604 | entries.push_back(t: { .name: entry.second.orig_name, .execution_model: entry.second.model }); |
2605 | return entries; |
2606 | } |
2607 | |
2608 | void Compiler::rename_entry_point(const std::string &old_name, const std::string &new_name, spv::ExecutionModel model) |
2609 | { |
2610 | auto &entry = get_entry_point(name: old_name, execution_model: model); |
2611 | entry.orig_name = new_name; |
2612 | entry.name = new_name; |
2613 | } |
2614 | |
2615 | void Compiler::set_entry_point(const std::string &name, spv::ExecutionModel model) |
2616 | { |
2617 | auto &entry = get_entry_point(name, execution_model: model); |
2618 | ir.default_entry_point = entry.self; |
2619 | } |
2620 | |
2621 | SPIREntryPoint &Compiler::get_first_entry_point(const std::string &name) |
2622 | { |
2623 | auto itr = find_if( |
2624 | first: begin(cont&: ir.entry_points), last: end(cont&: ir.entry_points), |
2625 | pred: [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool { return entry.second.orig_name == name; }); |
2626 | |
2627 | if (itr == end(cont&: ir.entry_points)) |
2628 | SPIRV_CROSS_THROW("Entry point does not exist."); |
2629 | |
2630 | return itr->second; |
2631 | } |
2632 | |
2633 | const SPIREntryPoint &Compiler::get_first_entry_point(const std::string &name) const |
2634 | { |
2635 | auto itr = find_if( |
2636 | first: begin(cont: ir.entry_points), last: end(cont: ir.entry_points), |
2637 | pred: [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool { return entry.second.orig_name == name; }); |
2638 | |
2639 | if (itr == end(cont: ir.entry_points)) |
2640 | SPIRV_CROSS_THROW("Entry point does not exist."); |
2641 | |
2642 | return itr->second; |
2643 | } |
2644 | |
2645 | SPIREntryPoint &Compiler::get_entry_point(const std::string &name, ExecutionModel model) |
2646 | { |
2647 | auto itr = find_if(first: begin(cont&: ir.entry_points), last: end(cont&: ir.entry_points), |
2648 | pred: [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool { |
2649 | return entry.second.orig_name == name && entry.second.model == model; |
2650 | }); |
2651 | |
2652 | if (itr == end(cont&: ir.entry_points)) |
2653 | SPIRV_CROSS_THROW("Entry point does not exist."); |
2654 | |
2655 | return itr->second; |
2656 | } |
2657 | |
2658 | const SPIREntryPoint &Compiler::get_entry_point(const std::string &name, ExecutionModel model) const |
2659 | { |
2660 | auto itr = find_if(first: begin(cont: ir.entry_points), last: end(cont: ir.entry_points), |
2661 | pred: [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool { |
2662 | return entry.second.orig_name == name && entry.second.model == model; |
2663 | }); |
2664 | |
2665 | if (itr == end(cont: ir.entry_points)) |
2666 | SPIRV_CROSS_THROW("Entry point does not exist."); |
2667 | |
2668 | return itr->second; |
2669 | } |
2670 | |
2671 | const string &Compiler::get_cleansed_entry_point_name(const std::string &name, ExecutionModel model) const |
2672 | { |
2673 | return get_entry_point(name, model).name; |
2674 | } |
2675 | |
2676 | const SPIREntryPoint &Compiler::get_entry_point() const |
2677 | { |
2678 | return ir.entry_points.find(x: ir.default_entry_point)->second; |
2679 | } |
2680 | |
2681 | SPIREntryPoint &Compiler::get_entry_point() |
2682 | { |
2683 | return ir.entry_points.find(x: ir.default_entry_point)->second; |
2684 | } |
2685 | |
2686 | bool Compiler::interface_variable_exists_in_entry_point(uint32_t id) const |
2687 | { |
2688 | auto &var = get<SPIRVariable>(id); |
2689 | |
2690 | if (ir.get_spirv_version() < 0x10400) |
2691 | { |
2692 | if (var.storage != StorageClassInput && var.storage != StorageClassOutput && |
2693 | var.storage != StorageClassUniformConstant) |
2694 | SPIRV_CROSS_THROW("Only Input, Output variables and Uniform constants are part of a shader linking interface."); |
2695 | |
2696 | // This is to avoid potential problems with very old glslang versions which did |
2697 | // not emit input/output interfaces properly. |
2698 | // We can assume they only had a single entry point, and single entry point |
2699 | // shaders could easily be assumed to use every interface variable anyways. |
2700 | if (ir.entry_points.size() <= 1) |
2701 | return true; |
2702 | } |
2703 | |
2704 | // In SPIR-V 1.4 and later, all global resource variables must be present. |
2705 | |
2706 | auto &execution = get_entry_point(); |
2707 | return find(first: begin(cont: execution.interface_variables), last: end(cont: execution.interface_variables), val: VariableID(id)) != |
2708 | end(cont: execution.interface_variables); |
2709 | } |
2710 | |
2711 | void Compiler::CombinedImageSamplerHandler::push_remap_parameters(const SPIRFunction &func, const uint32_t *args, |
2712 | uint32_t length) |
2713 | { |
2714 | // If possible, pipe through a remapping table so that parameters know |
2715 | // which variables they actually bind to in this scope. |
2716 | unordered_map<uint32_t, uint32_t> remapping; |
2717 | for (uint32_t i = 0; i < length; i++) |
2718 | remapping[func.arguments[i].id] = remap_parameter(id: args[i]); |
2719 | parameter_remapping.push(x: std::move(remapping)); |
2720 | } |
2721 | |
2722 | void Compiler::CombinedImageSamplerHandler::pop_remap_parameters() |
2723 | { |
2724 | parameter_remapping.pop(); |
2725 | } |
2726 | |
2727 | uint32_t Compiler::CombinedImageSamplerHandler::remap_parameter(uint32_t id) |
2728 | { |
2729 | auto *var = compiler.maybe_get_backing_variable(chain: id); |
2730 | if (var) |
2731 | id = var->self; |
2732 | |
2733 | if (parameter_remapping.empty()) |
2734 | return id; |
2735 | |
2736 | auto &remapping = parameter_remapping.top(); |
2737 | auto itr = remapping.find(x: id); |
2738 | if (itr != end(cont&: remapping)) |
2739 | return itr->second; |
2740 | else |
2741 | return id; |
2742 | } |
2743 | |
2744 | bool Compiler::CombinedImageSamplerHandler::begin_function_scope(const uint32_t *args, uint32_t length) |
2745 | { |
2746 | if (length < 3) |
2747 | return false; |
2748 | |
2749 | auto &callee = compiler.get<SPIRFunction>(id: args[2]); |
2750 | args += 3; |
2751 | length -= 3; |
2752 | push_remap_parameters(func: callee, args, length); |
2753 | functions.push(x: &callee); |
2754 | return true; |
2755 | } |
2756 | |
2757 | bool Compiler::CombinedImageSamplerHandler::end_function_scope(const uint32_t *args, uint32_t length) |
2758 | { |
2759 | if (length < 3) |
2760 | return false; |
2761 | |
2762 | auto &callee = compiler.get<SPIRFunction>(id: args[2]); |
2763 | args += 3; |
2764 | |
2765 | // There are two types of cases we have to handle, |
2766 | // a callee might call sampler2D(texture2D, sampler) directly where |
2767 | // one or more parameters originate from parameters. |
2768 | // Alternatively, we need to provide combined image samplers to our callees, |
2769 | // and in this case we need to add those as well. |
2770 | |
2771 | pop_remap_parameters(); |
2772 | |
2773 | // Our callee has now been processed at least once. |
2774 | // No point in doing it again. |
2775 | callee.do_combined_parameters = false; |
2776 | |
2777 | auto ¶ms = functions.top()->combined_parameters; |
2778 | functions.pop(); |
2779 | if (functions.empty()) |
2780 | return true; |
2781 | |
2782 | auto &caller = *functions.top(); |
2783 | if (caller.do_combined_parameters) |
2784 | { |
2785 | for (auto ¶m : params) |
2786 | { |
2787 | VariableID image_id = param.global_image ? param.image_id : VariableID(args[param.image_id]); |
2788 | VariableID sampler_id = param.global_sampler ? param.sampler_id : VariableID(args[param.sampler_id]); |
2789 | |
2790 | auto *i = compiler.maybe_get_backing_variable(chain: image_id); |
2791 | auto *s = compiler.maybe_get_backing_variable(chain: sampler_id); |
2792 | if (i) |
2793 | image_id = i->self; |
2794 | if (s) |
2795 | sampler_id = s->self; |
2796 | |
2797 | register_combined_image_sampler(caller, combined_id: 0, texture_id: image_id, sampler_id, depth: param.depth); |
2798 | } |
2799 | } |
2800 | |
2801 | return true; |
2802 | } |
2803 | |
2804 | void Compiler::CombinedImageSamplerHandler::register_combined_image_sampler(SPIRFunction &caller, |
2805 | VariableID combined_module_id, |
2806 | VariableID image_id, VariableID sampler_id, |
2807 | bool depth) |
2808 | { |
2809 | // We now have a texture ID and a sampler ID which will either be found as a global |
2810 | // or a parameter in our own function. If both are global, they will not need a parameter, |
2811 | // otherwise, add it to our list. |
2812 | SPIRFunction::CombinedImageSamplerParameter param = { |
2813 | .id: 0u, .image_id: image_id, .sampler_id: sampler_id, .global_image: true, .global_sampler: true, .depth: depth, |
2814 | }; |
2815 | |
2816 | auto texture_itr = find_if(first: begin(cont&: caller.arguments), last: end(cont&: caller.arguments), |
2817 | pred: [image_id](const SPIRFunction::Parameter &p) { return p.id == image_id; }); |
2818 | auto sampler_itr = find_if(first: begin(cont&: caller.arguments), last: end(cont&: caller.arguments), |
2819 | pred: [sampler_id](const SPIRFunction::Parameter &p) { return p.id == sampler_id; }); |
2820 | |
2821 | if (texture_itr != end(cont&: caller.arguments)) |
2822 | { |
2823 | param.global_image = false; |
2824 | param.image_id = uint32_t(texture_itr - begin(cont&: caller.arguments)); |
2825 | } |
2826 | |
2827 | if (sampler_itr != end(cont&: caller.arguments)) |
2828 | { |
2829 | param.global_sampler = false; |
2830 | param.sampler_id = uint32_t(sampler_itr - begin(cont&: caller.arguments)); |
2831 | } |
2832 | |
2833 | if (param.global_image && param.global_sampler) |
2834 | return; |
2835 | |
2836 | auto itr = find_if(first: begin(cont&: caller.combined_parameters), last: end(cont&: caller.combined_parameters), |
2837 | pred: [¶m](const SPIRFunction::CombinedImageSamplerParameter &p) { |
2838 | return param.image_id == p.image_id && param.sampler_id == p.sampler_id && |
2839 | param.global_image == p.global_image && param.global_sampler == p.global_sampler; |
2840 | }); |
2841 | |
2842 | if (itr == end(cont&: caller.combined_parameters)) |
2843 | { |
2844 | uint32_t id = compiler.ir.increase_bound_by(count: 3); |
2845 | auto type_id = id + 0; |
2846 | auto ptr_type_id = id + 1; |
2847 | auto combined_id = id + 2; |
2848 | auto &base = compiler.expression_type(id: image_id); |
2849 | auto &type = compiler.set<SPIRType>(id: type_id, args: OpTypeSampledImage); |
2850 | auto &ptr_type = compiler.set<SPIRType>(id: ptr_type_id, args: OpTypePointer); |
2851 | |
2852 | type = base; |
2853 | type.self = type_id; |
2854 | type.basetype = SPIRType::SampledImage; |
2855 | type.pointer = false; |
2856 | type.storage = StorageClassGeneric; |
2857 | type.image.depth = depth; |
2858 | |
2859 | ptr_type = type; |
2860 | ptr_type.pointer = true; |
2861 | ptr_type.storage = StorageClassUniformConstant; |
2862 | ptr_type.parent_type = type_id; |
2863 | |
2864 | // Build new variable. |
2865 | compiler.set<SPIRVariable>(id: combined_id, args&: ptr_type_id, args: StorageClassFunction, args: 0); |
2866 | |
2867 | // Inherit RelaxedPrecision. |
2868 | // If any of OpSampledImage, underlying image or sampler are marked, inherit the decoration. |
2869 | bool relaxed_precision = |
2870 | compiler.has_decoration(id: sampler_id, decoration: DecorationRelaxedPrecision) || |
2871 | compiler.has_decoration(id: image_id, decoration: DecorationRelaxedPrecision) || |
2872 | (combined_module_id && compiler.has_decoration(id: combined_module_id, decoration: DecorationRelaxedPrecision)); |
2873 | |
2874 | if (relaxed_precision) |
2875 | compiler.set_decoration(id: combined_id, decoration: DecorationRelaxedPrecision); |
2876 | |
2877 | param.id = combined_id; |
2878 | |
2879 | compiler.set_name(id: combined_id, |
2880 | name: join(ts: "SPIRV_Cross_Combined", ts: compiler.to_name(id: image_id), ts: compiler.to_name(id: sampler_id))); |
2881 | |
2882 | caller.combined_parameters.push_back(t: param); |
2883 | caller.shadow_arguments.push_back(t: { .type: ptr_type_id, .id: combined_id, .read_count: 0u, .write_count: 0u, .alias_global_variable: true }); |
2884 | } |
2885 | } |
2886 | |
2887 | bool Compiler::DummySamplerForCombinedImageHandler::handle(Op opcode, const uint32_t *args, uint32_t length) |
2888 | { |
2889 | if (need_dummy_sampler) |
2890 | { |
2891 | // No need to traverse further, we know the result. |
2892 | return false; |
2893 | } |
2894 | |
2895 | switch (opcode) |
2896 | { |
2897 | case OpLoad: |
2898 | { |
2899 | if (length < 3) |
2900 | return false; |
2901 | |
2902 | uint32_t result_type = args[0]; |
2903 | |
2904 | auto &type = compiler.get<SPIRType>(id: result_type); |
2905 | bool separate_image = |
2906 | type.basetype == SPIRType::Image && type.image.sampled == 1 && type.image.dim != DimBuffer; |
2907 | |
2908 | // If not separate image, don't bother. |
2909 | if (!separate_image) |
2910 | return true; |
2911 | |
2912 | uint32_t id = args[1]; |
2913 | uint32_t ptr = args[2]; |
2914 | compiler.set<SPIRExpression>(id, args: "", args&: result_type, args: true); |
2915 | compiler.register_read(expr: id, chain: ptr, forwarded: true); |
2916 | break; |
2917 | } |
2918 | |
2919 | case OpImageFetch: |
2920 | case OpImageQuerySizeLod: |
2921 | case OpImageQuerySize: |
2922 | case OpImageQueryLevels: |
2923 | case OpImageQuerySamples: |
2924 | { |
2925 | // If we are fetching or querying LOD from a plain OpTypeImage, we must pre-combine with our dummy sampler. |
2926 | auto *var = compiler.maybe_get_backing_variable(chain: args[2]); |
2927 | if (var) |
2928 | { |
2929 | auto &type = compiler.get<SPIRType>(id: var->basetype); |
2930 | if (type.basetype == SPIRType::Image && type.image.sampled == 1 && type.image.dim != DimBuffer) |
2931 | need_dummy_sampler = true; |
2932 | } |
2933 | |
2934 | break; |
2935 | } |
2936 | |
2937 | case OpInBoundsAccessChain: |
2938 | case OpAccessChain: |
2939 | case OpPtrAccessChain: |
2940 | { |
2941 | if (length < 3) |
2942 | return false; |
2943 | |
2944 | uint32_t result_type = args[0]; |
2945 | auto &type = compiler.get<SPIRType>(id: result_type); |
2946 | bool separate_image = |
2947 | type.basetype == SPIRType::Image && type.image.sampled == 1 && type.image.dim != DimBuffer; |
2948 | if (!separate_image) |
2949 | return true; |
2950 | |
2951 | uint32_t id = args[1]; |
2952 | uint32_t ptr = args[2]; |
2953 | compiler.set<SPIRExpression>(id, args: "", args&: result_type, args: true); |
2954 | compiler.register_read(expr: id, chain: ptr, forwarded: true); |
2955 | |
2956 | // Other backends might use SPIRAccessChain for this later. |
2957 | compiler.ir.ids[id].set_allow_type_rewrite(); |
2958 | break; |
2959 | } |
2960 | |
2961 | default: |
2962 | break; |
2963 | } |
2964 | |
2965 | return true; |
2966 | } |
2967 | |
2968 | bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *args, uint32_t length) |
2969 | { |
2970 | // We need to figure out where samplers and images are loaded from, so do only the bare bones compilation we need. |
2971 | bool is_fetch = false; |
2972 | |
2973 | switch (opcode) |
2974 | { |
2975 | case OpLoad: |
2976 | { |
2977 | if (length < 3) |
2978 | return false; |
2979 | |
2980 | uint32_t result_type = args[0]; |
2981 | |
2982 | auto &type = compiler.get<SPIRType>(id: result_type); |
2983 | bool separate_image = type.basetype == SPIRType::Image && type.image.sampled == 1; |
2984 | bool separate_sampler = type.basetype == SPIRType::Sampler; |
2985 | |
2986 | // If not separate image or sampler, don't bother. |
2987 | if (!separate_image && !separate_sampler) |
2988 | return true; |
2989 | |
2990 | uint32_t id = args[1]; |
2991 | uint32_t ptr = args[2]; |
2992 | compiler.set<SPIRExpression>(id, args: "", args&: result_type, args: true); |
2993 | compiler.register_read(expr: id, chain: ptr, forwarded: true); |
2994 | return true; |
2995 | } |
2996 | |
2997 | case OpInBoundsAccessChain: |
2998 | case OpAccessChain: |
2999 | case OpPtrAccessChain: |
3000 | { |
3001 | if (length < 3) |
3002 | return false; |
3003 | |
3004 | // Technically, it is possible to have arrays of textures and arrays of samplers and combine them, but this becomes essentially |
3005 | // impossible to implement, since we don't know which concrete sampler we are accessing. |
3006 | // One potential way is to create a combinatorial explosion where N textures and M samplers are combined into N * M sampler2Ds, |
3007 | // but this seems ridiculously complicated for a problem which is easy to work around. |
3008 | // Checking access chains like this assumes we don't have samplers or textures inside uniform structs, but this makes no sense. |
3009 | |
3010 | uint32_t result_type = args[0]; |
3011 | |
3012 | auto &type = compiler.get<SPIRType>(id: result_type); |
3013 | bool separate_image = type.basetype == SPIRType::Image && type.image.sampled == 1; |
3014 | bool separate_sampler = type.basetype == SPIRType::Sampler; |
3015 | if (separate_sampler) |
3016 | SPIRV_CROSS_THROW( |
3017 | "Attempting to use arrays or structs of separate samplers. This is not possible to statically " |
3018 | "remap to plain GLSL."); |
3019 | |
3020 | if (separate_image) |
3021 | { |
3022 | uint32_t id = args[1]; |
3023 | uint32_t ptr = args[2]; |
3024 | compiler.set<SPIRExpression>(id, args: "", args&: result_type, args: true); |
3025 | compiler.register_read(expr: id, chain: ptr, forwarded: true); |
3026 | } |
3027 | return true; |
3028 | } |
3029 | |
3030 | case OpImageFetch: |
3031 | case OpImageQuerySizeLod: |
3032 | case OpImageQuerySize: |
3033 | case OpImageQueryLevels: |
3034 | case OpImageQuerySamples: |
3035 | { |
3036 | // If we are fetching from a plain OpTypeImage or querying LOD, we must pre-combine with our dummy sampler. |
3037 | auto *var = compiler.maybe_get_backing_variable(chain: args[2]); |
3038 | if (!var) |
3039 | return true; |
3040 | |
3041 | auto &type = compiler.get<SPIRType>(id: var->basetype); |
3042 | if (type.basetype == SPIRType::Image && type.image.sampled == 1 && type.image.dim != DimBuffer) |
3043 | { |
3044 | if (compiler.dummy_sampler_id == 0) |
3045 | SPIRV_CROSS_THROW("texelFetch without sampler was found, but no dummy sampler has been created with " |
3046 | "build_dummy_sampler_for_combined_images()."); |
3047 | |
3048 | // Do it outside. |
3049 | is_fetch = true; |
3050 | break; |
3051 | } |
3052 | |
3053 | return true; |
3054 | } |
3055 | |
3056 | case OpSampledImage: |
3057 | // Do it outside. |
3058 | break; |
3059 | |
3060 | default: |
3061 | return true; |
3062 | } |
3063 | |
3064 | // Registers sampler2D calls used in case they are parameters so |
3065 | // that their callees know which combined image samplers to propagate down the call stack. |
3066 | if (!functions.empty()) |
3067 | { |
3068 | auto &callee = *functions.top(); |
3069 | if (callee.do_combined_parameters) |
3070 | { |
3071 | uint32_t image_id = args[2]; |
3072 | |
3073 | auto *image = compiler.maybe_get_backing_variable(chain: image_id); |
3074 | if (image) |
3075 | image_id = image->self; |
3076 | |
3077 | uint32_t sampler_id = is_fetch ? compiler.dummy_sampler_id : args[3]; |
3078 | auto *sampler = compiler.maybe_get_backing_variable(chain: sampler_id); |
3079 | if (sampler) |
3080 | sampler_id = sampler->self; |
3081 | |
3082 | uint32_t combined_id = args[1]; |
3083 | |
3084 | auto &combined_type = compiler.get<SPIRType>(id: args[0]); |
3085 | register_combined_image_sampler(caller&: callee, combined_module_id: combined_id, image_id, sampler_id, depth: combined_type.image.depth); |
3086 | } |
3087 | } |
3088 | |
3089 | // For function calls, we need to remap IDs which are function parameters into global variables. |
3090 | // This information is statically known from the current place in the call stack. |
3091 | // Function parameters are not necessarily pointers, so if we don't have a backing variable, remapping will know |
3092 | // which backing variable the image/sample came from. |
3093 | VariableID image_id = remap_parameter(id: args[2]); |
3094 | VariableID sampler_id = is_fetch ? compiler.dummy_sampler_id : remap_parameter(id: args[3]); |
3095 | |
3096 | auto itr = find_if(first: begin(cont&: compiler.combined_image_samplers), last: end(cont&: compiler.combined_image_samplers), |
3097 | pred: [image_id, sampler_id](const CombinedImageSampler &combined) { |
3098 | return combined.image_id == image_id && combined.sampler_id == sampler_id; |
3099 | }); |
3100 | |
3101 | if (itr == end(cont&: compiler.combined_image_samplers)) |
3102 | { |
3103 | uint32_t sampled_type; |
3104 | uint32_t combined_module_id; |
3105 | if (is_fetch) |
3106 | { |
3107 | // Have to invent the sampled image type. |
3108 | sampled_type = compiler.ir.increase_bound_by(count: 1); |
3109 | auto &type = compiler.set<SPIRType>(id: sampled_type, args: OpTypeSampledImage); |
3110 | type = compiler.expression_type(id: args[2]); |
3111 | type.self = sampled_type; |
3112 | type.basetype = SPIRType::SampledImage; |
3113 | type.image.depth = false; |
3114 | combined_module_id = 0; |
3115 | } |
3116 | else |
3117 | { |
3118 | sampled_type = args[0]; |
3119 | combined_module_id = args[1]; |
3120 | } |
3121 | |
3122 | auto id = compiler.ir.increase_bound_by(count: 2); |
3123 | auto type_id = id + 0; |
3124 | auto combined_id = id + 1; |
3125 | |
3126 | // Make a new type, pointer to OpTypeSampledImage, so we can make a variable of this type. |
3127 | // We will probably have this type lying around, but it doesn't hurt to make duplicates for internal purposes. |
3128 | auto &type = compiler.set<SPIRType>(id: type_id, args: OpTypePointer); |
3129 | auto &base = compiler.get<SPIRType>(id: sampled_type); |
3130 | type = base; |
3131 | type.pointer = true; |
3132 | type.storage = StorageClassUniformConstant; |
3133 | type.parent_type = type_id; |
3134 | |
3135 | // Build new variable. |
3136 | compiler.set<SPIRVariable>(id: combined_id, args&: type_id, args: StorageClassUniformConstant, args: 0); |
3137 | |
3138 | // Inherit RelaxedPrecision (and potentially other useful flags if deemed relevant). |
3139 | // If any of OpSampledImage, underlying image or sampler are marked, inherit the decoration. |
3140 | bool relaxed_precision = |
3141 | (sampler_id && compiler.has_decoration(id: sampler_id, decoration: DecorationRelaxedPrecision)) || |
3142 | (image_id && compiler.has_decoration(id: image_id, decoration: DecorationRelaxedPrecision)) || |
3143 | (combined_module_id && compiler.has_decoration(id: combined_module_id, decoration: DecorationRelaxedPrecision)); |
3144 | |
3145 | if (relaxed_precision) |
3146 | compiler.set_decoration(id: combined_id, decoration: DecorationRelaxedPrecision); |
3147 | |
3148 | // Propagate the array type for the original image as well. |
3149 | auto *var = compiler.maybe_get_backing_variable(chain: image_id); |
3150 | if (var) |
3151 | { |
3152 | auto &parent_type = compiler.get<SPIRType>(id: var->basetype); |
3153 | type.array = parent_type.array; |
3154 | type.array_size_literal = parent_type.array_size_literal; |
3155 | } |
3156 | |
3157 | compiler.combined_image_samplers.push_back(t: { .combined_id: combined_id, .image_id: image_id, .sampler_id: sampler_id }); |
3158 | } |
3159 | |
3160 | return true; |
3161 | } |
3162 | |
3163 | VariableID Compiler::build_dummy_sampler_for_combined_images() |
3164 | { |
3165 | DummySamplerForCombinedImageHandler handler(*this); |
3166 | traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler); |
3167 | if (handler.need_dummy_sampler) |
3168 | { |
3169 | uint32_t offset = ir.increase_bound_by(count: 3); |
3170 | auto type_id = offset + 0; |
3171 | auto ptr_type_id = offset + 1; |
3172 | auto var_id = offset + 2; |
3173 | |
3174 | auto &sampler = set<SPIRType>(id: type_id, args: OpTypeSampler); |
3175 | sampler.basetype = SPIRType::Sampler; |
3176 | |
3177 | auto &ptr_sampler = set<SPIRType>(id: ptr_type_id, args: OpTypePointer); |
3178 | ptr_sampler = sampler; |
3179 | ptr_sampler.self = type_id; |
3180 | ptr_sampler.storage = StorageClassUniformConstant; |
3181 | ptr_sampler.pointer = true; |
3182 | ptr_sampler.parent_type = type_id; |
3183 | |
3184 | set<SPIRVariable>(id: var_id, args&: ptr_type_id, args: StorageClassUniformConstant, args: 0); |
3185 | set_name(id: var_id, name: "SPIRV_Cross_DummySampler"); |
3186 | dummy_sampler_id = var_id; |
3187 | return var_id; |
3188 | } |
3189 | else |
3190 | return 0; |
3191 | } |
3192 | |
3193 | void Compiler::build_combined_image_samplers() |
3194 | { |
3195 | ir.for_each_typed_id<SPIRFunction>(op: [&](uint32_t, SPIRFunction &func) { |
3196 | func.combined_parameters.clear(); |
3197 | func.shadow_arguments.clear(); |
3198 | func.do_combined_parameters = true; |
3199 | }); |
3200 | |
3201 | combined_image_samplers.clear(); |
3202 | CombinedImageSamplerHandler handler(*this); |
3203 | traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler); |
3204 | } |
3205 | |
3206 | SmallVector<SpecializationConstant> Compiler::get_specialization_constants() const |
3207 | { |
3208 | SmallVector<SpecializationConstant> spec_consts; |
3209 | ir.for_each_typed_id<SPIRConstant>(op: [&](uint32_t, const SPIRConstant &c) { |
3210 | if (c.specialization && has_decoration(id: c.self, decoration: DecorationSpecId)) |
3211 | spec_consts.push_back(t: { .id: c.self, .constant_id: get_decoration(id: c.self, decoration: DecorationSpecId) }); |
3212 | }); |
3213 | return spec_consts; |
3214 | } |
3215 | |
3216 | SPIRConstant &Compiler::get_constant(ConstantID id) |
3217 | { |
3218 | return get<SPIRConstant>(id); |
3219 | } |
3220 | |
3221 | const SPIRConstant &Compiler::get_constant(ConstantID id) const |
3222 | { |
3223 | return get<SPIRConstant>(id); |
3224 | } |
3225 | |
3226 | static bool exists_unaccessed_path_to_return(const CFG &cfg, uint32_t block, const unordered_set<uint32_t> &blocks, |
3227 | unordered_set<uint32_t> &visit_cache) |
3228 | { |
3229 | // This block accesses the variable. |
3230 | if (blocks.find(x: block) != end(cont: blocks)) |
3231 | return false; |
3232 | |
3233 | // We are at the end of the CFG. |
3234 | if (cfg.get_succeeding_edges(block).empty()) |
3235 | return true; |
3236 | |
3237 | // If any of our successors have a path to the end, there exists a path from block. |
3238 | for (auto &succ : cfg.get_succeeding_edges(block)) |
3239 | { |
3240 | if (visit_cache.count(x: succ) == 0) |
3241 | { |
3242 | if (exists_unaccessed_path_to_return(cfg, block: succ, blocks, visit_cache)) |
3243 | return true; |
3244 | visit_cache.insert(x: succ); |
3245 | } |
3246 | } |
3247 | |
3248 | return false; |
3249 | } |
3250 | |
3251 | void Compiler::analyze_parameter_preservation( |
3252 | SPIRFunction &entry, const CFG &cfg, const unordered_map<uint32_t, unordered_set<uint32_t>> &variable_to_blocks, |
3253 | const unordered_map<uint32_t, unordered_set<uint32_t>> &complete_write_blocks) |
3254 | { |
3255 | for (auto &arg : entry.arguments) |
3256 | { |
3257 | // Non-pointers are always inputs. |
3258 | auto &type = get<SPIRType>(id: arg.type); |
3259 | if (!type.pointer) |
3260 | continue; |
3261 | |
3262 | // Opaque argument types are always in |
3263 | bool potential_preserve; |
3264 | switch (type.basetype) |
3265 | { |
3266 | case SPIRType::Sampler: |
3267 | case SPIRType::Image: |
3268 | case SPIRType::SampledImage: |
3269 | case SPIRType::AtomicCounter: |
3270 | potential_preserve = false; |
3271 | break; |
3272 | |
3273 | default: |
3274 | potential_preserve = true; |
3275 | break; |
3276 | } |
3277 | |
3278 | if (!potential_preserve) |
3279 | continue; |
3280 | |
3281 | auto itr = variable_to_blocks.find(x: arg.id); |
3282 | if (itr == end(cont: variable_to_blocks)) |
3283 | { |
3284 | // Variable is never accessed. |
3285 | continue; |
3286 | } |
3287 | |
3288 | // We have accessed a variable, but there was no complete writes to that variable. |
3289 | // We deduce that we must preserve the argument. |
3290 | itr = complete_write_blocks.find(x: arg.id); |
3291 | if (itr == end(cont: complete_write_blocks)) |
3292 | { |
3293 | arg.read_count++; |
3294 | continue; |
3295 | } |
3296 | |
3297 | // If there is a path through the CFG where no block completely writes to the variable, the variable will be in an undefined state |
3298 | // when the function returns. We therefore need to implicitly preserve the variable in case there are writers in the function. |
3299 | // Major case here is if a function is |
3300 | // void foo(int &var) { if (cond) var = 10; } |
3301 | // Using read/write counts, we will think it's just an out variable, but it really needs to be inout, |
3302 | // because if we don't write anything whatever we put into the function must return back to the caller. |
3303 | unordered_set<uint32_t> visit_cache; |
3304 | if (exists_unaccessed_path_to_return(cfg, block: entry.entry_block, blocks: itr->second, visit_cache)) |
3305 | arg.read_count++; |
3306 | } |
3307 | } |
3308 | |
3309 | Compiler::AnalyzeVariableScopeAccessHandler::AnalyzeVariableScopeAccessHandler(Compiler &compiler_, |
3310 | SPIRFunction &entry_) |
3311 | : compiler(compiler_) |
3312 | , entry(entry_) |
3313 | { |
3314 | } |
3315 | |
3316 | bool Compiler::AnalyzeVariableScopeAccessHandler::follow_function_call(const SPIRFunction &) |
3317 | { |
3318 | // Only analyze within this function. |
3319 | return false; |
3320 | } |
3321 | |
3322 | void Compiler::AnalyzeVariableScopeAccessHandler::set_current_block(const SPIRBlock &block) |
3323 | { |
3324 | current_block = █ |
3325 | |
3326 | // If we're branching to a block which uses OpPhi, in GLSL |
3327 | // this will be a variable write when we branch, |
3328 | // so we need to track access to these variables as well to |
3329 | // have a complete picture. |
3330 | const auto test_phi = [this, &block](uint32_t to) { |
3331 | auto &next = compiler.get<SPIRBlock>(id: to); |
3332 | for (auto &phi : next.phi_variables) |
3333 | { |
3334 | if (phi.parent == block.self) |
3335 | { |
3336 | accessed_variables_to_block[phi.function_variable].insert(x: block.self); |
3337 | // Phi variables are also accessed in our target branch block. |
3338 | accessed_variables_to_block[phi.function_variable].insert(x: next.self); |
3339 | |
3340 | notify_variable_access(id: phi.local_variable, block: block.self); |
3341 | } |
3342 | } |
3343 | }; |
3344 | |
3345 | switch (block.terminator) |
3346 | { |
3347 | case SPIRBlock::Direct: |
3348 | notify_variable_access(id: block.condition, block: block.self); |
3349 | test_phi(block.next_block); |
3350 | break; |
3351 | |
3352 | case SPIRBlock::Select: |
3353 | notify_variable_access(id: block.condition, block: block.self); |
3354 | test_phi(block.true_block); |
3355 | test_phi(block.false_block); |
3356 | break; |
3357 | |
3358 | case SPIRBlock::MultiSelect: |
3359 | { |
3360 | notify_variable_access(id: block.condition, block: block.self); |
3361 | auto &cases = compiler.get_case_list(block); |
3362 | for (auto &target : cases) |
3363 | test_phi(target.block); |
3364 | if (block.default_block) |
3365 | test_phi(block.default_block); |
3366 | break; |
3367 | } |
3368 | |
3369 | default: |
3370 | break; |
3371 | } |
3372 | } |
3373 | |
3374 | void Compiler::AnalyzeVariableScopeAccessHandler::notify_variable_access(uint32_t id, uint32_t block) |
3375 | { |
3376 | if (id == 0) |
3377 | return; |
3378 | |
3379 | // Access chains used in multiple blocks mean hoisting all the variables used to construct the access chain as not all backends can use pointers. |
3380 | auto itr = rvalue_forward_children.find(x: id); |
3381 | if (itr != end(cont&: rvalue_forward_children)) |
3382 | for (auto child_id : itr->second) |
3383 | notify_variable_access(id: child_id, block); |
3384 | |
3385 | if (id_is_phi_variable(id)) |
3386 | accessed_variables_to_block[id].insert(x: block); |
3387 | else if (id_is_potential_temporary(id)) |
3388 | accessed_temporaries_to_block[id].insert(x: block); |
3389 | } |
3390 | |
3391 | bool Compiler::AnalyzeVariableScopeAccessHandler::id_is_phi_variable(uint32_t id) const |
3392 | { |
3393 | if (id >= compiler.get_current_id_bound()) |
3394 | return false; |
3395 | auto *var = compiler.maybe_get<SPIRVariable>(id); |
3396 | return var && var->phi_variable; |
3397 | } |
3398 | |
3399 | bool Compiler::AnalyzeVariableScopeAccessHandler::id_is_potential_temporary(uint32_t id) const |
3400 | { |
3401 | if (id >= compiler.get_current_id_bound()) |
3402 | return false; |
3403 | |
3404 | // Temporaries are not created before we start emitting code. |
3405 | return compiler.ir.ids[id].empty() || (compiler.ir.ids[id].get_type() == TypeExpression); |
3406 | } |
3407 | |
3408 | bool Compiler::AnalyzeVariableScopeAccessHandler::handle_terminator(const SPIRBlock &block) |
3409 | { |
3410 | switch (block.terminator) |
3411 | { |
3412 | case SPIRBlock::Return: |
3413 | if (block.return_value) |
3414 | notify_variable_access(id: block.return_value, block: block.self); |
3415 | break; |
3416 | |
3417 | case SPIRBlock::Select: |
3418 | case SPIRBlock::MultiSelect: |
3419 | notify_variable_access(id: block.condition, block: block.self); |
3420 | break; |
3421 | |
3422 | default: |
3423 | break; |
3424 | } |
3425 | |
3426 | return true; |
3427 | } |
3428 | |
3429 | bool Compiler::AnalyzeVariableScopeAccessHandler::handle(spv::Op op, const uint32_t *args, uint32_t length) |
3430 | { |
3431 | // Keep track of the types of temporaries, so we can hoist them out as necessary. |
3432 | uint32_t result_type = 0, result_id = 0; |
3433 | if (compiler.instruction_to_result_type(result_type, result_id, op, args, length)) |
3434 | { |
3435 | // For some opcodes, we will need to override the result id. |
3436 | // If we need to hoist the temporary, the temporary type is the input, not the result. |
3437 | if (op == OpConvertUToAccelerationStructureKHR) |
3438 | { |
3439 | auto itr = result_id_to_type.find(x: args[2]); |
3440 | if (itr != result_id_to_type.end()) |
3441 | result_type = itr->second; |
3442 | } |
3443 | |
3444 | result_id_to_type[result_id] = result_type; |
3445 | } |
3446 | |
3447 | switch (op) |
3448 | { |
3449 | case OpStore: |
3450 | { |
3451 | if (length < 2) |
3452 | return false; |
3453 | |
3454 | ID ptr = args[0]; |
3455 | auto *var = compiler.maybe_get_backing_variable(chain: ptr); |
3456 | |
3457 | // If we store through an access chain, we have a partial write. |
3458 | if (var) |
3459 | { |
3460 | accessed_variables_to_block[var->self].insert(x: current_block->self); |
3461 | if (var->self == ptr) |
3462 | complete_write_variables_to_block[var->self].insert(x: current_block->self); |
3463 | else |
3464 | partial_write_variables_to_block[var->self].insert(x: current_block->self); |
3465 | } |
3466 | |
3467 | // args[0] might be an access chain we have to track use of. |
3468 | notify_variable_access(id: args[0], block: current_block->self); |
3469 | // Might try to store a Phi variable here. |
3470 | notify_variable_access(id: args[1], block: current_block->self); |
3471 | break; |
3472 | } |
3473 | |
3474 | case OpAccessChain: |
3475 | case OpInBoundsAccessChain: |
3476 | case OpPtrAccessChain: |
3477 | { |
3478 | if (length < 3) |
3479 | return false; |
3480 | |
3481 | // Access chains used in multiple blocks mean hoisting all the variables used to construct the access chain as not all backends can use pointers. |
3482 | uint32_t ptr = args[2]; |
3483 | auto *var = compiler.maybe_get<SPIRVariable>(id: ptr); |
3484 | if (var) |
3485 | { |
3486 | accessed_variables_to_block[var->self].insert(x: current_block->self); |
3487 | rvalue_forward_children[args[1]].insert(x: var->self); |
3488 | } |
3489 | |
3490 | // args[2] might be another access chain we have to track use of. |
3491 | for (uint32_t i = 2; i < length; i++) |
3492 | { |
3493 | notify_variable_access(id: args[i], block: current_block->self); |
3494 | rvalue_forward_children[args[1]].insert(x: args[i]); |
3495 | } |
3496 | |
3497 | // Also keep track of the access chain pointer itself. |
3498 | // In exceptionally rare cases, we can end up with a case where |
3499 | // the access chain is generated in the loop body, but is consumed in continue block. |
3500 | // This means we need complex loop workarounds, and we must detect this via CFG analysis. |
3501 | notify_variable_access(id: args[1], block: current_block->self); |
3502 | |
3503 | // The result of an access chain is a fixed expression and is not really considered a temporary. |
3504 | auto &e = compiler.set<SPIRExpression>(id: args[1], args: "", args: args[0], args: true); |
3505 | auto *backing_variable = compiler.maybe_get_backing_variable(chain: ptr); |
3506 | e.loaded_from = backing_variable ? VariableID(backing_variable->self) : VariableID(0); |
3507 | |
3508 | // Other backends might use SPIRAccessChain for this later. |
3509 | compiler.ir.ids[args[1]].set_allow_type_rewrite(); |
3510 | access_chain_expressions.insert(x: args[1]); |
3511 | break; |
3512 | } |
3513 | |
3514 | case OpCopyMemory: |
3515 | { |
3516 | if (length < 2) |
3517 | return false; |
3518 | |
3519 | ID lhs = args[0]; |
3520 | ID rhs = args[1]; |
3521 | auto *var = compiler.maybe_get_backing_variable(chain: lhs); |
3522 | |
3523 | // If we store through an access chain, we have a partial write. |
3524 | if (var) |
3525 | { |
3526 | accessed_variables_to_block[var->self].insert(x: current_block->self); |
3527 | if (var->self == lhs) |
3528 | complete_write_variables_to_block[var->self].insert(x: current_block->self); |
3529 | else |
3530 | partial_write_variables_to_block[var->self].insert(x: current_block->self); |
3531 | } |
3532 | |
3533 | // args[0:1] might be access chains we have to track use of. |
3534 | for (uint32_t i = 0; i < 2; i++) |
3535 | notify_variable_access(id: args[i], block: current_block->self); |
3536 | |
3537 | var = compiler.maybe_get_backing_variable(chain: rhs); |
3538 | if (var) |
3539 | accessed_variables_to_block[var->self].insert(x: current_block->self); |
3540 | break; |
3541 | } |
3542 | |
3543 | case OpCopyObject: |
3544 | { |
3545 | // OpCopyObject copies the underlying non-pointer type, |
3546 | // so any temp variable should be declared using the underlying type. |
3547 | // If the type is a pointer, get its base type and overwrite the result type mapping. |
3548 | auto &type = compiler.get<SPIRType>(id: result_type); |
3549 | if (type.pointer) |
3550 | result_id_to_type[result_id] = type.parent_type; |
3551 | |
3552 | if (length < 3) |
3553 | return false; |
3554 | |
3555 | auto *var = compiler.maybe_get_backing_variable(chain: args[2]); |
3556 | if (var) |
3557 | accessed_variables_to_block[var->self].insert(x: current_block->self); |
3558 | |
3559 | // Might be an access chain which we have to keep track of. |
3560 | notify_variable_access(id: args[1], block: current_block->self); |
3561 | if (access_chain_expressions.count(x: args[2])) |
3562 | access_chain_expressions.insert(x: args[1]); |
3563 | |
3564 | // Might try to copy a Phi variable here. |
3565 | notify_variable_access(id: args[2], block: current_block->self); |
3566 | break; |
3567 | } |
3568 | |
3569 | case OpLoad: |
3570 | { |
3571 | if (length < 3) |
3572 | return false; |
3573 | uint32_t ptr = args[2]; |
3574 | auto *var = compiler.maybe_get_backing_variable(chain: ptr); |
3575 | if (var) |
3576 | accessed_variables_to_block[var->self].insert(x: current_block->self); |
3577 | |
3578 | // Loaded value is a temporary. |
3579 | notify_variable_access(id: args[1], block: current_block->self); |
3580 | |
3581 | // Might be an access chain we have to track use of. |
3582 | notify_variable_access(id: args[2], block: current_block->self); |
3583 | |
3584 | // If we're loading an opaque type we cannot lower it to a temporary, |
3585 | // we must defer access of args[2] until it's used. |
3586 | auto &type = compiler.get<SPIRType>(id: args[0]); |
3587 | if (compiler.type_is_opaque_value(type)) |
3588 | rvalue_forward_children[args[1]].insert(x: args[2]); |
3589 | break; |
3590 | } |
3591 | |
3592 | case OpFunctionCall: |
3593 | { |
3594 | if (length < 3) |
3595 | return false; |
3596 | |
3597 | // Return value may be a temporary. |
3598 | if (compiler.get_type(id: args[0]).basetype != SPIRType::Void) |
3599 | notify_variable_access(id: args[1], block: current_block->self); |
3600 | |
3601 | length -= 3; |
3602 | args += 3; |
3603 | |
3604 | for (uint32_t i = 0; i < length; i++) |
3605 | { |
3606 | auto *var = compiler.maybe_get_backing_variable(chain: args[i]); |
3607 | if (var) |
3608 | { |
3609 | accessed_variables_to_block[var->self].insert(x: current_block->self); |
3610 | // Assume we can get partial writes to this variable. |
3611 | partial_write_variables_to_block[var->self].insert(x: current_block->self); |
3612 | } |
3613 | |
3614 | // Cannot easily prove if argument we pass to a function is completely written. |
3615 | // Usually, functions write to a dummy variable, |
3616 | // which is then copied to in full to the real argument. |
3617 | |
3618 | // Might try to copy a Phi variable here. |
3619 | notify_variable_access(id: args[i], block: current_block->self); |
3620 | } |
3621 | break; |
3622 | } |
3623 | |
3624 | case OpSelect: |
3625 | { |
3626 | // In case of variable pointers, we might access a variable here. |
3627 | // We cannot prove anything about these accesses however. |
3628 | for (uint32_t i = 1; i < length; i++) |
3629 | { |
3630 | if (i >= 3) |
3631 | { |
3632 | auto *var = compiler.maybe_get_backing_variable(chain: args[i]); |
3633 | if (var) |
3634 | { |
3635 | accessed_variables_to_block[var->self].insert(x: current_block->self); |
3636 | // Assume we can get partial writes to this variable. |
3637 | partial_write_variables_to_block[var->self].insert(x: current_block->self); |
3638 | } |
3639 | } |
3640 | |
3641 | // Might try to copy a Phi variable here. |
3642 | notify_variable_access(id: args[i], block: current_block->self); |
3643 | } |
3644 | break; |
3645 | } |
3646 | |
3647 | case OpExtInst: |
3648 | { |
3649 | for (uint32_t i = 4; i < length; i++) |
3650 | notify_variable_access(id: args[i], block: current_block->self); |
3651 | notify_variable_access(id: args[1], block: current_block->self); |
3652 | |
3653 | uint32_t extension_set = args[2]; |
3654 | if (compiler.get<SPIRExtension>(id: extension_set).ext == SPIRExtension::GLSL) |
3655 | { |
3656 | auto op_450 = static_cast<GLSLstd450>(args[3]); |
3657 | switch (op_450) |
3658 | { |
3659 | case GLSLstd450Modf: |
3660 | case GLSLstd450Frexp: |
3661 | { |
3662 | uint32_t ptr = args[5]; |
3663 | auto *var = compiler.maybe_get_backing_variable(chain: ptr); |
3664 | if (var) |
3665 | { |
3666 | accessed_variables_to_block[var->self].insert(x: current_block->self); |
3667 | if (var->self == ptr) |
3668 | complete_write_variables_to_block[var->self].insert(x: current_block->self); |
3669 | else |
3670 | partial_write_variables_to_block[var->self].insert(x: current_block->self); |
3671 | } |
3672 | break; |
3673 | } |
3674 | |
3675 | default: |
3676 | break; |
3677 | } |
3678 | } |
3679 | break; |
3680 | } |
3681 | |
3682 | case OpArrayLength: |
3683 | // Only result is a temporary. |
3684 | notify_variable_access(id: args[1], block: current_block->self); |
3685 | break; |
3686 | |
3687 | case OpLine: |
3688 | case OpNoLine: |
3689 | // Uses literals, but cannot be a phi variable or temporary, so ignore. |
3690 | break; |
3691 | |
3692 | // Atomics shouldn't be able to access function-local variables. |
3693 | // Some GLSL builtins access a pointer. |
3694 | |
3695 | case OpCompositeInsert: |
3696 | case OpVectorShuffle: |
3697 | // Specialize for opcode which contains literals. |
3698 | for (uint32_t i = 1; i < 4; i++) |
3699 | notify_variable_access(id: args[i], block: current_block->self); |
3700 | break; |
3701 | |
3702 | case OpCompositeExtract: |
3703 | // Specialize for opcode which contains literals. |
3704 | for (uint32_t i = 1; i < 3; i++) |
3705 | notify_variable_access(id: args[i], block: current_block->self); |
3706 | break; |
3707 | |
3708 | case OpImageWrite: |
3709 | for (uint32_t i = 0; i < length; i++) |
3710 | { |
3711 | // Argument 3 is a literal. |
3712 | if (i != 3) |
3713 | notify_variable_access(id: args[i], block: current_block->self); |
3714 | } |
3715 | break; |
3716 | |
3717 | case OpImageSampleImplicitLod: |
3718 | case OpImageSampleExplicitLod: |
3719 | case OpImageSparseSampleImplicitLod: |
3720 | case OpImageSparseSampleExplicitLod: |
3721 | case OpImageSampleProjImplicitLod: |
3722 | case OpImageSampleProjExplicitLod: |
3723 | case OpImageSparseSampleProjImplicitLod: |
3724 | case OpImageSparseSampleProjExplicitLod: |
3725 | case OpImageFetch: |
3726 | case OpImageSparseFetch: |
3727 | case OpImageRead: |
3728 | case OpImageSparseRead: |
3729 | for (uint32_t i = 1; i < length; i++) |
3730 | { |
3731 | // Argument 4 is a literal. |
3732 | if (i != 4) |
3733 | notify_variable_access(id: args[i], block: current_block->self); |
3734 | } |
3735 | break; |
3736 | |
3737 | case OpImageSampleDrefImplicitLod: |
3738 | case OpImageSampleDrefExplicitLod: |
3739 | case OpImageSparseSampleDrefImplicitLod: |
3740 | case OpImageSparseSampleDrefExplicitLod: |
3741 | case OpImageSampleProjDrefImplicitLod: |
3742 | case OpImageSampleProjDrefExplicitLod: |
3743 | case OpImageSparseSampleProjDrefImplicitLod: |
3744 | case OpImageSparseSampleProjDrefExplicitLod: |
3745 | case OpImageGather: |
3746 | case OpImageSparseGather: |
3747 | case OpImageDrefGather: |
3748 | case OpImageSparseDrefGather: |
3749 | for (uint32_t i = 1; i < length; i++) |
3750 | { |
3751 | // Argument 5 is a literal. |
3752 | if (i != 5) |
3753 | notify_variable_access(id: args[i], block: current_block->self); |
3754 | } |
3755 | break; |
3756 | |
3757 | default: |
3758 | { |
3759 | // Rather dirty way of figuring out where Phi variables are used. |
3760 | // As long as only IDs are used, we can scan through instructions and try to find any evidence that |
3761 | // the ID of a variable has been used. |
3762 | // There are potential false positives here where a literal is used in-place of an ID, |
3763 | // but worst case, it does not affect the correctness of the compile. |
3764 | // Exhaustive analysis would be better here, but it's not worth it for now. |
3765 | for (uint32_t i = 0; i < length; i++) |
3766 | notify_variable_access(id: args[i], block: current_block->self); |
3767 | break; |
3768 | } |
3769 | } |
3770 | return true; |
3771 | } |
3772 | |
3773 | Compiler::StaticExpressionAccessHandler::StaticExpressionAccessHandler(Compiler &compiler_, uint32_t variable_id_) |
3774 | : compiler(compiler_) |
3775 | , variable_id(variable_id_) |
3776 | { |
3777 | } |
3778 | |
3779 | bool Compiler::StaticExpressionAccessHandler::follow_function_call(const SPIRFunction &) |
3780 | { |
3781 | return false; |
3782 | } |
3783 | |
3784 | bool Compiler::StaticExpressionAccessHandler::handle(spv::Op op, const uint32_t *args, uint32_t length) |
3785 | { |
3786 | switch (op) |
3787 | { |
3788 | case OpStore: |
3789 | if (length < 2) |
3790 | return false; |
3791 | if (args[0] == variable_id) |
3792 | { |
3793 | static_expression = args[1]; |
3794 | write_count++; |
3795 | } |
3796 | break; |
3797 | |
3798 | case OpLoad: |
3799 | if (length < 3) |
3800 | return false; |
3801 | if (args[2] == variable_id && static_expression == 0) // Tried to read from variable before it was initialized. |
3802 | return false; |
3803 | break; |
3804 | |
3805 | case OpAccessChain: |
3806 | case OpInBoundsAccessChain: |
3807 | case OpPtrAccessChain: |
3808 | if (length < 3) |
3809 | return false; |
3810 | if (args[2] == variable_id) // If we try to access chain our candidate variable before we store to it, bail. |
3811 | return false; |
3812 | break; |
3813 | |
3814 | default: |
3815 | break; |
3816 | } |
3817 | |
3818 | return true; |
3819 | } |
3820 | |
3821 | void Compiler::find_function_local_luts(SPIRFunction &entry, const AnalyzeVariableScopeAccessHandler &handler, |
3822 | bool single_function) |
3823 | { |
3824 | auto &cfg = *function_cfgs.find(x: entry.self)->second; |
3825 | |
3826 | // For each variable which is statically accessed. |
3827 | for (auto &accessed_var : handler.accessed_variables_to_block) |
3828 | { |
3829 | auto &blocks = accessed_var.second; |
3830 | auto &var = get<SPIRVariable>(id: accessed_var.first); |
3831 | auto &type = expression_type(id: accessed_var.first); |
3832 | |
3833 | // First check if there are writes to the variable. Later, if there are none, we'll |
3834 | // reconsider it as globally accessed LUT. |
3835 | if (!var.is_written_to) |
3836 | { |
3837 | var.is_written_to = handler.complete_write_variables_to_block.count(x: var.self) != 0 || |
3838 | handler.partial_write_variables_to_block.count(x: var.self) != 0; |
3839 | } |
3840 | |
3841 | // Only consider function local variables here. |
3842 | // If we only have a single function in our CFG, private storage is also fine, |
3843 | // since it behaves like a function local variable. |
3844 | bool allow_lut = var.storage == StorageClassFunction || (single_function && var.storage == StorageClassPrivate); |
3845 | if (!allow_lut) |
3846 | continue; |
3847 | |
3848 | // We cannot be a phi variable. |
3849 | if (var.phi_variable) |
3850 | continue; |
3851 | |
3852 | // Only consider arrays here. |
3853 | if (type.array.empty()) |
3854 | continue; |
3855 | |
3856 | // If the variable has an initializer, make sure it is a constant expression. |
3857 | uint32_t static_constant_expression = 0; |
3858 | if (var.initializer) |
3859 | { |
3860 | if (ir.ids[var.initializer].get_type() != TypeConstant) |
3861 | continue; |
3862 | static_constant_expression = var.initializer; |
3863 | |
3864 | // There can be no stores to this variable, we have now proved we have a LUT. |
3865 | if (var.is_written_to) |
3866 | continue; |
3867 | } |
3868 | else |
3869 | { |
3870 | // We can have one, and only one write to the variable, and that write needs to be a constant. |
3871 | |
3872 | // No partial writes allowed. |
3873 | if (handler.partial_write_variables_to_block.count(x: var.self) != 0) |
3874 | continue; |
3875 | |
3876 | auto itr = handler.complete_write_variables_to_block.find(x: var.self); |
3877 | |
3878 | // No writes? |
3879 | if (itr == end(cont: handler.complete_write_variables_to_block)) |
3880 | continue; |
3881 | |
3882 | // We write to the variable in more than one block. |
3883 | auto &write_blocks = itr->second; |
3884 | if (write_blocks.size() != 1) |
3885 | continue; |
3886 | |
3887 | // The write needs to happen in the dominating block. |
3888 | DominatorBuilder builder(cfg); |
3889 | for (auto &block : blocks) |
3890 | builder.add_block(block); |
3891 | uint32_t dominator = builder.get_dominator(); |
3892 | |
3893 | // The complete write happened in a branch or similar, cannot deduce static expression. |
3894 | if (write_blocks.count(x: dominator) == 0) |
3895 | continue; |
3896 | |
3897 | // Find the static expression for this variable. |
3898 | StaticExpressionAccessHandler static_expression_handler(*this, var.self); |
3899 | traverse_all_reachable_opcodes(block: get<SPIRBlock>(id: dominator), handler&: static_expression_handler); |
3900 | |
3901 | // We want one, and exactly one write |
3902 | if (static_expression_handler.write_count != 1 || static_expression_handler.static_expression == 0) |
3903 | continue; |
3904 | |
3905 | // Is it a constant expression? |
3906 | if (ir.ids[static_expression_handler.static_expression].get_type() != TypeConstant) |
3907 | continue; |
3908 | |
3909 | // We found a LUT! |
3910 | static_constant_expression = static_expression_handler.static_expression; |
3911 | } |
3912 | |
3913 | get<SPIRConstant>(id: static_constant_expression).is_used_as_lut = true; |
3914 | var.static_expression = static_constant_expression; |
3915 | var.statically_assigned = true; |
3916 | var.remapped_variable = true; |
3917 | } |
3918 | } |
3919 | |
3920 | void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeAccessHandler &handler) |
3921 | { |
3922 | // First, we map out all variable access within a function. |
3923 | // Essentially a map of block -> { variables accessed in the basic block } |
3924 | traverse_all_reachable_opcodes(func: entry, handler); |
3925 | |
3926 | auto &cfg = *function_cfgs.find(x: entry.self)->second; |
3927 | |
3928 | // Analyze if there are parameters which need to be implicitly preserved with an "in" qualifier. |
3929 | analyze_parameter_preservation(entry, cfg, variable_to_blocks: handler.accessed_variables_to_block, |
3930 | complete_write_blocks: handler.complete_write_variables_to_block); |
3931 | |
3932 | unordered_map<uint32_t, uint32_t> potential_loop_variables; |
3933 | |
3934 | // Find the loop dominator block for each block. |
3935 | for (auto &block_id : entry.blocks) |
3936 | { |
3937 | auto &block = get<SPIRBlock>(id: block_id); |
3938 | |
3939 | auto itr = ir.continue_block_to_loop_header.find(x: block_id); |
3940 | if (itr != end(cont&: ir.continue_block_to_loop_header) && itr->second != block_id) |
3941 | { |
3942 | // Continue block might be unreachable in the CFG, but we still like to know the loop dominator. |
3943 | // Edge case is when continue block is also the loop header, don't set the dominator in this case. |
3944 | block.loop_dominator = itr->second; |
3945 | } |
3946 | else |
3947 | { |
3948 | uint32_t loop_dominator = cfg.find_loop_dominator(block: block_id); |
3949 | if (loop_dominator != block_id) |
3950 | block.loop_dominator = loop_dominator; |
3951 | else |
3952 | block.loop_dominator = SPIRBlock::NoDominator; |
3953 | } |
3954 | } |
3955 | |
3956 | // For each variable which is statically accessed. |
3957 | for (auto &var : handler.accessed_variables_to_block) |
3958 | { |
3959 | // Only deal with variables which are considered local variables in this function. |
3960 | if (find(first: begin(cont&: entry.local_variables), last: end(cont&: entry.local_variables), val: VariableID(var.first)) == |
3961 | end(cont&: entry.local_variables)) |
3962 | continue; |
3963 | |
3964 | DominatorBuilder builder(cfg); |
3965 | auto &blocks = var.second; |
3966 | auto &type = expression_type(id: var.first); |
3967 | BlockID potential_continue_block = 0; |
3968 | |
3969 | // Figure out which block is dominating all accesses of those variables. |
3970 | for (auto &block : blocks) |
3971 | { |
3972 | // If we're accessing a variable inside a continue block, this variable might be a loop variable. |
3973 | // We can only use loop variables with scalars, as we cannot track static expressions for vectors. |
3974 | if (is_continue(next: block)) |
3975 | { |
3976 | // Potentially awkward case to check for. |
3977 | // We might have a variable inside a loop, which is touched by the continue block, |
3978 | // but is not actually a loop variable. |
3979 | // The continue block is dominated by the inner part of the loop, which does not make sense in high-level |
3980 | // language output because it will be declared before the body, |
3981 | // so we will have to lift the dominator up to the relevant loop header instead. |
3982 | builder.add_block(block: ir.continue_block_to_loop_header[block]); |
3983 | |
3984 | // Arrays or structs cannot be loop variables. |
3985 | if (type.vecsize == 1 && type.columns == 1 && type.basetype != SPIRType::Struct && type.array.empty()) |
3986 | { |
3987 | // The variable is used in multiple continue blocks, this is not a loop |
3988 | // candidate, signal that by setting block to -1u. |
3989 | if (potential_continue_block == 0) |
3990 | potential_continue_block = block; |
3991 | else |
3992 | potential_continue_block = ~(0u); |
3993 | } |
3994 | } |
3995 | |
3996 | builder.add_block(block); |
3997 | } |
3998 | |
3999 | builder.lift_continue_block_dominator(); |
4000 | |
4001 | // Add it to a per-block list of variables. |
4002 | BlockID dominating_block = builder.get_dominator(); |
4003 | |
4004 | if (dominating_block && potential_continue_block != 0 && potential_continue_block != ~0u) |
4005 | { |
4006 | auto &inner_block = get<SPIRBlock>(id: dominating_block); |
4007 | |
4008 | BlockID merge_candidate = 0; |
4009 | |
4010 | // Analyze the dominator. If it lives in a different loop scope than the candidate continue |
4011 | // block, reject the loop variable candidate. |
4012 | if (inner_block.merge == SPIRBlock::MergeLoop) |
4013 | merge_candidate = inner_block.merge_block; |
4014 | else if (inner_block.loop_dominator != SPIRBlock::NoDominator) |
4015 | merge_candidate = get<SPIRBlock>(id: inner_block.loop_dominator).merge_block; |
4016 | |
4017 | if (merge_candidate != 0 && cfg.is_reachable(block: merge_candidate)) |
4018 | { |
4019 | // If the merge block has a higher post-visit order, we know that continue candidate |
4020 | // cannot reach the merge block, and we have two separate scopes. |
4021 | if (!cfg.is_reachable(block: potential_continue_block) || |
4022 | cfg.get_visit_order(block: merge_candidate) > cfg.get_visit_order(block: potential_continue_block)) |
4023 | { |
4024 | potential_continue_block = 0; |
4025 | } |
4026 | } |
4027 | } |
4028 | |
4029 | if (potential_continue_block != 0 && potential_continue_block != ~0u) |
4030 | potential_loop_variables[var.first] = potential_continue_block; |
4031 | |
4032 | // For variables whose dominating block is inside a loop, there is a risk that these variables |
4033 | // actually need to be preserved across loop iterations. We can express this by adding |
4034 | // a "read" access to the loop header. |
4035 | // In the dominating block, we must see an OpStore or equivalent as the first access of an OpVariable. |
4036 | // Should that fail, we look for the outermost loop header and tack on an access there. |
4037 | // Phi nodes cannot have this problem. |
4038 | if (dominating_block) |
4039 | { |
4040 | auto &variable = get<SPIRVariable>(id: var.first); |
4041 | if (!variable.phi_variable) |
4042 | { |
4043 | auto *block = &get<SPIRBlock>(id: dominating_block); |
4044 | bool preserve = may_read_undefined_variable_in_block(block: *block, var: var.first); |
4045 | if (preserve) |
4046 | { |
4047 | // Find the outermost loop scope. |
4048 | while (block->loop_dominator != BlockID(SPIRBlock::NoDominator)) |
4049 | block = &get<SPIRBlock>(id: block->loop_dominator); |
4050 | |
4051 | if (block->self != dominating_block) |
4052 | { |
4053 | builder.add_block(block: block->self); |
4054 | dominating_block = builder.get_dominator(); |
4055 | } |
4056 | } |
4057 | } |
4058 | } |
4059 | |
4060 | // If all blocks here are dead code, this will be 0, so the variable in question |
4061 | // will be completely eliminated. |
4062 | if (dominating_block) |
4063 | { |
4064 | auto &block = get<SPIRBlock>(id: dominating_block); |
4065 | block.dominated_variables.push_back(t: var.first); |
4066 | get<SPIRVariable>(id: var.first).dominator = dominating_block; |
4067 | } |
4068 | } |
4069 | |
4070 | for (auto &var : handler.accessed_temporaries_to_block) |
4071 | { |
4072 | auto itr = handler.result_id_to_type.find(x: var.first); |
4073 | |
4074 | if (itr == end(cont&: handler.result_id_to_type)) |
4075 | { |
4076 | // We found a false positive ID being used, ignore. |
4077 | // This should probably be an assert. |
4078 | continue; |
4079 | } |
4080 | |
4081 | // There is no point in doing domination analysis for opaque types. |
4082 | auto &type = get<SPIRType>(id: itr->second); |
4083 | if (type_is_opaque_value(type)) |
4084 | continue; |
4085 | |
4086 | DominatorBuilder builder(cfg); |
4087 | bool force_temporary = false; |
4088 | bool used_in_header_hoisted_continue_block = false; |
4089 | |
4090 | // Figure out which block is dominating all accesses of those temporaries. |
4091 | auto &blocks = var.second; |
4092 | for (auto &block : blocks) |
4093 | { |
4094 | builder.add_block(block); |
4095 | |
4096 | if (blocks.size() != 1 && is_continue(next: block)) |
4097 | { |
4098 | // The risk here is that inner loop can dominate the continue block. |
4099 | // Any temporary we access in the continue block must be declared before the loop. |
4100 | // This is moot for complex loops however. |
4101 | auto &loop_header_block = get<SPIRBlock>(id: ir.continue_block_to_loop_header[block]); |
4102 | assert(loop_header_block.merge == SPIRBlock::MergeLoop); |
4103 | builder.add_block(block: loop_header_block.self); |
4104 | used_in_header_hoisted_continue_block = true; |
4105 | } |
4106 | } |
4107 | |
4108 | uint32_t dominating_block = builder.get_dominator(); |
4109 | |
4110 | if (blocks.size() != 1 && is_single_block_loop(next: dominating_block)) |
4111 | { |
4112 | // Awkward case, because the loop header is also the continue block, |
4113 | // so hoisting to loop header does not help. |
4114 | force_temporary = true; |
4115 | } |
4116 | |
4117 | if (dominating_block) |
4118 | { |
4119 | // If we touch a variable in the dominating block, this is the expected setup. |
4120 | // SPIR-V normally mandates this, but we have extra cases for temporary use inside loops. |
4121 | bool first_use_is_dominator = blocks.count(x: dominating_block) != 0; |
4122 | |
4123 | if (!first_use_is_dominator || force_temporary) |
4124 | { |
4125 | if (handler.access_chain_expressions.count(x: var.first)) |
4126 | { |
4127 | // Exceptionally rare case. |
4128 | // We cannot declare temporaries of access chains (except on MSL perhaps with pointers). |
4129 | // Rather than do that, we force the indexing expressions to be declared in the right scope by |
4130 | // tracking their usage to that end. There is no temporary to hoist. |
4131 | // However, we still need to observe declaration order of the access chain. |
4132 | |
4133 | if (used_in_header_hoisted_continue_block) |
4134 | { |
4135 | // For this scenario, we used an access chain inside a continue block where we also registered an access to header block. |
4136 | // This is a problem as we need to declare an access chain properly first with full definition. |
4137 | // We cannot use temporaries for these expressions, |
4138 | // so we must make sure the access chain is declared ahead of time. |
4139 | // Force a complex for loop to deal with this. |
4140 | // TODO: Out-of-order declaring for loops where continue blocks are emitted last might be another option. |
4141 | auto &loop_header_block = get<SPIRBlock>(id: dominating_block); |
4142 | assert(loop_header_block.merge == SPIRBlock::MergeLoop); |
4143 | loop_header_block.complex_continue = true; |
4144 | } |
4145 | } |
4146 | else |
4147 | { |
4148 | // This should be very rare, but if we try to declare a temporary inside a loop, |
4149 | // and that temporary is used outside the loop as well (spirv-opt inliner likes this) |
4150 | // we should actually emit the temporary outside the loop. |
4151 | hoisted_temporaries.insert(x: var.first); |
4152 | forced_temporaries.insert(x: var.first); |
4153 | |
4154 | auto &block_temporaries = get<SPIRBlock>(id: dominating_block).declare_temporary; |
4155 | block_temporaries.emplace_back(ts&: handler.result_id_to_type[var.first], ts: var.first); |
4156 | } |
4157 | } |
4158 | else if (blocks.size() > 1) |
4159 | { |
4160 | // Keep track of the temporary as we might have to declare this temporary. |
4161 | // This can happen if the loop header dominates a temporary, but we have a complex fallback loop. |
4162 | // In this case, the header is actually inside the for (;;) {} block, and we have problems. |
4163 | // What we need to do is hoist the temporaries outside the for (;;) {} block in case the header block |
4164 | // declares the temporary. |
4165 | auto &block_temporaries = get<SPIRBlock>(id: dominating_block).potential_declare_temporary; |
4166 | block_temporaries.emplace_back(ts&: handler.result_id_to_type[var.first], ts: var.first); |
4167 | } |
4168 | } |
4169 | } |
4170 | |
4171 | unordered_set<uint32_t> seen_blocks; |
4172 | |
4173 | // Now, try to analyze whether or not these variables are actually loop variables. |
4174 | for (auto &loop_variable : potential_loop_variables) |
4175 | { |
4176 | auto &var = get<SPIRVariable>(id: loop_variable.first); |
4177 | auto dominator = var.dominator; |
4178 | BlockID block = loop_variable.second; |
4179 | |
4180 | // The variable was accessed in multiple continue blocks, ignore. |
4181 | if (block == BlockID(~(0u)) || block == BlockID(0)) |
4182 | continue; |
4183 | |
4184 | // Dead code. |
4185 | if (dominator == ID(0)) |
4186 | continue; |
4187 | |
4188 | BlockID header = 0; |
4189 | |
4190 | // Find the loop header for this block if we are a continue block. |
4191 | { |
4192 | auto itr = ir.continue_block_to_loop_header.find(x: block); |
4193 | if (itr != end(cont&: ir.continue_block_to_loop_header)) |
4194 | { |
4195 | header = itr->second; |
4196 | } |
4197 | else if (get<SPIRBlock>(id: block).continue_block == block) |
4198 | { |
4199 | // Also check for self-referential continue block. |
4200 | header = block; |
4201 | } |
4202 | } |
4203 | |
4204 | assert(header); |
4205 | auto &header_block = get<SPIRBlock>(id: header); |
4206 | auto &blocks = handler.accessed_variables_to_block[loop_variable.first]; |
4207 | |
4208 | // If a loop variable is not used before the loop, it's probably not a loop variable. |
4209 | bool has_accessed_variable = blocks.count(x: header) != 0; |
4210 | |
4211 | // Now, there are two conditions we need to meet for the variable to be a loop variable. |
4212 | // 1. The dominating block must have a branch-free path to the loop header, |
4213 | // this way we statically know which expression should be part of the loop variable initializer. |
4214 | |
4215 | // Walk from the dominator, if there is one straight edge connecting |
4216 | // dominator and loop header, we statically know the loop initializer. |
4217 | bool static_loop_init = true; |
4218 | while (dominator != header) |
4219 | { |
4220 | if (blocks.count(x: dominator) != 0) |
4221 | has_accessed_variable = true; |
4222 | |
4223 | auto &succ = cfg.get_succeeding_edges(block: dominator); |
4224 | if (succ.size() != 1) |
4225 | { |
4226 | static_loop_init = false; |
4227 | break; |
4228 | } |
4229 | |
4230 | auto &pred = cfg.get_preceding_edges(block: succ.front()); |
4231 | if (pred.size() != 1 || pred.front() != dominator) |
4232 | { |
4233 | static_loop_init = false; |
4234 | break; |
4235 | } |
4236 | |
4237 | dominator = succ.front(); |
4238 | } |
4239 | |
4240 | if (!static_loop_init || !has_accessed_variable) |
4241 | continue; |
4242 | |
4243 | // The second condition we need to meet is that no access after the loop |
4244 | // merge can occur. Walk the CFG to see if we find anything. |
4245 | |
4246 | seen_blocks.clear(); |
4247 | cfg.walk_from(seen_blocks, block: header_block.merge_block, op: [&](uint32_t walk_block) -> bool { |
4248 | // We found a block which accesses the variable outside the loop. |
4249 | if (blocks.find(x: walk_block) != end(cont&: blocks)) |
4250 | static_loop_init = false; |
4251 | return true; |
4252 | }); |
4253 | |
4254 | if (!static_loop_init) |
4255 | continue; |
4256 | |
4257 | // We have a loop variable. |
4258 | header_block.loop_variables.push_back(t: loop_variable.first); |
4259 | // Need to sort here as variables come from an unordered container, and pushing stuff in wrong order |
4260 | // will break reproducability in regression runs. |
4261 | sort(first: begin(cont&: header_block.loop_variables), last: end(cont&: header_block.loop_variables)); |
4262 | get<SPIRVariable>(id: loop_variable.first).loop_variable = true; |
4263 | } |
4264 | } |
4265 | |
4266 | bool Compiler::may_read_undefined_variable_in_block(const SPIRBlock &block, uint32_t var) |
4267 | { |
4268 | for (auto &op : block.ops) |
4269 | { |
4270 | auto *ops = stream(instr: op); |
4271 | switch (op.op) |
4272 | { |
4273 | case OpStore: |
4274 | case OpCopyMemory: |
4275 | if (ops[0] == var) |
4276 | return false; |
4277 | break; |
4278 | |
4279 | case OpAccessChain: |
4280 | case OpInBoundsAccessChain: |
4281 | case OpPtrAccessChain: |
4282 | // Access chains are generally used to partially read and write. It's too hard to analyze |
4283 | // if all constituents are written fully before continuing, so just assume it's preserved. |
4284 | // This is the same as the parameter preservation analysis. |
4285 | if (ops[2] == var) |
4286 | return true; |
4287 | break; |
4288 | |
4289 | case OpSelect: |
4290 | // Variable pointers. |
4291 | // We might read before writing. |
4292 | if (ops[3] == var || ops[4] == var) |
4293 | return true; |
4294 | break; |
4295 | |
4296 | case OpPhi: |
4297 | { |
4298 | // Variable pointers. |
4299 | // We might read before writing. |
4300 | if (op.length < 2) |
4301 | break; |
4302 | |
4303 | uint32_t count = op.length - 2; |
4304 | for (uint32_t i = 0; i < count; i += 2) |
4305 | if (ops[i + 2] == var) |
4306 | return true; |
4307 | break; |
4308 | } |
4309 | |
4310 | case OpCopyObject: |
4311 | case OpLoad: |
4312 | if (ops[2] == var) |
4313 | return true; |
4314 | break; |
4315 | |
4316 | case OpFunctionCall: |
4317 | { |
4318 | if (op.length < 3) |
4319 | break; |
4320 | |
4321 | // May read before writing. |
4322 | uint32_t count = op.length - 3; |
4323 | for (uint32_t i = 0; i < count; i++) |
4324 | if (ops[i + 3] == var) |
4325 | return true; |
4326 | break; |
4327 | } |
4328 | |
4329 | default: |
4330 | break; |
4331 | } |
4332 | } |
4333 | |
4334 | // Not accessed somehow, at least not in a usual fashion. |
4335 | // It's likely accessed in a branch, so assume we must preserve. |
4336 | return true; |
4337 | } |
4338 | |
4339 | Bitset Compiler::get_buffer_block_flags(VariableID id) const |
4340 | { |
4341 | return ir.get_buffer_block_flags(var: get<SPIRVariable>(id)); |
4342 | } |
4343 | |
4344 | bool Compiler::get_common_basic_type(const SPIRType &type, SPIRType::BaseType &base_type) |
4345 | { |
4346 | if (type.basetype == SPIRType::Struct) |
4347 | { |
4348 | base_type = SPIRType::Unknown; |
4349 | for (auto &member_type : type.member_types) |
4350 | { |
4351 | SPIRType::BaseType member_base; |
4352 | if (!get_common_basic_type(type: get<SPIRType>(id: member_type), base_type&: member_base)) |
4353 | return false; |
4354 | |
4355 | if (base_type == SPIRType::Unknown) |
4356 | base_type = member_base; |
4357 | else if (base_type != member_base) |
4358 | return false; |
4359 | } |
4360 | return true; |
4361 | } |
4362 | else |
4363 | { |
4364 | base_type = type.basetype; |
4365 | return true; |
4366 | } |
4367 | } |
4368 | |
4369 | void Compiler::ActiveBuiltinHandler::handle_builtin(const SPIRType &type, BuiltIn builtin, |
4370 | const Bitset &decoration_flags) |
4371 | { |
4372 | // If used, we will need to explicitly declare a new array size for these builtins. |
4373 | |
4374 | if (builtin == BuiltInClipDistance) |
4375 | { |
4376 | if (!type.array_size_literal[0]) |
4377 | SPIRV_CROSS_THROW("Array size for ClipDistance must be a literal."); |
4378 | uint32_t array_size = type.array[0]; |
4379 | if (array_size == 0) |
4380 | SPIRV_CROSS_THROW("Array size for ClipDistance must not be unsized."); |
4381 | compiler.clip_distance_count = array_size; |
4382 | } |
4383 | else if (builtin == BuiltInCullDistance) |
4384 | { |
4385 | if (!type.array_size_literal[0]) |
4386 | SPIRV_CROSS_THROW("Array size for CullDistance must be a literal."); |
4387 | uint32_t array_size = type.array[0]; |
4388 | if (array_size == 0) |
4389 | SPIRV_CROSS_THROW("Array size for CullDistance must not be unsized."); |
4390 | compiler.cull_distance_count = array_size; |
4391 | } |
4392 | else if (builtin == BuiltInPosition) |
4393 | { |
4394 | if (decoration_flags.get(bit: DecorationInvariant)) |
4395 | compiler.position_invariant = true; |
4396 | } |
4397 | } |
4398 | |
4399 | void Compiler::ActiveBuiltinHandler::add_if_builtin(uint32_t id, bool allow_blocks) |
4400 | { |
4401 | // Only handle plain variables here. |
4402 | // Builtins which are part of a block are handled in AccessChain. |
4403 | // If allow_blocks is used however, this is to handle initializers of blocks, |
4404 | // which implies that all members are written to. |
4405 | |
4406 | auto *var = compiler.maybe_get<SPIRVariable>(id); |
4407 | auto *m = compiler.ir.find_meta(id); |
4408 | if (var && m) |
4409 | { |
4410 | auto &type = compiler.get<SPIRType>(id: var->basetype); |
4411 | auto &decorations = m->decoration; |
4412 | auto &flags = type.storage == StorageClassInput ? |
4413 | compiler.active_input_builtins : compiler.active_output_builtins; |
4414 | if (decorations.builtin) |
4415 | { |
4416 | flags.set(decorations.builtin_type); |
4417 | handle_builtin(type, builtin: decorations.builtin_type, decoration_flags: decorations.decoration_flags); |
4418 | } |
4419 | else if (allow_blocks && compiler.has_decoration(id: type.self, decoration: DecorationBlock)) |
4420 | { |
4421 | uint32_t member_count = uint32_t(type.member_types.size()); |
4422 | for (uint32_t i = 0; i < member_count; i++) |
4423 | { |
4424 | if (compiler.has_member_decoration(id: type.self, index: i, decoration: DecorationBuiltIn)) |
4425 | { |
4426 | auto &member_type = compiler.get<SPIRType>(id: type.member_types[i]); |
4427 | BuiltIn builtin = BuiltIn(compiler.get_member_decoration(id: type.self, index: i, decoration: DecorationBuiltIn)); |
4428 | flags.set(builtin); |
4429 | handle_builtin(type: member_type, builtin, decoration_flags: compiler.get_member_decoration_bitset(id: type.self, index: i)); |
4430 | } |
4431 | } |
4432 | } |
4433 | } |
4434 | } |
4435 | |
4436 | void Compiler::ActiveBuiltinHandler::add_if_builtin(uint32_t id) |
4437 | { |
4438 | add_if_builtin(id, allow_blocks: false); |
4439 | } |
4440 | |
4441 | void Compiler::ActiveBuiltinHandler::add_if_builtin_or_block(uint32_t id) |
4442 | { |
4443 | add_if_builtin(id, allow_blocks: true); |
4444 | } |
4445 | |
4446 | bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args, uint32_t length) |
4447 | { |
4448 | switch (opcode) |
4449 | { |
4450 | case OpStore: |
4451 | if (length < 1) |
4452 | return false; |
4453 | |
4454 | add_if_builtin(id: args[0]); |
4455 | break; |
4456 | |
4457 | case OpCopyMemory: |
4458 | if (length < 2) |
4459 | return false; |
4460 | |
4461 | add_if_builtin(id: args[0]); |
4462 | add_if_builtin(id: args[1]); |
4463 | break; |
4464 | |
4465 | case OpCopyObject: |
4466 | case OpLoad: |
4467 | if (length < 3) |
4468 | return false; |
4469 | |
4470 | add_if_builtin(id: args[2]); |
4471 | break; |
4472 | |
4473 | case OpSelect: |
4474 | if (length < 5) |
4475 | return false; |
4476 | |
4477 | add_if_builtin(id: args[3]); |
4478 | add_if_builtin(id: args[4]); |
4479 | break; |
4480 | |
4481 | case OpPhi: |
4482 | { |
4483 | if (length < 2) |
4484 | return false; |
4485 | |
4486 | uint32_t count = length - 2; |
4487 | args += 2; |
4488 | for (uint32_t i = 0; i < count; i += 2) |
4489 | add_if_builtin(id: args[i]); |
4490 | break; |
4491 | } |
4492 | |
4493 | case OpFunctionCall: |
4494 | { |
4495 | if (length < 3) |
4496 | return false; |
4497 | |
4498 | uint32_t count = length - 3; |
4499 | args += 3; |
4500 | for (uint32_t i = 0; i < count; i++) |
4501 | add_if_builtin(id: args[i]); |
4502 | break; |
4503 | } |
4504 | |
4505 | case OpAccessChain: |
4506 | case OpInBoundsAccessChain: |
4507 | case OpPtrAccessChain: |
4508 | { |
4509 | if (length < 4) |
4510 | return false; |
4511 | |
4512 | // Only consider global variables, cannot consider variables in functions yet, or other |
4513 | // access chains as they have not been created yet. |
4514 | auto *var = compiler.maybe_get<SPIRVariable>(id: args[2]); |
4515 | if (!var) |
4516 | break; |
4517 | |
4518 | // Required if we access chain into builtins like gl_GlobalInvocationID. |
4519 | add_if_builtin(id: args[2]); |
4520 | |
4521 | // Start traversing type hierarchy at the proper non-pointer types. |
4522 | auto *type = &compiler.get_variable_data_type(var: *var); |
4523 | |
4524 | auto &flags = |
4525 | var->storage == StorageClassInput ? compiler.active_input_builtins : compiler.active_output_builtins; |
4526 | |
4527 | uint32_t count = length - 3; |
4528 | args += 3; |
4529 | for (uint32_t i = 0; i < count; i++) |
4530 | { |
4531 | // Pointers |
4532 | // PtrAccessChain functions more like a pointer offset. Type remains the same. |
4533 | if (opcode == OpPtrAccessChain && i == 0) |
4534 | continue; |
4535 | |
4536 | // Arrays |
4537 | if (!type->array.empty()) |
4538 | { |
4539 | type = &compiler.get<SPIRType>(id: type->parent_type); |
4540 | } |
4541 | // Structs |
4542 | else if (type->basetype == SPIRType::Struct) |
4543 | { |
4544 | uint32_t index = compiler.get<SPIRConstant>(id: args[i]).scalar(); |
4545 | |
4546 | if (index < uint32_t(compiler.ir.meta[type->self].members.size())) |
4547 | { |
4548 | auto &decorations = compiler.ir.meta[type->self].members[index]; |
4549 | if (decorations.builtin) |
4550 | { |
4551 | flags.set(decorations.builtin_type); |
4552 | handle_builtin(type: compiler.get<SPIRType>(id: type->member_types[index]), builtin: decorations.builtin_type, |
4553 | decoration_flags: decorations.decoration_flags); |
4554 | } |
4555 | } |
4556 | |
4557 | type = &compiler.get<SPIRType>(id: type->member_types[index]); |
4558 | } |
4559 | else |
4560 | { |
4561 | // No point in traversing further. We won't find any extra builtins. |
4562 | break; |
4563 | } |
4564 | } |
4565 | break; |
4566 | } |
4567 | |
4568 | default: |
4569 | break; |
4570 | } |
4571 | |
4572 | return true; |
4573 | } |
4574 | |
4575 | void Compiler::update_active_builtins() |
4576 | { |
4577 | active_input_builtins.reset(); |
4578 | active_output_builtins.reset(); |
4579 | cull_distance_count = 0; |
4580 | clip_distance_count = 0; |
4581 | ActiveBuiltinHandler handler(*this); |
4582 | traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler); |
4583 | |
4584 | ir.for_each_typed_id<SPIRVariable>(op: [&](uint32_t, const SPIRVariable &var) { |
4585 | if (var.storage != StorageClassOutput) |
4586 | return; |
4587 | if (!interface_variable_exists_in_entry_point(id: var.self)) |
4588 | return; |
4589 | |
4590 | // Also, make sure we preserve output variables which are only initialized, but never accessed by any code. |
4591 | if (var.initializer != ID(0)) |
4592 | handler.add_if_builtin_or_block(id: var.self); |
4593 | }); |
4594 | } |
4595 | |
4596 | // Returns whether this shader uses a builtin of the storage class |
4597 | bool Compiler::has_active_builtin(BuiltIn builtin, StorageClass storage) const |
4598 | { |
4599 | const Bitset *flags; |
4600 | switch (storage) |
4601 | { |
4602 | case StorageClassInput: |
4603 | flags = &active_input_builtins; |
4604 | break; |
4605 | case StorageClassOutput: |
4606 | flags = &active_output_builtins; |
4607 | break; |
4608 | |
4609 | default: |
4610 | return false; |
4611 | } |
4612 | return flags->get(bit: builtin); |
4613 | } |
4614 | |
4615 | void Compiler::analyze_image_and_sampler_usage() |
4616 | { |
4617 | CombinedImageSamplerDrefHandler dref_handler(*this); |
4618 | traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler&: dref_handler); |
4619 | |
4620 | CombinedImageSamplerUsageHandler handler(*this, dref_handler.dref_combined_samplers); |
4621 | traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler); |
4622 | |
4623 | // Need to run this traversal twice. First time, we propagate any comparison sampler usage from leaf functions |
4624 | // down to main(). |
4625 | // In the second pass, we can propagate up forced depth state coming from main() up into leaf functions. |
4626 | handler.dependency_hierarchy.clear(); |
4627 | traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler); |
4628 | |
4629 | comparison_ids = std::move(handler.comparison_ids); |
4630 | need_subpass_input = handler.need_subpass_input; |
4631 | need_subpass_input_ms = handler.need_subpass_input_ms; |
4632 | |
4633 | // Forward information from separate images and samplers into combined image samplers. |
4634 | for (auto &combined : combined_image_samplers) |
4635 | if (comparison_ids.count(x: combined.sampler_id)) |
4636 | comparison_ids.insert(x: combined.combined_id); |
4637 | } |
4638 | |
4639 | bool Compiler::CombinedImageSamplerDrefHandler::handle(spv::Op opcode, const uint32_t *args, uint32_t) |
4640 | { |
4641 | // Mark all sampled images which are used with Dref. |
4642 | switch (opcode) |
4643 | { |
4644 | case OpImageSampleDrefExplicitLod: |
4645 | case OpImageSampleDrefImplicitLod: |
4646 | case OpImageSampleProjDrefExplicitLod: |
4647 | case OpImageSampleProjDrefImplicitLod: |
4648 | case OpImageSparseSampleProjDrefImplicitLod: |
4649 | case OpImageSparseSampleDrefImplicitLod: |
4650 | case OpImageSparseSampleProjDrefExplicitLod: |
4651 | case OpImageSparseSampleDrefExplicitLod: |
4652 | case OpImageDrefGather: |
4653 | case OpImageSparseDrefGather: |
4654 | dref_combined_samplers.insert(x: args[2]); |
4655 | return true; |
4656 | |
4657 | default: |
4658 | break; |
4659 | } |
4660 | |
4661 | return true; |
4662 | } |
4663 | |
4664 | const CFG &Compiler::get_cfg_for_current_function() const |
4665 | { |
4666 | assert(current_function); |
4667 | return get_cfg_for_function(id: current_function->self); |
4668 | } |
4669 | |
4670 | const CFG &Compiler::get_cfg_for_function(uint32_t id) const |
4671 | { |
4672 | auto cfg_itr = function_cfgs.find(x: id); |
4673 | assert(cfg_itr != end(function_cfgs)); |
4674 | assert(cfg_itr->second); |
4675 | return *cfg_itr->second; |
4676 | } |
4677 | |
4678 | void Compiler::build_function_control_flow_graphs_and_analyze() |
4679 | { |
4680 | CFGBuilder handler(*this); |
4681 | handler.function_cfgs[ir.default_entry_point].reset(p: new CFG(*this, get<SPIRFunction>(id: ir.default_entry_point))); |
4682 | traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler); |
4683 | function_cfgs = std::move(handler.function_cfgs); |
4684 | bool single_function = function_cfgs.size() <= 1; |
4685 | |
4686 | for (auto &f : function_cfgs) |
4687 | { |
4688 | auto &func = get<SPIRFunction>(id: f.first); |
4689 | AnalyzeVariableScopeAccessHandler scope_handler(*this, func); |
4690 | analyze_variable_scope(entry&: func, handler&: scope_handler); |
4691 | find_function_local_luts(entry&: func, handler: scope_handler, single_function); |
4692 | |
4693 | // Check if we can actually use the loop variables we found in analyze_variable_scope. |
4694 | // To use multiple initializers, we need the same type and qualifiers. |
4695 | for (auto block : func.blocks) |
4696 | { |
4697 | auto &b = get<SPIRBlock>(id: block); |
4698 | if (b.loop_variables.size() < 2) |
4699 | continue; |
4700 | |
4701 | auto &flags = get_decoration_bitset(id: b.loop_variables.front()); |
4702 | uint32_t type = get<SPIRVariable>(id: b.loop_variables.front()).basetype; |
4703 | bool invalid_initializers = false; |
4704 | for (auto loop_variable : b.loop_variables) |
4705 | { |
4706 | if (flags != get_decoration_bitset(id: loop_variable) || |
4707 | type != get<SPIRVariable>(id: b.loop_variables.front()).basetype) |
4708 | { |
4709 | invalid_initializers = true; |
4710 | break; |
4711 | } |
4712 | } |
4713 | |
4714 | if (invalid_initializers) |
4715 | { |
4716 | for (auto loop_variable : b.loop_variables) |
4717 | get<SPIRVariable>(id: loop_variable).loop_variable = false; |
4718 | b.loop_variables.clear(); |
4719 | } |
4720 | } |
4721 | } |
4722 | |
4723 | // Find LUTs which are not function local. Only consider this case if the CFG is multi-function, |
4724 | // otherwise we treat Private as Function trivially. |
4725 | // Needs to be analyzed from the outside since we have to block the LUT optimization if at least |
4726 | // one function writes to it. |
4727 | if (!single_function) |
4728 | { |
4729 | for (auto &id : global_variables) |
4730 | { |
4731 | auto &var = get<SPIRVariable>(id); |
4732 | auto &type = get_variable_data_type(var); |
4733 | |
4734 | if (is_array(type) && var.storage == StorageClassPrivate && |
4735 | var.initializer && !var.is_written_to && |
4736 | ir.ids[var.initializer].get_type() == TypeConstant) |
4737 | { |
4738 | get<SPIRConstant>(id: var.initializer).is_used_as_lut = true; |
4739 | var.static_expression = var.initializer; |
4740 | var.statically_assigned = true; |
4741 | var.remapped_variable = true; |
4742 | } |
4743 | } |
4744 | } |
4745 | } |
4746 | |
4747 | Compiler::CFGBuilder::CFGBuilder(Compiler &compiler_) |
4748 | : compiler(compiler_) |
4749 | { |
4750 | } |
4751 | |
4752 | bool Compiler::CFGBuilder::handle(spv::Op, const uint32_t *, uint32_t) |
4753 | { |
4754 | return true; |
4755 | } |
4756 | |
4757 | bool Compiler::CFGBuilder::follow_function_call(const SPIRFunction &func) |
4758 | { |
4759 | if (function_cfgs.find(x: func.self) == end(cont&: function_cfgs)) |
4760 | { |
4761 | function_cfgs[func.self].reset(p: new CFG(compiler, func)); |
4762 | return true; |
4763 | } |
4764 | else |
4765 | return false; |
4766 | } |
4767 | |
4768 | void Compiler::CombinedImageSamplerUsageHandler::add_dependency(uint32_t dst, uint32_t src) |
4769 | { |
4770 | dependency_hierarchy[dst].insert(x: src); |
4771 | // Propagate up any comparison state if we're loading from one such variable. |
4772 | if (comparison_ids.count(x: src)) |
4773 | comparison_ids.insert(x: dst); |
4774 | } |
4775 | |
4776 | bool Compiler::CombinedImageSamplerUsageHandler::begin_function_scope(const uint32_t *args, uint32_t length) |
4777 | { |
4778 | if (length < 3) |
4779 | return false; |
4780 | |
4781 | auto &func = compiler.get<SPIRFunction>(id: args[2]); |
4782 | const auto *arg = &args[3]; |
4783 | length -= 3; |
4784 | |
4785 | for (uint32_t i = 0; i < length; i++) |
4786 | { |
4787 | auto &argument = func.arguments[i]; |
4788 | add_dependency(dst: argument.id, src: arg[i]); |
4789 | } |
4790 | |
4791 | return true; |
4792 | } |
4793 | |
4794 | void Compiler::CombinedImageSamplerUsageHandler::add_hierarchy_to_comparison_ids(uint32_t id) |
4795 | { |
4796 | // Traverse the variable dependency hierarchy and tag everything in its path with comparison ids. |
4797 | comparison_ids.insert(x: id); |
4798 | |
4799 | for (auto &dep_id : dependency_hierarchy[id]) |
4800 | add_hierarchy_to_comparison_ids(id: dep_id); |
4801 | } |
4802 | |
4803 | bool Compiler::CombinedImageSamplerUsageHandler::handle(Op opcode, const uint32_t *args, uint32_t length) |
4804 | { |
4805 | switch (opcode) |
4806 | { |
4807 | case OpAccessChain: |
4808 | case OpInBoundsAccessChain: |
4809 | case OpPtrAccessChain: |
4810 | case OpLoad: |
4811 | { |
4812 | if (length < 3) |
4813 | return false; |
4814 | |
4815 | add_dependency(dst: args[1], src: args[2]); |
4816 | |
4817 | // Ideally defer this to OpImageRead, but then we'd need to track loaded IDs. |
4818 | // If we load an image, we're going to use it and there is little harm in declaring an unused gl_FragCoord. |
4819 | auto &type = compiler.get<SPIRType>(id: args[0]); |
4820 | if (type.image.dim == DimSubpassData) |
4821 | { |
4822 | need_subpass_input = true; |
4823 | if (type.image.ms) |
4824 | need_subpass_input_ms = true; |
4825 | } |
4826 | |
4827 | // If we load a SampledImage and it will be used with Dref, propagate the state up. |
4828 | if (dref_combined_samplers.count(x: args[1]) != 0) |
4829 | add_hierarchy_to_comparison_ids(id: args[1]); |
4830 | break; |
4831 | } |
4832 | |
4833 | case OpSampledImage: |
4834 | { |
4835 | if (length < 4) |
4836 | return false; |
4837 | |
4838 | // If the underlying resource has been used for comparison then duplicate loads of that resource must be too. |
4839 | // This image must be a depth image. |
4840 | uint32_t result_id = args[1]; |
4841 | uint32_t image = args[2]; |
4842 | uint32_t sampler = args[3]; |
4843 | |
4844 | if (dref_combined_samplers.count(x: result_id) != 0) |
4845 | { |
4846 | add_hierarchy_to_comparison_ids(id: image); |
4847 | |
4848 | // This sampler must be a SamplerComparisonState, and not a regular SamplerState. |
4849 | add_hierarchy_to_comparison_ids(id: sampler); |
4850 | |
4851 | // Mark the OpSampledImage itself as being comparison state. |
4852 | comparison_ids.insert(x: result_id); |
4853 | } |
4854 | return true; |
4855 | } |
4856 | |
4857 | default: |
4858 | break; |
4859 | } |
4860 | |
4861 | return true; |
4862 | } |
4863 | |
4864 | bool Compiler::buffer_is_hlsl_counter_buffer(VariableID id) const |
4865 | { |
4866 | auto *m = ir.find_meta(id); |
4867 | return m && m->hlsl_is_magic_counter_buffer; |
4868 | } |
4869 | |
4870 | bool Compiler::buffer_get_hlsl_counter_buffer(VariableID id, uint32_t &counter_id) const |
4871 | { |
4872 | auto *m = ir.find_meta(id); |
4873 | |
4874 | // First, check for the proper decoration. |
4875 | if (m && m->hlsl_magic_counter_buffer != 0) |
4876 | { |
4877 | counter_id = m->hlsl_magic_counter_buffer; |
4878 | return true; |
4879 | } |
4880 | else |
4881 | return false; |
4882 | } |
4883 | |
4884 | void Compiler::make_constant_null(uint32_t id, uint32_t type) |
4885 | { |
4886 | auto &constant_type = get<SPIRType>(id: type); |
4887 | |
4888 | if (constant_type.pointer) |
4889 | { |
4890 | auto &constant = set<SPIRConstant>(id, args&: type); |
4891 | constant.make_null(constant_type_: constant_type); |
4892 | } |
4893 | else if (!constant_type.array.empty()) |
4894 | { |
4895 | assert(constant_type.parent_type); |
4896 | uint32_t parent_id = ir.increase_bound_by(count: 1); |
4897 | make_constant_null(id: parent_id, type: constant_type.parent_type); |
4898 | |
4899 | if (!constant_type.array_size_literal.back()) |
4900 | SPIRV_CROSS_THROW("Array size of OpConstantNull must be a literal."); |
4901 | |
4902 | SmallVector<uint32_t> elements(constant_type.array.back()); |
4903 | for (uint32_t i = 0; i < constant_type.array.back(); i++) |
4904 | elements[i] = parent_id; |
4905 | set<SPIRConstant>(id, args&: type, args: elements.data(), args: uint32_t(elements.size()), args: false); |
4906 | } |
4907 | else if (!constant_type.member_types.empty()) |
4908 | { |
4909 | uint32_t member_ids = ir.increase_bound_by(count: uint32_t(constant_type.member_types.size())); |
4910 | SmallVector<uint32_t> elements(constant_type.member_types.size()); |
4911 | for (uint32_t i = 0; i < constant_type.member_types.size(); i++) |
4912 | { |
4913 | make_constant_null(id: member_ids + i, type: constant_type.member_types[i]); |
4914 | elements[i] = member_ids + i; |
4915 | } |
4916 | set<SPIRConstant>(id, args&: type, args: elements.data(), args: uint32_t(elements.size()), args: false); |
4917 | } |
4918 | else |
4919 | { |
4920 | auto &constant = set<SPIRConstant>(id, args&: type); |
4921 | constant.make_null(constant_type_: constant_type); |
4922 | } |
4923 | } |
4924 | |
4925 | const SmallVector<spv::Capability> &Compiler::get_declared_capabilities() const |
4926 | { |
4927 | return ir.declared_capabilities; |
4928 | } |
4929 | |
4930 | const SmallVector<std::string> &Compiler::get_declared_extensions() const |
4931 | { |
4932 | return ir.declared_extensions; |
4933 | } |
4934 | |
4935 | std::string Compiler::get_remapped_declared_block_name(VariableID id) const |
4936 | { |
4937 | return get_remapped_declared_block_name(id, fallback_prefer_instance_name: false); |
4938 | } |
4939 | |
4940 | std::string Compiler::get_remapped_declared_block_name(uint32_t id, bool fallback_prefer_instance_name) const |
4941 | { |
4942 | auto itr = declared_block_names.find(x: id); |
4943 | if (itr != end(cont: declared_block_names)) |
4944 | { |
4945 | return itr->second; |
4946 | } |
4947 | else |
4948 | { |
4949 | auto &var = get<SPIRVariable>(id); |
4950 | |
4951 | if (fallback_prefer_instance_name) |
4952 | { |
4953 | return to_name(id: var.self); |
4954 | } |
4955 | else |
4956 | { |
4957 | auto &type = get<SPIRType>(id: var.basetype); |
4958 | auto *type_meta = ir.find_meta(id: type.self); |
4959 | auto *block_name = type_meta ? &type_meta->decoration.alias : nullptr; |
4960 | return (!block_name || block_name->empty()) ? get_block_fallback_name(id) : *block_name; |
4961 | } |
4962 | } |
4963 | } |
4964 | |
4965 | bool Compiler::reflection_ssbo_instance_name_is_significant() const |
4966 | { |
4967 | if (ir.source.known) |
4968 | { |
4969 | // UAVs from HLSL source tend to be declared in a way where the type is reused |
4970 | // but the instance name is significant, and that's the name we should report. |
4971 | // For GLSL, SSBOs each have their own block type as that's how GLSL is written. |
4972 | return ir.source.hlsl; |
4973 | } |
4974 | |
4975 | unordered_set<uint32_t> ssbo_type_ids; |
4976 | bool aliased_ssbo_types = false; |
4977 | |
4978 | // If we don't have any OpSource information, we need to perform some shaky heuristics. |
4979 | ir.for_each_typed_id<SPIRVariable>(op: [&](uint32_t, const SPIRVariable &var) { |
4980 | auto &type = this->get<SPIRType>(id: var.basetype); |
4981 | if (!type.pointer || var.storage == StorageClassFunction) |
4982 | return; |
4983 | |
4984 | bool ssbo = var.storage == StorageClassStorageBuffer || |
4985 | (var.storage == StorageClassUniform && has_decoration(id: type.self, decoration: DecorationBufferBlock)); |
4986 | |
4987 | if (ssbo) |
4988 | { |
4989 | if (ssbo_type_ids.count(x: type.self)) |
4990 | aliased_ssbo_types = true; |
4991 | else |
4992 | ssbo_type_ids.insert(x: type.self); |
4993 | } |
4994 | }); |
4995 | |
4996 | // If the block name is aliased, assume we have HLSL-style UAV declarations. |
4997 | return aliased_ssbo_types; |
4998 | } |
4999 | |
5000 | bool Compiler::instruction_to_result_type(uint32_t &result_type, uint32_t &result_id, spv::Op op, |
5001 | const uint32_t *args, uint32_t length) |
5002 | { |
5003 | if (length < 2) |
5004 | return false; |
5005 | |
5006 | bool has_result_id = false, has_result_type = false; |
5007 | HasResultAndType(opcode: op, hasResult: &has_result_id, hasResultType: &has_result_type); |
5008 | if (has_result_id && has_result_type) |
5009 | { |
5010 | result_type = args[0]; |
5011 | result_id = args[1]; |
5012 | return true; |
5013 | } |
5014 | else |
5015 | return false; |
5016 | } |
5017 | |
5018 | Bitset Compiler::combined_decoration_for_member(const SPIRType &type, uint32_t index) const |
5019 | { |
5020 | Bitset flags; |
5021 | auto *type_meta = ir.find_meta(id: type.self); |
5022 | |
5023 | if (type_meta) |
5024 | { |
5025 | auto &members = type_meta->members; |
5026 | if (index >= members.size()) |
5027 | return flags; |
5028 | auto &dec = members[index]; |
5029 | |
5030 | flags.merge_or(other: dec.decoration_flags); |
5031 | |
5032 | auto &member_type = get<SPIRType>(id: type.member_types[index]); |
5033 | |
5034 | // If our member type is a struct, traverse all the child members as well recursively. |
5035 | auto &member_childs = member_type.member_types; |
5036 | for (uint32_t i = 0; i < member_childs.size(); i++) |
5037 | { |
5038 | auto &child_member_type = get<SPIRType>(id: member_childs[i]); |
5039 | if (!child_member_type.pointer) |
5040 | flags.merge_or(other: combined_decoration_for_member(type: member_type, index: i)); |
5041 | } |
5042 | } |
5043 | |
5044 | return flags; |
5045 | } |
5046 | |
5047 | bool Compiler::is_desktop_only_format(spv::ImageFormat format) |
5048 | { |
5049 | switch (format) |
5050 | { |
5051 | // Desktop-only formats |
5052 | case ImageFormatR11fG11fB10f: |
5053 | case ImageFormatR16f: |
5054 | case ImageFormatRgb10A2: |
5055 | case ImageFormatR8: |
5056 | case ImageFormatRg8: |
5057 | case ImageFormatR16: |
5058 | case ImageFormatRg16: |
5059 | case ImageFormatRgba16: |
5060 | case ImageFormatR16Snorm: |
5061 | case ImageFormatRg16Snorm: |
5062 | case ImageFormatRgba16Snorm: |
5063 | case ImageFormatR8Snorm: |
5064 | case ImageFormatRg8Snorm: |
5065 | case ImageFormatR8ui: |
5066 | case ImageFormatRg8ui: |
5067 | case ImageFormatR16ui: |
5068 | case ImageFormatRgb10a2ui: |
5069 | case ImageFormatR8i: |
5070 | case ImageFormatRg8i: |
5071 | case ImageFormatR16i: |
5072 | return true; |
5073 | default: |
5074 | break; |
5075 | } |
5076 | |
5077 | return false; |
5078 | } |
5079 | |
5080 | // An image is determined to be a depth image if it is marked as a depth image and is not also |
5081 | // explicitly marked with a color format, or if there are any sample/gather compare operations on it. |
5082 | bool Compiler::is_depth_image(const SPIRType &type, uint32_t id) const |
5083 | { |
5084 | return (type.image.depth && type.image.format == ImageFormatUnknown) || comparison_ids.count(x: id); |
5085 | } |
5086 | |
5087 | bool Compiler::type_is_opaque_value(const SPIRType &type) const |
5088 | { |
5089 | return !type.pointer && (type.basetype == SPIRType::SampledImage || type.basetype == SPIRType::Image || |
5090 | type.basetype == SPIRType::Sampler); |
5091 | } |
5092 | |
5093 | // Make these member functions so we can easily break on any force_recompile events. |
5094 | void Compiler::force_recompile() |
5095 | { |
5096 | is_force_recompile = true; |
5097 | } |
5098 | |
5099 | void Compiler::force_recompile_guarantee_forward_progress() |
5100 | { |
5101 | force_recompile(); |
5102 | is_force_recompile_forward_progress = true; |
5103 | } |
5104 | |
5105 | bool Compiler::is_forcing_recompilation() const |
5106 | { |
5107 | return is_force_recompile; |
5108 | } |
5109 | |
5110 | void Compiler::clear_force_recompile() |
5111 | { |
5112 | is_force_recompile = false; |
5113 | is_force_recompile_forward_progress = false; |
5114 | } |
5115 | |
5116 | Compiler::PhysicalStorageBufferPointerHandler::PhysicalStorageBufferPointerHandler(Compiler &compiler_) |
5117 | : compiler(compiler_) |
5118 | { |
5119 | } |
5120 | |
5121 | Compiler::PhysicalBlockMeta *Compiler::PhysicalStorageBufferPointerHandler::find_block_meta(uint32_t id) const |
5122 | { |
5123 | auto chain_itr = access_chain_to_physical_block.find(x: id); |
5124 | if (chain_itr != access_chain_to_physical_block.end()) |
5125 | return chain_itr->second; |
5126 | else |
5127 | return nullptr; |
5128 | } |
5129 | |
5130 | void Compiler::PhysicalStorageBufferPointerHandler::mark_aligned_access(uint32_t id, const uint32_t *args, uint32_t length) |
5131 | { |
5132 | uint32_t mask = *args; |
5133 | args++; |
5134 | length--; |
5135 | if (length && (mask & MemoryAccessVolatileMask) != 0) |
5136 | { |
5137 | args++; |
5138 | length--; |
5139 | } |
5140 | |
5141 | if (length && (mask & MemoryAccessAlignedMask) != 0) |
5142 | { |
5143 | uint32_t alignment = *args; |
5144 | auto *meta = find_block_meta(id); |
5145 | |
5146 | // This makes the assumption that the application does not rely on insane edge cases like: |
5147 | // Bind buffer with ADDR = 8, use block offset of 8 bytes, load/store with 16 byte alignment. |
5148 | // If we emit the buffer with alignment = 16 here, the first element at offset = 0 should |
5149 | // actually have alignment of 8 bytes, but this is too theoretical and awkward to support. |
5150 | // We could potentially keep track of any offset in the access chain, but it's |
5151 | // practically impossible for high level compilers to emit code like that, |
5152 | // so deducing overall alignment requirement based on maximum observed Alignment value is probably fine. |
5153 | if (meta && alignment > meta->alignment) |
5154 | meta->alignment = alignment; |
5155 | } |
5156 | } |
5157 | |
5158 | bool Compiler::PhysicalStorageBufferPointerHandler::type_is_bda_block_entry(uint32_t type_id) const |
5159 | { |
5160 | auto &type = compiler.get<SPIRType>(id: type_id); |
5161 | return compiler.is_physical_pointer(type); |
5162 | } |
5163 | |
5164 | uint32_t Compiler::PhysicalStorageBufferPointerHandler::get_minimum_scalar_alignment(const SPIRType &type) const |
5165 | { |
5166 | if (type.storage == spv::StorageClassPhysicalStorageBufferEXT) |
5167 | return 8; |
5168 | else if (type.basetype == SPIRType::Struct) |
5169 | { |
5170 | uint32_t alignment = 0; |
5171 | for (auto &member_type : type.member_types) |
5172 | { |
5173 | uint32_t member_align = get_minimum_scalar_alignment(type: compiler.get<SPIRType>(id: member_type)); |
5174 | if (member_align > alignment) |
5175 | alignment = member_align; |
5176 | } |
5177 | return alignment; |
5178 | } |
5179 | else |
5180 | return type.width / 8; |
5181 | } |
5182 | |
5183 | void Compiler::PhysicalStorageBufferPointerHandler::setup_meta_chain(uint32_t type_id, uint32_t var_id) |
5184 | { |
5185 | if (type_is_bda_block_entry(type_id)) |
5186 | { |
5187 | auto &meta = physical_block_type_meta[type_id]; |
5188 | access_chain_to_physical_block[var_id] = &meta; |
5189 | |
5190 | auto &type = compiler.get<SPIRType>(id: type_id); |
5191 | |
5192 | if (!compiler.is_physical_pointer_to_buffer_block(type)) |
5193 | non_block_types.insert(x: type_id); |
5194 | |
5195 | if (meta.alignment == 0) |
5196 | meta.alignment = get_minimum_scalar_alignment(type: compiler.get_pointee_type(type)); |
5197 | } |
5198 | } |
5199 | |
5200 | bool Compiler::PhysicalStorageBufferPointerHandler::handle(Op op, const uint32_t *args, uint32_t length) |
5201 | { |
5202 | // When a BDA pointer comes to life, we need to keep a mapping of SSA ID -> type ID for the pointer type. |
5203 | // For every load and store, we'll need to be able to look up the type ID being accessed and mark any alignment |
5204 | // requirements. |
5205 | switch (op) |
5206 | { |
5207 | case OpConvertUToPtr: |
5208 | case OpBitcast: |
5209 | case OpCompositeExtract: |
5210 | // Extract can begin a new chain if we had a struct or array of pointers as input. |
5211 | // We don't begin chains before we have a pure scalar pointer. |
5212 | setup_meta_chain(type_id: args[0], var_id: args[1]); |
5213 | break; |
5214 | |
5215 | case OpAccessChain: |
5216 | case OpInBoundsAccessChain: |
5217 | case OpPtrAccessChain: |
5218 | case OpCopyObject: |
5219 | { |
5220 | auto itr = access_chain_to_physical_block.find(x: args[2]); |
5221 | if (itr != access_chain_to_physical_block.end()) |
5222 | access_chain_to_physical_block[args[1]] = itr->second; |
5223 | break; |
5224 | } |
5225 | |
5226 | case OpLoad: |
5227 | { |
5228 | setup_meta_chain(type_id: args[0], var_id: args[1]); |
5229 | if (length >= 4) |
5230 | mark_aligned_access(id: args[2], args: args + 3, length: length - 3); |
5231 | break; |
5232 | } |
5233 | |
5234 | case OpStore: |
5235 | { |
5236 | if (length >= 3) |
5237 | mark_aligned_access(id: args[0], args: args + 2, length: length - 2); |
5238 | break; |
5239 | } |
5240 | |
5241 | default: |
5242 | break; |
5243 | } |
5244 | |
5245 | return true; |
5246 | } |
5247 | |
5248 | uint32_t Compiler::PhysicalStorageBufferPointerHandler::get_base_non_block_type_id(uint32_t type_id) const |
5249 | { |
5250 | auto *type = &compiler.get<SPIRType>(id: type_id); |
5251 | while (compiler.is_physical_pointer(type: *type) && !type_is_bda_block_entry(type_id)) |
5252 | { |
5253 | type_id = type->parent_type; |
5254 | type = &compiler.get<SPIRType>(id: type_id); |
5255 | } |
5256 | |
5257 | assert(type_is_bda_block_entry(type_id)); |
5258 | return type_id; |
5259 | } |
5260 | |
5261 | void Compiler::PhysicalStorageBufferPointerHandler::analyze_non_block_types_from_block(const SPIRType &type) |
5262 | { |
5263 | for (auto &member : type.member_types) |
5264 | { |
5265 | auto &subtype = compiler.get<SPIRType>(id: member); |
5266 | |
5267 | if (compiler.is_physical_pointer(type: subtype) && !compiler.is_physical_pointer_to_buffer_block(type: subtype)) |
5268 | non_block_types.insert(x: get_base_non_block_type_id(type_id: member)); |
5269 | else if (subtype.basetype == SPIRType::Struct && !compiler.is_pointer(type: subtype)) |
5270 | analyze_non_block_types_from_block(type: subtype); |
5271 | } |
5272 | } |
5273 | |
5274 | void Compiler::analyze_non_block_pointer_types() |
5275 | { |
5276 | PhysicalStorageBufferPointerHandler handler(*this); |
5277 | traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler); |
5278 | |
5279 | // Analyze any block declaration we have to make. It might contain |
5280 | // physical pointers to POD types which we never used, and thus never added to the list. |
5281 | // We'll need to add those pointer types to the set of types we declare. |
5282 | ir.for_each_typed_id<SPIRType>(op: [&](uint32_t id, SPIRType &type) { |
5283 | // Only analyze the raw block struct, not any pointer-to-struct, since that's just redundant. |
5284 | if (type.self == id && |
5285 | (has_decoration(id: type.self, decoration: DecorationBlock) || |
5286 | has_decoration(id: type.self, decoration: DecorationBufferBlock))) |
5287 | { |
5288 | handler.analyze_non_block_types_from_block(type); |
5289 | } |
5290 | }); |
5291 | |
5292 | physical_storage_non_block_pointer_types.reserve(count: handler.non_block_types.size()); |
5293 | for (auto type : handler.non_block_types) |
5294 | physical_storage_non_block_pointer_types.push_back(t: type); |
5295 | sort(first: begin(cont&: physical_storage_non_block_pointer_types), last: end(cont&: physical_storage_non_block_pointer_types)); |
5296 | physical_storage_type_to_alignment = std::move(handler.physical_block_type_meta); |
5297 | } |
5298 | |
5299 | bool Compiler::InterlockedResourceAccessPrepassHandler::handle(Op op, const uint32_t *, uint32_t) |
5300 | { |
5301 | if (op == OpBeginInvocationInterlockEXT || op == OpEndInvocationInterlockEXT) |
5302 | { |
5303 | if (interlock_function_id != 0 && interlock_function_id != call_stack.back()) |
5304 | { |
5305 | // Most complex case, we have no sensible way of dealing with this |
5306 | // other than taking the 100% conservative approach, exit early. |
5307 | split_function_case = true; |
5308 | return false; |
5309 | } |
5310 | else |
5311 | { |
5312 | interlock_function_id = call_stack.back(); |
5313 | // If this call is performed inside control flow we have a problem. |
5314 | auto &cfg = compiler.get_cfg_for_function(id: interlock_function_id); |
5315 | |
5316 | uint32_t from_block_id = compiler.get<SPIRFunction>(id: interlock_function_id).entry_block; |
5317 | bool outside_control_flow = cfg.node_terminates_control_flow_in_sub_graph(from: from_block_id, to: current_block_id); |
5318 | if (!outside_control_flow) |
5319 | control_flow_interlock = true; |
5320 | } |
5321 | } |
5322 | return true; |
5323 | } |
5324 | |
5325 | void Compiler::InterlockedResourceAccessPrepassHandler::rearm_current_block(const SPIRBlock &block) |
5326 | { |
5327 | current_block_id = block.self; |
5328 | } |
5329 | |
5330 | bool Compiler::InterlockedResourceAccessPrepassHandler::begin_function_scope(const uint32_t *args, uint32_t length) |
5331 | { |
5332 | if (length < 3) |
5333 | return false; |
5334 | call_stack.push_back(t: args[2]); |
5335 | return true; |
5336 | } |
5337 | |
5338 | bool Compiler::InterlockedResourceAccessPrepassHandler::end_function_scope(const uint32_t *, uint32_t) |
5339 | { |
5340 | call_stack.pop_back(); |
5341 | return true; |
5342 | } |
5343 | |
5344 | bool Compiler::InterlockedResourceAccessHandler::begin_function_scope(const uint32_t *args, uint32_t length) |
5345 | { |
5346 | if (length < 3) |
5347 | return false; |
5348 | |
5349 | if (args[2] == interlock_function_id) |
5350 | call_stack_is_interlocked = true; |
5351 | |
5352 | call_stack.push_back(t: args[2]); |
5353 | return true; |
5354 | } |
5355 | |
5356 | bool Compiler::InterlockedResourceAccessHandler::end_function_scope(const uint32_t *, uint32_t) |
5357 | { |
5358 | if (call_stack.back() == interlock_function_id) |
5359 | call_stack_is_interlocked = false; |
5360 | |
5361 | call_stack.pop_back(); |
5362 | return true; |
5363 | } |
5364 | |
5365 | void Compiler::InterlockedResourceAccessHandler::access_potential_resource(uint32_t id) |
5366 | { |
5367 | if ((use_critical_section && in_crit_sec) || (control_flow_interlock && call_stack_is_interlocked) || |
5368 | split_function_case) |
5369 | { |
5370 | compiler.interlocked_resources.insert(x: id); |
5371 | } |
5372 | } |
5373 | |
5374 | bool Compiler::InterlockedResourceAccessHandler::handle(Op opcode, const uint32_t *args, uint32_t length) |
5375 | { |
5376 | // Only care about critical section analysis if we have simple case. |
5377 | if (use_critical_section) |
5378 | { |
5379 | if (opcode == OpBeginInvocationInterlockEXT) |
5380 | { |
5381 | in_crit_sec = true; |
5382 | return true; |
5383 | } |
5384 | |
5385 | if (opcode == OpEndInvocationInterlockEXT) |
5386 | { |
5387 | // End critical section--nothing more to do. |
5388 | return false; |
5389 | } |
5390 | } |
5391 | |
5392 | // We need to figure out where images and buffers are loaded from, so do only the bare bones compilation we need. |
5393 | switch (opcode) |
5394 | { |
5395 | case OpLoad: |
5396 | { |
5397 | if (length < 3) |
5398 | return false; |
5399 | |
5400 | uint32_t ptr = args[2]; |
5401 | auto *var = compiler.maybe_get_backing_variable(chain: ptr); |
5402 | |
5403 | // We're only concerned with buffer and image memory here. |
5404 | if (!var) |
5405 | break; |
5406 | |
5407 | switch (var->storage) |
5408 | { |
5409 | default: |
5410 | break; |
5411 | |
5412 | case StorageClassUniformConstant: |
5413 | { |
5414 | uint32_t result_type = args[0]; |
5415 | uint32_t id = args[1]; |
5416 | compiler.set<SPIRExpression>(id, args: "", args&: result_type, args: true); |
5417 | compiler.register_read(expr: id, chain: ptr, forwarded: true); |
5418 | break; |
5419 | } |
5420 | |
5421 | case StorageClassUniform: |
5422 | // Must have BufferBlock; we only care about SSBOs. |
5423 | if (!compiler.has_decoration(id: compiler.get<SPIRType>(id: var->basetype).self, decoration: DecorationBufferBlock)) |
5424 | break; |
5425 | // fallthrough |
5426 | case StorageClassStorageBuffer: |
5427 | access_potential_resource(id: var->self); |
5428 | break; |
5429 | } |
5430 | break; |
5431 | } |
5432 | |
5433 | case OpInBoundsAccessChain: |
5434 | case OpAccessChain: |
5435 | case OpPtrAccessChain: |
5436 | { |
5437 | if (length < 3) |
5438 | return false; |
5439 | |
5440 | uint32_t result_type = args[0]; |
5441 | |
5442 | auto &type = compiler.get<SPIRType>(id: result_type); |
5443 | if (type.storage == StorageClassUniform || type.storage == StorageClassUniformConstant || |
5444 | type.storage == StorageClassStorageBuffer) |
5445 | { |
5446 | uint32_t id = args[1]; |
5447 | uint32_t ptr = args[2]; |
5448 | compiler.set<SPIRExpression>(id, args: "", args&: result_type, args: true); |
5449 | compiler.register_read(expr: id, chain: ptr, forwarded: true); |
5450 | compiler.ir.ids[id].set_allow_type_rewrite(); |
5451 | } |
5452 | break; |
5453 | } |
5454 | |
5455 | case OpImageTexelPointer: |
5456 | { |
5457 | if (length < 3) |
5458 | return false; |
5459 | |
5460 | uint32_t result_type = args[0]; |
5461 | uint32_t id = args[1]; |
5462 | uint32_t ptr = args[2]; |
5463 | auto &e = compiler.set<SPIRExpression>(id, args: "", args&: result_type, args: true); |
5464 | auto *var = compiler.maybe_get_backing_variable(chain: ptr); |
5465 | if (var) |
5466 | e.loaded_from = var->self; |
5467 | break; |
5468 | } |
5469 | |
5470 | case OpStore: |
5471 | case OpImageWrite: |
5472 | case OpAtomicStore: |
5473 | { |
5474 | if (length < 1) |
5475 | return false; |
5476 | |
5477 | uint32_t ptr = args[0]; |
5478 | auto *var = compiler.maybe_get_backing_variable(chain: ptr); |
5479 | if (var && (var->storage == StorageClassUniform || var->storage == StorageClassUniformConstant || |
5480 | var->storage == StorageClassStorageBuffer)) |
5481 | { |
5482 | access_potential_resource(id: var->self); |
5483 | } |
5484 | |
5485 | break; |
5486 | } |
5487 | |
5488 | case OpCopyMemory: |
5489 | { |
5490 | if (length < 2) |
5491 | return false; |
5492 | |
5493 | uint32_t dst = args[0]; |
5494 | uint32_t src = args[1]; |
5495 | auto *dst_var = compiler.maybe_get_backing_variable(chain: dst); |
5496 | auto *src_var = compiler.maybe_get_backing_variable(chain: src); |
5497 | |
5498 | if (dst_var && (dst_var->storage == StorageClassUniform || dst_var->storage == StorageClassStorageBuffer)) |
5499 | access_potential_resource(id: dst_var->self); |
5500 | |
5501 | if (src_var) |
5502 | { |
5503 | if (src_var->storage != StorageClassUniform && src_var->storage != StorageClassStorageBuffer) |
5504 | break; |
5505 | |
5506 | if (src_var->storage == StorageClassUniform && |
5507 | !compiler.has_decoration(id: compiler.get<SPIRType>(id: src_var->basetype).self, decoration: DecorationBufferBlock)) |
5508 | { |
5509 | break; |
5510 | } |
5511 | |
5512 | access_potential_resource(id: src_var->self); |
5513 | } |
5514 | |
5515 | break; |
5516 | } |
5517 | |
5518 | case OpImageRead: |
5519 | case OpAtomicLoad: |
5520 | { |
5521 | if (length < 3) |
5522 | return false; |
5523 | |
5524 | uint32_t ptr = args[2]; |
5525 | auto *var = compiler.maybe_get_backing_variable(chain: ptr); |
5526 | |
5527 | // We're only concerned with buffer and image memory here. |
5528 | if (!var) |
5529 | break; |
5530 | |
5531 | switch (var->storage) |
5532 | { |
5533 | default: |
5534 | break; |
5535 | |
5536 | case StorageClassUniform: |
5537 | // Must have BufferBlock; we only care about SSBOs. |
5538 | if (!compiler.has_decoration(id: compiler.get<SPIRType>(id: var->basetype).self, decoration: DecorationBufferBlock)) |
5539 | break; |
5540 | // fallthrough |
5541 | case StorageClassUniformConstant: |
5542 | case StorageClassStorageBuffer: |
5543 | access_potential_resource(id: var->self); |
5544 | break; |
5545 | } |
5546 | break; |
5547 | } |
5548 | |
5549 | case OpAtomicExchange: |
5550 | case OpAtomicCompareExchange: |
5551 | case OpAtomicIIncrement: |
5552 | case OpAtomicIDecrement: |
5553 | case OpAtomicIAdd: |
5554 | case OpAtomicISub: |
5555 | case OpAtomicSMin: |
5556 | case OpAtomicUMin: |
5557 | case OpAtomicSMax: |
5558 | case OpAtomicUMax: |
5559 | case OpAtomicAnd: |
5560 | case OpAtomicOr: |
5561 | case OpAtomicXor: |
5562 | { |
5563 | if (length < 3) |
5564 | return false; |
5565 | |
5566 | uint32_t ptr = args[2]; |
5567 | auto *var = compiler.maybe_get_backing_variable(chain: ptr); |
5568 | if (var && (var->storage == StorageClassUniform || var->storage == StorageClassUniformConstant || |
5569 | var->storage == StorageClassStorageBuffer)) |
5570 | { |
5571 | access_potential_resource(id: var->self); |
5572 | } |
5573 | |
5574 | break; |
5575 | } |
5576 | |
5577 | default: |
5578 | break; |
5579 | } |
5580 | |
5581 | return true; |
5582 | } |
5583 | |
5584 | void Compiler::analyze_interlocked_resource_usage() |
5585 | { |
5586 | if (get_execution_model() == ExecutionModelFragment && |
5587 | (get_entry_point().flags.get(bit: ExecutionModePixelInterlockOrderedEXT) || |
5588 | get_entry_point().flags.get(bit: ExecutionModePixelInterlockUnorderedEXT) || |
5589 | get_entry_point().flags.get(bit: ExecutionModeSampleInterlockOrderedEXT) || |
5590 | get_entry_point().flags.get(bit: ExecutionModeSampleInterlockUnorderedEXT))) |
5591 | { |
5592 | InterlockedResourceAccessPrepassHandler prepass_handler(*this, ir.default_entry_point); |
5593 | traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler&: prepass_handler); |
5594 | |
5595 | InterlockedResourceAccessHandler handler(*this, ir.default_entry_point); |
5596 | handler.interlock_function_id = prepass_handler.interlock_function_id; |
5597 | handler.split_function_case = prepass_handler.split_function_case; |
5598 | handler.control_flow_interlock = prepass_handler.control_flow_interlock; |
5599 | handler.use_critical_section = !handler.split_function_case && !handler.control_flow_interlock; |
5600 | |
5601 | traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler); |
5602 | |
5603 | // For GLSL. If we hit any of these cases, we have to fall back to conservative approach. |
5604 | interlocked_is_complex = |
5605 | !handler.use_critical_section || handler.interlock_function_id != ir.default_entry_point; |
5606 | } |
5607 | } |
5608 | |
5609 | // Helper function |
5610 | bool Compiler::check_internal_recursion(const SPIRType &type, std::unordered_set<uint32_t> &checked_ids) |
5611 | { |
5612 | if (type.basetype != SPIRType::Struct) |
5613 | return false; |
5614 | |
5615 | if (checked_ids.count(x: type.self)) |
5616 | return true; |
5617 | |
5618 | // Recurse into struct members |
5619 | bool is_recursive = false; |
5620 | checked_ids.insert(x: type.self); |
5621 | uint32_t mbr_cnt = uint32_t(type.member_types.size()); |
5622 | for (uint32_t mbr_idx = 0; !is_recursive && mbr_idx < mbr_cnt; mbr_idx++) |
5623 | { |
5624 | uint32_t mbr_type_id = type.member_types[mbr_idx]; |
5625 | auto &mbr_type = get<SPIRType>(id: mbr_type_id); |
5626 | is_recursive |= check_internal_recursion(type: mbr_type, checked_ids); |
5627 | } |
5628 | checked_ids.erase(x: type.self); |
5629 | return is_recursive; |
5630 | } |
5631 | |
5632 | // Return whether the struct type contains a structural recursion nested somewhere within its content. |
5633 | bool Compiler::type_contains_recursion(const SPIRType &type) |
5634 | { |
5635 | std::unordered_set<uint32_t> checked_ids; |
5636 | return check_internal_recursion(type, checked_ids); |
5637 | } |
5638 | |
5639 | bool Compiler::type_is_array_of_pointers(const SPIRType &type) const |
5640 | { |
5641 | if (!is_array(type)) |
5642 | return false; |
5643 | |
5644 | // BDA types must have parent type hierarchy. |
5645 | if (!type.parent_type) |
5646 | return false; |
5647 | |
5648 | // Punch through all array layers. |
5649 | auto *parent = &get<SPIRType>(id: type.parent_type); |
5650 | while (is_array(type: *parent)) |
5651 | parent = &get<SPIRType>(id: parent->parent_type); |
5652 | |
5653 | return is_pointer(type: *parent); |
5654 | } |
5655 | |
5656 | bool Compiler::flush_phi_required(BlockID from, BlockID to) const |
5657 | { |
5658 | auto &child = get<SPIRBlock>(id: to); |
5659 | for (auto &phi : child.phi_variables) |
5660 | if (phi.parent == from) |
5661 | return true; |
5662 | return false; |
5663 | } |
5664 | |
5665 | void Compiler::add_loop_level() |
5666 | { |
5667 | current_loop_level++; |
5668 | } |
5669 |
Definitions
- Compiler
- Compiler
- Compiler
- Compiler
- set_ir
- set_ir
- compile
- variable_storage_is_aliased
- block_is_control_dependent
- block_is_pure
- to_name
- function_is_pure
- function_is_control_dependent
- register_global_read_dependencies
- register_global_read_dependencies
- maybe_get_backing_variable
- register_read
- register_write
- flush_dependees
- flush_all_aliased_variables
- flush_all_atomic_capable_variables
- flush_control_dependent_expressions
- flush_all_active_variables
- expression_type_id
- expression_type
- expression_is_lvalue
- is_immutable
- storage_class_is_interface
- is_hidden_variable
- is_builtin_type
- is_builtin_variable
- is_member_builtin
- is_scalar
- is_vector
- is_matrix
- is_array
- is_pointer
- is_physical_pointer
- is_physical_pointer_to_buffer_block
- is_runtime_size_array
- get_shader_resources
- get_shader_resources
- handle
- get_active_interface_variables
- set_enabled_interface_variables
- get_shader_resources
- type_is_top_level_block
- type_is_block_like
- parse_fixup
- update_name_cache
- update_name_cache
- set_name
- get_type
- get_type_from_variable
- get_pointee_type_id
- get_pointee_type
- get_pointee_type
- get_variable_data_type_id
- get_variable_data_type
- get_variable_data_type
- get_variable_element_type
- get_variable_element_type
- is_sampled_image_type
- set_member_decoration_string
- set_member_decoration
- set_member_name
- get_member_name
- set_qualified_name
- set_member_qualified_name
- get_member_qualified_name
- get_member_decoration
- get_member_decoration_bitset
- has_member_decoration
- unset_member_decoration
- set_decoration_string
- set_decoration
- set_extended_decoration
- set_extended_member_decoration
- get_default_extended_decoration
- get_extended_decoration
- get_extended_member_decoration
- has_extended_decoration
- has_extended_member_decoration
- unset_extended_decoration
- unset_extended_member_decoration
- get_storage_class
- get_name
- get_fallback_name
- get_block_fallback_name
- get_decoration_bitset
- has_decoration
- get_decoration_string
- get_member_decoration_string
- get_decoration
- unset_decoration
- get_binary_offset_for_decoration
- block_is_noop
- block_is_loop_candidate
- execution_is_noop
- execution_is_branchless
- execution_is_direct_branch
- continue_block_type
- get_case_list
- traverse_all_reachable_opcodes
- traverse_all_reachable_opcodes
- type_struct_member_offset
- type_struct_member_array_stride
- type_struct_member_matrix_stride
- get_declared_struct_size
- get_declared_struct_size_runtime_array
- evaluate_spec_constant_u32
- evaluate_constant_u32
- get_declared_struct_member_size
- handle
- get_active_buffer_ranges
- types_are_logically_equivalent
- get_execution_mode_bitset
- set_execution_mode
- unset_execution_mode
- get_work_group_size_specialization_constants
- get_execution_mode_argument
- get_execution_model
- is_tessellation_shader
- is_vertex_like_shader
- is_tessellation_shader
- is_tessellating_triangles
- set_remapped_variable_state
- get_remapped_variable_state
- set_subpass_input_remapped_components
- get_subpass_input_remapped_components
- add_implied_read_expression
- add_implied_read_expression
- add_active_interface_variable
- inherit_expression_dependencies
- get_entry_points_and_stages
- rename_entry_point
- set_entry_point
- get_first_entry_point
- get_first_entry_point
- get_entry_point
- get_entry_point
- get_cleansed_entry_point_name
- get_entry_point
- get_entry_point
- interface_variable_exists_in_entry_point
- push_remap_parameters
- pop_remap_parameters
- remap_parameter
- begin_function_scope
- end_function_scope
- register_combined_image_sampler
- handle
- handle
- build_dummy_sampler_for_combined_images
- build_combined_image_samplers
- get_specialization_constants
- get_constant
- get_constant
- exists_unaccessed_path_to_return
- analyze_parameter_preservation
- AnalyzeVariableScopeAccessHandler
- follow_function_call
- set_current_block
- notify_variable_access
- id_is_phi_variable
- id_is_potential_temporary
- handle_terminator
- handle
- StaticExpressionAccessHandler
- follow_function_call
- handle
- find_function_local_luts
- analyze_variable_scope
- may_read_undefined_variable_in_block
- get_buffer_block_flags
- get_common_basic_type
- handle_builtin
- add_if_builtin
- add_if_builtin
- add_if_builtin_or_block
- handle
- update_active_builtins
- has_active_builtin
- analyze_image_and_sampler_usage
- handle
- get_cfg_for_current_function
- get_cfg_for_function
- build_function_control_flow_graphs_and_analyze
- CFGBuilder
- handle
- follow_function_call
- add_dependency
- begin_function_scope
- add_hierarchy_to_comparison_ids
- handle
- buffer_is_hlsl_counter_buffer
- buffer_get_hlsl_counter_buffer
- make_constant_null
- get_declared_capabilities
- get_declared_extensions
- get_remapped_declared_block_name
- get_remapped_declared_block_name
- reflection_ssbo_instance_name_is_significant
- instruction_to_result_type
- combined_decoration_for_member
- is_desktop_only_format
- is_depth_image
- type_is_opaque_value
- force_recompile
- force_recompile_guarantee_forward_progress
- is_forcing_recompilation
- clear_force_recompile
- PhysicalStorageBufferPointerHandler
- find_block_meta
- mark_aligned_access
- type_is_bda_block_entry
- get_minimum_scalar_alignment
- setup_meta_chain
- handle
- get_base_non_block_type_id
- analyze_non_block_types_from_block
- analyze_non_block_pointer_types
- handle
- rearm_current_block
- begin_function_scope
- end_function_scope
- begin_function_scope
- end_function_scope
- access_potential_resource
- handle
- analyze_interlocked_resource_usage
- check_internal_recursion
- type_contains_recursion
- type_is_array_of_pointers
- flush_phi_required
Learn Advanced QML with KDAB
Find out more