blob: b5ce3ab87c2a95931ca4e796b15e550c966ecaf0 [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{
2877 if (id_is_phi_variable(id))
2878 accessed_variables_to_block[id].insert(block);
2879 else if (id_is_potential_temporary(id))
2880 accessed_temporaries_to_block[id].insert(block);
2881}
2882
2883bool Compiler::AnalyzeVariableScopeAccessHandler::id_is_phi_variable(uint32_t id) const
2884{
2885 if (id >= compiler.get_current_id_bound())
2886 return false;
2887 auto *var = compiler.maybe_get<SPIRVariable>(id);
2888 return var && var->phi_variable;
2889}
2890
2891bool Compiler::AnalyzeVariableScopeAccessHandler::id_is_potential_temporary(uint32_t id) const
2892{
2893 if (id >= compiler.get_current_id_bound())
2894 return false;
2895
2896 // Temporaries are not created before we start emitting code.
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002897 return compiler.ir.ids[id].empty() || (compiler.ir.ids[id].get_type() == TypeExpression);
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002898}
2899
2900bool Compiler::AnalyzeVariableScopeAccessHandler::handle(spv::Op op, const uint32_t *args, uint32_t length)
2901{
2902 // Keep track of the types of temporaries, so we can hoist them out as necessary.
2903 uint32_t result_type, result_id;
2904 if (compiler.instruction_to_result_type(result_type, result_id, op, args, length))
2905 result_id_to_type[result_id] = result_type;
2906
2907 switch (op)
2908 {
2909 case OpStore:
2910 {
2911 if (length < 2)
2912 return false;
2913
2914 uint32_t ptr = args[0];
2915 auto *var = compiler.maybe_get_backing_variable(ptr);
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002916
2917 // If we store through an access chain, we have a partial write.
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02002918 if (var)
2919 {
2920 accessed_variables_to_block[var->self].insert(current_block->self);
2921 if (var->self == ptr)
2922 complete_write_variables_to_block[var->self].insert(current_block->self);
2923 else
2924 partial_write_variables_to_block[var->self].insert(current_block->self);
2925 }
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002926
2927 // Might try to store a Phi variable here.
2928 notify_variable_access(args[1], current_block->self);
2929 break;
2930 }
2931
2932 case OpAccessChain:
2933 case OpInBoundsAccessChain:
Chip Davis3bfb2f92018-12-03 02:06:33 -06002934 case OpPtrAccessChain:
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002935 {
2936 if (length < 3)
2937 return false;
2938
2939 uint32_t ptr = args[2];
2940 auto *var = compiler.maybe_get<SPIRVariable>(ptr);
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02002941 if (var)
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002942 accessed_variables_to_block[var->self].insert(current_block->self);
2943
2944 for (uint32_t i = 3; i < length; i++)
2945 notify_variable_access(args[i], current_block->self);
2946
2947 // The result of an access chain is a fixed expression and is not really considered a temporary.
2948 auto &e = compiler.set<SPIRExpression>(args[1], "", args[0], true);
2949 auto *backing_variable = compiler.maybe_get_backing_variable(ptr);
2950 e.loaded_from = backing_variable ? backing_variable->self : 0;
2951
2952 // Other backends might use SPIRAccessChain for this later.
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02002953 compiler.ir.ids[args[1]].set_allow_type_rewrite();
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002954 break;
2955 }
2956
2957 case OpCopyMemory:
2958 {
2959 if (length < 2)
2960 return false;
2961
2962 uint32_t lhs = args[0];
2963 uint32_t rhs = args[1];
2964 auto *var = compiler.maybe_get_backing_variable(lhs);
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002965
2966 // If we store through an access chain, we have a partial write.
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02002967 if (var)
2968 {
2969 accessed_variables_to_block[var->self].insert(current_block->self);
2970 if (var->self == lhs)
2971 complete_write_variables_to_block[var->self].insert(current_block->self);
2972 else
2973 partial_write_variables_to_block[var->self].insert(current_block->self);
2974 }
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002975
2976 var = compiler.maybe_get_backing_variable(rhs);
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02002977 if (var)
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002978 accessed_variables_to_block[var->self].insert(current_block->self);
2979 break;
2980 }
2981
2982 case OpCopyObject:
2983 {
2984 if (length < 3)
2985 return false;
2986
2987 auto *var = compiler.maybe_get_backing_variable(args[2]);
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02002988 if (var)
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02002989 accessed_variables_to_block[var->self].insert(current_block->self);
2990
2991 // Might try to copy a Phi variable here.
2992 notify_variable_access(args[2], current_block->self);
2993 break;
2994 }
2995
2996 case OpLoad:
2997 {
2998 if (length < 3)
2999 return false;
3000 uint32_t ptr = args[2];
3001 auto *var = compiler.maybe_get_backing_variable(ptr);
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02003002 if (var)
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003003 accessed_variables_to_block[var->self].insert(current_block->self);
3004
3005 // Loaded value is a temporary.
3006 notify_variable_access(args[1], current_block->self);
3007 break;
3008 }
3009
3010 case OpFunctionCall:
3011 {
3012 if (length < 3)
3013 return false;
3014
3015 length -= 3;
3016 args += 3;
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02003017
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003018 for (uint32_t i = 0; i < length; i++)
3019 {
3020 auto *var = compiler.maybe_get_backing_variable(args[i]);
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02003021 if (var)
3022 {
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003023 accessed_variables_to_block[var->self].insert(current_block->self);
Hans-Kristian Arntzenb5ed7062018-07-05 10:42:05 +02003024 // Assume we can get partial writes to this variable.
3025 partial_write_variables_to_block[var->self].insert(current_block->self);
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02003026 }
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003027
3028 // Cannot easily prove if argument we pass to a function is completely written.
3029 // Usually, functions write to a dummy variable,
3030 // which is then copied to in full to the real argument.
3031
3032 // Might try to copy a Phi variable here.
3033 notify_variable_access(args[i], current_block->self);
3034 }
3035
3036 // Return value may be a temporary.
3037 notify_variable_access(args[1], current_block->self);
3038 break;
3039 }
3040
3041 case OpExtInst:
3042 {
3043 for (uint32_t i = 4; i < length; i++)
3044 notify_variable_access(args[i], current_block->self);
3045 notify_variable_access(args[1], current_block->self);
3046 break;
3047 }
3048
3049 case OpArrayLength:
Hans-Kristian Arntzend2cc43e2019-02-19 17:00:49 +01003050 // Uses literals, but cannot be a phi variable or temporary, so ignore.
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003051 break;
3052
3053 // Atomics shouldn't be able to access function-local variables.
3054 // Some GLSL builtins access a pointer.
3055
3056 case OpCompositeInsert:
3057 case OpVectorShuffle:
3058 // Specialize for opcode which contains literals.
3059 for (uint32_t i = 1; i < 4; i++)
3060 notify_variable_access(args[i], current_block->self);
3061 break;
3062
3063 case OpCompositeExtract:
3064 // Specialize for opcode which contains literals.
3065 for (uint32_t i = 1; i < 3; i++)
3066 notify_variable_access(args[i], current_block->self);
3067 break;
3068
Hans-Kristian Arntzend2cc43e2019-02-19 17:00:49 +01003069 case OpImageWrite:
3070 for (uint32_t i = 0; i < length; i++)
3071 {
3072 // Argument 3 is a literal.
3073 if (i != 3)
3074 notify_variable_access(args[i], current_block->self);
3075 }
3076 break;
3077
3078 case OpImageSampleImplicitLod:
3079 case OpImageSampleExplicitLod:
3080 case OpImageSparseSampleImplicitLod:
3081 case OpImageSparseSampleExplicitLod:
3082 case OpImageSampleProjImplicitLod:
3083 case OpImageSampleProjExplicitLod:
3084 case OpImageSparseSampleProjImplicitLod:
3085 case OpImageSparseSampleProjExplicitLod:
3086 case OpImageFetch:
3087 case OpImageSparseFetch:
3088 case OpImageRead:
3089 case OpImageSparseRead:
3090 for (uint32_t i = 1; i < length; i++)
3091 {
3092 // Argument 4 is a literal.
3093 if (i != 4)
3094 notify_variable_access(args[i], current_block->self);
3095 }
3096 break;
3097
3098 case OpImageSampleDrefImplicitLod:
3099 case OpImageSampleDrefExplicitLod:
3100 case OpImageSparseSampleDrefImplicitLod:
3101 case OpImageSparseSampleDrefExplicitLod:
3102 case OpImageSampleProjDrefImplicitLod:
3103 case OpImageSampleProjDrefExplicitLod:
3104 case OpImageSparseSampleProjDrefImplicitLod:
3105 case OpImageSparseSampleProjDrefExplicitLod:
3106 case OpImageGather:
3107 case OpImageSparseGather:
3108 case OpImageDrefGather:
3109 case OpImageSparseDrefGather:
3110 for (uint32_t i = 1; i < length; i++)
3111 {
3112 // Argument 5 is a literal.
3113 if (i != 5)
3114 notify_variable_access(args[i], current_block->self);
3115 }
3116 break;
3117
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003118 default:
3119 {
3120 // Rather dirty way of figuring out where Phi variables are used.
3121 // As long as only IDs are used, we can scan through instructions and try to find any evidence that
3122 // the ID of a variable has been used.
3123 // There are potential false positives here where a literal is used in-place of an ID,
3124 // but worst case, it does not affect the correctness of the compile.
3125 // Exhaustive analysis would be better here, but it's not worth it for now.
3126 for (uint32_t i = 0; i < length; i++)
3127 notify_variable_access(args[i], current_block->self);
3128 break;
3129 }
3130 }
3131 return true;
3132}
3133
Hans-Kristian Arntzen8c314112018-07-05 14:18:34 +02003134Compiler::StaticExpressionAccessHandler::StaticExpressionAccessHandler(Compiler &compiler_, uint32_t variable_id_)
3135 : compiler(compiler_)
3136 , variable_id(variable_id_)
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003137{
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003138}
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003139
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003140bool Compiler::StaticExpressionAccessHandler::follow_function_call(const SPIRFunction &)
3141{
3142 return false;
3143}
3144
3145bool Compiler::StaticExpressionAccessHandler::handle(spv::Op op, const uint32_t *args, uint32_t length)
3146{
3147 switch (op)
3148 {
3149 case OpStore:
3150 if (length < 2)
3151 return false;
3152 if (args[0] == variable_id)
3153 {
3154 static_expression = args[1];
3155 write_count++;
3156 }
3157 break;
3158
3159 case OpLoad:
3160 if (length < 3)
3161 return false;
3162 if (args[2] == variable_id && static_expression == 0) // Tried to read from variable before it was initialized.
3163 return false;
3164 break;
3165
3166 case OpAccessChain:
3167 case OpInBoundsAccessChain:
Chip Davis3bfb2f92018-12-03 02:06:33 -06003168 case OpPtrAccessChain:
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003169 if (length < 3)
3170 return false;
3171 if (args[2] == variable_id) // If we try to access chain our candidate variable before we store to it, bail.
3172 return false;
3173 break;
3174
3175 default:
3176 break;
3177 }
3178
3179 return true;
3180}
3181
Hans-Kristian Arntzen3e584f22019-02-06 10:38:18 +01003182void Compiler::find_function_local_luts(SPIRFunction &entry, const AnalyzeVariableScopeAccessHandler &handler,
3183 bool single_function)
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003184{
3185 auto &cfg = *function_cfgs.find(entry.self)->second;
3186
3187 // For each variable which is statically accessed.
3188 for (auto &accessed_var : handler.accessed_variables_to_block)
3189 {
3190 auto &blocks = accessed_var.second;
3191 auto &var = get<SPIRVariable>(accessed_var.first);
3192 auto &type = expression_type(accessed_var.first);
3193
3194 // Only consider function local variables here.
Hans-Kristian Arntzen3e584f22019-02-06 10:38:18 +01003195 // If we only have a single function in our CFG, private storage is also fine,
3196 // since it behaves like a function local variable.
3197 bool allow_lut = var.storage == StorageClassFunction || (single_function && var.storage == StorageClassPrivate);
3198 if (!allow_lut)
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003199 continue;
3200
3201 // We cannot be a phi variable.
3202 if (var.phi_variable)
3203 continue;
3204
3205 // Only consider arrays here.
3206 if (type.array.empty())
3207 continue;
3208
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003209 // If the variable has an initializer, make sure it is a constant expression.
3210 uint32_t static_constant_expression = 0;
3211 if (var.initializer)
3212 {
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003213 if (ir.ids[var.initializer].get_type() != TypeConstant)
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003214 continue;
3215 static_constant_expression = var.initializer;
3216
3217 // There can be no stores to this variable, we have now proved we have a LUT.
3218 if (handler.complete_write_variables_to_block.count(var.self) != 0 ||
3219 handler.partial_write_variables_to_block.count(var.self) != 0)
3220 continue;
3221 }
3222 else
3223 {
3224 // We can have one, and only one write to the variable, and that write needs to be a constant.
3225
3226 // No partial writes allowed.
3227 if (handler.partial_write_variables_to_block.count(var.self) != 0)
3228 continue;
3229
3230 auto itr = handler.complete_write_variables_to_block.find(var.self);
3231
3232 // No writes?
3233 if (itr == end(handler.complete_write_variables_to_block))
3234 continue;
3235
3236 // We write to the variable in more than one block.
3237 auto &write_blocks = itr->second;
3238 if (write_blocks.size() != 1)
3239 continue;
3240
3241 // The write needs to happen in the dominating block.
3242 DominatorBuilder builder(cfg);
3243 for (auto &block : blocks)
3244 builder.add_block(block);
3245 uint32_t dominator = builder.get_dominator();
3246
3247 // The complete write happened in a branch or similar, cannot deduce static expression.
3248 if (write_blocks.count(dominator) == 0)
3249 continue;
3250
3251 // Find the static expression for this variable.
3252 StaticExpressionAccessHandler static_expression_handler(*this, var.self);
3253 traverse_all_reachable_opcodes(get<SPIRBlock>(dominator), static_expression_handler);
3254
3255 // We want one, and exactly one write
3256 if (static_expression_handler.write_count != 1 || static_expression_handler.static_expression == 0)
3257 continue;
3258
3259 // Is it a constant expression?
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003260 if (ir.ids[static_expression_handler.static_expression].get_type() != TypeConstant)
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003261 continue;
3262
3263 // We found a LUT!
3264 static_constant_expression = static_expression_handler.static_expression;
3265 }
3266
3267 get<SPIRConstant>(static_constant_expression).is_used_as_lut = true;
3268 var.static_expression = static_constant_expression;
3269 var.statically_assigned = true;
3270 var.remapped_variable = true;
3271 }
3272}
3273
3274void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeAccessHandler &handler)
3275{
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003276 // First, we map out all variable access within a function.
3277 // Essentially a map of block -> { variables accessed in the basic block }
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003278 traverse_all_reachable_opcodes(entry, handler);
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003279
Hans-Kristian Arntzenc26c41b2018-07-04 17:26:53 +02003280 auto &cfg = *function_cfgs.find(entry.self)->second;
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003281
Hans-Kristian Arntzenb2c2e642017-03-25 16:25:30 +01003282 // Analyze if there are parameters which need to be implicitly preserved with an "in" qualifier.
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003283 analyze_parameter_preservation(entry, cfg, handler.accessed_variables_to_block,
3284 handler.complete_write_variables_to_block);
Hans-Kristian Arntzenb2c2e642017-03-25 16:25:30 +01003285
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003286 unordered_map<uint32_t, uint32_t> potential_loop_variables;
3287
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003288 // For each variable which is statically accessed.
3289 for (auto &var : handler.accessed_variables_to_block)
3290 {
Hans-Kristian Arntzen6fdadb92018-07-04 16:46:25 +02003291 // Only deal with variables which are considered local variables in this function.
3292 if (find(begin(entry.local_variables), end(entry.local_variables), var.first) == end(entry.local_variables))
3293 continue;
3294
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003295 DominatorBuilder builder(cfg);
3296 auto &blocks = var.second;
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003297 auto &type = expression_type(var.first);
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003298
3299 // Figure out which block is dominating all accesses of those variables.
3300 for (auto &block : blocks)
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003301 {
Hans-Kristian Arntzena714d422016-12-16 12:43:12 +01003302 // If we're accessing a variable inside a continue block, this variable might be a loop variable.
3303 // We can only use loop variables with scalars, as we cannot track static expressions for vectors.
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003304 if (is_continue(block))
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003305 {
Hans-Kristian Arntzenb737d2b2017-12-05 17:40:23 +01003306 // Potentially awkward case to check for.
3307 // We might have a variable inside a loop, which is touched by the continue block,
3308 // but is not actually a loop variable.
3309 // The continue block is dominated by the inner part of the loop, which does not make sense in high-level
3310 // language output because it will be declared before the body,
3311 // so we will have to lift the dominator up to the relevant loop header instead.
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003312 builder.add_block(ir.continue_block_to_loop_header[block]);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003313
Hans-Kristian Arntzen922420e2018-03-07 14:54:11 +01003314 // Arrays or structs cannot be loop variables.
3315 if (type.vecsize == 1 && type.columns == 1 && type.basetype != SPIRType::Struct && type.array.empty())
Hans-Kristian Arntzenb737d2b2017-12-05 17:40:23 +01003316 {
3317 // The variable is used in multiple continue blocks, this is not a loop
3318 // candidate, signal that by setting block to -1u.
3319 auto &potential = potential_loop_variables[var.first];
3320
3321 if (potential == 0)
3322 potential = block;
3323 else
3324 potential = ~(0u);
3325 }
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003326 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003327 builder.add_block(block);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003328 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003329
Hans-Kristian Arntzen5ff11cc2016-11-18 16:45:11 +01003330 builder.lift_continue_block_dominator();
3331
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003332 // Add it to a per-block list of variables.
3333 uint32_t dominating_block = builder.get_dominator();
Hans-Kristian Arntzenb737d2b2017-12-05 17:40:23 +01003334
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003335 // If all blocks here are dead code, this will be 0, so the variable in question
3336 // will be completely eliminated.
3337 if (dominating_block)
3338 {
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003339 auto &block = get<SPIRBlock>(dominating_block);
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003340 block.dominated_variables.push_back(var.first);
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003341 get<SPIRVariable>(var.first).dominator = dominating_block;
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003342 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003343 }
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003344
Hans-Kristian Arntzend4e470b2018-01-12 10:42:39 +01003345 for (auto &var : handler.accessed_temporaries_to_block)
3346 {
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003347 auto itr = handler.result_id_to_type.find(var.first);
3348
3349 if (itr == end(handler.result_id_to_type))
3350 {
3351 // We found a false positive ID being used, ignore.
3352 // This should probably be an assert.
3353 continue;
3354 }
3355
Hans-Kristian Arntzend2cc43e2019-02-19 17:00:49 +01003356 // There is no point in doing domination analysis for opaque types.
3357 auto &type = get<SPIRType>(itr->second);
3358 if (type_is_opaque_value(type))
3359 continue;
3360
Hans-Kristian Arntzend4e470b2018-01-12 10:42:39 +01003361 DominatorBuilder builder(cfg);
Hans-Kristian Arntzen9fbd8b72018-03-12 14:58:40 +01003362 bool force_temporary = false;
Hans-Kristian Arntzend4e470b2018-01-12 10:42:39 +01003363
3364 // Figure out which block is dominating all accesses of those temporaries.
3365 auto &blocks = var.second;
3366 for (auto &block : blocks)
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003367 {
Hans-Kristian Arntzend4e470b2018-01-12 10:42:39 +01003368 builder.add_block(block);
3369
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003370 // If a temporary is used in more than one block, we might have to lift continue block
3371 // access up to loop header like we did for variables.
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003372 if (blocks.size() != 1 && is_continue(block))
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003373 builder.add_block(ir.continue_block_to_loop_header[block]);
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003374 else if (blocks.size() != 1 && is_single_block_loop(block))
Hans-Kristian Arntzen9fbd8b72018-03-12 14:58:40 +01003375 {
3376 // Awkward case, because the loop header is also the continue block.
3377 force_temporary = true;
3378 }
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003379 }
Hans-Kristian Arntzend4e470b2018-01-12 10:42:39 +01003380
3381 uint32_t dominating_block = builder.get_dominator();
3382 if (dominating_block)
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003383 {
3384 // If we touch a variable in the dominating block, this is the expected setup.
3385 // SPIR-V normally mandates this, but we have extra cases for temporary use inside loops.
3386 bool first_use_is_dominator = blocks.count(dominating_block) != 0;
3387
Hans-Kristian Arntzen9fbd8b72018-03-12 14:58:40 +01003388 if (!first_use_is_dominator || force_temporary)
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003389 {
3390 // This should be very rare, but if we try to declare a temporary inside a loop,
3391 // and that temporary is used outside the loop as well (spirv-opt inliner likes this)
3392 // we should actually emit the temporary outside the loop.
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003393 hoisted_temporaries.insert(var.first);
3394 forced_temporaries.insert(var.first);
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003395
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003396 auto &block_temporaries = get<SPIRBlock>(dominating_block).declare_temporary;
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003397 block_temporaries.emplace_back(handler.result_id_to_type[var.first], var.first);
3398 }
Hans-Kristian Arntzenc1947aa2018-03-24 04:16:18 +01003399 else if (blocks.size() > 1)
3400 {
3401 // Keep track of the temporary as we might have to declare this temporary.
3402 // This can happen if the loop header dominates a temporary, but we have a complex fallback loop.
3403 // In this case, the header is actually inside the for (;;) {} block, and we have problems.
3404 // What we need to do is hoist the temporaries outside the for (;;) {} block in case the header block
3405 // declares the temporary.
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003406 auto &block_temporaries = get<SPIRBlock>(dominating_block).potential_declare_temporary;
Hans-Kristian Arntzenc1947aa2018-03-24 04:16:18 +01003407 block_temporaries.emplace_back(handler.result_id_to_type[var.first], var.first);
3408 }
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01003409 }
Hans-Kristian Arntzend4e470b2018-01-12 10:42:39 +01003410 }
3411
Hans-Kristian Arntzenf0c50a62017-07-25 18:22:15 +02003412 unordered_set<uint32_t> seen_blocks;
3413
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003414 // Now, try to analyze whether or not these variables are actually loop variables.
3415 for (auto &loop_variable : potential_loop_variables)
3416 {
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003417 auto &var = get<SPIRVariable>(loop_variable.first);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003418 auto dominator = var.dominator;
3419 auto block = loop_variable.second;
3420
3421 // The variable was accessed in multiple continue blocks, ignore.
Graham Wihlidalfadc1f92017-01-05 20:14:53 +01003422 if (block == ~(0u) || block == 0)
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003423 continue;
3424
3425 // Dead code.
3426 if (dominator == 0)
3427 continue;
3428
3429 uint32_t header = 0;
3430
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003431 // Find the loop header for this block if we are a continue block.
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003432 {
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003433 auto itr = ir.continue_block_to_loop_header.find(block);
3434 if (itr != end(ir.continue_block_to_loop_header))
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003435 {
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003436 header = itr->second;
3437 }
3438 else if (get<SPIRBlock>(block).continue_block == block)
3439 {
3440 // Also check for self-referential continue block.
3441 header = block;
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003442 }
3443 }
3444
3445 assert(header);
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003446 auto &header_block = get<SPIRBlock>(header);
Hans-Kristian Arntzen4a7a3722018-01-23 20:27:43 +01003447 auto &blocks = handler.accessed_variables_to_block[loop_variable.first];
3448
3449 // If a loop variable is not used before the loop, it's probably not a loop variable.
3450 bool has_accessed_variable = blocks.count(header) != 0;
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003451
3452 // Now, there are two conditions we need to meet for the variable to be a loop variable.
3453 // 1. The dominating block must have a branch-free path to the loop header,
3454 // this way we statically know which expression should be part of the loop variable initializer.
3455
3456 // Walk from the dominator, if there is one straight edge connecting
3457 // dominator and loop header, we statically know the loop initializer.
3458 bool static_loop_init = true;
3459 while (dominator != header)
3460 {
Hans-Kristian Arntzen4a7a3722018-01-23 20:27:43 +01003461 if (blocks.count(dominator) != 0)
3462 has_accessed_variable = true;
3463
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003464 auto &succ = cfg.get_succeeding_edges(dominator);
3465 if (succ.size() != 1)
3466 {
3467 static_loop_init = false;
3468 break;
3469 }
3470
3471 auto &pred = cfg.get_preceding_edges(succ.front());
3472 if (pred.size() != 1 || pred.front() != dominator)
3473 {
3474 static_loop_init = false;
3475 break;
3476 }
3477
3478 dominator = succ.front();
3479 }
3480
Hans-Kristian Arntzen4a7a3722018-01-23 20:27:43 +01003481 if (!static_loop_init || !has_accessed_variable)
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003482 continue;
3483
3484 // The second condition we need to meet is that no access after the loop
3485 // merge can occur. Walk the CFG to see if we find anything.
Hans-Kristian Arntzenf0c50a62017-07-25 18:22:15 +02003486
3487 seen_blocks.clear();
3488 cfg.walk_from(seen_blocks, header_block.merge_block, [&](uint32_t walk_block) {
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003489 // We found a block which accesses the variable outside the loop.
3490 if (blocks.find(walk_block) != end(blocks))
3491 static_loop_init = false;
3492 });
3493
3494 if (!static_loop_init)
3495 continue;
3496
3497 // We have a loop variable.
3498 header_block.loop_variables.push_back(loop_variable.first);
Hans-Kristian Arntzen44b32162016-12-16 14:01:09 +01003499 // Need to sort here as variables come from an unordered container, and pushing stuff in wrong order
3500 // will break reproducability in regression runs.
3501 sort(begin(header_block.loop_variables), end(header_block.loop_variables));
Hans-Kristian Arntzen72161292018-07-04 16:20:06 +02003502 get<SPIRVariable>(loop_variable.first).loop_variable = true;
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003503 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003504}
Hans-Kristian Arntzen016b1d82017-01-21 10:07:38 +01003505
Hans-Kristian Arntzen3a9b0452018-06-03 12:00:22 +02003506Bitset Compiler::get_buffer_block_flags(uint32_t id) const
3507{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003508 return ir.get_buffer_block_flags(get<SPIRVariable>(id));
Hans-Kristian Arntzen016b1d82017-01-21 10:07:38 +01003509}
Hans-Kristian Arntzen95409792017-01-21 12:29:20 +01003510
3511bool Compiler::get_common_basic_type(const SPIRType &type, SPIRType::BaseType &base_type)
3512{
3513 if (type.basetype == SPIRType::Struct)
3514 {
3515 base_type = SPIRType::Unknown;
3516 for (auto &member_type : type.member_types)
3517 {
3518 SPIRType::BaseType member_base;
3519 if (!get_common_basic_type(get<SPIRType>(member_type), member_base))
3520 return false;
3521
3522 if (base_type == SPIRType::Unknown)
3523 base_type = member_base;
3524 else if (base_type != member_base)
3525 return false;
3526 }
3527 return true;
3528 }
3529 else
3530 {
3531 base_type = type.basetype;
3532 return true;
3533 }
3534}
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003535
Hans-Kristian Arntzen719cf9d2018-03-13 14:05:33 +01003536void Compiler::ActiveBuiltinHandler::handle_builtin(const SPIRType &type, BuiltIn builtin,
3537 const Bitset &decoration_flags)
Hans-Kristian Arntzenfb3f92a2018-02-22 14:36:50 +01003538{
3539 // If used, we will need to explicitly declare a new array size for these builtins.
3540
3541 if (builtin == BuiltInClipDistance)
3542 {
3543 if (!type.array_size_literal[0])
3544 SPIRV_CROSS_THROW("Array size for ClipDistance must be a literal.");
3545 uint32_t array_size = type.array[0];
3546 if (array_size == 0)
3547 SPIRV_CROSS_THROW("Array size for ClipDistance must not be unsized.");
3548 compiler.clip_distance_count = array_size;
3549 }
3550 else if (builtin == BuiltInCullDistance)
3551 {
3552 if (!type.array_size_literal[0])
3553 SPIRV_CROSS_THROW("Array size for CullDistance must be a literal.");
3554 uint32_t array_size = type.array[0];
3555 if (array_size == 0)
3556 SPIRV_CROSS_THROW("Array size for CullDistance must not be unsized.");
3557 compiler.cull_distance_count = array_size;
3558 }
Hans-Kristian Arntzen3c1b1472018-03-01 12:30:55 +01003559 else if (builtin == BuiltInPosition)
3560 {
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003561 if (decoration_flags.get(DecorationInvariant))
Hans-Kristian Arntzen3c1b1472018-03-01 12:30:55 +01003562 compiler.position_invariant = true;
3563 }
Hans-Kristian Arntzenfb3f92a2018-02-22 14:36:50 +01003564}
3565
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003566bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args, uint32_t length)
3567{
3568 const auto add_if_builtin = [&](uint32_t id) {
3569 // Only handles variables here.
3570 // Builtins which are part of a block are handled in AccessChain.
3571 auto *var = compiler.maybe_get<SPIRVariable>(id);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003572 auto &decorations = compiler.ir.meta[id].decoration;
Hans-Kristian Arntzen3c1b1472018-03-01 12:30:55 +01003573 if (var && decorations.builtin)
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003574 {
3575 auto &type = compiler.get<SPIRType>(var->basetype);
3576 auto &flags =
Hans-Kristian Arntzen2ebe1a82017-03-06 16:50:46 +01003577 type.storage == StorageClassInput ? compiler.active_input_builtins : compiler.active_output_builtins;
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003578 flags.set(decorations.builtin_type);
Hans-Kristian Arntzen3c1b1472018-03-01 12:30:55 +01003579 handle_builtin(type, decorations.builtin_type, decorations.decoration_flags);
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003580 }
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003581 };
3582
3583 switch (opcode)
3584 {
3585 case OpStore:
3586 if (length < 1)
3587 return false;
3588
3589 add_if_builtin(args[0]);
3590 break;
3591
3592 case OpCopyMemory:
3593 if (length < 2)
3594 return false;
3595
3596 add_if_builtin(args[0]);
3597 add_if_builtin(args[1]);
3598 break;
3599
3600 case OpCopyObject:
3601 case OpLoad:
3602 if (length < 3)
3603 return false;
3604
3605 add_if_builtin(args[2]);
3606 break;
3607
Chip Davis3bfb2f92018-12-03 02:06:33 -06003608 case OpSelect:
3609 if (length < 5)
3610 return false;
3611
3612 add_if_builtin(args[3]);
3613 add_if_builtin(args[4]);
3614 break;
3615
3616 case OpPhi:
3617 {
3618 if (length < 2)
3619 return false;
3620
3621 uint32_t count = length - 2;
3622 args += 2;
3623 for (uint32_t i = 0; i < count; i += 2)
3624 add_if_builtin(args[i]);
3625 break;
3626 }
3627
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003628 case OpFunctionCall:
3629 {
3630 if (length < 3)
3631 return false;
3632
3633 uint32_t count = length - 3;
3634 args += 3;
3635 for (uint32_t i = 0; i < count; i++)
3636 add_if_builtin(args[i]);
3637 break;
3638 }
3639
3640 case OpAccessChain:
3641 case OpInBoundsAccessChain:
Chip Davis3bfb2f92018-12-03 02:06:33 -06003642 case OpPtrAccessChain:
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003643 {
3644 if (length < 4)
3645 return false;
3646
3647 // Only consider global variables, cannot consider variables in functions yet, or other
3648 // access chains as they have not been created yet.
3649 auto *var = compiler.maybe_get<SPIRVariable>(args[2]);
3650 if (!var)
3651 break;
3652
Hans-Kristian Arntzen945425e2017-08-15 10:23:04 +02003653 // Required if we access chain into builtins like gl_GlobalInvocationID.
3654 add_if_builtin(args[2]);
3655
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003656 // Start traversing type hierarchy at the proper non-pointer types.
Chip Davis3bfb2f92018-12-03 02:06:33 -06003657 auto *type = &compiler.get_variable_data_type(*var);
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003658
Hans-Kristian Arntzen2ebe1a82017-03-06 16:50:46 +01003659 auto &flags =
Chip Daviseb89c3a2019-02-03 23:58:46 -06003660 var->storage == StorageClassInput ? compiler.active_input_builtins : compiler.active_output_builtins;
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003661
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003662 uint32_t count = length - 3;
3663 args += 3;
3664 for (uint32_t i = 0; i < count; i++)
3665 {
Chip Davis3bfb2f92018-12-03 02:06:33 -06003666 // Pointers
3667 if (opcode == OpPtrAccessChain && i == 0)
3668 {
3669 type = &compiler.get<SPIRType>(type->parent_type);
3670 continue;
3671 }
3672
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003673 // Arrays
3674 if (!type->array.empty())
3675 {
3676 type = &compiler.get<SPIRType>(type->parent_type);
3677 }
3678 // Structs
3679 else if (type->basetype == SPIRType::Struct)
3680 {
3681 uint32_t index = compiler.get<SPIRConstant>(args[i]).scalar();
3682
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003683 if (index < uint32_t(compiler.ir.meta[type->self].members.size()))
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003684 {
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003685 auto &decorations = compiler.ir.meta[type->self].members[index];
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003686 if (decorations.builtin)
Hans-Kristian Arntzenfb3f92a2018-02-22 14:36:50 +01003687 {
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003688 flags.set(decorations.builtin_type);
Hans-Kristian Arntzen3c1b1472018-03-01 12:30:55 +01003689 handle_builtin(compiler.get<SPIRType>(type->member_types[index]), decorations.builtin_type,
3690 decorations.decoration_flags);
Hans-Kristian Arntzenfb3f92a2018-02-22 14:36:50 +01003691 }
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003692 }
3693
3694 type = &compiler.get<SPIRType>(type->member_types[index]);
3695 }
3696 else
3697 {
3698 // No point in traversing further. We won't find any extra builtins.
3699 break;
3700 }
3701 }
3702 break;
3703 }
3704
3705 default:
3706 break;
3707 }
3708
3709 return true;
3710}
3711
3712void Compiler::update_active_builtins()
3713{
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003714 active_input_builtins.reset();
3715 active_output_builtins.reset();
Hans-Kristian Arntzenfb3f92a2018-02-22 14:36:50 +01003716 cull_distance_count = 0;
3717 clip_distance_count = 0;
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003718 ActiveBuiltinHandler handler(*this);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003719 traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003720}
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003721
Bill Hollings192bdc92017-05-24 09:31:38 -04003722// Returns whether this shader uses a builtin of the storage class
3723bool Compiler::has_active_builtin(BuiltIn builtin, StorageClass storage)
3724{
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003725 const Bitset *flags;
Bill Hollings192bdc92017-05-24 09:31:38 -04003726 switch (storage)
3727 {
3728 case StorageClassInput:
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003729 flags = &active_input_builtins;
Bill Hollings192bdc92017-05-24 09:31:38 -04003730 break;
3731 case StorageClassOutput:
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003732 flags = &active_output_builtins;
Bill Hollings192bdc92017-05-24 09:31:38 -04003733 break;
3734
3735 default:
3736 return false;
3737 }
Hans-Kristian Arntzene8e58842018-03-12 13:09:25 +01003738 return flags->get(builtin);
Bill Hollings192bdc92017-05-24 09:31:38 -04003739}
3740
Hans-Kristian Arntzen18a594a2018-02-09 10:26:20 +01003741void Compiler::analyze_image_and_sampler_usage()
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003742{
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003743 CombinedImageSamplerDrefHandler dref_handler(*this);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003744 traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), dref_handler);
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003745
3746 CombinedImageSamplerUsageHandler handler(*this, dref_handler.dref_combined_samplers);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003747 traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003748 comparison_ids = move(handler.comparison_ids);
Hans-Kristian Arntzen18a594a2018-02-09 10:26:20 +01003749 need_subpass_input = handler.need_subpass_input;
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003750
3751 // Forward information from separate images and samplers into combined image samplers.
3752 for (auto &combined : combined_image_samplers)
3753 if (comparison_ids.count(combined.sampler_id))
3754 comparison_ids.insert(combined.combined_id);
3755}
3756
3757bool Compiler::CombinedImageSamplerDrefHandler::handle(spv::Op opcode, const uint32_t *args, uint32_t)
3758{
3759 // Mark all sampled images which are used with Dref.
3760 switch (opcode)
3761 {
3762 case OpImageSampleDrefExplicitLod:
3763 case OpImageSampleDrefImplicitLod:
3764 case OpImageSampleProjDrefExplicitLod:
3765 case OpImageSampleProjDrefImplicitLod:
3766 case OpImageSparseSampleProjDrefImplicitLod:
3767 case OpImageSparseSampleDrefImplicitLod:
3768 case OpImageSparseSampleProjDrefExplicitLod:
3769 case OpImageSparseSampleDrefExplicitLod:
3770 case OpImageDrefGather:
3771 case OpImageSparseDrefGather:
3772 dref_combined_samplers.insert(args[2]);
3773 return true;
3774
3775 default:
3776 break;
3777 }
3778
3779 return true;
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003780}
3781
Hans-Kristian Arntzenb5ed7062018-07-05 10:42:05 +02003782void Compiler::build_function_control_flow_graphs_and_analyze()
Hans-Kristian Arntzenc26c41b2018-07-04 17:26:53 +02003783{
3784 CFGBuilder handler(*this);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003785 handler.function_cfgs[ir.default_entry_point].reset(new CFG(*this, get<SPIRFunction>(ir.default_entry_point)));
3786 traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
Hans-Kristian Arntzenc26c41b2018-07-04 17:26:53 +02003787 function_cfgs = move(handler.function_cfgs);
Hans-Kristian Arntzen3e584f22019-02-06 10:38:18 +01003788 bool single_function = function_cfgs.size() <= 1;
Hans-Kristian Arntzenb5ed7062018-07-05 10:42:05 +02003789
3790 for (auto &f : function_cfgs)
3791 {
3792 auto &func = get<SPIRFunction>(f.first);
Hans-Kristian Arntzend29f48e2018-07-05 13:25:57 +02003793 AnalyzeVariableScopeAccessHandler scope_handler(*this, func);
3794 analyze_variable_scope(func, scope_handler);
Hans-Kristian Arntzen3e584f22019-02-06 10:38:18 +01003795 find_function_local_luts(func, scope_handler, single_function);
Hans-Kristian Arntzenb5ed7062018-07-05 10:42:05 +02003796
3797 // Check if we can actually use the loop variables we found in analyze_variable_scope.
3798 // To use multiple initializers, we need the same type and qualifiers.
3799 for (auto block : func.blocks)
3800 {
3801 auto &b = get<SPIRBlock>(block);
3802 if (b.loop_variables.size() < 2)
3803 continue;
3804
3805 auto &flags = get_decoration_bitset(b.loop_variables.front());
3806 uint32_t type = get<SPIRVariable>(b.loop_variables.front()).basetype;
3807 bool invalid_initializers = false;
3808 for (auto loop_variable : b.loop_variables)
3809 {
3810 if (flags != get_decoration_bitset(loop_variable) ||
3811 type != get<SPIRVariable>(b.loop_variables.front()).basetype)
3812 {
3813 invalid_initializers = true;
3814 break;
3815 }
3816 }
3817
3818 if (invalid_initializers)
3819 {
3820 for (auto loop_variable : b.loop_variables)
3821 get<SPIRVariable>(loop_variable).loop_variable = false;
3822 b.loop_variables.clear();
3823 }
3824 }
3825 }
Hans-Kristian Arntzenc26c41b2018-07-04 17:26:53 +02003826}
3827
Hans-Kristian Arntzen9b92e682019-03-29 10:29:44 +01003828Compiler::CFGBuilder::CFGBuilder(Compiler &compiler_)
Hans-Kristian Arntzen8c314112018-07-05 14:18:34 +02003829 : compiler(compiler_)
Hans-Kristian Arntzenc26c41b2018-07-04 17:26:53 +02003830{
3831}
3832
3833bool Compiler::CFGBuilder::handle(spv::Op, const uint32_t *, uint32_t)
3834{
3835 return true;
3836}
3837
3838bool Compiler::CFGBuilder::follow_function_call(const SPIRFunction &func)
3839{
3840 if (function_cfgs.find(func.self) == end(function_cfgs))
3841 {
3842 function_cfgs[func.self].reset(new CFG(compiler, func));
3843 return true;
3844 }
3845 else
3846 return false;
3847}
3848
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003849bool Compiler::CombinedImageSamplerUsageHandler::begin_function_scope(const uint32_t *args, uint32_t length)
3850{
3851 if (length < 3)
3852 return false;
3853
3854 auto &func = compiler.get<SPIRFunction>(args[2]);
3855 const auto *arg = &args[3];
3856 length -= 3;
3857
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003858 for (uint32_t i = 0; i < length; i++)
3859 {
3860 auto &argument = func.arguments[i];
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003861 dependency_hierarchy[argument.id].insert(arg[i]);
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003862 }
3863
3864 return true;
3865}
3866
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003867void Compiler::CombinedImageSamplerUsageHandler::add_hierarchy_to_comparison_ids(uint32_t id)
Hans-Kristian Arntzen18a594a2018-02-09 10:26:20 +01003868{
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003869 // Traverse the variable dependency hierarchy and tag everything in its path with comparison ids.
3870 comparison_ids.insert(id);
3871 for (auto &dep_id : dependency_hierarchy[id])
3872 add_hierarchy_to_comparison_ids(dep_id);
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003873}
3874
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003875bool Compiler::CombinedImageSamplerUsageHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
3876{
3877 switch (opcode)
3878 {
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003879 case OpAccessChain:
3880 case OpInBoundsAccessChain:
Chip Davis3bfb2f92018-12-03 02:06:33 -06003881 case OpPtrAccessChain:
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003882 case OpLoad:
3883 {
3884 if (length < 3)
3885 return false;
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003886 dependency_hierarchy[args[1]].insert(args[2]);
Hans-Kristian Arntzen18a594a2018-02-09 10:26:20 +01003887
3888 // Ideally defer this to OpImageRead, but then we'd need to track loaded IDs.
3889 // If we load an image, we're going to use it and there is little harm in declaring an unused gl_FragCoord.
3890 auto &type = compiler.get<SPIRType>(args[0]);
3891 if (type.image.dim == DimSubpassData)
3892 need_subpass_input = true;
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003893
3894 // If we load a SampledImage and it will be used with Dref, propagate the state up.
3895 if (dref_combined_samplers.count(args[1]) != 0)
3896 add_hierarchy_to_comparison_ids(args[1]);
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003897 break;
3898 }
3899
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003900 case OpSampledImage:
3901 {
3902 if (length < 4)
3903 return false;
3904
3905 uint32_t result_type = args[0];
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003906 uint32_t result_id = args[1];
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003907 auto &type = compiler.get<SPIRType>(result_type);
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003908 if (type.image.depth || dref_combined_samplers.count(result_id) != 0)
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003909 {
Hans-Kristian Arntzen18a594a2018-02-09 10:26:20 +01003910 // This image must be a depth image.
3911 uint32_t image = args[2];
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003912 add_hierarchy_to_comparison_ids(image);
Hans-Kristian Arntzen18a594a2018-02-09 10:26:20 +01003913
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003914 // This sampler must be a SamplerComparisonState, and not a regular SamplerState.
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003915 uint32_t sampler = args[3];
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02003916 add_hierarchy_to_comparison_ids(sampler);
3917
3918 // Mark the OpSampledImage itself as being comparison state.
3919 comparison_ids.insert(result_id);
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003920 }
3921 return true;
3922 }
3923
3924 default:
3925 break;
3926 }
3927
3928 return true;
3929}
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02003930
3931bool Compiler::buffer_is_hlsl_counter_buffer(uint32_t id) const
3932{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01003933 auto *m = ir.find_meta(id);
3934 return m && m->hlsl_is_magic_counter_buffer;
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02003935}
3936
3937bool Compiler::buffer_get_hlsl_counter_buffer(uint32_t id, uint32_t &counter_id) const
3938{
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01003939 auto *m = ir.find_meta(id);
3940
Hans-Kristian Arntzen215d3ca2018-03-20 20:04:12 +01003941 // First, check for the proper decoration.
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01003942 if (m && m->hlsl_magic_counter_buffer != 0)
Hans-Kristian Arntzen215d3ca2018-03-20 20:04:12 +01003943 {
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01003944 counter_id = m->hlsl_magic_counter_buffer;
Hans-Kristian Arntzen215d3ca2018-03-20 20:04:12 +01003945 return true;
3946 }
Hans-Kristian Arntzen9aa623a2018-11-22 10:23:58 +01003947 else
3948 return false;
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02003949}
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003950
3951void Compiler::make_constant_null(uint32_t id, uint32_t type)
3952{
3953 auto &constant_type = get<SPIRType>(type);
3954
Chip Davis3bfb2f92018-12-03 02:06:33 -06003955 if (constant_type.pointer)
3956 {
3957 auto &constant = set<SPIRConstant>(id, type);
3958 constant.make_null(constant_type);
3959 }
3960 else if (!constant_type.array.empty())
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003961 {
3962 assert(constant_type.parent_type);
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003963 uint32_t parent_id = ir.increase_bound_by(1);
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003964 make_constant_null(parent_id, constant_type.parent_type);
3965
3966 if (!constant_type.array_size_literal.back())
3967 SPIRV_CROSS_THROW("Array size of OpConstantNull must be a literal.");
3968
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02003969 SmallVector<uint32_t> elements(constant_type.array.back());
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003970 for (uint32_t i = 0; i < constant_type.array.back(); i++)
3971 elements[i] = parent_id;
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02003972 set<SPIRConstant>(id, type, elements.data(), uint32_t(elements.size()), false);
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003973 }
3974 else if (!constant_type.member_types.empty())
3975 {
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003976 uint32_t member_ids = ir.increase_bound_by(uint32_t(constant_type.member_types.size()));
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02003977 SmallVector<uint32_t> elements(constant_type.member_types.size());
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003978 for (uint32_t i = 0; i < constant_type.member_types.size(); i++)
3979 {
3980 make_constant_null(member_ids + i, constant_type.member_types[i]);
3981 elements[i] = member_ids + i;
3982 }
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02003983 set<SPIRConstant>(id, type, elements.data(), uint32_t(elements.size()), false);
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003984 }
3985 else
3986 {
3987 auto &constant = set<SPIRConstant>(id, type);
3988 constant.make_null(constant_type);
3989 }
3990}
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02003991
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02003992const SmallVector<spv::Capability> &Compiler::get_declared_capabilities() const
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02003993{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003994 return ir.declared_capabilities;
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02003995}
3996
Hans-Kristian Arntzena489ba72019-04-02 11:19:03 +02003997const SmallVector<std::string> &Compiler::get_declared_extensions() const
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02003998{
Hans-Kristian Arntzen5bcf02f2018-10-05 11:30:57 +02003999 return ir.declared_extensions;
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02004000}
Hans-Kristian Arntzen2c90ea32017-12-01 14:20:51 +01004001
4002std::string Compiler::get_remapped_declared_block_name(uint32_t id) const
4003{
4004 auto itr = declared_block_names.find(id);
4005 if (itr != end(declared_block_names))
4006 return itr->second;
4007 else
4008 {
4009 auto &var = get<SPIRVariable>(id);
4010 auto &type = get<SPIRType>(var.basetype);
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01004011
4012 auto *type_meta = ir.find_meta(type.self);
4013 auto *block_name = type_meta ? &type_meta->decoration.alias : nullptr;
4014 return (!block_name || block_name->empty()) ? get_block_fallback_name(id) : *block_name;
Hans-Kristian Arntzen2c90ea32017-12-01 14:20:51 +01004015 }
4016}
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01004017
4018bool Compiler::instruction_to_result_type(uint32_t &result_type, uint32_t &result_id, spv::Op op, const uint32_t *args,
4019 uint32_t length)
4020{
4021 // Most instructions follow the pattern of <result-type> <result-id> <arguments>.
4022 // There are some exceptions.
4023 switch (op)
4024 {
4025 case OpStore:
4026 case OpCopyMemory:
4027 case OpCopyMemorySized:
4028 case OpImageWrite:
4029 case OpAtomicStore:
4030 case OpAtomicFlagClear:
4031 case OpEmitStreamVertex:
4032 case OpEndStreamPrimitive:
4033 case OpControlBarrier:
4034 case OpMemoryBarrier:
4035 case OpGroupWaitEvents:
4036 case OpRetainEvent:
4037 case OpReleaseEvent:
4038 case OpSetUserEventStatus:
4039 case OpCaptureEventProfilingInfo:
4040 case OpCommitReadPipe:
4041 case OpCommitWritePipe:
4042 case OpGroupCommitReadPipe:
4043 case OpGroupCommitWritePipe:
4044 return false;
4045
4046 default:
Hans-Kristian Arntzend2cc43e2019-02-19 17:00:49 +01004047 if (length > 1 && maybe_get<SPIRType>(args[0]) != nullptr)
Hans-Kristian Arntzen7d223b82018-01-18 12:07:10 +01004048 {
4049 result_type = args[0];
4050 result_id = args[1];
4051 return true;
4052 }
4053 else
4054 return false;
4055 }
4056}
Brad Davis6c88b002018-06-18 09:30:16 -07004057
4058Bitset Compiler::combined_decoration_for_member(const SPIRType &type, uint32_t index) const
4059{
4060 Bitset flags;
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01004061 auto *type_meta = ir.find_meta(type.self);
Brad Davis6c88b002018-06-18 09:30:16 -07004062
Hans-Kristian Arntzenb6298782019-01-10 14:04:01 +01004063 if (type_meta)
4064 {
4065 auto &memb = type_meta->members;
4066 if (index >= memb.size())
4067 return flags;
4068 auto &dec = memb[index];
4069
4070 // If our type is a struct, traverse all the members as well recursively.
4071 flags.merge_or(dec.decoration_flags);
4072 for (uint32_t i = 0; i < type.member_types.size(); i++)
4073 flags.merge_or(combined_decoration_for_member(get<SPIRType>(type.member_types[i]), i));
4074 }
Brad Davis6c88b002018-06-18 09:30:16 -07004075
4076 return flags;
4077}
4078
Brad Davis76204002018-06-20 10:25:38 -07004079bool Compiler::is_desktop_only_format(spv::ImageFormat format)
Brad Davis6c88b002018-06-18 09:30:16 -07004080{
4081 switch (format)
4082 {
Brad Davis76204002018-06-20 10:25:38 -07004083 // Desktop-only formats
Brad Davis6c88b002018-06-18 09:30:16 -07004084 case ImageFormatR11fG11fB10f:
Brad Davis6c88b002018-06-18 09:30:16 -07004085 case ImageFormatR16f:
Brad Davis6c88b002018-06-18 09:30:16 -07004086 case ImageFormatRgb10A2:
Brad Davis6c88b002018-06-18 09:30:16 -07004087 case ImageFormatR8:
Brad Davis6c88b002018-06-18 09:30:16 -07004088 case ImageFormatRg8:
Brad Davis6c88b002018-06-18 09:30:16 -07004089 case ImageFormatR16:
Brad Davis6c88b002018-06-18 09:30:16 -07004090 case ImageFormatRg16:
Brad Davis6c88b002018-06-18 09:30:16 -07004091 case ImageFormatRgba16:
Brad Davis6c88b002018-06-18 09:30:16 -07004092 case ImageFormatR16Snorm:
Brad Davis6c88b002018-06-18 09:30:16 -07004093 case ImageFormatRg16Snorm:
Brad Davis6c88b002018-06-18 09:30:16 -07004094 case ImageFormatRgba16Snorm:
Brad Davis6c88b002018-06-18 09:30:16 -07004095 case ImageFormatR8Snorm:
Brad Davis6c88b002018-06-18 09:30:16 -07004096 case ImageFormatRg8Snorm:
Brad Davis6c88b002018-06-18 09:30:16 -07004097 case ImageFormatR8ui:
Brad Davis6c88b002018-06-18 09:30:16 -07004098 case ImageFormatRg8ui:
Brad Davis6c88b002018-06-18 09:30:16 -07004099 case ImageFormatR16ui:
Brad Davis6c88b002018-06-18 09:30:16 -07004100 case ImageFormatRgb10a2ui:
Brad Davis6c88b002018-06-18 09:30:16 -07004101 case ImageFormatR8i:
Brad Davis6c88b002018-06-18 09:30:16 -07004102 case ImageFormatRg8i:
Brad Davis6c88b002018-06-18 09:30:16 -07004103 case ImageFormatR16i:
Brad Davis76204002018-06-20 10:25:38 -07004104 return true;
Brad Davis6c88b002018-06-18 09:30:16 -07004105 default:
Brad Davis76204002018-06-20 10:25:38 -07004106 break;
Brad Davis6c88b002018-06-18 09:30:16 -07004107 }
Brad Davis76204002018-06-20 10:25:38 -07004108
4109 return false;
Brad Davis6c88b002018-06-18 09:30:16 -07004110}
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02004111
Hans-Kristian Arntzen9b92e682019-03-29 10:29:44 +01004112bool Compiler::image_is_comparison(const SPIRType &type, uint32_t id) const
Hans-Kristian Arntzene0447322018-07-04 14:25:10 +02004113{
4114 return type.image.depth || (comparison_ids.count(id) != 0);
4115}
Hans-Kristian Arntzend2cc43e2019-02-19 17:00:49 +01004116
Hans-Kristian Arntzen9b92e682019-03-29 10:29:44 +01004117bool Compiler::type_is_opaque_value(const SPIRType &type) const
Hans-Kristian Arntzend2cc43e2019-02-19 17:00:49 +01004118{
4119 return !type.pointer && (type.basetype == SPIRType::SampledImage || type.basetype == SPIRType::Image ||
4120 type.basetype == SPIRType::Sampler);
4121}
Hans-Kristian Arntzen317144a2019-04-05 12:06:10 +02004122
4123// Make these member functions so we can easily break on any force_recompile events.
4124void Compiler::force_recompile()
4125{
4126 is_force_recompile = true;
4127}
4128
4129bool Compiler::is_forcing_recompilation() const
4130{
4131 return is_force_recompile;
4132}
4133
4134void Compiler::clear_force_recompile()
4135{
4136 is_force_recompile = false;
4137}