9b4defe202
Support flattening StorageOutput & StorageInput matrices and arrays. No longer move matrix & array inputs to separate buffer. Add separate SPIRFunction::fixup_statements_in & SPIRFunction::fixup_statements_out instead of just SPIRFunction::fixup_statements. Emit SPIRFunction::fixup_statements at beginning of functions. CompilerMSL track vars_needing_early_declaration. Pass global output variables as variables to functions that access them. Sort input structs by location, same as output structs. Emit struct declarations in order output, input, uniforms. Regenerate reference shaders to new formats defined by above.
1285 lines
28 KiB
C++
1285 lines
28 KiB
C++
/*
|
|
* Copyright 2015-2018 ARM Limited
|
|
*
|
|
* 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.
|
|
*/
|
|
|
|
#ifndef SPIRV_CROSS_COMMON_HPP
|
|
#define SPIRV_CROSS_COMMON_HPP
|
|
|
|
#include "spirv.hpp"
|
|
|
|
#include <algorithm>
|
|
#include <cstdio>
|
|
#include <cstring>
|
|
#include <functional>
|
|
#include <locale>
|
|
#include <memory>
|
|
#include <sstream>
|
|
#include <stack>
|
|
#include <stdexcept>
|
|
#include <stdint.h>
|
|
#include <string>
|
|
#include <unordered_map>
|
|
#include <unordered_set>
|
|
#include <utility>
|
|
#include <vector>
|
|
|
|
namespace spirv_cross
|
|
{
|
|
|
|
#ifdef SPIRV_CROSS_EXCEPTIONS_TO_ASSERTIONS
|
|
#ifndef _MSC_VER
|
|
[[noreturn]]
|
|
#endif
|
|
inline void
|
|
report_and_abort(const std::string &msg)
|
|
{
|
|
#ifdef NDEBUG
|
|
(void)msg;
|
|
#else
|
|
fprintf(stderr, "There was a compiler error: %s\n", msg.c_str());
|
|
#endif
|
|
fflush(stderr);
|
|
abort();
|
|
}
|
|
|
|
#define SPIRV_CROSS_THROW(x) report_and_abort(x)
|
|
#else
|
|
class CompilerError : public std::runtime_error
|
|
{
|
|
public:
|
|
CompilerError(const std::string &str)
|
|
: std::runtime_error(str)
|
|
{
|
|
}
|
|
};
|
|
|
|
#define SPIRV_CROSS_THROW(x) throw CompilerError(x)
|
|
#endif
|
|
|
|
#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
|
|
|
|
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)...);
|
|
}
|
|
} // namespace inner
|
|
|
|
class Bitset
|
|
{
|
|
public:
|
|
Bitset() = default;
|
|
explicit inline Bitset(uint64_t lower_)
|
|
: lower(lower_)
|
|
{
|
|
}
|
|
|
|
inline bool get(uint32_t bit) const
|
|
{
|
|
if (bit < 64)
|
|
return (lower & (1ull << bit)) != 0;
|
|
else
|
|
return higher.count(bit) != 0;
|
|
}
|
|
|
|
inline void set(uint32_t bit)
|
|
{
|
|
if (bit < 64)
|
|
lower |= 1ull << bit;
|
|
else
|
|
higher.insert(bit);
|
|
}
|
|
|
|
inline void clear(uint32_t bit)
|
|
{
|
|
if (bit < 64)
|
|
lower &= ~(1ull << bit);
|
|
else
|
|
higher.erase(bit);
|
|
}
|
|
|
|
inline uint64_t get_lower() const
|
|
{
|
|
return lower;
|
|
}
|
|
|
|
inline void reset()
|
|
{
|
|
lower = 0;
|
|
higher.clear();
|
|
}
|
|
|
|
inline void merge_and(const Bitset &other)
|
|
{
|
|
lower &= other.lower;
|
|
std::unordered_set<uint32_t> tmp_set;
|
|
for (auto &v : higher)
|
|
if (other.higher.count(v) != 0)
|
|
tmp_set.insert(v);
|
|
higher = std::move(tmp_set);
|
|
}
|
|
|
|
inline void merge_or(const Bitset &other)
|
|
{
|
|
lower |= other.lower;
|
|
for (auto &v : other.higher)
|
|
higher.insert(v);
|
|
}
|
|
|
|
inline bool operator==(const Bitset &other) const
|
|
{
|
|
if (lower != other.lower)
|
|
return false;
|
|
|
|
if (higher.size() != other.higher.size())
|
|
return false;
|
|
|
|
for (auto &v : higher)
|
|
if (other.higher.count(v) == 0)
|
|
return false;
|
|
|
|
return true;
|
|
}
|
|
|
|
inline bool operator!=(const Bitset &other) const
|
|
{
|
|
return !(*this == other);
|
|
}
|
|
|
|
template <typename Op>
|
|
void for_each_bit(const Op &op) const
|
|
{
|
|
// TODO: Add ctz-based iteration.
|
|
for (uint32_t i = 0; i < 64; i++)
|
|
{
|
|
if (lower & (1ull << i))
|
|
op(i);
|
|
}
|
|
|
|
if (higher.empty())
|
|
return;
|
|
|
|
// Need to enforce an order here for reproducible results,
|
|
// but hitting this path should happen extremely rarely, so having this slow path is fine.
|
|
std::vector<uint32_t> bits;
|
|
bits.reserve(higher.size());
|
|
for (auto &v : higher)
|
|
bits.push_back(v);
|
|
std::sort(std::begin(bits), std::end(bits));
|
|
|
|
for (auto &v : bits)
|
|
op(v);
|
|
}
|
|
|
|
inline bool empty() const
|
|
{
|
|
return lower == 0 && higher.empty();
|
|
}
|
|
|
|
private:
|
|
// The most common bits to set are all lower than 64,
|
|
// so optimize for this case. Bits spilling outside 64 go into a slower data structure.
|
|
// In almost all cases, higher data structure will not be used.
|
|
uint64_t lower = 0;
|
|
std::unordered_set<uint32_t> higher;
|
|
};
|
|
|
|
// 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
|
|
#ifndef SPIRV_CROSS_FLT_FMT
|
|
#define SPIRV_CROSS_FLT_FMT "%.32g"
|
|
#endif
|
|
|
|
#ifdef _MSC_VER
|
|
// sprintf warning.
|
|
// We cannot rely on snprintf existing because, ..., MSVC.
|
|
#pragma warning(push)
|
|
#pragma warning(disable : 4996)
|
|
#endif
|
|
|
|
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;
|
|
}
|
|
|
|
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;
|
|
}
|
|
|
|
#ifdef _MSC_VER
|
|
#pragma warning(pop)
|
|
#endif
|
|
|
|
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,
|
|
TypeConstantOp,
|
|
TypeCombinedImageSampler,
|
|
TypeAccessChain,
|
|
TypeUndef
|
|
};
|
|
|
|
struct SPIRUndef : IVariant
|
|
{
|
|
enum
|
|
{
|
|
type = TypeUndef
|
|
};
|
|
SPIRUndef(uint32_t basetype_)
|
|
: basetype(basetype_)
|
|
{
|
|
}
|
|
uint32_t basetype;
|
|
};
|
|
|
|
// 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;
|
|
};
|
|
|
|
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;
|
|
};
|
|
|
|
struct SPIRType : IVariant
|
|
{
|
|
enum
|
|
{
|
|
type = TypeType
|
|
};
|
|
|
|
enum BaseType
|
|
{
|
|
Unknown,
|
|
Void,
|
|
Boolean,
|
|
Char,
|
|
Int,
|
|
UInt,
|
|
Int64,
|
|
UInt64,
|
|
AtomicCounter,
|
|
Half,
|
|
Float,
|
|
Double,
|
|
Struct,
|
|
Image,
|
|
SampledImage,
|
|
Sampler
|
|
};
|
|
|
|
// Scalar/vector/matrix support.
|
|
BaseType basetype = Unknown;
|
|
uint32_t width = 0;
|
|
uint32_t vecsize = 1;
|
|
uint32_t columns = 1;
|
|
|
|
// Arrays, support array of arrays by having a vector of array sizes.
|
|
std::vector<uint32_t> array;
|
|
|
|
// 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;
|
|
|
|
// Pointers
|
|
bool pointer = false;
|
|
spv::StorageClass storage = spv::StorageClassGeneric;
|
|
|
|
std::vector<uint32_t> member_types;
|
|
|
|
struct ImageType
|
|
{
|
|
uint32_t type;
|
|
spv::Dim dim;
|
|
bool depth;
|
|
bool arrayed;
|
|
bool ms;
|
|
uint32_t sampled;
|
|
spv::ImageFormat format;
|
|
spv::AccessQualifier access;
|
|
} image;
|
|
|
|
// 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;
|
|
|
|
// 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;
|
|
|
|
// Used in backends to avoid emitting members with conflicting names.
|
|
std::unordered_set<std::string> member_name_cache;
|
|
};
|
|
|
|
struct SPIRExtension : IVariant
|
|
{
|
|
enum
|
|
{
|
|
type = TypeExtension
|
|
};
|
|
|
|
enum Extension
|
|
{
|
|
Unsupported,
|
|
GLSL,
|
|
SPV_AMD_shader_ballot,
|
|
SPV_AMD_shader_explicit_vertex_parameter,
|
|
SPV_AMD_shader_trinary_minmax,
|
|
SPV_AMD_gcn_shader
|
|
};
|
|
|
|
SPIRExtension(Extension ext_)
|
|
: ext(ext_)
|
|
{
|
|
}
|
|
|
|
Extension ext;
|
|
};
|
|
|
|
// 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, const std::string &entry_name)
|
|
: self(self_)
|
|
, name(entry_name)
|
|
, orig_name(entry_name)
|
|
, model(execution_model)
|
|
{
|
|
}
|
|
SPIREntryPoint() = default;
|
|
|
|
uint32_t self = 0;
|
|
std::string name;
|
|
std::string orig_name;
|
|
std::vector<uint32_t> interface_variables;
|
|
|
|
Bitset flags;
|
|
struct
|
|
{
|
|
uint32_t x = 0, y = 0, z = 0;
|
|
uint32_t constant = 0; // Workgroup size can be expressed as a constant/spec-constant instead.
|
|
} workgroup_size;
|
|
uint32_t invocations = 0;
|
|
uint32_t output_vertices = 0;
|
|
spv::ExecutionModel model;
|
|
};
|
|
|
|
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.
|
|
// An expression being immutable can be speculative,
|
|
// it is assumed that this is true almost always.
|
|
bool immutable = false;
|
|
|
|
// Before use, this expression must be transposed.
|
|
// This is needed for targets which don't support row_major layouts.
|
|
bool need_transpose = false;
|
|
|
|
// A list of expressions which this expression depends on.
|
|
std::vector<uint32_t> expression_dependencies;
|
|
};
|
|
|
|
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,
|
|
MergeToSelectContinueForLoop
|
|
};
|
|
|
|
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;
|
|
|
|
// Declare these temporaries, but only conditionally if this block turns out to be
|
|
// a complex loop header.
|
|
std::vector<std::pair<uint32_t, uint32_t>> potential_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;
|
|
|
|
// 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;
|
|
|
|
// 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;
|
|
|
|
// Some expressions are control-flow dependent, i.e. any instruction which relies on derivatives or
|
|
// sub-group-like operations.
|
|
// Make sure that we only use these expressions in the original block.
|
|
std::vector<uint32_t> invalidate_expressions;
|
|
};
|
|
|
|
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;
|
|
|
|
// 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;
|
|
};
|
|
|
|
// 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
|
|
{
|
|
uint32_t id;
|
|
uint32_t image_id;
|
|
uint32_t sampler_id;
|
|
bool global_image;
|
|
bool global_sampler;
|
|
bool depth;
|
|
};
|
|
|
|
uint32_t return_type;
|
|
uint32_t function_type;
|
|
std::vector<Parameter> arguments;
|
|
|
|
// Can be used by backends to add magic arguments.
|
|
// Currently used by combined image/sampler implementation.
|
|
|
|
std::vector<Parameter> shadow_arguments;
|
|
std::vector<uint32_t> local_variables;
|
|
uint32_t entry_block = 0;
|
|
std::vector<uint32_t> blocks;
|
|
std::vector<CombinedImageSamplerParameter> combined_parameters;
|
|
|
|
void add_local_variable(uint32_t id)
|
|
{
|
|
local_variables.push_back(id);
|
|
}
|
|
|
|
void add_parameter(uint32_t parameter_type, uint32_t id, bool alias_global_variable = false)
|
|
{
|
|
// Arguments are read-only until proven otherwise.
|
|
arguments.push_back({ parameter_type, id, 0u, 0u, alias_global_variable });
|
|
}
|
|
|
|
// Statements to be emitted when the function returns.
|
|
// Mostly used for lowering internal data structures onto flattened structures.
|
|
std::vector<std::string> fixup_statements_out;
|
|
|
|
// Statements to be emitted when the function begins.
|
|
// Mostly used for populating internal data structures from flattened structures.
|
|
std::vector<std::string> fixup_statements_in;
|
|
|
|
bool active = false;
|
|
bool flush_undeclared = true;
|
|
bool do_combined_parameters = true;
|
|
bool analyzed_variable_scope = false;
|
|
};
|
|
|
|
struct SPIRAccessChain : IVariant
|
|
{
|
|
enum
|
|
{
|
|
type = TypeAccessChain
|
|
};
|
|
|
|
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_)
|
|
{
|
|
}
|
|
|
|
// 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;
|
|
uint32_t matrix_stride = 0;
|
|
bool row_major_matrix = false;
|
|
bool immutable = false;
|
|
};
|
|
|
|
struct SPIRVariable : IVariant
|
|
{
|
|
enum
|
|
{
|
|
type = TypeVariable
|
|
};
|
|
|
|
SPIRVariable() = default;
|
|
SPIRVariable(uint32_t basetype_, spv::StorageClass storage_, uint32_t initializer_ = 0, uint32_t basevariable_ = 0)
|
|
: basetype(basetype_)
|
|
, storage(storage_)
|
|
, initializer(initializer_)
|
|
, basevariable(basevariable_)
|
|
{
|
|
}
|
|
|
|
uint32_t basetype = 0;
|
|
spv::StorageClass storage = spv::StorageClassGeneric;
|
|
uint32_t decoration = 0;
|
|
uint32_t initializer = 0;
|
|
uint32_t basevariable = 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;
|
|
uint32_t remapped_components = 0;
|
|
|
|
// 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;
|
|
|
|
SPIRFunction::Parameter *parameter = nullptr;
|
|
};
|
|
|
|
struct SPIRConstant : IVariant
|
|
{
|
|
enum
|
|
{
|
|
type = TypeConstant
|
|
};
|
|
|
|
union Constant {
|
|
uint32_t u32;
|
|
int32_t i32;
|
|
float f32;
|
|
|
|
uint64_t u64;
|
|
int64_t i64;
|
|
double f64;
|
|
};
|
|
|
|
struct ConstantVector
|
|
{
|
|
Constant r[4];
|
|
// If != 0, this element is a specialization constant, and we should keep track of it as such.
|
|
uint32_t id[4];
|
|
uint32_t vecsize = 1;
|
|
|
|
// Workaround for MSVC 2013, initializing an array breaks.
|
|
ConstantVector()
|
|
{
|
|
memset(r, 0, sizeof(r));
|
|
for (unsigned i = 0; i < 4; i++)
|
|
id[i] = 0;
|
|
}
|
|
};
|
|
|
|
struct ConstantMatrix
|
|
{
|
|
ConstantVector c[4];
|
|
// If != 0, this column is a specialization constant, and we should keep track of it as such.
|
|
uint32_t id[4];
|
|
uint32_t columns = 1;
|
|
|
|
// Workaround for MSVC 2013, initializing an array breaks.
|
|
ConstantMatrix()
|
|
{
|
|
for (unsigned i = 0; i < 4; i++)
|
|
id[i] = 0;
|
|
}
|
|
};
|
|
|
|
static inline float f16_to_f32(uint16_t u16_value)
|
|
{
|
|
// Based on the GLM implementation.
|
|
int s = (u16_value >> 15) & 0x1;
|
|
int e = (u16_value >> 10) & 0x1f;
|
|
int m = (u16_value >> 0) & 0x3ff;
|
|
|
|
union {
|
|
float f32;
|
|
uint32_t u32;
|
|
} u;
|
|
|
|
if (e == 0)
|
|
{
|
|
if (m == 0)
|
|
{
|
|
u.u32 = uint32_t(s) << 31;
|
|
return u.f32;
|
|
}
|
|
else
|
|
{
|
|
while ((m & 0x400) == 0)
|
|
{
|
|
m <<= 1;
|
|
e--;
|
|
}
|
|
|
|
e++;
|
|
m &= ~0x400;
|
|
}
|
|
}
|
|
else if (e == 31)
|
|
{
|
|
if (m == 0)
|
|
{
|
|
u.u32 = (uint32_t(s) << 31) | 0x7f800000u;
|
|
return u.f32;
|
|
}
|
|
else
|
|
{
|
|
u.u32 = (uint32_t(s) << 31) | 0x7f800000u | (m << 13);
|
|
return u.f32;
|
|
}
|
|
}
|
|
|
|
e += 127 - 15;
|
|
m <<= 13;
|
|
u.u32 = (uint32_t(s) << 31) | (e << 23) | m;
|
|
return u.f32;
|
|
}
|
|
|
|
inline uint32_t specialization_constant_id(uint32_t col, uint32_t row) const
|
|
{
|
|
return m.c[col].id[row];
|
|
}
|
|
|
|
inline uint32_t specialization_constant_id(uint32_t col) const
|
|
{
|
|
return m.id[col];
|
|
}
|
|
|
|
inline uint32_t scalar(uint32_t col = 0, uint32_t row = 0) const
|
|
{
|
|
return m.c[col].r[row].u32;
|
|
}
|
|
|
|
inline uint16_t scalar_u16(uint32_t col = 0, uint32_t row = 0) const
|
|
{
|
|
return uint16_t(m.c[col].r[row].u32 & 0xffffu);
|
|
}
|
|
|
|
inline float scalar_f16(uint32_t col = 0, uint32_t row = 0) const
|
|
{
|
|
return f16_to_f32(scalar_u16(col, row));
|
|
}
|
|
|
|
inline float scalar_f32(uint32_t col = 0, uint32_t row = 0) const
|
|
{
|
|
return m.c[col].r[row].f32;
|
|
}
|
|
|
|
inline int32_t scalar_i32(uint32_t col = 0, uint32_t row = 0) const
|
|
{
|
|
return m.c[col].r[row].i32;
|
|
}
|
|
|
|
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;
|
|
}
|
|
|
|
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;
|
|
}
|
|
|
|
inline void make_null(const SPIRType &constant_type_)
|
|
{
|
|
m = {};
|
|
m.columns = constant_type_.columns;
|
|
for (auto &c : m.c)
|
|
c.vecsize = constant_type_.vecsize;
|
|
}
|
|
|
|
explicit SPIRConstant(uint32_t constant_type_)
|
|
: constant_type(constant_type_)
|
|
{
|
|
}
|
|
|
|
SPIRConstant() = default;
|
|
|
|
SPIRConstant(uint32_t constant_type_, const uint32_t *elements, uint32_t num_elements, bool specialized)
|
|
: constant_type(constant_type_)
|
|
, specialization(specialized)
|
|
{
|
|
subconstants.insert(end(subconstants), elements, elements + num_elements);
|
|
specialization = specialized;
|
|
}
|
|
|
|
// Construct scalar (32-bit).
|
|
SPIRConstant(uint32_t constant_type_, uint32_t v0, bool specialized)
|
|
: constant_type(constant_type_)
|
|
, specialization(specialized)
|
|
{
|
|
m.c[0].r[0].u32 = v0;
|
|
m.c[0].vecsize = 1;
|
|
m.columns = 1;
|
|
}
|
|
|
|
// Construct scalar (64-bit).
|
|
SPIRConstant(uint32_t constant_type_, uint64_t v0, bool specialized)
|
|
: constant_type(constant_type_)
|
|
, specialization(specialized)
|
|
{
|
|
m.c[0].r[0].u64 = v0;
|
|
m.c[0].vecsize = 1;
|
|
m.columns = 1;
|
|
}
|
|
|
|
// Construct vectors and matrices.
|
|
SPIRConstant(uint32_t constant_type_, const SPIRConstant *const *vector_elements, uint32_t num_elements,
|
|
bool specialized)
|
|
: constant_type(constant_type_)
|
|
, specialization(specialized)
|
|
{
|
|
bool matrix = vector_elements[0]->m.c[0].vecsize > 1;
|
|
|
|
if (matrix)
|
|
{
|
|
m.columns = num_elements;
|
|
|
|
for (uint32_t i = 0; i < num_elements; i++)
|
|
{
|
|
m.c[i] = vector_elements[i]->m.c[0];
|
|
if (vector_elements[i]->specialization)
|
|
m.id[i] = vector_elements[i]->self;
|
|
}
|
|
}
|
|
else
|
|
{
|
|
m.c[0].vecsize = num_elements;
|
|
m.columns = 1;
|
|
|
|
for (uint32_t i = 0; i < num_elements; i++)
|
|
{
|
|
m.c[0].r[i] = vector_elements[i]->m.c[0].r[0];
|
|
if (vector_elements[i]->specialization)
|
|
m.c[0].id[i] = vector_elements[i]->self;
|
|
}
|
|
}
|
|
}
|
|
|
|
uint32_t constant_type;
|
|
ConstantMatrix m;
|
|
|
|
// If this constant is a specialization constant (i.e. created with OpSpecConstant*).
|
|
bool specialization = false;
|
|
// If this constant is used as an array length which creates specialization restrictions on some backends.
|
|
bool is_used_as_array_length = false;
|
|
|
|
// 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 (!allow_type_rewrite && type != TypeNone && type != new_type)
|
|
SPIRV_CROSS_THROW("Overwriting a variant with new type.");
|
|
type = new_type;
|
|
allow_type_rewrite = false;
|
|
}
|
|
|
|
template <typename T>
|
|
T &get()
|
|
{
|
|
if (!holder)
|
|
SPIRV_CROSS_THROW("nullptr");
|
|
if (T::type != type)
|
|
SPIRV_CROSS_THROW("Bad cast");
|
|
return *static_cast<T *>(holder.get());
|
|
}
|
|
|
|
template <typename T>
|
|
const T &get() const
|
|
{
|
|
if (!holder)
|
|
SPIRV_CROSS_THROW("nullptr");
|
|
if (T::type != type)
|
|
SPIRV_CROSS_THROW("Bad cast");
|
|
return *static_cast<const T *>(holder.get());
|
|
}
|
|
|
|
uint32_t get_type() const
|
|
{
|
|
return type;
|
|
}
|
|
uint32_t get_id() const
|
|
{
|
|
return holder ? holder->self : 0;
|
|
}
|
|
bool empty() const
|
|
{
|
|
return !holder;
|
|
}
|
|
void reset()
|
|
{
|
|
holder.reset();
|
|
type = TypeNone;
|
|
}
|
|
|
|
void set_allow_type_rewrite()
|
|
{
|
|
allow_type_rewrite = true;
|
|
}
|
|
|
|
private:
|
|
std::unique_ptr<IVariant> holder;
|
|
uint32_t type = TypeNone;
|
|
bool allow_type_rewrite = false;
|
|
};
|
|
|
|
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;
|
|
std::string qualified_alias;
|
|
std::string hlsl_semantic;
|
|
Bitset decoration_flags;
|
|
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;
|
|
uint32_t matrix_stride = 0;
|
|
uint32_t input_attachment = 0;
|
|
uint32_t spec_id = 0;
|
|
uint32_t index = 0;
|
|
bool builtin = false;
|
|
};
|
|
|
|
Decoration decoration;
|
|
std::vector<Decoration> members;
|
|
uint32_t sampler = 0;
|
|
|
|
std::unordered_map<uint32_t, uint32_t> decoration_word_offset;
|
|
|
|
// 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;
|
|
|
|
// For SPV_GOOGLE_hlsl_functionality1, this avoids the workaround.
|
|
bool hlsl_is_magic_counter_buffer = false;
|
|
// ID for the sibling counter buffer.
|
|
uint32_t hlsl_magic_counter_buffer = 0;
|
|
};
|
|
|
|
// A user callback that remaps the type of any variable.
|
|
// 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.
|
|
using VariableTypeRemapCallback =
|
|
std::function<void(const SPIRType &type, const std::string &var_name, std::string &name_of_type)>;
|
|
|
|
class ClassicLocale
|
|
{
|
|
public:
|
|
ClassicLocale()
|
|
{
|
|
old = std::locale::global(std::locale::classic());
|
|
}
|
|
~ClassicLocale()
|
|
{
|
|
std::locale::global(old);
|
|
}
|
|
|
|
private:
|
|
std::locale old;
|
|
};
|
|
|
|
class Hasher
|
|
{
|
|
public:
|
|
inline void u32(uint32_t value)
|
|
{
|
|
h = (h * 0x100000001b3ull) ^ value;
|
|
}
|
|
|
|
inline uint64_t get() const
|
|
{
|
|
return h;
|
|
}
|
|
|
|
private:
|
|
uint64_t h = 0xcbf29ce484222325ull;
|
|
};
|
|
|
|
static inline bool type_is_floating_point(const SPIRType &type)
|
|
{
|
|
return type.basetype == SPIRType::Half || type.basetype == SPIRType::Float || type.basetype == SPIRType::Double;
|
|
}
|
|
} // namespace spirv_cross
|
|
|
|
#endif
|