AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Various nit fixes and improvements from review.

Browse Source
This commit is contained in:
Hans-Kristian Arntzen 2023-12-07 12:18:35 +01:00
parent e3d1b9afbb
commit 833c5936b0
8 changed files with 189 additions and 238 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
spirv_common.hpp
View File

@ -549,7 +549,7 @@ struct SPIRType : IVariant
}; };
spv::Op op = spv::Op::OpNop; spv::Op op = spv::Op::OpNop;
SPIRType(spv::Op op_) : op(op_) {} explicit SPIRType(spv::Op op_) : op(op_) {}
enum BaseType enum BaseType
{ {

67
spirv_cross.cpp
View File

@ -627,20 +627,22 @@ bool Compiler::is_matrix(const SPIRType &type) const
bool Compiler::is_array(const SPIRType &type) const bool Compiler::is_array(const SPIRType &type) const
{ {
return (type.op == OpTypeArray || type.op == OpTypeRuntimeArray); return type.op == OpTypeArray || type.op == OpTypeRuntimeArray;
} }
bool Compiler::is_pointer(const SPIRType &type) const bool Compiler::is_pointer(const SPIRType &type) const
{ {
return type.op == OpTypePointer; return type.op == OpTypePointer && type.basetype != SPIRType::Unknown; // Ignore function pointers.
}
bool Compiler::is_physical_pointer(const SPIRType &type) const
{
return type.op == OpTypePointer && type.storage == StorageClassPhysicalStorageBuffer;
} }
bool Compiler::is_runtime_size_array(const SPIRType &type) bool Compiler::is_runtime_size_array(const SPIRType &type)
{ {
if (type.array.empty()) return type.op == OpTypeRuntimeArray;
return false;
assert(type.array.size() == type.array_size_literal.size());
return type.array_size_literal.back() && type.array.back() == 0;
} }
ShaderResources Compiler::get_shader_resources() const ShaderResources Compiler::get_shader_resources() const
@ -2743,8 +2745,8 @@ void Compiler::CombinedImageSamplerHandler::register_combined_image_sampler(SPIR
auto ptr_type_id = id + 1; auto ptr_type_id = id + 1;
auto combined_id = id + 2; auto combined_id = id + 2;
auto &base = compiler.expression_type(image_id); auto &base = compiler.expression_type(image_id);
auto &type = compiler.set<SPIRType>(type_id, spv::Op::OpTypeSampledImage); auto &type = compiler.set<SPIRType>(type_id, OpTypeSampledImage);
auto &ptr_type = compiler.set<SPIRType>(ptr_type_id, spv::Op::OpTypePointer); auto &ptr_type = compiler.set<SPIRType>(ptr_type_id, OpTypePointer);
type = base; type = base;
type.self = type_id; type.self = type_id;
@ -3003,7 +3005,7 @@ bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *ar
{ {
// Have to invent the sampled image type. // Have to invent the sampled image type.
sampled_type = compiler.ir.increase_bound_by(1); sampled_type = compiler.ir.increase_bound_by(1);
auto &type = compiler.set<SPIRType>(sampled_type, spv::Op::OpTypeSampledImage); auto &type = compiler.set<SPIRType>(sampled_type, OpTypeSampledImage);
type = compiler.expression_type(args[2]); type = compiler.expression_type(args[2]);
type.self = sampled_type; type.self = sampled_type;
type.basetype = SPIRType::SampledImage; type.basetype = SPIRType::SampledImage;
@ -3022,7 +3024,7 @@ bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *ar
// Make a new type, pointer to OpTypeSampledImage, so we can make a variable of this type. // Make a new type, pointer to OpTypeSampledImage, so we can make a variable of this type.
// We will probably have this type lying around, but it doesn't hurt to make duplicates for internal purposes. // We will probably have this type lying around, but it doesn't hurt to make duplicates for internal purposes.
auto &type = compiler.set<SPIRType>(type_id, spv::Op::OpTypePointer); auto &type = compiler.set<SPIRType>(type_id, OpTypePointer);
auto &base = compiler.get<SPIRType>(sampled_type); auto &base = compiler.get<SPIRType>(sampled_type);
type = base; type = base;
type.pointer = true; type.pointer = true;
@ -3068,10 +3070,10 @@ VariableID Compiler::build_dummy_sampler_for_combined_images()
auto ptr_type_id = offset + 1; auto ptr_type_id = offset + 1;
auto var_id = offset + 2; auto var_id = offset + 2;
auto &sampler = set<SPIRType>(type_id, spv::Op::OpTypeSampler); auto &sampler = set<SPIRType>(type_id, OpTypeSampler);
sampler.basetype = SPIRType::Sampler; sampler.basetype = SPIRType::Sampler;
auto &ptr_sampler = set<SPIRType>(ptr_type_id, spv::Op::OpTypePointer); auto &ptr_sampler = set<SPIRType>(ptr_type_id, OpTypePointer);
ptr_sampler = sampler; ptr_sampler = sampler;
ptr_sampler.self = type_id; ptr_sampler.self = type_id;
ptr_sampler.storage = StorageClassUniformConstant; ptr_sampler.storage = StorageClassUniformConstant;
@ -5501,7 +5503,7 @@ bool Compiler::type_contains_recursion(const SPIRType &type)
bool Compiler::type_is_array_of_pointers(const SPIRType &type) const bool Compiler::type_is_array_of_pointers(const SPIRType &type) const
{ {
if (!type_is_top_level_array(type)) if (!is_array(type))
return false; return false;
// BDA types must have parent type hierarchy. // BDA types must have parent type hierarchy.
@ -5510,45 +5512,10 @@ bool Compiler::type_is_array_of_pointers(const SPIRType &type) const
// Punch through all array layers. // Punch through all array layers.
auto *parent = &get<SPIRType>(type.parent_type); auto *parent = &get<SPIRType>(type.parent_type);
while (type_is_top_level_array(*parent)) while (is_array(*parent))
parent = &get<SPIRType>(parent->parent_type); parent = &get<SPIRType>(parent->parent_type);
return type_is_top_level_pointer(*parent); return is_pointer(*parent);
}
bool Compiler::type_is_top_level_pointer(const SPIRType &type) const
{
if (!type.pointer)
return false;
// Function pointers, should not be hit by valid SPIR-V.
// Parent type will be SPIRFunction instead.
if (type.basetype == SPIRType::Unknown)
return false;
// Some types are synthesized in-place without complete type hierarchy and might not have parent types,
// but these types are never array-of-pointer or any complicated BDA type, infer reasonable defaults.
if (type.parent_type)
return type.pointer_depth > get<SPIRType>(type.parent_type).pointer_depth;
else
return true;
}
bool Compiler::type_is_top_level_physical_pointer(const SPIRType &type) const
{
return type_is_top_level_pointer(type) && type.storage == StorageClassPhysicalStorageBuffer;
}
bool Compiler::type_is_top_level_array(const SPIRType &type) const
{
if (type.array.empty())
return false;
// If we have pointer and array, we infer pointer-to-array as it's the only meaningful thing outside BDA.
if (type.parent_type)
return type.array.size() > get<SPIRType>(type.parent_type).array.size();
else
return !type.pointer;
} }
bool Compiler::flush_phi_required(BlockID from, BlockID to) const bool Compiler::flush_phi_required(BlockID from, BlockID to) const

4
spirv_cross.hpp
View File

@ -684,6 +684,7 @@ protected:
bool is_matrix(const SPIRType &type) const; bool is_matrix(const SPIRType &type) const;
bool is_array(const SPIRType &type) const; bool is_array(const SPIRType &type) const;
bool is_pointer(const SPIRType &type) const; bool is_pointer(const SPIRType &type) const;
bool is_physical_pointer(const SPIRType &type) const;
static bool is_runtime_size_array(const SPIRType &type); static bool is_runtime_size_array(const SPIRType &type);
uint32_t expression_type_id(uint32_t id) const; uint32_t expression_type_id(uint32_t id) const;
const SPIRType &expression_type(uint32_t id) const; const SPIRType &expression_type(uint32_t id) const;
@ -1149,9 +1150,6 @@ protected:
bool check_internal_recursion(const SPIRType &type, std::unordered_set<uint32_t> &checked_ids); bool check_internal_recursion(const SPIRType &type, std::unordered_set<uint32_t> &checked_ids);
bool type_contains_recursion(const SPIRType &type); bool type_contains_recursion(const SPIRType &type);
bool type_is_array_of_pointers(const SPIRType &type) const; bool type_is_array_of_pointers(const SPIRType &type) const;
bool type_is_top_level_physical_pointer(const SPIRType &type) const;
bool type_is_top_level_pointer(const SPIRType &type) const;
bool type_is_top_level_array(const SPIRType &type) const;
bool type_is_block_like(const SPIRType &type) const; bool type_is_block_like(const SPIRType &type) const;
bool type_is_top_level_block(const SPIRType &type) const; bool type_is_top_level_block(const SPIRType &type) const;
bool type_is_opaque_value(const SPIRType &type) const; bool type_is_opaque_value(const SPIRType &type) const;

70
spirv_glsl.cpp
View File

@ -223,7 +223,7 @@ static const char *to_pls_layout(PlsFormat format)
} }
} }
static std::tuple<spv::Op, SPIRType::BaseType> pls_format_to_basetype(PlsFormat format) static std::pair<spv::Op, SPIRType::BaseType> pls_format_to_basetype(PlsFormat format)
{ {
switch (format) switch (format)
{ {
@ -234,17 +234,17 @@ static std::tuple<spv::Op, SPIRType::BaseType> pls_format_to_basetype(PlsFormat
case PlsRGB10A2: case PlsRGB10A2:
case PlsRGBA8: case PlsRGBA8:
case PlsRG16: case PlsRG16:
return std::make_tuple(spv::OpTypeFloat, SPIRType::Float); return std::make_pair(spv::OpTypeFloat, SPIRType::Float);
case PlsRGBA8I: case PlsRGBA8I:
case PlsRG16I: case PlsRG16I:
return std::make_tuple(spv::OpTypeInt, SPIRType::Int); return std::make_pair(spv::OpTypeInt, SPIRType::Int);
case PlsRGB10A2UI: case PlsRGB10A2UI:
case PlsRGBA8UI: case PlsRGBA8UI:
case PlsRG16UI: case PlsRG16UI:
case PlsR32UI: case PlsR32UI:
return std::make_tuple(spv::OpTypeInt, SPIRType::UInt); return std::make_pair(spv::OpTypeInt, SPIRType::UInt);
} }
} }
@ -1529,7 +1529,7 @@ uint32_t CompilerGLSL::type_to_packed_alignment(const SPIRType &type, const Bits
{ {
// If using PhysicalStorageBufferEXT storage class, this is a pointer, // If using PhysicalStorageBufferEXT storage class, this is a pointer,
// and is 64-bit. // and is 64-bit.
if (type_is_top_level_physical_pointer(type)) if (is_physical_pointer(type))
{ {
if (!type.pointer) if (!type.pointer)
SPIRV_CROSS_THROW("Types in PhysicalStorageBufferEXT must be pointers."); SPIRV_CROSS_THROW("Types in PhysicalStorageBufferEXT must be pointers.");
@ -1544,7 +1544,7 @@ uint32_t CompilerGLSL::type_to_packed_alignment(const SPIRType &type, const Bits
else else
SPIRV_CROSS_THROW("AddressingModelPhysicalStorageBuffer64EXT must be used for PhysicalStorageBufferEXT."); SPIRV_CROSS_THROW("AddressingModelPhysicalStorageBuffer64EXT must be used for PhysicalStorageBufferEXT.");
} }
else if (type_is_top_level_array(type)) else if (is_array(type))
{ {
uint32_t minimum_alignment = 1; uint32_t minimum_alignment = 1;
if (packing_is_vec4_padded(packing)) if (packing_is_vec4_padded(packing))
@ -1652,7 +1652,7 @@ uint32_t CompilerGLSL::type_to_packed_size(const SPIRType &type, const Bitset &f
{ {
// If using PhysicalStorageBufferEXT storage class, this is a pointer, // If using PhysicalStorageBufferEXT storage class, this is a pointer,
// and is 64-bit. // and is 64-bit.
if (type_is_top_level_physical_pointer(type)) if (is_physical_pointer(type))
{ {
if (!type.pointer) if (!type.pointer)
SPIRV_CROSS_THROW("Types in PhysicalStorageBufferEXT must be pointers."); SPIRV_CROSS_THROW("Types in PhysicalStorageBufferEXT must be pointers.");
@ -1662,7 +1662,7 @@ uint32_t CompilerGLSL::type_to_packed_size(const SPIRType &type, const Bitset &f
else else
SPIRV_CROSS_THROW("AddressingModelPhysicalStorageBuffer64EXT must be used for PhysicalStorageBufferEXT."); SPIRV_CROSS_THROW("AddressingModelPhysicalStorageBuffer64EXT must be used for PhysicalStorageBufferEXT.");
} }
else if (type_is_top_level_array(type)) else if (is_array(type))
{ {
uint32_t packed_size = to_array_size_literal(type) * type_to_packed_array_stride(type, flags, packing); uint32_t packed_size = to_array_size_literal(type) * type_to_packed_array_stride(type, flags, packing);
@ -1840,7 +1840,7 @@ bool CompilerGLSL::buffer_is_packing_standard(const SPIRType &type, BufferPackin
} }
// Verify array stride rules. // Verify array stride rules.
if (type_is_top_level_array(memb_type) && if (is_array(memb_type) &&
type_to_packed_array_stride(memb_type, member_flags, packing) != type_to_packed_array_stride(memb_type, member_flags, packing) !=
type_struct_member_array_stride(type, i)) type_struct_member_array_stride(type, i))
{ {
@ -2489,7 +2489,7 @@ void CompilerGLSL::emit_buffer_block_flattened(const SPIRVariable &var)
SPIRType::BaseType basic_type; SPIRType::BaseType basic_type;
if (get_common_basic_type(type, basic_type)) if (get_common_basic_type(type, basic_type))
{ {
SPIRType tmp { spv::Op::OpTypeVector }; SPIRType tmp { OpTypeVector };
tmp.basetype = basic_type; tmp.basetype = basic_type;
tmp.vecsize = 4; tmp.vecsize = 4;
if (basic_type != SPIRType::Float && basic_type != SPIRType::Int && basic_type != SPIRType::UInt) if (basic_type != SPIRType::Float && basic_type != SPIRType::Int && basic_type != SPIRType::UInt)
@ -3926,7 +3926,7 @@ void CompilerGLSL::emit_output_variable_initializer(const SPIRVariable &var)
auto &member_type = get<SPIRType>(member_type_id); auto &member_type = get<SPIRType>(member_type_id);
auto array_type = member_type; auto array_type = member_type;
array_type.parent_type = member_type_id; array_type.parent_type = member_type_id;
array_type.op = spv::Op::OpTypeArray; array_type.op = OpTypeArray;
array_type.array.push_back(array_size); array_type.array.push_back(array_size);
array_type.array_size_literal.push_back(true); array_type.array_size_literal.push_back(true);
@ -3950,7 +3950,7 @@ void CompilerGLSL::emit_output_variable_initializer(const SPIRVariable &var)
if (is_control_point) if (is_control_point)
{ {
uint32_t ids = ir.increase_bound_by(3); uint32_t ids = ir.increase_bound_by(3);
auto& uint_type = set<SPIRType>(ids, spv::Op::OpTypeInt); auto &uint_type = set<SPIRType>(ids, OpTypeInt);
uint_type.basetype = SPIRType::UInt; uint_type.basetype = SPIRType::UInt;
uint_type.width = 32; uint_type.width = 32;
set<SPIRExpression>(ids + 1, builtin_to_glsl(BuiltInInvocationId, StorageClassInput), ids, true); set<SPIRExpression>(ids + 1, builtin_to_glsl(BuiltInInvocationId, StorageClassInput), ids, true);
@ -5148,7 +5148,7 @@ string CompilerGLSL::to_rerolled_array_expression(const SPIRType &parent_type,
type.basetype == SPIRType::Boolean && type.basetype == SPIRType::Boolean &&
backend.boolean_in_struct_remapped_type != SPIRType::Boolean; backend.boolean_in_struct_remapped_type != SPIRType::Boolean;
SPIRType tmp_type { spv::Op::OpNop }; SPIRType tmp_type { OpNop };
if (remapped_boolean) if (remapped_boolean)
{ {
tmp_type = get<SPIRType>(type.parent_type); tmp_type = get<SPIRType>(type.parent_type);
@ -5169,7 +5169,7 @@ string CompilerGLSL::to_rerolled_array_expression(const SPIRType &parent_type,
for (uint32_t i = 0; i < size; i++) for (uint32_t i = 0; i < size; i++)
{ {
auto subexpr = join(base_expr, "[", convert_to_string(i), "]"); auto subexpr = join(base_expr, "[", convert_to_string(i), "]");
if (!type_is_top_level_array(parent)) if (!is_array(parent))
{ {
if (remapped_boolean) if (remapped_boolean)
subexpr = join(type_to_glsl(tmp_type), "(", subexpr, ")"); subexpr = join(type_to_glsl(tmp_type), "(", subexpr, ")");
@ -5195,7 +5195,7 @@ string CompilerGLSL::to_composite_constructor_expression(const SPIRType &parent_
type.basetype == SPIRType::Boolean && type.basetype == SPIRType::Boolean &&
backend.boolean_in_struct_remapped_type != SPIRType::Boolean; backend.boolean_in_struct_remapped_type != SPIRType::Boolean;
if (type_is_top_level_array(type)) if (is_array(type))
{ {
reroll_array = !backend.array_is_value_type || reroll_array = !backend.array_is_value_type ||
(block_like_type && !backend.array_is_value_type_in_buffer_blocks); (block_like_type && !backend.array_is_value_type_in_buffer_blocks);
@ -5748,7 +5748,7 @@ string CompilerGLSL::constant_expression(const SPIRConstant &c,
{ {
auto &type = get<SPIRType>(c.constant_type); auto &type = get<SPIRType>(c.constant_type);
if (type_is_top_level_pointer(type)) if (is_pointer(type))
{ {
return backend.null_pointer_literal; return backend.null_pointer_literal;
} }
@ -5763,21 +5763,21 @@ string CompilerGLSL::constant_expression(const SPIRConstant &c,
// with Offset = 0, using no ArrayStride on the enclosed array type. // with Offset = 0, using no ArrayStride on the enclosed array type.
// A particular CTS test hits this scenario. // A particular CTS test hits this scenario.
bool array_type_decays = inside_block_like_struct_scope && bool array_type_decays = inside_block_like_struct_scope &&
type_is_top_level_array(type) && is_array(type) &&
!backend.array_is_value_type_in_buffer_blocks; !backend.array_is_value_type_in_buffer_blocks;
// Allow Metal to use the array<T> template to make arrays a value type // Allow Metal to use the array<T> template to make arrays a value type
bool needs_trailing_tracket = false; bool needs_trailing_tracket = false;
if (backend.use_initializer_list && backend.use_typed_initializer_list && type.basetype == SPIRType::Struct && if (backend.use_initializer_list && backend.use_typed_initializer_list && type.basetype == SPIRType::Struct &&
!type_is_top_level_array(type)) !is_array(type))
{ {
res = type_to_glsl_constructor(type) + "{ "; res = type_to_glsl_constructor(type) + "{ ";
} }
else if (backend.use_initializer_list && backend.use_typed_initializer_list && backend.array_is_value_type && else if (backend.use_initializer_list && backend.use_typed_initializer_list && backend.array_is_value_type &&
type_is_top_level_array(type) && !array_type_decays) is_array(type) && !array_type_decays)
{ {
const auto *p_type = &type; const auto *p_type = &type;
SPIRType tmp_type { spv::Op::OpNop }; SPIRType tmp_type { OpNop };
if (inside_struct_scope && if (inside_struct_scope &&
backend.boolean_in_struct_remapped_type != SPIRType::Boolean && backend.boolean_in_struct_remapped_type != SPIRType::Boolean &&
@ -5818,7 +5818,7 @@ string CompilerGLSL::constant_expression(const SPIRConstant &c,
res += to_name(elem); res += to_name(elem);
else else
{ {
if (!type_is_top_level_array(type) && type.basetype == SPIRType::Struct) if (!is_array(type) && type.basetype == SPIRType::Struct)
{ {
// When we get down to emitting struct members, override the block-like information. // When we get down to emitting struct members, override the block-like information.
// For constants, we can freely mix and match block-like state. // For constants, we can freely mix and match block-like state.
@ -5916,7 +5916,7 @@ string CompilerGLSL::convert_half_to_string(const SPIRConstant &c, uint32_t col,
// of complicated workarounds, just value-cast to the half type always. // of complicated workarounds, just value-cast to the half type always.
if (std::isnan(float_value) || std::isinf(float_value)) if (std::isnan(float_value) || std::isinf(float_value))
{ {
SPIRType type { spv::Op::OpTypeFloat }; SPIRType type { OpTypeFloat };
type.basetype = SPIRType::Half; type.basetype = SPIRType::Half;
type.vecsize = 1; type.vecsize = 1;
type.columns = 1; type.columns = 1;
@ -5932,7 +5932,7 @@ string CompilerGLSL::convert_half_to_string(const SPIRConstant &c, uint32_t col,
} }
else else
{ {
SPIRType type { spv::Op::OpTypeFloat }; SPIRType type { OpTypeFloat };
type.basetype = SPIRType::Half; type.basetype = SPIRType::Half;
type.vecsize = 1; type.vecsize = 1;
type.columns = 1; type.columns = 1;
@ -5952,8 +5952,8 @@ string CompilerGLSL::convert_float_to_string(const SPIRConstant &c, uint32_t col
// Use special representation. // Use special representation.
if (!is_legacy()) if (!is_legacy())
{ {
SPIRType out_type { spv::Op::OpTypeFloat }; SPIRType out_type { OpTypeFloat };
SPIRType in_type { spv::Op::OpTypeInt }; SPIRType in_type { OpTypeInt };
out_type.basetype = SPIRType::Float; out_type.basetype = SPIRType::Float;
in_type.basetype = SPIRType::UInt; in_type.basetype = SPIRType::UInt;
out_type.vecsize = 1; out_type.vecsize = 1;
@ -6022,8 +6022,8 @@ std::string CompilerGLSL::convert_double_to_string(const SPIRConstant &c, uint32
// Use special representation. // Use special representation.
if (!is_legacy()) if (!is_legacy())
{ {
SPIRType out_type { spv::Op::OpTypeFloat }; SPIRType out_type { OpTypeFloat };
SPIRType in_type { spv::Op::OpTypeInt }; SPIRType in_type { OpTypeInt };
out_type.basetype = SPIRType::Double; out_type.basetype = SPIRType::Double;
in_type.basetype = SPIRType::UInt64; in_type.basetype = SPIRType::UInt64;
out_type.vecsize = 1; out_type.vecsize = 1;
@ -6731,7 +6731,7 @@ SPIRType CompilerGLSL::binary_op_bitcast_helper(string &cast_op0, string &cast_o
// Create a fake type so we can bitcast to it. // Create a fake type so we can bitcast to it.
// We only deal with regular arithmetic types here like int, uints and so on. // We only deal with regular arithmetic types here like int, uints and so on.
SPIRType expected_type = type0.op; SPIRType expected_type{type0.op};
expected_type.basetype = input_type; expected_type.basetype = input_type;
expected_type.vecsize = type0.vecsize; expected_type.vecsize = type0.vecsize;
expected_type.columns = type0.columns; expected_type.columns = type0.columns;
@ -7086,7 +7086,7 @@ void CompilerGLSL::emit_bitfield_insert_op(uint32_t result_type, uint32_t result
auto op3_expr = to_unpacked_expression(op3); auto op3_expr = to_unpacked_expression(op3);
assert(offset_count_type == SPIRType::UInt || offset_count_type == SPIRType::Int); assert(offset_count_type == SPIRType::UInt || offset_count_type == SPIRType::Int);
SPIRType target_type { spv::Op::OpTypeInt }; SPIRType target_type { OpTypeInt };
target_type.width = 32; target_type.width = 32;
target_type.vecsize = 1; target_type.vecsize = 1;
target_type.basetype = offset_count_type; target_type.basetype = offset_count_type;
@ -7878,7 +7878,7 @@ bool CompilerGLSL::expression_is_constant_null(uint32_t id) const
bool CompilerGLSL::expression_is_non_value_type_array(uint32_t ptr) bool CompilerGLSL::expression_is_non_value_type_array(uint32_t ptr)
{ {
auto &type = expression_type(ptr); auto &type = expression_type(ptr);
if (!type_is_top_level_array(get_pointee_type(type))) if (!is_array(get_pointee_type(type)))
return false; return false;
if (!backend.array_is_value_type) if (!backend.array_is_value_type)
@ -10019,7 +10019,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
index &= 0x7fffffffu; index &= 0x7fffffffu;
} }
bool ptr_chain_array_entry = ptr_chain && i == 0 && type_is_top_level_array(*type); bool ptr_chain_array_entry = ptr_chain && i == 0 && is_array(*type);
if (ptr_chain_array_entry) if (ptr_chain_array_entry)
{ {
@ -15367,8 +15367,8 @@ string CompilerGLSL::pls_decl(const PlsRemap &var)
auto op_and_basetype = pls_format_to_basetype(var.format); auto op_and_basetype = pls_format_to_basetype(var.format);
SPIRType type { std::get<0>(op_and_basetype) }; SPIRType type { op_and_basetype.first };
type.basetype = std::get<1>(op_and_basetype); type.basetype = op_and_basetype.second;
auto vecsize = pls_format_to_components(var.format); auto vecsize = pls_format_to_components(var.format);
if (vecsize > 1) if (vecsize > 1)
{ {
@ -17685,7 +17685,7 @@ bool CompilerGLSL::unroll_array_to_complex_store(uint32_t target_id, uint32_t so
else else
array_expr = to_expression(type.array.back()); array_expr = to_expression(type.array.back());
SPIRType target_type { spv::Op::OpTypeInt }; SPIRType target_type { OpTypeInt };
target_type.basetype = SPIRType::Int; target_type.basetype = SPIRType::Int;
statement("for (int i = 0; i < int(", array_expr, "); i++)"); statement("for (int i = 0; i < int(", array_expr, "); i++)");
@ -17750,7 +17750,7 @@ void CompilerGLSL::unroll_array_from_complex_load(uint32_t target_id, uint32_t s
statement(new_expr, "[i] = gl_in[i].", expr, ";"); statement(new_expr, "[i] = gl_in[i].", expr, ";");
else if (is_sample_mask) else if (is_sample_mask)
{ {
SPIRType target_type { spv::Op::OpTypeInt }; SPIRType target_type { OpTypeInt };
target_type.basetype = SPIRType::Int; target_type.basetype = SPIRType::Int;
statement(new_expr, "[i] = ", bitcast_expression(target_type, type.basetype, join(expr, "[i]")), ";"); statement(new_expr, "[i] = ", bitcast_expression(target_type, type.basetype, join(expr, "[i]")), ";");
} }

12
spirv_hlsl.cpp
View File

@ -2432,7 +2432,7 @@ void CompilerHLSL::analyze_meshlet_writes()
uint32_t op_ptr = op_type + 2; uint32_t op_ptr = op_type + 2;
uint32_t op_var = op_type + 3; uint32_t op_var = op_type + 3;
auto &type = set<SPIRType>(op_type, spv::Op::OpTypeStruct); auto &type = set<SPIRType>(op_type, OpTypeStruct);
type.basetype = SPIRType::Struct; type.basetype = SPIRType::Struct;
set_name(op_type, block_name); set_name(op_type, block_name);
set_decoration(op_type, DecorationBlock); set_decoration(op_type, DecorationBlock);
@ -4508,7 +4508,7 @@ void CompilerHLSL::read_access_chain(string *expr, const string &lhs, const SPIR
{ {
auto &type = get<SPIRType>(chain.basetype); auto &type = get<SPIRType>(chain.basetype);
SPIRType target_type { is_scalar(type) ? spv::Op::OpTypeInt : type.op }; SPIRType target_type { is_scalar(type) ? OpTypeInt : type.op };
target_type.basetype = SPIRType::UInt; target_type.basetype = SPIRType::UInt;
target_type.vecsize = type.vecsize; target_type.vecsize = type.vecsize;
target_type.columns = type.columns; target_type.columns = type.columns;
@ -4755,7 +4755,7 @@ void CompilerHLSL::write_access_chain_array(const SPIRAccessChain &chain, uint32
uint32_t id = ir.increase_bound_by(2); uint32_t id = ir.increase_bound_by(2);
uint32_t int_type_id = id + 1; uint32_t int_type_id = id + 1;
SPIRType int_type { spv::Op::OpTypeInt }; SPIRType int_type { OpTypeInt };
int_type.basetype = SPIRType::Int; int_type.basetype = SPIRType::Int;
int_type.width = 32; int_type.width = 32;
set<SPIRType>(int_type_id, int_type); set<SPIRType>(int_type_id, int_type);
@ -4843,7 +4843,7 @@ void CompilerHLSL::write_access_chain(const SPIRAccessChain &chain, uint32_t val
// Make sure we trigger a read of the constituents in the access chain. // Make sure we trigger a read of the constituents in the access chain.
track_expression_read(chain.self); track_expression_read(chain.self);
SPIRType target_type { is_scalar(type) ? spv::Op::OpTypeInt : type.op }; SPIRType target_type { is_scalar(type) ? OpTypeInt : type.op };
target_type.basetype = SPIRType::UInt; target_type.basetype = SPIRType::UInt;
target_type.vecsize = type.vecsize; target_type.vecsize = type.vecsize;
target_type.columns = type.columns; target_type.columns = type.columns;
@ -6583,14 +6583,14 @@ VariableID CompilerHLSL::remap_num_workgroups_builtin()
uint32_t block_pointer_type_id = offset + 2; uint32_t block_pointer_type_id = offset + 2;
uint32_t variable_id = offset + 3; uint32_t variable_id = offset + 3;
SPIRType uint_type { spv::Op::OpTypeVector }; SPIRType uint_type { OpTypeVector };
uint_type.basetype = SPIRType::UInt; uint_type.basetype = SPIRType::UInt;
uint_type.width = 32; uint_type.width = 32;
uint_type.vecsize = 3; uint_type.vecsize = 3;
uint_type.columns = 1; uint_type.columns = 1;
set<SPIRType>(uint_type_id, uint_type); set<SPIRType>(uint_type_id, uint_type);
SPIRType block_type { spv::Op::OpTypeStruct }; SPIRType block_type { OpTypeStruct };
block_type.basetype = SPIRType::Struct; block_type.basetype = SPIRType::Struct;
block_type.member_types.push_back(uint_type_id); block_type.member_types.push_back(uint_type_id);
set<SPIRType>(block_type_id, block_type); set<SPIRType>(block_type_id, block_type);

241
spirv_msl.cpp
View File

@ -478,21 +478,20 @@ void CompilerMSL::build_implicit_builtins()
uint32_t var_id = offset + 2; uint32_t var_id = offset + 2;
// Create gl_FragCoord. // Create gl_FragCoord.
SPIRType vec4_type { spv::Op::OpTypeVector }; SPIRType vec4_type { OpTypeVector };
vec4_type.basetype = SPIRType::Float; vec4_type.basetype = SPIRType::Float;
vec4_type.width = 32; vec4_type.width = 32;
vec4_type.vecsize = 4; vec4_type.vecsize = 4;
set<SPIRType>(type_id, vec4_type); set<SPIRType>(type_id, vec4_type);
SPIRType vec4_type_ptr { spv::Op::OpTypePointer }; SPIRType vec4_type_ptr = vec4_type;
vec4_type_ptr = vec4_type; vec4_type_ptr.op = OpTypePointer;
vec4_type_ptr.op = spv::Op::OpTypePointer;
vec4_type_ptr.pointer = true; vec4_type_ptr.pointer = true;
vec4_type_ptr.pointer_depth++; vec4_type_ptr.pointer_depth++;
vec4_type_ptr.parent_type = type_id; vec4_type_ptr.parent_type = type_id;
vec4_type_ptr.storage = StorageClassInput; vec4_type_ptr.storage = StorageClassInput;
auto &ptr_type = set<SPIRType>(type_ptr_id, vec4_type_ptr); auto &ptr_type = set<SPIRType>(type_ptr_id, vec4_type_ptr);
ptr_type.self = type_ptr_id; ptr_type.self = type_id;
set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput); set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
set_decoration(var_id, DecorationBuiltIn, BuiltInFragCoord); set_decoration(var_id, DecorationBuiltIn, BuiltInFragCoord);
@ -507,15 +506,14 @@ void CompilerMSL::build_implicit_builtins()
uint32_t var_id = offset + 1; uint32_t var_id = offset + 1;
// Create gl_Layer. // Create gl_Layer.
SPIRType uint_type_ptr { spv::Op::OpTypePointer }; SPIRType uint_type_ptr = get_uint_type();
uint_type_ptr = get_uint_type(); uint_type_ptr.op = OpTypePointer;
uint_type_ptr.op = spv::Op::OpTypePointer;
uint_type_ptr.pointer = true; uint_type_ptr.pointer = true;
uint_type_ptr.pointer_depth++; uint_type_ptr.pointer_depth++;
uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.parent_type = get_uint_type_id();
uint_type_ptr.storage = StorageClassInput; uint_type_ptr.storage = StorageClassInput;
auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr); auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
ptr_type.self = type_ptr_id; ptr_type.self = get_uint_type_id();
set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput); set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
set_decoration(var_id, DecorationBuiltIn, BuiltInLayer); set_decoration(var_id, DecorationBuiltIn, BuiltInLayer);
@ -530,15 +528,14 @@ void CompilerMSL::build_implicit_builtins()
uint32_t var_id = offset + 1; uint32_t var_id = offset + 1;
// Create gl_ViewIndex. // Create gl_ViewIndex.
SPIRType uint_type_ptr { spv::Op::OpTypePointer }; SPIRType uint_type_ptr = get_uint_type();
uint_type_ptr = get_uint_type(); uint_type_ptr.op = OpTypePointer;
uint_type_ptr.op = spv::Op::OpTypePointer;
uint_type_ptr.pointer = true; uint_type_ptr.pointer = true;
uint_type_ptr.pointer_depth++; uint_type_ptr.pointer_depth++;
uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.parent_type = get_uint_type_id();
uint_type_ptr.storage = StorageClassInput; uint_type_ptr.storage = StorageClassInput;
auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr); auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
ptr_type.self = type_ptr_id; ptr_type.self = get_uint_type_id();
set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput); set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
set_decoration(var_id, DecorationBuiltIn, BuiltInViewIndex); set_decoration(var_id, DecorationBuiltIn, BuiltInViewIndex);
@ -554,15 +551,14 @@ void CompilerMSL::build_implicit_builtins()
uint32_t var_id = offset + 1; uint32_t var_id = offset + 1;
// Create gl_SampleID. // Create gl_SampleID.
SPIRType uint_type_ptr { spv::Op::OpTypePointer }; SPIRType uint_type_ptr = get_uint_type();
uint_type_ptr = get_uint_type(); uint_type_ptr.op = OpTypePointer;
uint_type_ptr.op = spv::Op::OpTypePointer;
uint_type_ptr.pointer = true; uint_type_ptr.pointer = true;
uint_type_ptr.pointer_depth++; uint_type_ptr.pointer_depth++;
uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.parent_type = get_uint_type_id();
uint_type_ptr.storage = StorageClassInput; uint_type_ptr.storage = StorageClassInput;
auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr); auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
ptr_type.self = type_ptr_id; ptr_type.self = get_uint_type_id();
set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput); set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
set_decoration(var_id, DecorationBuiltIn, BuiltInSampleId); set_decoration(var_id, DecorationBuiltIn, BuiltInSampleId);
@ -575,15 +571,14 @@ void CompilerMSL::build_implicit_builtins()
{ {
uint32_t type_ptr_id = ir.increase_bound_by(1); uint32_t type_ptr_id = ir.increase_bound_by(1);
SPIRType uint_type_ptr { spv::Op::OpTypePointer }; SPIRType uint_type_ptr = get_uint_type();
uint_type_ptr = get_uint_type(); uint_type_ptr.op = OpTypePointer;
uint_type_ptr.op = spv::Op::OpTypePointer;
uint_type_ptr.pointer = true; uint_type_ptr.pointer = true;
uint_type_ptr.pointer_depth++; uint_type_ptr.pointer_depth++;
uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.parent_type = get_uint_type_id();
uint_type_ptr.storage = StorageClassInput; uint_type_ptr.storage = StorageClassInput;
auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr); auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
ptr_type.self = type_ptr_id; ptr_type.self = get_uint_type_id();
if (need_vertex_params && !has_vertex_idx) if (need_vertex_params && !has_vertex_idx)
{ {
@ -636,9 +631,8 @@ void CompilerMSL::build_implicit_builtins()
// Note that we can't just abuse gl_ViewIndex for this purpose: it's an input, but // Note that we can't just abuse gl_ViewIndex for this purpose: it's an input, but
// gl_Layer is an output in vertex-pipeline shaders. // gl_Layer is an output in vertex-pipeline shaders.
uint32_t type_ptr_out_id = ir.increase_bound_by(2); uint32_t type_ptr_out_id = ir.increase_bound_by(2);
SPIRType uint_type_ptr_out { spv::Op::OpTypePointer }; SPIRType uint_type_ptr_out = get_uint_type();
uint_type_ptr_out = get_uint_type(); uint_type_ptr.op = OpTypePointer;
uint_type_ptr.op = spv::Op::OpTypePointer;
uint_type_ptr_out.pointer = true; uint_type_ptr_out.pointer = true;
uint_type_ptr_out.pointer_depth++; uint_type_ptr_out.pointer_depth++;
uint_type_ptr_out.parent_type = get_uint_type_id(); uint_type_ptr_out.parent_type = get_uint_type_id();
@ -669,15 +663,14 @@ void CompilerMSL::build_implicit_builtins()
{ {
uint32_t type_ptr_id = ir.increase_bound_by(1); uint32_t type_ptr_id = ir.increase_bound_by(1);
SPIRType uint_type_ptr { spv::Op::OpTypePointer }; SPIRType uint_type_ptr = get_uint_type();
uint_type_ptr = get_uint_type(); uint_type_ptr.op = OpTypePointer;
uint_type_ptr.op = spv::Op::OpTypePointer;
uint_type_ptr.pointer = true; uint_type_ptr.pointer = true;
uint_type_ptr.pointer_depth++; uint_type_ptr.pointer_depth++;
uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.parent_type = get_uint_type_id();
uint_type_ptr.storage = StorageClassInput; uint_type_ptr.storage = StorageClassInput;
auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr); auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
ptr_type.self = type_ptr_id; ptr_type.self = get_uint_type_id();
if ((need_tesc_params && msl_options.multi_patch_workgroup) || need_grid_params) if ((need_tesc_params && msl_options.multi_patch_workgroup) || need_grid_params)
{ {
@ -730,15 +723,14 @@ void CompilerMSL::build_implicit_builtins()
uint32_t var_id = offset + 1; uint32_t var_id = offset + 1;
// Create gl_SubgroupInvocationID. // Create gl_SubgroupInvocationID.
SPIRType uint_type_ptr { spv::Op::OpTypePointer }; SPIRType uint_type_ptr = get_uint_type();
uint_type_ptr = get_uint_type(); uint_type_ptr.op = OpTypePointer;
uint_type_ptr.op = spv::Op::OpTypePointer;
uint_type_ptr.pointer = true; uint_type_ptr.pointer = true;
uint_type_ptr.pointer_depth++; uint_type_ptr.pointer_depth++;
uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.parent_type = get_uint_type_id();
uint_type_ptr.storage = StorageClassInput; uint_type_ptr.storage = StorageClassInput;
auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr); auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
ptr_type.self = type_ptr_id; ptr_type.self = get_uint_type_id();
set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput); set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
set_decoration(var_id, DecorationBuiltIn, BuiltInSubgroupLocalInvocationId); set_decoration(var_id, DecorationBuiltIn, BuiltInSubgroupLocalInvocationId);
@ -753,15 +745,14 @@ void CompilerMSL::build_implicit_builtins()
uint32_t var_id = offset + 1; uint32_t var_id = offset + 1;
// Create gl_SubgroupSize. // Create gl_SubgroupSize.
SPIRType uint_type_ptr { spv::Op::OpTypePointer }; SPIRType uint_type_ptr = get_uint_type();
uint_type_ptr = get_uint_type(); uint_type_ptr.op = OpTypePointer;
uint_type_ptr.op = spv::Op::OpTypePointer;
uint_type_ptr.pointer = true; uint_type_ptr.pointer = true;
uint_type_ptr.pointer_depth++; uint_type_ptr.pointer_depth++;
uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.parent_type = get_uint_type_id();
uint_type_ptr.storage = StorageClassInput; uint_type_ptr.storage = StorageClassInput;
auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr); auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
ptr_type.self = type_ptr_id; ptr_type.self = get_uint_type_id();
set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput); set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
set_decoration(var_id, DecorationBuiltIn, BuiltInSubgroupSize); set_decoration(var_id, DecorationBuiltIn, BuiltInSubgroupSize);
@ -813,16 +804,15 @@ void CompilerMSL::build_implicit_builtins()
uint32_t var_id = offset + 1; uint32_t var_id = offset + 1;
// Create gl_SampleMask. // Create gl_SampleMask.
SPIRType uint_type_ptr_out { spv::Op::OpTypePointer }; SPIRType uint_type_ptr_out = get_uint_type();
uint_type_ptr_out = get_uint_type(); uint_type_ptr_out.op = OpTypePointer;
uint_type_ptr_out.op = spv::Op::OpTypePointer;
uint_type_ptr_out.pointer = true; uint_type_ptr_out.pointer = true;
uint_type_ptr_out.pointer_depth++; uint_type_ptr_out.pointer_depth++;
uint_type_ptr_out.parent_type = get_uint_type_id(); uint_type_ptr_out.parent_type = get_uint_type_id();
uint_type_ptr_out.storage = StorageClassOutput; uint_type_ptr_out.storage = StorageClassOutput;
auto &ptr_out_type = set<SPIRType>(offset, uint_type_ptr_out); auto &ptr_out_type = set<SPIRType>(offset, uint_type_ptr_out);
ptr_out_type.self = offset; ptr_out_type.self = get_uint_type_id();
set<SPIRVariable>(var_id, offset, StorageClassOutput); set<SPIRVariable>(var_id, offset, StorageClassOutput);
set_decoration(var_id, DecorationBuiltIn, BuiltInSampleMask); set_decoration(var_id, DecorationBuiltIn, BuiltInSampleMask);
builtin_sample_mask_id = var_id; builtin_sample_mask_id = var_id;
@ -837,14 +827,13 @@ void CompilerMSL::build_implicit_builtins()
uint32_t var_id = offset + 2; uint32_t var_id = offset + 2;
// Create gl_HelperInvocation. // Create gl_HelperInvocation.
SPIRType bool_type { spv::Op::OpTypeBool }; SPIRType bool_type { OpTypeBool };
bool_type.basetype = SPIRType::Boolean; bool_type.basetype = SPIRType::Boolean;
bool_type.width = 8; bool_type.width = 8;
bool_type.vecsize = 1; bool_type.vecsize = 1;
set<SPIRType>(type_id, bool_type); set<SPIRType>(type_id, bool_type);
SPIRType bool_type_ptr_in { spv::Op::OpTypePointer }; SPIRType bool_type_ptr_in = bool_type;
bool_type_ptr_in = bool_type;
bool_type_ptr_in.op = spv::OpTypePointer; bool_type_ptr_in.op = spv::OpTypePointer;
bool_type_ptr_in.pointer = true; bool_type_ptr_in.pointer = true;
bool_type_ptr_in.pointer_depth++; bool_type_ptr_in.pointer_depth++;
@ -852,7 +841,7 @@ void CompilerMSL::build_implicit_builtins()
bool_type_ptr_in.storage = StorageClassInput; bool_type_ptr_in.storage = StorageClassInput;
auto &ptr_in_type = set<SPIRType>(type_ptr_id, bool_type_ptr_in); auto &ptr_in_type = set<SPIRType>(type_ptr_id, bool_type_ptr_in);
ptr_in_type.self = type_ptr_id; ptr_in_type.self = type_id;
set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput); set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
set_decoration(var_id, DecorationBuiltIn, BuiltInHelperInvocation); set_decoration(var_id, DecorationBuiltIn, BuiltInHelperInvocation);
builtin_helper_invocation_id = var_id; builtin_helper_invocation_id = var_id;
@ -866,16 +855,15 @@ void CompilerMSL::build_implicit_builtins()
uint32_t var_id = offset + 1; uint32_t var_id = offset + 1;
// Create gl_LocalInvocationIndex. // Create gl_LocalInvocationIndex.
SPIRType uint_type_ptr { spv::Op::OpTypePointer }; SPIRType uint_type_ptr = get_uint_type();
uint_type_ptr = get_uint_type(); uint_type_ptr.op = OpTypePointer;
uint_type_ptr.op = spv::Op::OpTypePointer;
uint_type_ptr.pointer = true; uint_type_ptr.pointer = true;
uint_type_ptr.pointer_depth++; uint_type_ptr.pointer_depth++;
uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.parent_type = get_uint_type_id();
uint_type_ptr.storage = StorageClassInput; uint_type_ptr.storage = StorageClassInput;
auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr); auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
ptr_type.self = type_ptr_id; ptr_type.self = get_uint_type_id();
set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput); set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
set_decoration(var_id, DecorationBuiltIn, BuiltInLocalInvocationIndex); set_decoration(var_id, DecorationBuiltIn, BuiltInLocalInvocationIndex);
builtin_local_invocation_index_id = var_id; builtin_local_invocation_index_id = var_id;
@ -891,14 +879,14 @@ void CompilerMSL::build_implicit_builtins()
// Create gl_WorkgroupSize. // Create gl_WorkgroupSize.
uint32_t type_id = build_extended_vector_type(get_uint_type_id(), 3); uint32_t type_id = build_extended_vector_type(get_uint_type_id(), 3);
SPIRType uint_type_ptr = get<SPIRType>(type_id); SPIRType uint_type_ptr = get<SPIRType>(type_id);
uint_type_ptr.op = spv::Op::OpTypePointer; uint_type_ptr.op = OpTypePointer;
uint_type_ptr.pointer = true; uint_type_ptr.pointer = true;
uint_type_ptr.pointer_depth++; uint_type_ptr.pointer_depth++;
uint_type_ptr.parent_type = type_id; uint_type_ptr.parent_type = type_id;
uint_type_ptr.storage = StorageClassInput; uint_type_ptr.storage = StorageClassInput;
auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr); auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
ptr_type.self = type_ptr_id; ptr_type.self = type_id;
set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput); set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
set_decoration(var_id, DecorationBuiltIn, BuiltInWorkgroupSize); set_decoration(var_id, DecorationBuiltIn, BuiltInWorkgroupSize);
builtin_workgroup_size_id = var_id; builtin_workgroup_size_id = var_id;
@ -999,21 +987,20 @@ void CompilerMSL::build_implicit_builtins()
uint32_t var_id = offset + 2; uint32_t var_id = offset + 2;
// Create gl_Position. // Create gl_Position.
SPIRType vec4_type { spv::Op::OpTypeVector }; SPIRType vec4_type { OpTypeVector };
vec4_type.basetype = SPIRType::Float; vec4_type.basetype = SPIRType::Float;
vec4_type.width = 32; vec4_type.width = 32;
vec4_type.vecsize = 4; vec4_type.vecsize = 4;
set<SPIRType>(type_id, vec4_type); set<SPIRType>(type_id, vec4_type);
SPIRType vec4_type_ptr { spv::Op::OpTypePointer }; SPIRType vec4_type_ptr = vec4_type;
vec4_type_ptr = vec4_type; vec4_type_ptr.op = OpTypePointer;
vec4_type_ptr.op = spv::Op::OpTypePointer;
vec4_type_ptr.pointer = true; vec4_type_ptr.pointer = true;
vec4_type_ptr.pointer_depth++; vec4_type_ptr.pointer_depth++;
vec4_type_ptr.parent_type = type_id; vec4_type_ptr.parent_type = type_id;
vec4_type_ptr.storage = StorageClassOutput; vec4_type_ptr.storage = StorageClassOutput;
auto &ptr_type = set<SPIRType>(type_ptr_id, vec4_type_ptr); auto &ptr_type = set<SPIRType>(type_ptr_id, vec4_type_ptr);
ptr_type.self = type_ptr_id; ptr_type.self = type_id;
set<SPIRVariable>(var_id, type_ptr_id, StorageClassOutput); set<SPIRVariable>(var_id, type_ptr_id, StorageClassOutput);
set_decoration(var_id, DecorationBuiltIn, BuiltInPosition); set_decoration(var_id, DecorationBuiltIn, BuiltInPosition);
@ -1084,7 +1071,7 @@ uint32_t CompilerMSL::build_constant_uint_array_pointer()
// Create a buffer to hold extra data, including the swizzle constants. // Create a buffer to hold extra data, including the swizzle constants.
SPIRType uint_type_pointer = get_uint_type(); SPIRType uint_type_pointer = get_uint_type();
uint_type_pointer.op = spv::Op::OpTypePointer; uint_type_pointer.op = OpTypePointer;
uint_type_pointer.pointer = true; uint_type_pointer.pointer = true;
uint_type_pointer.pointer_depth++; uint_type_pointer.pointer_depth++;
uint_type_pointer.parent_type = get_uint_type_id(); uint_type_pointer.parent_type = get_uint_type_id();
@ -1163,7 +1150,7 @@ uint32_t CompilerMSL::get_uint_type_id()
uint_type_id = ir.increase_bound_by(1); uint_type_id = ir.increase_bound_by(1);
SPIRType type { spv::Op::OpTypeInt }; SPIRType type { OpTypeInt };
type.basetype = SPIRType::UInt; type.basetype = SPIRType::UInt;
type.width = 32; type.width = 32;
set<SPIRType>(uint_type_id, type); set<SPIRType>(uint_type_id, type);
@ -2186,7 +2173,7 @@ void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std::
// Make sure we have an actual pointer type, // Make sure we have an actual pointer type,
// so that we will get the appropriate address space when declaring these builtins. // so that we will get the appropriate address space when declaring these builtins.
auto &ptr = set<SPIRType>(ptr_type_id, get<SPIRType>(mbr_type_id)); auto &ptr = set<SPIRType>(ptr_type_id, get<SPIRType>(mbr_type_id));
ptr.self = ptr_type_id; ptr.self = mbr_type_id;
ptr.storage = var.storage; ptr.storage = var.storage;
ptr.pointer = true; ptr.pointer = true;
ptr.pointer_depth++; ptr.pointer_depth++;
@ -2348,23 +2335,25 @@ uint32_t CompilerMSL::build_extended_vector_type(uint32_t type_id, uint32_t comp
{ {
assert(components > 1); assert(components > 1);
uint32_t new_type_id = ir.increase_bound_by(1); uint32_t new_type_id = ir.increase_bound_by(1);
auto &old_type = get<SPIRType>(type_id); const auto *p_old_type = &get<SPIRType>(type_id);
SPIRType* old_ptr_t = nullptr; const SPIRType *old_ptr_t = nullptr;
SPIRType* old_array_t = nullptr; const SPIRType *old_array_t = nullptr;
if (is_pointer(old_type))
if (is_pointer(*p_old_type))
{ {
old_ptr_t = &old_type; old_ptr_t = p_old_type;
old_type = get_pointee_type(*old_ptr_t); p_old_type = &get_pointee_type(*old_ptr_t);
}
if (is_array(old_type))
{
old_array_t = &old_type;
old_type = get_type(old_array_t->parent_type);
} }
auto *type = &set<SPIRType>(new_type_id, old_type); if (is_array(*p_old_type))
{
old_array_t = p_old_type;
p_old_type = &get_type(old_array_t->parent_type);
}
auto *type = &set<SPIRType>(new_type_id, *p_old_type);
assert(is_scalar(*type) || is_vector(*type)); assert(is_scalar(*type) || is_vector(*type));
type->op = spv::Op::OpTypeVector; type->op = OpTypeVector;
type->vecsize = components; type->vecsize = components;
if (basetype != SPIRType::Unknown) if (basetype != SPIRType::Unknown)
type->basetype = basetype; type->basetype = basetype;
@ -2378,6 +2367,7 @@ uint32_t CompilerMSL::build_extended_vector_type(uint32_t type_id, uint32_t comp
{ {
uint32_t array_type_id = ir.increase_bound_by(1); uint32_t array_type_id = ir.increase_bound_by(1);
type = &set<SPIRType>(array_type_id, *type); type = &set<SPIRType>(array_type_id, *type);
type->op = OpTypeArray;
type->parent_type = new_type_id; type->parent_type = new_type_id;
type->array = old_array_t->array; type->array = old_array_t->array;
type->array_size_literal = old_array_t->array_size_literal; type->array_size_literal = old_array_t->array_size_literal;
@ -2388,7 +2378,7 @@ uint32_t CompilerMSL::build_extended_vector_type(uint32_t type_id, uint32_t comp
{ {
uint32_t ptr_type_id = ir.increase_bound_by(1); uint32_t ptr_type_id = ir.increase_bound_by(1);
type = &set<SPIRType>(ptr_type_id, *type); type = &set<SPIRType>(ptr_type_id, *type);
type->self = new_type_id; type->op = OpTypePointer;
type->parent_type = new_type_id; type->parent_type = new_type_id;
type->storage = old_ptr_t->storage; type->storage = old_ptr_t->storage;
type->pointer = true; type->pointer = true;
@ -3968,7 +3958,7 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
// declaraion is emitted, because it is cleared after each compilation pass. // declaraion is emitted, because it is cleared after each compilation pass.
uint32_t next_id = ir.increase_bound_by(3); uint32_t next_id = ir.increase_bound_by(3);
uint32_t ib_type_id = next_id++; uint32_t ib_type_id = next_id++;
auto &ib_type = set<SPIRType>(ib_type_id, spv::Op::OpTypeStruct); auto &ib_type = set<SPIRType>(ib_type_id, OpTypeStruct);
ib_type.basetype = SPIRType::Struct; ib_type.basetype = SPIRType::Struct;
ib_type.storage = storage; ib_type.storage = storage;
set_decoration(ib_type_id, DecorationBlock); set_decoration(ib_type_id, DecorationBlock);
@ -4198,7 +4188,7 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
uint32_t ptr_type_id = offset + 3; uint32_t ptr_type_id = offset + 3;
uint32_t var_id = offset + 4; uint32_t var_id = offset + 4;
SPIRType type { spv::Op::OpTypeInt }; SPIRType type { OpTypeInt };
switch (input.second.format) switch (input.second.format)
{ {
case MSL_SHADER_VARIABLE_FORMAT_UINT16: case MSL_SHADER_VARIABLE_FORMAT_UINT16:
@ -4215,26 +4205,26 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
set<SPIRType>(type_id, type); set<SPIRType>(type_id, type);
if (input.second.vecsize > 1) if (input.second.vecsize > 1)
{ {
type.op = spv::Op::OpTypeVector; type.op = OpTypeVector;
type.vecsize = input.second.vecsize; type.vecsize = input.second.vecsize;
set<SPIRType>(vec_type_id, type); set<SPIRType>(vec_type_id, type);
type_id = vec_type_id; type_id = vec_type_id;
} }
type.op = spv::Op::OpTypeArray; type.op = OpTypeArray;
type.array.push_back(0); type.array.push_back(0);
type.array_size_literal.push_back(true); type.array_size_literal.push_back(true);
type.parent_type = type_id; type.parent_type = type_id;
set<SPIRType>(array_type_id, type); set<SPIRType>(array_type_id, type);
type.self = array_type_id; type.self = type_id;
type.op = spv::Op::OpTypePointer; type.op = OpTypePointer;
type.pointer = true; type.pointer = true;
type.pointer_depth++; type.pointer_depth++;
type.parent_type = array_type_id; type.parent_type = array_type_id;
type.storage = storage; type.storage = storage;
auto &ptr_type = set<SPIRType>(ptr_type_id, type); auto &ptr_type = set<SPIRType>(ptr_type_id, type);
ptr_type.self = ptr_type_id; ptr_type.self = array_type_id;
auto &fake_var = set<SPIRVariable>(var_id, ptr_type_id, storage); auto &fake_var = set<SPIRVariable>(var_id, ptr_type_id, storage);
set_decoration(var_id, DecorationLocation, input.first.location); set_decoration(var_id, DecorationLocation, input.first.location);
@ -4266,7 +4256,7 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
uint32_t ptr_type_id = offset + 3; uint32_t ptr_type_id = offset + 3;
uint32_t var_id = offset + 4; uint32_t var_id = offset + 4;
SPIRType type { spv::Op::OpTypeInt }; SPIRType type { OpTypeInt };
switch (output.second.format) switch (output.second.format)
{ {
case MSL_SHADER_VARIABLE_FORMAT_UINT16: case MSL_SHADER_VARIABLE_FORMAT_UINT16:
@ -4283,7 +4273,7 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
set<SPIRType>(type_id, type); set<SPIRType>(type_id, type);
if (output.second.vecsize > 1) if (output.second.vecsize > 1)
{ {
type.op = spv::Op::OpTypeVector; type.op = OpTypeVector;
type.vecsize = output.second.vecsize; type.vecsize = output.second.vecsize;
set<SPIRType>(vec_type_id, type); set<SPIRType>(vec_type_id, type);
type_id = vec_type_id; type_id = vec_type_id;
@ -4291,14 +4281,14 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
if (is_tesc_shader()) if (is_tesc_shader())
{ {
type.op = spv::Op::OpTypeArray; type.op = OpTypeArray;
type.array.push_back(0); type.array.push_back(0);
type.array_size_literal.push_back(true); type.array_size_literal.push_back(true);
type.parent_type = vec_type_id; type.parent_type = type_id;
set<SPIRType>(array_type_id, type); set<SPIRType>(array_type_id, type);
} }
type.op = spv::Op::OpTypePointer; type.op = OpTypePointer;
type.pointer = true; type.pointer = true;
type.pointer_depth++; type.pointer_depth++;
type.parent_type = is_tesc_shader() ? array_type_id : type_id; type.parent_type = is_tesc_shader() ? array_type_id : type_id;
@ -4385,6 +4375,7 @@ uint32_t CompilerMSL::add_interface_block_pointer(uint32_t ib_var_id, StorageCla
// do the same with our struct here. // do the same with our struct here.
uint32_t ib_ptr_type_id = next_id++; uint32_t ib_ptr_type_id = next_id++;
auto &ib_ptr_type = set<SPIRType>(ib_ptr_type_id, ib_type); auto &ib_ptr_type = set<SPIRType>(ib_ptr_type_id, ib_type);
ib_ptr_type.op = OpTypePointer;
ib_ptr_type.parent_type = ib_ptr_type.type_alias = ib_type.self; ib_ptr_type.parent_type = ib_ptr_type.type_alias = ib_type.self;
ib_ptr_type.pointer = true; ib_ptr_type.pointer = true;
ib_ptr_type.pointer_depth++; ib_ptr_type.pointer_depth++;
@ -4438,27 +4429,24 @@ uint32_t CompilerMSL::add_interface_block_pointer(uint32_t ib_var_id, StorageCla
// Otherwise, create a new type, and return it's ID. // Otherwise, create a new type, and return it's ID.
uint32_t CompilerMSL::ensure_correct_builtin_type(uint32_t type_id, BuiltIn builtin) uint32_t CompilerMSL::ensure_correct_builtin_type(uint32_t type_id, BuiltIn builtin)
{ {
auto& type = get<SPIRType>(type_id); auto &type = get<SPIRType>(type_id);
auto pointee_type = type; auto &pointee_type = get_pointee_type(type);
if (is_pointer(type))
pointee_type = get<SPIRType>(type.parent_type);
if ((builtin == BuiltInSampleMask && is_array(pointee_type)) || if ((builtin == BuiltInSampleMask && is_array(pointee_type)) ||
((builtin == BuiltInLayer || builtin == BuiltInViewportIndex || builtin == BuiltInFragStencilRefEXT) && ((builtin == BuiltInLayer || builtin == BuiltInViewportIndex || builtin == BuiltInFragStencilRefEXT) &&
pointee_type.basetype != SPIRType::UInt)) pointee_type.basetype != SPIRType::UInt))
{ {
uint32_t next_id = ir.increase_bound_by(type.pointer ? 2 : 1); uint32_t next_id = ir.increase_bound_by(type_is_pointer(type) ? 2 : 1);
uint32_t base_type_id = next_id++; uint32_t base_type_id = next_id++;
auto &base_type = set<SPIRType>(base_type_id, spv::Op::OpTypeInt); auto &base_type = set<SPIRType>(base_type_id, OpTypeInt);
base_type.basetype = SPIRType::UInt; base_type.basetype = SPIRType::UInt;
base_type.width = 32; base_type.width = 32;
if (!type.pointer) if (!type_is_pointer(type))
return base_type_id; return base_type_id;
uint32_t ptr_type_id = next_id++; uint32_t ptr_type_id = next_id++;
auto &ptr_type = set<SPIRType>(ptr_type_id, spv::OpTypePointer); auto &ptr_type = set<SPIRType>(ptr_type_id, base_type);
ptr_type = base_type;
ptr_type.op = spv::OpTypePointer; ptr_type.op = spv::OpTypePointer;
ptr_type.pointer = true; ptr_type.pointer = true;
ptr_type.pointer_depth++; ptr_type.pointer_depth++;
@ -4949,7 +4937,7 @@ void CompilerMSL::ensure_member_packing_rules_msl(SPIRType &ib_type, uint32_t in
{ {
type.columns = 1; type.columns = 1;
assert(type.array.empty()); assert(type.array.empty());
type.op = spv::Op::OpTypeArray; type.op = OpTypeArray;
type.array.push_back(1); type.array.push_back(1);
type.array_size_literal.push_back(true); type.array_size_literal.push_back(true);
} }
@ -4966,7 +4954,7 @@ void CompilerMSL::ensure_member_packing_rules_msl(SPIRType &ib_type, uint32_t in
type.vecsize = type.columns; type.vecsize = type.columns;
type.columns = 1; type.columns = 1;
assert(type.array.empty()); assert(type.array.empty());
type.op = spv::Op::OpTypeArray; type.op = OpTypeArray;
type.array.push_back(1); type.array.push_back(1);
type.array_size_literal.push_back(true); type.array_size_literal.push_back(true);
} }
@ -4991,7 +4979,7 @@ void CompilerMSL::emit_store_statement(uint32_t lhs_expression, uint32_t rhs_exp
if (has_decoration(lhs_expression, DecorationBuiltIn) && if (has_decoration(lhs_expression, DecorationBuiltIn) &&
BuiltIn(get_decoration(lhs_expression, DecorationBuiltIn)) == BuiltInSampleMask && BuiltIn(get_decoration(lhs_expression, DecorationBuiltIn)) == BuiltInSampleMask &&
type_is_top_level_array(type)) is_array(type))
{ {
// Storing an array to SampleMask, have to remove the array-ness before storing. // Storing an array to SampleMask, have to remove the array-ness before storing.
statement(to_expression(lhs_expression), " = ", to_enclosed_unpacked_expression(rhs_expression), "[0];"); statement(to_expression(lhs_expression), " = ", to_enclosed_unpacked_expression(rhs_expression), "[0];");
@ -7507,7 +7495,7 @@ void CompilerMSL::declare_constant_arrays()
// FIXME: However, hoisting constants to main() means we need to pass down constant arrays to leaf functions if they are used there. // FIXME: However, hoisting constants to main() means we need to pass down constant arrays to leaf functions if they are used there.
// If there are multiple functions in the module, drop this case to avoid breaking use cases which do not need to // If there are multiple functions in the module, drop this case to avoid breaking use cases which do not need to
// link into Metal libraries. This is hacky. // link into Metal libraries. This is hacky.
if (type_is_top_level_array(type) && (!fully_inlined || is_scalar(type) || is_vector(type))) if (is_array(type) && (!fully_inlined || is_scalar(type) || is_vector(type)))
{ {
add_resource_name(c.self); add_resource_name(c.self);
auto name = to_name(c.self); auto name = to_name(c.self);
@ -7539,7 +7527,7 @@ void CompilerMSL::declare_complex_constant_arrays()
return; return;
auto &type = this->get<SPIRType>(c.constant_type); auto &type = this->get<SPIRType>(c.constant_type);
if (type_is_top_level_array(type) && !(is_scalar(type) || is_vector(type))) if (is_array(type) && !(is_scalar(type) || is_vector(type)))
{ {
add_resource_name(c.self); add_resource_name(c.self);
auto name = to_name(c.self); auto name = to_name(c.self);
@ -8278,7 +8266,7 @@ bool CompilerMSL::emit_tessellation_access_chain(const uint32_t *ops, uint32_t l
// We're not going to emit the actual member name, we let any further OpLoad take care of that. // We're not going to emit the actual member name, we let any further OpLoad take care of that.
// Tag the access chain with the member index we're referencing. // Tag the access chain with the member index we're referencing.
auto& result_pointee_type = get_pointee_type(result_ptr_type); auto &result_pointee_type = get_pointee_type(result_ptr_type);
bool defer_access_chain = flatten_composites && (is_matrix(result_pointee_type) || is_array(result_pointee_type) || bool defer_access_chain = flatten_composites && (is_matrix(result_pointee_type) || is_array(result_pointee_type) ||
result_pointee_type.basetype == SPIRType::Struct); result_pointee_type.basetype == SPIRType::Struct);
@ -9921,7 +9909,7 @@ bool CompilerMSL::maybe_emit_array_assignment(uint32_t id_lhs, uint32_t id_rhs)
{ {
// We only care about assignments of an entire array // We only care about assignments of an entire array
auto &type = expression_type(id_lhs); auto &type = expression_type(id_lhs);
if (!type_is_top_level_array(get_pointee_type(type))) if (!is_array(get_pointee_type(type)))
return false; return false;
auto *var = maybe_get<SPIRVariable>(id_lhs); auto *var = maybe_get<SPIRVariable>(id_lhs);
@ -10871,7 +10859,7 @@ string CompilerMSL::to_function_name(const TextureFunctionNameArguments &args)
string CompilerMSL::convert_to_f32(const string &expr, uint32_t components) string CompilerMSL::convert_to_f32(const string &expr, uint32_t components)
{ {
SPIRType t { components > 1 ? spv::Op::OpTypeVector : spv::Op::OpTypeFloat }; SPIRType t { components > 1 ? OpTypeVector : OpTypeFloat };
t.basetype = SPIRType::Float; t.basetype = SPIRType::Float;
t.vecsize = components; t.vecsize = components;
t.columns = 1; t.columns = 1;
@ -11980,7 +11968,7 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_
if (is_matrix(physical_type)) if (is_matrix(physical_type))
row_major = has_member_decoration(type.self, index, DecorationRowMajor); row_major = has_member_decoration(type.self, index, DecorationRowMajor);
SPIRType row_major_physical_type { spv::Op::OpTypeMatrix }; SPIRType row_major_physical_type { OpTypeMatrix };
const SPIRType *declared_type = &physical_type; const SPIRType *declared_type = &physical_type;
// If a struct is being declared with physical layout, // If a struct is being declared with physical layout,
@ -14375,7 +14363,6 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
auto &type = get_variable_data_type(var); auto &type = get_variable_data_type(var);
auto &var_type = get<SPIRType>(arg.type); auto &var_type = get<SPIRType>(arg.type);
StorageClass type_storage = var_type.storage; StorageClass type_storage = var_type.storage;
bool is_pointer = var_type.pointer;
// If we need to modify the name of the variable, make sure we use the original variable. // If we need to modify the name of the variable, make sure we use the original variable.
// Our alias is just a shadow variable. // Our alias is just a shadow variable.
@ -14383,7 +14370,7 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
if (arg.alias_global_variable && var.basevariable) if (arg.alias_global_variable && var.basevariable)
name_id = var.basevariable; name_id = var.basevariable;
bool constref = !arg.alias_global_variable && is_pointer && arg.write_count == 0; bool constref = !arg.alias_global_variable && is_pointer(var_type) && arg.write_count == 0;
// Framebuffer fetch is plain value, const looks out of place, but it is not wrong. // Framebuffer fetch is plain value, const looks out of place, but it is not wrong.
if (type_is_msl_framebuffer_fetch(type)) if (type_is_msl_framebuffer_fetch(type))
constref = false; constref = false;
@ -14481,7 +14468,7 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
} }
} }
if (!builtin && !is_pointer && if (!builtin && !is_pointer(var_type) &&
(type_storage == StorageClassFunction || type_storage == StorageClassGeneric)) (type_storage == StorageClassFunction || type_storage == StorageClassGeneric))
{ {
// If the argument is a pure value and not an opaque type, we will pass by value. // If the argument is a pure value and not an opaque type, we will pass by value.
@ -15060,7 +15047,7 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id, bool member)
string type_name; string type_name;
// Pointer? // Pointer?
if (type_is_top_level_pointer(type) || type_is_array_of_pointers(type)) if (is_pointer(type) || type_is_array_of_pointers(type))
{ {
assert(type.pointer_depth > 0); assert(type.pointer_depth > 0);
@ -15088,7 +15075,7 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id, bool member)
// the C-style nesting works right. // the C-style nesting works right.
// FIXME: This is somewhat of a hack. // FIXME: This is somewhat of a hack.
bool old_is_using_builtin_array = is_using_builtin_array; bool old_is_using_builtin_array = is_using_builtin_array;
if (type_is_top_level_physical_pointer(type)) if (is_physical_pointer(type))
is_using_builtin_array = false; is_using_builtin_array = false;
type_name = join(type_address_space, " ", type_to_glsl(*p_parent_type, id)); type_name = join(type_address_space, " ", type_to_glsl(*p_parent_type, id));
@ -16947,8 +16934,7 @@ bool CompilerMSL::OpCodePreprocessor::handle(Op opcode, const uint32_t *args, ui
auto *var = compiler.maybe_get_backing_variable(args[2]); auto *var = compiler.maybe_get_backing_variable(args[2]);
if (var != nullptr) if (var != nullptr)
{ {
auto &type = compiler.get<SPIRType>(var->basetype); if (!compiler.is_var_runtime_size_array(*var))
if (!is_runtime_size_array(type))
compiler.buffers_requiring_array_length.insert(var->self); compiler.buffers_requiring_array_length.insert(var->self);
} }
break; break;
@ -17361,7 +17347,7 @@ void CompilerMSL::cast_from_variable_load(uint32_t source_id, std::string &expr,
// Type fixups for workgroup variables if they are booleans. // Type fixups for workgroup variables if they are booleans.
if (rewrite_boolean_load) if (rewrite_boolean_load)
{ {
if (type_is_top_level_array(expr_type)) if (is_array(expr_type))
expr = to_rerolled_array_expression(expr_type, expr, expr_type); expr = to_rerolled_array_expression(expr_type, expr, expr_type);
else else
expr = join(type_to_glsl(expr_type), "(", expr, ")"); expr = join(type_to_glsl(expr_type), "(", expr, ")");
@ -17438,7 +17424,7 @@ void CompilerMSL::cast_from_variable_load(uint32_t source_id, std::string &expr,
break; break;
} }
if (type_is_top_level_array(expr_type) && builtin == BuiltInSampleMask) if (is_array(expr_type) && builtin == BuiltInSampleMask)
{ {
// Needs special handling. // Needs special handling.
auto wrap_expr = join(type_to_glsl(expr_type), "({ "); auto wrap_expr = join(type_to_glsl(expr_type), "({ ");
@ -17448,7 +17434,7 @@ void CompilerMSL::cast_from_variable_load(uint32_t source_id, std::string &expr,
} }
else if (expected_type != expr_type.basetype) else if (expected_type != expr_type.basetype)
{ {
if (type_is_top_level_array(expr_type) && (builtin == BuiltInTessLevelInner || builtin == BuiltInTessLevelOuter)) if (is_array(expr_type) && (builtin == BuiltInTessLevelInner || builtin == BuiltInTessLevelOuter))
{ {
// Triggers when loading TessLevel directly as an array. // Triggers when loading TessLevel directly as an array.
// Need explicit padding + cast. // Need explicit padding + cast.
@ -17507,7 +17493,7 @@ void CompilerMSL::cast_to_variable_store(uint32_t target_id, std::string &expr,
// Type fixups for workgroup variables or struct members if they are booleans. // Type fixups for workgroup variables or struct members if they are booleans.
if (rewrite_boolean_store) if (rewrite_boolean_store)
{ {
if (type_is_top_level_array(expr_type)) if (is_array(expr_type))
{ {
expr = to_rerolled_array_expression(*var_type, expr, expr_type); expr = to_rerolled_array_expression(*var_type, expr, expr_type);
} }
@ -17786,6 +17772,7 @@ void CompilerMSL::analyze_argument_buffers()
// Create a buffer to hold extra data, including the swizzle constants. // Create a buffer to hold extra data, including the swizzle constants.
SPIRType uint_type_pointer = get_uint_type(); SPIRType uint_type_pointer = get_uint_type();
uint_type_pointer.op = OpTypePointer;
uint_type_pointer.pointer = true; uint_type_pointer.pointer = true;
uint_type_pointer.pointer_depth++; uint_type_pointer.pointer_depth++;
uint_type_pointer.parent_type = get_uint_type_id(); uint_type_pointer.parent_type = get_uint_type_id();
@ -17841,7 +17828,7 @@ void CompilerMSL::analyze_argument_buffers()
uint32_t ptr_type_id = next_id + 2; uint32_t ptr_type_id = next_id + 2;
argument_buffer_ids[desc_set] = next_id; argument_buffer_ids[desc_set] = next_id;
auto &buffer_type = set<SPIRType>(type_id, spv::Op::OpTypeStruct); auto &buffer_type = set<SPIRType>(type_id, OpTypeStruct);
buffer_type.basetype = SPIRType::Struct; buffer_type.basetype = SPIRType::Struct;
@ -17858,7 +17845,7 @@ void CompilerMSL::analyze_argument_buffers()
set_name(type_id, join("spvDescriptorSetBuffer", desc_set)); set_name(type_id, join("spvDescriptorSetBuffer", desc_set));
auto &ptr_type = set<SPIRType>(ptr_type_id, spv::Op::OpTypePointer); auto &ptr_type = set<SPIRType>(ptr_type_id, OpTypePointer);
ptr_type = buffer_type; ptr_type = buffer_type;
ptr_type.op = spv::OpTypePointer; ptr_type.op = spv::OpTypePointer;
ptr_type.pointer = true; ptr_type.pointer = true;
@ -17943,14 +17930,14 @@ void CompilerMSL::analyze_argument_buffers()
bool type_is_array = !type.array.empty(); bool type_is_array = !type.array.empty();
uint32_t sampler_type_id = ir.increase_bound_by(type_is_array ? 2 : 1); uint32_t sampler_type_id = ir.increase_bound_by(type_is_array ? 2 : 1);
auto &new_sampler_type = set<SPIRType>(sampler_type_id, spv::Op::OpTypeSampler); auto &new_sampler_type = set<SPIRType>(sampler_type_id, OpTypeSampler);
new_sampler_type.basetype = SPIRType::Sampler; new_sampler_type.basetype = SPIRType::Sampler;
new_sampler_type.storage = StorageClassUniformConstant; new_sampler_type.storage = StorageClassUniformConstant;
if (type_is_array) if (type_is_array)
{ {
uint32_t sampler_type_array_id = sampler_type_id + 1; uint32_t sampler_type_array_id = sampler_type_id + 1;
auto &sampler_type_array = set<SPIRType>(sampler_type_array_id, spv::Op::OpTypeArray); auto &sampler_type_array = set<SPIRType>(sampler_type_array_id, OpTypeArray);
sampler_type_array = new_sampler_type; sampler_type_array = new_sampler_type;
sampler_type_array.array = type.array; sampler_type_array.array = type.array;
sampler_type_array.array_size_literal = type.array_size_literal; sampler_type_array.array_size_literal = type.array_size_literal;
@ -18001,18 +17988,19 @@ void CompilerMSL::analyze_argument_buffers()
uint32_t atomic_type_id = offset; uint32_t atomic_type_id = offset;
uint32_t type_ptr_id = offset + 1; uint32_t type_ptr_id = offset + 1;
SPIRType atomic_type { spv::Op::OpTypeInt }; SPIRType atomic_type { OpTypeInt };
atomic_type.basetype = SPIRType::AtomicCounter; atomic_type.basetype = SPIRType::AtomicCounter;
atomic_type.width = 32; atomic_type.width = 32;
atomic_type.vecsize = 1; atomic_type.vecsize = 1;
set<SPIRType>(atomic_type_id, atomic_type); set<SPIRType>(atomic_type_id, atomic_type);
atomic_type.op = OpTypePointer;
atomic_type.pointer = true; atomic_type.pointer = true;
atomic_type.pointer_depth++; atomic_type.pointer_depth++;
atomic_type.parent_type = atomic_type_id; atomic_type.parent_type = atomic_type_id;
atomic_type.storage = StorageClassStorageBuffer; atomic_type.storage = StorageClassStorageBuffer;
auto &atomic_ptr_type = set<SPIRType>(type_ptr_id, atomic_type); auto &atomic_ptr_type = set<SPIRType>(type_ptr_id, atomic_type);
atomic_ptr_type.self = type_ptr_id; atomic_ptr_type.self = atomic_type_id;
buffer_type.member_types.push_back(type_ptr_id); buffer_type.member_types.push_back(type_ptr_id);
} }
@ -18067,12 +18055,12 @@ void CompilerMSL::add_argument_buffer_padding_buffer_type(SPIRType &struct_type,
if (!argument_buffer_padding_buffer_type_id) if (!argument_buffer_padding_buffer_type_id)
{ {
uint32_t buff_type_id = ir.increase_bound_by(2); uint32_t buff_type_id = ir.increase_bound_by(2);
auto &buff_type = set<SPIRType>(buff_type_id, spv::Op::OpNop); auto &buff_type = set<SPIRType>(buff_type_id, OpNop);
buff_type.basetype = rez_bind.basetype; buff_type.basetype = rez_bind.basetype;
buff_type.storage = StorageClassUniformConstant; buff_type.storage = StorageClassUniformConstant;
uint32_t ptr_type_id = buff_type_id + 1; uint32_t ptr_type_id = buff_type_id + 1;
auto &ptr_type = set<SPIRType>(ptr_type_id, spv::Op::OpTypePointer); auto &ptr_type = set<SPIRType>(ptr_type_id, OpTypePointer);
ptr_type = buff_type; ptr_type = buff_type;
ptr_type.op = spv::OpTypePointer; ptr_type.op = spv::OpTypePointer;
ptr_type.pointer = true; ptr_type.pointer = true;
@ -18092,12 +18080,12 @@ void CompilerMSL::add_argument_buffer_padding_image_type(SPIRType &struct_type,
if (!argument_buffer_padding_image_type_id) if (!argument_buffer_padding_image_type_id)
{ {
uint32_t base_type_id = ir.increase_bound_by(2); uint32_t base_type_id = ir.increase_bound_by(2);
auto &base_type = set<SPIRType>(base_type_id, spv::Op::OpTypeFloat); auto &base_type = set<SPIRType>(base_type_id, OpTypeFloat);
base_type.basetype = SPIRType::Float; base_type.basetype = SPIRType::Float;
base_type.width = 32; base_type.width = 32;
uint32_t img_type_id = base_type_id + 1; uint32_t img_type_id = base_type_id + 1;
auto &img_type = set<SPIRType>(img_type_id, spv::Op::OpTypeImage); auto &img_type = set<SPIRType>(img_type_id, OpTypeImage);
img_type.basetype = SPIRType::Image; img_type.basetype = SPIRType::Image;
img_type.storage = StorageClassUniformConstant; img_type.storage = StorageClassUniformConstant;
@ -18123,7 +18111,7 @@ void CompilerMSL::add_argument_buffer_padding_sampler_type(SPIRType &struct_type
if (!argument_buffer_padding_sampler_type_id) if (!argument_buffer_padding_sampler_type_id)
{ {
uint32_t samp_type_id = ir.increase_bound_by(1); uint32_t samp_type_id = ir.increase_bound_by(1);
auto &samp_type = set<SPIRType>(samp_type_id, spv::Op::OpTypeSampler); auto &samp_type = set<SPIRType>(samp_type_id, OpTypeSampler);
samp_type.basetype = SPIRType::Sampler; samp_type.basetype = SPIRType::Sampler;
samp_type.storage = StorageClassUniformConstant; samp_type.storage = StorageClassUniformConstant;
@ -18142,9 +18130,8 @@ void CompilerMSL::add_argument_buffer_padding_type(uint32_t mbr_type_id, SPIRTyp
if (count > 1) if (count > 1)
{ {
uint32_t ary_type_id = ir.increase_bound_by(1); uint32_t ary_type_id = ir.increase_bound_by(1);
auto &ary_type = set<SPIRType>(ary_type_id, spv::Op::OpTypeArray); auto &ary_type = set<SPIRType>(ary_type_id, get<SPIRType>(type_id));
ary_type = get<SPIRType>(type_id); ary_type.op = OpTypeArray;
ary_type.op = spv::Op::OpTypeArray;
ary_type.array.push_back(count); ary_type.array.push_back(count);
ary_type.array_size_literal.push_back(true); ary_type.array_size_literal.push_back(true);
ary_type.parent_type = type_id; ary_type.parent_type = type_id;

23
spirv_parser.cpp
View File

@ -568,9 +568,8 @@ void Parser::parse(const Instruction &instruction)
uint32_t vecsize = ops[2]; uint32_t vecsize = ops[2];
auto &base = get<SPIRType>(ops[1]); auto &base = get<SPIRType>(ops[1]);
auto &vecbase = set<SPIRType>(id, op); auto &vecbase = set<SPIRType>(id, base);
vecbase = base;
vecbase.op = op; vecbase.op = op;
vecbase.vecsize = vecsize; vecbase.vecsize = vecsize;
vecbase.self = id; vecbase.self = id;
@ -584,9 +583,8 @@ void Parser::parse(const Instruction &instruction)
uint32_t colcount = ops[2]; uint32_t colcount = ops[2];
auto &base = get<SPIRType>(ops[1]); auto &base = get<SPIRType>(ops[1]);
auto &matrixbase = set<SPIRType>(id, op); auto &matrixbase = set<SPIRType>(id, base);
matrixbase = base;
matrixbase.op = op; matrixbase.op = op;
matrixbase.columns = colcount; matrixbase.columns = colcount;
matrixbase.self = id; matrixbase.self = id;
@ -597,12 +595,10 @@ void Parser::parse(const Instruction &instruction)
case OpTypeArray: case OpTypeArray:
{ {
uint32_t id = ops[0]; uint32_t id = ops[0];
auto &arraybase = set<SPIRType>(id, op);
uint32_t tid = ops[1]; uint32_t tid = ops[1];
auto &base = get<SPIRType>(tid); auto &base = get<SPIRType>(tid);
auto &arraybase = set<SPIRType>(id, base);
arraybase = base;
arraybase.op = op; arraybase.op = op;
arraybase.parent_type = tid; arraybase.parent_type = tid;
@ -618,7 +614,9 @@ void Parser::parse(const Instruction &instruction)
arraybase.array_size_literal.push_back(literal); arraybase.array_size_literal.push_back(literal);
arraybase.array.push_back(literal ? c->scalar() : cid); arraybase.array.push_back(literal ? c->scalar() : cid);
// Do NOT set arraybase.self!
// .self resolves down to non-array/non-pointer type.
arraybase.self = base.self;
break; break;
} }
@ -627,19 +625,20 @@ void Parser::parse(const Instruction &instruction)
uint32_t id = ops[0]; uint32_t id = ops[0];
auto &base = get<SPIRType>(ops[1]); auto &base = get<SPIRType>(ops[1]);
auto &arraybase = set<SPIRType>(id, op); auto &arraybase = set<SPIRType>(id, base);
// We're copying type information into Array types, so we'll need a fixup for any physical pointer // We're copying type information into Array types, so we'll need a fixup for any physical pointer
// references. // references.
if (base.forward_pointer) if (base.forward_pointer)
forward_pointer_fixups.push_back({ id, ops[1] }); forward_pointer_fixups.push_back({ id, ops[1] });
arraybase = base;
arraybase.op = op; arraybase.op = op;
arraybase.array.push_back(0); arraybase.array.push_back(0);
arraybase.array_size_literal.push_back(true); arraybase.array_size_literal.push_back(true);
arraybase.parent_type = ops[1]; arraybase.parent_type = ops[1];
// Do NOT set arraybase.self!
// .self resolves down to non-array/non-pointer type.
arraybase.self = base.self;
break; break;
} }
@ -1032,7 +1031,7 @@ void Parser::parse(const Instruction &instruction)
{ {
uint32_t ids = ir.increase_bound_by(2); uint32_t ids = ir.increase_bound_by(2);
auto& type = set<SPIRType>(ids, spv::Op::OpTypeInt); auto &type = set<SPIRType>(ids, OpTypeInt);
type.basetype = SPIRType::Int; type.basetype = SPIRType::Int;
type.width = 32; type.width = 32;
auto &c = set<SPIRConstant>(ids + 1, ids); auto &c = set<SPIRConstant>(ids + 1, ids);

8
spirv_reflect.cpp
View File

@ -291,7 +291,7 @@ static bool naturally_emit_type(const SPIRType &type)
bool CompilerReflection::type_is_reference(const SPIRType &type) const bool CompilerReflection::type_is_reference(const SPIRType &type) const
{ {
// Physical pointers and arrays of physical pointers need to refer to the pointee's type. // Physical pointers and arrays of physical pointers need to refer to the pointee's type.
return type_is_top_level_physical_pointer(type) || return is_physical_pointer(type) ||
(type_is_array_of_pointers(type) && type.storage == StorageClassPhysicalStorageBuffer); (type_is_array_of_pointers(type) && type.storage == StorageClassPhysicalStorageBuffer);
} }
@ -341,7 +341,7 @@ void CompilerReflection::emit_type(uint32_t type_id, bool &emitted_open_tag)
json_stream->emit_json_key_object("_" + std::to_string(type_id)); json_stream->emit_json_key_object("_" + std::to_string(type_id));
json_stream->emit_json_key_value("name", name); json_stream->emit_json_key_value("name", name);
if (type_is_top_level_physical_pointer(type)) if (is_physical_pointer(type))
{ {
json_stream->emit_json_key_value("type", "_" + std::to_string(type.parent_type)); json_stream->emit_json_key_value("type", "_" + std::to_string(type.parent_type));
json_stream->emit_json_key_value("physical_pointer", true); json_stream->emit_json_key_value("physical_pointer", true);
@ -404,7 +404,7 @@ void CompilerReflection::emit_type_member(const SPIRType &type, uint32_t index)
void CompilerReflection::emit_type_array(const SPIRType &type) void CompilerReflection::emit_type_array(const SPIRType &type)
{ {
if (!type_is_top_level_physical_pointer(type) && !type.array.empty()) if (!is_physical_pointer(type) && !type.array.empty())
{ {
json_stream->emit_json_key_array("array"); json_stream->emit_json_key_array("array");
// Note that we emit the zeros here as a means of identifying // Note that we emit the zeros here as a means of identifying
@ -444,7 +444,7 @@ void CompilerReflection::emit_type_member_qualifiers(const SPIRType &type, uint3
if (dec.decoration_flags.get(DecorationRowMajor)) if (dec.decoration_flags.get(DecorationRowMajor))
json_stream->emit_json_key_value("row_major", true); json_stream->emit_json_key_value("row_major", true);
if (type_is_top_level_physical_pointer(membertype)) if (is_physical_pointer(membertype))
json_stream->emit_json_key_value("physical_pointer", true); json_stream->emit_json_key_value("physical_pointer", true);
} }
} }