[spirv-reader] Start emitting functions
Build the SPIR-V tools IR context, which gives us a lot of information
about the SPIR-V module.
Add helper functions to emit Tint types from SPIR-V types and result
IDs. Currently only supports the void type.
Emit a function with its return type, and emit the contents of its
entry block. Currently only supports the OpReturn instruction.
Adds a basic test framework for parsing SPIR-V assembly and comparing
the result to an IR disassembly.
Bug: tint:1907
Change-Id: I4eb7441e9459320af11c883ed85aaa3d470b3a64
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/165540
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
diff --git a/src/tint/cmd/test/BUILD.bazel b/src/tint/cmd/test/BUILD.bazel
index e0e9b73..a44e992 100644
--- a/src/tint/cmd/test/BUILD.bazel
+++ b/src/tint/cmd/test/BUILD.bazel
@@ -121,6 +121,11 @@
],
"//conditions:default": [],
}) + select({
+ ":tint_build_spv_reader": [
+ "//src/tint/lang/spirv/reader/parser:test",
+ ],
+ "//conditions:default": [],
+ }) + select({
":tint_build_spv_reader_and_tint_build_wgsl_reader_and_tint_build_wgsl_writer": [
"//src/tint/lang/spirv/reader/ast_lower:test",
],
diff --git a/src/tint/cmd/test/BUILD.cmake b/src/tint/cmd/test/BUILD.cmake
index 909b3e8..e79c221 100644
--- a/src/tint/cmd/test/BUILD.cmake
+++ b/src/tint/cmd/test/BUILD.cmake
@@ -133,6 +133,12 @@
)
endif(TINT_BUILD_MSL_WRITER AND TINT_BUILD_WGSL_READER AND TINT_BUILD_WGSL_WRITER)
+if(TINT_BUILD_SPV_READER)
+ tint_target_add_dependencies(tint_cmd_test_test_cmd test_cmd
+ tint_lang_spirv_reader_parser_test
+ )
+endif(TINT_BUILD_SPV_READER)
+
if(TINT_BUILD_SPV_READER AND TINT_BUILD_WGSL_READER AND TINT_BUILD_WGSL_WRITER)
tint_target_add_dependencies(tint_cmd_test_test_cmd test_cmd
tint_lang_spirv_reader_ast_lower_test
diff --git a/src/tint/cmd/test/BUILD.gn b/src/tint/cmd/test/BUILD.gn
index 622cf0f..55c146b 100644
--- a/src/tint/cmd/test/BUILD.gn
+++ b/src/tint/cmd/test/BUILD.gn
@@ -127,6 +127,10 @@
deps += [ "${tint_src_dir}/lang/msl/writer/ast_raise:unittests" ]
}
+ if (tint_build_spv_reader) {
+ deps += [ "${tint_src_dir}/lang/spirv/reader/parser:unittests" ]
+ }
+
if (tint_build_spv_reader && tint_build_wgsl_reader &&
tint_build_wgsl_writer) {
deps += [ "${tint_src_dir}/lang/spirv/reader/ast_lower:unittests" ]
diff --git a/src/tint/lang/core/ir/builder.cc b/src/tint/lang/core/ir/builder.cc
index 9315f8a4..ba14605 100644
--- a/src/tint/lang/core/ir/builder.cc
+++ b/src/tint/lang/core/ir/builder.cc
@@ -51,17 +51,24 @@
return ir.blocks.Create<ir::MultiInBlock>();
}
-Function* Builder::Function(std::string_view name,
- const core::type::Type* return_type,
+Function* Builder::Function(const core::type::Type* return_type,
Function::PipelineStage stage,
std::optional<std::array<uint32_t, 3>> wg_size) {
auto* ir_func = ir.values.Create<ir::Function>(return_type, stage, wg_size);
ir_func->SetBlock(Block());
- ir.SetName(ir_func, name);
ir.functions.Push(ir_func);
return ir_func;
}
+Function* Builder::Function(std::string_view name,
+ const core::type::Type* return_type,
+ Function::PipelineStage stage,
+ std::optional<std::array<uint32_t, 3>> wg_size) {
+ auto* ir_func = Function(return_type, stage, wg_size);
+ ir.SetName(ir_func, name);
+ return ir_func;
+}
+
ir::Loop* Builder::Loop() {
return Append(ir.instructions.Create<ir::Loop>(Block(), MultiInBlock(), MultiInBlock()));
}
diff --git a/src/tint/lang/core/ir/builder.h b/src/tint/lang/core/ir/builder.h
index ebbd851..fc847c7 100644
--- a/src/tint/lang/core/ir/builder.h
+++ b/src/tint/lang/core/ir/builder.h
@@ -235,12 +235,21 @@
/// @returns a new multi-in block
ir::MultiInBlock* MultiInBlock();
- /// Creates a function instruction
+ /// Creates an unnamed function
+ /// @param return_type the function return type
+ /// @param stage the function stage
+ /// @param wg_size the workgroup_size
+ /// @returns the function
+ ir::Function* Function(const core::type::Type* return_type,
+ Function::PipelineStage stage = Function::PipelineStage::kUndefined,
+ std::optional<std::array<uint32_t, 3>> wg_size = {});
+
+ /// Creates a function
/// @param name the function name
/// @param return_type the function return type
/// @param stage the function stage
/// @param wg_size the workgroup_size
- /// @returns the instruction
+ /// @returns the function
ir::Function* Function(std::string_view name,
const core::type::Type* return_type,
Function::PipelineStage stage = Function::PipelineStage::kUndefined,
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.bazel b/src/tint/lang/spirv/reader/parser/BUILD.bazel
index 44f0586..ff01e93 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/parser/BUILD.bazel
@@ -67,6 +67,48 @@
] + select({
":tint_build_spv_reader_or_tint_build_spv_writer": [
"//src/tint/lang/spirv/validate",
+ "@spirv_tools//:spirv_tools_opt",
+ "@spirv_tools",
+ ],
+ "//conditions:default": [],
+ }),
+ copts = COPTS,
+ visibility = ["//visibility:public"],
+)
+cc_library(
+ name = "test",
+ alwayslink = True,
+ srcs = [
+ "function_test.cc",
+ "helper_test.h",
+ ],
+ deps = [
+ "//src/tint/api/common",
+ "//src/tint/lang/core",
+ "//src/tint/lang/core/constant",
+ "//src/tint/lang/core/ir",
+ "//src/tint/lang/core/type",
+ "//src/tint/utils/containers",
+ "//src/tint/utils/diagnostic",
+ "//src/tint/utils/ice",
+ "//src/tint/utils/id",
+ "//src/tint/utils/macros",
+ "//src/tint/utils/math",
+ "//src/tint/utils/memory",
+ "//src/tint/utils/reflection",
+ "//src/tint/utils/result",
+ "//src/tint/utils/rtti",
+ "//src/tint/utils/symbol",
+ "//src/tint/utils/text",
+ "//src/tint/utils/traits",
+ "@gtest",
+ ] + select({
+ ":tint_build_spv_reader": [
+ "//src/tint/lang/spirv/reader/parser",
+ ],
+ "//conditions:default": [],
+ }) + select({
+ ":tint_build_spv_reader_or_tint_build_spv_writer": [
"@spirv_tools",
],
"//conditions:default": [],
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.cmake b/src/tint/lang/spirv/reader/parser/BUILD.cmake
index 0aee0a6..2b6ba04 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/parser/BUILD.cmake
@@ -72,6 +72,56 @@
tint_lang_spirv_validate
)
tint_target_add_external_dependencies(tint_lang_spirv_reader_parser lib
+ "spirv-opt-internal"
+ "spirv-tools"
+ )
+endif(TINT_BUILD_SPV_READER OR TINT_BUILD_SPV_WRITER)
+
+endif(TINT_BUILD_SPV_READER)
+if(TINT_BUILD_SPV_READER)
+################################################################################
+# Target: tint_lang_spirv_reader_parser_test
+# Kind: test
+# Condition: TINT_BUILD_SPV_READER
+################################################################################
+tint_add_target(tint_lang_spirv_reader_parser_test test
+ lang/spirv/reader/parser/function_test.cc
+ lang/spirv/reader/parser/helper_test.h
+)
+
+tint_target_add_dependencies(tint_lang_spirv_reader_parser_test test
+ tint_api_common
+ tint_lang_core
+ tint_lang_core_constant
+ tint_lang_core_ir
+ tint_lang_core_type
+ tint_utils_containers
+ tint_utils_diagnostic
+ tint_utils_ice
+ tint_utils_id
+ tint_utils_macros
+ tint_utils_math
+ tint_utils_memory
+ tint_utils_reflection
+ tint_utils_result
+ tint_utils_rtti
+ tint_utils_symbol
+ tint_utils_text
+ tint_utils_traits
+)
+
+tint_target_add_external_dependencies(tint_lang_spirv_reader_parser_test test
+ "gtest"
+)
+
+if(TINT_BUILD_SPV_READER)
+ tint_target_add_dependencies(tint_lang_spirv_reader_parser_test test
+ tint_lang_spirv_reader_parser
+ )
+endif(TINT_BUILD_SPV_READER)
+
+if(TINT_BUILD_SPV_READER OR TINT_BUILD_SPV_WRITER)
+ tint_target_add_external_dependencies(tint_lang_spirv_reader_parser_test test
"spirv-tools"
)
endif(TINT_BUILD_SPV_READER OR TINT_BUILD_SPV_WRITER)
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.gn b/src/tint/lang/spirv/reader/parser/BUILD.gn
index ccce1da..f2cd647 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.gn
+++ b/src/tint/lang/spirv/reader/parser/BUILD.gn
@@ -37,6 +37,10 @@
import("../../../../../../scripts/tint_overrides_with_defaults.gni")
import("${tint_src_dir}/tint.gni")
+
+if (tint_build_unittests || tint_build_benchmarks) {
+ import("//testing/test.gni")
+}
if (tint_build_spv_reader) {
libtint_source_set("parser") {
sources = [
@@ -67,10 +71,56 @@
if (tint_build_spv_reader || tint_build_spv_writer) {
deps += [
+ "${tint_spirv_tools_dir}:spvtools",
"${tint_spirv_tools_dir}:spvtools_headers",
+ "${tint_spirv_tools_dir}:spvtools_opt",
+ "${tint_spirv_tools_dir}:spvtools_val",
"${tint_spirv_tools_dir}:spvtools_val",
"${tint_src_dir}/lang/spirv/validate",
]
}
+ public_configs = [ "${tint_spirv_tools_dir}/:spvtools_internal_config" ]
+ }
+}
+if (tint_build_unittests) {
+ if (tint_build_spv_reader) {
+ tint_unittests_source_set("unittests") {
+ sources = [
+ "function_test.cc",
+ "helper_test.h",
+ ]
+ deps = [
+ "${tint_src_dir}:gmock_and_gtest",
+ "${tint_src_dir}/api/common",
+ "${tint_src_dir}/lang/core",
+ "${tint_src_dir}/lang/core/constant",
+ "${tint_src_dir}/lang/core/ir",
+ "${tint_src_dir}/lang/core/type",
+ "${tint_src_dir}/utils/containers",
+ "${tint_src_dir}/utils/diagnostic",
+ "${tint_src_dir}/utils/ice",
+ "${tint_src_dir}/utils/id",
+ "${tint_src_dir}/utils/macros",
+ "${tint_src_dir}/utils/math",
+ "${tint_src_dir}/utils/memory",
+ "${tint_src_dir}/utils/reflection",
+ "${tint_src_dir}/utils/result",
+ "${tint_src_dir}/utils/rtti",
+ "${tint_src_dir}/utils/symbol",
+ "${tint_src_dir}/utils/text",
+ "${tint_src_dir}/utils/traits",
+ ]
+
+ if (tint_build_spv_reader) {
+ deps += [ "${tint_src_dir}/lang/spirv/reader/parser" ]
+ }
+
+ if (tint_build_spv_reader || tint_build_spv_writer) {
+ deps += [
+ "${tint_spirv_tools_dir}:spvtools_headers",
+ "${tint_spirv_tools_dir}:spvtools_val",
+ ]
+ }
+ }
}
}
diff --git a/src/tint/lang/spirv/reader/parser/function_test.cc b/src/tint/lang/spirv/reader/parser/function_test.cc
new file mode 100644
index 0000000..be6c28c
--- /dev/null
+++ b/src/tint/lang/spirv/reader/parser/function_test.cc
@@ -0,0 +1,56 @@
+// Copyright 2023 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/reader/parser/helper_test.h"
+
+namespace tint::spirv::reader {
+namespace {
+
+TEST_F(SpirvParserTest, EmptyEntryPoint) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%1 = func():void -> %b1 {
+ %b1 = block {
+ ret
+ }
+}
+)");
+}
+
+} // namespace
+} // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/helper_test.h b/src/tint/lang/spirv/reader/parser/helper_test.h
new file mode 100644
index 0000000..9b03490
--- /dev/null
+++ b/src/tint/lang/spirv/reader/parser/helper_test.h
@@ -0,0 +1,94 @@
+// Copyright 2023 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#ifndef SRC_TINT_LANG_SPIRV_READER_PARSER_HELPER_TEST_H_
+#define SRC_TINT_LANG_SPIRV_READER_PARSER_HELPER_TEST_H_
+
+#include <string>
+#include <vector>
+
+#include "gmock/gmock.h"
+#include "gtest/gtest.h"
+#include "spirv-tools/libspirv.hpp"
+#include "src/tint/lang/core/ir/disassembler.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/spirv/reader/parser/parser.h"
+
+namespace tint::spirv::reader {
+
+// Helper macro to run the parser and compare the disassembled IR to a string.
+// Automatically prefixes the IR disassembly with a newline to improve formatting of tests.
+#define EXPECT_IR(asm, ir) ASSERT_EQ("\n" + Run(asm), ir) << ir
+
+/// Base helper class for testing the SPIR-V parser implementation.
+template <typename BASE>
+class SpirvParserTestHelperBase : public BASE {
+ protected:
+ /// Run the parser on a SPIR-V module and return the Tint IR or an error string.
+ /// @param spirv_asm the SPIR-V assembly to parse
+ /// @returns the disassembled Tint IR
+ std::string Run(std::string spirv_asm) {
+ // Assemble the SPIR-V input.
+ StringStream err;
+ std::vector<uint32_t> binary;
+ spvtools::SpirvTools tools(SPV_ENV_UNIVERSAL_1_0);
+ tools.SetMessageConsumer(
+ [&err](spv_message_level_t, const char*, const spv_position_t& pos, const char* msg) {
+ err << "SPIR-V assembly failed:" << pos.line << ":" << pos.column << ": " << msg;
+ });
+ auto assembled =
+ tools.Assemble(spirv_asm, &binary, SPV_TEXT_TO_BINARY_OPTION_PRESERVE_NUMERIC_IDS);
+ if (!assembled) {
+ return err.str();
+ }
+
+ // Parse the SPIR-V to produce an IR module.
+ auto parsed = Parse(Slice(binary.data(), binary.size()));
+ if (!parsed) {
+ return parsed.Failure().reason.str();
+ }
+
+ // Validate the IR module.
+ auto validated = core::ir::Validate(parsed.Get());
+ if (!validated) {
+ return validated.Failure().reason.str();
+ }
+
+ // Return the disassembled IR module.
+ return core::ir::Disassemble(parsed.Get());
+ }
+};
+
+using SpirvParserTest = SpirvParserTestHelperBase<testing::Test>;
+
+template <typename T>
+using SpirvParserTestWithParam = SpirvParserTestHelperBase<testing::TestWithParam<T>>;
+
+} // namespace tint::spirv::reader
+
+#endif // SRC_TINT_LANG_SPIRV_READER_PARSER_HELPER_TEST_H_
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 11ad5b2..aacba85 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -27,9 +27,20 @@
#include "src/tint/lang/spirv/reader/parser/parser.h"
+#include <memory>
#include <utility>
#include <vector>
+TINT_BEGIN_DISABLE_WARNING(NEWLINE_EOF);
+TINT_BEGIN_DISABLE_WARNING(OLD_STYLE_CAST);
+TINT_BEGIN_DISABLE_WARNING(SIGN_CONVERSION);
+TINT_BEGIN_DISABLE_WARNING(WEAK_VTABLES);
+#include "source/opt/build_module.h"
+TINT_END_DISABLE_WARNING(WEAK_VTABLES);
+TINT_END_DISABLE_WARNING(SIGN_CONVERSION);
+TINT_END_DISABLE_WARNING(OLD_STYLE_CAST);
+TINT_END_DISABLE_WARNING(NEWLINE_EOF);
+
#include "src/tint/lang/core/ir/builder.h"
#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/spirv/validate/validate.h"
@@ -42,6 +53,7 @@
constexpr auto kTargetEnv = SPV_ENV_VULKAN_1_1;
/// PIMPL class for SPIR-V parser.
+/// Validates the SPIR-V module and then parses it to produce a Tint IR module.
class Parser {
public:
/// @param spirv the SPIR-V binary data
@@ -53,14 +65,83 @@
return result.Failure();
}
- // TODO(crbug.com/tint/1907): Parse the module.
+ // Build the SPIR-V tools internal representation of the SPIR-V module.
+ spvtools::Context context(kTargetEnv);
+ spirv_context_ =
+ spvtools::BuildModule(kTargetEnv, context.CContext()->consumer, spirv.data, spirv.len);
+ if (!spirv_context_) {
+ return Failure("failed to build the internal representation of the module");
+ }
+
+ // TODO(crbug.com/tint/1907): Emit module-scope variables.
+
+ EmitFunctions();
+
+ // TODO(crbug.com/tint/1907): Handle entry point declarations and execution modes.
+ // TODO(crbug.com/tint/1907): Handle entry point declarations and execution modes.
+ // TODO(crbug.com/tint/1907): Handle annotation instructions.
+ // TODO(crbug.com/tint/1907): Handle names.
return std::move(ir_);
}
+ /// @param type a SPIR-V type object
+ /// @returns a Tint type object
+ const core::type::Type* Type(const spvtools::opt::analysis::Type* type) {
+ switch (type->kind()) {
+ case spvtools::opt::analysis::Type::kVoid:
+ return ty_.void_();
+ default:
+ TINT_UNIMPLEMENTED() << "unhandled SPIR-V type: " << type->str();
+ return ty_.void_();
+ }
+ }
+
+ /// @param id a SPIR-V result ID for a type declaration instruction
+ /// @returns a Tint type object
+ const core::type::Type* Type(uint32_t id) {
+ return Type(spirv_context_->get_type_mgr()->GetType(id));
+ }
+
+ /// Emit the functions.
+ void EmitFunctions() {
+ for (auto& func : *spirv_context_->module()) {
+ // TODO(crbug.com/tint/1907): Emit function parameters as well.
+ current_function_ = b_.Function(
+ Type(func.type_id()), core::ir::Function::PipelineStage::kUndefined, std::nullopt);
+ EmitBlock(current_function_->Block(), *func.entry());
+ }
+ }
+
+ /// Emit the contents of SPIR-V block @p src into Tint IR block @p dst.
+ /// @param dst the Tint IR block to append to
+ /// @param src the SPIR-V block to emit
+ void EmitBlock(core::ir::Block* dst, const spvtools::opt::BasicBlock& src) {
+ for (auto& inst : src) {
+ switch (inst.opcode()) {
+ case spv::Op::OpReturn:
+ dst->Append(b_.Return(current_function_));
+ break;
+ default:
+ TINT_UNIMPLEMENTED()
+ << "unhandled SPIR-V instruction: " << static_cast<uint32_t>(inst.opcode());
+ }
+ }
+ }
+
private:
/// The generated IR module.
core::ir::Module ir_;
+ /// The Tint IR builder.
+ core::ir::Builder b_{ir_};
+ /// The Tint type manager.
+ core::type::Manager& ty_{ir_.Types()};
+
+ /// The Tint IR function that is currently being emitted.
+ core::ir::Function* current_function_ = nullptr;
+
+ /// The SPIR-V context containing the SPIR-V tools intermediate representation.
+ std::unique_ptr<spvtools::opt::IRContext> spirv_context_;
};
} // namespace
