blob: 4cf844f292e7ac4327a21f4c4aa53e6d70fdf88f [file] [log] [blame]
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001/*
Hans-Kristian Arntzen18c37bc2017-01-28 09:00:40 +01002 * Copyright 2015-2017 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 Arntzen75471fb2016-03-02 18:09:16 +010020#include <algorithm>
Hans-Kristian Arntzen5ea59bd2016-05-23 13:30:02 +020021#include <cstring>
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010022#include <utility>
23
24using namespace std;
25using namespace spv;
Hans-Kristian Arntzen147e53a2016-04-04 09:36:04 +020026using namespace spirv_cross;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010027
28#define log(...) fprintf(stderr, __VA_ARGS__)
29
Hans-Kristian Arntzen94ff3552017-10-10 17:32:26 +020030static string ensure_valid_identifier(const string &name, bool member)
Hans-Kristian Arntzenfd432f82017-03-18 10:52:41 +010031{
32 // Functions in glslangValidator are mangled with name(<mangled> stuff.
33 // Normally, we would never see '(' in any legal identifiers, so just strip them out.
34 auto str = name.substr(0, name.find('('));
35
36 for (uint32_t i = 0; i < str.size(); i++)
37 {
38 auto &c = str[i];
39
Hans-Kristian Arntzen94ff3552017-10-10 17:32:26 +020040 if (member)
41 {
42 // _m<num> variables are reserved by the internal implementation,
43 // otherwise, make sure the name is a valid identifier.
44 if (i == 0)
45 c = isalpha(c) ? c : '_';
46 else if (i == 2 && str[0] == '_' && str[1] == 'm')
47 c = isalpha(c) ? c : '_';
48 else
49 c = isalnum(c) ? c : '_';
50 }
Hans-Kristian Arntzenfd432f82017-03-18 10:52:41 +010051 else
Hans-Kristian Arntzen94ff3552017-10-10 17:32:26 +020052 {
53 // _<num> variables are reserved by the internal implementation,
54 // otherwise, make sure the name is a valid identifier.
55 if (i == 0 || (str[0] == '_' && i == 1))
56 c = isalpha(c) ? c : '_';
57 else
58 c = isalnum(c) ? c : '_';
59 }
Hans-Kristian Arntzenfd432f82017-03-18 10:52:41 +010060 }
61 return str;
62}
63
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010064Instruction::Instruction(const vector<uint32_t> &spirv, uint32_t &index)
65{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020066 op = spirv[index] & 0xffff;
67 count = (spirv[index] >> 16) & 0xffff;
Hans-Kristian Arntzen416566b2016-07-08 10:47:03 +020068
69 if (count == 0)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +010070 SPIRV_CROSS_THROW("SPIR-V instructions cannot consume 0 words. Invalid SPIR-V file.");
Hans-Kristian Arntzen416566b2016-07-08 10:47:03 +020071
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020072 offset = index + 1;
73 length = count - 1;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010074
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020075 index += count;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010076
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020077 if (index > spirv.size())
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +010078 SPIRV_CROSS_THROW("SPIR-V instruction goes out of bounds.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010079}
80
81Compiler::Compiler(vector<uint32_t> ir)
82 : spirv(move(ir))
83{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020084 parse();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010085}
86
Yuriy O'Donnellae8de512017-04-01 12:31:34 +020087Compiler::Compiler(const uint32_t *ir, size_t word_count)
88 : spirv(ir, ir + word_count)
89{
90 parse();
91}
92
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020093string Compiler::compile()
94{
Endre Oma6ad8b302017-01-11 15:57:05 +010095 // Force a classic "C" locale, reverts when function returns
96 ClassicLocale classic_locale;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020097 return "";
98}
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010099
100bool Compiler::variable_storage_is_aliased(const SPIRVariable &v)
101{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200102 auto &type = get<SPIRType>(v.basetype);
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +0200103 bool ssbo = v.storage == StorageClassStorageBuffer ||
104 ((meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)) != 0);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200105 bool image = type.basetype == SPIRType::Image;
106 bool counter = type.basetype == SPIRType::AtomicCounter;
Hans-Kristian Arntzen7d8add32016-07-12 15:00:10 +0200107 bool is_restrict = (meta[v.self].decoration.decoration_flags & (1ull << DecorationRestrict)) != 0;
108 return !is_restrict && (ssbo || image || counter);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100109}
110
111bool Compiler::block_is_pure(const SPIRBlock &block)
112{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200113 for (auto &i : block.ops)
114 {
115 auto ops = stream(i);
116 auto op = static_cast<Op>(i.op);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100117
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200118 switch (op)
119 {
120 case OpFunctionCall:
121 {
122 uint32_t func = ops[2];
123 if (!function_is_pure(get<SPIRFunction>(func)))
124 return false;
125 break;
126 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100127
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +0100128 case OpCopyMemory:
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200129 case OpStore:
130 {
131 auto &type = expression_type(ops[0]);
132 if (type.storage != StorageClassFunction)
133 return false;
134 break;
135 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100136
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200137 case OpImageWrite:
138 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100139
Hans-Kristian Arntzen5af1a512016-05-05 09:51:42 +0200140 // Atomics are impure.
141 case OpAtomicLoad:
142 case OpAtomicStore:
143 case OpAtomicExchange:
144 case OpAtomicCompareExchange:
Bill Hollings8f6df772017-05-19 18:14:08 -0400145 case OpAtomicCompareExchangeWeak:
Hans-Kristian Arntzen5af1a512016-05-05 09:51:42 +0200146 case OpAtomicIIncrement:
147 case OpAtomicIDecrement:
148 case OpAtomicIAdd:
149 case OpAtomicISub:
150 case OpAtomicSMin:
151 case OpAtomicUMin:
152 case OpAtomicSMax:
153 case OpAtomicUMax:
154 case OpAtomicAnd:
155 case OpAtomicOr:
156 case OpAtomicXor:
157 return false;
158
159 // Geometry shader builtins modify global state.
160 case OpEndPrimitive:
161 case OpEmitStreamVertex:
162 case OpEndStreamPrimitive:
163 case OpEmitVertex:
164 return false;
165
166 // Barriers disallow any reordering, so we should treat blocks with barrier as writing.
167 case OpControlBarrier:
168 case OpMemoryBarrier:
169 return false;
170
Hans-Kristian Arntzen1079e792017-10-10 10:22:40 +0200171 // OpExtInst is potentially impure depending on extension, but GLSL builtins are at least pure.
Hans-Kristian Arntzen5af1a512016-05-05 09:51:42 +0200172
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200173 default:
174 break;
175 }
176 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100177
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200178 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100179}
180
Hans-Kristian Arntzen61c31c62017-03-07 13:27:04 +0100181string Compiler::to_name(uint32_t id, bool allow_alias) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100182{
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +0200183 if (allow_alias && ids.at(id).get_type() == TypeType)
184 {
185 // If this type is a simple alias, emit the
186 // name of the original type instead.
187 // We don't want to override the meta alias
188 // as that can be overridden by the reflection APIs after parse.
189 auto &type = get<SPIRType>(id);
190 if (type.type_alias)
191 return to_name(type.type_alias);
192 }
193
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200194 if (meta[id].decoration.alias.empty())
195 return join("_", id);
196 else
197 return meta.at(id).decoration.alias;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100198}
199
200bool Compiler::function_is_pure(const SPIRFunction &func)
201{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200202 for (auto block : func.blocks)
203 {
204 if (!block_is_pure(get<SPIRBlock>(block)))
205 {
206 //fprintf(stderr, "Function %s is impure!\n", to_name(func.self).c_str());
207 return false;
208 }
209 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100210
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200211 //fprintf(stderr, "Function %s is pure!\n", to_name(func.self).c_str());
212 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100213}
214
215void Compiler::register_global_read_dependencies(const SPIRBlock &block, uint32_t id)
216{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200217 for (auto &i : block.ops)
218 {
219 auto ops = stream(i);
220 auto op = static_cast<Op>(i.op);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100221
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200222 switch (op)
223 {
224 case OpFunctionCall:
225 {
226 uint32_t func = ops[2];
227 register_global_read_dependencies(get<SPIRFunction>(func), id);
228 break;
229 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100230
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200231 case OpLoad:
232 case OpImageRead:
233 {
234 // If we're in a storage class which does not get invalidated, adding dependencies here is no big deal.
235 auto *var = maybe_get_backing_variable(ops[2]);
236 if (var && var->storage != StorageClassFunction)
237 {
238 auto &type = get<SPIRType>(var->basetype);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100239
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200240 // InputTargets are immutable.
241 if (type.basetype != SPIRType::Image && type.image.dim != DimSubpassData)
242 var->dependees.push_back(id);
243 }
244 break;
245 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100246
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200247 default:
248 break;
249 }
250 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100251}
252
253void Compiler::register_global_read_dependencies(const SPIRFunction &func, uint32_t id)
254{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200255 for (auto block : func.blocks)
256 register_global_read_dependencies(get<SPIRBlock>(block), id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100257}
258
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200259SPIRVariable *Compiler::maybe_get_backing_variable(uint32_t chain)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100260{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200261 auto *var = maybe_get<SPIRVariable>(chain);
262 if (!var)
263 {
264 auto *cexpr = maybe_get<SPIRExpression>(chain);
265 if (cexpr)
266 var = maybe_get<SPIRVariable>(cexpr->loaded_from);
Hans-Kristian Arntzen7d7f4b32017-08-10 17:12:48 +0200267
268 auto *access_chain = maybe_get<SPIRAccessChain>(chain);
269 if (access_chain)
270 var = maybe_get<SPIRVariable>(access_chain->loaded_from);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200271 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100272
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200273 return var;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100274}
275
276void Compiler::register_read(uint32_t expr, uint32_t chain, bool forwarded)
277{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200278 auto &e = get<SPIRExpression>(expr);
279 auto *var = maybe_get_backing_variable(chain);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100280
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200281 if (var)
282 {
283 e.loaded_from = var->self;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100284
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200285 // If the backing variable is immutable, we do not need to depend on the variable.
286 if (forwarded && !is_immutable(var->self))
287 var->dependees.push_back(e.self);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100288
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200289 // If we load from a parameter, make sure we create "inout" if we also write to the parameter.
290 // The default is "in" however, so we never invalidate our compilation by reading.
291 if (var && var->parameter)
292 var->parameter->read_count++;
293 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100294}
295
296void Compiler::register_write(uint32_t chain)
297{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200298 auto *var = maybe_get<SPIRVariable>(chain);
299 if (!var)
300 {
301 // If we're storing through an access chain, invalidate the backing variable instead.
302 auto *expr = maybe_get<SPIRExpression>(chain);
303 if (expr && expr->loaded_from)
304 var = maybe_get<SPIRVariable>(expr->loaded_from);
Hans-Kristian Arntzen7d7f4b32017-08-10 17:12:48 +0200305
306 auto *access_chain = maybe_get<SPIRAccessChain>(chain);
307 if (access_chain && access_chain->loaded_from)
308 var = maybe_get<SPIRVariable>(access_chain->loaded_from);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200309 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100310
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200311 if (var)
312 {
313 // If our variable is in a storage class which can alias with other buffers,
314 // invalidate all variables which depend on aliased variables.
315 if (variable_storage_is_aliased(*var))
316 flush_all_aliased_variables();
317 else if (var)
318 flush_dependees(*var);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100319
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200320 // We tried to write to a parameter which is not marked with out qualifier, force a recompile.
321 if (var->parameter && var->parameter->write_count == 0)
322 {
323 var->parameter->write_count++;
324 force_recompile = true;
325 }
326 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100327}
328
329void Compiler::flush_dependees(SPIRVariable &var)
330{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200331 for (auto expr : var.dependees)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200332 invalid_expressions.insert(expr);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200333 var.dependees.clear();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100334}
335
336void Compiler::flush_all_aliased_variables()
337{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200338 for (auto aliased : aliased_variables)
339 flush_dependees(get<SPIRVariable>(aliased));
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100340}
341
342void Compiler::flush_all_atomic_capable_variables()
343{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200344 for (auto global : global_variables)
345 flush_dependees(get<SPIRVariable>(global));
346 flush_all_aliased_variables();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100347}
348
349void Compiler::flush_all_active_variables()
350{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200351 // Invalidate all temporaries we read from variables in this block since they were forwarded.
352 // Invalidate all temporaries we read from globals.
353 for (auto &v : current_function->local_variables)
354 flush_dependees(get<SPIRVariable>(v));
355 for (auto &arg : current_function->arguments)
356 flush_dependees(get<SPIRVariable>(arg.id));
357 for (auto global : global_variables)
358 flush_dependees(get<SPIRVariable>(global));
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100359
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200360 flush_all_aliased_variables();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100361}
362
Bill Hollings1e84a372017-08-12 00:21:13 -0400363uint32_t Compiler::expression_type_id(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100364{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200365 switch (ids[id].get_type())
366 {
367 case TypeVariable:
Bill Hollings1e84a372017-08-12 00:21:13 -0400368 return get<SPIRVariable>(id).basetype;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100369
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200370 case TypeExpression:
Bill Hollings1e84a372017-08-12 00:21:13 -0400371 return get<SPIRExpression>(id).expression_type;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100372
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200373 case TypeConstant:
Bill Hollings1e84a372017-08-12 00:21:13 -0400374 return get<SPIRConstant>(id).constant_type;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100375
Hans-Kristian Arntzen7e8afa82016-10-03 15:54:02 +0200376 case TypeConstantOp:
Bill Hollings1e84a372017-08-12 00:21:13 -0400377 return get<SPIRConstantOp>(id).basetype;
Hans-Kristian Arntzen7e8afa82016-10-03 15:54:02 +0200378
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200379 case TypeUndef:
Bill Hollings1e84a372017-08-12 00:21:13 -0400380 return get<SPIRUndef>(id).basetype;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100381
Hans-Kristian Arntzen100e9d32017-04-25 10:44:55 +0200382 case TypeCombinedImageSampler:
Bill Hollings1e84a372017-08-12 00:21:13 -0400383 return get<SPIRCombinedImageSampler>(id).combined_type;
Hans-Kristian Arntzen100e9d32017-04-25 10:44:55 +0200384
Hans-Kristian Arntzen3cbdbec2017-08-10 15:36:30 +0200385 case TypeAccessChain:
Hans-Kristian Arntzene2bb5b82017-08-15 09:34:30 +0200386 return get<SPIRAccessChain>(id).basetype;
Hans-Kristian Arntzen3cbdbec2017-08-10 15:36:30 +0200387
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200388 default:
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100389 SPIRV_CROSS_THROW("Cannot resolve expression type.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200390 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100391}
392
Bill Hollings1e84a372017-08-12 00:21:13 -0400393const SPIRType &Compiler::expression_type(uint32_t id) const
394{
395 return get<SPIRType>(expression_type_id(id));
396}
397
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100398bool Compiler::expression_is_lvalue(uint32_t id) const
399{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200400 auto &type = expression_type(id);
401 switch (type.basetype)
402 {
403 case SPIRType::SampledImage:
404 case SPIRType::Image:
405 case SPIRType::Sampler:
406 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100407
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200408 default:
409 return true;
410 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100411}
412
413bool Compiler::is_immutable(uint32_t id) const
414{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200415 if (ids[id].get_type() == TypeVariable)
416 {
417 auto &var = get<SPIRVariable>(id);
Hans-Kristian Arntzen92134e42016-04-01 19:58:26 +0200418
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200419 // Anything we load from the UniformConstant address space is guaranteed to be immutable.
420 bool pointer_to_const = var.storage == StorageClassUniformConstant;
Hans-Kristian Arntzen36a0b632016-07-12 14:33:04 +0200421 return pointer_to_const || var.phi_variable || !expression_is_lvalue(id);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200422 }
Hans-Kristian Arntzen7d7f4b32017-08-10 17:12:48 +0200423 else if (ids[id].get_type() == TypeAccessChain)
424 return get<SPIRAccessChain>(id).immutable;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200425 else if (ids[id].get_type() == TypeExpression)
426 return get<SPIRExpression>(id).immutable;
Hans-Kristian Arntzen7e8afa82016-10-03 15:54:02 +0200427 else if (ids[id].get_type() == TypeConstant || ids[id].get_type() == TypeConstantOp ||
428 ids[id].get_type() == TypeUndef)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200429 return true;
430 else
431 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100432}
433
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200434static inline bool storage_class_is_interface(spv::StorageClass storage)
435{
436 switch (storage)
437 {
438 case StorageClassInput:
439 case StorageClassOutput:
440 case StorageClassUniform:
441 case StorageClassUniformConstant:
442 case StorageClassAtomicCounter:
443 case StorageClassPushConstant:
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +0200444 case StorageClassStorageBuffer:
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200445 return true;
446
447 default:
448 return false;
449 }
450}
451
452bool Compiler::is_hidden_variable(const SPIRVariable &var, bool include_builtins) const
453{
454 if ((is_builtin_variable(var) && !include_builtins) || var.remapped_variable)
455 return true;
456
Hans-Kristian Arntzen1b5ca8d2016-09-10 16:20:19 +0200457 // Combined image samplers are always considered active as they are "magic" variables.
458 if (find_if(begin(combined_image_samplers), end(combined_image_samplers), [&var](const CombinedImageSampler &samp) {
459 return samp.combined_id == var.self;
Hans-Kristian Arntzen1079e792017-10-10 10:22:40 +0200460 }) != end(combined_image_samplers))
Hans-Kristian Arntzen1b5ca8d2016-09-10 16:20:19 +0200461 {
462 return false;
463 }
464
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200465 bool hidden = false;
466 if (check_active_interface_variables && storage_class_is_interface(var.storage))
467 hidden = active_interface_variables.find(var.self) == end(active_interface_variables);
468 return hidden;
469}
470
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100471bool Compiler::is_builtin_variable(const SPIRVariable &var) const
472{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200473 if (var.compat_builtin || meta[var.self].decoration.builtin)
474 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100475
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200476 // We can have builtin structs as well. If one member of a struct is builtin, the struct must also be builtin.
477 for (auto &m : meta[get<SPIRType>(var.basetype).self].members)
478 if (m.builtin)
479 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100480
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200481 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100482}
483
484bool Compiler::is_member_builtin(const SPIRType &type, uint32_t index, BuiltIn *builtin) const
485{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200486 auto &memb = meta[type.self].members;
487 if (index < memb.size() && memb[index].builtin)
488 {
489 if (builtin)
490 *builtin = memb[index].builtin_type;
491 return true;
492 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100493
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200494 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100495}
496
Bill Hollings103aabf2016-04-06 17:42:27 -0400497bool Compiler::is_scalar(const SPIRType &type) const
498{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200499 return type.vecsize == 1 && type.columns == 1;
Bill Hollings103aabf2016-04-06 17:42:27 -0400500}
501
502bool Compiler::is_vector(const SPIRType &type) const
503{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200504 return type.vecsize > 1 && type.columns == 1;
Bill Hollings103aabf2016-04-06 17:42:27 -0400505}
506
507bool Compiler::is_matrix(const SPIRType &type) const
508{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200509 return type.vecsize > 1 && type.columns > 1;
Bill Hollings103aabf2016-04-06 17:42:27 -0400510}
511
Bill Hollingsf591bc02017-06-30 19:10:46 -0400512bool Compiler::is_array(const SPIRType &type) const
513{
514 return !type.array.empty();
515}
516
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100517ShaderResources Compiler::get_shader_resources() const
518{
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200519 return get_shader_resources(nullptr);
520}
521
522ShaderResources Compiler::get_shader_resources(const unordered_set<uint32_t> &active_variables) const
523{
524 return get_shader_resources(&active_variables);
525}
526
527bool Compiler::InterfaceVariableAccessHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
528{
529 uint32_t variable = 0;
530 switch (opcode)
531 {
532 // Need this first, otherwise, GCC complains about unhandled switch statements.
533 default:
534 break;
535
536 case OpFunctionCall:
537 {
538 // Invalid SPIR-V.
539 if (length < 3)
540 return false;
541
542 uint32_t count = length - 3;
543 args += 3;
544 for (uint32_t i = 0; i < count; i++)
545 {
546 auto *var = compiler.maybe_get<SPIRVariable>(args[i]);
547 if (var && storage_class_is_interface(var->storage))
548 variables.insert(args[i]);
549 }
550 break;
551 }
552
553 case OpAtomicStore:
554 case OpStore:
555 // Invalid SPIR-V.
556 if (length < 1)
557 return false;
558 variable = args[0];
559 break;
560
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +0100561 case OpCopyMemory:
562 {
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +0100563 if (length < 2)
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +0100564 return false;
565
566 auto *var = compiler.maybe_get<SPIRVariable>(args[0]);
567 if (var && storage_class_is_interface(var->storage))
568 variables.insert(variable);
569
570 var = compiler.maybe_get<SPIRVariable>(args[1]);
571 if (var && storage_class_is_interface(var->storage))
572 variables.insert(variable);
573 break;
574 }
575
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200576 case OpAccessChain:
577 case OpInBoundsAccessChain:
578 case OpLoad:
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +0100579 case OpCopyObject:
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200580 case OpImageTexelPointer:
581 case OpAtomicLoad:
582 case OpAtomicExchange:
583 case OpAtomicCompareExchange:
Bill Hollings8f6df772017-05-19 18:14:08 -0400584 case OpAtomicCompareExchangeWeak:
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200585 case OpAtomicIIncrement:
586 case OpAtomicIDecrement:
587 case OpAtomicIAdd:
588 case OpAtomicISub:
589 case OpAtomicSMin:
590 case OpAtomicUMin:
591 case OpAtomicSMax:
592 case OpAtomicUMax:
593 case OpAtomicAnd:
594 case OpAtomicOr:
595 case OpAtomicXor:
596 // Invalid SPIR-V.
597 if (length < 3)
598 return false;
599 variable = args[2];
600 break;
601 }
602
603 if (variable)
604 {
605 auto *var = compiler.maybe_get<SPIRVariable>(variable);
606 if (var && storage_class_is_interface(var->storage))
607 variables.insert(variable);
608 }
609 return true;
610}
611
612unordered_set<uint32_t> Compiler::get_active_interface_variables() const
613{
614 // Traverse the call graph and find all interface variables which are in use.
615 unordered_set<uint32_t> variables;
616 InterfaceVariableAccessHandler handler(*this, variables);
617 traverse_all_reachable_opcodes(get<SPIRFunction>(entry_point), handler);
618 return variables;
619}
620
621void Compiler::set_enabled_interface_variables(std::unordered_set<uint32_t> active_variables)
622{
623 active_interface_variables = move(active_variables);
624 check_active_interface_variables = true;
625}
626
627ShaderResources Compiler::get_shader_resources(const unordered_set<uint32_t> *active_variables) const
628{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200629 ShaderResources res;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100630
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200631 for (auto &id : ids)
632 {
633 if (id.get_type() != TypeVariable)
634 continue;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100635
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200636 auto &var = id.get<SPIRVariable>();
637 auto &type = get<SPIRType>(var.basetype);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100638
Hans-Kristian Arntzend5dc5f32016-07-05 13:21:26 +0200639 // It is possible for uniform storage classes to be passed as function parameters, so detect
640 // that. To detect function parameters, check of StorageClass of variable is function scope.
641 if (var.storage == StorageClassFunction || !type.pointer || is_builtin_variable(var))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200642 continue;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100643
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200644 if (active_variables && active_variables->find(var.self) == end(*active_variables))
645 continue;
646
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200647 // Input
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +0200648 if (var.storage == StorageClassInput && interface_variable_exists_in_entry_point(var.self))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200649 {
650 if (meta[type.self].decoration.decoration_flags & (1ull << DecorationBlock))
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200651 res.stage_inputs.push_back({ var.self, var.basetype, type.self, meta[type.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200652 else
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200653 res.stage_inputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200654 }
655 // Subpass inputs
656 else if (var.storage == StorageClassUniformConstant && type.image.dim == DimSubpassData)
657 {
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200658 res.subpass_inputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200659 }
660 // Outputs
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +0200661 else if (var.storage == StorageClassOutput && interface_variable_exists_in_entry_point(var.self))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200662 {
663 if (meta[type.self].decoration.decoration_flags & (1ull << DecorationBlock))
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200664 res.stage_outputs.push_back({ var.self, var.basetype, type.self, meta[type.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200665 else
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200666 res.stage_outputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200667 }
668 // UBOs
669 else if (type.storage == StorageClassUniform &&
670 (meta[type.self].decoration.decoration_flags & (1ull << DecorationBlock)))
671 {
Hans-Kristian Arntzenaab31072017-09-29 12:16:53 +0200672 auto &block_name = meta[type.self].decoration.alias;
673 res.uniform_buffers.push_back({ var.self, var.basetype, type.self,
674 block_name.empty() ? get_block_fallback_name(var.self) : block_name });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200675 }
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +0200676 // Old way to declare SSBOs.
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200677 else if (type.storage == StorageClassUniform &&
678 (meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)))
679 {
Hans-Kristian Arntzenaab31072017-09-29 12:16:53 +0200680 auto &block_name = meta[type.self].decoration.alias;
681 res.storage_buffers.push_back({ var.self, var.basetype, type.self,
682 block_name.empty() ? get_block_fallback_name(var.self) : block_name });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200683 }
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +0200684 // Modern way to declare SSBOs.
685 else if (type.storage == StorageClassStorageBuffer)
686 {
Hans-Kristian Arntzenaab31072017-09-29 12:16:53 +0200687 auto &block_name = meta[type.self].decoration.alias;
688 res.storage_buffers.push_back({ var.self, var.basetype, type.self,
689 block_name.empty() ? get_block_fallback_name(var.self) : block_name });
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +0200690 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200691 // Push constant blocks
692 else if (type.storage == StorageClassPushConstant)
693 {
694 // There can only be one push constant block, but keep the vector in case this restriction is lifted
695 // in the future.
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200696 res.push_constant_buffers.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200697 }
698 // Images
Hans-Kristian Arntzene9202082016-09-10 13:05:35 +0200699 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Image &&
700 type.image.sampled == 2)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200701 {
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200702 res.storage_images.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200703 }
Hans-Kristian Arntzene9202082016-09-10 13:05:35 +0200704 // Separate images
705 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Image &&
706 type.image.sampled == 1)
707 {
708 res.separate_images.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
709 }
710 // Separate samplers
711 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Sampler)
712 {
713 res.separate_samplers.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
714 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200715 // Textures
716 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::SampledImage)
717 {
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200718 res.sampled_images.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200719 }
720 // Atomic counters
721 else if (type.storage == StorageClassAtomicCounter)
722 {
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200723 res.atomic_counters.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200724 }
725 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100726
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200727 return res;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100728}
729
730static inline uint32_t swap_endian(uint32_t v)
731{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200732 return ((v >> 24) & 0x000000ffu) | ((v >> 8) & 0x0000ff00u) | ((v << 8) & 0x00ff0000u) | ((v << 24) & 0xff000000u);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100733}
734
735static string extract_string(const vector<uint32_t> &spirv, uint32_t offset)
736{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200737 string ret;
738 for (uint32_t i = offset; i < spirv.size(); i++)
739 {
740 uint32_t w = spirv[i];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100741
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200742 for (uint32_t j = 0; j < 4; j++, w >>= 8)
743 {
744 char c = w & 0xff;
745 if (c == '\0')
746 return ret;
747 ret += c;
748 }
749 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100750
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100751 SPIRV_CROSS_THROW("String was not terminated before EOF");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100752}
753
Hans-Kristian Arntzen45ad58a2016-05-10 23:39:41 +0200754static bool is_valid_spirv_version(uint32_t version)
755{
756 switch (version)
757 {
758 // Allow v99 since it tends to just work.
759 case 99:
760 case 0x10000: // SPIR-V 1.0
761 case 0x10100: // SPIR-V 1.1
Hans-Kristian Arntzene5595272017-05-22 13:59:58 +0200762 case 0x10200: // SPIR-V 1.2
Hans-Kristian Arntzen45ad58a2016-05-10 23:39:41 +0200763 return true;
764
765 default:
766 return false;
767 }
768}
769
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100770void Compiler::parse()
771{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200772 auto len = spirv.size();
773 if (len < 5)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100774 SPIRV_CROSS_THROW("SPIRV file too small.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100775
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200776 auto s = spirv.data();
Hans-Kristian Arntzen5ac88272016-04-11 13:38:18 +0200777
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200778 // Endian-swap if we need to.
779 if (s[0] == swap_endian(MagicNumber))
Hans-Kristian Arntzen5ea59bd2016-05-23 13:30:02 +0200780 transform(begin(spirv), end(spirv), begin(spirv), [](uint32_t c) { return swap_endian(c); });
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100781
Hans-Kristian Arntzen45ad58a2016-05-10 23:39:41 +0200782 if (s[0] != MagicNumber || !is_valid_spirv_version(s[1]))
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100783 SPIRV_CROSS_THROW("Invalid SPIRV format.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100784
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200785 uint32_t bound = s[3];
786 ids.resize(bound);
787 meta.resize(bound);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100788
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200789 uint32_t offset = 5;
790 while (offset < len)
791 inst.emplace_back(spirv, offset);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100792
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200793 for (auto &i : inst)
794 parse(i);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100795
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200796 if (current_function)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100797 SPIRV_CROSS_THROW("Function was not terminated.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200798 if (current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100799 SPIRV_CROSS_THROW("Block was not terminated.");
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +0200800
801 // Figure out specialization constants for work group sizes.
802 for (auto &id : ids)
803 {
804 if (id.get_type() == TypeConstant)
805 {
806 auto &c = id.get<SPIRConstant>();
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +0200807 if (meta[c.self].decoration.builtin && meta[c.self].decoration.builtin_type == BuiltInWorkgroupSize)
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +0200808 {
809 // In current SPIR-V, there can be just one constant like this.
810 // All entry points will receive the constant value.
811 for (auto &entry : entry_points)
812 {
813 entry.second.workgroup_size.constant = c.self;
814 entry.second.workgroup_size.x = c.scalar(0, 0);
815 entry.second.workgroup_size.y = c.scalar(0, 1);
816 entry.second.workgroup_size.z = c.scalar(0, 2);
817 }
818 }
819 }
820 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100821}
822
823void Compiler::flatten_interface_block(uint32_t id)
824{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200825 auto &var = get<SPIRVariable>(id);
826 auto &type = get<SPIRType>(var.basetype);
827 auto flags = meta.at(type.self).decoration.decoration_flags;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100828
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200829 if (!type.array.empty())
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100830 SPIRV_CROSS_THROW("Type is array of UBOs.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200831 if (type.basetype != SPIRType::Struct)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100832 SPIRV_CROSS_THROW("Type is not a struct.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200833 if ((flags & (1ull << DecorationBlock)) == 0)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100834 SPIRV_CROSS_THROW("Type is not a block.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200835 if (type.member_types.empty())
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100836 SPIRV_CROSS_THROW("Member list of struct is empty.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100837
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200838 uint32_t t = type.member_types[0];
839 for (auto &m : type.member_types)
840 if (t != m)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100841 SPIRV_CROSS_THROW("Types in block differ.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100842
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200843 auto &mtype = get<SPIRType>(t);
844 if (!mtype.array.empty())
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100845 SPIRV_CROSS_THROW("Member type cannot be arrays.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200846 if (mtype.basetype == SPIRType::Struct)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100847 SPIRV_CROSS_THROW("Member type cannot be struct.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100848
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200849 // Inherit variable name from interface block name.
850 meta.at(var.self).decoration.alias = meta.at(type.self).decoration.alias;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100851
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200852 auto storage = var.storage;
853 if (storage == StorageClassUniform)
854 storage = StorageClassUniformConstant;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100855
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200856 // Change type definition in-place into an array instead.
857 // Access chains will still work as-is.
858 uint32_t array_size = uint32_t(type.member_types.size());
859 type = mtype;
860 type.array.push_back(array_size);
861 type.pointer = true;
862 type.storage = storage;
863 var.storage = storage;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100864}
865
866void Compiler::update_name_cache(unordered_set<string> &cache, string &name)
867{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200868 if (name.empty())
869 return;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100870
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200871 if (cache.find(name) == end(cache))
872 {
873 cache.insert(name);
874 return;
875 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100876
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200877 uint32_t counter = 0;
878 auto tmpname = name;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100879
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200880 // If there is a collision (very rare),
881 // keep tacking on extra identifier until it's unique.
882 do
883 {
884 counter++;
885 name = tmpname + "_" + convert_to_string(counter);
886 } while (cache.find(name) != end(cache));
887 cache.insert(name);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100888}
889
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200890void Compiler::set_name(uint32_t id, const std::string &name)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100891{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200892 auto &str = meta.at(id).decoration.alias;
893 str.clear();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100894
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200895 if (name.empty())
896 return;
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +0200897
898 // glslang uses identifiers to pass along meaningful information
899 // about HLSL reflection.
900 auto &m = meta.at(id);
Hans-Kristian Arntzen08b3c672017-05-09 09:30:30 +0200901 if (source.hlsl && name.size() >= 6 && name.find("@count") == name.size() - 6)
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +0200902 {
903 m.hlsl_magic_counter_buffer_candidate = true;
904 m.hlsl_magic_counter_buffer_name = name.substr(0, name.find("@count"));
905 }
906 else
907 {
908 m.hlsl_magic_counter_buffer_candidate = false;
909 m.hlsl_magic_counter_buffer_name.clear();
910 }
911
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200912 // Reserved for temporaries.
Hans-Kristian Arntzenc9728942016-07-06 11:19:20 +0200913 if (name[0] == '_' && name.size() >= 2 && isdigit(name[1]))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200914 return;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100915
Hans-Kristian Arntzen94ff3552017-10-10 17:32:26 +0200916 str = ensure_valid_identifier(name, false);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100917}
918
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200919const SPIRType &Compiler::get_type(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100920{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200921 return get<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100922}
923
Robert Konrad451bdee2016-09-24 22:17:01 +0200924const SPIRType &Compiler::get_type_from_variable(uint32_t id) const
925{
926 return get<SPIRType>(get<SPIRVariable>(id).basetype);
927}
928
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100929void Compiler::set_member_decoration(uint32_t id, uint32_t index, Decoration decoration, uint32_t argument)
930{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200931 meta.at(id).members.resize(max(meta[id].members.size(), size_t(index) + 1));
932 auto &dec = meta.at(id).members[index];
933 dec.decoration_flags |= 1ull << decoration;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100934
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200935 switch (decoration)
936 {
937 case DecorationBuiltIn:
938 dec.builtin = true;
939 dec.builtin_type = static_cast<BuiltIn>(argument);
940 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100941
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200942 case DecorationLocation:
943 dec.location = argument;
944 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100945
Bill Hollings81757502017-01-29 13:28:20 -0500946 case DecorationBinding:
947 dec.binding = argument;
948 break;
949
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200950 case DecorationOffset:
951 dec.offset = argument;
952 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100953
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +0200954 case DecorationSpecId:
955 dec.spec_id = argument;
956 break;
957
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +0100958 case DecorationMatrixStride:
959 dec.matrix_stride = argument;
960 break;
961
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200962 default:
963 break;
964 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100965}
966
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200967void Compiler::set_member_name(uint32_t id, uint32_t index, const std::string &name)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100968{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200969 meta.at(id).members.resize(max(meta[id].members.size(), size_t(index) + 1));
Hans-Kristian Arntzenfd432f82017-03-18 10:52:41 +0100970
971 auto &str = meta.at(id).members[index].alias;
972 str.clear();
973 if (name.empty())
974 return;
975
976 // Reserved for unnamed members.
Hans-Kristian Arntzen94ff3552017-10-10 17:32:26 +0200977 if (name[0] == '_' && name.size() >= 3 && name[1] == 'm' && isdigit(name[2]))
Hans-Kristian Arntzenfd432f82017-03-18 10:52:41 +0100978 return;
979
Hans-Kristian Arntzen94ff3552017-10-10 17:32:26 +0200980 str = ensure_valid_identifier(name, true);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100981}
982
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200983const std::string &Compiler::get_member_name(uint32_t id, uint32_t index) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100984{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200985 auto &m = meta.at(id);
986 if (index >= m.members.size())
987 {
988 static string empty;
989 return empty;
990 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100991
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200992 return m.members[index].alias;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100993}
994
Bill Hollings1e84a372017-08-12 00:21:13 -0400995void Compiler::set_member_qualified_name(uint32_t type_id, uint32_t index, const std::string &name)
Bill Hollingsc45e74f2016-07-08 12:39:22 -0400996{
Bill Hollingsc1b81542017-05-22 21:41:19 -0400997 meta.at(type_id).members.resize(max(meta[type_id].members.size(), size_t(index) + 1));
998 meta.at(type_id).members[index].qualified_alias = name;
999}
1000
Bill Hollings1e84a372017-08-12 00:21:13 -04001001const std::string &Compiler::get_member_qualified_name(uint32_t type_id, uint32_t index) const
Bill Hollingsc1b81542017-05-22 21:41:19 -04001002{
Bill Hollings1e84a372017-08-12 00:21:13 -04001003 const static string empty;
Bill Hollingsc1b81542017-05-22 21:41:19 -04001004
1005 auto &m = meta.at(type_id);
Bill Hollings1e84a372017-08-12 00:21:13 -04001006 if (index < m.members.size())
1007 return m.members[index].qualified_alias;
1008 else
Bill Hollingsc1b81542017-05-22 21:41:19 -04001009 return empty;
Bill Hollingsc45e74f2016-07-08 12:39:22 -04001010}
1011
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001012uint32_t Compiler::get_member_decoration(uint32_t id, uint32_t index, Decoration decoration) const
1013{
Hans-Kristian Arntzen15679c72016-08-17 11:35:34 +02001014 auto &m = meta.at(id);
1015 if (index >= m.members.size())
1016 return 0;
1017
1018 auto &dec = m.members[index];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001019 if (!(dec.decoration_flags & (1ull << decoration)))
1020 return 0;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001021
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001022 switch (decoration)
1023 {
1024 case DecorationBuiltIn:
1025 return dec.builtin_type;
1026 case DecorationLocation:
1027 return dec.location;
Bill Hollings81757502017-01-29 13:28:20 -05001028 case DecorationBinding:
1029 return dec.binding;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001030 case DecorationOffset:
1031 return dec.offset;
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02001032 case DecorationSpecId:
1033 return dec.spec_id;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001034 default:
Graham Wihlidal9b1ee8f2017-01-05 21:01:49 +01001035 return 1;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001036 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001037}
1038
1039uint64_t Compiler::get_member_decoration_mask(uint32_t id, uint32_t index) const
1040{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001041 auto &m = meta.at(id);
1042 if (index >= m.members.size())
1043 return 0;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001044
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001045 return m.members[index].decoration_flags;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001046}
1047
Bill Hollingsdc694272017-03-11 12:17:22 -05001048bool Compiler::has_member_decoration(uint32_t id, uint32_t index, Decoration decoration) const
Bill Hollings484931d2017-02-28 21:44:36 -05001049{
Bill Hollingsdc694272017-03-11 12:17:22 -05001050 return get_member_decoration_mask(id, index) & (1ull << decoration);
Bill Hollings484931d2017-02-28 21:44:36 -05001051}
1052
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001053void Compiler::unset_member_decoration(uint32_t id, uint32_t index, Decoration decoration)
1054{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001055 auto &m = meta.at(id);
1056 if (index >= m.members.size())
1057 return;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001058
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001059 auto &dec = m.members[index];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001060
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001061 dec.decoration_flags &= ~(1ull << decoration);
1062 switch (decoration)
1063 {
1064 case DecorationBuiltIn:
1065 dec.builtin = false;
1066 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001067
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001068 case DecorationLocation:
1069 dec.location = 0;
1070 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001071
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001072 case DecorationOffset:
1073 dec.offset = 0;
1074 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001075
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02001076 case DecorationSpecId:
1077 dec.spec_id = 0;
1078 break;
1079
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001080 default:
1081 break;
1082 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001083}
1084
1085void Compiler::set_decoration(uint32_t id, Decoration decoration, uint32_t argument)
1086{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001087 auto &dec = meta.at(id).decoration;
1088 dec.decoration_flags |= 1ull << decoration;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001089
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001090 switch (decoration)
1091 {
1092 case DecorationBuiltIn:
1093 dec.builtin = true;
1094 dec.builtin_type = static_cast<BuiltIn>(argument);
1095 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001096
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001097 case DecorationLocation:
1098 dec.location = argument;
1099 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001100
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001101 case DecorationOffset:
1102 dec.offset = argument;
1103 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001104
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001105 case DecorationArrayStride:
1106 dec.array_stride = argument;
1107 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001108
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001109 case DecorationMatrixStride:
1110 dec.matrix_stride = argument;
1111 break;
1112
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001113 case DecorationBinding:
1114 dec.binding = argument;
1115 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001116
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001117 case DecorationDescriptorSet:
1118 dec.set = argument;
1119 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001120
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001121 case DecorationInputAttachmentIndex:
1122 dec.input_attachment = argument;
1123 break;
Hans-Kristian Arntzen12cfbb22016-05-04 13:41:04 +02001124
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02001125 case DecorationSpecId:
1126 dec.spec_id = argument;
1127 break;
1128
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001129 default:
1130 break;
1131 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001132}
1133
1134StorageClass Compiler::get_storage_class(uint32_t id) const
1135{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001136 return get<SPIRVariable>(id).storage;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001137}
1138
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001139const std::string &Compiler::get_name(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001140{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001141 return meta.at(id).decoration.alias;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001142}
1143
Hans-Kristian Arntzenaab31072017-09-29 12:16:53 +02001144const std::string Compiler::get_fallback_name(uint32_t id) const
1145{
1146 return join("_", id);
1147}
1148
1149const std::string Compiler::get_block_fallback_name(uint32_t id) const
1150{
1151 auto &var = get<SPIRVariable>(id);
1152 return join("_", get<SPIRType>(var.basetype).self, "_", id);
1153}
1154
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001155uint64_t Compiler::get_decoration_mask(uint32_t id) const
1156{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001157 auto &dec = meta.at(id).decoration;
1158 return dec.decoration_flags;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001159}
1160
Bill Hollingsdc694272017-03-11 12:17:22 -05001161bool Compiler::has_decoration(uint32_t id, Decoration decoration) const
Bill Hollings484931d2017-02-28 21:44:36 -05001162{
Bill Hollingsdc694272017-03-11 12:17:22 -05001163 return get_decoration_mask(id) & (1ull << decoration);
Bill Hollings484931d2017-02-28 21:44:36 -05001164}
1165
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001166uint32_t Compiler::get_decoration(uint32_t id, Decoration decoration) const
1167{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001168 auto &dec = meta.at(id).decoration;
1169 if (!(dec.decoration_flags & (1ull << decoration)))
1170 return 0;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001171
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001172 switch (decoration)
1173 {
1174 case DecorationBuiltIn:
1175 return dec.builtin_type;
1176 case DecorationLocation:
1177 return dec.location;
1178 case DecorationOffset:
1179 return dec.offset;
1180 case DecorationBinding:
1181 return dec.binding;
1182 case DecorationDescriptorSet:
1183 return dec.set;
1184 case DecorationInputAttachmentIndex:
1185 return dec.input_attachment;
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02001186 case DecorationSpecId:
1187 return dec.spec_id;
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001188 case DecorationArrayStride:
1189 return dec.array_stride;
1190 case DecorationMatrixStride:
1191 return dec.matrix_stride;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001192 default:
Graham Wihlidal9b1ee8f2017-01-05 21:01:49 +01001193 return 1;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001194 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001195}
1196
1197void Compiler::unset_decoration(uint32_t id, Decoration decoration)
1198{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001199 auto &dec = meta.at(id).decoration;
1200 dec.decoration_flags &= ~(1ull << decoration);
1201 switch (decoration)
1202 {
1203 case DecorationBuiltIn:
1204 dec.builtin = false;
1205 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001206
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001207 case DecorationLocation:
1208 dec.location = 0;
1209 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001210
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001211 case DecorationOffset:
1212 dec.offset = 0;
1213 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001214
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001215 case DecorationBinding:
1216 dec.binding = 0;
1217 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001218
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001219 case DecorationDescriptorSet:
1220 dec.set = 0;
1221 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001222
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02001223 case DecorationInputAttachmentIndex:
1224 dec.input_attachment = 0;
1225 break;
1226
1227 case DecorationSpecId:
1228 dec.spec_id = 0;
1229 break;
1230
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001231 default:
1232 break;
1233 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001234}
1235
Hans-Kristian Arntzen93fe19c2017-04-25 20:10:24 +02001236bool Compiler::get_binary_offset_for_decoration(uint32_t id, spv::Decoration decoration, uint32_t &word_offset) const
1237{
1238 auto &word_offsets = meta.at(id).decoration_word_offset;
1239 auto itr = word_offsets.find(decoration);
1240 if (itr == end(word_offsets))
1241 return false;
1242
1243 word_offset = itr->second;
1244 return true;
1245}
1246
Hans-Kristian Arntzen926916d2016-05-05 09:15:25 +02001247void Compiler::parse(const Instruction &instruction)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001248{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001249 auto ops = stream(instruction);
1250 auto op = static_cast<Op>(instruction.op);
1251 uint32_t length = instruction.length;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001252
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001253 switch (op)
1254 {
1255 case OpMemoryModel:
1256 case OpSourceExtension:
1257 case OpNop:
Robert Konrad8f7c1af2016-08-10 02:43:51 +02001258 case OpLine:
David Srbeckya08b2692017-06-27 15:34:45 +01001259 case OpString:
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001260 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001261
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001262 case OpSource:
1263 {
1264 auto lang = static_cast<SourceLanguage>(ops[0]);
1265 switch (lang)
1266 {
1267 case SourceLanguageESSL:
1268 source.es = true;
1269 source.version = ops[1];
1270 source.known = true;
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02001271 source.hlsl = false;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001272 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001273
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001274 case SourceLanguageGLSL:
1275 source.es = false;
1276 source.version = ops[1];
1277 source.known = true;
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02001278 source.hlsl = false;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001279 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001280
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02001281 case SourceLanguageHLSL:
1282 // For purposes of cross-compiling, this is GLSL 450.
1283 source.es = false;
1284 source.version = 450;
1285 source.known = true;
1286 source.hlsl = true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001287 break;
1288
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001289 default:
1290 source.known = false;
1291 break;
1292 }
1293 break;
1294 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001295
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001296 case OpUndef:
1297 {
1298 uint32_t result_type = ops[0];
1299 uint32_t id = ops[1];
1300 set<SPIRUndef>(id, result_type);
1301 break;
1302 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001303
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001304 case OpCapability:
1305 {
1306 uint32_t cap = ops[0];
1307 if (cap == CapabilityKernel)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001308 SPIRV_CROSS_THROW("Kernel capability not supported.");
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02001309
1310 declared_capabilities.push_back(static_cast<Capability>(ops[0]));
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001311 break;
1312 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001313
Hans-Kristian Arntzen299e19f2017-03-21 16:33:54 +01001314 case OpExtension:
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02001315 {
1316 auto ext = extract_string(spirv, instruction.offset);
1317 declared_extensions.push_back(move(ext));
Hans-Kristian Arntzen299e19f2017-03-21 16:33:54 +01001318 break;
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02001319 }
Hans-Kristian Arntzen299e19f2017-03-21 16:33:54 +01001320
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001321 case OpExtInstImport:
1322 {
1323 uint32_t id = ops[0];
1324 auto ext = extract_string(spirv, instruction.offset + 1);
1325 if (ext == "GLSL.std.450")
1326 set<SPIRExtension>(id, SPIRExtension::GLSL);
1327 else
Hans-Kristian Arntzen299e19f2017-03-21 16:33:54 +01001328 set<SPIRExtension>(id, SPIRExtension::Unsupported);
1329
1330 // Other SPIR-V extensions currently not supported.
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001331
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001332 break;
1333 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001334
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001335 case OpEntryPoint:
1336 {
Bill Hollingsac00c602016-10-24 09:24:24 -04001337 auto itr =
1338 entry_points.insert(make_pair(ops[1], SPIREntryPoint(ops[1], static_cast<ExecutionModel>(ops[0]),
1339 extract_string(spirv, instruction.offset + 2))));
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001340 auto &e = itr.first->second;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001341
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001342 // Strings need nul-terminator and consume the whole word.
Hans-Kristian Arntzen7630d3c2016-11-21 12:14:02 +01001343 uint32_t strlen_words = uint32_t((e.name.size() + 1 + 3) >> 2);
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001344 e.interface_variables.insert(end(e.interface_variables), ops + strlen_words + 2, ops + instruction.length);
1345
Bill Hollingsac00c602016-10-24 09:24:24 -04001346 // Set the name of the entry point in case OpName is not provided later
1347 set_name(ops[1], e.name);
Bill Hollings0943d9f2016-10-23 21:42:54 -04001348
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001349 // If we don't have an entry, make the first one our "default".
1350 if (!entry_point)
1351 entry_point = ops[1];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001352 break;
1353 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001354
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001355 case OpExecutionMode:
1356 {
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001357 auto &execution = entry_points[ops[0]];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001358 auto mode = static_cast<ExecutionMode>(ops[1]);
1359 execution.flags |= 1ull << mode;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001360
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001361 switch (mode)
1362 {
1363 case ExecutionModeInvocations:
1364 execution.invocations = ops[2];
1365 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001366
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001367 case ExecutionModeLocalSize:
1368 execution.workgroup_size.x = ops[2];
1369 execution.workgroup_size.y = ops[3];
1370 execution.workgroup_size.z = ops[4];
1371 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001372
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001373 case ExecutionModeOutputVertices:
1374 execution.output_vertices = ops[2];
1375 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001376
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001377 default:
1378 break;
1379 }
1380 break;
1381 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001382
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001383 case OpName:
1384 {
1385 uint32_t id = ops[0];
1386 set_name(id, extract_string(spirv, instruction.offset + 1));
1387 break;
1388 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001389
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001390 case OpMemberName:
1391 {
1392 uint32_t id = ops[0];
1393 uint32_t member = ops[1];
1394 set_member_name(id, member, extract_string(spirv, instruction.offset + 2));
1395 break;
1396 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001397
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001398 case OpDecorate:
1399 {
1400 uint32_t id = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001401
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001402 auto decoration = static_cast<Decoration>(ops[1]);
1403 if (length >= 3)
Hans-Kristian Arntzen93fe19c2017-04-25 20:10:24 +02001404 {
1405 meta[id].decoration_word_offset[decoration] = uint32_t(&ops[2] - spirv.data());
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001406 set_decoration(id, decoration, ops[2]);
Hans-Kristian Arntzen93fe19c2017-04-25 20:10:24 +02001407 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001408 else
1409 set_decoration(id, decoration);
Bill Hollings0943d9f2016-10-23 21:42:54 -04001410
Bill Hollingsac00c602016-10-24 09:24:24 -04001411 break;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001412 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001413
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001414 case OpMemberDecorate:
1415 {
1416 uint32_t id = ops[0];
1417 uint32_t member = ops[1];
1418 auto decoration = static_cast<Decoration>(ops[2]);
1419 if (length >= 4)
1420 set_member_decoration(id, member, decoration, ops[3]);
1421 else
1422 set_member_decoration(id, member, decoration);
1423 break;
1424 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001425
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001426 // Build up basic types.
1427 case OpTypeVoid:
1428 {
1429 uint32_t id = ops[0];
1430 auto &type = set<SPIRType>(id);
1431 type.basetype = SPIRType::Void;
1432 break;
1433 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001434
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001435 case OpTypeBool:
1436 {
1437 uint32_t id = ops[0];
1438 auto &type = set<SPIRType>(id);
Hans-Kristian Arntzen5f629272016-06-05 20:13:45 +02001439 type.basetype = SPIRType::Boolean;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001440 type.width = 1;
1441 break;
1442 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001443
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001444 case OpTypeFloat:
1445 {
1446 uint32_t id = ops[0];
1447 uint32_t width = ops[1];
1448 auto &type = set<SPIRType>(id);
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001449 type.basetype = width > 32 ? SPIRType::Double : SPIRType::Float;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001450 type.width = width;
1451 break;
1452 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001453
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001454 case OpTypeInt:
1455 {
1456 uint32_t id = ops[0];
1457 uint32_t width = ops[1];
1458 auto &type = set<SPIRType>(id);
Hans-Kristian Arntzenfc2230f2016-07-27 11:27:00 +02001459 type.basetype =
1460 ops[2] ? (width > 32 ? SPIRType::Int64 : SPIRType::Int) : (width > 32 ? SPIRType::UInt64 : SPIRType::UInt);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001461 type.width = width;
1462 break;
1463 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001464
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001465 // Build composite types by "inheriting".
1466 // NOTE: The self member is also copied! For pointers and array modifiers this is a good thing
1467 // since we can refer to decorations on pointee classes which is needed for UBO/SSBO, I/O blocks in geometry/tess etc.
1468 case OpTypeVector:
1469 {
1470 uint32_t id = ops[0];
1471 uint32_t vecsize = ops[2];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001472
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001473 auto &base = get<SPIRType>(ops[1]);
1474 auto &vecbase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001475
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001476 vecbase = base;
1477 vecbase.vecsize = vecsize;
1478 vecbase.self = id;
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001479 vecbase.parent_type = ops[1];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001480 break;
1481 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001482
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001483 case OpTypeMatrix:
1484 {
1485 uint32_t id = ops[0];
1486 uint32_t colcount = ops[2];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001487
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001488 auto &base = get<SPIRType>(ops[1]);
1489 auto &matrixbase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001490
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001491 matrixbase = base;
1492 matrixbase.columns = colcount;
1493 matrixbase.self = id;
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001494 matrixbase.parent_type = ops[1];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001495 break;
1496 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001497
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001498 case OpTypeArray:
1499 {
1500 uint32_t id = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001501
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001502 auto &base = get<SPIRType>(ops[1]);
1503 auto &arraybase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001504
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001505 arraybase = base;
Hans-Kristian Arntzen5d4bb682016-10-03 17:17:11 +02001506
1507 auto *c = maybe_get<SPIRConstant>(ops[2]);
1508 bool literal = c && !c->specialization;
1509
1510 arraybase.array_size_literal.push_back(literal);
1511 arraybase.array.push_back(literal ? c->scalar() : ops[2]);
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001512 arraybase.parent_type = ops[1];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001513 // Do NOT set arraybase.self!
1514 break;
1515 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001516
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001517 case OpTypeRuntimeArray:
1518 {
1519 uint32_t id = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001520
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001521 auto &base = get<SPIRType>(ops[1]);
1522 auto &arraybase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001523
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001524 arraybase = base;
1525 arraybase.array.push_back(0);
Hans-Kristian Arntzen5d4bb682016-10-03 17:17:11 +02001526 arraybase.array_size_literal.push_back(true);
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001527 arraybase.parent_type = ops[1];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001528 // Do NOT set arraybase.self!
1529 break;
1530 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001531
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001532 case OpTypeImage:
1533 {
1534 uint32_t id = ops[0];
1535 auto &type = set<SPIRType>(id);
1536 type.basetype = SPIRType::Image;
1537 type.image.type = ops[1];
1538 type.image.dim = static_cast<Dim>(ops[2]);
1539 type.image.depth = ops[3] != 0;
1540 type.image.arrayed = ops[4] != 0;
1541 type.image.ms = ops[5] != 0;
1542 type.image.sampled = ops[6];
1543 type.image.format = static_cast<ImageFormat>(ops[7]);
Bill Hollingsb41e1482017-05-29 20:45:05 -04001544 type.image.access = (length >= 9) ? static_cast<AccessQualifier>(ops[8]) : AccessQualifierMax;
Bill Hollings6f3381a2017-06-01 16:29:39 -04001545
1546 if (type.image.sampled == 0)
1547 SPIRV_CROSS_THROW("OpTypeImage Sampled parameter must not be zero.");
1548
Robert Konrad98028232017-01-15 22:42:22 +01001549 break;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001550 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001551
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001552 case OpTypeSampledImage:
1553 {
1554 uint32_t id = ops[0];
1555 uint32_t imagetype = ops[1];
1556 auto &type = set<SPIRType>(id);
1557 type = get<SPIRType>(imagetype);
1558 type.basetype = SPIRType::SampledImage;
1559 type.self = id;
1560 break;
1561 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001562
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001563 case OpTypeSampler:
1564 {
1565 uint32_t id = ops[0];
1566 auto &type = set<SPIRType>(id);
1567 type.basetype = SPIRType::Sampler;
1568 break;
1569 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001570
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001571 case OpTypePointer:
1572 {
1573 uint32_t id = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001574
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001575 auto &base = get<SPIRType>(ops[2]);
1576 auto &ptrbase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001577
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001578 ptrbase = base;
1579 if (ptrbase.pointer)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001580 SPIRV_CROSS_THROW("Cannot make pointer-to-pointer type.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001581 ptrbase.pointer = true;
1582 ptrbase.storage = static_cast<StorageClass>(ops[1]);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001583
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001584 if (ptrbase.storage == StorageClassAtomicCounter)
1585 ptrbase.basetype = SPIRType::AtomicCounter;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001586
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001587 ptrbase.parent_type = ops[2];
1588
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001589 // Do NOT set ptrbase.self!
1590 break;
1591 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001592
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001593 case OpTypeStruct:
1594 {
1595 uint32_t id = ops[0];
1596 auto &type = set<SPIRType>(id);
1597 type.basetype = SPIRType::Struct;
1598 for (uint32_t i = 1; i < length; i++)
1599 type.member_types.push_back(ops[i]);
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02001600
1601 // Check if we have seen this struct type before, with just different
1602 // decorations.
Hans-Kristian Arntzen5ad43402016-05-28 09:47:52 +02001603 //
1604 // Add workaround for issue #17 as well by looking at OpName for the struct
1605 // types, which we shouldn't normally do.
1606 // We should not normally have to consider type aliases like this to begin with
1607 // however ... glslang issues #304, #307 cover this.
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02001608
Hans-Kristian Arntzen818c1be2017-07-31 09:31:20 +02001609 // For stripped names, never consider struct type aliasing.
1610 // We risk declaring the same struct multiple times, but type-punning is not allowed
1611 // so this is safe.
1612 bool consider_aliasing = !get_name(type.self).empty();
1613 if (consider_aliasing)
1614 {
1615 for (auto &other : global_struct_cache)
1616 {
1617 if (get_name(type.self) == get_name(other) &&
1618 types_are_logically_equivalent(type, get<SPIRType>(other)))
1619 {
1620 type.type_alias = other;
1621 break;
1622 }
1623 }
1624
1625 if (type.type_alias == 0)
1626 global_struct_cache.push_back(id);
1627 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001628 break;
1629 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001630
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001631 case OpTypeFunction:
1632 {
1633 uint32_t id = ops[0];
1634 uint32_t ret = ops[1];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001635
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001636 auto &func = set<SPIRFunctionPrototype>(id, ret);
1637 for (uint32_t i = 2; i < length; i++)
1638 func.parameter_types.push_back(ops[i]);
1639 break;
1640 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001641
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001642 // Variable declaration
1643 // All variables are essentially pointers with a storage qualifier.
1644 case OpVariable:
1645 {
1646 uint32_t type = ops[0];
1647 uint32_t id = ops[1];
1648 auto storage = static_cast<StorageClass>(ops[2]);
1649 uint32_t initializer = length == 4 ? ops[3] : 0;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001650
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001651 if (storage == StorageClassFunction)
1652 {
1653 if (!current_function)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001654 SPIRV_CROSS_THROW("No function currently in scope");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001655 current_function->add_local_variable(id);
1656 }
1657 else if (storage == StorageClassPrivate || storage == StorageClassWorkgroup || storage == StorageClassOutput)
1658 {
1659 global_variables.push_back(id);
1660 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001661
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001662 auto &var = set<SPIRVariable>(id, type, storage, initializer);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001663
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001664 if (variable_storage_is_aliased(var))
1665 aliased_variables.push_back(var.self);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001666
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001667 break;
1668 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001669
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001670 // OpPhi
1671 // OpPhi is a fairly magical opcode.
1672 // It selects temporary variables based on which parent block we *came from*.
1673 // In high-level languages we can "de-SSA" by creating a function local, and flush out temporaries to this function-local
1674 // variable to emulate SSA Phi.
1675 case OpPhi:
1676 {
1677 if (!current_function)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001678 SPIRV_CROSS_THROW("No function currently in scope");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001679 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001680 SPIRV_CROSS_THROW("No block currently in scope");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001681
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001682 uint32_t result_type = ops[0];
1683 uint32_t id = ops[1];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001684
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001685 // Instead of a temporary, create a new function-wide temporary with this ID instead.
1686 auto &var = set<SPIRVariable>(id, result_type, spv::StorageClassFunction);
1687 var.phi_variable = true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001688
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001689 current_function->add_local_variable(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001690
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001691 for (uint32_t i = 2; i + 2 <= length; i += 2)
1692 current_block->phi_variables.push_back({ ops[i], ops[i + 1], id });
1693 break;
1694 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001695
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001696 // Constants
1697 case OpSpecConstant:
1698 case OpConstant:
1699 {
1700 uint32_t id = ops[1];
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001701 auto &type = get<SPIRType>(ops[0]);
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001702
Hans-Kristian Arntzenfc2230f2016-07-27 11:27:00 +02001703 if (type.width > 32)
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001704 set<SPIRConstant>(id, ops[0], ops[2] | (uint64_t(ops[3]) << 32), op == OpSpecConstant);
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001705 else
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001706 set<SPIRConstant>(id, ops[0], ops[2], op == OpSpecConstant);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001707 break;
1708 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001709
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001710 case OpSpecConstantFalse:
1711 case OpConstantFalse:
1712 {
1713 uint32_t id = ops[1];
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001714 set<SPIRConstant>(id, ops[0], uint32_t(0), op == OpSpecConstantFalse);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001715 break;
1716 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001717
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001718 case OpSpecConstantTrue:
1719 case OpConstantTrue:
1720 {
1721 uint32_t id = ops[1];
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001722 set<SPIRConstant>(id, ops[0], uint32_t(1), op == OpSpecConstantTrue);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001723 break;
1724 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001725
Hans-Kristian Arntzen2abdc132017-08-02 10:33:03 +02001726 case OpConstantNull:
1727 {
1728 uint32_t id = ops[1];
1729 uint32_t type = ops[0];
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02001730 make_constant_null(id, type);
Hans-Kristian Arntzen2abdc132017-08-02 10:33:03 +02001731 break;
1732 }
1733
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001734 case OpSpecConstantComposite:
1735 case OpConstantComposite:
1736 {
1737 uint32_t id = ops[1];
1738 uint32_t type = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001739
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001740 auto &ctype = get<SPIRType>(type);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001741
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001742 // We can have constants which are structs and arrays.
1743 // In this case, our SPIRConstant will be a list of other SPIRConstant ids which we
1744 // can refer to.
1745 if (ctype.basetype == SPIRType::Struct || !ctype.array.empty())
1746 {
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001747 set<SPIRConstant>(id, type, ops + 2, length - 2, op == OpSpecConstantComposite);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001748 }
1749 else
1750 {
Hans-Kristian Arntzenceefae52017-09-27 16:10:29 +02001751 uint32_t elements = length - 2;
1752 if (elements > 4)
1753 SPIRV_CROSS_THROW("OpConstantComposite only supports 1, 2, 3 and 4 elements.");
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001754
1755 const SPIRConstant *c[4];
Hans-Kristian Arntzenceefae52017-09-27 16:10:29 +02001756 for (uint32_t i = 0; i < elements; i++)
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001757 c[i] = &get<SPIRConstant>(ops[2 + i]);
Hans-Kristian Arntzenceefae52017-09-27 16:10:29 +02001758 set<SPIRConstant>(id, type, c, elements, op == OpSpecConstantComposite);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001759 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001760 break;
1761 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001762
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001763 // Functions
1764 case OpFunction:
1765 {
1766 uint32_t res = ops[0];
1767 uint32_t id = ops[1];
1768 // Control
1769 uint32_t type = ops[3];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001770
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001771 if (current_function)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001772 SPIRV_CROSS_THROW("Must end a function before starting a new one!");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001773
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001774 current_function = &set<SPIRFunction>(id, res, type);
1775 break;
1776 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001777
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001778 case OpFunctionParameter:
1779 {
1780 uint32_t type = ops[0];
1781 uint32_t id = ops[1];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001782
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001783 if (!current_function)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001784 SPIRV_CROSS_THROW("Must be in a function!");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001785
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001786 current_function->add_parameter(type, id);
1787 set<SPIRVariable>(id, type, StorageClassFunction);
1788 break;
1789 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001790
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001791 case OpFunctionEnd:
1792 {
Hans-Kristian Arntzen526d06d2016-11-12 10:03:18 +01001793 if (current_block)
1794 {
1795 // Very specific error message, but seems to come up quite often.
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001796 SPIRV_CROSS_THROW(
Hans-Kristian Arntzen526d06d2016-11-12 10:03:18 +01001797 "Cannot end a function before ending the current block.\n"
1798 "Likely cause: If this SPIR-V was created from glslang HLSL, make sure the entry point is valid.");
1799 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001800 current_function = nullptr;
1801 break;
1802 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001803
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001804 // Blocks
1805 case OpLabel:
1806 {
1807 // OpLabel always starts a block.
1808 if (!current_function)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001809 SPIRV_CROSS_THROW("Blocks cannot exist outside functions!");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001810
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001811 uint32_t id = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001812
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001813 current_function->blocks.push_back(id);
1814 if (!current_function->entry_block)
1815 current_function->entry_block = id;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001816
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001817 if (current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001818 SPIRV_CROSS_THROW("Cannot start a block before ending the current block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001819
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001820 current_block = &set<SPIRBlock>(id);
1821 break;
1822 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001823
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001824 // Branch instructions end blocks.
1825 case OpBranch:
1826 {
1827 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001828 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001829
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001830 uint32_t target = ops[0];
1831 current_block->terminator = SPIRBlock::Direct;
1832 current_block->next_block = target;
1833 current_block = nullptr;
1834 break;
1835 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001836
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001837 case OpBranchConditional:
1838 {
1839 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001840 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001841
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001842 current_block->condition = ops[0];
1843 current_block->true_block = ops[1];
1844 current_block->false_block = ops[2];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001845
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001846 current_block->terminator = SPIRBlock::Select;
1847 current_block = nullptr;
1848 break;
1849 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001850
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001851 case OpSwitch:
1852 {
1853 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001854 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001855
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001856 if (current_block->merge == SPIRBlock::MergeNone)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001857 SPIRV_CROSS_THROW("Switch statement is not structured");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001858
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001859 current_block->terminator = SPIRBlock::MultiSelect;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001860
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001861 current_block->condition = ops[0];
1862 current_block->default_block = ops[1];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001863
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001864 for (uint32_t i = 2; i + 2 <= length; i += 2)
1865 current_block->cases.push_back({ ops[i], ops[i + 1] });
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001866
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001867 // If we jump to next block, make it break instead since we're inside a switch case block at that point.
1868 multiselect_merge_targets.insert(current_block->next_block);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001869
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001870 current_block = nullptr;
1871 break;
1872 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001873
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001874 case OpKill:
1875 {
1876 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001877 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001878 current_block->terminator = SPIRBlock::Kill;
1879 current_block = nullptr;
1880 break;
1881 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001882
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001883 case OpReturn:
1884 {
1885 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001886 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001887 current_block->terminator = SPIRBlock::Return;
1888 current_block = nullptr;
1889 break;
1890 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001891
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001892 case OpReturnValue:
1893 {
1894 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001895 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001896 current_block->terminator = SPIRBlock::Return;
1897 current_block->return_value = ops[0];
1898 current_block = nullptr;
1899 break;
1900 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001901
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001902 case OpUnreachable:
1903 {
1904 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001905 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001906 current_block->terminator = SPIRBlock::Unreachable;
1907 current_block = nullptr;
1908 break;
1909 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001910
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001911 case OpSelectionMerge:
1912 {
1913 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001914 SPIRV_CROSS_THROW("Trying to modify a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001915
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001916 current_block->next_block = ops[0];
1917 current_block->merge = SPIRBlock::MergeSelection;
1918 selection_merge_targets.insert(current_block->next_block);
1919 break;
1920 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001921
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001922 case OpLoopMerge:
1923 {
1924 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001925 SPIRV_CROSS_THROW("Trying to modify a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001926
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001927 current_block->merge_block = ops[0];
1928 current_block->continue_block = ops[1];
1929 current_block->merge = SPIRBlock::MergeLoop;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001930
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001931 loop_blocks.insert(current_block->self);
1932 loop_merge_targets.insert(current_block->merge_block);
Hans-Kristian Arntzen97f81ba2016-04-01 12:37:29 +02001933
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001934 // Don't add loop headers to continue blocks,
1935 // which would make it impossible branch into the loop header since
1936 // they are treated as continues.
1937 if (current_block->continue_block != current_block->self)
1938 continue_blocks.insert(current_block->continue_block);
1939 break;
1940 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001941
Hans-Kristian Arntzen7e8afa82016-10-03 15:54:02 +02001942 case OpSpecConstantOp:
1943 {
1944 if (length < 3)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001945 SPIRV_CROSS_THROW("OpSpecConstantOp not enough arguments.");
Hans-Kristian Arntzen7e8afa82016-10-03 15:54:02 +02001946
1947 uint32_t result_type = ops[0];
1948 uint32_t id = ops[1];
1949 auto spec_op = static_cast<Op>(ops[2]);
1950
1951 set<SPIRConstantOp>(id, result_type, spec_op, ops + 3, length - 3);
1952 break;
1953 }
1954
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001955 // Actual opcodes.
1956 default:
1957 {
1958 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001959 SPIRV_CROSS_THROW("Currently no block to insert opcode.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001960
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001961 current_block->ops.push_back(instruction);
1962 break;
1963 }
1964 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001965}
1966
1967bool Compiler::block_is_loop_candidate(const SPIRBlock &block, SPIRBlock::Method method) const
1968{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001969 // Tried and failed.
1970 if (block.disable_block_optimization || block.complex_continue)
1971 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001972
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001973 if (method == SPIRBlock::MergeToSelectForLoop)
1974 {
1975 // Try to detect common for loop pattern
1976 // which the code backend can use to create cleaner code.
1977 // for(;;) { if (cond) { some_body; } else { break; } }
1978 // is the pattern we're looking for.
1979 bool ret = block.terminator == SPIRBlock::Select && block.merge == SPIRBlock::MergeLoop &&
1980 block.true_block != block.merge_block && block.true_block != block.self &&
1981 block.false_block == block.merge_block;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001982
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001983 // If we have OpPhi which depends on branches which came from our own block,
1984 // we need to flush phi variables in else block instead of a trivial break,
1985 // so we cannot assume this is a for loop candidate.
1986 if (ret)
1987 {
1988 for (auto &phi : block.phi_variables)
1989 if (phi.parent == block.self)
1990 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001991
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001992 auto *merge = maybe_get<SPIRBlock>(block.merge_block);
1993 if (merge)
1994 for (auto &phi : merge->phi_variables)
1995 if (phi.parent == block.self)
1996 return false;
1997 }
1998 return ret;
1999 }
2000 else if (method == SPIRBlock::MergeToDirectForLoop)
2001 {
2002 // Empty loop header that just sets up merge target
2003 // and branches to loop body.
2004 bool ret = block.terminator == SPIRBlock::Direct && block.merge == SPIRBlock::MergeLoop && block.ops.empty();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002005
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002006 if (!ret)
2007 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002008
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002009 auto &child = get<SPIRBlock>(block.next_block);
2010 ret = child.terminator == SPIRBlock::Select && child.merge == SPIRBlock::MergeNone &&
2011 child.false_block == block.merge_block && child.true_block != block.merge_block &&
2012 child.true_block != block.self;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002013
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002014 // If we have OpPhi which depends on branches which came from our own block,
2015 // we need to flush phi variables in else block instead of a trivial break,
2016 // so we cannot assume this is a for loop candidate.
2017 if (ret)
2018 {
2019 for (auto &phi : block.phi_variables)
2020 if (phi.parent == block.self || phi.parent == child.self)
2021 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002022
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002023 for (auto &phi : child.phi_variables)
2024 if (phi.parent == block.self)
2025 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002026
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002027 auto *merge = maybe_get<SPIRBlock>(block.merge_block);
2028 if (merge)
2029 for (auto &phi : merge->phi_variables)
2030 if (phi.parent == block.self || phi.parent == child.false_block)
2031 return false;
2032 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002033
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002034 return ret;
2035 }
2036 else
2037 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002038}
2039
2040bool Compiler::block_is_outside_flow_control_from_block(const SPIRBlock &from, const SPIRBlock &to)
2041{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002042 auto *start = &from;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002043
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002044 if (start->self == to.self)
2045 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002046
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002047 // Break cycles.
2048 if (is_continue(start->self))
2049 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002050
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002051 // If our select block doesn't merge, we must break or continue in these blocks,
2052 // so if continues occur branchless within these blocks, consider them branchless as well.
2053 // This is typically used for loop control.
2054 if (start->terminator == SPIRBlock::Select && start->merge == SPIRBlock::MergeNone &&
2055 (block_is_outside_flow_control_from_block(get<SPIRBlock>(start->true_block), to) ||
2056 block_is_outside_flow_control_from_block(get<SPIRBlock>(start->false_block), to)))
2057 {
2058 return true;
2059 }
2060 else if (start->merge_block && block_is_outside_flow_control_from_block(get<SPIRBlock>(start->merge_block), to))
2061 {
2062 return true;
2063 }
2064 else if (start->next_block && block_is_outside_flow_control_from_block(get<SPIRBlock>(start->next_block), to))
2065 {
2066 return true;
2067 }
2068 else
2069 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002070}
2071
2072bool Compiler::execution_is_noop(const SPIRBlock &from, const SPIRBlock &to) const
2073{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002074 if (!execution_is_branchless(from, to))
2075 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002076
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002077 auto *start = &from;
2078 for (;;)
2079 {
2080 if (start->self == to.self)
2081 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002082
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002083 if (!start->ops.empty())
2084 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002085
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002086 start = &get<SPIRBlock>(start->next_block);
2087 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002088}
2089
2090bool Compiler::execution_is_branchless(const SPIRBlock &from, const SPIRBlock &to) const
2091{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002092 auto *start = &from;
2093 for (;;)
2094 {
2095 if (start->self == to.self)
2096 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002097
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002098 if (start->terminator == SPIRBlock::Direct && start->merge == SPIRBlock::MergeNone)
2099 start = &get<SPIRBlock>(start->next_block);
2100 else
2101 return false;
2102 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002103}
2104
Hans-Kristian Arntzen926916d2016-05-05 09:15:25 +02002105SPIRBlock::ContinueBlockType Compiler::continue_block_type(const SPIRBlock &block) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002106{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002107 // The block was deemed too complex during code emit, pick conservative fallback paths.
2108 if (block.complex_continue)
2109 return SPIRBlock::ComplexLoop;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002110
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002111 // In older glslang output continue block can be equal to the loop header.
2112 // In this case, execution is clearly branchless, so just assume a while loop header here.
2113 if (block.merge == SPIRBlock::MergeLoop)
2114 return SPIRBlock::WhileLoop;
Hans-Kristian Arntzen97f81ba2016-04-01 12:37:29 +02002115
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002116 auto &dominator = get<SPIRBlock>(block.loop_dominator);
Hans-Kristian Arntzen97f81ba2016-04-01 12:37:29 +02002117
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002118 if (execution_is_noop(block, dominator))
2119 return SPIRBlock::WhileLoop;
2120 else if (execution_is_branchless(block, dominator))
2121 return SPIRBlock::ForLoop;
2122 else
2123 {
2124 if (block.merge == SPIRBlock::MergeNone && block.terminator == SPIRBlock::Select &&
2125 block.true_block == dominator.self && block.false_block == dominator.merge_block)
2126 {
2127 return SPIRBlock::DoWhileLoop;
2128 }
2129 else
2130 return SPIRBlock::ComplexLoop;
2131 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002132}
2133
2134bool Compiler::traverse_all_reachable_opcodes(const SPIRBlock &block, OpcodeHandler &handler) const
2135{
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01002136 handler.set_current_block(block);
2137
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002138 // Ideally, perhaps traverse the CFG instead of all blocks in order to eliminate dead blocks,
2139 // but this shouldn't be a problem in practice unless the SPIR-V is doing insane things like recursing
2140 // inside dead blocks ...
2141 for (auto &i : block.ops)
2142 {
2143 auto ops = stream(i);
2144 auto op = static_cast<Op>(i.op);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002145
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002146 if (!handler.handle(op, ops, i.length))
2147 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002148
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002149 if (op == OpFunctionCall)
2150 {
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01002151 auto &func = get<SPIRFunction>(ops[2]);
2152 if (handler.follow_function_call(func))
2153 {
2154 if (!handler.begin_function_scope(ops, i.length))
2155 return false;
2156 if (!traverse_all_reachable_opcodes(get<SPIRFunction>(ops[2]), handler))
2157 return false;
2158 if (!handler.end_function_scope(ops, i.length))
2159 return false;
2160 }
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002161 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002162 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002163
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002164 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002165}
2166
2167bool Compiler::traverse_all_reachable_opcodes(const SPIRFunction &func, OpcodeHandler &handler) const
2168{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002169 for (auto block : func.blocks)
2170 if (!traverse_all_reachable_opcodes(get<SPIRBlock>(block), handler))
2171 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002172
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002173 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002174}
2175
2176uint32_t Compiler::type_struct_member_offset(const SPIRType &type, uint32_t index) const
2177{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002178 // Decoration must be set in valid SPIR-V, otherwise throw.
2179 auto &dec = meta[type.self].members.at(index);
2180 if (dec.decoration_flags & (1ull << DecorationOffset))
2181 return dec.offset;
2182 else
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01002183 SPIRV_CROSS_THROW("Struct member does not have Offset set.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002184}
2185
2186uint32_t Compiler::type_struct_member_array_stride(const SPIRType &type, uint32_t index) const
2187{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002188 // Decoration must be set in valid SPIR-V, otherwise throw.
2189 // ArrayStride is part of the array type not OpMemberDecorate.
2190 auto &dec = meta[type.member_types[index]].decoration;
2191 if (dec.decoration_flags & (1ull << DecorationArrayStride))
2192 return dec.array_stride;
2193 else
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01002194 SPIRV_CROSS_THROW("Struct member does not have ArrayStride set.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002195}
2196
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01002197uint32_t Compiler::type_struct_member_matrix_stride(const SPIRType &type, uint32_t index) const
2198{
2199 // Decoration must be set in valid SPIR-V, otherwise throw.
2200 // MatrixStride is part of OpMemberDecorate.
2201 auto &dec = meta[type.self].members[index];
2202 if (dec.decoration_flags & (1ull << DecorationMatrixStride))
2203 return dec.matrix_stride;
2204 else
2205 SPIRV_CROSS_THROW("Struct member does not have MatrixStride set.");
2206}
2207
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002208size_t Compiler::get_declared_struct_size(const SPIRType &type) const
2209{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002210 uint32_t last = uint32_t(type.member_types.size() - 1);
2211 size_t offset = type_struct_member_offset(type, last);
2212 size_t size = get_declared_struct_member_size(type, last);
2213 return offset + size;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002214}
2215
2216size_t Compiler::get_declared_struct_member_size(const SPIRType &struct_type, uint32_t index) const
2217{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002218 auto flags = get_member_decoration_mask(struct_type.self, index);
2219 auto &type = get<SPIRType>(struct_type.member_types[index]);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002220
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08002221 switch (type.basetype)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002222 {
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08002223 case SPIRType::Unknown:
2224 case SPIRType::Void:
2225 case SPIRType::Boolean: // Bools are purely logical, and cannot be used for externally visible types.
2226 case SPIRType::AtomicCounter:
2227 case SPIRType::Image:
2228 case SPIRType::SampledImage:
2229 case SPIRType::Sampler:
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01002230 SPIRV_CROSS_THROW("Querying size for object with opaque size.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002231
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08002232 default:
2233 break;
2234 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002235
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08002236 if (!type.array.empty())
2237 {
2238 // For arrays, we can use ArrayStride to get an easy check.
2239 return type_struct_member_array_stride(struct_type, index) * type.array.back();
2240 }
2241 else if (type.basetype == SPIRType::Struct)
2242 {
2243 return get_declared_struct_size(type);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002244 }
2245 else
2246 {
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002247 unsigned vecsize = type.vecsize;
2248 unsigned columns = type.columns;
Hans-Kristian Arntzen2d79d362016-11-28 15:01:36 +01002249
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08002250 // Vectors.
2251 if (columns == 1)
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01002252 {
2253 size_t component_size = type.width / 8;
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08002254 return vecsize * component_size;
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01002255 }
Hans-Kristian Arntzen2d79d362016-11-28 15:01:36 +01002256 else
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002257 {
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01002258 uint32_t matrix_stride = type_struct_member_matrix_stride(struct_type, index);
Hans-Kristian Arntzen2d79d362016-11-28 15:01:36 +01002259
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01002260 // Per SPIR-V spec, matrices must be tightly packed and aligned up for vec3 accesses.
2261 if (flags & (1ull << DecorationRowMajor))
2262 return matrix_stride * vecsize;
2263 else if (flags & (1ull << DecorationColMajor))
2264 return matrix_stride * columns;
2265 else
2266 SPIRV_CROSS_THROW("Either row-major or column-major must be declared for matrices.");
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08002267 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002268 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002269}
2270
2271bool Compiler::BufferAccessHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
2272{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002273 if (opcode != OpAccessChain && opcode != OpInBoundsAccessChain)
2274 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002275
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002276 // Invalid SPIR-V.
2277 if (length < 4)
2278 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002279
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002280 if (args[2] != id)
2281 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002282
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002283 // Don't bother traversing the entire access chain tree yet.
2284 // If we access a struct member, assume we access the entire member.
2285 uint32_t index = compiler.get<SPIRConstant>(args[3]).scalar();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002286
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002287 // Seen this index already.
2288 if (seen.find(index) != end(seen))
2289 return true;
2290 seen.insert(index);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002291
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002292 auto &type = compiler.expression_type(id);
2293 uint32_t offset = compiler.type_struct_member_offset(type, index);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002294
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002295 size_t range;
2296 // If we have another member in the struct, deduce the range by looking at the next member.
2297 // This is okay since structs in SPIR-V can have padding, but Offset decoration must be
2298 // monotonically increasing.
2299 // Of course, this doesn't take into account if the SPIR-V for some reason decided to add
2300 // very large amounts of padding, but that's not really a big deal.
2301 if (index + 1 < type.member_types.size())
2302 {
2303 range = compiler.type_struct_member_offset(type, index + 1) - offset;
2304 }
2305 else
2306 {
2307 // No padding, so just deduce it from the size of the member directly.
2308 range = compiler.get_declared_struct_member_size(type, index);
2309 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002310
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002311 ranges.push_back({ index, offset, range });
2312 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002313}
2314
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002315std::vector<BufferRange> Compiler::get_active_buffer_ranges(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002316{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002317 std::vector<BufferRange> ranges;
2318 BufferAccessHandler handler(*this, ranges, id);
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002319 traverse_all_reachable_opcodes(get<SPIRFunction>(entry_point), handler);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002320 return ranges;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002321}
2322
Bill Hollings103aabf2016-04-06 17:42:27 -04002323// Increase the number of IDs by the specified incremental amount.
2324// Returns the value of the first ID available for use in the expanded bound.
2325uint32_t Compiler::increase_bound_by(uint32_t incr_amount)
2326{
Hans-Kristian Arntzen14bc1ff2016-09-10 18:07:52 +02002327 auto curr_bound = ids.size();
2328 auto new_bound = curr_bound + incr_amount;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002329 ids.resize(new_bound);
2330 meta.resize(new_bound);
Hans-Kristian Arntzen14bc1ff2016-09-10 18:07:52 +02002331 return uint32_t(curr_bound);
Bill Hollings103aabf2016-04-06 17:42:27 -04002332}
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02002333
2334bool Compiler::types_are_logically_equivalent(const SPIRType &a, const SPIRType &b) const
2335{
2336 if (a.basetype != b.basetype)
2337 return false;
2338 if (a.width != b.width)
2339 return false;
2340 if (a.vecsize != b.vecsize)
2341 return false;
2342 if (a.columns != b.columns)
2343 return false;
2344 if (a.array.size() != b.array.size())
2345 return false;
2346
Hans-Kristian Arntzen46892ab2016-05-23 13:45:20 +02002347 size_t array_count = a.array.size();
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02002348 if (array_count && memcmp(a.array.data(), b.array.data(), array_count * sizeof(uint32_t)) != 0)
2349 return false;
2350
2351 if (a.basetype == SPIRType::Image || a.basetype == SPIRType::SampledImage)
2352 {
2353 if (memcmp(&a.image, &b.image, sizeof(SPIRType::Image)) != 0)
2354 return false;
2355 }
2356
2357 if (a.member_types.size() != b.member_types.size())
2358 return false;
2359
Hans-Kristian Arntzen46892ab2016-05-23 13:45:20 +02002360 size_t member_types = a.member_types.size();
2361 for (size_t i = 0; i < member_types; i++)
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02002362 {
2363 if (!types_are_logically_equivalent(get<SPIRType>(a.member_types[i]), get<SPIRType>(b.member_types[i])))
2364 return false;
2365 }
2366
2367 return true;
2368}
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002369
2370uint64_t Compiler::get_execution_mode_mask() const
2371{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002372 return get_entry_point().flags;
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002373}
2374
2375void Compiler::set_execution_mode(ExecutionMode mode, uint32_t arg0, uint32_t arg1, uint32_t arg2)
2376{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002377 auto &execution = get_entry_point();
2378
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002379 execution.flags |= 1ull << mode;
2380 switch (mode)
2381 {
2382 case ExecutionModeLocalSize:
2383 execution.workgroup_size.x = arg0;
2384 execution.workgroup_size.y = arg1;
2385 execution.workgroup_size.z = arg2;
2386 break;
2387
2388 case ExecutionModeInvocations:
2389 execution.invocations = arg0;
2390 break;
2391
2392 case ExecutionModeOutputVertices:
2393 execution.output_vertices = arg0;
2394 break;
2395
2396 default:
2397 break;
2398 }
2399}
2400
2401void Compiler::unset_execution_mode(ExecutionMode mode)
2402{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002403 auto &execution = get_entry_point();
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002404 execution.flags &= ~(1ull << mode);
2405}
2406
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +02002407uint32_t Compiler::get_work_group_size_specialization_constants(SpecializationConstant &x, SpecializationConstant &y,
2408 SpecializationConstant &z) const
2409{
2410 auto &execution = get_entry_point();
Hans-Kristian Arntzenca69b612017-11-06 09:49:52 +01002411 x = { 0, 0 };
2412 y = { 0, 0 };
2413 z = { 0, 0 };
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +02002414
2415 if (execution.workgroup_size.constant != 0)
2416 {
2417 auto &c = get<SPIRConstant>(execution.workgroup_size.constant);
2418
2419 if (c.m.c[0].id[0] != 0)
2420 {
2421 x.id = c.m.c[0].id[0];
2422 x.constant_id = get_decoration(c.m.c[0].id[0], DecorationSpecId);
2423 }
2424
2425 if (c.m.c[0].id[1] != 0)
2426 {
2427 y.id = c.m.c[0].id[1];
2428 y.constant_id = get_decoration(c.m.c[0].id[1], DecorationSpecId);
2429 }
2430
2431 if (c.m.c[0].id[2] != 0)
2432 {
2433 z.id = c.m.c[0].id[2];
2434 z.constant_id = get_decoration(c.m.c[0].id[2], DecorationSpecId);
2435 }
2436 }
2437
2438 return execution.workgroup_size.constant;
2439}
2440
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002441uint32_t Compiler::get_execution_mode_argument(spv::ExecutionMode mode, uint32_t index) const
2442{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002443 auto &execution = get_entry_point();
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002444 switch (mode)
2445 {
2446 case ExecutionModeLocalSize:
2447 switch (index)
2448 {
2449 case 0:
2450 return execution.workgroup_size.x;
2451 case 1:
2452 return execution.workgroup_size.y;
2453 case 2:
2454 return execution.workgroup_size.z;
2455 default:
2456 return 0;
2457 }
2458
2459 case ExecutionModeInvocations:
2460 return execution.invocations;
2461
2462 case ExecutionModeOutputVertices:
2463 return execution.output_vertices;
2464
2465 default:
2466 return 0;
2467 }
2468}
2469
2470ExecutionModel Compiler::get_execution_model() const
2471{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002472 auto &execution = get_entry_point();
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002473 return execution.model;
2474}
Hans-Kristian Arntzen8e63c772016-07-06 09:58:01 +02002475
2476void Compiler::set_remapped_variable_state(uint32_t id, bool remap_enable)
2477{
2478 get<SPIRVariable>(id).remapped_variable = remap_enable;
2479}
2480
2481bool Compiler::get_remapped_variable_state(uint32_t id) const
2482{
2483 return get<SPIRVariable>(id).remapped_variable;
2484}
Hans-Kristian Arntzen078eec52016-07-06 11:04:06 +02002485
2486void Compiler::set_subpass_input_remapped_components(uint32_t id, uint32_t components)
2487{
2488 get<SPIRVariable>(id).remapped_components = components;
2489}
2490
2491uint32_t Compiler::get_subpass_input_remapped_components(uint32_t id) const
2492{
2493 return get<SPIRVariable>(id).remapped_components;
2494}
Hans-Kristian Arntzen36a0b632016-07-12 14:33:04 +02002495
2496void Compiler::inherit_expression_dependencies(uint32_t dst, uint32_t source_expression)
2497{
Hans-Kristian Arntzen75391f92017-03-20 22:38:05 +01002498 // Don't inherit any expression dependencies if the expression in dst
2499 // is not a forwarded temporary.
2500 if (forwarded_temporaries.find(dst) == end(forwarded_temporaries) ||
2501 forced_temporaries.find(dst) != end(forced_temporaries))
2502 {
2503 return;
2504 }
2505
Hans-Kristian Arntzen36a0b632016-07-12 14:33:04 +02002506 auto &e = get<SPIRExpression>(dst);
2507 auto *s = maybe_get<SPIRExpression>(source_expression);
2508 if (!s)
2509 return;
2510
2511 auto &e_deps = e.expression_dependencies;
2512 auto &s_deps = s->expression_dependencies;
2513
2514 // If we depend on a expression, we also depend on all sub-dependencies from source.
2515 e_deps.push_back(source_expression);
2516 e_deps.insert(end(e_deps), begin(s_deps), end(s_deps));
2517
2518 // Eliminate duplicated dependencies.
2519 e_deps.erase(unique(begin(e_deps), end(e_deps)), end(e_deps));
2520}
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002521
2522vector<string> Compiler::get_entry_points() const
2523{
2524 vector<string> entries;
2525 for (auto &entry : entry_points)
2526 entries.push_back(entry.second.name);
2527 return entries;
2528}
2529
2530void Compiler::set_entry_point(const std::string &name)
2531{
2532 auto &entry = get_entry_point(name);
2533 entry_point = entry.self;
2534}
2535
2536SPIREntryPoint &Compiler::get_entry_point(const std::string &name)
2537{
2538 auto itr =
2539 find_if(begin(entry_points), end(entry_points),
2540 [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool { return entry.second.name == name; });
2541
2542 if (itr == end(entry_points))
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01002543 SPIRV_CROSS_THROW("Entry point does not exist.");
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002544
2545 return itr->second;
2546}
2547
2548const SPIREntryPoint &Compiler::get_entry_point(const std::string &name) const
2549{
2550 auto itr =
2551 find_if(begin(entry_points), end(entry_points),
2552 [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool { return entry.second.name == name; });
2553
2554 if (itr == end(entry_points))
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01002555 SPIRV_CROSS_THROW("Entry point does not exist.");
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002556
2557 return itr->second;
2558}
2559
2560const SPIREntryPoint &Compiler::get_entry_point() const
2561{
2562 return entry_points.find(entry_point)->second;
2563}
2564
2565SPIREntryPoint &Compiler::get_entry_point()
2566{
2567 return entry_points.find(entry_point)->second;
2568}
2569
2570bool Compiler::interface_variable_exists_in_entry_point(uint32_t id) const
2571{
2572 auto &var = get<SPIRVariable>(id);
Robert Konrada778c362017-01-15 16:39:03 +01002573 if (var.storage != StorageClassInput && var.storage != StorageClassOutput &&
2574 var.storage != StorageClassUniformConstant)
Hans-Kristian Arntzen24df8f02017-02-04 10:26:26 +01002575 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 +02002576
2577 // This is to avoid potential problems with very old glslang versions which did
2578 // not emit input/output interfaces properly.
2579 // We can assume they only had a single entry point, and single entry point
2580 // shaders could easily be assumed to use every interface variable anyways.
2581 if (entry_points.size() <= 1)
2582 return true;
2583
2584 auto &execution = get_entry_point();
2585 return find(begin(execution.interface_variables), end(execution.interface_variables), id) !=
2586 end(execution.interface_variables);
2587}
Hans-Kristian Arntzenbcb55602016-09-10 13:56:36 +02002588
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002589void Compiler::CombinedImageSamplerHandler::push_remap_parameters(const SPIRFunction &func, const uint32_t *args,
2590 uint32_t length)
2591{
2592 // If possible, pipe through a remapping table so that parameters know
2593 // which variables they actually bind to in this scope.
2594 unordered_map<uint32_t, uint32_t> remapping;
2595 for (uint32_t i = 0; i < length; i++)
2596 remapping[func.arguments[i].id] = remap_parameter(args[i]);
2597 parameter_remapping.push(move(remapping));
2598}
2599
2600void Compiler::CombinedImageSamplerHandler::pop_remap_parameters()
2601{
2602 parameter_remapping.pop();
2603}
2604
2605uint32_t Compiler::CombinedImageSamplerHandler::remap_parameter(uint32_t id)
2606{
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002607 auto *var = compiler.maybe_get_backing_variable(id);
2608 if (var)
2609 id = var->self;
2610
Hans-Kristian Arntzen901b45e2016-09-10 22:21:57 +02002611 if (parameter_remapping.empty())
2612 return id;
2613
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002614 auto &remapping = parameter_remapping.top();
2615 auto itr = remapping.find(id);
2616 if (itr != end(remapping))
2617 return itr->second;
2618 else
2619 return id;
2620}
2621
2622bool Compiler::CombinedImageSamplerHandler::begin_function_scope(const uint32_t *args, uint32_t length)
2623{
2624 if (length < 3)
2625 return false;
2626
2627 auto &callee = compiler.get<SPIRFunction>(args[2]);
2628 args += 3;
2629 length -= 3;
2630 push_remap_parameters(callee, args, length);
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002631 functions.push(&callee);
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002632 return true;
2633}
2634
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002635bool Compiler::CombinedImageSamplerHandler::end_function_scope(const uint32_t *args, uint32_t length)
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002636{
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002637 if (length < 3)
2638 return false;
2639
2640 auto &callee = compiler.get<SPIRFunction>(args[2]);
2641 args += 3;
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002642
2643 // There are two types of cases we have to handle,
2644 // a callee might call sampler2D(texture2D, sampler) directly where
2645 // one or more parameters originate from parameters.
2646 // Alternatively, we need to provide combined image samplers to our callees,
2647 // and in this case we need to add those as well.
2648
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002649 pop_remap_parameters();
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002650
2651 // Our callee has now been processed at least once.
2652 // No point in doing it again.
2653 callee.do_combined_parameters = false;
2654
2655 auto &params = functions.top()->combined_parameters;
2656 functions.pop();
2657 if (functions.empty())
2658 return true;
2659
2660 auto &caller = *functions.top();
2661 if (caller.do_combined_parameters)
2662 {
2663 for (auto &param : params)
2664 {
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002665 uint32_t image_id = param.global_image ? param.image_id : args[param.image_id];
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002666 uint32_t sampler_id = param.global_sampler ? param.sampler_id : args[param.sampler_id];
2667
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002668 auto *i = compiler.maybe_get_backing_variable(image_id);
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002669 auto *s = compiler.maybe_get_backing_variable(sampler_id);
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002670 if (i)
2671 image_id = i->self;
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002672 if (s)
2673 sampler_id = s->self;
2674
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002675 register_combined_image_sampler(caller, image_id, sampler_id, param.depth);
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002676 }
2677 }
2678
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002679 return true;
2680}
2681
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002682void Compiler::CombinedImageSamplerHandler::register_combined_image_sampler(SPIRFunction &caller, uint32_t image_id,
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002683 uint32_t sampler_id, bool depth)
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002684{
2685 // We now have a texture ID and a sampler ID which will either be found as a global
2686 // or a parameter in our own function. If both are global, they will not need a parameter,
2687 // otherwise, add it to our list.
2688 SPIRFunction::CombinedImageSamplerParameter param = {
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002689 0u, image_id, sampler_id, true, true, depth,
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002690 };
2691
2692 auto texture_itr = find_if(begin(caller.arguments), end(caller.arguments),
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002693 [image_id](const SPIRFunction::Parameter &p) { return p.id == image_id; });
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002694 auto sampler_itr = find_if(begin(caller.arguments), end(caller.arguments),
2695 [sampler_id](const SPIRFunction::Parameter &p) { return p.id == sampler_id; });
2696
2697 if (texture_itr != end(caller.arguments))
2698 {
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002699 param.global_image = false;
Hans-Kristian Arntzen7630d3c2016-11-21 12:14:02 +01002700 param.image_id = uint32_t(texture_itr - begin(caller.arguments));
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002701 }
2702
2703 if (sampler_itr != end(caller.arguments))
2704 {
2705 param.global_sampler = false;
Hans-Kristian Arntzen7630d3c2016-11-21 12:14:02 +01002706 param.sampler_id = uint32_t(sampler_itr - begin(caller.arguments));
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002707 }
2708
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002709 if (param.global_image && param.global_sampler)
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002710 return;
2711
2712 auto itr = find_if(begin(caller.combined_parameters), end(caller.combined_parameters),
2713 [&param](const SPIRFunction::CombinedImageSamplerParameter &p) {
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002714 return param.image_id == p.image_id && param.sampler_id == p.sampler_id &&
2715 param.global_image == p.global_image && param.global_sampler == p.global_sampler;
Hans-Kristian Arntzen1079e792017-10-10 10:22:40 +02002716 });
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002717
2718 if (itr == end(caller.combined_parameters))
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002719 {
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002720 uint32_t id = compiler.increase_bound_by(3);
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002721 auto type_id = id + 0;
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002722 auto ptr_type_id = id + 1;
2723 auto combined_id = id + 2;
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002724 auto &base = compiler.expression_type(image_id);
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002725 auto &type = compiler.set<SPIRType>(type_id);
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002726 auto &ptr_type = compiler.set<SPIRType>(ptr_type_id);
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002727
2728 type = base;
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002729 type.self = type_id;
2730 type.basetype = SPIRType::SampledImage;
2731 type.pointer = false;
2732 type.storage = StorageClassGeneric;
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002733 type.image.depth = depth;
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002734
2735 ptr_type = type;
2736 ptr_type.pointer = true;
2737 ptr_type.storage = StorageClassUniformConstant;
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002738
2739 // Build new variable.
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002740 compiler.set<SPIRVariable>(combined_id, ptr_type_id, StorageClassFunction, 0);
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002741
2742 // Inherit RelaxedPrecision (and potentially other useful flags if deemed relevant).
2743 auto &new_flags = compiler.meta[combined_id].decoration.decoration_flags;
2744 auto old_flags = compiler.meta[sampler_id].decoration.decoration_flags;
2745 new_flags = old_flags & (1ull << DecorationRelaxedPrecision);
2746
2747 param.id = combined_id;
2748
2749 compiler.set_name(combined_id,
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002750 join("SPIRV_Cross_Combined", compiler.to_name(image_id), compiler.to_name(sampler_id)));
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002751
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002752 caller.combined_parameters.push_back(param);
Hans-Kristian Arntzen9cb86162017-02-05 10:50:14 +01002753 caller.shadow_arguments.push_back({ ptr_type_id, combined_id, 0u, 0u, true });
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002754 }
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002755}
2756
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002757bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
2758{
2759 // We need to figure out where samplers and images are loaded from, so do only the bare bones compilation we need.
2760 switch (opcode)
2761 {
2762 case OpLoad:
2763 {
2764 if (length < 3)
2765 return false;
2766
2767 uint32_t result_type = args[0];
2768
2769 auto &type = compiler.get<SPIRType>(result_type);
2770 bool separate_image = type.basetype == SPIRType::Image && type.image.sampled == 1;
2771 bool separate_sampler = type.basetype == SPIRType::Sampler;
2772
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002773 // If not separate image or sampler, don't bother.
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002774 if (!separate_image && !separate_sampler)
2775 return true;
2776
2777 uint32_t id = args[1];
2778 uint32_t ptr = args[2];
2779 compiler.set<SPIRExpression>(id, "", result_type, true);
2780 compiler.register_read(id, ptr, true);
2781 return true;
2782 }
2783
2784 case OpInBoundsAccessChain:
2785 case OpAccessChain:
2786 {
2787 if (length < 3)
2788 return false;
2789
2790 // Technically, it is possible to have arrays of textures and arrays of samplers and combine them, but this becomes essentially
2791 // impossible to implement, since we don't know which concrete sampler we are accessing.
2792 // One potential way is to create a combinatorial explosion where N textures and M samplers are combined into N * M sampler2Ds,
2793 // but this seems ridiculously complicated for a problem which is easy to work around.
2794 // Checking access chains like this assumes we don't have samplers or textures inside uniform structs, but this makes no sense.
2795
2796 auto &type = compiler.get<SPIRType>(args[0]);
2797 bool separate_image = type.basetype == SPIRType::Image && type.image.sampled == 1;
2798 bool separate_sampler = type.basetype == SPIRType::Sampler;
2799 if (separate_image)
Hans-Kristian Arntzen36e1c472017-05-06 13:59:00 +02002800 SPIRV_CROSS_THROW("Attempting to use arrays or structs of separate images. This is not possible to "
2801 "statically remap to plain GLSL.");
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002802 if (separate_sampler)
Hans-Kristian Arntzen36e1c472017-05-06 13:59:00 +02002803 SPIRV_CROSS_THROW(
2804 "Attempting to use arrays or structs of separate samplers. This is not possible to statically "
2805 "remap to plain GLSL.");
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002806 return true;
2807 }
2808
2809 case OpSampledImage:
2810 // Do it outside.
2811 break;
2812
2813 default:
2814 return true;
2815 }
2816
2817 if (length < 4)
2818 return false;
2819
Hans-Kristian Arntzendfb65972016-09-11 12:05:20 +02002820 // Registers sampler2D calls used in case they are parameters so
2821 // that their callees know which combined image samplers to propagate down the call stack.
2822 if (!functions.empty())
2823 {
2824 auto &callee = *functions.top();
2825 if (callee.do_combined_parameters)
2826 {
2827 uint32_t image_id = args[2];
2828
2829 auto *image = compiler.maybe_get_backing_variable(image_id);
2830 if (image)
2831 image_id = image->self;
2832
2833 uint32_t sampler_id = args[3];
2834 auto *sampler = compiler.maybe_get_backing_variable(sampler_id);
2835 if (sampler)
2836 sampler_id = sampler->self;
2837
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002838 auto &combined_type = compiler.get<SPIRType>(args[0]);
2839 register_combined_image_sampler(callee, image_id, sampler_id, combined_type.image.depth);
Hans-Kristian Arntzendfb65972016-09-11 12:05:20 +02002840 }
2841 }
2842
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002843 // For function calls, we need to remap IDs which are function parameters into global variables.
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002844 // This information is statically known from the current place in the call stack.
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002845 // Function parameters are not necessarily pointers, so if we don't have a backing variable, remapping will know
2846 // which backing variable the image/sample came from.
Hans-Kristian Arntzendfb65972016-09-11 12:05:20 +02002847 uint32_t image_id = remap_parameter(args[2]);
2848 uint32_t sampler_id = remap_parameter(args[3]);
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002849
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002850 auto itr = find_if(begin(compiler.combined_image_samplers), end(compiler.combined_image_samplers),
2851 [image_id, sampler_id](const CombinedImageSampler &combined) {
2852 return combined.image_id == image_id && combined.sampler_id == sampler_id;
Hans-Kristian Arntzen1079e792017-10-10 10:22:40 +02002853 });
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002854
2855 if (itr == end(compiler.combined_image_samplers))
2856 {
2857 auto id = compiler.increase_bound_by(2);
2858 auto type_id = id + 0;
2859 auto combined_id = id + 1;
2860 auto sampled_type = args[0];
2861
2862 // Make a new type, pointer to OpTypeSampledImage, so we can make a variable of this type.
2863 // We will probably have this type lying around, but it doesn't hurt to make duplicates for internal purposes.
2864 auto &type = compiler.set<SPIRType>(type_id);
2865 auto &base = compiler.get<SPIRType>(sampled_type);
2866 type = base;
2867 type.pointer = true;
2868 type.storage = StorageClassUniformConstant;
2869
2870 // Build new variable.
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002871 compiler.set<SPIRVariable>(combined_id, type_id, StorageClassUniformConstant, 0);
Hans-Kristian Arntzen14bc1ff2016-09-10 18:07:52 +02002872
2873 // Inherit RelaxedPrecision (and potentially other useful flags if deemed relevant).
2874 auto &new_flags = compiler.meta[combined_id].decoration.decoration_flags;
2875 auto old_flags = compiler.meta[sampler_id].decoration.decoration_flags;
2876 new_flags = old_flags & (1ull << DecorationRelaxedPrecision);
2877
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002878 compiler.combined_image_samplers.push_back({ combined_id, image_id, sampler_id });
2879 }
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002880
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002881 return true;
2882}
2883
Hans-Kristian Arntzenbcb55602016-09-10 13:56:36 +02002884void Compiler::build_combined_image_samplers()
2885{
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002886 for (auto &id : ids)
2887 {
2888 if (id.get_type() == TypeFunction)
2889 {
2890 auto &func = id.get<SPIRFunction>();
2891 func.combined_parameters.clear();
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002892 func.shadow_arguments.clear();
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002893 func.do_combined_parameters = true;
2894 }
2895 }
2896
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002897 combined_image_samplers.clear();
2898 CombinedImageSamplerHandler handler(*this);
2899 traverse_all_reachable_opcodes(get<SPIRFunction>(entry_point), handler);
Hans-Kristian Arntzenbcb55602016-09-10 13:56:36 +02002900}
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02002901
2902vector<SpecializationConstant> Compiler::get_specialization_constants() const
2903{
2904 vector<SpecializationConstant> spec_consts;
2905 for (auto &id : ids)
2906 {
2907 if (id.get_type() == TypeConstant)
2908 {
2909 auto &c = id.get<SPIRConstant>();
2910 if (c.specialization)
2911 {
2912 spec_consts.push_back({ c.self, get_decoration(c.self, DecorationSpecId) });
2913 }
2914 }
2915 }
2916 return spec_consts;
2917}
2918
2919SPIRConstant &Compiler::get_constant(uint32_t id)
2920{
2921 return get<SPIRConstant>(id);
2922}
2923
2924const SPIRConstant &Compiler::get_constant(uint32_t id) const
2925{
2926 return get<SPIRConstant>(id);
2927}
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01002928
Hans-Kristian Arntzenbf5c0752017-03-25 16:28:44 +01002929static 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 +01002930{
2931 // This block accesses the variable.
2932 if (blocks.find(block) != end(blocks))
2933 return false;
2934
2935 // We are at the end of the CFG.
2936 if (cfg.get_succeeding_edges(block).empty())
2937 return true;
2938
2939 // If any of our successors have a path to the end, there exists a path from block.
2940 for (auto &succ : cfg.get_succeeding_edges(block))
2941 if (exists_unaccessed_path_to_return(cfg, succ, blocks))
2942 return true;
2943
2944 return false;
2945}
2946
Hans-Kristian Arntzenbf5c0752017-03-25 16:28:44 +01002947void Compiler::analyze_parameter_preservation(
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02002948 SPIRFunction &entry, const CFG &cfg, const unordered_map<uint32_t, unordered_set<uint32_t>> &variable_to_blocks,
2949 const unordered_map<uint32_t, unordered_set<uint32_t>> &complete_write_blocks)
Hans-Kristian Arntzenb2c2e642017-03-25 16:25:30 +01002950{
2951 for (auto &arg : entry.arguments)
2952 {
2953 // Non-pointers are always inputs.
2954 auto &type = get<SPIRType>(arg.type);
2955 if (!type.pointer)
2956 continue;
2957
2958 // Opaque argument types are always in
2959 bool potential_preserve;
2960 switch (type.basetype)
2961 {
2962 case SPIRType::Sampler:
2963 case SPIRType::Image:
2964 case SPIRType::SampledImage:
2965 case SPIRType::AtomicCounter:
2966 potential_preserve = false;
2967 break;
2968
2969 default:
2970 potential_preserve = true;
2971 break;
2972 }
2973
2974 if (!potential_preserve)
2975 continue;
2976
2977 auto itr = variable_to_blocks.find(arg.id);
2978 if (itr == end(variable_to_blocks))
2979 {
2980 // Variable is never accessed.
2981 continue;
2982 }
2983
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02002984 // We have accessed a variable, but there was no complete writes to that variable.
2985 // We deduce that we must preserve the argument.
2986 itr = complete_write_blocks.find(arg.id);
2987 if (itr == end(complete_write_blocks))
2988 {
2989 arg.read_count++;
2990 continue;
2991 }
2992
2993 // 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 +01002994 // when the function returns. We therefore need to implicitly preserve the variable in case there are writers in the function.
2995 // Major case here is if a function is
2996 // void foo(int &var) { if (cond) var = 10; }
2997 // Using read/write counts, we will think it's just an out variable, but it really needs to be inout,
2998 // because if we don't write anything whatever we put into the function must return back to the caller.
2999 if (exists_unaccessed_path_to_return(cfg, entry.entry_block, itr->second))
3000 arg.read_count++;
3001 }
3002}
3003
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003004void Compiler::analyze_variable_scope(SPIRFunction &entry)
3005{
3006 struct AccessHandler : OpcodeHandler
3007 {
3008 public:
3009 AccessHandler(Compiler &compiler_)
3010 : compiler(compiler_)
3011 {
3012 }
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +01003013
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003014 bool follow_function_call(const SPIRFunction &)
3015 {
3016 // Only analyze within this function.
3017 return false;
3018 }
3019
3020 void set_current_block(const SPIRBlock &block)
3021 {
3022 current_block = &block;
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003023
3024 // If we're branching to a block which uses OpPhi, in GLSL
3025 // this will be a variable write when we branch,
3026 // so we need to track access to these variables as well to
3027 // have a complete picture.
3028 const auto test_phi = [this, &block](uint32_t to) {
3029 auto &next = compiler.get<SPIRBlock>(to);
3030 for (auto &phi : next.phi_variables)
Hans-Kristian Arntzena1c0ab62017-08-21 09:36:53 +02003031 {
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003032 if (phi.parent == block.self)
Hans-Kristian Arntzena1c0ab62017-08-21 09:36:53 +02003033 {
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003034 accessed_variables_to_block[phi.function_variable].insert(block.self);
Hans-Kristian Arntzena1c0ab62017-08-21 09:36:53 +02003035 // Phi variables are also accessed in our target branch block.
3036 accessed_variables_to_block[phi.function_variable].insert(next.self);
3037 }
3038 }
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003039 };
3040
3041 switch (block.terminator)
3042 {
3043 case SPIRBlock::Direct:
3044 test_phi(block.next_block);
3045 break;
3046
3047 case SPIRBlock::Select:
3048 test_phi(block.true_block);
3049 test_phi(block.false_block);
3050 break;
3051
3052 case SPIRBlock::MultiSelect:
3053 for (auto &target : block.cases)
3054 test_phi(target.block);
3055 if (block.default_block)
3056 test_phi(block.default_block);
3057 break;
3058
3059 default:
3060 break;
3061 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003062 }
3063
3064 bool handle(spv::Op op, const uint32_t *args, uint32_t length)
3065 {
3066 switch (op)
3067 {
3068 case OpStore:
3069 {
3070 if (length < 2)
3071 return false;
3072
3073 uint32_t ptr = args[0];
3074 auto *var = compiler.maybe_get_backing_variable(ptr);
3075 if (var && var->storage == StorageClassFunction)
3076 accessed_variables_to_block[var->self].insert(current_block->self);
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02003077
3078 // If we store through an access chain, we have a partial write.
3079 if (var && var->self == ptr && var->storage == StorageClassFunction)
3080 complete_write_variables_to_block[var->self].insert(current_block->self);
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003081 break;
3082 }
3083
3084 case OpAccessChain:
3085 case OpInBoundsAccessChain:
3086 {
3087 if (length < 3)
3088 return false;
3089
3090 uint32_t ptr = args[2];
3091 auto *var = compiler.maybe_get<SPIRVariable>(ptr);
3092 if (var && var->storage == StorageClassFunction)
3093 accessed_variables_to_block[var->self].insert(current_block->self);
3094 break;
3095 }
3096
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +01003097 case OpCopyMemory:
3098 {
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003099 if (length < 2)
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +01003100 return false;
3101
3102 uint32_t lhs = args[0];
3103 uint32_t rhs = args[1];
3104 auto *var = compiler.maybe_get_backing_variable(lhs);
3105 if (var && var->storage == StorageClassFunction)
3106 accessed_variables_to_block[var->self].insert(current_block->self);
3107
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02003108 // If we store through an access chain, we have a partial write.
3109 if (var->self == lhs)
3110 complete_write_variables_to_block[var->self].insert(current_block->self);
3111
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +01003112 var = compiler.maybe_get_backing_variable(rhs);
3113 if (var && var->storage == StorageClassFunction)
3114 accessed_variables_to_block[var->self].insert(current_block->self);
3115 break;
3116 }
3117
3118 case OpCopyObject:
3119 {
3120 if (length < 3)
3121 return false;
3122
3123 auto *var = compiler.maybe_get_backing_variable(args[2]);
3124 if (var && var->storage == StorageClassFunction)
3125 accessed_variables_to_block[var->self].insert(current_block->self);
3126 break;
3127 }
3128
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003129 case OpLoad:
3130 {
3131 if (length < 3)
3132 return false;
3133 uint32_t ptr = args[2];
3134 auto *var = compiler.maybe_get_backing_variable(ptr);
3135 if (var && var->storage == StorageClassFunction)
3136 accessed_variables_to_block[var->self].insert(current_block->self);
3137 break;
3138 }
3139
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +01003140 case OpFunctionCall:
3141 {
3142 if (length < 3)
3143 return false;
3144
3145 length -= 3;
3146 args += 3;
3147 for (uint32_t i = 0; i < length; i++)
3148 {
3149 auto *var = compiler.maybe_get_backing_variable(args[i]);
3150 if (var && var->storage == StorageClassFunction)
3151 accessed_variables_to_block[var->self].insert(current_block->self);
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02003152
3153 // Cannot easily prove if argument we pass to a function is completely written.
3154 // Usually, functions write to a dummy variable,
3155 // which is then copied to in full to the real argument.
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +01003156 }
3157 break;
3158 }
3159
3160 case OpPhi:
3161 {
3162 if (length < 2)
3163 return false;
3164
3165 // Phi nodes are implemented as function variables, so register an access here.
3166 accessed_variables_to_block[args[1]].insert(current_block->self);
3167 break;
3168 }
3169
Hans-Kristian Arntzen1079e792017-10-10 10:22:40 +02003170 // Atomics shouldn't be able to access function-local variables.
3171 // Some GLSL builtins access a pointer.
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003172
3173 default:
3174 break;
3175 }
3176 return true;
3177 }
3178
3179 Compiler &compiler;
3180 std::unordered_map<uint32_t, std::unordered_set<uint32_t>> accessed_variables_to_block;
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02003181 std::unordered_map<uint32_t, std::unordered_set<uint32_t>> complete_write_variables_to_block;
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003182 const SPIRBlock *current_block = nullptr;
3183 } handler(*this);
3184
3185 // First, we map out all variable access within a function.
3186 // Essentially a map of block -> { variables accessed in the basic block }
Hans-Kristian Arntzen7630d3c2016-11-21 12:14:02 +01003187 this->traverse_all_reachable_opcodes(entry, handler);
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003188
3189 // Compute the control flow graph for this function.
3190 CFG cfg(*this, entry);
3191
Hans-Kristian Arntzenb2c2e642017-03-25 16:25:30 +01003192 // Analyze if there are parameters which need to be implicitly preserved with an "in" qualifier.
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02003193 analyze_parameter_preservation(entry, cfg, handler.accessed_variables_to_block,
3194 handler.complete_write_variables_to_block);
Hans-Kristian Arntzenb2c2e642017-03-25 16:25:30 +01003195
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003196 unordered_map<uint32_t, uint32_t> potential_loop_variables;
3197
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003198 // For each variable which is statically accessed.
3199 for (auto &var : handler.accessed_variables_to_block)
3200 {
3201 DominatorBuilder builder(cfg);
3202 auto &blocks = var.second;
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003203 auto &type = this->expression_type(var.first);
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003204
3205 // Figure out which block is dominating all accesses of those variables.
3206 for (auto &block : blocks)
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003207 {
Hans-Kristian Arntzena714d422016-12-16 12:43:12 +01003208 // If we're accessing a variable inside a continue block, this variable might be a loop variable.
3209 // We can only use loop variables with scalars, as we cannot track static expressions for vectors.
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003210 if (this->is_continue(block) && type.vecsize == 1 && type.columns == 1)
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003211 {
3212 // The variable is used in multiple continue blocks, this is not a loop
3213 // candidate, signal that by setting block to -1u.
3214 auto &potential = potential_loop_variables[var.first];
3215
3216 if (potential == 0)
3217 potential = block;
3218 else
Graham Wihlidalfadc1f92017-01-05 20:14:53 +01003219 potential = ~(0u);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003220 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003221 builder.add_block(block);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003222 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003223
Hans-Kristian Arntzen5ff11cc2016-11-18 16:45:11 +01003224 builder.lift_continue_block_dominator();
3225
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003226 // Add it to a per-block list of variables.
3227 uint32_t dominating_block = builder.get_dominator();
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003228 // If all blocks here are dead code, this will be 0, so the variable in question
3229 // will be completely eliminated.
3230 if (dominating_block)
3231 {
Hans-Kristian Arntzen7630d3c2016-11-21 12:14:02 +01003232 auto &block = this->get<SPIRBlock>(dominating_block);
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003233 block.dominated_variables.push_back(var.first);
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003234 this->get<SPIRVariable>(var.first).dominator = dominating_block;
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003235 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003236 }
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003237
Hans-Kristian Arntzenf0c50a62017-07-25 18:22:15 +02003238 unordered_set<uint32_t> seen_blocks;
3239
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003240 // Now, try to analyze whether or not these variables are actually loop variables.
3241 for (auto &loop_variable : potential_loop_variables)
3242 {
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003243 auto &var = this->get<SPIRVariable>(loop_variable.first);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003244 auto dominator = var.dominator;
3245 auto block = loop_variable.second;
3246
3247 // The variable was accessed in multiple continue blocks, ignore.
Graham Wihlidalfadc1f92017-01-05 20:14:53 +01003248 if (block == ~(0u) || block == 0)
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003249 continue;
3250
3251 // Dead code.
3252 if (dominator == 0)
3253 continue;
3254
3255 uint32_t header = 0;
3256
3257 // Find the loop header for this block.
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003258 for (auto b : this->loop_blocks)
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003259 {
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003260 auto &potential_header = this->get<SPIRBlock>(b);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003261 if (potential_header.continue_block == block)
3262 {
3263 header = b;
3264 break;
3265 }
3266 }
3267
3268 assert(header);
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003269 auto &header_block = this->get<SPIRBlock>(header);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003270
3271 // Now, there are two conditions we need to meet for the variable to be a loop variable.
3272 // 1. The dominating block must have a branch-free path to the loop header,
3273 // this way we statically know which expression should be part of the loop variable initializer.
3274
3275 // Walk from the dominator, if there is one straight edge connecting
3276 // dominator and loop header, we statically know the loop initializer.
3277 bool static_loop_init = true;
3278 while (dominator != header)
3279 {
3280 auto &succ = cfg.get_succeeding_edges(dominator);
3281 if (succ.size() != 1)
3282 {
3283 static_loop_init = false;
3284 break;
3285 }
3286
3287 auto &pred = cfg.get_preceding_edges(succ.front());
3288 if (pred.size() != 1 || pred.front() != dominator)
3289 {
3290 static_loop_init = false;
3291 break;
3292 }
3293
3294 dominator = succ.front();
3295 }
3296
3297 if (!static_loop_init)
3298 continue;
3299
3300 // The second condition we need to meet is that no access after the loop
3301 // merge can occur. Walk the CFG to see if we find anything.
3302 auto &blocks = handler.accessed_variables_to_block[loop_variable.first];
Hans-Kristian Arntzenf0c50a62017-07-25 18:22:15 +02003303
3304 seen_blocks.clear();
3305 cfg.walk_from(seen_blocks, header_block.merge_block, [&](uint32_t walk_block) {
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003306 // We found a block which accesses the variable outside the loop.
3307 if (blocks.find(walk_block) != end(blocks))
3308 static_loop_init = false;
3309 });
3310
3311 if (!static_loop_init)
3312 continue;
3313
3314 // We have a loop variable.
3315 header_block.loop_variables.push_back(loop_variable.first);
Hans-Kristian Arntzen44b32162016-12-16 14:01:09 +01003316 // Need to sort here as variables come from an unordered container, and pushing stuff in wrong order
3317 // will break reproducability in regression runs.
3318 sort(begin(header_block.loop_variables), end(header_block.loop_variables));
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003319 this->get<SPIRVariable>(loop_variable.first).loop_variable = true;
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003320 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003321}
Hans-Kristian Arntzen016b1d82017-01-21 10:07:38 +01003322
3323uint64_t Compiler::get_buffer_block_flags(const SPIRVariable &var)
3324{
3325 auto &type = get<SPIRType>(var.basetype);
Hans-Kristian Arntzen8a80e622017-01-22 08:51:24 +01003326 assert(type.basetype == SPIRType::Struct);
Hans-Kristian Arntzen016b1d82017-01-21 10:07:38 +01003327
3328 // Some flags like non-writable, non-readable are actually found
3329 // as member decorations. If all members have a decoration set, propagate
3330 // the decoration up as a regular variable decoration.
3331 uint64_t base_flags = meta[var.self].decoration.decoration_flags;
Hans-Kristian Arntzen016b1d82017-01-21 10:07:38 +01003332
Hans-Kristian Arntzen8a80e622017-01-22 08:51:24 +01003333 if (type.member_types.empty())
3334 return base_flags;
3335
3336 uint64_t all_members_flag_mask = ~(0ull);
3337 for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++)
3338 all_members_flag_mask &= get_member_decoration_mask(type.self, i);
3339
3340 return base_flags | all_members_flag_mask;
Hans-Kristian Arntzen016b1d82017-01-21 10:07:38 +01003341}
Hans-Kristian Arntzen95409792017-01-21 12:29:20 +01003342
3343bool Compiler::get_common_basic_type(const SPIRType &type, SPIRType::BaseType &base_type)
3344{
3345 if (type.basetype == SPIRType::Struct)
3346 {
3347 base_type = SPIRType::Unknown;
3348 for (auto &member_type : type.member_types)
3349 {
3350 SPIRType::BaseType member_base;
3351 if (!get_common_basic_type(get<SPIRType>(member_type), member_base))
3352 return false;
3353
3354 if (base_type == SPIRType::Unknown)
3355 base_type = member_base;
3356 else if (base_type != member_base)
3357 return false;
3358 }
3359 return true;
3360 }
3361 else
3362 {
3363 base_type = type.basetype;
3364 return true;
3365 }
3366}
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003367
3368bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args, uint32_t length)
3369{
3370 const auto add_if_builtin = [&](uint32_t id) {
3371 // Only handles variables here.
3372 // Builtins which are part of a block are handled in AccessChain.
3373 auto *var = compiler.maybe_get<SPIRVariable>(id);
3374 if (var && compiler.meta[id].decoration.builtin)
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003375 {
3376 auto &type = compiler.get<SPIRType>(var->basetype);
3377 auto &flags =
Hans-Kristian Arntzen2ebe1a82017-03-06 16:50:46 +01003378 type.storage == StorageClassInput ? compiler.active_input_builtins : compiler.active_output_builtins;
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003379 flags |= 1ull << compiler.meta[id].decoration.builtin_type;
3380 }
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003381 };
3382
3383 switch (opcode)
3384 {
3385 case OpStore:
3386 if (length < 1)
3387 return false;
3388
3389 add_if_builtin(args[0]);
3390 break;
3391
3392 case OpCopyMemory:
3393 if (length < 2)
3394 return false;
3395
3396 add_if_builtin(args[0]);
3397 add_if_builtin(args[1]);
3398 break;
3399
3400 case OpCopyObject:
3401 case OpLoad:
3402 if (length < 3)
3403 return false;
3404
3405 add_if_builtin(args[2]);
3406 break;
3407
3408 case OpFunctionCall:
3409 {
3410 if (length < 3)
3411 return false;
3412
3413 uint32_t count = length - 3;
3414 args += 3;
3415 for (uint32_t i = 0; i < count; i++)
3416 add_if_builtin(args[i]);
3417 break;
3418 }
3419
3420 case OpAccessChain:
3421 case OpInBoundsAccessChain:
3422 {
3423 if (length < 4)
3424 return false;
3425
3426 // Only consider global variables, cannot consider variables in functions yet, or other
3427 // access chains as they have not been created yet.
3428 auto *var = compiler.maybe_get<SPIRVariable>(args[2]);
3429 if (!var)
3430 break;
3431
Hans-Kristian Arntzen945425e2017-08-15 10:23:04 +02003432 // Required if we access chain into builtins like gl_GlobalInvocationID.
3433 add_if_builtin(args[2]);
3434
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003435 auto *type = &compiler.get<SPIRType>(var->basetype);
3436
3437 // Start traversing type hierarchy at the proper non-pointer types.
3438 while (type->pointer)
3439 {
3440 assert(type->parent_type);
3441 type = &compiler.get<SPIRType>(type->parent_type);
3442 }
3443
Hans-Kristian Arntzen2ebe1a82017-03-06 16:50:46 +01003444 auto &flags =
3445 type->storage == StorageClassInput ? compiler.active_input_builtins : compiler.active_output_builtins;
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003446
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003447 uint32_t count = length - 3;
3448 args += 3;
3449 for (uint32_t i = 0; i < count; i++)
3450 {
3451 // Arrays
3452 if (!type->array.empty())
3453 {
3454 type = &compiler.get<SPIRType>(type->parent_type);
3455 }
3456 // Structs
3457 else if (type->basetype == SPIRType::Struct)
3458 {
3459 uint32_t index = compiler.get<SPIRConstant>(args[i]).scalar();
3460
3461 if (index < uint32_t(compiler.meta[type->self].members.size()))
3462 {
3463 auto &decorations = compiler.meta[type->self].members[index];
3464 if (decorations.builtin)
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003465 flags |= 1ull << decorations.builtin_type;
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003466 }
3467
3468 type = &compiler.get<SPIRType>(type->member_types[index]);
3469 }
3470 else
3471 {
3472 // No point in traversing further. We won't find any extra builtins.
3473 break;
3474 }
3475 }
3476 break;
3477 }
3478
3479 default:
3480 break;
3481 }
3482
3483 return true;
3484}
3485
3486void Compiler::update_active_builtins()
3487{
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003488 active_input_builtins = 0;
3489 active_output_builtins = 0;
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003490 ActiveBuiltinHandler handler(*this);
3491 traverse_all_reachable_opcodes(get<SPIRFunction>(entry_point), handler);
3492}
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003493
Bill Hollings192bdc92017-05-24 09:31:38 -04003494// Returns whether this shader uses a builtin of the storage class
3495bool Compiler::has_active_builtin(BuiltIn builtin, StorageClass storage)
3496{
3497 uint64_t flags;
3498 switch (storage)
3499 {
3500 case StorageClassInput:
3501 flags = active_input_builtins;
3502 break;
3503 case StorageClassOutput:
3504 flags = active_output_builtins;
3505 break;
3506
3507 default:
3508 return false;
3509 }
3510 return flags & (1ull << builtin);
3511}
3512
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003513void Compiler::analyze_sampler_comparison_states()
3514{
3515 CombinedImageSamplerUsageHandler handler(*this);
3516 traverse_all_reachable_opcodes(get<SPIRFunction>(entry_point), handler);
3517 comparison_samplers = move(handler.comparison_samplers);
3518}
3519
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003520bool Compiler::CombinedImageSamplerUsageHandler::begin_function_scope(const uint32_t *args, uint32_t length)
3521{
3522 if (length < 3)
3523 return false;
3524
3525 auto &func = compiler.get<SPIRFunction>(args[2]);
3526 const auto *arg = &args[3];
3527 length -= 3;
3528
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003529 for (uint32_t i = 0; i < length; i++)
3530 {
3531 auto &argument = func.arguments[i];
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003532 dependency_hierarchy[argument.id].insert(arg[i]);
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003533 }
3534
3535 return true;
3536}
3537
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003538void Compiler::CombinedImageSamplerUsageHandler::add_hierarchy_to_comparison_samplers(uint32_t sampler)
3539{
3540 // Traverse the variable dependency hierarchy and tag everything in its path with comparison samplers.
3541 comparison_samplers.insert(sampler);
3542 for (auto &samp : dependency_hierarchy[sampler])
3543 add_hierarchy_to_comparison_samplers(samp);
3544}
3545
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003546bool Compiler::CombinedImageSamplerUsageHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
3547{
3548 switch (opcode)
3549 {
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003550 case OpAccessChain:
3551 case OpInBoundsAccessChain:
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003552 case OpLoad:
3553 {
3554 if (length < 3)
3555 return false;
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003556 dependency_hierarchy[args[1]].insert(args[2]);
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003557 break;
3558 }
3559
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003560 case OpSampledImage:
3561 {
3562 if (length < 4)
3563 return false;
3564
3565 uint32_t result_type = args[0];
3566 auto &type = compiler.get<SPIRType>(result_type);
3567 if (type.image.depth)
3568 {
3569 // This sampler must be a SamplerComparisionState, and not a regular SamplerState.
3570 uint32_t sampler = args[3];
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003571 add_hierarchy_to_comparison_samplers(sampler);
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003572 }
3573 return true;
3574 }
3575
3576 default:
3577 break;
3578 }
3579
3580 return true;
3581}
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02003582
3583bool Compiler::buffer_is_hlsl_counter_buffer(uint32_t id) const
3584{
3585 if (meta.at(id).hlsl_magic_counter_buffer_candidate)
3586 {
3587 auto *var = maybe_get<SPIRVariable>(id);
Hans-Kristian Arntzen9c9e2672017-05-09 09:38:33 +02003588 // Ensure that this is actually a buffer object.
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +02003589 return var && (var->storage == StorageClassStorageBuffer ||
3590 has_decoration(get<SPIRType>(var->basetype).self, DecorationBufferBlock));
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02003591 }
3592 else
3593 return false;
3594}
3595
3596bool Compiler::buffer_get_hlsl_counter_buffer(uint32_t id, uint32_t &counter_id) const
3597{
3598 auto &name = get_name(id);
3599 uint32_t id_bound = get_current_id_bound();
3600 for (uint32_t i = 0; i < id_bound; i++)
3601 {
3602 if (meta[i].hlsl_magic_counter_buffer_candidate && meta[i].hlsl_magic_counter_buffer_name == name)
3603 {
3604 auto *var = maybe_get<SPIRVariable>(i);
Hans-Kristian Arntzen9c9e2672017-05-09 09:38:33 +02003605 // Ensure that this is actually a buffer object.
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +02003606 if (var && (var->storage == StorageClassStorageBuffer ||
3607 has_decoration(get<SPIRType>(var->basetype).self, DecorationBufferBlock)))
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02003608 {
3609 counter_id = i;
3610 return true;
3611 }
3612 }
3613 }
3614 return false;
3615}
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003616
3617void Compiler::make_constant_null(uint32_t id, uint32_t type)
3618{
3619 auto &constant_type = get<SPIRType>(type);
3620
3621 if (!constant_type.array.empty())
3622 {
3623 assert(constant_type.parent_type);
3624 uint32_t parent_id = increase_bound_by(1);
3625 make_constant_null(parent_id, constant_type.parent_type);
3626
3627 if (!constant_type.array_size_literal.back())
3628 SPIRV_CROSS_THROW("Array size of OpConstantNull must be a literal.");
3629
3630 vector<uint32_t> elements(constant_type.array.back());
3631 for (uint32_t i = 0; i < constant_type.array.back(); i++)
3632 elements[i] = parent_id;
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02003633 set<SPIRConstant>(id, type, elements.data(), uint32_t(elements.size()), false);
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003634 }
3635 else if (!constant_type.member_types.empty())
3636 {
Bill Hollings77f58122017-08-11 14:54:58 -04003637 uint32_t member_ids = increase_bound_by(uint32_t(constant_type.member_types.size()));
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003638 vector<uint32_t> elements(constant_type.member_types.size());
3639 for (uint32_t i = 0; i < constant_type.member_types.size(); i++)
3640 {
3641 make_constant_null(member_ids + i, constant_type.member_types[i]);
3642 elements[i] = member_ids + i;
3643 }
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02003644 set<SPIRConstant>(id, type, elements.data(), uint32_t(elements.size()), false);
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003645 }
3646 else
3647 {
3648 auto &constant = set<SPIRConstant>(id, type);
3649 constant.make_null(constant_type);
3650 }
3651}
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02003652
3653const std::vector<spv::Capability> &Compiler::get_declared_capabilities() const
3654{
3655 return declared_capabilities;
3656}
3657
3658const std::vector<std::string> &Compiler::get_declared_extensions() const
3659{
3660 return declared_extensions;
3661}