2016-03-02 17:09:16 +00:00
|
|
|
/*
|
2017-01-28 08:00:40 +00:00
|
|
|
* Copyright 2015-2017 ARM Limited
|
2016-03-02 17:09:16 +00:00
|
|
|
*
|
|
|
|
* Licensed under the Apache License, Version 2.0 (the "License");
|
|
|
|
* you may not use this file except in compliance with the License.
|
|
|
|
* You may obtain a copy of the License at
|
|
|
|
*
|
|
|
|
* http://www.apache.org/licenses/LICENSE-2.0
|
|
|
|
*
|
|
|
|
* Unless required by applicable law or agreed to in writing, software
|
|
|
|
* distributed under the License is distributed on an "AS IS" BASIS,
|
|
|
|
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
|
|
* See the License for the specific language governing permissions and
|
|
|
|
* limitations under the License.
|
|
|
|
*/
|
|
|
|
|
2016-11-11 17:04:14 +00:00
|
|
|
#ifndef SPIRV_CROSS_COMMON_HPP
|
|
|
|
#define SPIRV_CROSS_COMMON_HPP
|
2016-03-02 17:09:16 +00:00
|
|
|
|
2017-03-25 15:28:44 +00:00
|
|
|
#include "spirv.hpp"
|
2017-03-25 15:43:26 +00:00
|
|
|
|
2016-12-12 21:33:22 +00:00
|
|
|
#include <cstdio>
|
|
|
|
#include <cstring>
|
2016-09-20 08:17:41 +00:00
|
|
|
#include <functional>
|
2017-01-11 14:57:05 +00:00
|
|
|
#include <locale>
|
2017-03-25 15:43:26 +00:00
|
|
|
#include <memory>
|
2016-03-02 17:09:16 +00:00
|
|
|
#include <sstream>
|
2017-03-25 15:43:26 +00:00
|
|
|
#include <stack>
|
|
|
|
#include <stdexcept>
|
|
|
|
#include <string>
|
|
|
|
#include <unordered_map>
|
2017-03-25 15:25:30 +00:00
|
|
|
#include <unordered_set>
|
2017-03-25 15:43:26 +00:00
|
|
|
#include <utility>
|
2017-03-25 15:28:44 +00:00
|
|
|
#include <vector>
|
2016-03-02 17:09:16 +00:00
|
|
|
|
2016-04-04 07:36:04 +00:00
|
|
|
namespace spirv_cross
|
2016-03-02 17:09:16 +00:00
|
|
|
{
|
2016-12-12 21:33:22 +00:00
|
|
|
|
|
|
|
#ifdef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
|
2016-12-15 19:46:10 +00:00
|
|
|
#ifndef _MSC_VER
|
|
|
|
[[noreturn]]
|
|
|
|
#endif
|
|
|
|
inline void
|
|
|
|
report_and_abort(const std::string &msg)
|
2016-12-12 21:33:22 +00:00
|
|
|
{
|
|
|
|
#ifdef NDEBUG
|
|
|
|
(void)msg;
|
|
|
|
#else
|
2016-12-15 19:46:10 +00:00
|
|
|
fprintf(stderr, "There was a compiler error: %s\n", msg.c_str());
|
2016-12-12 21:33:22 +00:00
|
|
|
#endif
|
2017-09-08 07:33:34 +00:00
|
|
|
fflush(stderr);
|
2016-12-12 21:33:22 +00:00
|
|
|
abort();
|
|
|
|
}
|
|
|
|
|
|
|
|
#define SPIRV_CROSS_THROW(x) report_and_abort(x)
|
|
|
|
#else
|
2016-05-05 07:33:18 +00:00
|
|
|
class CompilerError : public std::runtime_error
|
|
|
|
{
|
|
|
|
public:
|
|
|
|
CompilerError(const std::string &str)
|
|
|
|
: std::runtime_error(str)
|
|
|
|
{
|
|
|
|
}
|
|
|
|
};
|
|
|
|
|
2016-12-12 21:33:22 +00:00
|
|
|
#define SPIRV_CROSS_THROW(x) throw CompilerError(x)
|
|
|
|
#endif
|
|
|
|
|
2017-03-15 09:31:19 +00:00
|
|
|
#if __cplusplus >= 201402l
|
|
|
|
#define SPIRV_CROSS_DEPRECATED(reason) [[deprecated(reason)]]
|
|
|
|
#elif defined(__GNUC__)
|
|
|
|
#define SPIRV_CROSS_DEPRECATED(reason) __attribute__((deprecated))
|
|
|
|
#elif defined(_MSC_VER)
|
|
|
|
#define SPIRV_CROSS_DEPRECATED(reason) __declspec(deprecated(reason))
|
|
|
|
#else
|
|
|
|
#define SPIRV_CROSS_DEPRECATED(reason)
|
|
|
|
#endif
|
|
|
|
|
2016-05-05 07:33:18 +00:00
|
|
|
namespace inner
|
|
|
|
{
|
|
|
|
template <typename T>
|
|
|
|
void join_helper(std::ostringstream &stream, T &&t)
|
|
|
|
{
|
|
|
|
stream << std::forward<T>(t);
|
|
|
|
}
|
|
|
|
|
|
|
|
template <typename T, typename... Ts>
|
|
|
|
void join_helper(std::ostringstream &stream, T &&t, Ts &&... ts)
|
|
|
|
{
|
|
|
|
stream << std::forward<T>(t);
|
|
|
|
join_helper(stream, std::forward<Ts>(ts)...);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
// Helper template to avoid lots of nasty string temporary munging.
|
|
|
|
template <typename... Ts>
|
|
|
|
std::string join(Ts &&... ts)
|
|
|
|
{
|
|
|
|
std::ostringstream stream;
|
|
|
|
inner::join_helper(stream, std::forward<Ts>(ts)...);
|
|
|
|
return stream.str();
|
|
|
|
}
|
|
|
|
|
|
|
|
inline std::string merge(const std::vector<std::string> &list)
|
|
|
|
{
|
|
|
|
std::string s;
|
|
|
|
for (auto &elem : list)
|
|
|
|
{
|
|
|
|
s += elem;
|
|
|
|
if (&elem != &list.back())
|
|
|
|
s += ", ";
|
|
|
|
}
|
|
|
|
return s;
|
|
|
|
}
|
|
|
|
|
|
|
|
template <typename T>
|
|
|
|
inline std::string convert_to_string(T &&t)
|
|
|
|
{
|
|
|
|
return std::to_string(std::forward<T>(t));
|
|
|
|
}
|
|
|
|
|
|
|
|
// Allow implementations to set a convenient standard precision
|
2016-04-08 19:12:40 +00:00
|
|
|
#ifndef SPIRV_CROSS_FLT_FMT
|
2016-05-05 07:33:18 +00:00
|
|
|
#define SPIRV_CROSS_FLT_FMT "%.32g"
|
2016-04-06 21:42:27 +00:00
|
|
|
#endif
|
|
|
|
|
2017-01-12 09:57:44 +00:00
|
|
|
#ifdef _MSC_VER
|
|
|
|
#pragma warning(push)
|
2017-01-13 15:32:54 +00:00
|
|
|
#pragma warning(disable : 4996)
|
2017-01-12 09:57:44 +00:00
|
|
|
#endif
|
|
|
|
|
2016-05-05 07:33:18 +00:00
|
|
|
inline std::string convert_to_string(float t)
|
|
|
|
{
|
|
|
|
// std::to_string for floating point values is broken.
|
|
|
|
// Fallback to something more sane.
|
|
|
|
char buf[64];
|
|
|
|
sprintf(buf, SPIRV_CROSS_FLT_FMT, t);
|
|
|
|
// Ensure that the literal is float.
|
|
|
|
if (!strchr(buf, '.') && !strchr(buf, 'e'))
|
|
|
|
strcat(buf, ".0");
|
|
|
|
return buf;
|
2016-03-02 17:09:16 +00:00
|
|
|
}
|
|
|
|
|
2016-05-05 07:33:18 +00:00
|
|
|
inline std::string convert_to_string(double t)
|
|
|
|
{
|
|
|
|
// std::to_string for floating point values is broken.
|
|
|
|
// Fallback to something more sane.
|
|
|
|
char buf[64];
|
|
|
|
sprintf(buf, SPIRV_CROSS_FLT_FMT, t);
|
|
|
|
// Ensure that the literal is float.
|
|
|
|
if (!strchr(buf, '.') && !strchr(buf, 'e'))
|
|
|
|
strcat(buf, ".0");
|
|
|
|
return buf;
|
|
|
|
}
|
|
|
|
|
2017-01-12 09:57:44 +00:00
|
|
|
#ifdef _MSC_VER
|
|
|
|
#pragma warning(pop)
|
|
|
|
#endif
|
|
|
|
|
2016-05-05 07:33:18 +00:00
|
|
|
struct Instruction
|
|
|
|
{
|
|
|
|
Instruction(const std::vector<uint32_t> &spirv, uint32_t &index);
|
|
|
|
|
|
|
|
uint16_t op;
|
|
|
|
uint16_t count;
|
|
|
|
uint32_t offset;
|
|
|
|
uint32_t length;
|
|
|
|
};
|
|
|
|
|
|
|
|
// Helper for Variant interface.
|
|
|
|
struct IVariant
|
|
|
|
{
|
|
|
|
virtual ~IVariant() = default;
|
|
|
|
uint32_t self = 0;
|
|
|
|
};
|
|
|
|
|
|
|
|
enum Types
|
|
|
|
{
|
|
|
|
TypeNone,
|
|
|
|
TypeType,
|
|
|
|
TypeVariable,
|
|
|
|
TypeConstant,
|
|
|
|
TypeFunction,
|
|
|
|
TypeFunctionPrototype,
|
|
|
|
TypePointer,
|
|
|
|
TypeBlock,
|
|
|
|
TypeExtension,
|
|
|
|
TypeExpression,
|
2016-10-03 13:54:02 +00:00
|
|
|
TypeConstantOp,
|
2017-04-25 08:44:55 +00:00
|
|
|
TypeCombinedImageSampler,
|
2017-08-10 13:36:30 +00:00
|
|
|
TypeAccessChain,
|
2016-05-05 07:33:18 +00:00
|
|
|
TypeUndef
|
|
|
|
};
|
|
|
|
|
|
|
|
struct SPIRUndef : IVariant
|
|
|
|
{
|
|
|
|
enum
|
|
|
|
{
|
|
|
|
type = TypeUndef
|
|
|
|
};
|
|
|
|
SPIRUndef(uint32_t basetype_)
|
|
|
|
: basetype(basetype_)
|
|
|
|
{
|
|
|
|
}
|
|
|
|
uint32_t basetype;
|
|
|
|
};
|
|
|
|
|
2017-04-25 08:44:55 +00:00
|
|
|
// This type is only used by backends which need to access the combined image and sampler IDs separately after
|
|
|
|
// the OpSampledImage opcode.
|
|
|
|
struct SPIRCombinedImageSampler : IVariant
|
|
|
|
{
|
|
|
|
enum
|
|
|
|
{
|
|
|
|
type = TypeCombinedImageSampler
|
|
|
|
};
|
|
|
|
SPIRCombinedImageSampler(uint32_t type_, uint32_t image_, uint32_t sampler_)
|
|
|
|
: combined_type(type_)
|
|
|
|
, image(image_)
|
|
|
|
, sampler(sampler_)
|
|
|
|
{
|
|
|
|
}
|
|
|
|
uint32_t combined_type;
|
|
|
|
uint32_t image;
|
|
|
|
uint32_t sampler;
|
|
|
|
};
|
|
|
|
|
2016-10-03 13:54:02 +00:00
|
|
|
struct SPIRConstantOp : IVariant
|
|
|
|
{
|
|
|
|
enum
|
|
|
|
{
|
|
|
|
type = TypeConstantOp
|
|
|
|
};
|
|
|
|
|
|
|
|
SPIRConstantOp(uint32_t result_type, spv::Op op, const uint32_t *args, uint32_t length)
|
|
|
|
: opcode(op)
|
|
|
|
, arguments(args, args + length)
|
|
|
|
, basetype(result_type)
|
|
|
|
{
|
|
|
|
}
|
|
|
|
|
|
|
|
spv::Op opcode;
|
|
|
|
std::vector<uint32_t> arguments;
|
|
|
|
uint32_t basetype;
|
|
|
|
};
|
|
|
|
|
2016-05-05 07:33:18 +00:00
|
|
|
struct SPIRType : IVariant
|
|
|
|
{
|
|
|
|
enum
|
|
|
|
{
|
|
|
|
type = TypeType
|
|
|
|
};
|
|
|
|
|
|
|
|
enum BaseType
|
|
|
|
{
|
|
|
|
Unknown,
|
|
|
|
Void,
|
2016-06-05 18:13:45 +00:00
|
|
|
Boolean,
|
2016-05-05 07:33:18 +00:00
|
|
|
Char,
|
|
|
|
Int,
|
|
|
|
UInt,
|
2016-07-27 09:27:00 +00:00
|
|
|
Int64,
|
|
|
|
UInt64,
|
2016-05-05 07:33:18 +00:00
|
|
|
AtomicCounter,
|
|
|
|
Float,
|
2016-07-27 08:59:00 +00:00
|
|
|
Double,
|
2016-05-05 07:33:18 +00:00
|
|
|
Struct,
|
|
|
|
Image,
|
|
|
|
SampledImage,
|
|
|
|
Sampler
|
|
|
|
};
|
|
|
|
|
|
|
|
// Scalar/vector/matrix support.
|
|
|
|
BaseType basetype = Unknown;
|
|
|
|
uint32_t width = 0;
|
|
|
|
uint32_t vecsize = 1;
|
|
|
|
uint32_t columns = 1;
|
|
|
|
|
2016-10-03 15:17:11 +00:00
|
|
|
// Arrays, support array of arrays by having a vector of array sizes.
|
2016-05-05 07:33:18 +00:00
|
|
|
std::vector<uint32_t> array;
|
|
|
|
|
2016-10-03 15:17:11 +00:00
|
|
|
// Array elements can be either specialization constants or specialization ops.
|
|
|
|
// This array determines how to interpret the array size.
|
|
|
|
// If an element is true, the element is a literal,
|
|
|
|
// otherwise, it's an expression, which must be resolved on demand.
|
|
|
|
// The actual size is not really known until runtime.
|
|
|
|
std::vector<bool> array_size_literal;
|
|
|
|
|
2016-05-05 07:33:18 +00:00
|
|
|
// Pointers
|
|
|
|
bool pointer = false;
|
|
|
|
spv::StorageClass storage = spv::StorageClassGeneric;
|
|
|
|
|
|
|
|
std::vector<uint32_t> member_types;
|
|
|
|
|
|
|
|
struct Image
|
|
|
|
{
|
|
|
|
uint32_t type;
|
|
|
|
spv::Dim dim;
|
|
|
|
bool depth;
|
|
|
|
bool arrayed;
|
|
|
|
bool ms;
|
|
|
|
uint32_t sampled;
|
|
|
|
spv::ImageFormat format;
|
2017-05-30 00:45:05 +00:00
|
|
|
spv::AccessQualifier access;
|
2016-05-05 07:33:18 +00:00
|
|
|
} image;
|
2016-05-23 08:57:22 +00:00
|
|
|
|
|
|
|
// Structs can be declared multiple times if they are used as part of interface blocks.
|
|
|
|
// We want to detect this so that we only emit the struct definition once.
|
|
|
|
// Since we cannot rely on OpName to be equal, we need to figure out aliases.
|
|
|
|
uint32_t type_alias = 0;
|
2016-05-23 10:25:09 +00:00
|
|
|
|
2017-01-21 10:30:33 +00:00
|
|
|
// Denotes the type which this type is based on.
|
|
|
|
// Allows the backend to traverse how a complex type is built up during access chains.
|
|
|
|
uint32_t parent_type = 0;
|
|
|
|
|
2016-05-23 10:25:09 +00:00
|
|
|
// Used in backends to avoid emitting members with conflicting names.
|
|
|
|
std::unordered_set<std::string> member_name_cache;
|
2016-05-05 07:33:18 +00:00
|
|
|
};
|
|
|
|
|
|
|
|
struct SPIRExtension : IVariant
|
|
|
|
{
|
|
|
|
enum
|
|
|
|
{
|
|
|
|
type = TypeExtension
|
|
|
|
};
|
|
|
|
|
|
|
|
enum Extension
|
|
|
|
{
|
2017-03-21 15:33:54 +00:00
|
|
|
Unsupported,
|
2016-05-05 07:33:18 +00:00
|
|
|
GLSL
|
|
|
|
};
|
2016-03-02 17:09:16 +00:00
|
|
|
|
2016-05-05 07:33:18 +00:00
|
|
|
SPIRExtension(Extension ext_)
|
|
|
|
: ext(ext_)
|
|
|
|
{
|
|
|
|
}
|
|
|
|
|
|
|
|
Extension ext;
|
|
|
|
};
|
|
|
|
|
2016-07-28 09:16:02 +00:00
|
|
|
// SPIREntryPoint is not a variant since its IDs are used to decorate OpFunction,
|
|
|
|
// so in order to avoid conflicts, we can't stick them in the ids array.
|
|
|
|
struct SPIREntryPoint
|
|
|
|
{
|
|
|
|
SPIREntryPoint(uint32_t self_, spv::ExecutionModel execution_model, std::string entry_name)
|
|
|
|
: self(self_)
|
|
|
|
, name(std::move(entry_name))
|
|
|
|
, model(execution_model)
|
|
|
|
{
|
|
|
|
}
|
|
|
|
SPIREntryPoint() = default;
|
|
|
|
|
|
|
|
uint32_t self = 0;
|
|
|
|
std::string name;
|
|
|
|
std::vector<uint32_t> interface_variables;
|
|
|
|
|
|
|
|
uint64_t flags = 0;
|
|
|
|
struct
|
|
|
|
{
|
|
|
|
uint32_t x = 0, y = 0, z = 0;
|
|
|
|
} workgroup_size;
|
|
|
|
uint32_t invocations = 0;
|
|
|
|
uint32_t output_vertices = 0;
|
2016-09-19 11:17:04 +00:00
|
|
|
spv::ExecutionModel model;
|
2016-07-28 09:16:02 +00:00
|
|
|
};
|
|
|
|
|
2016-05-05 07:33:18 +00:00
|
|
|
struct SPIRExpression : IVariant
|
|
|
|
{
|
|
|
|
enum
|
|
|
|
{
|
|
|
|
type = TypeExpression
|
|
|
|
};
|
|
|
|
|
|
|
|
// Only created by the backend target to avoid creating tons of temporaries.
|
|
|
|
SPIRExpression(std::string expr, uint32_t expression_type_, bool immutable_)
|
|
|
|
: expression(move(expr))
|
|
|
|
, expression_type(expression_type_)
|
|
|
|
, immutable(immutable_)
|
|
|
|
{
|
|
|
|
}
|
|
|
|
|
|
|
|
// If non-zero, prepend expression with to_expression(base_expression).
|
|
|
|
// Used in amortizing multiple calls to to_expression()
|
|
|
|
// where in certain cases that would quickly force a temporary when not needed.
|
|
|
|
uint32_t base_expression = 0;
|
|
|
|
|
|
|
|
std::string expression;
|
|
|
|
uint32_t expression_type = 0;
|
|
|
|
|
|
|
|
// If this expression is a forwarded load,
|
|
|
|
// allow us to reference the original variable.
|
|
|
|
uint32_t loaded_from = 0;
|
|
|
|
|
|
|
|
// If this expression will never change, we can avoid lots of temporaries
|
|
|
|
// in high level source.
|
2016-07-12 12:33:04 +00:00
|
|
|
// An expression being immutable can be speculative,
|
|
|
|
// it is assumed that this is true almost always.
|
2016-05-05 07:33:18 +00:00
|
|
|
bool immutable = false;
|
|
|
|
|
2017-01-13 15:31:13 +00:00
|
|
|
// Before use, this expression must be transposed.
|
|
|
|
// This is needed for targets which don't support row_major layouts.
|
|
|
|
bool need_transpose = false;
|
|
|
|
|
2016-07-12 12:33:04 +00:00
|
|
|
// A list of expressions which this expression depends on.
|
|
|
|
std::vector<uint32_t> expression_dependencies;
|
2016-05-05 07:33:18 +00:00
|
|
|
};
|
|
|
|
|
|
|
|
struct SPIRFunctionPrototype : IVariant
|
|
|
|
{
|
|
|
|
enum
|
|
|
|
{
|
|
|
|
type = TypeFunctionPrototype
|
|
|
|
};
|
|
|
|
|
|
|
|
SPIRFunctionPrototype(uint32_t return_type_)
|
|
|
|
: return_type(return_type_)
|
|
|
|
{
|
|
|
|
}
|
|
|
|
|
|
|
|
uint32_t return_type;
|
|
|
|
std::vector<uint32_t> parameter_types;
|
|
|
|
};
|
|
|
|
|
|
|
|
struct SPIRBlock : IVariant
|
|
|
|
{
|
|
|
|
enum
|
|
|
|
{
|
|
|
|
type = TypeBlock
|
|
|
|
};
|
|
|
|
|
|
|
|
enum Terminator
|
|
|
|
{
|
|
|
|
Unknown,
|
|
|
|
Direct, // Emit next block directly without a particular condition.
|
|
|
|
|
|
|
|
Select, // Block ends with an if/else block.
|
|
|
|
MultiSelect, // Block ends with switch statement.
|
|
|
|
|
|
|
|
Return, // Block ends with return.
|
|
|
|
Unreachable, // Noop
|
|
|
|
Kill // Discard
|
|
|
|
};
|
|
|
|
|
|
|
|
enum Merge
|
|
|
|
{
|
|
|
|
MergeNone,
|
|
|
|
MergeLoop,
|
|
|
|
MergeSelection
|
|
|
|
};
|
|
|
|
|
|
|
|
enum Method
|
|
|
|
{
|
|
|
|
MergeToSelectForLoop,
|
|
|
|
MergeToDirectForLoop
|
|
|
|
};
|
|
|
|
|
|
|
|
enum ContinueBlockType
|
|
|
|
{
|
|
|
|
ContinueNone,
|
|
|
|
|
|
|
|
// Continue block is branchless and has at least one instruction.
|
|
|
|
ForLoop,
|
|
|
|
|
|
|
|
// Noop continue block.
|
|
|
|
WhileLoop,
|
|
|
|
|
|
|
|
// Continue block is conditional.
|
|
|
|
DoWhileLoop,
|
|
|
|
|
|
|
|
// Highly unlikely that anything will use this,
|
|
|
|
// since it is really awkward/impossible to express in GLSL.
|
|
|
|
ComplexLoop
|
|
|
|
};
|
|
|
|
|
|
|
|
enum
|
|
|
|
{
|
|
|
|
NoDominator = 0xffffffffu
|
|
|
|
};
|
|
|
|
|
|
|
|
Terminator terminator = Unknown;
|
|
|
|
Merge merge = MergeNone;
|
|
|
|
uint32_t next_block = 0;
|
|
|
|
uint32_t merge_block = 0;
|
|
|
|
uint32_t continue_block = 0;
|
|
|
|
|
|
|
|
uint32_t return_value = 0; // If 0, return nothing (void).
|
|
|
|
uint32_t condition = 0;
|
|
|
|
uint32_t true_block = 0;
|
|
|
|
uint32_t false_block = 0;
|
|
|
|
uint32_t default_block = 0;
|
|
|
|
|
|
|
|
std::vector<Instruction> ops;
|
|
|
|
|
|
|
|
struct Phi
|
|
|
|
{
|
|
|
|
uint32_t local_variable; // flush local variable ...
|
|
|
|
uint32_t parent; // If we're in from_block and want to branch into this block ...
|
|
|
|
uint32_t function_variable; // to this function-global "phi" variable first.
|
|
|
|
};
|
|
|
|
|
|
|
|
// Before entering this block flush out local variables to magical "phi" variables.
|
|
|
|
std::vector<Phi> phi_variables;
|
|
|
|
|
|
|
|
// Declare these temporaries before beginning the block.
|
|
|
|
// Used for handling complex continue blocks which have side effects.
|
|
|
|
std::vector<std::pair<uint32_t, uint32_t>> declare_temporary;
|
|
|
|
|
|
|
|
struct Case
|
|
|
|
{
|
|
|
|
uint32_t value;
|
|
|
|
uint32_t block;
|
|
|
|
};
|
|
|
|
std::vector<Case> cases;
|
|
|
|
|
|
|
|
// If we have tried to optimize code for this block but failed,
|
|
|
|
// keep track of this.
|
|
|
|
bool disable_block_optimization = false;
|
|
|
|
|
|
|
|
// If the continue block is complex, fallback to "dumb" for loops.
|
|
|
|
bool complex_continue = false;
|
|
|
|
|
|
|
|
// The dominating block which this block might be within.
|
|
|
|
// Used in continue; blocks to determine if we really need to write continue.
|
|
|
|
uint32_t loop_dominator = 0;
|
2016-11-11 17:04:14 +00:00
|
|
|
|
|
|
|
// All access to these variables are dominated by this block,
|
|
|
|
// so before branching anywhere we need to make sure that we declare these variables.
|
|
|
|
std::vector<uint32_t> dominated_variables;
|
2016-12-15 16:14:47 +00:00
|
|
|
|
|
|
|
// These are variables which should be declared in a for loop header, if we
|
|
|
|
// fail to use a classic for-loop,
|
|
|
|
// we remove these variables, and fall back to regular variables outside the loop.
|
|
|
|
std::vector<uint32_t> loop_variables;
|
2016-05-05 07:33:18 +00:00
|
|
|
};
|
|
|
|
|
|
|
|
struct SPIRFunction : IVariant
|
|
|
|
{
|
|
|
|
enum
|
|
|
|
{
|
|
|
|
type = TypeFunction
|
|
|
|
};
|
|
|
|
|
|
|
|
SPIRFunction(uint32_t return_type_, uint32_t function_type_)
|
|
|
|
: return_type(return_type_)
|
|
|
|
, function_type(function_type_)
|
|
|
|
{
|
|
|
|
}
|
|
|
|
|
|
|
|
struct Parameter
|
|
|
|
{
|
|
|
|
uint32_t type;
|
|
|
|
uint32_t id;
|
|
|
|
uint32_t read_count;
|
|
|
|
uint32_t write_count;
|
2017-02-05 09:50:14 +00:00
|
|
|
|
|
|
|
// Set to true if this parameter aliases a global variable,
|
|
|
|
// used mostly in Metal where global variables
|
|
|
|
// have to be passed down to functions as regular arguments.
|
|
|
|
// However, for this kind of variable, we should not care about
|
|
|
|
// read and write counts as access to the function arguments
|
|
|
|
// is not local to the function in question.
|
|
|
|
bool alias_global_variable;
|
2016-05-05 07:33:18 +00:00
|
|
|
};
|
|
|
|
|
2016-09-11 09:39:20 +00:00
|
|
|
// When calling a function, and we're remapping separate image samplers,
|
|
|
|
// resolve these arguments into combined image samplers and pass them
|
|
|
|
// as additional arguments in this order.
|
|
|
|
// It gets more complicated as functions can pull in their own globals
|
|
|
|
// and combine them with parameters,
|
|
|
|
// so we need to distinguish if something is local parameter index
|
|
|
|
// or a global ID.
|
|
|
|
struct CombinedImageSamplerParameter
|
|
|
|
{
|
2016-09-11 10:36:12 +00:00
|
|
|
uint32_t id;
|
2016-09-11 11:47:06 +00:00
|
|
|
uint32_t image_id;
|
2016-09-11 09:39:20 +00:00
|
|
|
uint32_t sampler_id;
|
2016-09-11 11:47:06 +00:00
|
|
|
bool global_image;
|
2016-09-11 09:39:20 +00:00
|
|
|
bool global_sampler;
|
2017-05-06 11:21:35 +00:00
|
|
|
bool depth;
|
2016-09-11 09:39:20 +00:00
|
|
|
};
|
|
|
|
|
2016-05-05 07:33:18 +00:00
|
|
|
uint32_t return_type;
|
|
|
|
uint32_t function_type;
|
|
|
|
std::vector<Parameter> arguments;
|
2016-09-11 10:36:12 +00:00
|
|
|
|
|
|
|
// Can be used by backends to add magic arguments.
|
|
|
|
// Currently used by combined image/sampler implementation.
|
|
|
|
|
|
|
|
std::vector<Parameter> shadow_arguments;
|
2016-05-05 07:33:18 +00:00
|
|
|
std::vector<uint32_t> local_variables;
|
|
|
|
uint32_t entry_block = 0;
|
|
|
|
std::vector<uint32_t> blocks;
|
2016-09-11 09:39:20 +00:00
|
|
|
std::vector<CombinedImageSamplerParameter> combined_parameters;
|
2016-05-05 07:33:18 +00:00
|
|
|
|
|
|
|
void add_local_variable(uint32_t id)
|
|
|
|
{
|
|
|
|
local_variables.push_back(id);
|
|
|
|
}
|
|
|
|
|
2017-02-05 09:50:14 +00:00
|
|
|
void add_parameter(uint32_t parameter_type, uint32_t id, bool alias_global_variable = false)
|
2016-05-05 07:33:18 +00:00
|
|
|
{
|
|
|
|
// Arguments are read-only until proven otherwise.
|
2017-02-05 09:50:14 +00:00
|
|
|
arguments.push_back({ parameter_type, id, 0u, 0u, alias_global_variable });
|
2016-05-05 07:33:18 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
bool active = false;
|
|
|
|
bool flush_undeclared = true;
|
2016-09-11 09:39:20 +00:00
|
|
|
bool do_combined_parameters = true;
|
2016-11-18 16:06:49 +00:00
|
|
|
bool analyzed_variable_scope = false;
|
2016-05-05 07:33:18 +00:00
|
|
|
};
|
|
|
|
|
2017-08-10 13:36:30 +00:00
|
|
|
struct SPIRAccessChain : IVariant
|
|
|
|
{
|
|
|
|
enum
|
|
|
|
{
|
|
|
|
type = TypeAccessChain
|
|
|
|
};
|
|
|
|
|
2017-08-28 07:01:03 +00:00
|
|
|
SPIRAccessChain(uint32_t basetype_, spv::StorageClass storage_, std::string base_, std::string dynamic_index_,
|
|
|
|
int32_t static_index_)
|
|
|
|
: basetype(basetype_)
|
|
|
|
, storage(storage_)
|
|
|
|
, base(base_)
|
|
|
|
, dynamic_index(std::move(dynamic_index_))
|
|
|
|
, static_index(static_index_)
|
2017-08-10 13:36:30 +00:00
|
|
|
{
|
|
|
|
}
|
|
|
|
|
|
|
|
// The access chain represents an offset into a buffer.
|
|
|
|
// Some backends need more complicated handling of access chains to be able to use buffers, like HLSL
|
|
|
|
// which has no usable buffer type ala GLSL SSBOs.
|
|
|
|
// StructuredBuffer is too limited, so our only option is to deal with ByteAddressBuffer which works with raw addresses.
|
|
|
|
|
|
|
|
uint32_t basetype;
|
|
|
|
spv::StorageClass storage;
|
|
|
|
std::string base;
|
|
|
|
std::string dynamic_index;
|
|
|
|
int32_t static_index;
|
|
|
|
|
|
|
|
uint32_t loaded_from = 0;
|
|
|
|
bool need_transpose = false;
|
2017-08-10 15:12:48 +00:00
|
|
|
bool immutable = false;
|
2017-08-10 13:36:30 +00:00
|
|
|
};
|
|
|
|
|
2016-05-05 07:33:18 +00:00
|
|
|
struct SPIRVariable : IVariant
|
|
|
|
{
|
|
|
|
enum
|
|
|
|
{
|
|
|
|
type = TypeVariable
|
|
|
|
};
|
|
|
|
|
|
|
|
SPIRVariable() = default;
|
|
|
|
SPIRVariable(uint32_t basetype_, spv::StorageClass storage_, uint32_t initializer_ = 0)
|
|
|
|
: basetype(basetype_)
|
|
|
|
, storage(storage_)
|
|
|
|
, initializer(initializer_)
|
|
|
|
{
|
|
|
|
}
|
|
|
|
|
|
|
|
uint32_t basetype = 0;
|
|
|
|
spv::StorageClass storage = spv::StorageClassGeneric;
|
|
|
|
uint32_t decoration = 0;
|
|
|
|
uint32_t initializer = 0;
|
|
|
|
|
|
|
|
std::vector<uint32_t> dereference_chain;
|
|
|
|
bool compat_builtin = false;
|
|
|
|
|
|
|
|
// If a variable is shadowed, we only statically assign to it
|
|
|
|
// and never actually emit a statement for it.
|
|
|
|
// When we read the variable as an expression, just forward
|
|
|
|
// shadowed_id as the expression.
|
|
|
|
bool statically_assigned = false;
|
|
|
|
uint32_t static_expression = 0;
|
|
|
|
|
|
|
|
// Temporaries which can remain forwarded as long as this variable is not modified.
|
|
|
|
std::vector<uint32_t> dependees;
|
|
|
|
bool forwardable = true;
|
|
|
|
|
|
|
|
bool deferred_declaration = false;
|
|
|
|
bool phi_variable = false;
|
|
|
|
bool remapped_variable = false;
|
2016-07-06 09:04:06 +00:00
|
|
|
uint32_t remapped_components = 0;
|
2016-05-05 07:33:18 +00:00
|
|
|
|
2016-12-15 16:14:47 +00:00
|
|
|
// The block which dominates all access to this variable.
|
|
|
|
uint32_t dominator = 0;
|
|
|
|
// If true, this variable is a loop variable, when accessing the variable
|
|
|
|
// outside a loop,
|
|
|
|
// we should statically forward it.
|
|
|
|
bool loop_variable = false;
|
|
|
|
// Set to true while we're inside the for loop.
|
|
|
|
bool loop_variable_enable = false;
|
|
|
|
|
2016-05-05 07:33:18 +00:00
|
|
|
SPIRFunction::Parameter *parameter = nullptr;
|
|
|
|
};
|
|
|
|
|
|
|
|
struct SPIRConstant : IVariant
|
|
|
|
{
|
|
|
|
enum
|
|
|
|
{
|
|
|
|
type = TypeConstant
|
|
|
|
};
|
|
|
|
|
2016-05-23 11:30:02 +00:00
|
|
|
union Constant {
|
2016-05-05 07:33:18 +00:00
|
|
|
uint32_t u32;
|
|
|
|
int32_t i32;
|
|
|
|
float f32;
|
2016-07-27 08:59:00 +00:00
|
|
|
|
|
|
|
uint64_t u64;
|
|
|
|
int64_t i64;
|
|
|
|
double f64;
|
2016-05-05 07:33:18 +00:00
|
|
|
};
|
|
|
|
|
|
|
|
struct ConstantVector
|
|
|
|
{
|
|
|
|
Constant r[4];
|
2017-09-27 13:16:33 +00:00
|
|
|
uint32_t id[4] = {}; // If != 0, this is a specialization constant, and we should keep track of it as such.
|
|
|
|
uint32_t vecsize = 1;
|
2016-05-05 07:33:18 +00:00
|
|
|
};
|
|
|
|
|
|
|
|
struct ConstantMatrix
|
|
|
|
{
|
|
|
|
ConstantVector c[4];
|
2017-09-27 13:16:33 +00:00
|
|
|
uint32_t columns = 1;
|
2016-05-05 07:33:18 +00:00
|
|
|
};
|
|
|
|
|
|
|
|
inline uint32_t scalar(uint32_t col = 0, uint32_t row = 0) const
|
|
|
|
{
|
|
|
|
return m.c[col].r[row].u32;
|
|
|
|
}
|
|
|
|
|
|
|
|
inline float scalar_f32(uint32_t col = 0, uint32_t row = 0) const
|
|
|
|
{
|
|
|
|
return m.c[col].r[row].f32;
|
|
|
|
}
|
|
|
|
|
2016-07-27 08:59:00 +00:00
|
|
|
inline int32_t scalar_i32(uint32_t col = 0, uint32_t row = 0) const
|
2016-05-05 07:33:18 +00:00
|
|
|
{
|
|
|
|
return m.c[col].r[row].i32;
|
|
|
|
}
|
|
|
|
|
2016-07-27 08:59:00 +00:00
|
|
|
inline double scalar_f64(uint32_t col = 0, uint32_t row = 0) const
|
|
|
|
{
|
|
|
|
return m.c[col].r[row].f64;
|
|
|
|
}
|
|
|
|
|
|
|
|
inline int64_t scalar_i64(uint32_t col = 0, uint32_t row = 0) const
|
|
|
|
{
|
|
|
|
return m.c[col].r[row].i64;
|
|
|
|
}
|
|
|
|
|
|
|
|
inline uint64_t scalar_u64(uint32_t col = 0, uint32_t row = 0) const
|
|
|
|
{
|
|
|
|
return m.c[col].r[row].u64;
|
|
|
|
}
|
|
|
|
|
2016-05-05 07:33:18 +00:00
|
|
|
inline const ConstantVector &vector() const
|
|
|
|
{
|
|
|
|
return m.c[0];
|
|
|
|
}
|
|
|
|
inline uint32_t vector_size() const
|
|
|
|
{
|
|
|
|
return m.c[0].vecsize;
|
|
|
|
}
|
|
|
|
inline uint32_t columns() const
|
|
|
|
{
|
|
|
|
return m.columns;
|
|
|
|
}
|
|
|
|
|
2017-08-03 12:32:07 +00:00
|
|
|
inline void make_null(const SPIRType &constant_type_)
|
|
|
|
{
|
|
|
|
std::memset(&m, 0, sizeof(m));
|
|
|
|
m.columns = constant_type_.columns;
|
|
|
|
for (auto &c : m.c)
|
|
|
|
c.vecsize = constant_type_.vecsize;
|
|
|
|
}
|
|
|
|
|
2017-08-02 08:33:03 +00:00
|
|
|
SPIRConstant(uint32_t constant_type_)
|
|
|
|
: constant_type(constant_type_)
|
|
|
|
{
|
|
|
|
}
|
|
|
|
|
2017-09-27 13:16:33 +00:00
|
|
|
SPIRConstant(uint32_t constant_type_, const uint32_t *elements, uint32_t num_elements, bool specialized)
|
|
|
|
: constant_type(constant_type_), specialization(specialized)
|
2016-05-05 07:33:18 +00:00
|
|
|
{
|
|
|
|
subconstants.insert(end(subconstants), elements, elements + num_elements);
|
2017-09-27 13:16:33 +00:00
|
|
|
specialization = specialized;
|
2016-05-05 07:33:18 +00:00
|
|
|
}
|
|
|
|
|
2017-09-27 13:16:33 +00:00
|
|
|
// Construct scalar (32-bit).
|
|
|
|
SPIRConstant(uint32_t constant_type_, uint32_t v0, bool specialized)
|
|
|
|
: constant_type(constant_type_), specialization(specialized)
|
2016-05-05 07:33:18 +00:00
|
|
|
{
|
|
|
|
m.c[0].r[0].u32 = v0;
|
|
|
|
m.c[0].vecsize = 1;
|
|
|
|
m.columns = 1;
|
|
|
|
|
2017-09-27 13:16:33 +00:00
|
|
|
if (specialized)
|
|
|
|
m.c[0].id[0] = self;
|
2016-05-05 07:33:18 +00:00
|
|
|
}
|
|
|
|
|
2017-09-27 13:16:33 +00:00
|
|
|
// Construct scalar (64-bit).
|
|
|
|
SPIRConstant(uint32_t constant_type_, uint64_t v0, bool specialized)
|
|
|
|
: constant_type(constant_type_), specialization(specialized)
|
2016-07-27 08:59:00 +00:00
|
|
|
{
|
|
|
|
m.c[0].r[0].u64 = v0;
|
|
|
|
m.c[0].vecsize = 1;
|
|
|
|
m.columns = 1;
|
|
|
|
|
2017-09-27 13:16:33 +00:00
|
|
|
if (specialized)
|
|
|
|
m.c[0].id[0] = self;
|
2016-07-27 08:59:00 +00:00
|
|
|
}
|
|
|
|
|
2017-09-27 13:16:33 +00:00
|
|
|
// Construct vector.
|
|
|
|
SPIRConstant(uint32_t constant_type_, const SPIRConstant * const *vector_elements, uint32_t num_elements, bool specialized)
|
|
|
|
: constant_type(constant_type_), specialization(specialized)
|
2016-07-27 08:59:00 +00:00
|
|
|
{
|
2017-09-27 13:16:33 +00:00
|
|
|
m.c[0].vecsize = num_elements;
|
2016-07-27 08:59:00 +00:00
|
|
|
m.columns = 1;
|
|
|
|
|
2017-09-27 13:16:33 +00:00
|
|
|
for (uint32_t i = 0; i < num_elements; i++)
|
|
|
|
{
|
|
|
|
m.c[0].r[i] = vector_elements[i]->m.c[0].r[0];
|
|
|
|
m.c[0].id[i] = vector_elements[i]->m.c[0].id[0];
|
|
|
|
}
|
2016-05-05 07:33:18 +00:00
|
|
|
}
|
|
|
|
|
2017-09-27 13:16:33 +00:00
|
|
|
// Construct matrix.
|
|
|
|
SPIRConstant(uint32_t constant_type_, const ConstantVector *vectors, uint32_t num_vectors, bool specialized)
|
|
|
|
: constant_type(constant_type_), specialization(specialized)
|
2016-05-05 07:33:18 +00:00
|
|
|
{
|
2017-09-27 13:16:33 +00:00
|
|
|
m.columns = num_vectors;
|
|
|
|
memcpy(m.c, vectors, num_vectors * sizeof(*vectors));
|
2016-05-05 07:33:18 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
uint32_t constant_type;
|
|
|
|
ConstantMatrix m;
|
2017-09-27 13:16:33 +00:00
|
|
|
bool specialization = false; // If the constant is a specialization constant (i.e. created with OpSpecConstant*).
|
2016-05-05 07:33:18 +00:00
|
|
|
|
|
|
|
// For composites which are constant arrays, etc.
|
|
|
|
std::vector<uint32_t> subconstants;
|
|
|
|
};
|
|
|
|
|
|
|
|
class Variant
|
|
|
|
{
|
|
|
|
public:
|
|
|
|
// MSVC 2013 workaround, we shouldn't need these constructors.
|
|
|
|
Variant() = default;
|
|
|
|
Variant(Variant &&other)
|
|
|
|
{
|
|
|
|
*this = std::move(other);
|
|
|
|
}
|
|
|
|
Variant &operator=(Variant &&other)
|
|
|
|
{
|
|
|
|
if (this != &other)
|
|
|
|
{
|
|
|
|
holder = move(other.holder);
|
|
|
|
type = other.type;
|
|
|
|
other.type = TypeNone;
|
|
|
|
}
|
|
|
|
return *this;
|
|
|
|
}
|
|
|
|
|
|
|
|
void set(std::unique_ptr<IVariant> val, uint32_t new_type)
|
|
|
|
{
|
|
|
|
holder = std::move(val);
|
|
|
|
if (type != TypeNone && type != new_type)
|
2016-12-12 21:33:22 +00:00
|
|
|
SPIRV_CROSS_THROW("Overwriting a variant with new type.");
|
2016-05-05 07:33:18 +00:00
|
|
|
type = new_type;
|
|
|
|
}
|
|
|
|
|
|
|
|
template <typename T>
|
|
|
|
T &get()
|
|
|
|
{
|
|
|
|
if (!holder)
|
2016-12-12 21:33:22 +00:00
|
|
|
SPIRV_CROSS_THROW("nullptr");
|
2016-05-05 07:33:18 +00:00
|
|
|
if (T::type != type)
|
2016-12-12 21:33:22 +00:00
|
|
|
SPIRV_CROSS_THROW("Bad cast");
|
2016-05-05 07:33:18 +00:00
|
|
|
return *static_cast<T *>(holder.get());
|
|
|
|
}
|
|
|
|
|
|
|
|
template <typename T>
|
|
|
|
const T &get() const
|
|
|
|
{
|
|
|
|
if (!holder)
|
2016-12-12 21:33:22 +00:00
|
|
|
SPIRV_CROSS_THROW("nullptr");
|
2016-05-05 07:33:18 +00:00
|
|
|
if (T::type != type)
|
2016-12-12 21:33:22 +00:00
|
|
|
SPIRV_CROSS_THROW("Bad cast");
|
2016-05-05 07:33:18 +00:00
|
|
|
return *static_cast<const T *>(holder.get());
|
|
|
|
}
|
|
|
|
|
|
|
|
uint32_t get_type() const
|
|
|
|
{
|
|
|
|
return type;
|
|
|
|
}
|
|
|
|
bool empty() const
|
|
|
|
{
|
|
|
|
return !holder;
|
|
|
|
}
|
|
|
|
void reset()
|
|
|
|
{
|
|
|
|
holder.reset();
|
|
|
|
type = TypeNone;
|
|
|
|
}
|
|
|
|
|
|
|
|
private:
|
|
|
|
std::unique_ptr<IVariant> holder;
|
|
|
|
uint32_t type = TypeNone;
|
|
|
|
};
|
|
|
|
|
|
|
|
template <typename T>
|
|
|
|
T &variant_get(Variant &var)
|
|
|
|
{
|
|
|
|
return var.get<T>();
|
|
|
|
}
|
|
|
|
|
|
|
|
template <typename T>
|
|
|
|
const T &variant_get(const Variant &var)
|
|
|
|
{
|
|
|
|
return var.get<T>();
|
|
|
|
}
|
|
|
|
|
|
|
|
template <typename T, typename... P>
|
|
|
|
T &variant_set(Variant &var, P &&... args)
|
|
|
|
{
|
|
|
|
auto uptr = std::unique_ptr<T>(new T(std::forward<P>(args)...));
|
|
|
|
auto ptr = uptr.get();
|
|
|
|
var.set(std::move(uptr), T::type);
|
|
|
|
return *ptr;
|
|
|
|
}
|
|
|
|
|
|
|
|
struct Meta
|
|
|
|
{
|
|
|
|
struct Decoration
|
|
|
|
{
|
|
|
|
std::string alias;
|
2016-10-24 13:24:24 +00:00
|
|
|
std::string qualified_alias;
|
2016-05-05 07:33:18 +00:00
|
|
|
uint64_t decoration_flags = 0;
|
|
|
|
spv::BuiltIn builtin_type;
|
|
|
|
uint32_t location = 0;
|
|
|
|
uint32_t set = 0;
|
|
|
|
uint32_t binding = 0;
|
|
|
|
uint32_t offset = 0;
|
|
|
|
uint32_t array_stride = 0;
|
2017-01-21 10:30:33 +00:00
|
|
|
uint32_t matrix_stride = 0;
|
2016-05-05 07:33:18 +00:00
|
|
|
uint32_t input_attachment = 0;
|
2016-09-17 13:16:07 +00:00
|
|
|
uint32_t spec_id = 0;
|
2016-05-05 07:33:18 +00:00
|
|
|
bool builtin = false;
|
|
|
|
};
|
|
|
|
|
|
|
|
Decoration decoration;
|
|
|
|
std::vector<Decoration> members;
|
|
|
|
uint32_t sampler = 0;
|
2017-04-25 18:10:24 +00:00
|
|
|
|
2017-04-26 07:25:28 +00:00
|
|
|
std::unordered_map<uint32_t, uint32_t> decoration_word_offset;
|
2017-04-19 15:33:14 +00:00
|
|
|
|
|
|
|
// Used when the parser has detected a candidate identifier which matches
|
|
|
|
// known "magic" counter buffers as emitted by HLSL frontends.
|
|
|
|
// We will need to match the identifiers by name later when reflecting resources.
|
|
|
|
// We could use the regular alias later, but the alias will be mangled when parsing SPIR-V because the identifier
|
|
|
|
// is not a valid identifier in any high-level language.
|
|
|
|
std::string hlsl_magic_counter_buffer_name;
|
|
|
|
bool hlsl_magic_counter_buffer_candidate = false;
|
2016-05-05 07:33:18 +00:00
|
|
|
};
|
2016-09-20 08:17:41 +00:00
|
|
|
|
|
|
|
// A user callback that remaps the type of any variable.
|
2016-09-20 08:55:09 +00:00
|
|
|
// var_name is the declared name of the variable.
|
|
|
|
// name_of_type is the textual name of the type which will be used in the code unless written to by the callback.
|
2016-09-20 08:17:41 +00:00
|
|
|
using VariableTypeRemapCallback =
|
|
|
|
std::function<void(const SPIRType &type, const std::string &var_name, std::string &name_of_type)>;
|
2017-01-11 14:57:05 +00:00
|
|
|
|
|
|
|
class ClassicLocale
|
|
|
|
{
|
|
|
|
public:
|
|
|
|
ClassicLocale()
|
|
|
|
{
|
|
|
|
old = std::locale::global(std::locale::classic());
|
|
|
|
}
|
|
|
|
~ClassicLocale()
|
|
|
|
{
|
|
|
|
std::locale::global(old);
|
|
|
|
}
|
|
|
|
|
|
|
|
private:
|
|
|
|
std::locale old;
|
|
|
|
};
|
2016-05-05 07:33:18 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
#endif
|