1	/*
2	* Copyright 2015-2021 Arm Limited
3	* SPDX-License-Identifier: Apache-2.0 OR MIT
4	*
5	* Licensed under the Apache License, Version 2.0 (the "License");
6	* you may not use this file except in compliance with the License.
7	* You may obtain a copy of the License at
8	*
9	* http://www.apache.org/licenses/LICENSE-2.0
10	*
11	* Unless required by applicable law or agreed to in writing, software
12	* distributed under the License is distributed on an "AS IS" BASIS,
13	* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
14	* See the License for the specific language governing permissions and
15	* limitations under the License.
16	*/
17
18	/*
19	* At your option, you may choose to accept this material under either:
20	* 1. The Apache License, Version 2.0, found at <http://www.apache.org/licenses/LICENSE-2.0>, or
21	* 2. The MIT License, found at <http://opensource.org/licenses/MIT>.
22	*/
23
24	#include "spirv_cross.hpp"
25	#include "GLSL.std.450.h"
26	#include "spirv_cfg.hpp"
27	#include "spirv_common.hpp"
28	#include "spirv_parser.hpp"
29	#include <algorithm>
30	#include <cstring>
31	#include <utility>
32
33	using namespace std;
34	using namespace spv;
35	using namespace SPIRV_CROSS_NAMESPACE;
36
37	Compiler::Compiler(vector<uint32_t> ir_)
38	{
39	Parser parser(std::move(ir_));
40	parser.parse();
41	set_ir(std::move(parser.get_parsed_ir()));
42	}
43
44	Compiler::Compiler(const uint32_t *ir_, size_t word_count)
45	{
46	Parser parser(ir_, word_count);
47	parser.parse();
48	set_ir(std::move(parser.get_parsed_ir()));
49	}
50
51	Compiler::Compiler(const ParsedIR &ir_)
52	{
53	set_ir(ir_);
54	}
55
56	Compiler::Compiler(ParsedIR &&ir_)
57	{
58	set_ir(std::move(ir_));
59	}
60
61	void Compiler::set_ir(ParsedIR &&ir_)
62	{
63	ir = std::move(ir_);
64	parse_fixup();
65	}
66
67	void Compiler::set_ir(const ParsedIR &ir_)
68	{
69	ir = ir_;
70	parse_fixup();
71	}
72
73	string Compiler::compile()
74	{
75	return "";
76	}
77
78	bool Compiler::variable_storage_is_aliased(const SPIRVariable &v)
79	{
80	auto &type = get<SPIRType>(id: v.basetype);
81	bool ssbo = v.storage == StorageClassStorageBuffer ||
82	ir.meta[type.self].decoration.decoration_flags.get(bit: DecorationBufferBlock);
83	bool image = type.basetype == SPIRType::Image;
84	bool counter = type.basetype == SPIRType::AtomicCounter;
85	bool buffer_reference = type.storage == StorageClassPhysicalStorageBufferEXT;
86
87	bool is_restrict;
88	if (ssbo)
89	is_restrict = ir.get_buffer_block_flags(var: v).get(bit: DecorationRestrict);
90	else
91	is_restrict = has_decoration(id: v.self, decoration: DecorationRestrict);
92
93	return !is_restrict && (ssbo || image || counter || buffer_reference);
94	}
95
96	bool Compiler::block_is_control_dependent(const SPIRBlock &block)
97	{
98	for (auto &i : block.ops)
99	{
100	auto ops = stream(instr: i);
101	auto op = static_cast<Op>(i.op);
102
103	switch (op)
104	{
105	case OpFunctionCall:
106	{
107	uint32_t func = ops[2];
108	if (function_is_control_dependent(func: get<SPIRFunction>(id: func)))
109	return true;
110	break;
111	}
112
113	// Derivatives
114	case OpDPdx:
115	case OpDPdxCoarse:
116	case OpDPdxFine:
117	case OpDPdy:
118	case OpDPdyCoarse:
119	case OpDPdyFine:
120	case OpFwidth:
121	case OpFwidthCoarse:
122	case OpFwidthFine:
123
124	// Anything implicit LOD
125	case OpImageSampleImplicitLod:
126	case OpImageSampleDrefImplicitLod:
127	case OpImageSampleProjImplicitLod:
128	case OpImageSampleProjDrefImplicitLod:
129	case OpImageSparseSampleImplicitLod:
130	case OpImageSparseSampleDrefImplicitLod:
131	case OpImageSparseSampleProjImplicitLod:
132	case OpImageSparseSampleProjDrefImplicitLod:
133	case OpImageQueryLod:
134	case OpImageDrefGather:
135	case OpImageGather:
136	case OpImageSparseDrefGather:
137	case OpImageSparseGather:
138
139	// Anything subgroups
140	case OpGroupNonUniformElect:
141	case OpGroupNonUniformAll:
142	case OpGroupNonUniformAny:
143	case OpGroupNonUniformAllEqual:
144	case OpGroupNonUniformBroadcast:
145	case OpGroupNonUniformBroadcastFirst:
146	case OpGroupNonUniformBallot:
147	case OpGroupNonUniformInverseBallot:
148	case OpGroupNonUniformBallotBitExtract:
149	case OpGroupNonUniformBallotBitCount:
150	case OpGroupNonUniformBallotFindLSB:
151	case OpGroupNonUniformBallotFindMSB:
152	case OpGroupNonUniformShuffle:
153	case OpGroupNonUniformShuffleXor:
154	case OpGroupNonUniformShuffleUp:
155	case OpGroupNonUniformShuffleDown:
156	case OpGroupNonUniformIAdd:
157	case OpGroupNonUniformFAdd:
158	case OpGroupNonUniformIMul:
159	case OpGroupNonUniformFMul:
160	case OpGroupNonUniformSMin:
161	case OpGroupNonUniformUMin:
162	case OpGroupNonUniformFMin:
163	case OpGroupNonUniformSMax:
164	case OpGroupNonUniformUMax:
165	case OpGroupNonUniformFMax:
166	case OpGroupNonUniformBitwiseAnd:
167	case OpGroupNonUniformBitwiseOr:
168	case OpGroupNonUniformBitwiseXor:
169	case OpGroupNonUniformLogicalAnd:
170	case OpGroupNonUniformLogicalOr:
171	case OpGroupNonUniformLogicalXor:
172	case OpGroupNonUniformQuadBroadcast:
173	case OpGroupNonUniformQuadSwap:
174
175	// Control barriers
176	case OpControlBarrier:
177	return true;
178
179	default:
180	break;
181	}
182	}
183
184	return false;
185	}
186
187	bool Compiler::block_is_pure(const SPIRBlock &block)
188	{
189	// This is a global side effect of the function.
190	if (block.terminator == SPIRBlock::Kill ||
191	block.terminator == SPIRBlock::TerminateRay ||
192	block.terminator == SPIRBlock::IgnoreIntersection ||
193	block.terminator == SPIRBlock::EmitMeshTasks)
194	return false;
195
196	for (auto &i : block.ops)
197	{
198	auto ops = stream(instr: i);
199	auto op = static_cast<Op>(i.op);
200
201	switch (op)
202	{
203	case OpFunctionCall:
204	{
205	uint32_t func = ops[2];
206	if (!function_is_pure(func: get<SPIRFunction>(id: func)))
207	return false;
208	break;
209	}
210
211	case OpCopyMemory:
212	case OpStore:
213	{
214	auto &type = expression_type(id: ops[0]);
215	if (type.storage != StorageClassFunction)
216	return false;
217	break;
218	}
219
220	case OpImageWrite:
221	return false;
222
223	// Atomics are impure.
224	case OpAtomicLoad:
225	case OpAtomicStore:
226	case OpAtomicExchange:
227	case OpAtomicCompareExchange:
228	case OpAtomicCompareExchangeWeak:
229	case OpAtomicIIncrement:
230	case OpAtomicIDecrement:
231	case OpAtomicIAdd:
232	case OpAtomicISub:
233	case OpAtomicSMin:
234	case OpAtomicUMin:
235	case OpAtomicSMax:
236	case OpAtomicUMax:
237	case OpAtomicAnd:
238	case OpAtomicOr:
239	case OpAtomicXor:
240	return false;
241
242	// Geometry shader builtins modify global state.
243	case OpEndPrimitive:
244	case OpEmitStreamVertex:
245	case OpEndStreamPrimitive:
246	case OpEmitVertex:
247	return false;
248
249	// Mesh shader functions modify global state.
250	// (EmitMeshTasks is a terminator).
251	case OpSetMeshOutputsEXT:
252	return false;
253
254	// Barriers disallow any reordering, so we should treat blocks with barrier as writing.
255	case OpControlBarrier:
256	case OpMemoryBarrier:
257	return false;
258
259	// Ray tracing builtins are impure.
260	case OpReportIntersectionKHR:
261	case OpIgnoreIntersectionNV:
262	case OpTerminateRayNV:
263	case OpTraceNV:
264	case OpTraceRayKHR:
265	case OpExecuteCallableNV:
266	case OpExecuteCallableKHR:
267	case OpRayQueryInitializeKHR:
268	case OpRayQueryTerminateKHR:
269	case OpRayQueryGenerateIntersectionKHR:
270	case OpRayQueryConfirmIntersectionKHR:
271	case OpRayQueryProceedKHR:
272	// There are various getters in ray query, but they are considered pure.
273	return false;
274
275	// OpExtInst is potentially impure depending on extension, but GLSL builtins are at least pure.
276
277	case OpDemoteToHelperInvocationEXT:
278	// This is a global side effect of the function.
279	return false;
280
281	case OpExtInst:
282	{
283	uint32_t extension_set = ops[2];
284	if (get<SPIRExtension>(id: extension_set).ext == SPIRExtension::GLSL)
285	{
286	auto op_450 = static_cast<GLSLstd450>(ops[3]);
287	switch (op_450)
288	{
289	case GLSLstd450Modf:
290	case GLSLstd450Frexp:
291	{
292	auto &type = expression_type(id: ops[5]);
293	if (type.storage != StorageClassFunction)
294	return false;
295	break;
296	}
297
298	default:
299	break;
300	}
301	}
302	break;
303	}
304
305	default:
306	break;
307	}
308	}
309
310	return true;
311	}
312
313	string Compiler::to_name(uint32_t id, bool allow_alias) const
314	{
315	if (allow_alias && ir.ids[id].get_type() == TypeType)
316	{
317	// If this type is a simple alias, emit the
318	// name of the original type instead.
319	// We don't want to override the meta alias
320	// as that can be overridden by the reflection APIs after parse.
321	auto &type = get<SPIRType>(id);
322	if (type.type_alias)
323	{
324	// If the alias master has been specially packed, we will have emitted a clean variant as well,
325	// so skip the name aliasing here.
326	if (!has_extended_decoration(id: type.type_alias, decoration: SPIRVCrossDecorationBufferBlockRepacked))
327	return to_name(id: type.type_alias);
328	}
329	}
330
331	auto &alias = ir.get_name(id);
332	if (alias.empty())
333	return join(ts: "_", ts&: id);
334	else
335	return alias;
336	}
337
338	bool Compiler::function_is_pure(const SPIRFunction &func)
339	{
340	for (auto block : func.blocks)
341	if (!block_is_pure(block: get<SPIRBlock>(id: block)))
342	return false;
343
344	return true;
345	}
346
347	bool Compiler::function_is_control_dependent(const SPIRFunction &func)
348	{
349	for (auto block : func.blocks)
350	if (block_is_control_dependent(block: get<SPIRBlock>(id: block)))
351	return true;
352
353	return false;
354	}
355
356	void Compiler::register_global_read_dependencies(const SPIRBlock &block, uint32_t id)
357	{
358	for (auto &i : block.ops)
359	{
360	auto ops = stream(instr: i);
361	auto op = static_cast<Op>(i.op);
362
363	switch (op)
364	{
365	case OpFunctionCall:
366	{
367	uint32_t func = ops[2];
368	register_global_read_dependencies(func: get<SPIRFunction>(id: func), id);
369	break;
370	}
371
372	case OpLoad:
373	case OpImageRead:
374	{
375	// If we're in a storage class which does not get invalidated, adding dependencies here is no big deal.
376	auto *var = maybe_get_backing_variable(chain: ops[2]);
377	if (var && var->storage != StorageClassFunction)
378	{
379	auto &type = get<SPIRType>(id: var->basetype);
380
381	// InputTargets are immutable.
382	if (type.basetype != SPIRType::Image && type.image.dim != DimSubpassData)
383	var->dependees.push_back(t: id);
384	}
385	break;
386	}
387
388	default:
389	break;
390	}
391	}
392	}
393
394	void Compiler::register_global_read_dependencies(const SPIRFunction &func, uint32_t id)
395	{
396	for (auto block : func.blocks)
397	register_global_read_dependencies(block: get<SPIRBlock>(id: block), id);
398	}
399
400	SPIRVariable *Compiler::maybe_get_backing_variable(uint32_t chain)
401	{
402	auto *var = maybe_get<SPIRVariable>(id: chain);
403	if (!var)
404	{
405	auto *cexpr = maybe_get<SPIRExpression>(id: chain);
406	if (cexpr)
407	var = maybe_get<SPIRVariable>(id: cexpr->loaded_from);
408
409	auto *access_chain = maybe_get<SPIRAccessChain>(id: chain);
410	if (access_chain)
411	var = maybe_get<SPIRVariable>(id: access_chain->loaded_from);
412	}
413
414	return var;
415	}
416
417	void Compiler::register_read(uint32_t expr, uint32_t chain, bool forwarded)
418	{
419	auto &e = get<SPIRExpression>(id: expr);
420	auto *var = maybe_get_backing_variable(chain);
421
422	if (var)
423	{
424	e.loaded_from = var->self;
425
426	// If the backing variable is immutable, we do not need to depend on the variable.
427	if (forwarded && !is_immutable(id: var->self))
428	var->dependees.push_back(t: e.self);
429
430	// If we load from a parameter, make sure we create "inout" if we also write to the parameter.
431	// The default is "in" however, so we never invalidate our compilation by reading.
432	if (var && var->parameter)
433	var->parameter->read_count++;
434	}
435	}
436
437	void Compiler::register_write(uint32_t chain)
438	{
439	auto *var = maybe_get<SPIRVariable>(id: chain);
440	if (!var)
441	{
442	// If we're storing through an access chain, invalidate the backing variable instead.
443	auto *expr = maybe_get<SPIRExpression>(id: chain);
444	if (expr && expr->loaded_from)
445	var = maybe_get<SPIRVariable>(id: expr->loaded_from);
446
447	auto *access_chain = maybe_get<SPIRAccessChain>(id: chain);
448	if (access_chain && access_chain->loaded_from)
449	var = maybe_get<SPIRVariable>(id: access_chain->loaded_from);
450	}
451
452	auto &chain_type = expression_type(id: chain);
453
454	if (var)
455	{
456	bool check_argument_storage_qualifier = true;
457	auto &type = expression_type(id: chain);
458
459	// If our variable is in a storage class which can alias with other buffers,
460	// invalidate all variables which depend on aliased variables. And if this is a
461	// variable pointer, then invalidate all variables regardless.
462	if (get_variable_data_type(var: *var).pointer)
463	{
464	flush_all_active_variables();
465
466	if (type.pointer_depth == 1)
467	{
468	// We have a backing variable which is a pointer-to-pointer type.
469	// We are storing some data through a pointer acquired through that variable,
470	// but we are not writing to the value of the variable itself,
471	// i.e., we are not modifying the pointer directly.
472	// If we are storing a non-pointer type (pointer_depth == 1),
473	// we know that we are storing some unrelated data.
474	// A case here would be
475	// void foo(Foo * const *arg) {
476	// Foo *bar = *arg;
477	// bar->unrelated = 42;
478	// }
479	// arg, the argument is constant.
480	check_argument_storage_qualifier = false;
481	}
482	}
483
484	if (type.storage == StorageClassPhysicalStorageBufferEXT || variable_storage_is_aliased(v: *var))
485	flush_all_aliased_variables();
486	else if (var)
487	flush_dependees(var&: *var);
488
489	// We tried to write to a parameter which is not marked with out qualifier, force a recompile.
490	if (check_argument_storage_qualifier && var->parameter && var->parameter->write_count == 0)
491	{
492	var->parameter->write_count++;
493	force_recompile();
494	}
495	}
496	else if (chain_type.pointer)
497	{
498	// If we stored through a variable pointer, then we don't know which
499	// variable we stored to. So *all* expressions after this point need to
500	// be invalidated.
501	// FIXME: If we can prove that the variable pointer will point to
502	// only certain variables, we can invalidate only those.
503	flush_all_active_variables();
504	}
505
506	// If chain_type.pointer is false, we're not writing to memory backed variables, but temporaries instead.
507	// This can happen in copy_logical_type where we unroll complex reads and writes to temporaries.
508	}
509
510	void Compiler::flush_dependees(SPIRVariable &var)
511	{
512	for (auto expr : var.dependees)
513	invalid_expressions.insert(x: expr);
514	var.dependees.clear();
515	}
516
517	void Compiler::flush_all_aliased_variables()
518	{
519	for (auto aliased : aliased_variables)
520	flush_dependees(var&: get<SPIRVariable>(id: aliased));
521	}
522
523	void Compiler::flush_all_atomic_capable_variables()
524	{
525	for (auto global : global_variables)
526	flush_dependees(var&: get<SPIRVariable>(id: global));
527	flush_all_aliased_variables();
528	}
529
530	void Compiler::flush_control_dependent_expressions(uint32_t block_id)
531	{
532	auto &block = get<SPIRBlock>(id: block_id);
533	for (auto &expr : block.invalidate_expressions)
534	invalid_expressions.insert(x: expr);
535	block.invalidate_expressions.clear();
536	}
537
538	void Compiler::flush_all_active_variables()
539	{
540	// Invalidate all temporaries we read from variables in this block since they were forwarded.
541	// Invalidate all temporaries we read from globals.
542	for (auto &v : current_function->local_variables)
543	flush_dependees(var&: get<SPIRVariable>(id: v));
544	for (auto &arg : current_function->arguments)
545	flush_dependees(var&: get<SPIRVariable>(id: arg.id));
546	for (auto global : global_variables)
547	flush_dependees(var&: get<SPIRVariable>(id: global));
548
549	flush_all_aliased_variables();
550	}
551
552	uint32_t Compiler::expression_type_id(uint32_t id) const
553	{
554	switch (ir.ids[id].get_type())
555	{
556	case TypeVariable:
557	return get<SPIRVariable>(id).basetype;
558
559	case TypeExpression:
560	return get<SPIRExpression>(id).expression_type;
561
562	case TypeConstant:
563	return get<SPIRConstant>(id).constant_type;
564
565	case TypeConstantOp:
566	return get<SPIRConstantOp>(id).basetype;
567
568	case TypeUndef:
569	return get<SPIRUndef>(id).basetype;
570
571	case TypeCombinedImageSampler:
572	return get<SPIRCombinedImageSampler>(id).combined_type;
573
574	case TypeAccessChain:
575	return get<SPIRAccessChain>(id).basetype;
576
577	default:
578	SPIRV_CROSS_THROW("Cannot resolve expression type.");
579	}
580	}
581
582	const SPIRType &Compiler::expression_type(uint32_t id) const
583	{
584	return get<SPIRType>(id: expression_type_id(id));
585	}
586
587	bool Compiler::expression_is_lvalue(uint32_t id) const
588	{
589	auto &type = expression_type(id);
590	switch (type.basetype)
591	{
592	case SPIRType::SampledImage:
593	case SPIRType::Image:
594	case SPIRType::Sampler:
595	return false;
596
597	default:
598	return true;
599	}
600	}
601
602	bool Compiler::is_immutable(uint32_t id) const
603	{
604	if (ir.ids[id].get_type() == TypeVariable)
605	{
606	auto &var = get<SPIRVariable>(id);
607
608	// Anything we load from the UniformConstant address space is guaranteed to be immutable.
609	bool pointer_to_const = var.storage == StorageClassUniformConstant;
610	return pointer_to_const || var.phi_variable || !expression_is_lvalue(id);
611	}
612	else if (ir.ids[id].get_type() == TypeAccessChain)
613	return get<SPIRAccessChain>(id).immutable;
614	else if (ir.ids[id].get_type() == TypeExpression)
615	return get<SPIRExpression>(id).immutable;
616	else if (ir.ids[id].get_type() == TypeConstant || ir.ids[id].get_type() == TypeConstantOp ||
617	ir.ids[id].get_type() == TypeUndef)
618	return true;
619	else
620	return false;
621	}
622
623	static inline bool storage_class_is_interface(spv::StorageClass storage)
624	{
625	switch (storage)
626	{
627	case StorageClassInput:
628	case StorageClassOutput:
629	case StorageClassUniform:
630	case StorageClassUniformConstant:
631	case StorageClassAtomicCounter:
632	case StorageClassPushConstant:
633	case StorageClassStorageBuffer:
634	return true;
635
636	default:
637	return false;
638	}
639	}
640
641	bool Compiler::is_hidden_variable(const SPIRVariable &var, bool include_builtins) const
642	{
643	if ((is_builtin_variable(var) && !include_builtins) || var.remapped_variable)
644	return true;
645
646	// Combined image samplers are always considered active as they are "magic" variables.
647	if (find_if(first: begin(cont: combined_image_samplers), last: end(cont: combined_image_samplers), pred: [&var](const CombinedImageSampler &samp) {
648	return samp.combined_id == var.self;
649	}) != end(cont: combined_image_samplers))
650	{
651	return false;
652	}
653
654	// In SPIR-V 1.4 and up we must also use the active variable interface to disable global variables
655	// which are not part of the entry point.
656	if (ir.get_spirv_version() >= 0x10400 && var.storage != spv::StorageClassGeneric &&
657	var.storage != spv::StorageClassFunction && !interface_variable_exists_in_entry_point(id: var.self))
658	{
659	return true;
660	}
661
662	return check_active_interface_variables && storage_class_is_interface(storage: var.storage) &&
663	active_interface_variables.find(x: var.self) == end(cont: active_interface_variables);
664	}
665
666	bool Compiler::is_builtin_type(const SPIRType &type) const
667	{
668	auto *type_meta = ir.find_meta(id: type.self);
669
670	// We can have builtin structs as well. If one member of a struct is builtin, the struct must also be builtin.
671	if (type_meta)
672	for (auto &m : type_meta->members)
673	if (m.builtin)
674	return true;
675
676	return false;
677	}
678
679	bool Compiler::is_builtin_variable(const SPIRVariable &var) const
680	{
681	auto *m = ir.find_meta(id: var.self);
682
683	if (var.compat_builtin || (m && m->decoration.builtin))
684	return true;
685	else
686	return is_builtin_type(type: get<SPIRType>(id: var.basetype));
687	}
688
689	bool Compiler::is_member_builtin(const SPIRType &type, uint32_t index, BuiltIn *builtin) const
690	{
691	auto *type_meta = ir.find_meta(id: type.self);
692
693	if (type_meta)
694	{
695	auto &memb = type_meta->members;
696	if (index < memb.size() && memb[index].builtin)
697	{
698	if (builtin)
699	*builtin = memb[index].builtin_type;
700	return true;
701	}
702	}
703
704	return false;
705	}
706
707	bool Compiler::is_scalar(const SPIRType &type) const
708	{
709	return type.basetype != SPIRType::Struct && type.vecsize == 1 && type.columns == 1;
710	}
711
712	bool Compiler::is_vector(const SPIRType &type) const
713	{
714	return type.vecsize > 1 && type.columns == 1;
715	}
716
717	bool Compiler::is_matrix(const SPIRType &type) const
718	{
719	return type.vecsize > 1 && type.columns > 1;
720	}
721
722	bool Compiler::is_array(const SPIRType &type) const
723	{
724	return type.op == OpTypeArray || type.op == OpTypeRuntimeArray;
725	}
726
727	bool Compiler::is_pointer(const SPIRType &type) const
728	{
729	return type.op == OpTypePointer && type.basetype != SPIRType::Unknown; // Ignore function pointers.
730	}
731
732	bool Compiler::is_physical_pointer(const SPIRType &type) const
733	{
734	return type.op == OpTypePointer && type.storage == StorageClassPhysicalStorageBuffer;
735	}
736
737	bool Compiler::is_physical_pointer_to_buffer_block(const SPIRType &type) const
738	{
739	return is_physical_pointer(type) && get_pointee_type(type).self == type.parent_type &&
740	(has_decoration(id: type.self, decoration: DecorationBlock) ||
741	has_decoration(id: type.self, decoration: DecorationBufferBlock));
742	}
743
744	bool Compiler::is_runtime_size_array(const SPIRType &type)
745	{
746	return type.op == OpTypeRuntimeArray;
747	}
748
749	ShaderResources Compiler::get_shader_resources() const
750	{
751	return get_shader_resources(active_variables: nullptr);
752	}
753
754	ShaderResources Compiler::get_shader_resources(const unordered_set<VariableID> &active_variables) const
755	{
756	return get_shader_resources(active_variables: &active_variables);
757	}
758
759	bool Compiler::InterfaceVariableAccessHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
760	{
761	uint32_t variable = 0;
762	switch (opcode)
763	{
764	// Need this first, otherwise, GCC complains about unhandled switch statements.
765	default:
766	break;
767
768	case OpFunctionCall:
769	{
770	// Invalid SPIR-V.
771	if (length < 3)
772	return false;
773
774	uint32_t count = length - 3;
775	args += 3;
776	for (uint32_t i = 0; i < count; i++)
777	{
778	auto *var = compiler.maybe_get<SPIRVariable>(id: args[i]);
779	if (var && storage_class_is_interface(storage: var->storage))
780	variables.insert(x: args[i]);
781	}
782	break;
783	}
784
785	case OpSelect:
786	{
787	// Invalid SPIR-V.
788	if (length < 5)
789	return false;
790
791	uint32_t count = length - 3;
792	args += 3;
793	for (uint32_t i = 0; i < count; i++)
794	{
795	auto *var = compiler.maybe_get<SPIRVariable>(id: args[i]);
796	if (var && storage_class_is_interface(storage: var->storage))
797	variables.insert(x: args[i]);
798	}
799	break;
800	}
801
802	case OpPhi:
803	{
804	// Invalid SPIR-V.
805	if (length < 2)
806	return false;
807
808	uint32_t count = length - 2;
809	args += 2;
810	for (uint32_t i = 0; i < count; i += 2)
811	{
812	auto *var = compiler.maybe_get<SPIRVariable>(id: args[i]);
813	if (var && storage_class_is_interface(storage: var->storage))
814	variables.insert(x: args[i]);
815	}
816	break;
817	}
818
819	case OpAtomicStore:
820	case OpStore:
821	// Invalid SPIR-V.
822	if (length < 1)
823	return false;
824	variable = args[0];
825	break;
826
827	case OpCopyMemory:
828	{
829	if (length < 2)
830	return false;
831
832	auto *var = compiler.maybe_get<SPIRVariable>(id: args[0]);
833	if (var && storage_class_is_interface(storage: var->storage))
834	variables.insert(x: args[0]);
835
836	var = compiler.maybe_get<SPIRVariable>(id: args[1]);
837	if (var && storage_class_is_interface(storage: var->storage))
838	variables.insert(x: args[1]);
839	break;
840	}
841
842	case OpExtInst:
843	{
844	if (length < 3)
845	return false;
846	auto &extension_set = compiler.get<SPIRExtension>(id: args[2]);
847	switch (extension_set.ext)
848	{
849	case SPIRExtension::GLSL:
850	{
851	auto op = static_cast<GLSLstd450>(args[3]);
852
853	switch (op)
854	{
855	case GLSLstd450InterpolateAtCentroid:
856	case GLSLstd450InterpolateAtSample:
857	case GLSLstd450InterpolateAtOffset:
858	{
859	auto *var = compiler.maybe_get<SPIRVariable>(id: args[4]);
860	if (var && storage_class_is_interface(storage: var->storage))
861	variables.insert(x: args[4]);
862	break;
863	}
864
865	case GLSLstd450Modf:
866	case GLSLstd450Fract:
867	{
868	auto *var = compiler.maybe_get<SPIRVariable>(id: args[5]);
869	if (var && storage_class_is_interface(storage: var->storage))
870	variables.insert(x: args[5]);
871	break;
872	}
873
874	default:
875	break;
876	}
877	break;
878	}
879	case SPIRExtension::SPV_AMD_shader_explicit_vertex_parameter:
880	{
881	enum AMDShaderExplicitVertexParameter
882	{
883	InterpolateAtVertexAMD = 1
884	};
885
886	auto op = static_cast<AMDShaderExplicitVertexParameter>(args[3]);
887
888	switch (op)
889	{
890	case InterpolateAtVertexAMD:
891	{
892	auto *var = compiler.maybe_get<SPIRVariable>(id: args[4]);
893	if (var && storage_class_is_interface(storage: var->storage))
894	variables.insert(x: args[4]);
895	break;
896	}
897
898	default:
899	break;
900	}
901	break;
902	}
903	default:
904	break;
905	}
906	break;
907	}
908
909	case OpAccessChain:
910	case OpInBoundsAccessChain:
911	case OpPtrAccessChain:
912	case OpLoad:
913	case OpCopyObject:
914	case OpImageTexelPointer:
915	case OpAtomicLoad:
916	case OpAtomicExchange:
917	case OpAtomicCompareExchange:
918	case OpAtomicCompareExchangeWeak:
919	case OpAtomicIIncrement:
920	case OpAtomicIDecrement:
921	case OpAtomicIAdd:
922	case OpAtomicISub:
923	case OpAtomicSMin:
924	case OpAtomicUMin:
925	case OpAtomicSMax:
926	case OpAtomicUMax:
927	case OpAtomicAnd:
928	case OpAtomicOr:
929	case OpAtomicXor:
930	case OpArrayLength:
931	// Invalid SPIR-V.
932	if (length < 3)
933	return false;
934	variable = args[2];
935	break;
936	}
937
938	if (variable)
939	{
940	auto *var = compiler.maybe_get<SPIRVariable>(id: variable);
941	if (var && storage_class_is_interface(storage: var->storage))
942	variables.insert(x: variable);
943	}
944	return true;
945	}
946
947	unordered_set<VariableID> Compiler::get_active_interface_variables() const
948	{
949	// Traverse the call graph and find all interface variables which are in use.
950	unordered_set<VariableID> variables;
951	InterfaceVariableAccessHandler handler(*this, variables);
952	traverse_all_reachable_opcodes(block: get<SPIRFunction>(id: ir.default_entry_point), handler);
953
954	ir.for_each_typed_id<SPIRVariable>(op: [&](uint32_t, const SPIRVariable &var) {
955	if (var.storage != StorageClassOutput)
956	return;
957	if (!interface_variable_exists_in_entry_point(id: var.self))
958	return;
959
960	// An output variable which is just declared (but uninitialized) might be read by subsequent stages
961	// so we should force-enable these outputs,
962	// since compilation will fail if a subsequent stage attempts to read from the variable in question.
963	// Also, make sure we preserve output variables which are only initialized, but never accessed by any code.
964	if (var.initializer != ID(0) || get_execution_model() != ExecutionModelFragment)
965	variables.insert(x: var.self);
966	});
967
968	// If we needed to create one, we'll need it.
969	if (dummy_sampler_id)
970	variables.insert(x: dummy_sampler_id);
971
972	return variables;
973	}
974
975	void Compiler::set_enabled_interface_variables(std::unordered_set<VariableID> active_variables)
976	{
977	active_interface_variables = std::move(active_variables);
978	check_active_interface_variables = true;
979	}
980
981	ShaderResources Compiler::get_shader_resources(const unordered_set<VariableID> *active_variables) const
982	{
983	ShaderResources res;
984
985	bool ssbo_instance_name = reflection_ssbo_instance_name_is_significant();
986
987	ir.for_each_typed_id<SPIRVariable>(op: [&](uint32_t, const SPIRVariable &var) {
988	auto &type = this->get<SPIRType>(id: var.basetype);
989
990	// It is possible for uniform storage classes to be passed as function parameters, so detect
991	// that. To detect function parameters, check of StorageClass of variable is function scope.
992	if (var.storage == StorageClassFunction || !type.pointer)
993	return;
994
995	if (active_variables && active_variables->find(x: var.self) == end(cont: *active_variables))
996	return;
997
998	// In SPIR-V 1.4 and up, every global must be present in the entry point interface list,
999	// not just IO variables.
1000	bool active_in_entry_point = true;
1001	if (ir.get_spirv_version() < 0x10400)
1002	{
1003	if (var.storage == StorageClassInput || var.storage == StorageClassOutput)
1004	active_in_entry_point = interface_variable_exists_in_entry_point(id: var.self);
1005	}
1006	else
1007	active_in_entry_point = interface_variable_exists_in_entry_point(id: var.self);
1008
1009	if (!active_in_entry_point)
1010	return;
1011
1012	bool is_builtin = is_builtin_variable(var);
1013
1014	if (is_builtin)
1015	{
1016	if (var.storage != StorageClassInput && var.storage != StorageClassOutput)
1017	return;
1018
1019	auto &list = var.storage == StorageClassInput ? res.builtin_inputs : res.builtin_outputs;
1020	BuiltInResource resource;
1021
1022	if (has_decoration(id: type.self, decoration: DecorationBlock))
1023	{
1024	resource.resource = { .id: var.self, .type_id: var.basetype, .base_type_id: type.self,
1025	.name: get_remapped_declared_block_name(id: var.self, fallback_prefer_instance_name: false) };
1026
1027	for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++)
1028	{
1029	resource.value_type_id = type.member_types[i];
1030	resource.builtin = BuiltIn(get_member_decoration(id: type.self, index: i, decoration: DecorationBuiltIn));
1031	list.push_back(t: resource);
1032	}
1033	}
1034	else
1035	{
1036	bool strip_array =
1037	!has_decoration(id: var.self, decoration: DecorationPatch) && (
1038	get_execution_model() == ExecutionModelTessellationControl ||
1039	(get_execution_model() == ExecutionModelTessellationEvaluation &&
1040	var.storage == StorageClassInput));
1041
1042	resource.resource = { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) };
1043
1044	if (strip_array && !type.array.empty())
1045	resource.value_type_id = get_variable_data_type(var).parent_type;
1046	else
1047	resource.value_type_id = get_variable_data_type_id(var);
1048
1049	assert(resource.value_type_id);
1050
1051	resource.builtin = BuiltIn(get_decoration(id: var.self, decoration: DecorationBuiltIn));
1052	list.push_back(t: std::move(resource));
1053	}
1054	return;
1055	}
1056
1057	// Input
1058	if (var.storage == StorageClassInput)
1059	{
1060	if (has_decoration(id: type.self, decoration: DecorationBlock))
1061	{
1062	res.stage_inputs.push_back(
1063	t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self,
1064	.name: get_remapped_declared_block_name(id: var.self, fallback_prefer_instance_name: false) });
1065	}
1066	else
1067	res.stage_inputs.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) });
1068	}
1069	// Subpass inputs
1070	else if (var.storage == StorageClassUniformConstant && type.image.dim == DimSubpassData)
1071	{
1072	res.subpass_inputs.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) });
1073	}
1074	// Outputs
1075	else if (var.storage == StorageClassOutput)
1076	{
1077	if (has_decoration(id: type.self, decoration: DecorationBlock))
1078	{
1079	res.stage_outputs.push_back(
1080	t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_remapped_declared_block_name(id: var.self, fallback_prefer_instance_name: false) });
1081	}
1082	else
1083	res.stage_outputs.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) });
1084	}
1085	// UBOs
1086	else if (type.storage == StorageClassUniform && has_decoration(id: type.self, decoration: DecorationBlock))
1087	{
1088	res.uniform_buffers.push_back(
1089	t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_remapped_declared_block_name(id: var.self, fallback_prefer_instance_name: false) });
1090	}
1091	// Old way to declare SSBOs.
1092	else if (type.storage == StorageClassUniform && has_decoration(id: type.self, decoration: DecorationBufferBlock))
1093	{
1094	res.storage_buffers.push_back(
1095	t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_remapped_declared_block_name(id: var.self, fallback_prefer_instance_name: ssbo_instance_name) });
1096	}
1097	// Modern way to declare SSBOs.
1098	else if (type.storage == StorageClassStorageBuffer)
1099	{
1100	res.storage_buffers.push_back(
1101	t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_remapped_declared_block_name(id: var.self, fallback_prefer_instance_name: ssbo_instance_name) });
1102	}
1103	// Push constant blocks
1104	else if (type.storage == StorageClassPushConstant)
1105	{
1106	// There can only be one push constant block, but keep the vector in case this restriction is lifted
1107	// in the future.
1108	res.push_constant_buffers.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) });
1109	}
1110	else if (type.storage == StorageClassShaderRecordBufferKHR)
1111	{
1112	res.shader_record_buffers.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_remapped_declared_block_name(id: var.self, fallback_prefer_instance_name: ssbo_instance_name) });
1113	}
1114	// Atomic counters
1115	else if (type.storage == StorageClassAtomicCounter)
1116	{
1117	res.atomic_counters.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) });
1118	}
1119	else if (type.storage == StorageClassUniformConstant)
1120	{
1121	if (type.basetype == SPIRType::Image)
1122	{
1123	// Images
1124	if (type.image.sampled == 2)
1125	{
1126	res.storage_images.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) });
1127	}
1128	// Separate images
1129	else if (type.image.sampled == 1)
1130	{
1131	res.separate_images.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) });
1132	}
1133	}
1134	// Separate samplers
1135	else if (type.basetype == SPIRType::Sampler)
1136	{
1137	res.separate_samplers.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) });
1138	}
1139	// Textures
1140	else if (type.basetype == SPIRType::SampledImage)
1141	{
1142	res.sampled_images.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) });
1143	}
1144	// Acceleration structures
1145	else if (type.basetype == SPIRType::AccelerationStructure)
1146	{
1147	res.acceleration_structures.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) });
1148	}
1149	else
1150	{
1151	res.gl_plain_uniforms.push_back(t: { .id: var.self, .type_id: var.basetype, .base_type_id: type.self, .name: get_name(id: var.self) });
1152	}
1153	}
1154	});
1155
1156	return res;
1157	}
1158
1159	bool Compiler::type_is_top_level_block(const SPIRType &type) const
1160	{
1161	if (type.basetype != SPIRType::Struct)
1162	return false;
1163	return has_decoration(id: type.self, decoration: DecorationBlock) || has_decoration(id: type.self, decoration: DecorationBufferBlock);
1164	}
1165
1166	bool Compiler::type_is_block_like(const SPIRType &type) const
1167	{
1168	if (type_is_top_level_block(type))
1169	return true;
1170
1171	if (type.basetype == SPIRType::Struct)
1172	{
1173	// Block-like types may have Offset decorations.
1174	for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++)
1175	if (has_member_decoration(id: type.self, index: i, decoration: DecorationOffset))
1176	return true;
1177	}
1178
1179	return false;
1180	}
1181
1182	void Compiler::parse_fixup()
1183	{
1184	// Figure out specialization constants for work group sizes.
1185	for (auto id_ : ir.ids_for_constant_or_variable)
1186	{
1187	auto &id = ir.ids[id_];
1188
1189	if (id.get_type() == TypeConstant)
1190	{
1191	auto &c = id.get<SPIRConstant>();
1192	if (has_decoration(id: c.self, decoration: DecorationBuiltIn) &&
1193	BuiltIn(get_decoration(id: c.self, decoration: DecorationBuiltIn)) == BuiltInWorkgroupSize)
1194	{
1195	// In current SPIR-V, there can be just one constant like this.
1196	// All entry points will receive the constant value.
1197	// WorkgroupSize take precedence over LocalSizeId.
1198	for (auto &entry : ir.entry_points)
1199	{
1200	entry.second.workgroup_size.constant = c.self;
1201	entry.second.workgroup_size.x = c.scalar(col: 0, row: 0);
1202	entry.second.workgroup_size.y = c.scalar(col: 0, row: 1);
1203	entry.second.workgroup_size.z = c.scalar(col: 0, row: 2);
1204	}
1205	}
1206	}
1207	else if (id.get_type() == TypeVariable)
1208	{
1209	auto &var = id.get<SPIRVariable>();
1210	if (var.storage == StorageClassPrivate || var.storage == StorageClassWorkgroup ||
1211	var.storage == StorageClassTaskPayloadWorkgroupEXT ||
1212	var.storage == StorageClassOutput)
1213	{
1214	global_variables.push_back(t: var.self);
1215	}
1216	if (variable_storage_is_aliased(v: var))
1217	aliased_variables.push_back(t: var.self);
1218	}
1219	}
1220	}
1221
1222	void Compiler::update_name_cache(unordered_set<string> &cache_primary, const unordered_set<string> &cache_secondary,
1223	string &name)
1224	{
1225	if (name.empty())
1226	return;
1227
1228	const auto find_name = [&](const string &n) -> bool {
1229	if (cache_primary.find(x: n) != end(cont&: cache_primary))
1230	return true;
1231
1232	if (&cache_primary != &cache_secondary)
1233	if (cache_secondary.find(x: n) != end(cont: cache_secondary))
1234	return true;
1235
1236	return false;
1237	};
1238
1239	const auto insert_name = [&](const string &n) { cache_primary.insert(x: n); };
1240
1241	if (!find_name(name))
1242	{
1243	insert_name(name);
1244	return;
1245	}
1246
1247	uint32_t counter = 0;
1248	auto tmpname = name;
1249
1250	bool use_linked_underscore = true;
1251
1252	if (tmpname == "_")
1253	{
1254	// We cannot just append numbers, as we will end up creating internally reserved names.
1255	// Make it like _0_<counter> instead.
1256	tmpname += "0";
1257	}
1258	else if (tmpname.back() == '_')
1259	{
1260	// The last_character is an underscore, so we don't need to link in underscore.
1261	// This would violate double underscore rules.
1262	use_linked_underscore = false;
1263	}
1264
1265	// If there is a collision (very rare),
1266	// keep tacking on extra identifier until it's unique.
1267	do
1268	{
1269	counter++;
1270	name = tmpname + (use_linked_underscore ? "_" : "") + convert_to_string(t: counter);
1271	} while (find_name(name));
1272	insert_name(name);
1273	}
1274
1275	void Compiler::update_name_cache(unordered_set<string> &cache, string &name)
1276	{
1277	update_name_cache(cache_primary&: cache, cache_secondary: cache, name);
1278	}
1279
1280	void Compiler::set_name(ID id, const std::string &name)
1281	{
1282	ir.set_name(id, name);
1283	}
1284
1285	const SPIRType &Compiler::get_type(TypeID id) const
1286	{
1287	return get<SPIRType>(id);
1288	}
1289
1290	const SPIRType &Compiler::get_type_from_variable(VariableID id) const
1291	{
1292	return get<SPIRType>(id: get<SPIRVariable>(id).basetype);
1293	}
1294
1295	uint32_t Compiler::get_pointee_type_id(uint32_t type_id) const
1296	{
1297	auto *p_type = &get<SPIRType>(id: type_id);
1298	if (p_type->pointer)
1299	{
1300	assert(p_type->parent_type);
1301	type_id = p_type->parent_type;
1302	}
1303	return type_id;
1304	}
1305
1306	const SPIRType &Compiler::get_pointee_type(const SPIRType &type) const
1307	{
1308	auto *p_type = &type;
1309	if (p_type->pointer)
1310	{
1311	assert(p_type->parent_type);
1312	p_type = &get<SPIRType>(id: p_type->parent_type);
1313	}
1314	return *p_type;
1315	}
1316
1317	const SPIRType &Compiler::get_pointee_type(uint32_t type_id) const
1318	{
1319	return get_pointee_type(type: get<SPIRType>(id: type_id));
1320	}
1321
1322	uint32_t Compiler::get_variable_data_type_id(const SPIRVariable &var) const
1323	{
1324	if (var.phi_variable || var.storage == spv::StorageClass::StorageClassAtomicCounter)
1325	return var.basetype;
1326	return get_pointee_type_id(type_id: var.basetype);
1327	}
1328
1329	SPIRType &Compiler::get_variable_data_type(const SPIRVariable &var)
1330	{
1331	return get<SPIRType>(id: get_variable_data_type_id(var));
1332	}
1333
1334	const SPIRType &Compiler::get_variable_data_type(const SPIRVariable &var) const
1335	{
1336	return get<SPIRType>(id: get_variable_data_type_id(var));
1337	}
1338
1339	SPIRType &Compiler::get_variable_element_type(const SPIRVariable &var)
1340	{
1341	SPIRType *type = &get_variable_data_type(var);
1342	if (is_array(type: *type))
1343	type = &get<SPIRType>(id: type->parent_type);
1344	return *type;
1345	}
1346
1347	const SPIRType &Compiler::get_variable_element_type(const SPIRVariable &var) const
1348	{
1349	const SPIRType *type = &get_variable_data_type(var);
1350	if (is_array(type: *type))
1351	type = &get<SPIRType>(id: type->parent_type);
1352	return *type;
1353	}
1354
1355	bool Compiler::is_sampled_image_type(const SPIRType &type)
1356	{
1357	return (type.basetype == SPIRType::Image || type.basetype == SPIRType::SampledImage) && type.image.sampled == 1 &&
1358	type.image.dim != DimBuffer;
1359	}
1360
1361	void Compiler::set_member_decoration_string(TypeID id, uint32_t index, spv::Decoration decoration,
1362	const std::string &argument)
1363	{
1364	ir.set_member_decoration_string(id, index, decoration, argument);
1365	}
1366
1367	void Compiler::set_member_decoration(TypeID id, uint32_t index, Decoration decoration, uint32_t argument)
1368	{
1369	ir.set_member_decoration(id, index, decoration, argument);
1370	}
1371
1372	void Compiler::set_member_name(TypeID id, uint32_t index, const std::string &name)
1373	{
1374	ir.set_member_name(id, index, name);
1375	}
1376
1377	const std::string &Compiler::get_member_name(TypeID id, uint32_t index) const
1378	{
1379	return ir.get_member_name(id, index);
1380	}
1381
1382	void Compiler::set_qualified_name(uint32_t id, const string &name)
1383	{
1384	ir.meta[id].decoration.qualified_alias = name;
1385	}
1386
1387	void Compiler::set_member_qualified_name(uint32_t type_id, uint32_t index, const std::string &name)
1388	{
1389	ir.meta[type_id].members.resize(new_size: max(a: ir.meta[type_id].members.size(), b: size_t(index) + 1));
1390	ir.meta[type_id].members[index].qualified_alias = name;
1391	}
1392
1393	const string &Compiler::get_member_qualified_name(TypeID type_id, uint32_t index) const
1394	{
1395	auto *m = ir.find_meta(id: type_id);
1396	if (m && index < m->members.size())
1397	return m->members[index].qualified_alias;
1398	else
1399	return ir.get_empty_string();
1400	}
1401
1402	uint32_t Compiler::get_member_decoration(TypeID id, uint32_t index, Decoration decoration) const
1403	{
1404	return ir.get_member_decoration(id, index, decoration);
1405	}
1406
1407	const Bitset &Compiler::get_member_decoration_bitset(TypeID id, uint32_t index) const
1408	{
1409	return ir.get_member_decoration_bitset(id, index);
1410	}
1411
1412	bool Compiler::has_member_decoration(TypeID id, uint32_t index, Decoration decoration) const
1413	{
1414	return ir.has_member_decoration(id, index, decoration);
1415	}
1416
1417	void Compiler::unset_member_decoration(TypeID id, uint32_t index, Decoration decoration)
1418	{
1419	ir.unset_member_decoration(id, index, decoration);
1420	}
1421
1422	void Compiler::set_decoration_string(ID id, spv::Decoration decoration, const std::string &argument)
1423	{
1424	ir.set_decoration_string(id, decoration, argument);
1425	}
1426
1427	void Compiler::set_decoration(ID id, Decoration decoration, uint32_t argument)
1428	{
1429	ir.set_decoration(id, decoration, argument);
1430	}
1431
1432	void Compiler::set_extended_decoration(uint32_t id, ExtendedDecorations decoration, uint32_t value)
1433	{
1434	auto &dec = ir.meta[id].decoration;
1435	dec.extended.flags.set(decoration);
1436	dec.extended.values[decoration] = value;
1437	}
1438
1439	void Compiler::set_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration,
1440	uint32_t value)
1441	{
1442	ir.meta[type].members.resize(new_size: max(a: ir.meta[type].members.size(), b: size_t(index) + 1));
1443	auto &dec = ir.meta[type].members[index];
1444	dec.extended.flags.set(decoration);
1445	dec.extended.values[decoration] = value;
1446	}
1447
1448	static uint32_t get_default_extended_decoration(ExtendedDecorations decoration)
1449	{
1450	switch (decoration)
1451	{
1452	case SPIRVCrossDecorationResourceIndexPrimary:
1453	case SPIRVCrossDecorationResourceIndexSecondary:
1454	case SPIRVCrossDecorationResourceIndexTertiary:
1455	case SPIRVCrossDecorationResourceIndexQuaternary:
1456	case SPIRVCrossDecorationInterfaceMemberIndex:
1457	return ~(0u);
1458
1459	default:
1460	return 0;
1461	}
1462	}
1463
1464	uint32_t Compiler::get_extended_decoration(uint32_t id, ExtendedDecorations decoration) const
1465	{
1466	auto *m = ir.find_meta(id);
1467	if (!m)
1468	return 0;
1469
1470	auto &dec = m->decoration;
1471
1472	if (!dec.extended.flags.get(bit: decoration))
1473	return get_default_extended_decoration(decoration);
1474
1475	return dec.extended.values[decoration];
1476	}
1477
1478	uint32_t Compiler::get_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const
1479	{
1480	auto *m = ir.find_meta(id: type);
1481	if (!m)
1482	return 0;
1483
1484	if (index >= m->members.size())
1485	return 0;
1486
1487	auto &dec = m->members[index];
1488	if (!dec.extended.flags.get(bit: decoration))
1489	return get_default_extended_decoration(decoration);
1490	return dec.extended.values[decoration];
1491	}
1492
1493	bool Compiler::has_extended_decoration(uint32_t id, ExtendedDecorations decoration) const
1494	{
1495	auto *m = ir.find_meta(id);
1496	if (!m)
1497	return false;
1498
1499	auto &dec = m->decoration;
1500	return dec.extended.flags.get(bit: decoration);
1501	}
1502
1503	bool Compiler::has_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const
1504	{
1505	auto *m = ir.find_meta(id: type);
1506	if (!m)
1507	return false;
1508
1509	if (index >= m->members.size())
1510	return false;
1511
1512	auto &dec = m->members[index];
1513	return dec.extended.flags.get(bit: decoration);
1514	}
1515
1516	void Compiler::unset_extended_decoration(uint32_t id, ExtendedDecorations decoration)
1517	{
1518	auto &dec = ir.meta[id].decoration;
1519	dec.extended.flags.clear(bit: decoration);
1520	dec.extended.values[decoration] = 0;
1521	}
1522
1523	void Compiler::unset_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration)
1524	{
1525	ir.meta[type].members.resize(new_size: max(a: ir.meta[type].members.size(), b: size_t(index) + 1));
1526	auto &dec = ir.meta[type].members[index];
1527	dec.extended.flags.clear(bit: decoration);
1528	dec.extended.values[decoration] = 0;
1529	}
1530
1531	StorageClass Compiler::get_storage_class(VariableID id) const
1532	{
1533	return get<SPIRVariable>(id).storage;
1534	}
1535
1536	const std::string &Compiler::get_name(ID id) const
1537	{
1538	return ir.get_name(id);
1539	}
1540
1541	const std::string Compiler::get_fallback_name(ID id) const
1542	{
1543	return join(ts: "_", ts&: id);
1544	}
1545
1546	const std::string Compiler::get_block_fallback_name(VariableID id) const
1547	{
1548	auto &var = get<SPIRVariable>(id);
1549	if (get_name(id).empty())
1550	return join(ts: "_", ts: get<SPIRType>(id: var.basetype).self, ts: "_", ts&: id);
1551	else
1552	return get_name(id);
1553	}
1554
1555	const Bitset &Compiler::get_decoration_bitset(ID id) const
1556	{
1557	return ir.get_decoration_bitset(id);
1558	}
1559
1560	bool Compiler::has_decoration(ID id, Decoration decoration) const
1561	{
1562	return ir.has_decoration(id, decoration);
1563	}
1564
1565	const string &Compiler::get_decoration_string(ID id, Decoration decoration) const
1566	{
1567	return ir.get_decoration_string(id, decoration);
1568	}
1569
1570	const string &Compiler::get_member_decoration_string(TypeID id, uint32_t index, Decoration decoration) const
1571	{
1572	return ir.get_member_decoration_string(id, index, decoration);
1573	}
1574
1575	uint32_t Compiler::get_decoration(ID id, Decoration decoration) const
1576	{
1577	return ir.get_decoration(id, decoration);
1578	}
1579
1580	void Compiler::unset_decoration(ID id, Decoration decoration)
1581	{
1582	ir.unset_decoration(id, decoration);
1583	}
1584
1585	bool Compiler::get_binary_offset_for_decoration(VariableID id, spv::Decoration decoration, uint32_t &word_offset) const
1586	{
1587	auto *m = ir.find_meta(id);
1588	if (!m)
1589	return false;
1590
1591	auto &word_offsets = m->decoration_word_offset;
1592	auto itr = word_offsets.find(x: decoration);
1593	if (itr == end(cont: word_offsets))
1594	return false;
1595
1596	word_offset = itr->second;
1597	return true;
1598	}
1599
1600	bool Compiler::block_is_noop(const SPIRBlock &block) const
1601	{
1602	if (block.terminator != SPIRBlock::Direct)
1603	return false;
1604
1605	auto &child = get<SPIRBlock>(id: block.next_block);
1606
1607	// If this block participates in PHI, the block isn't really noop.
1608	for (auto &phi : block.phi_variables)
1609	if (phi.parent == block.self || phi.parent == child.self)
1610	return false;
1611
1612	for (auto &phi : child.phi_variables)
1613	if (phi.parent == block.self)
1614	return false;
1615
1616	// Verify all instructions have no semantic impact.
1617	for (auto &i : block.ops)
1618	{
1619	auto op = static_cast<Op>(i.op);
1620
1621	switch (op)
1622	{
1623	// Non-Semantic instructions.
1624	case OpLine:
1625	case OpNoLine:
1626	break;
1627
1628	case OpExtInst:
1629	{
1630	auto *ops = stream(instr: i);
1631	auto ext = get<SPIRExtension>(id: ops[2]).ext;
1632
1633	bool ext_is_nonsemantic_only =
1634	ext == SPIRExtension::NonSemanticShaderDebugInfo ||
1635	ext == SPIRExtension::SPV_debug_info ||
1636	ext == SPIRExtension::NonSemanticGeneric;
1637
1638	if (!ext_is_nonsemantic_only)
1639	return false;
1640
1641	break;
1642	}
1643
1644	default:
1645	return false;
1646	}
1647	}
1648
1649	return true;
1650	}
1651
1652	bool Compiler::block_is_loop_candidate(const SPIRBlock &block, SPIRBlock::Method method) const
1653	{
1654	// Tried and failed.
1655	if (block.disable_block_optimization || block.complex_continue)
1656	return false;
1657
1658	if (method == SPIRBlock::MergeToSelectForLoop || method == SPIRBlock::MergeToSelectContinueForLoop)
1659	{
1660	// Try to detect common for loop pattern
1661	// which the code backend can use to create cleaner code.
1662	// for(;;) { if (cond) { some_body; } else { break; } }
1663	// is the pattern we're looking for.
1664	const auto *false_block = maybe_get<SPIRBlock>(id: block.false_block);
1665	const auto *true_block = maybe_get<SPIRBlock>(id: block.true_block);
1666	const auto *merge_block = maybe_get<SPIRBlock>(id: block.merge_block);
1667
1668	bool false_block_is_merge = block.false_block == block.merge_block ||
1669	(false_block && merge_block && execution_is_noop(from: *false_block, to: *merge_block));
1670
1671	bool true_block_is_merge = block.true_block == block.merge_block ||
1672	(true_block && merge_block && execution_is_noop(from: *true_block, to: *merge_block));
1673
1674	bool positive_candidate =
1675	block.true_block != block.merge_block && block.true_block != block.self && false_block_is_merge;
1676
1677	bool negative_candidate =
1678	block.false_block != block.merge_block && block.false_block != block.self && true_block_is_merge;
1679
1680	bool ret = block.terminator == SPIRBlock::Select && block.merge == SPIRBlock::MergeLoop &&
1681	(positive_candidate || negative_candidate);
1682
1683	if (ret && positive_candidate && method == SPIRBlock::MergeToSelectContinueForLoop)
1684	ret = block.true_block == block.continue_block;
1685	else if (ret && negative_candidate && method == SPIRBlock::MergeToSelectContinueForLoop)
1686	ret = block.false_block == block.continue_block;
1687
1688	// If we have OpPhi which depends on branches which came from our own block,
1689	// we need to flush phi variables in else block instead of a trivial break,
1690	// so we cannot assume this is a for loop candidate.
1691	if (ret)
1692	{
1693	for (auto &phi : block.phi_variables)
1694	if (phi.parent == block.self)
1695	return false;
1696
1697	auto *merge = maybe_get<SPIRBlock>(id: block.merge_block);
1698	if (merge)
1699	for (auto &phi : merge->phi_variables)
1700	if (phi.parent == block.self)
1701	return false;
1702	}
1703	return ret;
1704	}
1705	else if (method == SPIRBlock::MergeToDirectForLoop)
1706	{
1707	// Empty loop header that just sets up merge target
1708	// and branches to loop body.
1709	bool ret = block.terminator == SPIRBlock::Direct && block.merge == SPIRBlock::MergeLoop && block_is_noop(block);
1710
1711	if (!ret)
1712	return false;
1713
1714	auto &child = get<SPIRBlock>(id: block.next_block);
1715
1716	const auto *false_block = maybe_get<SPIRBlock>(id: child.false_block);
1717	const auto *true_block = maybe_get<SPIRBlock>(id: child.true_block);
1718	const auto *merge_block = maybe_get<SPIRBlock>(id: block.merge_block);
1719
1720	bool false_block_is_merge = child.false_block == block.merge_block ||
1721	(false_block && merge_block && execution_is_noop(from: *false_block, to: *merge_block));
1722
1723	bool true_block_is_merge = child.true_block == block.merge_block ||
1724	(true_block && merge_block && execution_is_noop(from: *true_block, to: *merge_block));
1725
1726	bool positive_candidate =
1727	child.true_block != block.merge_block && child.true_block != block.self && false_block_is_merge;
1728
1729	bool negative_candidate =
1730	child.false_block != block.merge_block && child.false_block != block.self && true_block_is_merge;
1731
1732	ret = child.terminator == SPIRBlock::Select && child.merge == SPIRBlock::MergeNone &&
1733	(positive_candidate || negative_candidate);
1734
1735	if (ret)
1736	{
1737	auto *merge = maybe_get<SPIRBlock>(id: block.merge_block);
1738	if (merge)
1739	for (auto &phi : merge->phi_variables)
1740	if (phi.parent == block.self || phi.parent == child.false_block)
1741	return false;
1742	}
1743
1744	return ret;
1745	}
1746	else
1747	return false;
1748	}
1749
1750	bool Compiler::execution_is_noop(const SPIRBlock &from, const SPIRBlock &to) const
1751	{
1752	if (!execution_is_branchless(from, to))
1753	return false;
1754
1755	auto *start = &from;
1756	for (;;)
1757	{
1758	if (start->self == to.self)
1759	return true;
1760
1761	if (!block_is_noop(block: *start))
1762	return false;
1763
1764	auto &next = get<SPIRBlock>(id: start->next_block);
1765	start = &next;
1766	}
1767	}
1768
1769	bool Compiler::execution_is_branchless(const SPIRBlock &from, const SPIRBlock &to) const
1770	{
1771	auto *start = &from;
1772	for (;;)
1773	{
1774	if (start->self == to.self)
1775	return true;
1776
1777	if (start->terminator == SPIRBlock::Direct && start->merge == SPIRBlock::MergeNone)
1778	start = &get<SPIRBlock>(id: start->next_block);
1779	else
1780	return false;
1781	}
1782	}
1783
1784	bool Compiler::execution_is_direct_branch(const SPIRBlock &from, const SPIRBlock &to) const
1785	{
1786	return from.terminator == SPIRBlock::Direct && from.merge == SPIRBlock::MergeNone && from.next_block == to.self;
1787	}
1788
1789	SPIRBlock::ContinueBlockType Compiler::continue_block_type(const SPIRBlock &block) const
1790	{
1791	// The block was deemed too complex during code emit, pick conservative fallback paths.
1792	if (block.complex_continue)
1793	return SPIRBlock::ComplexLoop;
1794
1795	// In older glslang output continue block can be equal to the loop header.
1796	// In this case, execution is clearly branchless, so just assume a while loop header here.
1797	if (block.merge == SPIRBlock::MergeLoop)
1798	return SPIRBlock::WhileLoop;
1799
1800	if (block.loop_dominator == BlockID(SPIRBlock::NoDominator))
1801	{
1802	// Continue block is never reached from CFG.
1803	return SPIRBlock::ComplexLoop;
1804	}
1805
1806	auto &dominator = get<SPIRBlock>(id: block.loop_dominator);
1807
1808	if (execution_is_noop(from: block, to: dominator))
1809	return SPIRBlock::WhileLoop;
1810	else if (execution_is_branchless(from: block, to: dominator))
1811	return SPIRBlock::ForLoop;
1812	else
1813	{
1814	const auto *false_block = maybe_get<SPIRBlock>(id: block.false_block);
1815	const auto *true_block = maybe_get<SPIRBlock>(id: block.true_block);
1816	const auto *merge_block = maybe_get<SPIRBlock>(id: dominator.merge_block);
1817
1818	// If we need to flush Phi in this block, we cannot have a DoWhile loop.
1819	bool flush_phi_to_false = false_block && flush_phi_required(from: block.self, to: block.false_block);
1820	bool flush_phi_to_true = true_block && flush_phi_required(from: block.self, to: block.true_block);
1821	if (flush_phi_to_false || flush_phi_to_true)
1822	return SPIRBlock::ComplexLoop;
1823
1824	bool positive_do_while = block.true_block == dominator.self &&
1825	(block.false_block == dominator.merge_block ||
1826	(false_block && merge_block && execution_is_noop(from: *false_block, to: *merge_block)));
1827
1828	bool negative_do_while = block.false_block == dominator.self &&
1829	(block.true_block == dominator.merge_block ||
1830	(true_block && merge_block && execution_is_noop(from: *true_block, to: *merge_block)));
1831
1832	if (block.merge == SPIRBlock::MergeNone && block.terminator == SPIRBlock::Select &&
1833	(positive_do_while || negative_do_while))
1834	{
1835	return SPIRBlock::DoWhileLoop;
1836	}
1837	else
1838	return SPIRBlock::ComplexLoop;
1839	}
1840	}
1841
1842	const SmallVector<SPIRBlock::Case> &Compiler::get_case_list(const SPIRBlock &block) const
1843	{
1844	uint32_t width = 0;
1845
1846	// First we check if we can get the type directly from the block.condition
1847	// since it can be a SPIRConstant or a SPIRVariable.
1848	if (const auto *constant = maybe_get<SPIRConstant>(id: block.condition))
1849	{
1850	const auto &type = get<SPIRType>(id: constant->constant_type);
1851	width = type.width;
1852	}
1853	else if (const auto *op = maybe_get<SPIRConstantOp>(id: block.condition))
1854	{
1855	const auto &type = get<SPIRType>(id: op->basetype);
1856	width = type.width;
1857	}
1858	else if (const auto *var = maybe_get<SPIRVariable>(id: block.condition))
1859	{
1860	const auto &type = get<SPIRType>(id: var->basetype);
1861	width = type.width;
1862	}
1863	else if (const auto *undef = maybe_get<SPIRUndef>(id: block.condition))
1864	{
1865	const auto &type = get<SPIRType>(id: undef->basetype);
1866	width = type.width;
1867	}
1868	else
1869	{
1870	auto search = ir.load_type_width.find(x: block.condition);
1871	if (search == ir.load_type_width.end())
1872	{
1873	SPIRV_CROSS_THROW("Use of undeclared variable on a switch statement.");
1874	}
1875
1876	width = search->second;
1877	}
1878
1879	if (width > 32)
1880	return block.cases_64bit;
1881
1882	return block.cases_32bit;
1883	}
1884
1885	bool Compiler::traverse_all_reachable_opcodes(const SPIRBlock &block, OpcodeHandler &handler) const
1886	{
1887	handler.set_current_block(block);
1888	handler.rearm_current_block(block);
1889
1890	// Ideally, perhaps traverse the CFG instead of all blocks in order to eliminate dead blocks,
1891	// but this shouldn't be a problem in practice unless the SPIR-V is doing insane things like recursing
1892	// inside dead blocks ...
1893	for (auto &i : block.ops)
1894	{
1895	auto ops = stream(instr: i);
1896	auto op = static_cast<Op>(i.op);
1897
1898	if (!handler.handle(opcode: op, args: ops, length: i.length))
1899	return false;
1900
1901	if (op == OpFunctionCall)
1902	{
1903	auto &func = get<SPIRFunction>(id: ops[2]);
1904	if (handler.follow_function_call(func))
1905	{
1906	if (!handler.begin_function_scope(ops, i.length))
1907	return false;
1908	if (!traverse_all_reachable_opcodes(block: get<SPIRFunction>(id: ops[2]), handler))
1909	return false;
1910	if (!handler.end_function_scope(ops, i.length))
1911	return false;
1912
1913	handler.rearm_current_block(block);
1914	}
1915	}
1916	}
1917
1918	if (!handler.handle_terminator(block))
1919	return false;
1920
1921	return true;
1922	}
1923
1924	bool Compiler::traverse_all_reachable_opcodes(const SPIRFunction &func, OpcodeHandler &handler) const
1925	{
1926	for (auto block : func.blocks)
1927	if (!traverse_all_reachable_opcodes(block: get<SPIRBlock>(id: block), handler))
1928	return false;
1929
1930	return true;
1931	}
1932
1933	uint32_t Compiler::type_struct_member_offset(const SPIRType &type, uint32_t index) const
1934	{
1935	auto *type_meta = ir.find_meta(id: type.self);
1936	if (type_meta)
1937	{
1938	// Decoration must be set in valid SPIR-V, otherwise throw.
1939	auto &dec = type_meta->members[index];
1940	if (dec.decoration_flags.get(bit: DecorationOffset))
1941	return dec.offset;
1942	else
1943	SPIRV_CROSS_THROW("Struct member does not have Offset set.");
1944	}
1945	else
1946	SPIRV_CROSS_THROW("Struct member does not have Offset set.");
1947	}
1948
1949	uint32_t Compiler::type_struct_member_array_stride(const SPIRType &type, uint32_t index) const
1950	{
1951	auto *type_meta = ir.find_meta(id: type.member_types[index]);
1952	if (type_meta)
1953	{
1954	// Decoration must be set in valid SPIR-V, otherwise throw.
1955	// ArrayStride is part of the array type not OpMemberDecorate.
1956	auto &dec = type_meta->decoration;
1957	if (dec.decoration_flags.get(bit: DecorationArrayStride))
1958	return dec.array_stride;
1959	else
1960	SPIRV_CROSS_THROW("Struct member does not have ArrayStride set.");
1961	}
1962	else
1963	SPIRV_CROSS_THROW("Struct member does not have ArrayStride set.");
1964	}
1965
1966	uint32_t Compiler::type_struct_member_matrix_stride(const SPIRType &type, uint32_t index) const
1967	{
1968	auto *type_meta = ir.find_meta(id: type.self);
1969	if (type_meta)
1970	{
1971	// Decoration must be set in valid SPIR-V, otherwise throw.
1972	// MatrixStride is part of OpMemberDecorate.
1973	auto &dec = type_meta->members[index];
1974	if (dec.decoration_flags.get(bit: DecorationMatrixStride))
1975	return dec.matrix_stride;
1976	else
1977	SPIRV_CROSS_THROW("Struct member does not have MatrixStride set.");
1978	}
1979	else
1980	SPIRV_CROSS_THROW("Struct member does not have MatrixStride set.");
1981	}
1982
1983	size_t Compiler::get_declared_struct_size(const SPIRType &type) const
1984	{
1985	if (type.member_types.empty())
1986	SPIRV_CROSS_THROW("Declared struct in block cannot be empty.");
1987
1988	// Offsets can be declared out of order, so we need to deduce the actual size
1989	// based on last member instead.
1990	uint32_t member_index = 0;
1991	size_t highest_offset = 0;
1992	for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++)
1993	{
1994	size_t offset = type_struct_member_offset(type, index: i);
1995	if (offset > highest_offset)
1996	{
1997	highest_offset = offset;
1998	member_index = i;
1999	}
2000	}
2001
2002	size_t size = get_declared_struct_member_size(struct_type: type, index: member_index);
2003	return highest_offset + size;
2004	}
2005
2006	size_t Compiler::get_declared_struct_size_runtime_array(const SPIRType &type, size_t array_size) const
2007	{
2008	if (type.member_types.empty())
2009	SPIRV_CROSS_THROW("Declared struct in block cannot be empty.");
2010
2011	size_t size = get_declared_struct_size(type);
2012	auto &last_type = get<SPIRType>(id: type.member_types.back());
2013	if (!last_type.array.empty() && last_type.array_size_literal[0] && last_type.array[0] == 0) // Runtime array
2014	size += array_size * type_struct_member_array_stride(type, index: uint32_t(type.member_types.size() - 1));
2015
2016	return size;
2017	}
2018
2019	uint32_t Compiler::evaluate_spec_constant_u32(const SPIRConstantOp &spec) const
2020	{
2021	auto &result_type = get<SPIRType>(id: spec.basetype);
2022	if (result_type.basetype != SPIRType::UInt && result_type.basetype != SPIRType::Int &&
2023	result_type.basetype != SPIRType::Boolean)
2024	{
2025	SPIRV_CROSS_THROW(
2026	"Only 32-bit integers and booleans are currently supported when evaluating specialization constants.\n");
2027	}
2028
2029	if (!is_scalar(type: result_type))
2030	SPIRV_CROSS_THROW("Spec constant evaluation must be a scalar.\n");
2031
2032	uint32_t value = 0;
2033
2034	const auto eval_u32 = [&](uint32_t id) -> uint32_t {
2035	auto &type = expression_type(id);
2036	if (type.basetype != SPIRType::UInt && type.basetype != SPIRType::Int && type.basetype != SPIRType::Boolean)
2037	{
2038	SPIRV_CROSS_THROW("Only 32-bit integers and booleans are currently supported when evaluating "
2039	"specialization constants.\n");
2040	}
2041
2042	if (!is_scalar(type))
2043	SPIRV_CROSS_THROW("Spec constant evaluation must be a scalar.\n");
2044	if (const auto *c = this->maybe_get<SPIRConstant>(id))
2045	return c->scalar();
2046	else
2047	return evaluate_spec_constant_u32(spec: this->get<SPIRConstantOp>(id));
2048	};
2049
2050	#define binary_spec_op(op, binary_op) \
2051	case Op##op: \
2052	value = eval_u32(spec.arguments[0]) binary_op eval_u32(spec.arguments[1]); \
2053	break
2054	#define binary_spec_op_cast(op, binary_op, type) \
2055	case Op##op: \
2056	value = uint32_t(type(eval_u32(spec.arguments[0])) binary_op type(eval_u32(spec.arguments[1]))); \
2057	break
2058
2059	// Support the basic opcodes which are typically used when computing array sizes.
2060	switch (spec.opcode)
2061	{
2062	binary_spec_op(IAdd, +);
2063	binary_spec_op(ISub, -);
2064	binary_spec_op(IMul, *);
2065	binary_spec_op(BitwiseAnd, &);
2066	binary_spec_op(BitwiseOr, |);
2067	binary_spec_op(BitwiseXor, ^);
2068	binary_spec_op(LogicalAnd, &);
2069	binary_spec_op(LogicalOr, |);
2070	binary_spec_op(ShiftLeftLogical, <<);
2071	binary_spec_op(ShiftRightLogical, >>);
2072	binary_spec_op_cast(ShiftRightArithmetic, >>, int32_t);
2073	binary_spec_op(LogicalEqual, ==);
2074	binary_spec_op(LogicalNotEqual, !=);
2075	binary_spec_op(IEqual, ==);
2076	binary_spec_op(INotEqual, !=);
2077	binary_spec_op(ULessThan, <);
2078	binary_spec_op(ULessThanEqual, <=);
2079	binary_spec_op(UGreaterThan, >);
2080	binary_spec_op(UGreaterThanEqual, >=);
2081	binary_spec_op_cast(SLessThan, <, int32_t);
2082	binary_spec_op_cast(SLessThanEqual, <=, int32_t);
2083	binary_spec_op_cast(SGreaterThan, >, int32_t);
2084	binary_spec_op_cast(SGreaterThanEqual, >=, int32_t);
2085	#undef binary_spec_op
2086	#undef binary_spec_op_cast
2087
2088	case OpLogicalNot:
2089	value = uint32_t(!eval_u32(spec.arguments[0]));
2090	break;
2091
2092	case OpNot:
2093	value = ~eval_u32(spec.arguments[0]);
2094	break;
2095
2096	case OpSNegate:
2097	value = uint32_t(-int32_t(eval_u32(spec.arguments[0])));
2098	break;
2099
2100	case OpSelect:
2101	value = eval_u32(spec.arguments[0]) ? eval_u32(spec.arguments[1]) : eval_u32(spec.arguments[2]);
2102	break;
2103
2104	case OpUMod:
2105	{
2106	uint32_t a = eval_u32(spec.arguments[0]);
2107	uint32_t b = eval_u32(spec.arguments[1]);
2108	if (b == 0)
2109	SPIRV_CROSS_THROW("Undefined behavior in UMod, b == 0.\n");
2110	value = a % b;
2111	break;
2112	}
2113
2114	case OpSRem:
2115	{
2116	auto a = int32_t(eval_u32(spec.arguments[0]));
2117	auto b = int32_t(eval_u32(spec.arguments[1]));
2118	if (b == 0)
2119	SPIRV_CROSS_THROW("Undefined behavior in SRem, b == 0.\n");
2120	value = a % b;
2121	break;
2122	}
2123
2124	case OpSMod:
2125	{
2126	auto a = int32_t(eval_u32(spec.arguments[0]));
2127	auto b = int32_t(eval_u32(spec.arguments[1]));
2128	if (b == 0)
2129	SPIRV_CROSS_THROW("Undefined behavior in SMod, b == 0.\n");
2130	auto v = a % b;
2131
2132	// Makes sure we match the sign of b, not a.
2133	if ((b < 0 && v > 0) || (b > 0 && v < 0))
2134	v += b;
2135	value = v;
2136	break;
2137	}
2138
2139	case OpUDiv:
2140	{
2141	uint32_t a = eval_u32(spec.arguments[0]);
2142	uint32_t b = eval_u32(spec.arguments[1]);
2143	if (b == 0)
2144	SPIRV_CROSS_THROW("Undefined behavior in UDiv, b == 0.\n");
2145	value = a / b;
2146	break;
2147	}
2148
2149	case OpSDiv:
2150	{
2151	auto a = int32_t(eval_u32(spec.arguments[0]));
2152	auto b = int32_t(eval_u32(spec.arguments[1]));
2153	if (b == 0)
2154	SPIRV_CROSS_THROW("Undefined behavior in SDiv, b == 0.\n");
2155	value = a / b;
2156	break;
2157	}
2158
2159	default:
2160	SPIRV_CROSS_THROW("Unsupported spec constant opcode for evaluation.\n");
2161	}
2162
2163	return value;
2164	}
2165
2166	uint32_t Compiler::evaluate_constant_u32(uint32_t id) const
2167	{
2168	if (const auto *c = maybe_get<SPIRConstant>(id))
2169	return c->scalar();
2170	else
2171	return evaluate_spec_constant_u32(spec: get<SPIRConstantOp>(id));
2172	}
2173
2174	size_t Compiler::get_declared_struct_member_size(const SPIRType &struct_type, uint32_t index) const
2175	{
2176	if (struct_type.member_types.empty())
2177	SPIRV_CROSS_THROW("Declared struct in block cannot be empty.");
2178
2179	auto &flags = get_member_decoration_bitset(id: struct_type.self, index);
2180	auto &type = get<SPIRType>(id: struct_type.member_types[index]);
2181
2182	switch (type.basetype)
2183	{
2184	case SPIRType::Unknown:
2185	case SPIRType::Void:
2186	case SPIRType::Boolean: // Bools are purely logical, and cannot be used for externally visible types.
2187	case SPIRType::AtomicCounter:
2188	case SPIRType::Image:
2189	case SPIRType::SampledImage:
2190	case SPIRType::Sampler:
2191	SPIRV_CROSS_THROW("Querying size for object with opaque size.");
2192
2193	default:
2194	break;
2195	}
2196
2197	if (type.pointer && type.storage == StorageClassPhysicalStorageBuffer)
2198	{
2199	// Check if this is a top-level pointer type, and not an array of pointers.
2200	if (type.pointer_depth > get<SPIRType>(id: type.parent_type).pointer_depth)
2201	return 8;
2202	}
2203
2204	if (!type.array.empty())
2205	{
2206	// For arrays, we can use ArrayStride to get an easy check.
2207	bool array_size_literal = type.array_size_literal.back();
2208	uint32_t array_size = array_size_literal ? type.array.back() : evaluate_constant_u32(id: type.array.back());
2209	return type_struct_member_array_stride(type: struct_type, index) * array_size;
2210	}
2211	else if (type.basetype == SPIRType::Struct)
2212	{
2213	return get_declared_struct_size(type);
2214	}
2215	else
2216	{
2217	unsigned vecsize = type.vecsize;
2218	unsigned columns = type.columns;
2219
2220	// Vectors.
2221	if (columns == 1)
2222	{
2223	size_t component_size = type.width / 8;
2224	return vecsize * component_size;
2225	}
2226	else
2227	{
2228	uint32_t matrix_stride = type_struct_member_matrix_stride(type: struct_type, index);
2229
2230	// Per SPIR-V spec, matrices must be tightly packed and aligned up for vec3 accesses.
2231	if (flags.get(bit: DecorationRowMajor))
2232	return matrix_stride * vecsize;
2233	else if (flags.get(bit: DecorationColMajor))
2234	return matrix_stride * columns;
2235	else
2236	SPIRV_CROSS_THROW("Either row-major or column-major must be declared for matrices.");
2237	}
2238	}
2239	}
2240
2241	bool Compiler::BufferAccessHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
2242	{
2243	if (opcode != OpAccessChain && opcode != OpInBoundsAccessChain && opcode != OpPtrAccessChain)
2244	return true;
2245
2246	bool ptr_chain = (opcode == OpPtrAccessChain);
2247
2248	// Invalid SPIR-V.
2249	if (length < (ptr_chain ? 5u : 4u))
2250	return false;
2251
2252	if (args[2] != id)
2253	return true;
2254
2255	// Don't bother traversing the entire access chain tree yet.
2256	// If we access a struct member, assume we access the entire member.
2257	uint32_t index = compiler.get<SPIRConstant>(id: args[ptr_chain ? 4 : 3]).scalar();
2258
2259	// Seen this index already.
2260	if (seen.find(x: index) != end(cont&: seen))
2261	return true;
2262	seen.insert(x: index);
2263
2264	auto &type = compiler.expression_type(id);
2265	uint32_t offset = compiler.type_struct_member_offset(type, index);
2266
2267	size_t range;
2268	// If we have another member in the struct, deduce the range by looking at the next member.
2269	// This is okay since structs in SPIR-V can have padding, but Offset decoration must be
2270	// monotonically increasing.
2271	// Of course, this doesn't take into account if the SPIR-V for some reason decided to add
2272	// very large amounts of padding, but that's not really a big deal.
2273	if (index + 1 < type.member_types.size())
2274	{
2275	range = compiler.type_struct_member_offset(type, index: index + 1) - offset;
2276	}
2277	else
2278	{
2279	// No padding, so just deduce it from the size of the member directly.
2280	range = compiler.get_declared_struct_member_size(struct_type: type, index);
2281	}
2282
2283	ranges.push_back(t: { .index: index, .offset: offset, .range: range });
2284	return true;
2285	}
2286
2287	SmallVector<BufferRange> Compiler::get_active_buffer_ranges(VariableID id) const
2288	{
2289	SmallVector<BufferRange> ranges;
2290	BufferAccessHandler handler(*this, ranges, id);
2291	traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler);
2292	return ranges;
2293	}
2294
2295	bool Compiler::types_are_logically_equivalent(const SPIRType &a, const SPIRType &b) const
2296	{
2297	if (a.basetype != b.basetype)
2298	return false;
2299	if (a.width != b.width)
2300	return false;
2301	if (a.vecsize != b.vecsize)
2302	return false;
2303	if (a.columns != b.columns)
2304	return false;
2305	if (a.array.size() != b.array.size())
2306	return false;
2307
2308	size_t array_count = a.array.size();
2309	if (array_count && memcmp(s1: a.array.data(), s2: b.array.data(), n: array_count * sizeof(uint32_t)) != 0)
2310	return false;
2311
2312	if (a.basetype == SPIRType::Image || a.basetype == SPIRType::SampledImage)
2313	{
2314	if (memcmp(s1: &a.image, s2: &b.image, n: sizeof(SPIRType::Image)) != 0)
2315	return false;
2316	}
2317
2318	if (a.member_types.size() != b.member_types.size())
2319	return false;
2320
2321	size_t member_types = a.member_types.size();
2322	for (size_t i = 0; i < member_types; i++)
2323	{
2324	if (!types_are_logically_equivalent(a: get<SPIRType>(id: a.member_types[i]), b: get<SPIRType>(id: b.member_types[i])))
2325	return false;
2326	}
2327
2328	return true;
2329	}
2330
2331	const Bitset &Compiler::get_execution_mode_bitset() const
2332	{
2333	return get_entry_point().flags;
2334	}
2335
2336	void Compiler::set_execution_mode(ExecutionMode mode, uint32_t arg0, uint32_t arg1, uint32_t arg2)
2337	{
2338	auto &execution = get_entry_point();
2339
2340	execution.flags.set(mode);
2341	switch (mode)
2342	{
2343	case ExecutionModeLocalSize:
2344	execution.workgroup_size.x = arg0;
2345	execution.workgroup_size.y = arg1;
2346	execution.workgroup_size.z = arg2;
2347	break;
2348
2349	case ExecutionModeLocalSizeId:
2350	execution.workgroup_size.id_x = arg0;
2351	execution.workgroup_size.id_y = arg1;
2352	execution.workgroup_size.id_z = arg2;
2353	break;
2354
2355	case ExecutionModeInvocations:
2356	execution.invocations = arg0;
2357	break;
2358
2359	case ExecutionModeOutputVertices:
2360	execution.output_vertices = arg0;
2361	break;
2362
2363	case ExecutionModeOutputPrimitivesEXT:
2364	execution.output_primitives = arg0;
2365	break;
2366
2367	default:
2368	break;
2369	}
2370	}
2371
2372	void Compiler::unset_execution_mode(ExecutionMode mode)
2373	{
2374	auto &execution = get_entry_point();
2375	execution.flags.clear(bit: mode);
2376	}
2377
2378	uint32_t Compiler::get_work_group_size_specialization_constants(SpecializationConstant &x, SpecializationConstant &y,
2379	SpecializationConstant &z) const
2380	{
2381	auto &execution = get_entry_point();
2382	x = { .id: 0, .constant_id: 0 };
2383	y = { .id: 0, .constant_id: 0 };
2384	z = { .id: 0, .constant_id: 0 };
2385
2386	// WorkgroupSize builtin takes precedence over LocalSize / LocalSizeId.
2387	if (execution.workgroup_size.constant != 0)
2388	{
2389	auto &c = get<SPIRConstant>(id: execution.workgroup_size.constant);
2390
2391	if (c.m.c[0].id[0] != ID(0))
2392	{
2393	x.id = c.m.c[0].id[0];
2394	x.constant_id = get_decoration(id: c.m.c[0].id[0], decoration: DecorationSpecId);
2395	}
2396
2397	if (c.m.c[0].id[1] != ID(0))
2398	{
2399	y.id = c.m.c[0].id[1];
2400	y.constant_id = get_decoration(id: c.m.c[0].id[1], decoration: DecorationSpecId);
2401	}
2402
2403	if (c.m.c[0].id[2] != ID(0))
2404	{
2405	z.id = c.m.c[0].id[2];
2406	z.constant_id = get_decoration(id: c.m.c[0].id[2], decoration: DecorationSpecId);
2407	}
2408	}
2409	else if (execution.flags.get(bit: ExecutionModeLocalSizeId))
2410	{
2411	auto &cx = get<SPIRConstant>(id: execution.workgroup_size.id_x);
2412	if (cx.specialization)
2413	{
2414	x.id = execution.workgroup_size.id_x;
2415	x.constant_id = get_decoration(id: execution.workgroup_size.id_x, decoration: DecorationSpecId);
2416	}
2417
2418	auto &cy = get<SPIRConstant>(id: execution.workgroup_size.id_y);
2419	if (cy.specialization)
2420	{
2421	y.id = execution.workgroup_size.id_y;
2422	y.constant_id = get_decoration(id: execution.workgroup_size.id_y, decoration: DecorationSpecId);
2423	}
2424
2425	auto &cz = get<SPIRConstant>(id: execution.workgroup_size.id_z);
2426	if (cz.specialization)
2427	{
2428	z.id = execution.workgroup_size.id_z;
2429	z.constant_id = get_decoration(id: execution.workgroup_size.id_z, decoration: DecorationSpecId);
2430	}
2431	}
2432
2433	return execution.workgroup_size.constant;
2434	}
2435
2436	uint32_t Compiler::get_execution_mode_argument(spv::ExecutionMode mode, uint32_t index) const
2437	{
2438	auto &execution = get_entry_point();
2439	switch (mode)
2440	{
2441	case ExecutionModeLocalSizeId:
2442	if (execution.flags.get(bit: ExecutionModeLocalSizeId))
2443	{
2444	switch (index)
2445	{
2446	case 0:
2447	return execution.workgroup_size.id_x;
2448	case 1:
2449	return execution.workgroup_size.id_y;
2450	case 2:
2451	return execution.workgroup_size.id_z;
2452	default:
2453	return 0;
2454	}
2455	}
2456	else
2457	return 0;
2458
2459	case ExecutionModeLocalSize:
2460	switch (index)
2461	{
2462	case 0:
2463	if (execution.flags.get(bit: ExecutionModeLocalSizeId) && execution.workgroup_size.id_x != 0)
2464	return get<SPIRConstant>(id: execution.workgroup_size.id_x).scalar();
2465	else
2466	return execution.workgroup_size.x;
2467	case 1:
2468	if (execution.flags.get(bit: ExecutionModeLocalSizeId) && execution.workgroup_size.id_y != 0)
2469	return get<SPIRConstant>(id: execution.workgroup_size.id_y).scalar();
2470	else
2471	return execution.workgroup_size.y;
2472	case 2:
2473	if (execution.flags.get(bit: ExecutionModeLocalSizeId) && execution.workgroup_size.id_z != 0)
2474	return get<SPIRConstant>(id: execution.workgroup_size.id_z).scalar();
2475	else
2476	return execution.workgroup_size.z;
2477	default:
2478	return 0;
2479	}
2480
2481	case ExecutionModeInvocations:
2482	return execution.invocations;
2483
2484	case ExecutionModeOutputVertices:
2485	return execution.output_vertices;
2486
2487	case ExecutionModeOutputPrimitivesEXT:
2488	return execution.output_primitives;
2489
2490	default:
2491	return 0;
2492	}
2493	}
2494
2495	ExecutionModel Compiler::get_execution_model() const
2496	{
2497	auto &execution = get_entry_point();
2498	return execution.model;
2499	}
2500
2501	bool Compiler::is_tessellation_shader(ExecutionModel model)
2502	{
2503	return model == ExecutionModelTessellationControl || model == ExecutionModelTessellationEvaluation;
2504	}
2505
2506	bool Compiler::is_vertex_like_shader() const
2507	{
2508	auto model = get_execution_model();
2509	return model == ExecutionModelVertex || model == ExecutionModelGeometry ||
2510	model == ExecutionModelTessellationControl || model == ExecutionModelTessellationEvaluation;
2511	}
2512
2513	bool Compiler::is_tessellation_shader() const
2514	{
2515	return is_tessellation_shader(model: get_execution_model());
2516	}
2517
2518	bool Compiler::is_tessellating_triangles() const
2519	{
2520	return get_execution_mode_bitset().get(bit: ExecutionModeTriangles);
2521	}
2522
2523	void Compiler::set_remapped_variable_state(VariableID id, bool remap_enable)
2524	{
2525	get<SPIRVariable>(id).remapped_variable = remap_enable;
2526	}
2527
2528	bool Compiler::get_remapped_variable_state(VariableID id) const
2529	{
2530	return get<SPIRVariable>(id).remapped_variable;
2531	}
2532
2533	void Compiler::set_subpass_input_remapped_components(VariableID id, uint32_t components)
2534	{
2535	get<SPIRVariable>(id).remapped_components = components;
2536	}
2537
2538	uint32_t Compiler::get_subpass_input_remapped_components(VariableID id) const
2539	{
2540	return get<SPIRVariable>(id).remapped_components;
2541	}
2542
2543	void Compiler::add_implied_read_expression(SPIRExpression &e, uint32_t source)
2544	{
2545	auto itr = find(first: begin(cont&: e.implied_read_expressions), last: end(cont&: e.implied_read_expressions), val: ID(source));
2546	if (itr == end(cont&: e.implied_read_expressions))
2547	e.implied_read_expressions.push_back(t: source);
2548	}
2549
2550	void Compiler::add_implied_read_expression(SPIRAccessChain &e, uint32_t source)
2551	{
2552	auto itr = find(first: begin(cont&: e.implied_read_expressions), last: end(cont&: e.implied_read_expressions), val: ID(source));
2553	if (itr == end(cont&: e.implied_read_expressions))
2554	e.implied_read_expressions.push_back(t: source);
2555	}
2556
2557	void Compiler::add_active_interface_variable(uint32_t var_id)
2558	{
2559	active_interface_variables.insert(x: var_id);
2560
2561	// In SPIR-V 1.4 and up we must also track the interface variable in the entry point.
2562	if (ir.get_spirv_version() >= 0x10400)
2563	{
2564	auto &vars = get_entry_point().interface_variables;
2565	if (find(first: begin(cont&: vars), last: end(cont&: vars), val: VariableID(var_id)) == end(cont&: vars))
2566	vars.push_back(t: var_id);
2567	}
2568	}
2569
2570	void Compiler::inherit_expression_dependencies(uint32_t dst, uint32_t source_expression)
2571	{
2572	auto *ptr_e = maybe_get<SPIRExpression>(id: dst);
2573
2574	if (is_position_invariant() && ptr_e && maybe_get<SPIRExpression>(id: source_expression))
2575	{
2576	auto &deps = ptr_e->invariance_dependencies;
2577	if (std::find(first: deps.begin(), last: deps.end(), val: source_expression) == deps.end())
2578	deps.push_back(t: source_expression);
2579	}
2580
2581	// Don't inherit any expression dependencies if the expression in dst
2582	// is not a forwarded temporary.
2583	if (forwarded_temporaries.find(x: dst) == end(cont&: forwarded_temporaries) ||
2584	forced_temporaries.find(x: dst) != end(cont&: forced_temporaries))
2585	{
2586	return;
2587	}
2588
2589	auto &e = *ptr_e;
2590	auto *phi = maybe_get<SPIRVariable>(id: source_expression);
2591	if (phi && phi->phi_variable)
2592	{
2593	// We have used a phi variable, which can change at the end of the block,
2594	// so make sure we take a dependency on this phi variable.
2595	phi->dependees.push_back(t: dst);
2596	}
2597
2598	auto *s = maybe_get<SPIRExpression>(id: source_expression);
2599	if (!s)
2600	return;
2601
2602	auto &e_deps = e.expression_dependencies;
2603	auto &s_deps = s->expression_dependencies;
2604
2605	// If we depend on a expression, we also depend on all sub-dependencies from source.
2606	e_deps.push_back(t: source_expression);
2607	e_deps.insert(itr: end(cont&: e_deps), insert_begin: begin(cont&: s_deps), insert_end: end(cont&: s_deps));
2608
2609	// Eliminate duplicated dependencies.
2610	sort(first: begin(cont&: e_deps), last: end(cont&: e_deps));
2611	e_deps.erase(start_erase: unique(first: begin(cont&: e_deps), last: end(cont&: e_deps)), end_erase: end(cont&: e_deps));
2612	}
2613
2614	SmallVector<EntryPoint> Compiler::get_entry_points_and_stages() const
2615	{
2616	SmallVector<EntryPoint> entries;
2617	for (auto &entry : ir.entry_points)
2618	entries.push_back(t: { .name: entry.second.orig_name, .execution_model: entry.second.model });
2619	return entries;
2620	}
2621
2622	void Compiler::rename_entry_point(const std::string &old_name, const std::string &new_name, spv::ExecutionModel model)
2623	{
2624	auto &entry = get_entry_point(name: old_name, execution_model: model);
2625	entry.orig_name = new_name;
2626	entry.name = new_name;
2627	}
2628
2629	void Compiler::set_entry_point(const std::string &name, spv::ExecutionModel model)
2630	{
2631	auto &entry = get_entry_point(name, execution_model: model);
2632	ir.default_entry_point = entry.self;
2633	}
2634
2635	SPIREntryPoint &Compiler::get_first_entry_point(const std::string &name)
2636	{
2637	auto itr = find_if(
2638	first: begin(cont&: ir.entry_points), last: end(cont&: ir.entry_points),
2639	pred: [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool { return entry.second.orig_name == name; });
2640
2641	if (itr == end(cont&: ir.entry_points))
2642	SPIRV_CROSS_THROW("Entry point does not exist.");
2643
2644	return itr->second;
2645	}
2646
2647	const SPIREntryPoint &Compiler::get_first_entry_point(const std::string &name) const
2648	{
2649	auto itr = find_if(
2650	first: begin(cont: ir.entry_points), last: end(cont: ir.entry_points),
2651	pred: [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool { return entry.second.orig_name == name; });
2652
2653	if (itr == end(cont: ir.entry_points))
2654	SPIRV_CROSS_THROW("Entry point does not exist.");
2655
2656	return itr->second;
2657	}
2658
2659	SPIREntryPoint &Compiler::get_entry_point(const std::string &name, ExecutionModel model)
2660	{
2661	auto itr = find_if(first: begin(cont&: ir.entry_points), last: end(cont&: ir.entry_points),
2662	pred: [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool {
2663	return entry.second.orig_name == name && entry.second.model == model;
2664	});
2665
2666	if (itr == end(cont&: ir.entry_points))
2667	SPIRV_CROSS_THROW("Entry point does not exist.");
2668
2669	return itr->second;
2670	}
2671
2672	const SPIREntryPoint &Compiler::get_entry_point(const std::string &name, ExecutionModel model) const
2673	{
2674	auto itr = find_if(first: begin(cont: ir.entry_points), last: end(cont: ir.entry_points),
2675	pred: [&](const std::pair<uint32_t, SPIREntryPoint> &entry) -> bool {
2676	return entry.second.orig_name == name && entry.second.model == model;
2677	});
2678
2679	if (itr == end(cont: ir.entry_points))
2680	SPIRV_CROSS_THROW("Entry point does not exist.");
2681
2682	return itr->second;
2683	}
2684
2685	const string &Compiler::get_cleansed_entry_point_name(const std::string &name, ExecutionModel model) const
2686	{
2687	return get_entry_point(name, model).name;
2688	}
2689
2690	const SPIREntryPoint &Compiler::get_entry_point() const
2691	{
2692	return ir.entry_points.find(x: ir.default_entry_point)->second;
2693	}
2694
2695	SPIREntryPoint &Compiler::get_entry_point()
2696	{
2697	return ir.entry_points.find(x: ir.default_entry_point)->second;
2698	}
2699
2700	bool Compiler::interface_variable_exists_in_entry_point(uint32_t id) const
2701	{
2702	auto &var = get<SPIRVariable>(id);
2703
2704	if (ir.get_spirv_version() < 0x10400)
2705	{
2706	if (var.storage != StorageClassInput && var.storage != StorageClassOutput &&
2707	var.storage != StorageClassUniformConstant)
2708	SPIRV_CROSS_THROW("Only Input, Output variables and Uniform constants are part of a shader linking interface.");
2709
2710	// This is to avoid potential problems with very old glslang versions which did
2711	// not emit input/output interfaces properly.
2712	// We can assume they only had a single entry point, and single entry point
2713	// shaders could easily be assumed to use every interface variable anyways.
2714	if (ir.entry_points.size() <= 1)
2715	return true;
2716	}
2717
2718	// In SPIR-V 1.4 and later, all global resource variables must be present.
2719
2720	auto &execution = get_entry_point();
2721	return find(first: begin(cont: execution.interface_variables), last: end(cont: execution.interface_variables), val: VariableID(id)) !=
2722	end(cont: execution.interface_variables);
2723	}
2724
2725	void Compiler::CombinedImageSamplerHandler::push_remap_parameters(const SPIRFunction &func, const uint32_t *args,
2726	uint32_t length)
2727	{
2728	// If possible, pipe through a remapping table so that parameters know
2729	// which variables they actually bind to in this scope.
2730	unordered_map<uint32_t, uint32_t> remapping;
2731	for (uint32_t i = 0; i < length; i++)
2732	remapping[func.arguments[i].id] = remap_parameter(id: args[i]);
2733	parameter_remapping.push(x: std::move(remapping));
2734	}
2735
2736	void Compiler::CombinedImageSamplerHandler::pop_remap_parameters()
2737	{
2738	parameter_remapping.pop();
2739	}
2740
2741	uint32_t Compiler::CombinedImageSamplerHandler::remap_parameter(uint32_t id)
2742	{
2743	auto *var = compiler.maybe_get_backing_variable(chain: id);
2744	if (var)
2745	id = var->self;
2746
2747	if (parameter_remapping.empty())
2748	return id;
2749
2750	auto &remapping = parameter_remapping.top();
2751	auto itr = remapping.find(x: id);
2752	if (itr != end(cont&: remapping))
2753	return itr->second;
2754	else
2755	return id;
2756	}
2757
2758	bool Compiler::CombinedImageSamplerHandler::begin_function_scope(const uint32_t *args, uint32_t length)
2759	{
2760	if (length < 3)
2761	return false;
2762
2763	auto &callee = compiler.get<SPIRFunction>(id: args[2]);
2764	args += 3;
2765	length -= 3;
2766	push_remap_parameters(func: callee, args, length);
2767	functions.push(x: &callee);
2768	return true;
2769	}
2770
2771	bool Compiler::CombinedImageSamplerHandler::end_function_scope(const uint32_t *args, uint32_t length)
2772	{
2773	if (length < 3)
2774	return false;
2775
2776	auto &callee = compiler.get<SPIRFunction>(id: args[2]);
2777	args += 3;
2778
2779	// There are two types of cases we have to handle,
2780	// a callee might call sampler2D(texture2D, sampler) directly where
2781	// one or more parameters originate from parameters.
2782	// Alternatively, we need to provide combined image samplers to our callees,
2783	// and in this case we need to add those as well.
2784
2785	pop_remap_parameters();
2786
2787	// Our callee has now been processed at least once.
2788	// No point in doing it again.
2789	callee.do_combined_parameters = false;
2790
2791	auto &params = functions.top()->combined_parameters;
2792	functions.pop();
2793	if (functions.empty())
2794	return true;
2795
2796	auto &caller = *functions.top();
2797	if (caller.do_combined_parameters)
2798	{
2799	for (auto &param : params)
2800	{
2801	VariableID image_id = param.global_image ? param.image_id : VariableID(args[param.image_id]);
2802	VariableID sampler_id = param.global_sampler ? param.sampler_id : VariableID(args[param.sampler_id]);
2803
2804	auto *i = compiler.maybe_get_backing_variable(chain: image_id);
2805	auto *s = compiler.maybe_get_backing_variable(chain: sampler_id);
2806	if (i)
2807	image_id = i->self;
2808	if (s)
2809	sampler_id = s->self;
2810
2811	register_combined_image_sampler(caller, combined_id: 0, texture_id: image_id, sampler_id, depth: param.depth);
2812	}
2813	}
2814
2815	return true;
2816	}
2817
2818	void Compiler::CombinedImageSamplerHandler::register_combined_image_sampler(SPIRFunction &caller,
2819	VariableID combined_module_id,
2820	VariableID image_id, VariableID sampler_id,
2821	bool depth)
2822	{
2823	// We now have a texture ID and a sampler ID which will either be found as a global
2824	// or a parameter in our own function. If both are global, they will not need a parameter,
2825	// otherwise, add it to our list.
2826	SPIRFunction::CombinedImageSamplerParameter param = {
2827	.id: 0u, .image_id: image_id, .sampler_id: sampler_id, .global_image: true, .global_sampler: true, .depth: depth,
2828	};
2829
2830	auto texture_itr = find_if(first: begin(cont&: caller.arguments), last: end(cont&: caller.arguments),
2831	pred: [image_id](const SPIRFunction::Parameter &p) { return p.id == image_id; });
2832	auto sampler_itr = find_if(first: begin(cont&: caller.arguments), last: end(cont&: caller.arguments),
2833	pred: [sampler_id](const SPIRFunction::Parameter &p) { return p.id == sampler_id; });
2834
2835	if (texture_itr != end(cont&: caller.arguments))
2836	{
2837	param.global_image = false;
2838	param.image_id = uint32_t(texture_itr - begin(cont&: caller.arguments));
2839	}
2840
2841	if (sampler_itr != end(cont&: caller.arguments))
2842	{
2843	param.global_sampler = false;
2844	param.sampler_id = uint32_t(sampler_itr - begin(cont&: caller.arguments));
2845	}
2846
2847	if (param.global_image && param.global_sampler)
2848	return;
2849
2850	auto itr = find_if(first: begin(cont&: caller.combined_parameters), last: end(cont&: caller.combined_parameters),
2851	pred: [&param](const SPIRFunction::CombinedImageSamplerParameter &p) {
2852	return param.image_id == p.image_id && param.sampler_id == p.sampler_id &&
2853	param.global_image == p.global_image && param.global_sampler == p.global_sampler;
2854	});
2855
2856	if (itr == end(cont&: caller.combined_parameters))
2857	{
2858	uint32_t id = compiler.ir.increase_bound_by(count: 3);
2859	auto type_id = id + 0;
2860	auto ptr_type_id = id + 1;
2861	auto combined_id = id + 2;
2862	auto &base = compiler.expression_type(id: image_id);
2863	auto &type = compiler.set<SPIRType>(id: type_id, args: OpTypeSampledImage);
2864	auto &ptr_type = compiler.set<SPIRType>(id: ptr_type_id, args: OpTypePointer);
2865
2866	type = base;
2867	type.self = type_id;
2868	type.basetype = SPIRType::SampledImage;
2869	type.pointer = false;
2870	type.storage = StorageClassGeneric;
2871	type.image.depth = depth;
2872
2873	ptr_type = type;
2874	ptr_type.pointer = true;
2875	ptr_type.storage = StorageClassUniformConstant;
2876	ptr_type.parent_type = type_id;
2877
2878	// Build new variable.
2879	compiler.set<SPIRVariable>(id: combined_id, args&: ptr_type_id, args: StorageClassFunction, args: 0);
2880
2881	// Inherit RelaxedPrecision.
2882	// If any of OpSampledImage, underlying image or sampler are marked, inherit the decoration.
2883	bool relaxed_precision =
2884	compiler.has_decoration(id: sampler_id, decoration: DecorationRelaxedPrecision) ||
2885	compiler.has_decoration(id: image_id, decoration: DecorationRelaxedPrecision) ||
2886	(combined_module_id && compiler.has_decoration(id: combined_module_id, decoration: DecorationRelaxedPrecision));
2887
2888	if (relaxed_precision)
2889	compiler.set_decoration(id: combined_id, decoration: DecorationRelaxedPrecision);
2890
2891	param.id = combined_id;
2892
2893	compiler.set_name(id: combined_id,
2894	name: join(ts: "SPIRV_Cross_Combined", ts: compiler.to_name(id: image_id), ts: compiler.to_name(id: sampler_id)));
2895
2896	caller.combined_parameters.push_back(t: param);
2897	caller.shadow_arguments.push_back(t: { .type: ptr_type_id, .id: combined_id, .read_count: 0u, .write_count: 0u, .alias_global_variable: true });
2898	}
2899	}
2900
2901	bool Compiler::DummySamplerForCombinedImageHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
2902	{
2903	if (need_dummy_sampler)
2904	{
2905	// No need to traverse further, we know the result.
2906	return false;
2907	}
2908
2909	switch (opcode)
2910	{
2911	case OpLoad:
2912	{
2913	if (length < 3)
2914	return false;
2915
2916	uint32_t result_type = args[0];
2917
2918	auto &type = compiler.get<SPIRType>(id: result_type);
2919	bool separate_image =
2920	type.basetype == SPIRType::Image && type.image.sampled == 1 && type.image.dim != DimBuffer;
2921
2922	// If not separate image, don't bother.
2923	if (!separate_image)
2924	return true;
2925
2926	uint32_t id = args[1];
2927	uint32_t ptr = args[2];
2928	compiler.set<SPIRExpression>(id, args: "", args&: result_type, args: true);
2929	compiler.register_read(expr: id, chain: ptr, forwarded: true);
2930	break;
2931	}
2932
2933	case OpImageFetch:
2934	case OpImageQuerySizeLod:
2935	case OpImageQuerySize:
2936	case OpImageQueryLevels:
2937	case OpImageQuerySamples:
2938	{
2939	// If we are fetching or querying LOD from a plain OpTypeImage, we must pre-combine with our dummy sampler.
2940	auto *var = compiler.maybe_get_backing_variable(chain: args[2]);
2941	if (var)
2942	{
2943	auto &type = compiler.get<SPIRType>(id: var->basetype);
2944	if (type.basetype == SPIRType::Image && type.image.sampled == 1 && type.image.dim != DimBuffer)
2945	need_dummy_sampler = true;
2946	}
2947
2948	break;
2949	}
2950
2951	case OpInBoundsAccessChain:
2952	case OpAccessChain:
2953	case OpPtrAccessChain:
2954	{
2955	if (length < 3)
2956	return false;
2957
2958	uint32_t result_type = args[0];
2959	auto &type = compiler.get<SPIRType>(id: result_type);
2960	bool separate_image =
2961	type.basetype == SPIRType::Image && type.image.sampled == 1 && type.image.dim != DimBuffer;
2962	if (!separate_image)
2963	return true;
2964
2965	uint32_t id = args[1];
2966	uint32_t ptr = args[2];
2967	compiler.set<SPIRExpression>(id, args: "", args&: result_type, args: true);
2968	compiler.register_read(expr: id, chain: ptr, forwarded: true);
2969
2970	// Other backends might use SPIRAccessChain for this later.
2971	compiler.ir.ids[id].set_allow_type_rewrite();
2972	break;
2973	}
2974
2975	default:
2976	break;
2977	}
2978
2979	return true;
2980	}
2981
2982	bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
2983	{
2984	// We need to figure out where samplers and images are loaded from, so do only the bare bones compilation we need.
2985	bool is_fetch = false;
2986
2987	switch (opcode)
2988	{
2989	case OpLoad:
2990	{
2991	if (length < 3)
2992	return false;
2993
2994	uint32_t result_type = args[0];
2995
2996	auto &type = compiler.get<SPIRType>(id: result_type);
2997	bool separate_image = type.basetype == SPIRType::Image && type.image.sampled == 1;
2998	bool separate_sampler = type.basetype == SPIRType::Sampler;
2999
3000	// If not separate image or sampler, don't bother.
3001	if (!separate_image && !separate_sampler)
3002	return true;
3003
3004	uint32_t id = args[1];
3005	uint32_t ptr = args[2];
3006	compiler.set<SPIRExpression>(id, args: "", args&: result_type, args: true);
3007	compiler.register_read(expr: id, chain: ptr, forwarded: true);
3008	return true;
3009	}
3010
3011	case OpInBoundsAccessChain:
3012	case OpAccessChain:
3013	case OpPtrAccessChain:
3014	{
3015	if (length < 3)
3016	return false;
3017
3018	// Technically, it is possible to have arrays of textures and arrays of samplers and combine them, but this becomes essentially
3019	// impossible to implement, since we don't know which concrete sampler we are accessing.
3020	// One potential way is to create a combinatorial explosion where N textures and M samplers are combined into N * M sampler2Ds,
3021	// but this seems ridiculously complicated for a problem which is easy to work around.
3022	// Checking access chains like this assumes we don't have samplers or textures inside uniform structs, but this makes no sense.
3023
3024	uint32_t result_type = args[0];
3025
3026	auto &type = compiler.get<SPIRType>(id: result_type);
3027	bool separate_image = type.basetype == SPIRType::Image && type.image.sampled == 1;
3028	bool separate_sampler = type.basetype == SPIRType::Sampler;
3029	if (separate_sampler)
3030	SPIRV_CROSS_THROW(
3031	"Attempting to use arrays or structs of separate samplers. This is not possible to statically "
3032	"remap to plain GLSL.");
3033
3034	if (separate_image)
3035	{
3036	uint32_t id = args[1];
3037	uint32_t ptr = args[2];
3038	compiler.set<SPIRExpression>(id, args: "", args&: result_type, args: true);
3039	compiler.register_read(expr: id, chain: ptr, forwarded: true);
3040	}
3041	return true;
3042	}
3043
3044	case OpImageFetch:
3045	case OpImageQuerySizeLod:
3046	case OpImageQuerySize:
3047	case OpImageQueryLevels:
3048	case OpImageQuerySamples:
3049	{
3050	// If we are fetching from a plain OpTypeImage or querying LOD, we must pre-combine with our dummy sampler.
3051	auto *var = compiler.maybe_get_backing_variable(chain: args[2]);
3052	if (!var)
3053	return true;
3054
3055	auto &type = compiler.get<SPIRType>(id: var->basetype);
3056	if (type.basetype == SPIRType::Image && type.image.sampled == 1 && type.image.dim != DimBuffer)
3057	{
3058	if (compiler.dummy_sampler_id == 0)
3059	SPIRV_CROSS_THROW("texelFetch without sampler was found, but no dummy sampler has been created with "
3060	"build_dummy_sampler_for_combined_images().");
3061
3062	// Do it outside.
3063	is_fetch = true;
3064	break;
3065	}
3066
3067	return true;
3068	}
3069
3070	case OpSampledImage:
3071	// Do it outside.
3072	break;
3073
3074	default:
3075	return true;
3076	}
3077
3078	// Registers sampler2D calls used in case they are parameters so
3079	// that their callees know which combined image samplers to propagate down the call stack.
3080	if (!functions.empty())
3081	{
3082	auto &callee = *functions.top();
3083	if (callee.do_combined_parameters)
3084	{
3085	uint32_t image_id = args[2];
3086
3087	auto *image = compiler.maybe_get_backing_variable(chain: image_id);
3088	if (image)
3089	image_id = image->self;
3090
3091	uint32_t sampler_id = is_fetch ? compiler.dummy_sampler_id : args[3];
3092	auto *sampler = compiler.maybe_get_backing_variable(chain: sampler_id);
3093	if (sampler)
3094	sampler_id = sampler->self;
3095
3096	uint32_t combined_id = args[1];
3097
3098	auto &combined_type = compiler.get<SPIRType>(id: args[0]);
3099	register_combined_image_sampler(caller&: callee, combined_module_id: combined_id, image_id, sampler_id, depth: combined_type.image.depth);
3100	}
3101	}
3102
3103	// For function calls, we need to remap IDs which are function parameters into global variables.
3104	// This information is statically known from the current place in the call stack.
3105	// Function parameters are not necessarily pointers, so if we don't have a backing variable, remapping will know
3106	// which backing variable the image/sample came from.
3107	VariableID image_id = remap_parameter(id: args[2]);
3108	VariableID sampler_id = is_fetch ? compiler.dummy_sampler_id : remap_parameter(id: args[3]);
3109
3110	auto itr = find_if(first: begin(cont&: compiler.combined_image_samplers), last: end(cont&: compiler.combined_image_samplers),
3111	pred: [image_id, sampler_id](const CombinedImageSampler &combined) {
3112	return combined.image_id == image_id && combined.sampler_id == sampler_id;
3113	});
3114
3115	if (itr == end(cont&: compiler.combined_image_samplers))
3116	{
3117	uint32_t sampled_type;
3118	uint32_t combined_module_id;
3119	if (is_fetch)
3120	{
3121	// Have to invent the sampled image type.
3122	sampled_type = compiler.ir.increase_bound_by(count: 1);
3123	auto &type = compiler.set<SPIRType>(id: sampled_type, args: OpTypeSampledImage);
3124	type = compiler.expression_type(id: args[2]);
3125	type.self = sampled_type;
3126	type.basetype = SPIRType::SampledImage;
3127	type.image.depth = false;
3128	combined_module_id = 0;
3129	}
3130	else
3131	{
3132	sampled_type = args[0];
3133	combined_module_id = args[1];
3134	}
3135
3136	auto id = compiler.ir.increase_bound_by(count: 2);
3137	auto type_id = id + 0;
3138	auto combined_id = id + 1;
3139
3140	// Make a new type, pointer to OpTypeSampledImage, so we can make a variable of this type.
3141	// We will probably have this type lying around, but it doesn't hurt to make duplicates for internal purposes.
3142	auto &type = compiler.set<SPIRType>(id: type_id, args: OpTypePointer);
3143	auto &base = compiler.get<SPIRType>(id: sampled_type);
3144	type = base;
3145	type.pointer = true;
3146	type.storage = StorageClassUniformConstant;
3147	type.parent_type = type_id;
3148
3149	// Build new variable.
3150	compiler.set<SPIRVariable>(id: combined_id, args&: type_id, args: StorageClassUniformConstant, args: 0);
3151
3152	// Inherit RelaxedPrecision (and potentially other useful flags if deemed relevant).
3153	// If any of OpSampledImage, underlying image or sampler are marked, inherit the decoration.
3154	bool relaxed_precision =
3155	(sampler_id && compiler.has_decoration(id: sampler_id, decoration: DecorationRelaxedPrecision)) ||
3156	(image_id && compiler.has_decoration(id: image_id, decoration: DecorationRelaxedPrecision)) ||
3157	(combined_module_id && compiler.has_decoration(id: combined_module_id, decoration: DecorationRelaxedPrecision));
3158
3159	if (relaxed_precision)
3160	compiler.set_decoration(id: combined_id, decoration: DecorationRelaxedPrecision);
3161
3162	// Propagate the array type for the original image as well.
3163	auto *var = compiler.maybe_get_backing_variable(chain: image_id);
3164	if (var)
3165	{
3166	auto &parent_type = compiler.get<SPIRType>(id: var->basetype);
3167	type.array = parent_type.array;
3168	type.array_size_literal = parent_type.array_size_literal;
3169	}
3170
3171	compiler.combined_image_samplers.push_back(t: { .combined_id: combined_id, .image_id: image_id, .sampler_id: sampler_id });
3172	}
3173
3174	return true;
3175	}
3176
3177	VariableID Compiler::build_dummy_sampler_for_combined_images()
3178	{
3179	DummySamplerForCombinedImageHandler handler(*this);
3180	traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler);
3181	if (handler.need_dummy_sampler)
3182	{
3183	uint32_t offset = ir.increase_bound_by(count: 3);
3184	auto type_id = offset + 0;
3185	auto ptr_type_id = offset + 1;
3186	auto var_id = offset + 2;
3187
3188	auto &sampler = set<SPIRType>(id: type_id, args: OpTypeSampler);
3189	sampler.basetype = SPIRType::Sampler;
3190
3191	auto &ptr_sampler = set<SPIRType>(id: ptr_type_id, args: OpTypePointer);
3192	ptr_sampler = sampler;
3193	ptr_sampler.self = type_id;
3194	ptr_sampler.storage = StorageClassUniformConstant;
3195	ptr_sampler.pointer = true;
3196	ptr_sampler.parent_type = type_id;
3197
3198	set<SPIRVariable>(id: var_id, args&: ptr_type_id, args: StorageClassUniformConstant, args: 0);
3199	set_name(id: var_id, name: "SPIRV_Cross_DummySampler");
3200	dummy_sampler_id = var_id;
3201	return var_id;
3202	}
3203	else
3204	return 0;
3205	}
3206
3207	void Compiler::build_combined_image_samplers()
3208	{
3209	ir.for_each_typed_id<SPIRFunction>(op: [&](uint32_t, SPIRFunction &func) {
3210	func.combined_parameters.clear();
3211	func.shadow_arguments.clear();
3212	func.do_combined_parameters = true;
3213	});
3214
3215	combined_image_samplers.clear();
3216	CombinedImageSamplerHandler handler(*this);
3217	traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler);
3218	}
3219
3220	SmallVector<SpecializationConstant> Compiler::get_specialization_constants() const
3221	{
3222	SmallVector<SpecializationConstant> spec_consts;
3223	ir.for_each_typed_id<SPIRConstant>(op: [&](uint32_t, const SPIRConstant &c) {
3224	if (c.specialization && has_decoration(id: c.self, decoration: DecorationSpecId))
3225	spec_consts.push_back(t: { .id: c.self, .constant_id: get_decoration(id: c.self, decoration: DecorationSpecId) });
3226	});
3227	return spec_consts;
3228	}
3229
3230	SPIRConstant &Compiler::get_constant(ConstantID id)
3231	{
3232	return get<SPIRConstant>(id);
3233	}
3234
3235	const SPIRConstant &Compiler::get_constant(ConstantID id) const
3236	{
3237	return get<SPIRConstant>(id);
3238	}
3239
3240	static bool exists_unaccessed_path_to_return(const CFG &cfg, uint32_t block, const unordered_set<uint32_t> &blocks,
3241	unordered_set<uint32_t> &visit_cache)
3242	{
3243	// This block accesses the variable.
3244	if (blocks.find(x: block) != end(cont: blocks))
3245	return false;
3246
3247	// We are at the end of the CFG.
3248	if (cfg.get_succeeding_edges(block).empty())
3249	return true;
3250
3251	// If any of our successors have a path to the end, there exists a path from block.
3252	for (auto &succ : cfg.get_succeeding_edges(block))
3253	{
3254	if (visit_cache.count(x: succ) == 0)
3255	{
3256	if (exists_unaccessed_path_to_return(cfg, block: succ, blocks, visit_cache))
3257	return true;
3258	visit_cache.insert(x: succ);
3259	}
3260	}
3261
3262	return false;
3263	}
3264
3265	void Compiler::analyze_parameter_preservation(
3266	SPIRFunction &entry, const CFG &cfg, const unordered_map<uint32_t, unordered_set<uint32_t>> &variable_to_blocks,
3267	const unordered_map<uint32_t, unordered_set<uint32_t>> &complete_write_blocks)
3268	{
3269	for (auto &arg : entry.arguments)
3270	{
3271	// Non-pointers are always inputs.
3272	auto &type = get<SPIRType>(id: arg.type);
3273	if (!type.pointer)
3274	continue;
3275
3276	// Opaque argument types are always in
3277	bool potential_preserve;
3278	switch (type.basetype)
3279	{
3280	case SPIRType::Sampler:
3281	case SPIRType::Image:
3282	case SPIRType::SampledImage:
3283	case SPIRType::AtomicCounter:
3284	potential_preserve = false;
3285	break;
3286
3287	default:
3288	potential_preserve = true;
3289	break;
3290	}
3291
3292	if (!potential_preserve)
3293	continue;
3294
3295	auto itr = variable_to_blocks.find(x: arg.id);
3296	if (itr == end(cont: variable_to_blocks))
3297	{
3298	// Variable is never accessed.
3299	continue;
3300	}
3301
3302	// We have accessed a variable, but there was no complete writes to that variable.
3303	// We deduce that we must preserve the argument.
3304	itr = complete_write_blocks.find(x: arg.id);
3305	if (itr == end(cont: complete_write_blocks))
3306	{
3307	arg.read_count++;
3308	continue;
3309	}
3310
3311	// If there is a path through the CFG where no block completely writes to the variable, the variable will be in an undefined state
3312	// when the function returns. We therefore need to implicitly preserve the variable in case there are writers in the function.
3313	// Major case here is if a function is
3314	// void foo(int &var) { if (cond) var = 10; }
3315	// Using read/write counts, we will think it's just an out variable, but it really needs to be inout,
3316	// because if we don't write anything whatever we put into the function must return back to the caller.
3317	unordered_set<uint32_t> visit_cache;
3318	if (exists_unaccessed_path_to_return(cfg, block: entry.entry_block, blocks: itr->second, visit_cache))
3319	arg.read_count++;
3320	}
3321	}
3322
3323	Compiler::AnalyzeVariableScopeAccessHandler::AnalyzeVariableScopeAccessHandler(Compiler &compiler_,
3324	SPIRFunction &entry_)
3325	: compiler(compiler_)
3326	, entry(entry_)
3327	{
3328	}
3329
3330	bool Compiler::AnalyzeVariableScopeAccessHandler::follow_function_call(const SPIRFunction &)
3331	{
3332	// Only analyze within this function.
3333	return false;
3334	}
3335
3336	void Compiler::AnalyzeVariableScopeAccessHandler::set_current_block(const SPIRBlock &block)
3337	{
3338	current_block = &block;
3339
3340	// If we're branching to a block which uses OpPhi, in GLSL
3341	// this will be a variable write when we branch,
3342	// so we need to track access to these variables as well to
3343	// have a complete picture.
3344	const auto test_phi = [this, &block](uint32_t to) {
3345	auto &next = compiler.get<SPIRBlock>(id: to);
3346	for (auto &phi : next.phi_variables)
3347	{
3348	if (phi.parent == block.self)
3349	{
3350	accessed_variables_to_block[phi.function_variable].insert(x: block.self);
3351	// Phi variables are also accessed in our target branch block.
3352	accessed_variables_to_block[phi.function_variable].insert(x: next.self);
3353
3354	notify_variable_access(id: phi.local_variable, block: block.self);
3355	}
3356	}
3357	};
3358
3359	switch (block.terminator)
3360	{
3361	case SPIRBlock::Direct:
3362	notify_variable_access(id: block.condition, block: block.self);
3363	test_phi(block.next_block);
3364	break;
3365
3366	case SPIRBlock::Select:
3367	notify_variable_access(id: block.condition, block: block.self);
3368	test_phi(block.true_block);
3369	test_phi(block.false_block);
3370	break;
3371
3372	case SPIRBlock::MultiSelect:
3373	{
3374	notify_variable_access(id: block.condition, block: block.self);
3375	auto &cases = compiler.get_case_list(block);
3376	for (auto &target : cases)
3377	test_phi(target.block);
3378	if (block.default_block)
3379	test_phi(block.default_block);
3380	break;
3381	}
3382
3383	default:
3384	break;
3385	}
3386	}
3387
3388	void Compiler::AnalyzeVariableScopeAccessHandler::notify_variable_access(uint32_t id, uint32_t block)
3389	{
3390	if (id == 0)
3391	return;
3392
3393	// Access chains used in multiple blocks mean hoisting all the variables used to construct the access chain as not all backends can use pointers.
3394	auto itr = rvalue_forward_children.find(x: id);
3395	if (itr != end(cont&: rvalue_forward_children))
3396	for (auto child_id : itr->second)
3397	notify_variable_access(id: child_id, block);
3398
3399	if (id_is_phi_variable(id))
3400	accessed_variables_to_block[id].insert(x: block);
3401	else if (id_is_potential_temporary(id))
3402	accessed_temporaries_to_block[id].insert(x: block);
3403	}
3404
3405	bool Compiler::AnalyzeVariableScopeAccessHandler::id_is_phi_variable(uint32_t id) const
3406	{
3407	if (id >= compiler.get_current_id_bound())
3408	return false;
3409	auto *var = compiler.maybe_get<SPIRVariable>(id);
3410	return var && var->phi_variable;
3411	}
3412
3413	bool Compiler::AnalyzeVariableScopeAccessHandler::id_is_potential_temporary(uint32_t id) const
3414	{
3415	if (id >= compiler.get_current_id_bound())
3416	return false;
3417
3418	// Temporaries are not created before we start emitting code.
3419	return compiler.ir.ids[id].empty() || (compiler.ir.ids[id].get_type() == TypeExpression);
3420	}
3421
3422	bool Compiler::AnalyzeVariableScopeAccessHandler::handle_terminator(const SPIRBlock &block)
3423	{
3424	switch (block.terminator)
3425	{
3426	case SPIRBlock::Return:
3427	if (block.return_value)
3428	notify_variable_access(id: block.return_value, block: block.self);
3429	break;
3430
3431	case SPIRBlock::Select:
3432	case SPIRBlock::MultiSelect:
3433	notify_variable_access(id: block.condition, block: block.self);
3434	break;
3435
3436	default:
3437	break;
3438	}
3439
3440	return true;
3441	}
3442
3443	bool Compiler::AnalyzeVariableScopeAccessHandler::handle(spv::Op op, const uint32_t *args, uint32_t length)
3444	{
3445	// Keep track of the types of temporaries, so we can hoist them out as necessary.
3446	uint32_t result_type = 0, result_id = 0;
3447	if (compiler.instruction_to_result_type(result_type, result_id, op, args, length))
3448	{
3449	// For some opcodes, we will need to override the result id.
3450	// If we need to hoist the temporary, the temporary type is the input, not the result.
3451	if (op == OpConvertUToAccelerationStructureKHR)
3452	{
3453	auto itr = result_id_to_type.find(x: args[2]);
3454	if (itr != result_id_to_type.end())
3455	result_type = itr->second;
3456	}
3457
3458	result_id_to_type[result_id] = result_type;
3459	}
3460
3461	switch (op)
3462	{
3463	case OpStore:
3464	{
3465	if (length < 2)
3466	return false;
3467
3468	ID ptr = args[0];
3469	auto *var = compiler.maybe_get_backing_variable(chain: ptr);
3470
3471	// If we store through an access chain, we have a partial write.
3472	if (var)
3473	{
3474	accessed_variables_to_block[var->self].insert(x: current_block->self);
3475	if (var->self == ptr)
3476	complete_write_variables_to_block[var->self].insert(x: current_block->self);
3477	else
3478	partial_write_variables_to_block[var->self].insert(x: current_block->self);
3479	}
3480
3481	// args[0] might be an access chain we have to track use of.
3482	notify_variable_access(id: args[0], block: current_block->self);
3483	// Might try to store a Phi variable here.
3484	notify_variable_access(id: args[1], block: current_block->self);
3485	break;
3486	}
3487
3488	case OpAccessChain:
3489	case OpInBoundsAccessChain:
3490	case OpPtrAccessChain:
3491	{
3492	if (length < 3)
3493	return false;
3494
3495	// Access chains used in multiple blocks mean hoisting all the variables used to construct the access chain as not all backends can use pointers.
3496	uint32_t ptr = args[2];
3497	auto *var = compiler.maybe_get<SPIRVariable>(id: ptr);
3498	if (var)
3499	{
3500	accessed_variables_to_block[var->self].insert(x: current_block->self);
3501	rvalue_forward_children[args[1]].insert(x: var->self);
3502	}
3503
3504	// args[2] might be another access chain we have to track use of.
3505	for (uint32_t i = 2; i < length; i++)
3506	{
3507	notify_variable_access(id: args[i], block: current_block->self);
3508	rvalue_forward_children[args[1]].insert(x: args[i]);
3509	}
3510
3511	// Also keep track of the access chain pointer itself.
3512	// In exceptionally rare cases, we can end up with a case where
3513	// the access chain is generated in the loop body, but is consumed in continue block.
3514	// This means we need complex loop workarounds, and we must detect this via CFG analysis.
3515	notify_variable_access(id: args[1], block: current_block->self);
3516
3517	// The result of an access chain is a fixed expression and is not really considered a temporary.
3518	auto &e = compiler.set<SPIRExpression>(id: args[1], args: "", args: args[0], args: true);
3519	auto *backing_variable = compiler.maybe_get_backing_variable(chain: ptr);
3520	e.loaded_from = backing_variable ? VariableID(backing_variable->self) : VariableID(0);
3521
3522	// Other backends might use SPIRAccessChain for this later.
3523	compiler.ir.ids[args[1]].set_allow_type_rewrite();
3524	access_chain_expressions.insert(x: args[1]);
3525	break;
3526	}
3527
3528	case OpCopyMemory:
3529	{
3530	if (length < 2)
3531	return false;
3532
3533	ID lhs = args[0];
3534	ID rhs = args[1];
3535	auto *var = compiler.maybe_get_backing_variable(chain: lhs);
3536
3537	// If we store through an access chain, we have a partial write.
3538	if (var)
3539	{
3540	accessed_variables_to_block[var->self].insert(x: current_block->self);
3541	if (var->self == lhs)
3542	complete_write_variables_to_block[var->self].insert(x: current_block->self);
3543	else
3544	partial_write_variables_to_block[var->self].insert(x: current_block->self);
3545	}
3546
3547	// args[0:1] might be access chains we have to track use of.
3548	for (uint32_t i = 0; i < 2; i++)
3549	notify_variable_access(id: args[i], block: current_block->self);
3550
3551	var = compiler.maybe_get_backing_variable(chain: rhs);
3552	if (var)
3553	accessed_variables_to_block[var->self].insert(x: current_block->self);
3554	break;
3555	}
3556
3557	case OpCopyObject:
3558	{
3559	// OpCopyObject copies the underlying non-pointer type,
3560	// so any temp variable should be declared using the underlying type.
3561	// If the type is a pointer, get its base type and overwrite the result type mapping.
3562	auto &type = compiler.get<SPIRType>(id: result_type);
3563	if (type.pointer)
3564	result_id_to_type[result_id] = type.parent_type;
3565
3566	if (length < 3)
3567	return false;
3568
3569	auto *var = compiler.maybe_get_backing_variable(chain: args[2]);
3570	if (var)
3571	accessed_variables_to_block[var->self].insert(x: current_block->self);
3572
3573	// Might be an access chain which we have to keep track of.
3574	notify_variable_access(id: args[1], block: current_block->self);
3575	if (access_chain_expressions.count(x: args[2]))
3576	access_chain_expressions.insert(x: args[1]);
3577
3578	// Might try to copy a Phi variable here.
3579	notify_variable_access(id: args[2], block: current_block->self);
3580	break;
3581	}
3582
3583	case OpLoad:
3584	{
3585	if (length < 3)
3586	return false;
3587	uint32_t ptr = args[2];
3588	auto *var = compiler.maybe_get_backing_variable(chain: ptr);
3589	if (var)
3590	accessed_variables_to_block[var->self].insert(x: current_block->self);
3591
3592	// Loaded value is a temporary.
3593	notify_variable_access(id: args[1], block: current_block->self);
3594
3595	// Might be an access chain we have to track use of.
3596	notify_variable_access(id: args[2], block: current_block->self);
3597
3598	// If we're loading an opaque type we cannot lower it to a temporary,
3599	// we must defer access of args[2] until it's used.
3600	auto &type = compiler.get<SPIRType>(id: args[0]);
3601	if (compiler.type_is_opaque_value(type))
3602	rvalue_forward_children[args[1]].insert(x: args[2]);
3603	break;
3604	}
3605
3606	case OpFunctionCall:
3607	{
3608	if (length < 3)
3609	return false;
3610
3611	// Return value may be a temporary.
3612	if (compiler.get_type(id: args[0]).basetype != SPIRType::Void)
3613	notify_variable_access(id: args[1], block: current_block->self);
3614
3615	length -= 3;
3616	args += 3;
3617
3618	for (uint32_t i = 0; i < length; i++)
3619	{
3620	auto *var = compiler.maybe_get_backing_variable(chain: args[i]);
3621	if (var)
3622	{
3623	accessed_variables_to_block[var->self].insert(x: current_block->self);
3624	// Assume we can get partial writes to this variable.
3625	partial_write_variables_to_block[var->self].insert(x: current_block->self);
3626	}
3627
3628	// Cannot easily prove if argument we pass to a function is completely written.
3629	// Usually, functions write to a dummy variable,
3630	// which is then copied to in full to the real argument.
3631
3632	// Might try to copy a Phi variable here.
3633	notify_variable_access(id: args[i], block: current_block->self);
3634	}
3635	break;
3636	}
3637
3638	case OpSelect:
3639	{
3640	// In case of variable pointers, we might access a variable here.
3641	// We cannot prove anything about these accesses however.
3642	for (uint32_t i = 1; i < length; i++)
3643	{
3644	if (i >= 3)
3645	{
3646	auto *var = compiler.maybe_get_backing_variable(chain: args[i]);
3647	if (var)
3648	{
3649	accessed_variables_to_block[var->self].insert(x: current_block->self);
3650	// Assume we can get partial writes to this variable.
3651	partial_write_variables_to_block[var->self].insert(x: current_block->self);
3652	}
3653	}
3654
3655	// Might try to copy a Phi variable here.
3656	notify_variable_access(id: args[i], block: current_block->self);
3657	}
3658	break;
3659	}
3660
3661	case OpExtInst:
3662	{
3663	for (uint32_t i = 4; i < length; i++)
3664	notify_variable_access(id: args[i], block: current_block->self);
3665	notify_variable_access(id: args[1], block: current_block->self);
3666
3667	uint32_t extension_set = args[2];
3668	if (compiler.get<SPIRExtension>(id: extension_set).ext == SPIRExtension::GLSL)
3669	{
3670	auto op_450 = static_cast<GLSLstd450>(args[3]);
3671	switch (op_450)
3672	{
3673	case GLSLstd450Modf:
3674	case GLSLstd450Frexp:
3675	{
3676	uint32_t ptr = args[5];
3677	auto *var = compiler.maybe_get_backing_variable(chain: ptr);
3678	if (var)
3679	{
3680	accessed_variables_to_block[var->self].insert(x: current_block->self);
3681	if (var->self == ptr)
3682	complete_write_variables_to_block[var->self].insert(x: current_block->self);
3683	else
3684	partial_write_variables_to_block[var->self].insert(x: current_block->self);
3685	}
3686	break;
3687	}
3688
3689	default:
3690	break;
3691	}
3692	}
3693	break;
3694	}
3695
3696	case OpArrayLength:
3697	// Only result is a temporary.
3698	notify_variable_access(id: args[1], block: current_block->self);
3699	break;
3700
3701	case OpLine:
3702	case OpNoLine:
3703	// Uses literals, but cannot be a phi variable or temporary, so ignore.
3704	break;
3705
3706	// Atomics shouldn't be able to access function-local variables.
3707	// Some GLSL builtins access a pointer.
3708
3709	case OpCompositeInsert:
3710	case OpVectorShuffle:
3711	// Specialize for opcode which contains literals.
3712	for (uint32_t i = 1; i < 4; i++)
3713	notify_variable_access(id: args[i], block: current_block->self);
3714	break;
3715
3716	case OpCompositeExtract:
3717	// Specialize for opcode which contains literals.
3718	for (uint32_t i = 1; i < 3; i++)
3719	notify_variable_access(id: args[i], block: current_block->self);
3720	break;
3721
3722	case OpImageWrite:
3723	for (uint32_t i = 0; i < length; i++)
3724	{
3725	// Argument 3 is a literal.
3726	if (i != 3)
3727	notify_variable_access(id: args[i], block: current_block->self);
3728	}
3729	break;
3730
3731	case OpImageSampleImplicitLod:
3732	case OpImageSampleExplicitLod:
3733	case OpImageSparseSampleImplicitLod:
3734	case OpImageSparseSampleExplicitLod:
3735	case OpImageSampleProjImplicitLod:
3736	case OpImageSampleProjExplicitLod:
3737	case OpImageSparseSampleProjImplicitLod:
3738	case OpImageSparseSampleProjExplicitLod:
3739	case OpImageFetch:
3740	case OpImageSparseFetch:
3741	case OpImageRead:
3742	case OpImageSparseRead:
3743	for (uint32_t i = 1; i < length; i++)
3744	{
3745	// Argument 4 is a literal.
3746	if (i != 4)
3747	notify_variable_access(id: args[i], block: current_block->self);
3748	}
3749	break;
3750
3751	case OpImageSampleDrefImplicitLod:
3752	case OpImageSampleDrefExplicitLod:
3753	case OpImageSparseSampleDrefImplicitLod:
3754	case OpImageSparseSampleDrefExplicitLod:
3755	case OpImageSampleProjDrefImplicitLod:
3756	case OpImageSampleProjDrefExplicitLod:
3757	case OpImageSparseSampleProjDrefImplicitLod:
3758	case OpImageSparseSampleProjDrefExplicitLod:
3759	case OpImageGather:
3760	case OpImageSparseGather:
3761	case OpImageDrefGather:
3762	case OpImageSparseDrefGather:
3763	for (uint32_t i = 1; i < length; i++)
3764	{
3765	// Argument 5 is a literal.
3766	if (i != 5)
3767	notify_variable_access(id: args[i], block: current_block->self);
3768	}
3769	break;
3770
3771	default:
3772	{
3773	// Rather dirty way of figuring out where Phi variables are used.
3774	// As long as only IDs are used, we can scan through instructions and try to find any evidence that
3775	// the ID of a variable has been used.
3776	// There are potential false positives here where a literal is used in-place of an ID,
3777	// but worst case, it does not affect the correctness of the compile.
3778	// Exhaustive analysis would be better here, but it's not worth it for now.
3779	for (uint32_t i = 0; i < length; i++)
3780	notify_variable_access(id: args[i], block: current_block->self);
3781	break;
3782	}
3783	}
3784	return true;
3785	}
3786
3787	Compiler::StaticExpressionAccessHandler::StaticExpressionAccessHandler(Compiler &compiler_, uint32_t variable_id_)
3788	: compiler(compiler_)
3789	, variable_id(variable_id_)
3790	{
3791	}
3792
3793	bool Compiler::StaticExpressionAccessHandler::follow_function_call(const SPIRFunction &)
3794	{
3795	return false;
3796	}
3797
3798	bool Compiler::StaticExpressionAccessHandler::handle(spv::Op op, const uint32_t *args, uint32_t length)
3799	{
3800	switch (op)
3801	{
3802	case OpStore:
3803	if (length < 2)
3804	return false;
3805	if (args[0] == variable_id)
3806	{
3807	static_expression = args[1];
3808	write_count++;
3809	}
3810	break;
3811
3812	case OpLoad:
3813	if (length < 3)
3814	return false;
3815	if (args[2] == variable_id && static_expression == 0) // Tried to read from variable before it was initialized.
3816	return false;
3817	break;
3818
3819	case OpAccessChain:
3820	case OpInBoundsAccessChain:
3821	case OpPtrAccessChain:
3822	if (length < 3)
3823	return false;
3824	if (args[2] == variable_id) // If we try to access chain our candidate variable before we store to it, bail.
3825	return false;
3826	break;
3827
3828	default:
3829	break;
3830	}
3831
3832	return true;
3833	}
3834
3835	void Compiler::find_function_local_luts(SPIRFunction &entry, const AnalyzeVariableScopeAccessHandler &handler,
3836	bool single_function)
3837	{
3838	auto &cfg = *function_cfgs.find(x: entry.self)->second;
3839
3840	// For each variable which is statically accessed.
3841	for (auto &accessed_var : handler.accessed_variables_to_block)
3842	{
3843	auto &blocks = accessed_var.second;
3844	auto &var = get<SPIRVariable>(id: accessed_var.first);
3845	auto &type = expression_type(id: accessed_var.first);
3846
3847	// First check if there are writes to the variable. Later, if there are none, we'll
3848	// reconsider it as globally accessed LUT.
3849	if (!var.is_written_to)
3850	{
3851	var.is_written_to = handler.complete_write_variables_to_block.count(x: var.self) != 0 ||
3852	handler.partial_write_variables_to_block.count(x: var.self) != 0;
3853	}
3854
3855	// Only consider function local variables here.
3856	// If we only have a single function in our CFG, private storage is also fine,
3857	// since it behaves like a function local variable.
3858	bool allow_lut = var.storage == StorageClassFunction || (single_function && var.storage == StorageClassPrivate);
3859	if (!allow_lut)
3860	continue;
3861
3862	// We cannot be a phi variable.
3863	if (var.phi_variable)
3864	continue;
3865
3866	// Only consider arrays here.
3867	if (type.array.empty())
3868	continue;
3869
3870	// If the variable has an initializer, make sure it is a constant expression.
3871	uint32_t static_constant_expression = 0;
3872	if (var.initializer)
3873	{
3874	if (ir.ids[var.initializer].get_type() != TypeConstant)
3875	continue;
3876	static_constant_expression = var.initializer;
3877
3878	// There can be no stores to this variable, we have now proved we have a LUT.
3879	if (var.is_written_to)
3880	continue;
3881	}
3882	else
3883	{
3884	// We can have one, and only one write to the variable, and that write needs to be a constant.
3885
3886	// No partial writes allowed.
3887	if (handler.partial_write_variables_to_block.count(x: var.self) != 0)
3888	continue;
3889
3890	auto itr = handler.complete_write_variables_to_block.find(x: var.self);
3891
3892	// No writes?
3893	if (itr == end(cont: handler.complete_write_variables_to_block))
3894	continue;
3895
3896	// We write to the variable in more than one block.
3897	auto &write_blocks = itr->second;
3898	if (write_blocks.size() != 1)
3899	continue;
3900
3901	// The write needs to happen in the dominating block.
3902	DominatorBuilder builder(cfg);
3903	for (auto &block : blocks)
3904	builder.add_block(block);
3905	uint32_t dominator = builder.get_dominator();
3906
3907	// The complete write happened in a branch or similar, cannot deduce static expression.
3908	if (write_blocks.count(x: dominator) == 0)
3909	continue;
3910
3911	// Find the static expression for this variable.
3912	StaticExpressionAccessHandler static_expression_handler(*this, var.self);
3913	traverse_all_reachable_opcodes(block: get<SPIRBlock>(id: dominator), handler&: static_expression_handler);
3914
3915	// We want one, and exactly one write
3916	if (static_expression_handler.write_count != 1 || static_expression_handler.static_expression == 0)
3917	continue;
3918
3919	// Is it a constant expression?
3920	if (ir.ids[static_expression_handler.static_expression].get_type() != TypeConstant)
3921	continue;
3922
3923	// We found a LUT!
3924	static_constant_expression = static_expression_handler.static_expression;
3925	}
3926
3927	get<SPIRConstant>(id: static_constant_expression).is_used_as_lut = true;
3928	var.static_expression = static_constant_expression;
3929	var.statically_assigned = true;
3930	var.remapped_variable = true;
3931	}
3932	}
3933
3934	void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeAccessHandler &handler)
3935	{
3936	// First, we map out all variable access within a function.
3937	// Essentially a map of block -> { variables accessed in the basic block }
3938	traverse_all_reachable_opcodes(func: entry, handler);
3939
3940	auto &cfg = *function_cfgs.find(x: entry.self)->second;
3941
3942	// Analyze if there are parameters which need to be implicitly preserved with an "in" qualifier.
3943	analyze_parameter_preservation(entry, cfg, variable_to_blocks: handler.accessed_variables_to_block,
3944	complete_write_blocks: handler.complete_write_variables_to_block);
3945
3946	unordered_map<uint32_t, uint32_t> potential_loop_variables;
3947
3948	// Find the loop dominator block for each block.
3949	for (auto &block_id : entry.blocks)
3950	{
3951	auto &block = get<SPIRBlock>(id: block_id);
3952
3953	auto itr = ir.continue_block_to_loop_header.find(x: block_id);
3954	if (itr != end(cont&: ir.continue_block_to_loop_header) && itr->second != block_id)
3955	{
3956	// Continue block might be unreachable in the CFG, but we still like to know the loop dominator.
3957	// Edge case is when continue block is also the loop header, don't set the dominator in this case.
3958	block.loop_dominator = itr->second;
3959	}
3960	else
3961	{
3962	uint32_t loop_dominator = cfg.find_loop_dominator(block: block_id);
3963	if (loop_dominator != block_id)
3964	block.loop_dominator = loop_dominator;
3965	else
3966	block.loop_dominator = SPIRBlock::NoDominator;
3967	}
3968	}
3969
3970	// For each variable which is statically accessed.
3971	for (auto &var : handler.accessed_variables_to_block)
3972	{
3973	// Only deal with variables which are considered local variables in this function.
3974	if (find(first: begin(cont&: entry.local_variables), last: end(cont&: entry.local_variables), val: VariableID(var.first)) ==
3975	end(cont&: entry.local_variables))
3976	continue;
3977
3978	DominatorBuilder builder(cfg);
3979	auto &blocks = var.second;
3980	auto &type = expression_type(id: var.first);
3981	BlockID potential_continue_block = 0;
3982
3983	// Figure out which block is dominating all accesses of those variables.
3984	for (auto &block : blocks)
3985	{
3986	// If we're accessing a variable inside a continue block, this variable might be a loop variable.
3987	// We can only use loop variables with scalars, as we cannot track static expressions for vectors.
3988	if (is_continue(next: block))
3989	{
3990	// Potentially awkward case to check for.
3991	// We might have a variable inside a loop, which is touched by the continue block,
3992	// but is not actually a loop variable.
3993	// The continue block is dominated by the inner part of the loop, which does not make sense in high-level
3994	// language output because it will be declared before the body,
3995	// so we will have to lift the dominator up to the relevant loop header instead.
3996	builder.add_block(block: ir.continue_block_to_loop_header[block]);
3997
3998	// Arrays or structs cannot be loop variables.
3999	if (type.vecsize == 1 && type.columns == 1 && type.basetype != SPIRType::Struct && type.array.empty())
4000	{
4001	// The variable is used in multiple continue blocks, this is not a loop
4002	// candidate, signal that by setting block to -1u.
4003	if (potential_continue_block == 0)
4004	potential_continue_block = block;
4005	else
4006	potential_continue_block = ~(0u);
4007	}
4008	}
4009
4010	builder.add_block(block);
4011	}
4012
4013	builder.lift_continue_block_dominator();
4014
4015	// Add it to a per-block list of variables.
4016	BlockID dominating_block = builder.get_dominator();
4017
4018	if (dominating_block && potential_continue_block != 0 && potential_continue_block != ~0u)
4019	{
4020	auto &inner_block = get<SPIRBlock>(id: dominating_block);
4021
4022	BlockID merge_candidate = 0;
4023
4024	// Analyze the dominator. If it lives in a different loop scope than the candidate continue
4025	// block, reject the loop variable candidate.
4026	if (inner_block.merge == SPIRBlock::MergeLoop)
4027	merge_candidate = inner_block.merge_block;
4028	else if (inner_block.loop_dominator != SPIRBlock::NoDominator)
4029	merge_candidate = get<SPIRBlock>(id: inner_block.loop_dominator).merge_block;
4030
4031	if (merge_candidate != 0 && cfg.is_reachable(block: merge_candidate))
4032	{
4033	// If the merge block has a higher post-visit order, we know that continue candidate
4034	// cannot reach the merge block, and we have two separate scopes.
4035	if (!cfg.is_reachable(block: potential_continue_block) ||
4036	cfg.get_visit_order(block: merge_candidate) > cfg.get_visit_order(block: potential_continue_block))
4037	{
4038	potential_continue_block = 0;
4039	}
4040	}
4041	}
4042
4043	if (potential_continue_block != 0 && potential_continue_block != ~0u)
4044	potential_loop_variables[var.first] = potential_continue_block;
4045
4046	// For variables whose dominating block is inside a loop, there is a risk that these variables
4047	// actually need to be preserved across loop iterations. We can express this by adding
4048	// a "read" access to the loop header.
4049	// In the dominating block, we must see an OpStore or equivalent as the first access of an OpVariable.
4050	// Should that fail, we look for the outermost loop header and tack on an access there.
4051	// Phi nodes cannot have this problem.
4052	if (dominating_block)
4053	{
4054	auto &variable = get<SPIRVariable>(id: var.first);
4055	if (!variable.phi_variable)
4056	{
4057	auto *block = &get<SPIRBlock>(id: dominating_block);
4058	bool preserve = may_read_undefined_variable_in_block(block: *block, var: var.first);
4059	if (preserve)
4060	{
4061	// Find the outermost loop scope.
4062	while (block->loop_dominator != BlockID(SPIRBlock::NoDominator))
4063	block = &get<SPIRBlock>(id: block->loop_dominator);
4064
4065	if (block->self != dominating_block)
4066	{
4067	builder.add_block(block: block->self);
4068	dominating_block = builder.get_dominator();
4069	}
4070	}
4071	}
4072	}
4073
4074	// If all blocks here are dead code, this will be 0, so the variable in question
4075	// will be completely eliminated.
4076	if (dominating_block)
4077	{
4078	auto &block = get<SPIRBlock>(id: dominating_block);
4079	block.dominated_variables.push_back(t: var.first);
4080	get<SPIRVariable>(id: var.first).dominator = dominating_block;
4081	}
4082	}
4083
4084	for (auto &var : handler.accessed_temporaries_to_block)
4085	{
4086	auto itr = handler.result_id_to_type.find(x: var.first);
4087
4088	if (itr == end(cont&: handler.result_id_to_type))
4089	{
4090	// We found a false positive ID being used, ignore.
4091	// This should probably be an assert.
4092	continue;
4093	}
4094
4095	// There is no point in doing domination analysis for opaque types.
4096	auto &type = get<SPIRType>(id: itr->second);
4097	if (type_is_opaque_value(type))
4098	continue;
4099
4100	DominatorBuilder builder(cfg);
4101	bool force_temporary = false;
4102	bool used_in_header_hoisted_continue_block = false;
4103
4104	// Figure out which block is dominating all accesses of those temporaries.
4105	auto &blocks = var.second;
4106	for (auto &block : blocks)
4107	{
4108	builder.add_block(block);
4109
4110	if (blocks.size() != 1 && is_continue(next: block))
4111	{
4112	// The risk here is that inner loop can dominate the continue block.
4113	// Any temporary we access in the continue block must be declared before the loop.
4114	// This is moot for complex loops however.
4115	auto &loop_header_block = get<SPIRBlock>(id: ir.continue_block_to_loop_header[block]);
4116	assert(loop_header_block.merge == SPIRBlock::MergeLoop);
4117	builder.add_block(block: loop_header_block.self);
4118	used_in_header_hoisted_continue_block = true;
4119	}
4120	}
4121
4122	uint32_t dominating_block = builder.get_dominator();
4123
4124	if (blocks.size() != 1 && is_single_block_loop(next: dominating_block))
4125	{
4126	// Awkward case, because the loop header is also the continue block,
4127	// so hoisting to loop header does not help.
4128	force_temporary = true;
4129	}
4130
4131	if (dominating_block)
4132	{
4133	// If we touch a variable in the dominating block, this is the expected setup.
4134	// SPIR-V normally mandates this, but we have extra cases for temporary use inside loops.
4135	bool first_use_is_dominator = blocks.count(x: dominating_block) != 0;
4136
4137	if (!first_use_is_dominator || force_temporary)
4138	{
4139	if (handler.access_chain_expressions.count(x: var.first))
4140	{
4141	// Exceptionally rare case.
4142	// We cannot declare temporaries of access chains (except on MSL perhaps with pointers).
4143	// Rather than do that, we force the indexing expressions to be declared in the right scope by
4144	// tracking their usage to that end. There is no temporary to hoist.
4145	// However, we still need to observe declaration order of the access chain.
4146
4147	if (used_in_header_hoisted_continue_block)
4148	{
4149	// For this scenario, we used an access chain inside a continue block where we also registered an access to header block.
4150	// This is a problem as we need to declare an access chain properly first with full definition.
4151	// We cannot use temporaries for these expressions,
4152	// so we must make sure the access chain is declared ahead of time.
4153	// Force a complex for loop to deal with this.
4154	// TODO: Out-of-order declaring for loops where continue blocks are emitted last might be another option.
4155	auto &loop_header_block = get<SPIRBlock>(id: dominating_block);
4156	assert(loop_header_block.merge == SPIRBlock::MergeLoop);
4157	loop_header_block.complex_continue = true;
4158	}
4159	}
4160	else
4161	{
4162	// This should be very rare, but if we try to declare a temporary inside a loop,
4163	// and that temporary is used outside the loop as well (spirv-opt inliner likes this)
4164	// we should actually emit the temporary outside the loop.
4165	hoisted_temporaries.insert(x: var.first);
4166	forced_temporaries.insert(x: var.first);
4167
4168	auto &block_temporaries = get<SPIRBlock>(id: dominating_block).declare_temporary;
4169	block_temporaries.emplace_back(ts&: handler.result_id_to_type[var.first], ts: var.first);
4170	}
4171	}
4172	else if (blocks.size() > 1)
4173	{
4174	// Keep track of the temporary as we might have to declare this temporary.
4175	// This can happen if the loop header dominates a temporary, but we have a complex fallback loop.
4176	// In this case, the header is actually inside the for (;;) {} block, and we have problems.
4177	// What we need to do is hoist the temporaries outside the for (;;) {} block in case the header block
4178	// declares the temporary.
4179	auto &block_temporaries = get<SPIRBlock>(id: dominating_block).potential_declare_temporary;
4180	block_temporaries.emplace_back(ts&: handler.result_id_to_type[var.first], ts: var.first);
4181	}
4182	}
4183	}
4184
4185	unordered_set<uint32_t> seen_blocks;
4186
4187	// Now, try to analyze whether or not these variables are actually loop variables.
4188	for (auto &loop_variable : potential_loop_variables)
4189	{
4190	auto &var = get<SPIRVariable>(id: loop_variable.first);
4191	auto dominator = var.dominator;
4192	BlockID block = loop_variable.second;
4193
4194	// The variable was accessed in multiple continue blocks, ignore.
4195	if (block == BlockID(~(0u)) || block == BlockID(0))
4196	continue;
4197
4198	// Dead code.
4199	if (dominator == ID(0))
4200	continue;
4201
4202	BlockID header = 0;
4203
4204	// Find the loop header for this block if we are a continue block.
4205	{
4206	auto itr = ir.continue_block_to_loop_header.find(x: block);
4207	if (itr != end(cont&: ir.continue_block_to_loop_header))
4208	{
4209	header = itr->second;
4210	}
4211	else if (get<SPIRBlock>(id: block).continue_block == block)
4212	{
4213	// Also check for self-referential continue block.
4214	header = block;
4215	}
4216	}
4217
4218	assert(header);
4219	auto &header_block = get<SPIRBlock>(id: header);
4220	auto &blocks = handler.accessed_variables_to_block[loop_variable.first];
4221
4222	// If a loop variable is not used before the loop, it's probably not a loop variable.
4223	bool has_accessed_variable = blocks.count(x: header) != 0;
4224
4225	// Now, there are two conditions we need to meet for the variable to be a loop variable.
4226	// 1. The dominating block must have a branch-free path to the loop header,
4227	// this way we statically know which expression should be part of the loop variable initializer.
4228
4229	// Walk from the dominator, if there is one straight edge connecting
4230	// dominator and loop header, we statically know the loop initializer.
4231	bool static_loop_init = true;
4232	while (dominator != header)
4233	{
4234	if (blocks.count(x: dominator) != 0)
4235	has_accessed_variable = true;
4236
4237	auto &succ = cfg.get_succeeding_edges(block: dominator);
4238	if (succ.size() != 1)
4239	{
4240	static_loop_init = false;
4241	break;
4242	}
4243
4244	auto &pred = cfg.get_preceding_edges(block: succ.front());
4245	if (pred.size() != 1 || pred.front() != dominator)
4246	{
4247	static_loop_init = false;
4248	break;
4249	}
4250
4251	dominator = succ.front();
4252	}
4253
4254	if (!static_loop_init || !has_accessed_variable)
4255	continue;
4256
4257	// The second condition we need to meet is that no access after the loop
4258	// merge can occur. Walk the CFG to see if we find anything.
4259
4260	seen_blocks.clear();
4261	cfg.walk_from(seen_blocks, block: header_block.merge_block, op: [&](uint32_t walk_block) -> bool {
4262	// We found a block which accesses the variable outside the loop.
4263	if (blocks.find(x: walk_block) != end(cont&: blocks))
4264	static_loop_init = false;
4265	return true;
4266	});
4267
4268	if (!static_loop_init)
4269	continue;
4270
4271	// We have a loop variable.
4272	header_block.loop_variables.push_back(t: loop_variable.first);
4273	// Need to sort here as variables come from an unordered container, and pushing stuff in wrong order
4274	// will break reproducability in regression runs.
4275	sort(first: begin(cont&: header_block.loop_variables), last: end(cont&: header_block.loop_variables));
4276	get<SPIRVariable>(id: loop_variable.first).loop_variable = true;
4277	}
4278	}
4279
4280	bool Compiler::may_read_undefined_variable_in_block(const SPIRBlock &block, uint32_t var)
4281	{
4282	for (auto &op : block.ops)
4283	{
4284	auto *ops = stream(instr: op);
4285	switch (op.op)
4286	{
4287	case OpStore:
4288	case OpCopyMemory:
4289	if (ops[0] == var)
4290	return false;
4291	break;
4292
4293	case OpAccessChain:
4294	case OpInBoundsAccessChain:
4295	case OpPtrAccessChain:
4296	// Access chains are generally used to partially read and write. It's too hard to analyze
4297	// if all constituents are written fully before continuing, so just assume it's preserved.
4298	// This is the same as the parameter preservation analysis.
4299	if (ops[2] == var)
4300	return true;
4301	break;
4302
4303	case OpSelect:
4304	// Variable pointers.
4305	// We might read before writing.
4306	if (ops[3] == var || ops[4] == var)
4307	return true;
4308	break;
4309
4310	case OpPhi:
4311	{
4312	// Variable pointers.
4313	// We might read before writing.
4314	if (op.length < 2)
4315	break;
4316
4317	uint32_t count = op.length - 2;
4318	for (uint32_t i = 0; i < count; i += 2)
4319	if (ops[i + 2] == var)
4320	return true;
4321	break;
4322	}
4323
4324	case OpCopyObject:
4325	case OpLoad:
4326	if (ops[2] == var)
4327	return true;
4328	break;
4329
4330	case OpFunctionCall:
4331	{
4332	if (op.length < 3)
4333	break;
4334
4335	// May read before writing.
4336	uint32_t count = op.length - 3;
4337	for (uint32_t i = 0; i < count; i++)
4338	if (ops[i + 3] == var)
4339	return true;
4340	break;
4341	}
4342
4343	default:
4344	break;
4345	}
4346	}
4347
4348	// Not accessed somehow, at least not in a usual fashion.
4349	// It's likely accessed in a branch, so assume we must preserve.
4350	return true;
4351	}
4352
4353	Bitset Compiler::get_buffer_block_flags(VariableID id) const
4354	{
4355	return ir.get_buffer_block_flags(var: get<SPIRVariable>(id));
4356	}
4357
4358	bool Compiler::get_common_basic_type(const SPIRType &type, SPIRType::BaseType &base_type)
4359	{
4360	if (type.basetype == SPIRType::Struct)
4361	{
4362	base_type = SPIRType::Unknown;
4363	for (auto &member_type : type.member_types)
4364	{
4365	SPIRType::BaseType member_base;
4366	if (!get_common_basic_type(type: get<SPIRType>(id: member_type), base_type&: member_base))
4367	return false;
4368
4369	if (base_type == SPIRType::Unknown)
4370	base_type = member_base;
4371	else if (base_type != member_base)
4372	return false;
4373	}
4374	return true;
4375	}
4376	else
4377	{
4378	base_type = type.basetype;
4379	return true;
4380	}
4381	}
4382
4383	void Compiler::ActiveBuiltinHandler::handle_builtin(const SPIRType &type, BuiltIn builtin,
4384	const Bitset &decoration_flags)
4385	{
4386	// If used, we will need to explicitly declare a new array size for these builtins.
4387
4388	if (builtin == BuiltInClipDistance)
4389	{
4390	if (!type.array_size_literal[0])
4391	SPIRV_CROSS_THROW("Array size for ClipDistance must be a literal.");
4392	uint32_t array_size = type.array[0];
4393	if (array_size == 0)
4394	SPIRV_CROSS_THROW("Array size for ClipDistance must not be unsized.");
4395	compiler.clip_distance_count = array_size;
4396	}
4397	else if (builtin == BuiltInCullDistance)
4398	{
4399	if (!type.array_size_literal[0])
4400	SPIRV_CROSS_THROW("Array size for CullDistance must be a literal.");
4401	uint32_t array_size = type.array[0];
4402	if (array_size == 0)
4403	SPIRV_CROSS_THROW("Array size for CullDistance must not be unsized.");
4404	compiler.cull_distance_count = array_size;
4405	}
4406	else if (builtin == BuiltInPosition)
4407	{
4408	if (decoration_flags.get(bit: DecorationInvariant))
4409	compiler.position_invariant = true;
4410	}
4411	}
4412
4413	void Compiler::ActiveBuiltinHandler::add_if_builtin(uint32_t id, bool allow_blocks)
4414	{
4415	// Only handle plain variables here.
4416	// Builtins which are part of a block are handled in AccessChain.
4417	// If allow_blocks is used however, this is to handle initializers of blocks,
4418	// which implies that all members are written to.
4419
4420	auto *var = compiler.maybe_get<SPIRVariable>(id);
4421	auto *m = compiler.ir.find_meta(id);
4422	if (var && m)
4423	{
4424	auto &type = compiler.get<SPIRType>(id: var->basetype);
4425	auto &decorations = m->decoration;
4426	auto &flags = type.storage == StorageClassInput ?
4427	compiler.active_input_builtins : compiler.active_output_builtins;
4428	if (decorations.builtin)
4429	{
4430	flags.set(decorations.builtin_type);
4431	handle_builtin(type, builtin: decorations.builtin_type, decoration_flags: decorations.decoration_flags);
4432	}
4433	else if (allow_blocks && compiler.has_decoration(id: type.self, decoration: DecorationBlock))
4434	{
4435	uint32_t member_count = uint32_t(type.member_types.size());
4436	for (uint32_t i = 0; i < member_count; i++)
4437	{
4438	if (compiler.has_member_decoration(id: type.self, index: i, decoration: DecorationBuiltIn))
4439	{
4440	auto &member_type = compiler.get<SPIRType>(id: type.member_types[i]);
4441	BuiltIn builtin = BuiltIn(compiler.get_member_decoration(id: type.self, index: i, decoration: DecorationBuiltIn));
4442	flags.set(builtin);
4443	handle_builtin(type: member_type, builtin, decoration_flags: compiler.get_member_decoration_bitset(id: type.self, index: i));
4444	}
4445	}
4446	}
4447	}
4448	}
4449
4450	void Compiler::ActiveBuiltinHandler::add_if_builtin(uint32_t id)
4451	{
4452	add_if_builtin(id, allow_blocks: false);
4453	}
4454
4455	void Compiler::ActiveBuiltinHandler::add_if_builtin_or_block(uint32_t id)
4456	{
4457	add_if_builtin(id, allow_blocks: true);
4458	}
4459
4460	bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args, uint32_t length)
4461	{
4462	switch (opcode)
4463	{
4464	case OpStore:
4465	if (length < 1)
4466	return false;
4467
4468	add_if_builtin(id: args[0]);
4469	break;
4470
4471	case OpCopyMemory:
4472	if (length < 2)
4473	return false;
4474
4475	add_if_builtin(id: args[0]);
4476	add_if_builtin(id: args[1]);
4477	break;
4478
4479	case OpCopyObject:
4480	case OpLoad:
4481	if (length < 3)
4482	return false;
4483
4484	add_if_builtin(id: args[2]);
4485	break;
4486
4487	case OpSelect:
4488	if (length < 5)
4489	return false;
4490
4491	add_if_builtin(id: args[3]);
4492	add_if_builtin(id: args[4]);
4493	break;
4494
4495	case OpPhi:
4496	{
4497	if (length < 2)
4498	return false;
4499
4500	uint32_t count = length - 2;
4501	args += 2;
4502	for (uint32_t i = 0; i < count; i += 2)
4503	add_if_builtin(id: args[i]);
4504	break;
4505	}
4506
4507	case OpFunctionCall:
4508	{
4509	if (length < 3)
4510	return false;
4511
4512	uint32_t count = length - 3;
4513	args += 3;
4514	for (uint32_t i = 0; i < count; i++)
4515	add_if_builtin(id: args[i]);
4516	break;
4517	}
4518
4519	case OpAccessChain:
4520	case OpInBoundsAccessChain:
4521	case OpPtrAccessChain:
4522	{
4523	if (length < 4)
4524	return false;
4525
4526	// Only consider global variables, cannot consider variables in functions yet, or other
4527	// access chains as they have not been created yet.
4528	auto *var = compiler.maybe_get<SPIRVariable>(id: args[2]);
4529	if (!var)
4530	break;
4531
4532	// Required if we access chain into builtins like gl_GlobalInvocationID.
4533	add_if_builtin(id: args[2]);
4534
4535	// Start traversing type hierarchy at the proper non-pointer types.
4536	auto *type = &compiler.get_variable_data_type(var: *var);
4537
4538	auto &flags =
4539	var->storage == StorageClassInput ? compiler.active_input_builtins : compiler.active_output_builtins;
4540
4541	uint32_t count = length - 3;
4542	args += 3;
4543	for (uint32_t i = 0; i < count; i++)
4544	{
4545	// Pointers
4546	// PtrAccessChain functions more like a pointer offset. Type remains the same.
4547	if (opcode == OpPtrAccessChain && i == 0)
4548	continue;
4549
4550	// Arrays
4551	if (!type->array.empty())
4552	{
4553	type = &compiler.get<SPIRType>(id: type->parent_type);
4554	}
4555	// Structs
4556	else if (type->basetype == SPIRType::Struct)
4557	{
4558	uint32_t index = compiler.get<SPIRConstant>(id: args[i]).scalar();
4559
4560	if (index < uint32_t(compiler.ir.meta[type->self].members.size()))
4561	{
4562	auto &decorations = compiler.ir.meta[type->self].members[index];
4563	if (decorations.builtin)
4564	{
4565	flags.set(decorations.builtin_type);
4566	handle_builtin(type: compiler.get<SPIRType>(id: type->member_types[index]), builtin: decorations.builtin_type,
4567	decoration_flags: decorations.decoration_flags);
4568	}
4569	}
4570
4571	type = &compiler.get<SPIRType>(id: type->member_types[index]);
4572	}
4573	else
4574	{
4575	// No point in traversing further. We won't find any extra builtins.
4576	break;
4577	}
4578	}
4579	break;
4580	}
4581
4582	default:
4583	break;
4584	}
4585
4586	return true;
4587	}
4588
4589	void Compiler::update_active_builtins()
4590	{
4591	active_input_builtins.reset();
4592	active_output_builtins.reset();
4593	cull_distance_count = 0;
4594	clip_distance_count = 0;
4595	ActiveBuiltinHandler handler(*this);
4596	traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler);
4597
4598	ir.for_each_typed_id<SPIRVariable>(op: [&](uint32_t, const SPIRVariable &var) {
4599	if (var.storage != StorageClassOutput)
4600	return;
4601	if (!interface_variable_exists_in_entry_point(id: var.self))
4602	return;
4603
4604	// Also, make sure we preserve output variables which are only initialized, but never accessed by any code.
4605	if (var.initializer != ID(0))
4606	handler.add_if_builtin_or_block(id: var.self);
4607	});
4608	}
4609
4610	// Returns whether this shader uses a builtin of the storage class
4611	bool Compiler::has_active_builtin(BuiltIn builtin, StorageClass storage) const
4612	{
4613	const Bitset *flags;
4614	switch (storage)
4615	{
4616	case StorageClassInput:
4617	flags = &active_input_builtins;
4618	break;
4619	case StorageClassOutput:
4620	flags = &active_output_builtins;
4621	break;
4622
4623	default:
4624	return false;
4625	}
4626	return flags->get(bit: builtin);
4627	}
4628
4629	void Compiler::analyze_image_and_sampler_usage()
4630	{
4631	CombinedImageSamplerDrefHandler dref_handler(*this);
4632	traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler&: dref_handler);
4633
4634	CombinedImageSamplerUsageHandler handler(*this, dref_handler.dref_combined_samplers);
4635	traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler);
4636
4637	// Need to run this traversal twice. First time, we propagate any comparison sampler usage from leaf functions
4638	// down to main().
4639	// In the second pass, we can propagate up forced depth state coming from main() up into leaf functions.
4640	handler.dependency_hierarchy.clear();
4641	traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler);
4642
4643	comparison_ids = std::move(handler.comparison_ids);
4644	need_subpass_input = handler.need_subpass_input;
4645	need_subpass_input_ms = handler.need_subpass_input_ms;
4646
4647	// Forward information from separate images and samplers into combined image samplers.
4648	for (auto &combined : combined_image_samplers)
4649	if (comparison_ids.count(x: combined.sampler_id))
4650	comparison_ids.insert(x: combined.combined_id);
4651	}
4652
4653	bool Compiler::CombinedImageSamplerDrefHandler::handle(spv::Op opcode, const uint32_t *args, uint32_t)
4654	{
4655	// Mark all sampled images which are used with Dref.
4656	switch (opcode)
4657	{
4658	case OpImageSampleDrefExplicitLod:
4659	case OpImageSampleDrefImplicitLod:
4660	case OpImageSampleProjDrefExplicitLod:
4661	case OpImageSampleProjDrefImplicitLod:
4662	case OpImageSparseSampleProjDrefImplicitLod:
4663	case OpImageSparseSampleDrefImplicitLod:
4664	case OpImageSparseSampleProjDrefExplicitLod:
4665	case OpImageSparseSampleDrefExplicitLod:
4666	case OpImageDrefGather:
4667	case OpImageSparseDrefGather:
4668	dref_combined_samplers.insert(x: args[2]);
4669	return true;
4670
4671	default:
4672	break;
4673	}
4674
4675	return true;
4676	}
4677
4678	const CFG &Compiler::get_cfg_for_current_function() const
4679	{
4680	assert(current_function);
4681	return get_cfg_for_function(id: current_function->self);
4682	}
4683
4684	const CFG &Compiler::get_cfg_for_function(uint32_t id) const
4685	{
4686	auto cfg_itr = function_cfgs.find(x: id);
4687	assert(cfg_itr != end(function_cfgs));
4688	assert(cfg_itr->second);
4689	return *cfg_itr->second;
4690	}
4691
4692	void Compiler::build_function_control_flow_graphs_and_analyze()
4693	{
4694	CFGBuilder handler(*this);
4695	handler.function_cfgs[ir.default_entry_point].reset(p: new CFG(*this, get<SPIRFunction>(id: ir.default_entry_point)));
4696	traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler);
4697	function_cfgs = std::move(handler.function_cfgs);
4698	bool single_function = function_cfgs.size() <= 1;
4699
4700	for (auto &f : function_cfgs)
4701	{
4702	auto &func = get<SPIRFunction>(id: f.first);
4703	AnalyzeVariableScopeAccessHandler scope_handler(*this, func);
4704	analyze_variable_scope(entry&: func, handler&: scope_handler);
4705	find_function_local_luts(entry&: func, handler: scope_handler, single_function);
4706
4707	// Check if we can actually use the loop variables we found in analyze_variable_scope.
4708	// To use multiple initializers, we need the same type and qualifiers.
4709	for (auto block : func.blocks)
4710	{
4711	auto &b = get<SPIRBlock>(id: block);
4712	if (b.loop_variables.size() < 2)
4713	continue;
4714
4715	auto &flags = get_decoration_bitset(id: b.loop_variables.front());
4716	uint32_t type = get<SPIRVariable>(id: b.loop_variables.front()).basetype;
4717	bool invalid_initializers = false;
4718	for (auto loop_variable : b.loop_variables)
4719	{
4720	if (flags != get_decoration_bitset(id: loop_variable) ||
4721	type != get<SPIRVariable>(id: b.loop_variables.front()).basetype)
4722	{
4723	invalid_initializers = true;
4724	break;
4725	}
4726	}
4727
4728	if (invalid_initializers)
4729	{
4730	for (auto loop_variable : b.loop_variables)
4731	get<SPIRVariable>(id: loop_variable).loop_variable = false;
4732	b.loop_variables.clear();
4733	}
4734	}
4735	}
4736
4737	// Find LUTs which are not function local. Only consider this case if the CFG is multi-function,
4738	// otherwise we treat Private as Function trivially.
4739	// Needs to be analyzed from the outside since we have to block the LUT optimization if at least
4740	// one function writes to it.
4741	if (!single_function)
4742	{
4743	for (auto &id : global_variables)
4744	{
4745	auto &var = get<SPIRVariable>(id);
4746	auto &type = get_variable_data_type(var);
4747
4748	if (is_array(type) && var.storage == StorageClassPrivate &&
4749	var.initializer && !var.is_written_to &&
4750	ir.ids[var.initializer].get_type() == TypeConstant)
4751	{
4752	get<SPIRConstant>(id: var.initializer).is_used_as_lut = true;
4753	var.static_expression = var.initializer;
4754	var.statically_assigned = true;
4755	var.remapped_variable = true;
4756	}
4757	}
4758	}
4759	}
4760
4761	Compiler::CFGBuilder::CFGBuilder(Compiler &compiler_)
4762	: compiler(compiler_)
4763	{
4764	}
4765
4766	bool Compiler::CFGBuilder::handle(spv::Op, const uint32_t *, uint32_t)
4767	{
4768	return true;
4769	}
4770
4771	bool Compiler::CFGBuilder::follow_function_call(const SPIRFunction &func)
4772	{
4773	if (function_cfgs.find(x: func.self) == end(cont&: function_cfgs))
4774	{
4775	function_cfgs[func.self].reset(p: new CFG(compiler, func));
4776	return true;
4777	}
4778	else
4779	return false;
4780	}
4781
4782	void Compiler::CombinedImageSamplerUsageHandler::add_dependency(uint32_t dst, uint32_t src)
4783	{
4784	dependency_hierarchy[dst].insert(x: src);
4785	// Propagate up any comparison state if we're loading from one such variable.
4786	if (comparison_ids.count(x: src))
4787	comparison_ids.insert(x: dst);
4788	}
4789
4790	bool Compiler::CombinedImageSamplerUsageHandler::begin_function_scope(const uint32_t *args, uint32_t length)
4791	{
4792	if (length < 3)
4793	return false;
4794
4795	auto &func = compiler.get<SPIRFunction>(id: args[2]);
4796	const auto *arg = &args[3];
4797	length -= 3;
4798
4799	for (uint32_t i = 0; i < length; i++)
4800	{
4801	auto &argument = func.arguments[i];
4802	add_dependency(dst: argument.id, src: arg[i]);
4803	}
4804
4805	return true;
4806	}
4807
4808	void Compiler::CombinedImageSamplerUsageHandler::add_hierarchy_to_comparison_ids(uint32_t id)
4809	{
4810	// Traverse the variable dependency hierarchy and tag everything in its path with comparison ids.
4811	comparison_ids.insert(x: id);
4812
4813	for (auto &dep_id : dependency_hierarchy[id])
4814	add_hierarchy_to_comparison_ids(id: dep_id);
4815	}
4816
4817	bool Compiler::CombinedImageSamplerUsageHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
4818	{
4819	switch (opcode)
4820	{
4821	case OpAccessChain:
4822	case OpInBoundsAccessChain:
4823	case OpPtrAccessChain:
4824	case OpLoad:
4825	{
4826	if (length < 3)
4827	return false;
4828
4829	add_dependency(dst: args[1], src: args[2]);
4830
4831	// Ideally defer this to OpImageRead, but then we'd need to track loaded IDs.
4832	// If we load an image, we're going to use it and there is little harm in declaring an unused gl_FragCoord.
4833	auto &type = compiler.get<SPIRType>(id: args[0]);
4834	if (type.image.dim == DimSubpassData)
4835	{
4836	need_subpass_input = true;
4837	if (type.image.ms)
4838	need_subpass_input_ms = true;
4839	}
4840
4841	// If we load a SampledImage and it will be used with Dref, propagate the state up.
4842	if (dref_combined_samplers.count(x: args[1]) != 0)
4843	add_hierarchy_to_comparison_ids(id: args[1]);
4844	break;
4845	}
4846
4847	case OpSampledImage:
4848	{
4849	if (length < 4)
4850	return false;
4851
4852	// If the underlying resource has been used for comparison then duplicate loads of that resource must be too.
4853	// This image must be a depth image.
4854	uint32_t result_id = args[1];
4855	uint32_t image = args[2];
4856	uint32_t sampler = args[3];
4857
4858	if (dref_combined_samplers.count(x: result_id) != 0)
4859	{
4860	add_hierarchy_to_comparison_ids(id: image);
4861
4862	// This sampler must be a SamplerComparisonState, and not a regular SamplerState.
4863	add_hierarchy_to_comparison_ids(id: sampler);
4864
4865	// Mark the OpSampledImage itself as being comparison state.
4866	comparison_ids.insert(x: result_id);
4867	}
4868	return true;
4869	}
4870
4871	default:
4872	break;
4873	}
4874
4875	return true;
4876	}
4877
4878	bool Compiler::buffer_is_hlsl_counter_buffer(VariableID id) const
4879	{
4880	auto *m = ir.find_meta(id);
4881	return m && m->hlsl_is_magic_counter_buffer;
4882	}
4883
4884	bool Compiler::buffer_get_hlsl_counter_buffer(VariableID id, uint32_t &counter_id) const
4885	{
4886	auto *m = ir.find_meta(id);
4887
4888	// First, check for the proper decoration.
4889	if (m && m->hlsl_magic_counter_buffer != 0)
4890	{
4891	counter_id = m->hlsl_magic_counter_buffer;
4892	return true;
4893	}
4894	else
4895	return false;
4896	}
4897
4898	void Compiler::make_constant_null(uint32_t id, uint32_t type)
4899	{
4900	auto &constant_type = get<SPIRType>(id: type);
4901
4902	if (constant_type.pointer)
4903	{
4904	auto &constant = set<SPIRConstant>(id, args&: type);
4905	constant.make_null(constant_type_: constant_type);
4906	}
4907	else if (!constant_type.array.empty())
4908	{
4909	assert(constant_type.parent_type);
4910	uint32_t parent_id = ir.increase_bound_by(count: 1);
4911	make_constant_null(id: parent_id, type: constant_type.parent_type);
4912
4913	if (!constant_type.array_size_literal.back())
4914	SPIRV_CROSS_THROW("Array size of OpConstantNull must be a literal.");
4915
4916	SmallVector<uint32_t> elements(constant_type.array.back());
4917	for (uint32_t i = 0; i < constant_type.array.back(); i++)
4918	elements[i] = parent_id;
4919	set<SPIRConstant>(id, args&: type, args: elements.data(), args: uint32_t(elements.size()), args: false);
4920	}
4921	else if (!constant_type.member_types.empty())
4922	{
4923	uint32_t member_ids = ir.increase_bound_by(count: uint32_t(constant_type.member_types.size()));
4924	SmallVector<uint32_t> elements(constant_type.member_types.size());
4925	for (uint32_t i = 0; i < constant_type.member_types.size(); i++)
4926	{
4927	make_constant_null(id: member_ids + i, type: constant_type.member_types[i]);
4928	elements[i] = member_ids + i;
4929	}
4930	set<SPIRConstant>(id, args&: type, args: elements.data(), args: uint32_t(elements.size()), args: false);
4931	}
4932	else
4933	{
4934	auto &constant = set<SPIRConstant>(id, args&: type);
4935	constant.make_null(constant_type_: constant_type);
4936	}
4937	}
4938
4939	const SmallVector<spv::Capability> &Compiler::get_declared_capabilities() const
4940	{
4941	return ir.declared_capabilities;
4942	}
4943
4944	const SmallVector<std::string> &Compiler::get_declared_extensions() const
4945	{
4946	return ir.declared_extensions;
4947	}
4948
4949	std::string Compiler::get_remapped_declared_block_name(VariableID id) const
4950	{
4951	return get_remapped_declared_block_name(id, fallback_prefer_instance_name: false);
4952	}
4953
4954	std::string Compiler::get_remapped_declared_block_name(uint32_t id, bool fallback_prefer_instance_name) const
4955	{
4956	auto itr = declared_block_names.find(x: id);
4957	if (itr != end(cont: declared_block_names))
4958	{
4959	return itr->second;
4960	}
4961	else
4962	{
4963	auto &var = get<SPIRVariable>(id);
4964
4965	if (fallback_prefer_instance_name)
4966	{
4967	return to_name(id: var.self);
4968	}
4969	else
4970	{
4971	auto &type = get<SPIRType>(id: var.basetype);
4972	auto *type_meta = ir.find_meta(id: type.self);
4973	auto *block_name = type_meta ? &type_meta->decoration.alias : nullptr;
4974	return (!block_name || block_name->empty()) ? get_block_fallback_name(id) : *block_name;
4975	}
4976	}
4977	}
4978
4979	bool Compiler::reflection_ssbo_instance_name_is_significant() const
4980	{
4981	if (ir.source.known)
4982	{
4983	// UAVs from HLSL source tend to be declared in a way where the type is reused
4984	// but the instance name is significant, and that's the name we should report.
4985	// For GLSL, SSBOs each have their own block type as that's how GLSL is written.
4986	return ir.source.hlsl;
4987	}
4988
4989	unordered_set<uint32_t> ssbo_type_ids;
4990	bool aliased_ssbo_types = false;
4991
4992	// If we don't have any OpSource information, we need to perform some shaky heuristics.
4993	ir.for_each_typed_id<SPIRVariable>(op: [&](uint32_t, const SPIRVariable &var) {
4994	auto &type = this->get<SPIRType>(id: var.basetype);
4995	if (!type.pointer || var.storage == StorageClassFunction)
4996	return;
4997
4998	bool ssbo = var.storage == StorageClassStorageBuffer ||
4999	(var.storage == StorageClassUniform && has_decoration(id: type.self, decoration: DecorationBufferBlock));
5000
5001	if (ssbo)
5002	{
5003	if (ssbo_type_ids.count(x: type.self))
5004	aliased_ssbo_types = true;
5005	else
5006	ssbo_type_ids.insert(x: type.self);
5007	}
5008	});
5009
5010	// If the block name is aliased, assume we have HLSL-style UAV declarations.
5011	return aliased_ssbo_types;
5012	}
5013
5014	bool Compiler::instruction_to_result_type(uint32_t &result_type, uint32_t &result_id, spv::Op op,
5015	const uint32_t *args, uint32_t length)
5016	{
5017	if (length < 2)
5018	return false;
5019
5020	bool has_result_id = false, has_result_type = false;
5021	HasResultAndType(opcode: op, hasResult: &has_result_id, hasResultType: &has_result_type);
5022	if (has_result_id && has_result_type)
5023	{
5024	result_type = args[0];
5025	result_id = args[1];
5026	return true;
5027	}
5028	else
5029	return false;
5030	}
5031
5032	Bitset Compiler::combined_decoration_for_member(const SPIRType &type, uint32_t index) const
5033	{
5034	Bitset flags;
5035	auto *type_meta = ir.find_meta(id: type.self);
5036
5037	if (type_meta)
5038	{
5039	auto &members = type_meta->members;
5040	if (index >= members.size())
5041	return flags;
5042	auto &dec = members[index];
5043
5044	flags.merge_or(other: dec.decoration_flags);
5045
5046	auto &member_type = get<SPIRType>(id: type.member_types[index]);
5047
5048	// If our member type is a struct, traverse all the child members as well recursively.
5049	auto &member_childs = member_type.member_types;
5050	for (uint32_t i = 0; i < member_childs.size(); i++)
5051	{
5052	auto &child_member_type = get<SPIRType>(id: member_childs[i]);
5053	if (!child_member_type.pointer)
5054	flags.merge_or(other: combined_decoration_for_member(type: member_type, index: i));
5055	}
5056	}
5057
5058	return flags;
5059	}
5060
5061	bool Compiler::is_desktop_only_format(spv::ImageFormat format)
5062	{
5063	switch (format)
5064	{
5065	// Desktop-only formats
5066	case ImageFormatR11fG11fB10f:
5067	case ImageFormatR16f:
5068	case ImageFormatRgb10A2:
5069	case ImageFormatR8:
5070	case ImageFormatRg8:
5071	case ImageFormatR16:
5072	case ImageFormatRg16:
5073	case ImageFormatRgba16:
5074	case ImageFormatR16Snorm:
5075	case ImageFormatRg16Snorm:
5076	case ImageFormatRgba16Snorm:
5077	case ImageFormatR8Snorm:
5078	case ImageFormatRg8Snorm:
5079	case ImageFormatR8ui:
5080	case ImageFormatRg8ui:
5081	case ImageFormatR16ui:
5082	case ImageFormatRgb10a2ui:
5083	case ImageFormatR8i:
5084	case ImageFormatRg8i:
5085	case ImageFormatR16i:
5086	return true;
5087	default:
5088	break;
5089	}
5090
5091	return false;
5092	}
5093
5094	// An image is determined to be a depth image if it is marked as a depth image and is not also
5095	// explicitly marked with a color format, or if there are any sample/gather compare operations on it.
5096	bool Compiler::is_depth_image(const SPIRType &type, uint32_t id) const
5097	{
5098	return (type.image.depth && type.image.format == ImageFormatUnknown) || comparison_ids.count(x: id);
5099	}
5100
5101	bool Compiler::type_is_opaque_value(const SPIRType &type) const
5102	{
5103	return !type.pointer && (type.basetype == SPIRType::SampledImage || type.basetype == SPIRType::Image ||
5104	type.basetype == SPIRType::Sampler);
5105	}
5106
5107	// Make these member functions so we can easily break on any force_recompile events.
5108	void Compiler::force_recompile()
5109	{
5110	is_force_recompile = true;
5111	}
5112
5113	void Compiler::force_recompile_guarantee_forward_progress()
5114	{
5115	force_recompile();
5116	is_force_recompile_forward_progress = true;
5117	}
5118
5119	bool Compiler::is_forcing_recompilation() const
5120	{
5121	return is_force_recompile;
5122	}
5123
5124	void Compiler::clear_force_recompile()
5125	{
5126	is_force_recompile = false;
5127	is_force_recompile_forward_progress = false;
5128	}
5129
5130	Compiler::PhysicalStorageBufferPointerHandler::PhysicalStorageBufferPointerHandler(Compiler &compiler_)
5131	: compiler(compiler_)
5132	{
5133	}
5134
5135	Compiler::PhysicalBlockMeta *Compiler::PhysicalStorageBufferPointerHandler::find_block_meta(uint32_t id) const
5136	{
5137	auto chain_itr = access_chain_to_physical_block.find(x: id);
5138	if (chain_itr != access_chain_to_physical_block.end())
5139	return chain_itr->second;
5140	else
5141	return nullptr;
5142	}
5143
5144	void Compiler::PhysicalStorageBufferPointerHandler::mark_aligned_access(uint32_t id, const uint32_t *args, uint32_t length)
5145	{
5146	uint32_t mask = *args;
5147	args++;
5148	length--;
5149	if (length && (mask & MemoryAccessVolatileMask) != 0)
5150	{
5151	args++;
5152	length--;
5153	}
5154
5155	if (length && (mask & MemoryAccessAlignedMask) != 0)
5156	{
5157	uint32_t alignment = *args;
5158	auto *meta = find_block_meta(id);
5159
5160	// This makes the assumption that the application does not rely on insane edge cases like:
5161	// Bind buffer with ADDR = 8, use block offset of 8 bytes, load/store with 16 byte alignment.
5162	// If we emit the buffer with alignment = 16 here, the first element at offset = 0 should
5163	// actually have alignment of 8 bytes, but this is too theoretical and awkward to support.
5164	// We could potentially keep track of any offset in the access chain, but it's
5165	// practically impossible for high level compilers to emit code like that,
5166	// so deducing overall alignment requirement based on maximum observed Alignment value is probably fine.
5167	if (meta && alignment > meta->alignment)
5168	meta->alignment = alignment;
5169	}
5170	}
5171
5172	bool Compiler::PhysicalStorageBufferPointerHandler::type_is_bda_block_entry(uint32_t type_id) const
5173	{
5174	auto &type = compiler.get<SPIRType>(id: type_id);
5175	return compiler.is_physical_pointer(type);
5176	}
5177
5178	uint32_t Compiler::PhysicalStorageBufferPointerHandler::get_minimum_scalar_alignment(const SPIRType &type) const
5179	{
5180	if (type.storage == spv::StorageClassPhysicalStorageBufferEXT)
5181	return 8;
5182	else if (type.basetype == SPIRType::Struct)
5183	{
5184	uint32_t alignment = 0;
5185	for (auto &member_type : type.member_types)
5186	{
5187	uint32_t member_align = get_minimum_scalar_alignment(type: compiler.get<SPIRType>(id: member_type));
5188	if (member_align > alignment)
5189	alignment = member_align;
5190	}
5191	return alignment;
5192	}
5193	else
5194	return type.width / 8;
5195	}
5196
5197	void Compiler::PhysicalStorageBufferPointerHandler::setup_meta_chain(uint32_t type_id, uint32_t var_id)
5198	{
5199	if (type_is_bda_block_entry(type_id))
5200	{
5201	auto &meta = physical_block_type_meta[type_id];
5202	access_chain_to_physical_block[var_id] = &meta;
5203
5204	auto &type = compiler.get<SPIRType>(id: type_id);
5205
5206	if (!compiler.is_physical_pointer_to_buffer_block(type))
5207	non_block_types.insert(x: type_id);
5208
5209	if (meta.alignment == 0)
5210	meta.alignment = get_minimum_scalar_alignment(type: compiler.get_pointee_type(type));
5211	}
5212	}
5213
5214	bool Compiler::PhysicalStorageBufferPointerHandler::handle(Op op, const uint32_t *args, uint32_t length)
5215	{
5216	// When a BDA pointer comes to life, we need to keep a mapping of SSA ID -> type ID for the pointer type.
5217	// For every load and store, we'll need to be able to look up the type ID being accessed and mark any alignment
5218	// requirements.
5219	switch (op)
5220	{
5221	case OpConvertUToPtr:
5222	case OpBitcast:
5223	case OpCompositeExtract:
5224	// Extract can begin a new chain if we had a struct or array of pointers as input.
5225	// We don't begin chains before we have a pure scalar pointer.
5226	setup_meta_chain(type_id: args[0], var_id: args[1]);
5227	break;
5228
5229	case OpAccessChain:
5230	case OpInBoundsAccessChain:
5231	case OpPtrAccessChain:
5232	case OpCopyObject:
5233	{
5234	auto itr = access_chain_to_physical_block.find(x: args[2]);
5235	if (itr != access_chain_to_physical_block.end())
5236	access_chain_to_physical_block[args[1]] = itr->second;
5237	break;
5238	}
5239
5240	case OpLoad:
5241	{
5242	setup_meta_chain(type_id: args[0], var_id: args[1]);
5243	if (length >= 4)
5244	mark_aligned_access(id: args[2], args: args + 3, length: length - 3);
5245	break;
5246	}
5247
5248	case OpStore:
5249	{
5250	if (length >= 3)
5251	mark_aligned_access(id: args[0], args: args + 2, length: length - 2);
5252	break;
5253	}
5254
5255	default:
5256	break;
5257	}
5258
5259	return true;
5260	}
5261
5262	uint32_t Compiler::PhysicalStorageBufferPointerHandler::get_base_non_block_type_id(uint32_t type_id) const
5263	{
5264	auto *type = &compiler.get<SPIRType>(id: type_id);
5265	while (compiler.is_physical_pointer(type: *type) && !type_is_bda_block_entry(type_id))
5266	{
5267	type_id = type->parent_type;
5268	type = &compiler.get<SPIRType>(id: type_id);
5269	}
5270
5271	assert(type_is_bda_block_entry(type_id));
5272	return type_id;
5273	}
5274
5275	void Compiler::PhysicalStorageBufferPointerHandler::analyze_non_block_types_from_block(const SPIRType &type)
5276	{
5277	for (auto &member : type.member_types)
5278	{
5279	auto &subtype = compiler.get<SPIRType>(id: member);
5280
5281	if (compiler.is_physical_pointer(type: subtype) && !compiler.is_physical_pointer_to_buffer_block(type: subtype))
5282	non_block_types.insert(x: get_base_non_block_type_id(type_id: member));
5283	else if (subtype.basetype == SPIRType::Struct && !compiler.is_pointer(type: subtype))
5284	analyze_non_block_types_from_block(type: subtype);
5285	}
5286	}
5287
5288	void Compiler::analyze_non_block_pointer_types()
5289	{
5290	PhysicalStorageBufferPointerHandler handler(*this);
5291	traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler);
5292
5293	// Analyze any block declaration we have to make. It might contain
5294	// physical pointers to POD types which we never used, and thus never added to the list.
5295	// We'll need to add those pointer types to the set of types we declare.
5296	ir.for_each_typed_id<SPIRType>(op: [&](uint32_t id, SPIRType &type) {
5297	// Only analyze the raw block struct, not any pointer-to-struct, since that's just redundant.
5298	if (type.self == id &&
5299	(has_decoration(id: type.self, decoration: DecorationBlock) ||
5300	has_decoration(id: type.self, decoration: DecorationBufferBlock)))
5301	{
5302	handler.analyze_non_block_types_from_block(type);
5303	}
5304	});
5305
5306	physical_storage_non_block_pointer_types.reserve(count: handler.non_block_types.size());
5307	for (auto type : handler.non_block_types)
5308	physical_storage_non_block_pointer_types.push_back(t: type);
5309	sort(first: begin(cont&: physical_storage_non_block_pointer_types), last: end(cont&: physical_storage_non_block_pointer_types));
5310	physical_storage_type_to_alignment = std::move(handler.physical_block_type_meta);
5311	}
5312
5313	bool Compiler::InterlockedResourceAccessPrepassHandler::handle(Op op, const uint32_t *, uint32_t)
5314	{
5315	if (op == OpBeginInvocationInterlockEXT || op == OpEndInvocationInterlockEXT)
5316	{
5317	if (interlock_function_id != 0 && interlock_function_id != call_stack.back())
5318	{
5319	// Most complex case, we have no sensible way of dealing with this
5320	// other than taking the 100% conservative approach, exit early.
5321	split_function_case = true;
5322	return false;
5323	}
5324	else
5325	{
5326	interlock_function_id = call_stack.back();
5327	// If this call is performed inside control flow we have a problem.
5328	auto &cfg = compiler.get_cfg_for_function(id: interlock_function_id);
5329
5330	uint32_t from_block_id = compiler.get<SPIRFunction>(id: interlock_function_id).entry_block;
5331	bool outside_control_flow = cfg.node_terminates_control_flow_in_sub_graph(from: from_block_id, to: current_block_id);
5332	if (!outside_control_flow)
5333	control_flow_interlock = true;
5334	}
5335	}
5336	return true;
5337	}
5338
5339	void Compiler::InterlockedResourceAccessPrepassHandler::rearm_current_block(const SPIRBlock &block)
5340	{
5341	current_block_id = block.self;
5342	}
5343
5344	bool Compiler::InterlockedResourceAccessPrepassHandler::begin_function_scope(const uint32_t *args, uint32_t length)
5345	{
5346	if (length < 3)
5347	return false;
5348	call_stack.push_back(t: args[2]);
5349	return true;
5350	}
5351
5352	bool Compiler::InterlockedResourceAccessPrepassHandler::end_function_scope(const uint32_t *, uint32_t)
5353	{
5354	call_stack.pop_back();
5355	return true;
5356	}
5357
5358	bool Compiler::InterlockedResourceAccessHandler::begin_function_scope(const uint32_t *args, uint32_t length)
5359	{
5360	if (length < 3)
5361	return false;
5362
5363	if (args[2] == interlock_function_id)
5364	call_stack_is_interlocked = true;
5365
5366	call_stack.push_back(t: args[2]);
5367	return true;
5368	}
5369
5370	bool Compiler::InterlockedResourceAccessHandler::end_function_scope(const uint32_t *, uint32_t)
5371	{
5372	if (call_stack.back() == interlock_function_id)
5373	call_stack_is_interlocked = false;
5374
5375	call_stack.pop_back();
5376	return true;
5377	}
5378
5379	void Compiler::InterlockedResourceAccessHandler::access_potential_resource(uint32_t id)
5380	{
5381	if ((use_critical_section && in_crit_sec) || (control_flow_interlock && call_stack_is_interlocked) ||
5382	split_function_case)
5383	{
5384	compiler.interlocked_resources.insert(x: id);
5385	}
5386	}
5387
5388	bool Compiler::InterlockedResourceAccessHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
5389	{
5390	// Only care about critical section analysis if we have simple case.
5391	if (use_critical_section)
5392	{
5393	if (opcode == OpBeginInvocationInterlockEXT)
5394	{
5395	in_crit_sec = true;
5396	return true;
5397	}
5398
5399	if (opcode == OpEndInvocationInterlockEXT)
5400	{
5401	// End critical section--nothing more to do.
5402	return false;
5403	}
5404	}
5405
5406	// We need to figure out where images and buffers are loaded from, so do only the bare bones compilation we need.
5407	switch (opcode)
5408	{
5409	case OpLoad:
5410	{
5411	if (length < 3)
5412	return false;
5413
5414	uint32_t ptr = args[2];
5415	auto *var = compiler.maybe_get_backing_variable(chain: ptr);
5416
5417	// We're only concerned with buffer and image memory here.
5418	if (!var)
5419	break;
5420
5421	switch (var->storage)
5422	{
5423	default:
5424	break;
5425
5426	case StorageClassUniformConstant:
5427	{
5428	uint32_t result_type = args[0];
5429	uint32_t id = args[1];
5430	compiler.set<SPIRExpression>(id, args: "", args&: result_type, args: true);
5431	compiler.register_read(expr: id, chain: ptr, forwarded: true);
5432	break;
5433	}
5434
5435	case StorageClassUniform:
5436	// Must have BufferBlock; we only care about SSBOs.
5437	if (!compiler.has_decoration(id: compiler.get<SPIRType>(id: var->basetype).self, decoration: DecorationBufferBlock))
5438	break;
5439	// fallthrough
5440	case StorageClassStorageBuffer:
5441	access_potential_resource(id: var->self);
5442	break;
5443	}
5444	break;
5445	}
5446
5447	case OpInBoundsAccessChain:
5448	case OpAccessChain:
5449	case OpPtrAccessChain:
5450	{
5451	if (length < 3)
5452	return false;
5453
5454	uint32_t result_type = args[0];
5455
5456	auto &type = compiler.get<SPIRType>(id: result_type);
5457	if (type.storage == StorageClassUniform || type.storage == StorageClassUniformConstant ||
5458	type.storage == StorageClassStorageBuffer)
5459	{
5460	uint32_t id = args[1];
5461	uint32_t ptr = args[2];
5462	compiler.set<SPIRExpression>(id, args: "", args&: result_type, args: true);
5463	compiler.register_read(expr: id, chain: ptr, forwarded: true);
5464	compiler.ir.ids[id].set_allow_type_rewrite();
5465	}
5466	break;
5467	}
5468
5469	case OpImageTexelPointer:
5470	{
5471	if (length < 3)
5472	return false;
5473
5474	uint32_t result_type = args[0];
5475	uint32_t id = args[1];
5476	uint32_t ptr = args[2];
5477	auto &e = compiler.set<SPIRExpression>(id, args: "", args&: result_type, args: true);
5478	auto *var = compiler.maybe_get_backing_variable(chain: ptr);
5479	if (var)
5480	e.loaded_from = var->self;
5481	break;
5482	}
5483
5484	case OpStore:
5485	case OpImageWrite:
5486	case OpAtomicStore:
5487	{
5488	if (length < 1)
5489	return false;
5490
5491	uint32_t ptr = args[0];
5492	auto *var = compiler.maybe_get_backing_variable(chain: ptr);
5493	if (var && (var->storage == StorageClassUniform || var->storage == StorageClassUniformConstant ||
5494	var->storage == StorageClassStorageBuffer))
5495	{
5496	access_potential_resource(id: var->self);
5497	}
5498
5499	break;
5500	}
5501
5502	case OpCopyMemory:
5503	{
5504	if (length < 2)
5505	return false;
5506
5507	uint32_t dst = args[0];
5508	uint32_t src = args[1];
5509	auto *dst_var = compiler.maybe_get_backing_variable(chain: dst);
5510	auto *src_var = compiler.maybe_get_backing_variable(chain: src);
5511
5512	if (dst_var && (dst_var->storage == StorageClassUniform || dst_var->storage == StorageClassStorageBuffer))
5513	access_potential_resource(id: dst_var->self);
5514
5515	if (src_var)
5516	{
5517	if (src_var->storage != StorageClassUniform && src_var->storage != StorageClassStorageBuffer)
5518	break;
5519
5520	if (src_var->storage == StorageClassUniform &&
5521	!compiler.has_decoration(id: compiler.get<SPIRType>(id: src_var->basetype).self, decoration: DecorationBufferBlock))
5522	{
5523	break;
5524	}
5525
5526	access_potential_resource(id: src_var->self);
5527	}
5528
5529	break;
5530	}
5531
5532	case OpImageRead:
5533	case OpAtomicLoad:
5534	{
5535	if (length < 3)
5536	return false;
5537
5538	uint32_t ptr = args[2];
5539	auto *var = compiler.maybe_get_backing_variable(chain: ptr);
5540
5541	// We're only concerned with buffer and image memory here.
5542	if (!var)
5543	break;
5544
5545	switch (var->storage)
5546	{
5547	default:
5548	break;
5549
5550	case StorageClassUniform:
5551	// Must have BufferBlock; we only care about SSBOs.
5552	if (!compiler.has_decoration(id: compiler.get<SPIRType>(id: var->basetype).self, decoration: DecorationBufferBlock))
5553	break;
5554	// fallthrough
5555	case StorageClassUniformConstant:
5556	case StorageClassStorageBuffer:
5557	access_potential_resource(id: var->self);
5558	break;
5559	}
5560	break;
5561	}
5562
5563	case OpAtomicExchange:
5564	case OpAtomicCompareExchange:
5565	case OpAtomicIIncrement:
5566	case OpAtomicIDecrement:
5567	case OpAtomicIAdd:
5568	case OpAtomicISub:
5569	case OpAtomicSMin:
5570	case OpAtomicUMin:
5571	case OpAtomicSMax:
5572	case OpAtomicUMax:
5573	case OpAtomicAnd:
5574	case OpAtomicOr:
5575	case OpAtomicXor:
5576	{
5577	if (length < 3)
5578	return false;
5579
5580	uint32_t ptr = args[2];
5581	auto *var = compiler.maybe_get_backing_variable(chain: ptr);
5582	if (var && (var->storage == StorageClassUniform || var->storage == StorageClassUniformConstant ||
5583	var->storage == StorageClassStorageBuffer))
5584	{
5585	access_potential_resource(id: var->self);
5586	}
5587
5588	break;
5589	}
5590
5591	default:
5592	break;
5593	}
5594
5595	return true;
5596	}
5597
5598	void Compiler::analyze_interlocked_resource_usage()
5599	{
5600	if (get_execution_model() == ExecutionModelFragment &&
5601	(get_entry_point().flags.get(bit: ExecutionModePixelInterlockOrderedEXT) ||
5602	get_entry_point().flags.get(bit: ExecutionModePixelInterlockUnorderedEXT) ||
5603	get_entry_point().flags.get(bit: ExecutionModeSampleInterlockOrderedEXT) ||
5604	get_entry_point().flags.get(bit: ExecutionModeSampleInterlockUnorderedEXT)))
5605	{
5606	InterlockedResourceAccessPrepassHandler prepass_handler(*this, ir.default_entry_point);
5607	traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler&: prepass_handler);
5608
5609	InterlockedResourceAccessHandler handler(*this, ir.default_entry_point);
5610	handler.interlock_function_id = prepass_handler.interlock_function_id;
5611	handler.split_function_case = prepass_handler.split_function_case;
5612	handler.control_flow_interlock = prepass_handler.control_flow_interlock;
5613	handler.use_critical_section = !handler.split_function_case && !handler.control_flow_interlock;
5614
5615	traverse_all_reachable_opcodes(func: get<SPIRFunction>(id: ir.default_entry_point), handler);
5616
5617	// For GLSL. If we hit any of these cases, we have to fall back to conservative approach.
5618	interlocked_is_complex =
5619	!handler.use_critical_section || handler.interlock_function_id != ir.default_entry_point;
5620	}
5621	}
5622
5623	// Helper function
5624	bool Compiler::check_internal_recursion(const SPIRType &type, std::unordered_set<uint32_t> &checked_ids)
5625	{
5626	if (type.basetype != SPIRType::Struct)
5627	return false;
5628
5629	if (checked_ids.count(x: type.self))
5630	return true;
5631
5632	// Recurse into struct members
5633	bool is_recursive = false;
5634	checked_ids.insert(x: type.self);
5635	uint32_t mbr_cnt = uint32_t(type.member_types.size());
5636	for (uint32_t mbr_idx = 0; !is_recursive && mbr_idx < mbr_cnt; mbr_idx++)
5637	{
5638	uint32_t mbr_type_id = type.member_types[mbr_idx];
5639	auto &mbr_type = get<SPIRType>(id: mbr_type_id);
5640	is_recursive |= check_internal_recursion(type: mbr_type, checked_ids);
5641	}
5642	checked_ids.erase(x: type.self);
5643	return is_recursive;
5644	}
5645
5646	// Return whether the struct type contains a structural recursion nested somewhere within its content.
5647	bool Compiler::type_contains_recursion(const SPIRType &type)
5648	{
5649	std::unordered_set<uint32_t> checked_ids;
5650	return check_internal_recursion(type, checked_ids);
5651	}
5652
5653	bool Compiler::type_is_array_of_pointers(const SPIRType &type) const
5654	{
5655	if (!is_array(type))
5656	return false;
5657
5658	// BDA types must have parent type hierarchy.
5659	if (!type.parent_type)
5660	return false;
5661
5662	// Punch through all array layers.
5663	auto *parent = &get<SPIRType>(id: type.parent_type);
5664	while (is_array(type: *parent))
5665	parent = &get<SPIRType>(id: parent->parent_type);
5666
5667	return is_pointer(type: *parent);
5668	}
5669
5670	bool Compiler::flush_phi_required(BlockID from, BlockID to) const
5671	{
5672	auto &child = get<SPIRBlock>(id: to);
5673	for (auto &phi : child.phi_variables)
5674	if (phi.parent == from)
5675	return true;
5676	return false;
5677	}
5678
5679	void Compiler::add_loop_level()
5680	{
5681	current_loop_level++;
5682	}
5683