[tint][ir] Begin building proto-based serializer
Change-Id: Idf01b213b642afb4614ee0b0dbc4d1a70979cc70
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/155702
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 208a231..9a4eba5 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -41,6 +41,12 @@
endif()
endif()
+if (TARGET libprotobuf)
+ set(TINT_BUILD_IR_BINARY 1)
+else()
+ set(TINT_BUILD_IR_BINARY 0)
+endif()
+
################################################################################
# Helper functions
################################################################################
@@ -50,6 +56,7 @@
target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_GLSL_VALIDATOR=$<BOOL:${TINT_BUILD_GLSL_VALIDATOR}>)
target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_GLSL_WRITER=$<BOOL:${TINT_BUILD_GLSL_WRITER}>)
target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_HLSL_WRITER=$<BOOL:${TINT_BUILD_HLSL_WRITER}>)
+ target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_IR_BINARY=$<BOOL:${TINT_BUILD_IR_BINARY}>)
target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_IS_LINUX=$<BOOL:${TINT_BUILD_IS_LINUX}>)
target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_IS_MAC=$<BOOL:${TINT_BUILD_IS_MAC}>)
target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_IS_WIN=$<BOOL:${TINT_BUILD_IS_WIN}>)
@@ -60,7 +67,6 @@
target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_WGSL_READER=$<BOOL:${TINT_BUILD_WGSL_READER}>)
target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_WGSL_WRITER=$<BOOL:${TINT_BUILD_WGSL_WRITER}>)
-
if(TINT_BUILD_FUZZERS)
target_compile_options(${TARGET} PRIVATE "-fsanitize=fuzzer")
endif()
diff --git a/src/tint/cmd/test/BUILD.bazel b/src/tint/cmd/test/BUILD.bazel
index 080bcb9..e0e9b73 100644
--- a/src/tint/cmd/test/BUILD.bazel
+++ b/src/tint/cmd/test/BUILD.bazel
@@ -103,6 +103,11 @@
],
"//conditions:default": [],
}) + select({
+ ":tint_build_ir_binary": [
+ "//src/tint/lang/core/ir/binary:test",
+ ],
+ "//conditions:default": [],
+ }) + select({
":tint_build_msl_writer": [
"//src/tint/lang/msl/writer/ast_printer:test",
"//src/tint/lang/msl/writer/common:test",
@@ -176,6 +181,11 @@
)
alias(
+ name = "tint_build_ir_binary",
+ actual = "//src/tint:tint_build_ir_binary_true",
+)
+
+alias(
name = "tint_build_msl_writer",
actual = "//src/tint:tint_build_msl_writer_true",
)
diff --git a/src/tint/cmd/test/BUILD.cmake b/src/tint/cmd/test/BUILD.cmake
index 3709145..909b3e8 100644
--- a/src/tint/cmd/test/BUILD.cmake
+++ b/src/tint/cmd/test/BUILD.cmake
@@ -112,6 +112,12 @@
)
endif(TINT_BUILD_HLSL_WRITER AND TINT_BUILD_WGSL_READER AND TINT_BUILD_WGSL_WRITER)
+if(TINT_BUILD_IR_BINARY)
+ tint_target_add_dependencies(tint_cmd_test_test_cmd test_cmd
+ tint_lang_core_ir_binary_test
+ )
+endif(TINT_BUILD_IR_BINARY)
+
if(TINT_BUILD_MSL_WRITER)
tint_target_add_dependencies(tint_cmd_test_test_cmd test_cmd
tint_lang_msl_writer_ast_printer_test
diff --git a/src/tint/cmd/test/BUILD.gn b/src/tint/cmd/test/BUILD.gn
index e307f93..622cf0f 100644
--- a/src/tint/cmd/test/BUILD.gn
+++ b/src/tint/cmd/test/BUILD.gn
@@ -109,6 +109,10 @@
deps += [ "${tint_src_dir}/lang/hlsl/writer/ast_raise:unittests" ]
}
+ if (tint_build_ir_binary) {
+ deps += [ "${tint_src_dir}/lang/core/ir/binary:unittests" ]
+ }
+
if (tint_build_msl_writer) {
deps += [
"${tint_src_dir}/lang/msl/writer/ast_printer:unittests",
diff --git a/src/tint/lang/core/ir/BUILD.cmake b/src/tint/lang/core/ir/BUILD.cmake
index 19c6a8a..c901636 100644
--- a/src/tint/lang/core/ir/BUILD.cmake
+++ b/src/tint/lang/core/ir/BUILD.cmake
@@ -34,6 +34,7 @@
# Do not modify this file directly
################################################################################
+include(lang/core/ir/binary/BUILD.cmake)
include(lang/core/ir/transform/BUILD.cmake)
################################################################################
diff --git a/src/tint/lang/core/ir/binary/BUILD.bazel b/src/tint/lang/core/ir/binary/BUILD.bazel
new file mode 100644
index 0000000..e781ef7
--- /dev/null
+++ b/src/tint/lang/core/ir/binary/BUILD.bazel
@@ -0,0 +1,120 @@
+# 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.
+
+################################################################################
+# File generated by 'tools/src/cmd/gen' using the template:
+# tools/src/cmd/gen/build/BUILD.bazel.tmpl
+#
+# To regenerate run: './tools/run gen'
+#
+# Do not modify this file directly
+################################################################################
+
+load("//src/tint:flags.bzl", "COPTS")
+load("@bazel_skylib//lib:selects.bzl", "selects")
+cc_library(
+ name = "binary",
+ srcs = [
+ "decode.cc",
+ "encode.cc",
+ ],
+ hdrs = [
+ "decode.h",
+ "encode.h",
+ ],
+ deps = [
+ "//src/tint/api/common",
+ "//src/tint/lang/core",
+ "//src/tint/lang/core/constant",
+ "//src/tint/lang/core/intrinsic",
+ "//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",
+ ] + select({
+ ":tint_build_ir_binary": [
+ "",
+ ],
+ "//conditions:default": [],
+ }),
+ copts = COPTS,
+ visibility = ["//visibility:public"],
+)
+cc_library(
+ name = "test",
+ alwayslink = True,
+ srcs = [
+ "roundtrip_test.cc",
+ ],
+ deps = [
+ "//src/tint/api/common",
+ "//src/tint/lang/core",
+ "//src/tint/lang/core/constant",
+ "//src/tint/lang/core/intrinsic",
+ "//src/tint/lang/core/ir",
+ "//src/tint/lang/core/ir:test",
+ "//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_ir_binary": [
+ "//src/tint/lang/core/ir/binary",
+ ],
+ "//conditions:default": [],
+ }),
+ copts = COPTS,
+ visibility = ["//visibility:public"],
+)
+
+alias(
+ name = "tint_build_ir_binary",
+ actual = "//src/tint:tint_build_ir_binary_true",
+)
+
diff --git a/src/tint/lang/core/ir/binary/BUILD.cfg b/src/tint/lang/core/ir/binary/BUILD.cfg
new file mode 100644
index 0000000..deb7ad2
--- /dev/null
+++ b/src/tint/lang/core/ir/binary/BUILD.cfg
@@ -0,0 +1,3 @@
+{
+ "condition": "tint_build_ir_binary"
+}
diff --git a/src/tint/lang/core/ir/binary/BUILD.cmake b/src/tint/lang/core/ir/binary/BUILD.cmake
new file mode 100644
index 0000000..61650da
--- /dev/null
+++ b/src/tint/lang/core/ir/binary/BUILD.cmake
@@ -0,0 +1,133 @@
+# 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.
+
+################################################################################
+# File generated by 'tools/src/cmd/gen' using the template:
+# tools/src/cmd/gen/build/BUILD.cmake.tmpl
+#
+# To regenerate run: './tools/run gen'
+#
+# Do not modify this file directly
+################################################################################
+
+if(TINT_BUILD_IR_BINARY)
+################################################################################
+# Target: tint_lang_core_ir_binary
+# Kind: lib
+# Condition: TINT_BUILD_IR_BINARY
+################################################################################
+tint_add_target(tint_lang_core_ir_binary lib
+ lang/core/ir/binary/decode.cc
+ lang/core/ir/binary/decode.h
+ lang/core/ir/binary/encode.cc
+ lang/core/ir/binary/encode.h
+)
+
+tint_target_add_dependencies(tint_lang_core_ir_binary lib
+ tint_api_common
+ tint_lang_core
+ tint_lang_core_constant
+ tint_lang_core_intrinsic
+ 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
+)
+
+if(TINT_BUILD_IR_BINARY)
+ tint_target_add_dependencies(tint_lang_core_ir_binary lib
+ tint_lang_core_ir_binary_proto
+ )
+endif(TINT_BUILD_IR_BINARY)
+
+endif(TINT_BUILD_IR_BINARY)
+if(TINT_BUILD_IR_BINARY)
+################################################################################
+# Target: tint_lang_core_ir_binary_proto
+# Kind: proto
+# Condition: TINT_BUILD_IR_BINARY
+################################################################################
+tint_add_target(tint_lang_core_ir_binary_proto proto
+ lang/core/ir/binary/ir.proto
+)
+
+endif(TINT_BUILD_IR_BINARY)
+if(TINT_BUILD_IR_BINARY)
+################################################################################
+# Target: tint_lang_core_ir_binary_test
+# Kind: test
+# Condition: TINT_BUILD_IR_BINARY
+################################################################################
+tint_add_target(tint_lang_core_ir_binary_test test
+ lang/core/ir/binary/roundtrip_test.cc
+)
+
+tint_target_add_dependencies(tint_lang_core_ir_binary_test test
+ tint_api_common
+ tint_lang_core
+ tint_lang_core_constant
+ tint_lang_core_intrinsic
+ tint_lang_core_ir
+ tint_lang_core_ir_test
+ 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_core_ir_binary_test test
+ "gtest"
+)
+
+if(TINT_BUILD_IR_BINARY)
+ tint_target_add_dependencies(tint_lang_core_ir_binary_test test
+ tint_lang_core_ir_binary
+ )
+endif(TINT_BUILD_IR_BINARY)
+
+endif(TINT_BUILD_IR_BINARY)
\ No newline at end of file
diff --git a/src/tint/lang/core/ir/binary/BUILD.gn b/src/tint/lang/core/ir/binary/BUILD.gn
new file mode 100644
index 0000000..d0333c8
--- /dev/null
+++ b/src/tint/lang/core/ir/binary/BUILD.gn
@@ -0,0 +1,118 @@
+# 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.
+
+################################################################################
+# File generated by 'tools/src/cmd/gen' using the template:
+# tools/src/cmd/gen/build/BUILD.gn.tmpl
+#
+# To regenerate run: './tools/run gen'
+#
+# Do not modify this file directly
+################################################################################
+
+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_ir_binary) {
+ libtint_source_set("binary") {
+ sources = [
+ "decode.cc",
+ "decode.h",
+ "encode.cc",
+ "encode.h",
+ ]
+ deps = [
+ "${tint_src_dir}/api/common",
+ "${tint_src_dir}/lang/core",
+ "${tint_src_dir}/lang/core/constant",
+ "${tint_src_dir}/lang/core/intrinsic",
+ "${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_ir_binary) {
+ deps += [ "${tint_src_dir}/lang/core/ir/binary:proto" ]
+ }
+ }
+}
+if (tint_build_ir_binary) {
+ tint_proto_library("proto") {
+ sources = [ "ir.proto" ]
+ deps = []
+ }
+}
+if (tint_build_unittests) {
+ if (tint_build_ir_binary) {
+ tint_unittests_source_set("unittests") {
+ sources = [ "roundtrip_test.cc" ]
+ 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/intrinsic",
+ "${tint_src_dir}/lang/core/ir",
+ "${tint_src_dir}/lang/core/ir:unittests",
+ "${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_ir_binary) {
+ deps += [ "${tint_src_dir}/lang/core/ir/binary" ]
+ }
+ }
+ }
+}
diff --git a/src/tint/lang/core/ir/binary/decode.cc b/src/tint/lang/core/ir/binary/decode.cc
new file mode 100644
index 0000000..cf9576f
--- /dev/null
+++ b/src/tint/lang/core/ir/binary/decode.cc
@@ -0,0 +1,235 @@
+// 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/core/ir/binary/decode.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/utils/macros/compiler.h"
+
+TINT_BEGIN_DISABLE_PROTOBUF_WARNINGS();
+#include "src/tint/lang/core/ir/binary/ir.pb.h"
+TINT_END_DISABLE_PROTOBUF_WARNINGS();
+
+using namespace tint::core::fluent_types; // NOLINT
+
+namespace tint::core::ir::binary {
+namespace {
+
+struct Decoder {
+ pb::Module& mod_in_;
+ Module& mod_out_;
+ Hashmap<uint32_t, ir::Block*, 32> blocks_{};
+ Hashmap<uint32_t, const type::Type*, 32> types_{};
+ Hashmap<uint32_t, ir::Value*, 32> values_{};
+ Builder b{mod_out_};
+
+ void Decode() {
+ // Build all the functions in a separate pass, before we decode them.
+ // This allows for forward references, while preserving function declaration order.
+ for (size_t i = 0, n = static_cast<size_t>(mod_in_.functions().size()); i < n; i++) {
+ b.ir.functions.Push(b.ir.values.Create<ir::Function>());
+ }
+ for (size_t i = 0, n = static_cast<size_t>(mod_in_.functions().size()); i < n; i++) {
+ Function(b.ir.functions[i], mod_in_.functions()[static_cast<int>(i)]);
+ }
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Functions
+ ////////////////////////////////////////////////////////////////////////////
+ void Function(ir::Function* fn_out, const pb::Function& fn_in) {
+ if (!fn_in.name().empty()) {
+ b.ir.SetName(fn_out, fn_in.name());
+ }
+ fn_out->SetReturnType(Type(fn_in.return_type()));
+ if (fn_in.has_pipeline_stage()) {
+ fn_out->SetStage(PipelineStage(fn_in.pipeline_stage()));
+ }
+ if (fn_in.has_workgroup_size()) {
+ auto& wg_size_in = fn_in.workgroup_size();
+ fn_out->SetWorkgroupSize(wg_size_in.x(), wg_size_in.y(), wg_size_in.z());
+ }
+
+ Vector<FunctionParam*, 8> params_out;
+ for (auto param_in : fn_in.parameters()) {
+ params_out.Push(ValueAs<ir::FunctionParam>(param_in));
+ }
+ fn_out->SetParams(std::move(params_out));
+ fn_out->SetBlock(Block(fn_in.block()));
+ }
+
+ Function::PipelineStage PipelineStage(pb::PipelineStage stage) {
+ switch (stage) {
+ case pb::PipelineStage::Compute:
+ return Function::PipelineStage::kCompute;
+ case pb::PipelineStage::Fragment:
+ return Function::PipelineStage::kFragment;
+ case pb::PipelineStage::Vertex:
+ return Function::PipelineStage::kVertex;
+ default:
+ TINT_ICE() << "unhandled PipelineStage: " << stage;
+ return Function::PipelineStage::kCompute;
+ }
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Blocks
+ ////////////////////////////////////////////////////////////////////////////
+ ir::Block* Block(uint32_t id) {
+ if (id == 0) {
+ return nullptr;
+ }
+ return blocks_.GetOrCreate(id, [&] {
+ auto& block_in = mod_in_.blocks().at(static_cast<int>(id) - 1);
+ auto* block_out = b.Block();
+ for (auto& inst : block_in.instructions()) {
+ block_out->Append(Instruction(inst));
+ }
+ return block_out;
+ });
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Instructions
+ ////////////////////////////////////////////////////////////////////////////
+ ir::Instruction* Instruction(const pb::Instruction& inst_in) {
+ ir::Instruction* inst_out = nullptr;
+ switch (inst_in.kind()) {
+ case pb::InstructionKind::Return:
+ inst_out = b.ir.instructions.Create<ir::Return>();
+ break;
+ default:
+ TINT_UNIMPLEMENTED() << inst_in.kind();
+ break;
+ }
+ TINT_ASSERT_OR_RETURN_VALUE(inst_out, nullptr);
+
+ return inst_out;
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Types
+ ////////////////////////////////////////////////////////////////////////////
+ const type::Type* Type(uint32_t id) {
+ if (id == 0) {
+ return nullptr;
+ }
+ return types_.GetOrCreate(id, [&]() -> const type::Type* {
+ auto& ty_in = mod_in_.types().at(static_cast<int>(id) - 1);
+ switch (ty_in.kind_case()) {
+ case pb::TypeDecl::KindCase::kBasic:
+ switch (ty_in.basic()) {
+ case pb::BasicType::void_:
+ return mod_out_.Types().Get<void>();
+ case pb::BasicType::bool_:
+ return mod_out_.Types().Get<bool>();
+ case pb::BasicType::i32:
+ return mod_out_.Types().Get<i32>();
+ case pb::BasicType::u32:
+ return mod_out_.Types().Get<u32>();
+ case pb::BasicType::f32:
+ return mod_out_.Types().Get<f32>();
+ default:
+ TINT_ICE() << "invalid BasicType: " << ty_in.basic();
+ return nullptr;
+ }
+ case pb::TypeDecl::KindCase::kVector:
+ case pb::TypeDecl::KindCase::kMatrix:
+ case pb::TypeDecl::KindCase::kArray:
+ case pb::TypeDecl::KindCase::kAtomic:
+ TINT_UNIMPLEMENTED() << ty_in.kind_case();
+ return nullptr;
+
+ case pb::TypeDecl::KindCase::KIND_NOT_SET:
+ break;
+ }
+ TINT_ICE() << "invalid TypeDecl.kind";
+ return nullptr;
+ });
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Values
+ ////////////////////////////////////////////////////////////////////////////
+ ir::Value* Value(uint32_t id) {
+ if (id == 0) {
+ return nullptr;
+ }
+ return values_.GetOrCreate(id, [&]() -> ir::Value* {
+ auto& val_in = mod_in_.values().at(static_cast<int>(id) - 1);
+ auto* type = Type(val_in.type());
+ ir::Value* val_out = nullptr;
+ switch (val_in.kind()) {
+ case pb::ValueKind::instruction_result:
+ val_out = b.InstructionResult(type);
+ break;
+ case pb::ValueKind::function_parameter:
+ val_out = b.FunctionParam(type);
+ break;
+ default:
+ TINT_ICE() << "invalid TypeDecl.kind";
+ return nullptr;
+ }
+ if (val_in.has_name()) {
+ mod_out_.SetName(val_out, val_in.name());
+ }
+ return val_out;
+ });
+ }
+
+ template <typename T>
+ T* ValueAs(uint32_t id) {
+ auto* value = Value(id);
+ if (auto cast = value->As<T>(); TINT_LIKELY(cast)) {
+ return cast;
+ }
+ TINT_ICE() << "Value " << id << " is " << value->TypeInfo().name << " expected "
+ << TypeInfo::Of<T>().name;
+ return nullptr;
+ }
+};
+
+} // namespace
+
+Result<Module> Decode(Slice<const std::byte> encoded) {
+ GOOGLE_PROTOBUF_VERIFY_VERSION;
+
+ pb::Module mod_in;
+ if (!mod_in.ParseFromArray(encoded.data, static_cast<int>(encoded.len))) {
+ return Failure{"failed to deserialize protobuf"};
+ }
+
+ Module mod_out;
+ Decoder{mod_in, mod_out}.Decode();
+
+ return mod_out;
+}
+
+} // namespace tint::core::ir::binary
diff --git a/src/tint/lang/core/ir/binary/decode.h b/src/tint/lang/core/ir/binary/decode.h
new file mode 100644
index 0000000..d4810b3
--- /dev/null
+++ b/src/tint/lang/core/ir/binary/decode.h
@@ -0,0 +1,44 @@
+// 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_CORE_IR_BINARY_DECODE_H_
+#define SRC_TINT_LANG_CORE_IR_BINARY_DECODE_H_
+
+#include "src/tint/utils/result/result.h"
+
+// Forward declarartion
+namespace tint::core::ir {
+class Module;
+} // namespace tint::core::ir
+
+namespace tint::core::ir::binary {
+
+Result<Module> Decode(Slice<const std::byte> encoded);
+
+} // namespace tint::core::ir::binary
+
+#endif // SRC_TINT_LANG_CORE_IR_BINARY_DECODE_H_
diff --git a/src/tint/lang/core/ir/binary/encode.cc b/src/tint/lang/core/ir/binary/encode.cc
new file mode 100644
index 0000000..340ef66
--- /dev/null
+++ b/src/tint/lang/core/ir/binary/encode.cc
@@ -0,0 +1,192 @@
+// 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/core/ir/binary/encode.h"
+
+#include "src/tint/lang/core/ir/function_param.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/core/ir/return.h"
+#include "src/tint/lang/core/type/bool.h"
+#include "src/tint/lang/core/type/f32.h"
+#include "src/tint/lang/core/type/i32.h"
+#include "src/tint/lang/core/type/u32.h"
+#include "src/tint/lang/core/type/void.h"
+#include "src/tint/utils/macros/compiler.h"
+#include "src/tint/utils/rtti/switch.h"
+
+TINT_BEGIN_DISABLE_PROTOBUF_WARNINGS();
+#include "src/tint/lang/core/ir/binary/ir.pb.h"
+TINT_END_DISABLE_PROTOBUF_WARNINGS();
+
+namespace tint::core::ir::binary {
+namespace {
+struct Encoder {
+ const Module& mod_in_;
+ pb::Module& mod_out_;
+ Hashmap<const core::ir::Function*, uint32_t, 32> functions_{};
+ Hashmap<const core::ir::Block*, uint32_t, 32> blocks_{};
+ Hashmap<const core::type::Type*, uint32_t, 32> types_{};
+ Hashmap<const core::ir::Value*, uint32_t, 32> values_{};
+
+ void Encode() {
+ Vector<pb::Function*, 8> fns_out;
+ for (auto& fn_in : mod_in_.functions) {
+ uint32_t id = static_cast<uint32_t>(fns_out.Length() + 1);
+ fns_out.Push(mod_out_.add_functions());
+ functions_.Add(fn_in, id);
+ }
+ for (size_t i = 0, n = mod_in_.functions.Length(); i < n; i++) {
+ Function(fns_out[i], mod_in_.functions[i]);
+ }
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Functions
+ ////////////////////////////////////////////////////////////////////////////
+ void Function(pb::Function* fn_out, const ir::Function* fn_in) {
+ if (auto name = mod_in_.NameOf(fn_in)) {
+ fn_out->set_name(name.Name());
+ }
+ fn_out->set_return_type(Type(fn_in->ReturnType()));
+ if (fn_in->Stage() != Function::PipelineStage::kUndefined) {
+ fn_out->set_pipeline_stage(PipelineStage(fn_in->Stage()));
+ }
+ if (auto wg_size_in = fn_in->WorkgroupSize()) {
+ auto& wg_size_out = *fn_out->mutable_workgroup_size();
+ wg_size_out.set_x((*wg_size_in)[0]);
+ wg_size_out.set_y((*wg_size_in)[1]);
+ wg_size_out.set_z((*wg_size_in)[2]);
+ }
+ for (auto* param_in : fn_in->Params()) {
+ fn_out->add_parameters(FunctionParam(param_in));
+ }
+ fn_out->set_block(Block(fn_in->Block()));
+ }
+
+ pb::PipelineStage PipelineStage(Function::PipelineStage stage) {
+ switch (stage) {
+ case Function::PipelineStage::kCompute:
+ return pb::PipelineStage::Compute;
+ case Function::PipelineStage::kFragment:
+ return pb::PipelineStage::Fragment;
+ case Function::PipelineStage::kVertex:
+ return pb::PipelineStage::Vertex;
+ default:
+ TINT_ICE() << "unhandled PipelineStage: " << stage;
+ return pb::PipelineStage::Compute;
+ }
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Blocks
+ ////////////////////////////////////////////////////////////////////////////
+ uint32_t Block(const ir::Block* block_in) {
+ if (block_in == nullptr) {
+ return 0;
+ }
+ return blocks_.GetOrCreate(block_in, [&]() -> uint32_t {
+ auto& block_out = *mod_out_.add_blocks();
+ for (auto* inst : *block_in) {
+ Instruction(block_out.add_instructions(), inst);
+ }
+ return static_cast<uint32_t>(blocks_.Count());
+ });
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Instructions
+ ////////////////////////////////////////////////////////////////////////////
+ void Instruction(pb::Instruction* inst_out, const ir::Instruction* inst_in) {
+ auto kind = Switch(
+ inst_in, //
+ [&](const ir::Return*) { return pb::InstructionKind::Return; }, //
+ TINT_ICE_ON_NO_MATCH);
+ inst_out->set_kind(kind);
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Types
+ ////////////////////////////////////////////////////////////////////////////
+ uint32_t Type(const core::type::Type* type) {
+ if (type == nullptr) {
+ return 0;
+ }
+ return types_.GetOrCreate(type, [&]() -> uint32_t {
+ auto basic = tint::Switch<pb::BasicType>(
+ type, //
+ [&](const core::type::Void*) { return pb::BasicType::void_; },
+ [&](const core::type::Bool*) { return pb::BasicType::bool_; },
+ [&](const core::type::I32*) { return pb::BasicType::i32; },
+ [&](const core::type::U32*) { return pb::BasicType::u32; },
+ [&](const core::type::F32*) { return pb::BasicType::f32; }, //
+ TINT_ICE_ON_NO_MATCH);
+ mod_out_.add_types()->set_basic(basic);
+ return static_cast<uint32_t>(types_.Count());
+ });
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Values
+ ////////////////////////////////////////////////////////////////////////////
+ // uint32_t Value(const ir::Value* value) {
+ // return Switch(value, //
+ // [&](const ir::FunctionParam* p) { return FunctionParam(p); });
+ // }
+
+ uint32_t FunctionParam(const ir::FunctionParam* param) {
+ return values_.GetOrCreate(param, [&] {
+ auto& val_out = *mod_out_.add_values();
+ val_out.set_kind(pb::ValueKind::function_parameter);
+ val_out.set_type(Type(param->Type()));
+ if (auto name = mod_in_.NameOf(param); name.IsValid()) {
+ val_out.set_name(name.Name());
+ }
+ return static_cast<uint32_t>(values_.Count());
+ });
+ }
+};
+
+} // namespace
+
+Result<Vector<std::byte, 0>> Encode(const Module& mod_in) {
+ GOOGLE_PROTOBUF_VERIFY_VERSION;
+
+ pb::Module mod_out;
+ Encoder{mod_in, mod_out}.Encode();
+
+ Vector<std::byte, 0> buffer;
+ size_t len = mod_out.ByteSizeLong();
+ buffer.Resize(len);
+ if (len > 0) {
+ if (!mod_out.SerializeToArray(&buffer[0], static_cast<int>(len))) {
+ return Failure{"failed to serialize protobuf"};
+ }
+ }
+ return buffer;
+}
+
+} // namespace tint::core::ir::binary
diff --git a/src/tint/lang/core/ir/binary/encode.h b/src/tint/lang/core/ir/binary/encode.h
new file mode 100644
index 0000000..3cdd520
--- /dev/null
+++ b/src/tint/lang/core/ir/binary/encode.h
@@ -0,0 +1,45 @@
+// 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_CORE_IR_BINARY_ENCODE_H_
+#define SRC_TINT_LANG_CORE_IR_BINARY_ENCODE_H_
+
+#include "src/tint/utils/containers/vector.h"
+#include "src/tint/utils/result/result.h"
+
+// Forward declarartion
+namespace tint::core::ir {
+class Module;
+} // namespace tint::core::ir
+
+namespace tint::core::ir::binary {
+
+Result<Vector<std::byte, 0>> Encode(const Module& module);
+
+} // namespace tint::core::ir::binary
+
+#endif // SRC_TINT_LANG_CORE_IR_BINARY_ENCODE_H_
diff --git a/src/tint/lang/core/ir/binary/ir.proto b/src/tint/lang/core/ir/binary/ir.proto
new file mode 100644
index 0000000..c85bb8e
--- /dev/null
+++ b/src/tint/lang/core/ir/binary/ir.proto
@@ -0,0 +1,128 @@
+// Copyright 2023 The Tint Authors
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+syntax = "proto3";
+
+package tint.core.ir.binary.pb;
+
+message Module {
+ repeated TypeDecl types = 1;
+ repeated Value values = 2;
+ Block root_block = 3;
+ repeated Function functions = 4;
+ repeated Block blocks = 5;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Types
+////////////////////////////////////////////////////////////////////////////////
+message TypeDecl {
+ oneof kind {
+ BasicType basic = 1;
+ VectorType vector = 2;
+ MatrixType matrix = 3;
+ ArrayType array = 4;
+ uint32 atomic = 5; // Module.types
+ // TODO: textures, samplers
+ }
+}
+
+// Non-compound types
+enum BasicType {
+ void = 0;
+ bool = 1;
+ i32 = 2;
+ u32 = 3;
+ f32 = 4;
+ f16 = 5;
+}
+
+message VectorType {
+ uint32 width = 1;
+ uint32 element_type = 2; // Module.types
+}
+
+message MatrixType {
+ uint32 num_columns = 1;
+ uint32 num_width = 2;
+ uint32 element_type = 3; // Module.types
+}
+
+message ArrayType {
+ uint32 count = 1;
+ uint32 element_type = 2; // Module.types
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Values
+////////////////////////////////////////////////////////////////////////////////
+message Value {
+ ValueKind kind = 1;
+ uint32 type = 2; // Module.types
+ optional string name = 3;
+}
+
+enum ValueKind {
+ instruction_result = 0;
+ function_parameter = 1;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Functions
+////////////////////////////////////////////////////////////////////////////////
+message Function {
+ optional string name = 1;
+ uint32 return_type = 2;
+ optional PipelineStage pipeline_stage = 3;
+ optional WorkgroupSize workgroup_size = 4;
+ repeated uint32 parameters = 5; // Module.values
+ uint32 block = 6; // Module.blocks
+}
+
+enum PipelineStage {
+ Compute = 0;
+ Fragment = 1;
+ Vertex = 2;
+}
+
+message WorkgroupSize {
+ uint32 x = 1;
+ uint32 y = 2;
+ uint32 z = 3;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Blocks
+////////////////////////////////////////////////////////////////////////////////
+message Block {
+ repeated uint32 parameters = 1; // Module.values
+ repeated Instruction instructions = 2;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Instructions
+////////////////////////////////////////////////////////////////////////////////
+message Instruction {
+ InstructionKind kind = 1;
+ repeated uint32 operands = 2;
+ repeated uint32 results = 3;
+}
+
+enum InstructionKind {
+ Return = 0;
+ UnaryOp = 1;
+ BinaryOp = 2;
+ Builtin = 3;
+ Constructor = 4;
+}
diff --git a/src/tint/lang/core/ir/binary/roundtrip_test.cc b/src/tint/lang/core/ir/binary/roundtrip_test.cc
new file mode 100644
index 0000000..90a5309
--- /dev/null
+++ b/src/tint/lang/core/ir/binary/roundtrip_test.cc
@@ -0,0 +1,114 @@
+// 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/core/ir/ir_helper_test.h"
+
+#include "src/tint/lang/core/ir/binary/decode.h"
+#include "src/tint/lang/core/ir/binary/encode.h"
+#include "src/tint/lang/core/ir/disassembler.h"
+
+namespace tint::core::ir::binary {
+namespace {
+
+using namespace tint::core::number_suffixes; // NOLINT
+
+template <typename T = testing::Test>
+class IRBinaryRoundtripTestBase : public IRTestParamHelper<T> {
+ public:
+ std::pair<std::string, std::string> Roundtrip() {
+ auto pre = Disassemble(this->mod);
+ auto encoded = Encode(this->mod);
+ if (!encoded) {
+ return {pre, encoded.Failure().reason.str()};
+ }
+ auto decoded = Decode(encoded->Slice());
+ if (!decoded) {
+ return {pre, decoded.Failure().reason.str()};
+ }
+ auto post = Disassemble(decoded.Get());
+ return {pre, post};
+ }
+};
+
+#define RUN_TEST() \
+ { \
+ auto [pre, post] = Roundtrip(); \
+ EXPECT_EQ(pre, post); \
+ } \
+ TINT_REQUIRE_SEMICOLON
+
+using IRBinaryRoundtripTest = IRBinaryRoundtripTestBase<>;
+TEST_F(IRBinaryRoundtripTest, EmptyModule) {
+ RUN_TEST();
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Functions
+////////////////////////////////////////////////////////////////////////////////
+TEST_F(IRBinaryRoundtripTest, Fn_i32_ret) {
+ b.Function("Function", ty.i32());
+ RUN_TEST();
+}
+
+using IRBinaryRoundtripTest_FnPipelineStage = IRBinaryRoundtripTestBase<Function::PipelineStage>;
+TEST_P(IRBinaryRoundtripTest_FnPipelineStage, Test) {
+ b.Function("Function", ty.i32(), GetParam());
+ RUN_TEST();
+}
+INSTANTIATE_TEST_SUITE_P(,
+ IRBinaryRoundtripTest_FnPipelineStage,
+ testing::Values(Function::PipelineStage::kCompute,
+ Function::PipelineStage::kFragment,
+ Function::PipelineStage::kVertex));
+
+TEST_F(IRBinaryRoundtripTest, Fn_WorkgroupSize) {
+ b.Function("Function", ty.i32(), Function::PipelineStage::kCompute,
+ std::array<uint32_t, 3>{1, 2, 3});
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Fn_Parameters) {
+ auto* fn = b.Function("Function", ty.void_());
+ auto* p0 = b.FunctionParam(ty.i32());
+ auto* p1 = b.FunctionParam(ty.u32());
+ auto* p2 = b.FunctionParam(ty.f32());
+ b.ir.SetName(p1, "p1");
+ fn->SetParams({p0, p1, p2});
+ RUN_TEST();
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Instructions
+////////////////////////////////////////////////////////////////////////////////
+TEST_F(IRBinaryRoundtripTest, Return) {
+ auto* fn = b.Function("Function", ty.void_());
+ b.Append(fn->Block(), [&] { b.Return(fn); });
+ RUN_TEST();
+}
+
+} // namespace
+} // namespace tint::core::ir::binary
diff --git a/src/tint/lang/core/ir/function.cc b/src/tint/lang/core/ir/function.cc
index d3edbf1..8c095e0 100644
--- a/src/tint/lang/core/ir/function.cc
+++ b/src/tint/lang/core/ir/function.cc
@@ -36,6 +36,8 @@
namespace tint::core::ir {
+Function::Function() = default;
+
Function::Function(const core::type::Type* rt,
PipelineStage stage,
std::optional<std::array<uint32_t, 3>> wg_size)
diff --git a/src/tint/lang/core/ir/function.h b/src/tint/lang/core/ir/function.h
index 1937673..0193884 100644
--- a/src/tint/lang/core/ir/function.h
+++ b/src/tint/lang/core/ir/function.h
@@ -73,6 +73,9 @@
};
/// Constructor
+ Function();
+
+ /// Constructor
/// @param rt the function return type
/// @param stage the function stage
/// @param wg_size the workgroup_size
@@ -103,6 +106,9 @@
/// @returns the workgroup size information
std::optional<std::array<uint32_t, 3>> WorkgroupSize() const { return workgroup_size_; }
+ /// @param type the return type for the function
+ void SetReturnType(const core::type::Type* type) { return_.type = type; }
+
/// @returns the return type for the function
const core::type::Type* ReturnType() const { return return_.type; }
@@ -165,7 +171,7 @@
void Destroy() override;
private:
- PipelineStage pipeline_stage_;
+ PipelineStage pipeline_stage_ = PipelineStage::kUndefined;
std::optional<std::array<uint32_t, 3>> workgroup_size_;
struct {
diff --git a/src/tint/lang/core/ir/return.cc b/src/tint/lang/core/ir/return.cc
index f105afc..fb6cd8a 100644
--- a/src/tint/lang/core/ir/return.cc
+++ b/src/tint/lang/core/ir/return.cc
@@ -37,6 +37,8 @@
namespace tint::core::ir {
+Return::Return() = default;
+
Return::Return(Function* func) {
AddOperand(Return::kFunctionOperandOffset, func);
}
diff --git a/src/tint/lang/core/ir/return.h b/src/tint/lang/core/ir/return.h
index 4fcdc86..4e8827b 100644
--- a/src/tint/lang/core/ir/return.h
+++ b/src/tint/lang/core/ir/return.h
@@ -49,6 +49,9 @@
/// The offset in Operands() for the return argument
static constexpr size_t kArgsOperandOffset = 1;
+ /// Constructor (no operands)
+ Return();
+
/// Constructor (no return value)
/// @param func the function being returned
explicit Return(Function* func);
diff --git a/src/tint/lang/core/type/manager.h b/src/tint/lang/core/type/manager.h
index ada8c7e..8f361da 100644
--- a/src/tint/lang/core/type/manager.h
+++ b/src/tint/lang/core/type/manager.h
@@ -121,6 +121,8 @@
return Get<core::type::F16>(std::forward<ARGS>(args)...);
} else if constexpr (std::is_same_v<T, bool>) {
return Get<core::type::Bool>(std::forward<ARGS>(args)...);
+ } else if constexpr (std::is_same_v<T, void>) {
+ return Get<core::type::Void>(std::forward<ARGS>(args)...);
} else if constexpr (core::fluent_types::IsVector<T>) {
return vec<typename T::type, T::width>(std::forward<ARGS>(args)...);
} else if constexpr (core::fluent_types::IsMatrix<T>) {
diff --git a/src/tint/utils/macros/compiler.h b/src/tint/utils/macros/compiler.h
index b1fc791..623e4c4 100644
--- a/src/tint/utils/macros/compiler.h
+++ b/src/tint/utils/macros/compiler.h
@@ -45,8 +45,14 @@
#define TINT_DISABLE_WARNING_WEAK_VTABLES /* currently no-op */
#define TINT_DISABLE_WARNING_FLOAT_EQUAL /* currently no-op */
#define TINT_DISABLE_WARNING_DEPRECATED __pragma(warning(disable : 4996))
-#define TINT_DISABLE_WARNING_RESERVED_IDENTIFIER /* currently no-op */
-#define TINT_DISABLE_WARNING_UNUSED_VALUE /* currently no-op */
+#define TINT_DISABLE_WARNING_RESERVED_IDENTIFIER /* currently no-op */
+#define TINT_DISABLE_WARNING_RESERVED_MACRO_IDENTIFIER /* currently no-op */
+#define TINT_DISABLE_WARNING_UNUSED_VALUE /* currently no-op */
+#define TINT_DISABLE_WARNING_UNUSED_PARAMETER __pragma(warning(disable : 4100))
+#define TINT_DISABLE_WARNING_SHADOW_FIELD_IN_CONSTRUCTOR /* currently no-op */
+#define TINT_DISABLE_WARNING_EXTRA_SEMICOLON /* currently no-op */
+#define TINT_DISABLE_WARNING_ZERO_AS_NULLPTR /* currently no-op */
+#define TINT_DISABLE_WARNING_MISSING_DESTRUCTOR_OVERRIDE /* currently no-op */
// clang-format off
#define TINT_BEGIN_DISABLE_WARNING(name) \
@@ -58,14 +64,20 @@
TINT_REQUIRE_SEMICOLON
// clang-format on
+#define TINT_BEGIN_DISABLE_PROTOBUF_WARNINGS() \
+ __pragma(warning(push)) TINT_DISABLE_WARNING_UNUSED_PARAMETER TINT_REQUIRE_SEMICOLON
+#define TINT_END_DISABLE_PROTOBUF_WARNINGS() __pragma(warning(pop)) TINT_REQUIRE_SEMICOLON
+
#define TINT_UNLIKELY(x) x /* currently no-op */
#define TINT_LIKELY(x) x /* currently no-op */
+
#elif defined(__clang__)
////////////////////////////////////////////////////////////////////////////////
// Clang
////////////////////////////////////////////////////////////////////////////////
-#define TINT_DISABLE_WARNING_CONSTANT_OVERFLOW /* currently no-op */
-#define TINT_DISABLE_WARNING_MAYBE_UNINITIALIZED /* currently no-op */
+#define TINT_DISABLE_WARNING_CONSTANT_OVERFLOW /* currently no-op */
+#define TINT_DISABLE_WARNING_MAYBE_UNINITIALIZED \
+ _Pragma("clang diagnostic ignored \"-Wconditional-uninitialized\"")
#define TINT_DISABLE_WARNING_NEWLINE_EOF _Pragma("clang diagnostic ignored \"-Wnewline-eof\"")
#define TINT_DISABLE_WARNING_OLD_STYLE_CAST _Pragma("clang diagnostic ignored \"-Wold-style-cast\"")
#define TINT_DISABLE_WARNING_SIGN_CONVERSION \
@@ -76,13 +88,45 @@
#define TINT_DISABLE_WARNING_DEPRECATED /* currently no-op */
#define TINT_DISABLE_WARNING_RESERVED_IDENTIFIER \
_Pragma("clang diagnostic ignored \"-Wreserved-identifier\"")
+#define TINT_DISABLE_WARNING_RESERVED_MACRO_IDENTIFIER \
+ _Pragma("clang diagnostic ignored \"-Wreserved-macro-identifier\"")
#define TINT_DISABLE_WARNING_UNUSED_VALUE _Pragma("clang diagnostic ignored \"-Wunused-value\"")
+#define TINT_DISABLE_WARNING_UNUSED_PARAMETER \
+ _Pragma("clang diagnostic ignored \"-Wunused-parameter\"")
+#define TINT_DISABLE_WARNING_SHADOW_FIELD_IN_CONSTRUCTOR \
+ _Pragma("clang diagnostic ignored \"-Wshadow-field-in-constructor\"")
+#define TINT_DISABLE_WARNING_EXTRA_SEMICOLON \
+ _Pragma("clang diagnostic ignored \"-Wextra-semi-stmt\"")
+#define TINT_DISABLE_WARNING_ZERO_AS_NULLPTR \
+ _Pragma("clang diagnostic ignored \"-Wzero-as-null-pointer-constant\"")
+#define TINT_DISABLE_WARNING_MISSING_DESTRUCTOR_OVERRIDE \
+ _Pragma("clang diagnostic ignored \"-Wsuggest-destructor-override\"") \
+ _Pragma("clang diagnostic ignored \"-Winconsistent-missing-destructor-override\"")
// clang-format off
+#define TINT_BEGIN_DISABLE_PROTOBUF_WARNINGS() \
+ _Pragma("clang diagnostic push") \
+ TINT_DISABLE_WARNING_EXTRA_SEMICOLON \
+ TINT_DISABLE_WARNING_MAYBE_UNINITIALIZED \
+ TINT_DISABLE_WARNING_MISSING_DESTRUCTOR_OVERRIDE \
+ TINT_DISABLE_WARNING_RESERVED_IDENTIFIER \
+ TINT_DISABLE_WARNING_RESERVED_MACRO_IDENTIFIER \
+ TINT_DISABLE_WARNING_SHADOW_FIELD_IN_CONSTRUCTOR \
+ TINT_DISABLE_WARNING_SIGN_CONVERSION \
+ TINT_DISABLE_WARNING_UNUSED_PARAMETER \
+ TINT_DISABLE_WARNING_WEAK_VTABLES \
+ TINT_DISABLE_WARNING_ZERO_AS_NULLPTR \
+ TINT_REQUIRE_SEMICOLON
+
+#define TINT_END_DISABLE_PROTOBUF_WARNINGS() \
+ _Pragma("clang diagnostic pop") \
+ TINT_REQUIRE_SEMICOLON
+
#define TINT_BEGIN_DISABLE_WARNING(name) \
_Pragma("clang diagnostic push") \
TINT_CONCAT(TINT_DISABLE_WARNING_, name) \
TINT_REQUIRE_SEMICOLON
+
#define TINT_END_DISABLE_WARNING(name) \
_Pragma("clang diagnostic pop") \
TINT_REQUIRE_SEMICOLON
@@ -97,15 +141,26 @@
#define TINT_DISABLE_WARNING_CONSTANT_OVERFLOW /* currently no-op */
#define TINT_DISABLE_WARNING_MAYBE_UNINITIALIZED \
_Pragma("GCC diagnostic ignored \"-Wmaybe-uninitialized\"")
-#define TINT_DISABLE_WARNING_NEWLINE_EOF /* currently no-op */
-#define TINT_DISABLE_WARNING_OLD_STYLE_CAST /* currently no-op */
-#define TINT_DISABLE_WARNING_SIGN_CONVERSION /* currently no-op */
-#define TINT_DISABLE_WARNING_UNREACHABLE_CODE /* currently no-op */
-#define TINT_DISABLE_WARNING_WEAK_VTABLES /* currently no-op */
-#define TINT_DISABLE_WARNING_FLOAT_EQUAL /* currently no-op */
-#define TINT_DISABLE_WARNING_DEPRECATED /* currently no-op */
-#define TINT_DISABLE_WARNING_RESERVED_IDENTIFIER /* currently no-op */
+#define TINT_DISABLE_WARNING_NEWLINE_EOF /* currently no-op */
+#define TINT_DISABLE_WARNING_OLD_STYLE_CAST /* currently no-op */
+#define TINT_DISABLE_WARNING_SIGN_CONVERSION /* currently no-op */
+#define TINT_DISABLE_WARNING_UNREACHABLE_CODE /* currently no-op */
+#define TINT_DISABLE_WARNING_WEAK_VTABLES /* currently no-op */
+#define TINT_DISABLE_WARNING_FLOAT_EQUAL /* currently no-op */
+#define TINT_DISABLE_WARNING_DEPRECATED /* currently no-op */
+#define TINT_DISABLE_WARNING_RESERVED_IDENTIFIER /* currently no-op */
+#define TINT_DISABLE_WARNING_RESERVED_MACRO_IDENTIFIER /* currently no-op */
#define TINT_DISABLE_WARNING_UNUSED_VALUE _Pragma("GCC diagnostic ignored \"-Wunused-value\"")
+#define TINT_DISABLE_WARNING_UNUSED_PARAMETER \
+ _Pragma("GCC diagnostic ignored \"-Wunused-parameter\"")
+#define TINT_DISABLE_WARNING_SHADOW_FIELD_IN_CONSTRUCTOR /* currently no-op */
+#define TINT_DISABLE_WARNING_EXTRA_SEMICOLON /* currently no-op */
+#define TINT_DISABLE_WARNING_ZERO_AS_NULLPTR /* currently no-op */
+#define TINT_DISABLE_WARNING_MISSING_DESTRUCTOR_OVERRIDE /* currently no-op */
+
+#define TINT_BEGIN_DISABLE_PROTOBUF_WARNINGS() \
+ _Pragma("GCC diagnostic push") TINT_DISABLE_WARNING_UNUSED_PARAMETER TINT_REQUIRE_SEMICOLON
+#define TINT_END_DISABLE_PROTOBUF_WARNINGS() _Pragma("GCC diagnostic pop") TINT_REQUIRE_SEMICOLON
// clang-format off
#define TINT_BEGIN_DISABLE_WARNING(name) \
@@ -125,6 +180,8 @@
////////////////////////////////////////////////////////////////////////////////
#define TINT_BEGIN_DISABLE_WARNING(name) TINT_REQUIRE_SEMICOLON
#define TINT_END_DISABLE_WARNING(name) TINT_REQUIRE_SEMICOLON
+#define TINT_BEGIN_DISABLE_PROTOBUF_WARNINGS() TINT_REQUIRE_SEMICOLON
+#define TINT_END_DISABLE_PROTOBUF_WARNINGS() TINT_REQUIRE_SEMICOLON
#define TINT_UNLIKELY(x) x
#define TINT_LIKELY(x) x
diff --git a/third_party/protobuf.cmake b/third_party/protobuf.cmake
index 731068f..f1686e8 100644
--- a/third_party/protobuf.cmake
+++ b/third_party/protobuf.cmake
@@ -158,6 +158,8 @@
set(COMMENT "${COMMENT}, plugin-options: ${PLUGIN_OPTIONS}")
endif()
+ file(MAKE_DIRECTORY "${ARGS_PROTOC_OUT_DIR}/${REL_DIR}")
+
add_custom_command(
OUTPUT ${GENERATED_SRCS}
COMMAND protobuf::protoc
