blob: cbe9b26ac02e76dc32c32cfb7e037dadd24974b5 [file] [log] [blame]
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001/*
Hans-Kristian Arntzen318c17c2019-01-04 12:38:35 +01002 * Copyright 2015-2019 Arm Limited
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01003 *
4 * Licensed under the Apache License, Version 2.0 (the "License");
5 * you may not use this file except in compliance with the License.
6 * You may obtain a copy of the License at
7 *
8 * http://www.apache.org/licenses/LICENSE-2.0
9 *
10 * Unless required by applicable law or agreed to in writing, software
11 * distributed under the License is distributed on an "AS IS" BASIS,
12 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 * See the License for the specific language governing permissions and
14 * limitations under the License.
15 */
16
Hans-Kristian Arntzen147e53a2016-04-04 09:36:04 +020017#include "spirv_cross.hpp"
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010018#include "GLSL.std.450.h"
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +010019#include "spirv_cfg.hpp"
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020020#include "spirv_parser.hpp"
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010021#include <algorithm>
Hans-Kristian Arntzen5ea59bd2016-05-23 13:30:02 +020022#include <cstring>
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010023#include <utility>
24
25using namespace std;
26using namespace spv;
Hans-Kristian Arntzen9b92e682019-03-29 10:29:44 +010027using namespace SPIRV_CROSS_NAMESPACE;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010028
Hans-Kristian Arntzen3fe57d32019-04-09 12:46:23 +020029Compiler::Compiler(vector<uint32_t> ir_)
Hans-Kristian Arntzenfd432f82017-03-18 10:52:41 +010030{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020031 Parser parser(move(ir_));
32 parser.parse();
33 set_ir(move(parser.get_parsed_ir()));
Hans-Kristian Arntzenfd432f82017-03-18 10:52:41 +010034}
35
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020036Compiler::Compiler(const uint32_t *ir_, size_t word_count)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010037{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020038 Parser parser(ir_, word_count);
39 parser.parse();
40 set_ir(move(parser.get_parsed_ir()));
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010041}
42
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020043Compiler::Compiler(const ParsedIR &ir_)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010044{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020045 set_ir(ir_);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010046}
47
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020048Compiler::Compiler(ParsedIR &&ir_)
Yuriy O'Donnellae8de512017-04-01 12:31:34 +020049{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020050 set_ir(move(ir_));
51}
52
53void Compiler::set_ir(ParsedIR &&ir_)
54{
55 ir = move(ir_);
56 parse_fixup();
57}
58
59void Compiler::set_ir(const ParsedIR &ir_)
60{
61 ir = ir_;
62 parse_fixup();
Yuriy O'Donnellae8de512017-04-01 12:31:34 +020063}
64
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020065string Compiler::compile()
66{
67 return "";
68}
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010069
70bool Compiler::variable_storage_is_aliased(const SPIRVariable &v)
71{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020072 auto &type = get<SPIRType>(v.basetype);
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +020073 bool ssbo = v.storage == StorageClassStorageBuffer ||
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020074 ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020075 bool image = type.basetype == SPIRType::Image;
76 bool counter = type.basetype == SPIRType::AtomicCounter;
Hans-Kristian Arntzen7eba2472018-05-11 10:14:20 +020077
78 bool is_restrict;
79 if (ssbo)
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +020080 is_restrict = ir.get_buffer_block_flags(v).get(DecorationRestrict);
Hans-Kristian Arntzen7eba2472018-05-11 10:14:20 +020081 else
82 is_restrict = has_decoration(v.self, DecorationRestrict);
83
Hans-Kristian Arntzen7d8add32016-07-12 15:00:10 +020084 return !is_restrict && (ssbo || image || counter);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010085}
86
87bool Compiler::block_is_pure(const SPIRBlock &block)
88{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020089 for (auto &i : block.ops)
90 {
91 auto ops = stream(i);
92 auto op = static_cast<Op>(i.op);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010093
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020094 switch (op)
95 {
96 case OpFunctionCall:
97 {
98 uint32_t func = ops[2];
99 if (!function_is_pure(get<SPIRFunction>(func)))
100 return false;
101 break;
102 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100103
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +0100104 case OpCopyMemory:
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200105 case OpStore:
106 {
107 auto &type = expression_type(ops[0]);
108 if (type.storage != StorageClassFunction)
109 return false;
110 break;
111 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100112
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200113 case OpImageWrite:
114 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100115
Hans-Kristian Arntzen5af1a512016-05-05 09:51:42 +0200116 // Atomics are impure.
117 case OpAtomicLoad:
118 case OpAtomicStore:
119 case OpAtomicExchange:
120 case OpAtomicCompareExchange:
Bill Hollings8f6df772017-05-19 18:14:08 -0400121 case OpAtomicCompareExchangeWeak:
Hans-Kristian Arntzen5af1a512016-05-05 09:51:42 +0200122 case OpAtomicIIncrement:
123 case OpAtomicIDecrement:
124 case OpAtomicIAdd:
125 case OpAtomicISub:
126 case OpAtomicSMin:
127 case OpAtomicUMin:
128 case OpAtomicSMax:
129 case OpAtomicUMax:
130 case OpAtomicAnd:
131 case OpAtomicOr:
132 case OpAtomicXor:
133 return false;
134
135 // Geometry shader builtins modify global state.
136 case OpEndPrimitive:
137 case OpEmitStreamVertex:
138 case OpEndStreamPrimitive:
139 case OpEmitVertex:
140 return false;
141
142 // Barriers disallow any reordering, so we should treat blocks with barrier as writing.
143 case OpControlBarrier:
144 case OpMemoryBarrier:
145 return false;
146
Patrick Mours90c91e42019-03-26 14:16:38 +0100147 // Ray tracing builtins are impure.
148 case OpReportIntersectionNV:
149 case OpIgnoreIntersectionNV:
150 case OpTerminateRayNV:
151 case OpTraceNV:
152 case OpExecuteCallableNV:
153 return false;
154
Hans-Kristian Arntzence18d4c2017-11-17 13:38:29 +0100155 // OpExtInst is potentially impure depending on extension, but GLSL builtins are at least pure.
Hans-Kristian Arntzen5af1a512016-05-05 09:51:42 +0200156
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200157 default:
158 break;
159 }
160 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100161
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200162 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100163}
164
Hans-Kristian Arntzen61c31c62017-03-07 13:27:04 +0100165string Compiler::to_name(uint32_t id, bool allow_alias) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100166{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100167 if (allow_alias && ir.ids[id].get_type() == TypeType)
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +0200168 {
169 // If this type is a simple alias, emit the
170 // name of the original type instead.
171 // We don't want to override the meta alias
172 // as that can be overridden by the reflection APIs after parse.
173 auto &type = get<SPIRType>(id);
174 if (type.type_alias)
Hans-Kristian Arntzen294259e2018-03-05 16:27:04 +0100175 {
176 // If the alias master has been specially packed, we will have emitted a clean variant as well,
177 // so skip the name aliasing here.
Hans-Kristian Arntzende7e5cc2019-01-17 11:22:24 +0100178 if (!has_extended_decoration(type.type_alias, SPIRVCrossDecorationPacked))
Hans-Kristian Arntzen294259e2018-03-05 16:27:04 +0100179 return to_name(type.type_alias);
180 }
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +0200181 }
182
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100183 auto &alias = ir.get_name(id);
184 if (alias.empty())
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200185 return join("_", id);
186 else
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100187 return alias;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100188}
189
190bool Compiler::function_is_pure(const SPIRFunction &func)
191{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200192 for (auto block : func.blocks)
193 {
194 if (!block_is_pure(get<SPIRBlock>(block)))
195 {
196 //fprintf(stderr, "Function %s is impure!\n", to_name(func.self).c_str());
197 return false;
198 }
199 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100200
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200201 //fprintf(stderr, "Function %s is pure!\n", to_name(func.self).c_str());
202 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100203}
204
205void Compiler::register_global_read_dependencies(const SPIRBlock &block, uint32_t id)
206{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200207 for (auto &i : block.ops)
208 {
209 auto ops = stream(i);
210 auto op = static_cast<Op>(i.op);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100211
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200212 switch (op)
213 {
214 case OpFunctionCall:
215 {
216 uint32_t func = ops[2];
217 register_global_read_dependencies(get<SPIRFunction>(func), id);
218 break;
219 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100220
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200221 case OpLoad:
222 case OpImageRead:
223 {
224 // If we're in a storage class which does not get invalidated, adding dependencies here is no big deal.
225 auto *var = maybe_get_backing_variable(ops[2]);
226 if (var && var->storage != StorageClassFunction)
227 {
228 auto &type = get<SPIRType>(var->basetype);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100229
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200230 // InputTargets are immutable.
231 if (type.basetype != SPIRType::Image && type.image.dim != DimSubpassData)
232 var->dependees.push_back(id);
233 }
234 break;
235 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100236
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200237 default:
238 break;
239 }
240 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100241}
242
243void Compiler::register_global_read_dependencies(const SPIRFunction &func, uint32_t id)
244{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200245 for (auto block : func.blocks)
246 register_global_read_dependencies(get<SPIRBlock>(block), id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100247}
248
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200249SPIRVariable *Compiler::maybe_get_backing_variable(uint32_t chain)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100250{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200251 auto *var = maybe_get<SPIRVariable>(chain);
252 if (!var)
253 {
254 auto *cexpr = maybe_get<SPIRExpression>(chain);
255 if (cexpr)
256 var = maybe_get<SPIRVariable>(cexpr->loaded_from);
Hans-Kristian Arntzen7d7f4b32017-08-10 17:12:48 +0200257
258 auto *access_chain = maybe_get<SPIRAccessChain>(chain);
259 if (access_chain)
260 var = maybe_get<SPIRVariable>(access_chain->loaded_from);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200261 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100262
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200263 return var;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100264}
265
266void Compiler::register_read(uint32_t expr, uint32_t chain, bool forwarded)
267{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200268 auto &e = get<SPIRExpression>(expr);
269 auto *var = maybe_get_backing_variable(chain);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100270
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200271 if (var)
272 {
273 e.loaded_from = var->self;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100274
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200275 // If the backing variable is immutable, we do not need to depend on the variable.
276 if (forwarded && !is_immutable(var->self))
277 var->dependees.push_back(e.self);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100278
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200279 // If we load from a parameter, make sure we create "inout" if we also write to the parameter.
280 // The default is "in" however, so we never invalidate our compilation by reading.
281 if (var && var->parameter)
282 var->parameter->read_count++;
283 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100284}
285
286void Compiler::register_write(uint32_t chain)
287{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200288 auto *var = maybe_get<SPIRVariable>(chain);
289 if (!var)
290 {
291 // If we're storing through an access chain, invalidate the backing variable instead.
292 auto *expr = maybe_get<SPIRExpression>(chain);
293 if (expr && expr->loaded_from)
294 var = maybe_get<SPIRVariable>(expr->loaded_from);
Hans-Kristian Arntzen7d7f4b32017-08-10 17:12:48 +0200295
296 auto *access_chain = maybe_get<SPIRAccessChain>(chain);
297 if (access_chain && access_chain->loaded_from)
298 var = maybe_get<SPIRVariable>(access_chain->loaded_from);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200299 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100300
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200301 if (var)
302 {
303 // If our variable is in a storage class which can alias with other buffers,
Chip Davisd6aa9112019-01-08 15:16:17 -0600304 // invalidate all variables which depend on aliased variables. And if this is a
305 // variable pointer, then invalidate all variables regardless.
306 if (get_variable_data_type(*var).pointer)
307 flush_all_active_variables();
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200308 if (variable_storage_is_aliased(*var))
309 flush_all_aliased_variables();
310 else if (var)
311 flush_dependees(*var);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100312
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200313 // We tried to write to a parameter which is not marked with out qualifier, force a recompile.
314 if (var->parameter && var->parameter->write_count == 0)
315 {
316 var->parameter->write_count++;
Hans-Kristian Arntzen317144a2019-04-05 12:06:10 +0200317 force_recompile();
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200318 }
319 }
Chip Davisd6aa9112019-01-08 15:16:17 -0600320 else
321 {
322 // If we stored through a variable pointer, then we don't know which
323 // variable we stored to. So *all* expressions after this point need to
324 // be invalidated.
325 // FIXME: If we can prove that the variable pointer will point to
326 // only certain variables, we can invalidate only those.
327 flush_all_active_variables();
328 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100329}
330
331void Compiler::flush_dependees(SPIRVariable &var)
332{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200333 for (auto expr : var.dependees)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200334 invalid_expressions.insert(expr);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200335 var.dependees.clear();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100336}
337
338void Compiler::flush_all_aliased_variables()
339{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200340 for (auto aliased : aliased_variables)
341 flush_dependees(get<SPIRVariable>(aliased));
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100342}
343
344void Compiler::flush_all_atomic_capable_variables()
345{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200346 for (auto global : global_variables)
347 flush_dependees(get<SPIRVariable>(global));
348 flush_all_aliased_variables();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100349}
350
Hans-Kristian Arntzen938c7de2018-03-12 17:34:54 +0100351void Compiler::flush_control_dependent_expressions(uint32_t block_id)
352{
353 auto &block = get<SPIRBlock>(block_id);
354 for (auto &expr : block.invalidate_expressions)
355 invalid_expressions.insert(expr);
356 block.invalidate_expressions.clear();
357}
358
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100359void Compiler::flush_all_active_variables()
360{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200361 // Invalidate all temporaries we read from variables in this block since they were forwarded.
362 // Invalidate all temporaries we read from globals.
363 for (auto &v : current_function->local_variables)
364 flush_dependees(get<SPIRVariable>(v));
365 for (auto &arg : current_function->arguments)
366 flush_dependees(get<SPIRVariable>(arg.id));
367 for (auto global : global_variables)
368 flush_dependees(get<SPIRVariable>(global));
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100369
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200370 flush_all_aliased_variables();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100371}
372
Bill Hollings1e84a372017-08-12 00:21:13 -0400373uint32_t Compiler::expression_type_id(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100374{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200375 switch (ir.ids[id].get_type())
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200376 {
377 case TypeVariable:
Bill Hollings1e84a372017-08-12 00:21:13 -0400378 return get<SPIRVariable>(id).basetype;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100379
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200380 case TypeExpression:
Bill Hollings1e84a372017-08-12 00:21:13 -0400381 return get<SPIRExpression>(id).expression_type;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100382
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200383 case TypeConstant:
Bill Hollings1e84a372017-08-12 00:21:13 -0400384 return get<SPIRConstant>(id).constant_type;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100385
Hans-Kristian Arntzen7e8afa82016-10-03 15:54:02 +0200386 case TypeConstantOp:
Bill Hollings1e84a372017-08-12 00:21:13 -0400387 return get<SPIRConstantOp>(id).basetype;
Hans-Kristian Arntzen7e8afa82016-10-03 15:54:02 +0200388
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200389 case TypeUndef:
Bill Hollings1e84a372017-08-12 00:21:13 -0400390 return get<SPIRUndef>(id).basetype;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100391
Hans-Kristian Arntzen100e9d32017-04-25 10:44:55 +0200392 case TypeCombinedImageSampler:
Bill Hollings1e84a372017-08-12 00:21:13 -0400393 return get<SPIRCombinedImageSampler>(id).combined_type;
Hans-Kristian Arntzen100e9d32017-04-25 10:44:55 +0200394
Hans-Kristian Arntzen3cbdbec2017-08-10 15:36:30 +0200395 case TypeAccessChain:
Hans-Kristian Arntzene2bb5b82017-08-15 09:34:30 +0200396 return get<SPIRAccessChain>(id).basetype;
Hans-Kristian Arntzen3cbdbec2017-08-10 15:36:30 +0200397
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200398 default:
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100399 SPIRV_CROSS_THROW("Cannot resolve expression type.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200400 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100401}
402
Bill Hollings1e84a372017-08-12 00:21:13 -0400403const SPIRType &Compiler::expression_type(uint32_t id) const
404{
405 return get<SPIRType>(expression_type_id(id));
406}
407
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100408bool Compiler::expression_is_lvalue(uint32_t id) const
409{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200410 auto &type = expression_type(id);
411 switch (type.basetype)
412 {
413 case SPIRType::SampledImage:
414 case SPIRType::Image:
415 case SPIRType::Sampler:
416 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100417
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200418 default:
419 return true;
420 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100421}
422
423bool Compiler::is_immutable(uint32_t id) const
424{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200425 if (ir.ids[id].get_type() == TypeVariable)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200426 {
427 auto &var = get<SPIRVariable>(id);
Hans-Kristian Arntzen92134e42016-04-01 19:58:26 +0200428
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200429 // Anything we load from the UniformConstant address space is guaranteed to be immutable.
430 bool pointer_to_const = var.storage == StorageClassUniformConstant;
Hans-Kristian Arntzen36a0b632016-07-12 14:33:04 +0200431 return pointer_to_const || var.phi_variable || !expression_is_lvalue(id);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200432 }
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200433 else if (ir.ids[id].get_type() == TypeAccessChain)
Hans-Kristian Arntzen7d7f4b32017-08-10 17:12:48 +0200434 return get<SPIRAccessChain>(id).immutable;
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200435 else if (ir.ids[id].get_type() == TypeExpression)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200436 return get<SPIRExpression>(id).immutable;
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200437 else if (ir.ids[id].get_type() == TypeConstant || ir.ids[id].get_type() == TypeConstantOp ||
438 ir.ids[id].get_type() == TypeUndef)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200439 return true;
440 else
441 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100442}
443
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200444static inline bool storage_class_is_interface(spv::StorageClass storage)
445{
446 switch (storage)
447 {
448 case StorageClassInput:
449 case StorageClassOutput:
450 case StorageClassUniform:
451 case StorageClassUniformConstant:
452 case StorageClassAtomicCounter:
453 case StorageClassPushConstant:
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +0200454 case StorageClassStorageBuffer:
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200455 return true;
456
457 default:
458 return false;
459 }
460}
461
462bool Compiler::is_hidden_variable(const SPIRVariable &var, bool include_builtins) const
463{
464 if ((is_builtin_variable(var) && !include_builtins) || var.remapped_variable)
465 return true;
466
Hans-Kristian Arntzen1b5ca8d2016-09-10 16:20:19 +0200467 // Combined image samplers are always considered active as they are "magic" variables.
468 if (find_if(begin(combined_image_samplers), end(combined_image_samplers), [&var](const CombinedImageSampler &samp) {
469 return samp.combined_id == var.self;
Hans-Kristian Arntzence18d4c2017-11-17 13:38:29 +0100470 }) != end(combined_image_samplers))
Hans-Kristian Arntzen1b5ca8d2016-09-10 16:20:19 +0200471 {
472 return false;
473 }
474
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200475 bool hidden = false;
476 if (check_active_interface_variables && storage_class_is_interface(var.storage))
477 hidden = active_interface_variables.find(var.self) == end(active_interface_variables);
478 return hidden;
479}
480
Hans-Kristian Arntzend3100602018-09-13 14:42:05 +0200481bool Compiler::is_builtin_type(const SPIRType &type) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100482{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100483 auto *type_meta = ir.find_meta(type.self);
484
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200485 // We can have builtin structs as well. If one member of a struct is builtin, the struct must also be builtin.
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100486 if (type_meta)
487 for (auto &m : type_meta->members)
488 if (m.builtin)
489 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100490
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200491 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100492}
493
Hans-Kristian Arntzend3100602018-09-13 14:42:05 +0200494bool Compiler::is_builtin_variable(const SPIRVariable &var) const
495{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100496 auto *m = ir.find_meta(var.self);
497
498 if (var.compat_builtin || (m && m->decoration.builtin))
Hans-Kristian Arntzend3100602018-09-13 14:42:05 +0200499 return true;
500 else
501 return is_builtin_type(get<SPIRType>(var.basetype));
502}
503
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100504bool Compiler::is_member_builtin(const SPIRType &type, uint32_t index, BuiltIn *builtin) const
505{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100506 auto *type_meta = ir.find_meta(type.self);
507
508 if (type_meta)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200509 {
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100510 auto &memb = type_meta->members;
511 if (index < memb.size() && memb[index].builtin)
512 {
513 if (builtin)
514 *builtin = memb[index].builtin_type;
515 return true;
516 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200517 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100518
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200519 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100520}
521
Bill Hollings103aabf2016-04-06 17:42:27 -0400522bool Compiler::is_scalar(const SPIRType &type) const
523{
Hans-Kristian Arntzen64ca1ec2019-01-16 16:16:39 +0100524 return type.basetype != SPIRType::Struct && type.vecsize == 1 && type.columns == 1;
Bill Hollings103aabf2016-04-06 17:42:27 -0400525}
526
527bool Compiler::is_vector(const SPIRType &type) const
528{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200529 return type.vecsize > 1 && type.columns == 1;
Bill Hollings103aabf2016-04-06 17:42:27 -0400530}
531
532bool Compiler::is_matrix(const SPIRType &type) const
533{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200534 return type.vecsize > 1 && type.columns > 1;
Bill Hollings103aabf2016-04-06 17:42:27 -0400535}
536
Bill Hollingsf591bc02017-06-30 19:10:46 -0400537bool Compiler::is_array(const SPIRType &type) const
538{
539 return !type.array.empty();
540}
541
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100542ShaderResources Compiler::get_shader_resources() const
543{
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200544 return get_shader_resources(nullptr);
545}
546
547ShaderResources Compiler::get_shader_resources(const unordered_set<uint32_t> &active_variables) const
548{
549 return get_shader_resources(&active_variables);
550}
551
552bool Compiler::InterfaceVariableAccessHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
553{
554 uint32_t variable = 0;
555 switch (opcode)
556 {
557 // Need this first, otherwise, GCC complains about unhandled switch statements.
558 default:
559 break;
560
561 case OpFunctionCall:
562 {
563 // Invalid SPIR-V.
564 if (length < 3)
565 return false;
566
567 uint32_t count = length - 3;
568 args += 3;
569 for (uint32_t i = 0; i < count; i++)
570 {
571 auto *var = compiler.maybe_get<SPIRVariable>(args[i]);
572 if (var && storage_class_is_interface(var->storage))
573 variables.insert(args[i]);
574 }
575 break;
576 }
577
Chip Davis3bfb2f92018-12-03 02:06:33 -0600578 case OpSelect:
579 {
580 // Invalid SPIR-V.
581 if (length < 5)
582 return false;
583
584 uint32_t count = length - 3;
585 args += 3;
586 for (uint32_t i = 0; i < count; i++)
587 {
588 auto *var = compiler.maybe_get<SPIRVariable>(args[i]);
589 if (var && storage_class_is_interface(var->storage))
590 variables.insert(args[i]);
591 }
592 break;
593 }
594
595 case OpPhi:
596 {
597 // Invalid SPIR-V.
598 if (length < 2)
599 return false;
600
601 uint32_t count = length - 2;
602 args += 2;
603 for (uint32_t i = 0; i < count; i += 2)
604 {
605 auto *var = compiler.maybe_get<SPIRVariable>(args[i]);
606 if (var && storage_class_is_interface(var->storage))
607 variables.insert(args[i]);
608 }
609 break;
610 }
611
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200612 case OpAtomicStore:
613 case OpStore:
614 // Invalid SPIR-V.
615 if (length < 1)
616 return false;
617 variable = args[0];
618 break;
619
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +0100620 case OpCopyMemory:
621 {
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +0100622 if (length < 2)
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +0100623 return false;
624
625 auto *var = compiler.maybe_get<SPIRVariable>(args[0]);
626 if (var && storage_class_is_interface(var->storage))
627 variables.insert(variable);
628
629 var = compiler.maybe_get<SPIRVariable>(args[1]);
630 if (var && storage_class_is_interface(var->storage))
631 variables.insert(variable);
632 break;
633 }
634
Hans-Kristian Arntzenbcdff2d2017-11-22 19:27:03 +0100635 case OpExtInst:
636 {
637 if (length < 5)
638 return false;
639 uint32_t extension_set = args[2];
640 if (compiler.get<SPIRExtension>(extension_set).ext == SPIRExtension::SPV_AMD_shader_explicit_vertex_parameter)
641 {
642 enum AMDShaderExplicitVertexParameter
643 {
644 InterpolateAtVertexAMD = 1
645 };
646
647 auto op = static_cast<AMDShaderExplicitVertexParameter>(args[3]);
648
649 switch (op)
650 {
651 case InterpolateAtVertexAMD:
652 {
653 auto *var = compiler.maybe_get<SPIRVariable>(args[4]);
654 if (var && storage_class_is_interface(var->storage))
655 variables.insert(args[4]);
656 break;
657 }
658
659 default:
660 break;
661 }
662 }
663 break;
664 }
665
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200666 case OpAccessChain:
667 case OpInBoundsAccessChain:
Chip Davis3bfb2f92018-12-03 02:06:33 -0600668 case OpPtrAccessChain:
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200669 case OpLoad:
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +0100670 case OpCopyObject:
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200671 case OpImageTexelPointer:
672 case OpAtomicLoad:
673 case OpAtomicExchange:
674 case OpAtomicCompareExchange:
Bill Hollings8f6df772017-05-19 18:14:08 -0400675 case OpAtomicCompareExchangeWeak:
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200676 case OpAtomicIIncrement:
677 case OpAtomicIDecrement:
678 case OpAtomicIAdd:
679 case OpAtomicISub:
680 case OpAtomicSMin:
681 case OpAtomicUMin:
682 case OpAtomicSMax:
683 case OpAtomicUMax:
684 case OpAtomicAnd:
685 case OpAtomicOr:
686 case OpAtomicXor:
687 // Invalid SPIR-V.
688 if (length < 3)
689 return false;
690 variable = args[2];
691 break;
692 }
693
694 if (variable)
695 {
696 auto *var = compiler.maybe_get<SPIRVariable>(variable);
697 if (var && storage_class_is_interface(var->storage))
698 variables.insert(variable);
699 }
700 return true;
701}
702
703unordered_set<uint32_t> Compiler::get_active_interface_variables() const
704{
705 // Traverse the call graph and find all interface variables which are in use.
706 unordered_set<uint32_t> variables;
707 InterfaceVariableAccessHandler handler(*this, variables);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200708 traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
Hans-Kristian Arntzen1a2e4de2018-02-21 13:43:16 +0100709
Hans-Kristian Arntzen3e098792019-01-30 10:29:08 +0100710 // Make sure we preserve output variables which are only initialized, but never accessed by any code.
711 ir.for_each_typed_id<SPIRVariable>([&](uint32_t, const SPIRVariable &var) {
712 if (var.storage == StorageClassOutput && var.initializer != 0)
713 variables.insert(var.self);
714 });
715
Hans-Kristian Arntzen1a2e4de2018-02-21 13:43:16 +0100716 // If we needed to create one, we'll need it.
717 if (dummy_sampler_id)
718 variables.insert(dummy_sampler_id);
719
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200720 return variables;
721}
722
723void Compiler::set_enabled_interface_variables(std::unordered_set<uint32_t> active_variables)
724{
725 active_interface_variables = move(active_variables);
726 check_active_interface_variables = true;
727}
728
729ShaderResources Compiler::get_shader_resources(const unordered_set<uint32_t> *active_variables) const
730{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200731 ShaderResources res;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100732
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +0100733 ir.for_each_typed_id<SPIRVariable>([&](uint32_t, const SPIRVariable &var) {
Hans-Kristian Arntzen2fb9aa22019-01-11 09:29:28 +0100734 auto &type = this->get<SPIRType>(var.basetype);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100735
Hans-Kristian Arntzend5dc5f32016-07-05 13:21:26 +0200736 // It is possible for uniform storage classes to be passed as function parameters, so detect
737 // that. To detect function parameters, check of StorageClass of variable is function scope.
738 if (var.storage == StorageClassFunction || !type.pointer || is_builtin_variable(var))
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +0100739 return;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100740
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200741 if (active_variables && active_variables->find(var.self) == end(*active_variables))
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +0100742 return;
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200743
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200744 // Input
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +0200745 if (var.storage == StorageClassInput && interface_variable_exists_in_entry_point(var.self))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200746 {
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100747 if (has_decoration(type.self, DecorationBlock))
748 {
Hans-Kristian Arntzenaa2557c2017-12-05 09:58:12 +0100749 res.stage_inputs.push_back(
Hans-Kristian Arntzen6e1c3cc2019-01-11 12:56:00 +0100750 { var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100751 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200752 else
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100753 res.stage_inputs.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200754 }
755 // Subpass inputs
756 else if (var.storage == StorageClassUniformConstant && type.image.dim == DimSubpassData)
757 {
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100758 res.subpass_inputs.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200759 }
760 // Outputs
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +0200761 else if (var.storage == StorageClassOutput && interface_variable_exists_in_entry_point(var.self))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200762 {
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100763 if (has_decoration(type.self, DecorationBlock))
764 {
Hans-Kristian Arntzenaa2557c2017-12-05 09:58:12 +0100765 res.stage_outputs.push_back(
Hans-Kristian Arntzen6e1c3cc2019-01-11 12:56:00 +0100766 { var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100767 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200768 else
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100769 res.stage_outputs.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200770 }
771 // UBOs
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100772 else if (type.storage == StorageClassUniform && has_decoration(type.self, DecorationBlock))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200773 {
Hans-Kristian Arntzenaa2557c2017-12-05 09:58:12 +0100774 res.uniform_buffers.push_back(
775 { var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200776 }
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +0200777 // Old way to declare SSBOs.
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100778 else if (type.storage == StorageClassUniform && has_decoration(type.self, DecorationBufferBlock))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200779 {
Hans-Kristian Arntzenaa2557c2017-12-05 09:58:12 +0100780 res.storage_buffers.push_back(
781 { var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200782 }
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +0200783 // Modern way to declare SSBOs.
784 else if (type.storage == StorageClassStorageBuffer)
785 {
Hans-Kristian Arntzenaa2557c2017-12-05 09:58:12 +0100786 res.storage_buffers.push_back(
787 { var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +0200788 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200789 // Push constant blocks
790 else if (type.storage == StorageClassPushConstant)
791 {
792 // There can only be one push constant block, but keep the vector in case this restriction is lifted
793 // in the future.
Hans-Kristian Arntzen6e1c3cc2019-01-11 12:56:00 +0100794 res.push_constant_buffers.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200795 }
796 // Images
Hans-Kristian Arntzene9202082016-09-10 13:05:35 +0200797 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Image &&
798 type.image.sampled == 2)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200799 {
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100800 res.storage_images.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200801 }
Hans-Kristian Arntzene9202082016-09-10 13:05:35 +0200802 // Separate images
803 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Image &&
804 type.image.sampled == 1)
805 {
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100806 res.separate_images.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
Hans-Kristian Arntzene9202082016-09-10 13:05:35 +0200807 }
808 // Separate samplers
809 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Sampler)
810 {
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100811 res.separate_samplers.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
Hans-Kristian Arntzene9202082016-09-10 13:05:35 +0200812 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200813 // Textures
814 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::SampledImage)
815 {
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100816 res.sampled_images.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200817 }
818 // Atomic counters
819 else if (type.storage == StorageClassAtomicCounter)
820 {
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +0100821 res.atomic_counters.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200822 }
Patrick Moursb2a66752019-03-26 15:02:00 +0100823 // Acceleration structures
824 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::AccelerationStructureNV)
825 {
826 res.acceleration_structures.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
827 }
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +0100828 });
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100829
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200830 return res;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100831}
832
Hans-Kristian Arntzen294259e2018-03-05 16:27:04 +0100833bool Compiler::type_is_block_like(const SPIRType &type) const
834{
835 if (type.basetype != SPIRType::Struct)
836 return false;
837
Hans-Kristian Arntzenac0e93f2018-03-07 10:29:20 +0100838 if (has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock))
Hans-Kristian Arntzen294259e2018-03-05 16:27:04 +0100839 {
840 return true;
841 }
842
843 // Block-like types may have Offset decorations.
844 for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++)
845 if (has_member_decoration(type.self, i, DecorationOffset))
846 return true;
847
848 return false;
849}
850
851void Compiler::fixup_type_alias()
852{
853 // Due to how some backends work, the "master" type of type_alias must be a block-like type if it exists.
854 // FIXME: Multiple alias types which are both block-like will be awkward, for now, it's best to just drop the type
855 // alias if the slave type is a block type.
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +0100856 ir.for_each_typed_id<SPIRType>([&](uint32_t self, SPIRType &type) {
Hans-Kristian Arntzen294259e2018-03-05 16:27:04 +0100857 if (type.type_alias && type_is_block_like(type))
858 {
859 // Become the master.
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +0100860 ir.for_each_typed_id<SPIRType>([&](uint32_t other_id, SPIRType &other_type) {
861 if (other_id == type.self)
862 return;
Hans-Kristian Arntzen294259e2018-03-05 16:27:04 +0100863
Hans-Kristian Arntzen294259e2018-03-05 16:27:04 +0100864 if (other_type.type_alias == type.type_alias)
865 other_type.type_alias = type.self;
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +0100866 });
Hans-Kristian Arntzen294259e2018-03-05 16:27:04 +0100867
Hans-Kristian Arntzen2fb9aa22019-01-11 09:29:28 +0100868 this->get<SPIRType>(type.type_alias).type_alias = self;
Hans-Kristian Arntzen294259e2018-03-05 16:27:04 +0100869 type.type_alias = 0;
870 }
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +0100871 });
Hans-Kristian Arntzen294259e2018-03-05 16:27:04 +0100872
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +0100873 ir.for_each_typed_id<SPIRType>([&](uint32_t, SPIRType &type) {
Hans-Kristian Arntzen294259e2018-03-05 16:27:04 +0100874 if (type.type_alias && type_is_block_like(type))
875 {
876 // This is not allowed, drop the type_alias.
877 type.type_alias = 0;
878 }
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +0100879 });
880
881 // Reorder declaration of types so that the master of the type alias is always emitted first.
882 // We need this in case a type B depends on type A (A must come before in the vector), but A is an alias of a type Abuffer, which
883 // means declaration of A doesn't happen (yet), and order would be B, ABuffer and not ABuffer, B. Fix this up here.
884 auto &type_ids = ir.ids_for_type[TypeType];
885 for (auto alias_itr = begin(type_ids); alias_itr != end(type_ids); ++alias_itr)
886 {
887 auto &type = get<SPIRType>(*alias_itr);
Hans-Kristian Arntzende7e5cc2019-01-17 11:22:24 +0100888 if (type.type_alias != 0 && !has_extended_decoration(type.type_alias, SPIRVCrossDecorationPacked))
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +0100889 {
890 // We will skip declaring this type, so make sure the type_alias type comes before.
891 auto master_itr = find(begin(type_ids), end(type_ids), type.type_alias);
892 assert(master_itr != end(type_ids));
893
894 if (alias_itr < master_itr)
895 {
896 // Must also swap the type order for the constant-type joined array.
897 auto &joined_types = ir.ids_for_constant_or_type;
898 auto alt_alias_itr = find(begin(joined_types), end(joined_types), *alias_itr);
899 auto alt_master_itr = find(begin(joined_types), end(joined_types), *master_itr);
900 assert(alt_alias_itr != end(joined_types));
901 assert(alt_master_itr != end(joined_types));
902
903 swap(*alias_itr, *master_itr);
904 swap(*alt_alias_itr, *alt_master_itr);
905 }
906 }
Hans-Kristian Arntzen294259e2018-03-05 16:27:04 +0100907 }
908}
909
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200910void Compiler::parse_fixup()
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100911{
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +0200912 // Figure out specialization constants for work group sizes.
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +0100913 for (auto id_ : ir.ids_for_constant_or_variable)
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +0200914 {
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +0100915 auto &id = ir.ids[id_];
916
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +0200917 if (id.get_type() == TypeConstant)
918 {
919 auto &c = id.get<SPIRConstant>();
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200920 if (ir.meta[c.self].decoration.builtin && ir.meta[c.self].decoration.builtin_type == BuiltInWorkgroupSize)
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +0200921 {
922 // In current SPIR-V, there can be just one constant like this.
923 // All entry points will receive the constant value.
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200924 for (auto &entry : ir.entry_points)
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +0200925 {
926 entry.second.workgroup_size.constant = c.self;
927 entry.second.workgroup_size.x = c.scalar(0, 0);
928 entry.second.workgroup_size.y = c.scalar(0, 1);
929 entry.second.workgroup_size.z = c.scalar(0, 2);
930 }
931 }
932 }
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +0200933 else if (id.get_type() == TypeVariable)
934 {
935 auto &var = id.get<SPIRVariable>();
936 if (var.storage == StorageClassPrivate || var.storage == StorageClassWorkgroup ||
937 var.storage == StorageClassOutput)
938 global_variables.push_back(var.self);
939 if (variable_storage_is_aliased(var))
940 aliased_variables.push_back(var.self);
941 }
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +0200942 }
Hans-Kristian Arntzen294259e2018-03-05 16:27:04 +0100943
944 fixup_type_alias();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100945}
946
Hans-Kristian Arntzen5b876222019-01-07 10:01:28 +0100947void Compiler::update_name_cache(unordered_set<string> &cache_primary, const unordered_set<string> &cache_secondary,
948 string &name)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100949{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200950 if (name.empty())
951 return;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100952
Hans-Kristian Arntzen9728f9c2019-01-04 12:15:43 +0100953 const auto find_name = [&](const string &n) -> bool {
954 if (cache_primary.find(n) != end(cache_primary))
955 return true;
956
957 if (&cache_primary != &cache_secondary)
958 if (cache_secondary.find(n) != end(cache_secondary))
959 return true;
960
961 return false;
962 };
963
Hans-Kristian Arntzen5b876222019-01-07 10:01:28 +0100964 const auto insert_name = [&](const string &n) { cache_primary.insert(n); };
Hans-Kristian Arntzen9728f9c2019-01-04 12:15:43 +0100965
966 if (!find_name(name))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200967 {
Hans-Kristian Arntzen9728f9c2019-01-04 12:15:43 +0100968 insert_name(name);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200969 return;
970 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100971
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200972 uint32_t counter = 0;
973 auto tmpname = name;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100974
Hans-Kristian Arntzen6bcc8902018-06-04 10:13:57 +0200975 bool use_linked_underscore = true;
976
977 if (tmpname == "_")
978 {
979 // We cannot just append numbers, as we will end up creating internally reserved names.
980 // Make it like _0_<counter> instead.
981 tmpname += "0";
982 }
983 else if (tmpname.back() == '_')
984 {
985 // The last_character is an underscore, so we don't need to link in underscore.
986 // This would violate double underscore rules.
987 use_linked_underscore = false;
988 }
989
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200990 // If there is a collision (very rare),
991 // keep tacking on extra identifier until it's unique.
992 do
993 {
994 counter++;
Hans-Kristian Arntzen6bcc8902018-06-04 10:13:57 +0200995 name = tmpname + (use_linked_underscore ? "_" : "") + convert_to_string(counter);
Hans-Kristian Arntzen9728f9c2019-01-04 12:15:43 +0100996 } while (find_name(name));
997 insert_name(name);
998}
999
1000void Compiler::update_name_cache(unordered_set<string> &cache, string &name)
1001{
1002 update_name_cache(cache, cache, name);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001003}
1004
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001005void Compiler::set_name(uint32_t id, const std::string &name)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001006{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001007 ir.set_name(id, name);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001008}
1009
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001010const SPIRType &Compiler::get_type(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001011{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001012 return get<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001013}
1014
Robert Konrad451bdee2016-09-24 22:17:01 +02001015const SPIRType &Compiler::get_type_from_variable(uint32_t id) const
1016{
1017 return get<SPIRType>(get<SPIRVariable>(id).basetype);
1018}
1019
Chip Davisfc02b3d2019-01-08 12:54:40 -06001020uint32_t Compiler::get_pointee_type_id(uint32_t type_id) const
Bill Hollingse0910312018-06-24 15:06:12 -04001021{
1022 auto *p_type = &get<SPIRType>(type_id);
Chip Davis3bfb2f92018-12-03 02:06:33 -06001023 if (p_type->pointer)
Bill Hollingse0910312018-06-24 15:06:12 -04001024 {
1025 assert(p_type->parent_type);
1026 type_id = p_type->parent_type;
Bill Hollingse0910312018-06-24 15:06:12 -04001027 }
1028 return type_id;
1029}
1030
Chip Davisfc02b3d2019-01-08 12:54:40 -06001031const SPIRType &Compiler::get_pointee_type(const SPIRType &type) const
Bill Hollingse0910312018-06-24 15:06:12 -04001032{
1033 auto *p_type = &type;
Chip Davis3bfb2f92018-12-03 02:06:33 -06001034 if (p_type->pointer)
Bill Hollingse0910312018-06-24 15:06:12 -04001035 {
1036 assert(p_type->parent_type);
1037 p_type = &get<SPIRType>(p_type->parent_type);
1038 }
1039 return *p_type;
1040}
1041
Chip Davisfc02b3d2019-01-08 12:54:40 -06001042const SPIRType &Compiler::get_pointee_type(uint32_t type_id) const
Bill Hollingse0910312018-06-24 15:06:12 -04001043{
Chip Davisfc02b3d2019-01-08 12:54:40 -06001044 return get_pointee_type(get<SPIRType>(type_id));
Bill Hollingse0910312018-06-24 15:06:12 -04001045}
1046
Chip Davis3bfb2f92018-12-03 02:06:33 -06001047uint32_t Compiler::get_variable_data_type_id(const SPIRVariable &var) const
1048{
1049 if (var.phi_variable)
1050 return var.basetype;
Chip Davisfc02b3d2019-01-08 12:54:40 -06001051 return get_pointee_type_id(var.basetype);
Chip Davis3bfb2f92018-12-03 02:06:33 -06001052}
1053
1054SPIRType &Compiler::get_variable_data_type(const SPIRVariable &var)
1055{
1056 return get<SPIRType>(get_variable_data_type_id(var));
1057}
1058
1059const SPIRType &Compiler::get_variable_data_type(const SPIRVariable &var) const
1060{
1061 return get<SPIRType>(get_variable_data_type_id(var));
1062}
1063
Chip Daviseb89c3a2019-02-03 23:58:46 -06001064SPIRType &Compiler::get_variable_element_type(const SPIRVariable &var)
1065{
1066 SPIRType *type = &get_variable_data_type(var);
1067 if (is_array(*type))
1068 type = &get<SPIRType>(type->parent_type);
1069 return *type;
1070}
1071
1072const SPIRType &Compiler::get_variable_element_type(const SPIRVariable &var) const
1073{
1074 const SPIRType *type = &get_variable_data_type(var);
1075 if (is_array(*type))
1076 type = &get<SPIRType>(type->parent_type);
1077 return *type;
1078}
1079
Chip Davis8855ea02018-09-24 12:24:58 -05001080bool Compiler::is_sampled_image_type(const SPIRType &type)
1081{
Hans-Kristian Arntzenc07c3032018-09-27 13:36:38 +02001082 return (type.basetype == SPIRType::Image || type.basetype == SPIRType::SampledImage) && type.image.sampled == 1 &&
1083 type.image.dim != DimBuffer;
Chip Davis8855ea02018-09-24 12:24:58 -05001084}
1085
Hans-Kristian Arntzen215d3ca2018-03-20 20:04:12 +01001086void Compiler::set_member_decoration_string(uint32_t id, uint32_t index, spv::Decoration decoration,
1087 const std::string &argument)
1088{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001089 ir.set_member_decoration_string(id, index, decoration, argument);
Hans-Kristian Arntzen215d3ca2018-03-20 20:04:12 +01001090}
1091
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001092void Compiler::set_member_decoration(uint32_t id, uint32_t index, Decoration decoration, uint32_t argument)
1093{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001094 ir.set_member_decoration(id, index, decoration, argument);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001095}
1096
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001097void Compiler::set_member_name(uint32_t id, uint32_t index, const std::string &name)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001098{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001099 ir.set_member_name(id, index, name);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001100}
1101
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001102const std::string &Compiler::get_member_name(uint32_t id, uint32_t index) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001103{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001104 return ir.get_member_name(id, index);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001105}
1106
Hans-Kristian Arntzene47a77d2019-03-14 10:29:34 +01001107void Compiler::set_qualified_name(uint32_t id, const string &name)
1108{
1109 ir.meta[id].decoration.qualified_alias = name;
1110}
1111
Bill Hollings1e84a372017-08-12 00:21:13 -04001112void Compiler::set_member_qualified_name(uint32_t type_id, uint32_t index, const std::string &name)
Bill Hollingsc45e74f2016-07-08 12:39:22 -04001113{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01001114 ir.meta[type_id].members.resize(max(ir.meta[type_id].members.size(), size_t(index) + 1));
1115 ir.meta[type_id].members[index].qualified_alias = name;
Bill Hollingsc1b81542017-05-22 21:41:19 -04001116}
1117
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01001118const string &Compiler::get_member_qualified_name(uint32_t type_id, uint32_t index) const
Bill Hollingsc1b81542017-05-22 21:41:19 -04001119{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01001120 auto *m = ir.find_meta(type_id);
1121 if (m && index < m->members.size())
1122 return m->members[index].qualified_alias;
Bill Hollings1e84a372017-08-12 00:21:13 -04001123 else
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01001124 return ir.get_empty_string();
Bill Hollingsc45e74f2016-07-08 12:39:22 -04001125}
1126
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001127uint32_t Compiler::get_member_decoration(uint32_t id, uint32_t index, Decoration decoration) const
1128{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001129 return ir.get_member_decoration(id, index, decoration);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001130}
1131
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01001132const Bitset &Compiler::get_member_decoration_bitset(uint32_t id, uint32_t index) const
1133{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001134 return ir.get_member_decoration_bitset(id, index);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001135}
1136
Bill Hollingsdc694272017-03-11 12:17:22 -05001137bool Compiler::has_member_decoration(uint32_t id, uint32_t index, Decoration decoration) const
Bill Hollings484931d2017-02-28 21:44:36 -05001138{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001139 return ir.has_member_decoration(id, index, decoration);
Bill Hollings484931d2017-02-28 21:44:36 -05001140}
1141
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001142void Compiler::unset_member_decoration(uint32_t id, uint32_t index, Decoration decoration)
1143{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001144 ir.unset_member_decoration(id, index, decoration);
Hans-Kristian Arntzen215d3ca2018-03-20 20:04:12 +01001145}
1146
1147void Compiler::set_decoration_string(uint32_t id, spv::Decoration decoration, const std::string &argument)
1148{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001149 ir.set_decoration_string(id, decoration, argument);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001150}
1151
1152void Compiler::set_decoration(uint32_t id, Decoration decoration, uint32_t argument)
1153{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001154 ir.set_decoration(id, decoration, argument);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001155}
1156
Hans-Kristian Arntzende7e5cc2019-01-17 11:22:24 +01001157void Compiler::set_extended_decoration(uint32_t id, ExtendedDecorations decoration, uint32_t value)
1158{
1159 auto &dec = ir.meta[id].decoration;
1160 switch (decoration)
1161 {
1162 case SPIRVCrossDecorationPacked:
1163 dec.extended.packed = true;
1164 break;
1165
1166 case SPIRVCrossDecorationPackedType:
1167 dec.extended.packed_type = value;
1168 break;
Chip Daviseb89c3a2019-02-03 23:58:46 -06001169
1170 case SPIRVCrossDecorationInterfaceMemberIndex:
1171 dec.extended.ib_member_index = value;
1172 break;
1173
1174 case SPIRVCrossDecorationInterfaceOrigID:
1175 dec.extended.ib_orig_id = value;
1176 break;
Hans-Kristian Arntzene47a77d2019-03-14 10:29:34 +01001177
1178 case SPIRVCrossDecorationArgumentBufferID:
1179 dec.extended.argument_buffer_id = value;
1180 break;
Hans-Kristian Arntzende7e5cc2019-01-17 11:22:24 +01001181 }
1182}
1183
Hans-Kristian Arntzen40e77232019-01-17 11:29:50 +01001184void Compiler::set_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration,
1185 uint32_t value)
Hans-Kristian Arntzende7e5cc2019-01-17 11:22:24 +01001186{
1187 ir.meta[type].members.resize(max(ir.meta[type].members.size(), size_t(index) + 1));
1188 auto &dec = ir.meta[type].members[index];
1189
1190 switch (decoration)
1191 {
1192 case SPIRVCrossDecorationPacked:
1193 dec.extended.packed = true;
1194 break;
1195
1196 case SPIRVCrossDecorationPackedType:
1197 dec.extended.packed_type = value;
1198 break;
Chip Daviseb89c3a2019-02-03 23:58:46 -06001199
1200 case SPIRVCrossDecorationInterfaceMemberIndex:
1201 dec.extended.ib_member_index = value;
1202 break;
1203
1204 case SPIRVCrossDecorationInterfaceOrigID:
1205 dec.extended.ib_orig_id = value;
1206 break;
Hans-Kristian Arntzene47a77d2019-03-14 10:29:34 +01001207
1208 case SPIRVCrossDecorationArgumentBufferID:
1209 dec.extended.argument_buffer_id = value;
1210 break;
Hans-Kristian Arntzende7e5cc2019-01-17 11:22:24 +01001211 }
1212}
1213
1214uint32_t Compiler::get_extended_decoration(uint32_t id, ExtendedDecorations decoration) const
1215{
1216 auto *m = ir.find_meta(id);
1217 if (!m)
1218 return 0;
1219
1220 auto &dec = m->decoration;
1221 switch (decoration)
1222 {
1223 case SPIRVCrossDecorationPacked:
1224 return uint32_t(dec.extended.packed);
1225
1226 case SPIRVCrossDecorationPackedType:
1227 return dec.extended.packed_type;
Chip Daviseb89c3a2019-02-03 23:58:46 -06001228
1229 case SPIRVCrossDecorationInterfaceMemberIndex:
1230 return dec.extended.ib_member_index;
1231
1232 case SPIRVCrossDecorationInterfaceOrigID:
1233 return dec.extended.ib_orig_id;
Hans-Kristian Arntzene47a77d2019-03-14 10:29:34 +01001234
1235 case SPIRVCrossDecorationArgumentBufferID:
1236 return dec.extended.argument_buffer_id;
Hans-Kristian Arntzende7e5cc2019-01-17 11:22:24 +01001237 }
1238
1239 return 0;
1240}
1241
1242uint32_t Compiler::get_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const
1243{
1244 auto *m = ir.find_meta(type);
1245 if (!m)
1246 return 0;
1247
1248 if (index >= m->members.size())
1249 return 0;
1250
1251 auto &dec = m->members[index];
1252 switch (decoration)
1253 {
1254 case SPIRVCrossDecorationPacked:
1255 return uint32_t(dec.extended.packed);
1256
1257 case SPIRVCrossDecorationPackedType:
1258 return dec.extended.packed_type;
Chip Daviseb89c3a2019-02-03 23:58:46 -06001259
1260 case SPIRVCrossDecorationInterfaceMemberIndex:
1261 return dec.extended.ib_member_index;
1262
1263 case SPIRVCrossDecorationInterfaceOrigID:
1264 return dec.extended.ib_orig_id;
Hans-Kristian Arntzene47a77d2019-03-14 10:29:34 +01001265
1266 case SPIRVCrossDecorationArgumentBufferID:
1267 return dec.extended.argument_buffer_id;
Hans-Kristian Arntzende7e5cc2019-01-17 11:22:24 +01001268 }
1269
1270 return 0;
1271}
1272
1273bool Compiler::has_extended_decoration(uint32_t id, ExtendedDecorations decoration) const
1274{
1275 auto *m = ir.find_meta(id);
1276 if (!m)
1277 return false;
1278
1279 auto &dec = m->decoration;
1280 switch (decoration)
1281 {
1282 case SPIRVCrossDecorationPacked:
1283 return dec.extended.packed;
1284
1285 case SPIRVCrossDecorationPackedType:
1286 return dec.extended.packed_type != 0;
Chip Daviseb89c3a2019-02-03 23:58:46 -06001287
1288 case SPIRVCrossDecorationInterfaceMemberIndex:
Chip Davis8860a972019-02-11 16:14:00 -06001289 return dec.extended.ib_member_index != uint32_t(-1);
Chip Daviseb89c3a2019-02-03 23:58:46 -06001290
1291 case SPIRVCrossDecorationInterfaceOrigID:
1292 return dec.extended.ib_orig_id != 0;
Hans-Kristian Arntzene47a77d2019-03-14 10:29:34 +01001293
1294 case SPIRVCrossDecorationArgumentBufferID:
1295 return dec.extended.argument_buffer_id != 0;
Hans-Kristian Arntzende7e5cc2019-01-17 11:22:24 +01001296 }
1297
1298 return false;
1299}
1300
1301bool Compiler::has_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const
1302{
1303 auto *m = ir.find_meta(type);
1304 if (!m)
1305 return false;
1306
1307 if (index >= m->members.size())
1308 return false;
1309
1310 auto &dec = m->members[index];
1311 switch (decoration)
1312 {
1313 case SPIRVCrossDecorationPacked:
1314 return dec.extended.packed;
1315
1316 case SPIRVCrossDecorationPackedType:
1317 return dec.extended.packed_type != 0;
Chip Daviseb89c3a2019-02-03 23:58:46 -06001318
1319 case SPIRVCrossDecorationInterfaceMemberIndex:
Chip Davis8860a972019-02-11 16:14:00 -06001320 return dec.extended.ib_member_index != uint32_t(-1);
Chip Daviseb89c3a2019-02-03 23:58:46 -06001321
1322 case SPIRVCrossDecorationInterfaceOrigID:
1323 return dec.extended.ib_orig_id != 0;
Hans-Kristian Arntzene47a77d2019-03-14 10:29:34 +01001324
1325 case SPIRVCrossDecorationArgumentBufferID:
1326 return dec.extended.argument_buffer_id != uint32_t(-1);
Hans-Kristian Arntzende7e5cc2019-01-17 11:22:24 +01001327 }
1328
1329 return false;
1330}
1331
1332void Compiler::unset_extended_decoration(uint32_t id, ExtendedDecorations decoration)
1333{
1334 auto &dec = ir.meta[id].decoration;
1335 switch (decoration)
1336 {
1337 case SPIRVCrossDecorationPacked:
1338 dec.extended.packed = false;
1339 break;
1340
1341 case SPIRVCrossDecorationPackedType:
1342 dec.extended.packed_type = 0;
1343 break;
Chip Daviseb89c3a2019-02-03 23:58:46 -06001344
1345 case SPIRVCrossDecorationInterfaceMemberIndex:
Hans-Kristian Arntzene47a77d2019-03-14 10:29:34 +01001346 dec.extended.ib_member_index = ~(0u);
Chip Daviseb89c3a2019-02-03 23:58:46 -06001347 break;
1348
1349 case SPIRVCrossDecorationInterfaceOrigID:
1350 dec.extended.ib_orig_id = 0;
1351 break;
Hans-Kristian Arntzene47a77d2019-03-14 10:29:34 +01001352
1353 case SPIRVCrossDecorationArgumentBufferID:
1354 dec.extended.argument_buffer_id = 0;
1355 break;
Hans-Kristian Arntzende7e5cc2019-01-17 11:22:24 +01001356 }
1357}
1358
1359void Compiler::unset_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration)
1360{
1361 ir.meta[type].members.resize(max(ir.meta[type].members.size(), size_t(index) + 1));
1362 auto &dec = ir.meta[type].members[index];
1363
1364 switch (decoration)
1365 {
1366 case SPIRVCrossDecorationPacked:
1367 dec.extended.packed = false;
1368 break;
1369
1370 case SPIRVCrossDecorationPackedType:
1371 dec.extended.packed_type = 0;
1372 break;
Chip Daviseb89c3a2019-02-03 23:58:46 -06001373
1374 case SPIRVCrossDecorationInterfaceMemberIndex:
Hans-Kristian Arntzene47a77d2019-03-14 10:29:34 +01001375 dec.extended.ib_member_index = ~(0u);
Chip Daviseb89c3a2019-02-03 23:58:46 -06001376 break;
1377
1378 case SPIRVCrossDecorationInterfaceOrigID:
1379 dec.extended.ib_orig_id = 0;
1380 break;
Hans-Kristian Arntzene47a77d2019-03-14 10:29:34 +01001381
1382 case SPIRVCrossDecorationArgumentBufferID:
1383 dec.extended.argument_buffer_id = 0;
1384 break;
Hans-Kristian Arntzende7e5cc2019-01-17 11:22:24 +01001385 }
1386}
1387
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001388StorageClass Compiler::get_storage_class(uint32_t id) const
1389{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001390 return get<SPIRVariable>(id).storage;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001391}
1392
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001393const std::string &Compiler::get_name(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001394{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01001395 return ir.get_name(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001396}
1397
Hans-Kristian Arntzenaab31072017-09-29 12:16:53 +02001398const std::string Compiler::get_fallback_name(uint32_t id) const
1399{
1400 return join("_", id);
1401}
1402
1403const std::string Compiler::get_block_fallback_name(uint32_t id) const
1404{
1405 auto &var = get<SPIRVariable>(id);
Hans-Kristian Arntzen2c90ea32017-12-01 14:20:51 +01001406 if (get_name(id).empty())
1407 return join("_", get<SPIRType>(var.basetype).self, "_", id);
1408 else
1409 return get_name(id);
Hans-Kristian Arntzenaab31072017-09-29 12:16:53 +02001410}
1411
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01001412const Bitset &Compiler::get_decoration_bitset(uint32_t id) const
1413{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001414 return ir.get_decoration_bitset(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001415}
1416
Bill Hollingsdc694272017-03-11 12:17:22 -05001417bool Compiler::has_decoration(uint32_t id, Decoration decoration) const
Bill Hollings484931d2017-02-28 21:44:36 -05001418{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001419 return ir.has_decoration(id, decoration);
Bill Hollings484931d2017-02-28 21:44:36 -05001420}
1421
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001422const string &Compiler::get_decoration_string(uint32_t id, Decoration decoration) const
Hans-Kristian Arntzen215d3ca2018-03-20 20:04:12 +01001423{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001424 return ir.get_decoration_string(id, decoration);
1425}
Hans-Kristian Arntzen215d3ca2018-03-20 20:04:12 +01001426
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001427const string &Compiler::get_member_decoration_string(uint32_t id, uint32_t index, Decoration decoration) const
1428{
1429 return ir.get_member_decoration_string(id, index, decoration);
Hans-Kristian Arntzen215d3ca2018-03-20 20:04:12 +01001430}
1431
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001432uint32_t Compiler::get_decoration(uint32_t id, Decoration decoration) const
1433{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001434 return ir.get_decoration(id, decoration);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001435}
1436
1437void Compiler::unset_decoration(uint32_t id, Decoration decoration)
1438{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001439 ir.unset_decoration(id, decoration);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001440}
1441
Hans-Kristian Arntzen93fe19c2017-04-25 20:10:24 +02001442bool Compiler::get_binary_offset_for_decoration(uint32_t id, spv::Decoration decoration, uint32_t &word_offset) const
1443{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01001444 auto *m = ir.find_meta(id);
1445 if (!m)
1446 return false;
1447
1448 auto &word_offsets = m->decoration_word_offset;
Hans-Kristian Arntzen93fe19c2017-04-25 20:10:24 +02001449 auto itr = word_offsets.find(decoration);
1450 if (itr == end(word_offsets))
1451 return false;
1452
1453 word_offset = itr->second;
1454 return true;
1455}
1456
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001457bool Compiler::block_is_loop_candidate(const SPIRBlock &block, SPIRBlock::Method method) const
1458{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001459 // Tried and failed.
1460 if (block.disable_block_optimization || block.complex_continue)
1461 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001462
Hans-Kristian Arntzen28cccc32018-03-08 17:51:55 +01001463 if (method == SPIRBlock::MergeToSelectForLoop || method == SPIRBlock::MergeToSelectContinueForLoop)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001464 {
1465 // Try to detect common for loop pattern
1466 // which the code backend can use to create cleaner code.
1467 // for(;;) { if (cond) { some_body; } else { break; } }
1468 // is the pattern we're looking for.
Hans-Kristian Arntzen70ff96b2019-03-06 11:24:43 +01001469 const auto *false_block = maybe_get<SPIRBlock>(block.false_block);
1470 const auto *true_block = maybe_get<SPIRBlock>(block.true_block);
1471 const auto *merge_block = maybe_get<SPIRBlock>(block.merge_block);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001472
Hans-Kristian Arntzen8bfb04d2019-03-06 12:20:13 +01001473 bool false_block_is_merge = block.false_block == block.merge_block ||
1474 (false_block && merge_block && execution_is_noop(*false_block, *merge_block));
Hans-Kristian Arntzen70ff96b2019-03-06 11:24:43 +01001475
Hans-Kristian Arntzen8bfb04d2019-03-06 12:20:13 +01001476 bool true_block_is_merge = block.true_block == block.merge_block ||
1477 (true_block && merge_block && execution_is_noop(*true_block, *merge_block));
Hans-Kristian Arntzen70ff96b2019-03-06 11:24:43 +01001478
Hans-Kristian Arntzen8bfb04d2019-03-06 12:20:13 +01001479 bool positive_candidate =
1480 block.true_block != block.merge_block && block.true_block != block.self && false_block_is_merge;
Hans-Kristian Arntzen70ff96b2019-03-06 11:24:43 +01001481
Hans-Kristian Arntzen8bfb04d2019-03-06 12:20:13 +01001482 bool negative_candidate =
1483 block.false_block != block.merge_block && block.false_block != block.self && true_block_is_merge;
Hans-Kristian Arntzen70ff96b2019-03-06 11:24:43 +01001484
1485 bool ret = block.terminator == SPIRBlock::Select && block.merge == SPIRBlock::MergeLoop &&
1486 (positive_candidate || negative_candidate);
1487
1488 if (ret && positive_candidate && method == SPIRBlock::MergeToSelectContinueForLoop)
Hans-Kristian Arntzen28cccc32018-03-08 17:51:55 +01001489 ret = block.true_block == block.continue_block;
Hans-Kristian Arntzen70ff96b2019-03-06 11:24:43 +01001490 else if (ret && negative_candidate && method == SPIRBlock::MergeToSelectContinueForLoop)
1491 ret = block.false_block == block.continue_block;
Hans-Kristian Arntzen28cccc32018-03-08 17:51:55 +01001492
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001493 // If we have OpPhi which depends on branches which came from our own block,
1494 // we need to flush phi variables in else block instead of a trivial break,
1495 // so we cannot assume this is a for loop candidate.
1496 if (ret)
1497 {
1498 for (auto &phi : block.phi_variables)
1499 if (phi.parent == block.self)
1500 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001501
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001502 auto *merge = maybe_get<SPIRBlock>(block.merge_block);
1503 if (merge)
1504 for (auto &phi : merge->phi_variables)
1505 if (phi.parent == block.self)
1506 return false;
1507 }
1508 return ret;
1509 }
1510 else if (method == SPIRBlock::MergeToDirectForLoop)
1511 {
1512 // Empty loop header that just sets up merge target
1513 // and branches to loop body.
1514 bool ret = block.terminator == SPIRBlock::Direct && block.merge == SPIRBlock::MergeLoop && block.ops.empty();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001515
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001516 if (!ret)
1517 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001518
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001519 auto &child = get<SPIRBlock>(block.next_block);
Hans-Kristian Arntzen70ff96b2019-03-06 11:24:43 +01001520
1521 const auto *false_block = maybe_get<SPIRBlock>(child.false_block);
1522 const auto *true_block = maybe_get<SPIRBlock>(child.true_block);
1523 const auto *merge_block = maybe_get<SPIRBlock>(block.merge_block);
1524
Hans-Kristian Arntzen8bfb04d2019-03-06 12:20:13 +01001525 bool false_block_is_merge = child.false_block == block.merge_block ||
1526 (false_block && merge_block && execution_is_noop(*false_block, *merge_block));
Hans-Kristian Arntzen70ff96b2019-03-06 11:24:43 +01001527
Hans-Kristian Arntzen8bfb04d2019-03-06 12:20:13 +01001528 bool true_block_is_merge = child.true_block == block.merge_block ||
1529 (true_block && merge_block && execution_is_noop(*true_block, *merge_block));
Hans-Kristian Arntzen70ff96b2019-03-06 11:24:43 +01001530
Hans-Kristian Arntzen8bfb04d2019-03-06 12:20:13 +01001531 bool positive_candidate =
1532 child.true_block != block.merge_block && child.true_block != block.self && false_block_is_merge;
Hans-Kristian Arntzen70ff96b2019-03-06 11:24:43 +01001533
Hans-Kristian Arntzen8bfb04d2019-03-06 12:20:13 +01001534 bool negative_candidate =
1535 child.false_block != block.merge_block && child.false_block != block.self && true_block_is_merge;
Hans-Kristian Arntzen70ff96b2019-03-06 11:24:43 +01001536
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001537 ret = child.terminator == SPIRBlock::Select && child.merge == SPIRBlock::MergeNone &&
Hans-Kristian Arntzen70ff96b2019-03-06 11:24:43 +01001538 (positive_candidate || negative_candidate);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001539
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001540 // If we have OpPhi which depends on branches which came from our own block,
1541 // we need to flush phi variables in else block instead of a trivial break,
1542 // so we cannot assume this is a for loop candidate.
1543 if (ret)
1544 {
1545 for (auto &phi : block.phi_variables)
1546 if (phi.parent == block.self || phi.parent == child.self)
1547 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001548
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001549 for (auto &phi : child.phi_variables)
1550 if (phi.parent == block.self)
1551 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001552
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001553 auto *merge = maybe_get<SPIRBlock>(block.merge_block);
1554 if (merge)
1555 for (auto &phi : merge->phi_variables)
1556 if (phi.parent == block.self || phi.parent == child.false_block)
1557 return false;
1558 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001559
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001560 return ret;
1561 }
1562 else
1563 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001564}
1565
1566bool Compiler::block_is_outside_flow_control_from_block(const SPIRBlock &from, const SPIRBlock &to)
1567{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001568 auto *start = &from;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001569
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001570 if (start->self == to.self)
1571 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001572
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001573 // Break cycles.
1574 if (is_continue(start->self))
1575 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001576
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001577 // If our select block doesn't merge, we must break or continue in these blocks,
1578 // so if continues occur branchless within these blocks, consider them branchless as well.
1579 // This is typically used for loop control.
1580 if (start->terminator == SPIRBlock::Select && start->merge == SPIRBlock::MergeNone &&
1581 (block_is_outside_flow_control_from_block(get<SPIRBlock>(start->true_block), to) ||
1582 block_is_outside_flow_control_from_block(get<SPIRBlock>(start->false_block), to)))
1583 {
1584 return true;
1585 }
1586 else if (start->merge_block && block_is_outside_flow_control_from_block(get<SPIRBlock>(start->merge_block), to))
1587 {
1588 return true;
1589 }
1590 else if (start->next_block && block_is_outside_flow_control_from_block(get<SPIRBlock>(start->next_block), to))
1591 {
1592 return true;
1593 }
1594 else
1595 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001596}
1597
1598bool Compiler::execution_is_noop(const SPIRBlock &from, const SPIRBlock &to) const
1599{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001600 if (!execution_is_branchless(from, to))
1601 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001602
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001603 auto *start = &from;
1604 for (;;)
1605 {
1606 if (start->self == to.self)
1607 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001608
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001609 if (!start->ops.empty())
1610 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001611
Hans-Kristian Arntzen40bb42f2018-08-06 13:28:01 +02001612 auto &next = get<SPIRBlock>(start->next_block);
1613 // Flushing phi variables does not count as noop.
1614 for (auto &phi : next.phi_variables)
1615 if (phi.parent == start->self)
1616 return false;
1617
1618 start = &next;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001619 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001620}
1621
1622bool Compiler::execution_is_branchless(const SPIRBlock &from, const SPIRBlock &to) const
1623{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001624 auto *start = &from;
1625 for (;;)
1626 {
1627 if (start->self == to.self)
1628 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001629
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001630 if (start->terminator == SPIRBlock::Direct && start->merge == SPIRBlock::MergeNone)
1631 start = &get<SPIRBlock>(start->next_block);
1632 else
1633 return false;
1634 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001635}
1636
Hans-Kristian Arntzen926916d2016-05-05 09:15:25 +02001637SPIRBlock::ContinueBlockType Compiler::continue_block_type(const SPIRBlock &block) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001638{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001639 // The block was deemed too complex during code emit, pick conservative fallback paths.
1640 if (block.complex_continue)
1641 return SPIRBlock::ComplexLoop;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001642
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001643 // In older glslang output continue block can be equal to the loop header.
1644 // In this case, execution is clearly branchless, so just assume a while loop header here.
1645 if (block.merge == SPIRBlock::MergeLoop)
1646 return SPIRBlock::WhileLoop;
Hans-Kristian Arntzen97f81ba2016-04-01 12:37:29 +02001647
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001648 auto &dominator = get<SPIRBlock>(block.loop_dominator);
Hans-Kristian Arntzen97f81ba2016-04-01 12:37:29 +02001649
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001650 if (execution_is_noop(block, dominator))
1651 return SPIRBlock::WhileLoop;
1652 else if (execution_is_branchless(block, dominator))
1653 return SPIRBlock::ForLoop;
1654 else
1655 {
Hans-Kristian Arntzenef243372019-03-06 12:17:38 +01001656 const auto *false_block = maybe_get<SPIRBlock>(block.false_block);
1657 const auto *true_block = maybe_get<SPIRBlock>(block.true_block);
1658 const auto *merge_block = maybe_get<SPIRBlock>(dominator.merge_block);
1659
1660 bool positive_do_while = block.true_block == dominator.self &&
1661 (block.false_block == dominator.merge_block ||
1662 (false_block && merge_block && execution_is_noop(*false_block, *merge_block)));
1663
1664 bool negative_do_while = block.false_block == dominator.self &&
1665 (block.true_block == dominator.merge_block ||
1666 (true_block && merge_block && execution_is_noop(*true_block, *merge_block)));
1667
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001668 if (block.merge == SPIRBlock::MergeNone && block.terminator == SPIRBlock::Select &&
Hans-Kristian Arntzenef243372019-03-06 12:17:38 +01001669 (positive_do_while || negative_do_while))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001670 {
1671 return SPIRBlock::DoWhileLoop;
1672 }
1673 else
1674 return SPIRBlock::ComplexLoop;
1675 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001676}
1677
1678bool Compiler::traverse_all_reachable_opcodes(const SPIRBlock &block, OpcodeHandler &handler) const
1679{
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01001680 handler.set_current_block(block);
1681
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001682 // Ideally, perhaps traverse the CFG instead of all blocks in order to eliminate dead blocks,
1683 // but this shouldn't be a problem in practice unless the SPIR-V is doing insane things like recursing
1684 // inside dead blocks ...
1685 for (auto &i : block.ops)
1686 {
1687 auto ops = stream(i);
1688 auto op = static_cast<Op>(i.op);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001689
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001690 if (!handler.handle(op, ops, i.length))
1691 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001692
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02001693 if (op == OpFunctionCall)
1694 {
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01001695 auto &func = get<SPIRFunction>(ops[2]);
1696 if (handler.follow_function_call(func))
1697 {
1698 if (!handler.begin_function_scope(ops, i.length))
1699 return false;
1700 if (!traverse_all_reachable_opcodes(get<SPIRFunction>(ops[2]), handler))
1701 return false;
1702 if (!handler.end_function_scope(ops, i.length))
1703 return false;
1704 }
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02001705 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001706 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001707
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001708 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001709}
1710
1711bool Compiler::traverse_all_reachable_opcodes(const SPIRFunction &func, OpcodeHandler &handler) const
1712{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001713 for (auto block : func.blocks)
1714 if (!traverse_all_reachable_opcodes(get<SPIRBlock>(block), handler))
1715 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001716
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001717 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001718}
1719
1720uint32_t Compiler::type_struct_member_offset(const SPIRType &type, uint32_t index) const
1721{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01001722 auto *type_meta = ir.find_meta(type.self);
1723 if (type_meta)
1724 {
1725 // Decoration must be set in valid SPIR-V, otherwise throw.
1726 auto &dec = type_meta->members[index];
1727 if (dec.decoration_flags.get(DecorationOffset))
1728 return dec.offset;
1729 else
1730 SPIRV_CROSS_THROW("Struct member does not have Offset set.");
1731 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001732 else
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001733 SPIRV_CROSS_THROW("Struct member does not have Offset set.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001734}
1735
1736uint32_t Compiler::type_struct_member_array_stride(const SPIRType &type, uint32_t index) const
1737{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01001738 auto *type_meta = ir.find_meta(type.member_types[index]);
1739 if (type_meta)
1740 {
1741 // Decoration must be set in valid SPIR-V, otherwise throw.
1742 // ArrayStride is part of the array type not OpMemberDecorate.
1743 auto &dec = type_meta->decoration;
1744 if (dec.decoration_flags.get(DecorationArrayStride))
1745 return dec.array_stride;
1746 else
1747 SPIRV_CROSS_THROW("Struct member does not have ArrayStride set.");
1748 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001749 else
Hans-Kristian Arntzenfc37c522019-04-02 19:18:13 +02001750 SPIRV_CROSS_THROW("Struct member does not have ArrayStride set.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001751}
1752
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001753uint32_t Compiler::type_struct_member_matrix_stride(const SPIRType &type, uint32_t index) const
1754{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01001755 auto *type_meta = ir.find_meta(type.self);
1756 if (type_meta)
1757 {
1758 // Decoration must be set in valid SPIR-V, otherwise throw.
1759 // MatrixStride is part of OpMemberDecorate.
1760 auto &dec = type_meta->members[index];
1761 if (dec.decoration_flags.get(DecorationMatrixStride))
1762 return dec.matrix_stride;
1763 else
1764 SPIRV_CROSS_THROW("Struct member does not have MatrixStride set.");
1765 }
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001766 else
1767 SPIRV_CROSS_THROW("Struct member does not have MatrixStride set.");
1768}
1769
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001770size_t Compiler::get_declared_struct_size(const SPIRType &type) const
1771{
Hans-Kristian Arntzen694b3142018-04-05 16:26:54 +02001772 if (type.member_types.empty())
1773 SPIRV_CROSS_THROW("Declared struct in block cannot be empty.");
1774
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001775 uint32_t last = uint32_t(type.member_types.size() - 1);
1776 size_t offset = type_struct_member_offset(type, last);
1777 size_t size = get_declared_struct_member_size(type, last);
1778 return offset + size;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001779}
1780
Hans-Kristian Arntzene86018f2018-09-10 11:08:47 +02001781size_t Compiler::get_declared_struct_size_runtime_array(const SPIRType &type, size_t array_size) const
1782{
1783 if (type.member_types.empty())
1784 SPIRV_CROSS_THROW("Declared struct in block cannot be empty.");
1785
1786 size_t size = get_declared_struct_size(type);
1787 auto &last_type = get<SPIRType>(type.member_types.back());
1788 if (!last_type.array.empty() && last_type.array_size_literal[0] && last_type.array[0] == 0) // Runtime array
1789 size += array_size * type_struct_member_array_stride(type, uint32_t(type.member_types.size() - 1));
1790
1791 return size;
1792}
1793
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001794size_t Compiler::get_declared_struct_member_size(const SPIRType &struct_type, uint32_t index) const
1795{
Hans-Kristian Arntzen694b3142018-04-05 16:26:54 +02001796 if (struct_type.member_types.empty())
1797 SPIRV_CROSS_THROW("Declared struct in block cannot be empty.");
1798
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01001799 auto &flags = get_member_decoration_bitset(struct_type.self, index);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001800 auto &type = get<SPIRType>(struct_type.member_types[index]);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001801
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08001802 switch (type.basetype)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001803 {
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08001804 case SPIRType::Unknown:
1805 case SPIRType::Void:
1806 case SPIRType::Boolean: // Bools are purely logical, and cannot be used for externally visible types.
1807 case SPIRType::AtomicCounter:
1808 case SPIRType::Image:
1809 case SPIRType::SampledImage:
1810 case SPIRType::Sampler:
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001811 SPIRV_CROSS_THROW("Querying size for object with opaque size.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001812
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08001813 default:
1814 break;
1815 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001816
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08001817 if (!type.array.empty())
1818 {
1819 // For arrays, we can use ArrayStride to get an easy check.
Hans-Kristian Arntzendd603ea2018-02-23 15:09:28 +01001820 bool array_size_literal = type.array_size_literal.back();
1821 uint32_t array_size = array_size_literal ? type.array.back() : get<SPIRConstant>(type.array.back()).scalar();
1822 return type_struct_member_array_stride(struct_type, index) * array_size;
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08001823 }
1824 else if (type.basetype == SPIRType::Struct)
1825 {
1826 return get_declared_struct_size(type);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001827 }
1828 else
1829 {
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001830 unsigned vecsize = type.vecsize;
1831 unsigned columns = type.columns;
Hans-Kristian Arntzen2d79d362016-11-28 15:01:36 +01001832
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08001833 // Vectors.
1834 if (columns == 1)
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001835 {
1836 size_t component_size = type.width / 8;
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08001837 return vecsize * component_size;
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001838 }
Hans-Kristian Arntzen2d79d362016-11-28 15:01:36 +01001839 else
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001840 {
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001841 uint32_t matrix_stride = type_struct_member_matrix_stride(struct_type, index);
Hans-Kristian Arntzen2d79d362016-11-28 15:01:36 +01001842
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001843 // Per SPIR-V spec, matrices must be tightly packed and aligned up for vec3 accesses.
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01001844 if (flags.get(DecorationRowMajor))
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001845 return matrix_stride * vecsize;
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01001846 else if (flags.get(DecorationColMajor))
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001847 return matrix_stride * columns;
1848 else
1849 SPIRV_CROSS_THROW("Either row-major or column-major must be declared for matrices.");
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08001850 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001851 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001852}
1853
1854bool Compiler::BufferAccessHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
1855{
Chip Davis3bfb2f92018-12-03 02:06:33 -06001856 if (opcode != OpAccessChain && opcode != OpInBoundsAccessChain && opcode != OpPtrAccessChain)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001857 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001858
Chip Davis3bfb2f92018-12-03 02:06:33 -06001859 bool ptr_chain = (opcode == OpPtrAccessChain);
1860
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001861 // Invalid SPIR-V.
Hans-Kristian Arntzen2fb9aa22019-01-11 09:29:28 +01001862 if (length < (ptr_chain ? 5u : 4u))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001863 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001864
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001865 if (args[2] != id)
1866 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001867
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001868 // Don't bother traversing the entire access chain tree yet.
1869 // If we access a struct member, assume we access the entire member.
Chip Davis3bfb2f92018-12-03 02:06:33 -06001870 uint32_t index = compiler.get<SPIRConstant>(args[ptr_chain ? 4 : 3]).scalar();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001871
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001872 // Seen this index already.
1873 if (seen.find(index) != end(seen))
1874 return true;
1875 seen.insert(index);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001876
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001877 auto &type = compiler.expression_type(id);
1878 uint32_t offset = compiler.type_struct_member_offset(type, index);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001879
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001880 size_t range;
1881 // If we have another member in the struct, deduce the range by looking at the next member.
1882 // This is okay since structs in SPIR-V can have padding, but Offset decoration must be
1883 // monotonically increasing.
1884 // Of course, this doesn't take into account if the SPIR-V for some reason decided to add
1885 // very large amounts of padding, but that's not really a big deal.
1886 if (index + 1 < type.member_types.size())
1887 {
1888 range = compiler.type_struct_member_offset(type, index + 1) - offset;
1889 }
1890 else
1891 {
1892 // No padding, so just deduce it from the size of the member directly.
1893 range = compiler.get_declared_struct_member_size(type, index);
1894 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001895
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001896 ranges.push_back({ index, offset, range });
1897 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001898}
1899
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02001900SmallVector<BufferRange> Compiler::get_active_buffer_ranges(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001901{
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02001902 SmallVector<BufferRange> ranges;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001903 BufferAccessHandler handler(*this, ranges, id);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02001904 traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001905 return ranges;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001906}
1907
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02001908bool Compiler::types_are_logically_equivalent(const SPIRType &a, const SPIRType &b) const
1909{
1910 if (a.basetype != b.basetype)
1911 return false;
1912 if (a.width != b.width)
1913 return false;
1914 if (a.vecsize != b.vecsize)
1915 return false;
1916 if (a.columns != b.columns)
1917 return false;
1918 if (a.array.size() != b.array.size())
1919 return false;
1920
Hans-Kristian Arntzen46892ab2016-05-23 13:45:20 +02001921 size_t array_count = a.array.size();
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02001922 if (array_count && memcmp(a.array.data(), b.array.data(), array_count * sizeof(uint32_t)) != 0)
1923 return false;
1924
1925 if (a.basetype == SPIRType::Image || a.basetype == SPIRType::SampledImage)
1926 {
1927 if (memcmp(&a.image, &b.image, sizeof(SPIRType::Image)) != 0)
1928 return false;
1929 }
1930
1931 if (a.member_types.size() != b.member_types.size())
1932 return false;
1933
Hans-Kristian Arntzen46892ab2016-05-23 13:45:20 +02001934 size_t member_types = a.member_types.size();
1935 for (size_t i = 0; i < member_types; i++)
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02001936 {
1937 if (!types_are_logically_equivalent(get<SPIRType>(a.member_types[i]), get<SPIRType>(b.member_types[i])))
1938 return false;
1939 }
1940
1941 return true;
1942}
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02001943
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01001944const Bitset &Compiler::get_execution_mode_bitset() const
1945{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001946 return get_entry_point().flags;
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02001947}
1948
1949void Compiler::set_execution_mode(ExecutionMode mode, uint32_t arg0, uint32_t arg1, uint32_t arg2)
1950{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001951 auto &execution = get_entry_point();
1952
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01001953 execution.flags.set(mode);
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02001954 switch (mode)
1955 {
1956 case ExecutionModeLocalSize:
1957 execution.workgroup_size.x = arg0;
1958 execution.workgroup_size.y = arg1;
1959 execution.workgroup_size.z = arg2;
1960 break;
1961
1962 case ExecutionModeInvocations:
1963 execution.invocations = arg0;
1964 break;
1965
1966 case ExecutionModeOutputVertices:
1967 execution.output_vertices = arg0;
1968 break;
1969
1970 default:
1971 break;
1972 }
1973}
1974
1975void Compiler::unset_execution_mode(ExecutionMode mode)
1976{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001977 auto &execution = get_entry_point();
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01001978 execution.flags.clear(mode);
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02001979}
1980
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +02001981uint32_t Compiler::get_work_group_size_specialization_constants(SpecializationConstant &x, SpecializationConstant &y,
1982 SpecializationConstant &z) const
1983{
1984 auto &execution = get_entry_point();
Hans-Kristian Arntzenca69b612017-11-06 09:49:52 +01001985 x = { 0, 0 };
1986 y = { 0, 0 };
1987 z = { 0, 0 };
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +02001988
1989 if (execution.workgroup_size.constant != 0)
1990 {
1991 auto &c = get<SPIRConstant>(execution.workgroup_size.constant);
1992
1993 if (c.m.c[0].id[0] != 0)
1994 {
1995 x.id = c.m.c[0].id[0];
1996 x.constant_id = get_decoration(c.m.c[0].id[0], DecorationSpecId);
1997 }
1998
1999 if (c.m.c[0].id[1] != 0)
2000 {
2001 y.id = c.m.c[0].id[1];
2002 y.constant_id = get_decoration(c.m.c[0].id[1], DecorationSpecId);
2003 }
2004
2005 if (c.m.c[0].id[2] != 0)
2006 {
2007 z.id = c.m.c[0].id[2];
2008 z.constant_id = get_decoration(c.m.c[0].id[2], DecorationSpecId);
2009 }
2010 }
2011
2012 return execution.workgroup_size.constant;
2013}
2014
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002015uint32_t Compiler::get_execution_mode_argument(spv::ExecutionMode mode, uint32_t index) const
2016{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002017 auto &execution = get_entry_point();
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002018 switch (mode)
2019 {
2020 case ExecutionModeLocalSize:
2021 switch (index)
2022 {
2023 case 0:
2024 return execution.workgroup_size.x;
2025 case 1:
2026 return execution.workgroup_size.y;
2027 case 2:
2028 return execution.workgroup_size.z;
2029 default:
2030 return 0;
2031 }
2032
2033 case ExecutionModeInvocations:
2034 return execution.invocations;
2035
2036 case ExecutionModeOutputVertices:
2037 return execution.output_vertices;
2038
2039 default:
2040 return 0;
2041 }
2042}
2043
2044ExecutionModel Compiler::get_execution_model() const
2045{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002046 auto &execution = get_entry_point();
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002047 return execution.model;
2048}
Hans-Kristian Arntzen8e63c772016-07-06 09:58:01 +02002049
Chip Davise75add42019-02-05 18:13:26 -06002050bool Compiler::is_tessellation_shader(ExecutionModel model)
2051{
2052 return model == ExecutionModelTessellationControl || model == ExecutionModelTessellationEvaluation;
2053}
2054
2055bool Compiler::is_tessellation_shader() const
2056{
2057 return is_tessellation_shader(get_execution_model());
2058}
2059
Hans-Kristian Arntzen8e63c772016-07-06 09:58:01 +02002060void Compiler::set_remapped_variable_state(uint32_t id, bool remap_enable)
2061{
2062 get<SPIRVariable>(id).remapped_variable = remap_enable;
2063}
2064
2065bool Compiler::get_remapped_variable_state(uint32_t id) const
2066{
2067 return get<SPIRVariable>(id).remapped_variable;
2068}
Hans-Kristian Arntzen078eec52016-07-06 11:04:06 +02002069
2070void Compiler::set_subpass_input_remapped_components(uint32_t id, uint32_t components)
2071{
2072 get<SPIRVariable>(id).remapped_components = components;
2073}
2074
2075uint32_t Compiler::get_subpass_input_remapped_components(uint32_t id) const
2076{
2077 return get<SPIRVariable>(id).remapped_components;
2078}
Hans-Kristian Arntzen36a0b632016-07-12 14:33:04 +02002079
Hans-Kristian Arntzenacae6072019-01-04 13:19:50 +01002080void Compiler::add_implied_read_expression(SPIRExpression &e, uint32_t source)
2081{
2082 auto itr = find(begin(e.implied_read_expressions), end(e.implied_read_expressions), source);
2083 if (itr == end(e.implied_read_expressions))
2084 e.implied_read_expressions.push_back(source);
2085}
2086
2087void Compiler::add_implied_read_expression(SPIRAccessChain &e, uint32_t source)
2088{
2089 auto itr = find(begin(e.implied_read_expressions), end(e.implied_read_expressions), source);
2090 if (itr == end(e.implied_read_expressions))
2091 e.implied_read_expressions.push_back(source);
2092}
2093
Hans-Kristian Arntzen36a0b632016-07-12 14:33:04 +02002094void Compiler::inherit_expression_dependencies(uint32_t dst, uint32_t source_expression)
2095{
Hans-Kristian Arntzen75391f92017-03-20 22:38:05 +01002096 // Don't inherit any expression dependencies if the expression in dst
2097 // is not a forwarded temporary.
2098 if (forwarded_temporaries.find(dst) == end(forwarded_temporaries) ||
2099 forced_temporaries.find(dst) != end(forced_temporaries))
2100 {
2101 return;
2102 }
2103
Hans-Kristian Arntzen36a0b632016-07-12 14:33:04 +02002104 auto &e = get<SPIRExpression>(dst);
Hans-Kristian Arntzene0efa732018-03-09 13:21:38 +01002105 auto *phi = maybe_get<SPIRVariable>(source_expression);
2106 if (phi && phi->phi_variable)
2107 {
2108 // We have used a phi variable, which can change at the end of the block,
2109 // so make sure we take a dependency on this phi variable.
2110 phi->dependees.push_back(dst);
2111 }
2112
Hans-Kristian Arntzen36a0b632016-07-12 14:33:04 +02002113 auto *s = maybe_get<SPIRExpression>(source_expression);
2114 if (!s)
2115 return;
2116
2117 auto &e_deps = e.expression_dependencies;
2118 auto &s_deps = s->expression_dependencies;
2119
2120 // If we depend on a expression, we also depend on all sub-dependencies from source.
2121 e_deps.push_back(source_expression);
2122 e_deps.insert(end(e_deps), begin(s_deps), end(s_deps));
2123
2124 // Eliminate duplicated dependencies.
Hans-Kristian Arntzene0efa732018-03-09 13:21:38 +01002125 sort(begin(e_deps), end(e_deps));
Hans-Kristian Arntzen36a0b632016-07-12 14:33:04 +02002126 e_deps.erase(unique(begin(e_deps), end(e_deps)), end(e_deps));
2127}
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002128
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02002129SmallVector<EntryPoint> Compiler::get_entry_points_and_stages() const
Hans-Kristian Arntzeneecbeaa2018-03-01 14:00:04 +01002130{
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02002131 SmallVector<EntryPoint> entries;
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002132 for (auto &entry : ir.entry_points)
Hans-Kristian Arntzeneecbeaa2018-03-01 14:00:04 +01002133 entries.push_back({ entry.second.orig_name, entry.second.model });
2134 return entries;
2135}
2136
Hans-Kristian Arntzeneecbeaa2018-03-01 14:00:04 +01002137void Compiler::rename_entry_point(const std::string &old_name, const std::string &new_name, spv::ExecutionModel model)
2138{
2139 auto &entry = get_entry_point(old_name, model);
Hans-Kristian Arntzen4427cb92017-11-13 13:49:11 +01002140 entry.orig_name = new_name;
2141 entry.name = new_name;
2142}
2143
Hans-Kristian Arntzeneecbeaa2018-03-01 14:00:04 +01002144void Compiler::set_entry_point(const std::string &name, spv::ExecutionModel model)
2145{
2146 auto &entry = get_entry_point(name, model);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002147 ir.default_entry_point = entry.self;
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002148}
2149
Hans-Kristian Arntzeneecbeaa2018-03-01 14:00:04 +01002150SPIREntryPoint &Compiler::get_first_entry_point(const std::string &name)
2151{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002152 auto itr = find_if(
2153 begin(ir.entry_points), end(ir.entry_points),
2154 [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool { return entry.second.orig_name == name; });
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002155
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002156 if (itr == end(ir.entry_points))
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01002157 SPIRV_CROSS_THROW("Entry point does not exist.");
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002158
2159 return itr->second;
2160}
2161
Hans-Kristian Arntzeneecbeaa2018-03-01 14:00:04 +01002162const SPIREntryPoint &Compiler::get_first_entry_point(const std::string &name) const
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002163{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002164 auto itr = find_if(
2165 begin(ir.entry_points), end(ir.entry_points),
2166 [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool { return entry.second.orig_name == name; });
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002167
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002168 if (itr == end(ir.entry_points))
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01002169 SPIRV_CROSS_THROW("Entry point does not exist.");
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002170
2171 return itr->second;
2172}
2173
Hans-Kristian Arntzeneecbeaa2018-03-01 14:00:04 +01002174SPIREntryPoint &Compiler::get_entry_point(const std::string &name, ExecutionModel model)
2175{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002176 auto itr = find_if(begin(ir.entry_points), end(ir.entry_points),
2177 [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool {
2178 return entry.second.orig_name == name && entry.second.model == model;
2179 });
Hans-Kristian Arntzeneecbeaa2018-03-01 14:00:04 +01002180
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002181 if (itr == end(ir.entry_points))
Hans-Kristian Arntzeneecbeaa2018-03-01 14:00:04 +01002182 SPIRV_CROSS_THROW("Entry point does not exist.");
2183
2184 return itr->second;
2185}
2186
2187const SPIREntryPoint &Compiler::get_entry_point(const std::string &name, ExecutionModel model) const
2188{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002189 auto itr = find_if(begin(ir.entry_points), end(ir.entry_points),
2190 [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool {
2191 return entry.second.orig_name == name && entry.second.model == model;
2192 });
Hans-Kristian Arntzeneecbeaa2018-03-01 14:00:04 +01002193
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002194 if (itr == end(ir.entry_points))
Hans-Kristian Arntzeneecbeaa2018-03-01 14:00:04 +01002195 SPIRV_CROSS_THROW("Entry point does not exist.");
2196
2197 return itr->second;
2198}
2199
Hans-Kristian Arntzeneecbeaa2018-03-01 14:00:04 +01002200const string &Compiler::get_cleansed_entry_point_name(const std::string &name, ExecutionModel model) const
2201{
2202 return get_entry_point(name, model).name;
Bill Hollings10148472017-11-10 16:40:33 -05002203}
2204
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002205const SPIREntryPoint &Compiler::get_entry_point() const
2206{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002207 return ir.entry_points.find(ir.default_entry_point)->second;
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002208}
2209
2210SPIREntryPoint &Compiler::get_entry_point()
2211{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002212 return ir.entry_points.find(ir.default_entry_point)->second;
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002213}
2214
2215bool Compiler::interface_variable_exists_in_entry_point(uint32_t id) const
2216{
2217 auto &var = get<SPIRVariable>(id);
Robert Konrada778c362017-01-15 16:39:03 +01002218 if (var.storage != StorageClassInput && var.storage != StorageClassOutput &&
2219 var.storage != StorageClassUniformConstant)
Hans-Kristian Arntzen24df8f02017-02-04 10:26:26 +01002220 SPIRV_CROSS_THROW("Only Input, Output variables and Uniform constants are part of a shader linking interface.");
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002221
2222 // This is to avoid potential problems with very old glslang versions which did
2223 // not emit input/output interfaces properly.
2224 // We can assume they only had a single entry point, and single entry point
2225 // shaders could easily be assumed to use every interface variable anyways.
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002226 if (ir.entry_points.size() <= 1)
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002227 return true;
2228
2229 auto &execution = get_entry_point();
2230 return find(begin(execution.interface_variables), end(execution.interface_variables), id) !=
2231 end(execution.interface_variables);
2232}
Hans-Kristian Arntzenbcb55602016-09-10 13:56:36 +02002233
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002234void Compiler::CombinedImageSamplerHandler::push_remap_parameters(const SPIRFunction &func, const uint32_t *args,
2235 uint32_t length)
2236{
2237 // If possible, pipe through a remapping table so that parameters know
2238 // which variables they actually bind to in this scope.
2239 unordered_map<uint32_t, uint32_t> remapping;
2240 for (uint32_t i = 0; i < length; i++)
2241 remapping[func.arguments[i].id] = remap_parameter(args[i]);
2242 parameter_remapping.push(move(remapping));
2243}
2244
2245void Compiler::CombinedImageSamplerHandler::pop_remap_parameters()
2246{
2247 parameter_remapping.pop();
2248}
2249
2250uint32_t Compiler::CombinedImageSamplerHandler::remap_parameter(uint32_t id)
2251{
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002252 auto *var = compiler.maybe_get_backing_variable(id);
2253 if (var)
2254 id = var->self;
2255
Hans-Kristian Arntzen901b45e2016-09-10 22:21:57 +02002256 if (parameter_remapping.empty())
2257 return id;
2258
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002259 auto &remapping = parameter_remapping.top();
2260 auto itr = remapping.find(id);
2261 if (itr != end(remapping))
2262 return itr->second;
2263 else
2264 return id;
2265}
2266
2267bool Compiler::CombinedImageSamplerHandler::begin_function_scope(const uint32_t *args, uint32_t length)
2268{
2269 if (length < 3)
2270 return false;
2271
2272 auto &callee = compiler.get<SPIRFunction>(args[2]);
2273 args += 3;
2274 length -= 3;
2275 push_remap_parameters(callee, args, length);
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002276 functions.push(&callee);
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002277 return true;
2278}
2279
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002280bool Compiler::CombinedImageSamplerHandler::end_function_scope(const uint32_t *args, uint32_t length)
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002281{
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002282 if (length < 3)
2283 return false;
2284
2285 auto &callee = compiler.get<SPIRFunction>(args[2]);
2286 args += 3;
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002287
2288 // There are two types of cases we have to handle,
2289 // a callee might call sampler2D(texture2D, sampler) directly where
2290 // one or more parameters originate from parameters.
2291 // Alternatively, we need to provide combined image samplers to our callees,
2292 // and in this case we need to add those as well.
2293
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002294 pop_remap_parameters();
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002295
2296 // Our callee has now been processed at least once.
2297 // No point in doing it again.
2298 callee.do_combined_parameters = false;
2299
2300 auto &params = functions.top()->combined_parameters;
2301 functions.pop();
2302 if (functions.empty())
2303 return true;
2304
2305 auto &caller = *functions.top();
2306 if (caller.do_combined_parameters)
2307 {
2308 for (auto &param : params)
2309 {
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002310 uint32_t image_id = param.global_image ? param.image_id : args[param.image_id];
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002311 uint32_t sampler_id = param.global_sampler ? param.sampler_id : args[param.sampler_id];
2312
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002313 auto *i = compiler.maybe_get_backing_variable(image_id);
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002314 auto *s = compiler.maybe_get_backing_variable(sampler_id);
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002315 if (i)
2316 image_id = i->self;
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002317 if (s)
2318 sampler_id = s->self;
2319
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002320 register_combined_image_sampler(caller, image_id, sampler_id, param.depth);
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002321 }
2322 }
2323
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002324 return true;
2325}
2326
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002327void Compiler::CombinedImageSamplerHandler::register_combined_image_sampler(SPIRFunction &caller, uint32_t image_id,
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002328 uint32_t sampler_id, bool depth)
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002329{
2330 // We now have a texture ID and a sampler ID which will either be found as a global
2331 // or a parameter in our own function. If both are global, they will not need a parameter,
2332 // otherwise, add it to our list.
2333 SPIRFunction::CombinedImageSamplerParameter param = {
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002334 0u, image_id, sampler_id, true, true, depth,
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002335 };
2336
2337 auto texture_itr = find_if(begin(caller.arguments), end(caller.arguments),
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002338 [image_id](const SPIRFunction::Parameter &p) { return p.id == image_id; });
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002339 auto sampler_itr = find_if(begin(caller.arguments), end(caller.arguments),
2340 [sampler_id](const SPIRFunction::Parameter &p) { return p.id == sampler_id; });
2341
2342 if (texture_itr != end(caller.arguments))
2343 {
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002344 param.global_image = false;
Hans-Kristian Arntzen7630d3c2016-11-21 12:14:02 +01002345 param.image_id = uint32_t(texture_itr - begin(caller.arguments));
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002346 }
2347
2348 if (sampler_itr != end(caller.arguments))
2349 {
2350 param.global_sampler = false;
Hans-Kristian Arntzen7630d3c2016-11-21 12:14:02 +01002351 param.sampler_id = uint32_t(sampler_itr - begin(caller.arguments));
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002352 }
2353
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002354 if (param.global_image && param.global_sampler)
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002355 return;
2356
2357 auto itr = find_if(begin(caller.combined_parameters), end(caller.combined_parameters),
2358 [&param](const SPIRFunction::CombinedImageSamplerParameter &p) {
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002359 return param.image_id == p.image_id && param.sampler_id == p.sampler_id &&
2360 param.global_image == p.global_image && param.global_sampler == p.global_sampler;
Hans-Kristian Arntzence18d4c2017-11-17 13:38:29 +01002361 });
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002362
2363 if (itr == end(caller.combined_parameters))
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002364 {
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002365 uint32_t id = compiler.ir.increase_bound_by(3);
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002366 auto type_id = id + 0;
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002367 auto ptr_type_id = id + 1;
2368 auto combined_id = id + 2;
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002369 auto &base = compiler.expression_type(image_id);
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002370 auto &type = compiler.set<SPIRType>(type_id);
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002371 auto &ptr_type = compiler.set<SPIRType>(ptr_type_id);
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002372
2373 type = base;
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002374 type.self = type_id;
2375 type.basetype = SPIRType::SampledImage;
2376 type.pointer = false;
2377 type.storage = StorageClassGeneric;
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002378 type.image.depth = depth;
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002379
2380 ptr_type = type;
2381 ptr_type.pointer = true;
2382 ptr_type.storage = StorageClassUniformConstant;
Chip Davis3bfb2f92018-12-03 02:06:33 -06002383 ptr_type.parent_type = type_id;
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002384
2385 // Build new variable.
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002386 compiler.set<SPIRVariable>(combined_id, ptr_type_id, StorageClassFunction, 0);
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002387
2388 // Inherit RelaxedPrecision (and potentially other useful flags if deemed relevant).
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002389 auto &new_flags = compiler.ir.meta[combined_id].decoration.decoration_flags;
2390 auto &old_flags = compiler.ir.meta[sampler_id].decoration.decoration_flags;
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01002391 new_flags.reset();
2392 if (old_flags.get(DecorationRelaxedPrecision))
2393 new_flags.set(DecorationRelaxedPrecision);
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002394
2395 param.id = combined_id;
2396
2397 compiler.set_name(combined_id,
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002398 join("SPIRV_Cross_Combined", compiler.to_name(image_id), compiler.to_name(sampler_id)));
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002399
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002400 caller.combined_parameters.push_back(param);
Hans-Kristian Arntzen9cb86162017-02-05 10:50:14 +01002401 caller.shadow_arguments.push_back({ ptr_type_id, combined_id, 0u, 0u, true });
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002402 }
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002403}
2404
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002405bool Compiler::DummySamplerForCombinedImageHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
2406{
2407 if (need_dummy_sampler)
2408 {
2409 // No need to traverse further, we know the result.
2410 return false;
2411 }
2412
2413 switch (opcode)
2414 {
2415 case OpLoad:
2416 {
2417 if (length < 3)
2418 return false;
2419
2420 uint32_t result_type = args[0];
2421
2422 auto &type = compiler.get<SPIRType>(result_type);
Hans-Kristian Arntzen47b37422018-02-21 13:46:16 +01002423 bool separate_image =
2424 type.basetype == SPIRType::Image && type.image.sampled == 1 && type.image.dim != DimBuffer;
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002425
2426 // If not separate image, don't bother.
2427 if (!separate_image)
2428 return true;
2429
2430 uint32_t id = args[1];
2431 uint32_t ptr = args[2];
2432 compiler.set<SPIRExpression>(id, "", result_type, true);
2433 compiler.register_read(id, ptr, true);
2434 break;
2435 }
2436
2437 case OpImageFetch:
Hans-Kristian Arntzen40bbf6b2018-04-30 11:18:18 +02002438 case OpImageQuerySizeLod:
2439 case OpImageQuerySize:
2440 case OpImageQueryLevels:
2441 case OpImageQuerySamples:
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002442 {
Hans-Kristian Arntzen40bbf6b2018-04-30 11:18:18 +02002443 // If we are fetching or querying LOD from a plain OpTypeImage, we must pre-combine with our dummy sampler.
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002444 auto *var = compiler.maybe_get_backing_variable(args[2]);
2445 if (var)
2446 {
2447 auto &type = compiler.get<SPIRType>(var->basetype);
2448 if (type.basetype == SPIRType::Image && type.image.sampled == 1 && type.image.dim != DimBuffer)
2449 need_dummy_sampler = true;
2450 }
2451
2452 break;
2453 }
2454
2455 case OpInBoundsAccessChain:
2456 case OpAccessChain:
Chip Davis3bfb2f92018-12-03 02:06:33 -06002457 case OpPtrAccessChain:
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002458 {
2459 if (length < 3)
2460 return false;
2461
Hans-Kristian Arntzena39eb482018-04-23 11:52:05 +02002462 uint32_t result_type = args[0];
2463 auto &type = compiler.get<SPIRType>(result_type);
Hans-Kristian Arntzen47b37422018-02-21 13:46:16 +01002464 bool separate_image =
2465 type.basetype == SPIRType::Image && type.image.sampled == 1 && type.image.dim != DimBuffer;
Hans-Kristian Arntzena39eb482018-04-23 11:52:05 +02002466 if (!separate_image)
2467 return true;
2468
2469 uint32_t id = args[1];
2470 uint32_t ptr = args[2];
2471 compiler.set<SPIRExpression>(id, "", result_type, true);
2472 compiler.register_read(id, ptr, true);
Hans-Kristian Arntzen7eba2472018-05-11 10:14:20 +02002473
2474 // Other backends might use SPIRAccessChain for this later.
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002475 compiler.ir.ids[id].set_allow_type_rewrite();
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002476 break;
2477 }
2478
2479 default:
2480 break;
2481 }
2482
2483 return true;
2484}
2485
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002486bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
2487{
2488 // We need to figure out where samplers and images are loaded from, so do only the bare bones compilation we need.
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002489 bool is_fetch = false;
2490
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002491 switch (opcode)
2492 {
2493 case OpLoad:
2494 {
2495 if (length < 3)
2496 return false;
2497
2498 uint32_t result_type = args[0];
2499
2500 auto &type = compiler.get<SPIRType>(result_type);
2501 bool separate_image = type.basetype == SPIRType::Image && type.image.sampled == 1;
2502 bool separate_sampler = type.basetype == SPIRType::Sampler;
2503
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002504 // If not separate image or sampler, don't bother.
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002505 if (!separate_image && !separate_sampler)
2506 return true;
2507
2508 uint32_t id = args[1];
2509 uint32_t ptr = args[2];
2510 compiler.set<SPIRExpression>(id, "", result_type, true);
2511 compiler.register_read(id, ptr, true);
2512 return true;
2513 }
2514
2515 case OpInBoundsAccessChain:
2516 case OpAccessChain:
Chip Davis3bfb2f92018-12-03 02:06:33 -06002517 case OpPtrAccessChain:
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002518 {
2519 if (length < 3)
2520 return false;
2521
2522 // Technically, it is possible to have arrays of textures and arrays of samplers and combine them, but this becomes essentially
2523 // impossible to implement, since we don't know which concrete sampler we are accessing.
2524 // One potential way is to create a combinatorial explosion where N textures and M samplers are combined into N * M sampler2Ds,
2525 // but this seems ridiculously complicated for a problem which is easy to work around.
2526 // Checking access chains like this assumes we don't have samplers or textures inside uniform structs, but this makes no sense.
2527
Hans-Kristian Arntzena39eb482018-04-23 11:52:05 +02002528 uint32_t result_type = args[0];
2529
2530 auto &type = compiler.get<SPIRType>(result_type);
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002531 bool separate_image = type.basetype == SPIRType::Image && type.image.sampled == 1;
2532 bool separate_sampler = type.basetype == SPIRType::Sampler;
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002533 if (separate_sampler)
Hans-Kristian Arntzen36e1c472017-05-06 13:59:00 +02002534 SPIRV_CROSS_THROW(
2535 "Attempting to use arrays or structs of separate samplers. This is not possible to statically "
2536 "remap to plain GLSL.");
Hans-Kristian Arntzena39eb482018-04-23 11:52:05 +02002537
2538 if (separate_image)
2539 {
2540 uint32_t id = args[1];
2541 uint32_t ptr = args[2];
2542 compiler.set<SPIRExpression>(id, "", result_type, true);
2543 compiler.register_read(id, ptr, true);
2544 }
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002545 return true;
2546 }
2547
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002548 case OpImageFetch:
Hans-Kristian Arntzen40bbf6b2018-04-30 11:18:18 +02002549 case OpImageQuerySizeLod:
2550 case OpImageQuerySize:
2551 case OpImageQueryLevels:
2552 case OpImageQuerySamples:
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002553 {
Hans-Kristian Arntzen40bbf6b2018-04-30 11:18:18 +02002554 // If we are fetching from a plain OpTypeImage or querying LOD, we must pre-combine with our dummy sampler.
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002555 auto *var = compiler.maybe_get_backing_variable(args[2]);
2556 if (!var)
2557 return true;
2558
2559 auto &type = compiler.get<SPIRType>(var->basetype);
2560 if (type.basetype == SPIRType::Image && type.image.sampled == 1 && type.image.dim != DimBuffer)
2561 {
2562 if (compiler.dummy_sampler_id == 0)
Hans-Kristian Arntzen47b37422018-02-21 13:46:16 +01002563 SPIRV_CROSS_THROW("texelFetch without sampler was found, but no dummy sampler has been created with "
2564 "build_dummy_sampler_for_combined_images().");
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002565
2566 // Do it outside.
2567 is_fetch = true;
2568 break;
2569 }
2570
2571 return true;
2572 }
2573
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002574 case OpSampledImage:
2575 // Do it outside.
2576 break;
2577
2578 default:
2579 return true;
2580 }
2581
Hans-Kristian Arntzendfb65972016-09-11 12:05:20 +02002582 // Registers sampler2D calls used in case they are parameters so
2583 // that their callees know which combined image samplers to propagate down the call stack.
2584 if (!functions.empty())
2585 {
2586 auto &callee = *functions.top();
2587 if (callee.do_combined_parameters)
2588 {
2589 uint32_t image_id = args[2];
2590
2591 auto *image = compiler.maybe_get_backing_variable(image_id);
2592 if (image)
2593 image_id = image->self;
2594
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002595 uint32_t sampler_id = is_fetch ? compiler.dummy_sampler_id : args[3];
Hans-Kristian Arntzendfb65972016-09-11 12:05:20 +02002596 auto *sampler = compiler.maybe_get_backing_variable(sampler_id);
2597 if (sampler)
2598 sampler_id = sampler->self;
2599
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002600 auto &combined_type = compiler.get<SPIRType>(args[0]);
2601 register_combined_image_sampler(callee, image_id, sampler_id, combined_type.image.depth);
Hans-Kristian Arntzendfb65972016-09-11 12:05:20 +02002602 }
2603 }
2604
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002605 // For function calls, we need to remap IDs which are function parameters into global variables.
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002606 // This information is statically known from the current place in the call stack.
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002607 // Function parameters are not necessarily pointers, so if we don't have a backing variable, remapping will know
2608 // which backing variable the image/sample came from.
Hans-Kristian Arntzendfb65972016-09-11 12:05:20 +02002609 uint32_t image_id = remap_parameter(args[2]);
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002610 uint32_t sampler_id = is_fetch ? compiler.dummy_sampler_id : remap_parameter(args[3]);
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002611
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002612 auto itr = find_if(begin(compiler.combined_image_samplers), end(compiler.combined_image_samplers),
2613 [image_id, sampler_id](const CombinedImageSampler &combined) {
2614 return combined.image_id == image_id && combined.sampler_id == sampler_id;
Hans-Kristian Arntzence18d4c2017-11-17 13:38:29 +01002615 });
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002616
2617 if (itr == end(compiler.combined_image_samplers))
2618 {
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002619 uint32_t sampled_type;
2620 if (is_fetch)
2621 {
2622 // Have to invent the sampled image type.
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002623 sampled_type = compiler.ir.increase_bound_by(1);
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002624 auto &type = compiler.set<SPIRType>(sampled_type);
2625 type = compiler.expression_type(args[2]);
2626 type.self = sampled_type;
2627 type.basetype = SPIRType::SampledImage;
Hans-Kristian Arntzen40bbf6b2018-04-30 11:18:18 +02002628 type.image.depth = false;
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002629 }
2630 else
2631 {
2632 sampled_type = args[0];
2633 }
2634
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002635 auto id = compiler.ir.increase_bound_by(2);
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002636 auto type_id = id + 0;
2637 auto combined_id = id + 1;
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002638
2639 // Make a new type, pointer to OpTypeSampledImage, so we can make a variable of this type.
2640 // We will probably have this type lying around, but it doesn't hurt to make duplicates for internal purposes.
2641 auto &type = compiler.set<SPIRType>(type_id);
2642 auto &base = compiler.get<SPIRType>(sampled_type);
2643 type = base;
2644 type.pointer = true;
2645 type.storage = StorageClassUniformConstant;
Chip Davis3bfb2f92018-12-03 02:06:33 -06002646 type.parent_type = type_id;
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002647
2648 // Build new variable.
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002649 compiler.set<SPIRVariable>(combined_id, type_id, StorageClassUniformConstant, 0);
Hans-Kristian Arntzen14bc1ff2016-09-10 18:07:52 +02002650
2651 // Inherit RelaxedPrecision (and potentially other useful flags if deemed relevant).
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002652 auto &new_flags = compiler.ir.meta[combined_id].decoration.decoration_flags;
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002653 // Fetch inherits precision from the image, not sampler (there is no sampler).
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002654 auto &old_flags = compiler.ir.meta[is_fetch ? image_id : sampler_id].decoration.decoration_flags;
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01002655 new_flags.reset();
2656 if (old_flags.get(DecorationRelaxedPrecision))
2657 new_flags.set(DecorationRelaxedPrecision);
Hans-Kristian Arntzen14bc1ff2016-09-10 18:07:52 +02002658
Hans-Kristian Arntzena39eb482018-04-23 11:52:05 +02002659 // Propagate the array type for the original image as well.
2660 auto *var = compiler.maybe_get_backing_variable(image_id);
2661 if (var)
2662 {
2663 auto &parent_type = compiler.get<SPIRType>(var->basetype);
2664 type.array = parent_type.array;
2665 type.array_size_literal = parent_type.array_size_literal;
2666 }
2667
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002668 compiler.combined_image_samplers.push_back({ combined_id, image_id, sampler_id });
2669 }
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002670
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002671 return true;
2672}
2673
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002674uint32_t Compiler::build_dummy_sampler_for_combined_images()
2675{
2676 DummySamplerForCombinedImageHandler handler(*this);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002677 traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002678 if (handler.need_dummy_sampler)
2679 {
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002680 uint32_t offset = ir.increase_bound_by(3);
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002681 auto type_id = offset + 0;
2682 auto ptr_type_id = offset + 1;
2683 auto var_id = offset + 2;
2684
2685 SPIRType sampler_type;
2686 auto &sampler = set<SPIRType>(type_id);
2687 sampler.basetype = SPIRType::Sampler;
2688
2689 auto &ptr_sampler = set<SPIRType>(ptr_type_id);
2690 ptr_sampler = sampler;
2691 ptr_sampler.self = type_id;
2692 ptr_sampler.storage = StorageClassUniformConstant;
Hans-Kristian Arntzen1a2e4de2018-02-21 13:43:16 +01002693 ptr_sampler.pointer = true;
Chip Davis3bfb2f92018-12-03 02:06:33 -06002694 ptr_sampler.parent_type = type_id;
Hans-Kristian Arntzen4db70612018-02-21 13:08:30 +01002695
2696 set<SPIRVariable>(var_id, ptr_type_id, StorageClassUniformConstant, 0);
2697 set_name(var_id, "SPIRV_Cross_DummySampler");
2698 dummy_sampler_id = var_id;
2699 return var_id;
2700 }
2701 else
2702 return 0;
2703}
2704
Hans-Kristian Arntzenbcb55602016-09-10 13:56:36 +02002705void Compiler::build_combined_image_samplers()
2706{
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +01002707 ir.for_each_typed_id<SPIRFunction>([&](uint32_t, SPIRFunction &func) {
2708 func.combined_parameters.clear();
2709 func.shadow_arguments.clear();
2710 func.do_combined_parameters = true;
2711 });
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002712
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002713 combined_image_samplers.clear();
2714 CombinedImageSamplerHandler handler(*this);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002715 traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
Hans-Kristian Arntzenbcb55602016-09-10 13:56:36 +02002716}
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02002717
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02002718SmallVector<SpecializationConstant> Compiler::get_specialization_constants() const
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02002719{
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02002720 SmallVector<SpecializationConstant> spec_consts;
Hans-Kristian Arntzend92de002019-01-10 09:49:33 +01002721 ir.for_each_typed_id<SPIRConstant>([&](uint32_t, const SPIRConstant &c) {
2722 if (c.specialization && has_decoration(c.self, DecorationSpecId))
2723 spec_consts.push_back({ c.self, get_decoration(c.self, DecorationSpecId) });
2724 });
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02002725 return spec_consts;
2726}
2727
2728SPIRConstant &Compiler::get_constant(uint32_t id)
2729{
2730 return get<SPIRConstant>(id);
2731}
2732
2733const SPIRConstant &Compiler::get_constant(uint32_t id) const
2734{
2735 return get<SPIRConstant>(id);
2736}
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01002737
Hans-Kristian Arntzenbf5c0752017-03-25 16:28:44 +01002738static bool exists_unaccessed_path_to_return(const CFG &cfg, uint32_t block, const unordered_set<uint32_t> &blocks)
Hans-Kristian Arntzenb2c2e642017-03-25 16:25:30 +01002739{
2740 // This block accesses the variable.
2741 if (blocks.find(block) != end(blocks))
2742 return false;
2743
2744 // We are at the end of the CFG.
2745 if (cfg.get_succeeding_edges(block).empty())
2746 return true;
2747
2748 // If any of our successors have a path to the end, there exists a path from block.
2749 for (auto &succ : cfg.get_succeeding_edges(block))
2750 if (exists_unaccessed_path_to_return(cfg, succ, blocks))
2751 return true;
2752
2753 return false;
2754}
2755
Hans-Kristian Arntzenbf5c0752017-03-25 16:28:44 +01002756void Compiler::analyze_parameter_preservation(
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02002757 SPIRFunction &entry, const CFG &cfg, const unordered_map<uint32_t, unordered_set<uint32_t>> &variable_to_blocks,
2758 const unordered_map<uint32_t, unordered_set<uint32_t>> &complete_write_blocks)
Hans-Kristian Arntzenb2c2e642017-03-25 16:25:30 +01002759{
2760 for (auto &arg : entry.arguments)
2761 {
2762 // Non-pointers are always inputs.
2763 auto &type = get<SPIRType>(arg.type);
2764 if (!type.pointer)
2765 continue;
2766
2767 // Opaque argument types are always in
2768 bool potential_preserve;
2769 switch (type.basetype)
2770 {
2771 case SPIRType::Sampler:
2772 case SPIRType::Image:
2773 case SPIRType::SampledImage:
2774 case SPIRType::AtomicCounter:
2775 potential_preserve = false;
2776 break;
2777
2778 default:
2779 potential_preserve = true;
2780 break;
2781 }
2782
2783 if (!potential_preserve)
2784 continue;
2785
2786 auto itr = variable_to_blocks.find(arg.id);
2787 if (itr == end(variable_to_blocks))
2788 {
2789 // Variable is never accessed.
2790 continue;
2791 }
2792
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02002793 // We have accessed a variable, but there was no complete writes to that variable.
2794 // We deduce that we must preserve the argument.
2795 itr = complete_write_blocks.find(arg.id);
2796 if (itr == end(complete_write_blocks))
2797 {
2798 arg.read_count++;
2799 continue;
2800 }
2801
2802 // If there is a path through the CFG where no block completely writes to the variable, the variable will be in an undefined state
Hans-Kristian Arntzenb2c2e642017-03-25 16:25:30 +01002803 // when the function returns. We therefore need to implicitly preserve the variable in case there are writers in the function.
2804 // Major case here is if a function is
2805 // void foo(int &var) { if (cond) var = 10; }
2806 // Using read/write counts, we will think it's just an out variable, but it really needs to be inout,
2807 // because if we don't write anything whatever we put into the function must return back to the caller.
2808 if (exists_unaccessed_path_to_return(cfg, entry.entry_block, itr->second))
2809 arg.read_count++;
2810 }
2811}
2812
Hans-Kristian Arntzenb5ed7062018-07-05 10:42:05 +02002813Compiler::AnalyzeVariableScopeAccessHandler::AnalyzeVariableScopeAccessHandler(Compiler &compiler_,
2814 SPIRFunction &entry_)
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002815 : compiler(compiler_)
2816 , entry(entry_)
2817{
2818}
2819
Hans-Kristian Arntzenb5ed7062018-07-05 10:42:05 +02002820bool Compiler::AnalyzeVariableScopeAccessHandler::follow_function_call(const SPIRFunction &)
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002821{
2822 // Only analyze within this function.
2823 return false;
2824}
2825
Hans-Kristian Arntzenb5ed7062018-07-05 10:42:05 +02002826void Compiler::AnalyzeVariableScopeAccessHandler::set_current_block(const SPIRBlock &block)
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002827{
2828 current_block = &block;
2829
2830 // If we're branching to a block which uses OpPhi, in GLSL
2831 // this will be a variable write when we branch,
2832 // so we need to track access to these variables as well to
2833 // have a complete picture.
2834 const auto test_phi = [this, &block](uint32_t to) {
2835 auto &next = compiler.get<SPIRBlock>(to);
2836 for (auto &phi : next.phi_variables)
2837 {
2838 if (phi.parent == block.self)
2839 {
2840 accessed_variables_to_block[phi.function_variable].insert(block.self);
2841 // Phi variables are also accessed in our target branch block.
2842 accessed_variables_to_block[phi.function_variable].insert(next.self);
2843
2844 notify_variable_access(phi.local_variable, block.self);
2845 }
2846 }
2847 };
2848
2849 switch (block.terminator)
2850 {
2851 case SPIRBlock::Direct:
2852 notify_variable_access(block.condition, block.self);
2853 test_phi(block.next_block);
2854 break;
2855
2856 case SPIRBlock::Select:
2857 notify_variable_access(block.condition, block.self);
2858 test_phi(block.true_block);
2859 test_phi(block.false_block);
2860 break;
2861
2862 case SPIRBlock::MultiSelect:
2863 notify_variable_access(block.condition, block.self);
2864 for (auto &target : block.cases)
2865 test_phi(target.block);
2866 if (block.default_block)
2867 test_phi(block.default_block);
2868 break;
2869
2870 default:
2871 break;
2872 }
2873}
2874
2875void Compiler::AnalyzeVariableScopeAccessHandler::notify_variable_access(uint32_t id, uint32_t block)
2876{
Hans-Kristian Arntzene23c9ea2019-04-10 15:46:48 +02002877 if (id == 0)
2878 return;
2879
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002880 if (id_is_phi_variable(id))
2881 accessed_variables_to_block[id].insert(block);
2882 else if (id_is_potential_temporary(id))
2883 accessed_temporaries_to_block[id].insert(block);
2884}
2885
2886bool Compiler::AnalyzeVariableScopeAccessHandler::id_is_phi_variable(uint32_t id) const
2887{
2888 if (id >= compiler.get_current_id_bound())
2889 return false;
2890 auto *var = compiler.maybe_get<SPIRVariable>(id);
2891 return var && var->phi_variable;
2892}
2893
2894bool Compiler::AnalyzeVariableScopeAccessHandler::id_is_potential_temporary(uint32_t id) const
2895{
2896 if (id >= compiler.get_current_id_bound())
2897 return false;
2898
2899 // Temporaries are not created before we start emitting code.
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002900 return compiler.ir.ids[id].empty() || (compiler.ir.ids[id].get_type() == TypeExpression);
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002901}
2902
2903bool Compiler::AnalyzeVariableScopeAccessHandler::handle(spv::Op op, const uint32_t *args, uint32_t length)
2904{
2905 // Keep track of the types of temporaries, so we can hoist them out as necessary.
2906 uint32_t result_type, result_id;
2907 if (compiler.instruction_to_result_type(result_type, result_id, op, args, length))
2908 result_id_to_type[result_id] = result_type;
2909
2910 switch (op)
2911 {
2912 case OpStore:
2913 {
2914 if (length < 2)
2915 return false;
2916
2917 uint32_t ptr = args[0];
2918 auto *var = compiler.maybe_get_backing_variable(ptr);
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002919
2920 // If we store through an access chain, we have a partial write.
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02002921 if (var)
2922 {
2923 accessed_variables_to_block[var->self].insert(current_block->self);
2924 if (var->self == ptr)
2925 complete_write_variables_to_block[var->self].insert(current_block->self);
2926 else
2927 partial_write_variables_to_block[var->self].insert(current_block->self);
2928 }
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002929
Hans-Kristian Arntzene23c9ea2019-04-10 15:46:48 +02002930 // args[0] might be an access chain we have to track use of.
2931 notify_variable_access(args[0], current_block->self);
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002932 // Might try to store a Phi variable here.
2933 notify_variable_access(args[1], current_block->self);
2934 break;
2935 }
2936
2937 case OpAccessChain:
2938 case OpInBoundsAccessChain:
Chip Davis3bfb2f92018-12-03 02:06:33 -06002939 case OpPtrAccessChain:
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002940 {
2941 if (length < 3)
2942 return false;
2943
2944 uint32_t ptr = args[2];
2945 auto *var = compiler.maybe_get<SPIRVariable>(ptr);
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02002946 if (var)
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002947 accessed_variables_to_block[var->self].insert(current_block->self);
2948
Hans-Kristian Arntzene23c9ea2019-04-10 15:46:48 +02002949 // args[2] might be another access chain we have to track use of.
2950 for (uint32_t i = 2; i < length; i++)
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002951 notify_variable_access(args[i], current_block->self);
2952
Hans-Kristian Arntzene23c9ea2019-04-10 15:46:48 +02002953 // Also keep track of the access chain pointer itself.
2954 // In exceptionally rare cases, we can end up with a case where
2955 // the access chain is generated in the loop body, but is consumed in continue block.
2956 // This means we need complex loop workarounds, and we must detect this via CFG analysis.
2957 notify_variable_access(args[1], current_block->self);
2958
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002959 // The result of an access chain is a fixed expression and is not really considered a temporary.
2960 auto &e = compiler.set<SPIRExpression>(args[1], "", args[0], true);
2961 auto *backing_variable = compiler.maybe_get_backing_variable(ptr);
2962 e.loaded_from = backing_variable ? backing_variable->self : 0;
2963
2964 // Other backends might use SPIRAccessChain for this later.
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002965 compiler.ir.ids[args[1]].set_allow_type_rewrite();
Hans-Kristian Arntzene23c9ea2019-04-10 15:46:48 +02002966 access_chain_expressions.insert(args[1]);
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002967 break;
2968 }
2969
2970 case OpCopyMemory:
2971 {
2972 if (length < 2)
2973 return false;
2974
2975 uint32_t lhs = args[0];
2976 uint32_t rhs = args[1];
2977 auto *var = compiler.maybe_get_backing_variable(lhs);
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002978
2979 // If we store through an access chain, we have a partial write.
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02002980 if (var)
2981 {
2982 accessed_variables_to_block[var->self].insert(current_block->self);
2983 if (var->self == lhs)
2984 complete_write_variables_to_block[var->self].insert(current_block->self);
2985 else
2986 partial_write_variables_to_block[var->self].insert(current_block->self);
2987 }
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002988
Hans-Kristian Arntzene23c9ea2019-04-10 15:46:48 +02002989 // args[0:1] might be access chains we have to track use of.
2990 for (uint32_t i = 0; i < 2; i++)
2991 notify_variable_access(args[i], current_block->self);
2992
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002993 var = compiler.maybe_get_backing_variable(rhs);
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02002994 if (var)
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002995 accessed_variables_to_block[var->self].insert(current_block->self);
2996 break;
2997 }
2998
2999 case OpCopyObject:
3000 {
3001 if (length < 3)
3002 return false;
3003
3004 auto *var = compiler.maybe_get_backing_variable(args[2]);
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02003005 if (var)
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003006 accessed_variables_to_block[var->self].insert(current_block->self);
3007
Hans-Kristian Arntzene23c9ea2019-04-10 15:46:48 +02003008 // Might be an access chain which we have to keep track of.
3009 notify_variable_access(args[1], current_block->self);
3010 if (access_chain_expressions.count(args[2]))
3011 access_chain_expressions.insert(args[1]);
3012
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003013 // Might try to copy a Phi variable here.
3014 notify_variable_access(args[2], current_block->self);
3015 break;
3016 }
3017
3018 case OpLoad:
3019 {
3020 if (length < 3)
3021 return false;
3022 uint32_t ptr = args[2];
3023 auto *var = compiler.maybe_get_backing_variable(ptr);
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02003024 if (var)
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003025 accessed_variables_to_block[var->self].insert(current_block->self);
3026
3027 // Loaded value is a temporary.
3028 notify_variable_access(args[1], current_block->self);
Hans-Kristian Arntzene23c9ea2019-04-10 15:46:48 +02003029
3030 // Might be an access chain we have to track use of.
3031 notify_variable_access(args[2], current_block->self);
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003032 break;
3033 }
3034
3035 case OpFunctionCall:
3036 {
3037 if (length < 3)
3038 return false;
3039
3040 length -= 3;
3041 args += 3;
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02003042
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003043 for (uint32_t i = 0; i < length; i++)
3044 {
3045 auto *var = compiler.maybe_get_backing_variable(args[i]);
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02003046 if (var)
3047 {
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003048 accessed_variables_to_block[var->self].insert(current_block->self);
Hans-Kristian Arntzenb5ed7062018-07-05 10:42:05 +02003049 // Assume we can get partial writes to this variable.
3050 partial_write_variables_to_block[var->self].insert(current_block->self);
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02003051 }
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003052
3053 // Cannot easily prove if argument we pass to a function is completely written.
3054 // Usually, functions write to a dummy variable,
3055 // which is then copied to in full to the real argument.
3056
3057 // Might try to copy a Phi variable here.
3058 notify_variable_access(args[i], current_block->self);
3059 }
3060
3061 // Return value may be a temporary.
3062 notify_variable_access(args[1], current_block->self);
3063 break;
3064 }
3065
3066 case OpExtInst:
3067 {
3068 for (uint32_t i = 4; i < length; i++)
3069 notify_variable_access(args[i], current_block->self);
3070 notify_variable_access(args[1], current_block->self);
3071 break;
3072 }
3073
3074 case OpArrayLength:
Hans-Kristian Arntzend2cc43e2019-02-19 17:00:49 +01003075 // Uses literals, but cannot be a phi variable or temporary, so ignore.
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003076 break;
3077
3078 // Atomics shouldn't be able to access function-local variables.
3079 // Some GLSL builtins access a pointer.
3080
3081 case OpCompositeInsert:
3082 case OpVectorShuffle:
3083 // Specialize for opcode which contains literals.
3084 for (uint32_t i = 1; i < 4; i++)
3085 notify_variable_access(args[i], current_block->self);
3086 break;
3087
3088 case OpCompositeExtract:
3089 // Specialize for opcode which contains literals.
3090 for (uint32_t i = 1; i < 3; i++)
3091 notify_variable_access(args[i], current_block->self);
3092 break;
3093
Hans-Kristian Arntzend2cc43e2019-02-19 17:00:49 +01003094 case OpImageWrite:
3095 for (uint32_t i = 0; i < length; i++)
3096 {
3097 // Argument 3 is a literal.
3098 if (i != 3)
3099 notify_variable_access(args[i], current_block->self);
3100 }
3101 break;
3102
3103 case OpImageSampleImplicitLod:
3104 case OpImageSampleExplicitLod:
3105 case OpImageSparseSampleImplicitLod:
3106 case OpImageSparseSampleExplicitLod:
3107 case OpImageSampleProjImplicitLod:
3108 case OpImageSampleProjExplicitLod:
3109 case OpImageSparseSampleProjImplicitLod:
3110 case OpImageSparseSampleProjExplicitLod:
3111 case OpImageFetch:
3112 case OpImageSparseFetch:
3113 case OpImageRead:
3114 case OpImageSparseRead:
3115 for (uint32_t i = 1; i < length; i++)
3116 {
3117 // Argument 4 is a literal.
3118 if (i != 4)
3119 notify_variable_access(args[i], current_block->self);
3120 }
3121 break;
3122
3123 case OpImageSampleDrefImplicitLod:
3124 case OpImageSampleDrefExplicitLod:
3125 case OpImageSparseSampleDrefImplicitLod:
3126 case OpImageSparseSampleDrefExplicitLod:
3127 case OpImageSampleProjDrefImplicitLod:
3128 case OpImageSampleProjDrefExplicitLod:
3129 case OpImageSparseSampleProjDrefImplicitLod:
3130 case OpImageSparseSampleProjDrefExplicitLod:
3131 case OpImageGather:
3132 case OpImageSparseGather:
3133 case OpImageDrefGather:
3134 case OpImageSparseDrefGather:
3135 for (uint32_t i = 1; i < length; i++)
3136 {
3137 // Argument 5 is a literal.
3138 if (i != 5)
3139 notify_variable_access(args[i], current_block->self);
3140 }
3141 break;
3142
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003143 default:
3144 {
3145 // Rather dirty way of figuring out where Phi variables are used.
3146 // As long as only IDs are used, we can scan through instructions and try to find any evidence that
3147 // the ID of a variable has been used.
3148 // There are potential false positives here where a literal is used in-place of an ID,
3149 // but worst case, it does not affect the correctness of the compile.
3150 // Exhaustive analysis would be better here, but it's not worth it for now.
3151 for (uint32_t i = 0; i < length; i++)
3152 notify_variable_access(args[i], current_block->self);
3153 break;
3154 }
3155 }
3156 return true;
3157}
3158
Hans-Kristian Arntzen8c314112018-07-05 14:18:34 +02003159Compiler::StaticExpressionAccessHandler::StaticExpressionAccessHandler(Compiler &compiler_, uint32_t variable_id_)
3160 : compiler(compiler_)
3161 , variable_id(variable_id_)
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003162{
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003163}
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003164
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003165bool Compiler::StaticExpressionAccessHandler::follow_function_call(const SPIRFunction &)
3166{
3167 return false;
3168}
3169
3170bool Compiler::StaticExpressionAccessHandler::handle(spv::Op op, const uint32_t *args, uint32_t length)
3171{
3172 switch (op)
3173 {
3174 case OpStore:
3175 if (length < 2)
3176 return false;
3177 if (args[0] == variable_id)
3178 {
3179 static_expression = args[1];
3180 write_count++;
3181 }
3182 break;
3183
3184 case OpLoad:
3185 if (length < 3)
3186 return false;
3187 if (args[2] == variable_id && static_expression == 0) // Tried to read from variable before it was initialized.
3188 return false;
3189 break;
3190
3191 case OpAccessChain:
3192 case OpInBoundsAccessChain:
Chip Davis3bfb2f92018-12-03 02:06:33 -06003193 case OpPtrAccessChain:
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003194 if (length < 3)
3195 return false;
3196 if (args[2] == variable_id) // If we try to access chain our candidate variable before we store to it, bail.
3197 return false;
3198 break;
3199
3200 default:
3201 break;
3202 }
3203
3204 return true;
3205}
3206
Hans-Kristian Arntzen3e584f22019-02-06 10:38:18 +01003207void Compiler::find_function_local_luts(SPIRFunction &entry, const AnalyzeVariableScopeAccessHandler &handler,
3208 bool single_function)
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003209{
3210 auto &cfg = *function_cfgs.find(entry.self)->second;
3211
3212 // For each variable which is statically accessed.
3213 for (auto &accessed_var : handler.accessed_variables_to_block)
3214 {
3215 auto &blocks = accessed_var.second;
3216 auto &var = get<SPIRVariable>(accessed_var.first);
3217 auto &type = expression_type(accessed_var.first);
3218
3219 // Only consider function local variables here.
Hans-Kristian Arntzen3e584f22019-02-06 10:38:18 +01003220 // If we only have a single function in our CFG, private storage is also fine,
3221 // since it behaves like a function local variable.
3222 bool allow_lut = var.storage == StorageClassFunction || (single_function && var.storage == StorageClassPrivate);
3223 if (!allow_lut)
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003224 continue;
3225
3226 // We cannot be a phi variable.
3227 if (var.phi_variable)
3228 continue;
3229
3230 // Only consider arrays here.
3231 if (type.array.empty())
3232 continue;
3233
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003234 // If the variable has an initializer, make sure it is a constant expression.
3235 uint32_t static_constant_expression = 0;
3236 if (var.initializer)
3237 {
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003238 if (ir.ids[var.initializer].get_type() != TypeConstant)
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003239 continue;
3240 static_constant_expression = var.initializer;
3241
3242 // There can be no stores to this variable, we have now proved we have a LUT.
3243 if (handler.complete_write_variables_to_block.count(var.self) != 0 ||
3244 handler.partial_write_variables_to_block.count(var.self) != 0)
3245 continue;
3246 }
3247 else
3248 {
3249 // We can have one, and only one write to the variable, and that write needs to be a constant.
3250
3251 // No partial writes allowed.
3252 if (handler.partial_write_variables_to_block.count(var.self) != 0)
3253 continue;
3254
3255 auto itr = handler.complete_write_variables_to_block.find(var.self);
3256
3257 // No writes?
3258 if (itr == end(handler.complete_write_variables_to_block))
3259 continue;
3260
3261 // We write to the variable in more than one block.
3262 auto &write_blocks = itr->second;
3263 if (write_blocks.size() != 1)
3264 continue;
3265
3266 // The write needs to happen in the dominating block.
3267 DominatorBuilder builder(cfg);
3268 for (auto &block : blocks)
3269 builder.add_block(block);
3270 uint32_t dominator = builder.get_dominator();
3271
3272 // The complete write happened in a branch or similar, cannot deduce static expression.
3273 if (write_blocks.count(dominator) == 0)
3274 continue;
3275
3276 // Find the static expression for this variable.
3277 StaticExpressionAccessHandler static_expression_handler(*this, var.self);
3278 traverse_all_reachable_opcodes(get<SPIRBlock>(dominator), static_expression_handler);
3279
3280 // We want one, and exactly one write
3281 if (static_expression_handler.write_count != 1 || static_expression_handler.static_expression == 0)
3282 continue;
3283
3284 // Is it a constant expression?
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003285 if (ir.ids[static_expression_handler.static_expression].get_type() != TypeConstant)
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003286 continue;
3287
3288 // We found a LUT!
3289 static_constant_expression = static_expression_handler.static_expression;
3290 }
3291
3292 get<SPIRConstant>(static_constant_expression).is_used_as_lut = true;
3293 var.static_expression = static_constant_expression;
3294 var.statically_assigned = true;
3295 var.remapped_variable = true;
3296 }
3297}
3298
3299void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeAccessHandler &handler)
3300{
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003301 // First, we map out all variable access within a function.
3302 // Essentially a map of block -> { variables accessed in the basic block }
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003303 traverse_all_reachable_opcodes(entry, handler);
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003304
Hans-Kristian Arntzenc26c41b2018-07-04 17:26:53 +02003305 auto &cfg = *function_cfgs.find(entry.self)->second;
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003306
Hans-Kristian Arntzenb2c2e642017-03-25 16:25:30 +01003307 // Analyze if there are parameters which need to be implicitly preserved with an "in" qualifier.
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003308 analyze_parameter_preservation(entry, cfg, handler.accessed_variables_to_block,
3309 handler.complete_write_variables_to_block);
Hans-Kristian Arntzenb2c2e642017-03-25 16:25:30 +01003310
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003311 unordered_map<uint32_t, uint32_t> potential_loop_variables;
3312
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003313 // For each variable which is statically accessed.
3314 for (auto &var : handler.accessed_variables_to_block)
3315 {
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02003316 // Only deal with variables which are considered local variables in this function.
3317 if (find(begin(entry.local_variables), end(entry.local_variables), var.first) == end(entry.local_variables))
3318 continue;
3319
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003320 DominatorBuilder builder(cfg);
3321 auto &blocks = var.second;
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003322 auto &type = expression_type(var.first);
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003323
3324 // Figure out which block is dominating all accesses of those variables.
3325 for (auto &block : blocks)
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003326 {
Hans-Kristian Arntzena714d422016-12-16 12:43:12 +01003327 // If we're accessing a variable inside a continue block, this variable might be a loop variable.
3328 // We can only use loop variables with scalars, as we cannot track static expressions for vectors.
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003329 if (is_continue(block))
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003330 {
Hans-Kristian Arntzenb737d2b2017-12-05 17:40:23 +01003331 // Potentially awkward case to check for.
3332 // We might have a variable inside a loop, which is touched by the continue block,
3333 // but is not actually a loop variable.
3334 // The continue block is dominated by the inner part of the loop, which does not make sense in high-level
3335 // language output because it will be declared before the body,
3336 // so we will have to lift the dominator up to the relevant loop header instead.
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003337 builder.add_block(ir.continue_block_to_loop_header[block]);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003338
Hans-Kristian Arntzen922420e2018-03-07 14:54:11 +01003339 // Arrays or structs cannot be loop variables.
3340 if (type.vecsize == 1 && type.columns == 1 && type.basetype != SPIRType::Struct && type.array.empty())
Hans-Kristian Arntzenb737d2b2017-12-05 17:40:23 +01003341 {
3342 // The variable is used in multiple continue blocks, this is not a loop
3343 // candidate, signal that by setting block to -1u.
3344 auto &potential = potential_loop_variables[var.first];
3345
3346 if (potential == 0)
3347 potential = block;
3348 else
3349 potential = ~(0u);
3350 }
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003351 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003352 builder.add_block(block);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003353 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003354
Hans-Kristian Arntzen5ff11cc2016-11-18 16:45:11 +01003355 builder.lift_continue_block_dominator();
3356
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003357 // Add it to a per-block list of variables.
3358 uint32_t dominating_block = builder.get_dominator();
Hans-Kristian Arntzenb737d2b2017-12-05 17:40:23 +01003359
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003360 // If all blocks here are dead code, this will be 0, so the variable in question
3361 // will be completely eliminated.
3362 if (dominating_block)
3363 {
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003364 auto &block = get<SPIRBlock>(dominating_block);
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003365 block.dominated_variables.push_back(var.first);
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003366 get<SPIRVariable>(var.first).dominator = dominating_block;
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003367 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003368 }
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003369
Hans-Kristian Arntzend4e470b2018-01-12 10:42:39 +01003370 for (auto &var : handler.accessed_temporaries_to_block)
3371 {
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003372 auto itr = handler.result_id_to_type.find(var.first);
3373
3374 if (itr == end(handler.result_id_to_type))
3375 {
3376 // We found a false positive ID being used, ignore.
3377 // This should probably be an assert.
3378 continue;
3379 }
3380
Hans-Kristian Arntzend2cc43e2019-02-19 17:00:49 +01003381 // There is no point in doing domination analysis for opaque types.
3382 auto &type = get<SPIRType>(itr->second);
3383 if (type_is_opaque_value(type))
3384 continue;
3385
Hans-Kristian Arntzend4e470b2018-01-12 10:42:39 +01003386 DominatorBuilder builder(cfg);
Hans-Kristian Arntzen9fbd8b72018-03-12 14:58:40 +01003387 bool force_temporary = false;
Hans-Kristian Arntzend4e470b2018-01-12 10:42:39 +01003388
3389 // Figure out which block is dominating all accesses of those temporaries.
3390 auto &blocks = var.second;
3391 for (auto &block : blocks)
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003392 {
Hans-Kristian Arntzend4e470b2018-01-12 10:42:39 +01003393 builder.add_block(block);
3394
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003395 // If a temporary is used in more than one block, we might have to lift continue block
3396 // access up to loop header like we did for variables.
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003397 if (blocks.size() != 1 && is_continue(block))
Hans-Kristian Arntzene23c9ea2019-04-10 15:46:48 +02003398 {
3399 auto &loop_header_block = get<SPIRBlock>(ir.continue_block_to_loop_header[block]);
3400 assert(loop_header_block.merge == SPIRBlock::MergeLoop);
3401
3402 // Only relevant if the loop is not marked as complex.
3403 if (!loop_header_block.complex_continue)
3404 builder.add_block(loop_header_block.self);
3405 }
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003406 else if (blocks.size() != 1 && is_single_block_loop(block))
Hans-Kristian Arntzen9fbd8b72018-03-12 14:58:40 +01003407 {
3408 // Awkward case, because the loop header is also the continue block.
3409 force_temporary = true;
3410 }
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003411 }
Hans-Kristian Arntzend4e470b2018-01-12 10:42:39 +01003412
3413 uint32_t dominating_block = builder.get_dominator();
3414 if (dominating_block)
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003415 {
3416 // If we touch a variable in the dominating block, this is the expected setup.
3417 // SPIR-V normally mandates this, but we have extra cases for temporary use inside loops.
3418 bool first_use_is_dominator = blocks.count(dominating_block) != 0;
3419
Hans-Kristian Arntzen9fbd8b72018-03-12 14:58:40 +01003420 if (!first_use_is_dominator || force_temporary)
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003421 {
Hans-Kristian Arntzene23c9ea2019-04-10 15:46:48 +02003422 if (handler.access_chain_expressions.count(var.first))
3423 {
3424 // Exceptionally rare case.
3425 // We cannot declare temporaries of access chains (except on MSL perhaps with pointers).
3426 // Rather than do that, we force a complex loop to make sure access chains are created and consumed
3427 // in expected order.
3428 auto &loop_header_block = get<SPIRBlock>(dominating_block);
3429 assert(loop_header_block.merge == SPIRBlock::MergeLoop);
3430 loop_header_block.complex_continue = true;
3431 }
3432 else
3433 {
3434 // This should be very rare, but if we try to declare a temporary inside a loop,
3435 // and that temporary is used outside the loop as well (spirv-opt inliner likes this)
3436 // we should actually emit the temporary outside the loop.
3437 hoisted_temporaries.insert(var.first);
3438 forced_temporaries.insert(var.first);
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003439
Hans-Kristian Arntzene23c9ea2019-04-10 15:46:48 +02003440 auto &block_temporaries = get<SPIRBlock>(dominating_block).declare_temporary;
3441 block_temporaries.emplace_back(handler.result_id_to_type[var.first], var.first);
3442 }
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003443 }
Hans-Kristian Arntzenc1947aa2018-03-24 04:16:18 +01003444 else if (blocks.size() > 1)
3445 {
3446 // Keep track of the temporary as we might have to declare this temporary.
3447 // This can happen if the loop header dominates a temporary, but we have a complex fallback loop.
3448 // In this case, the header is actually inside the for (;;) {} block, and we have problems.
3449 // What we need to do is hoist the temporaries outside the for (;;) {} block in case the header block
3450 // declares the temporary.
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003451 auto &block_temporaries = get<SPIRBlock>(dominating_block).potential_declare_temporary;
Hans-Kristian Arntzenc1947aa2018-03-24 04:16:18 +01003452 block_temporaries.emplace_back(handler.result_id_to_type[var.first], var.first);
3453 }
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003454 }
Hans-Kristian Arntzend4e470b2018-01-12 10:42:39 +01003455 }
3456
Hans-Kristian Arntzenf0c50a62017-07-25 18:22:15 +02003457 unordered_set<uint32_t> seen_blocks;
3458
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003459 // Now, try to analyze whether or not these variables are actually loop variables.
3460 for (auto &loop_variable : potential_loop_variables)
3461 {
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003462 auto &var = get<SPIRVariable>(loop_variable.first);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003463 auto dominator = var.dominator;
3464 auto block = loop_variable.second;
3465
3466 // The variable was accessed in multiple continue blocks, ignore.
Graham Wihlidalfadc1f92017-01-05 20:14:53 +01003467 if (block == ~(0u) || block == 0)
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003468 continue;
3469
3470 // Dead code.
3471 if (dominator == 0)
3472 continue;
3473
3474 uint32_t header = 0;
3475
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003476 // Find the loop header for this block if we are a continue block.
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003477 {
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003478 auto itr = ir.continue_block_to_loop_header.find(block);
3479 if (itr != end(ir.continue_block_to_loop_header))
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003480 {
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003481 header = itr->second;
3482 }
3483 else if (get<SPIRBlock>(block).continue_block == block)
3484 {
3485 // Also check for self-referential continue block.
3486 header = block;
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003487 }
3488 }
3489
3490 assert(header);
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003491 auto &header_block = get<SPIRBlock>(header);
Hans-Kristian Arntzen4a7a3722018-01-23 20:27:43 +01003492 auto &blocks = handler.accessed_variables_to_block[loop_variable.first];
3493
3494 // If a loop variable is not used before the loop, it's probably not a loop variable.
3495 bool has_accessed_variable = blocks.count(header) != 0;
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003496
3497 // Now, there are two conditions we need to meet for the variable to be a loop variable.
3498 // 1. The dominating block must have a branch-free path to the loop header,
3499 // this way we statically know which expression should be part of the loop variable initializer.
3500
3501 // Walk from the dominator, if there is one straight edge connecting
3502 // dominator and loop header, we statically know the loop initializer.
3503 bool static_loop_init = true;
3504 while (dominator != header)
3505 {
Hans-Kristian Arntzen4a7a3722018-01-23 20:27:43 +01003506 if (blocks.count(dominator) != 0)
3507 has_accessed_variable = true;
3508
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003509 auto &succ = cfg.get_succeeding_edges(dominator);
3510 if (succ.size() != 1)
3511 {
3512 static_loop_init = false;
3513 break;
3514 }
3515
3516 auto &pred = cfg.get_preceding_edges(succ.front());
3517 if (pred.size() != 1 || pred.front() != dominator)
3518 {
3519 static_loop_init = false;
3520 break;
3521 }
3522
3523 dominator = succ.front();
3524 }
3525
Hans-Kristian Arntzen4a7a3722018-01-23 20:27:43 +01003526 if (!static_loop_init || !has_accessed_variable)
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003527 continue;
3528
3529 // The second condition we need to meet is that no access after the loop
3530 // merge can occur. Walk the CFG to see if we find anything.
Hans-Kristian Arntzenf0c50a62017-07-25 18:22:15 +02003531
3532 seen_blocks.clear();
3533 cfg.walk_from(seen_blocks, header_block.merge_block, [&](uint32_t walk_block) {
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003534 // We found a block which accesses the variable outside the loop.
3535 if (blocks.find(walk_block) != end(blocks))
3536 static_loop_init = false;
3537 });
3538
3539 if (!static_loop_init)
3540 continue;
3541
3542 // We have a loop variable.
3543 header_block.loop_variables.push_back(loop_variable.first);
Hans-Kristian Arntzen44b32162016-12-16 14:01:09 +01003544 // Need to sort here as variables come from an unordered container, and pushing stuff in wrong order
3545 // will break reproducability in regression runs.
3546 sort(begin(header_block.loop_variables), end(header_block.loop_variables));
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003547 get<SPIRVariable>(loop_variable.first).loop_variable = true;
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003548 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003549}
Hans-Kristian Arntzen016b1d82017-01-21 10:07:38 +01003550
Hans-Kristian Arntzen3a9b0452018-06-03 12:00:22 +02003551Bitset Compiler::get_buffer_block_flags(uint32_t id) const
3552{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003553 return ir.get_buffer_block_flags(get<SPIRVariable>(id));
Hans-Kristian Arntzen016b1d82017-01-21 10:07:38 +01003554}
Hans-Kristian Arntzen95409792017-01-21 12:29:20 +01003555
3556bool Compiler::get_common_basic_type(const SPIRType &type, SPIRType::BaseType &base_type)
3557{
3558 if (type.basetype == SPIRType::Struct)
3559 {
3560 base_type = SPIRType::Unknown;
3561 for (auto &member_type : type.member_types)
3562 {
3563 SPIRType::BaseType member_base;
3564 if (!get_common_basic_type(get<SPIRType>(member_type), member_base))
3565 return false;
3566
3567 if (base_type == SPIRType::Unknown)
3568 base_type = member_base;
3569 else if (base_type != member_base)
3570 return false;
3571 }
3572 return true;
3573 }
3574 else
3575 {
3576 base_type = type.basetype;
3577 return true;
3578 }
3579}
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003580
Hans-Kristian Arntzen719cf9d2018-03-13 14:05:33 +01003581void Compiler::ActiveBuiltinHandler::handle_builtin(const SPIRType &type, BuiltIn builtin,
3582 const Bitset &decoration_flags)
Hans-Kristian Arntzenfb3f92a2018-02-22 14:36:50 +01003583{
3584 // If used, we will need to explicitly declare a new array size for these builtins.
3585
3586 if (builtin == BuiltInClipDistance)
3587 {
3588 if (!type.array_size_literal[0])
3589 SPIRV_CROSS_THROW("Array size for ClipDistance must be a literal.");
3590 uint32_t array_size = type.array[0];
3591 if (array_size == 0)
3592 SPIRV_CROSS_THROW("Array size for ClipDistance must not be unsized.");
3593 compiler.clip_distance_count = array_size;
3594 }
3595 else if (builtin == BuiltInCullDistance)
3596 {
3597 if (!type.array_size_literal[0])
3598 SPIRV_CROSS_THROW("Array size for CullDistance must be a literal.");
3599 uint32_t array_size = type.array[0];
3600 if (array_size == 0)
3601 SPIRV_CROSS_THROW("Array size for CullDistance must not be unsized.");
3602 compiler.cull_distance_count = array_size;
3603 }
Hans-Kristian Arntzen3c1b1472018-03-01 12:30:55 +01003604 else if (builtin == BuiltInPosition)
3605 {
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003606 if (decoration_flags.get(DecorationInvariant))
Hans-Kristian Arntzen3c1b1472018-03-01 12:30:55 +01003607 compiler.position_invariant = true;
3608 }
Hans-Kristian Arntzenfb3f92a2018-02-22 14:36:50 +01003609}
3610
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003611bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args, uint32_t length)
3612{
3613 const auto add_if_builtin = [&](uint32_t id) {
3614 // Only handles variables here.
3615 // Builtins which are part of a block are handled in AccessChain.
3616 auto *var = compiler.maybe_get<SPIRVariable>(id);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003617 auto &decorations = compiler.ir.meta[id].decoration;
Hans-Kristian Arntzen3c1b1472018-03-01 12:30:55 +01003618 if (var && decorations.builtin)
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003619 {
3620 auto &type = compiler.get<SPIRType>(var->basetype);
3621 auto &flags =
Hans-Kristian Arntzen2ebe1a82017-03-06 16:50:46 +01003622 type.storage == StorageClassInput ? compiler.active_input_builtins : compiler.active_output_builtins;
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003623 flags.set(decorations.builtin_type);
Hans-Kristian Arntzen3c1b1472018-03-01 12:30:55 +01003624 handle_builtin(type, decorations.builtin_type, decorations.decoration_flags);
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003625 }
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003626 };
3627
3628 switch (opcode)
3629 {
3630 case OpStore:
3631 if (length < 1)
3632 return false;
3633
3634 add_if_builtin(args[0]);
3635 break;
3636
3637 case OpCopyMemory:
3638 if (length < 2)
3639 return false;
3640
3641 add_if_builtin(args[0]);
3642 add_if_builtin(args[1]);
3643 break;
3644
3645 case OpCopyObject:
3646 case OpLoad:
3647 if (length < 3)
3648 return false;
3649
3650 add_if_builtin(args[2]);
3651 break;
3652
Chip Davis3bfb2f92018-12-03 02:06:33 -06003653 case OpSelect:
3654 if (length < 5)
3655 return false;
3656
3657 add_if_builtin(args[3]);
3658 add_if_builtin(args[4]);
3659 break;
3660
3661 case OpPhi:
3662 {
3663 if (length < 2)
3664 return false;
3665
3666 uint32_t count = length - 2;
3667 args += 2;
3668 for (uint32_t i = 0; i < count; i += 2)
3669 add_if_builtin(args[i]);
3670 break;
3671 }
3672
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003673 case OpFunctionCall:
3674 {
3675 if (length < 3)
3676 return false;
3677
3678 uint32_t count = length - 3;
3679 args += 3;
3680 for (uint32_t i = 0; i < count; i++)
3681 add_if_builtin(args[i]);
3682 break;
3683 }
3684
3685 case OpAccessChain:
3686 case OpInBoundsAccessChain:
Chip Davis3bfb2f92018-12-03 02:06:33 -06003687 case OpPtrAccessChain:
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003688 {
3689 if (length < 4)
3690 return false;
3691
3692 // Only consider global variables, cannot consider variables in functions yet, or other
3693 // access chains as they have not been created yet.
3694 auto *var = compiler.maybe_get<SPIRVariable>(args[2]);
3695 if (!var)
3696 break;
3697
Hans-Kristian Arntzen945425e2017-08-15 10:23:04 +02003698 // Required if we access chain into builtins like gl_GlobalInvocationID.
3699 add_if_builtin(args[2]);
3700
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003701 // Start traversing type hierarchy at the proper non-pointer types.
Chip Davis3bfb2f92018-12-03 02:06:33 -06003702 auto *type = &compiler.get_variable_data_type(*var);
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003703
Hans-Kristian Arntzen2ebe1a82017-03-06 16:50:46 +01003704 auto &flags =
Chip Daviseb89c3a2019-02-03 23:58:46 -06003705 var->storage == StorageClassInput ? compiler.active_input_builtins : compiler.active_output_builtins;
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003706
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003707 uint32_t count = length - 3;
3708 args += 3;
3709 for (uint32_t i = 0; i < count; i++)
3710 {
Chip Davis3bfb2f92018-12-03 02:06:33 -06003711 // Pointers
3712 if (opcode == OpPtrAccessChain && i == 0)
3713 {
3714 type = &compiler.get<SPIRType>(type->parent_type);
3715 continue;
3716 }
3717
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003718 // Arrays
3719 if (!type->array.empty())
3720 {
3721 type = &compiler.get<SPIRType>(type->parent_type);
3722 }
3723 // Structs
3724 else if (type->basetype == SPIRType::Struct)
3725 {
3726 uint32_t index = compiler.get<SPIRConstant>(args[i]).scalar();
3727
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003728 if (index < uint32_t(compiler.ir.meta[type->self].members.size()))
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003729 {
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003730 auto &decorations = compiler.ir.meta[type->self].members[index];
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003731 if (decorations.builtin)
Hans-Kristian Arntzenfb3f92a2018-02-22 14:36:50 +01003732 {
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003733 flags.set(decorations.builtin_type);
Hans-Kristian Arntzen3c1b1472018-03-01 12:30:55 +01003734 handle_builtin(compiler.get<SPIRType>(type->member_types[index]), decorations.builtin_type,
3735 decorations.decoration_flags);
Hans-Kristian Arntzenfb3f92a2018-02-22 14:36:50 +01003736 }
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003737 }
3738
3739 type = &compiler.get<SPIRType>(type->member_types[index]);
3740 }
3741 else
3742 {
3743 // No point in traversing further. We won't find any extra builtins.
3744 break;
3745 }
3746 }
3747 break;
3748 }
3749
3750 default:
3751 break;
3752 }
3753
3754 return true;
3755}
3756
3757void Compiler::update_active_builtins()
3758{
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003759 active_input_builtins.reset();
3760 active_output_builtins.reset();
Hans-Kristian Arntzenfb3f92a2018-02-22 14:36:50 +01003761 cull_distance_count = 0;
3762 clip_distance_count = 0;
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003763 ActiveBuiltinHandler handler(*this);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003764 traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003765}
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003766
Bill Hollings192bdc92017-05-24 09:31:38 -04003767// Returns whether this shader uses a builtin of the storage class
3768bool Compiler::has_active_builtin(BuiltIn builtin, StorageClass storage)
3769{
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003770 const Bitset *flags;
Bill Hollings192bdc92017-05-24 09:31:38 -04003771 switch (storage)
3772 {
3773 case StorageClassInput:
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003774 flags = &active_input_builtins;
Bill Hollings192bdc92017-05-24 09:31:38 -04003775 break;
3776 case StorageClassOutput:
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003777 flags = &active_output_builtins;
Bill Hollings192bdc92017-05-24 09:31:38 -04003778 break;
3779
3780 default:
3781 return false;
3782 }
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003783 return flags->get(builtin);
Bill Hollings192bdc92017-05-24 09:31:38 -04003784}
3785
Hans-Kristian Arntzen18a594a2018-02-09 10:26:20 +01003786void Compiler::analyze_image_and_sampler_usage()
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003787{
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003788 CombinedImageSamplerDrefHandler dref_handler(*this);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003789 traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), dref_handler);
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003790
3791 CombinedImageSamplerUsageHandler handler(*this, dref_handler.dref_combined_samplers);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003792 traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003793 comparison_ids = move(handler.comparison_ids);
Hans-Kristian Arntzen18a594a2018-02-09 10:26:20 +01003794 need_subpass_input = handler.need_subpass_input;
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003795
3796 // Forward information from separate images and samplers into combined image samplers.
3797 for (auto &combined : combined_image_samplers)
3798 if (comparison_ids.count(combined.sampler_id))
3799 comparison_ids.insert(combined.combined_id);
3800}
3801
3802bool Compiler::CombinedImageSamplerDrefHandler::handle(spv::Op opcode, const uint32_t *args, uint32_t)
3803{
3804 // Mark all sampled images which are used with Dref.
3805 switch (opcode)
3806 {
3807 case OpImageSampleDrefExplicitLod:
3808 case OpImageSampleDrefImplicitLod:
3809 case OpImageSampleProjDrefExplicitLod:
3810 case OpImageSampleProjDrefImplicitLod:
3811 case OpImageSparseSampleProjDrefImplicitLod:
3812 case OpImageSparseSampleDrefImplicitLod:
3813 case OpImageSparseSampleProjDrefExplicitLod:
3814 case OpImageSparseSampleDrefExplicitLod:
3815 case OpImageDrefGather:
3816 case OpImageSparseDrefGather:
3817 dref_combined_samplers.insert(args[2]);
3818 return true;
3819
3820 default:
3821 break;
3822 }
3823
3824 return true;
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003825}
3826
Hans-Kristian Arntzenb5ed7062018-07-05 10:42:05 +02003827void Compiler::build_function_control_flow_graphs_and_analyze()
Hans-Kristian Arntzenc26c41b2018-07-04 17:26:53 +02003828{
3829 CFGBuilder handler(*this);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003830 handler.function_cfgs[ir.default_entry_point].reset(new CFG(*this, get<SPIRFunction>(ir.default_entry_point)));
3831 traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
Hans-Kristian Arntzenc26c41b2018-07-04 17:26:53 +02003832 function_cfgs = move(handler.function_cfgs);
Hans-Kristian Arntzen3e584f22019-02-06 10:38:18 +01003833 bool single_function = function_cfgs.size() <= 1;
Hans-Kristian Arntzenb5ed7062018-07-05 10:42:05 +02003834
3835 for (auto &f : function_cfgs)
3836 {
3837 auto &func = get<SPIRFunction>(f.first);
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003838 AnalyzeVariableScopeAccessHandler scope_handler(*this, func);
3839 analyze_variable_scope(func, scope_handler);
Hans-Kristian Arntzen3e584f22019-02-06 10:38:18 +01003840 find_function_local_luts(func, scope_handler, single_function);
Hans-Kristian Arntzenb5ed7062018-07-05 10:42:05 +02003841
3842 // Check if we can actually use the loop variables we found in analyze_variable_scope.
3843 // To use multiple initializers, we need the same type and qualifiers.
3844 for (auto block : func.blocks)
3845 {
3846 auto &b = get<SPIRBlock>(block);
3847 if (b.loop_variables.size() < 2)
3848 continue;
3849
3850 auto &flags = get_decoration_bitset(b.loop_variables.front());
3851 uint32_t type = get<SPIRVariable>(b.loop_variables.front()).basetype;
3852 bool invalid_initializers = false;
3853 for (auto loop_variable : b.loop_variables)
3854 {
3855 if (flags != get_decoration_bitset(loop_variable) ||
3856 type != get<SPIRVariable>(b.loop_variables.front()).basetype)
3857 {
3858 invalid_initializers = true;
3859 break;
3860 }
3861 }
3862
3863 if (invalid_initializers)
3864 {
3865 for (auto loop_variable : b.loop_variables)
3866 get<SPIRVariable>(loop_variable).loop_variable = false;
3867 b.loop_variables.clear();
3868 }
3869 }
3870 }
Hans-Kristian Arntzenc26c41b2018-07-04 17:26:53 +02003871}
3872
Hans-Kristian Arntzen9b92e682019-03-29 10:29:44 +01003873Compiler::CFGBuilder::CFGBuilder(Compiler &compiler_)
Hans-Kristian Arntzen8c314112018-07-05 14:18:34 +02003874 : compiler(compiler_)
Hans-Kristian Arntzenc26c41b2018-07-04 17:26:53 +02003875{
3876}
3877
3878bool Compiler::CFGBuilder::handle(spv::Op, const uint32_t *, uint32_t)
3879{
3880 return true;
3881}
3882
3883bool Compiler::CFGBuilder::follow_function_call(const SPIRFunction &func)
3884{
3885 if (function_cfgs.find(func.self) == end(function_cfgs))
3886 {
3887 function_cfgs[func.self].reset(new CFG(compiler, func));
3888 return true;
3889 }
3890 else
3891 return false;
3892}
3893
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003894bool Compiler::CombinedImageSamplerUsageHandler::begin_function_scope(const uint32_t *args, uint32_t length)
3895{
3896 if (length < 3)
3897 return false;
3898
3899 auto &func = compiler.get<SPIRFunction>(args[2]);
3900 const auto *arg = &args[3];
3901 length -= 3;
3902
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003903 for (uint32_t i = 0; i < length; i++)
3904 {
3905 auto &argument = func.arguments[i];
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003906 dependency_hierarchy[argument.id].insert(arg[i]);
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003907 }
3908
3909 return true;
3910}
3911
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003912void Compiler::CombinedImageSamplerUsageHandler::add_hierarchy_to_comparison_ids(uint32_t id)
Hans-Kristian Arntzen18a594a2018-02-09 10:26:20 +01003913{
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003914 // Traverse the variable dependency hierarchy and tag everything in its path with comparison ids.
3915 comparison_ids.insert(id);
3916 for (auto &dep_id : dependency_hierarchy[id])
3917 add_hierarchy_to_comparison_ids(dep_id);
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003918}
3919
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003920bool Compiler::CombinedImageSamplerUsageHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
3921{
3922 switch (opcode)
3923 {
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003924 case OpAccessChain:
3925 case OpInBoundsAccessChain:
Chip Davis3bfb2f92018-12-03 02:06:33 -06003926 case OpPtrAccessChain:
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003927 case OpLoad:
3928 {
3929 if (length < 3)
3930 return false;
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003931 dependency_hierarchy[args[1]].insert(args[2]);
Hans-Kristian Arntzen18a594a2018-02-09 10:26:20 +01003932
3933 // Ideally defer this to OpImageRead, but then we'd need to track loaded IDs.
3934 // If we load an image, we're going to use it and there is little harm in declaring an unused gl_FragCoord.
3935 auto &type = compiler.get<SPIRType>(args[0]);
3936 if (type.image.dim == DimSubpassData)
3937 need_subpass_input = true;
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003938
3939 // If we load a SampledImage and it will be used with Dref, propagate the state up.
3940 if (dref_combined_samplers.count(args[1]) != 0)
3941 add_hierarchy_to_comparison_ids(args[1]);
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003942 break;
3943 }
3944
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003945 case OpSampledImage:
3946 {
3947 if (length < 4)
3948 return false;
3949
3950 uint32_t result_type = args[0];
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003951 uint32_t result_id = args[1];
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003952 auto &type = compiler.get<SPIRType>(result_type);
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003953 if (type.image.depth || dref_combined_samplers.count(result_id) != 0)
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003954 {
Hans-Kristian Arntzen18a594a2018-02-09 10:26:20 +01003955 // This image must be a depth image.
3956 uint32_t image = args[2];
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003957 add_hierarchy_to_comparison_ids(image);
Hans-Kristian Arntzen18a594a2018-02-09 10:26:20 +01003958
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003959 // This sampler must be a SamplerComparisonState, and not a regular SamplerState.
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003960 uint32_t sampler = args[3];
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003961 add_hierarchy_to_comparison_ids(sampler);
3962
3963 // Mark the OpSampledImage itself as being comparison state.
3964 comparison_ids.insert(result_id);
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003965 }
3966 return true;
3967 }
3968
3969 default:
3970 break;
3971 }
3972
3973 return true;
3974}
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02003975
3976bool Compiler::buffer_is_hlsl_counter_buffer(uint32_t id) const
3977{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01003978 auto *m = ir.find_meta(id);
3979 return m && m->hlsl_is_magic_counter_buffer;
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02003980}
3981
3982bool Compiler::buffer_get_hlsl_counter_buffer(uint32_t id, uint32_t &counter_id) const
3983{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01003984 auto *m = ir.find_meta(id);
3985
Hans-Kristian Arntzen215d3ca2018-03-20 20:04:12 +01003986 // First, check for the proper decoration.
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01003987 if (m && m->hlsl_magic_counter_buffer != 0)
Hans-Kristian Arntzen215d3ca2018-03-20 20:04:12 +01003988 {
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01003989 counter_id = m->hlsl_magic_counter_buffer;
Hans-Kristian Arntzen215d3ca2018-03-20 20:04:12 +01003990 return true;
3991 }
Hans-Kristian Arntzen9aa623a2018-11-22 10:23:58 +01003992 else
3993 return false;
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02003994}
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003995
3996void Compiler::make_constant_null(uint32_t id, uint32_t type)
3997{
3998 auto &constant_type = get<SPIRType>(type);
3999
Chip Davis3bfb2f92018-12-03 02:06:33 -06004000 if (constant_type.pointer)
4001 {
4002 auto &constant = set<SPIRConstant>(id, type);
4003 constant.make_null(constant_type);
4004 }
4005 else if (!constant_type.array.empty())
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02004006 {
4007 assert(constant_type.parent_type);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02004008 uint32_t parent_id = ir.increase_bound_by(1);
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02004009 make_constant_null(parent_id, constant_type.parent_type);
4010
4011 if (!constant_type.array_size_literal.back())
4012 SPIRV_CROSS_THROW("Array size of OpConstantNull must be a literal.");
4013
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02004014 SmallVector<uint32_t> elements(constant_type.array.back());
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02004015 for (uint32_t i = 0; i < constant_type.array.back(); i++)
4016 elements[i] = parent_id;
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02004017 set<SPIRConstant>(id, type, elements.data(), uint32_t(elements.size()), false);
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02004018 }
4019 else if (!constant_type.member_types.empty())
4020 {
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02004021 uint32_t member_ids = ir.increase_bound_by(uint32_t(constant_type.member_types.size()));
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02004022 SmallVector<uint32_t> elements(constant_type.member_types.size());
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02004023 for (uint32_t i = 0; i < constant_type.member_types.size(); i++)
4024 {
4025 make_constant_null(member_ids + i, constant_type.member_types[i]);
4026 elements[i] = member_ids + i;
4027 }
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02004028 set<SPIRConstant>(id, type, elements.data(), uint32_t(elements.size()), false);
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02004029 }
4030 else
4031 {
4032 auto &constant = set<SPIRConstant>(id, type);
4033 constant.make_null(constant_type);
4034 }
4035}
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02004036
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02004037const SmallVector<spv::Capability> &Compiler::get_declared_capabilities() const
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02004038{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02004039 return ir.declared_capabilities;
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02004040}
4041
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02004042const SmallVector<std::string> &Compiler::get_declared_extensions() const
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02004043{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02004044 return ir.declared_extensions;
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02004045}
Hans-Kristian Arntzen2c90ea32017-12-01 14:20:51 +01004046
4047std::string Compiler::get_remapped_declared_block_name(uint32_t id) const
4048{
4049 auto itr = declared_block_names.find(id);
4050 if (itr != end(declared_block_names))
4051 return itr->second;
4052 else
4053 {
4054 auto &var = get<SPIRVariable>(id);
4055 auto &type = get<SPIRType>(var.basetype);
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01004056
4057 auto *type_meta = ir.find_meta(type.self);
4058 auto *block_name = type_meta ? &type_meta->decoration.alias : nullptr;
4059 return (!block_name || block_name->empty()) ? get_block_fallback_name(id) : *block_name;
Hans-Kristian Arntzen2c90ea32017-12-01 14:20:51 +01004060 }
4061}
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01004062
4063bool Compiler::instruction_to_result_type(uint32_t &result_type, uint32_t &result_id, spv::Op op, const uint32_t *args,
4064 uint32_t length)
4065{
4066 // Most instructions follow the pattern of <result-type> <result-id> <arguments>.
4067 // There are some exceptions.
4068 switch (op)
4069 {
4070 case OpStore:
4071 case OpCopyMemory:
4072 case OpCopyMemorySized:
4073 case OpImageWrite:
4074 case OpAtomicStore:
4075 case OpAtomicFlagClear:
4076 case OpEmitStreamVertex:
4077 case OpEndStreamPrimitive:
4078 case OpControlBarrier:
4079 case OpMemoryBarrier:
4080 case OpGroupWaitEvents:
4081 case OpRetainEvent:
4082 case OpReleaseEvent:
4083 case OpSetUserEventStatus:
4084 case OpCaptureEventProfilingInfo:
4085 case OpCommitReadPipe:
4086 case OpCommitWritePipe:
4087 case OpGroupCommitReadPipe:
4088 case OpGroupCommitWritePipe:
4089 return false;
4090
4091 default:
Hans-Kristian Arntzend2cc43e2019-02-19 17:00:49 +01004092 if (length > 1 && maybe_get<SPIRType>(args[0]) != nullptr)
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01004093 {
4094 result_type = args[0];
4095 result_id = args[1];
4096 return true;
4097 }
4098 else
4099 return false;
4100 }
4101}
Brad Davis6c88b002018-06-18 09:30:16 -07004102
4103Bitset Compiler::combined_decoration_for_member(const SPIRType &type, uint32_t index) const
4104{
4105 Bitset flags;
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01004106 auto *type_meta = ir.find_meta(type.self);
Brad Davis6c88b002018-06-18 09:30:16 -07004107
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01004108 if (type_meta)
4109 {
4110 auto &memb = type_meta->members;
4111 if (index >= memb.size())
4112 return flags;
4113 auto &dec = memb[index];
4114
4115 // If our type is a struct, traverse all the members as well recursively.
4116 flags.merge_or(dec.decoration_flags);
4117 for (uint32_t i = 0; i < type.member_types.size(); i++)
4118 flags.merge_or(combined_decoration_for_member(get<SPIRType>(type.member_types[i]), i));
4119 }
Brad Davis6c88b002018-06-18 09:30:16 -07004120
4121 return flags;
4122}
4123
Brad Davis76204002018-06-20 10:25:38 -07004124bool Compiler::is_desktop_only_format(spv::ImageFormat format)
Brad Davis6c88b002018-06-18 09:30:16 -07004125{
4126 switch (format)
4127 {
Brad Davis76204002018-06-20 10:25:38 -07004128 // Desktop-only formats
Brad Davis6c88b002018-06-18 09:30:16 -07004129 case ImageFormatR11fG11fB10f:
Brad Davis6c88b002018-06-18 09:30:16 -07004130 case ImageFormatR16f:
Brad Davis6c88b002018-06-18 09:30:16 -07004131 case ImageFormatRgb10A2:
Brad Davis6c88b002018-06-18 09:30:16 -07004132 case ImageFormatR8:
Brad Davis6c88b002018-06-18 09:30:16 -07004133 case ImageFormatRg8:
Brad Davis6c88b002018-06-18 09:30:16 -07004134 case ImageFormatR16:
Brad Davis6c88b002018-06-18 09:30:16 -07004135 case ImageFormatRg16:
Brad Davis6c88b002018-06-18 09:30:16 -07004136 case ImageFormatRgba16:
Brad Davis6c88b002018-06-18 09:30:16 -07004137 case ImageFormatR16Snorm:
Brad Davis6c88b002018-06-18 09:30:16 -07004138 case ImageFormatRg16Snorm:
Brad Davis6c88b002018-06-18 09:30:16 -07004139 case ImageFormatRgba16Snorm:
Brad Davis6c88b002018-06-18 09:30:16 -07004140 case ImageFormatR8Snorm:
Brad Davis6c88b002018-06-18 09:30:16 -07004141 case ImageFormatRg8Snorm:
Brad Davis6c88b002018-06-18 09:30:16 -07004142 case ImageFormatR8ui:
Brad Davis6c88b002018-06-18 09:30:16 -07004143 case ImageFormatRg8ui:
Brad Davis6c88b002018-06-18 09:30:16 -07004144 case ImageFormatR16ui:
Brad Davis6c88b002018-06-18 09:30:16 -07004145 case ImageFormatRgb10a2ui:
Brad Davis6c88b002018-06-18 09:30:16 -07004146 case ImageFormatR8i:
Brad Davis6c88b002018-06-18 09:30:16 -07004147 case ImageFormatRg8i:
Brad Davis6c88b002018-06-18 09:30:16 -07004148 case ImageFormatR16i:
Brad Davis76204002018-06-20 10:25:38 -07004149 return true;
Brad Davis6c88b002018-06-18 09:30:16 -07004150 default:
Brad Davis76204002018-06-20 10:25:38 -07004151 break;
Brad Davis6c88b002018-06-18 09:30:16 -07004152 }
Brad Davis76204002018-06-20 10:25:38 -07004153
4154 return false;
Brad Davis6c88b002018-06-18 09:30:16 -07004155}
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02004156
Hans-Kristian Arntzen9b92e682019-03-29 10:29:44 +01004157bool Compiler::image_is_comparison(const SPIRType &type, uint32_t id) const
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02004158{
4159 return type.image.depth || (comparison_ids.count(id) != 0);
4160}
Hans-Kristian Arntzend2cc43e2019-02-19 17:00:49 +01004161
Hans-Kristian Arntzen9b92e682019-03-29 10:29:44 +01004162bool Compiler::type_is_opaque_value(const SPIRType &type) const
Hans-Kristian Arntzend2cc43e2019-02-19 17:00:49 +01004163{
4164 return !type.pointer && (type.basetype == SPIRType::SampledImage || type.basetype == SPIRType::Image ||
4165 type.basetype == SPIRType::Sampler);
4166}
Hans-Kristian Arntzen317144a2019-04-05 12:06:10 +02004167
4168// Make these member functions so we can easily break on any force_recompile events.
4169void Compiler::force_recompile()
4170{
4171 is_force_recompile = true;
4172}
4173
4174bool Compiler::is_forcing_recompilation() const
4175{
4176 return is_force_recompile;
4177}
4178
4179void Compiler::clear_force_recompile()
4180{
4181 is_force_recompile = false;
4182}