Import Tint changes from Dawn
Contains manual fixes for:
* CMakeLists.txt
* DEPS
* scripts/tint_overrides_with_defaults.gni
And one new files:
* third_party/protobuf.cmake
Also bumps Kokoro Clang version to 13.0.1, to match the Dawn bots.
Changes:
- d86e9a765cb57b9e3c8476731a4de964e2dd25a3 [spirv-reader] Pass workgroup_id as argument by James Price <jrprice@google.com>
- 3f11e1e7dff3c0a0a0ecc6d3ae23f9070f4c0aa8 [tint][ir] Serialize Swizzle instruction by Ben Clayton <bclayton@google.com>
- 59bc89f8fe3967eee8eef4465eea31dc22fe1a79 Revert "clang-format" by Stephen White <senorblanco@google.com>
- f91e88ad4a2306f6b65dd97849d33a2c515af6ac clang-format by Ben Clayton <bclayton@google.com>
- 84a7cbf8ebc415970472d2b32eaef27a417b59bc [tint][ir] Serialize [Load|Store]VectorElement instructions by Ben Clayton <bclayton@google.com>
- 53c0fa9a80310e69b5b1b6cae96f9469f20b654e Enable language extension `packed_4x8_integer_dot_product` by Jiawei Shao <jiawei.shao@intel.com>
- c668c1690ed23db80238aabbe42002a15ee2eb0e [tint][ir] Serialize Unary and Binary instructions by Ben Clayton <bclayton@google.com>
- dacc2637282217ac5d9a2c2a45a7ca2c54dbac61 [tint][ir] Serialize Var binding points by Ben Clayton <bclayton@google.com>
- 7d9a66e3d2a562819b750d4aa1bdbb912de770b5 [tint][ir] Serialize the root block by Ben Clayton <bclayton@google.com>
- cba30d2b886939bc7b84f7b5687236ff9ca83046 [tint][ir] Serialize Load and Store instructions by Ben Clayton <bclayton@google.com>
- 112549bea523b3cfc1c2c0e0f37b2e2f442f2ca5 [tint][ir] Serialize UserCall by Ben Clayton <bclayton@google.com>
- 0bc302670162c75ef9417f4e1c3d775bc96cdaa0 [tint][ir] Serialize fixed-size array types by Ben Clayton <bclayton@google.com>
- 7f9721daadcf26d752f56847ff319644d1c5913e [tint][ir] Serialize var instructions and pointer types. by Ben Clayton <bclayton@google.com>
- 3f620ab5d809ad4c2b90a165a1ec4241dbbb7913 [tint][ir] Serialize Access instructions by Ben Clayton <bclayton@google.com>
- fc5664ab75a62d969cdc4d3870a073bf4ba5e0c4 [tint][ir] Serialize vectors and matrices by Ben Clayton <bclayton@google.com>
- 069238f0d8b1eee5daaff62a7f5978f69dcab954 [tint][ir][validator] 'source' -> 'instruction' by Ben Clayton <bclayton@google.com>
- 4eff8bec72c2900d4a8b9219535433016ceef279 [tint][ir] Serialize constant scalars and lets by Ben Clayton <bclayton@google.com>
- 85b673578773a7104615183fce70b0dbf6462d58 [tint][ir] Serialize Discard instructions by Ben Clayton <bclayton@google.com>
- 2388287c0ecacdb37a09ab2b73f10a7c66eff7bc [tint][fuzz] Don't attempt to memcpy from nullptr by Ben Clayton <bclayton@google.com>
- 3ee81db189e8e5cebbec8c6f33f9d15e9124ed67 Fix license: it's BSD-3 not Apache 2 by David Neto <dneto@google.com>
- 0a8dc81b6184478c1b4c83bda1ecb71da660bb04 [tint][utils] Make bytes::Reader an interface by Ben Clayton <bclayton@google.com>
- 44a29b5e3743ad38a8d5147119471d28fabd44e5 [tint][gn] Fix include_dirs in tint_common_config by Ben Clayton <bclayton@google.com>
- ae7240f3642136f636a49d07dfc1c87cd29cf0b9 Vulkan: Use the polyfill of dot4{I|U}8Packed() when neces... by Jiawei Shao <jiawei.shao@intel.com>
- 648a2b32c94e2bbda7f12076ff57ace9699e602e OpenGL: Support dot4I8Packed() and dot4U8Packed() by Jiawei Shao <jiawei.shao@intel.com>
- 90c7cc2304a6b629e5c5fdef18b4e1a1b3a2fbaa [tint][ir] Begin building proto-based serializer by Ben Clayton <bclayton@google.com>
GitOrigin-RevId: d86e9a765cb57b9e3c8476731a4de964e2dd25a3
Change-Id: Ib656c7a3768d7a278d640f35ce2534b8c50f73d2
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/165140
Commit-Queue: Ben Clayton <bclayton@google.com>
Auto-Submit: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/CMakeLists.txt b/CMakeLists.txt
index fc95d64..0deacd3 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -26,6 +26,8 @@
set(CMAKE_BUILD_TYPE "Debug")
endif()
+set(DAWN_BUILD_GEN_DIR "${CMAKE_CURRENT_BINARY_DIR}/gen")
+
# TINT_IS_SUBPROJECT is 1 if added via add_subdirectory() from another project.
get_directory_property(TINT_IS_SUBPROJECT PARENT_DIRECTORY)
if(TINT_IS_SUBPROJECT)
@@ -74,6 +76,7 @@
set_if_not_defined(TINT_MARKUPSAFE_DIR "${TINT_THIRD_PARTY_DIR}/markupsafe" "Directory in which to find MarkupSafe")
set_if_not_defined(TINT_KHRONOS_DIR "${TINT_THIRD_PARTY_DIR}/khronos" "Directory in which to find Khronos GL headers")
set_if_not_defined(TINT_SWIFTSHADER_DIR "${TINT_THIRD_PARTY_DIR}/swiftshader" "Directory in which to find swiftshader")
+set_if_not_defined(TINT_PROTOBUF_DIR "${TINT_THIRD_PARTY_DIR}/protobuf" "Directory in which to find protobuf")
set_if_not_defined(TINT_SPIRV_TOOLS_DIR "${TINT_VULKAN_DEPS_DIR}/spirv-tools/src" "Directory in which to find SPIRV-Tools")
set_if_not_defined(TINT_SPIRV_HEADERS_DIR "${TINT_VULKAN_DEPS_DIR}/spirv-headers/src" "Directory in which to find SPIRV-Headers")
@@ -392,6 +395,11 @@
endif(TINT_EMIT_COVERAGE)
endfunction()
+if (EXISTS "${TINT_PROTOBUF_DIR}/cmake")
+ # Needs to come before SPIR-V Tools
+ include("third_party/protobuf.cmake")
+endif()
+
add_subdirectory(third_party)
add_subdirectory(src/tint)
diff --git a/DEPS b/DEPS
index 4138b73..55f27cc 100644
--- a/DEPS
+++ b/DEPS
@@ -129,7 +129,7 @@
},
'third_party/protobuf': {
- 'url': '{chromium_git}/external/github.com/protocolbuffers/protobuf.git@2b673bbb57e34fe1bd4570f726fc86b769a3a3d2',
+ 'url': '{chromium_git}/chromium/src/third_party/protobuf@41759e11ec427e29e1a72b9401d2af3f6e02d839',
},
}
diff --git a/kokoro/linux/docker.sh b/kokoro/linux/docker.sh
index a15ff37..eb0f1e9 100755
--- a/kokoro/linux/docker.sh
+++ b/kokoro/linux/docker.sh
@@ -99,7 +99,7 @@
COMMON_CMAKE_FLAGS+=" -DTINT_BUILD_BENCHMARKS=1"
if [ "$BUILD_TOOLCHAIN" == "clang" ]; then
- using clang-10.0.0
+ using clang-13.0.1
COMMON_CMAKE_FLAGS+=" -DTINT_BUILD_FUZZERS=1"
COMMON_CMAKE_FLAGS+=" -DTINT_BUILD_SPIRV_TOOLS_FUZZER=1"
COMMON_CMAKE_FLAGS+=" -DTINT_BUILD_AST_FUZZER=1"
diff --git a/scripts/tint_overrides_with_defaults.gni b/scripts/tint_overrides_with_defaults.gni
index 0edfb87..4e33a50 100644
--- a/scripts/tint_overrides_with_defaults.gni
+++ b/scripts/tint_overrides_with_defaults.gni
@@ -101,6 +101,11 @@
tint_build_syntax_tree_writer = false
}
+ # Build the IR binary serializer
+ if (!defined(tint_build_ir_binary)) {
+ tint_build_ir_binary = tint_has_protobuf
+ }
+
# Build unittests
if (!defined(tint_build_unittests)) {
tint_build_unittests = true
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 6f78959..f8fe3f2 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -42,13 +42,13 @@
# Common - Configs, etc. shared across targets
###############################################################################
+# tint_gen_dir is the relative path of tint_root_dir to '//', under root_gen_dir
+tint_gen_dir = "${root_gen_dir}/" + rebase_path("${tint_root_dir}", "//")
+
config("tint_common_config") {
include_dirs = [
- "${target_gen_dir}",
- "${tint_root_dir}/",
- "${tint_spirv_headers_dir}/include",
- "${tint_spirv_tools_dir}/",
- "${tint_spirv_tools_dir}/include",
+ "${tint_root_dir}",
+ "${tint_gen_dir}",
]
}
@@ -129,7 +129,6 @@
include_dirs = [
"${tint_root_dir}/",
"${tint_root_dir}/include/",
- "${tint_spirv_headers_dir}/include",
]
}
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/fuzz/ir/fuzz.cc b/src/tint/cmd/fuzz/ir/fuzz.cc
index fc96dcc..8934bf2 100644
--- a/src/tint/cmd/fuzz/ir/fuzz.cc
+++ b/src/tint/cmd/fuzz/ir/fuzz.cc
@@ -46,7 +46,6 @@
bool IsUnsupported(const ast::Enable* enable) {
for (auto ext : enable->extensions) {
switch (ext->name) {
- case tint::wgsl::Extension::kChromiumExperimentalDp4A:
case tint::wgsl::Extension::kChromiumExperimentalFullPtrParameters:
case tint::wgsl::Extension::kChromiumExperimentalPixelLocal:
case tint::wgsl::Extension::kChromiumExperimentalPushConstant:
diff --git a/src/tint/cmd/fuzz/ir/fuzz.h b/src/tint/cmd/fuzz/ir/fuzz.h
index 897f6e1..36d8c02 100644
--- a/src/tint/cmd/fuzz/ir/fuzz.h
+++ b/src/tint/cmd/fuzz/ir/fuzz.h
@@ -53,7 +53,10 @@
static IRFuzzer Create(std::string_view name, void (*fn)(core::ir::Module&, ARGS...)) {
if constexpr (sizeof...(ARGS) > 0) {
auto fn_with_decode = [fn](core::ir::Module& module, Slice<const std::byte> data) {
- bytes::Reader reader{data};
+ if (!data.data) {
+ return;
+ }
+ bytes::BufferReader reader{data};
if (auto data_args = bytes::Decode<std::tuple<std::decay_t<ARGS>...>>(reader)) {
auto all_args =
std::tuple_cat(std::tuple<core::ir::Module&>{module}, data_args.Get());
diff --git a/src/tint/cmd/fuzz/wgsl/fuzz.h b/src/tint/cmd/fuzz/wgsl/fuzz.h
index 744ceb5..da2efe4 100644
--- a/src/tint/cmd/fuzz/wgsl/fuzz.h
+++ b/src/tint/cmd/fuzz/wgsl/fuzz.h
@@ -50,7 +50,10 @@
static ProgramFuzzer Create(std::string_view name, void (*fn)(const Program&, ARGS...)) {
if constexpr (sizeof...(ARGS) > 0) {
auto fn_with_decode = [fn](const Program& program, Slice<const std::byte> data) {
- bytes::Reader reader{data};
+ if (!data.data) {
+ return;
+ }
+ bytes::BufferReader reader{data};
if (auto data_args = bytes::Decode<std::tuple<std::decay_t<ARGS>...>>(reader)) {
auto all_args =
std::tuple_cat(std::tuple<const Program&>{program}, data_args.Get());
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/builtin_fn.cc b/src/tint/lang/core/builtin_fn.cc
index 97c1baf..0109b7e 100644
--- a/src/tint/lang/core/builtin_fn.cc
+++ b/src/tint/lang/core/builtin_fn.cc
@@ -689,7 +689,7 @@
f == BuiltinFn::kAtomicCompareExchangeWeak;
}
-bool IsDP4a(BuiltinFn f) {
+bool IsPacked4x8IntegerDotProductBuiltin(BuiltinFn f) {
return f == BuiltinFn::kDot4I8Packed || f == BuiltinFn::kDot4U8Packed;
}
diff --git a/src/tint/lang/core/builtin_fn.cc.tmpl b/src/tint/lang/core/builtin_fn.cc.tmpl
index 376b33f..0a39072 100644
--- a/src/tint/lang/core/builtin_fn.cc.tmpl
+++ b/src/tint/lang/core/builtin_fn.cc.tmpl
@@ -100,7 +100,7 @@
f == BuiltinFn::kAtomicCompareExchangeWeak;
}
-bool IsDP4a(BuiltinFn f) {
+bool IsPacked4x8IntegerDotProductBuiltin(BuiltinFn f) {
return f == BuiltinFn::kDot4I8Packed || f == BuiltinFn::kDot4U8Packed;
}
diff --git a/src/tint/lang/core/builtin_fn.h b/src/tint/lang/core/builtin_fn.h
index 78ba924..badcd4c 100644
--- a/src/tint/lang/core/builtin_fn.h
+++ b/src/tint/lang/core/builtin_fn.h
@@ -468,10 +468,12 @@
/// @returns true if the given `f` is an atomic builtin
bool IsAtomic(BuiltinFn f);
-/// Determines if the given `f` is a DP4a builtin.
+/// Determines if the given `f` is a builtin defined in the language extension
+/// `packed_4x8_integer_dot_product`.
/// @param f the builtin type
-/// @returns true if the given `f` is a DP4a builtin
-bool IsDP4a(BuiltinFn f);
+/// @returns true if the given `f` is a builtin defined in the language extension
+/// `packed_4x8_integer_dot_product`.
+bool IsPacked4x8IntegerDotProductBuiltin(BuiltinFn f);
/// Determines if the given `f` is a subgroup builtin.
/// @param f the builtin type
diff --git a/src/tint/lang/core/builtin_fn.h.tmpl b/src/tint/lang/core/builtin_fn.h.tmpl
index d5198fa..bb08fff 100644
--- a/src/tint/lang/core/builtin_fn.h.tmpl
+++ b/src/tint/lang/core/builtin_fn.h.tmpl
@@ -108,10 +108,12 @@
/// @returns true if the given `f` is an atomic builtin
bool IsAtomic(BuiltinFn f);
-/// Determines if the given `f` is a DP4a builtin.
+/// Determines if the given `f` is a builtin defined in the language extension
+/// `packed_4x8_integer_dot_product`.
/// @param f the builtin type
-/// @returns true if the given `f` is a DP4a builtin
-bool IsDP4a(BuiltinFn f);
+/// @returns true if the given `f` is a builtin defined in the language extension
+/// `packed_4x8_integer_dot_product`.
+bool IsPacked4x8IntegerDotProductBuiltin(BuiltinFn f);
/// Determines if the given `f` is a subgroup builtin.
/// @param f the builtin type
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/access.cc b/src/tint/lang/core/ir/access.cc
index c469a5d..b4434e6 100644
--- a/src/tint/lang/core/ir/access.cc
+++ b/src/tint/lang/core/ir/access.cc
@@ -37,6 +37,8 @@
namespace tint::core::ir {
//! @cond Doxygen_Suppress
+Access::Access() = default;
+
Access::Access(InstructionResult* result, Value* object, VectorRef<Value*> indices) {
AddOperand(Access::kObjectOperandOffset, object);
AddOperands(Access::kIndicesOperandOffset, std::move(indices));
diff --git a/src/tint/lang/core/ir/access.h b/src/tint/lang/core/ir/access.h
index 0bf2caa..6ee0eb6 100644
--- a/src/tint/lang/core/ir/access.h
+++ b/src/tint/lang/core/ir/access.h
@@ -44,6 +44,9 @@
/// The base offset in Operands() for the access indices
static constexpr size_t kIndicesOperandOffset = 1;
+ /// Constructor (no results, no operands)
+ Access();
+
/// Constructor
/// @param result the result value
/// @param object the accessor object
diff --git a/src/tint/lang/core/ir/binary.cc b/src/tint/lang/core/ir/binary.cc
index 41c9972..732f029 100644
--- a/src/tint/lang/core/ir/binary.cc
+++ b/src/tint/lang/core/ir/binary.cc
@@ -34,6 +34,8 @@
namespace tint::core::ir {
+Binary::Binary() = default;
+
Binary::Binary(InstructionResult* result, BinaryOp op, Value* lhs, Value* rhs) : op_(op) {
AddOperand(Binary::kLhsOperandOffset, lhs);
AddOperand(Binary::kRhsOperandOffset, rhs);
diff --git a/src/tint/lang/core/ir/binary.h b/src/tint/lang/core/ir/binary.h
index 6490be9..5f8c26a 100644
--- a/src/tint/lang/core/ir/binary.h
+++ b/src/tint/lang/core/ir/binary.h
@@ -67,6 +67,9 @@
/// The offset in Operands() for the RHS
static constexpr size_t kRhsOperandOffset = 1;
+ /// Constructor (no results, no operands)
+ Binary();
+
/// Constructor
/// @param result the result value
/// @param op the binary operator
@@ -81,6 +84,9 @@
/// @returns the binary operator
BinaryOp Op() const { return op_; }
+ /// @param op the new binary operator
+ void SetOp(BinaryOp op) { op_ = op; }
+
/// @returns the left-hand-side value for the instruction
Value* LHS() { return operands_[kLhsOperandOffset]; }
@@ -97,7 +103,7 @@
std::string FriendlyName() const override { return "binary"; }
private:
- BinaryOp op_;
+ BinaryOp op_ = BinaryOp::kAdd;
};
/// @param kind the enum value
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..6d23588
--- /dev/null
+++ b/src/tint/lang/core/ir/binary/decode.cc
@@ -0,0 +1,597 @@
+// 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/containers/transform.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_;
+ Vector<ir::Block*, 32> blocks_{};
+ Vector<const type::Type*, 32> types_{};
+ Vector<const core::constant::Value*, 32> constant_values_{};
+ Vector<ir::Value*, 32> values_{};
+ Builder b{mod_out_};
+
+ void Decode() {
+ {
+ const size_t n = static_cast<size_t>(mod_in_.types().size());
+ types_.Reserve(n);
+ for (auto& type_in : mod_in_.types()) {
+ types_.Push(CreateType(type_in));
+ }
+ }
+ {
+ const size_t n = static_cast<size_t>(mod_in_.functions().size());
+ mod_out_.functions.Reserve(n);
+ for (auto& fn_in : mod_in_.functions()) {
+ mod_out_.functions.Push(CreateFunction(fn_in));
+ }
+ }
+ {
+ const size_t n = static_cast<size_t>(mod_in_.blocks().size());
+ blocks_.Reserve(n);
+ for (size_t i = 0; i < n; i++) {
+ auto id = static_cast<uint32_t>(i + 1);
+ if (id == mod_in_.root_block()) {
+ blocks_.Push(mod_out_.root_block);
+ } else {
+ auto& block_in = mod_in_.blocks()[static_cast<int>(i)];
+ blocks_.Push(CreateBlock(block_in));
+ }
+ }
+ }
+ {
+ const size_t n = static_cast<size_t>(mod_in_.constant_values().size());
+ constant_values_.Reserve(n);
+ for (auto& value_in : mod_in_.constant_values()) {
+ constant_values_.Push(CreateConstantValue(value_in));
+ }
+ }
+ {
+ const size_t n = static_cast<size_t>(mod_in_.values().size());
+ values_.Reserve(n);
+ for (auto& value_in : mod_in_.values()) {
+ values_.Push(CreateValue(value_in));
+ }
+ }
+ for (size_t i = 0, n = static_cast<size_t>(mod_in_.functions().size()); i < n; i++) {
+ PopulateFunction(mod_out_.functions[i], mod_in_.functions()[static_cast<int>(i)]);
+ }
+ for (size_t i = 0, n = static_cast<size_t>(mod_in_.blocks().size()); i < n; i++) {
+ PopulateBlock(blocks_[i], mod_in_.blocks()[static_cast<int>(i)]);
+ }
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Functions
+ ////////////////////////////////////////////////////////////////////////////
+ ir::Function* CreateFunction(const pb::Function&) {
+ return mod_out_.values.Create<ir::Function>();
+ }
+
+ void PopulateFunction(ir::Function* fn_out, const pb::Function& fn_in) {
+ if (!fn_in.name().empty()) {
+ mod_out_.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()));
+ }
+
+ ir::Function* Function(uint32_t id) { return id > 0 ? mod_out_.functions[id - 1] : nullptr; }
+
+ 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* CreateBlock(const pb::Block&) { return b.Block(); }
+
+ ir::Block* PopulateBlock(ir::Block* block_out, const pb::Block& block_in) {
+ for (auto& inst : block_in.instructions()) {
+ block_out->Append(Instruction(inst));
+ }
+ return block_out;
+ }
+
+ ir::Block* Block(uint32_t id) { return id > 0 ? blocks_[id - 1] : nullptr; }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Instructions
+ ////////////////////////////////////////////////////////////////////////////
+ ir::Instruction* Instruction(const pb::Instruction& inst_in) {
+ ir::Instruction* inst_out = nullptr;
+ switch (inst_in.kind_case()) {
+ case pb::Instruction::KindCase::kAccess:
+ inst_out = CreateInstructionAccess(inst_in.access());
+ break;
+ case pb::Instruction::KindCase::kBinary:
+ inst_out = CreateInstructionBinary(inst_in.binary());
+ break;
+ case pb::Instruction::KindCase::kConstruct:
+ inst_out = CreateInstructionConstruct(inst_in.construct());
+ break;
+ case pb::Instruction::KindCase::kDiscard:
+ inst_out = CreateInstructionDiscard(inst_in.discard());
+ break;
+ case pb::Instruction::KindCase::kLet:
+ inst_out = CreateInstructionLet(inst_in.let());
+ break;
+ case pb::Instruction::KindCase::kLoad:
+ inst_out = CreateInstructionLoad(inst_in.load());
+ break;
+ case pb::Instruction::KindCase::kLoadVectorElement:
+ inst_out = CreateInstructionLoadVectorElement(inst_in.load_vector_element());
+ break;
+ case pb::Instruction::KindCase::kReturn:
+ inst_out = CreateInstructionReturn(inst_in.return_());
+ break;
+ case pb::Instruction::KindCase::kStore:
+ inst_out = CreateInstructionStore(inst_in.store());
+ break;
+ case pb::Instruction::KindCase::kStoreVectorElement:
+ inst_out = CreateInstructionStoreVectorElement(inst_in.store_vector_element());
+ break;
+ case pb::Instruction::KindCase::kSwizzle:
+ inst_out = CreateInstructionSwizzle(inst_in.swizzle());
+ break;
+ case pb::Instruction::KindCase::kUnary:
+ inst_out = CreateInstructionUnary(inst_in.unary());
+ break;
+ case pb::Instruction::KindCase::kUserCall:
+ inst_out = CreateInstructionUserCall(inst_in.user_call());
+ break;
+ case pb::Instruction::KindCase::kVar:
+ inst_out = CreateInstructionVar(inst_in.var());
+ break;
+ default:
+ TINT_UNIMPLEMENTED() << inst_in.kind_case();
+ break;
+ }
+ TINT_ASSERT_OR_RETURN_VALUE(inst_out, nullptr);
+
+ Vector<ir::Value*, 4> operands;
+ for (auto id : inst_in.operands()) {
+ operands.Push(Value(id));
+ }
+ inst_out->SetOperands(std::move(operands));
+
+ Vector<ir::InstructionResult*, 4> results;
+ for (auto id : inst_in.results()) {
+ results.Push(ValueAs<ir::InstructionResult>(id));
+ }
+ inst_out->SetResults(std::move(results));
+
+ return inst_out;
+ }
+
+ ir::Access* CreateInstructionAccess(const pb::InstructionAccess&) {
+ return mod_out_.instructions.Create<ir::Access>();
+ }
+
+ ir::Binary* CreateInstructionBinary(const pb::InstructionBinary& binary_in) {
+ auto* binary_out = mod_out_.instructions.Create<ir::Binary>();
+ binary_out->SetOp(BinaryOp(binary_in.op()));
+ return binary_out;
+ }
+
+ ir::Construct* CreateInstructionConstruct(const pb::InstructionConstruct&) {
+ return mod_out_.instructions.Create<ir::Construct>();
+ }
+
+ ir::Discard* CreateInstructionDiscard(const pb::InstructionDiscard&) {
+ return mod_out_.instructions.Create<ir::Discard>();
+ }
+
+ ir::Let* CreateInstructionLet(const pb::InstructionLet&) {
+ return mod_out_.instructions.Create<ir::Let>();
+ }
+
+ ir::Load* CreateInstructionLoad(const pb::InstructionLoad&) {
+ return mod_out_.instructions.Create<ir::Load>();
+ }
+
+ ir::LoadVectorElement* CreateInstructionLoadVectorElement(
+ const pb::InstructionLoadVectorElement&) {
+ return mod_out_.instructions.Create<ir::LoadVectorElement>();
+ }
+
+ ir::Return* CreateInstructionReturn(const pb::InstructionReturn&) {
+ return mod_out_.instructions.Create<ir::Return>();
+ }
+
+ ir::Store* CreateInstructionStore(const pb::InstructionStore&) {
+ return mod_out_.instructions.Create<ir::Store>();
+ }
+
+ ir::StoreVectorElement* CreateInstructionStoreVectorElement(
+ const pb::InstructionStoreVectorElement&) {
+ return mod_out_.instructions.Create<ir::StoreVectorElement>();
+ }
+
+ ir::Swizzle* CreateInstructionSwizzle(const pb::InstructionSwizzle& swizzle_in) {
+ auto* swizzle_out = mod_out_.instructions.Create<ir::Swizzle>();
+ Vector<uint32_t, 4> indices;
+ for (auto idx : swizzle_in.indices()) {
+ indices.Push(idx);
+ }
+ swizzle_out->SetIndices(indices);
+ return swizzle_out;
+ }
+
+ ir::Unary* CreateInstructionUnary(const pb::InstructionUnary& unary_in) {
+ auto* unary_out = mod_out_.instructions.Create<ir::Unary>();
+ unary_out->SetOp(UnaryOp(unary_in.op()));
+ return unary_out;
+ }
+
+ ir::UserCall* CreateInstructionUserCall(const pb::InstructionUserCall&) {
+ return mod_out_.instructions.Create<ir::UserCall>();
+ }
+
+ ir::Var* CreateInstructionVar(const pb::InstructionVar& var_in) {
+ auto* var_out = mod_out_.instructions.Create<ir::Var>();
+ if (var_in.has_binding_point()) {
+ auto& bp_in = var_in.binding_point();
+ var_out->SetBindingPoint(bp_in.group(), bp_in.binding());
+ }
+ return var_out;
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Types
+ ////////////////////////////////////////////////////////////////////////////
+ const type::Type* CreateType(const pb::Type type_in) {
+ switch (type_in.kind_case()) {
+ case pb::Type::KindCase::kBasic:
+ return CreateTypeBasic(type_in.basic());
+ case pb::Type::KindCase::kVector:
+ return CreateTypeVector(type_in.vector());
+ case pb::Type::KindCase::kMatrix:
+ return CreateTypeMatrix(type_in.matrix());
+ case pb::Type::KindCase::kPointer:
+ return CreateTypePointer(type_in.pointer());
+ case pb::Type::KindCase::kArray:
+ return CreateTypeArray(type_in.array());
+ case pb::Type::KindCase::kAtomic:
+ TINT_UNIMPLEMENTED() << type_in.kind_case();
+ return nullptr;
+
+ case pb::Type::KindCase::KIND_NOT_SET:
+ break;
+ }
+ TINT_ICE() << "invalid TypeDecl.kind";
+ return nullptr;
+ }
+
+ const type::Type* CreateTypeBasic(pb::BasicType basic_in) {
+ switch (basic_in) {
+ 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>();
+ case pb::BasicType::f16:
+ return mod_out_.Types().Get<f16>();
+ default:
+ TINT_ICE() << "invalid BasicType: " << basic_in;
+ return nullptr;
+ }
+ }
+
+ const type::Vector* CreateTypeVector(const pb::VectorType& vector_in) {
+ auto* el_ty = Type(vector_in.element_type());
+ return mod_out_.Types().vec(el_ty, vector_in.width());
+ }
+
+ const type::Matrix* CreateTypeMatrix(const pb::MatrixType& matrix_in) {
+ auto* el_ty = Type(matrix_in.element_type());
+ auto* column_ty = mod_out_.Types().vec(el_ty, matrix_in.num_rows());
+ return mod_out_.Types().mat(column_ty, matrix_in.num_columns());
+ }
+
+ const type::Pointer* CreateTypePointer(const pb::PointerType& pointer_in) {
+ auto address_space = AddressSpace(pointer_in.address_space());
+ auto* store_ty = Type(pointer_in.store_type());
+ auto access = Access(pointer_in.access());
+ return mod_out_.Types().ptr(address_space, store_ty, access);
+ }
+
+ const type::Array* CreateTypeArray(const pb::ArrayType& array_in) {
+ auto* element = Type(array_in.element());
+ uint32_t stride = static_cast<uint32_t>(array_in.stride());
+ uint32_t count = static_cast<uint32_t>(array_in.count());
+ return mod_out_.Types().array(element, count, stride);
+ }
+
+ const type::Type* Type(size_t id) { return id > 0 ? types_[id - 1] : nullptr; }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Values
+ ////////////////////////////////////////////////////////////////////////////
+ ir::Value* CreateValue(const pb::Value& value_in) {
+ ir::Value* value_out = nullptr;
+ switch (value_in.kind_case()) {
+ case pb::Value::KindCase::kFunction: {
+ value_out = Function(value_in.function());
+ break;
+ }
+ case pb::Value::KindCase::kInstructionResult: {
+ auto& res_in = value_in.instruction_result();
+ auto* type = Type(res_in.type());
+ value_out = b.InstructionResult(type);
+ if (res_in.has_name()) {
+ mod_out_.SetName(value_out, res_in.name());
+ }
+ break;
+ }
+ case pb::Value::KindCase::kFunctionParameter: {
+ auto& param_in = value_in.function_parameter();
+ auto* type = Type(param_in.type());
+ value_out = b.FunctionParam(type);
+ if (param_in.has_name()) {
+ mod_out_.SetName(value_out, param_in.name());
+ }
+ break;
+ }
+ case pb::Value::KindCase::kConstant: {
+ value_out = b.Constant(ConstantValue(value_in.constant()));
+ break;
+ }
+ default:
+ TINT_ICE() << "invalid TypeDecl.kind: " << value_in.kind_case();
+ return nullptr;
+ }
+ return value_out;
+ }
+
+ ir::Value* Value(uint32_t id) { return id > 0 ? values_[id - 1] : nullptr; }
+
+ 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 ? value->TypeInfo().name : "<null>")
+ << " expected " << TypeInfo::Of<T>().name;
+ return nullptr;
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // ConstantValues
+ ////////////////////////////////////////////////////////////////////////////
+ const core::constant::Value* CreateConstantValue(const pb::ConstantValue& value_in) {
+ switch (value_in.kind_case()) {
+ case pb::ConstantValue::KindCase::kScalar:
+ return CreateConstantScalar(value_in.scalar());
+ case pb::ConstantValue::KindCase::kComposite:
+ return CreateConstantComposite(value_in.composite());
+ case pb::ConstantValue::KindCase::kSplat:
+ return CreateConstantSplat(value_in.splat());
+ default:
+ TINT_ICE() << "invalid ConstantValue.kind: " << value_in.kind_case();
+ return nullptr;
+ }
+ }
+
+ const core::constant::Value* CreateConstantScalar(const pb::ConstantValueScalar& value_in) {
+ switch (value_in.kind_case()) {
+ case pb::ConstantValueScalar::KindCase::kBool:
+ return b.ConstantValue(value_in.bool_());
+ case pb::ConstantValueScalar::KindCase::kI32:
+ return b.ConstantValue(i32(value_in.i32()));
+ case pb::ConstantValueScalar::KindCase::kU32:
+ return b.ConstantValue(u32(value_in.u32()));
+ case pb::ConstantValueScalar::KindCase::kF32:
+ return b.ConstantValue(f32(value_in.f32()));
+ case pb::ConstantValueScalar::KindCase::kF16:
+ return b.ConstantValue(f16(value_in.f16()));
+ default:
+ TINT_ICE() << "invalid ConstantValueScalar.kind: " << value_in.kind_case();
+ return nullptr;
+ }
+ }
+
+ const core::constant::Value* CreateConstantComposite(
+ const pb::ConstantValueComposite& composite_in) {
+ auto* type = Type(composite_in.type());
+ Vector<const core::constant::Value*, 8> elements_out;
+ for (auto element_id : composite_in.elements()) {
+ elements_out.Push(ConstantValue(element_id));
+ }
+ return mod_out_.constant_values.Composite(type, std::move(elements_out));
+ }
+
+ const core::constant::Value* CreateConstantSplat(const pb::ConstantValueSplat& splat_in) {
+ auto* type = Type(splat_in.type());
+ auto* elem = ConstantValue(splat_in.elements());
+ return mod_out_.constant_values.Splat(type, elem, splat_in.count());
+ }
+
+ const core::constant::Value* ConstantValue(uint32_t id) {
+ return id > 0 ? constant_values_[id - 1] : nullptr;
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Enums
+ ////////////////////////////////////////////////////////////////////////////
+ core::AddressSpace AddressSpace(pb::AddressSpace in) {
+ switch (in) {
+ case pb::AddressSpace::function:
+ return core::AddressSpace::kFunction;
+ case pb::AddressSpace::handle:
+ return core::AddressSpace::kHandle;
+ case pb::AddressSpace::pixel_local:
+ return core::AddressSpace::kPixelLocal;
+ case pb::AddressSpace::private_:
+ return core::AddressSpace::kPrivate;
+ case pb::AddressSpace::push_constant:
+ return core::AddressSpace::kPushConstant;
+ case pb::AddressSpace::storage:
+ return core::AddressSpace::kStorage;
+ case pb::AddressSpace::uniform:
+ return core::AddressSpace::kUniform;
+ case pb::AddressSpace::workgroup:
+ return core::AddressSpace::kWorkgroup;
+ default:
+ TINT_ICE() << "invalid AddressSpace: " << in;
+ return core::AddressSpace::kUndefined;
+ }
+ }
+
+ core::Access Access(pb::AccessControl in) {
+ switch (in) {
+ case pb::AccessControl::read:
+ return core::Access::kRead;
+ case pb::AccessControl::write:
+ return core::Access::kWrite;
+ case pb::AccessControl::read_write:
+ return core::Access::kReadWrite;
+ default:
+ TINT_ICE() << "invalid Access: " << in;
+ return core::Access::kUndefined;
+ }
+ }
+
+ core::ir::UnaryOp UnaryOp(pb::UnaryOp in) {
+ switch (in) {
+ case pb::UnaryOp::complement:
+ return core::ir::UnaryOp::kComplement;
+ case pb::UnaryOp::negation:
+ return core::ir::UnaryOp::kNegation;
+
+ default:
+ TINT_ICE() << "invalid UnaryOp: " << in;
+ return core::ir::UnaryOp::kComplement;
+ }
+ }
+
+ core::ir::BinaryOp BinaryOp(pb::BinaryOp in) {
+ switch (in) {
+ case pb::BinaryOp::add_:
+ return core::ir::BinaryOp::kAdd;
+ case pb::BinaryOp::subtract:
+ return core::ir::BinaryOp::kSubtract;
+ case pb::BinaryOp::multiply:
+ return core::ir::BinaryOp::kMultiply;
+ case pb::BinaryOp::divide:
+ return core::ir::BinaryOp::kDivide;
+ case pb::BinaryOp::modulo:
+ return core::ir::BinaryOp::kModulo;
+ case pb::BinaryOp::and_:
+ return core::ir::BinaryOp::kAnd;
+ case pb::BinaryOp::or_:
+ return core::ir::BinaryOp::kOr;
+ case pb::BinaryOp::xor_:
+ return core::ir::BinaryOp::kXor;
+ case pb::BinaryOp::equal:
+ return core::ir::BinaryOp::kEqual;
+ case pb::BinaryOp::not_equal:
+ return core::ir::BinaryOp::kNotEqual;
+ case pb::BinaryOp::less_than:
+ return core::ir::BinaryOp::kLessThan;
+ case pb::BinaryOp::greater_than:
+ return core::ir::BinaryOp::kGreaterThan;
+ case pb::BinaryOp::less_than_equal:
+ return core::ir::BinaryOp::kLessThanEqual;
+ case pb::BinaryOp::greater_than_equal:
+ return core::ir::BinaryOp::kGreaterThanEqual;
+ case pb::BinaryOp::shift_left:
+ return core::ir::BinaryOp::kShiftLeft;
+ case pb::BinaryOp::shift_right:
+ return core::ir::BinaryOp::kShiftRight;
+
+ default:
+ TINT_ICE() << "invalid BinaryOp: " << in;
+ return core::ir::BinaryOp::kAdd;
+ }
+ }
+};
+
+} // 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..7846a4a
--- /dev/null
+++ b/src/tint/lang/core/ir/binary/encode.cc
@@ -0,0 +1,482 @@
+// 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 <utility>
+
+#include "src/tint/lang/core/constant/composite.h"
+#include "src/tint/lang/core/constant/scalar.h"
+#include "src/tint/lang/core/constant/splat.h"
+#include "src/tint/lang/core/ir/access.h"
+#include "src/tint/lang/core/ir/binary.h"
+#include "src/tint/lang/core/ir/construct.h"
+#include "src/tint/lang/core/ir/discard.h"
+#include "src/tint/lang/core/ir/function_param.h"
+#include "src/tint/lang/core/ir/let.h"
+#include "src/tint/lang/core/ir/load.h"
+#include "src/tint/lang/core/ir/load_vector_element.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/core/ir/return.h"
+#include "src/tint/lang/core/ir/store.h"
+#include "src/tint/lang/core/ir/store_vector_element.h"
+#include "src/tint/lang/core/ir/swizzle.h"
+#include "src/tint/lang/core/ir/unary.h"
+#include "src/tint/lang/core/ir/user_call.h"
+#include "src/tint/lang/core/ir/var.h"
+#include "src/tint/lang/core/type/array.h"
+#include "src/tint/lang/core/type/bool.h"
+#include "src/tint/lang/core/type/f16.h"
+#include "src/tint/lang/core/type/f32.h"
+#include "src/tint/lang/core/type/i32.h"
+#include "src/tint/lang/core/type/matrix.h"
+#include "src/tint/lang/core/type/pointer.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_{};
+ Hashmap<const core::constant::Value*, uint32_t, 32> constant_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++) {
+ PopulateFunction(fns_out[i], mod_in_.functions[i]);
+ }
+ mod_out_.set_root_block(Block(mod_in_.root_block));
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Functions
+ ////////////////////////////////////////////////////////////////////////////
+ void PopulateFunction(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(Value(param_in));
+ }
+ fn_out->set_block(Block(fn_in->Block()));
+ }
+
+ uint32_t Function(const ir::Function* fn_in) { return fn_in ? *functions_.Get(fn_in) : 0; }
+
+ 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) {
+ Switch(
+ inst_in, //
+ [&](const ir::Access* i) { InstructionAccess(*inst_out.mutable_access(), i); },
+ [&](const ir::Binary* i) { InstructionBinary(*inst_out.mutable_binary(), i); },
+ [&](const ir::Construct* i) { InstructionConstruct(*inst_out.mutable_construct(), i); },
+ [&](const ir::Discard* i) { InstructionDiscard(*inst_out.mutable_discard(), i); },
+ [&](const ir::Let* i) { InstructionLet(*inst_out.mutable_let(), i); },
+ [&](const ir::Load* i) { InstructionLoad(*inst_out.mutable_load(), i); },
+ [&](const ir::LoadVectorElement* i) {
+ InstructionLoadVectorElement(*inst_out.mutable_load_vector_element(), i);
+ },
+ [&](const ir::Return* i) { InstructionReturn(*inst_out.mutable_return_(), i); },
+ [&](const ir::Store* i) { InstructionStore(*inst_out.mutable_store(), i); },
+ [&](const ir::StoreVectorElement* i) {
+ InstructionStoreVectorElement(*inst_out.mutable_store_vector_element(), i);
+ },
+ [&](const ir::Swizzle* i) { InstructionSwizzle(*inst_out.mutable_swizzle(), i); },
+ [&](const ir::Unary* i) { InstructionUnary(*inst_out.mutable_unary(), i); },
+ [&](const ir::UserCall* i) { InstructionUserCall(*inst_out.mutable_user_call(), i); },
+ [&](const ir::Var* i) { InstructionVar(*inst_out.mutable_var(), i); },
+ TINT_ICE_ON_NO_MATCH);
+ for (auto* operand : inst_in->Operands()) {
+ inst_out.add_operands(Value(operand));
+ }
+ for (auto* result : inst_in->Results()) {
+ inst_out.add_results(Value(result));
+ }
+ }
+
+ void InstructionAccess(pb::InstructionAccess&, const ir::Access*) {}
+
+ void InstructionBinary(pb::InstructionBinary& binary_out, const ir::Binary* binary_in) {
+ binary_out.set_op(BinaryOp(binary_in->Op()));
+ }
+
+ void InstructionConstruct(pb::InstructionConstruct&, const ir::Construct*) {}
+
+ void InstructionDiscard(pb::InstructionDiscard&, const ir::Discard*) {}
+
+ void InstructionLet(pb::InstructionLet&, const ir::Let*) {}
+
+ void InstructionLoad(pb::InstructionLoad&, const ir::Load*) {}
+
+ void InstructionLoadVectorElement(pb::InstructionLoadVectorElement&,
+ const ir::LoadVectorElement*) {}
+
+ void InstructionReturn(pb::InstructionReturn&, const ir::Return*) {}
+
+ void InstructionStore(pb::InstructionStore&, const ir::Store*) {}
+
+ void InstructionStoreVectorElement(pb::InstructionStoreVectorElement&,
+ const ir::StoreVectorElement*) {}
+
+ void InstructionSwizzle(pb::InstructionSwizzle& swizzle_out, const ir::Swizzle* swizzle_in) {
+ for (auto idx : swizzle_in->Indices()) {
+ swizzle_out.add_indices(idx);
+ }
+ }
+
+ void InstructionUnary(pb::InstructionUnary& unary_out, const ir::Unary* unary_in) {
+ unary_out.set_op(UnaryOp(unary_in->Op()));
+ }
+
+ void InstructionUserCall(pb::InstructionUserCall&, const ir::UserCall*) {}
+
+ void InstructionVar(pb::InstructionVar& var_out, const ir::Var* var_in) {
+ if (auto bp_in = var_in->BindingPoint()) {
+ auto& bp_out = *var_out.mutable_binding_point();
+ bp_out.set_group(bp_in->group);
+ bp_out.set_binding(bp_in->binding);
+ }
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Types
+ ////////////////////////////////////////////////////////////////////////////
+ uint32_t Type(const core::type::Type* type_in) {
+ if (type_in == nullptr) {
+ return 0;
+ }
+ return types_.GetOrCreate(type_in, [&]() -> uint32_t {
+ pb::Type type_out;
+ Switch(
+ type_in, //
+ [&](const core::type::Void*) { type_out.set_basic(pb::BasicType::void_); },
+ [&](const core::type::Bool*) { type_out.set_basic(pb::BasicType::bool_); },
+ [&](const core::type::I32*) { type_out.set_basic(pb::BasicType::i32); },
+ [&](const core::type::U32*) { type_out.set_basic(pb::BasicType::u32); },
+ [&](const core::type::F32*) { type_out.set_basic(pb::BasicType::f32); },
+ [&](const core::type::F16*) { type_out.set_basic(pb::BasicType::f16); },
+ [&](const core::type::Vector* v) { VectorType(*type_out.mutable_vector(), v); },
+ [&](const core::type::Matrix* m) { MatrixType(*type_out.mutable_matrix(), m); },
+ [&](const core::type::Pointer* m) { PointerType(*type_out.mutable_pointer(), m); },
+ [&](const core::type::Array* m) { ArrayType(*type_out.mutable_array(), m); },
+ TINT_ICE_ON_NO_MATCH);
+
+ mod_out_.mutable_types()->Add(std::move(type_out));
+ return static_cast<uint32_t>(mod_out_.types().size());
+ });
+ }
+
+ void VectorType(pb::VectorType& vector_out, const core::type::Vector* vector_in) {
+ vector_out.set_width(vector_in->Width());
+ vector_out.set_element_type(Type(vector_in->type()));
+ }
+
+ void MatrixType(pb::MatrixType& matrix_out, const core::type::Matrix* matrix_in) {
+ matrix_out.set_num_columns(matrix_in->columns());
+ matrix_out.set_num_rows(matrix_in->rows());
+ matrix_out.set_element_type(Type(matrix_in->type()));
+ }
+
+ void PointerType(pb::PointerType& pointer_out, const core::type::Pointer* pointer_in) {
+ pointer_out.set_address_space(AddressSpace(pointer_in->AddressSpace()));
+ pointer_out.set_store_type(Type(pointer_in->StoreType()));
+ pointer_out.set_access(Access(pointer_in->Access()));
+ }
+
+ void ArrayType(pb::ArrayType& array_out, const core::type::Array* array_in) {
+ array_out.set_element(Type(array_in->ElemType()));
+ array_out.set_stride(array_in->Stride());
+ Switch(
+ array_in->Count(), //
+ [&](const core::type::ConstantArrayCount* c) { array_out.set_count(c->value); },
+ TINT_ICE_ON_NO_MATCH);
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Values
+ ////////////////////////////////////////////////////////////////////////////
+ uint32_t Value(const ir::Value* value_in) {
+ if (!value_in) {
+ return 0;
+ }
+ return values_.GetOrCreate(value_in, [&] {
+ auto& value_out = *mod_out_.add_values();
+ Switch(
+ value_in,
+ [&](const ir::InstructionResult* v) {
+ InstructionResult(*value_out.mutable_instruction_result(), v);
+ },
+ [&](const ir::FunctionParam* v) {
+ FunctionParameter(*value_out.mutable_function_parameter(), v);
+ },
+ [&](const ir::Function* v) { value_out.set_function(Function(v)); },
+ [&](const ir::Constant* v) { value_out.set_constant(ConstantValue(v->Value())); },
+ TINT_ICE_ON_NO_MATCH);
+
+ return static_cast<uint32_t>(mod_out_.values().size());
+ });
+ }
+
+ void InstructionResult(pb::InstructionResult& res_out, const ir::InstructionResult* res_in) {
+ res_out.set_type(Type(res_in->Type()));
+ if (auto name = mod_in_.NameOf(res_in); name.IsValid()) {
+ res_out.set_name(name.Name());
+ }
+ }
+
+ void FunctionParameter(pb::FunctionParameter& param_out, const ir::FunctionParam* param_in) {
+ param_out.set_type(Type(param_in->Type()));
+ if (auto name = mod_in_.NameOf(param_in); name.IsValid()) {
+ param_out.set_name(name.Name());
+ }
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // ConstantValues
+ ////////////////////////////////////////////////////////////////////////////
+ uint32_t ConstantValue(const core::constant::Value* constant_in) {
+ if (!constant_in) {
+ return 0;
+ }
+ return constant_values_.GetOrCreate(constant_in, [&] {
+ pb::ConstantValue constant_out;
+ Switch(
+ constant_in, //
+ [&](const core::constant::Scalar<bool>* b) {
+ constant_out.mutable_scalar()->set_bool_(b->value);
+ },
+ [&](const core::constant::Scalar<core::i32>* i32) {
+ constant_out.mutable_scalar()->set_i32(i32->value);
+ },
+ [&](const core::constant::Scalar<core::u32>* u32) {
+ constant_out.mutable_scalar()->set_u32(u32->value);
+ },
+ [&](const core::constant::Scalar<core::f32>* f32) {
+ constant_out.mutable_scalar()->set_f32(f32->value);
+ },
+ [&](const core::constant::Scalar<core::f16>* f16) {
+ constant_out.mutable_scalar()->set_f16(f16->value);
+ },
+ [&](const core::constant::Composite* composite) {
+ ConstantValueComposite(*constant_out.mutable_composite(), composite);
+ },
+ [&](const core::constant::Splat* splat) {
+ ConstantValueSplat(*constant_out.mutable_splat(), splat);
+ },
+ TINT_ICE_ON_NO_MATCH);
+
+ mod_out_.mutable_constant_values()->Add(std::move(constant_out));
+ return static_cast<uint32_t>(mod_out_.constant_values().size());
+ });
+ }
+
+ void ConstantValueComposite(pb::ConstantValueComposite& composite_out,
+ const core::constant::Composite* composite_in) {
+ composite_out.set_type(Type(composite_in->type));
+ for (auto* el : composite_in->elements) {
+ composite_out.add_elements(ConstantValue(el));
+ }
+ }
+
+ void ConstantValueSplat(pb::ConstantValueSplat& splat_out,
+ const core::constant::Splat* splat_in) {
+ splat_out.set_type(Type(splat_in->type));
+ splat_out.set_elements(ConstantValue(splat_in->el));
+ splat_out.set_count(static_cast<uint32_t>(splat_in->count));
+ }
+
+ ////////////////////////////////////////////////////////////////////////////
+ // Enums
+ ////////////////////////////////////////////////////////////////////////////
+ pb::AddressSpace AddressSpace(core::AddressSpace in) {
+ switch (in) {
+ case core::AddressSpace::kFunction:
+ return pb::AddressSpace::function;
+ case core::AddressSpace::kHandle:
+ return pb::AddressSpace::handle;
+ case core::AddressSpace::kPixelLocal:
+ return pb::AddressSpace::pixel_local;
+ case core::AddressSpace::kPrivate:
+ return pb::AddressSpace::private_;
+ case core::AddressSpace::kPushConstant:
+ return pb::AddressSpace::push_constant;
+ case core::AddressSpace::kStorage:
+ return pb::AddressSpace::storage;
+ case core::AddressSpace::kUniform:
+ return pb::AddressSpace::uniform;
+ case core::AddressSpace::kWorkgroup:
+ return pb::AddressSpace::workgroup;
+ default:
+ TINT_ICE() << "invalid AddressSpace: " << in;
+ return pb::AddressSpace::function;
+ }
+ }
+
+ pb::AccessControl Access(core::Access in) {
+ switch (in) {
+ case core::Access::kRead:
+ return pb::AccessControl::read;
+ case core::Access::kWrite:
+ return pb::AccessControl::write;
+ case core::Access::kReadWrite:
+ return pb::AccessControl::read_write;
+ default:
+ TINT_ICE() << "invalid Access: " << in;
+ return pb::AccessControl::read;
+ }
+ }
+
+ pb::UnaryOp UnaryOp(core::ir::UnaryOp in) {
+ switch (in) {
+ case core::ir::UnaryOp::kComplement:
+ return pb::UnaryOp::complement;
+ case core::ir::UnaryOp::kNegation:
+ return pb::UnaryOp::negation;
+ }
+ TINT_ICE() << "invalid UnaryOp: " << in;
+ return pb::UnaryOp::complement;
+ }
+
+ pb::BinaryOp BinaryOp(core::ir::BinaryOp in) {
+ switch (in) {
+ case core::ir::BinaryOp::kAdd:
+ return pb::BinaryOp::add_;
+ case core::ir::BinaryOp::kSubtract:
+ return pb::BinaryOp::subtract;
+ case core::ir::BinaryOp::kMultiply:
+ return pb::BinaryOp::multiply;
+ case core::ir::BinaryOp::kDivide:
+ return pb::BinaryOp::divide;
+ case core::ir::BinaryOp::kModulo:
+ return pb::BinaryOp::modulo;
+ case core::ir::BinaryOp::kAnd:
+ return pb::BinaryOp::and_;
+ case core::ir::BinaryOp::kOr:
+ return pb::BinaryOp::or_;
+ case core::ir::BinaryOp::kXor:
+ return pb::BinaryOp::xor_;
+ case core::ir::BinaryOp::kEqual:
+ return pb::BinaryOp::equal;
+ case core::ir::BinaryOp::kNotEqual:
+ return pb::BinaryOp::not_equal;
+ case core::ir::BinaryOp::kLessThan:
+ return pb::BinaryOp::less_than;
+ case core::ir::BinaryOp::kGreaterThan:
+ return pb::BinaryOp::greater_than;
+ case core::ir::BinaryOp::kLessThanEqual:
+ return pb::BinaryOp::less_than_equal;
+ case core::ir::BinaryOp::kGreaterThanEqual:
+ return pb::BinaryOp::greater_than_equal;
+ case core::ir::BinaryOp::kShiftLeft:
+ return pb::BinaryOp::shift_left;
+ case core::ir::BinaryOp::kShiftRight:
+ return pb::BinaryOp::shift_right;
+ }
+
+ TINT_ICE() << "invalid BinaryOp: " << in;
+ return pb::BinaryOp::add_;
+ }
+};
+
+} // 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..3db27a5
--- /dev/null
+++ b/src/tint/lang/core/ir/binary/ir.proto
@@ -0,0 +1,288 @@
+// 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.
+
+syntax = "proto3";
+
+package tint.core.ir.binary.pb;
+
+message Module {
+ repeated Type types = 1;
+ repeated Value values = 2;
+ repeated ConstantValue constant_values = 3;
+ repeated Function functions = 4;
+ repeated Block blocks = 5;
+ uint32 root_block = 6; // Module.blocks
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Types
+////////////////////////////////////////////////////////////////////////////////
+message Type {
+ oneof kind {
+ BasicType basic = 1;
+ VectorType vector = 2;
+ MatrixType matrix = 3;
+ ArrayType array = 4;
+ PointerType pointer = 5;
+ uint32 atomic = 6; // 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_rows = 2;
+ uint32 element_type = 3; // Module.types
+}
+
+message ArrayType {
+ uint32 element = 1; // Module.types
+ uint32 stride = 2;
+ uint32 count = 3;
+}
+
+message PointerType {
+ AddressSpace address_space = 1;
+ uint32 store_type = 2; // Module.types
+ AccessControl access = 3;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Values
+////////////////////////////////////////////////////////////////////////////////
+message Value {
+ oneof kind {
+ uint32 function = 1; // Module.functions
+ InstructionResult instruction_result = 2;
+ FunctionParameter function_parameter = 3;
+ uint32 constant = 4; // Module.constant_values
+ }
+}
+
+message InstructionResult {
+ uint32 type = 1; // Module.types
+ optional string name = 2;
+}
+
+message FunctionParameter {
+ uint32 type = 1; // Module.types
+ optional string name = 2;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// ConstantValues
+////////////////////////////////////////////////////////////////////////////////
+message ConstantValue {
+ oneof kind {
+ ConstantValueScalar scalar = 1;
+ ConstantValueComposite composite = 2;
+ ConstantValueSplat splat = 3;
+ }
+}
+
+message ConstantValueScalar {
+ oneof kind {
+ bool bool = 1;
+ int32 i32 = 2;
+ uint32 u32 = 3;
+ float f32 = 4;
+ float f16 = 5;
+ }
+}
+
+message ConstantValueComposite {
+ uint32 type = 1; // Module.types
+ repeated uint32 elements = 2; // Module.constant_values
+}
+
+message ConstantValueSplat {
+ uint32 type = 1; // Module.types
+ uint32 elements = 2; // Module.constant_values
+ uint32 count = 3; // splat count
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Functions
+////////////////////////////////////////////////////////////////////////////////
+message Function {
+ uint32 return_type = 1; // Module.types
+ uint32 block = 2; // Module.blocks
+ optional string name = 3;
+ optional PipelineStage pipeline_stage = 4;
+ optional WorkgroupSize workgroup_size = 5;
+ repeated uint32 parameters = 6; // Module.values
+}
+
+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 {
+ repeated uint32 operands = 1; // Module.values
+ repeated uint32 results = 2; // Module.values
+ oneof kind {
+ InstructionReturn return = 3;
+ InstructionUnary unary = 4;
+ InstructionBinary binary = 5;
+ InstructionBuiltin builtin = 6;
+ InstructionConstructor constructor = 7;
+ InstructionDiscard discard = 8;
+ InstructionLet let = 9;
+ InstructionVar var = 10;
+ InstructionConstruct construct = 11;
+ InstructionAccess access = 12;
+ InstructionUserCall user_call = 13;
+ InstructionLoad load = 14;
+ InstructionStore store = 15;
+ InstructionLoadVectorElement load_vector_element = 16;
+ InstructionStoreVectorElement store_vector_element = 17;
+ InstructionSwizzle swizzle = 18;
+ }
+}
+
+message InstructionReturn {}
+
+message InstructionUnary {
+ UnaryOp op = 1;
+}
+
+message InstructionBinary {
+ BinaryOp op = 1;
+}
+
+message InstructionBuiltin {}
+
+message InstructionConstructor {}
+
+message InstructionDiscard {}
+
+message InstructionLet {}
+
+message InstructionVar {
+ optional BindingPoint binding_point = 1;
+}
+
+message InstructionConstruct {}
+
+message InstructionAccess {}
+
+message InstructionUserCall {}
+
+message InstructionLoad {}
+
+message InstructionStore {}
+
+message InstructionLoadVectorElement {}
+
+message InstructionStoreVectorElement {}
+
+message InstructionSwizzle {
+ repeated uint32 indices = 1;
+}
+
+message BindingPoint {
+ uint32 group = 1;
+ uint32 binding = 2;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Enums
+////////////////////////////////////////////////////////////////////////////////
+enum AddressSpace {
+ function = 0;
+ handle = 1;
+ pixel_local = 2;
+ private = 3;
+ push_constant = 4;
+ storage = 5;
+ uniform = 6;
+ workgroup = 7;
+}
+
+enum AccessControl {
+ read = 0;
+ write = 1;
+ read_write = 2;
+}
+
+enum UnaryOp {
+ complement = 0;
+ negation = 1;
+}
+
+enum BinaryOp {
+ add_ = 0;
+ subtract = 1;
+ multiply = 2;
+ divide = 3;
+ modulo = 4;
+ and = 5;
+ or_ = 6;
+ xor_ = 7;
+ equal = 8;
+ not_equal = 9;
+ less_than = 10;
+ greater_than = 11;
+ less_than_equal = 12;
+ greater_than_equal = 13;
+ shift_left = 14;
+ shift_right = 15;
+}
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..57ade4e
--- /dev/null
+++ b/src/tint/lang/core/ir/binary/roundtrip_test.cc
@@ -0,0 +1,318 @@
+// 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
+using namespace tint::core::fluent_types; // 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();
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Root block
+////////////////////////////////////////////////////////////////////////////////
+TEST_F(IRBinaryRoundtripTest, RootBlock_Var_private_i32_Unnamed) {
+ b.Append(b.ir.root_block, [&] { b.Var<private_, i32>(); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, RootBlock_Var_workgroup_f32_Named) {
+ b.Append(b.ir.root_block, [&] { b.Var<workgroup, f32>("WG"); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, RootBlock_Var_storage_binding) {
+ b.Append(b.ir.root_block, [&] {
+ auto* v = b.Var<storage, f32>();
+ v->SetBindingPoint(10, 20);
+ });
+ 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();
+}
+
+TEST_F(IRBinaryRoundtripTest, Return_bool) {
+ auto* fn = b.Function("Function", ty.bool_());
+ b.Append(fn->Block(), [&] { b.Return(fn, true); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Return_i32) {
+ auto* fn = b.Function("Function", ty.i32());
+ b.Append(fn->Block(), [&] { b.Return(fn, 42_i); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Return_u32) {
+ auto* fn = b.Function("Function", ty.u32());
+ b.Append(fn->Block(), [&] { b.Return(fn, 42_u); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Return_f32) {
+ auto* fn = b.Function("Function", ty.f32());
+ b.Append(fn->Block(), [&] { b.Return(fn, 42_f); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Return_f16) {
+ auto* fn = b.Function("Function", ty.f16());
+ b.Append(fn->Block(), [&] { b.Return(fn, 42_h); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Return_vec3f_Composite) {
+ auto* fn = b.Function("Function", ty.vec3<f32>());
+ b.Append(fn->Block(), [&] { b.Return(fn, b.Composite<vec3<f32>>(1_f, 2_f, 3_f)); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Return_vec3f_Splat) {
+ auto* fn = b.Function("Function", ty.vec3<f32>());
+ b.Append(fn->Block(), [&] { b.Return(fn, b.Splat<vec3<f32>>(1_f, 3)); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Return_mat2x3f_Composite) {
+ auto* fn = b.Function("Function", ty.mat2x3<f32>());
+ b.Append(fn->Block(),
+ [&] { b.Return(fn, b.Composite<mat2x3<f32>>(1_f, 2_f, 3_f, 4_f, 5_f, 6_f)); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Return_mat2x3f_Splat) {
+ auto* fn = b.Function("Function", ty.mat2x3<f32>());
+ b.Append(fn->Block(), [&] { b.Return(fn, b.Splat<mat2x3<f32>>(1_f, 6)); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Return_array_f32_Composite) {
+ auto* fn = b.Function("Function", ty.array<f32, 3>());
+ b.Append(fn->Block(), [&] { b.Return(fn, b.Composite<array<f32, 3>>(1_i, 2_i, 3_i)); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Return_array_f32_Splat) {
+ auto* fn = b.Function("Function", ty.array<f32, 3>());
+ b.Append(fn->Block(), [&] { b.Return(fn, b.Splat<array<f32, 3>>(1_i, 3)); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Construct) {
+ auto* fn = b.Function("Function", ty.void_());
+ b.Append(fn->Block(), [&] {
+ b.Construct<vec3<f32>>(1_f, 2_f, 3_f);
+ b.Return(fn);
+ });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Discard) {
+ auto* fn = b.Function("Function", ty.void_());
+ b.Append(fn->Block(), [&] {
+ b.Discard();
+ b.Return(fn);
+ });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Let) {
+ auto* fn = b.Function("Function", ty.void_());
+ b.Append(fn->Block(), [&] {
+ b.Let("Let", b.Constant(42_i));
+ b.Return(fn);
+ });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Var) {
+ auto* fn = b.Function("Function", ty.void_());
+ b.Append(fn->Block(), [&] {
+ b.Var<function>("Var", b.Constant(42_i));
+ b.Return(fn);
+ });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Access) {
+ auto* fn = b.Function("Function", ty.f32());
+ b.Append(fn->Block(),
+ [&] { b.Return(fn, b.Access<f32>(b.Construct<mat4x4<f32>>(), 1_u, 2_u)); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, UserCall) {
+ auto* fn_a = b.Function("A", ty.f32());
+ b.Append(fn_a->Block(), [&] { b.Return(fn_a, 42_f); });
+ auto* fn_b = b.Function("B", ty.f32());
+ b.Append(fn_b->Block(), [&] { b.Return(fn_b, b.Call(fn_a)); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Load) {
+ auto p = b.FunctionParam<ptr<function, f32, read_write>>("p");
+ auto* fn = b.Function("Function", ty.f32());
+ fn->SetParams({p});
+ b.Append(fn->Block(), [&] { b.Return(fn, b.Load(p)); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Store) {
+ auto p = b.FunctionParam<ptr<function, f32, read_write>>("p");
+ auto* fn = b.Function("Function", ty.void_());
+ fn->SetParams({p});
+ b.Append(fn->Block(), [&] {
+ b.Store(p, 42_f);
+ b.Return(fn);
+ });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, LoadVectorElement) {
+ auto p = b.FunctionParam<ptr<function, vec3<f32>, read_write>>("p");
+ auto* fn = b.Function("Function", ty.f32());
+ fn->SetParams({p});
+ b.Append(fn->Block(), [&] { b.Return(fn, b.LoadVectorElement(p, 1_i)); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, StoreVectorElement) {
+ auto p = b.FunctionParam<ptr<function, vec3<f32>, read_write>>("p");
+ auto* fn = b.Function("Function", ty.void_());
+ fn->SetParams({p});
+ b.Append(fn->Block(), [&] {
+ b.StoreVectorElement(p, 1_u, 42_f);
+ b.Return(fn);
+ });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, UnaryOp) {
+ auto x = b.FunctionParam<bool>("x");
+ auto* fn = b.Function("Function", ty.bool_());
+ fn->SetParams({x});
+ b.Append(fn->Block(), [&] { b.Return(fn, b.Not<bool>(x)); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, BinaryOp) {
+ auto x = b.FunctionParam<f32>("x");
+ auto y = b.FunctionParam<f32>("y");
+ auto* fn = b.Function("Function", ty.f32());
+ fn->SetParams({x, y});
+ b.Append(fn->Block(), [&] { b.Return(fn, b.Add<f32>(x, y)); });
+ RUN_TEST();
+}
+
+TEST_F(IRBinaryRoundtripTest, Swizzle) {
+ auto x = b.FunctionParam<vec4<f32>>("x");
+ auto* fn = b.Function("Function", ty.vec3<f32>());
+ fn->SetParams({x});
+ b.Append(fn->Block(), [&] {
+ b.Return(fn, b.Swizzle<vec3<f32>>(x, Vector<uint32_t, 3>{1, 0, 2}));
+ });
+ RUN_TEST();
+}
+
+} // namespace
+} // namespace tint::core::ir::binary
diff --git a/src/tint/lang/core/ir/builder.h b/src/tint/lang/core/ir/builder.h
index 2a89fa9..bdbd507 100644
--- a/src/tint/lang/core/ir/builder.h
+++ b/src/tint/lang/core/ir/builder.h
@@ -365,6 +365,17 @@
}
/// Creates a new ir::Constant
+ /// @tparam TYPE the splat type
+ /// @param value the splat value
+ /// @param size the number of items
+ /// @returns the new constant
+ template <typename TYPE, typename ARG>
+ ir::Constant* Splat(ARG&& value, size_t size) {
+ auto* type = ir.Types().Get<TYPE>();
+ return Splat(type, std::forward<ARG>(value), size);
+ }
+
+ /// Creates a new ir::Constant
/// @param ty the constant type
/// @param values the composite values
/// @returns the new constant
@@ -374,6 +385,16 @@
ir.constant_values.Composite(ty, Vector{ConstantValue(std::forward<ARGS>(values))...}));
}
+ /// Creates a new ir::Constant
+ /// @tparam TYPE the constant type
+ /// @param values the composite values
+ /// @returns the new constant
+ template <typename TYPE, typename... ARGS, typename = DisableIfVectorLike<ARGS...>>
+ ir::Constant* Composite(ARGS&&... values) {
+ auto* type = ir.Types().Get<TYPE>();
+ return Composite(type, std::forward<ARGS>(values)...);
+ }
+
/// Creates a new zero-value ir::Constant
/// @param ty the constant type
/// @returns the new constant
@@ -1270,6 +1291,17 @@
Values(std::forward<ARGS>(indices)...)));
}
+ /// Creates a new `Access`
+ /// @tparam TYPE the return type
+ /// @param object the object being accessed
+ /// @param indices the access indices
+ /// @returns the instruction
+ template <typename TYPE, typename OBJ, typename... ARGS>
+ ir::Access* Access(OBJ&& object, ARGS&&... indices) {
+ auto* type = ir.Types().Get<TYPE>();
+ return Access(type, std::forward<OBJ>(object), std::forward<ARGS>(indices)...);
+ }
+
/// Creates a new `Swizzle`
/// @param type the return type
/// @param object the object being swizzled
@@ -1283,6 +1315,17 @@
}
/// Creates a new `Swizzle`
+ /// @tparam TYPE the return type
+ /// @param object the object being swizzled
+ /// @param indices the swizzle indices
+ /// @returns the instruction
+ template <typename TYPE, typename OBJ>
+ ir::Swizzle* Swizzle(OBJ&& object, VectorRef<uint32_t> indices) {
+ auto* type = ir.Types().Get<TYPE>();
+ return Swizzle(type, std::forward<OBJ>(object), std::move(indices));
+ }
+
+ /// Creates a new `Swizzle`
/// @param type the return type
/// @param object the object being swizzled
/// @param indices the swizzle indices
diff --git a/src/tint/lang/core/ir/construct.cc b/src/tint/lang/core/ir/construct.cc
index fd67aee..8840b1f 100644
--- a/src/tint/lang/core/ir/construct.cc
+++ b/src/tint/lang/core/ir/construct.cc
@@ -36,6 +36,8 @@
namespace tint::core::ir {
+Construct::Construct() = default;
+
Construct::Construct(InstructionResult* result, VectorRef<Value*> arguments) {
AddOperands(Construct::kArgsOperandOffset, std::move(arguments));
AddResult(result);
diff --git a/src/tint/lang/core/ir/construct.h b/src/tint/lang/core/ir/construct.h
index ff9ca22..d0bbeb8 100644
--- a/src/tint/lang/core/ir/construct.h
+++ b/src/tint/lang/core/ir/construct.h
@@ -41,6 +41,9 @@
/// The base offset in Operands() for the args
static constexpr size_t kArgsOperandOffset = 0;
+ /// Constructor (no result, no operands)
+ Construct();
+
/// Constructor
/// @param result the result value
/// @param args the constructor arguments
diff --git a/src/tint/lang/core/ir/control_instruction.h b/src/tint/lang/core/ir/control_instruction.h
index cb172db..0f98763 100644
--- a/src/tint/lang/core/ir/control_instruction.h
+++ b/src/tint/lang/core/ir/control_instruction.h
@@ -54,31 +54,6 @@
/// @param cb the function to call once for each block
virtual void ForeachBlock(const std::function<void(ir::Block*)>& cb) = 0;
- /// Sets the results of the control instruction
- /// @param values the new result values
- void SetResults(VectorRef<InstructionResult*> values) {
- for (auto* value : results_) {
- if (value) {
- value->SetInstruction(nullptr);
- }
- }
- results_ = std::move(values);
- for (auto* value : results_) {
- if (value) {
- value->SetInstruction(this);
- }
- }
- }
-
- /// Sets the results of the control instruction
- /// @param values the new result values
- template <typename... ARGS,
- typename = std::enable_if_t<!tint::IsVectorLike<
- tint::traits::Decay<tint::traits::NthTypeOf<0, ARGS..., void>>>>>
- void SetResults(ARGS&&... values) {
- SetResults(Vector{std::forward<ARGS>(values)...});
- }
-
/// @return All the exits for the flow control instruction
const Hashset<Exit*, 2>& Exits() const { return exits_; }
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/instruction.h b/src/tint/lang/core/ir/instruction.h
index 886df26..f8c72d5 100644
--- a/src/tint/lang/core/ir/instruction.h
+++ b/src/tint/lang/core/ir/instruction.h
@@ -61,6 +61,14 @@
/// @returns the operands of the instruction
virtual VectorRef<const ir::Value*> Operands() const = 0;
+ /// Replaces the operands of the instruction
+ /// @param operands the new operands of the instruction
+ virtual void SetOperands(VectorRef<ir::Value*> operands) = 0;
+
+ /// Replaces the results of the instruction
+ /// @param results the new results of the instruction
+ virtual void SetResults(VectorRef<ir::InstructionResult*> results) = 0;
+
/// @returns the result values for this instruction
virtual VectorRef<InstructionResult*> Results() = 0;
diff --git a/src/tint/lang/core/ir/let.cc b/src/tint/lang/core/ir/let.cc
index ddbf9a3..aeb185e 100644
--- a/src/tint/lang/core/ir/let.cc
+++ b/src/tint/lang/core/ir/let.cc
@@ -35,6 +35,8 @@
namespace tint::core::ir {
+Let::Let() = default;
+
Let::Let(InstructionResult* result, ir::Value* value) {
AddOperand(Let::kValueOperandOffset, value);
AddResult(result);
diff --git a/src/tint/lang/core/ir/let.h b/src/tint/lang/core/ir/let.h
index a2fe222..96a318d 100644
--- a/src/tint/lang/core/ir/let.h
+++ b/src/tint/lang/core/ir/let.h
@@ -40,10 +40,14 @@
/// The offset in Operands() for the value
static constexpr size_t kValueOperandOffset = 0;
+ /// Constructor (no result, no operands)
+ Let();
+
/// Constructor
/// @param result the result value
/// @param value the let's value
Let(InstructionResult* result, Value* value);
+
~Let() override;
/// @copydoc Instruction::Clone()
diff --git a/src/tint/lang/core/ir/load.cc b/src/tint/lang/core/ir/load.cc
index ba9f6d3..6563ea1 100644
--- a/src/tint/lang/core/ir/load.cc
+++ b/src/tint/lang/core/ir/load.cc
@@ -36,6 +36,10 @@
namespace tint::core::ir {
+Load::Load() {
+ flags_.Add(Flag::kSequenced);
+}
+
Load::Load(InstructionResult* result, Value* from) {
flags_.Add(Flag::kSequenced);
diff --git a/src/tint/lang/core/ir/load.h b/src/tint/lang/core/ir/load.h
index 167fa87..987b4b1 100644
--- a/src/tint/lang/core/ir/load.h
+++ b/src/tint/lang/core/ir/load.h
@@ -41,6 +41,9 @@
/// The offset in Operands() for the from value
static constexpr size_t kFromOperandOffset = 0;
+ /// Constructor (no results, no operands)
+ Load();
+
/// Constructor (infers type)
/// @param result the result value
/// @param from the value being loaded from
diff --git a/src/tint/lang/core/ir/load_vector_element.cc b/src/tint/lang/core/ir/load_vector_element.cc
index 4db0a53..750a72e 100644
--- a/src/tint/lang/core/ir/load_vector_element.cc
+++ b/src/tint/lang/core/ir/load_vector_element.cc
@@ -34,6 +34,10 @@
namespace tint::core::ir {
+LoadVectorElement::LoadVectorElement() {
+ flags_.Add(Flag::kSequenced);
+}
+
LoadVectorElement::LoadVectorElement(InstructionResult* result, ir::Value* from, ir::Value* index) {
flags_.Add(Flag::kSequenced);
diff --git a/src/tint/lang/core/ir/load_vector_element.h b/src/tint/lang/core/ir/load_vector_element.h
index f8cbb37..f94ddbf 100644
--- a/src/tint/lang/core/ir/load_vector_element.h
+++ b/src/tint/lang/core/ir/load_vector_element.h
@@ -44,6 +44,9 @@
/// The offset in Operands() for the `index` value
static constexpr size_t kIndexOperandOffset = 1;
+ /// Constructor (no results, no operands)
+ LoadVectorElement();
+
/// Constructor
/// @param result the result value
/// @param from the vector pointer
diff --git a/src/tint/lang/core/ir/operand_instruction.h b/src/tint/lang/core/ir/operand_instruction.h
index 2b3f83a..cab7667 100644
--- a/src/tint/lang/core/ir/operand_instruction.h
+++ b/src/tint/lang/core/ir/operand_instruction.h
@@ -65,9 +65,9 @@
}
}
- /// Sets the operands to @p operands
- /// @param operands the new operands for the instruction
- void SetOperands(VectorRef<ir::Value*> operands) {
+ /// Replaces the operands of the instruction
+ /// @param operands the new operands of the instruction
+ void SetOperands(VectorRef<ir::Value*> operands) override {
ClearOperands();
operands_ = std::move(operands);
for (size_t i = 0; i < operands_.Length(); i++) {
@@ -88,8 +88,36 @@
operands_.Clear();
}
+ /// Replaces the results of the instruction
+ /// @param results the new results of the instruction
+ void SetResults(VectorRef<ir::InstructionResult*> results) override {
+ ClearResults();
+ results_ = std::move(results);
+ for (auto* result : results_) {
+ if (result) {
+ result->SetInstruction(this);
+ }
+ }
+ }
+
+ /// Sets the results of the instruction
+ /// @param values the new result values
+ template <typename... ARGS,
+ typename = std::enable_if_t<!tint::IsVectorLike<
+ tint::traits::Decay<tint::traits::NthTypeOf<0, ARGS..., void>>>>>
+ void SetResults(ARGS&&... values) {
+ SetResults(Vector{std::forward<ARGS>(values)...});
+ }
+
/// Removes all results from the instruction.
- void ClearResults() { results_.Clear(); }
+ void ClearResults() {
+ for (auto* result : results_) {
+ if (result && result->Instruction() == this) {
+ result->SetInstruction(nullptr);
+ }
+ }
+ results_.Clear();
+ }
/// @returns the operands of the instruction
VectorRef<ir::Value*> Operands() override { return operands_; }
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/ir/store.cc b/src/tint/lang/core/ir/store.cc
index 7acfd51..2f6ab39 100644
--- a/src/tint/lang/core/ir/store.cc
+++ b/src/tint/lang/core/ir/store.cc
@@ -34,6 +34,10 @@
namespace tint::core::ir {
+Store::Store() {
+ flags_.Add(Flag::kSequenced);
+}
+
Store::Store(Value* to, Value* from) {
flags_.Add(Flag::kSequenced);
diff --git a/src/tint/lang/core/ir/store.h b/src/tint/lang/core/ir/store.h
index 9ed7cc5..2ba7936 100644
--- a/src/tint/lang/core/ir/store.h
+++ b/src/tint/lang/core/ir/store.h
@@ -44,6 +44,9 @@
/// The offset in Operands() for the `from` value
static constexpr size_t kFromOperandOffset = 1;
+ /// Constructor (no results, no operands)
+ Store();
+
/// Constructor
/// @param to the value to store too
/// @param from the value being stored from
diff --git a/src/tint/lang/core/ir/store_vector_element.cc b/src/tint/lang/core/ir/store_vector_element.cc
index c429c81..947c28e 100644
--- a/src/tint/lang/core/ir/store_vector_element.cc
+++ b/src/tint/lang/core/ir/store_vector_element.cc
@@ -34,6 +34,10 @@
namespace tint::core::ir {
+StoreVectorElement::StoreVectorElement() {
+ flags_.Add(Flag::kSequenced);
+}
+
StoreVectorElement::StoreVectorElement(ir::Value* to, ir::Value* index, ir::Value* value) {
flags_.Add(Flag::kSequenced);
diff --git a/src/tint/lang/core/ir/store_vector_element.h b/src/tint/lang/core/ir/store_vector_element.h
index 2ef3895..91e8697 100644
--- a/src/tint/lang/core/ir/store_vector_element.h
+++ b/src/tint/lang/core/ir/store_vector_element.h
@@ -47,6 +47,9 @@
/// The offset in Operands() for the `value` value
static constexpr size_t kValueOperandOffset = 2;
+ /// Constructor (no operands)
+ StoreVectorElement();
+
/// Constructor
/// @param to the vector pointer
/// @param index the new vector element index
diff --git a/src/tint/lang/core/ir/swizzle.cc b/src/tint/lang/core/ir/swizzle.cc
index 1a1e6dc..184ddf6 100644
--- a/src/tint/lang/core/ir/swizzle.cc
+++ b/src/tint/lang/core/ir/swizzle.cc
@@ -37,6 +37,8 @@
namespace tint::core::ir {
+Swizzle::Swizzle() = default;
+
Swizzle::Swizzle(InstructionResult* result, Value* object, VectorRef<uint32_t> indices)
: indices_(std::move(indices)) {
TINT_ASSERT(!indices.IsEmpty());
diff --git a/src/tint/lang/core/ir/swizzle.h b/src/tint/lang/core/ir/swizzle.h
index 1e5b0fc..88740d0 100644
--- a/src/tint/lang/core/ir/swizzle.h
+++ b/src/tint/lang/core/ir/swizzle.h
@@ -29,6 +29,7 @@
#define SRC_TINT_LANG_CORE_IR_SWIZZLE_H_
#include <string>
+#include <utility>
#include "src/tint/lang/core/ir/operand_instruction.h"
#include "src/tint/utils/rtti/castable.h"
@@ -41,6 +42,9 @@
/// The offset in Operands() for the object being swizzled
static constexpr size_t kObjectOperandOffset = 0;
+ /// Constructor (no results, no operands)
+ Swizzle();
+
/// Constructor
/// @param result the result value
/// @param object the object being swizzled
@@ -60,6 +64,9 @@
/// @returns the swizzle indices
VectorRef<uint32_t> Indices() const { return indices_; }
+ /// @param indices the new swizzle indices
+ void SetIndices(VectorRef<uint32_t> indices) { indices_ = std::move(indices); }
+
/// @returns the friendly name for the instruction
std::string FriendlyName() const override { return "swizzle"; }
diff --git a/src/tint/lang/core/ir/unary.cc b/src/tint/lang/core/ir/unary.cc
index b79a23c..d4de4e1 100644
--- a/src/tint/lang/core/ir/unary.cc
+++ b/src/tint/lang/core/ir/unary.cc
@@ -34,6 +34,8 @@
namespace tint::core::ir {
+Unary::Unary() = default;
+
Unary::Unary(InstructionResult* result, UnaryOp op, Value* val) : op_(op) {
AddOperand(Unary::kValueOperandOffset, val);
AddResult(result);
@@ -47,4 +49,13 @@
return ctx.ir.instructions.Create<Unary>(new_result, op_, val);
}
+std::string_view ToString(enum UnaryOp op) {
+ switch (op) {
+ case UnaryOp::kComplement:
+ return "complement";
+ case UnaryOp::kNegation:
+ return "negation";
+ }
+ return "<unknown>";
+}
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/unary.h b/src/tint/lang/core/ir/unary.h
index af03473..41a4875 100644
--- a/src/tint/lang/core/ir/unary.h
+++ b/src/tint/lang/core/ir/unary.h
@@ -47,6 +47,9 @@
/// The offset in Operands() for the value
static constexpr size_t kValueOperandOffset = 0;
+ /// Constructor (no results, no operands)
+ Unary();
+
/// Constructor
/// @param result the result value
/// @param op the unary operator
@@ -66,13 +69,26 @@
/// @returns the unary operator
UnaryOp Op() const { return op_; }
+ /// @param op the new unary operator
+ void SetOp(UnaryOp op) { op_ = op; }
+
/// @returns the friendly name for the instruction
std::string FriendlyName() const override { return "unary"; }
private:
- UnaryOp op_;
+ UnaryOp op_ = UnaryOp::kComplement;
};
+/// @param kind the enum value
+/// @returns the string for the given enum value
+std::string_view ToString(UnaryOp kind);
+
+/// Emits the name of the intrinsic type.
+template <typename STREAM, typename = traits::EnableIfIsOStream<STREAM>>
+auto& operator<<(STREAM& out, UnaryOp kind) {
+ return out << ToString(kind);
+}
+
} // namespace tint::core::ir
#endif // SRC_TINT_LANG_CORE_IR_UNARY_H_
diff --git a/src/tint/lang/core/ir/user_call.cc b/src/tint/lang/core/ir/user_call.cc
index 698cb62..6cadb32 100644
--- a/src/tint/lang/core/ir/user_call.cc
+++ b/src/tint/lang/core/ir/user_call.cc
@@ -36,6 +36,10 @@
namespace tint::core::ir {
+UserCall::UserCall() {
+ flags_.Add(Flag::kSequenced);
+}
+
UserCall::UserCall(InstructionResult* result, Function* func, VectorRef<Value*> arguments) {
flags_.Add(Flag::kSequenced);
AddOperand(UserCall::kFunctionOperandOffset, func);
diff --git a/src/tint/lang/core/ir/user_call.h b/src/tint/lang/core/ir/user_call.h
index 4036208..a7a9b97 100644
--- a/src/tint/lang/core/ir/user_call.h
+++ b/src/tint/lang/core/ir/user_call.h
@@ -45,6 +45,9 @@
/// The base offset in Operands() for the call arguments
static constexpr size_t kArgsOperandOffset = 1;
+ /// Constructor (no results, no operands)
+ UserCall();
+
/// Constructor
/// @param result the result value
/// @param func the function being called
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index a3a83df..9c24c34 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -484,15 +484,15 @@
for (size_t i = 0; i < results.Length(); ++i) {
auto* res = results[i];
if (!res) {
- AddResultError(inst, i, InstError(inst, "instruction result is undefined"));
+ AddResultError(inst, i, InstError(inst, "result is undefined"));
continue;
}
if (res->Instruction() == nullptr) {
- AddResultError(inst, i, InstError(inst, "instruction result source is undefined"));
+ AddResultError(inst, i, InstError(inst, "instruction of result is undefined"));
} else if (res->Instruction() != inst) {
AddResultError(inst, i,
- InstError(inst, "instruction result source has wrong instruction"));
+ InstError(inst, "instruction of result is a different instruction"));
}
}
@@ -506,14 +506,11 @@
// Note, a `nullptr` is a valid operand in some cases, like `var` so we can't just check
// for `nullptr` here.
if (!op->Alive()) {
- AddError(inst, i,
- InstError(inst, "instruction operand " + std::to_string(i) + " is not alive"));
+ AddError(inst, i, InstError(inst, "operand is not alive"));
}
if (!op->HasUsage(inst, i)) {
- AddError(
- inst, i,
- InstError(inst, "instruction operand " + std::to_string(i) + " missing usage"));
+ AddError(inst, i, InstError(inst, "operand missing usage"));
}
}
diff --git a/src/tint/lang/core/ir/validator_test.cc b/src/tint/lang/core/ir/validator_test.cc
index 37cbbcc..bb43fe3 100644
--- a/src/tint/lang/core/ir/validator_test.cc
+++ b/src/tint/lang/core/ir/validator_test.cc
@@ -833,7 +833,7 @@
auto res = ir::Validate(mod);
ASSERT_FALSE(res);
- EXPECT_EQ(res.Failure().reason.str(), R"(:3:5 error: if: instruction result is undefined
+ EXPECT_EQ(res.Failure().reason.str(), R"(:3:5 error: if: result is undefined
undef = if true [t: %b2, f: %b3] { # if_1
^^^^^
@@ -905,7 +905,7 @@
auto res = ir::Validate(mod);
ASSERT_FALSE(res);
- EXPECT_EQ(res.Failure().reason.str(), R"(:2:3 error: var: instruction result is undefined
+ EXPECT_EQ(res.Failure().reason.str(), R"(:2:3 error: var: result is undefined
undef = var
^^^^^
@@ -932,7 +932,7 @@
auto res = ir::Validate(mod);
ASSERT_FALSE(res);
- EXPECT_EQ(res.Failure().reason.str(), R"(:3:5 error: var: instruction result is undefined
+ EXPECT_EQ(res.Failure().reason.str(), R"(:3:5 error: var: result is undefined
undef = var
^^^^^
@@ -991,7 +991,7 @@
auto res = ir::Validate(mod);
ASSERT_FALSE(res);
- EXPECT_EQ(res.Failure().reason.str(), R"(:3:5 error: let: instruction result is undefined
+ EXPECT_EQ(res.Failure().reason.str(), R"(:3:5 error: let: result is undefined
undef = let 1i
^^^^^
@@ -1105,7 +1105,7 @@
EXPECT_EQ(res.Failure().reason.str(), expected);
}
-TEST_F(IR_ValidatorTest, Instruction_NullSource) {
+TEST_F(IR_ValidatorTest, Instruction_NullInstruction) {
auto* f = b.Function("my_func", ty.void_());
auto sb = b.Append(f->Block());
@@ -1116,7 +1116,8 @@
auto res = ir::Validate(mod);
ASSERT_FALSE(res);
- EXPECT_EQ(res.Failure().reason.str(), R"(:3:5 error: var: instruction result source is undefined
+ EXPECT_EQ(res.Failure().reason.str(),
+ R"(:3:5 error: var: instruction of result is undefined
%2:ptr<function, f32, read_write> = var
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -1147,7 +1148,7 @@
auto res = ir::Validate(mod);
ASSERT_FALSE(res);
- EXPECT_EQ(res.Failure().reason.str(), R"(:3:46 error: var: instruction operand 0 is not alive
+ EXPECT_EQ(res.Failure().reason.str(), R"(:3:46 error: var: operand is not alive
%2:ptr<function, f32, read_write> = var, %3
^^
@@ -1178,7 +1179,7 @@
auto res = ir::Validate(mod);
ASSERT_FALSE(res);
- EXPECT_EQ(res.Failure().reason.str(), R"(:3:46 error: var: instruction operand 0 missing usage
+ EXPECT_EQ(res.Failure().reason.str(), R"(:3:46 error: var: operand missing usage
%2:ptr<function, f32, read_write> = var, %3
^^
@@ -1285,7 +1286,7 @@
auto res = ir::Validate(mod);
ASSERT_FALSE(res);
- EXPECT_EQ(res.Failure().reason.str(), R"(:3:5 error: binary: instruction result is undefined
+ EXPECT_EQ(res.Failure().reason.str(), R"(:3:5 error: binary: result is undefined
undef = add 3i, 2i
^^^^^
@@ -1342,7 +1343,7 @@
auto res = ir::Validate(mod);
ASSERT_FALSE(res);
- EXPECT_EQ(res.Failure().reason.str(), R"(:3:5 error: unary: instruction result is undefined
+ EXPECT_EQ(res.Failure().reason.str(), R"(:3:5 error: unary: result is undefined
undef = negation 2i
^^^^^
@@ -2985,7 +2986,7 @@
auto res = ir::Validate(mod);
ASSERT_FALSE(res);
EXPECT_EQ(res.Failure().reason.str(),
- R"(:4:5 error: load_vector_element: instruction result is undefined
+ R"(:4:5 error: load_vector_element: result is undefined
undef = load_vector_element %2, 1i
^^^^^
diff --git a/src/tint/lang/core/ir/var.cc b/src/tint/lang/core/ir/var.cc
index b9a8f91..32b1c83 100644
--- a/src/tint/lang/core/ir/var.cc
+++ b/src/tint/lang/core/ir/var.cc
@@ -37,6 +37,8 @@
namespace tint::core::ir {
+Var::Var() = default;
+
Var::Var(InstructionResult* result) {
if (result && result->Type()) {
TINT_ASSERT(result->Type()->Is<core::type::Pointer>());
diff --git a/src/tint/lang/core/ir/var.h b/src/tint/lang/core/ir/var.h
index cd72271..4220a4c 100644
--- a/src/tint/lang/core/ir/var.h
+++ b/src/tint/lang/core/ir/var.h
@@ -58,6 +58,9 @@
/// The offset in Operands() for the initializer
static constexpr size_t kInitializerOperandOffset = 0;
+ /// Constructor (no results, no operands)
+ Var();
+
/// Constructor
/// @param result the result value
explicit Var(InstructionResult* result);
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/lang/glsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
index 9694635..fbc7d34 100644
--- a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
@@ -189,6 +189,7 @@
polyfills.saturate = true;
polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
polyfills.workgroup_uniform_load = true;
+ polyfills.dot_4x8_packed = true;
data.Add<ast::transform::BuiltinPolyfill::Config>(polyfills);
manager.Add<ast::transform::BuiltinPolyfill>(); // Must come before DirectVariableAccess
}
@@ -264,7 +265,6 @@
"GLSL", builder_.AST(), diagnostics_,
Vector{
wgsl::Extension::kChromiumDisableUniformityAnalysis,
- wgsl::Extension::kChromiumExperimentalDp4A,
wgsl::Extension::kChromiumExperimentalFullPtrParameters,
wgsl::Extension::kChromiumInternalDualSourceBlending,
wgsl::Extension::kChromiumExperimentalPushConstant,
diff --git a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
index c668179..0ffbb1d 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
@@ -380,7 +380,6 @@
"HLSL", builder_.AST(), diagnostics_,
Vector{
wgsl::Extension::kChromiumDisableUniformityAnalysis,
- wgsl::Extension::kChromiumExperimentalDp4A,
wgsl::Extension::kChromiumExperimentalFullPtrParameters,
wgsl::Extension::kChromiumExperimentalPushConstant,
wgsl::Extension::kChromiumExperimentalSubgroups,
@@ -1246,8 +1245,8 @@
if (builtin->IsAtomic()) {
return EmitWorkgroupAtomicCall(out, expr, builtin);
}
- if (builtin->IsDP4a()) {
- return EmitDP4aCall(out, expr, builtin);
+ if (builtin->IsPacked4x8IntegerDotProductBuiltin()) {
+ return EmitPacked4x8IntegerDotProductBuiltinCall(out, expr, builtin);
}
if (builtin->IsSubgroup()) {
if (builtin->Fn() == wgsl::BuiltinFn::kSubgroupBroadcast) {
@@ -2521,10 +2520,9 @@
});
}
-bool ASTPrinter::EmitDP4aCall(StringStream& out,
- const ast::CallExpression* expr,
- const sem::BuiltinFn* builtin) {
- // TODO(crbug.com/tint/1497): support the polyfill version of DP4a functions.
+bool ASTPrinter::EmitPacked4x8IntegerDotProductBuiltinCall(StringStream& out,
+ const ast::CallExpression* expr,
+ const sem::BuiltinFn* builtin) {
return CallBuiltinHelper(
out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
std::string functionName;
diff --git a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.h b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.h
index 0c449b2..633180c 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.h
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.h
@@ -286,14 +286,15 @@
bool EmitTruncCall(StringStream& out,
const ast::CallExpression* expr,
const sem::BuiltinFn* builtin);
- /// Handles generating a call to DP4a builtins (dot4I8Packed and dot4U8Packed)
+ /// Handles generating a call to the builtins defined in the language extension
+ /// `packed_4x8_integer_dot_product`.
/// @param out the output stream
/// @param expr the call expression
/// @param builtin the semantic information for the builtin
/// @returns true if the call expression is emitted
- bool EmitDP4aCall(StringStream& out,
- const ast::CallExpression* expr,
- const sem::BuiltinFn* builtin);
+ bool EmitPacked4x8IntegerDotProductBuiltinCall(StringStream& out,
+ const ast::CallExpression* expr,
+ const sem::BuiltinFn* builtin);
/// Handles generating a call to subgroup builtins.
/// @param out the output stream
/// @param expr the call expression
diff --git a/src/tint/lang/hlsl/writer/ast_printer/builtin_test.cc b/src/tint/lang/hlsl/writer/ast_printer/builtin_test.cc
index 0c5f91e..d9d6d89 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/builtin_test.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/builtin_test.cc
@@ -1449,7 +1449,7 @@
}
TEST_F(HlslASTPrinterTest_Builtin, Dot4I8Packed) {
- Enable(wgsl::Extension::kChromiumExperimentalDp4A);
+ Require(wgsl::LanguageFeature::kPacked4X8IntegerDotProduct);
auto* val1 = Var("val1", ty.u32());
auto* val2 = Var("val2", ty.u32());
@@ -1475,7 +1475,7 @@
}
TEST_F(HlslASTPrinterTest_Builtin, Dot4U8Packed) {
- Enable(wgsl::Extension::kChromiumExperimentalDp4A);
+ Require(wgsl::LanguageFeature::kPacked4X8IntegerDotProduct);
auto* val1 = Var("val1", ty.u32());
auto* val2 = Var("val2", ty.u32());
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
index 712eddb..a662d59 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -275,7 +275,6 @@
"MSL", builder_.AST(), diagnostics_,
Vector{
wgsl::Extension::kChromiumDisableUniformityAnalysis,
- wgsl::Extension::kChromiumExperimentalDp4A,
wgsl::Extension::kChromiumExperimentalFullPtrParameters,
wgsl::Extension::kChromiumExperimentalPixelLocal,
wgsl::Extension::kChromiumExperimentalSubgroups,
diff --git a/src/tint/lang/spirv/reader/ast_lower/BUILD.bazel b/src/tint/lang/spirv/reader/ast_lower/BUILD.bazel
index 4bb7e6a..0fdbcd1 100644
--- a/src/tint/lang/spirv/reader/ast_lower/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/ast_lower/BUILD.bazel
@@ -43,12 +43,14 @@
"decompose_strided_array.cc",
"decompose_strided_matrix.cc",
"fold_trivial_lets.cc",
+ "pass_workgroup_id_as_argument.cc",
],
hdrs = [
"atomics.h",
"decompose_strided_array.h",
"decompose_strided_matrix.h",
"fold_trivial_lets.h",
+ "pass_workgroup_id_as_argument.h",
],
deps = [
"//src/tint/api/common",
@@ -88,6 +90,7 @@
"decompose_strided_array_test.cc",
"decompose_strided_matrix_test.cc",
"fold_trivial_lets_test.cc",
+ "pass_workgroup_id_as_argument_test.cc",
],
deps = [
"//src/tint/api/common",
diff --git a/src/tint/lang/spirv/reader/ast_lower/BUILD.cmake b/src/tint/lang/spirv/reader/ast_lower/BUILD.cmake
index 6d3870b..1f1a04c 100644
--- a/src/tint/lang/spirv/reader/ast_lower/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/ast_lower/BUILD.cmake
@@ -49,6 +49,8 @@
lang/spirv/reader/ast_lower/decompose_strided_matrix.h
lang/spirv/reader/ast_lower/fold_trivial_lets.cc
lang/spirv/reader/ast_lower/fold_trivial_lets.h
+ lang/spirv/reader/ast_lower/pass_workgroup_id_as_argument.cc
+ lang/spirv/reader/ast_lower/pass_workgroup_id_as_argument.h
)
tint_target_add_dependencies(tint_lang_spirv_reader_ast_lower lib
@@ -91,6 +93,7 @@
lang/spirv/reader/ast_lower/decompose_strided_array_test.cc
lang/spirv/reader/ast_lower/decompose_strided_matrix_test.cc
lang/spirv/reader/ast_lower/fold_trivial_lets_test.cc
+ lang/spirv/reader/ast_lower/pass_workgroup_id_as_argument_test.cc
)
tint_target_add_dependencies(tint_lang_spirv_reader_ast_lower_test test
diff --git a/src/tint/lang/spirv/reader/ast_lower/BUILD.gn b/src/tint/lang/spirv/reader/ast_lower/BUILD.gn
index c2ded34..9381efc 100644
--- a/src/tint/lang/spirv/reader/ast_lower/BUILD.gn
+++ b/src/tint/lang/spirv/reader/ast_lower/BUILD.gn
@@ -52,6 +52,8 @@
"decompose_strided_matrix.h",
"fold_trivial_lets.cc",
"fold_trivial_lets.h",
+ "pass_workgroup_id_as_argument.cc",
+ "pass_workgroup_id_as_argument.h",
]
deps = [
"${tint_src_dir}/api/common",
@@ -91,6 +93,7 @@
"decompose_strided_array_test.cc",
"decompose_strided_matrix_test.cc",
"fold_trivial_lets_test.cc",
+ "pass_workgroup_id_as_argument_test.cc",
]
deps = [
"${tint_src_dir}:gmock_and_gtest",
diff --git a/src/tint/lang/spirv/reader/ast_lower/pass_workgroup_id_as_argument.cc b/src/tint/lang/spirv/reader/ast_lower/pass_workgroup_id_as_argument.cc
new file mode 100644
index 0000000..cfaf300
--- /dev/null
+++ b/src/tint/lang/spirv/reader/ast_lower/pass_workgroup_id_as_argument.cc
@@ -0,0 +1,163 @@
+// 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/ast_lower/pass_workgroup_id_as_argument.h"
+
+#include <utility>
+
+#include "src/tint/lang/wgsl/program/clone_context.h"
+#include "src/tint/lang/wgsl/program/program_builder.h"
+#include "src/tint/lang/wgsl/resolver/resolve.h"
+#include "src/tint/lang/wgsl/sem/function.h"
+#include "src/tint/lang/wgsl/sem/statement.h"
+#include "src/tint/utils/containers/hashmap.h"
+
+using namespace tint::core::fluent_types; // NOLINT
+
+TINT_INSTANTIATE_TYPEINFO(tint::spirv::reader::PassWorkgroupIdAsArgument);
+
+namespace tint::spirv::reader {
+
+/// PIMPL state for the transform.
+struct PassWorkgroupIdAsArgument::State {
+ /// The source program
+ const Program& src;
+ /// The target program builder
+ ProgramBuilder b;
+ /// The clone context
+ program::CloneContext ctx = {&b, &src, /* auto_clone_symbols */ true};
+ /// The semantic info.
+ const sem::Info& sem = src.Sem();
+
+ /// Map from function to the name of its workgroup_id parameter.
+ Hashmap<const ast::Function*, Symbol, 8> func_to_param;
+
+ /// Constructor
+ /// @param program the source program
+ explicit State(const Program& program) : src(program) {}
+
+ /// Runs the transform.
+ /// @returns the new program
+ ApplyResult Run() {
+ // Process all entry points in the module, looking for workgroup_id builtin parameters.
+ bool made_changes = false;
+ for (auto* func : src.AST().Functions()) {
+ if (func->IsEntryPoint()) {
+ for (auto* param : func->params) {
+ if (auto* builtin =
+ ast::GetAttribute<ast::BuiltinAttribute>(param->attributes)) {
+ if (sem.Get(builtin)->Value() == core::BuiltinValue::kWorkgroupId) {
+ ProcessBuiltin(func, param);
+ made_changes = true;
+ }
+ }
+ }
+ }
+ }
+ if (!made_changes) {
+ return SkipTransform;
+ }
+
+ ctx.Clone();
+ return resolver::Resolve(b);
+ }
+
+ /// Process a workgroup_id builtin.
+ /// @param ep the entry point function
+ /// @param builtin the builtin parameter
+ void ProcessBuiltin(const ast::Function* ep, const ast::Parameter* builtin) {
+ // Record the name of the parameter for the entry point function.
+ func_to_param.Add(ep, ctx.Clone(builtin->name->symbol));
+
+ // The reader should only produce a single use of the parameter which assigns to a global.
+ const auto& users = sem.Get(builtin)->Users();
+ TINT_ASSERT_OR_RETURN(users.Length() == 1u);
+ auto* assign = users[0]->Stmt()->Declaration()->As<ast::AssignmentStatement>();
+ auto& stmts =
+ sem.Get(assign)->Parent()->Declaration()->As<ast::BlockStatement>()->statements;
+ auto* rhs = assign->rhs;
+ if (auto* bitcast = rhs->As<ast::BitcastExpression>()) {
+ // The RHS may be bitcast to a signed integer, so we capture that bitcast.
+ auto let = b.Symbols().New("tint_wgid_bitcast");
+ ctx.InsertBefore(stmts, assign, b.Decl(b.Let(let, ctx.Clone(bitcast))));
+ func_to_param.Replace(ep, let);
+ rhs = bitcast->expr;
+ }
+ TINT_ASSERT_OR_RETURN(assign && rhs == users[0]->Declaration());
+ auto* lhs = sem.GetVal(assign->lhs)->As<sem::VariableUser>();
+ TINT_ASSERT_OR_RETURN(lhs &&
+ lhs->Variable()->AddressSpace() == core::AddressSpace::kPrivate);
+
+ // Replace all references to the global variable with a function parameter.
+ for (auto* user : lhs->Variable()->Users()) {
+ if (user == lhs) {
+ // Skip the assignment, which will be removed.
+ continue;
+ }
+ auto param = GetParameter(user->Stmt()->Function()->Declaration(),
+ lhs->Variable()->Declaration()->type);
+ ctx.Replace(user->Declaration(), b.Expr(param));
+ }
+
+ // Remove the global variable and the assignment to it.
+ ctx.Remove(src.AST().GlobalDeclarations(), lhs->Variable()->Declaration());
+ ctx.Remove(stmts, assign);
+ }
+
+ /// Get the workgroup_id parameter for a function, creating it and updating callsites if needed.
+ /// @param func the function
+ /// @param type the type of the parameter
+ /// @returns the name of the parameter
+ Symbol GetParameter(const ast::Function* func, const ast::Type& type) {
+ return func_to_param.GetOrCreate(func, [&] {
+ // Append a new parameter to the function.
+ auto name = b.Symbols().New("tint_wgid");
+ ctx.InsertBack(func->params, b.Param(name, ctx.Clone(type)));
+
+ // Recursively update all callsites to pass the workgroup_id as an argument.
+ for (auto* callsite : sem.Get(func)->CallSites()) {
+ auto param = GetParameter(callsite->Stmt()->Function()->Declaration(), type);
+ ctx.InsertBack(callsite->Declaration()->args, b.Expr(param));
+ }
+
+ return name;
+ });
+ }
+};
+
+PassWorkgroupIdAsArgument::PassWorkgroupIdAsArgument() = default;
+
+PassWorkgroupIdAsArgument::~PassWorkgroupIdAsArgument() = default;
+
+ast::transform::Transform::ApplyResult PassWorkgroupIdAsArgument::Apply(
+ const Program& src,
+ const ast::transform::DataMap&,
+ ast::transform::DataMap&) const {
+ return State(src).Run();
+}
+
+} // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/ast_lower/pass_workgroup_id_as_argument.h b/src/tint/lang/spirv/reader/ast_lower/pass_workgroup_id_as_argument.h
new file mode 100644
index 0000000..e680352
--- /dev/null
+++ b/src/tint/lang/spirv/reader/ast_lower/pass_workgroup_id_as_argument.h
@@ -0,0 +1,59 @@
+// 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_AST_LOWER_PASS_WORKGROUP_ID_AS_ARGUMENT_H_
+#define SRC_TINT_LANG_SPIRV_READER_AST_LOWER_PASS_WORKGROUP_ID_AS_ARGUMENT_H_
+
+#include "src/tint/lang/wgsl/ast/transform/transform.h"
+
+namespace tint::spirv::reader {
+
+/// PassWorkgroupIdAsArgument is a transform that passes the workgroup_id builtin as an argument to
+/// functions that need it, instead of using a module-scope private variable. This allows the
+/// uniformity analysis to see that it is uniform, enabling shaders that use barriers in control
+/// flow guarded by this builtin.
+class PassWorkgroupIdAsArgument final
+ : public Castable<PassWorkgroupIdAsArgument, ast::transform::Transform> {
+ public:
+ /// Constructor
+ PassWorkgroupIdAsArgument();
+
+ /// Destructor
+ ~PassWorkgroupIdAsArgument() override;
+
+ /// @copydoc ast::transform::Transform::Apply
+ ApplyResult Apply(const Program& program,
+ const ast::transform::DataMap& inputs,
+ ast::transform::DataMap& outputs) const override;
+
+ private:
+ struct State;
+};
+
+} // namespace tint::spirv::reader
+
+#endif // SRC_TINT_LANG_SPIRV_READER_AST_LOWER_PASS_WORKGROUP_ID_AS_ARGUMENT_H_
diff --git a/src/tint/lang/spirv/reader/ast_lower/pass_workgroup_id_as_argument_test.cc b/src/tint/lang/spirv/reader/ast_lower/pass_workgroup_id_as_argument_test.cc
new file mode 100644
index 0000000..3c1bbcd
--- /dev/null
+++ b/src/tint/lang/spirv/reader/ast_lower/pass_workgroup_id_as_argument_test.cc
@@ -0,0 +1,408 @@
+// 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/ast_lower/pass_workgroup_id_as_argument.h"
+
+#include "src/tint/lang/wgsl/ast/transform/helper_test.h"
+
+namespace tint::spirv::reader {
+namespace {
+
+using PassWorkgroupIdAsArgumentTest = ast::transform::TransformTest;
+
+TEST_F(PassWorkgroupIdAsArgumentTest, Basic) {
+ auto* src = R"(
+enable chromium_disable_uniformity_analysis;
+
+var<private> wgid : vec3u;
+
+fn inner() {
+ if (wgid.x == 0) {
+ workgroupBarrier();
+ }
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wgid_param : vec3u) {
+ wgid = wgid_param;
+ inner();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_disable_uniformity_analysis;
+
+fn inner(tint_wgid : vec3u) {
+ if ((tint_wgid.x == 0)) {
+ workgroupBarrier();
+ }
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wgid_param : vec3u) {
+ inner(wgid_param);
+}
+)";
+
+ auto got = Run<PassWorkgroupIdAsArgument>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PassWorkgroupIdAsArgumentTest, MultipleUses) {
+ auto* src = R"(
+enable chromium_disable_uniformity_analysis;
+
+var<private> wgid : vec3u;
+
+fn inner() {
+ if (wgid.x == 0) {
+ workgroupBarrier();
+ }
+ if (wgid.y == 0) {
+ workgroupBarrier();
+ }
+ if (wgid.z == 0) {
+ workgroupBarrier();
+ }
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wgid_param : vec3u) {
+ wgid = wgid_param;
+ inner();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_disable_uniformity_analysis;
+
+fn inner(tint_wgid : vec3u) {
+ if ((tint_wgid.x == 0)) {
+ workgroupBarrier();
+ }
+ if ((tint_wgid.y == 0)) {
+ workgroupBarrier();
+ }
+ if ((tint_wgid.z == 0)) {
+ workgroupBarrier();
+ }
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wgid_param : vec3u) {
+ inner(wgid_param);
+}
+)";
+
+ auto got = Run<PassWorkgroupIdAsArgument>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PassWorkgroupIdAsArgumentTest, NestedCall) {
+ auto* src = R"(
+enable chromium_disable_uniformity_analysis;
+
+var<private> wgid : vec3u;
+
+fn inner_2() {
+ if (wgid.x == 0) {
+ workgroupBarrier();
+ }
+}
+
+fn inner_1() {
+ inner_2();
+}
+
+fn inner() {
+ inner_1();
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wgid_param : vec3u) {
+ wgid = wgid_param;
+ inner();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_disable_uniformity_analysis;
+
+fn inner_2(tint_wgid : vec3u) {
+ if ((tint_wgid.x == 0)) {
+ workgroupBarrier();
+ }
+}
+
+fn inner_1(tint_wgid_1 : vec3u) {
+ inner_2(tint_wgid_1);
+}
+
+fn inner(tint_wgid_2 : vec3u) {
+ inner_1(tint_wgid_2);
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wgid_param : vec3u) {
+ inner(wgid_param);
+}
+)";
+
+ auto got = Run<PassWorkgroupIdAsArgument>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PassWorkgroupIdAsArgumentTest, NestedCall_UsesAtEachLevel) {
+ auto* src = R"(
+enable chromium_disable_uniformity_analysis;
+
+var<private> wgid : vec3u;
+
+fn inner_2() {
+ if (wgid.x == 0) {
+ workgroupBarrier();
+ }
+}
+
+fn inner_1() {
+ inner_2();
+ if (wgid.y == 0) {
+ workgroupBarrier();
+ }
+}
+
+fn inner() {
+ inner_1();
+ if (wgid.z == 0) {
+ workgroupBarrier();
+ }
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wgid_param : vec3u) {
+ wgid = wgid_param;
+ inner();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_disable_uniformity_analysis;
+
+fn inner_2(tint_wgid : vec3u) {
+ if ((tint_wgid.x == 0)) {
+ workgroupBarrier();
+ }
+}
+
+fn inner_1(tint_wgid_1 : vec3u) {
+ inner_2(tint_wgid_1);
+ if ((tint_wgid_1.y == 0)) {
+ workgroupBarrier();
+ }
+}
+
+fn inner(tint_wgid_2 : vec3u) {
+ inner_1(tint_wgid_2);
+ if ((tint_wgid_2.z == 0)) {
+ workgroupBarrier();
+ }
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wgid_param : vec3u) {
+ inner(wgid_param);
+}
+)";
+
+ auto got = Run<PassWorkgroupIdAsArgument>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PassWorkgroupIdAsArgumentTest, NestedCall_MultipleCallsites) {
+ auto* src = R"(
+enable chromium_disable_uniformity_analysis;
+
+var<private> wgid : vec3u;
+
+fn inner_2() {
+ if (wgid.x == 0) {
+ workgroupBarrier();
+ }
+}
+
+fn inner_1() {
+ inner_2();
+ inner_2();
+ inner_2();
+}
+
+fn inner() {
+ inner_1();
+ inner_2();
+ inner_1();
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wgid_param : vec3u) {
+ wgid = wgid_param;
+ inner();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_disable_uniformity_analysis;
+
+fn inner_2(tint_wgid : vec3u) {
+ if ((tint_wgid.x == 0)) {
+ workgroupBarrier();
+ }
+}
+
+fn inner_1(tint_wgid_1 : vec3u) {
+ inner_2(tint_wgid_1);
+ inner_2(tint_wgid_1);
+ inner_2(tint_wgid_1);
+}
+
+fn inner(tint_wgid_2 : vec3u) {
+ inner_1(tint_wgid_2);
+ inner_2(tint_wgid_2);
+ inner_1(tint_wgid_2);
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wgid_param : vec3u) {
+ inner(wgid_param);
+}
+)";
+
+ auto got = Run<PassWorkgroupIdAsArgument>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PassWorkgroupIdAsArgumentTest, NestedCall_OtherParameters) {
+ auto* src = R"(
+enable chromium_disable_uniformity_analysis;
+
+var<private> wgid : vec3u;
+
+fn inner_2(a : u32, b : u32) {
+ if (wgid.x + a == b) {
+ workgroupBarrier();
+ }
+}
+
+fn inner_1(a : u32) {
+ inner_2(a, 1);
+}
+
+fn inner() {
+ inner_1(2);
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wgid_param : vec3u) {
+ wgid = wgid_param;
+ inner();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_disable_uniformity_analysis;
+
+fn inner_2(a : u32, b : u32, tint_wgid : vec3u) {
+ if (((tint_wgid.x + a) == b)) {
+ workgroupBarrier();
+ }
+}
+
+fn inner_1(a : u32, tint_wgid_1 : vec3u) {
+ inner_2(a, 1, tint_wgid_1);
+}
+
+fn inner(tint_wgid_2 : vec3u) {
+ inner_1(2, tint_wgid_2);
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wgid_param : vec3u) {
+ inner(wgid_param);
+}
+)";
+
+ auto got = Run<PassWorkgroupIdAsArgument>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PassWorkgroupIdAsArgumentTest, BitcastToI32) {
+ auto* src = R"(
+enable chromium_disable_uniformity_analysis;
+
+var<private> wgid : vec3i;
+
+fn inner() {
+ if (wgid.x == 0i) {
+ workgroupBarrier();
+ }
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wgid_param : vec3u) {
+ wgid = bitcast<vec3i>(wgid_param);
+ inner();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_disable_uniformity_analysis;
+
+fn inner(tint_wgid : vec3i) {
+ if ((tint_wgid.x == 0i)) {
+ workgroupBarrier();
+ }
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(workgroup_id) wgid_param : vec3u) {
+ let tint_wgid_bitcast = bitcast<vec3i>(wgid_param);
+ inner(tint_wgid_bitcast);
+}
+)";
+
+ auto got = Run<PassWorkgroupIdAsArgument>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+} // namespace
+} // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/ast_parser/parse.cc b/src/tint/lang/spirv/reader/ast_parser/parse.cc
index e3f8cf6..c1b6af7 100644
--- a/src/tint/lang/spirv/reader/ast_parser/parse.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/parse.cc
@@ -33,16 +33,49 @@
#include "src/tint/lang/spirv/reader/ast_lower/decompose_strided_array.h"
#include "src/tint/lang/spirv/reader/ast_lower/decompose_strided_matrix.h"
#include "src/tint/lang/spirv/reader/ast_lower/fold_trivial_lets.h"
+#include "src/tint/lang/spirv/reader/ast_lower/pass_workgroup_id_as_argument.h"
#include "src/tint/lang/spirv/reader/ast_parser/ast_parser.h"
#include "src/tint/lang/wgsl/ast/transform/manager.h"
#include "src/tint/lang/wgsl/ast/transform/remove_unreachable_statements.h"
#include "src/tint/lang/wgsl/ast/transform/simplify_pointers.h"
#include "src/tint/lang/wgsl/ast/transform/unshadow.h"
+#include "src/tint/lang/wgsl/extension.h"
#include "src/tint/lang/wgsl/program/clone_context.h"
#include "src/tint/lang/wgsl/resolver/resolve.h"
namespace tint::spirv::reader::ast_parser {
+namespace {
+
+/// Trivial transform that removes the enable directive that disables the uniformity analysis.
+class ReenableUniformityAnalysis final
+ : public Castable<ReenableUniformityAnalysis, ast::transform::Transform> {
+ public:
+ ReenableUniformityAnalysis() {}
+ ~ReenableUniformityAnalysis() override {}
+
+ /// @copydoc ast::transform::Transform::Apply
+ ApplyResult Apply(const Program& src,
+ const ast::transform::DataMap&,
+ ast::transform::DataMap&) const override {
+ ProgramBuilder b;
+ program::CloneContext ctx = {&b, &src, /* auto_clone_symbols */ true};
+
+ // Remove the extension that disables the uniformity analysis.
+ for (auto* enable : src.AST().Enables()) {
+ if (enable->HasExtension(wgsl::Extension::kChromiumDisableUniformityAnalysis) &&
+ enable->extensions.Length() == 1u) {
+ ctx.Remove(src.AST().GlobalDeclarations(), enable);
+ }
+ }
+
+ ctx.Clone();
+ return resolver::Resolve(b);
+ }
+};
+
+} // namespace
+
Program Parse(const std::vector<uint32_t>& input, const Options& options) {
ASTParser parser(input);
bool parsed = parser.Parse();
@@ -60,13 +93,19 @@
builder.DiagnosticDirective(wgsl::DiagnosticSeverity::kOff, "derivative_uniformity");
}
+ // Disable the uniformity analysis temporarily.
+ // We will run transforms that attempt to change the AST to satisfy the analysis.
+ auto allowed_features = options.allowed_features;
+ allowed_features.extensions.insert(wgsl::Extension::kChromiumDisableUniformityAnalysis);
+ builder.Enable(wgsl::Extension::kChromiumDisableUniformityAnalysis);
+
// The SPIR-V parser can construct disjoint AST nodes, which is invalid for
// the Resolver. Clone the Program to clean these up.
Program program_with_disjoint_ast(std::move(builder));
ProgramBuilder output;
program::CloneContext(&output, &program_with_disjoint_ast, false).Clone();
- auto program = Program(resolver::Resolve(output, options.allowed_features));
+ auto program = Program(resolver::Resolve(output, allowed_features));
if (!program.IsValid()) {
return program;
}
@@ -76,11 +115,15 @@
manager.Add<ast::transform::Unshadow>();
manager.Add<ast::transform::SimplifyPointers>();
manager.Add<FoldTrivialLets>();
+ manager.Add<PassWorkgroupIdAsArgument>();
manager.Add<DecomposeStridedMatrix>();
manager.Add<DecomposeStridedArray>();
manager.Add<ast::transform::RemoveUnreachableStatements>();
manager.Add<Atomics>();
+ manager.Add<ReenableUniformityAnalysis>();
return manager.Run(program, {}, outputs);
}
} // namespace tint::spirv::reader::ast_parser
+
+TINT_INSTANTIATE_TYPEINFO(tint::spirv::reader::ast_parser::ReenableUniformityAnalysis);
diff --git a/src/tint/lang/spirv/reader/ast_parser/parser_test.cc b/src/tint/lang/spirv/reader/ast_parser/parser_test.cc
index 7375186..404e925 100644
--- a/src/tint/lang/spirv/reader/ast_parser/parser_test.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/parser_test.cc
@@ -91,6 +91,43 @@
EXPECT_EQ(program.Diagnostics().count(), 0u) << errs;
}
+TEST_F(ParserTest, WorkgroupIdGuardingBarrier) {
+ auto spv = test::Assemble(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %foo "foo" %wgid
+ OpExecutionMode %foo LocalSize 1 1 1
+ OpDecorate %wgid BuiltIn WorkgroupId
+ %uint = OpTypeInt 32 0
+ %vec3u = OpTypeVector %uint 3
+%_ptr_Input_vec3u = OpTypePointer Input %vec3u
+ %uint_0 = OpConstant %uint 0
+ %uint_2 = OpConstant %uint 2
+ %uint_8 = OpConstant %uint 8
+ %wgid = OpVariable %_ptr_Input_vec3u Input
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %func_type = OpTypeFunction %void
+ %foo = OpFunction %void None %func_type
+ %foo_start = OpLabel
+ %wgid_value = OpLoad %vec3u %wgid
+ %wgid_x = OpCompositeExtract %uint %wgid_value 0
+ %condition = OpIEqual %bool %wgid_x %uint_0
+ OpSelectionMerge %merge None
+ OpBranchConditional %condition %true_branch %merge
+%true_branch = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_8
+ OpBranch %merge
+ %merge = OpLabel
+ OpReturn
+ OpFunctionEnd
+)");
+ auto program = Parse(spv, {});
+ auto errs = program.Diagnostics().str();
+ EXPECT_TRUE(program.IsValid()) << errs;
+ EXPECT_EQ(program.Diagnostics().count(), 0u) << errs;
+}
+
// TODO(dneto): uint32 vec, valid SPIR-V
// TODO(dneto): uint32 vec, invalid SPIR-V
diff --git a/src/tint/lang/spirv/writer/ast_printer/ast_builtin_test.cc b/src/tint/lang/spirv/writer/ast_printer/ast_builtin_test.cc
index 6d11ca9..d9eee77 100644
--- a/src/tint/lang/spirv/writer/ast_printer/ast_builtin_test.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/ast_builtin_test.cc
@@ -4168,11 +4168,11 @@
} // namespace synchronization_builtin_tests
-// Tests for DP4A builtins, tint:1497
-namespace DP4A_builtin_tests {
+// Tests for `packed_4x8_integer_dot_product` builtins, tint:1497
+namespace Packed_4x8_integer_dot_product_builtin_tests {
TEST_F(BuiltinSpirvASTPrinterTest, Call_Dot4I8Packed) {
- Enable(wgsl::Extension::kChromiumExperimentalDp4A);
+ Require(wgsl::LanguageFeature::kPacked4X8IntegerDotProduct);
auto* val1 = Var("val1", ty.u32());
auto* val2 = Var("val2", ty.u32());
@@ -4184,7 +4184,10 @@
ASSERT_TRUE(b.GenerateFunction(func)) << b.Diagnostics();
auto got = DumpModule(b.Module());
- auto expect = R"(OpEntryPoint GLCompute %3 "test_function"
+ auto expect = R"(OpCapability DotProduct
+OpCapability DotProductInput4x8BitPacked
+OpExtension "SPV_KHR_integer_dot_product"
+OpEntryPoint GLCompute %3 "test_function"
OpExecutionMode %3 LocalSize 1 1 1
OpName %3 "test_function"
OpName %5 "val1"
@@ -4209,7 +4212,7 @@
}
TEST_F(BuiltinSpirvASTPrinterTest, Call_Dot4U8Packed) {
- Enable(wgsl::Extension::kChromiumExperimentalDp4A);
+ Require(wgsl::LanguageFeature::kPacked4X8IntegerDotProduct);
auto* val1 = Var("val1", ty.u32());
auto* val2 = Var("val2", ty.u32());
@@ -4221,7 +4224,10 @@
ASSERT_TRUE(b.GenerateFunction(func)) << b.Diagnostics();
auto got = DumpModule(b.Module());
- auto expect = R"(OpEntryPoint GLCompute %3 "test_function"
+ auto expect = R"(OpCapability DotProduct
+OpCapability DotProductInput4x8BitPacked
+OpExtension "SPV_KHR_integer_dot_product"
+OpEntryPoint GLCompute %3 "test_function"
OpExecutionMode %3 LocalSize 1 1 1
OpName %3 "test_function"
OpName %5 "val1"
@@ -4244,7 +4250,7 @@
EXPECT_EQ(got, expect);
}
-} // namespace DP4A_builtin_tests
+} // namespace Packed_4x8_integer_dot_product_builtin_tests
} // namespace
} // namespace tint::spirv::writer
diff --git a/src/tint/lang/spirv/writer/ast_printer/ast_printer.cc b/src/tint/lang/spirv/writer/ast_printer/ast_printer.cc
index 61ceab1..2fe2ab2 100644
--- a/src/tint/lang/spirv/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/ast_printer.cc
@@ -145,6 +145,7 @@
polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
polyfills.quantize_to_vec_f16 = true; // crbug.com/tint/1741
polyfills.workgroup_uniform_load = true;
+ polyfills.dot_4x8_packed = options.polyfill_dot_4x8_packed;
data.Add<ast::transform::BuiltinPolyfill::Config>(polyfills);
manager.Add<ast::transform::BuiltinPolyfill>(); // Must come before DirectVariableAccess
}
diff --git a/src/tint/lang/spirv/writer/ast_printer/builder.cc b/src/tint/lang/spirv/writer/ast_printer/builder.cc
index 2e1271e..26c4d3a 100644
--- a/src/tint/lang/spirv/writer/ast_printer/builder.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/builder.cc
@@ -277,7 +277,6 @@
"SPIR-V", builder_.AST(), builder_.Diagnostics(),
Vector{
wgsl::Extension::kChromiumDisableUniformityAnalysis,
- wgsl::Extension::kChromiumExperimentalDp4A,
wgsl::Extension::kChromiumExperimentalFullPtrParameters,
wgsl::Extension::kChromiumExperimentalPushConstant,
wgsl::Extension::kChromiumExperimentalSubgroups,
@@ -348,11 +347,6 @@
bool Builder::GenerateExtension(wgsl::Extension extension) {
switch (extension) {
- case wgsl::Extension::kChromiumExperimentalDp4A:
- module_.PushExtension("SPV_KHR_integer_dot_product");
- module_.PushCapability(SpvCapabilityDotProductKHR);
- module_.PushCapability(SpvCapabilityDotProductInput4x8BitPackedKHR);
- break;
case wgsl::Extension::kF16:
module_.PushCapability(SpvCapabilityFloat16);
module_.PushCapability(SpvCapabilityUniformAndStorageBuffer16BitAccess);
@@ -2534,6 +2528,7 @@
}
break;
case wgsl::BuiltinFn::kDot4I8Packed: {
+ DeclarePacked4x8IntegerDotProductCapabilitiesAndExtensions();
auto first_param_id = get_arg_as_value_id(0);
auto second_param_id = get_arg_as_value_id(1);
if (!push_function_inst(spv::Op::OpSDotKHR,
@@ -2546,6 +2541,7 @@
return result_id;
}
case wgsl::BuiltinFn::kDot4U8Packed: {
+ DeclarePacked4x8IntegerDotProductCapabilitiesAndExtensions();
auto first_param_id = get_arg_as_value_id(0);
auto second_param_id = get_arg_as_value_id(1);
if (!push_function_inst(spv::Op::OpUDotKHR,
@@ -4180,6 +4176,12 @@
return true;
}
+void Builder::DeclarePacked4x8IntegerDotProductCapabilitiesAndExtensions() {
+ module_.PushExtension("SPV_KHR_integer_dot_product");
+ module_.PushCapability(SpvCapabilityDotProductKHR);
+ module_.PushCapability(SpvCapabilityDotProductInput4x8BitPackedKHR);
+}
+
Builder::ContinuingInfo::ContinuingInfo(const ast::Statement* the_last_statement,
uint32_t loop_id,
uint32_t break_id)
diff --git a/src/tint/lang/spirv/writer/ast_printer/builder.h b/src/tint/lang/spirv/writer/ast_printer/builder.h
index 6737973..21df95e 100644
--- a/src/tint/lang/spirv/writer/ast_printer/builder.h
+++ b/src/tint/lang/spirv/writer/ast_printer/builder.h
@@ -528,6 +528,10 @@
/// Pops the top-most scope
void PopScope();
+ /// Declare all the extensions and capabilities required by `OpSDot` and `OpUDot` using 4x8
+ // packed integer vectors as input.
+ void DeclarePacked4x8IntegerDotProductCapabilitiesAndExtensions();
+
ProgramBuilder builder_;
writer::Module module_;
Function current_function_;
diff --git a/src/tint/lang/spirv/writer/common/options.h b/src/tint/lang/spirv/writer/common/options.h
index f5c7897..2c8a117 100644
--- a/src/tint/lang/spirv/writer/common/options.h
+++ b/src/tint/lang/spirv/writer/common/options.h
@@ -143,6 +143,9 @@
/// SPIRV module. Issue: dawn:464
bool experimental_require_subgroup_uniform_control_flow = false;
+ /// Set to `true` to generate polyfill for `dot4I8Packed` and `dot4U8Packed` builtins
+ bool polyfill_dot_4x8_packed = false;
+
/// The bindings
Bindings bindings;
@@ -155,6 +158,7 @@
emit_vertex_point_size,
clamp_frag_depth,
experimental_require_subgroup_uniform_control_flow,
+ polyfill_dot_4x8_packed,
bindings);
};
diff --git a/src/tint/lang/wgsl/ast/transform/builtin_polyfill_test.cc b/src/tint/lang/wgsl/ast/transform/builtin_polyfill_test.cc
index 7045762..b33c31a 100644
--- a/src/tint/lang/wgsl/ast/transform/builtin_polyfill_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/builtin_polyfill_test.cc
@@ -4067,8 +4067,6 @@
TEST_F(BuiltinPolyfillTest, Dot4I8Packed) {
auto* src = R"(
-enable chromium_experimental_dp4a;
-
fn f() {
let v1 = 0x01020304u;
let v2 = 0xF1F2F3F4u;
@@ -4077,8 +4075,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_dp4a;
-
fn tint_dot4_i8_packed(a : u32, b : u32) -> i32 {
const n = vec4<u32>(24, 16, 8, 0);
let a_i8 = (bitcast<vec4<i32>>((vec4<u32>(a) << n)) >> vec4<u32>(24));
@@ -4100,8 +4096,6 @@
TEST_F(BuiltinPolyfillTest, Dot4U8Packed) {
auto* src = R"(
-enable chromium_experimental_dp4a;
-
fn f() {
let v1 = 0x01020304u;
let v2 = 0xF1F2F3F4u;
@@ -4110,8 +4104,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_dp4a;
-
fn tint_dot4_u8_packed(a : u32, b : u32) -> u32 {
const n = vec4<u32>(24, 16, 8, 0);
let a_u8 = ((vec4<u32>(a) >> n) & vec4<u32>(255));
diff --git a/src/tint/lang/wgsl/builtin_fn.cc b/src/tint/lang/wgsl/builtin_fn.cc
index d392fcb..c43c067 100644
--- a/src/tint/lang/wgsl/builtin_fn.cc
+++ b/src/tint/lang/wgsl/builtin_fn.cc
@@ -694,7 +694,7 @@
f == BuiltinFn::kAtomicCompareExchangeWeak;
}
-bool IsDP4a(BuiltinFn f) {
+bool IsPacked4x8IntegerDotProductBuiltin(BuiltinFn f) {
return f == BuiltinFn::kDot4I8Packed || f == BuiltinFn::kDot4U8Packed;
}
diff --git a/src/tint/lang/wgsl/builtin_fn.cc.tmpl b/src/tint/lang/wgsl/builtin_fn.cc.tmpl
index cf79bee..a255971 100644
--- a/src/tint/lang/wgsl/builtin_fn.cc.tmpl
+++ b/src/tint/lang/wgsl/builtin_fn.cc.tmpl
@@ -100,7 +100,7 @@
f == BuiltinFn::kAtomicCompareExchangeWeak;
}
-bool IsDP4a(BuiltinFn f) {
+bool IsPacked4x8IntegerDotProductBuiltin(BuiltinFn f) {
return f == BuiltinFn::kDot4I8Packed || f == BuiltinFn::kDot4U8Packed;
}
diff --git a/src/tint/lang/wgsl/builtin_fn.h b/src/tint/lang/wgsl/builtin_fn.h
index 6829f67..d9c68da 100644
--- a/src/tint/lang/wgsl/builtin_fn.h
+++ b/src/tint/lang/wgsl/builtin_fn.h
@@ -471,10 +471,12 @@
/// @returns true if the given `f` is an atomic builtin
bool IsAtomic(BuiltinFn f);
-/// Determines if the given `f` is a DP4a builtin.
+/// Determines if the given `f` is a builtin defined in the language extension
+/// `packed_4x8_integer_dot_product`.
/// @param f the builtin type
-/// @returns true if the given `f` is a DP4a builtin
-bool IsDP4a(BuiltinFn f);
+/// @returns true if the given `f` is a builtin defined in the language extension
+/// `packed_4x8_integer_dot_product`.
+bool IsPacked4x8IntegerDotProductBuiltin(BuiltinFn f);
/// Determines if the given `f` is a subgroup builtin.
/// @param f the builtin type
diff --git a/src/tint/lang/wgsl/builtin_fn.h.tmpl b/src/tint/lang/wgsl/builtin_fn.h.tmpl
index e69eda2..57d0fe8 100644
--- a/src/tint/lang/wgsl/builtin_fn.h.tmpl
+++ b/src/tint/lang/wgsl/builtin_fn.h.tmpl
@@ -108,10 +108,12 @@
/// @returns true if the given `f` is an atomic builtin
bool IsAtomic(BuiltinFn f);
-/// Determines if the given `f` is a DP4a builtin.
+/// Determines if the given `f` is a builtin defined in the language extension
+/// `packed_4x8_integer_dot_product`.
/// @param f the builtin type
-/// @returns true if the given `f` is a DP4a builtin
-bool IsDP4a(BuiltinFn f);
+/// @returns true if the given `f` is a builtin defined in the language extension
+/// `packed_4x8_integer_dot_product`.
+bool IsPacked4x8IntegerDotProductBuiltin(BuiltinFn f);
/// Determines if the given `f` is a subgroup builtin.
/// @param f the builtin type
diff --git a/src/tint/lang/wgsl/extension.cc b/src/tint/lang/wgsl/extension.cc
index 6b001ef..c2f1762 100644
--- a/src/tint/lang/wgsl/extension.cc
+++ b/src/tint/lang/wgsl/extension.cc
@@ -45,9 +45,6 @@
if (str == "chromium_disable_uniformity_analysis") {
return Extension::kChromiumDisableUniformityAnalysis;
}
- if (str == "chromium_experimental_dp4a") {
- return Extension::kChromiumExperimentalDp4A;
- }
if (str == "chromium_experimental_framebuffer_fetch") {
return Extension::kChromiumExperimentalFramebufferFetch;
}
@@ -81,8 +78,6 @@
return "undefined";
case Extension::kChromiumDisableUniformityAnalysis:
return "chromium_disable_uniformity_analysis";
- case Extension::kChromiumExperimentalDp4A:
- return "chromium_experimental_dp4a";
case Extension::kChromiumExperimentalFramebufferFetch:
return "chromium_experimental_framebuffer_fetch";
case Extension::kChromiumExperimentalFullPtrParameters:
diff --git a/src/tint/lang/wgsl/extension.h b/src/tint/lang/wgsl/extension.h
index 259c925..41edcba 100644
--- a/src/tint/lang/wgsl/extension.h
+++ b/src/tint/lang/wgsl/extension.h
@@ -47,7 +47,6 @@
enum class Extension : uint8_t {
kUndefined,
kChromiumDisableUniformityAnalysis,
- kChromiumExperimentalDp4A,
kChromiumExperimentalFramebufferFetch,
kChromiumExperimentalFullPtrParameters,
kChromiumExperimentalPixelLocal,
@@ -76,17 +75,20 @@
Extension ParseExtension(std::string_view str);
constexpr std::string_view kExtensionStrings[] = {
- "chromium_disable_uniformity_analysis", "chromium_experimental_dp4a",
- "chromium_experimental_framebuffer_fetch", "chromium_experimental_full_ptr_parameters",
- "chromium_experimental_pixel_local", "chromium_experimental_push_constant",
- "chromium_experimental_subgroups", "chromium_internal_dual_source_blending",
- "chromium_internal_relaxed_uniform_layout", "f16",
+ "chromium_disable_uniformity_analysis",
+ "chromium_experimental_framebuffer_fetch",
+ "chromium_experimental_full_ptr_parameters",
+ "chromium_experimental_pixel_local",
+ "chromium_experimental_push_constant",
+ "chromium_experimental_subgroups",
+ "chromium_internal_dual_source_blending",
+ "chromium_internal_relaxed_uniform_layout",
+ "f16",
};
/// All extensions
static constexpr Extension kAllExtensions[] = {
Extension::kChromiumDisableUniformityAnalysis,
- Extension::kChromiumExperimentalDp4A,
Extension::kChromiumExperimentalFramebufferFetch,
Extension::kChromiumExperimentalFullPtrParameters,
Extension::kChromiumExperimentalPixelLocal,
diff --git a/src/tint/lang/wgsl/extension_bench.cc b/src/tint/lang/wgsl/extension_bench.cc
index 63c8636..c2e6bd2 100644
--- a/src/tint/lang/wgsl/extension_bench.cc
+++ b/src/tint/lang/wgsl/extension_bench.cc
@@ -52,69 +52,62 @@
"chromium_dis1ble_uniformity_analysis",
"chromium_qqisable_unifomity_anaJysis",
"chrollium_disable_uniformity_analysi77",
- "chromippHm_experqqmetal_dp4a",
- "chrmium_expecimntal_dp4",
- "chrmiumGexpebimental_dp4a",
- "chromium_experimental_dp4a",
- "chromium_exverimentiil_dp4a",
- "chro8ium_experimenWWal_dp4a",
- "chromiMm_eperimxxntal_dp4a",
- "cXromium_expermental_framggbuffer_fetch",
- "chVomiu_experimntal_fXauebuffer_fetch",
- "chromium_experimental_fr3mebuffer_fetch",
+ "cqqromium_eppperimental_framebuffe_fetcHH",
+ "chrmium_experimvntal_frcmebufer_ftch",
+ "chromium_expebimental_framGbufer_fetch",
"chromium_experimental_framebuffer_fetch",
- "chromium_experimental_fraEebuffer_fetch",
- "chromTTum_experimental_fraebuffePP_fetch",
- "chromum_experiddental_framebxxffer_fetch",
- "chromium_experimental_full_ptr_p44rameters",
- "SShromium_experimental_full_ptr_parameVVers",
- "chroRium_expRrimental_fu22l_ptr_prameters",
+ "chromium_experimental_vramebuffeii_fetch",
+ "chro8WWum_experimental_framebuffer_fetch",
+ "chromium_eperimenxxMl_framebuffer_fetch",
+ "chromium_expeggimeXtal_full_ptr_paraeters",
+ "chromium_expVrimental_full_ptr_puraXeer",
+ "chromium_experimental_full_ptr3parameters",
"chromium_experimental_full_ptr_parameters",
- "chromium_experimFntal_full_ptr_paramet9r",
- "chromium_experimentl_full_ptr_parameters",
- "chromiuHexperimental_fulOO_pVr_paramRRters",
- "chromium_experimenya_pixel_local",
- "chrromium_exp77rimentnnl_pixellllocGl",
- "chromium_exper4mental_pixel_lo00al",
+ "chromium_experimentalEfull_ptr_parameters",
+ "chromium_experimentalfull_ptr_PPaTTameters",
+ "chromium_ddxperimental_fullptrxxparameters",
+ "chromium_experi44ental_pixel_local",
+ "chromium_experimental_VVSixel_local",
+ "chroRium_experimental_pix22Rlocal",
"chromium_experimental_pixel_local",
- "chooomium_eperimenal_pxel_local",
- "chromium_experzzental_pixel_ocal",
- "chiiomiu_epperimental_pix11l_local",
- "chromXXum_experimental_push_constant",
- "chromII9um_experinnental_push_c55nstant",
- "chSSomium_experiHHental_pusrraaconstaYt",
+ "chromiuF_experiment9lpixel_local",
+ "chromium_experimental_pixel_loca",
+ "Vhromium_expeOOimentalHpixRRl_lcal",
+ "chromiym_experimental_push_contant",
+ "nnhro77ium_experimenGal_push_conrrllant",
+ "chromium_experimental_push_c4nstan00",
"chromium_experimental_push_constant",
- "chromium_epHrimentkk_psh_constant",
- "chromium_expegimenja_puRRh_costant",
- "chromium_bxperimental_push_contan",
- "chromium_experimental_sjbgroups",
- "chromium_experimental_sbgroups",
- "cromum_experimentalqsubgroups",
+ "chooomum_experimental_ush_constat",
+ "chromium_xperimntal_zzush_constant",
+ "chromi11m_experimepptal_psh_ciistant",
+ "chromium_experimental_subgroXXps",
+ "chromium55eIIperimental_subgnno99ps",
+ "chraamiuSS_experimentaHHr_subgrouYs",
"chromium_experimental_subgroups",
- "chromium_expNNrimental_subgoups",
- "chromium_experimetal_svvbgrous",
- "chromium_experiQental_subgroups",
- "chrorum_internal_dal_source_bleffding",
- "chromium_internal_dual_source_jlending",
- "chromiNNm_internal_dua8_sourwwe_blening",
+ "chkkomium_eperimntal_subgroup",
+ "jhromium_experRmental_subgogps",
+ "chromiubexperiental_subgroups",
+ "chromium_internal_dujl_source_blending",
+ "chromium_intenal_dual_source_blending",
+ "chqomium_internal_dual_source_beding",
"chromium_internal_dual_source_blending",
- "chromium_internal_dual_soure_blending",
- "chromium_irrternal_dual_source_blending",
- "chromium_internal_duaG_source_blending",
- "chromium_internalFFrelaxed_uniform_layout",
- "chromEum_internal_relaxed_unifrmlyout",
- "chromium_internalrrrelaxd_uniform_layout",
+ "chroium_NNnternal_dual_source_blending",
+ "chovvium_internal_dual_source_lending",
+ "chromium_internQQl_dual_sorce_blending",
+ "chromirm_intenal_rfflaxed_unifrm_layout",
+ "chromium_internal_jelaxed_uniform_layout",
+ "chromium_interna_relNNxed_uwwiform_lay82t",
"chromium_internal_relaxed_uniform_layout",
- "chromiuminternal_relaxed_uniform_layut",
- "cXroDium_internal_rJJlaed_uniform_layout",
- "chromium_int8nal_relaed_uniform_layut",
- "k",
- "16",
- "J1",
+ "chromium_internal_relaxed_uniform_layut",
+ "chromium_internal_relaxed_rrniform_layout",
+ "chromium_internal_relaxedGuniform_layout",
+ "FF16",
+ "",
+ "rr1",
"f16",
- "c16",
- "fO6",
- "_KKttvv",
+ "1",
+ "DJ1",
+ "",
};
for (auto _ : state) {
for (auto* str : kStrings) {
diff --git a/src/tint/lang/wgsl/extension_test.cc b/src/tint/lang/wgsl/extension_test.cc
index c40bacf..08cda9b 100644
--- a/src/tint/lang/wgsl/extension_test.cc
+++ b/src/tint/lang/wgsl/extension_test.cc
@@ -58,7 +58,6 @@
static constexpr Case kValidCases[] = {
{"chromium_disable_uniformity_analysis", Extension::kChromiumDisableUniformityAnalysis},
- {"chromium_experimental_dp4a", Extension::kChromiumExperimentalDp4A},
{"chromium_experimental_framebuffer_fetch", Extension::kChromiumExperimentalFramebufferFetch},
{"chromium_experimental_full_ptr_parameters",
Extension::kChromiumExperimentalFullPtrParameters},
@@ -74,33 +73,30 @@
{"chromium_disableuniformiccy_analysis", Extension::kUndefined},
{"chromil3_disable_unifority_analss", Extension::kUndefined},
{"chromium_disable_Vniformity_analysis", Extension::kUndefined},
- {"chro1ium_experimental_dp4a", Extension::kUndefined},
- {"chrJmium_experiqqetal_dp4a", Extension::kUndefined},
- {"chromium_experimenll77l_dp4a", Extension::kUndefined},
- {"cqqromium_eppperimental_framebuffe_fetcHH", Extension::kUndefined},
- {"chrmium_experimvntal_frcmebufer_ftch", Extension::kUndefined},
- {"chromium_expebimental_framGbufer_fetch", Extension::kUndefined},
- {"chvomium_exiierimental_full_ptr_parameters", Extension::kUndefined},
- {"chromium_WWxperimental_full8ptr_parameters", Extension::kUndefined},
- {"chromxxum_Mperimental_full_ptr_parameters", Extension::kUndefined},
- {"chromum_experimental_pixeX_loggal", Extension::kUndefined},
- {"chromium_expVrXmntal_ixel_local", Extension::kUndefined},
- {"3hromium_experimental_pixel_local", Extension::kUndefined},
- {"chromium_experEmental_push_constant", Extension::kUndefined},
- {"chPPomiumexperimental_push_conTTtant", Extension::kUndefined},
- {"chromixxm_experimentddl_push_constnt", Extension::kUndefined},
- {"chromium_experimental_44ubgroups", Extension::kUndefined},
- {"cSSromVVum_experimental_subgroups", Extension::kUndefined},
- {"chrmium_e22perimental_suRgrRups", Extension::kUndefined},
- {"chroFium_internal_dual_source_bl9ndig", Extension::kUndefined},
- {"chrmium_internal_dual_source_blending", Extension::kUndefined},
- {"cVromium_interHal_dualOOsouRRce_blening", Extension::kUndefined},
- {"chromium_internl_relaxyd_uniform_layout", Extension::kUndefined},
- {"chromnnum_internrr77_Gelaxell_uniform_layout", Extension::kUndefined},
- {"chromium_intern4l_relaxe00_uniform_layout", Extension::kUndefined},
- {"5", Extension::kUndefined},
- {"u16", Extension::kUndefined},
- {"f", Extension::kUndefined},
+ {"chromium_experimental_framebuf1er_fetch", Extension::kUndefined},
+ {"chromium_experiqqntal_framebuffer_fetch", Extension::kUndefined},
+ {"chromium_experimental_framebuffll77_fetch", Extension::kUndefined},
+ {"chroium_experimental_full_ptr_paqqppmetHHrs", Extension::kUndefined},
+ {"chrium_evperiental_full_ptr_paraceters", Extension::kUndefined},
+ {"chromium_expGimental_fullbptr_parameters", Extension::kUndefined},
+ {"vhromium_experimental_pixel_liical", Extension::kUndefined},
+ {"chromium_experiment8l_pixel_lWWcal", Extension::kUndefined},
+ {"chromium_expeimentMl_xxixel_local", Extension::kUndefined},
+ {"chrXmium_experimeggtal_ush_constant", Extension::kUndefined},
+ {"chromiu_experVmentalpusX_constant", Extension::kUndefined},
+ {"chro3ium_experimental_push_constant", Extension::kUndefined},
+ {"cEromium_experimental_subgroups", Extension::kUndefined},
+ {"TThromium_experiPPental_sugroups", Extension::kUndefined},
+ {"chddomium_experimental_subgroxxs", Extension::kUndefined},
+ {"chromium_internal_44ual_source_blending", Extension::kUndefined},
+ {"chromium_inteSSnal_dual_source_blendinVV", Extension::kUndefined},
+ {"chromiuR_interna22dual_source_blenRing", Extension::kUndefined},
+ {"chromium_int9rnal_relaxed_Fnifor_layout", Extension::kUndefined},
+ {"chrmium_internal_relaxed_uniform_layout", Extension::kUndefined},
+ {"VRhHomium_internal_relaxd_uniform_OOayout", Extension::kUndefined},
+ {"y1", Extension::kUndefined},
+ {"l77rrn6", Extension::kUndefined},
+ {"4016", Extension::kUndefined},
};
using ExtensionParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/lang/wgsl/features/language_feature.cc b/src/tint/lang/wgsl/features/language_feature.cc
index 708697b..0dc4b3a 100644
--- a/src/tint/lang/wgsl/features/language_feature.cc
+++ b/src/tint/lang/wgsl/features/language_feature.cc
@@ -57,6 +57,9 @@
if (str == "chromium_testing_unsafe_experimental") {
return LanguageFeature::kChromiumTestingUnsafeExperimental;
}
+ if (str == "packed_4x8_integer_dot_product") {
+ return LanguageFeature::kPacked4X8IntegerDotProduct;
+ }
if (str == "readonly_and_readwrite_storage_textures") {
return LanguageFeature::kReadonlyAndReadwriteStorageTextures;
}
@@ -77,6 +80,8 @@
return "chromium_testing_unimplemented";
case LanguageFeature::kChromiumTestingUnsafeExperimental:
return "chromium_testing_unsafe_experimental";
+ case LanguageFeature::kPacked4X8IntegerDotProduct:
+ return "packed_4x8_integer_dot_product";
case LanguageFeature::kReadonlyAndReadwriteStorageTextures:
return "readonly_and_readwrite_storage_textures";
}
diff --git a/src/tint/lang/wgsl/features/language_feature.h b/src/tint/lang/wgsl/features/language_feature.h
index 4788b2d..0f1d2c9 100644
--- a/src/tint/lang/wgsl/features/language_feature.h
+++ b/src/tint/lang/wgsl/features/language_feature.h
@@ -51,6 +51,7 @@
kChromiumTestingShippedWithKillswitch,
kChromiumTestingUnimplemented,
kChromiumTestingUnsafeExperimental,
+ kPacked4X8IntegerDotProduct,
kReadonlyAndReadwriteStorageTextures,
};
@@ -69,6 +70,7 @@
"chromium_testing_shipped_with_killswitch",
"chromium_testing_unimplemented",
"chromium_testing_unsafe_experimental",
+ "packed_4x8_integer_dot_product",
"readonly_and_readwrite_storage_textures",
};
@@ -79,6 +81,7 @@
LanguageFeature::kChromiumTestingShippedWithKillswitch,
LanguageFeature::kChromiumTestingUnimplemented,
LanguageFeature::kChromiumTestingUnsafeExperimental,
+ LanguageFeature::kPacked4X8IntegerDotProduct,
LanguageFeature::kReadonlyAndReadwriteStorageTextures,
};
diff --git a/src/tint/lang/wgsl/features/status.cc b/src/tint/lang/wgsl/features/status.cc
index 1520bbf..b0c4bb3 100644
--- a/src/tint/lang/wgsl/features/status.cc
+++ b/src/tint/lang/wgsl/features/status.cc
@@ -34,6 +34,7 @@
FeatureStatus GetLanguageFeatureStatus(LanguageFeature f) {
switch (f) {
case LanguageFeature::kReadonlyAndReadwriteStorageTextures:
+ case LanguageFeature::kPacked4X8IntegerDotProduct:
return FeatureStatus::kExperimental;
case LanguageFeature::kUndefined:
return FeatureStatus::kUnknown;
diff --git a/src/tint/lang/wgsl/helpers/check_supported_extensions_test.cc b/src/tint/lang/wgsl/helpers/check_supported_extensions_test.cc
index ed40d5a..415f32c 100644
--- a/src/tint/lang/wgsl/helpers/check_supported_extensions_test.cc
+++ b/src/tint/lang/wgsl/helpers/check_supported_extensions_test.cc
@@ -42,7 +42,7 @@
ASSERT_TRUE(CheckSupportedExtensions("writer", AST(), Diagnostics(),
Vector{
wgsl::Extension::kF16,
- wgsl::Extension::kChromiumExperimentalDp4A,
+ wgsl::Extension::kChromiumExperimentalSubgroups,
}));
}
@@ -51,7 +51,7 @@
ASSERT_FALSE(CheckSupportedExtensions("writer", AST(), Diagnostics(),
Vector{
- wgsl::Extension::kChromiumExperimentalDp4A,
+ wgsl::Extension::kChromiumExperimentalSubgroups,
}));
EXPECT_EQ(Diagnostics().str(), "12:34 error: writer backend does not support extension 'f16'");
}
diff --git a/src/tint/lang/wgsl/language_feature_test.cc b/src/tint/lang/wgsl/language_feature_test.cc
index 045f462..6a103ad 100644
--- a/src/tint/lang/wgsl/language_feature_test.cc
+++ b/src/tint/lang/wgsl/language_feature_test.cc
@@ -63,6 +63,7 @@
LanguageFeature::kChromiumTestingShippedWithKillswitch},
{"chromium_testing_unimplemented", LanguageFeature::kChromiumTestingUnimplemented},
{"chromium_testing_unsafe_experimental", LanguageFeature::kChromiumTestingUnsafeExperimental},
+ {"packed_4x8_integer_dot_product", LanguageFeature::kPacked4X8IntegerDotProduct},
{"readonly_and_readwrite_storage_textures",
LanguageFeature::kReadonlyAndReadwriteStorageTextures},
};
@@ -83,9 +84,12 @@
{"chXggmium_testing_unsafe_expermental", LanguageFeature::kUndefined},
{"Xhomiuu_testng_unsafe_experimental", LanguageFeature::kUndefined},
{"chromium_3esting_unsafe_experimental", LanguageFeature::kUndefined},
- {"readonly_and_readwrite_stErage_textures", LanguageFeature::kUndefined},
- {"readoTTly_and_readwrite_strage_tPPxtures", LanguageFeature::kUndefined},
- {"readoly_and_redddwrite_storaxxe_textures", LanguageFeature::kUndefined},
+ {"packed_4x8_integer_Eot_product", LanguageFeature::kUndefined},
+ {"paked_4x8_integePP_dTTt_product", LanguageFeature::kUndefined},
+ {"packed_4x8_integxxrdot_pddoduct", LanguageFeature::kUndefined},
+ {"readon44y_and_readwrite_storage_textures", LanguageFeature::kUndefined},
+ {"readonly_and_readwrite_storageVVSSextures", LanguageFeature::kUndefined},
+ {"rRadonly_an_rea22write_storRge_textures", LanguageFeature::kUndefined},
};
using LanguageFeatureParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc b/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc
index 6536271..0af6015 100644
--- a/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc
@@ -82,8 +82,8 @@
// Test a valid enable directive with multiple extensions.
TEST_F(EnableDirectiveTest, Multiple) {
- auto p =
- parser("enable f16, chromium_disable_uniformity_analysis, chromium_experimental_dp4a;");
+ auto p = parser(
+ "enable f16, chromium_disable_uniformity_analysis, chromium_experimental_subgroups;");
p->enable_directive();
EXPECT_FALSE(p->has_error()) << p->error();
auto program = p->program();
@@ -101,19 +101,19 @@
EXPECT_EQ(enable->extensions[1]->source.range.begin.column, 13u);
EXPECT_EQ(enable->extensions[1]->source.range.end.line, 1u);
EXPECT_EQ(enable->extensions[1]->source.range.end.column, 49u);
- EXPECT_EQ(enable->extensions[2]->name, wgsl::Extension::kChromiumExperimentalDp4A);
+ EXPECT_EQ(enable->extensions[2]->name, wgsl::Extension::kChromiumExperimentalSubgroups);
EXPECT_EQ(enable->extensions[2]->source.range.begin.line, 1u);
EXPECT_EQ(enable->extensions[2]->source.range.begin.column, 51u);
EXPECT_EQ(enable->extensions[2]->source.range.end.line, 1u);
- EXPECT_EQ(enable->extensions[2]->source.range.end.column, 77u);
+ EXPECT_EQ(enable->extensions[2]->source.range.end.column, 82u);
ASSERT_EQ(ast.GlobalDeclarations().Length(), 1u);
EXPECT_EQ(ast.GlobalDeclarations()[0], enable);
}
// Test a valid enable directive with multiple extensions.
TEST_F(EnableDirectiveTest, MultipleTrailingComma) {
- auto p =
- parser("enable f16, chromium_disable_uniformity_analysis, chromium_experimental_dp4a,;");
+ auto p = parser(
+ "enable f16, chromium_disable_uniformity_analysis, chromium_experimental_subgroups,;");
p->enable_directive();
EXPECT_FALSE(p->has_error()) << p->error();
auto program = p->program();
@@ -131,11 +131,11 @@
EXPECT_EQ(enable->extensions[1]->source.range.begin.column, 13u);
EXPECT_EQ(enable->extensions[1]->source.range.end.line, 1u);
EXPECT_EQ(enable->extensions[1]->source.range.end.column, 49u);
- EXPECT_EQ(enable->extensions[2]->name, wgsl::Extension::kChromiumExperimentalDp4A);
+ EXPECT_EQ(enable->extensions[2]->name, wgsl::Extension::kChromiumExperimentalSubgroups);
EXPECT_EQ(enable->extensions[2]->source.range.begin.line, 1u);
EXPECT_EQ(enable->extensions[2]->source.range.begin.column, 51u);
EXPECT_EQ(enable->extensions[2]->source.range.end.line, 1u);
- EXPECT_EQ(enable->extensions[2]->source.range.end.column, 77u);
+ EXPECT_EQ(enable->extensions[2]->source.range.end.column, 82u);
ASSERT_EQ(ast.GlobalDeclarations().Length(), 1u);
EXPECT_EQ(ast.GlobalDeclarations()[0], enable);
}
@@ -205,7 +205,7 @@
// Error when unknown extension found
EXPECT_TRUE(p->has_error());
EXPECT_EQ(p->error(), R"(1:8: expected extension
-Possible values: 'chromium_disable_uniformity_analysis', 'chromium_experimental_dp4a', 'chromium_experimental_framebuffer_fetch', 'chromium_experimental_full_ptr_parameters', 'chromium_experimental_pixel_local', 'chromium_experimental_push_constant', 'chromium_experimental_subgroups', 'chromium_internal_dual_source_blending', 'chromium_internal_relaxed_uniform_layout', 'f16')");
+Possible values: 'chromium_disable_uniformity_analysis', 'chromium_experimental_framebuffer_fetch', 'chromium_experimental_full_ptr_parameters', 'chromium_experimental_pixel_local', 'chromium_experimental_push_constant', 'chromium_experimental_subgroups', 'chromium_internal_dual_source_blending', 'chromium_internal_relaxed_uniform_layout', 'f16')");
auto program = p->program();
auto& ast = program.AST();
EXPECT_EQ(ast.Enables().Length(), 0u);
diff --git a/src/tint/lang/wgsl/resolver/builtin_validation_test.cc b/src/tint/lang/wgsl/resolver/builtin_validation_test.cc
index 60898ce..7a062b5 100644
--- a/src/tint/lang/wgsl/resolver/builtin_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/builtin_validation_test.cc
@@ -545,14 +545,11 @@
} // namespace texture_constexpr_args
-// TODO(crbug.com/tint/1497): Update or remove ResolverDP4aExtensionValidationTest when the
-// experimental extension chromium_experimental_dp4a is not needed.
-using ResolverDP4aExtensionValidationTest = ResolverTest;
+using ResolverPacked4x8IntegerDotProductExtensionValidationTest = ResolverTest;
-TEST_F(ResolverDP4aExtensionValidationTest, Dot4I8PackedWithExtension) {
- // enable chromium_experimental_dp4a;
+TEST_F(ResolverPacked4x8IntegerDotProductExtensionValidationTest, Dot4I8Packed) {
// fn func { return dot4I8Packed(1u, 2u); }
- Enable(wgsl::Extension::kChromiumExperimentalDp4A);
+ Require(wgsl::LanguageFeature::kPacked4X8IntegerDotProduct);
Func("func", tint::Empty, ty.i32(),
Vector{
@@ -563,7 +560,7 @@
EXPECT_TRUE(r()->Resolve());
}
-TEST_F(ResolverDP4aExtensionValidationTest, Dot4I8PackedWithoutExtension) {
+TEST_F(ResolverPacked4x8IntegerDotProductExtensionValidationTest, Dot4I8Packed_FeatureDisallowed) {
// fn func { return dot4I8Packed(1u, 2u); }
Func("func", tint::Empty, ty.i32(),
Vector{
@@ -571,16 +568,17 @@
Vector{Expr(1_u), Expr(2_u)})),
});
- EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(
- r()->error(),
- R"(12:34 error: cannot call built-in function 'dot4I8Packed' without extension chromium_experimental_dp4a)");
+ auto resolver = Resolver(this, {});
+ EXPECT_FALSE(resolver.Resolve());
+ EXPECT_EQ(resolver.error(),
+ "12:34 error: built-in function 'dot4I8Packed' requires the "
+ "packed_4x8_integer_dot_product language feature, which is not allowed in the "
+ "current environment");
}
-TEST_F(ResolverDP4aExtensionValidationTest, Dot4U8PackedWithExtension) {
- // enable chromium_experimental_dp4a;
+TEST_F(ResolverPacked4x8IntegerDotProductExtensionValidationTest, Dot4U8Packed) {
// fn func { return dot4U8Packed(1u, 2u); }
- Enable(wgsl::Extension::kChromiumExperimentalDp4A);
+ Require(wgsl::LanguageFeature::kPacked4X8IntegerDotProduct);
Func("func", tint::Empty, ty.u32(),
Vector{
@@ -591,7 +589,7 @@
EXPECT_TRUE(r()->Resolve());
}
-TEST_F(ResolverDP4aExtensionValidationTest, Dot4U8PackedWithoutExtension) {
+TEST_F(ResolverPacked4x8IntegerDotProductExtensionValidationTest, Dot4U8Packed_FeatureDisallowed) {
// fn func { return dot4U8Packed(1u, 2u); }
Func("func", tint::Empty, ty.u32(),
Vector{
@@ -599,10 +597,12 @@
Vector{Expr(1_u), Expr(2_u)})),
});
- EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(
- r()->error(),
- R"(12:34 error: cannot call built-in function 'dot4U8Packed' without extension chromium_experimental_dp4a)");
+ auto resolver = Resolver(this, {});
+ EXPECT_FALSE(resolver.Resolve());
+ EXPECT_EQ(resolver.error(),
+ "12:34 error: built-in function 'dot4U8Packed' requires the "
+ "packed_4x8_integer_dot_product language feature, which is not allowed in the "
+ "current environment");
}
TEST_F(ResolverBuiltinValidationTest, WorkgroupUniformLoad_WrongAddressSpace) {
diff --git a/src/tint/lang/wgsl/resolver/side_effects_test.cc b/src/tint/lang/wgsl/resolver/side_effects_test.cc
index 23cd725..b9fa38b 100644
--- a/src/tint/lang/wgsl/resolver/side_effects_test.cc
+++ b/src/tint/lang/wgsl/resolver/side_effects_test.cc
@@ -180,7 +180,6 @@
using SideEffectsBuiltinTest = resolver::ResolverTestWithParam<Case>;
TEST_P(SideEffectsBuiltinTest, Test) {
- Enable(tint::wgsl::Extension::kChromiumExperimentalDp4A);
auto& c = GetParam();
uint32_t next_binding = 0;
diff --git a/src/tint/lang/wgsl/sem/builtin_fn.cc b/src/tint/lang/wgsl/sem/builtin_fn.cc
index e180378..0e48dfc 100644
--- a/src/tint/lang/wgsl/sem/builtin_fn.cc
+++ b/src/tint/lang/wgsl/sem/builtin_fn.cc
@@ -93,8 +93,8 @@
return wgsl::IsAtomic(fn_);
}
-bool BuiltinFn::IsDP4a() const {
- return wgsl::IsDP4a(fn_);
+bool BuiltinFn::IsPacked4x8IntegerDotProductBuiltin() const {
+ return wgsl::IsPacked4x8IntegerDotProductBuiltin(fn_);
}
bool BuiltinFn::IsSubgroup() const {
@@ -106,9 +106,6 @@
}
wgsl::Extension BuiltinFn::RequiredExtension() const {
- if (IsDP4a()) {
- return wgsl::Extension::kChromiumExperimentalDp4A;
- }
if (IsSubgroup()) {
return wgsl::Extension::kChromiumExperimentalSubgroups;
}
@@ -119,6 +116,9 @@
if (fn_ == wgsl::BuiltinFn::kTextureBarrier) {
return wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures;
}
+ if (IsPacked4x8IntegerDotProductBuiltin()) {
+ return wgsl::LanguageFeature::kPacked4X8IntegerDotProduct;
+ }
return wgsl::LanguageFeature::kUndefined;
}
diff --git a/src/tint/lang/wgsl/sem/builtin_fn.h b/src/tint/lang/wgsl/sem/builtin_fn.h
index 83d62bf..7c5d12c 100644
--- a/src/tint/lang/wgsl/sem/builtin_fn.h
+++ b/src/tint/lang/wgsl/sem/builtin_fn.h
@@ -102,9 +102,9 @@
/// @returns true if builtin is a atomic builtin
bool IsAtomic() const;
- /// @returns true if builtin is a DP4a builtin (defined in the extension
- /// chromium_experimental_DP4a)
- bool IsDP4a() const;
+ /// @returns true if builtin is a builtin defined in the language extension
+ /// `packed_4x8_integer_dot_product`.
+ bool IsPacked4x8IntegerDotProductBuiltin() const;
/// @returns true if builtin is a subgroup builtin (defined in the extension
/// chromium_experimental_subgroups)
diff --git a/src/tint/lang/wgsl/wgsl.def b/src/tint/lang/wgsl/wgsl.def
index 7969b9f..1a3bbee 100644
--- a/src/tint/lang/wgsl/wgsl.def
+++ b/src/tint/lang/wgsl/wgsl.def
@@ -68,9 +68,6 @@
enum extension {
// WGSL Extension "f16"
f16
- // An extension for the experimental feature "chromium_experimental_dp4a".
- // See crbug.com/tint/1497 for more details
- chromium_experimental_dp4a
// A Chromium-specific extension for disabling uniformity analysis.
chromium_disable_uniformity_analysis
// A Chromium-specific extension for push constants
@@ -95,6 +92,7 @@
// https://gpuweb.github.io/gpuweb/wgsl/#language-extensions-sec
enum language_feature {
readonly_and_readwrite_storage_textures
+ packed_4x8_integer_dot_product
// Language features used only for testing whose status will never change.
chromium_testing_unimplemented
diff --git a/src/tint/utils/bytes/BUILD.bazel b/src/tint/utils/bytes/BUILD.bazel
index 1b60b1d..2bfa708 100644
--- a/src/tint/utils/bytes/BUILD.bazel
+++ b/src/tint/utils/bytes/BUILD.bazel
@@ -40,6 +40,7 @@
name = "bytes",
srcs = [
"bytes.cc",
+ "reader.cc",
],
hdrs = [
"decoder.h",
diff --git a/src/tint/utils/bytes/BUILD.cmake b/src/tint/utils/bytes/BUILD.cmake
index 1012db6..a3997fd 100644
--- a/src/tint/utils/bytes/BUILD.cmake
+++ b/src/tint/utils/bytes/BUILD.cmake
@@ -42,6 +42,7 @@
utils/bytes/bytes.cc
utils/bytes/decoder.h
utils/bytes/endianness.h
+ utils/bytes/reader.cc
utils/bytes/reader.h
utils/bytes/swap.h
)
diff --git a/src/tint/utils/bytes/BUILD.gn b/src/tint/utils/bytes/BUILD.gn
index 97931f1..c10512f 100644
--- a/src/tint/utils/bytes/BUILD.gn
+++ b/src/tint/utils/bytes/BUILD.gn
@@ -47,6 +47,7 @@
"bytes.cc",
"decoder.h",
"endianness.h",
+ "reader.cc",
"reader.h",
"swap.h",
]
diff --git a/src/tint/utils/bytes/decoder.h b/src/tint/utils/bytes/decoder.h
index 591f255..4496ad9 100644
--- a/src/tint/utils/bytes/decoder.h
+++ b/src/tint/utils/bytes/decoder.h
@@ -34,7 +34,6 @@
#include <utility>
#include "src/tint/utils/bytes/reader.h"
-#include "src/tint/utils/result/result.h"
namespace tint::bytes {
@@ -43,10 +42,11 @@
/// Decodes T from @p reader.
/// @param reader the byte reader
+/// @param args additional arguments used by Decoder<T>::Decode()
/// @returns the decoded object
-template <typename T>
-Result<T> Decode(Reader& reader) {
- return Decoder<T>::Decode(reader);
+template <typename T, typename... ARGS>
+Result<T> Decode(Reader& reader, ARGS&&... args) {
+ return Decoder<T>::Decode(reader, std::forward<ARGS>(args)...);
}
/// Decoder specialization for integer types
@@ -54,12 +54,10 @@
struct Decoder<T, std::enable_if_t<std::is_integral_v<T>>> {
/// Decode decodes the integer type from @p reader.
/// @param reader the reader to decode from
+ /// @param endianness the endianness of the integer
/// @returns the decoded integer type, or an error if the stream is too short.
- static Result<T> Decode(Reader& reader) {
- if (reader.BytesRemaining() < sizeof(T)) {
- return Failure{"EOF"};
- }
- return reader.Int<T>();
+ static Result<T> Decode(Reader& reader, Endianness endianness = Endianness::kLittle) {
+ return reader.Int<T>(endianness);
}
};
@@ -69,44 +67,31 @@
/// Decode decodes the floating point type from @p reader.
/// @param reader the reader to decode from
/// @returns the decoded floating point type, or an error if the stream is too short.
- static Result<T> Decode(Reader& reader) {
- if (reader.BytesRemaining() < sizeof(T)) {
- return Failure{"EOF"};
- }
- return reader.Float<T>();
- }
+ static Result<T> Decode(Reader& reader) { return reader.Float<T>(); }
};
/// Decoder specialization for a uint16_t length prefixed string.
-template <typename T>
-struct Decoder<T, std::enable_if_t<std::is_same_v<T, std::string>>> {
+template <>
+struct Decoder<std::string, void> {
/// Decode decodes the string from @p reader.
/// @param reader the reader to decode from
/// @returns the decoded string, or an error if the stream is too short.
- static Result<T> Decode(Reader& reader) {
- if (reader.BytesRemaining() < sizeof(uint16_t)) {
- return Failure{"EOF"};
- }
+ static Result<std::string> Decode(Reader& reader) {
auto len = reader.Int<uint16_t>();
- if (reader.BytesRemaining() < len) {
- return Failure{"EOF"};
+ if (!len) {
+ return len.Failure();
}
- return reader.String(len);
+ return reader.String(len.Get());
}
};
/// Decoder specialization for bool types
template <>
struct Decoder<bool, void> {
- static Result<bool> Decode(Reader& reader) {
- /// Decode decodes the boolean from @p reader.
- /// @param reader the reader to decode from
- /// @returns the decoded boolean, or an error if the stream is too short.
- if (reader.IsEOF()) {
- return Failure{"EOF"};
- }
- return reader.Bool();
- }
+ /// Decode decodes the boolean from @p reader.
+ /// @param reader the reader to decode from
+ /// @returns the decoded boolean, or an error if the stream is too short.
+ static Result<bool> Decode(Reader& reader) { return reader.Bool(); }
};
/// Decoder specialization for types that use TINT_REFLECT
@@ -143,10 +128,11 @@
std::unordered_map<K, V> out;
while (true) {
- if (reader.IsEOF()) {
- return Failure{"EOF"};
+ auto stop = bytes::Decode<bool>(reader);
+ if (!stop) {
+ return stop.Failure();
}
- if (reader.Bool()) {
+ if (stop.Get()) {
break;
}
auto key = bytes::Decode<K>(reader);
diff --git a/src/tint/utils/bytes/decoder_test.cc b/src/tint/utils/bytes/decoder_test.cc
index cce8e5a..12640b1 100644
--- a/src/tint/utils/bytes/decoder_test.cc
+++ b/src/tint/utils/bytes/decoder_test.cc
@@ -44,7 +44,7 @@
TEST(BytesDecoderTest, Uint8) {
auto data = Data(0x10, 0x20, 0x30, 0x40, 0x50, 0x60, 0x70, 0x80);
- auto reader = Reader{Slice{data}, 0, Endianness::kLittle};
+ auto reader = BufferReader{Slice{data}};
EXPECT_EQ(Decode<uint8_t>(reader).Get(), 0x10u);
EXPECT_EQ(Decode<uint8_t>(reader).Get(), 0x20u);
EXPECT_EQ(Decode<uint8_t>(reader).Get(), 0x30u);
@@ -58,7 +58,7 @@
TEST(BytesDecoderTest, Uint16) {
auto data = Data(0x10, 0x20, 0x30, 0x40, 0x50, 0x60, 0x70, 0x80);
- auto reader = Reader{Slice{data}, 0, Endianness::kLittle};
+ auto reader = BufferReader{Slice{data}};
EXPECT_EQ(Decode<uint16_t>(reader).Get(), 0x2010u);
EXPECT_EQ(Decode<uint16_t>(reader).Get(), 0x4030u);
EXPECT_EQ(Decode<uint16_t>(reader).Get(), 0x6050u);
@@ -68,22 +68,22 @@
TEST(BytesDecoderTest, Uint32) {
auto data = Data(0x10, 0x20, 0x30, 0x40, 0x50, 0x60, 0x70, 0x80);
- auto reader = Reader{Slice{data}, 0, Endianness::kBig};
- EXPECT_EQ(Decode<uint32_t>(reader).Get(), 0x10203040u);
- EXPECT_EQ(Decode<uint32_t>(reader).Get(), 0x50607080u);
+ auto reader = BufferReader{Slice{data}};
+ EXPECT_EQ(Decode<uint32_t>(reader, Endianness::kBig).Get(), 0x10203040u);
+ EXPECT_EQ(Decode<uint32_t>(reader, Endianness::kBig).Get(), 0x50607080u);
EXPECT_FALSE(Decode<uint32_t>(reader));
}
TEST(BytesDecoderTest, Float) {
auto data = Data(0x00, 0x00, 0x08, 0x41);
- auto reader = Reader{Slice{data}};
+ auto reader = BufferReader{Slice{data}};
EXPECT_EQ(Decode<float>(reader).Get(), 8.5f);
EXPECT_FALSE(Decode<float>(reader));
}
TEST(BytesDecoderTest, Bool) {
auto data = Data(0x0, 0x1, 0x2, 0x1, 0x0);
- auto reader = Reader{Slice{data}};
+ auto reader = BufferReader{Slice{data}};
EXPECT_EQ(Decode<bool>(reader).Get(), false);
EXPECT_EQ(Decode<bool>(reader).Get(), true);
EXPECT_EQ(Decode<bool>(reader).Get(), true);
@@ -93,8 +93,8 @@
}
TEST(BytesDecoderTest, String) {
- auto data = Data(0x0, 0x5, 'h', 'e', 'l', 'l', 'o', 0x0, 0x5, 'w', 'o', 'r', 'l', 'd');
- auto reader = Reader{Slice{data}, 0, Endianness::kBig};
+ auto data = Data(0x5, 0x0, 'h', 'e', 'l', 'l', 'o', 0x5, 0x0, 'w', 'o', 'r', 'l', 'd');
+ auto reader = BufferReader{Slice{data}};
EXPECT_EQ(Decode<std::string>(reader).Get(), "hello");
EXPECT_EQ(Decode<std::string>(reader).Get(), "world");
EXPECT_FALSE(Decode<std::string>(reader));
@@ -109,11 +109,11 @@
TEST(BytesDecoderTest, ReflectedObject) {
auto data = Data(0x10, 0x20, 0x30, 0x40, 0x50, 0x60, 0x70, 0x80);
- auto reader = Reader{Slice{data}, 0, Endianness::kBig};
+ auto reader = BufferReader{Slice{data}};
auto got = Decode<S>(reader);
EXPECT_EQ(got->a, 0x10u);
- EXPECT_EQ(got->b, 0x2030u);
- EXPECT_EQ(got->c, 0x40506070u);
+ EXPECT_EQ(got->b, 0x3020u);
+ EXPECT_EQ(got->c, 0x70605040u);
EXPECT_FALSE(Decode<S>(reader));
}
@@ -124,13 +124,13 @@
0x00, 0x50, 0x06, 0x60, //
0x00, 0x70, 0x08, 0x80, //
0x01);
- auto reader = Reader{Slice{data}, 0, Endianness::kBig};
+ auto reader = BufferReader{Slice{data}};
auto got = Decode<M>(reader);
EXPECT_THAT(got.Get(), testing::ContainerEq(M{
- std::pair<uint8_t, uint32_t>(0x10u, 0x0220u),
- std::pair<uint8_t, uint32_t>(0x30u, 0x0440u),
- std::pair<uint8_t, uint32_t>(0x50u, 0x0660u),
- std::pair<uint8_t, uint32_t>(0x70u, 0x0880u),
+ std::pair<uint8_t, uint32_t>(0x10u, 0x2002u),
+ std::pair<uint8_t, uint32_t>(0x30u, 0x4004u),
+ std::pair<uint8_t, uint32_t>(0x50u, 0x6006u),
+ std::pair<uint8_t, uint32_t>(0x70u, 0x8008u),
}));
EXPECT_FALSE(Decode<M>(reader));
}
@@ -141,8 +141,8 @@
0x20, 0x30, //
0x40, 0x50, 0x60, 0x70, //
0x80);
- auto reader = Reader{Slice{data}, 0, Endianness::kBig};
- EXPECT_THAT(Decode<T>(reader).Get(), (T{0x10u, 0x2030u, 0x40506070u}));
+ auto reader = BufferReader{Slice{data}};
+ EXPECT_THAT(Decode<T>(reader).Get(), (T{0x10u, 0x3020u, 0x70605040u}));
EXPECT_FALSE(Decode<T>(reader));
}
diff --git a/src/tint/utils/bytes/reader.cc b/src/tint/utils/bytes/reader.cc
new file mode 100644
index 0000000..42e0176
--- /dev/null
+++ b/src/tint/utils/bytes/reader.cc
@@ -0,0 +1,35 @@
+// 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/utils/bytes/reader.h"
+
+namespace tint::bytes {
+
+Reader::~Reader() = default;
+BufferReader::~BufferReader() = default;
+
+} // namespace tint::bytes
diff --git a/src/tint/utils/bytes/reader.h b/src/tint/utils/bytes/reader.h
index 38bbf4e..9e58c15 100644
--- a/src/tint/utils/bytes/reader.h
+++ b/src/tint/utils/bytes/reader.h
@@ -36,104 +36,118 @@
#include "src/tint/utils/bytes/swap.h"
#include "src/tint/utils/containers/slice.h"
#include "src/tint/utils/reflection/reflection.h"
+#include "src/tint/utils/result/result.h"
namespace tint::bytes {
-/// A binary stream reader.
-struct Reader {
- /// @returns true if there are no more bytes remaining
- bool IsEOF() const { return offset >= bytes.len; }
+/// A binary stream reader interface
+class Reader {
+ public:
+ /// Read reads bytes from the stream, blocking until there are @p count bytes available, or the
+ /// end of the stream has been reached.
+ /// @param out a pointer to the byte buffer that will be filled with the read data. Must be at
+ /// least @p count size.
+ /// @param count the number of bytes to read. Must be greater than 0.
+ /// @returns the number of bytes read from the stream. If Read() returns less than @p count,
+ /// then the end of the stream has been reached.
+ virtual size_t Read(std::byte* out, size_t count) = 0;
- /// @returns the number of bytes remaining in the stream
- size_t BytesRemaining() const { return IsEOF() ? 0 : bytes.len - offset; }
+ // Destructor
+ virtual ~Reader();
/// Reads an integer from the stream, performing byte swapping if the stream's endianness
- /// differs from the native endianness. If there are too few bytes remaining in the stream, then
- /// the missing data will be substituted with zeros.
+ /// differs from the native endianness.
+ /// If there are too few bytes remaining in the stream, then a failure is returned.
+ /// @param endianness the encoded endianness of the integer
/// @return the deserialized integer
template <typename T>
- T Int() {
+ Result<T> Int(Endianness endianness = Endianness::kLittle) {
static_assert(std::is_integral_v<T>);
T out = 0;
- if (!IsEOF()) {
- size_t n = std::min(sizeof(T), BytesRemaining());
- memcpy(&out, &bytes[offset], n);
- offset += n;
- if (NativeEndianness() != endianness) {
- out = Swap(out);
- }
+ if (size_t n = Read(reinterpret_cast<std::byte*>(&out), sizeof(T)); n != sizeof(T)) {
+ return Failure{"EOF"};
+ }
+ if (NativeEndianness() != endianness) {
+ out = Swap(out);
}
return out;
}
- /// Reads a float from the stream. If there are too few bytes remaining in the stream, then
- /// the missing data will be substituted with zeros.
+ /// Reads a float from the stream.
+ /// If there are too few bytes remaining in the stream, then a failure is returned.
/// @return the deserialized floating point number
template <typename T>
- T Float() {
+ Result<T> Float() {
static_assert(std::is_floating_point_v<T>);
T out = 0;
- if (!IsEOF()) {
- size_t n = std::min(sizeof(T), BytesRemaining());
- memcpy(&out, &bytes[offset], n);
- offset += n;
+ if (size_t n = Read(reinterpret_cast<std::byte*>(&out), sizeof(T)); n != sizeof(T)) {
+ return Failure{"EOF"};
}
return out;
}
/// Reads a boolean from the stream
+ /// If there are too few bytes remaining in the stream, then a failure is returned.
/// @returns true if the next byte is non-zero
- bool Bool() {
- if (IsEOF()) {
- return false;
+ Result<bool> Bool() {
+ std::byte b{0};
+ if (size_t n = Read(&b, 1); n != 1) {
+ return Failure{"EOF"};
}
- return bytes[offset++] != std::byte{0};
+ return b != std::byte{0};
}
- /// Reads a string of @p len bytes from the stream. If there are too few bytes remaining in the
- /// stream, then the returned string will be truncated.
+ /// Reads a string of @p len bytes from the stream.
+ /// If there are too few bytes remaining in the stream, then a failure is returned.
/// @param len the length of the returned string in bytes
/// @return the deserialized string
- std::string String(size_t len) {
- if (IsEOF()) {
- return "";
+ Result<std::string> String(size_t len) {
+ std::string out;
+ out.resize(len);
+ if (size_t n = Read(reinterpret_cast<std::byte*>(out.data()), sizeof(char) * len);
+ n != len) {
+ return Failure{"EOF"};
}
- size_t n = std::min(len, BytesRemaining());
- std::string out(reinterpret_cast<const char*>(&bytes[offset]), n);
- offset += n;
return out;
}
-
- /// The data to read from
- Slice<const std::byte> bytes;
-
- /// The current byte offset
- size_t offset = 0;
-
- /// The endianness of integers serialized in the stream
- Endianness endianness = Endianness::kLittle;
};
-/// Reads the templated type from the reader and assigns it to @p out
-/// @note This function does not
-template <typename T>
-Reader& operator>>(Reader& reader, T& out) {
- constexpr bool is_numeric = std::is_integral_v<T> || std::is_floating_point_v<T>;
- static_assert(is_numeric);
+/// BufferReader is an implementation of the Reader interface backed by a buffer.
+class BufferReader final : public Reader {
+ public:
+ // Destructor
+ ~BufferReader() override;
- if constexpr (std::is_integral_v<T>) {
- out = reader.Int<T>();
- return reader;
+ /// Constructor
+ /// @param data the data to read from
+ /// @param size the number of bytes in the buffer
+ BufferReader(const std::byte* data, size_t size) : data_(data), bytes_remaining_(size) {
+ TINT_ASSERT(data);
}
- if constexpr (std::is_floating_point_v<T>) {
- out = reader.Float<T>();
- return reader;
+ /// Constructor
+ /// @param slice the byte slice to read from
+ explicit BufferReader(Slice<const std::byte> slice)
+ : data_(slice.data), bytes_remaining_(slice.len) {
+ TINT_ASSERT(slice.data);
}
- // Unreachable
- return reader;
-}
+ /// @copydoc Reader::Read
+ size_t Read(std::byte* out, size_t count) override {
+ size_t n = std::min(count, bytes_remaining_);
+ memcpy(out, data_, n);
+ data_ += n;
+ bytes_remaining_ -= n;
+ return n;
+ }
+
+ private:
+ /// The data to read from
+ const std::byte* data_ = nullptr;
+
+ /// The number of bytes remaining
+ size_t bytes_remaining_ = 0;
+};
} // namespace tint::bytes
diff --git a/src/tint/utils/bytes/reader_test.cc b/src/tint/utils/bytes/reader_test.cc
index e8af39d..6b76e22 100644
--- a/src/tint/utils/bytes/reader_test.cc
+++ b/src/tint/utils/bytes/reader_test.cc
@@ -37,70 +37,49 @@
return std::array{std::byte{static_cast<uint8_t>(args)}...};
}
-TEST(BytesReaderTest, IntegerBigEndian) {
+TEST(BufferReaderTest, IntegerBigEndian) {
auto data = Data(0x10, 0x20, 0x30, 0x40);
- auto u32 = Reader{Slice{data}, 0, Endianness::kBig}.Int<uint32_t>();
+ auto u32 = BufferReader{Slice{data}}.Int<uint32_t>(Endianness::kBig);
EXPECT_EQ(u32, 0x10203040u);
- auto i32 = Reader{Slice{data}, 0, Endianness::kBig}.Int<int32_t>();
+ auto i32 = BufferReader{Slice{data}}.Int<int32_t>(Endianness::kBig);
EXPECT_EQ(i32, 0x10203040);
}
-TEST(BytesReaderTest, IntegerBigEndian_Offset) {
- auto data = Data(0x10, 0x20, 0x30, 0x40, 0x50, 0x60);
- auto u32 = Reader{Slice{data}, 2, Endianness::kBig}.Int<uint32_t>();
- EXPECT_EQ(u32, 0x30405060u);
- auto i32 = Reader{Slice{data}, 2, Endianness::kBig}.Int<int32_t>();
- EXPECT_EQ(i32, 0x30405060);
+TEST(BufferReaderTest, IntegerBigEndian_TooShort) {
+ auto data = Data(0x10, 0x20);
+ auto u32 = BufferReader{Slice{data}}.Int<uint32_t>(Endianness::kBig);
+ EXPECT_FALSE(u32);
+ auto i32 = BufferReader{Slice{data}}.Int<int32_t>(Endianness::kBig);
+ EXPECT_FALSE(i32);
}
-TEST(BytesReaderTest, IntegerBigEndian_Clipped) {
+TEST(BufferReaderTest, IntegerLittleEndian) {
auto data = Data(0x10, 0x20, 0x30, 0x40);
- auto u32 = Reader{Slice{data}, 2, Endianness::kBig}.Int<uint32_t>();
- EXPECT_EQ(u32, 0x30400000u);
- auto i32 = Reader{Slice{data}, 2, Endianness::kBig}.Int<int32_t>();
- EXPECT_EQ(i32, 0x30400000);
-}
-
-TEST(BytesReaderTest, IntegerLittleEndian) {
- auto data = Data(0x10, 0x20, 0x30, 0x40);
- auto u32 = Reader{Slice{data}, 0, Endianness::kLittle}.Int<uint32_t>();
+ auto u32 = BufferReader{Slice{data}}.Int<uint32_t>(Endianness::kLittle);
EXPECT_EQ(u32, 0x40302010u);
- auto i32 = Reader{Slice{data}, 0, Endianness::kLittle}.Int<int32_t>();
+ auto i32 = BufferReader{Slice{data}}.Int<int32_t>(Endianness::kLittle);
EXPECT_EQ(i32, 0x40302010);
}
-TEST(BytesReaderTest, IntegerLittleEndian_Offset) {
- auto data = Data(0x10, 0x20, 0x30, 0x40, 0x50, 0x60);
- auto u32 = Reader{Slice{data}, 2, Endianness::kLittle}.Int<uint32_t>();
- EXPECT_EQ(u32, 0x60504030u);
- auto i32 = Reader{Slice{data}, 2, Endianness::kLittle}.Int<int32_t>();
- EXPECT_EQ(i32, 0x60504030);
+TEST(BufferReaderTest, IntegerLittleEndian_TooShort) {
+ auto data = Data(0x30, 0x40);
+ auto u32 = BufferReader{Slice{data}}.Int<uint32_t>(Endianness::kLittle);
+ EXPECT_FALSE(u32);
+ auto i32 = BufferReader{Slice{data}}.Int<int32_t>(Endianness::kLittle);
+ EXPECT_FALSE(i32);
}
-TEST(BytesReaderTest, IntegerLittleEndian_Clipped) {
- auto data = Data(0x10, 0x20, 0x30, 0x40);
- auto u32 = Reader{Slice{data}, 2, Endianness::kLittle}.Int<uint32_t>();
- EXPECT_EQ(u32, 0x00004030u);
- auto i32 = Reader{Slice{data}, 2, Endianness::kLittle}.Int<int32_t>();
- EXPECT_EQ(i32, 0x00004030);
-}
-
-TEST(BytesReaderTest, Float) {
+TEST(BufferReaderTest, Float) {
auto data = Data(0x00, 0x00, 0x08, 0x41);
- float f32 = Reader{Slice{data}}.Float<float>();
- EXPECT_EQ(f32, 8.5f);
+ auto f32 = BufferReader{Slice{data}}.Float<float>();
+ ASSERT_TRUE(f32);
+ EXPECT_EQ(f32.Get(), 8.5f);
}
-TEST(BytesReaderTest, Float_Offset) {
- auto data = Data(0x00, 0x00, 0x08, 0x41, 0x80, 0x3e);
- float f32 = Reader{Slice{data}, 2}.Float<float>();
- EXPECT_EQ(f32, 0.25049614f);
-}
-
-TEST(BytesReaderTest, Float_Clipped) {
- auto data = Data(0x00, 0x00, 0x08, 0x41);
- float f32 = Reader{Slice{data}, 2}.Float<float>();
- EXPECT_EQ(f32, 2.3329e-41f);
+TEST(BufferReaderTest, Float_TooShort) {
+ auto data = Data(0x08, 0x41);
+ auto f32 = BufferReader{Slice{data}}.Float<float>();
+ EXPECT_FALSE(f32);
}
} // namespace
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
new file mode 100644
index 0000000..2d060d3
--- /dev/null
+++ b/third_party/protobuf.cmake
@@ -0,0 +1,176 @@
+# 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.
+
+set(protobuf_INSTALL OFF CACHE BOOL "Install protobuf binaries and files" FORCE)
+set(protobuf_BUILD_CONFORMANCE OFF CACHE BOOL "Build conformance tests" FORCE)
+set(protobuf_BUILD_EXAMPLES OFF CACHE BOOL "Build examples" FORCE)
+set(protobuf_BUILD_LIBPROTOC OFF CACHE BOOL "Build libprotoc" FORCE)
+set(protobuf_BUILD_TESTS OFF CACHE BOOL "Controls whether protobuf tests are built" FORCE)
+set(protobuf_MSVC_STATIC_RUNTIME OFF CACHE BOOL "Controls whether a protobuf static runtime is built" FORCE)
+
+set(protobuf_BUILD_PROTOC_BINARIES ON CACHE BOOL "Build libprotoc and protoc compiler" FORCE)
+set(protobuf_DISABLE_RTTI ON CACHE BOOL "Remove runtime type information in the binaries" FORCE)
+
+add_subdirectory("${TINT_PROTOBUF_DIR}/cmake")
+
+target_compile_definitions(libprotobuf PUBLIC "-DGOOGLE_PROTOBUF_INTERNAL_DONATE_STEAL_INLINE=0")
+
+# A simplified version of protobuf_generate()
+function(generate_protos)
+ set(OPTIONS APPEND_PATH)
+ set(SINGLE_ARGS TARGET LANGUAGE EXPORT_MACRO PROTOC_OUT_DIR PLUGIN PLUGIN_OPTIONS)
+ set(MULTI_ARGS IMPORT_DIRS GENERATE_EXTENSIONS PROTOC_OPTIONS)
+ cmake_parse_arguments(ARGS "${OPTIONS}" "${SINGLE_ARGS}" "${MULTI_ARGS}" "${ARGN}")
+
+ if(NOT ARGS_TARGET)
+ message(FATAL_ERROR "generate_protos called without a target")
+ endif()
+
+ if(NOT ARGS_LANGUAGE)
+ set(ARGS_LANGUAGE cpp)
+ endif()
+ string(TOLOWER ${ARGS_LANGUAGE} ARGS_LANGUAGE)
+
+ if(NOT ARGS_PROTOC_OUT_DIR)
+ set(ARGS_PROTOC_OUT_DIR ${CMAKE_CURRENT_BINARY_DIR})
+ endif()
+
+ foreach(OPTION ${ARGS_PLUGIN_OPTIONS})
+ # append comma - not using CMake lists and string replacement as users
+ # might have semicolons in options
+ if(PLUGIN_OPTIONS)
+ set( PLUGIN_OPTIONS "${PLUGIN_OPTIONS},")
+ endif()
+ set(PLUGIN_OPTIONS "${PLUGIN_OPTIONS}${OPTION}")
+ endforeach()
+
+ if(ARGS_PLUGIN)
+ set(_plugin "--plugin=${ARGS_PLUGIN}")
+ endif()
+
+ if(NOT ARGS_GENERATE_EXTENSIONS)
+ if(ARGS_LANGUAGE STREQUAL cpp)
+ set(ARGS_GENERATE_EXTENSIONS .pb.h .pb.cc)
+ elseif(ARGS_LANGUAGE STREQUAL python)
+ set(ARGS_GENERATE_EXTENSIONS _pb2.py)
+ else()
+ message(FATAL_ERROR "generate_protos given unknown Language ${LANGUAGE}, please provide a value for GENERATE_EXTENSIONS")
+ endif()
+ endif()
+
+ if(ARGS_TARGET)
+ get_target_property(SOURCE_LIST ${ARGS_TARGET} SOURCES)
+ foreach(FILE ${SOURCE_LIST})
+ if(FILE MATCHES ".proto$")
+ list(APPEND PROTO_FILES ${FILE})
+ endif()
+ endforeach()
+ endif()
+
+ if(NOT PROTO_FILES)
+ message(FATAL_ERROR "generate_protos could not find any .proto files")
+ endif()
+
+ if(ARGS_APPEND_PATH)
+ # Create an include path for each file specified
+ foreach(FILE ${PROTO_FILES})
+ get_filename_component(ABS_FILE ${FILE} ABSOLUTE)
+ get_filename_component(ABS_PATH ${ABS_FILE} PATH)
+ list(FIND PROTOBUF_INCLUDE_PATH ${ABS_PATH} FOUND)
+ if(${FOUND} EQUAL -1)
+ list(APPEND PROTOBUF_INCLUDE_PATH -I ${ABS_PATH})
+ endif()
+ endforeach()
+ endif()
+
+ foreach(DIR ${ARGS_IMPORT_DIRS})
+ get_filename_component(ABS_PATH ${DIR} ABSOLUTE)
+ list(FIND PROTOBUF_INCLUDE_PATH ${ABS_PATH} FOUND)
+ if(${FOUND} EQUAL -1)
+ list(APPEND PROTOBUF_INCLUDE_PATH -I ${ABS_PATH})
+ endif()
+ endforeach()
+
+ if(NOT PROTOBUF_INCLUDE_PATH)
+ set(PROTOBUF_INCLUDE_PATH -I ${CMAKE_CURRENT_SOURCE_DIR})
+ endif()
+
+ set(ALL_GENERATED_SRCS)
+ foreach(PROTO_FILE ${PROTO_FILES})
+ get_filename_component(ABS_FILE ${PROTO_FILE} ABSOLUTE)
+ get_filename_component(ABS_DIR ${ABS_FILE} DIRECTORY)
+
+ get_filename_component(FILE_FULL_NAME ${PROTO_FILE} NAME)
+ string(FIND "${FILE_FULL_NAME}" "." FILE_LAST_EXT_POS REVERSE)
+ string(SUBSTRING "${FILE_FULL_NAME}" 0 ${FILE_LAST_EXT_POS} BASENAME)
+
+ set(SUITABLE_INCLUDE_FOUND FALSE)
+ foreach(DIR ${PROTOBUF_INCLUDE_PATH})
+ if(NOT DIR STREQUAL "-I")
+ file(RELATIVE_PATH REL_DIR ${DIR} ${ABS_DIR})
+ string(FIND "${REL_DIR}" "../" IS_IN_PARENT_FOLDER)
+ if (NOT ${IS_IN_PARENT_FOLDER} EQUAL 0)
+ set(SUITABLE_INCLUDE_FOUND TRUE)
+ break()
+ endif()
+ endif()
+ endforeach()
+
+ if(NOT SUITABLE_INCLUDE_FOUND)
+ message(FATAL_ERROR "generate_protos could not find any correct proto include directory.")
+ endif()
+
+ set(GENERATED_SRCS)
+ foreach(EXT ${ARGS_GENERATE_EXTENSIONS})
+ list(APPEND GENERATED_SRCS "${ARGS_PROTOC_OUT_DIR}/${REL_DIR}/${BASENAME}${EXT}")
+ endforeach()
+ list(APPEND ALL_GENERATED_SRCS ${GENERATED_SRCS})
+
+ set(COMMENT "Running ${ARGS_LANGUAGE} protocol buffer compiler on ${PROTO_FILE}")
+ if(ARGS_PROTOC_OPTIONS)
+ set(COMMENT "${COMMENT}, protoc-options: ${ARGS_PROTOC_OPTIONS}")
+ endif()
+ if(PLUGIN_OPTIONS)
+ 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
+ ARGS ${ARGS_PROTOC_OPTIONS} --${ARGS_LANGUAGE}_out ${_plugin_options}:${ARGS_PROTOC_OUT_DIR} ${_plugin} ${PROTOBUF_INCLUDE_PATH} ${ABS_FILE}
+ DEPENDS ${ABS_FILE} protobuf::protoc
+ COMMENT ${COMMENT}
+ VERBATIM)
+ endforeach()
+
+ set_source_files_properties(${ALL_GENERATED_SRCS} PROPERTIES GENERATED TRUE)
+ if(ARGS_TARGET)
+ target_sources(${ARGS_TARGET} PRIVATE ${ALL_GENERATED_SRCS})
+ endif()
+endfunction()
