From 833c5936b0e8d1e7f0e5b3c83eccfa35e98f1e06 Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Thu, 7 Dec 2023 12:18:35 +0100 Subject: [PATCH] Various nit fixes and improvements from review. --- spirv_common.hpp | 2 +- spirv_cross.cpp | 67 ++++--------- spirv_cross.hpp | 4 +- spirv_glsl.cpp | 70 +++++++------- spirv_hlsl.cpp | 12 +-- spirv_msl.cpp | 241 ++++++++++++++++++++++------------------------ spirv_parser.cpp | 23 +++-- spirv_reflect.cpp | 8 +- 8 files changed, 189 insertions(+), 238 deletions(-) diff --git a/spirv_common.hpp b/spirv_common.hpp index 77cb63b7..7fce7ae7 100644 --- a/spirv_common.hpp +++ b/spirv_common.hpp @@ -549,7 +549,7 @@ struct SPIRType : IVariant }; spv::Op op = spv::Op::OpNop; - SPIRType(spv::Op op_) : op(op_) {} + explicit SPIRType(spv::Op op_) : op(op_) {} enum BaseType { diff --git a/spirv_cross.cpp b/spirv_cross.cpp index e0a6398b..4ab985ef 100644 --- a/spirv_cross.cpp +++ b/spirv_cross.cpp @@ -627,20 +627,22 @@ bool Compiler::is_matrix(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 { - 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) { - if (type.array.empty()) - return false; - assert(type.array.size() == type.array_size_literal.size()); - return type.array_size_literal.back() && type.array.back() == 0; + return type.op == OpTypeRuntimeArray; } 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 combined_id = id + 2; auto &base = compiler.expression_type(image_id); - auto &type = compiler.set(type_id, spv::Op::OpTypeSampledImage); - auto &ptr_type = compiler.set(ptr_type_id, spv::Op::OpTypePointer); + auto &type = compiler.set(type_id, OpTypeSampledImage); + auto &ptr_type = compiler.set(ptr_type_id, OpTypePointer); type = base; 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. sampled_type = compiler.ir.increase_bound_by(1); - auto &type = compiler.set(sampled_type, spv::Op::OpTypeSampledImage); + auto &type = compiler.set(sampled_type, OpTypeSampledImage); type = compiler.expression_type(args[2]); type.self = sampled_type; 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. // We will probably have this type lying around, but it doesn't hurt to make duplicates for internal purposes. - auto &type = compiler.set(type_id, spv::Op::OpTypePointer); + auto &type = compiler.set(type_id, OpTypePointer); auto &base = compiler.get(sampled_type); type = base; type.pointer = true; @@ -3068,10 +3070,10 @@ VariableID Compiler::build_dummy_sampler_for_combined_images() auto ptr_type_id = offset + 1; auto var_id = offset + 2; - auto &sampler = set(type_id, spv::Op::OpTypeSampler); + auto &sampler = set(type_id, OpTypeSampler); sampler.basetype = SPIRType::Sampler; - auto &ptr_sampler = set(ptr_type_id, spv::Op::OpTypePointer); + auto &ptr_sampler = set(ptr_type_id, OpTypePointer); ptr_sampler = sampler; ptr_sampler.self = type_id; 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 { - if (!type_is_top_level_array(type)) + if (!is_array(type)) return false; // 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. auto *parent = &get(type.parent_type); - while (type_is_top_level_array(*parent)) + while (is_array(*parent)) parent = &get(parent->parent_type); - return type_is_top_level_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(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(type.parent_type).array.size(); - else - return !type.pointer; + return is_pointer(*parent); } bool Compiler::flush_phi_required(BlockID from, BlockID to) const diff --git a/spirv_cross.hpp b/spirv_cross.hpp index e1613321..9fe6c41c 100644 --- a/spirv_cross.hpp +++ b/spirv_cross.hpp @@ -684,6 +684,7 @@ protected: bool is_matrix(const SPIRType &type) const; bool is_array(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); uint32_t expression_type_id(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 &checked_ids); bool type_contains_recursion(const SPIRType &type); 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_top_level_block(const SPIRType &type) const; bool type_is_opaque_value(const SPIRType &type) const; diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index adfe3287..9b8c7d79 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -223,7 +223,7 @@ static const char *to_pls_layout(PlsFormat format) } } -static std::tuple pls_format_to_basetype(PlsFormat format) +static std::pair pls_format_to_basetype(PlsFormat format) { switch (format) { @@ -234,17 +234,17 @@ static std::tuple pls_format_to_basetype(PlsFormat case PlsRGB10A2: case PlsRGBA8: case PlsRG16: - return std::make_tuple(spv::OpTypeFloat, SPIRType::Float); + return std::make_pair(spv::OpTypeFloat, SPIRType::Float); case PlsRGBA8I: case PlsRG16I: - return std::make_tuple(spv::OpTypeInt, SPIRType::Int); + return std::make_pair(spv::OpTypeInt, SPIRType::Int); case PlsRGB10A2UI: case PlsRGBA8UI: case PlsRG16UI: 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, // and is 64-bit. - if (type_is_top_level_physical_pointer(type)) + if (is_physical_pointer(type)) { if (!type.pointer) 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 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; 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, // and is 64-bit. - if (type_is_top_level_physical_pointer(type)) + if (is_physical_pointer(type)) { if (!type.pointer) 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 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); @@ -1840,7 +1840,7 @@ bool CompilerGLSL::buffer_is_packing_standard(const SPIRType &type, BufferPackin } // 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_struct_member_array_stride(type, i)) { @@ -2489,7 +2489,7 @@ void CompilerGLSL::emit_buffer_block_flattened(const SPIRVariable &var) SPIRType::BaseType basic_type; if (get_common_basic_type(type, basic_type)) { - SPIRType tmp { spv::Op::OpTypeVector }; + SPIRType tmp { OpTypeVector }; tmp.basetype = basic_type; tmp.vecsize = 4; 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(member_type_id); auto array_type = member_type; 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_size_literal.push_back(true); @@ -3950,7 +3950,7 @@ void CompilerGLSL::emit_output_variable_initializer(const SPIRVariable &var) if (is_control_point) { uint32_t ids = ir.increase_bound_by(3); - auto& uint_type = set(ids, spv::Op::OpTypeInt); + auto &uint_type = set(ids, OpTypeInt); uint_type.basetype = SPIRType::UInt; uint_type.width = 32; set(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 && backend.boolean_in_struct_remapped_type != SPIRType::Boolean; - SPIRType tmp_type { spv::Op::OpNop }; + SPIRType tmp_type { OpNop }; if (remapped_boolean) { tmp_type = get(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++) { auto subexpr = join(base_expr, "[", convert_to_string(i), "]"); - if (!type_is_top_level_array(parent)) + if (!is_array(parent)) { if (remapped_boolean) 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 && 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 || (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(c.constant_type); - if (type_is_top_level_pointer(type)) + if (is_pointer(type)) { 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. // A particular CTS test hits this scenario. 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; // Allow Metal to use the array template to make arrays a value type bool needs_trailing_tracket = false; 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) + "{ "; } 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; - SPIRType tmp_type { spv::Op::OpNop }; + SPIRType tmp_type { OpNop }; if (inside_struct_scope && backend.boolean_in_struct_remapped_type != SPIRType::Boolean && @@ -5818,7 +5818,7 @@ string CompilerGLSL::constant_expression(const SPIRConstant &c, res += to_name(elem); 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. // 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. if (std::isnan(float_value) || std::isinf(float_value)) { - SPIRType type { spv::Op::OpTypeFloat }; + SPIRType type { OpTypeFloat }; type.basetype = SPIRType::Half; type.vecsize = 1; type.columns = 1; @@ -5932,7 +5932,7 @@ string CompilerGLSL::convert_half_to_string(const SPIRConstant &c, uint32_t col, } else { - SPIRType type { spv::Op::OpTypeFloat }; + SPIRType type { OpTypeFloat }; type.basetype = SPIRType::Half; type.vecsize = 1; type.columns = 1; @@ -5952,8 +5952,8 @@ string CompilerGLSL::convert_float_to_string(const SPIRConstant &c, uint32_t col // Use special representation. if (!is_legacy()) { - SPIRType out_type { spv::Op::OpTypeFloat }; - SPIRType in_type { spv::Op::OpTypeInt }; + SPIRType out_type { OpTypeFloat }; + SPIRType in_type { OpTypeInt }; out_type.basetype = SPIRType::Float; in_type.basetype = SPIRType::UInt; out_type.vecsize = 1; @@ -6022,8 +6022,8 @@ std::string CompilerGLSL::convert_double_to_string(const SPIRConstant &c, uint32 // Use special representation. if (!is_legacy()) { - SPIRType out_type { spv::Op::OpTypeFloat }; - SPIRType in_type { spv::Op::OpTypeInt }; + SPIRType out_type { OpTypeFloat }; + SPIRType in_type { OpTypeInt }; out_type.basetype = SPIRType::Double; in_type.basetype = SPIRType::UInt64; 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. // 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.vecsize = type0.vecsize; 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); 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.vecsize = 1; 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) { auto &type = expression_type(ptr); - if (!type_is_top_level_array(get_pointee_type(type))) + if (!is_array(get_pointee_type(type))) return false; if (!backend.array_is_value_type) @@ -10019,7 +10019,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice 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) { @@ -15367,8 +15367,8 @@ string CompilerGLSL::pls_decl(const PlsRemap &var) auto op_and_basetype = pls_format_to_basetype(var.format); - SPIRType type { std::get<0>(op_and_basetype) }; - type.basetype = std::get<1>(op_and_basetype); + SPIRType type { op_and_basetype.first }; + type.basetype = op_and_basetype.second; auto vecsize = pls_format_to_components(var.format); if (vecsize > 1) { @@ -17685,7 +17685,7 @@ bool CompilerGLSL::unroll_array_to_complex_store(uint32_t target_id, uint32_t so else array_expr = to_expression(type.array.back()); - SPIRType target_type { spv::Op::OpTypeInt }; + SPIRType target_type { OpTypeInt }; target_type.basetype = SPIRType::Int; 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, ";"); else if (is_sample_mask) { - SPIRType target_type { spv::Op::OpTypeInt }; + SPIRType target_type { OpTypeInt }; target_type.basetype = SPIRType::Int; statement(new_expr, "[i] = ", bitcast_expression(target_type, type.basetype, join(expr, "[i]")), ";"); } diff --git a/spirv_hlsl.cpp b/spirv_hlsl.cpp index 5e30324d..9b834936 100644 --- a/spirv_hlsl.cpp +++ b/spirv_hlsl.cpp @@ -2432,7 +2432,7 @@ void CompilerHLSL::analyze_meshlet_writes() uint32_t op_ptr = op_type + 2; uint32_t op_var = op_type + 3; - auto &type = set(op_type, spv::Op::OpTypeStruct); + auto &type = set(op_type, OpTypeStruct); type.basetype = SPIRType::Struct; set_name(op_type, block_name); set_decoration(op_type, DecorationBlock); @@ -4508,7 +4508,7 @@ void CompilerHLSL::read_access_chain(string *expr, const string &lhs, const SPIR { auto &type = get(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.vecsize = type.vecsize; 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 int_type_id = id + 1; - SPIRType int_type { spv::Op::OpTypeInt }; + SPIRType int_type { OpTypeInt }; int_type.basetype = SPIRType::Int; int_type.width = 32; set(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. 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.vecsize = type.vecsize; 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 variable_id = offset + 3; - SPIRType uint_type { spv::Op::OpTypeVector }; + SPIRType uint_type { OpTypeVector }; uint_type.basetype = SPIRType::UInt; uint_type.width = 32; uint_type.vecsize = 3; uint_type.columns = 1; set(uint_type_id, uint_type); - SPIRType block_type { spv::Op::OpTypeStruct }; + SPIRType block_type { OpTypeStruct }; block_type.basetype = SPIRType::Struct; block_type.member_types.push_back(uint_type_id); set(block_type_id, block_type); diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 0a8769b0..40c73a0c 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -478,21 +478,20 @@ void CompilerMSL::build_implicit_builtins() uint32_t var_id = offset + 2; // Create gl_FragCoord. - SPIRType vec4_type { spv::Op::OpTypeVector }; + SPIRType vec4_type { OpTypeVector }; vec4_type.basetype = SPIRType::Float; vec4_type.width = 32; vec4_type.vecsize = 4; set(type_id, vec4_type); - SPIRType vec4_type_ptr { spv::Op::OpTypePointer }; - vec4_type_ptr = vec4_type; - vec4_type_ptr.op = spv::Op::OpTypePointer; + SPIRType vec4_type_ptr = vec4_type; + vec4_type_ptr.op = OpTypePointer; vec4_type_ptr.pointer = true; vec4_type_ptr.pointer_depth++; vec4_type_ptr.parent_type = type_id; vec4_type_ptr.storage = StorageClassInput; auto &ptr_type = set(type_ptr_id, vec4_type_ptr); - ptr_type.self = type_ptr_id; + ptr_type.self = type_id; set(var_id, type_ptr_id, StorageClassInput); set_decoration(var_id, DecorationBuiltIn, BuiltInFragCoord); @@ -507,15 +506,14 @@ void CompilerMSL::build_implicit_builtins() uint32_t var_id = offset + 1; // Create gl_Layer. - SPIRType uint_type_ptr { spv::Op::OpTypePointer }; - uint_type_ptr = get_uint_type(); - uint_type_ptr.op = spv::Op::OpTypePointer; + SPIRType uint_type_ptr = get_uint_type(); + uint_type_ptr.op = OpTypePointer; uint_type_ptr.pointer = true; uint_type_ptr.pointer_depth++; uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.storage = StorageClassInput; auto &ptr_type = set(type_ptr_id, uint_type_ptr); - ptr_type.self = type_ptr_id; + ptr_type.self = get_uint_type_id(); set(var_id, type_ptr_id, StorageClassInput); set_decoration(var_id, DecorationBuiltIn, BuiltInLayer); @@ -530,15 +528,14 @@ void CompilerMSL::build_implicit_builtins() uint32_t var_id = offset + 1; // Create gl_ViewIndex. - SPIRType uint_type_ptr { spv::Op::OpTypePointer }; - uint_type_ptr = get_uint_type(); - uint_type_ptr.op = spv::Op::OpTypePointer; + SPIRType uint_type_ptr = get_uint_type(); + uint_type_ptr.op = OpTypePointer; uint_type_ptr.pointer = true; uint_type_ptr.pointer_depth++; uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.storage = StorageClassInput; auto &ptr_type = set(type_ptr_id, uint_type_ptr); - ptr_type.self = type_ptr_id; + ptr_type.self = get_uint_type_id(); set(var_id, type_ptr_id, StorageClassInput); set_decoration(var_id, DecorationBuiltIn, BuiltInViewIndex); @@ -554,15 +551,14 @@ void CompilerMSL::build_implicit_builtins() uint32_t var_id = offset + 1; // Create gl_SampleID. - SPIRType uint_type_ptr { spv::Op::OpTypePointer }; - uint_type_ptr = get_uint_type(); - uint_type_ptr.op = spv::Op::OpTypePointer; + SPIRType uint_type_ptr = get_uint_type(); + uint_type_ptr.op = OpTypePointer; uint_type_ptr.pointer = true; uint_type_ptr.pointer_depth++; uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.storage = StorageClassInput; auto &ptr_type = set(type_ptr_id, uint_type_ptr); - ptr_type.self = type_ptr_id; + ptr_type.self = get_uint_type_id(); set(var_id, type_ptr_id, StorageClassInput); 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); - SPIRType uint_type_ptr { spv::Op::OpTypePointer }; - uint_type_ptr = get_uint_type(); - uint_type_ptr.op = spv::Op::OpTypePointer; + SPIRType uint_type_ptr = get_uint_type(); + uint_type_ptr.op = OpTypePointer; uint_type_ptr.pointer = true; uint_type_ptr.pointer_depth++; uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.storage = StorageClassInput; auto &ptr_type = set(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) { @@ -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 // gl_Layer is an output in vertex-pipeline shaders. uint32_t type_ptr_out_id = ir.increase_bound_by(2); - SPIRType uint_type_ptr_out { spv::Op::OpTypePointer }; - uint_type_ptr_out = get_uint_type(); - uint_type_ptr.op = spv::Op::OpTypePointer; + SPIRType uint_type_ptr_out = get_uint_type(); + uint_type_ptr.op = OpTypePointer; uint_type_ptr_out.pointer = true; uint_type_ptr_out.pointer_depth++; 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); - SPIRType uint_type_ptr { spv::Op::OpTypePointer }; - uint_type_ptr = get_uint_type(); - uint_type_ptr.op = spv::Op::OpTypePointer; + SPIRType uint_type_ptr = get_uint_type(); + uint_type_ptr.op = OpTypePointer; uint_type_ptr.pointer = true; uint_type_ptr.pointer_depth++; uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.storage = StorageClassInput; auto &ptr_type = set(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) { @@ -730,15 +723,14 @@ void CompilerMSL::build_implicit_builtins() uint32_t var_id = offset + 1; // Create gl_SubgroupInvocationID. - SPIRType uint_type_ptr { spv::Op::OpTypePointer }; - uint_type_ptr = get_uint_type(); - uint_type_ptr.op = spv::Op::OpTypePointer; + SPIRType uint_type_ptr = get_uint_type(); + uint_type_ptr.op = OpTypePointer; uint_type_ptr.pointer = true; uint_type_ptr.pointer_depth++; uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.storage = StorageClassInput; auto &ptr_type = set(type_ptr_id, uint_type_ptr); - ptr_type.self = type_ptr_id; + ptr_type.self = get_uint_type_id(); set(var_id, type_ptr_id, StorageClassInput); set_decoration(var_id, DecorationBuiltIn, BuiltInSubgroupLocalInvocationId); @@ -753,15 +745,14 @@ void CompilerMSL::build_implicit_builtins() uint32_t var_id = offset + 1; // Create gl_SubgroupSize. - SPIRType uint_type_ptr { spv::Op::OpTypePointer }; - uint_type_ptr = get_uint_type(); - uint_type_ptr.op = spv::Op::OpTypePointer; + SPIRType uint_type_ptr = get_uint_type(); + uint_type_ptr.op = OpTypePointer; uint_type_ptr.pointer = true; uint_type_ptr.pointer_depth++; uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.storage = StorageClassInput; auto &ptr_type = set(type_ptr_id, uint_type_ptr); - ptr_type.self = type_ptr_id; + ptr_type.self = get_uint_type_id(); set(var_id, type_ptr_id, StorageClassInput); set_decoration(var_id, DecorationBuiltIn, BuiltInSubgroupSize); @@ -813,16 +804,15 @@ void CompilerMSL::build_implicit_builtins() uint32_t var_id = offset + 1; // Create gl_SampleMask. - SPIRType uint_type_ptr_out { spv::Op::OpTypePointer }; - uint_type_ptr_out = get_uint_type(); - uint_type_ptr_out.op = spv::Op::OpTypePointer; + SPIRType uint_type_ptr_out = get_uint_type(); + uint_type_ptr_out.op = OpTypePointer; uint_type_ptr_out.pointer = true; uint_type_ptr_out.pointer_depth++; uint_type_ptr_out.parent_type = get_uint_type_id(); uint_type_ptr_out.storage = StorageClassOutput; auto &ptr_out_type = set(offset, uint_type_ptr_out); - ptr_out_type.self = offset; + ptr_out_type.self = get_uint_type_id(); set(var_id, offset, StorageClassOutput); set_decoration(var_id, DecorationBuiltIn, BuiltInSampleMask); builtin_sample_mask_id = var_id; @@ -837,14 +827,13 @@ void CompilerMSL::build_implicit_builtins() uint32_t var_id = offset + 2; // Create gl_HelperInvocation. - SPIRType bool_type { spv::Op::OpTypeBool }; + SPIRType bool_type { OpTypeBool }; bool_type.basetype = SPIRType::Boolean; bool_type.width = 8; bool_type.vecsize = 1; set(type_id, bool_type); - SPIRType bool_type_ptr_in { spv::Op::OpTypePointer }; - bool_type_ptr_in = bool_type; + SPIRType bool_type_ptr_in = bool_type; bool_type_ptr_in.op = spv::OpTypePointer; bool_type_ptr_in.pointer = true; bool_type_ptr_in.pointer_depth++; @@ -852,7 +841,7 @@ void CompilerMSL::build_implicit_builtins() bool_type_ptr_in.storage = StorageClassInput; auto &ptr_in_type = set(type_ptr_id, bool_type_ptr_in); - ptr_in_type.self = type_ptr_id; + ptr_in_type.self = type_id; set(var_id, type_ptr_id, StorageClassInput); set_decoration(var_id, DecorationBuiltIn, BuiltInHelperInvocation); builtin_helper_invocation_id = var_id; @@ -866,16 +855,15 @@ void CompilerMSL::build_implicit_builtins() uint32_t var_id = offset + 1; // Create gl_LocalInvocationIndex. - SPIRType uint_type_ptr { spv::Op::OpTypePointer }; - uint_type_ptr = get_uint_type(); - uint_type_ptr.op = spv::Op::OpTypePointer; + SPIRType uint_type_ptr = get_uint_type(); + uint_type_ptr.op = OpTypePointer; uint_type_ptr.pointer = true; uint_type_ptr.pointer_depth++; uint_type_ptr.parent_type = get_uint_type_id(); uint_type_ptr.storage = StorageClassInput; auto &ptr_type = set(type_ptr_id, uint_type_ptr); - ptr_type.self = type_ptr_id; + ptr_type.self = get_uint_type_id(); set(var_id, type_ptr_id, StorageClassInput); set_decoration(var_id, DecorationBuiltIn, BuiltInLocalInvocationIndex); builtin_local_invocation_index_id = var_id; @@ -891,14 +879,14 @@ void CompilerMSL::build_implicit_builtins() // Create gl_WorkgroupSize. uint32_t type_id = build_extended_vector_type(get_uint_type_id(), 3); SPIRType uint_type_ptr = get(type_id); - uint_type_ptr.op = spv::Op::OpTypePointer; + uint_type_ptr.op = OpTypePointer; uint_type_ptr.pointer = true; uint_type_ptr.pointer_depth++; uint_type_ptr.parent_type = type_id; uint_type_ptr.storage = StorageClassInput; auto &ptr_type = set(type_ptr_id, uint_type_ptr); - ptr_type.self = type_ptr_id; + ptr_type.self = type_id; set(var_id, type_ptr_id, StorageClassInput); set_decoration(var_id, DecorationBuiltIn, BuiltInWorkgroupSize); builtin_workgroup_size_id = var_id; @@ -999,21 +987,20 @@ void CompilerMSL::build_implicit_builtins() uint32_t var_id = offset + 2; // Create gl_Position. - SPIRType vec4_type { spv::Op::OpTypeVector }; + SPIRType vec4_type { OpTypeVector }; vec4_type.basetype = SPIRType::Float; vec4_type.width = 32; vec4_type.vecsize = 4; set(type_id, vec4_type); - SPIRType vec4_type_ptr { spv::Op::OpTypePointer }; - vec4_type_ptr = vec4_type; - vec4_type_ptr.op = spv::Op::OpTypePointer; + SPIRType vec4_type_ptr = vec4_type; + vec4_type_ptr.op = OpTypePointer; vec4_type_ptr.pointer = true; vec4_type_ptr.pointer_depth++; vec4_type_ptr.parent_type = type_id; vec4_type_ptr.storage = StorageClassOutput; auto &ptr_type = set(type_ptr_id, vec4_type_ptr); - ptr_type.self = type_ptr_id; + ptr_type.self = type_id; set(var_id, type_ptr_id, StorageClassOutput); 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. 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_depth++; 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); - SPIRType type { spv::Op::OpTypeInt }; + SPIRType type { OpTypeInt }; type.basetype = SPIRType::UInt; type.width = 32; set(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, // so that we will get the appropriate address space when declaring these builtins. auto &ptr = set(ptr_type_id, get(mbr_type_id)); - ptr.self = ptr_type_id; + ptr.self = mbr_type_id; ptr.storage = var.storage; ptr.pointer = true; ptr.pointer_depth++; @@ -2348,23 +2335,25 @@ uint32_t CompilerMSL::build_extended_vector_type(uint32_t type_id, uint32_t comp { assert(components > 1); uint32_t new_type_id = ir.increase_bound_by(1); - auto &old_type = get(type_id); - SPIRType* old_ptr_t = nullptr; - SPIRType* old_array_t = nullptr; - if (is_pointer(old_type)) + const auto *p_old_type = &get(type_id); + const SPIRType *old_ptr_t = nullptr; + const SPIRType *old_array_t = nullptr; + + if (is_pointer(*p_old_type)) { - old_ptr_t = &old_type; - 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); + old_ptr_t = p_old_type; + p_old_type = &get_pointee_type(*old_ptr_t); } - auto *type = &set(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(new_type_id, *p_old_type); assert(is_scalar(*type) || is_vector(*type)); - type->op = spv::Op::OpTypeVector; + type->op = OpTypeVector; type->vecsize = components; if (basetype != SPIRType::Unknown) 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); type = &set(array_type_id, *type); + type->op = OpTypeArray; type->parent_type = new_type_id; type->array = old_array_t->array; 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); type = &set(ptr_type_id, *type); - type->self = new_type_id; + type->op = OpTypePointer; type->parent_type = new_type_id; type->storage = old_ptr_t->storage; 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. uint32_t next_id = ir.increase_bound_by(3); uint32_t ib_type_id = next_id++; - auto &ib_type = set(ib_type_id, spv::Op::OpTypeStruct); + auto &ib_type = set(ib_type_id, OpTypeStruct); ib_type.basetype = SPIRType::Struct; ib_type.storage = storage; 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 var_id = offset + 4; - SPIRType type { spv::Op::OpTypeInt }; + SPIRType type { OpTypeInt }; switch (input.second.format) { case MSL_SHADER_VARIABLE_FORMAT_UINT16: @@ -4215,26 +4205,26 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch) set(type_id, type); if (input.second.vecsize > 1) { - type.op = spv::Op::OpTypeVector; + type.op = OpTypeVector; type.vecsize = input.second.vecsize; set(vec_type_id, type); type_id = vec_type_id; } - type.op = spv::Op::OpTypeArray; + type.op = OpTypeArray; type.array.push_back(0); type.array_size_literal.push_back(true); type.parent_type = type_id; set(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_depth++; type.parent_type = array_type_id; type.storage = storage; auto &ptr_type = set(ptr_type_id, type); - ptr_type.self = ptr_type_id; + ptr_type.self = array_type_id; auto &fake_var = set(var_id, ptr_type_id, storage); 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 var_id = offset + 4; - SPIRType type { spv::Op::OpTypeInt }; + SPIRType type { OpTypeInt }; switch (output.second.format) { case MSL_SHADER_VARIABLE_FORMAT_UINT16: @@ -4283,7 +4273,7 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch) set(type_id, type); if (output.second.vecsize > 1) { - type.op = spv::Op::OpTypeVector; + type.op = OpTypeVector; type.vecsize = output.second.vecsize; set(vec_type_id, type); type_id = vec_type_id; @@ -4291,14 +4281,14 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch) if (is_tesc_shader()) { - type.op = spv::Op::OpTypeArray; + type.op = OpTypeArray; type.array.push_back(0); type.array_size_literal.push_back(true); - type.parent_type = vec_type_id; + type.parent_type = type_id; set(array_type_id, type); } - type.op = spv::Op::OpTypePointer; + type.op = OpTypePointer; type.pointer = true; type.pointer_depth++; 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. uint32_t ib_ptr_type_id = next_id++; auto &ib_ptr_type = set(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.pointer = true; 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. uint32_t CompilerMSL::ensure_correct_builtin_type(uint32_t type_id, BuiltIn builtin) { - auto& type = get(type_id); - auto pointee_type = type; - if (is_pointer(type)) - pointee_type = get(type.parent_type); + auto &type = get(type_id); + auto &pointee_type = get_pointee_type(type); if ((builtin == BuiltInSampleMask && is_array(pointee_type)) || ((builtin == BuiltInLayer || builtin == BuiltInViewportIndex || builtin == BuiltInFragStencilRefEXT) && 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++; - auto &base_type = set(base_type_id, spv::Op::OpTypeInt); + auto &base_type = set(base_type_id, OpTypeInt); base_type.basetype = SPIRType::UInt; base_type.width = 32; - if (!type.pointer) + if (!type_is_pointer(type)) return base_type_id; uint32_t ptr_type_id = next_id++; - auto &ptr_type = set(ptr_type_id, spv::OpTypePointer); - ptr_type = base_type; + auto &ptr_type = set(ptr_type_id, base_type); ptr_type.op = spv::OpTypePointer; ptr_type.pointer = true; ptr_type.pointer_depth++; @@ -4949,7 +4937,7 @@ void CompilerMSL::ensure_member_packing_rules_msl(SPIRType &ib_type, uint32_t in { type.columns = 1; assert(type.array.empty()); - type.op = spv::Op::OpTypeArray; + type.op = OpTypeArray; type.array.push_back(1); 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.columns = 1; assert(type.array.empty()); - type.op = spv::Op::OpTypeArray; + type.op = OpTypeArray; type.array.push_back(1); 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) && 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. 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. // 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. - 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); auto name = to_name(c.self); @@ -7539,7 +7527,7 @@ void CompilerMSL::declare_complex_constant_arrays() return; auto &type = this->get(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); 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. // 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) || 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 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; auto *var = maybe_get(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) { - SPIRType t { components > 1 ? spv::Op::OpTypeVector : spv::Op::OpTypeFloat }; + SPIRType t { components > 1 ? OpTypeVector : OpTypeFloat }; t.basetype = SPIRType::Float; t.vecsize = components; t.columns = 1; @@ -11980,7 +11968,7 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_ if (is_matrix(physical_type)) 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; // 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 &var_type = get(arg.type); 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. // 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) 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. if (type_is_msl_framebuffer_fetch(type)) 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)) { // 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; // 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); @@ -15088,7 +15075,7 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id, bool member) // the C-style nesting works right. // FIXME: This is somewhat of a hack. 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; 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]); if (var != nullptr) { - auto &type = compiler.get(var->basetype); - if (!is_runtime_size_array(type)) + if (!compiler.is_var_runtime_size_array(*var)) compiler.buffers_requiring_array_length.insert(var->self); } 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. 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); else 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; } - if (type_is_top_level_array(expr_type) && builtin == BuiltInSampleMask) + if (is_array(expr_type) && builtin == BuiltInSampleMask) { // Needs special handling. 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) { - 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. // 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. 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); } @@ -17786,6 +17772,7 @@ void CompilerMSL::analyze_argument_buffers() // Create a buffer to hold extra data, including the swizzle constants. SPIRType uint_type_pointer = get_uint_type(); + uint_type_pointer.op = OpTypePointer; uint_type_pointer.pointer = true; uint_type_pointer.pointer_depth++; 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; argument_buffer_ids[desc_set] = next_id; - auto &buffer_type = set(type_id, spv::Op::OpTypeStruct); + auto &buffer_type = set(type_id, OpTypeStruct); buffer_type.basetype = SPIRType::Struct; @@ -17858,7 +17845,7 @@ void CompilerMSL::analyze_argument_buffers() set_name(type_id, join("spvDescriptorSetBuffer", desc_set)); - auto &ptr_type = set(ptr_type_id, spv::Op::OpTypePointer); + auto &ptr_type = set(ptr_type_id, OpTypePointer); ptr_type = buffer_type; ptr_type.op = spv::OpTypePointer; ptr_type.pointer = true; @@ -17943,14 +17930,14 @@ void CompilerMSL::analyze_argument_buffers() bool type_is_array = !type.array.empty(); uint32_t sampler_type_id = ir.increase_bound_by(type_is_array ? 2 : 1); - auto &new_sampler_type = set(sampler_type_id, spv::Op::OpTypeSampler); + auto &new_sampler_type = set(sampler_type_id, OpTypeSampler); new_sampler_type.basetype = SPIRType::Sampler; new_sampler_type.storage = StorageClassUniformConstant; if (type_is_array) { uint32_t sampler_type_array_id = sampler_type_id + 1; - auto &sampler_type_array = set(sampler_type_array_id, spv::Op::OpTypeArray); + auto &sampler_type_array = set(sampler_type_array_id, OpTypeArray); sampler_type_array = new_sampler_type; sampler_type_array.array = type.array; 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 type_ptr_id = offset + 1; - SPIRType atomic_type { spv::Op::OpTypeInt }; + SPIRType atomic_type { OpTypeInt }; atomic_type.basetype = SPIRType::AtomicCounter; atomic_type.width = 32; atomic_type.vecsize = 1; set(atomic_type_id, atomic_type); + atomic_type.op = OpTypePointer; atomic_type.pointer = true; atomic_type.pointer_depth++; atomic_type.parent_type = atomic_type_id; atomic_type.storage = StorageClassStorageBuffer; auto &atomic_ptr_type = set(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); } @@ -18067,12 +18055,12 @@ void CompilerMSL::add_argument_buffer_padding_buffer_type(SPIRType &struct_type, if (!argument_buffer_padding_buffer_type_id) { uint32_t buff_type_id = ir.increase_bound_by(2); - auto &buff_type = set(buff_type_id, spv::Op::OpNop); + auto &buff_type = set(buff_type_id, OpNop); buff_type.basetype = rez_bind.basetype; buff_type.storage = StorageClassUniformConstant; uint32_t ptr_type_id = buff_type_id + 1; - auto &ptr_type = set(ptr_type_id, spv::Op::OpTypePointer); + auto &ptr_type = set(ptr_type_id, OpTypePointer); ptr_type = buff_type; ptr_type.op = spv::OpTypePointer; 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) { uint32_t base_type_id = ir.increase_bound_by(2); - auto &base_type = set(base_type_id, spv::Op::OpTypeFloat); + auto &base_type = set(base_type_id, OpTypeFloat); base_type.basetype = SPIRType::Float; base_type.width = 32; uint32_t img_type_id = base_type_id + 1; - auto &img_type = set(img_type_id, spv::Op::OpTypeImage); + auto &img_type = set(img_type_id, OpTypeImage); img_type.basetype = SPIRType::Image; 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) { uint32_t samp_type_id = ir.increase_bound_by(1); - auto &samp_type = set(samp_type_id, spv::Op::OpTypeSampler); + auto &samp_type = set(samp_type_id, OpTypeSampler); samp_type.basetype = SPIRType::Sampler; samp_type.storage = StorageClassUniformConstant; @@ -18142,9 +18130,8 @@ void CompilerMSL::add_argument_buffer_padding_type(uint32_t mbr_type_id, SPIRTyp if (count > 1) { uint32_t ary_type_id = ir.increase_bound_by(1); - auto &ary_type = set(ary_type_id, spv::Op::OpTypeArray); - ary_type = get(type_id); - ary_type.op = spv::Op::OpTypeArray; + auto &ary_type = set(ary_type_id, get(type_id)); + ary_type.op = OpTypeArray; ary_type.array.push_back(count); ary_type.array_size_literal.push_back(true); ary_type.parent_type = type_id; diff --git a/spirv_parser.cpp b/spirv_parser.cpp index d403bbd9..6108dbb6 100644 --- a/spirv_parser.cpp +++ b/spirv_parser.cpp @@ -568,9 +568,8 @@ void Parser::parse(const Instruction &instruction) uint32_t vecsize = ops[2]; auto &base = get(ops[1]); - auto &vecbase = set(id, op); + auto &vecbase = set(id, base); - vecbase = base; vecbase.op = op; vecbase.vecsize = vecsize; vecbase.self = id; @@ -584,9 +583,8 @@ void Parser::parse(const Instruction &instruction) uint32_t colcount = ops[2]; auto &base = get(ops[1]); - auto &matrixbase = set(id, op); + auto &matrixbase = set(id, base); - matrixbase = base; matrixbase.op = op; matrixbase.columns = colcount; matrixbase.self = id; @@ -597,12 +595,10 @@ void Parser::parse(const Instruction &instruction) case OpTypeArray: { uint32_t id = ops[0]; - auto &arraybase = set(id, op); - uint32_t tid = ops[1]; auto &base = get(tid); + auto &arraybase = set(id, base); - arraybase = base; arraybase.op = op; arraybase.parent_type = tid; @@ -618,7 +614,9 @@ void Parser::parse(const Instruction &instruction) arraybase.array_size_literal.push_back(literal); 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; } @@ -627,19 +625,20 @@ void Parser::parse(const Instruction &instruction) uint32_t id = ops[0]; auto &base = get(ops[1]); - auto &arraybase = set(id, op); + auto &arraybase = set(id, base); // We're copying type information into Array types, so we'll need a fixup for any physical pointer // references. if (base.forward_pointer) forward_pointer_fixups.push_back({ id, ops[1] }); - arraybase = base; arraybase.op = op; arraybase.array.push_back(0); arraybase.array_size_literal.push_back(true); arraybase.parent_type = ops[1]; - // Do NOT set arraybase.self! + + // .self resolves down to non-array/non-pointer type. + arraybase.self = base.self; break; } @@ -1032,7 +1031,7 @@ void Parser::parse(const Instruction &instruction) { uint32_t ids = ir.increase_bound_by(2); - auto& type = set(ids, spv::Op::OpTypeInt); + auto &type = set(ids, OpTypeInt); type.basetype = SPIRType::Int; type.width = 32; auto &c = set(ids + 1, ids); diff --git a/spirv_reflect.cpp b/spirv_reflect.cpp index 9fcd3bc0..b0277372 100644 --- a/spirv_reflect.cpp +++ b/spirv_reflect.cpp @@ -291,7 +291,7 @@ static bool naturally_emit_type(const SPIRType &type) bool CompilerReflection::type_is_reference(const SPIRType &type) const { // 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); } @@ -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_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("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) { - 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"); // 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)) 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); } }