blob: 30a189c61635945734e66ae4b5fb4b2710e1df38 [file] [log] [blame]
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001/*
2 * Copyright 2015-2016 ARM Limited
3 *
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 Arntzen75471fb2016-03-02 18:09:16 +010019#include <algorithm>
Hans-Kristian Arntzen5ea59bd2016-05-23 13:30:02 +020020#include <cstring>
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010021#include <utility>
22
23using namespace std;
24using namespace spv;
Hans-Kristian Arntzen147e53a2016-04-04 09:36:04 +020025using namespace spirv_cross;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010026
27#define log(...) fprintf(stderr, __VA_ARGS__)
28
29Instruction::Instruction(const vector<uint32_t> &spirv, uint32_t &index)
30{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020031 op = spirv[index] & 0xffff;
32 count = (spirv[index] >> 16) & 0xffff;
Hans-Kristian Arntzen416566b2016-07-08 10:47:03 +020033
34 if (count == 0)
35 throw CompilerError("SPIR-V instructions cannot consume 0 words. Invalid SPIR-V file.");
36
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020037 offset = index + 1;
38 length = count - 1;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010039
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020040 index += count;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010041
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020042 if (index > spirv.size())
43 throw CompilerError("SPIR-V instruction goes out of bounds.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010044}
45
46Compiler::Compiler(vector<uint32_t> ir)
47 : spirv(move(ir))
48{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020049 parse();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010050}
51
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020052string Compiler::compile()
53{
54 return "";
55}
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010056
57bool Compiler::variable_storage_is_aliased(const SPIRVariable &v)
58{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020059 auto &type = get<SPIRType>(v.basetype);
60 bool ssbo = (meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)) != 0;
61 bool image = type.basetype == SPIRType::Image;
62 bool counter = type.basetype == SPIRType::AtomicCounter;
Hans-Kristian Arntzen7d8add32016-07-12 15:00:10 +020063 bool is_restrict = (meta[v.self].decoration.decoration_flags & (1ull << DecorationRestrict)) != 0;
64 return !is_restrict && (ssbo || image || counter);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010065}
66
67bool Compiler::block_is_pure(const SPIRBlock &block)
68{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020069 for (auto &i : block.ops)
70 {
71 auto ops = stream(i);
72 auto op = static_cast<Op>(i.op);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010073
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020074 switch (op)
75 {
76 case OpFunctionCall:
77 {
78 uint32_t func = ops[2];
79 if (!function_is_pure(get<SPIRFunction>(func)))
80 return false;
81 break;
82 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010083
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020084 case OpStore:
85 {
86 auto &type = expression_type(ops[0]);
87 if (type.storage != StorageClassFunction)
88 return false;
89 break;
90 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010091
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +020092 case OpImageWrite:
93 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +010094
Hans-Kristian Arntzen5af1a512016-05-05 09:51:42 +020095 // Atomics are impure.
96 case OpAtomicLoad:
97 case OpAtomicStore:
98 case OpAtomicExchange:
99 case OpAtomicCompareExchange:
100 case OpAtomicIIncrement:
101 case OpAtomicIDecrement:
102 case OpAtomicIAdd:
103 case OpAtomicISub:
104 case OpAtomicSMin:
105 case OpAtomicUMin:
106 case OpAtomicSMax:
107 case OpAtomicUMax:
108 case OpAtomicAnd:
109 case OpAtomicOr:
110 case OpAtomicXor:
111 return false;
112
113 // Geometry shader builtins modify global state.
114 case OpEndPrimitive:
115 case OpEmitStreamVertex:
116 case OpEndStreamPrimitive:
117 case OpEmitVertex:
118 return false;
119
120 // Barriers disallow any reordering, so we should treat blocks with barrier as writing.
121 case OpControlBarrier:
122 case OpMemoryBarrier:
123 return false;
124
125 // OpExtInst is potentially impure depending on extension, but GLSL builtins are at least pure.
126
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200127 default:
128 break;
129 }
130 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100131
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200132 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100133}
134
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +0200135string Compiler::to_name(uint32_t id, bool allow_alias)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100136{
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +0200137 if (allow_alias && ids.at(id).get_type() == TypeType)
138 {
139 // If this type is a simple alias, emit the
140 // name of the original type instead.
141 // We don't want to override the meta alias
142 // as that can be overridden by the reflection APIs after parse.
143 auto &type = get<SPIRType>(id);
144 if (type.type_alias)
145 return to_name(type.type_alias);
146 }
147
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200148 if (meta[id].decoration.alias.empty())
149 return join("_", id);
150 else
151 return meta.at(id).decoration.alias;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100152}
153
154bool Compiler::function_is_pure(const SPIRFunction &func)
155{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200156 for (auto block : func.blocks)
157 {
158 if (!block_is_pure(get<SPIRBlock>(block)))
159 {
160 //fprintf(stderr, "Function %s is impure!\n", to_name(func.self).c_str());
161 return false;
162 }
163 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100164
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200165 //fprintf(stderr, "Function %s is pure!\n", to_name(func.self).c_str());
166 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100167}
168
169void Compiler::register_global_read_dependencies(const SPIRBlock &block, uint32_t id)
170{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200171 for (auto &i : block.ops)
172 {
173 auto ops = stream(i);
174 auto op = static_cast<Op>(i.op);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100175
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200176 switch (op)
177 {
178 case OpFunctionCall:
179 {
180 uint32_t func = ops[2];
181 register_global_read_dependencies(get<SPIRFunction>(func), id);
182 break;
183 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100184
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200185 case OpLoad:
186 case OpImageRead:
187 {
188 // If we're in a storage class which does not get invalidated, adding dependencies here is no big deal.
189 auto *var = maybe_get_backing_variable(ops[2]);
190 if (var && var->storage != StorageClassFunction)
191 {
192 auto &type = get<SPIRType>(var->basetype);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100193
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200194 // InputTargets are immutable.
195 if (type.basetype != SPIRType::Image && type.image.dim != DimSubpassData)
196 var->dependees.push_back(id);
197 }
198 break;
199 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100200
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200201 default:
202 break;
203 }
204 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100205}
206
207void Compiler::register_global_read_dependencies(const SPIRFunction &func, uint32_t id)
208{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200209 for (auto block : func.blocks)
210 register_global_read_dependencies(get<SPIRBlock>(block), id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100211}
212
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200213SPIRVariable *Compiler::maybe_get_backing_variable(uint32_t chain)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100214{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200215 auto *var = maybe_get<SPIRVariable>(chain);
216 if (!var)
217 {
218 auto *cexpr = maybe_get<SPIRExpression>(chain);
219 if (cexpr)
220 var = maybe_get<SPIRVariable>(cexpr->loaded_from);
221 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100222
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200223 return var;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100224}
225
226void Compiler::register_read(uint32_t expr, uint32_t chain, bool forwarded)
227{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200228 auto &e = get<SPIRExpression>(expr);
229 auto *var = maybe_get_backing_variable(chain);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100230
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200231 if (var)
232 {
233 e.loaded_from = var->self;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100234
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200235 // If the backing variable is immutable, we do not need to depend on the variable.
236 if (forwarded && !is_immutable(var->self))
237 var->dependees.push_back(e.self);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100238
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200239 // If we load from a parameter, make sure we create "inout" if we also write to the parameter.
240 // The default is "in" however, so we never invalidate our compilation by reading.
241 if (var && var->parameter)
242 var->parameter->read_count++;
243 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100244}
245
246void Compiler::register_write(uint32_t chain)
247{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200248 auto *var = maybe_get<SPIRVariable>(chain);
249 if (!var)
250 {
251 // If we're storing through an access chain, invalidate the backing variable instead.
252 auto *expr = maybe_get<SPIRExpression>(chain);
253 if (expr && expr->loaded_from)
254 var = maybe_get<SPIRVariable>(expr->loaded_from);
255 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100256
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200257 if (var)
258 {
259 // If our variable is in a storage class which can alias with other buffers,
260 // invalidate all variables which depend on aliased variables.
261 if (variable_storage_is_aliased(*var))
262 flush_all_aliased_variables();
263 else if (var)
264 flush_dependees(*var);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100265
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200266 // We tried to write to a parameter which is not marked with out qualifier, force a recompile.
267 if (var->parameter && var->parameter->write_count == 0)
268 {
269 var->parameter->write_count++;
270 force_recompile = true;
271 }
272 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100273}
274
275void Compiler::flush_dependees(SPIRVariable &var)
276{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200277 for (auto expr : var.dependees)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200278 invalid_expressions.insert(expr);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200279 var.dependees.clear();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100280}
281
282void Compiler::flush_all_aliased_variables()
283{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200284 for (auto aliased : aliased_variables)
285 flush_dependees(get<SPIRVariable>(aliased));
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100286}
287
288void Compiler::flush_all_atomic_capable_variables()
289{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200290 for (auto global : global_variables)
291 flush_dependees(get<SPIRVariable>(global));
292 flush_all_aliased_variables();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100293}
294
295void Compiler::flush_all_active_variables()
296{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200297 // Invalidate all temporaries we read from variables in this block since they were forwarded.
298 // Invalidate all temporaries we read from globals.
299 for (auto &v : current_function->local_variables)
300 flush_dependees(get<SPIRVariable>(v));
301 for (auto &arg : current_function->arguments)
302 flush_dependees(get<SPIRVariable>(arg.id));
303 for (auto global : global_variables)
304 flush_dependees(get<SPIRVariable>(global));
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100305
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200306 flush_all_aliased_variables();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100307}
308
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200309const SPIRType &Compiler::expression_type(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100310{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200311 switch (ids[id].get_type())
312 {
313 case TypeVariable:
314 return get<SPIRType>(get<SPIRVariable>(id).basetype);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100315
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200316 case TypeExpression:
317 return get<SPIRType>(get<SPIRExpression>(id).expression_type);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100318
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200319 case TypeConstant:
320 return get<SPIRType>(get<SPIRConstant>(id).constant_type);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100321
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200322 case TypeUndef:
323 return get<SPIRType>(get<SPIRUndef>(id).basetype);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100324
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200325 default:
326 throw CompilerError("Cannot resolve expression type.");
327 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100328}
329
330bool Compiler::expression_is_lvalue(uint32_t id) const
331{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200332 auto &type = expression_type(id);
333 switch (type.basetype)
334 {
335 case SPIRType::SampledImage:
336 case SPIRType::Image:
337 case SPIRType::Sampler:
338 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100339
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200340 default:
341 return true;
342 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100343}
344
345bool Compiler::is_immutable(uint32_t id) const
346{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200347 if (ids[id].get_type() == TypeVariable)
348 {
349 auto &var = get<SPIRVariable>(id);
Hans-Kristian Arntzen92134e42016-04-01 19:58:26 +0200350
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200351 // Anything we load from the UniformConstant address space is guaranteed to be immutable.
352 bool pointer_to_const = var.storage == StorageClassUniformConstant;
Hans-Kristian Arntzen36a0b632016-07-12 14:33:04 +0200353 return pointer_to_const || var.phi_variable || !expression_is_lvalue(id);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200354 }
355 else if (ids[id].get_type() == TypeExpression)
356 return get<SPIRExpression>(id).immutable;
357 else if (ids[id].get_type() == TypeConstant || ids[id].get_type() == TypeUndef)
358 return true;
359 else
360 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100361}
362
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200363static inline bool storage_class_is_interface(spv::StorageClass storage)
364{
365 switch (storage)
366 {
367 case StorageClassInput:
368 case StorageClassOutput:
369 case StorageClassUniform:
370 case StorageClassUniformConstant:
371 case StorageClassAtomicCounter:
372 case StorageClassPushConstant:
373 return true;
374
375 default:
376 return false;
377 }
378}
379
380bool Compiler::is_hidden_variable(const SPIRVariable &var, bool include_builtins) const
381{
382 if ((is_builtin_variable(var) && !include_builtins) || var.remapped_variable)
383 return true;
384
Hans-Kristian Arntzen1b5ca8d2016-09-10 16:20:19 +0200385 // Combined image samplers are always considered active as they are "magic" variables.
386 if (find_if(begin(combined_image_samplers), end(combined_image_samplers), [&var](const CombinedImageSampler &samp) {
387 return samp.combined_id == var.self;
388 }) != end(combined_image_samplers))
389 {
390 return false;
391 }
392
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200393 bool hidden = false;
394 if (check_active_interface_variables && storage_class_is_interface(var.storage))
395 hidden = active_interface_variables.find(var.self) == end(active_interface_variables);
396 return hidden;
397}
398
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100399bool Compiler::is_builtin_variable(const SPIRVariable &var) const
400{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200401 if (var.compat_builtin || meta[var.self].decoration.builtin)
402 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100403
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200404 // We can have builtin structs as well. If one member of a struct is builtin, the struct must also be builtin.
405 for (auto &m : meta[get<SPIRType>(var.basetype).self].members)
406 if (m.builtin)
407 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100408
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200409 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100410}
411
412bool Compiler::is_member_builtin(const SPIRType &type, uint32_t index, BuiltIn *builtin) const
413{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200414 auto &memb = meta[type.self].members;
415 if (index < memb.size() && memb[index].builtin)
416 {
417 if (builtin)
418 *builtin = memb[index].builtin_type;
419 return true;
420 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100421
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200422 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100423}
424
Bill Hollings103aabf2016-04-06 17:42:27 -0400425bool Compiler::is_scalar(const SPIRType &type) const
426{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200427 return type.vecsize == 1 && type.columns == 1;
Bill Hollings103aabf2016-04-06 17:42:27 -0400428}
429
430bool Compiler::is_vector(const SPIRType &type) const
431{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200432 return type.vecsize > 1 && type.columns == 1;
Bill Hollings103aabf2016-04-06 17:42:27 -0400433}
434
435bool Compiler::is_matrix(const SPIRType &type) const
436{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200437 return type.vecsize > 1 && type.columns > 1;
Bill Hollings103aabf2016-04-06 17:42:27 -0400438}
439
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100440ShaderResources Compiler::get_shader_resources() const
441{
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200442 return get_shader_resources(nullptr);
443}
444
445ShaderResources Compiler::get_shader_resources(const unordered_set<uint32_t> &active_variables) const
446{
447 return get_shader_resources(&active_variables);
448}
449
450bool Compiler::InterfaceVariableAccessHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
451{
452 uint32_t variable = 0;
453 switch (opcode)
454 {
455 // Need this first, otherwise, GCC complains about unhandled switch statements.
456 default:
457 break;
458
459 case OpFunctionCall:
460 {
461 // Invalid SPIR-V.
462 if (length < 3)
463 return false;
464
465 uint32_t count = length - 3;
466 args += 3;
467 for (uint32_t i = 0; i < count; i++)
468 {
469 auto *var = compiler.maybe_get<SPIRVariable>(args[i]);
470 if (var && storage_class_is_interface(var->storage))
471 variables.insert(args[i]);
472 }
473 break;
474 }
475
476 case OpAtomicStore:
477 case OpStore:
478 // Invalid SPIR-V.
479 if (length < 1)
480 return false;
481 variable = args[0];
482 break;
483
484 case OpAccessChain:
485 case OpInBoundsAccessChain:
486 case OpLoad:
487 case OpImageTexelPointer:
488 case OpAtomicLoad:
489 case OpAtomicExchange:
490 case OpAtomicCompareExchange:
491 case OpAtomicIIncrement:
492 case OpAtomicIDecrement:
493 case OpAtomicIAdd:
494 case OpAtomicISub:
495 case OpAtomicSMin:
496 case OpAtomicUMin:
497 case OpAtomicSMax:
498 case OpAtomicUMax:
499 case OpAtomicAnd:
500 case OpAtomicOr:
501 case OpAtomicXor:
502 // Invalid SPIR-V.
503 if (length < 3)
504 return false;
505 variable = args[2];
506 break;
507 }
508
509 if (variable)
510 {
511 auto *var = compiler.maybe_get<SPIRVariable>(variable);
512 if (var && storage_class_is_interface(var->storage))
513 variables.insert(variable);
514 }
515 return true;
516}
517
518unordered_set<uint32_t> Compiler::get_active_interface_variables() const
519{
520 // Traverse the call graph and find all interface variables which are in use.
521 unordered_set<uint32_t> variables;
522 InterfaceVariableAccessHandler handler(*this, variables);
523 traverse_all_reachable_opcodes(get<SPIRFunction>(entry_point), handler);
524 return variables;
525}
526
527void Compiler::set_enabled_interface_variables(std::unordered_set<uint32_t> active_variables)
528{
529 active_interface_variables = move(active_variables);
530 check_active_interface_variables = true;
531}
532
533ShaderResources Compiler::get_shader_resources(const unordered_set<uint32_t> *active_variables) const
534{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200535 ShaderResources res;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100536
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200537 for (auto &id : ids)
538 {
539 if (id.get_type() != TypeVariable)
540 continue;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100541
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200542 auto &var = id.get<SPIRVariable>();
543 auto &type = get<SPIRType>(var.basetype);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100544
Hans-Kristian Arntzend5dc5f32016-07-05 13:21:26 +0200545 // It is possible for uniform storage classes to be passed as function parameters, so detect
546 // that. To detect function parameters, check of StorageClass of variable is function scope.
547 if (var.storage == StorageClassFunction || !type.pointer || is_builtin_variable(var))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200548 continue;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100549
Hans-Kristian Arntzenf61a5d12016-08-26 12:58:50 +0200550 if (active_variables && active_variables->find(var.self) == end(*active_variables))
551 continue;
552
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200553 // Input
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +0200554 if (var.storage == StorageClassInput && interface_variable_exists_in_entry_point(var.self))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200555 {
556 if (meta[type.self].decoration.decoration_flags & (1ull << DecorationBlock))
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200557 res.stage_inputs.push_back({ var.self, var.basetype, type.self, meta[type.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200558 else
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200559 res.stage_inputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200560 }
561 // Subpass inputs
562 else if (var.storage == StorageClassUniformConstant && type.image.dim == DimSubpassData)
563 {
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200564 res.subpass_inputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200565 }
566 // Outputs
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +0200567 else if (var.storage == StorageClassOutput && interface_variable_exists_in_entry_point(var.self))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200568 {
569 if (meta[type.self].decoration.decoration_flags & (1ull << DecorationBlock))
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200570 res.stage_outputs.push_back({ var.self, var.basetype, type.self, meta[type.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200571 else
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200572 res.stage_outputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200573 }
574 // UBOs
575 else if (type.storage == StorageClassUniform &&
576 (meta[type.self].decoration.decoration_flags & (1ull << DecorationBlock)))
577 {
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200578 res.uniform_buffers.push_back({ var.self, var.basetype, type.self, meta[type.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200579 }
580 // SSBOs
581 else if (type.storage == StorageClassUniform &&
582 (meta[type.self].decoration.decoration_flags & (1ull << DecorationBufferBlock)))
583 {
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200584 res.storage_buffers.push_back({ var.self, var.basetype, type.self, meta[type.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200585 }
586 // Push constant blocks
587 else if (type.storage == StorageClassPushConstant)
588 {
589 // There can only be one push constant block, but keep the vector in case this restriction is lifted
590 // in the future.
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200591 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 +0200592 }
593 // Images
Hans-Kristian Arntzene9202082016-09-10 13:05:35 +0200594 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Image &&
595 type.image.sampled == 2)
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200596 {
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200597 res.storage_images.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200598 }
Hans-Kristian Arntzene9202082016-09-10 13:05:35 +0200599 // Separate images
600 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Image &&
601 type.image.sampled == 1)
602 {
603 res.separate_images.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
604 }
605 // Separate samplers
606 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Sampler)
607 {
608 res.separate_samplers.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
609 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200610 // Textures
611 else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::SampledImage)
612 {
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200613 res.sampled_images.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200614 }
615 // Atomic counters
616 else if (type.storage == StorageClassAtomicCounter)
617 {
Hans-Kristian Arntzen5c24d992016-07-12 21:20:18 +0200618 res.atomic_counters.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias });
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200619 }
620 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100621
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200622 return res;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100623}
624
625static inline uint32_t swap_endian(uint32_t v)
626{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200627 return ((v >> 24) & 0x000000ffu) | ((v >> 8) & 0x0000ff00u) | ((v << 8) & 0x00ff0000u) | ((v << 24) & 0xff000000u);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100628}
629
630static string extract_string(const vector<uint32_t> &spirv, uint32_t offset)
631{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200632 string ret;
633 for (uint32_t i = offset; i < spirv.size(); i++)
634 {
635 uint32_t w = spirv[i];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100636
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200637 for (uint32_t j = 0; j < 4; j++, w >>= 8)
638 {
639 char c = w & 0xff;
640 if (c == '\0')
641 return ret;
642 ret += c;
643 }
644 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100645
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200646 throw CompilerError("String was not terminated before EOF");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100647}
648
Hans-Kristian Arntzen45ad58a2016-05-10 23:39:41 +0200649static bool is_valid_spirv_version(uint32_t version)
650{
651 switch (version)
652 {
653 // Allow v99 since it tends to just work.
654 case 99:
655 case 0x10000: // SPIR-V 1.0
656 case 0x10100: // SPIR-V 1.1
657 return true;
658
659 default:
660 return false;
661 }
662}
663
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100664void Compiler::parse()
665{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200666 auto len = spirv.size();
667 if (len < 5)
668 throw CompilerError("SPIRV file too small.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100669
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200670 auto s = spirv.data();
Hans-Kristian Arntzen5ac88272016-04-11 13:38:18 +0200671
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200672 // Endian-swap if we need to.
673 if (s[0] == swap_endian(MagicNumber))
Hans-Kristian Arntzen5ea59bd2016-05-23 13:30:02 +0200674 transform(begin(spirv), end(spirv), begin(spirv), [](uint32_t c) { return swap_endian(c); });
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100675
Hans-Kristian Arntzen45ad58a2016-05-10 23:39:41 +0200676 if (s[0] != MagicNumber || !is_valid_spirv_version(s[1]))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200677 throw CompilerError("Invalid SPIRV format.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100678
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200679 uint32_t bound = s[3];
680 ids.resize(bound);
681 meta.resize(bound);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100682
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200683 uint32_t offset = 5;
684 while (offset < len)
685 inst.emplace_back(spirv, offset);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100686
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200687 for (auto &i : inst)
688 parse(i);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100689
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200690 if (current_function)
691 throw CompilerError("Function was not terminated.");
692 if (current_block)
693 throw CompilerError("Block was not terminated.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100694}
695
696void Compiler::flatten_interface_block(uint32_t id)
697{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200698 auto &var = get<SPIRVariable>(id);
699 auto &type = get<SPIRType>(var.basetype);
700 auto flags = meta.at(type.self).decoration.decoration_flags;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100701
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200702 if (!type.array.empty())
703 throw CompilerError("Type is array of UBOs.");
704 if (type.basetype != SPIRType::Struct)
705 throw CompilerError("Type is not a struct.");
706 if ((flags & (1ull << DecorationBlock)) == 0)
707 throw CompilerError("Type is not a block.");
708 if (type.member_types.empty())
709 throw CompilerError("Member list of struct is empty.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100710
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200711 uint32_t t = type.member_types[0];
712 for (auto &m : type.member_types)
713 if (t != m)
714 throw CompilerError("Types in block differ.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100715
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200716 auto &mtype = get<SPIRType>(t);
717 if (!mtype.array.empty())
718 throw CompilerError("Member type cannot be arrays.");
719 if (mtype.basetype == SPIRType::Struct)
720 throw CompilerError("Member type cannot be struct.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100721
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200722 // Inherit variable name from interface block name.
723 meta.at(var.self).decoration.alias = meta.at(type.self).decoration.alias;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100724
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200725 auto storage = var.storage;
726 if (storage == StorageClassUniform)
727 storage = StorageClassUniformConstant;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100728
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200729 // Change type definition in-place into an array instead.
730 // Access chains will still work as-is.
731 uint32_t array_size = uint32_t(type.member_types.size());
732 type = mtype;
733 type.array.push_back(array_size);
734 type.pointer = true;
735 type.storage = storage;
736 var.storage = storage;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100737}
738
739void Compiler::update_name_cache(unordered_set<string> &cache, string &name)
740{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200741 if (name.empty())
742 return;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100743
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200744 if (cache.find(name) == end(cache))
745 {
746 cache.insert(name);
747 return;
748 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100749
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200750 uint32_t counter = 0;
751 auto tmpname = name;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100752
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200753 // If there is a collision (very rare),
754 // keep tacking on extra identifier until it's unique.
755 do
756 {
757 counter++;
758 name = tmpname + "_" + convert_to_string(counter);
759 } while (cache.find(name) != end(cache));
760 cache.insert(name);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100761}
762
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200763void Compiler::set_name(uint32_t id, const std::string &name)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100764{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200765 auto &str = meta.at(id).decoration.alias;
766 str.clear();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100767
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200768 if (name.empty())
769 return;
770 // Reserved for temporaries.
Hans-Kristian Arntzenc9728942016-07-06 11:19:20 +0200771 if (name[0] == '_' && name.size() >= 2 && isdigit(name[1]))
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200772 return;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100773
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200774 // Functions in glslangValidator are mangled with name(<mangled> stuff.
775 // Normally, we would never see '(' in any legal indentifiers, so just strip them out.
776 str = name.substr(0, name.find('('));
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100777
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200778 for (uint32_t i = 0; i < str.size(); i++)
779 {
780 auto &c = str[i];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100781
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200782 // _<num> variables are reserved by the internal implementation,
783 // otherwise, make sure the name is a valid identifier.
784 if (i == 0 || (str[0] == '_' && i == 1))
785 c = isalpha(c) ? c : '_';
786 else
787 c = isalnum(c) ? c : '_';
788 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100789}
790
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200791const SPIRType &Compiler::get_type(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100792{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200793 return get<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100794}
795
796void Compiler::set_member_decoration(uint32_t id, uint32_t index, Decoration decoration, uint32_t argument)
797{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200798 meta.at(id).members.resize(max(meta[id].members.size(), size_t(index) + 1));
799 auto &dec = meta.at(id).members[index];
800 dec.decoration_flags |= 1ull << decoration;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100801
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200802 switch (decoration)
803 {
804 case DecorationBuiltIn:
805 dec.builtin = true;
806 dec.builtin_type = static_cast<BuiltIn>(argument);
807 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100808
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200809 case DecorationLocation:
810 dec.location = argument;
811 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100812
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200813 case DecorationOffset:
814 dec.offset = argument;
815 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100816
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200817 default:
818 break;
819 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100820}
821
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200822void Compiler::set_member_name(uint32_t id, uint32_t index, const std::string &name)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100823{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200824 meta.at(id).members.resize(max(meta[id].members.size(), size_t(index) + 1));
825 meta.at(id).members[index].alias = name;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100826}
827
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200828const std::string &Compiler::get_member_name(uint32_t id, uint32_t index) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100829{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200830 auto &m = meta.at(id);
831 if (index >= m.members.size())
832 {
833 static string empty;
834 return empty;
835 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100836
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200837 return m.members[index].alias;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100838}
839
840uint32_t Compiler::get_member_decoration(uint32_t id, uint32_t index, Decoration decoration) const
841{
Hans-Kristian Arntzen15679c72016-08-17 11:35:34 +0200842 auto &m = meta.at(id);
843 if (index >= m.members.size())
844 return 0;
845
846 auto &dec = m.members[index];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200847 if (!(dec.decoration_flags & (1ull << decoration)))
848 return 0;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100849
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200850 switch (decoration)
851 {
852 case DecorationBuiltIn:
853 return dec.builtin_type;
854 case DecorationLocation:
855 return dec.location;
856 case DecorationOffset:
857 return dec.offset;
858 default:
859 return 0;
860 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100861}
862
863uint64_t Compiler::get_member_decoration_mask(uint32_t id, uint32_t index) const
864{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200865 auto &m = meta.at(id);
866 if (index >= m.members.size())
867 return 0;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100868
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200869 return m.members[index].decoration_flags;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100870}
871
872void Compiler::unset_member_decoration(uint32_t id, uint32_t index, Decoration decoration)
873{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200874 auto &m = meta.at(id);
875 if (index >= m.members.size())
876 return;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100877
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200878 auto &dec = m.members[index];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100879
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200880 dec.decoration_flags &= ~(1ull << decoration);
881 switch (decoration)
882 {
883 case DecorationBuiltIn:
884 dec.builtin = false;
885 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100886
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200887 case DecorationLocation:
888 dec.location = 0;
889 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100890
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200891 case DecorationOffset:
892 dec.offset = 0;
893 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100894
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200895 default:
896 break;
897 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100898}
899
900void Compiler::set_decoration(uint32_t id, Decoration decoration, uint32_t argument)
901{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200902 auto &dec = meta.at(id).decoration;
903 dec.decoration_flags |= 1ull << decoration;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100904
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200905 switch (decoration)
906 {
907 case DecorationBuiltIn:
908 dec.builtin = true;
909 dec.builtin_type = static_cast<BuiltIn>(argument);
910 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100911
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200912 case DecorationLocation:
913 dec.location = argument;
914 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100915
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200916 case DecorationOffset:
917 dec.offset = argument;
918 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100919
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200920 case DecorationArrayStride:
921 dec.array_stride = argument;
922 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100923
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200924 case DecorationBinding:
925 dec.binding = argument;
926 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100927
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200928 case DecorationDescriptorSet:
929 dec.set = argument;
930 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100931
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200932 case DecorationInputAttachmentIndex:
933 dec.input_attachment = argument;
934 break;
Hans-Kristian Arntzen12cfbb22016-05-04 13:41:04 +0200935
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200936 default:
937 break;
938 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100939}
940
941StorageClass Compiler::get_storage_class(uint32_t id) const
942{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200943 return get<SPIRVariable>(id).storage;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100944}
945
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200946const std::string &Compiler::get_name(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100947{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200948 return meta.at(id).decoration.alias;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100949}
950
951uint64_t Compiler::get_decoration_mask(uint32_t id) const
952{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200953 auto &dec = meta.at(id).decoration;
954 return dec.decoration_flags;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100955}
956
957uint32_t Compiler::get_decoration(uint32_t id, Decoration decoration) const
958{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200959 auto &dec = meta.at(id).decoration;
960 if (!(dec.decoration_flags & (1ull << decoration)))
961 return 0;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100962
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200963 switch (decoration)
964 {
965 case DecorationBuiltIn:
966 return dec.builtin_type;
967 case DecorationLocation:
968 return dec.location;
969 case DecorationOffset:
970 return dec.offset;
971 case DecorationBinding:
972 return dec.binding;
973 case DecorationDescriptorSet:
974 return dec.set;
975 case DecorationInputAttachmentIndex:
976 return dec.input_attachment;
977 default:
978 return 0;
979 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100980}
981
982void Compiler::unset_decoration(uint32_t id, Decoration decoration)
983{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200984 auto &dec = meta.at(id).decoration;
985 dec.decoration_flags &= ~(1ull << decoration);
986 switch (decoration)
987 {
988 case DecorationBuiltIn:
989 dec.builtin = false;
990 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100991
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200992 case DecorationLocation:
993 dec.location = 0;
994 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100995
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +0200996 case DecorationOffset:
997 dec.offset = 0;
998 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +0100999
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001000 case DecorationBinding:
1001 dec.binding = 0;
1002 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001003
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001004 case DecorationDescriptorSet:
1005 dec.set = 0;
1006 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001007
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001008 default:
1009 break;
1010 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001011}
1012
Hans-Kristian Arntzen926916d2016-05-05 09:15:25 +02001013void Compiler::parse(const Instruction &instruction)
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001014{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001015 auto ops = stream(instruction);
1016 auto op = static_cast<Op>(instruction.op);
1017 uint32_t length = instruction.length;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001018
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001019 switch (op)
1020 {
1021 case OpMemoryModel:
1022 case OpSourceExtension:
1023 case OpNop:
Robert Konrad8f7c1af2016-08-10 02:43:51 +02001024 case OpLine:
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001025 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001026
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001027 case OpSource:
1028 {
1029 auto lang = static_cast<SourceLanguage>(ops[0]);
1030 switch (lang)
1031 {
1032 case SourceLanguageESSL:
1033 source.es = true;
1034 source.version = ops[1];
1035 source.known = true;
1036 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001037
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001038 case SourceLanguageGLSL:
1039 source.es = false;
1040 source.version = ops[1];
1041 source.known = true;
1042 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001043
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001044 default:
1045 source.known = false;
1046 break;
1047 }
1048 break;
1049 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001050
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001051 case OpUndef:
1052 {
1053 uint32_t result_type = ops[0];
1054 uint32_t id = ops[1];
1055 set<SPIRUndef>(id, result_type);
1056 break;
1057 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001058
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001059 case OpCapability:
1060 {
1061 uint32_t cap = ops[0];
1062 if (cap == CapabilityKernel)
1063 throw CompilerError("Kernel capability not supported.");
1064 break;
1065 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001066
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001067 case OpExtInstImport:
1068 {
1069 uint32_t id = ops[0];
1070 auto ext = extract_string(spirv, instruction.offset + 1);
1071 if (ext == "GLSL.std.450")
1072 set<SPIRExtension>(id, SPIRExtension::GLSL);
1073 else
1074 throw CompilerError("Only GLSL.std.450 extension interface supported.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001075
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001076 break;
1077 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001078
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001079 case OpEntryPoint:
1080 {
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001081 auto itr = entry_points.emplace(ops[1], SPIREntryPoint(ops[1], static_cast<ExecutionModel>(ops[0]),
1082 extract_string(spirv, instruction.offset + 2)));
1083 auto &e = itr.first->second;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001084
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001085 // Strings need nul-terminator and consume the whole word.
1086 uint32_t strlen_words = (e.name.size() + 1 + 3) >> 2;
1087 e.interface_variables.insert(end(e.interface_variables), ops + strlen_words + 2, ops + instruction.length);
1088
1089 // If we don't have an entry, make the first one our "default".
1090 if (!entry_point)
1091 entry_point = ops[1];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001092 break;
1093 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001094
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001095 case OpExecutionMode:
1096 {
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02001097 auto &execution = entry_points[ops[0]];
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001098 auto mode = static_cast<ExecutionMode>(ops[1]);
1099 execution.flags |= 1ull << mode;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001100
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001101 switch (mode)
1102 {
1103 case ExecutionModeInvocations:
1104 execution.invocations = ops[2];
1105 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001106
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001107 case ExecutionModeLocalSize:
1108 execution.workgroup_size.x = ops[2];
1109 execution.workgroup_size.y = ops[3];
1110 execution.workgroup_size.z = ops[4];
1111 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001112
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001113 case ExecutionModeOutputVertices:
1114 execution.output_vertices = ops[2];
1115 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001116
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001117 default:
1118 break;
1119 }
1120 break;
1121 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001122
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001123 case OpName:
1124 {
1125 uint32_t id = ops[0];
1126 set_name(id, extract_string(spirv, instruction.offset + 1));
1127 break;
1128 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001129
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001130 case OpMemberName:
1131 {
1132 uint32_t id = ops[0];
1133 uint32_t member = ops[1];
1134 set_member_name(id, member, extract_string(spirv, instruction.offset + 2));
1135 break;
1136 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001137
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001138 case OpDecorate:
1139 {
1140 uint32_t id = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001141
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001142 auto decoration = static_cast<Decoration>(ops[1]);
1143 if (length >= 3)
1144 set_decoration(id, decoration, ops[2]);
1145 else
1146 set_decoration(id, decoration);
1147 break;
1148 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001149
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001150 case OpMemberDecorate:
1151 {
1152 uint32_t id = ops[0];
1153 uint32_t member = ops[1];
1154 auto decoration = static_cast<Decoration>(ops[2]);
1155 if (length >= 4)
1156 set_member_decoration(id, member, decoration, ops[3]);
1157 else
1158 set_member_decoration(id, member, decoration);
1159 break;
1160 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001161
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001162 // Build up basic types.
1163 case OpTypeVoid:
1164 {
1165 uint32_t id = ops[0];
1166 auto &type = set<SPIRType>(id);
1167 type.basetype = SPIRType::Void;
1168 break;
1169 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001170
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001171 case OpTypeBool:
1172 {
1173 uint32_t id = ops[0];
1174 auto &type = set<SPIRType>(id);
Hans-Kristian Arntzen5f629272016-06-05 20:13:45 +02001175 type.basetype = SPIRType::Boolean;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001176 type.width = 1;
1177 break;
1178 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001179
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001180 case OpTypeFloat:
1181 {
1182 uint32_t id = ops[0];
1183 uint32_t width = ops[1];
1184 auto &type = set<SPIRType>(id);
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001185 type.basetype = width > 32 ? SPIRType::Double : SPIRType::Float;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001186 type.width = width;
1187 break;
1188 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001189
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001190 case OpTypeInt:
1191 {
1192 uint32_t id = ops[0];
1193 uint32_t width = ops[1];
1194 auto &type = set<SPIRType>(id);
Hans-Kristian Arntzenfc2230f2016-07-27 11:27:00 +02001195 type.basetype =
1196 ops[2] ? (width > 32 ? SPIRType::Int64 : SPIRType::Int) : (width > 32 ? SPIRType::UInt64 : SPIRType::UInt);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001197 type.width = width;
1198 break;
1199 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001200
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001201 // Build composite types by "inheriting".
1202 // NOTE: The self member is also copied! For pointers and array modifiers this is a good thing
1203 // since we can refer to decorations on pointee classes which is needed for UBO/SSBO, I/O blocks in geometry/tess etc.
1204 case OpTypeVector:
1205 {
1206 uint32_t id = ops[0];
1207 uint32_t vecsize = ops[2];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001208
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001209 auto &base = get<SPIRType>(ops[1]);
1210 auto &vecbase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001211
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001212 vecbase = base;
1213 vecbase.vecsize = vecsize;
1214 vecbase.self = id;
1215 break;
1216 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001217
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001218 case OpTypeMatrix:
1219 {
1220 uint32_t id = ops[0];
1221 uint32_t colcount = ops[2];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001222
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001223 auto &base = get<SPIRType>(ops[1]);
1224 auto &matrixbase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001225
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001226 matrixbase = base;
1227 matrixbase.columns = colcount;
1228 matrixbase.self = id;
1229 break;
1230 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001231
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001232 case OpTypeArray:
1233 {
1234 uint32_t id = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001235
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001236 auto &base = get<SPIRType>(ops[1]);
1237 auto &arraybase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001238
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001239 arraybase = base;
1240 arraybase.array.push_back(get<SPIRConstant>(ops[2]).scalar());
1241 // Do NOT set arraybase.self!
1242 break;
1243 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001244
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001245 case OpTypeRuntimeArray:
1246 {
1247 uint32_t id = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001248
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001249 auto &base = get<SPIRType>(ops[1]);
1250 auto &arraybase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001251
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001252 arraybase = base;
1253 arraybase.array.push_back(0);
1254 // Do NOT set arraybase.self!
1255 break;
1256 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001257
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001258 case OpTypeImage:
1259 {
1260 uint32_t id = ops[0];
1261 auto &type = set<SPIRType>(id);
1262 type.basetype = SPIRType::Image;
1263 type.image.type = ops[1];
1264 type.image.dim = static_cast<Dim>(ops[2]);
1265 type.image.depth = ops[3] != 0;
1266 type.image.arrayed = ops[4] != 0;
1267 type.image.ms = ops[5] != 0;
1268 type.image.sampled = ops[6];
1269 type.image.format = static_cast<ImageFormat>(ops[7]);
1270 break;
1271 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001272
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001273 case OpTypeSampledImage:
1274 {
1275 uint32_t id = ops[0];
1276 uint32_t imagetype = ops[1];
1277 auto &type = set<SPIRType>(id);
1278 type = get<SPIRType>(imagetype);
1279 type.basetype = SPIRType::SampledImage;
1280 type.self = id;
1281 break;
1282 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001283
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001284 // Not really used.
1285 case OpTypeSampler:
1286 {
1287 uint32_t id = ops[0];
1288 auto &type = set<SPIRType>(id);
1289 type.basetype = SPIRType::Sampler;
1290 break;
1291 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001292
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001293 case OpTypePointer:
1294 {
1295 uint32_t id = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001296
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001297 auto &base = get<SPIRType>(ops[2]);
1298 auto &ptrbase = set<SPIRType>(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001299
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001300 ptrbase = base;
1301 if (ptrbase.pointer)
1302 throw CompilerError("Cannot make pointer-to-pointer type.");
1303 ptrbase.pointer = true;
1304 ptrbase.storage = static_cast<StorageClass>(ops[1]);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001305
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001306 if (ptrbase.storage == StorageClassAtomicCounter)
1307 ptrbase.basetype = SPIRType::AtomicCounter;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001308
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001309 // Do NOT set ptrbase.self!
1310 break;
1311 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001312
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001313 case OpTypeStruct:
1314 {
1315 uint32_t id = ops[0];
1316 auto &type = set<SPIRType>(id);
1317 type.basetype = SPIRType::Struct;
1318 for (uint32_t i = 1; i < length; i++)
1319 type.member_types.push_back(ops[i]);
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02001320
1321 // Check if we have seen this struct type before, with just different
1322 // decorations.
Hans-Kristian Arntzen5ad43402016-05-28 09:47:52 +02001323 //
1324 // Add workaround for issue #17 as well by looking at OpName for the struct
1325 // types, which we shouldn't normally do.
1326 // We should not normally have to consider type aliases like this to begin with
1327 // however ... glslang issues #304, #307 cover this.
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02001328 for (auto &other : global_struct_cache)
1329 {
Hans-Kristian Arntzen5ad43402016-05-28 09:47:52 +02001330 if (get_name(type.self) == get_name(other) && types_are_logically_equivalent(type, get<SPIRType>(other)))
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02001331 {
1332 type.type_alias = other;
1333 break;
1334 }
1335 }
1336
1337 if (type.type_alias == 0)
1338 global_struct_cache.push_back(id);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001339 break;
1340 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001341
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001342 case OpTypeFunction:
1343 {
1344 uint32_t id = ops[0];
1345 uint32_t ret = ops[1];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001346
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001347 auto &func = set<SPIRFunctionPrototype>(id, ret);
1348 for (uint32_t i = 2; i < length; i++)
1349 func.parameter_types.push_back(ops[i]);
1350 break;
1351 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001352
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001353 // Variable declaration
1354 // All variables are essentially pointers with a storage qualifier.
1355 case OpVariable:
1356 {
1357 uint32_t type = ops[0];
1358 uint32_t id = ops[1];
1359 auto storage = static_cast<StorageClass>(ops[2]);
1360 uint32_t initializer = length == 4 ? ops[3] : 0;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001361
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001362 if (storage == StorageClassFunction)
1363 {
1364 if (!current_function)
1365 throw CompilerError("No function currently in scope");
1366 current_function->add_local_variable(id);
1367 }
1368 else if (storage == StorageClassPrivate || storage == StorageClassWorkgroup || storage == StorageClassOutput)
1369 {
1370 global_variables.push_back(id);
1371 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001372
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001373 auto &var = set<SPIRVariable>(id, type, storage, initializer);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001374
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001375 if (variable_storage_is_aliased(var))
1376 aliased_variables.push_back(var.self);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001377
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001378 // glslangValidator does not emit required qualifiers here.
1379 // Solve this by making the image access as restricted as possible
1380 // and loosen up if we need to.
1381 auto &vartype = expression_type(id);
1382 if (vartype.basetype == SPIRType::Image)
1383 {
1384 auto &flags = meta.at(id).decoration.decoration_flags;
1385 flags |= 1ull << DecorationNonWritable;
1386 flags |= 1ull << DecorationNonReadable;
1387 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001388
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001389 break;
1390 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001391
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001392 // OpPhi
1393 // OpPhi is a fairly magical opcode.
1394 // It selects temporary variables based on which parent block we *came from*.
1395 // In high-level languages we can "de-SSA" by creating a function local, and flush out temporaries to this function-local
1396 // variable to emulate SSA Phi.
1397 case OpPhi:
1398 {
1399 if (!current_function)
1400 throw CompilerError("No function currently in scope");
1401 if (!current_block)
1402 throw CompilerError("No block currently in scope");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001403
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001404 uint32_t result_type = ops[0];
1405 uint32_t id = ops[1];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001406
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001407 // Instead of a temporary, create a new function-wide temporary with this ID instead.
1408 auto &var = set<SPIRVariable>(id, result_type, spv::StorageClassFunction);
1409 var.phi_variable = true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001410
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001411 current_function->add_local_variable(id);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001412
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001413 for (uint32_t i = 2; i + 2 <= length; i += 2)
1414 current_block->phi_variables.push_back({ ops[i], ops[i + 1], id });
1415 break;
1416 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001417
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001418 // Constants
1419 case OpSpecConstant:
1420 case OpConstant:
1421 {
1422 uint32_t id = ops[1];
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001423 auto &type = get<SPIRType>(ops[0]);
Hans-Kristian Arntzenfc2230f2016-07-27 11:27:00 +02001424 if (type.width > 32)
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001425 set<SPIRConstant>(id, ops[0], ops[2] | (uint64_t(ops[3]) << 32)).specialization = op == OpSpecConstant;
1426 else
1427 set<SPIRConstant>(id, ops[0], ops[2]).specialization = op == OpSpecConstant;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001428 break;
1429 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001430
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001431 case OpSpecConstantFalse:
1432 case OpConstantFalse:
1433 {
1434 uint32_t id = ops[1];
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001435 set<SPIRConstant>(id, ops[0], uint32_t(0)).specialization = op == OpSpecConstantFalse;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001436 break;
1437 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001438
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001439 case OpSpecConstantTrue:
1440 case OpConstantTrue:
1441 {
1442 uint32_t id = ops[1];
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001443 set<SPIRConstant>(id, ops[0], uint32_t(1)).specialization = op == OpSpecConstantTrue;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001444 break;
1445 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001446
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001447 case OpSpecConstantComposite:
1448 case OpConstantComposite:
1449 {
1450 uint32_t id = ops[1];
1451 uint32_t type = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001452
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001453 auto &ctype = get<SPIRType>(type);
1454 SPIRConstant *constant = nullptr;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001455
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001456 // We can have constants which are structs and arrays.
1457 // In this case, our SPIRConstant will be a list of other SPIRConstant ids which we
1458 // can refer to.
1459 if (ctype.basetype == SPIRType::Struct || !ctype.array.empty())
1460 {
1461 constant = &set<SPIRConstant>(id, type, ops + 2, length - 2);
1462 constant->specialization = op == OpSpecConstantComposite;
1463 break;
1464 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001465
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001466 bool type_64bit = ctype.width > 32;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001467 bool matrix = ctype.columns > 1;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001468
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001469 if (matrix)
1470 {
1471 switch (length - 2)
1472 {
1473 case 1:
1474 constant = &set<SPIRConstant>(id, type, get<SPIRConstant>(ops[2]).vector());
1475 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001476
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001477 case 2:
1478 constant = &set<SPIRConstant>(id, type, get<SPIRConstant>(ops[2]).vector(),
1479 get<SPIRConstant>(ops[3]).vector());
1480 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001481
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001482 case 3:
1483 constant = &set<SPIRConstant>(id, type, get<SPIRConstant>(ops[2]).vector(),
1484 get<SPIRConstant>(ops[3]).vector(), get<SPIRConstant>(ops[4]).vector());
1485 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001486
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001487 case 4:
1488 constant =
1489 &set<SPIRConstant>(id, type, get<SPIRConstant>(ops[2]).vector(), get<SPIRConstant>(ops[3]).vector(),
1490 get<SPIRConstant>(ops[4]).vector(), get<SPIRConstant>(ops[5]).vector());
1491 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001492
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001493 default:
1494 throw CompilerError("OpConstantComposite only supports 1, 2, 3 and 4 columns.");
1495 }
1496 }
1497 else
1498 {
1499 switch (length - 2)
1500 {
1501 case 1:
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001502 if (type_64bit)
1503 constant = &set<SPIRConstant>(id, type, get<SPIRConstant>(ops[2]).scalar_u64());
1504 else
1505 constant = &set<SPIRConstant>(id, type, get<SPIRConstant>(ops[2]).scalar());
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001506 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001507
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001508 case 2:
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001509 if (type_64bit)
1510 {
1511 constant = &set<SPIRConstant>(id, type, get<SPIRConstant>(ops[2]).scalar_u64(),
1512 get<SPIRConstant>(ops[3]).scalar_u64());
1513 }
1514 else
1515 {
1516 constant = &set<SPIRConstant>(id, type, get<SPIRConstant>(ops[2]).scalar(),
1517 get<SPIRConstant>(ops[3]).scalar());
1518 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001519 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001520
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001521 case 3:
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001522 if (type_64bit)
1523 {
1524 constant = &set<SPIRConstant>(id, type, get<SPIRConstant>(ops[2]).scalar_u64(),
1525 get<SPIRConstant>(ops[3]).scalar_u64(),
1526 get<SPIRConstant>(ops[4]).scalar_u64());
1527 }
1528 else
1529 {
1530 constant =
1531 &set<SPIRConstant>(id, type, get<SPIRConstant>(ops[2]).scalar(),
1532 get<SPIRConstant>(ops[3]).scalar(), get<SPIRConstant>(ops[4]).scalar());
1533 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001534 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001535
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001536 case 4:
Hans-Kristian Arntzenfa0255c2016-07-27 10:59:00 +02001537 if (type_64bit)
1538 {
1539 constant = &set<SPIRConstant>(
1540 id, type, get<SPIRConstant>(ops[2]).scalar_u64(), get<SPIRConstant>(ops[3]).scalar_u64(),
1541 get<SPIRConstant>(ops[4]).scalar_u64(), get<SPIRConstant>(ops[5]).scalar_u64());
1542 }
1543 else
1544 {
1545 constant = &set<SPIRConstant>(
1546 id, type, get<SPIRConstant>(ops[2]).scalar(), get<SPIRConstant>(ops[3]).scalar(),
1547 get<SPIRConstant>(ops[4]).scalar(), get<SPIRConstant>(ops[5]).scalar());
1548 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001549 break;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001550
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001551 default:
1552 throw CompilerError("OpConstantComposite only supports 1, 2, 3 and 4 components.");
1553 }
1554 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001555
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001556 constant->specialization = op == OpSpecConstantComposite;
1557 break;
1558 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001559
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001560 // Functions
1561 case OpFunction:
1562 {
1563 uint32_t res = ops[0];
1564 uint32_t id = ops[1];
1565 // Control
1566 uint32_t type = ops[3];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001567
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001568 if (current_function)
1569 throw CompilerError("Must end a function before starting a new one!");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001570
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001571 current_function = &set<SPIRFunction>(id, res, type);
1572 break;
1573 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001574
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001575 case OpFunctionParameter:
1576 {
1577 uint32_t type = ops[0];
1578 uint32_t id = ops[1];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001579
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001580 if (!current_function)
1581 throw CompilerError("Must be in a function!");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001582
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001583 current_function->add_parameter(type, id);
1584 set<SPIRVariable>(id, type, StorageClassFunction);
1585 break;
1586 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001587
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001588 case OpFunctionEnd:
1589 {
1590 current_function = nullptr;
1591 break;
1592 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001593
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001594 // Blocks
1595 case OpLabel:
1596 {
1597 // OpLabel always starts a block.
1598 if (!current_function)
1599 throw CompilerError("Blocks cannot exist outside functions!");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001600
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001601 uint32_t id = ops[0];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001602
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001603 current_function->blocks.push_back(id);
1604 if (!current_function->entry_block)
1605 current_function->entry_block = id;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001606
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001607 if (current_block)
1608 throw CompilerError("Cannot start a block before ending the current block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001609
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001610 current_block = &set<SPIRBlock>(id);
1611 break;
1612 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001613
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001614 // Branch instructions end blocks.
1615 case OpBranch:
1616 {
1617 if (!current_block)
1618 throw CompilerError("Trying to end a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001619
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001620 uint32_t target = ops[0];
1621 current_block->terminator = SPIRBlock::Direct;
1622 current_block->next_block = target;
1623 current_block = nullptr;
1624 break;
1625 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001626
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001627 case OpBranchConditional:
1628 {
1629 if (!current_block)
1630 throw CompilerError("Trying to end a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001631
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001632 current_block->condition = ops[0];
1633 current_block->true_block = ops[1];
1634 current_block->false_block = ops[2];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001635
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001636 current_block->terminator = SPIRBlock::Select;
1637 current_block = nullptr;
1638 break;
1639 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001640
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001641 case OpSwitch:
1642 {
1643 if (!current_block)
1644 throw CompilerError("Trying to end a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001645
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001646 if (current_block->merge == SPIRBlock::MergeNone)
1647 throw CompilerError("Switch statement is not structured");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001648
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001649 current_block->terminator = SPIRBlock::MultiSelect;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001650
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001651 current_block->condition = ops[0];
1652 current_block->default_block = ops[1];
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001653
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001654 for (uint32_t i = 2; i + 2 <= length; i += 2)
1655 current_block->cases.push_back({ ops[i], ops[i + 1] });
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001656
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001657 // If we jump to next block, make it break instead since we're inside a switch case block at that point.
1658 multiselect_merge_targets.insert(current_block->next_block);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001659
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001660 current_block = nullptr;
1661 break;
1662 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001663
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001664 case OpKill:
1665 {
1666 if (!current_block)
1667 throw CompilerError("Trying to end a non-existing block.");
1668 current_block->terminator = SPIRBlock::Kill;
1669 current_block = nullptr;
1670 break;
1671 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001672
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001673 case OpReturn:
1674 {
1675 if (!current_block)
1676 throw CompilerError("Trying to end a non-existing block.");
1677 current_block->terminator = SPIRBlock::Return;
1678 current_block = nullptr;
1679 break;
1680 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001681
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001682 case OpReturnValue:
1683 {
1684 if (!current_block)
1685 throw CompilerError("Trying to end a non-existing block.");
1686 current_block->terminator = SPIRBlock::Return;
1687 current_block->return_value = ops[0];
1688 current_block = nullptr;
1689 break;
1690 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001691
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001692 case OpUnreachable:
1693 {
1694 if (!current_block)
1695 throw CompilerError("Trying to end a non-existing block.");
1696 current_block->terminator = SPIRBlock::Unreachable;
1697 current_block = nullptr;
1698 break;
1699 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001700
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001701 case OpSelectionMerge:
1702 {
1703 if (!current_block)
1704 throw CompilerError("Trying to modify a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001705
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001706 current_block->next_block = ops[0];
1707 current_block->merge = SPIRBlock::MergeSelection;
1708 selection_merge_targets.insert(current_block->next_block);
1709 break;
1710 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001711
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001712 case OpLoopMerge:
1713 {
1714 if (!current_block)
1715 throw CompilerError("Trying to modify a non-existing block.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001716
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001717 current_block->merge_block = ops[0];
1718 current_block->continue_block = ops[1];
1719 current_block->merge = SPIRBlock::MergeLoop;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001720
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001721 loop_blocks.insert(current_block->self);
1722 loop_merge_targets.insert(current_block->merge_block);
Hans-Kristian Arntzen97f81ba2016-04-01 12:37:29 +02001723
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001724 // Don't add loop headers to continue blocks,
1725 // which would make it impossible branch into the loop header since
1726 // they are treated as continues.
1727 if (current_block->continue_block != current_block->self)
1728 continue_blocks.insert(current_block->continue_block);
1729 break;
1730 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001731
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001732 // Actual opcodes.
1733 default:
1734 {
1735 if (!current_block)
1736 throw CompilerError("Currently no block to insert opcode.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001737
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001738 current_block->ops.push_back(instruction);
1739 break;
1740 }
1741 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001742}
1743
1744bool Compiler::block_is_loop_candidate(const SPIRBlock &block, SPIRBlock::Method method) const
1745{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001746 // Tried and failed.
1747 if (block.disable_block_optimization || block.complex_continue)
1748 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001749
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001750 if (method == SPIRBlock::MergeToSelectForLoop)
1751 {
1752 // Try to detect common for loop pattern
1753 // which the code backend can use to create cleaner code.
1754 // for(;;) { if (cond) { some_body; } else { break; } }
1755 // is the pattern we're looking for.
1756 bool ret = block.terminator == SPIRBlock::Select && block.merge == SPIRBlock::MergeLoop &&
1757 block.true_block != block.merge_block && block.true_block != block.self &&
1758 block.false_block == block.merge_block;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001759
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001760 // If we have OpPhi which depends on branches which came from our own block,
1761 // we need to flush phi variables in else block instead of a trivial break,
1762 // so we cannot assume this is a for loop candidate.
1763 if (ret)
1764 {
1765 for (auto &phi : block.phi_variables)
1766 if (phi.parent == block.self)
1767 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001768
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001769 auto *merge = maybe_get<SPIRBlock>(block.merge_block);
1770 if (merge)
1771 for (auto &phi : merge->phi_variables)
1772 if (phi.parent == block.self)
1773 return false;
1774 }
1775 return ret;
1776 }
1777 else if (method == SPIRBlock::MergeToDirectForLoop)
1778 {
1779 // Empty loop header that just sets up merge target
1780 // and branches to loop body.
1781 bool ret = block.terminator == SPIRBlock::Direct && block.merge == SPIRBlock::MergeLoop && block.ops.empty();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001782
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001783 if (!ret)
1784 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001785
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001786 auto &child = get<SPIRBlock>(block.next_block);
1787 ret = child.terminator == SPIRBlock::Select && child.merge == SPIRBlock::MergeNone &&
1788 child.false_block == block.merge_block && child.true_block != block.merge_block &&
1789 child.true_block != block.self;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001790
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001791 // If we have OpPhi which depends on branches which came from our own block,
1792 // we need to flush phi variables in else block instead of a trivial break,
1793 // so we cannot assume this is a for loop candidate.
1794 if (ret)
1795 {
1796 for (auto &phi : block.phi_variables)
1797 if (phi.parent == block.self || phi.parent == child.self)
1798 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001799
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001800 for (auto &phi : child.phi_variables)
1801 if (phi.parent == block.self)
1802 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001803
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001804 auto *merge = maybe_get<SPIRBlock>(block.merge_block);
1805 if (merge)
1806 for (auto &phi : merge->phi_variables)
1807 if (phi.parent == block.self || phi.parent == child.false_block)
1808 return false;
1809 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001810
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001811 return ret;
1812 }
1813 else
1814 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001815}
1816
1817bool Compiler::block_is_outside_flow_control_from_block(const SPIRBlock &from, const SPIRBlock &to)
1818{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001819 auto *start = &from;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001820
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001821 if (start->self == to.self)
1822 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001823
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001824 // Break cycles.
1825 if (is_continue(start->self))
1826 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001827
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001828 // If our select block doesn't merge, we must break or continue in these blocks,
1829 // so if continues occur branchless within these blocks, consider them branchless as well.
1830 // This is typically used for loop control.
1831 if (start->terminator == SPIRBlock::Select && start->merge == SPIRBlock::MergeNone &&
1832 (block_is_outside_flow_control_from_block(get<SPIRBlock>(start->true_block), to) ||
1833 block_is_outside_flow_control_from_block(get<SPIRBlock>(start->false_block), to)))
1834 {
1835 return true;
1836 }
1837 else if (start->merge_block && block_is_outside_flow_control_from_block(get<SPIRBlock>(start->merge_block), to))
1838 {
1839 return true;
1840 }
1841 else if (start->next_block && block_is_outside_flow_control_from_block(get<SPIRBlock>(start->next_block), to))
1842 {
1843 return true;
1844 }
1845 else
1846 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001847}
1848
1849bool Compiler::execution_is_noop(const SPIRBlock &from, const SPIRBlock &to) const
1850{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001851 if (!execution_is_branchless(from, to))
1852 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001853
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001854 auto *start = &from;
1855 for (;;)
1856 {
1857 if (start->self == to.self)
1858 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001859
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001860 if (!start->ops.empty())
1861 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001862
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001863 start = &get<SPIRBlock>(start->next_block);
1864 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001865}
1866
1867bool Compiler::execution_is_branchless(const SPIRBlock &from, const SPIRBlock &to) const
1868{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001869 auto *start = &from;
1870 for (;;)
1871 {
1872 if (start->self == to.self)
1873 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001874
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001875 if (start->terminator == SPIRBlock::Direct && start->merge == SPIRBlock::MergeNone)
1876 start = &get<SPIRBlock>(start->next_block);
1877 else
1878 return false;
1879 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001880}
1881
Hans-Kristian Arntzen926916d2016-05-05 09:15:25 +02001882SPIRBlock::ContinueBlockType Compiler::continue_block_type(const SPIRBlock &block) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001883{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001884 // The block was deemed too complex during code emit, pick conservative fallback paths.
1885 if (block.complex_continue)
1886 return SPIRBlock::ComplexLoop;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001887
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001888 // In older glslang output continue block can be equal to the loop header.
1889 // In this case, execution is clearly branchless, so just assume a while loop header here.
1890 if (block.merge == SPIRBlock::MergeLoop)
1891 return SPIRBlock::WhileLoop;
Hans-Kristian Arntzen97f81ba2016-04-01 12:37:29 +02001892
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001893 auto &dominator = get<SPIRBlock>(block.loop_dominator);
Hans-Kristian Arntzen97f81ba2016-04-01 12:37:29 +02001894
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001895 if (execution_is_noop(block, dominator))
1896 return SPIRBlock::WhileLoop;
1897 else if (execution_is_branchless(block, dominator))
1898 return SPIRBlock::ForLoop;
1899 else
1900 {
1901 if (block.merge == SPIRBlock::MergeNone && block.terminator == SPIRBlock::Select &&
1902 block.true_block == dominator.self && block.false_block == dominator.merge_block)
1903 {
1904 return SPIRBlock::DoWhileLoop;
1905 }
1906 else
1907 return SPIRBlock::ComplexLoop;
1908 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001909}
1910
1911bool Compiler::traverse_all_reachable_opcodes(const SPIRBlock &block, OpcodeHandler &handler) const
1912{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001913 // Ideally, perhaps traverse the CFG instead of all blocks in order to eliminate dead blocks,
1914 // but this shouldn't be a problem in practice unless the SPIR-V is doing insane things like recursing
1915 // inside dead blocks ...
1916 for (auto &i : block.ops)
1917 {
1918 auto ops = stream(i);
1919 auto op = static_cast<Op>(i.op);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001920
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001921 if (!handler.handle(op, ops, i.length))
1922 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001923
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02001924 if (op == OpFunctionCall)
1925 {
1926 if (!handler.begin_function_scope(ops, i.length))
1927 return false;
1928 if (!traverse_all_reachable_opcodes(get<SPIRFunction>(ops[2]), handler))
1929 return false;
1930 if (!handler.end_function_scope())
1931 return false;
1932 }
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001933 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001934
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001935 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001936}
1937
1938bool Compiler::traverse_all_reachable_opcodes(const SPIRFunction &func, OpcodeHandler &handler) const
1939{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001940 for (auto block : func.blocks)
1941 if (!traverse_all_reachable_opcodes(get<SPIRBlock>(block), handler))
1942 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001943
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001944 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001945}
1946
1947uint32_t Compiler::type_struct_member_offset(const SPIRType &type, uint32_t index) const
1948{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001949 // Decoration must be set in valid SPIR-V, otherwise throw.
1950 auto &dec = meta[type.self].members.at(index);
1951 if (dec.decoration_flags & (1ull << DecorationOffset))
1952 return dec.offset;
1953 else
1954 throw CompilerError("Struct member does not have Offset set.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001955}
1956
1957uint32_t Compiler::type_struct_member_array_stride(const SPIRType &type, uint32_t index) const
1958{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001959 // Decoration must be set in valid SPIR-V, otherwise throw.
1960 // ArrayStride is part of the array type not OpMemberDecorate.
1961 auto &dec = meta[type.member_types[index]].decoration;
1962 if (dec.decoration_flags & (1ull << DecorationArrayStride))
1963 return dec.array_stride;
1964 else
1965 throw CompilerError("Struct member does not have ArrayStride set.");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001966}
1967
1968size_t Compiler::get_declared_struct_size(const SPIRType &type) const
1969{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001970 uint32_t last = uint32_t(type.member_types.size() - 1);
1971 size_t offset = type_struct_member_offset(type, last);
1972 size_t size = get_declared_struct_member_size(type, last);
1973 return offset + size;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001974}
1975
1976size_t Compiler::get_declared_struct_member_size(const SPIRType &struct_type, uint32_t index) const
1977{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001978 auto flags = get_member_decoration_mask(struct_type.self, index);
1979 auto &type = get<SPIRType>(struct_type.member_types[index]);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001980
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001981 if (type.basetype != SPIRType::Struct)
1982 {
1983 switch (type.basetype)
1984 {
1985 case SPIRType::Unknown:
1986 case SPIRType::Void:
Hans-Kristian Arntzen5f629272016-06-05 20:13:45 +02001987 case SPIRType::Boolean: // Bools are purely logical, and cannot be used for externally visible types.
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001988 case SPIRType::AtomicCounter:
1989 case SPIRType::Image:
1990 case SPIRType::SampledImage:
1991 case SPIRType::Sampler:
1992 throw CompilerError("Querying size for object with opaque size.\n");
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001993
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001994 default:
1995 break;
1996 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01001997
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02001998 size_t component_size = type.width / 8;
1999 unsigned vecsize = type.vecsize;
2000 unsigned columns = type.columns;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002001
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002002 if (type.array.empty())
2003 {
2004 // Vectors.
2005 if (columns == 1)
2006 return vecsize * component_size;
2007 else
2008 {
2009 // Per SPIR-V spec, matrices must be tightly packed and aligned up for vec3 accesses.
2010 if ((flags & (1ull << DecorationRowMajor)) && columns == 3)
2011 columns = 4;
2012 else if ((flags & (1ull << DecorationColMajor)) && vecsize == 3)
2013 vecsize = 4;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002014
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002015 return vecsize * columns * component_size;
2016 }
2017 }
2018 else
2019 {
2020 // For arrays, we can use ArrayStride to get an easy check.
2021 return type_struct_member_array_stride(struct_type, index) * type.array.back();
2022 }
2023 }
2024 else
2025 {
2026 // Recurse.
2027 uint32_t last = uint32_t(struct_type.member_types.size() - 1);
2028 uint32_t offset = type_struct_member_offset(struct_type, last);
2029 size_t size = get_declared_struct_size(get<SPIRType>(struct_type.member_types.back()));
2030 return offset + size;
2031 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002032}
2033
2034bool Compiler::BufferAccessHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
2035{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002036 if (opcode != OpAccessChain && opcode != OpInBoundsAccessChain)
2037 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002038
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002039 // Invalid SPIR-V.
2040 if (length < 4)
2041 return false;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002042
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002043 if (args[2] != id)
2044 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002045
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002046 // Don't bother traversing the entire access chain tree yet.
2047 // If we access a struct member, assume we access the entire member.
2048 uint32_t index = compiler.get<SPIRConstant>(args[3]).scalar();
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002049
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002050 // Seen this index already.
2051 if (seen.find(index) != end(seen))
2052 return true;
2053 seen.insert(index);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002054
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002055 auto &type = compiler.expression_type(id);
2056 uint32_t offset = compiler.type_struct_member_offset(type, index);
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002057
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002058 size_t range;
2059 // If we have another member in the struct, deduce the range by looking at the next member.
2060 // This is okay since structs in SPIR-V can have padding, but Offset decoration must be
2061 // monotonically increasing.
2062 // Of course, this doesn't take into account if the SPIR-V for some reason decided to add
2063 // very large amounts of padding, but that's not really a big deal.
2064 if (index + 1 < type.member_types.size())
2065 {
2066 range = compiler.type_struct_member_offset(type, index + 1) - offset;
2067 }
2068 else
2069 {
2070 // No padding, so just deduce it from the size of the member directly.
2071 range = compiler.get_declared_struct_member_size(type, index);
2072 }
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002073
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002074 ranges.push_back({ index, offset, range });
2075 return true;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002076}
2077
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002078std::vector<BufferRange> Compiler::get_active_buffer_ranges(uint32_t id) const
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002079{
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002080 std::vector<BufferRange> ranges;
2081 BufferAccessHandler handler(*this, ranges, id);
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002082 traverse_all_reachable_opcodes(get<SPIRFunction>(entry_point), handler);
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002083 return ranges;
Hans-Kristian Arntzen75471fb2016-03-02 18:09:16 +01002084}
2085
Bill Hollings103aabf2016-04-06 17:42:27 -04002086// Increase the number of IDs by the specified incremental amount.
2087// Returns the value of the first ID available for use in the expanded bound.
2088uint32_t Compiler::increase_bound_by(uint32_t incr_amount)
2089{
Hans-Kristian Arntzen14bc1ff2016-09-10 18:07:52 +02002090 auto curr_bound = ids.size();
2091 auto new_bound = curr_bound + incr_amount;
Hans-Kristian Arntzen4b8ed532016-05-05 09:33:18 +02002092 ids.resize(new_bound);
2093 meta.resize(new_bound);
Hans-Kristian Arntzen14bc1ff2016-09-10 18:07:52 +02002094 return uint32_t(curr_bound);
Bill Hollings103aabf2016-04-06 17:42:27 -04002095}
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02002096
2097bool Compiler::types_are_logically_equivalent(const SPIRType &a, const SPIRType &b) const
2098{
2099 if (a.basetype != b.basetype)
2100 return false;
2101 if (a.width != b.width)
2102 return false;
2103 if (a.vecsize != b.vecsize)
2104 return false;
2105 if (a.columns != b.columns)
2106 return false;
2107 if (a.array.size() != b.array.size())
2108 return false;
2109
Hans-Kristian Arntzen46892ab2016-05-23 13:45:20 +02002110 size_t array_count = a.array.size();
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02002111 if (array_count && memcmp(a.array.data(), b.array.data(), array_count * sizeof(uint32_t)) != 0)
2112 return false;
2113
2114 if (a.basetype == SPIRType::Image || a.basetype == SPIRType::SampledImage)
2115 {
2116 if (memcmp(&a.image, &b.image, sizeof(SPIRType::Image)) != 0)
2117 return false;
2118 }
2119
2120 if (a.member_types.size() != b.member_types.size())
2121 return false;
2122
Hans-Kristian Arntzen46892ab2016-05-23 13:45:20 +02002123 size_t member_types = a.member_types.size();
2124 for (size_t i = 0; i < member_types; i++)
Hans-Kristian Arntzenf05373b2016-05-23 10:57:22 +02002125 {
2126 if (!types_are_logically_equivalent(get<SPIRType>(a.member_types[i]), get<SPIRType>(b.member_types[i])))
2127 return false;
2128 }
2129
2130 return true;
2131}
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002132
2133uint64_t Compiler::get_execution_mode_mask() const
2134{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002135 return get_entry_point().flags;
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002136}
2137
2138void Compiler::set_execution_mode(ExecutionMode mode, uint32_t arg0, uint32_t arg1, uint32_t arg2)
2139{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002140 auto &execution = get_entry_point();
2141
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002142 execution.flags |= 1ull << mode;
2143 switch (mode)
2144 {
2145 case ExecutionModeLocalSize:
2146 execution.workgroup_size.x = arg0;
2147 execution.workgroup_size.y = arg1;
2148 execution.workgroup_size.z = arg2;
2149 break;
2150
2151 case ExecutionModeInvocations:
2152 execution.invocations = arg0;
2153 break;
2154
2155 case ExecutionModeOutputVertices:
2156 execution.output_vertices = arg0;
2157 break;
2158
2159 default:
2160 break;
2161 }
2162}
2163
2164void Compiler::unset_execution_mode(ExecutionMode mode)
2165{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002166 auto &execution = get_entry_point();
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002167 execution.flags &= ~(1ull << mode);
2168}
2169
2170uint32_t Compiler::get_execution_mode_argument(spv::ExecutionMode mode, uint32_t index) const
2171{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002172 auto &execution = get_entry_point();
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002173 switch (mode)
2174 {
2175 case ExecutionModeLocalSize:
2176 switch (index)
2177 {
2178 case 0:
2179 return execution.workgroup_size.x;
2180 case 1:
2181 return execution.workgroup_size.y;
2182 case 2:
2183 return execution.workgroup_size.z;
2184 default:
2185 return 0;
2186 }
2187
2188 case ExecutionModeInvocations:
2189 return execution.invocations;
2190
2191 case ExecutionModeOutputVertices:
2192 return execution.output_vertices;
2193
2194 default:
2195 return 0;
2196 }
2197}
2198
2199ExecutionModel Compiler::get_execution_model() const
2200{
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002201 auto &execution = get_entry_point();
Hans-Kristian Arntzen3c285a12016-07-04 13:30:05 +02002202 return execution.model;
2203}
Hans-Kristian Arntzen8e63c772016-07-06 09:58:01 +02002204
2205void Compiler::set_remapped_variable_state(uint32_t id, bool remap_enable)
2206{
2207 get<SPIRVariable>(id).remapped_variable = remap_enable;
2208}
2209
2210bool Compiler::get_remapped_variable_state(uint32_t id) const
2211{
2212 return get<SPIRVariable>(id).remapped_variable;
2213}
Hans-Kristian Arntzen078eec52016-07-06 11:04:06 +02002214
2215void Compiler::set_subpass_input_remapped_components(uint32_t id, uint32_t components)
2216{
2217 get<SPIRVariable>(id).remapped_components = components;
2218}
2219
2220uint32_t Compiler::get_subpass_input_remapped_components(uint32_t id) const
2221{
2222 return get<SPIRVariable>(id).remapped_components;
2223}
Hans-Kristian Arntzen36a0b632016-07-12 14:33:04 +02002224
2225void Compiler::inherit_expression_dependencies(uint32_t dst, uint32_t source_expression)
2226{
2227 auto &e = get<SPIRExpression>(dst);
2228 auto *s = maybe_get<SPIRExpression>(source_expression);
2229 if (!s)
2230 return;
2231
2232 auto &e_deps = e.expression_dependencies;
2233 auto &s_deps = s->expression_dependencies;
2234
2235 // If we depend on a expression, we also depend on all sub-dependencies from source.
2236 e_deps.push_back(source_expression);
2237 e_deps.insert(end(e_deps), begin(s_deps), end(s_deps));
2238
2239 // Eliminate duplicated dependencies.
2240 e_deps.erase(unique(begin(e_deps), end(e_deps)), end(e_deps));
2241}
Hans-Kristian Arntzen042475e2016-07-28 11:16:02 +02002242
2243vector<string> Compiler::get_entry_points() const
2244{
2245 vector<string> entries;
2246 for (auto &entry : entry_points)
2247 entries.push_back(entry.second.name);
2248 return entries;
2249}
2250
2251void Compiler::set_entry_point(const std::string &name)
2252{
2253 auto &entry = get_entry_point(name);
2254 entry_point = entry.self;
2255}
2256
2257SPIREntryPoint &Compiler::get_entry_point(const std::string &name)
2258{
2259 auto itr =
2260 find_if(begin(entry_points), end(entry_points),
2261 [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool { return entry.second.name == name; });
2262
2263 if (itr == end(entry_points))
2264 throw CompilerError("Entry point does not exist.");
2265
2266 return itr->second;
2267}
2268
2269const SPIREntryPoint &Compiler::get_entry_point(const std::string &name) const
2270{
2271 auto itr =
2272 find_if(begin(entry_points), end(entry_points),
2273 [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool { return entry.second.name == name; });
2274
2275 if (itr == end(entry_points))
2276 throw CompilerError("Entry point does not exist.");
2277
2278 return itr->second;
2279}
2280
2281const SPIREntryPoint &Compiler::get_entry_point() const
2282{
2283 return entry_points.find(entry_point)->second;
2284}
2285
2286SPIREntryPoint &Compiler::get_entry_point()
2287{
2288 return entry_points.find(entry_point)->second;
2289}
2290
2291bool Compiler::interface_variable_exists_in_entry_point(uint32_t id) const
2292{
2293 auto &var = get<SPIRVariable>(id);
2294 if (var.storage != StorageClassInput && var.storage != StorageClassOutput)
2295 throw CompilerError("Only Input and Output variables are part of a shader linking interface.");
2296
2297 // This is to avoid potential problems with very old glslang versions which did
2298 // not emit input/output interfaces properly.
2299 // We can assume they only had a single entry point, and single entry point
2300 // shaders could easily be assumed to use every interface variable anyways.
2301 if (entry_points.size() <= 1)
2302 return true;
2303
2304 auto &execution = get_entry_point();
2305 return find(begin(execution.interface_variables), end(execution.interface_variables), id) !=
2306 end(execution.interface_variables);
2307}
Hans-Kristian Arntzenbcb55602016-09-10 13:56:36 +02002308
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002309void Compiler::CombinedImageSamplerHandler::push_remap_parameters(const SPIRFunction &func, const uint32_t *args,
2310 uint32_t length)
2311{
2312 // If possible, pipe through a remapping table so that parameters know
2313 // which variables they actually bind to in this scope.
2314 unordered_map<uint32_t, uint32_t> remapping;
2315 for (uint32_t i = 0; i < length; i++)
2316 remapping[func.arguments[i].id] = remap_parameter(args[i]);
2317 parameter_remapping.push(move(remapping));
2318}
2319
2320void Compiler::CombinedImageSamplerHandler::pop_remap_parameters()
2321{
2322 parameter_remapping.pop();
2323}
2324
2325uint32_t Compiler::CombinedImageSamplerHandler::remap_parameter(uint32_t id)
2326{
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002327 auto *var = compiler.maybe_get_backing_variable(id);
2328 if (var)
2329 id = var->self;
2330
Hans-Kristian Arntzen901b45e2016-09-10 22:21:57 +02002331 if (parameter_remapping.empty())
2332 return id;
2333
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002334 auto &remapping = parameter_remapping.top();
2335 auto itr = remapping.find(id);
2336 if (itr != end(remapping))
2337 return itr->second;
2338 else
2339 return id;
2340}
2341
2342bool Compiler::CombinedImageSamplerHandler::begin_function_scope(const uint32_t *args, uint32_t length)
2343{
2344 if (length < 3)
2345 return false;
2346
2347 auto &callee = compiler.get<SPIRFunction>(args[2]);
2348 args += 3;
2349 length -= 3;
2350 push_remap_parameters(callee, args, length);
2351 return true;
2352}
2353
2354bool Compiler::CombinedImageSamplerHandler::end_function_scope()
2355{
2356 pop_remap_parameters();
2357 return true;
2358}
2359
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002360bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
2361{
2362 // We need to figure out where samplers and images are loaded from, so do only the bare bones compilation we need.
2363 switch (opcode)
2364 {
2365 case OpLoad:
2366 {
2367 if (length < 3)
2368 return false;
2369
2370 uint32_t result_type = args[0];
2371
2372 auto &type = compiler.get<SPIRType>(result_type);
2373 bool separate_image = type.basetype == SPIRType::Image && type.image.sampled == 1;
2374 bool separate_sampler = type.basetype == SPIRType::Sampler;
2375
2376 // If not separate texture or sampler, don't bother.
2377 if (!separate_image && !separate_sampler)
2378 return true;
2379
2380 uint32_t id = args[1];
2381 uint32_t ptr = args[2];
2382 compiler.set<SPIRExpression>(id, "", result_type, true);
2383 compiler.register_read(id, ptr, true);
2384 return true;
2385 }
2386
2387 case OpInBoundsAccessChain:
2388 case OpAccessChain:
2389 {
2390 if (length < 3)
2391 return false;
2392
2393 // Technically, it is possible to have arrays of textures and arrays of samplers and combine them, but this becomes essentially
2394 // impossible to implement, since we don't know which concrete sampler we are accessing.
2395 // One potential way is to create a combinatorial explosion where N textures and M samplers are combined into N * M sampler2Ds,
2396 // but this seems ridiculously complicated for a problem which is easy to work around.
2397 // Checking access chains like this assumes we don't have samplers or textures inside uniform structs, but this makes no sense.
2398
2399 auto &type = compiler.get<SPIRType>(args[0]);
2400 bool separate_image = type.basetype == SPIRType::Image && type.image.sampled == 1;
2401 bool separate_sampler = type.basetype == SPIRType::Sampler;
2402 if (separate_image)
2403 throw CompilerError(
2404 "Attempting to use arrays of separate images. This is not possible to statically remap to plain GLSL.");
2405 if (separate_sampler)
2406 throw CompilerError("Attempting to use arrays of separate samplers. This is not possible to statically "
2407 "remap to plain GLSL.");
2408 return true;
2409 }
2410
2411 case OpSampledImage:
2412 // Do it outside.
2413 break;
2414
2415 default:
2416 return true;
2417 }
2418
2419 if (length < 4)
2420 return false;
2421
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002422 // For function calls, we need to remap IDs which are function parameters into global variables.
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002423 // This information is statically known from the current place in the call stack.
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002424 // Function parameters are not necessarily pointers, so if we don't have a backing variable, remapping will know
2425 // which backing variable the image/sample came from.
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002426 auto image_id = remap_parameter(args[2]);
2427 auto sampler_id = remap_parameter(args[3]);
2428
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002429 auto itr = find_if(begin(compiler.combined_image_samplers), end(compiler.combined_image_samplers),
2430 [image_id, sampler_id](const CombinedImageSampler &combined) {
2431 return combined.image_id == image_id && combined.sampler_id == sampler_id;
2432 });
2433
2434 if (itr == end(compiler.combined_image_samplers))
2435 {
2436 auto id = compiler.increase_bound_by(2);
2437 auto type_id = id + 0;
2438 auto combined_id = id + 1;
2439 auto sampled_type = args[0];
2440
2441 // Make a new type, pointer to OpTypeSampledImage, so we can make a variable of this type.
2442 // We will probably have this type lying around, but it doesn't hurt to make duplicates for internal purposes.
2443 auto &type = compiler.set<SPIRType>(type_id);
2444 auto &base = compiler.get<SPIRType>(sampled_type);
2445 type = base;
2446 type.pointer = true;
2447 type.storage = StorageClassUniformConstant;
2448
2449 // Build new variable.
Hans-Kristian Arntzen14bc1ff2016-09-10 18:07:52 +02002450 auto &var = compiler.set<SPIRVariable>(combined_id, type_id, StorageClassUniformConstant, 0);
2451 var.storage = StorageClassUniformConstant;
2452
2453 // Inherit RelaxedPrecision (and potentially other useful flags if deemed relevant).
2454 auto &new_flags = compiler.meta[combined_id].decoration.decoration_flags;
2455 auto old_flags = compiler.meta[sampler_id].decoration.decoration_flags;
2456 new_flags = old_flags & (1ull << DecorationRelaxedPrecision);
2457
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002458 compiler.combined_image_samplers.push_back({ combined_id, image_id, sampler_id });
2459 }
Hans-Kristian Arntzendd1513b2016-09-10 21:52:22 +02002460
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002461 return true;
2462}
2463
Hans-Kristian Arntzenbcb55602016-09-10 13:56:36 +02002464void Compiler::build_combined_image_samplers()
2465{
Hans-Kristian Arntzen71bacc42016-09-10 17:48:52 +02002466 combined_image_samplers.clear();
2467 CombinedImageSamplerHandler handler(*this);
2468 traverse_all_reachable_opcodes(get<SPIRFunction>(entry_point), handler);
Hans-Kristian Arntzenbcb55602016-09-10 13:56:36 +02002469}