blob: 4dfd1b232d04f3a9abb776c1ad78cdc8cd1020c4 [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 Arntzence18d4c2017-11-17 13:38:29 +0100171 // 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 Arntzence18d4c2017-11-17 13:38:29 +0100460 }) != 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 Arntzenbcdff2d2017-11-22 19:27:03 +0100576 case OpExtInst:
577 {
578 if (length < 5)
579 return false;
580 uint32_t extension_set = args[2];
581 if (compiler.get<SPIRExtension>(extension_set).ext == SPIRExtension::SPV_AMD_shader_explicit_vertex_parameter)
582 {
583 enum AMDShaderExplicitVertexParameter
584 {
585 InterpolateAtVertexAMD = 1
586 };
587
588 auto op = static_cast<AMDShaderExplicitVertexParameter>(args[3]);
589
590 switch (op)
591 {
592 case InterpolateAtVertexAMD:
593 {
594 auto *var = compiler.maybe_get<SPIRVariable>(args[4]);
595 if (var && storage_class_is_interface(var->storage))
596 variables.insert(args[4]);
597 break;
598 }
599
600 default:
601 break;
602 }
603 }
604 break;
605 }
606
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200607 case OpAccessChain:
608 case OpInBoundsAccessChain:
609 case OpLoad:
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +0100610 case OpCopyObject:
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200611 case OpImageTexelPointer:
612 case OpAtomicLoad:
613 case OpAtomicExchange:
614 case OpAtomicCompareExchange:
Bill Hollings8f6df772017-05-19 18:14:08 -0400615 case OpAtomicCompareExchangeWeak:
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200616 case OpAtomicIIncrement:
617 case OpAtomicIDecrement:
618 case OpAtomicIAdd:
619 case OpAtomicISub:
620 case OpAtomicSMin:
621 case OpAtomicUMin:
622 case OpAtomicSMax:
623 case OpAtomicUMax:
624 case OpAtomicAnd:
625 case OpAtomicOr:
626 case OpAtomicXor:
627 // Invalid SPIR-V.
628 if (length < 3)
629 return false;
630 variable = args[2];
631 break;
632 }
633
634 if (variable)
635 {
636 auto *var = compiler.maybe_get<SPIRVariable>(variable);
637 if (var && storage_class_is_interface(var->storage))
638 variables.insert(variable);
639 }
640 return true;
641}
642
643unordered_set<uint32_t> Compiler::get_active_interface_variables() const
644{
645 // Traverse the call graph and find all interface variables which are in use.
646 unordered_set<uint32_t> variables;
647 InterfaceVariableAccessHandler handler(*this, variables);
648 traverse_all_reachable_opcodes(get<SPIRFunction>(entry_point), handler);
649 return variables;
650}
651
652void Compiler::set_enabled_interface_variables(std::unordered_set<uint32_t> active_variables)
653{
654 active_interface_variables = move(active_variables);
655 check_active_interface_variables = true;
656}
657
658ShaderResources Compiler::get_shader_resources(const unordered_set<uint32_t> *active_variables) const
659{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200660 ShaderResources res;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100661
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200662 for (auto &id : ids)
663 {
664 if (id.get_type() != TypeVariable)
665 continue;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100666
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200667 auto &var = id.get<SPIRVariable>();
668 auto &type = get<SPIRType>(var.basetype);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100669
Hans-Kristian Arntzend5dc5f32016-07-05 13:21:26 +0200670 // It is possible for uniform storage classes to be passed as function parameters, so detect
671 // that. To detect function parameters, check of StorageClass of variable is function scope.
672 if (var.storage == StorageClassFunction || !type.pointer || is_builtin_variable(var))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200673 continue;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100674
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200675 if (active_variables && active_variables->find(var.self) == end(*active_variables))
676 continue;
677
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200678 // Input
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +0200679 if (var.storage == StorageClassInput && interface_variable_exists_in_entry_point(var.self))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200680 {
681 if (meta[type.self].decoration.decoration_flags & (1ull << DecorationBlock))
Hans-Kristian Arntzenaa2557c2017-12-05 09:58:12 +0100682 res.stage_inputs.push_back(
683 { var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200684 else
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200685 res.stage_inputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200686 }
687 // Subpass inputs
688 else if (var.storage == StorageClassUniformConstant && type.image.dim == DimSubpassData)
689 {
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200690 res.subpass_inputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200691 }
692 // Outputs
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +0200693 else if (var.storage == StorageClassOutput && interface_variable_exists_in_entry_point(var.self))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200694 {
695 if (meta[type.self].decoration.decoration_flags & (1ull << DecorationBlock))
Hans-Kristian Arntzenaa2557c2017-12-05 09:58:12 +0100696 res.stage_outputs.push_back(
697 { var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200698 else
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200699 res.stage_outputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200700 }
701 // UBOs
702 else if (type.storage == StorageClassUniform &&
703 (meta[type.self].decoration.decoration_flags & (1ull << DecorationBlock)))
704 {
Hans-Kristian Arntzenaa2557c2017-12-05 09:58:12 +0100705 res.uniform_buffers.push_back(
706 { var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200707 }
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +0200708 // Old way to declare SSBOs.
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200709 else if (type.storage == StorageClassUniform &&
710 (meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)))
711 {
Hans-Kristian Arntzenaa2557c2017-12-05 09:58:12 +0100712 res.storage_buffers.push_back(
713 { var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200714 }
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +0200715 // Modern way to declare SSBOs.
716 else if (type.storage == StorageClassStorageBuffer)
717 {
Hans-Kristian Arntzenaa2557c2017-12-05 09:58:12 +0100718 res.storage_buffers.push_back(
719 { var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) });
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +0200720 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200721 // Push constant blocks
722 else if (type.storage == StorageClassPushConstant)
723 {
724 // There can only be one push constant block, but keep the vector in case this restriction is lifted
725 // in the future.
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200726 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 +0200727 }
728 // Images
Hans-Kristian Arntzene9202082016-09-10 13:05:35 +0200729 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Image &&
730 type.image.sampled == 2)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200731 {
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200732 res.storage_images.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200733 }
Hans-Kristian Arntzene9202082016-09-10 13:05:35 +0200734 // Separate images
735 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Image &&
736 type.image.sampled == 1)
737 {
738 res.separate_images.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
739 }
740 // Separate samplers
741 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Sampler)
742 {
743 res.separate_samplers.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
744 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200745 // Textures
746 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::SampledImage)
747 {
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200748 res.sampled_images.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200749 }
750 // Atomic counters
751 else if (type.storage == StorageClassAtomicCounter)
752 {
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200753 res.atomic_counters.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200754 }
755 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100756
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200757 return res;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100758}
759
760static inline uint32_t swap_endian(uint32_t v)
761{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200762 return ((v >> 24) & 0x000000ffu) | ((v >> 8) & 0x0000ff00u) | ((v << 8) & 0x00ff0000u) | ((v << 24) & 0xff000000u);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100763}
764
765static string extract_string(const vector<uint32_t> &spirv, uint32_t offset)
766{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200767 string ret;
768 for (uint32_t i = offset; i < spirv.size(); i++)
769 {
770 uint32_t w = spirv[i];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100771
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200772 for (uint32_t j = 0; j < 4; j++, w >>= 8)
773 {
774 char c = w & 0xff;
775 if (c == '\0')
776 return ret;
777 ret += c;
778 }
779 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100780
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100781 SPIRV_CROSS_THROW("String was not terminated before EOF");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100782}
783
Hans-Kristian Arntzen45ad58a2016-05-10 23:39:41 +0200784static bool is_valid_spirv_version(uint32_t version)
785{
786 switch (version)
787 {
788 // Allow v99 since it tends to just work.
789 case 99:
790 case 0x10000: // SPIR-V 1.0
791 case 0x10100: // SPIR-V 1.1
Hans-Kristian Arntzene5595272017-05-22 13:59:58 +0200792 case 0x10200: // SPIR-V 1.2
Hans-Kristian Arntzen45ad58a2016-05-10 23:39:41 +0200793 return true;
794
795 default:
796 return false;
797 }
798}
799
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100800void Compiler::parse()
801{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200802 auto len = spirv.size();
803 if (len < 5)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100804 SPIRV_CROSS_THROW("SPIRV file too small.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100805
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200806 auto s = spirv.data();
Hans-Kristian Arntzen5ac88272016-04-11 13:38:18 +0200807
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200808 // Endian-swap if we need to.
809 if (s[0] == swap_endian(MagicNumber))
Hans-Kristian Arntzen5ea59bd2016-05-23 13:30:02 +0200810 transform(begin(spirv), end(spirv), begin(spirv), [](uint32_t c) { return swap_endian(c); });
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100811
Hans-Kristian Arntzen45ad58a2016-05-10 23:39:41 +0200812 if (s[0] != MagicNumber || !is_valid_spirv_version(s[1]))
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100813 SPIRV_CROSS_THROW("Invalid SPIRV format.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100814
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200815 uint32_t bound = s[3];
816 ids.resize(bound);
817 meta.resize(bound);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100818
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200819 uint32_t offset = 5;
820 while (offset < len)
821 inst.emplace_back(spirv, offset);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100822
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200823 for (auto &i : inst)
824 parse(i);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100825
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200826 if (current_function)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100827 SPIRV_CROSS_THROW("Function was not terminated.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200828 if (current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100829 SPIRV_CROSS_THROW("Block was not terminated.");
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +0200830
831 // Figure out specialization constants for work group sizes.
832 for (auto &id : ids)
833 {
834 if (id.get_type() == TypeConstant)
835 {
836 auto &c = id.get<SPIRConstant>();
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +0200837 if (meta[c.self].decoration.builtin && meta[c.self].decoration.builtin_type == BuiltInWorkgroupSize)
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +0200838 {
839 // In current SPIR-V, there can be just one constant like this.
840 // All entry points will receive the constant value.
841 for (auto &entry : entry_points)
842 {
843 entry.second.workgroup_size.constant = c.self;
844 entry.second.workgroup_size.x = c.scalar(0, 0);
845 entry.second.workgroup_size.y = c.scalar(0, 1);
846 entry.second.workgroup_size.z = c.scalar(0, 2);
847 }
848 }
849 }
850 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100851}
852
853void Compiler::flatten_interface_block(uint32_t id)
854{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200855 auto &var = get<SPIRVariable>(id);
856 auto &type = get<SPIRType>(var.basetype);
857 auto flags = meta.at(type.self).decoration.decoration_flags;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100858
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200859 if (!type.array.empty())
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100860 SPIRV_CROSS_THROW("Type is array of UBOs.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200861 if (type.basetype != SPIRType::Struct)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100862 SPIRV_CROSS_THROW("Type is not a struct.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200863 if ((flags & (1ull << DecorationBlock)) == 0)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100864 SPIRV_CROSS_THROW("Type is not a block.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200865 if (type.member_types.empty())
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100866 SPIRV_CROSS_THROW("Member list of struct is empty.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100867
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200868 uint32_t t = type.member_types[0];
869 for (auto &m : type.member_types)
870 if (t != m)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100871 SPIRV_CROSS_THROW("Types in block differ.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100872
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200873 auto &mtype = get<SPIRType>(t);
874 if (!mtype.array.empty())
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100875 SPIRV_CROSS_THROW("Member type cannot be arrays.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200876 if (mtype.basetype == SPIRType::Struct)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +0100877 SPIRV_CROSS_THROW("Member type cannot be struct.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100878
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200879 // Inherit variable name from interface block name.
880 meta.at(var.self).decoration.alias = meta.at(type.self).decoration.alias;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100881
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200882 auto storage = var.storage;
883 if (storage == StorageClassUniform)
884 storage = StorageClassUniformConstant;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100885
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200886 // Change type definition in-place into an array instead.
887 // Access chains will still work as-is.
888 uint32_t array_size = uint32_t(type.member_types.size());
889 type = mtype;
890 type.array.push_back(array_size);
891 type.pointer = true;
892 type.storage = storage;
893 var.storage = storage;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100894}
895
896void Compiler::update_name_cache(unordered_set<string> &cache, string &name)
897{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200898 if (name.empty())
899 return;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100900
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200901 if (cache.find(name) == end(cache))
902 {
903 cache.insert(name);
904 return;
905 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100906
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200907 uint32_t counter = 0;
908 auto tmpname = name;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100909
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200910 // If there is a collision (very rare),
911 // keep tacking on extra identifier until it's unique.
912 do
913 {
914 counter++;
915 name = tmpname + "_" + convert_to_string(counter);
916 } while (cache.find(name) != end(cache));
917 cache.insert(name);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100918}
919
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200920void Compiler::set_name(uint32_t id, const std::string &name)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100921{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200922 auto &str = meta.at(id).decoration.alias;
923 str.clear();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100924
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200925 if (name.empty())
926 return;
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +0200927
928 // glslang uses identifiers to pass along meaningful information
929 // about HLSL reflection.
930 auto &m = meta.at(id);
Hans-Kristian Arntzen08b3c672017-05-09 09:30:30 +0200931 if (source.hlsl && name.size() >= 6 && name.find("@count") == name.size() - 6)
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +0200932 {
933 m.hlsl_magic_counter_buffer_candidate = true;
934 m.hlsl_magic_counter_buffer_name = name.substr(0, name.find("@count"));
935 }
936 else
937 {
938 m.hlsl_magic_counter_buffer_candidate = false;
939 m.hlsl_magic_counter_buffer_name.clear();
940 }
941
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200942 // Reserved for temporaries.
Hans-Kristian Arntzenc9728942016-07-06 11:19:20 +0200943 if (name[0] == '_' && name.size() >= 2 && isdigit(name[1]))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200944 return;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100945
Hans-Kristian Arntzen94ff3552017-10-10 17:32:26 +0200946 str = ensure_valid_identifier(name, false);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100947}
948
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200949const SPIRType &Compiler::get_type(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100950{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200951 return get<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100952}
953
Robert Konrad451bdee2016-09-24 22:17:01 +0200954const SPIRType &Compiler::get_type_from_variable(uint32_t id) const
955{
956 return get<SPIRType>(get<SPIRVariable>(id).basetype);
957}
958
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100959void Compiler::set_member_decoration(uint32_t id, uint32_t index, Decoration decoration, uint32_t argument)
960{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200961 meta.at(id).members.resize(max(meta[id].members.size(), size_t(index) + 1));
962 auto &dec = meta.at(id).members[index];
963 dec.decoration_flags |= 1ull << decoration;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100964
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200965 switch (decoration)
966 {
967 case DecorationBuiltIn:
968 dec.builtin = true;
969 dec.builtin_type = static_cast<BuiltIn>(argument);
970 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100971
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200972 case DecorationLocation:
973 dec.location = argument;
974 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100975
Bill Hollings81757502017-01-29 13:28:20 -0500976 case DecorationBinding:
977 dec.binding = argument;
978 break;
979
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200980 case DecorationOffset:
981 dec.offset = argument;
982 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100983
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +0200984 case DecorationSpecId:
985 dec.spec_id = argument;
986 break;
987
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +0100988 case DecorationMatrixStride:
989 dec.matrix_stride = argument;
990 break;
991
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200992 default:
993 break;
994 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100995}
996
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200997void Compiler::set_member_name(uint32_t id, uint32_t index, const std::string &name)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100998{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200999 meta.at(id).members.resize(max(meta[id].members.size(), size_t(index) + 1));
Hans-Kristian Arntzenfd432f82017-03-18 10:52:41 +01001000
1001 auto &str = meta.at(id).members[index].alias;
1002 str.clear();
1003 if (name.empty())
1004 return;
1005
1006 // Reserved for unnamed members.
Hans-Kristian Arntzen94ff3552017-10-10 17:32:26 +02001007 if (name[0] == '_' && name.size() >= 3 && name[1] == 'm' && isdigit(name[2]))
Hans-Kristian Arntzenfd432f82017-03-18 10:52:41 +01001008 return;
1009
Hans-Kristian Arntzen94ff3552017-10-10 17:32:26 +02001010 str = ensure_valid_identifier(name, true);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001011}
1012
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001013const std::string &Compiler::get_member_name(uint32_t id, uint32_t index) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001014{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001015 auto &m = meta.at(id);
1016 if (index >= m.members.size())
1017 {
1018 static string empty;
1019 return empty;
1020 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001021
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001022 return m.members[index].alias;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001023}
1024
Bill Hollings1e84a372017-08-12 00:21:13 -04001025void Compiler::set_member_qualified_name(uint32_t type_id, uint32_t index, const std::string &name)
Bill Hollingsc45e74f2016-07-08 12:39:22 -04001026{
Bill Hollingsc1b81542017-05-22 21:41:19 -04001027 meta.at(type_id).members.resize(max(meta[type_id].members.size(), size_t(index) + 1));
1028 meta.at(type_id).members[index].qualified_alias = name;
1029}
1030
Bill Hollings1e84a372017-08-12 00:21:13 -04001031const std::string &Compiler::get_member_qualified_name(uint32_t type_id, uint32_t index) const
Bill Hollingsc1b81542017-05-22 21:41:19 -04001032{
Bill Hollings1e84a372017-08-12 00:21:13 -04001033 const static string empty;
Bill Hollingsc1b81542017-05-22 21:41:19 -04001034
1035 auto &m = meta.at(type_id);
Bill Hollings1e84a372017-08-12 00:21:13 -04001036 if (index < m.members.size())
1037 return m.members[index].qualified_alias;
1038 else
Bill Hollingsc1b81542017-05-22 21:41:19 -04001039 return empty;
Bill Hollingsc45e74f2016-07-08 12:39:22 -04001040}
1041
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001042uint32_t Compiler::get_member_decoration(uint32_t id, uint32_t index, Decoration decoration) const
1043{
Hans-Kristian Arntzen15679c72016-08-17 11:35:34 +02001044 auto &m = meta.at(id);
1045 if (index >= m.members.size())
1046 return 0;
1047
1048 auto &dec = m.members[index];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001049 if (!(dec.decoration_flags & (1ull << decoration)))
1050 return 0;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001051
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001052 switch (decoration)
1053 {
1054 case DecorationBuiltIn:
1055 return dec.builtin_type;
1056 case DecorationLocation:
1057 return dec.location;
Bill Hollings81757502017-01-29 13:28:20 -05001058 case DecorationBinding:
1059 return dec.binding;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001060 case DecorationOffset:
1061 return dec.offset;
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02001062 case DecorationSpecId:
1063 return dec.spec_id;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001064 default:
Graham Wihlidal9b1ee8f2017-01-05 21:01:49 +01001065 return 1;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001066 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001067}
1068
1069uint64_t Compiler::get_member_decoration_mask(uint32_t id, uint32_t index) const
1070{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001071 auto &m = meta.at(id);
1072 if (index >= m.members.size())
1073 return 0;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001074
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001075 return m.members[index].decoration_flags;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001076}
1077
Bill Hollingsdc694272017-03-11 12:17:22 -05001078bool Compiler::has_member_decoration(uint32_t id, uint32_t index, Decoration decoration) const
Bill Hollings484931d2017-02-28 21:44:36 -05001079{
Bill Hollingsdc694272017-03-11 12:17:22 -05001080 return get_member_decoration_mask(id, index) & (1ull << decoration);
Bill Hollings484931d2017-02-28 21:44:36 -05001081}
1082
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001083void Compiler::unset_member_decoration(uint32_t id, uint32_t index, Decoration decoration)
1084{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001085 auto &m = meta.at(id);
1086 if (index >= m.members.size())
1087 return;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001088
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001089 auto &dec = m.members[index];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001090
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001091 dec.decoration_flags &= ~(1ull << decoration);
1092 switch (decoration)
1093 {
1094 case DecorationBuiltIn:
1095 dec.builtin = false;
1096 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001097
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001098 case DecorationLocation:
1099 dec.location = 0;
1100 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001101
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001102 case DecorationOffset:
1103 dec.offset = 0;
1104 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001105
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02001106 case DecorationSpecId:
1107 dec.spec_id = 0;
1108 break;
1109
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001110 default:
1111 break;
1112 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001113}
1114
1115void Compiler::set_decoration(uint32_t id, Decoration decoration, uint32_t argument)
1116{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001117 auto &dec = meta.at(id).decoration;
1118 dec.decoration_flags |= 1ull << decoration;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001119
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001120 switch (decoration)
1121 {
1122 case DecorationBuiltIn:
1123 dec.builtin = true;
1124 dec.builtin_type = static_cast<BuiltIn>(argument);
1125 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001126
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001127 case DecorationLocation:
1128 dec.location = argument;
1129 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001130
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001131 case DecorationOffset:
1132 dec.offset = argument;
1133 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001134
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001135 case DecorationArrayStride:
1136 dec.array_stride = argument;
1137 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001138
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001139 case DecorationMatrixStride:
1140 dec.matrix_stride = argument;
1141 break;
1142
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001143 case DecorationBinding:
1144 dec.binding = argument;
1145 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001146
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001147 case DecorationDescriptorSet:
1148 dec.set = argument;
1149 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001150
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001151 case DecorationInputAttachmentIndex:
1152 dec.input_attachment = argument;
1153 break;
Hans-Kristian Arntzen12cfbb22016-05-04 13:41:04 +02001154
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02001155 case DecorationSpecId:
1156 dec.spec_id = argument;
1157 break;
1158
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001159 default:
1160 break;
1161 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001162}
1163
1164StorageClass Compiler::get_storage_class(uint32_t id) const
1165{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001166 return get<SPIRVariable>(id).storage;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001167}
1168
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001169const std::string &Compiler::get_name(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001170{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001171 return meta.at(id).decoration.alias;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001172}
1173
Hans-Kristian Arntzenaab31072017-09-29 12:16:53 +02001174const std::string Compiler::get_fallback_name(uint32_t id) const
1175{
1176 return join("_", id);
1177}
1178
1179const std::string Compiler::get_block_fallback_name(uint32_t id) const
1180{
1181 auto &var = get<SPIRVariable>(id);
Hans-Kristian Arntzen2c90ea32017-12-01 14:20:51 +01001182 if (get_name(id).empty())
1183 return join("_", get<SPIRType>(var.basetype).self, "_", id);
1184 else
1185 return get_name(id);
Hans-Kristian Arntzenaab31072017-09-29 12:16:53 +02001186}
1187
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001188uint64_t Compiler::get_decoration_mask(uint32_t id) const
1189{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001190 auto &dec = meta.at(id).decoration;
1191 return dec.decoration_flags;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001192}
1193
Bill Hollingsdc694272017-03-11 12:17:22 -05001194bool Compiler::has_decoration(uint32_t id, Decoration decoration) const
Bill Hollings484931d2017-02-28 21:44:36 -05001195{
Bill Hollingsdc694272017-03-11 12:17:22 -05001196 return get_decoration_mask(id) & (1ull << decoration);
Bill Hollings484931d2017-02-28 21:44:36 -05001197}
1198
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001199uint32_t Compiler::get_decoration(uint32_t id, Decoration decoration) const
1200{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001201 auto &dec = meta.at(id).decoration;
1202 if (!(dec.decoration_flags & (1ull << decoration)))
1203 return 0;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001204
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001205 switch (decoration)
1206 {
1207 case DecorationBuiltIn:
1208 return dec.builtin_type;
1209 case DecorationLocation:
1210 return dec.location;
1211 case DecorationOffset:
1212 return dec.offset;
1213 case DecorationBinding:
1214 return dec.binding;
1215 case DecorationDescriptorSet:
1216 return dec.set;
1217 case DecorationInputAttachmentIndex:
1218 return dec.input_attachment;
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02001219 case DecorationSpecId:
1220 return dec.spec_id;
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001221 case DecorationArrayStride:
1222 return dec.array_stride;
1223 case DecorationMatrixStride:
1224 return dec.matrix_stride;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001225 default:
Graham Wihlidal9b1ee8f2017-01-05 21:01:49 +01001226 return 1;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001227 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001228}
1229
1230void Compiler::unset_decoration(uint32_t id, Decoration decoration)
1231{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001232 auto &dec = meta.at(id).decoration;
1233 dec.decoration_flags &= ~(1ull << decoration);
1234 switch (decoration)
1235 {
1236 case DecorationBuiltIn:
1237 dec.builtin = false;
1238 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001239
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001240 case DecorationLocation:
1241 dec.location = 0;
1242 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001243
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001244 case DecorationOffset:
1245 dec.offset = 0;
1246 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001247
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001248 case DecorationBinding:
1249 dec.binding = 0;
1250 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001251
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001252 case DecorationDescriptorSet:
1253 dec.set = 0;
1254 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001255
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02001256 case DecorationInputAttachmentIndex:
1257 dec.input_attachment = 0;
1258 break;
1259
1260 case DecorationSpecId:
1261 dec.spec_id = 0;
1262 break;
1263
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001264 default:
1265 break;
1266 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001267}
1268
Hans-Kristian Arntzen93fe19c2017-04-25 20:10:24 +02001269bool Compiler::get_binary_offset_for_decoration(uint32_t id, spv::Decoration decoration, uint32_t &word_offset) const
1270{
1271 auto &word_offsets = meta.at(id).decoration_word_offset;
1272 auto itr = word_offsets.find(decoration);
1273 if (itr == end(word_offsets))
1274 return false;
1275
1276 word_offset = itr->second;
1277 return true;
1278}
1279
Hans-Kristian Arntzen926916d2016-05-05 09:15:25 +02001280void Compiler::parse(const Instruction &instruction)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001281{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001282 auto ops = stream(instruction);
1283 auto op = static_cast<Op>(instruction.op);
1284 uint32_t length = instruction.length;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001285
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001286 switch (op)
1287 {
1288 case OpMemoryModel:
1289 case OpSourceExtension:
1290 case OpNop:
Robert Konrad8f7c1af2016-08-10 02:43:51 +02001291 case OpLine:
Hans-Kristian Arntzen5b057f12017-11-21 10:11:40 +01001292 case OpNoLine:
David Srbeckya08b2692017-06-27 15:34:45 +01001293 case OpString:
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001294 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001295
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001296 case OpSource:
1297 {
1298 auto lang = static_cast<SourceLanguage>(ops[0]);
1299 switch (lang)
1300 {
1301 case SourceLanguageESSL:
1302 source.es = true;
1303 source.version = ops[1];
1304 source.known = true;
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02001305 source.hlsl = false;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001306 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001307
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001308 case SourceLanguageGLSL:
1309 source.es = false;
1310 source.version = ops[1];
1311 source.known = true;
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02001312 source.hlsl = false;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001313 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001314
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02001315 case SourceLanguageHLSL:
1316 // For purposes of cross-compiling, this is GLSL 450.
1317 source.es = false;
1318 source.version = 450;
1319 source.known = true;
1320 source.hlsl = true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001321 break;
1322
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001323 default:
1324 source.known = false;
1325 break;
1326 }
1327 break;
1328 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001329
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001330 case OpUndef:
1331 {
1332 uint32_t result_type = ops[0];
1333 uint32_t id = ops[1];
1334 set<SPIRUndef>(id, result_type);
1335 break;
1336 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001337
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001338 case OpCapability:
1339 {
1340 uint32_t cap = ops[0];
1341 if (cap == CapabilityKernel)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001342 SPIRV_CROSS_THROW("Kernel capability not supported.");
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02001343
1344 declared_capabilities.push_back(static_cast<Capability>(ops[0]));
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001345 break;
1346 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001347
Hans-Kristian Arntzen299e19f2017-03-21 16:33:54 +01001348 case OpExtension:
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02001349 {
1350 auto ext = extract_string(spirv, instruction.offset);
1351 declared_extensions.push_back(move(ext));
Hans-Kristian Arntzen299e19f2017-03-21 16:33:54 +01001352 break;
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02001353 }
Hans-Kristian Arntzen299e19f2017-03-21 16:33:54 +01001354
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001355 case OpExtInstImport:
1356 {
1357 uint32_t id = ops[0];
1358 auto ext = extract_string(spirv, instruction.offset + 1);
1359 if (ext == "GLSL.std.450")
1360 set<SPIRExtension>(id, SPIRExtension::GLSL);
Lou Kramer6671f522017-11-21 14:04:57 +01001361 else if (ext == "SPV_AMD_shader_ballot")
1362 set<SPIRExtension>(id, SPIRExtension::SPV_AMD_shader_ballot);
1363 else if (ext == "SPV_AMD_shader_explicit_vertex_parameter")
1364 set<SPIRExtension>(id, SPIRExtension::SPV_AMD_shader_explicit_vertex_parameter);
1365 else if (ext == "SPV_AMD_shader_trinary_minmax")
1366 set<SPIRExtension>(id, SPIRExtension::SPV_AMD_shader_trinary_minmax);
1367 else if (ext == "SPV_AMD_gcn_shader")
1368 set<SPIRExtension>(id, SPIRExtension::SPV_AMD_gcn_shader);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001369 else
Hans-Kristian Arntzen299e19f2017-03-21 16:33:54 +01001370 set<SPIRExtension>(id, SPIRExtension::Unsupported);
1371
1372 // Other SPIR-V extensions currently not supported.
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001373
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001374 break;
1375 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001376
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001377 case OpEntryPoint:
1378 {
Bill Hollingsac00c602016-10-24 09:24:24 -04001379 auto itr =
1380 entry_points.insert(make_pair(ops[1], SPIREntryPoint(ops[1], static_cast<ExecutionModel>(ops[0]),
1381 extract_string(spirv, instruction.offset + 2))));
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001382 auto &e = itr.first->second;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001383
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001384 // Strings need nul-terminator and consume the whole word.
Hans-Kristian Arntzen7630d3c2016-11-21 12:14:02 +01001385 uint32_t strlen_words = uint32_t((e.name.size() + 1 + 3) >> 2);
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001386 e.interface_variables.insert(end(e.interface_variables), ops + strlen_words + 2, ops + instruction.length);
1387
Bill Hollingsac00c602016-10-24 09:24:24 -04001388 // Set the name of the entry point in case OpName is not provided later
1389 set_name(ops[1], e.name);
Bill Hollings0943d9f2016-10-23 21:42:54 -04001390
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001391 // If we don't have an entry, make the first one our "default".
1392 if (!entry_point)
1393 entry_point = ops[1];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001394 break;
1395 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001396
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001397 case OpExecutionMode:
1398 {
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001399 auto &execution = entry_points[ops[0]];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001400 auto mode = static_cast<ExecutionMode>(ops[1]);
1401 execution.flags |= 1ull << mode;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001402
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001403 switch (mode)
1404 {
1405 case ExecutionModeInvocations:
1406 execution.invocations = ops[2];
1407 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001408
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001409 case ExecutionModeLocalSize:
1410 execution.workgroup_size.x = ops[2];
1411 execution.workgroup_size.y = ops[3];
1412 execution.workgroup_size.z = ops[4];
1413 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001414
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001415 case ExecutionModeOutputVertices:
1416 execution.output_vertices = ops[2];
1417 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001418
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001419 default:
1420 break;
1421 }
1422 break;
1423 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001424
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001425 case OpName:
1426 {
1427 uint32_t id = ops[0];
1428 set_name(id, extract_string(spirv, instruction.offset + 1));
1429 break;
1430 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001431
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001432 case OpMemberName:
1433 {
1434 uint32_t id = ops[0];
1435 uint32_t member = ops[1];
1436 set_member_name(id, member, extract_string(spirv, instruction.offset + 2));
1437 break;
1438 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001439
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001440 case OpDecorate:
1441 {
1442 uint32_t id = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001443
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001444 auto decoration = static_cast<Decoration>(ops[1]);
1445 if (length >= 3)
Hans-Kristian Arntzen93fe19c2017-04-25 20:10:24 +02001446 {
1447 meta[id].decoration_word_offset[decoration] = uint32_t(&ops[2] - spirv.data());
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001448 set_decoration(id, decoration, ops[2]);
Hans-Kristian Arntzen93fe19c2017-04-25 20:10:24 +02001449 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001450 else
1451 set_decoration(id, decoration);
Bill Hollings0943d9f2016-10-23 21:42:54 -04001452
Bill Hollingsac00c602016-10-24 09:24:24 -04001453 break;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001454 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001455
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001456 case OpMemberDecorate:
1457 {
1458 uint32_t id = ops[0];
1459 uint32_t member = ops[1];
1460 auto decoration = static_cast<Decoration>(ops[2]);
1461 if (length >= 4)
1462 set_member_decoration(id, member, decoration, ops[3]);
1463 else
1464 set_member_decoration(id, member, decoration);
1465 break;
1466 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001467
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001468 // Build up basic types.
1469 case OpTypeVoid:
1470 {
1471 uint32_t id = ops[0];
1472 auto &type = set<SPIRType>(id);
1473 type.basetype = SPIRType::Void;
1474 break;
1475 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001476
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001477 case OpTypeBool:
1478 {
1479 uint32_t id = ops[0];
1480 auto &type = set<SPIRType>(id);
Hans-Kristian Arntzen5f629272016-06-05 20:13:45 +02001481 type.basetype = SPIRType::Boolean;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001482 type.width = 1;
1483 break;
1484 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001485
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001486 case OpTypeFloat:
1487 {
1488 uint32_t id = ops[0];
1489 uint32_t width = ops[1];
1490 auto &type = set<SPIRType>(id);
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001491 type.basetype = width > 32 ? SPIRType::Double : SPIRType::Float;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001492 type.width = width;
1493 break;
1494 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001495
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001496 case OpTypeInt:
1497 {
1498 uint32_t id = ops[0];
1499 uint32_t width = ops[1];
1500 auto &type = set<SPIRType>(id);
Hans-Kristian Arntzenfc2230f2016-07-27 11:27:00 +02001501 type.basetype =
1502 ops[2] ? (width > 32 ? SPIRType::Int64 : SPIRType::Int) : (width > 32 ? SPIRType::UInt64 : SPIRType::UInt);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001503 type.width = width;
1504 break;
1505 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001506
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001507 // Build composite types by "inheriting".
1508 // NOTE: The self member is also copied! For pointers and array modifiers this is a good thing
1509 // since we can refer to decorations on pointee classes which is needed for UBO/SSBO, I/O blocks in geometry/tess etc.
1510 case OpTypeVector:
1511 {
1512 uint32_t id = ops[0];
1513 uint32_t vecsize = ops[2];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001514
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001515 auto &base = get<SPIRType>(ops[1]);
1516 auto &vecbase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001517
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001518 vecbase = base;
1519 vecbase.vecsize = vecsize;
1520 vecbase.self = id;
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001521 vecbase.parent_type = ops[1];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001522 break;
1523 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001524
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001525 case OpTypeMatrix:
1526 {
1527 uint32_t id = ops[0];
1528 uint32_t colcount = ops[2];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001529
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001530 auto &base = get<SPIRType>(ops[1]);
1531 auto &matrixbase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001532
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001533 matrixbase = base;
1534 matrixbase.columns = colcount;
1535 matrixbase.self = id;
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001536 matrixbase.parent_type = ops[1];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001537 break;
1538 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001539
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001540 case OpTypeArray:
1541 {
1542 uint32_t id = ops[0];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001543 auto &arraybase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001544
Bill Hollings1c180782017-11-05 21:34:42 -05001545 uint32_t tid = ops[1];
1546 auto &base = get<SPIRType>(tid);
Hans-Kristian Arntzen5d4bb682016-10-03 17:17:11 +02001547
Bill Hollings1c180782017-11-05 21:34:42 -05001548 arraybase = base;
1549 arraybase.parent_type = tid;
1550
1551 uint32_t cid = ops[2];
1552 mark_used_as_array_length(cid);
1553 auto *c = maybe_get<SPIRConstant>(cid);
Hans-Kristian Arntzen5d4bb682016-10-03 17:17:11 +02001554 bool literal = c && !c->specialization;
1555
1556 arraybase.array_size_literal.push_back(literal);
Bill Hollings1c180782017-11-05 21:34:42 -05001557 arraybase.array.push_back(literal ? c->scalar() : cid);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001558 // Do NOT set arraybase.self!
1559 break;
1560 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001561
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001562 case OpTypeRuntimeArray:
1563 {
1564 uint32_t id = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001565
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001566 auto &base = get<SPIRType>(ops[1]);
1567 auto &arraybase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001568
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001569 arraybase = base;
1570 arraybase.array.push_back(0);
Hans-Kristian Arntzen5d4bb682016-10-03 17:17:11 +02001571 arraybase.array_size_literal.push_back(true);
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001572 arraybase.parent_type = ops[1];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001573 // Do NOT set arraybase.self!
1574 break;
1575 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001576
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001577 case OpTypeImage:
1578 {
1579 uint32_t id = ops[0];
1580 auto &type = set<SPIRType>(id);
1581 type.basetype = SPIRType::Image;
1582 type.image.type = ops[1];
1583 type.image.dim = static_cast<Dim>(ops[2]);
1584 type.image.depth = ops[3] != 0;
1585 type.image.arrayed = ops[4] != 0;
1586 type.image.ms = ops[5] != 0;
1587 type.image.sampled = ops[6];
1588 type.image.format = static_cast<ImageFormat>(ops[7]);
Bill Hollingsb41e1482017-05-29 20:45:05 -04001589 type.image.access = (length >= 9) ? static_cast<AccessQualifier>(ops[8]) : AccessQualifierMax;
Bill Hollings6f3381a2017-06-01 16:29:39 -04001590
1591 if (type.image.sampled == 0)
1592 SPIRV_CROSS_THROW("OpTypeImage Sampled parameter must not be zero.");
1593
Robert Konrad98028232017-01-15 22:42:22 +01001594 break;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001595 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001596
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001597 case OpTypeSampledImage:
1598 {
1599 uint32_t id = ops[0];
1600 uint32_t imagetype = ops[1];
1601 auto &type = set<SPIRType>(id);
1602 type = get<SPIRType>(imagetype);
1603 type.basetype = SPIRType::SampledImage;
1604 type.self = id;
1605 break;
1606 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001607
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001608 case OpTypeSampler:
1609 {
1610 uint32_t id = ops[0];
1611 auto &type = set<SPIRType>(id);
1612 type.basetype = SPIRType::Sampler;
1613 break;
1614 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001615
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001616 case OpTypePointer:
1617 {
1618 uint32_t id = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001619
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001620 auto &base = get<SPIRType>(ops[2]);
1621 auto &ptrbase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001622
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001623 ptrbase = base;
1624 if (ptrbase.pointer)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001625 SPIRV_CROSS_THROW("Cannot make pointer-to-pointer type.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001626 ptrbase.pointer = true;
1627 ptrbase.storage = static_cast<StorageClass>(ops[1]);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001628
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001629 if (ptrbase.storage == StorageClassAtomicCounter)
1630 ptrbase.basetype = SPIRType::AtomicCounter;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001631
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01001632 ptrbase.parent_type = ops[2];
1633
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001634 // Do NOT set ptrbase.self!
1635 break;
1636 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001637
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001638 case OpTypeStruct:
1639 {
1640 uint32_t id = ops[0];
1641 auto &type = set<SPIRType>(id);
1642 type.basetype = SPIRType::Struct;
1643 for (uint32_t i = 1; i < length; i++)
1644 type.member_types.push_back(ops[i]);
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02001645
1646 // Check if we have seen this struct type before, with just different
1647 // decorations.
Hans-Kristian Arntzen5ad43402016-05-28 09:47:52 +02001648 //
1649 // Add workaround for issue #17 as well by looking at OpName for the struct
1650 // types, which we shouldn't normally do.
1651 // We should not normally have to consider type aliases like this to begin with
1652 // however ... glslang issues #304, #307 cover this.
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02001653
Hans-Kristian Arntzen818c1be2017-07-31 09:31:20 +02001654 // For stripped names, never consider struct type aliasing.
1655 // We risk declaring the same struct multiple times, but type-punning is not allowed
1656 // so this is safe.
1657 bool consider_aliasing = !get_name(type.self).empty();
1658 if (consider_aliasing)
1659 {
1660 for (auto &other : global_struct_cache)
1661 {
1662 if (get_name(type.self) == get_name(other) &&
1663 types_are_logically_equivalent(type, get<SPIRType>(other)))
1664 {
1665 type.type_alias = other;
1666 break;
1667 }
1668 }
1669
1670 if (type.type_alias == 0)
1671 global_struct_cache.push_back(id);
1672 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001673 break;
1674 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001675
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001676 case OpTypeFunction:
1677 {
1678 uint32_t id = ops[0];
1679 uint32_t ret = ops[1];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001680
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001681 auto &func = set<SPIRFunctionPrototype>(id, ret);
1682 for (uint32_t i = 2; i < length; i++)
1683 func.parameter_types.push_back(ops[i]);
1684 break;
1685 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001686
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001687 // Variable declaration
1688 // All variables are essentially pointers with a storage qualifier.
1689 case OpVariable:
1690 {
1691 uint32_t type = ops[0];
1692 uint32_t id = ops[1];
1693 auto storage = static_cast<StorageClass>(ops[2]);
1694 uint32_t initializer = length == 4 ? ops[3] : 0;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001695
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001696 if (storage == StorageClassFunction)
1697 {
1698 if (!current_function)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001699 SPIRV_CROSS_THROW("No function currently in scope");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001700 current_function->add_local_variable(id);
1701 }
1702 else if (storage == StorageClassPrivate || storage == StorageClassWorkgroup || storage == StorageClassOutput)
1703 {
1704 global_variables.push_back(id);
1705 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001706
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001707 auto &var = set<SPIRVariable>(id, type, storage, initializer);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001708
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001709 if (variable_storage_is_aliased(var))
1710 aliased_variables.push_back(var.self);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001711
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001712 break;
1713 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001714
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001715 // OpPhi
1716 // OpPhi is a fairly magical opcode.
1717 // It selects temporary variables based on which parent block we *came from*.
1718 // In high-level languages we can "de-SSA" by creating a function local, and flush out temporaries to this function-local
1719 // variable to emulate SSA Phi.
1720 case OpPhi:
1721 {
1722 if (!current_function)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001723 SPIRV_CROSS_THROW("No function currently in scope");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001724 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001725 SPIRV_CROSS_THROW("No block currently in scope");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001726
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001727 uint32_t result_type = ops[0];
1728 uint32_t id = ops[1];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001729
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001730 // Instead of a temporary, create a new function-wide temporary with this ID instead.
1731 auto &var = set<SPIRVariable>(id, result_type, spv::StorageClassFunction);
1732 var.phi_variable = true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001733
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001734 current_function->add_local_variable(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001735
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001736 for (uint32_t i = 2; i + 2 <= length; i += 2)
1737 current_block->phi_variables.push_back({ ops[i], ops[i + 1], id });
1738 break;
1739 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001740
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001741 // Constants
1742 case OpSpecConstant:
1743 case OpConstant:
1744 {
1745 uint32_t id = ops[1];
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001746 auto &type = get<SPIRType>(ops[0]);
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001747
Hans-Kristian Arntzenfc2230f2016-07-27 11:27:00 +02001748 if (type.width > 32)
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001749 set<SPIRConstant>(id, ops[0], ops[2] | (uint64_t(ops[3]) << 32), op == OpSpecConstant);
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001750 else
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001751 set<SPIRConstant>(id, ops[0], ops[2], op == OpSpecConstant);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001752 break;
1753 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001754
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001755 case OpSpecConstantFalse:
1756 case OpConstantFalse:
1757 {
1758 uint32_t id = ops[1];
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001759 set<SPIRConstant>(id, ops[0], uint32_t(0), op == OpSpecConstantFalse);
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 case OpSpecConstantTrue:
1764 case OpConstantTrue:
1765 {
1766 uint32_t id = ops[1];
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001767 set<SPIRConstant>(id, ops[0], uint32_t(1), op == OpSpecConstantTrue);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001768 break;
1769 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001770
Hans-Kristian Arntzen2abdc132017-08-02 10:33:03 +02001771 case OpConstantNull:
1772 {
1773 uint32_t id = ops[1];
1774 uint32_t type = ops[0];
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02001775 make_constant_null(id, type);
Hans-Kristian Arntzen2abdc132017-08-02 10:33:03 +02001776 break;
1777 }
1778
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001779 case OpSpecConstantComposite:
1780 case OpConstantComposite:
1781 {
1782 uint32_t id = ops[1];
1783 uint32_t type = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001784
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001785 auto &ctype = get<SPIRType>(type);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001786
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001787 // We can have constants which are structs and arrays.
1788 // In this case, our SPIRConstant will be a list of other SPIRConstant ids which we
1789 // can refer to.
1790 if (ctype.basetype == SPIRType::Struct || !ctype.array.empty())
1791 {
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001792 set<SPIRConstant>(id, type, ops + 2, length - 2, op == OpSpecConstantComposite);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001793 }
1794 else
1795 {
Hans-Kristian Arntzenceefae52017-09-27 16:10:29 +02001796 uint32_t elements = length - 2;
1797 if (elements > 4)
1798 SPIRV_CROSS_THROW("OpConstantComposite only supports 1, 2, 3 and 4 elements.");
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001799
1800 const SPIRConstant *c[4];
Hans-Kristian Arntzenceefae52017-09-27 16:10:29 +02001801 for (uint32_t i = 0; i < elements; i++)
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02001802 c[i] = &get<SPIRConstant>(ops[2 + i]);
Hans-Kristian Arntzenceefae52017-09-27 16:10:29 +02001803 set<SPIRConstant>(id, type, c, elements, op == OpSpecConstantComposite);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001804 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001805 break;
1806 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001807
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001808 // Functions
1809 case OpFunction:
1810 {
1811 uint32_t res = ops[0];
1812 uint32_t id = ops[1];
1813 // Control
1814 uint32_t type = ops[3];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001815
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001816 if (current_function)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001817 SPIRV_CROSS_THROW("Must end a function before starting a new one!");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001818
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001819 current_function = &set<SPIRFunction>(id, res, type);
1820 break;
1821 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001822
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001823 case OpFunctionParameter:
1824 {
1825 uint32_t type = ops[0];
1826 uint32_t id = ops[1];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001827
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001828 if (!current_function)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001829 SPIRV_CROSS_THROW("Must be in a function!");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001830
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001831 current_function->add_parameter(type, id);
1832 set<SPIRVariable>(id, type, StorageClassFunction);
1833 break;
1834 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001835
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001836 case OpFunctionEnd:
1837 {
Hans-Kristian Arntzen526d06d2016-11-12 10:03:18 +01001838 if (current_block)
1839 {
1840 // Very specific error message, but seems to come up quite often.
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001841 SPIRV_CROSS_THROW(
Hans-Kristian Arntzen526d06d2016-11-12 10:03:18 +01001842 "Cannot end a function before ending the current block.\n"
1843 "Likely cause: If this SPIR-V was created from glslang HLSL, make sure the entry point is valid.");
1844 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001845 current_function = nullptr;
1846 break;
1847 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001848
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001849 // Blocks
1850 case OpLabel:
1851 {
1852 // OpLabel always starts a block.
1853 if (!current_function)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001854 SPIRV_CROSS_THROW("Blocks cannot exist outside functions!");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001855
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001856 uint32_t id = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001857
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001858 current_function->blocks.push_back(id);
1859 if (!current_function->entry_block)
1860 current_function->entry_block = id;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001861
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001862 if (current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001863 SPIRV_CROSS_THROW("Cannot start a block before ending the current block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001864
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001865 current_block = &set<SPIRBlock>(id);
1866 break;
1867 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001868
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001869 // Branch instructions end blocks.
1870 case OpBranch:
1871 {
1872 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001873 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001874
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001875 uint32_t target = ops[0];
1876 current_block->terminator = SPIRBlock::Direct;
1877 current_block->next_block = target;
1878 current_block = nullptr;
1879 break;
1880 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001881
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001882 case OpBranchConditional:
1883 {
1884 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001885 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001886
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001887 current_block->condition = ops[0];
1888 current_block->true_block = ops[1];
1889 current_block->false_block = ops[2];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001890
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001891 current_block->terminator = SPIRBlock::Select;
1892 current_block = nullptr;
1893 break;
1894 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001895
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001896 case OpSwitch:
1897 {
1898 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001899 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001900
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001901 if (current_block->merge == SPIRBlock::MergeNone)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001902 SPIRV_CROSS_THROW("Switch statement is not structured");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001903
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001904 current_block->terminator = SPIRBlock::MultiSelect;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001905
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001906 current_block->condition = ops[0];
1907 current_block->default_block = ops[1];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001908
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001909 for (uint32_t i = 2; i + 2 <= length; i += 2)
1910 current_block->cases.push_back({ ops[i], ops[i + 1] });
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001911
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001912 // If we jump to next block, make it break instead since we're inside a switch case block at that point.
1913 multiselect_merge_targets.insert(current_block->next_block);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001914
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001915 current_block = nullptr;
1916 break;
1917 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001918
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001919 case OpKill:
1920 {
1921 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001922 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001923 current_block->terminator = SPIRBlock::Kill;
1924 current_block = nullptr;
1925 break;
1926 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001927
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001928 case OpReturn:
1929 {
1930 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001931 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001932 current_block->terminator = SPIRBlock::Return;
1933 current_block = nullptr;
1934 break;
1935 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001936
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001937 case OpReturnValue:
1938 {
1939 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001940 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001941 current_block->terminator = SPIRBlock::Return;
1942 current_block->return_value = ops[0];
1943 current_block = nullptr;
1944 break;
1945 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001946
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001947 case OpUnreachable:
1948 {
1949 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001950 SPIRV_CROSS_THROW("Trying to end a non-existing block.");
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001951 current_block->terminator = SPIRBlock::Unreachable;
1952 current_block = nullptr;
1953 break;
1954 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001955
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001956 case OpSelectionMerge:
1957 {
1958 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001959 SPIRV_CROSS_THROW("Trying to modify a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001960
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001961 current_block->next_block = ops[0];
1962 current_block->merge = SPIRBlock::MergeSelection;
1963 selection_merge_targets.insert(current_block->next_block);
1964 break;
1965 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001966
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001967 case OpLoopMerge:
1968 {
1969 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001970 SPIRV_CROSS_THROW("Trying to modify a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001971
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001972 current_block->merge_block = ops[0];
1973 current_block->continue_block = ops[1];
1974 current_block->merge = SPIRBlock::MergeLoop;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001975
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001976 loop_blocks.insert(current_block->self);
1977 loop_merge_targets.insert(current_block->merge_block);
Hans-Kristian Arntzen97f81ba2016-04-01 12:37:29 +02001978
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001979 // Don't add loop headers to continue blocks,
1980 // which would make it impossible branch into the loop header since
1981 // they are treated as continues.
1982 if (current_block->continue_block != current_block->self)
1983 continue_blocks.insert(current_block->continue_block);
1984 break;
1985 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001986
Hans-Kristian Arntzen7e8afa82016-10-03 15:54:02 +02001987 case OpSpecConstantOp:
1988 {
1989 if (length < 3)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01001990 SPIRV_CROSS_THROW("OpSpecConstantOp not enough arguments.");
Hans-Kristian Arntzen7e8afa82016-10-03 15:54:02 +02001991
1992 uint32_t result_type = ops[0];
1993 uint32_t id = ops[1];
1994 auto spec_op = static_cast<Op>(ops[2]);
1995
1996 set<SPIRConstantOp>(id, result_type, spec_op, ops + 3, length - 3);
1997 break;
1998 }
1999
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002000 // Actual opcodes.
2001 default:
2002 {
2003 if (!current_block)
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01002004 SPIRV_CROSS_THROW("Currently no block to insert opcode.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002005
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002006 current_block->ops.push_back(instruction);
2007 break;
2008 }
2009 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002010}
2011
2012bool Compiler::block_is_loop_candidate(const SPIRBlock &block, SPIRBlock::Method method) const
2013{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002014 // Tried and failed.
2015 if (block.disable_block_optimization || block.complex_continue)
2016 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002017
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002018 if (method == SPIRBlock::MergeToSelectForLoop)
2019 {
2020 // Try to detect common for loop pattern
2021 // which the code backend can use to create cleaner code.
2022 // for(;;) { if (cond) { some_body; } else { break; } }
2023 // is the pattern we're looking for.
2024 bool ret = block.terminator == SPIRBlock::Select && block.merge == SPIRBlock::MergeLoop &&
2025 block.true_block != block.merge_block && block.true_block != block.self &&
2026 block.false_block == block.merge_block;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002027
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002028 // If we have OpPhi which depends on branches which came from our own block,
2029 // we need to flush phi variables in else block instead of a trivial break,
2030 // so we cannot assume this is a for loop candidate.
2031 if (ret)
2032 {
2033 for (auto &phi : block.phi_variables)
2034 if (phi.parent == block.self)
2035 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002036
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002037 auto *merge = maybe_get<SPIRBlock>(block.merge_block);
2038 if (merge)
2039 for (auto &phi : merge->phi_variables)
2040 if (phi.parent == block.self)
2041 return false;
2042 }
2043 return ret;
2044 }
2045 else if (method == SPIRBlock::MergeToDirectForLoop)
2046 {
2047 // Empty loop header that just sets up merge target
2048 // and branches to loop body.
2049 bool ret = block.terminator == SPIRBlock::Direct && block.merge == SPIRBlock::MergeLoop && block.ops.empty();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002050
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002051 if (!ret)
2052 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002053
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002054 auto &child = get<SPIRBlock>(block.next_block);
2055 ret = child.terminator == SPIRBlock::Select && child.merge == SPIRBlock::MergeNone &&
2056 child.false_block == block.merge_block && child.true_block != block.merge_block &&
2057 child.true_block != block.self;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002058
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002059 // If we have OpPhi which depends on branches which came from our own block,
2060 // we need to flush phi variables in else block instead of a trivial break,
2061 // so we cannot assume this is a for loop candidate.
2062 if (ret)
2063 {
2064 for (auto &phi : block.phi_variables)
2065 if (phi.parent == block.self || phi.parent == child.self)
2066 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002067
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002068 for (auto &phi : child.phi_variables)
2069 if (phi.parent == block.self)
2070 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002071
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002072 auto *merge = maybe_get<SPIRBlock>(block.merge_block);
2073 if (merge)
2074 for (auto &phi : merge->phi_variables)
2075 if (phi.parent == block.self || phi.parent == child.false_block)
2076 return false;
2077 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002078
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002079 return ret;
2080 }
2081 else
2082 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002083}
2084
2085bool Compiler::block_is_outside_flow_control_from_block(const SPIRBlock &from, const SPIRBlock &to)
2086{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002087 auto *start = &from;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002088
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002089 if (start->self == to.self)
2090 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002091
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002092 // Break cycles.
2093 if (is_continue(start->self))
2094 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002095
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002096 // If our select block doesn't merge, we must break or continue in these blocks,
2097 // so if continues occur branchless within these blocks, consider them branchless as well.
2098 // This is typically used for loop control.
2099 if (start->terminator == SPIRBlock::Select && start->merge == SPIRBlock::MergeNone &&
2100 (block_is_outside_flow_control_from_block(get<SPIRBlock>(start->true_block), to) ||
2101 block_is_outside_flow_control_from_block(get<SPIRBlock>(start->false_block), to)))
2102 {
2103 return true;
2104 }
2105 else if (start->merge_block && block_is_outside_flow_control_from_block(get<SPIRBlock>(start->merge_block), to))
2106 {
2107 return true;
2108 }
2109 else if (start->next_block && block_is_outside_flow_control_from_block(get<SPIRBlock>(start->next_block), to))
2110 {
2111 return true;
2112 }
2113 else
2114 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002115}
2116
2117bool Compiler::execution_is_noop(const SPIRBlock &from, const SPIRBlock &to) const
2118{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002119 if (!execution_is_branchless(from, to))
2120 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002121
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002122 auto *start = &from;
2123 for (;;)
2124 {
2125 if (start->self == to.self)
2126 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002127
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002128 if (!start->ops.empty())
2129 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002130
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002131 start = &get<SPIRBlock>(start->next_block);
2132 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002133}
2134
2135bool Compiler::execution_is_branchless(const SPIRBlock &from, const SPIRBlock &to) const
2136{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002137 auto *start = &from;
2138 for (;;)
2139 {
2140 if (start->self == to.self)
2141 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002142
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002143 if (start->terminator == SPIRBlock::Direct && start->merge == SPIRBlock::MergeNone)
2144 start = &get<SPIRBlock>(start->next_block);
2145 else
2146 return false;
2147 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002148}
2149
Hans-Kristian Arntzen926916d2016-05-05 09:15:25 +02002150SPIRBlock::ContinueBlockType Compiler::continue_block_type(const SPIRBlock &block) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002151{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002152 // The block was deemed too complex during code emit, pick conservative fallback paths.
2153 if (block.complex_continue)
2154 return SPIRBlock::ComplexLoop;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002155
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002156 // In older glslang output continue block can be equal to the loop header.
2157 // In this case, execution is clearly branchless, so just assume a while loop header here.
2158 if (block.merge == SPIRBlock::MergeLoop)
2159 return SPIRBlock::WhileLoop;
Hans-Kristian Arntzen97f81ba2016-04-01 12:37:29 +02002160
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002161 auto &dominator = get<SPIRBlock>(block.loop_dominator);
Hans-Kristian Arntzen97f81ba2016-04-01 12:37:29 +02002162
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002163 if (execution_is_noop(block, dominator))
2164 return SPIRBlock::WhileLoop;
2165 else if (execution_is_branchless(block, dominator))
2166 return SPIRBlock::ForLoop;
2167 else
2168 {
2169 if (block.merge == SPIRBlock::MergeNone && block.terminator == SPIRBlock::Select &&
2170 block.true_block == dominator.self && block.false_block == dominator.merge_block)
2171 {
2172 return SPIRBlock::DoWhileLoop;
2173 }
2174 else
2175 return SPIRBlock::ComplexLoop;
2176 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002177}
2178
2179bool Compiler::traverse_all_reachable_opcodes(const SPIRBlock &block, OpcodeHandler &handler) const
2180{
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01002181 handler.set_current_block(block);
2182
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002183 // Ideally, perhaps traverse the CFG instead of all blocks in order to eliminate dead blocks,
2184 // but this shouldn't be a problem in practice unless the SPIR-V is doing insane things like recursing
2185 // inside dead blocks ...
2186 for (auto &i : block.ops)
2187 {
2188 auto ops = stream(i);
2189 auto op = static_cast<Op>(i.op);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002190
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002191 if (!handler.handle(op, ops, i.length))
2192 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002193
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002194 if (op == OpFunctionCall)
2195 {
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01002196 auto &func = get<SPIRFunction>(ops[2]);
2197 if (handler.follow_function_call(func))
2198 {
2199 if (!handler.begin_function_scope(ops, i.length))
2200 return false;
2201 if (!traverse_all_reachable_opcodes(get<SPIRFunction>(ops[2]), handler))
2202 return false;
2203 if (!handler.end_function_scope(ops, i.length))
2204 return false;
2205 }
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002206 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002207 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002208
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002209 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002210}
2211
2212bool Compiler::traverse_all_reachable_opcodes(const SPIRFunction &func, OpcodeHandler &handler) const
2213{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002214 for (auto block : func.blocks)
2215 if (!traverse_all_reachable_opcodes(get<SPIRBlock>(block), handler))
2216 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002217
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002218 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002219}
2220
2221uint32_t Compiler::type_struct_member_offset(const SPIRType &type, uint32_t index) const
2222{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002223 // Decoration must be set in valid SPIR-V, otherwise throw.
2224 auto &dec = meta[type.self].members.at(index);
2225 if (dec.decoration_flags & (1ull << DecorationOffset))
2226 return dec.offset;
2227 else
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01002228 SPIRV_CROSS_THROW("Struct member does not have Offset set.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002229}
2230
2231uint32_t Compiler::type_struct_member_array_stride(const SPIRType &type, uint32_t index) const
2232{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002233 // Decoration must be set in valid SPIR-V, otherwise throw.
2234 // ArrayStride is part of the array type not OpMemberDecorate.
2235 auto &dec = meta[type.member_types[index]].decoration;
2236 if (dec.decoration_flags & (1ull << DecorationArrayStride))
2237 return dec.array_stride;
2238 else
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01002239 SPIRV_CROSS_THROW("Struct member does not have ArrayStride set.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002240}
2241
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01002242uint32_t Compiler::type_struct_member_matrix_stride(const SPIRType &type, uint32_t index) const
2243{
2244 // Decoration must be set in valid SPIR-V, otherwise throw.
2245 // MatrixStride is part of OpMemberDecorate.
2246 auto &dec = meta[type.self].members[index];
2247 if (dec.decoration_flags & (1ull << DecorationMatrixStride))
2248 return dec.matrix_stride;
2249 else
2250 SPIRV_CROSS_THROW("Struct member does not have MatrixStride set.");
2251}
2252
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002253size_t Compiler::get_declared_struct_size(const SPIRType &type) const
2254{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002255 uint32_t last = uint32_t(type.member_types.size() - 1);
2256 size_t offset = type_struct_member_offset(type, last);
2257 size_t size = get_declared_struct_member_size(type, last);
2258 return offset + size;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002259}
2260
2261size_t Compiler::get_declared_struct_member_size(const SPIRType &struct_type, uint32_t index) const
2262{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002263 auto flags = get_member_decoration_mask(struct_type.self, index);
2264 auto &type = get<SPIRType>(struct_type.member_types[index]);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002265
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08002266 switch (type.basetype)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002267 {
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08002268 case SPIRType::Unknown:
2269 case SPIRType::Void:
2270 case SPIRType::Boolean: // Bools are purely logical, and cannot be used for externally visible types.
2271 case SPIRType::AtomicCounter:
2272 case SPIRType::Image:
2273 case SPIRType::SampledImage:
2274 case SPIRType::Sampler:
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01002275 SPIRV_CROSS_THROW("Querying size for object with opaque size.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002276
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08002277 default:
2278 break;
2279 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002280
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08002281 if (!type.array.empty())
2282 {
2283 // For arrays, we can use ArrayStride to get an easy check.
2284 return type_struct_member_array_stride(struct_type, index) * type.array.back();
2285 }
2286 else if (type.basetype == SPIRType::Struct)
2287 {
2288 return get_declared_struct_size(type);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002289 }
2290 else
2291 {
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002292 unsigned vecsize = type.vecsize;
2293 unsigned columns = type.columns;
Hans-Kristian Arntzen2d79d362016-11-28 15:01:36 +01002294
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08002295 // Vectors.
2296 if (columns == 1)
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01002297 {
2298 size_t component_size = type.width / 8;
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08002299 return vecsize * component_size;
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01002300 }
Hans-Kristian Arntzen2d79d362016-11-28 15:01:36 +01002301 else
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002302 {
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01002303 uint32_t matrix_stride = type_struct_member_matrix_stride(struct_type, index);
Hans-Kristian Arntzen2d79d362016-11-28 15:01:36 +01002304
Hans-Kristian Arntzend3cad992017-01-21 11:30:33 +01002305 // Per SPIR-V spec, matrices must be tightly packed and aligned up for vec3 accesses.
2306 if (flags & (1ull << DecorationRowMajor))
2307 return matrix_stride * vecsize;
2308 else if (flags & (1ull << DecorationColMajor))
2309 return matrix_stride * columns;
2310 else
2311 SPIRV_CROSS_THROW("Either row-major or column-major must be declared for matrices.");
Arseny Kapoulkinef63e7c52017-01-17 01:52:12 -08002312 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002313 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002314}
2315
2316bool Compiler::BufferAccessHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
2317{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002318 if (opcode != OpAccessChain && opcode != OpInBoundsAccessChain)
2319 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002320
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002321 // Invalid SPIR-V.
2322 if (length < 4)
2323 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002324
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002325 if (args[2] != id)
2326 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002327
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002328 // Don't bother traversing the entire access chain tree yet.
2329 // If we access a struct member, assume we access the entire member.
2330 uint32_t index = compiler.get<SPIRConstant>(args[3]).scalar();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002331
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002332 // Seen this index already.
2333 if (seen.find(index) != end(seen))
2334 return true;
2335 seen.insert(index);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002336
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002337 auto &type = compiler.expression_type(id);
2338 uint32_t offset = compiler.type_struct_member_offset(type, index);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002339
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002340 size_t range;
2341 // If we have another member in the struct, deduce the range by looking at the next member.
2342 // This is okay since structs in SPIR-V can have padding, but Offset decoration must be
2343 // monotonically increasing.
2344 // Of course, this doesn't take into account if the SPIR-V for some reason decided to add
2345 // very large amounts of padding, but that's not really a big deal.
2346 if (index + 1 < type.member_types.size())
2347 {
2348 range = compiler.type_struct_member_offset(type, index + 1) - offset;
2349 }
2350 else
2351 {
2352 // No padding, so just deduce it from the size of the member directly.
2353 range = compiler.get_declared_struct_member_size(type, index);
2354 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002355
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002356 ranges.push_back({ index, offset, range });
2357 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002358}
2359
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002360std::vector<BufferRange> Compiler::get_active_buffer_ranges(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002361{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002362 std::vector<BufferRange> ranges;
2363 BufferAccessHandler handler(*this, ranges, id);
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002364 traverse_all_reachable_opcodes(get<SPIRFunction>(entry_point), handler);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002365 return ranges;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002366}
2367
Bill Hollings103aabf2016-04-06 17:42:27 -04002368// Increase the number of IDs by the specified incremental amount.
2369// Returns the value of the first ID available for use in the expanded bound.
2370uint32_t Compiler::increase_bound_by(uint32_t incr_amount)
2371{
Hans-Kristian Arntzen14bc1ff2016-09-10 18:07:52 +02002372 auto curr_bound = ids.size();
2373 auto new_bound = curr_bound + incr_amount;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002374 ids.resize(new_bound);
2375 meta.resize(new_bound);
Hans-Kristian Arntzen14bc1ff2016-09-10 18:07:52 +02002376 return uint32_t(curr_bound);
Bill Hollings103aabf2016-04-06 17:42:27 -04002377}
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02002378
2379bool Compiler::types_are_logically_equivalent(const SPIRType &a, const SPIRType &b) const
2380{
2381 if (a.basetype != b.basetype)
2382 return false;
2383 if (a.width != b.width)
2384 return false;
2385 if (a.vecsize != b.vecsize)
2386 return false;
2387 if (a.columns != b.columns)
2388 return false;
2389 if (a.array.size() != b.array.size())
2390 return false;
2391
Hans-Kristian Arntzen46892ab2016-05-23 13:45:20 +02002392 size_t array_count = a.array.size();
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02002393 if (array_count && memcmp(a.array.data(), b.array.data(), array_count * sizeof(uint32_t)) != 0)
2394 return false;
2395
2396 if (a.basetype == SPIRType::Image || a.basetype == SPIRType::SampledImage)
2397 {
2398 if (memcmp(&a.image, &b.image, sizeof(SPIRType::Image)) != 0)
2399 return false;
2400 }
2401
2402 if (a.member_types.size() != b.member_types.size())
2403 return false;
2404
Hans-Kristian Arntzen46892ab2016-05-23 13:45:20 +02002405 size_t member_types = a.member_types.size();
2406 for (size_t i = 0; i < member_types; i++)
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02002407 {
2408 if (!types_are_logically_equivalent(get<SPIRType>(a.member_types[i]), get<SPIRType>(b.member_types[i])))
2409 return false;
2410 }
2411
2412 return true;
2413}
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002414
2415uint64_t Compiler::get_execution_mode_mask() const
2416{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002417 return get_entry_point().flags;
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002418}
2419
2420void Compiler::set_execution_mode(ExecutionMode mode, uint32_t arg0, uint32_t arg1, uint32_t arg2)
2421{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002422 auto &execution = get_entry_point();
2423
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002424 execution.flags |= 1ull << mode;
2425 switch (mode)
2426 {
2427 case ExecutionModeLocalSize:
2428 execution.workgroup_size.x = arg0;
2429 execution.workgroup_size.y = arg1;
2430 execution.workgroup_size.z = arg2;
2431 break;
2432
2433 case ExecutionModeInvocations:
2434 execution.invocations = arg0;
2435 break;
2436
2437 case ExecutionModeOutputVertices:
2438 execution.output_vertices = arg0;
2439 break;
2440
2441 default:
2442 break;
2443 }
2444}
2445
2446void Compiler::unset_execution_mode(ExecutionMode mode)
2447{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002448 auto &execution = get_entry_point();
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002449 execution.flags &= ~(1ull << mode);
2450}
2451
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +02002452uint32_t Compiler::get_work_group_size_specialization_constants(SpecializationConstant &x, SpecializationConstant &y,
2453 SpecializationConstant &z) const
2454{
2455 auto &execution = get_entry_point();
Hans-Kristian Arntzenca69b612017-11-06 09:49:52 +01002456 x = { 0, 0 };
2457 y = { 0, 0 };
2458 z = { 0, 0 };
Hans-Kristian Arntzen86eb8742017-09-28 11:33:30 +02002459
2460 if (execution.workgroup_size.constant != 0)
2461 {
2462 auto &c = get<SPIRConstant>(execution.workgroup_size.constant);
2463
2464 if (c.m.c[0].id[0] != 0)
2465 {
2466 x.id = c.m.c[0].id[0];
2467 x.constant_id = get_decoration(c.m.c[0].id[0], DecorationSpecId);
2468 }
2469
2470 if (c.m.c[0].id[1] != 0)
2471 {
2472 y.id = c.m.c[0].id[1];
2473 y.constant_id = get_decoration(c.m.c[0].id[1], DecorationSpecId);
2474 }
2475
2476 if (c.m.c[0].id[2] != 0)
2477 {
2478 z.id = c.m.c[0].id[2];
2479 z.constant_id = get_decoration(c.m.c[0].id[2], DecorationSpecId);
2480 }
2481 }
2482
2483 return execution.workgroup_size.constant;
2484}
2485
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002486uint32_t Compiler::get_execution_mode_argument(spv::ExecutionMode mode, uint32_t index) const
2487{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002488 auto &execution = get_entry_point();
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002489 switch (mode)
2490 {
2491 case ExecutionModeLocalSize:
2492 switch (index)
2493 {
2494 case 0:
2495 return execution.workgroup_size.x;
2496 case 1:
2497 return execution.workgroup_size.y;
2498 case 2:
2499 return execution.workgroup_size.z;
2500 default:
2501 return 0;
2502 }
2503
2504 case ExecutionModeInvocations:
2505 return execution.invocations;
2506
2507 case ExecutionModeOutputVertices:
2508 return execution.output_vertices;
2509
2510 default:
2511 return 0;
2512 }
2513}
2514
2515ExecutionModel Compiler::get_execution_model() const
2516{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002517 auto &execution = get_entry_point();
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002518 return execution.model;
2519}
Hans-Kristian Arntzen8e63c772016-07-06 09:58:01 +02002520
2521void Compiler::set_remapped_variable_state(uint32_t id, bool remap_enable)
2522{
2523 get<SPIRVariable>(id).remapped_variable = remap_enable;
2524}
2525
2526bool Compiler::get_remapped_variable_state(uint32_t id) const
2527{
2528 return get<SPIRVariable>(id).remapped_variable;
2529}
Hans-Kristian Arntzen078eec52016-07-06 11:04:06 +02002530
2531void Compiler::set_subpass_input_remapped_components(uint32_t id, uint32_t components)
2532{
2533 get<SPIRVariable>(id).remapped_components = components;
2534}
2535
2536uint32_t Compiler::get_subpass_input_remapped_components(uint32_t id) const
2537{
2538 return get<SPIRVariable>(id).remapped_components;
2539}
Hans-Kristian Arntzen36a0b632016-07-12 14:33:04 +02002540
2541void Compiler::inherit_expression_dependencies(uint32_t dst, uint32_t source_expression)
2542{
Hans-Kristian Arntzen75391f92017-03-20 22:38:05 +01002543 // Don't inherit any expression dependencies if the expression in dst
2544 // is not a forwarded temporary.
2545 if (forwarded_temporaries.find(dst) == end(forwarded_temporaries) ||
2546 forced_temporaries.find(dst) != end(forced_temporaries))
2547 {
2548 return;
2549 }
2550
Hans-Kristian Arntzen36a0b632016-07-12 14:33:04 +02002551 auto &e = get<SPIRExpression>(dst);
2552 auto *s = maybe_get<SPIRExpression>(source_expression);
2553 if (!s)
2554 return;
2555
2556 auto &e_deps = e.expression_dependencies;
2557 auto &s_deps = s->expression_dependencies;
2558
2559 // If we depend on a expression, we also depend on all sub-dependencies from source.
2560 e_deps.push_back(source_expression);
2561 e_deps.insert(end(e_deps), begin(s_deps), end(s_deps));
2562
2563 // Eliminate duplicated dependencies.
2564 e_deps.erase(unique(begin(e_deps), end(e_deps)), end(e_deps));
2565}
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002566
2567vector<string> Compiler::get_entry_points() const
2568{
2569 vector<string> entries;
2570 for (auto &entry : entry_points)
Bill Hollings10148472017-11-10 16:40:33 -05002571 entries.push_back(entry.second.orig_name);
Bill Hollings1c180782017-11-05 21:34:42 -05002572 return entries;
2573}
2574
Hans-Kristian Arntzen4427cb92017-11-13 13:49:11 +01002575void Compiler::rename_entry_point(const std::string &old_name, const std::string &new_name)
2576{
2577 auto &entry = get_entry_point(old_name);
2578 entry.orig_name = new_name;
2579 entry.name = new_name;
2580}
2581
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002582void Compiler::set_entry_point(const std::string &name)
2583{
2584 auto &entry = get_entry_point(name);
2585 entry_point = entry.self;
2586}
2587
2588SPIREntryPoint &Compiler::get_entry_point(const std::string &name)
2589{
2590 auto itr =
Bill Hollings10148472017-11-10 16:40:33 -05002591 find_if(begin(entry_points), end(entry_points), [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool {
2592 return entry.second.orig_name == name;
Hans-Kristian Arntzence18d4c2017-11-17 13:38:29 +01002593 });
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002594
2595 if (itr == end(entry_points))
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01002596 SPIRV_CROSS_THROW("Entry point does not exist.");
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002597
2598 return itr->second;
2599}
2600
2601const SPIREntryPoint &Compiler::get_entry_point(const std::string &name) const
2602{
2603 auto itr =
Bill Hollings10148472017-11-10 16:40:33 -05002604 find_if(begin(entry_points), end(entry_points), [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool {
2605 return entry.second.orig_name == name;
Hans-Kristian Arntzence18d4c2017-11-17 13:38:29 +01002606 });
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002607
2608 if (itr == end(entry_points))
Panagiotis Christopoulos Charitos946f7792016-12-12 22:33:22 +01002609 SPIRV_CROSS_THROW("Entry point does not exist.");
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002610
2611 return itr->second;
2612}
2613
Bill Hollings10148472017-11-10 16:40:33 -05002614const string &Compiler::get_cleansed_entry_point_name(const std::string &name) const
2615{
2616 return get_entry_point(name).name;
2617}
2618
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002619const SPIREntryPoint &Compiler::get_entry_point() const
2620{
2621 return entry_points.find(entry_point)->second;
2622}
2623
2624SPIREntryPoint &Compiler::get_entry_point()
2625{
2626 return entry_points.find(entry_point)->second;
2627}
2628
2629bool Compiler::interface_variable_exists_in_entry_point(uint32_t id) const
2630{
2631 auto &var = get<SPIRVariable>(id);
Robert Konrada778c362017-01-15 16:39:03 +01002632 if (var.storage != StorageClassInput && var.storage != StorageClassOutput &&
2633 var.storage != StorageClassUniformConstant)
Hans-Kristian Arntzen24df8f02017-02-04 10:26:26 +01002634 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 +02002635
2636 // This is to avoid potential problems with very old glslang versions which did
2637 // not emit input/output interfaces properly.
2638 // We can assume they only had a single entry point, and single entry point
2639 // shaders could easily be assumed to use every interface variable anyways.
2640 if (entry_points.size() <= 1)
2641 return true;
2642
2643 auto &execution = get_entry_point();
2644 return find(begin(execution.interface_variables), end(execution.interface_variables), id) !=
2645 end(execution.interface_variables);
2646}
Hans-Kristian Arntzenbcb55602016-09-10 13:56:36 +02002647
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002648void Compiler::CombinedImageSamplerHandler::push_remap_parameters(const SPIRFunction &func, const uint32_t *args,
2649 uint32_t length)
2650{
2651 // If possible, pipe through a remapping table so that parameters know
2652 // which variables they actually bind to in this scope.
2653 unordered_map<uint32_t, uint32_t> remapping;
2654 for (uint32_t i = 0; i < length; i++)
2655 remapping[func.arguments[i].id] = remap_parameter(args[i]);
2656 parameter_remapping.push(move(remapping));
2657}
2658
2659void Compiler::CombinedImageSamplerHandler::pop_remap_parameters()
2660{
2661 parameter_remapping.pop();
2662}
2663
2664uint32_t Compiler::CombinedImageSamplerHandler::remap_parameter(uint32_t id)
2665{
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002666 auto *var = compiler.maybe_get_backing_variable(id);
2667 if (var)
2668 id = var->self;
2669
Hans-Kristian Arntzen901b45e2016-09-10 22:21:57 +02002670 if (parameter_remapping.empty())
2671 return id;
2672
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002673 auto &remapping = parameter_remapping.top();
2674 auto itr = remapping.find(id);
2675 if (itr != end(remapping))
2676 return itr->second;
2677 else
2678 return id;
2679}
2680
2681bool Compiler::CombinedImageSamplerHandler::begin_function_scope(const uint32_t *args, uint32_t length)
2682{
2683 if (length < 3)
2684 return false;
2685
2686 auto &callee = compiler.get<SPIRFunction>(args[2]);
2687 args += 3;
2688 length -= 3;
2689 push_remap_parameters(callee, args, length);
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002690 functions.push(&callee);
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002691 return true;
2692}
2693
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002694bool Compiler::CombinedImageSamplerHandler::end_function_scope(const uint32_t *args, uint32_t length)
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002695{
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002696 if (length < 3)
2697 return false;
2698
2699 auto &callee = compiler.get<SPIRFunction>(args[2]);
2700 args += 3;
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002701
2702 // There are two types of cases we have to handle,
2703 // a callee might call sampler2D(texture2D, sampler) directly where
2704 // one or more parameters originate from parameters.
2705 // Alternatively, we need to provide combined image samplers to our callees,
2706 // and in this case we need to add those as well.
2707
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002708 pop_remap_parameters();
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002709
2710 // Our callee has now been processed at least once.
2711 // No point in doing it again.
2712 callee.do_combined_parameters = false;
2713
2714 auto &params = functions.top()->combined_parameters;
2715 functions.pop();
2716 if (functions.empty())
2717 return true;
2718
2719 auto &caller = *functions.top();
2720 if (caller.do_combined_parameters)
2721 {
2722 for (auto &param : params)
2723 {
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002724 uint32_t image_id = param.global_image ? param.image_id : args[param.image_id];
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002725 uint32_t sampler_id = param.global_sampler ? param.sampler_id : args[param.sampler_id];
2726
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002727 auto *i = compiler.maybe_get_backing_variable(image_id);
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002728 auto *s = compiler.maybe_get_backing_variable(sampler_id);
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002729 if (i)
2730 image_id = i->self;
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002731 if (s)
2732 sampler_id = s->self;
2733
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002734 register_combined_image_sampler(caller, image_id, sampler_id, param.depth);
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002735 }
2736 }
2737
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002738 return true;
2739}
2740
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002741void Compiler::CombinedImageSamplerHandler::register_combined_image_sampler(SPIRFunction &caller, uint32_t image_id,
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002742 uint32_t sampler_id, bool depth)
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002743{
2744 // We now have a texture ID and a sampler ID which will either be found as a global
2745 // or a parameter in our own function. If both are global, they will not need a parameter,
2746 // otherwise, add it to our list.
2747 SPIRFunction::CombinedImageSamplerParameter param = {
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002748 0u, image_id, sampler_id, true, true, depth,
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002749 };
2750
2751 auto texture_itr = find_if(begin(caller.arguments), end(caller.arguments),
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002752 [image_id](const SPIRFunction::Parameter &p) { return p.id == image_id; });
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002753 auto sampler_itr = find_if(begin(caller.arguments), end(caller.arguments),
2754 [sampler_id](const SPIRFunction::Parameter &p) { return p.id == sampler_id; });
2755
2756 if (texture_itr != end(caller.arguments))
2757 {
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002758 param.global_image = false;
Hans-Kristian Arntzen7630d3c2016-11-21 12:14:02 +01002759 param.image_id = uint32_t(texture_itr - begin(caller.arguments));
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002760 }
2761
2762 if (sampler_itr != end(caller.arguments))
2763 {
2764 param.global_sampler = false;
Hans-Kristian Arntzen7630d3c2016-11-21 12:14:02 +01002765 param.sampler_id = uint32_t(sampler_itr - begin(caller.arguments));
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002766 }
2767
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002768 if (param.global_image && param.global_sampler)
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002769 return;
2770
2771 auto itr = find_if(begin(caller.combined_parameters), end(caller.combined_parameters),
2772 [&param](const SPIRFunction::CombinedImageSamplerParameter &p) {
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002773 return param.image_id == p.image_id && param.sampler_id == p.sampler_id &&
2774 param.global_image == p.global_image && param.global_sampler == p.global_sampler;
Hans-Kristian Arntzence18d4c2017-11-17 13:38:29 +01002775 });
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002776
2777 if (itr == end(caller.combined_parameters))
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002778 {
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002779 uint32_t id = compiler.increase_bound_by(3);
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002780 auto type_id = id + 0;
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002781 auto ptr_type_id = id + 1;
2782 auto combined_id = id + 2;
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002783 auto &base = compiler.expression_type(image_id);
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002784 auto &type = compiler.set<SPIRType>(type_id);
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002785 auto &ptr_type = compiler.set<SPIRType>(ptr_type_id);
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002786
2787 type = base;
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002788 type.self = type_id;
2789 type.basetype = SPIRType::SampledImage;
2790 type.pointer = false;
2791 type.storage = StorageClassGeneric;
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002792 type.image.depth = depth;
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002793
2794 ptr_type = type;
2795 ptr_type.pointer = true;
2796 ptr_type.storage = StorageClassUniformConstant;
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002797
2798 // Build new variable.
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002799 compiler.set<SPIRVariable>(combined_id, ptr_type_id, StorageClassFunction, 0);
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002800
2801 // Inherit RelaxedPrecision (and potentially other useful flags if deemed relevant).
2802 auto &new_flags = compiler.meta[combined_id].decoration.decoration_flags;
2803 auto old_flags = compiler.meta[sampler_id].decoration.decoration_flags;
2804 new_flags = old_flags & (1ull << DecorationRelaxedPrecision);
2805
2806 param.id = combined_id;
2807
2808 compiler.set_name(combined_id,
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002809 join("SPIRV_Cross_Combined", compiler.to_name(image_id), compiler.to_name(sampler_id)));
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002810
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002811 caller.combined_parameters.push_back(param);
Hans-Kristian Arntzen9cb86162017-02-05 10:50:14 +01002812 caller.shadow_arguments.push_back({ ptr_type_id, combined_id, 0u, 0u, true });
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002813 }
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002814}
2815
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002816bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
2817{
2818 // We need to figure out where samplers and images are loaded from, so do only the bare bones compilation we need.
2819 switch (opcode)
2820 {
2821 case OpLoad:
2822 {
2823 if (length < 3)
2824 return false;
2825
2826 uint32_t result_type = args[0];
2827
2828 auto &type = compiler.get<SPIRType>(result_type);
2829 bool separate_image = type.basetype == SPIRType::Image && type.image.sampled == 1;
2830 bool separate_sampler = type.basetype == SPIRType::Sampler;
2831
Hans-Kristian Arntzen378fbe82016-09-11 13:47:06 +02002832 // If not separate image or sampler, don't bother.
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002833 if (!separate_image && !separate_sampler)
2834 return true;
2835
2836 uint32_t id = args[1];
2837 uint32_t ptr = args[2];
2838 compiler.set<SPIRExpression>(id, "", result_type, true);
2839 compiler.register_read(id, ptr, true);
2840 return true;
2841 }
2842
2843 case OpInBoundsAccessChain:
2844 case OpAccessChain:
2845 {
2846 if (length < 3)
2847 return false;
2848
2849 // Technically, it is possible to have arrays of textures and arrays of samplers and combine them, but this becomes essentially
2850 // impossible to implement, since we don't know which concrete sampler we are accessing.
2851 // One potential way is to create a combinatorial explosion where N textures and M samplers are combined into N * M sampler2Ds,
2852 // but this seems ridiculously complicated for a problem which is easy to work around.
2853 // Checking access chains like this assumes we don't have samplers or textures inside uniform structs, but this makes no sense.
2854
2855 auto &type = compiler.get<SPIRType>(args[0]);
2856 bool separate_image = type.basetype == SPIRType::Image && type.image.sampled == 1;
2857 bool separate_sampler = type.basetype == SPIRType::Sampler;
2858 if (separate_image)
Hans-Kristian Arntzen36e1c472017-05-06 13:59:00 +02002859 SPIRV_CROSS_THROW("Attempting to use arrays or structs of separate images. This is not possible to "
2860 "statically remap to plain GLSL.");
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002861 if (separate_sampler)
Hans-Kristian Arntzen36e1c472017-05-06 13:59:00 +02002862 SPIRV_CROSS_THROW(
2863 "Attempting to use arrays or structs of separate samplers. This is not possible to statically "
2864 "remap to plain GLSL.");
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002865 return true;
2866 }
2867
2868 case OpSampledImage:
2869 // Do it outside.
2870 break;
2871
2872 default:
2873 return true;
2874 }
2875
2876 if (length < 4)
2877 return false;
2878
Hans-Kristian Arntzendfb65972016-09-11 12:05:20 +02002879 // Registers sampler2D calls used in case they are parameters so
2880 // that their callees know which combined image samplers to propagate down the call stack.
2881 if (!functions.empty())
2882 {
2883 auto &callee = *functions.top();
2884 if (callee.do_combined_parameters)
2885 {
2886 uint32_t image_id = args[2];
2887
2888 auto *image = compiler.maybe_get_backing_variable(image_id);
2889 if (image)
2890 image_id = image->self;
2891
2892 uint32_t sampler_id = args[3];
2893 auto *sampler = compiler.maybe_get_backing_variable(sampler_id);
2894 if (sampler)
2895 sampler_id = sampler->self;
2896
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02002897 auto &combined_type = compiler.get<SPIRType>(args[0]);
2898 register_combined_image_sampler(callee, image_id, sampler_id, combined_type.image.depth);
Hans-Kristian Arntzendfb65972016-09-11 12:05:20 +02002899 }
2900 }
2901
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002902 // For function calls, we need to remap IDs which are function parameters into global variables.
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002903 // This information is statically known from the current place in the call stack.
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002904 // Function parameters are not necessarily pointers, so if we don't have a backing variable, remapping will know
2905 // which backing variable the image/sample came from.
Hans-Kristian Arntzendfb65972016-09-11 12:05:20 +02002906 uint32_t image_id = remap_parameter(args[2]);
2907 uint32_t sampler_id = remap_parameter(args[3]);
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002908
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002909 auto itr = find_if(begin(compiler.combined_image_samplers), end(compiler.combined_image_samplers),
2910 [image_id, sampler_id](const CombinedImageSampler &combined) {
2911 return combined.image_id == image_id && combined.sampler_id == sampler_id;
Hans-Kristian Arntzence18d4c2017-11-17 13:38:29 +01002912 });
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002913
2914 if (itr == end(compiler.combined_image_samplers))
2915 {
2916 auto id = compiler.increase_bound_by(2);
2917 auto type_id = id + 0;
2918 auto combined_id = id + 1;
2919 auto sampled_type = args[0];
2920
2921 // Make a new type, pointer to OpTypeSampledImage, so we can make a variable of this type.
2922 // We will probably have this type lying around, but it doesn't hurt to make duplicates for internal purposes.
2923 auto &type = compiler.set<SPIRType>(type_id);
2924 auto &base = compiler.get<SPIRType>(sampled_type);
2925 type = base;
2926 type.pointer = true;
2927 type.storage = StorageClassUniformConstant;
2928
2929 // Build new variable.
Hans-Kristian Arntzen948930b2016-09-11 12:36:12 +02002930 compiler.set<SPIRVariable>(combined_id, type_id, StorageClassUniformConstant, 0);
Hans-Kristian Arntzen14bc1ff2016-09-10 18:07:52 +02002931
2932 // Inherit RelaxedPrecision (and potentially other useful flags if deemed relevant).
2933 auto &new_flags = compiler.meta[combined_id].decoration.decoration_flags;
2934 auto old_flags = compiler.meta[sampler_id].decoration.decoration_flags;
2935 new_flags = old_flags & (1ull << DecorationRelaxedPrecision);
2936
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002937 compiler.combined_image_samplers.push_back({ combined_id, image_id, sampler_id });
2938 }
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002939
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002940 return true;
2941}
2942
Hans-Kristian Arntzenbcb55602016-09-10 13:56:36 +02002943void Compiler::build_combined_image_samplers()
2944{
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002945 for (auto &id : ids)
2946 {
2947 if (id.get_type() == TypeFunction)
2948 {
2949 auto &func = id.get<SPIRFunction>();
2950 func.combined_parameters.clear();
Hans-Kristian Arntzen313cb5f2016-09-11 12:54:08 +02002951 func.shadow_arguments.clear();
Hans-Kristian Arntzened98a8e2016-09-11 11:39:20 +02002952 func.do_combined_parameters = true;
2953 }
2954 }
2955
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002956 combined_image_samplers.clear();
2957 CombinedImageSamplerHandler handler(*this);
2958 traverse_all_reachable_opcodes(get<SPIRFunction>(entry_point), handler);
Hans-Kristian Arntzenbcb55602016-09-10 13:56:36 +02002959}
Hans-Kristian Arntzen6bd545b2016-09-17 15:16:07 +02002960
2961vector<SpecializationConstant> Compiler::get_specialization_constants() const
2962{
2963 vector<SpecializationConstant> spec_consts;
2964 for (auto &id : ids)
2965 {
2966 if (id.get_type() == TypeConstant)
2967 {
2968 auto &c = id.get<SPIRConstant>();
2969 if (c.specialization)
2970 {
2971 spec_consts.push_back({ c.self, get_decoration(c.self, DecorationSpecId) });
2972 }
2973 }
2974 }
2975 return spec_consts;
2976}
2977
2978SPIRConstant &Compiler::get_constant(uint32_t id)
2979{
2980 return get<SPIRConstant>(id);
2981}
2982
2983const SPIRConstant &Compiler::get_constant(uint32_t id) const
2984{
2985 return get<SPIRConstant>(id);
2986}
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01002987
Bill Hollingsbac657d2017-11-07 15:38:13 -05002988// Recursively marks any constants referenced by the specified constant instruction as being used
2989// as an array length. The id must be a constant instruction (SPIRConstant or SPIRConstantOp).
Bill Hollings1c180782017-11-05 21:34:42 -05002990void Compiler::mark_used_as_array_length(uint32_t id)
2991{
2992 switch (ids[id].get_type())
2993 {
2994 case TypeConstant:
2995 get<SPIRConstant>(id).is_used_as_array_length = true;
2996 break;
2997
2998 case TypeConstantOp:
2999 {
3000 auto &cop = get<SPIRConstantOp>(id);
3001 for (uint32_t arg_id : cop.arguments)
3002 mark_used_as_array_length(arg_id);
3003 }
3004
3005 case TypeUndef:
3006 return;
3007
3008 default:
3009 SPIRV_CROSS_THROW("Array lengths must be a constant instruction (OpConstant.. or OpSpecConstant...).");
3010 }
3011}
3012
Hans-Kristian Arntzenbf5c0752017-03-25 16:28:44 +01003013static 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 +01003014{
3015 // This block accesses the variable.
3016 if (blocks.find(block) != end(blocks))
3017 return false;
3018
3019 // We are at the end of the CFG.
3020 if (cfg.get_succeeding_edges(block).empty())
3021 return true;
3022
3023 // If any of our successors have a path to the end, there exists a path from block.
3024 for (auto &succ : cfg.get_succeeding_edges(block))
3025 if (exists_unaccessed_path_to_return(cfg, succ, blocks))
3026 return true;
3027
3028 return false;
3029}
3030
Hans-Kristian Arntzenbf5c0752017-03-25 16:28:44 +01003031void Compiler::analyze_parameter_preservation(
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02003032 SPIRFunction &entry, const CFG &cfg, const unordered_map<uint32_t, unordered_set<uint32_t>> &variable_to_blocks,
3033 const unordered_map<uint32_t, unordered_set<uint32_t>> &complete_write_blocks)
Hans-Kristian Arntzenb2c2e642017-03-25 16:25:30 +01003034{
3035 for (auto &arg : entry.arguments)
3036 {
3037 // Non-pointers are always inputs.
3038 auto &type = get<SPIRType>(arg.type);
3039 if (!type.pointer)
3040 continue;
3041
3042 // Opaque argument types are always in
3043 bool potential_preserve;
3044 switch (type.basetype)
3045 {
3046 case SPIRType::Sampler:
3047 case SPIRType::Image:
3048 case SPIRType::SampledImage:
3049 case SPIRType::AtomicCounter:
3050 potential_preserve = false;
3051 break;
3052
3053 default:
3054 potential_preserve = true;
3055 break;
3056 }
3057
3058 if (!potential_preserve)
3059 continue;
3060
3061 auto itr = variable_to_blocks.find(arg.id);
3062 if (itr == end(variable_to_blocks))
3063 {
3064 // Variable is never accessed.
3065 continue;
3066 }
3067
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02003068 // We have accessed a variable, but there was no complete writes to that variable.
3069 // We deduce that we must preserve the argument.
3070 itr = complete_write_blocks.find(arg.id);
3071 if (itr == end(complete_write_blocks))
3072 {
3073 arg.read_count++;
3074 continue;
3075 }
3076
3077 // 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 +01003078 // when the function returns. We therefore need to implicitly preserve the variable in case there are writers in the function.
3079 // Major case here is if a function is
3080 // void foo(int &var) { if (cond) var = 10; }
3081 // Using read/write counts, we will think it's just an out variable, but it really needs to be inout,
3082 // because if we don't write anything whatever we put into the function must return back to the caller.
3083 if (exists_unaccessed_path_to_return(cfg, entry.entry_block, itr->second))
3084 arg.read_count++;
3085 }
3086}
3087
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003088void Compiler::analyze_variable_scope(SPIRFunction &entry)
3089{
3090 struct AccessHandler : OpcodeHandler
3091 {
3092 public:
Hans-Kristian Arntzen07584282017-11-21 18:51:11 +01003093 AccessHandler(Compiler &compiler_, SPIRFunction &entry_)
Hans-Kristian Arntzen1d63d142017-11-22 11:30:09 +01003094 : compiler(compiler_)
3095 , entry(entry_)
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003096 {
3097 }
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +01003098
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003099 bool follow_function_call(const SPIRFunction &)
3100 {
3101 // Only analyze within this function.
3102 return false;
3103 }
3104
3105 void set_current_block(const SPIRBlock &block)
3106 {
3107 current_block = &block;
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003108
3109 // If we're branching to a block which uses OpPhi, in GLSL
3110 // this will be a variable write when we branch,
3111 // so we need to track access to these variables as well to
3112 // have a complete picture.
3113 const auto test_phi = [this, &block](uint32_t to) {
3114 auto &next = compiler.get<SPIRBlock>(to);
3115 for (auto &phi : next.phi_variables)
Hans-Kristian Arntzena1c0ab62017-08-21 09:36:53 +02003116 {
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003117 if (phi.parent == block.self)
Hans-Kristian Arntzena1c0ab62017-08-21 09:36:53 +02003118 {
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003119 accessed_variables_to_block[phi.function_variable].insert(block.self);
Hans-Kristian Arntzena1c0ab62017-08-21 09:36:53 +02003120 // Phi variables are also accessed in our target branch block.
3121 accessed_variables_to_block[phi.function_variable].insert(next.self);
3122 }
3123 }
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003124 };
3125
Hans-Kristian Arntzencdca1922017-11-21 16:52:01 +01003126 // A Phi node might be reading other phi variables as input, so check for this as well.
3127 for (auto &phi : block.phi_variables)
Hans-Kristian Arntzen0fd02812017-11-21 18:19:51 +01003128 {
Hans-Kristian Arntzencdca1922017-11-21 16:52:01 +01003129 if (id_is_phi_variable(phi.local_variable))
3130 accessed_variables_to_block[phi.local_variable].insert(block.self);
Hans-Kristian Arntzen0fd02812017-11-21 18:19:51 +01003131 else
3132 {
3133 // Temporary variable, due to potential issues with scoping,
3134 // always declare these variables up-front in the entry block.
3135 if (!compiler.hoisted_temporaries.count(phi.local_variable))
3136 {
Hans-Kristian Arntzen65cd4172017-11-22 14:25:00 +01003137 auto *undef = compiler.maybe_get<SPIRUndef>(phi.local_variable);
3138 // Undef variables are declared as global variables without initializer.
3139 // Never declare these variables.
3140 if (!undef)
3141 {
3142 auto &var = compiler.get<SPIRVariable>(phi.function_variable);
3143 auto &entry_block = compiler.get<SPIRBlock>(entry.entry_block);
3144 entry_block.declare_temporary.emplace_back(var.basetype, phi.local_variable);
3145 compiler.hoisted_temporaries.insert(phi.local_variable);
3146 compiler.forced_temporaries.insert(phi.local_variable);
3147 }
Hans-Kristian Arntzen0fd02812017-11-21 18:19:51 +01003148 }
3149 }
3150 }
Hans-Kristian Arntzencdca1922017-11-21 16:52:01 +01003151
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003152 switch (block.terminator)
3153 {
3154 case SPIRBlock::Direct:
3155 test_phi(block.next_block);
3156 break;
3157
3158 case SPIRBlock::Select:
3159 test_phi(block.true_block);
3160 test_phi(block.false_block);
3161 break;
3162
3163 case SPIRBlock::MultiSelect:
3164 for (auto &target : block.cases)
3165 test_phi(target.block);
3166 if (block.default_block)
3167 test_phi(block.default_block);
3168 break;
3169
3170 default:
3171 break;
3172 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003173 }
3174
Hans-Kristian Arntzen5b057f12017-11-21 10:11:40 +01003175 bool id_is_phi_variable(uint32_t id)
3176 {
3177 if (id >= compiler.get_current_id_bound())
3178 return false;
Hans-Kristian Arntzencdca1922017-11-21 16:52:01 +01003179 auto *var = compiler.maybe_get<SPIRVariable>(id);
Hans-Kristian Arntzen5b057f12017-11-21 10:11:40 +01003180 return var && var->phi_variable;
3181 }
3182
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003183 bool handle(spv::Op op, const uint32_t *args, uint32_t length)
3184 {
3185 switch (op)
3186 {
3187 case OpStore:
3188 {
3189 if (length < 2)
3190 return false;
3191
3192 uint32_t ptr = args[0];
3193 auto *var = compiler.maybe_get_backing_variable(ptr);
3194 if (var && var->storage == StorageClassFunction)
3195 accessed_variables_to_block[var->self].insert(current_block->self);
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02003196
3197 // If we store through an access chain, we have a partial write.
3198 if (var && var->self == ptr && var->storage == StorageClassFunction)
3199 complete_write_variables_to_block[var->self].insert(current_block->self);
Hans-Kristian Arntzen7e02f7f2017-11-22 11:04:29 +01003200
3201 // Might try to store a Phi variable here.
3202 if (id_is_phi_variable(args[1]))
3203 accessed_variables_to_block[args[1]].insert(current_block->self);
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003204 break;
3205 }
3206
3207 case OpAccessChain:
3208 case OpInBoundsAccessChain:
3209 {
3210 if (length < 3)
3211 return false;
3212
3213 uint32_t ptr = args[2];
3214 auto *var = compiler.maybe_get<SPIRVariable>(ptr);
3215 if (var && var->storage == StorageClassFunction)
3216 accessed_variables_to_block[var->self].insert(current_block->self);
3217 break;
3218 }
3219
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +01003220 case OpCopyMemory:
3221 {
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003222 if (length < 2)
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +01003223 return false;
3224
3225 uint32_t lhs = args[0];
3226 uint32_t rhs = args[1];
3227 auto *var = compiler.maybe_get_backing_variable(lhs);
3228 if (var && var->storage == StorageClassFunction)
3229 accessed_variables_to_block[var->self].insert(current_block->self);
3230
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02003231 // If we store through an access chain, we have a partial write.
3232 if (var->self == lhs)
3233 complete_write_variables_to_block[var->self].insert(current_block->self);
3234
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +01003235 var = compiler.maybe_get_backing_variable(rhs);
3236 if (var && var->storage == StorageClassFunction)
3237 accessed_variables_to_block[var->self].insert(current_block->self);
3238 break;
3239 }
3240
3241 case OpCopyObject:
3242 {
3243 if (length < 3)
3244 return false;
3245
3246 auto *var = compiler.maybe_get_backing_variable(args[2]);
3247 if (var && var->storage == StorageClassFunction)
3248 accessed_variables_to_block[var->self].insert(current_block->self);
Hans-Kristian Arntzen7e02f7f2017-11-22 11:04:29 +01003249
3250 // Might try to copy a Phi variable here.
3251 if (id_is_phi_variable(args[2]))
3252 accessed_variables_to_block[args[2]].insert(current_block->self);
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +01003253 break;
3254 }
3255
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003256 case OpLoad:
3257 {
3258 if (length < 3)
3259 return false;
3260 uint32_t ptr = args[2];
3261 auto *var = compiler.maybe_get_backing_variable(ptr);
3262 if (var && var->storage == StorageClassFunction)
3263 accessed_variables_to_block[var->self].insert(current_block->self);
3264 break;
3265 }
3266
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +01003267 case OpFunctionCall:
3268 {
3269 if (length < 3)
3270 return false;
3271
3272 length -= 3;
3273 args += 3;
3274 for (uint32_t i = 0; i < length; i++)
3275 {
3276 auto *var = compiler.maybe_get_backing_variable(args[i]);
3277 if (var && var->storage == StorageClassFunction)
3278 accessed_variables_to_block[var->self].insert(current_block->self);
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02003279
3280 // Cannot easily prove if argument we pass to a function is completely written.
3281 // Usually, functions write to a dummy variable,
3282 // which is then copied to in full to the real argument.
Hans-Kristian Arntzen7e02f7f2017-11-22 11:04:29 +01003283
3284 // Might try to copy a Phi variable here.
3285 if (id_is_phi_variable(args[i]))
3286 accessed_variables_to_block[args[i]].insert(current_block->self);
Hans-Kristian Arntzen0c9683c2016-11-18 09:59:54 +01003287 }
3288 break;
3289 }
3290
Hans-Kristian Arntzen5b057f12017-11-21 10:11:40 +01003291 case OpExtInst:
3292 {
3293 for (uint32_t i = 4; i < length; i++)
3294 if (id_is_phi_variable(args[i]))
3295 accessed_variables_to_block[args[i]].insert(current_block->self);
3296 break;
3297 }
3298
3299 case OpArrayLength:
3300 // Uses literals, but cannot be a phi variable, so ignore.
3301 break;
3302
Hans-Kristian Arntzence18d4c2017-11-17 13:38:29 +01003303 // Atomics shouldn't be able to access function-local variables.
3304 // Some GLSL builtins access a pointer.
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003305
3306 default:
Hans-Kristian Arntzen5b057f12017-11-21 10:11:40 +01003307 {
3308 // Rather dirty way of figuring out where Phi variables are used.
3309 // As long as only IDs are used, we can scan through instructions and try to find any evidence that
3310 // the ID of a variable has been used.
3311 // There are potential false positives here where a literal is used in-place of an ID,
3312 // but worst case, it does not affect the correctness of the compile.
3313 // Exhaustive analysis would be better here, but it's not worth it for now.
3314 for (uint32_t i = 0; i < length; i++)
3315 if (id_is_phi_variable(args[i]))
3316 accessed_variables_to_block[args[i]].insert(current_block->self);
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003317 break;
3318 }
Hans-Kristian Arntzen5b057f12017-11-21 10:11:40 +01003319 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003320 return true;
3321 }
3322
3323 Compiler &compiler;
Hans-Kristian Arntzen07584282017-11-21 18:51:11 +01003324 SPIRFunction &entry;
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003325 std::unordered_map<uint32_t, std::unordered_set<uint32_t>> accessed_variables_to_block;
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02003326 std::unordered_map<uint32_t, std::unordered_set<uint32_t>> complete_write_variables_to_block;
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003327 const SPIRBlock *current_block = nullptr;
Hans-Kristian Arntzen07584282017-11-21 18:51:11 +01003328 } handler(*this, entry);
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003329
3330 // First, we map out all variable access within a function.
3331 // Essentially a map of block -> { variables accessed in the basic block }
Hans-Kristian Arntzen7630d3c2016-11-21 12:14:02 +01003332 this->traverse_all_reachable_opcodes(entry, handler);
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003333
3334 // Compute the control flow graph for this function.
3335 CFG cfg(*this, entry);
3336
Hans-Kristian Arntzenb2c2e642017-03-25 16:25:30 +01003337 // Analyze if there are parameters which need to be implicitly preserved with an "in" qualifier.
Hans-Kristian Arntzen744d0402017-08-09 17:05:51 +02003338 analyze_parameter_preservation(entry, cfg, handler.accessed_variables_to_block,
3339 handler.complete_write_variables_to_block);
Hans-Kristian Arntzenb2c2e642017-03-25 16:25:30 +01003340
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003341 unordered_map<uint32_t, uint32_t> potential_loop_variables;
3342
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003343 // For each variable which is statically accessed.
3344 for (auto &var : handler.accessed_variables_to_block)
3345 {
3346 DominatorBuilder builder(cfg);
3347 auto &blocks = var.second;
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003348 auto &type = this->expression_type(var.first);
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003349
3350 // Figure out which block is dominating all accesses of those variables.
3351 for (auto &block : blocks)
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003352 {
Hans-Kristian Arntzena714d422016-12-16 12:43:12 +01003353 // If we're accessing a variable inside a continue block, this variable might be a loop variable.
3354 // We can only use loop variables with scalars, as we cannot track static expressions for vectors.
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003355 if (this->is_continue(block) && type.vecsize == 1 && type.columns == 1)
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003356 {
3357 // The variable is used in multiple continue blocks, this is not a loop
3358 // candidate, signal that by setting block to -1u.
3359 auto &potential = potential_loop_variables[var.first];
3360
3361 if (potential == 0)
3362 potential = block;
3363 else
Graham Wihlidalfadc1f92017-01-05 20:14:53 +01003364 potential = ~(0u);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003365 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003366 builder.add_block(block);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003367 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003368
Hans-Kristian Arntzen5ff11cc2016-11-18 16:45:11 +01003369 builder.lift_continue_block_dominator();
3370
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003371 // Add it to a per-block list of variables.
3372 uint32_t dominating_block = builder.get_dominator();
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003373 // If all blocks here are dead code, this will be 0, so the variable in question
3374 // will be completely eliminated.
3375 if (dominating_block)
3376 {
Hans-Kristian Arntzen7630d3c2016-11-21 12:14:02 +01003377 auto &block = this->get<SPIRBlock>(dominating_block);
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003378 block.dominated_variables.push_back(var.first);
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003379 this->get<SPIRVariable>(var.first).dominator = dominating_block;
Hans-Kristian Arntzenedbe8672016-11-17 22:15:07 +01003380 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003381 }
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003382
Hans-Kristian Arntzenf0c50a62017-07-25 18:22:15 +02003383 unordered_set<uint32_t> seen_blocks;
3384
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003385 // Now, try to analyze whether or not these variables are actually loop variables.
3386 for (auto &loop_variable : potential_loop_variables)
3387 {
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003388 auto &var = this->get<SPIRVariable>(loop_variable.first);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003389 auto dominator = var.dominator;
3390 auto block = loop_variable.second;
3391
3392 // The variable was accessed in multiple continue blocks, ignore.
Graham Wihlidalfadc1f92017-01-05 20:14:53 +01003393 if (block == ~(0u) || block == 0)
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003394 continue;
3395
3396 // Dead code.
3397 if (dominator == 0)
3398 continue;
3399
3400 uint32_t header = 0;
3401
3402 // Find the loop header for this block.
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003403 for (auto b : this->loop_blocks)
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003404 {
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003405 auto &potential_header = this->get<SPIRBlock>(b);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003406 if (potential_header.continue_block == block)
3407 {
3408 header = b;
3409 break;
3410 }
3411 }
3412
3413 assert(header);
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003414 auto &header_block = this->get<SPIRBlock>(header);
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003415
3416 // Now, there are two conditions we need to meet for the variable to be a loop variable.
3417 // 1. The dominating block must have a branch-free path to the loop header,
3418 // this way we statically know which expression should be part of the loop variable initializer.
3419
3420 // Walk from the dominator, if there is one straight edge connecting
3421 // dominator and loop header, we statically know the loop initializer.
3422 bool static_loop_init = true;
3423 while (dominator != header)
3424 {
3425 auto &succ = cfg.get_succeeding_edges(dominator);
3426 if (succ.size() != 1)
3427 {
3428 static_loop_init = false;
3429 break;
3430 }
3431
3432 auto &pred = cfg.get_preceding_edges(succ.front());
3433 if (pred.size() != 1 || pred.front() != dominator)
3434 {
3435 static_loop_init = false;
3436 break;
3437 }
3438
3439 dominator = succ.front();
3440 }
3441
3442 if (!static_loop_init)
3443 continue;
3444
3445 // The second condition we need to meet is that no access after the loop
3446 // merge can occur. Walk the CFG to see if we find anything.
3447 auto &blocks = handler.accessed_variables_to_block[loop_variable.first];
Hans-Kristian Arntzenf0c50a62017-07-25 18:22:15 +02003448
3449 seen_blocks.clear();
3450 cfg.walk_from(seen_blocks, header_block.merge_block, [&](uint32_t walk_block) {
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003451 // We found a block which accesses the variable outside the loop.
3452 if (blocks.find(walk_block) != end(blocks))
3453 static_loop_init = false;
3454 });
3455
3456 if (!static_loop_init)
3457 continue;
3458
3459 // We have a loop variable.
3460 header_block.loop_variables.push_back(loop_variable.first);
Hans-Kristian Arntzen44b32162016-12-16 14:01:09 +01003461 // Need to sort here as variables come from an unordered container, and pushing stuff in wrong order
3462 // will break reproducability in regression runs.
3463 sort(begin(header_block.loop_variables), end(header_block.loop_variables));
Hans-Kristian Arntzence3fe292017-01-12 10:57:44 +01003464 this->get<SPIRVariable>(loop_variable.first).loop_variable = true;
Hans-Kristian Arntzen4f07a322016-12-15 17:14:47 +01003465 }
Hans-Kristian Arntzendad4a342016-11-11 18:04:14 +01003466}
Hans-Kristian Arntzen016b1d82017-01-21 10:07:38 +01003467
3468uint64_t Compiler::get_buffer_block_flags(const SPIRVariable &var)
3469{
3470 auto &type = get<SPIRType>(var.basetype);
Hans-Kristian Arntzen8a80e622017-01-22 08:51:24 +01003471 assert(type.basetype == SPIRType::Struct);
Hans-Kristian Arntzen016b1d82017-01-21 10:07:38 +01003472
3473 // Some flags like non-writable, non-readable are actually found
3474 // as member decorations. If all members have a decoration set, propagate
3475 // the decoration up as a regular variable decoration.
3476 uint64_t base_flags = meta[var.self].decoration.decoration_flags;
Hans-Kristian Arntzen016b1d82017-01-21 10:07:38 +01003477
Hans-Kristian Arntzen8a80e622017-01-22 08:51:24 +01003478 if (type.member_types.empty())
3479 return base_flags;
3480
3481 uint64_t all_members_flag_mask = ~(0ull);
3482 for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++)
3483 all_members_flag_mask &= get_member_decoration_mask(type.self, i);
3484
3485 return base_flags | all_members_flag_mask;
Hans-Kristian Arntzen016b1d82017-01-21 10:07:38 +01003486}
Hans-Kristian Arntzen95409792017-01-21 12:29:20 +01003487
3488bool Compiler::get_common_basic_type(const SPIRType &type, SPIRType::BaseType &base_type)
3489{
3490 if (type.basetype == SPIRType::Struct)
3491 {
3492 base_type = SPIRType::Unknown;
3493 for (auto &member_type : type.member_types)
3494 {
3495 SPIRType::BaseType member_base;
3496 if (!get_common_basic_type(get<SPIRType>(member_type), member_base))
3497 return false;
3498
3499 if (base_type == SPIRType::Unknown)
3500 base_type = member_base;
3501 else if (base_type != member_base)
3502 return false;
3503 }
3504 return true;
3505 }
3506 else
3507 {
3508 base_type = type.basetype;
3509 return true;
3510 }
3511}
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003512
3513bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args, uint32_t length)
3514{
3515 const auto add_if_builtin = [&](uint32_t id) {
3516 // Only handles variables here.
3517 // Builtins which are part of a block are handled in AccessChain.
3518 auto *var = compiler.maybe_get<SPIRVariable>(id);
3519 if (var && compiler.meta[id].decoration.builtin)
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003520 {
3521 auto &type = compiler.get<SPIRType>(var->basetype);
3522 auto &flags =
Hans-Kristian Arntzen2ebe1a82017-03-06 16:50:46 +01003523 type.storage == StorageClassInput ? compiler.active_input_builtins : compiler.active_output_builtins;
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003524 flags |= 1ull << compiler.meta[id].decoration.builtin_type;
3525 }
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003526 };
3527
3528 switch (opcode)
3529 {
3530 case OpStore:
3531 if (length < 1)
3532 return false;
3533
3534 add_if_builtin(args[0]);
3535 break;
3536
3537 case OpCopyMemory:
3538 if (length < 2)
3539 return false;
3540
3541 add_if_builtin(args[0]);
3542 add_if_builtin(args[1]);
3543 break;
3544
3545 case OpCopyObject:
3546 case OpLoad:
3547 if (length < 3)
3548 return false;
3549
3550 add_if_builtin(args[2]);
3551 break;
3552
3553 case OpFunctionCall:
3554 {
3555 if (length < 3)
3556 return false;
3557
3558 uint32_t count = length - 3;
3559 args += 3;
3560 for (uint32_t i = 0; i < count; i++)
3561 add_if_builtin(args[i]);
3562 break;
3563 }
3564
3565 case OpAccessChain:
3566 case OpInBoundsAccessChain:
3567 {
3568 if (length < 4)
3569 return false;
3570
3571 // Only consider global variables, cannot consider variables in functions yet, or other
3572 // access chains as they have not been created yet.
3573 auto *var = compiler.maybe_get<SPIRVariable>(args[2]);
3574 if (!var)
3575 break;
3576
Hans-Kristian Arntzen945425e2017-08-15 10:23:04 +02003577 // Required if we access chain into builtins like gl_GlobalInvocationID.
3578 add_if_builtin(args[2]);
3579
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003580 auto *type = &compiler.get<SPIRType>(var->basetype);
3581
3582 // Start traversing type hierarchy at the proper non-pointer types.
3583 while (type->pointer)
3584 {
3585 assert(type->parent_type);
3586 type = &compiler.get<SPIRType>(type->parent_type);
3587 }
3588
Hans-Kristian Arntzen2ebe1a82017-03-06 16:50:46 +01003589 auto &flags =
3590 type->storage == StorageClassInput ? compiler.active_input_builtins : compiler.active_output_builtins;
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003591
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003592 uint32_t count = length - 3;
3593 args += 3;
3594 for (uint32_t i = 0; i < count; i++)
3595 {
3596 // Arrays
3597 if (!type->array.empty())
3598 {
3599 type = &compiler.get<SPIRType>(type->parent_type);
3600 }
3601 // Structs
3602 else if (type->basetype == SPIRType::Struct)
3603 {
3604 uint32_t index = compiler.get<SPIRConstant>(args[i]).scalar();
3605
3606 if (index < uint32_t(compiler.meta[type->self].members.size()))
3607 {
3608 auto &decorations = compiler.meta[type->self].members[index];
3609 if (decorations.builtin)
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003610 flags |= 1ull << decorations.builtin_type;
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003611 }
3612
3613 type = &compiler.get<SPIRType>(type->member_types[index]);
3614 }
3615 else
3616 {
3617 // No point in traversing further. We won't find any extra builtins.
3618 break;
3619 }
3620 }
3621 break;
3622 }
3623
3624 default:
3625 break;
3626 }
3627
3628 return true;
3629}
3630
3631void Compiler::update_active_builtins()
3632{
Hans-Kristian Arntzenbdea1a42017-03-06 16:50:15 +01003633 active_input_builtins = 0;
3634 active_output_builtins = 0;
Hans-Kristian Arntzen099f3072017-03-06 15:21:00 +01003635 ActiveBuiltinHandler handler(*this);
3636 traverse_all_reachable_opcodes(get<SPIRFunction>(entry_point), handler);
3637}
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003638
Bill Hollings192bdc92017-05-24 09:31:38 -04003639// Returns whether this shader uses a builtin of the storage class
3640bool Compiler::has_active_builtin(BuiltIn builtin, StorageClass storage)
3641{
3642 uint64_t flags;
3643 switch (storage)
3644 {
3645 case StorageClassInput:
3646 flags = active_input_builtins;
3647 break;
3648 case StorageClassOutput:
3649 flags = active_output_builtins;
3650 break;
3651
3652 default:
3653 return false;
3654 }
3655 return flags & (1ull << builtin);
3656}
3657
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003658void Compiler::analyze_sampler_comparison_states()
3659{
3660 CombinedImageSamplerUsageHandler handler(*this);
3661 traverse_all_reachable_opcodes(get<SPIRFunction>(entry_point), handler);
3662 comparison_samplers = move(handler.comparison_samplers);
3663}
3664
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003665bool Compiler::CombinedImageSamplerUsageHandler::begin_function_scope(const uint32_t *args, uint32_t length)
3666{
3667 if (length < 3)
3668 return false;
3669
3670 auto &func = compiler.get<SPIRFunction>(args[2]);
3671 const auto *arg = &args[3];
3672 length -= 3;
3673
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003674 for (uint32_t i = 0; i < length; i++)
3675 {
3676 auto &argument = func.arguments[i];
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003677 dependency_hierarchy[argument.id].insert(arg[i]);
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003678 }
3679
3680 return true;
3681}
3682
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003683void Compiler::CombinedImageSamplerUsageHandler::add_hierarchy_to_comparison_samplers(uint32_t sampler)
3684{
3685 // Traverse the variable dependency hierarchy and tag everything in its path with comparison samplers.
3686 comparison_samplers.insert(sampler);
3687 for (auto &samp : dependency_hierarchy[sampler])
3688 add_hierarchy_to_comparison_samplers(samp);
3689}
3690
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003691bool Compiler::CombinedImageSamplerUsageHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
3692{
3693 switch (opcode)
3694 {
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003695 case OpAccessChain:
3696 case OpInBoundsAccessChain:
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003697 case OpLoad:
3698 {
3699 if (length < 3)
3700 return false;
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003701 dependency_hierarchy[args[1]].insert(args[2]);
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003702 break;
3703 }
3704
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003705 case OpSampledImage:
3706 {
3707 if (length < 4)
3708 return false;
3709
3710 uint32_t result_type = args[0];
3711 auto &type = compiler.get<SPIRType>(result_type);
3712 if (type.image.depth)
3713 {
3714 // This sampler must be a SamplerComparisionState, and not a regular SamplerState.
3715 uint32_t sampler = args[3];
Hans-Kristian Arntzen07ee7d02017-05-06 13:53:06 +02003716 add_hierarchy_to_comparison_samplers(sampler);
Hans-Kristian Arntzenf4d72682017-05-06 13:21:35 +02003717 }
3718 return true;
3719 }
3720
3721 default:
3722 break;
3723 }
3724
3725 return true;
3726}
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02003727
3728bool Compiler::buffer_is_hlsl_counter_buffer(uint32_t id) const
3729{
3730 if (meta.at(id).hlsl_magic_counter_buffer_candidate)
3731 {
3732 auto *var = maybe_get<SPIRVariable>(id);
Hans-Kristian Arntzen9c9e2672017-05-09 09:38:33 +02003733 // Ensure that this is actually a buffer object.
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +02003734 return var && (var->storage == StorageClassStorageBuffer ||
3735 has_decoration(get<SPIRType>(var->basetype).self, DecorationBufferBlock));
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02003736 }
3737 else
3738 return false;
3739}
3740
3741bool Compiler::buffer_get_hlsl_counter_buffer(uint32_t id, uint32_t &counter_id) const
3742{
3743 auto &name = get_name(id);
3744 uint32_t id_bound = get_current_id_bound();
3745 for (uint32_t i = 0; i < id_bound; i++)
3746 {
3747 if (meta[i].hlsl_magic_counter_buffer_candidate && meta[i].hlsl_magic_counter_buffer_name == name)
3748 {
3749 auto *var = maybe_get<SPIRVariable>(i);
Hans-Kristian Arntzen9c9e2672017-05-09 09:38:33 +02003750 // Ensure that this is actually a buffer object.
Hans-Kristian Arntzen153fed02017-09-28 13:28:44 +02003751 if (var && (var->storage == StorageClassStorageBuffer ||
3752 has_decoration(get<SPIRType>(var->basetype).self, DecorationBufferBlock)))
Hans-Kristian Arntzenec45c9e2017-04-19 17:33:14 +02003753 {
3754 counter_id = i;
3755 return true;
3756 }
3757 }
3758 }
3759 return false;
3760}
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003761
3762void Compiler::make_constant_null(uint32_t id, uint32_t type)
3763{
3764 auto &constant_type = get<SPIRType>(type);
3765
3766 if (!constant_type.array.empty())
3767 {
3768 assert(constant_type.parent_type);
3769 uint32_t parent_id = increase_bound_by(1);
3770 make_constant_null(parent_id, constant_type.parent_type);
3771
3772 if (!constant_type.array_size_literal.back())
3773 SPIRV_CROSS_THROW("Array size of OpConstantNull must be a literal.");
3774
3775 vector<uint32_t> elements(constant_type.array.back());
3776 for (uint32_t i = 0; i < constant_type.array.back(); i++)
3777 elements[i] = parent_id;
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02003778 set<SPIRConstant>(id, type, elements.data(), uint32_t(elements.size()), false);
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003779 }
3780 else if (!constant_type.member_types.empty())
3781 {
Bill Hollings77f58122017-08-11 14:54:58 -04003782 uint32_t member_ids = increase_bound_by(uint32_t(constant_type.member_types.size()));
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003783 vector<uint32_t> elements(constant_type.member_types.size());
3784 for (uint32_t i = 0; i < constant_type.member_types.size(); i++)
3785 {
3786 make_constant_null(member_ids + i, constant_type.member_types[i]);
3787 elements[i] = member_ids + i;
3788 }
Hans-Kristian Arntzen5e1d6fb2017-09-27 15:16:33 +02003789 set<SPIRConstant>(id, type, elements.data(), uint32_t(elements.size()), false);
Hans-Kristian Arntzen48ccde32017-08-03 14:32:07 +02003790 }
3791 else
3792 {
3793 auto &constant = set<SPIRConstant>(id, type);
3794 constant.make_null(constant_type);
3795 }
3796}
Hans-Kristian Arntzen8d7a9092017-08-15 15:27:53 +02003797
3798const std::vector<spv::Capability> &Compiler::get_declared_capabilities() const
3799{
3800 return declared_capabilities;
3801}
3802
3803const std::vector<std::string> &Compiler::get_declared_extensions() const
3804{
3805 return declared_extensions;
3806}
Hans-Kristian Arntzen2c90ea32017-12-01 14:20:51 +01003807
3808std::string Compiler::get_remapped_declared_block_name(uint32_t id) const
3809{
3810 auto itr = declared_block_names.find(id);
3811 if (itr != end(declared_block_names))
3812 return itr->second;
3813 else
3814 {
3815 auto &var = get<SPIRVariable>(id);
3816 auto &type = get<SPIRType>(var.basetype);
3817 auto &block_name = meta[type.self].decoration.alias;
3818 return block_name.empty() ? get_block_fallback_name(id) : block_name;
3819 }
3820}