From c67a28fbe844810f4d395059bd00ea413ba85cef Mon Sep 17 00:00:00 2001 From: Jacob Lifshay Date: Mon, 14 Aug 2017 16:44:34 -0700 Subject: [PATCH] working on implementing vertex shader wrapper still need to implement copying the shader output to the output buffer. running the generated vertex shader should work. --- src/pipeline/pipeline.h | 2 +- src/spirv_to_llvm/spirv_to_llvm.cpp | 667 +++++++++++++++++++++++++++- src/spirv_to_llvm/spirv_to_llvm.h | 9 +- 3 files changed, 660 insertions(+), 18 deletions(-) diff --git a/src/pipeline/pipeline.h b/src/pipeline/pipeline.h index c31a8dd..c50bc08 100644 --- a/src/pipeline/pipeline.h +++ b/src/pipeline/pipeline.h @@ -69,7 +69,7 @@ class Graphics_pipeline final : public Pipeline public: #warning finish adding draw function parameters typedef void (*Vertex_shader_function)(std::uint32_t vertex_start_index, - std::uint32_t vertex_count, + std::uint32_t vertex_end_index, std::uint32_t instance_id, void *output_buffer); diff --git a/src/spirv_to_llvm/spirv_to_llvm.cpp b/src/spirv_to_llvm/spirv_to_llvm.cpp index 09ea92e..d23c946 100644 --- a/src/spirv_to_llvm/spirv_to_llvm.cpp +++ b/src/spirv_to_llvm/spirv_to_llvm.cpp @@ -798,33 +798,651 @@ public: no_instruction_index); outputs_member = io_struct->add_member(Struct_type_descriptor::Member({}, outputs_struct)); } - std::string generate_entry_function(Op_entry_point_state &entry_point) + std::string generate_entry_function(Op_entry_point_state &entry_point, + ::LLVMValueRef main_function) { - ::LLVMValueRef function = nullptr; + ::LLVMValueRef entry_function = nullptr; switch(execution_model) { case spirv::Execution_model::vertex: { - typedef void (*Vertex_shader_function)(std::uint32_t vertex_start_index, - std::uint32_t vertex_count, + typedef std::uint32_t Vertex_index_type; + auto llvm_vertex_index_type = + llvm_wrapper::Create_llvm_type()(context); + typedef void (*Vertex_shader_function)(Vertex_index_type vertex_start_index, + Vertex_index_type vertex_end_index, std::uint32_t instance_id, void *output_buffer); constexpr std::size_t arg_vertex_start_index = 0; - constexpr std::size_t arg_vertex_count = 1; + constexpr std::size_t arg_vertex_end_index = 1; constexpr std::size_t arg_instance_id = 2; constexpr std::size_t arg_output_buffer = 3; static_assert(std::is_same::value, "vertex shader function signature mismatch"); auto function_type = llvm_wrapper::Create_llvm_type()(context); - function = ::LLVMAddFunction( + entry_function = ::LLVMAddFunction( module.get(), get_prefixed_name("vertex_entry_point", true).c_str(), function_type); - llvm_wrapper::Module::set_function_target_machine(function, target_machine); - static_cast(arg_vertex_start_index); - static_cast(arg_vertex_count); - static_cast(arg_instance_id); - static_cast(arg_output_buffer); -#warning finish implementing vertex execution model + llvm_wrapper::Module::set_function_target_machine(entry_function, target_machine); + ::LLVMSetValueName(::LLVMGetParam(entry_function, arg_vertex_start_index), + "vertex_start_index"); + ::LLVMSetValueName(::LLVMGetParam(entry_function, arg_vertex_end_index), + "vertex_end_index"); + ::LLVMSetValueName(::LLVMGetParam(entry_function, arg_instance_id), "instance_id"); + ::LLVMSetValueName(::LLVMGetParam(entry_function, arg_output_buffer), "output_buffer"); + auto entry_block = ::LLVMAppendBasicBlockInContext(context, entry_function, "entry"); + auto loop_block = ::LLVMAppendBasicBlockInContext(context, entry_function, "loop"); + auto exit_block = ::LLVMAppendBasicBlockInContext(context, entry_function, "exit"); + ::LLVMPositionBuilderAtEnd(builder.get(), entry_block); + auto io_struct_type = io_struct->get_or_make_type(); + auto io_struct_pointer = + ::LLVMBuildAlloca(builder.get(), io_struct_type.type, "io_struct"); + ::LLVMSetAlignment( + ::LLVMBuildStore( + builder.get(), ::LLVMConstNull(io_struct_type.type), io_struct_pointer), + io_struct_type.alignment); + auto start_loop_condition = + ::LLVMBuildICmp(builder.get(), + ::LLVMIntULT, + ::LLVMGetParam(entry_function, arg_vertex_start_index), + ::LLVMGetParam(entry_function, arg_vertex_end_index), + "start_loop_condition"); + ::LLVMBuildCondBr(builder.get(), start_loop_condition, loop_block, exit_block); + ::LLVMPositionBuilderAtEnd(builder.get(), loop_block); + auto vertex_index = + ::LLVMBuildPhi(builder.get(), + llvm_wrapper::Create_llvm_type()(context), + "vertex_index"); + auto next_vertex_index = + ::LLVMBuildNUWAdd(builder.get(), + vertex_index, + ::LLVMConstInt(llvm_vertex_index_type, 1, false), + "next_vertex_index"); + constexpr std::size_t vertex_index_incoming_count = 2; + ::LLVMValueRef vertex_index_incoming_values[vertex_index_incoming_count] = { + next_vertex_index, ::LLVMGetParam(entry_function, arg_vertex_start_index), + }; + ::LLVMBasicBlockRef vertex_index_incoming_blocks[vertex_index_incoming_count] = { + loop_block, entry_block, + }; + ::LLVMAddIncoming(vertex_index, + vertex_index_incoming_values, + vertex_index_incoming_blocks, + vertex_index_incoming_count); + for(auto &member : io_struct->get_members(true)) + { + if(member.type == inputs_struct) + { + auto inputs_struct_pointer = ::LLVMBuildStructGEP( + builder.get(), io_struct_pointer, member.llvm_member_index, "inputs"); + for(auto &input_member : inputs_struct->get_members(true)) + { + auto input_pointer = ::LLVMBuildStructGEP(builder.get(), + inputs_struct_pointer, + input_member.llvm_member_index, + "input"); + util::optional built_in; + static_cast(input_pointer); + for(auto &decoration : input_member.decorations) + { + switch(decoration.value) + { + case Decoration::relaxed_precision: +#warning finish implementing Decoration::relaxed_precision + break; + case Decoration::spec_id: +#warning finish implementing Decoration::spec_id + break; + case Decoration::block: +#warning finish implementing Decoration::block + break; + case Decoration::buffer_block: +#warning finish implementing Decoration::buffer_block + break; + case Decoration::row_major: +#warning finish implementing Decoration::row_major + break; + case Decoration::col_major: +#warning finish implementing Decoration::col_major + break; + case Decoration::array_stride: +#warning finish implementing Decoration::array_stride + break; + case Decoration::matrix_stride: +#warning finish implementing Decoration::matrix_stride + break; + case Decoration::glsl_shared: +#warning finish implementing Decoration::glsl_shared + break; + case Decoration::glsl_packed: +#warning finish implementing Decoration::glsl_packed + break; + case Decoration::c_packed: +#warning finish implementing Decoration::c_packed + break; + case Decoration::built_in: + if(built_in) + throw Parser_error( + 0, 0, "multiple BuiltIn decorations on the same variable"); + built_in = util::get( + decoration.parameters) + .built_in; + continue; + case Decoration::no_perspective: +#warning finish implementing Decoration::no_perspective + break; + case Decoration::flat: +#warning finish implementing Decoration::flat + break; + case Decoration::patch: +#warning finish implementing Decoration::patch + break; + case Decoration::centroid: +#warning finish implementing Decoration::centroid + break; + case Decoration::sample: +#warning finish implementing Decoration::sample + break; + case Decoration::invariant: +#warning finish implementing Decoration::invariant + break; + case Decoration::restrict: +#warning finish implementing Decoration::restrict + break; + case Decoration::aliased: +#warning finish implementing Decoration::aliased + break; + case Decoration::volatile_: +#warning finish implementing Decoration::volatile_ + break; + case Decoration::constant: +#warning finish implementing Decoration::constant + break; + case Decoration::coherent: +#warning finish implementing Decoration::coherent + break; + case Decoration::non_writable: +#warning finish implementing Decoration::non_writable + break; + case Decoration::non_readable: +#warning finish implementing Decoration::non_readable + break; + case Decoration::uniform: +#warning finish implementing Decoration::uniform + break; + case Decoration::saturated_conversion: +#warning finish implementing Decoration::saturated_conversion + break; + case Decoration::stream: +#warning finish implementing Decoration::stream + break; + case Decoration::location: +#warning finish implementing Decoration::location + break; + case Decoration::component: +#warning finish implementing Decoration::component + break; + case Decoration::index: +#warning finish implementing Decoration::index + break; + case Decoration::binding: +#warning finish implementing Decoration::binding + break; + case Decoration::descriptor_set: +#warning finish implementing Decoration::descriptor_set + break; + case Decoration::offset: +#warning finish implementing Decoration::offset + break; + case Decoration::xfb_buffer: +#warning finish implementing Decoration::xfb_buffer + break; + case Decoration::xfb_stride: +#warning finish implementing Decoration::xfb_stride + break; + case Decoration::func_param_attr: +#warning finish implementing Decoration::func_param_attr + break; + case Decoration::fp_rounding_mode: +#warning finish implementing Decoration::fp_rounding_mode + break; + case Decoration::fp_fast_math_mode: +#warning finish implementing Decoration::fp_fast_math_mode + break; + case Decoration::linkage_attributes: +#warning finish implementing Decoration::linkage_attributes + break; + case Decoration::no_contraction: +#warning finish implementing Decoration::no_contraction + break; + case Decoration::input_attachment_index: +#warning finish implementing Decoration::input_attachment_index + break; + case Decoration::alignment: +#warning finish implementing Decoration::alignment + break; + case Decoration::max_byte_offset: +#warning finish implementing Decoration::max_byte_offset + break; + case Decoration::alignment_id: +#warning finish implementing Decoration::alignment_id + break; + case Decoration::max_byte_offset_id: +#warning finish implementing Decoration::max_byte_offset_id + break; + case Decoration::override_coverage_nv: +#warning finish implementing Decoration::override_coverage_nv + break; + case Decoration::passthrough_nv: +#warning finish implementing Decoration::passthrough_nv + break; + case Decoration::viewport_relative_nv: +#warning finish implementing Decoration::viewport_relative_nv + break; + case Decoration::secondary_viewport_relative_nv: +#warning finish implementing Decoration::secondary_viewport_relative_nv + break; + } + throw Parser_error( + 0, + 0, + "unimplemented member decoration on shader input variable: " + + std::string(get_enumerant_name(decoration.value))); + } + if(!built_in) + throw Parser_error( + 0, 0, "non-built-in shader input variables are not implemented"); + do + { + switch(*built_in) + { + case Built_in::position: +#warning finish implementing Built_in::position + break; + case Built_in::point_size: +#warning finish implementing Built_in::point_size + break; + case Built_in::clip_distance: +#warning finish implementing Built_in::clip_distance + break; + case Built_in::cull_distance: +#warning finish implementing Built_in::cull_distance + break; + case Built_in::vertex_id: +#warning finish implementing Built_in::vertex_id + break; + case Built_in::instance_id: +#warning finish implementing Built_in::instance_id + break; + case Built_in::primitive_id: +#warning finish implementing Built_in::primitive_id + break; + case Built_in::invocation_id: +#warning finish implementing Built_in::invocation_id + break; + case Built_in::layer: +#warning finish implementing Built_in::layer + break; + case Built_in::viewport_index: +#warning finish implementing Built_in::viewport_index + break; + case Built_in::tess_level_outer: +#warning finish implementing Built_in::tess_level_outer + break; + case Built_in::tess_level_inner: +#warning finish implementing Built_in::tess_level_inner + break; + case Built_in::tess_coord: +#warning finish implementing Built_in::tess_coord + break; + case Built_in::patch_vertices: +#warning finish implementing Built_in::patch_vertices + break; + case Built_in::frag_coord: +#warning finish implementing Built_in::frag_coord + break; + case Built_in::point_coord: +#warning finish implementing Built_in::point_coord + break; + case Built_in::front_facing: +#warning finish implementing Built_in::front_facing + break; + case Built_in::sample_id: +#warning finish implementing Built_in::sample_id + break; + case Built_in::sample_position: +#warning finish implementing Built_in::sample_position + break; + case Built_in::sample_mask: +#warning finish implementing Built_in::sample_mask + break; + case Built_in::frag_depth: +#warning finish implementing Built_in::frag_depth + break; + case Built_in::helper_invocation: +#warning finish implementing Built_in::helper_invocation + break; + case Built_in::num_workgroups: +#warning finish implementing Built_in::num_workgroups + break; + case Built_in::workgroup_size: +#warning finish implementing Built_in::workgroup_size + break; + case Built_in::workgroup_id: +#warning finish implementing Built_in::workgroup_id + break; + case Built_in::local_invocation_id: +#warning finish implementing Built_in::local_invocation_id + break; + case Built_in::global_invocation_id: +#warning finish implementing Built_in::global_invocation_id + break; + case Built_in::local_invocation_index: +#warning finish implementing Built_in::local_invocation_index + break; + case Built_in::work_dim: +#warning finish implementing Built_in::work_dim + break; + case Built_in::global_size: +#warning finish implementing Built_in::global_size + break; + case Built_in::enqueued_workgroup_size: +#warning finish implementing Built_in::enqueued_workgroup_size + break; + case Built_in::global_offset: +#warning finish implementing Built_in::global_offset + break; + case Built_in::global_linear_id: +#warning finish implementing Built_in::global_linear_id + break; + case Built_in::subgroup_size: +#warning finish implementing Built_in::subgroup_size + break; + case Built_in::subgroup_max_size: +#warning finish implementing Built_in::subgroup_max_size + break; + case Built_in::num_subgroups: +#warning finish implementing Built_in::num_subgroups + break; + case Built_in::num_enqueued_subgroups: +#warning finish implementing Built_in::num_enqueued_subgroups + break; + case Built_in::subgroup_id: +#warning finish implementing Built_in::subgroup_id + break; + case Built_in::subgroup_local_invocation_id: +#warning finish implementing Built_in::subgroup_local_invocation_id + break; + case Built_in::vertex_index: + { + if(::LLVMGetElementType(::LLVMTypeOf(input_pointer)) + != llvm_vertex_index_type) + throw Parser_error( + 0, 0, "invalid type for vertex index built-in variable"); + ::LLVMBuildStore(builder.get(), vertex_index, input_pointer); + continue; + } + case Built_in::instance_index: +#warning finish implementing Built_in::instance_index + break; + case Built_in::subgroup_eq_mask_khr: +#warning finish implementing Built_in::subgroup_eq_mask_khr + break; + case Built_in::subgroup_ge_mask_khr: +#warning finish implementing Built_in::subgroup_ge_mask_khr + break; + case Built_in::subgroup_gt_mask_khr: +#warning finish implementing Built_in::subgroup_gt_mask_khr + break; + case Built_in::subgroup_le_mask_khr: +#warning finish implementing Built_in::subgroup_le_mask_khr + break; + case Built_in::subgroup_lt_mask_khr: +#warning finish implementing Built_in::subgroup_lt_mask_khr + break; + case Built_in::base_vertex: +#warning finish implementing Built_in::base_vertex + break; + case Built_in::base_instance: +#warning finish implementing Built_in::base_instance + break; + case Built_in::draw_index: +#warning finish implementing Built_in::draw_index + break; + case Built_in::device_index: +#warning finish implementing Built_in::device_index + break; + case Built_in::view_index: +#warning finish implementing Built_in::view_index + break; + case Built_in::viewport_mask_nv: +#warning finish implementing Built_in::viewport_mask_nv + break; + case Built_in::secondary_position_nv: +#warning finish implementing Built_in::secondary_position_nv + break; + case Built_in::secondary_viewport_mask_nv: +#warning finish implementing Built_in::secondary_viewport_mask_nv + break; + case Built_in::position_per_view_nv: +#warning finish implementing Built_in::position_per_view_nv + break; + case Built_in::viewport_mask_per_view_nv: +#warning finish implementing Built_in::viewport_mask_per_view_nv + break; + } + throw Parser_error(0, + 0, + "unimplemented built in shader input variable: " + + std::string(get_enumerant_name(*built_in))); + } while(false); + } + } + else if(member.type == outputs_struct) + { + auto outputs_struct_pointer = ::LLVMBuildStructGEP( + builder.get(), io_struct_pointer, member.llvm_member_index, "outputs"); + for(auto &output_member : outputs_struct->get_members(true)) + { + auto output_pointer = ::LLVMBuildStructGEP(builder.get(), + outputs_struct_pointer, + output_member.llvm_member_index, + "output"); + static_cast(output_pointer); + for(auto &decoration : output_member.decorations) + { + switch(decoration.value) + { + case Decoration::relaxed_precision: +#warning finish implementing Decoration::relaxed_precision + break; + case Decoration::spec_id: +#warning finish implementing Decoration::spec_id + break; + case Decoration::block: +#warning finish implementing Decoration::block + break; + case Decoration::buffer_block: +#warning finish implementing Decoration::buffer_block + break; + case Decoration::row_major: +#warning finish implementing Decoration::row_major + break; + case Decoration::col_major: +#warning finish implementing Decoration::col_major + break; + case Decoration::array_stride: +#warning finish implementing Decoration::array_stride + break; + case Decoration::matrix_stride: +#warning finish implementing Decoration::matrix_stride + break; + case Decoration::glsl_shared: +#warning finish implementing Decoration::glsl_shared + break; + case Decoration::glsl_packed: +#warning finish implementing Decoration::glsl_packed + break; + case Decoration::c_packed: +#warning finish implementing Decoration::c_packed + break; + case Decoration::built_in: +#warning finish implementing Decoration::built_in + break; + case Decoration::no_perspective: +#warning finish implementing Decoration::no_perspective + break; + case Decoration::flat: +#warning finish implementing Decoration::flat + break; + case Decoration::patch: +#warning finish implementing Decoration::patch + break; + case Decoration::centroid: +#warning finish implementing Decoration::centroid + break; + case Decoration::sample: +#warning finish implementing Decoration::sample + break; + case Decoration::invariant: +#warning finish implementing Decoration::invariant + break; + case Decoration::restrict: +#warning finish implementing Decoration::restrict + break; + case Decoration::aliased: +#warning finish implementing Decoration::aliased + break; + case Decoration::volatile_: +#warning finish implementing Decoration::volatile_ + break; + case Decoration::constant: +#warning finish implementing Decoration::constant + break; + case Decoration::coherent: +#warning finish implementing Decoration::coherent + break; + case Decoration::non_writable: +#warning finish implementing Decoration::non_writable + break; + case Decoration::non_readable: +#warning finish implementing Decoration::non_readable + break; + case Decoration::uniform: +#warning finish implementing Decoration::uniform + break; + case Decoration::saturated_conversion: +#warning finish implementing Decoration::saturated_conversion + break; + case Decoration::stream: +#warning finish implementing Decoration::stream + break; + case Decoration::location: +#warning finish implementing Decoration::location + break; + case Decoration::component: +#warning finish implementing Decoration::component + break; + case Decoration::index: +#warning finish implementing Decoration::index + break; + case Decoration::binding: +#warning finish implementing Decoration::binding + break; + case Decoration::descriptor_set: +#warning finish implementing Decoration::descriptor_set + break; + case Decoration::offset: +#warning finish implementing Decoration::offset + break; + case Decoration::xfb_buffer: +#warning finish implementing Decoration::xfb_buffer + break; + case Decoration::xfb_stride: +#warning finish implementing Decoration::xfb_stride + break; + case Decoration::func_param_attr: +#warning finish implementing Decoration::func_param_attr + break; + case Decoration::fp_rounding_mode: +#warning finish implementing Decoration::fp_rounding_mode + break; + case Decoration::fp_fast_math_mode: +#warning finish implementing Decoration::fp_fast_math_mode + break; + case Decoration::linkage_attributes: +#warning finish implementing Decoration::linkage_attributes + break; + case Decoration::no_contraction: +#warning finish implementing Decoration::no_contraction + break; + case Decoration::input_attachment_index: +#warning finish implementing Decoration::input_attachment_index + break; + case Decoration::alignment: +#warning finish implementing Decoration::alignment + break; + case Decoration::max_byte_offset: +#warning finish implementing Decoration::max_byte_offset + break; + case Decoration::alignment_id: +#warning finish implementing Decoration::alignment_id + break; + case Decoration::max_byte_offset_id: +#warning finish implementing Decoration::max_byte_offset_id + break; + case Decoration::override_coverage_nv: +#warning finish implementing Decoration::override_coverage_nv + break; + case Decoration::passthrough_nv: +#warning finish implementing Decoration::passthrough_nv + break; + case Decoration::viewport_relative_nv: +#warning finish implementing Decoration::viewport_relative_nv + break; + case Decoration::secondary_viewport_relative_nv: +#warning finish implementing Decoration::secondary_viewport_relative_nv + break; + } + throw Parser_error( + 0, + 0, + "unimplemented member decoration on shader output variable: " + + std::string(get_enumerant_name(decoration.value))); + } + } + } + else + { + throw Parser_error(0, 0, "internal error: unhandled Io_struct member"); + } + } + { + constexpr std::size_t arg_count = 1; + assert(implicit_function_arguments.size() == arg_count); + assert(implicit_function_arguments[0]->get_or_make_type().type + == ::LLVMTypeOf(io_struct_pointer)); + ::LLVMValueRef args[arg_count] = { + io_struct_pointer, + }; + assert(::LLVMGetReturnType(::LLVMGetElementType(::LLVMTypeOf(main_function))) + == llvm_wrapper::Create_llvm_type()(context)); + ::LLVMBuildCall(builder.get(), main_function, args, arg_count, ""); + } +#warning add output copy + auto next_iteration_condition = + ::LLVMBuildICmp(builder.get(), + ::LLVMIntULT, + next_vertex_index, + ::LLVMGetParam(entry_function, arg_vertex_end_index), + "next_iteration_condition"); + ::LLVMBuildCondBr(builder.get(), next_iteration_condition, loop_block, exit_block); + ::LLVMPositionBuilderAtEnd(builder.get(), exit_block); + static_assert( + std::is_same()(0, 0, 0, nullptr)), + void>::value, + ""); + ::LLVMBuildRetVoid(builder.get()); break; } case spirv::Execution_model::tessellation_control: @@ -864,8 +1482,8 @@ public: "unimplemented execution model: " + std::string(spirv::get_enumerant_name(execution_model))); } - assert(function); - return ::LLVMGetValueName(function); + assert(entry_function); + return ::LLVMGetValueName(entry_function); } Converted_module run(const Word *shader_words, std::size_t shader_size) { @@ -899,7 +1517,8 @@ public: + " \"" + std::string(entry_point_name) + "\""); - entry_function_name = generate_entry_function(entry_point); + entry_function_name = + generate_entry_function(entry_point, id_state.function->function); } } if(entry_function_name.empty()) @@ -2601,20 +3220,32 @@ void Spirv_to_llvm::handle_instruction_op_type_function(Op_type_function instruc + instruction.parameter_0_type_parameter_1_type.size()); for(auto &arg : implicit_function_arguments) args.push_back(arg); + bool return_type_is_void = false; + auto return_type = get_type(instruction.return_type, instruction_start_index); + if(auto *simple_return_type = dynamic_cast(return_type.get())) + if(simple_return_type->get_or_make_type().type == ::LLVMVoidTypeInContext(context)) + return_type_is_void = true; + bool valid_for_entry_point = + instruction.parameter_0_type_parameter_1_type.empty() && return_type_is_void; for(Id_ref type : instruction.parameter_0_type_parameter_1_type) + { args.push_back(get_type(type, instruction_start_index)); + } auto &state = get_id_state(instruction.result); if(!state.decorations.empty()) throw Parser_error(instruction_start_index, instruction_start_index, "decorations on instruction not implemented: " + std::string(get_enumerant_name(instruction.get_operation()))); + constexpr bool is_var_arg = false; state.type = std::make_shared( state.decorations, get_type(instruction.return_type, instruction_start_index), std::move(args), instruction_start_index, - target_data); + target_data, + valid_for_entry_point, + is_var_arg); break; } case Stage::generate_code: @@ -2949,6 +3580,10 @@ void Spirv_to_llvm::handle_instruction_op_function(Op_function instruction, auto function_name = get_name(current_function_id); if(function_name.empty() && state.op_entry_points.size() == 1) function_name = std::string(state.op_entry_points[0].entry_point.name); + if(!state.op_entry_points.empty() && !function_type->is_valid_for_entry_point()) + throw Parser_error(instruction_start_index, + instruction_start_index, + "invalid function type for entry point"); function_name = get_or_make_prefixed_name(std::move(function_name), false); auto function = ::LLVMAddFunction( module.get(), function_name.c_str(), function_type->get_or_make_type().type); diff --git a/src/spirv_to_llvm/spirv_to_llvm.h b/src/spirv_to_llvm/spirv_to_llvm.h index 602ed09..8d57b8f 100644 --- a/src/spirv_to_llvm/spirv_to_llvm.h +++ b/src/spirv_to_llvm/spirv_to_llvm.h @@ -399,6 +399,7 @@ private: LLVM_type_and_alignment type; Recursion_checker_state recursion_checker_state; std::size_t instruction_start_index; + bool valid_for_entry_point; bool is_var_arg; public: @@ -407,12 +408,14 @@ public: std::vector> args, std::size_t instruction_start_index, ::LLVMTargetDataRef target_data, - bool is_var_arg = false) noexcept + bool valid_for_entry_point, + bool is_var_arg) noexcept : Type_descriptor(std::move(decorations)), return_type(std::move(return_type)), args(std::move(args)), type(nullptr, llvm_wrapper::Target_data::get_pointer_alignment(target_data)), instruction_start_index(instruction_start_index), + valid_for_entry_point(valid_for_entry_point), is_var_arg(is_var_arg) { } @@ -435,6 +438,10 @@ public: { type_visitor.visit(*this); } + bool is_valid_for_entry_point() const noexcept + { + return valid_for_entry_point; + } }; class Struct_type_descriptor final : public Type_descriptor -- 2.30.2