Mirrors/bgfx
1
0
Fork 0
mirror of https://github.com/bkaradzic/bgfx.git synced 2026-02-17 20:52:36 +01:00
Code Issues Packages Projects Releases Wiki Activity

Updated spirv-cross.

Browse Source
This commit is contained in:
Бранимир Караџић
2025-03-14 22:04:08 -07:00
parent ea634e3b11
commit 2fb9eec9a4
11 changed files with 5017 additions and 136 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2296
3rdparty/spirv-cross/spirv.h vendored
View File

File diff suppressed because it is too large Load Diff

2344
3rdparty/spirv-cross/spirv.hpp vendored
View File

File diff suppressed because it is too large Load Diff

4
3rdparty/spirv-cross/spirv_common.hpp vendored
View File

@@ -1410,6 +1410,10 @@ struct SPIRConstant : IVariant
// If true, this is a LUT, and should always be declared in the outer scope.
bool is_used_as_lut = false;
// If this is a null constant of array type with specialized length.
// May require special handling in initializer
bool is_null_array_specialized_length = false;
// For composites which are constant arrays, etc.
SmallVector<ConstantID> subconstants;

16
3rdparty/spirv-cross/spirv_cross.cpp vendored
View File

@@ -587,6 +587,7 @@ const SPIRType &Compiler::expression_type(uint32_t id) const
bool Compiler::expression_is_lvalue(uint32_t id) const
{
auto &type = expression_type(id);
switch (type.basetype)
{
case SPIRType::SampledImage:
@@ -4910,13 +4911,16 @@ void Compiler::make_constant_null(uint32_t id, uint32_t type)
uint32_t parent_id = ir.increase_bound_by(1);
make_constant_null(parent_id, constant_type.parent_type);
if (!constant_type.array_size_literal.back())
SPIRV_CROSS_THROW("Array size of OpConstantNull must be a literal.");
SmallVector<uint32_t> elements(constant_type.array.back());
for (uint32_t i = 0; i < constant_type.array.back(); i++)
// The array size of OpConstantNull can be either literal or specialization constant.
// In the latter case, we cannot take the value as-is, as it can be changed to anything.
// Rather, we assume it to be *one* for the sake of initializer.
bool is_literal_array_size = constant_type.array_size_literal.back();
uint32_t count = is_literal_array_size ? constant_type.array.back() : 1;
SmallVector<uint32_t> elements(count);
for (uint32_t i = 0; i < count; i++)
elements[i] = parent_id;
set<SPIRConstant>(id, type, elements.data(), uint32_t(elements.size()), false);
auto &constant = set<SPIRConstant>(id, type, elements.data(), uint32_t(elements.size()), false);
constant.is_null_array_specialized_length = !is_literal_array_size;
}
else if (!constant_type.member_types.empty())
{

15
3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp vendored
View File

@@ -1050,16 +1050,21 @@ void ParsedIR::make_constant_null(uint32_t id, uint32_t type, bool add_to_typed_
uint32_t parent_id = increase_bound_by(1);
make_constant_null(parent_id, constant_type.parent_type, add_to_typed_id_set);
if (!constant_type.array_size_literal.back())
SPIRV_CROSS_THROW("Array size of OpConstantNull must be a literal.");
// The array size of OpConstantNull can be either literal or specialization constant.
// In the latter case, we cannot take the value as-is, as it can be changed to anything.
// Rather, we assume it to be *one* for the sake of initializer.
bool is_literal_array_size = constant_type.array_size_literal.back();
uint32_t count = is_literal_array_size ? constant_type.array.back() : 1;
SmallVector<uint32_t> elements(constant_type.array.back());
for (uint32_t i = 0; i < constant_type.array.back(); i++)
SmallVector<uint32_t> elements(count);
for (uint32_t i = 0; i < count; i++)
elements[i] = parent_id;
if (add_to_typed_id_set)
add_typed_id(TypeConstant, id);
variant_set<SPIRConstant>(ids[id], type, elements.data(), uint32_t(elements.size()), false).self = id;
auto& constant = variant_set<SPIRConstant>(ids[id], type, elements.data(), uint32_t(elements.size()), false);
constant.self = id;
constant.is_null_array_specialized_length = !is_literal_array_size;
}
else if (!constant_type.member_types.empty())
{

118
3rdparty/spirv-cross/spirv_glsl.cpp vendored
View File

@@ -681,6 +681,8 @@ string CompilerGLSL::compile()
backend.requires_relaxed_precision_analysis = options.es || options.vulkan_semantics;
backend.support_precise_qualifier =
(!options.es && options.version >= 400) || (options.es && options.version >= 320);
backend.constant_null_initializer = "{ }";
backend.requires_matching_array_initializer = true;
if (is_legacy_es())
backend.support_case_fallthrough = false;
@@ -2841,7 +2843,7 @@ void CompilerGLSL::emit_uniform(const SPIRVariable &var)
statement(layout_for_variable(var), variable_decl(var), ";");
}
string CompilerGLSL::constant_value_macro_name(uint32_t id)
string CompilerGLSL::constant_value_macro_name(uint32_t id) const
{
return join("SPIRV_CROSS_CONSTANT_ID_", id);
}
@@ -4955,12 +4957,16 @@ void CompilerGLSL::emit_polyfills(uint32_t polyfills, bool relaxed)
// Subclasses may override to modify the return value.
string CompilerGLSL::to_func_call_arg(const SPIRFunction::Parameter &, uint32_t id)
{
// BDA expects pointers through function interface.
if (is_physical_pointer(expression_type(id)))
return to_pointer_expression(id);
// Make sure that we use the name of the original variable, and not the parameter alias.
uint32_t name_id = id;
auto *var = maybe_get<SPIRVariable>(id);
if (var && var->basevariable)
name_id = var->basevariable;
return to_expression(name_id);
return to_unpacked_expression(name_id);
}
void CompilerGLSL::force_temporary_and_recompile(uint32_t id)
@@ -5391,6 +5397,15 @@ string CompilerGLSL::to_non_uniform_aware_expression(uint32_t id)
return expr;
}
string CompilerGLSL::to_atomic_ptr_expression(uint32_t id)
{
string expr = to_non_uniform_aware_expression(id);
// If we have naked pointer to POD, we need to dereference to get the proper ".value" resolve.
if (should_dereference(id))
expr = dereference_expression(expression_type(id), expr);
return expr;
}
string CompilerGLSL::to_expression(uint32_t id, bool register_expression_read)
{
auto itr = invalid_expressions.find(id);
@@ -5898,6 +5913,11 @@ string CompilerGLSL::constant_expression(const SPIRConstant &c,
{
return backend.null_pointer_literal;
}
else if (c.is_null_array_specialized_length && backend.requires_matching_array_initializer)
{
require_extension_internal("GL_EXT_null_initializer");
return backend.constant_null_initializer;
}
else if (!c.subconstants.empty())
{
// Handles Arrays and structures.
@@ -6988,9 +7008,12 @@ void CompilerGLSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id,
require_extension_internal("GL_EXT_shader_atomic_float");
}
if (type.basetype == SPIRType::UInt64 || type.basetype == SPIRType::Int64)
require_extension_internal("GL_EXT_shader_atomic_int64");
forced_temporaries.insert(result_id);
emit_op(result_type, result_id,
join(op, "(", to_non_uniform_aware_expression(op0), ", ",
join(op, "(", to_atomic_ptr_expression(op0), ", ",
to_unpacked_expression(op1), ")"), false);
flush_all_atomic_capable_variables();
}
@@ -11249,7 +11272,7 @@ bool CompilerGLSL::should_dereference(uint32_t id)
{
const auto &type = expression_type(id);
// Non-pointer expressions don't need to be dereferenced.
if (!type.pointer)
if (!is_pointer(type))
return false;
// Handles shouldn't be dereferenced either.
@@ -11257,8 +11280,9 @@ bool CompilerGLSL::should_dereference(uint32_t id)
return false;
// If id is a variable but not a phi variable, we should not dereference it.
// BDA passed around as parameters are always pointers.
if (auto *var = maybe_get<SPIRVariable>(id))
return var->phi_variable;
return (var->parameter && is_physical_pointer(type)) || var->phi_variable;
if (auto *expr = maybe_get<SPIRExpression>(id))
{
@@ -11291,6 +11315,16 @@ bool CompilerGLSL::should_dereference(uint32_t id)
return true;
}
bool CompilerGLSL::should_dereference_caller_param(uint32_t id)
{
const auto &type = expression_type(id);
// BDA is always passed around as pointers.
if (is_physical_pointer(type))
return false;
return should_dereference(id);
}
bool CompilerGLSL::should_forward(uint32_t id) const
{
// If id is a variable we will try to forward it regardless of force_temporary check below
@@ -13853,8 +13887,11 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
const char *increment = unsigned_type ? "0u" : "0";
emit_op(ops[0], ops[1],
join(op, "(",
to_non_uniform_aware_expression(ops[2]), ", ", increment, ")"), false);
to_atomic_ptr_expression(ops[2]), ", ", increment, ")"), false);
flush_all_atomic_capable_variables();
if (type.basetype == SPIRType::UInt64 || type.basetype == SPIRType::Int64)
require_extension_internal("GL_EXT_shader_atomic_int64");
break;
}
@@ -13866,8 +13903,12 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
// Ignore semantics for now, probably only relevant to CL.
uint32_t val = ops[3];
const char *op = check_atomic_image(ptr) ? "imageAtomicExchange" : "atomicExchange";
statement(op, "(", to_non_uniform_aware_expression(ptr), ", ", to_expression(val), ");");
statement(op, "(", to_atomic_ptr_expression(ptr), ", ", to_expression(val), ");");
flush_all_atomic_capable_variables();
auto &type = expression_type(ptr);
if (type.basetype == SPIRType::UInt64 || type.basetype == SPIRType::Int64)
require_extension_internal("GL_EXT_shader_atomic_int64");
break;
}
@@ -13902,7 +13943,10 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
increment = "-1";
emit_op(ops[0], ops[1],
join(op, "(", to_non_uniform_aware_expression(ops[2]), ", ", increment, ")"), false);
join(op, "(", to_atomic_ptr_expression(ops[2]), ", ", increment, ")"), false);
if (type.basetype == SPIRType::UInt64 || type.basetype == SPIRType::Int64)
require_extension_internal("GL_EXT_shader_atomic_int64");
}
flush_all_atomic_capable_variables();
@@ -13921,9 +13965,13 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
{
const char *op = check_atomic_image(ops[2]) ? "imageAtomicAdd" : "atomicAdd";
forced_temporaries.insert(ops[1]);
auto expr = join(op, "(", to_non_uniform_aware_expression(ops[2]), ", -", to_enclosed_expression(ops[5]), ")");
auto expr = join(op, "(", to_atomic_ptr_expression(ops[2]), ", -", to_enclosed_expression(ops[5]), ")");
emit_op(ops[0], ops[1], expr, should_forward(ops[2]) && should_forward(ops[5]));
flush_all_atomic_capable_variables();
auto &type = get<SPIRType>(ops[0]);
if (type.basetype == SPIRType::UInt64 || type.basetype == SPIRType::Int64)
require_extension_internal("GL_EXT_shader_atomic_int64");
break;
}
@@ -14727,6 +14775,20 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
break;
}
case OpExtInstWithForwardRefsKHR:
{
uint32_t extension_set = ops[2];
auto ext = get<SPIRExtension>(extension_set).ext;
if (ext != SPIRExtension::SPV_debug_info &&
ext != SPIRExtension::NonSemanticShaderDebugInfo &&
ext != SPIRExtension::NonSemanticGeneric)
{
SPIRV_CROSS_THROW("Unexpected use of ExtInstWithForwardRefsKHR.");
}
break;
}
case OpExtInst:
{
uint32_t extension_set = ops[2];
@@ -15699,7 +15761,10 @@ string CompilerGLSL::argument_decl(const SPIRFunction::Parameter &arg)
auto &type = expression_type(arg.id);
const char *direction = "";
if (type.pointer)
if (is_pointer(type) &&
(type.storage == StorageClassFunction ||
type.storage == StorageClassPrivate ||
type.storage == StorageClassOutput))
{
// If we're passing around block types to function, we really mean reference in a pointer sense,
// but DXC does not like inout for mesh blocks, so workaround that. out is technically not correct,
@@ -15773,13 +15838,24 @@ string CompilerGLSL::variable_decl(const SPIRVariable &variable)
else if (options.force_zero_initialized_variables && type_can_zero_initialize(type))
res += join(" = ", to_zero_initialized_expression(get_variable_data_type_id(variable)));
}
else if (variable.initializer && !variable_decl_is_remapped_storage(variable, StorageClassWorkgroup))
else if (variable.initializer)
{
uint32_t expr = variable.initializer;
if (ir.ids[expr].get_type() != TypeUndef)
res += join(" = ", to_initializer_expression(variable));
else if (options.force_zero_initialized_variables && type_can_zero_initialize(type))
res += join(" = ", to_zero_initialized_expression(get_variable_data_type_id(variable)));
if (!variable_decl_is_remapped_storage(variable, StorageClassWorkgroup))
{
uint32_t expr = variable.initializer;
if (ir.ids[expr].get_type() != TypeUndef)
res += join(" = ", to_initializer_expression(variable));
else if (options.force_zero_initialized_variables && type_can_zero_initialize(type))
res += join(" = ", to_zero_initialized_expression(get_variable_data_type_id(variable)));
}
else
{
// Workgroup memory requires special handling. First, it can only be Null-Initialized.
// GLSL will handle this with null initializer, while others require more work after the decl
require_extension_internal("GL_EXT_null_initializer");
if (!backend.constant_null_initializer.empty())
res += join(" = ", backend.constant_null_initializer);
}
}
return res;
@@ -16540,6 +16616,12 @@ void CompilerGLSL::emit_function(SPIRFunction &func, const Bitset &return_flags)
// Comes from MSL which can push global variables as local variables in main function.
add_local_variable_name(var.self);
statement(variable_decl(var), ";");
// "Real" workgroup variables in compute shaders needs extra caretaking.
// They need to be initialized with an extra routine as they come in arbitrary form.
if (var.storage == StorageClassWorkgroup && var.initializer)
emit_workgroup_initialization(var);
var.deferred_declaration = false;
}
else if (var.storage == StorageClassPrivate)
@@ -16646,6 +16728,10 @@ void CompilerGLSL::emit_fixup()
}
}
void CompilerGLSL::emit_workgroup_initialization(const SPIRVariable &)
{
}
void CompilerGLSL::flush_phi(BlockID from, BlockID to)
{
auto &child = get<SPIRBlock>(to);

9
3rdparty/spirv-cross/spirv_glsl.hpp vendored
View File

@@ -297,6 +297,9 @@ public:
float_formatter = formatter;
}
// Returns the macro name corresponding to constant id
std::string constant_value_macro_name(uint32_t id) const;
protected:
struct ShaderSubgroupSupportHelper
{
@@ -450,6 +453,7 @@ protected:
virtual std::string variable_decl(const SPIRType &type, const std::string &name, uint32_t id = 0);
virtual bool variable_decl_is_remapped_storage(const SPIRVariable &var, spv::StorageClass storage) const;
virtual std::string to_func_call_arg(const SPIRFunction::Parameter &arg, uint32_t id);
virtual void emit_workgroup_initialization(const SPIRVariable &var);
struct TextureFunctionBaseArguments
{
@@ -622,6 +626,7 @@ protected:
const char *uint16_t_literal_suffix = "us";
const char *nonuniform_qualifier = "nonuniformEXT";
const char *boolean_mix_function = "mix";
std::string constant_null_initializer = "";
SPIRType::BaseType boolean_in_struct_remapped_type = SPIRType::Boolean;
bool swizzle_is_function = false;
bool shared_is_implied = false;
@@ -629,6 +634,7 @@ protected:
bool explicit_struct_type = false;
bool use_initializer_list = false;
bool use_typed_initializer_list = false;
bool requires_matching_array_initializer = false;
bool can_declare_struct_inline = true;
bool can_declare_arrays_inline = true;
bool native_row_major_matrix = true;
@@ -679,7 +685,6 @@ protected:
const SmallVector<uint32_t> &indices);
void emit_block_chain(SPIRBlock &block);
void emit_hoisted_temporaries(SmallVector<std::pair<TypeID, ID>> &temporaries);
std::string constant_value_macro_name(uint32_t id);
int get_constant_mapping_to_workgroup_component(const SPIRConstant &constant) const;
void emit_constant(const SPIRConstant &constant);
void emit_specialization_constant_op(const SPIRConstantOp &constant);
@@ -695,6 +700,7 @@ protected:
void emit_variable_temporary_copies(const SPIRVariable &var);
bool should_dereference(uint32_t id);
bool should_dereference_caller_param(uint32_t id);
bool should_forward(uint32_t id) const;
bool should_suppress_usage_tracking(uint32_t id) const;
void emit_mix_op(uint32_t result_type, uint32_t id, uint32_t left, uint32_t right, uint32_t lerp);
@@ -794,6 +800,7 @@ protected:
SPIRExpression &emit_uninitialized_temporary_expression(uint32_t type, uint32_t id);
void append_global_func_args(const SPIRFunction &func, uint32_t index, SmallVector<std::string> &arglist);
std::string to_non_uniform_aware_expression(uint32_t id);
std::string to_atomic_ptr_expression(uint32_t id);
std::string to_expression(uint32_t id, bool register_expression_read = true);
std::string to_composite_constructor_expression(const SPIRType &parent_type, uint32_t id, bool block_like_type);
std::string to_rerolled_array_expression(const SPIRType &parent_type, const std::string &expr, const SPIRType &type);

2
3rdparty/spirv-cross/spirv_hlsl.cpp vendored
View File

@@ -2476,12 +2476,14 @@ void CompilerHLSL::analyze_meshlet_writes()
set_decoration(op_type, DecorationPerPrimitiveEXT);
auto &arr = set<SPIRType>(op_arr, type);
arr.op = OpTypeArray;
arr.parent_type = type.self;
arr.array.push_back(per_primitive ? execution.output_primitives : execution.output_vertices);
arr.array_size_literal.push_back(true);
auto &ptr = set<SPIRType>(op_ptr, arr);
ptr.parent_type = arr.self;
ptr.op = OpTypePointer;
ptr.pointer = true;
ptr.pointer_depth++;
ptr.storage = StorageClassOutput;

338
3rdparty/spirv-cross/spirv_msl.cpp vendored
View File

@@ -272,7 +272,9 @@ void CompilerMSL::build_implicit_builtins()
(active_input_builtins.get(BuiltInVertexId) || active_input_builtins.get(BuiltInVertexIndex) ||
active_input_builtins.get(BuiltInBaseVertex) || active_input_builtins.get(BuiltInInstanceId) ||
active_input_builtins.get(BuiltInInstanceIndex) || active_input_builtins.get(BuiltInBaseInstance));
bool need_local_invocation_index = (msl_options.emulate_subgroups && active_input_builtins.get(BuiltInSubgroupId)) || is_mesh_shader();
bool need_local_invocation_index =
(msl_options.emulate_subgroups && active_input_builtins.get(BuiltInSubgroupId)) || is_mesh_shader() ||
needs_workgroup_zero_init;
bool need_workgroup_size = msl_options.emulate_subgroups && active_input_builtins.get(BuiltInNumSubgroups);
bool force_frag_depth_passthrough =
get_execution_model() == ExecutionModelFragment && !uses_explicit_early_fragment_test() && need_subpass_input &&
@@ -1611,7 +1613,7 @@ string CompilerMSL::compile()
backend.nonuniform_qualifier = "";
backend.support_small_type_sampling_result = true;
backend.force_merged_mesh_block = false;
backend.force_gl_in_out_block = get_execution_model() == ExecutionModelMeshEXT;
backend.force_gl_in_out_block = false;
backend.supports_empty_struct = true;
backend.support_64bit_switch = true;
backend.boolean_in_struct_remapped_type = SPIRType::Short;
@@ -1649,6 +1651,7 @@ string CompilerMSL::compile()
analyze_image_and_sampler_usage();
analyze_sampled_image_usage();
analyze_interlocked_resource_usage();
analyze_workgroup_variables();
preprocess_op_codes();
build_implicit_builtins();
@@ -2322,7 +2325,14 @@ void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std::
v.storage = StorageClassWorkgroup;
// Ensure the existing variable has a valid name and the new variable has all the same meta info
set_name(arg_id, ensure_valid_name(to_name(arg_id), "v"));
if (ir.meta[arg_id].decoration.builtin)
{
set_name(arg_id, builtin_to_glsl(bi_type, var.storage));
}
else
{
set_name(arg_id, ensure_valid_name(to_name(arg_id), "v"));
}
ir.meta[next_id] = ir.meta[arg_id];
}
else if (is_builtin && has_decoration(p_type->self, DecorationBlock))
@@ -3187,41 +3197,58 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass
string mbr_name = ensure_valid_name(append_member_name(mbr_name_qual, var_type, mbr_idx) + (mbr_is_indexable ? join("_", i) : ""), "m");
set_member_name(ib_type.self, ib_mbr_idx, mbr_name);
// The SPIRV location of interface variable, used to obtain the initial
// MSL location (the location variable) and interface matching
uint32_t ir_location = UINT32_MAX;
bool has_member_loc_decor = has_member_decoration(var_type.self, mbr_idx, DecorationLocation);
bool has_var_loc_decor = has_decoration(var.self, DecorationLocation);
uint32_t orig_vecsize = UINT32_MAX;
if (has_member_loc_decor)
ir_location = get_member_decoration(var_type.self, mbr_idx, DecorationLocation);
else if (has_var_loc_decor)
ir_location = get_accumulated_member_location(var, mbr_idx, meta.strip_array);
else if (is_builtin)
{
if (is_tessellation_shader() && storage == StorageClassInput && inputs_by_builtin.count(builtin))
ir_location = inputs_by_builtin[builtin].location;
else if (capture_output_to_buffer && storage == StorageClassOutput && outputs_by_builtin.count(builtin))
ir_location = outputs_by_builtin[builtin].location;
}
// Once we determine the location of the first member within nested structures,
// from a var of the topmost structure, the remaining flattened members of
// the nested structures will have consecutive location values. At this point,
// we've recursively tunnelled into structs, arrays, and matrices, and are
// down to a single location for each member now.
if (!is_builtin && location != UINT32_MAX)
if (location == UINT32_MAX && ir_location != UINT32_MAX)
location = ir_location + i;
if (storage == StorageClassInput && (has_member_loc_decor || has_var_loc_decor))
{
set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
mark_location_as_used_by_shader(location, *usable_type, storage);
location++;
uint32_t component = 0;
uint32_t orig_mbr_type_id = usable_type->self;
if (has_member_loc_decor)
component = get_member_decoration(var_type.self, mbr_idx, DecorationComponent);
var.basetype = ensure_correct_input_type(var.basetype, location, component, 0, meta.strip_array);
mbr_type_id = ensure_correct_input_type(usable_type->self, location, component, 0, meta.strip_array);
// For members of the composite interface block, we only change the interface block type
// when interface matching happens. In the meantime, we store the original vector size
// and insert a swizzle when loading from metal interface block (see fixup below)
if (mbr_type_id != orig_mbr_type_id)
orig_vecsize = get<SPIRType>(orig_mbr_type_id).vecsize;
if (storage == StorageClassInput && pull_model_inputs.count(var.self))
ib_type.member_types[ib_mbr_idx] = build_msl_interpolant_type(mbr_type_id, is_noperspective);
else
ib_type.member_types[ib_mbr_idx] = mbr_type_id;
}
else if (has_member_decoration(var_type.self, mbr_idx, DecorationLocation))
if ((!is_builtin && location != UINT32_MAX) || (is_builtin && ir_location != UINT32_MAX))
{
location = get_member_decoration(var_type.self, mbr_idx, DecorationLocation) + i;
set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
mark_location_as_used_by_shader(location, *usable_type, storage);
location++;
}
else if (has_decoration(var.self, DecorationLocation))
{
location = get_accumulated_member_location(var, mbr_idx, meta.strip_array) + i;
set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
mark_location_as_used_by_shader(location, *usable_type, storage);
location++;
}
else if (is_builtin && is_tessellation_shader() && storage == StorageClassInput && inputs_by_builtin.count(builtin))
{
location = inputs_by_builtin[builtin].location + i;
set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
mark_location_as_used_by_shader(location, *usable_type, storage);
location++;
}
else if (is_builtin && capture_output_to_buffer && storage == StorageClassOutput && outputs_by_builtin.count(builtin))
{
location = outputs_by_builtin[builtin].location + i;
set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
mark_location_as_used_by_shader(location, *usable_type, storage);
location++;
@@ -3261,6 +3288,7 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass
case StorageClassInput:
entry_func.fixup_hooks_in.push_back([=, &var]() {
string lerp_call;
string swizzle;
if (pull_model_inputs.count(var.self))
{
if (is_centroid)
@@ -3270,7 +3298,9 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass
else
lerp_call = ".interpolate_at_center()";
}
statement(var_chain, " = ", ib_var_ref, ".", mbr_name, lerp_call, ";");
if (orig_vecsize != UINT32_MAX)
swizzle = vector_swizzle(orig_vecsize, 0);
statement(var_chain, " = ", ib_var_ref, ".", mbr_name, lerp_call, swizzle, ";");
});
break;
@@ -3338,6 +3368,55 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor
qual_var_name += ".interpolate_at_center()";
}
// The SPIRV location of interface variable, used to obtain the initial
// MSL location (the location variable) and interface matching
uint32_t ir_location = UINT32_MAX;
bool has_member_loc_decor = has_member_decoration(var_type.self, mbr_idx, DecorationLocation);
bool has_var_loc_decor = has_decoration(var.self, DecorationLocation);
uint32_t orig_vecsize = UINT32_MAX;
if (has_member_loc_decor)
ir_location = get_member_decoration(var_type.self, mbr_idx, DecorationLocation);
else if (has_var_loc_decor)
ir_location = get_accumulated_member_location(var, mbr_idx, meta.strip_array);
else if (is_builtin)
{
if (is_tessellation_shader() && storage == StorageClassInput && inputs_by_builtin.count(builtin))
ir_location = inputs_by_builtin[builtin].location;
else if (capture_output_to_buffer && storage == StorageClassOutput && outputs_by_builtin.count(builtin))
ir_location = outputs_by_builtin[builtin].location;
}
// Once we determine the location of the first member within nested structures,
// from a var of the topmost structure, the remaining flattened members of
// the nested structures will have consecutive location values. At this point,
// we've recursively tunnelled into structs, arrays, and matrices, and are
// down to a single location for each member now.
if (location == UINT32_MAX && ir_location != UINT32_MAX)
location = ir_location;
if (storage == StorageClassInput && (has_member_loc_decor || has_var_loc_decor))
{
uint32_t component = 0;
uint32_t orig_mbr_type_id = mbr_type_id;
if (has_member_loc_decor)
component = get_member_decoration(var_type.self, mbr_idx, DecorationComponent);
mbr_type_id = ensure_correct_input_type(mbr_type_id, location, component, 0, meta.strip_array);
// For members of the composite interface block, we only change the interface block type
// when interface matching happens. In the meantime, we store the original vector size
// and insert a swizzle when loading from metal interface block (see fixup below)
if (mbr_type_id != orig_mbr_type_id)
orig_vecsize = get<SPIRType>(orig_mbr_type_id).vecsize;
if (storage == StorageClassInput && pull_model_inputs.count(var.self))
ib_type.member_types[ib_mbr_idx] = build_msl_interpolant_type(mbr_type_id, is_noperspective);
else
ib_type.member_types[ib_mbr_idx] = mbr_type_id;
}
bool flatten_stage_out = false;
string var_chain = var_chain_qual + "." + to_member_name(var_type, mbr_idx);
if (is_builtin && !meta.strip_array)
@@ -3353,7 +3432,11 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor
{
case StorageClassInput:
entry_func.fixup_hooks_in.push_back([=]() {
statement(var_chain, " = ", qual_var_name, ";");
string swizzle;
// Insert swizzle for widened interface block vector from interface matching
if (orig_vecsize != UINT32_MAX)
swizzle = vector_swizzle(orig_vecsize, 0);
statement(var_chain, " = ", qual_var_name, swizzle, ";");
});
break;
@@ -3369,64 +3452,12 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor
}
}
// Once we determine the location of the first member within nested structures,
// from a var of the topmost structure, the remaining flattened members of
// the nested structures will have consecutive location values. At this point,
// we've recursively tunnelled into structs, arrays, and matrices, and are
// down to a single location for each member now.
if (!is_builtin && location != UINT32_MAX)
if ((!is_builtin && location != UINT32_MAX) || (is_builtin && ir_location != UINT32_MAX))
{
set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
mark_location_as_used_by_shader(location, get<SPIRType>(mbr_type_id), storage);
location += type_to_location_count(get<SPIRType>(mbr_type_id));
}
else if (has_member_decoration(var_type.self, mbr_idx, DecorationLocation))
{
location = get_member_decoration(var_type.self, mbr_idx, DecorationLocation);
uint32_t comp = get_member_decoration(var_type.self, mbr_idx, DecorationComponent);
if (storage == StorageClassInput)
{
mbr_type_id = ensure_correct_input_type(mbr_type_id, location, comp, 0, meta.strip_array);
var_type.member_types[mbr_idx] = mbr_type_id;
if (storage == StorageClassInput && pull_model_inputs.count(var.self))
ib_type.member_types[ib_mbr_idx] = build_msl_interpolant_type(mbr_type_id, is_noperspective);
else
ib_type.member_types[ib_mbr_idx] = mbr_type_id;
}
set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
mark_location_as_used_by_shader(location, get<SPIRType>(mbr_type_id), storage);
location += type_to_location_count(get<SPIRType>(mbr_type_id));
}
else if (has_decoration(var.self, DecorationLocation))
{
location = get_accumulated_member_location(var, mbr_idx, meta.strip_array);
if (storage == StorageClassInput)
{
mbr_type_id = ensure_correct_input_type(mbr_type_id, location, 0, 0, meta.strip_array);
var_type.member_types[mbr_idx] = mbr_type_id;
if (storage == StorageClassInput && pull_model_inputs.count(var.self))
ib_type.member_types[ib_mbr_idx] = build_msl_interpolant_type(mbr_type_id, is_noperspective);
else
ib_type.member_types[ib_mbr_idx] = mbr_type_id;
}
set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
mark_location_as_used_by_shader(location, get<SPIRType>(mbr_type_id), storage);
location += type_to_location_count(get<SPIRType>(mbr_type_id));
}
else if (is_builtin && is_tessellation_shader() && storage == StorageClassInput && inputs_by_builtin.count(builtin))
{
location = inputs_by_builtin[builtin].location;
set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
mark_location_as_used_by_shader(location, get<SPIRType>(mbr_type_id), storage);
location += type_to_location_count(get<SPIRType>(mbr_type_id));
}
else if (is_builtin && capture_output_to_buffer && storage == StorageClassOutput && outputs_by_builtin.count(builtin))
{
location = outputs_by_builtin[builtin].location;
set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
mark_location_as_used_by_shader(location, get<SPIRType>(mbr_type_id), storage);
location += type_to_location_count(get<SPIRType>(mbr_type_id));
}
// Copy the component location, if present.
if (has_member_decoration(var_type.self, mbr_idx, DecorationComponent))
@@ -5550,6 +5581,10 @@ void CompilerMSL::emit_header()
if (suppress_incompatible_pointer_types_discard_qualifiers)
statement("#pragma clang diagnostic ignored \"-Wincompatible-pointer-types-discards-qualifiers\"");
// Disable warning about "sometimes unitialized" when zero-initializing simple threadgroup variables
if (suppress_sometimes_unitialized)
statement("#pragma clang diagnostic ignored \"-Wsometimes-uninitialized\"");
// Disable warning about missing braces for array<T> template to make arrays a value type
if (spv_function_implementations.count(SPVFuncImplUnsafeArray) != 0)
statement("#pragma clang diagnostic ignored \"-Wmissing-braces\"");
@@ -7916,8 +7951,16 @@ void CompilerMSL::emit_specialization_constants_and_structs()
{
SpecializationConstant wg_x, wg_y, wg_z;
ID workgroup_size_id = get_work_group_size_specialization_constants(wg_x, wg_y, wg_z);
bool emitted = false;
if (workgroup_size_id == 0 && is_mesh_shader())
{
auto &execution = get_entry_point();
statement("constant uint3 ", builtin_to_glsl(BuiltInWorkgroupSize, StorageClassWorkgroup),
" [[maybe_unused]] = ", "uint3(", execution.workgroup_size.x, ", ", execution.workgroup_size.y, ", ",
execution.workgroup_size.z, ");");
statement("");
}
bool emitted = false;
unordered_set<uint32_t> declared_structs;
unordered_set<uint32_t> aligned_structs;
@@ -8022,14 +8065,18 @@ void CompilerMSL::emit_specialization_constants_and_structs()
else if (has_decoration(c.self, DecorationSpecId))
{
// Fallback to macro overrides.
uint32_t constant_id = get_decoration(c.self, DecorationSpecId);
c.specialization_constant_macro_name =
constant_value_macro_name(get_decoration(c.self, DecorationSpecId));
constant_value_macro_name(constant_id);
statement("#ifndef ", c.specialization_constant_macro_name);
statement("#define ", c.specialization_constant_macro_name, " ", constant_expression(c));
statement("#endif");
statement("constant ", sc_type_name, " ", sc_name, " = ", c.specialization_constant_macro_name,
";");
// Record the usage of macro
constant_macro_ids.insert(constant_id);
}
else
{
@@ -9978,11 +10025,11 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
auto &type = get<SPIRType>(ops[0]);
auto &input_type = expression_type(ops[2]);
if (opcode != OpBitcast || type.pointer || input_type.pointer)
if (opcode != OpBitcast || is_pointer(type) || is_pointer(input_type))
{
string op;
if (type.vecsize == 1 && input_type.vecsize == 1)
if ((type.vecsize == 1 || is_pointer(type)) && (input_type.vecsize == 1 || is_pointer(input_type)))
op = join("reinterpret_cast<", type_to_glsl(type), ">(", to_unpacked_expression(ops[2]), ")");
else if (input_type.vecsize == 2)
op = join("reinterpret_cast<", type_to_glsl(type), ">(as_type<ulong>(", to_unpacked_expression(ops[2]), "))");
@@ -12317,7 +12364,7 @@ string CompilerMSL::to_func_call_arg(const SPIRFunction::Parameter &arg, uint32_
}
// Dereference pointer variables where needed.
// FIXME: This dereference is actually backwards. We should really just support passing pointer variables between functions.
else if (should_dereference(id))
else if (should_dereference_caller_param(id))
arg_str += dereference_expression(type, CompilerGLSL::to_func_call_arg(arg, id));
else
arg_str += CompilerGLSL::to_func_call_arg(arg, id);
@@ -13405,6 +13452,9 @@ bool CompilerMSL::uses_explicit_early_fragment_test()
string CompilerMSL::get_argument_address_space(const SPIRVariable &argument)
{
const auto &type = get<SPIRType>(argument.basetype);
// BDA is always passed around by value. There is no storage class for the argument itself.
if (is_physical_pointer(type))
return "";
return get_type_address_space(type, argument.self, true);
}
@@ -15140,17 +15190,21 @@ const char *CompilerMSL::descriptor_address_space(uint32_t id, StorageClass stor
string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
{
auto &var = get<SPIRVariable>(arg.id);
auto &type = get_variable_data_type(var);
auto &var_type = get<SPIRType>(arg.type);
StorageClass type_storage = var_type.storage;
// Physical pointer types are passed by pointer, not reference.
auto &data_type = get_variable_data_type(var);
bool passed_by_value = is_physical_pointer(var_type);
auto &type = passed_by_value ? var_type : data_type;
// If we need to modify the name of the variable, make sure we use the original variable.
// Our alias is just a shadow variable.
uint32_t name_id = var.self;
if (arg.alias_global_variable && var.basevariable)
name_id = var.basevariable;
bool constref = !arg.alias_global_variable && is_pointer(var_type) && arg.write_count == 0;
bool constref = !arg.alias_global_variable && !passed_by_value && is_pointer(var_type) && arg.write_count == 0;
// Framebuffer fetch is plain value, const looks out of place, but it is not wrong.
if (type_is_msl_framebuffer_fetch(type))
constref = false;
@@ -15181,8 +15235,7 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
if (var.basevariable && (var.basevariable == stage_in_ptr_var_id || var.basevariable == stage_out_ptr_var_id))
decl = join(cv_qualifier, type_to_glsl(type, arg.id));
else if (builtin && builtin_type != spv::BuiltInPrimitiveTriangleIndicesEXT &&
builtin_type != spv::BuiltInPrimitiveLineIndicesEXT && builtin_type != spv::BuiltInPrimitivePointIndicesEXT)
else if (builtin && !is_mesh_shader())
{
// Only use templated array for Clip/Cull distance when feasible.
// In other scenarios, we need need to override array length for tess levels (if used as outputs),
@@ -15234,7 +15287,7 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
else
{
// The type is a pointer type we need to emit cv_qualifier late.
if (is_pointer(type))
if (is_pointer(data_type))
{
decl = type_to_glsl(type, arg.id);
if (*cv_qualifier != '\0')
@@ -15246,8 +15299,8 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
}
}
if (!builtin && !is_pointer(var_type) &&
(type_storage == StorageClassFunction || type_storage == StorageClassGeneric))
if (passed_by_value || (!builtin && !is_pointer(var_type) &&
(type_storage == StorageClassFunction || type_storage == StorageClassGeneric)))
{
// If the argument is a pure value and not an opaque type, we will pass by value.
if (msl_options.force_native_arrays && is_array(type))
@@ -15348,7 +15401,7 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
// for the reference has to go before the '&', but after the '*'.
if (!address_space.empty())
{
if (is_pointer(type))
if (is_pointer(data_type))
{
if (*cv_qualifier == '\0')
decl += ' ';
@@ -15357,6 +15410,7 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
else
decl = join(address_space, " ", decl);
}
decl += "&";
decl += " ";
decl += to_restrict(name_id, true);
@@ -17608,6 +17662,23 @@ void CompilerMSL::analyze_sampled_image_usage()
}
}
void CompilerMSL::analyze_workgroup_variables()
{
ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
// If workgroup variables have initializer, it can only be ConstantNull (zero init)
if (var.storage == StorageClassWorkgroup && var.initializer)
{
needs_workgroup_zero_init = true;
// MSL compiler does not like the routine to initialize simple threadgroup variables,
// falsely claiming it is "sometimes uninitialized". Suppress it.
auto &type = get_variable_data_type(var);
if (type.array.empty() && type.member_types.empty())
suppress_sometimes_unitialized = true;
}
});
}
bool CompilerMSL::SampledImageScanner::handle(spv::Op opcode, const uint32_t *args, uint32_t length)
{
switch (opcode)
@@ -19007,7 +19078,7 @@ void CompilerMSL::analyze_argument_buffers()
set_extended_member_decoration(buffer_type.self, member_index, SPIRVCrossDecorationOverlappingBinding);
member_index++;
}
if (msl_options.replace_recursive_inputs && type_contains_recursion(buffer_type))
{
recursive_inputs.insert(type_id);
@@ -19167,6 +19238,11 @@ const char *CompilerMSL::get_combined_sampler_suffix() const
return sampler_name_suffix.c_str();
}
bool CompilerMSL::specialization_constant_is_macro(uint32_t const_id) const
{
return constant_macro_ids.find(const_id) != constant_macro_ids.end();
}
void CompilerMSL::emit_block_hints(const SPIRBlock &)
{
}
@@ -19389,6 +19465,70 @@ void CompilerMSL::emit_mesh_tasks(SPIRBlock &block)
statement("return;");
}
void CompilerMSL::emit_workgroup_initialization(const SPIRVariable &var)
{
auto &type = get_variable_data_type(var);
begin_scope();
if (type.array.empty() && type.member_types.empty())
{
// For simple shared variables, we just initialize it in thread 0 of the block
// We use short to represent bool for threadgroup variable to workaround compiler bug,
// so we do a temporary fixup here. Alas. (see the type_to_glsl method)
bool is_boolean = type.basetype == SPIRType::Boolean;
if (is_boolean)
type.basetype = SPIRType::Short;
statement("if (gl_LocalInvocationIndex == 0)");
begin_scope();
statement(to_name(var.self), " = ", to_initializer_expression(var), ";");
end_scope();
if (is_boolean)
type.basetype = SPIRType::Boolean;
}
else
{
// Otherwise, we use a loop to cooperatively initialize the memory within the group
// First, we define a few variable names;
string var_name = to_name(var.self);
string var_ptr_name = join(var_name, "_ptr");
string var_size_name = join(var_name, "_sz");
string var_pos_name = join(var_name, "_pos");
string var_stride_name = join(var_name, "_stride");
string var_ptr2_name = join(var_name, "_ptr2");
statement("threadgroup uint *", var_ptr_name, " = (threadgroup uint *)&", var_name, ";");
statement("uint ", var_size_name, " = ", "sizeof(", var_name, ");");
statement("uint ", var_pos_name, " = gl_LocalInvocationIndex;");
statement("uint ", var_stride_name, " = gl_WorkGroupSize.x * gl_WorkGroupSize.y * gl_WorkGroupSize.z;");
statement("while (sizeof(uint) * ", var_pos_name, " < ", var_size_name, ")");
begin_scope();
statement(var_ptr_name, "[", var_pos_name, "] = 0u;");
statement(var_pos_name, " += ", var_stride_name, ";");
end_scope();
statement("if (gl_LocalInvocationIndex == 0)");
begin_scope();
statement(var_pos_name, " = (", var_size_name, " / sizeof(uint)) * sizeof(uint);");
statement("threadgroup uchar *", var_ptr2_name, " = (threadgroup uchar *)&", var_name, ";");
statement("while (", var_pos_name, " < ", var_size_name, ")");
begin_scope();
statement(var_ptr2_name, "[", var_pos_name, "] = '\\0';");
statement(var_pos_name, "++;");
end_scope();
end_scope();
}
statement("threadgroup_barrier(mem_flags::mem_threadgroup);");
end_scope();
}
string CompilerMSL::additional_fixed_sample_mask_str() const
{
char print_buffer[32];

10
3rdparty/spirv-cross/spirv_msl.hpp vendored
View File

@@ -760,6 +760,11 @@ public:
void set_combined_sampler_suffix(const char *suffix);
const char *get_combined_sampler_suffix() const;
// Information about specialization constants that are translated into MSL macros
// instead of using function constant
// These must only be called after a successful call to CompilerMSL::compile().
bool specialization_constant_is_macro(uint32_t constant_id) const;
protected:
// An enum of SPIR-V functions that are implemented in additional
// source code that is added to the shader if necessary.
@@ -876,6 +881,7 @@ protected:
void emit_mesh_entry_point();
void emit_mesh_outputs();
void emit_mesh_tasks(SPIRBlock &block) override;
void emit_workgroup_initialization(const SPIRVariable &var) override;
// Allow Metal to use the array<T> template to make arrays a value type
std::string type_to_array_glsl(const SPIRType &type, uint32_t variable_id) override;
@@ -1137,6 +1143,7 @@ protected:
void emit_store_statement(uint32_t lhs_expression, uint32_t rhs_expression) override;
void analyze_sampled_image_usage();
void analyze_workgroup_variables();
bool access_chain_needs_stage_io_builtin_translation(uint32_t base) override;
bool prepare_access_chain_for_scalar_access(std::string &expr, const SPIRType &type, spv::StorageClass storage,
@@ -1171,6 +1178,7 @@ protected:
std::set<std::string> pragma_lines;
std::set<std::string> typedef_lines;
SmallVector<uint32_t> vars_needing_early_declaration;
std::unordered_set<uint32_t> constant_macro_ids;
std::unordered_map<StageSetBinding, std::pair<MSLResourceBinding, bool>, InternalHasher> resource_bindings;
std::unordered_map<StageSetBinding, uint32_t, InternalHasher> resource_arg_buff_idx_to_binding_number;
@@ -1218,6 +1226,7 @@ protected:
bool needs_subgroup_size = false;
bool needs_sample_id = false;
bool needs_helper_invocation = false;
bool needs_workgroup_zero_init = false;
bool writes_to_depth = false;
std::string qual_pos_var_name;
std::string stage_in_var_name = "in";
@@ -1280,6 +1289,7 @@ protected:
bool suppress_missing_prototypes = false;
bool suppress_incompatible_pointer_types_discard_qualifiers = false;
bool suppress_sometimes_unitialized = false;
void add_spv_func_and_recompile(SPVFuncImpl spv_func);

1
3rdparty/spirv-cross/spirv_parser.cpp vendored
View File

@@ -305,6 +305,7 @@ void Parser::parse(const Instruction &instruction)
}
case OpExtInst:
case OpExtInstWithForwardRefsKHR:
{
// The SPIR-V debug information extended instructions might come at global scope.
if (current_block)