[ir][msl] Add msl def file
Add a msl.def file and put the `threadgroup_barrier` method into it.
This is what the various barriers translate into but also carries the
barrier type required by MSL.
Bug: tint:1967
Change-Id: I6d822df111d57f64d7ac9023249cbe971c97ac6d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/163340
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/cmd/test/BUILD.bazel b/src/tint/cmd/test/BUILD.bazel
index 6f1e558..302b464 100644
--- a/src/tint/cmd/test/BUILD.bazel
+++ b/src/tint/cmd/test/BUILD.bazel
@@ -50,6 +50,7 @@
"//src/tint/lang/core/ir:test",
"//src/tint/lang/core/type:test",
"//src/tint/lang/core:test",
+ "//src/tint/lang/msl/ir:test",
"//src/tint/lang/spirv/ir:test",
"//src/tint/lang/wgsl/ast:test",
"//src/tint/lang/wgsl/features:test",
@@ -107,6 +108,7 @@
"//src/tint/lang/msl/writer/ast_printer:test",
"//src/tint/lang/msl/writer/common:test",
"//src/tint/lang/msl/writer/printer:test",
+ "//src/tint/lang/msl/writer/raise:test",
],
"//conditions:default": [],
}) + select({
diff --git a/src/tint/cmd/test/BUILD.cmake b/src/tint/cmd/test/BUILD.cmake
index 670c21f..a231ddd 100644
--- a/src/tint/cmd/test/BUILD.cmake
+++ b/src/tint/cmd/test/BUILD.cmake
@@ -51,6 +51,7 @@
tint_lang_core_ir_test
tint_lang_core_type_test
tint_lang_core_test
+ tint_lang_msl_ir_test
tint_lang_spirv_ir_test
tint_lang_wgsl_ast_test
tint_lang_wgsl_features_test
@@ -117,6 +118,7 @@
tint_lang_msl_writer_ast_printer_test
tint_lang_msl_writer_common_test
tint_lang_msl_writer_printer_test
+ tint_lang_msl_writer_raise_test
)
endif(TINT_BUILD_MSL_WRITER)
diff --git a/src/tint/cmd/test/BUILD.gn b/src/tint/cmd/test/BUILD.gn
index 17410ae..4fcb365 100644
--- a/src/tint/cmd/test/BUILD.gn
+++ b/src/tint/cmd/test/BUILD.gn
@@ -56,6 +56,7 @@
"${tint_src_dir}/lang/core/ir:unittests",
"${tint_src_dir}/lang/core/ir/transform:unittests",
"${tint_src_dir}/lang/core/type:unittests",
+ "${tint_src_dir}/lang/msl/ir:unittests",
"${tint_src_dir}/lang/spirv/ir:unittests",
"${tint_src_dir}/lang/wgsl:unittests",
"${tint_src_dir}/lang/wgsl/ast:unittests",
@@ -114,6 +115,7 @@
"${tint_src_dir}/lang/msl/writer/ast_printer:unittests",
"${tint_src_dir}/lang/msl/writer/common:unittests",
"${tint_src_dir}/lang/msl/writer/printer:unittests",
+ "${tint_src_dir}/lang/msl/writer/raise:unittests",
]
}
diff --git a/src/tint/lang/msl/BUILD.bazel b/src/tint/lang/msl/BUILD.bazel
index 1a114c6..af698ff 100644
--- a/src/tint/lang/msl/BUILD.bazel
+++ b/src/tint/lang/msl/BUILD.bazel
@@ -36,4 +36,19 @@
load("//src/tint:flags.bzl", "COPTS")
load("@bazel_skylib//lib:selects.bzl", "selects")
+cc_library(
+ name = "msl",
+ srcs = [
+ "builtin_fn.cc",
+ ],
+ hdrs = [
+ "barrier_type.h",
+ "builtin_fn.h",
+ ],
+ deps = [
+ "//src/tint/utils/traits",
+ ],
+ copts = COPTS,
+ visibility = ["//visibility:public"],
+)
diff --git a/src/tint/lang/msl/BUILD.cmake b/src/tint/lang/msl/BUILD.cmake
index bbcbda1..40cbc8e 100644
--- a/src/tint/lang/msl/BUILD.cmake
+++ b/src/tint/lang/msl/BUILD.cmake
@@ -34,5 +34,21 @@
# Do not modify this file directly
################################################################################
+include(lang/msl/intrinsic/BUILD.cmake)
+include(lang/msl/ir/BUILD.cmake)
include(lang/msl/validate/BUILD.cmake)
include(lang/msl/writer/BUILD.cmake)
+
+################################################################################
+# Target: tint_lang_msl
+# Kind: lib
+################################################################################
+tint_add_target(tint_lang_msl lib
+ lang/msl/barrier_type.h
+ lang/msl/builtin_fn.cc
+ lang/msl/builtin_fn.h
+)
+
+tint_target_add_dependencies(tint_lang_msl lib
+ tint_utils_traits
+)
diff --git a/src/tint/lang/msl/BUILD.gn b/src/tint/lang/msl/BUILD.gn
index 6caa255..a092431 100644
--- a/src/tint/lang/msl/BUILD.gn
+++ b/src/tint/lang/msl/BUILD.gn
@@ -37,3 +37,12 @@
import("../../../../scripts/tint_overrides_with_defaults.gni")
import("${tint_src_dir}/tint.gni")
+
+libtint_source_set("msl") {
+ sources = [
+ "barrier_type.h",
+ "builtin_fn.cc",
+ "builtin_fn.h",
+ ]
+ deps = [ "${tint_src_dir}/utils/traits" ]
+}
diff --git a/src/tint/lang/msl/barrier_type.h b/src/tint/lang/msl/barrier_type.h
new file mode 100644
index 0000000..c3c13b4
--- /dev/null
+++ b/src/tint/lang/msl/barrier_type.h
@@ -0,0 +1,49 @@
+// 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_MSL_BARRIER_TYPE_H_
+#define SRC_TINT_LANG_MSL_BARRIER_TYPE_H_
+
+#include <cstdint>
+
+namespace tint::msl {
+
+/// A barrier type
+enum BarrierType : uint8_t {
+ /// No type set
+ kNone = 0x0,
+ /// Device type
+ kDevice = 0x1,
+ /// Texture type
+ kTexture = 0x2,
+ /// Thead group type
+ kThreadGroup = 0x4,
+};
+
+} // namespace tint::msl
+
+#endif // SRC_TINT_LANG_MSL_BARRIER_TYPE_H_
diff --git a/src/tint/lang/msl/builtin_fn.cc b/src/tint/lang/msl/builtin_fn.cc
new file mode 100644
index 0000000..18d96de
--- /dev/null
+++ b/src/tint/lang/msl/builtin_fn.cc
@@ -0,0 +1,51 @@
+// 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:
+// src/tint/lang/msl/builtin_fn.cc.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+#include "src/tint/lang/msl/builtin_fn.h"
+
+namespace tint::msl {
+
+const char* str(BuiltinFn i) {
+ switch (i) {
+ case BuiltinFn::kNone:
+ return "<none>";
+ case BuiltinFn::kThreadgroupBarrier:
+ return "threadgroup_barrier";
+ }
+ return "<unknown>";
+}
+
+} // namespace tint::msl
diff --git a/src/tint/lang/msl/builtin_fn.cc.tmpl b/src/tint/lang/msl/builtin_fn.cc.tmpl
new file mode 100644
index 0000000..139045f
--- /dev/null
+++ b/src/tint/lang/msl/builtin_fn.cc.tmpl
@@ -0,0 +1,31 @@
+{{- /*
+--------------------------------------------------------------------------------
+Template file for use with tools/src/cmd/gen to generate builtin_fn.cc
+
+To update the generated file, run:
+ ./tools/run gen
+
+See:
+* tools/src/cmd/gen for structures used by this template
+* https://golang.org/pkg/text/template/ for documentation on the template syntax
+--------------------------------------------------------------------------------
+*/ -}}
+
+{{- $I := LoadIntrinsics "src/tint/lang/msl/msl.def" -}}
+#include "src/tint/lang/msl/builtin_fn.h"
+
+namespace tint::msl {
+
+const char* str(BuiltinFn i) {
+ switch (i) {
+ case BuiltinFn::kNone:
+ return "<none>";
+{{- range $I.Sem.Builtins }}
+ case BuiltinFn::k{{PascalCase .Name}}:
+ return "{{.Name}}";
+{{- end }}
+ }
+ return "<unknown>";
+}
+
+} // namespace tint::msl
diff --git a/src/tint/lang/msl/builtin_fn.h b/src/tint/lang/msl/builtin_fn.h
new file mode 100644
index 0000000..c652488
--- /dev/null
+++ b/src/tint/lang/msl/builtin_fn.h
@@ -0,0 +1,66 @@
+// 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:
+// src/tint/lang/msl/builtin_fn.h.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+#ifndef SRC_TINT_LANG_MSL_BUILTIN_FN_H_
+#define SRC_TINT_LANG_MSL_BUILTIN_FN_H_
+
+#include <cstdint>
+#include <string>
+
+#include "src/tint/utils/traits/traits.h"
+
+// \cond DO_NOT_DOCUMENT
+namespace tint::msl {
+
+/// Enumerator of all builtin functions
+enum class BuiltinFn : uint8_t {
+ kThreadgroupBarrier,
+ kNone,
+};
+
+/// @returns the name of the builtin function type.
+const char* str(BuiltinFn i);
+
+/// Emits the name of the builtin function type.
+template <typename STREAM, typename = traits::EnableIfIsOStream<STREAM>>
+auto& operator<<(STREAM& o, BuiltinFn i) {
+ return o << str(i);
+}
+
+} // namespace tint::msl
+// \endcond
+
+#endif // SRC_TINT_LANG_MSL_BUILTIN_FN_H_
diff --git a/src/tint/lang/msl/builtin_fn.h.tmpl b/src/tint/lang/msl/builtin_fn.h.tmpl
new file mode 100644
index 0000000..491524f
--- /dev/null
+++ b/src/tint/lang/msl/builtin_fn.h.tmpl
@@ -0,0 +1,47 @@
+{{- /*
+--------------------------------------------------------------------------------
+Template file for use with tools/src/cmd/gen to generate builtin_fn.h
+
+To update the generated file, run:
+ ./tools/run gen
+
+See:
+* tools/src/cmd/gen for structures used by this template
+* https://golang.org/pkg/text/template/ for documentation on the template syntax
+--------------------------------------------------------------------------------
+*/ -}}
+
+{{- $I := LoadIntrinsics "src/tint/lang/msl/msl.def" -}}
+
+#ifndef SRC_TINT_LANG_MSL_BUILTIN_FN_H_
+#define SRC_TINT_LANG_MSL_BUILTIN_FN_H_
+
+#include <cstdint>
+#include <string>
+
+#include "src/tint/utils/traits/traits.h"
+
+// \cond DO_NOT_DOCUMENT
+namespace tint::msl {
+
+/// Enumerator of all builtin functions
+enum class BuiltinFn : uint8_t {
+{{- range $I.Sem.Builtins }}
+ k{{PascalCase .Name}},
+{{- end }}
+ kNone,
+};
+
+/// @returns the name of the builtin function type.
+const char* str(BuiltinFn i);
+
+/// Emits the name of the builtin function type.
+template <typename STREAM, typename = traits::EnableIfIsOStream<STREAM>>
+auto& operator<<(STREAM& o, BuiltinFn i) {
+ return o << str(i);
+}
+
+} // namespace tint::msl
+// \endcond
+
+#endif // SRC_TINT_LANG_MSL_BUILTIN_FN_H_
diff --git a/src/tint/lang/msl/intrinsic/BUILD.bazel b/src/tint/lang/msl/intrinsic/BUILD.bazel
new file mode 100644
index 0000000..de60b66
--- /dev/null
+++ b/src/tint/lang/msl/intrinsic/BUILD.bazel
@@ -0,0 +1,69 @@
+# 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 = "intrinsic",
+ srcs = [
+ "data.cc",
+ ],
+ hdrs = [
+ "dialect.h",
+ ],
+ deps = [
+ "//src/tint/lang/core",
+ "//src/tint/lang/core/constant",
+ "//src/tint/lang/core/intrinsic",
+ "//src/tint/lang/core/type",
+ "//src/tint/lang/msl",
+ "//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/result",
+ "//src/tint/utils/rtti",
+ "//src/tint/utils/symbol",
+ "//src/tint/utils/text",
+ "//src/tint/utils/traits",
+ ],
+ copts = COPTS,
+ visibility = ["//visibility:public"],
+)
+
diff --git a/src/tint/lang/msl/intrinsic/BUILD.cmake b/src/tint/lang/msl/intrinsic/BUILD.cmake
new file mode 100644
index 0000000..60cca31
--- /dev/null
+++ b/src/tint/lang/msl/intrinsic/BUILD.cmake
@@ -0,0 +1,64 @@
+# 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
+################################################################################
+
+################################################################################
+# Target: tint_lang_msl_intrinsic
+# Kind: lib
+################################################################################
+tint_add_target(tint_lang_msl_intrinsic lib
+ lang/msl/intrinsic/data.cc
+ lang/msl/intrinsic/dialect.h
+)
+
+tint_target_add_dependencies(tint_lang_msl_intrinsic lib
+ tint_lang_core
+ tint_lang_core_constant
+ tint_lang_core_intrinsic
+ tint_lang_core_type
+ tint_lang_msl
+ tint_utils_containers
+ tint_utils_diagnostic
+ tint_utils_ice
+ tint_utils_id
+ tint_utils_macros
+ tint_utils_math
+ tint_utils_memory
+ tint_utils_result
+ tint_utils_rtti
+ tint_utils_symbol
+ tint_utils_text
+ tint_utils_traits
+)
diff --git a/src/tint/lang/msl/intrinsic/BUILD.gn b/src/tint/lang/msl/intrinsic/BUILD.gn
new file mode 100644
index 0000000..61d0771
--- /dev/null
+++ b/src/tint/lang/msl/intrinsic/BUILD.gn
@@ -0,0 +1,65 @@
+# 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")
+
+libtint_source_set("intrinsic") {
+ sources = [
+ "data.cc",
+ "dialect.h",
+ ]
+ deps = [
+ "${tint_src_dir}/lang/core",
+ "${tint_src_dir}/lang/core/constant",
+ "${tint_src_dir}/lang/core/intrinsic",
+ "${tint_src_dir}/lang/core/type",
+ "${tint_src_dir}/lang/msl",
+ "${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/result",
+ "${tint_src_dir}/utils/rtti",
+ "${tint_src_dir}/utils/symbol",
+ "${tint_src_dir}/utils/text",
+ "${tint_src_dir}/utils/traits",
+ ]
+}
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
new file mode 100644
index 0000000..484b078
--- /dev/null
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -0,0 +1,188 @@
+// 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:
+// src/tint/lang/msl/intrinsic/data.cc.tmpl
+//
+// To regenerate run: './tools/run gen'
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+#include <limits>
+#include <string>
+
+#include "src/tint/lang/core/intrinsic/type_matchers.h"
+#include "src/tint/lang/msl/intrinsic/dialect.h"
+#include "src/tint/utils/text/string_stream.h"
+
+namespace tint::msl::intrinsic {
+
+using namespace tint::core::intrinsic; // NOLINT(build/namespaces)
+
+namespace {
+
+using ConstEvalFunctionIndex = tint::core::intrinsic::ConstEvalFunctionIndex;
+using IntrinsicInfo = tint::core::intrinsic::IntrinsicInfo;
+using MatchState = tint::core::intrinsic::MatchState;
+using Number = tint::core::intrinsic::Number;
+using NumberMatcher = tint::core::intrinsic::NumberMatcher;
+using NumberMatcherIndex = tint::core::intrinsic::NumberMatcherIndex;
+using NumberMatcherIndicesIndex = tint::core::intrinsic::NumberMatcherIndicesIndex;
+using OverloadFlag = tint::core::intrinsic::OverloadFlag;
+using OverloadFlags = tint::core::intrinsic::OverloadFlags;
+using OverloadIndex = tint::core::intrinsic::OverloadIndex;
+using OverloadInfo = tint::core::intrinsic::OverloadInfo;
+using ParameterIndex = tint::core::intrinsic::ParameterIndex;
+using ParameterInfo = tint::core::intrinsic::ParameterInfo;
+using StringStream = tint::StringStream;
+using TemplateNumberIndex = tint::core::intrinsic::TemplateNumberIndex;
+using TemplateNumberInfo = tint::core::intrinsic::TemplateNumberInfo;
+using TemplateTypeIndex = tint::core::intrinsic::TemplateTypeIndex;
+using TemplateTypeInfo = tint::core::intrinsic::TemplateTypeInfo;
+using Type = tint::core::type::Type;
+using TypeMatcher = tint::core::intrinsic::TypeMatcher;
+using TypeMatcherIndex = tint::core::intrinsic::TypeMatcherIndex;
+using TypeMatcherIndicesIndex = tint::core::intrinsic::TypeMatcherIndicesIndex;
+
+template <size_t N>
+using TemplateNumberMatcher = tint::core::intrinsic::TemplateNumberMatcher<N>;
+
+template <size_t N>
+using TemplateTypeMatcher = tint::core::intrinsic::TemplateTypeMatcher<N>;
+
+// clang-format off
+
+/// TypeMatcher for 'type u32'
+constexpr TypeMatcher kU32Matcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+ if (!MatchU32(state, ty)) {
+ return nullptr;
+ }
+ return BuildU32(state, ty);
+ },
+/* string */ [](MatchState*) -> std::string {
+ return "u32";
+ }
+};
+
+
+/// Type and number matchers
+
+/// The template types, types, and type matchers
+constexpr TypeMatcher kTypeMatchers[] = {
+ /* [0] */ kU32Matcher,
+};
+
+constexpr TypeMatcherIndex kTypeMatcherIndices[] = {
+ /* [0] */ TypeMatcherIndex(0),
+};
+
+static_assert(TypeMatcherIndex::CanIndex(kTypeMatcherIndices),
+ "TypeMatcherIndex is not large enough to index kTypeMatcherIndices");
+
+constexpr ParameterInfo kParameters[] = {
+ {
+ /* [0] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* type_matcher_indices */ TypeMatcherIndicesIndex(0),
+ /* number_matcher_indices */ NumberMatcherIndicesIndex(/* invalid */),
+ },
+};
+
+static_assert(ParameterIndex::CanIndex(kParameters),
+ "ParameterIndex is not large enough to index kParameters");
+
+constexpr OverloadInfo kOverloads[] = {
+ {
+ /* [0] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_template_types */ 0,
+ /* num_template_numbers */ 0,
+ /* template_types */ TemplateTypeIndex(/* invalid */),
+ /* template_numbers */ TemplateNumberIndex(/* invalid */),
+ /* parameters */ ParameterIndex(0),
+ /* return_type_matcher_indices */ TypeMatcherIndicesIndex(/* invalid */),
+ /* return_number_matcher_indices */ NumberMatcherIndicesIndex(/* invalid */),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+};
+
+static_assert(OverloadIndex::CanIndex(kOverloads),
+ "OverloadIndex is not large enough to index kOverloads");
+
+constexpr IntrinsicInfo kBuiltins[] = {
+ {
+ /* [0] */
+ /* fn threadgroup_barrier(u32) */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(0),
+ },
+};
+
+// clang-format on
+
+} // anonymous namespace
+
+const core::intrinsic::TableData Dialect::kData{
+ /* template_types */ Empty,
+ /* template_numbers */ Empty,
+ /* type_matcher_indices */ kTypeMatcherIndices,
+ /* number_matcher_indices */ Empty,
+ /* type_matchers */ kTypeMatchers,
+ /* number_matchers */ Empty,
+ /* parameters */ kParameters,
+ /* overloads */ kOverloads,
+ /* const_eval_functions */ Empty,
+ /* ctor_conv */ Empty,
+ /* builtins */ kBuiltins,
+ /* binary '+' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '-' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '*' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '/' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '%' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '^' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '&' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '|' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '&&' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '||' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '==' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '!=' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '<' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '>' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '<=' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '>=' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '<<' */ tint::core::intrinsic::kNoOverloads,
+ /* binary '>>' */ tint::core::intrinsic::kNoOverloads,
+ /* unary '!' */ tint::core::intrinsic::kNoOverloads,
+ /* unary '~' */ tint::core::intrinsic::kNoOverloads,
+ /* unary '-' */ tint::core::intrinsic::kNoOverloads,
+};
+
+} // namespace tint::msl::intrinsic
diff --git a/src/tint/lang/msl/intrinsic/data.cc.tmpl b/src/tint/lang/msl/intrinsic/data.cc.tmpl
new file mode 100644
index 0000000..2314b5a
--- /dev/null
+++ b/src/tint/lang/msl/intrinsic/data.cc.tmpl
@@ -0,0 +1,34 @@
+{{- /*
+--------------------------------------------------------------------------------
+Template file for use with tools/src/cmd/gen to generate intrinsic_table.inl
+Used by BuiltinTable.cc for builtin overload resolution.
+
+To update the generated file, run:
+ ./tools/run gen
+
+See:
+* tools/src/cmd/gen for structures used by this template
+* https://golang.org/pkg/text/template/ for documentation on the template syntax
+--------------------------------------------------------------------------------
+*/ -}}
+
+{{- Import "src/tint/utils/templates/intrinsic_table_data.tmpl.inc" -}}
+
+{{- $I := LoadIntrinsics "src/tint/lang/msl/msl.def" -}}
+
+#include <limits>
+#include <string>
+
+#include "src/tint/lang/core/intrinsic/type_matchers.h"
+#include "src/tint/lang/msl/intrinsic/dialect.h"
+#include "src/tint/utils/text/string_stream.h"
+
+namespace tint::msl::intrinsic {
+
+using namespace tint::core::intrinsic; // NOLINT(build/namespaces)
+
+{{ Eval "Data"
+ "Intrinsics" $I
+ "Name" "Dialect::kData" -}}
+
+} // namespace tint::msl::intrinsic
diff --git a/src/tint/lang/msl/intrinsic/dialect.h b/src/tint/lang/msl/intrinsic/dialect.h
new file mode 100644
index 0000000..700df35
--- /dev/null
+++ b/src/tint/lang/msl/intrinsic/dialect.h
@@ -0,0 +1,51 @@
+// 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_MSL_INTRINSIC_DIALECT_H_
+#define SRC_TINT_LANG_MSL_INTRINSIC_DIALECT_H_
+
+#include "src/tint/lang/core/intrinsic/table_data.h"
+#include "src/tint/lang/msl/builtin_fn.h"
+
+namespace tint::msl::intrinsic {
+
+/// Dialect holds the intrinsic table data and types for the MSL dialect
+struct Dialect {
+ /// The dialect's intrinsic table data
+ static const core::intrinsic::TableData kData;
+
+ /// The dialect's builtin function enumerator
+ using BuiltinFn = msl::BuiltinFn;
+
+ /// @returns the name of the builtin function @p fn
+ /// @param fn the builtin function
+ static std::string_view ToString(BuiltinFn fn) { return str(fn); }
+};
+
+} // namespace tint::msl::intrinsic
+
+#endif // SRC_TINT_LANG_MSL_INTRINSIC_DIALECT_H_
diff --git a/src/tint/lang/msl/ir/BUILD.bazel b/src/tint/lang/msl/ir/BUILD.bazel
new file mode 100644
index 0000000..cff7707
--- /dev/null
+++ b/src/tint/lang/msl/ir/BUILD.bazel
@@ -0,0 +1,108 @@
+# 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 = "ir",
+ srcs = [
+ "builtin_call.cc",
+ ],
+ hdrs = [
+ "builtin_call.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/lang/msl",
+ "//src/tint/lang/msl/intrinsic",
+ "//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",
+ ],
+ copts = COPTS,
+ visibility = ["//visibility:public"],
+)
+cc_library(
+ name = "test",
+ alwayslink = True,
+ srcs = [
+ "builtin_call_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/lang/msl",
+ "//src/tint/lang/msl/intrinsic",
+ "//src/tint/lang/msl/ir",
+ "//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",
+ ],
+ copts = COPTS,
+ visibility = ["//visibility:public"],
+)
+
diff --git a/src/tint/lang/msl/ir/BUILD.cmake b/src/tint/lang/msl/ir/BUILD.cmake
new file mode 100644
index 0000000..1f4769d
--- /dev/null
+++ b/src/tint/lang/msl/ir/BUILD.cmake
@@ -0,0 +1,106 @@
+# 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
+################################################################################
+
+################################################################################
+# Target: tint_lang_msl_ir
+# Kind: lib
+################################################################################
+tint_add_target(tint_lang_msl_ir lib
+ lang/msl/ir/builtin_call.cc
+ lang/msl/ir/builtin_call.h
+)
+
+tint_target_add_dependencies(tint_lang_msl_ir lib
+ tint_api_common
+ tint_lang_core
+ tint_lang_core_constant
+ tint_lang_core_intrinsic
+ tint_lang_core_ir
+ tint_lang_core_type
+ tint_lang_msl
+ tint_lang_msl_intrinsic
+ 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
+)
+
+################################################################################
+# Target: tint_lang_msl_ir_test
+# Kind: test
+################################################################################
+tint_add_target(tint_lang_msl_ir_test test
+ lang/msl/ir/builtin_call_test.cc
+)
+
+tint_target_add_dependencies(tint_lang_msl_ir_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_lang_msl
+ tint_lang_msl_intrinsic
+ tint_lang_msl_ir
+ 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_msl_ir_test test
+ "gtest"
+)
diff --git a/src/tint/lang/msl/ir/BUILD.gn b/src/tint/lang/msl/ir/BUILD.gn
new file mode 100644
index 0000000..a072d1c
--- /dev/null
+++ b/src/tint/lang/msl/ir/BUILD.gn
@@ -0,0 +1,104 @@
+# 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")
+}
+
+libtint_source_set("ir") {
+ sources = [
+ "builtin_call.cc",
+ "builtin_call.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}/lang/msl",
+ "${tint_src_dir}/lang/msl/intrinsic",
+ "${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_unittests) {
+ tint_unittests_source_set("unittests") {
+ sources = [ "builtin_call_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}/lang/msl",
+ "${tint_src_dir}/lang/msl/intrinsic",
+ "${tint_src_dir}/lang/msl/ir",
+ "${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",
+ ]
+ }
+}
diff --git a/src/tint/lang/msl/ir/builtin_call.cc b/src/tint/lang/msl/ir/builtin_call.cc
new file mode 100644
index 0000000..8e20da3
--- /dev/null
+++ b/src/tint/lang/msl/ir/builtin_call.cc
@@ -0,0 +1,56 @@
+// Copyright 2023 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/msl/ir/builtin_call.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/utils/ice/ice.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::msl::ir::BuiltinCall);
+
+namespace tint::msl::ir {
+
+BuiltinCall::BuiltinCall(core::ir::InstructionResult* result,
+ BuiltinFn func,
+ VectorRef<core::ir::Value*> arguments)
+ : Base(result, arguments), func_(func) {
+ flags_.Add(Flag::kSequenced);
+ TINT_ASSERT(func != BuiltinFn::kNone);
+}
+
+BuiltinCall::~BuiltinCall() = default;
+
+BuiltinCall* BuiltinCall::Clone(core::ir::CloneContext& ctx) {
+ auto* new_result = ctx.Clone(Result(0));
+ auto new_args = ctx.Clone<BuiltinCall::kDefaultNumOperands>(Args());
+ return ctx.ir.instructions.Create<BuiltinCall>(new_result, func_, new_args);
+}
+
+} // namespace tint::msl::ir
diff --git a/src/tint/lang/msl/ir/builtin_call.h b/src/tint/lang/msl/ir/builtin_call.h
new file mode 100644
index 0000000..de63d4d
--- /dev/null
+++ b/src/tint/lang/msl/ir/builtin_call.h
@@ -0,0 +1,76 @@
+// 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_MSL_IR_BUILTIN_CALL_H_
+#define SRC_TINT_LANG_MSL_IR_BUILTIN_CALL_H_
+
+#include <string>
+
+#include "src/tint/lang/core/intrinsic/table_data.h"
+#include "src/tint/lang/core/ir/builtin_call.h"
+#include "src/tint/lang/msl/builtin_fn.h"
+#include "src/tint/lang/msl/intrinsic/dialect.h"
+#include "src/tint/utils/rtti/castable.h"
+
+namespace tint::msl::ir {
+
+/// A msl builtin call instruction in the IR.
+class BuiltinCall final : public Castable<BuiltinCall, core::ir::BuiltinCall> {
+ public:
+ /// Constructor
+ /// @param result the result value
+ /// @param func the builtin function
+ /// @param args the conversion arguments
+ BuiltinCall(core::ir::InstructionResult* result,
+ BuiltinFn func,
+ VectorRef<core::ir::Value*> args = tint::Empty);
+ ~BuiltinCall() override;
+
+ /// @copydoc core::ir::Instruction::Clone()
+ BuiltinCall* Clone(core::ir::CloneContext& ctx) override;
+
+ /// @returns the builtin function
+ BuiltinFn Func() const { return func_; }
+
+ /// @returns the identifier for the function
+ size_t FuncId() const override { return static_cast<size_t>(func_); }
+
+ /// @returns the friendly name for the instruction
+ std::string FriendlyName() const override { return std::string("msl.") + str(func_); }
+
+ /// @returns the table data to validate this builtin
+ const core::intrinsic::TableData& TableData() const override {
+ return msl::intrinsic::Dialect::kData;
+ }
+
+ private:
+ BuiltinFn func_;
+};
+
+} // namespace tint::msl::ir
+
+#endif // SRC_TINT_LANG_MSL_IR_BUILTIN_CALL_H_
diff --git a/src/tint/lang/msl/ir/builtin_call_test.cc b/src/tint/lang/msl/ir/builtin_call_test.cc
new file mode 100644
index 0000000..cbe47ba
--- /dev/null
+++ b/src/tint/lang/msl/ir/builtin_call_test.cc
@@ -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.
+
+#include "src/tint/lang/msl/ir/builtin_call.h"
+#include "gmock/gmock.h"
+#include "gtest/gtest-spi.h"
+#include "src/tint/lang/core/ir/ir_helper_test.h"
+
+namespace tint::msl::ir {
+namespace {
+
+using namespace tint::core::number_suffixes; // NOLINT
+ //
+using IR_MslBuiltinCallTest = core::ir::IRTestHelper;
+
+TEST_F(IR_MslBuiltinCallTest, Clone) {
+ auto* builtin = b.Call<BuiltinCall>(mod.Types().void_(), BuiltinFn::kThreadgroupBarrier, 0_u);
+
+ auto* new_b = clone_ctx.Clone(builtin);
+
+ EXPECT_NE(builtin, new_b);
+ EXPECT_NE(builtin->Result(0), new_b->Result(0));
+ EXPECT_EQ(mod.Types().void_(), new_b->Result(0)->Type());
+
+ EXPECT_EQ(BuiltinFn::kThreadgroupBarrier, new_b->Func());
+
+ auto args = new_b->Args();
+ EXPECT_EQ(1u, args.Length());
+
+ auto* val0 = args[0]->As<core::ir::Constant>()->Value();
+ EXPECT_EQ(0_u, val0->As<core::constant::Scalar<core::u32>>()->ValueAs<core::u32>());
+}
+
+} // namespace
+} // namespace tint::msl::ir
diff --git a/src/tint/lang/msl/msl.def b/src/tint/lang/msl/msl.def
new file mode 100644
index 0000000..7709699
--- /dev/null
+++ b/src/tint/lang/msl/msl.def
@@ -0,0 +1,38 @@
+// 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.
+
+////////////////////////////////////////////////////////////////////////////////
+// MSL builtin definition file //
+////////////////////////////////////////////////////////////////////////////////
+
+type u32
+
+////////////////////////////////////////////////////////////////////////////////
+// Builtin Functions //
+////////////////////////////////////////////////////////////////////////////////
+@stage("compute") fn threadgroup_barrier(u32)
+
diff --git a/src/tint/lang/msl/writer/printer/BUILD.bazel b/src/tint/lang/msl/writer/printer/BUILD.bazel
index 7de0377..ef25cfe 100644
--- a/src/tint/lang/msl/writer/printer/BUILD.bazel
+++ b/src/tint/lang/msl/writer/printer/BUILD.bazel
@@ -51,6 +51,9 @@
"//src/tint/lang/core/intrinsic",
"//src/tint/lang/core/ir",
"//src/tint/lang/core/type",
+ "//src/tint/lang/msl",
+ "//src/tint/lang/msl/intrinsic",
+ "//src/tint/lang/msl/ir",
"//src/tint/utils/containers",
"//src/tint/utils/diagnostic",
"//src/tint/utils/generator",
diff --git a/src/tint/lang/msl/writer/printer/BUILD.cmake b/src/tint/lang/msl/writer/printer/BUILD.cmake
index 37d0fcb..b922ea6 100644
--- a/src/tint/lang/msl/writer/printer/BUILD.cmake
+++ b/src/tint/lang/msl/writer/printer/BUILD.cmake
@@ -52,6 +52,9 @@
tint_lang_core_intrinsic
tint_lang_core_ir
tint_lang_core_type
+ tint_lang_msl
+ tint_lang_msl_intrinsic
+ tint_lang_msl_ir
tint_utils_containers
tint_utils_diagnostic
tint_utils_generator
diff --git a/src/tint/lang/msl/writer/printer/BUILD.gn b/src/tint/lang/msl/writer/printer/BUILD.gn
index ce460cc..3db1a8f 100644
--- a/src/tint/lang/msl/writer/printer/BUILD.gn
+++ b/src/tint/lang/msl/writer/printer/BUILD.gn
@@ -54,6 +54,9 @@
"${tint_src_dir}/lang/core/intrinsic",
"${tint_src_dir}/lang/core/ir",
"${tint_src_dir}/lang/core/type",
+ "${tint_src_dir}/lang/msl",
+ "${tint_src_dir}/lang/msl/intrinsic",
+ "${tint_src_dir}/lang/msl/ir",
"${tint_src_dir}/utils/containers",
"${tint_src_dir}/utils/diagnostic",
"${tint_src_dir}/utils/generator",
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index 2ededba..24b4b29 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -83,6 +83,8 @@
#include "src/tint/lang/core/type/u32.h"
#include "src/tint/lang/core/type/vector.h"
#include "src/tint/lang/core/type/void.h"
+#include "src/tint/lang/msl/barrier_type.h"
+#include "src/tint/lang/msl/ir/builtin_call.h"
#include "src/tint/lang/msl/writer/common/printer_support.h"
#include "src/tint/utils/containers/map.h"
#include "src/tint/utils/generator/text_generator.h"
@@ -372,6 +374,7 @@
[&](const core::ir::Var* var) { out << NameOf(var->Result(0)); }, //
[&](const core::ir::Bitcast* b) { EmitBitcast(out, b); }, //
[&](const core::ir::Access* a) { EmitAccess(out, a); }, //
+ [&](const msl::ir::BuiltinCall* c) { EmitMslBuiltinCall(out, c); }, //
[&](const core::ir::CoreBuiltinCall* c) { EmitCoreBuiltinCall(out, c); }, //
[&](const core::ir::UserCall* c) { EmitUserCall(out, c); }, //
[&](const core::ir::LoadVectorElement* e) {
@@ -753,6 +756,37 @@
}
}
+ void EmitMslBuiltinCall(StringStream& out, const msl::ir::BuiltinCall* c) {
+ switch (c->Func()) {
+ case msl::BuiltinFn::kThreadgroupBarrier: {
+ auto flags = c->Args()[0]->As<core::ir::Constant>()->Value()->ValueAs<uint8_t>();
+ out << "threadgroup_barrier(";
+ bool emitted_flag = false;
+
+ auto emit = [&](BarrierType type, const std::string& name) {
+ if ((flags & type) != type) {
+ return;
+ }
+
+ if (emitted_flag) {
+ out << " | ";
+ }
+ emitted_flag = true;
+ out << "mem_flags::mem_" << name;
+ };
+ emit(BarrierType::kDevice, "device");
+ emit(BarrierType::kThreadGroup, "threadgroup");
+ emit(BarrierType::kTexture, "texture");
+
+ out << ")";
+ return;
+ }
+ default:
+ TINT_ICE() << "undefined MSL ir function";
+ return;
+ }
+ }
+
void EmitCoreBuiltinCall(StringStream& out, const core::ir::CoreBuiltinCall* c) {
EmitCoreBuiltinName(out, c->Func());
out << "(";
diff --git a/src/tint/lang/msl/writer/raise/BUILD.bazel b/src/tint/lang/msl/writer/raise/BUILD.bazel
index a192956..c906eaa 100644
--- a/src/tint/lang/msl/writer/raise/BUILD.bazel
+++ b/src/tint/lang/msl/writer/raise/BUILD.bazel
@@ -39,24 +39,36 @@
cc_library(
name = "raise",
srcs = [
+ "builtin_polyfill.cc",
"raise.cc",
],
hdrs = [
+ "builtin_polyfill.h",
"raise.h",
],
deps = [
"//src/tint/api/common",
"//src/tint/api/options",
+ "//src/tint/lang/core",
+ "//src/tint/lang/core/constant",
+ "//src/tint/lang/core/intrinsic",
+ "//src/tint/lang/core/ir",
"//src/tint/lang/core/ir/transform",
+ "//src/tint/lang/core/type",
+ "//src/tint/lang/msl",
+ "//src/tint/lang/msl/intrinsic",
+ "//src/tint/lang/msl/ir",
"//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({
@@ -68,6 +80,43 @@
copts = COPTS,
visibility = ["//visibility:public"],
)
+cc_library(
+ name = "test",
+ alwayslink = True,
+ srcs = [
+ "builtin_polyfill_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/transform: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_msl_writer": [
+ "//src/tint/lang/msl/writer/raise",
+ ],
+ "//conditions:default": [],
+ }),
+ copts = COPTS,
+ visibility = ["//visibility:public"],
+)
alias(
name = "tint_build_msl_writer",
diff --git a/src/tint/lang/msl/writer/raise/BUILD.cmake b/src/tint/lang/msl/writer/raise/BUILD.cmake
index 0e41a3c..06e15b9 100644
--- a/src/tint/lang/msl/writer/raise/BUILD.cmake
+++ b/src/tint/lang/msl/writer/raise/BUILD.cmake
@@ -41,6 +41,8 @@
# Condition: TINT_BUILD_MSL_WRITER
################################################################################
tint_add_target(tint_lang_msl_writer_raise lib
+ lang/msl/writer/raise/builtin_polyfill.cc
+ lang/msl/writer/raise/builtin_polyfill.h
lang/msl/writer/raise/raise.cc
lang/msl/writer/raise/raise.h
)
@@ -48,16 +50,26 @@
tint_target_add_dependencies(tint_lang_msl_writer_raise lib
tint_api_common
tint_api_options
+ tint_lang_core
+ tint_lang_core_constant
+ tint_lang_core_intrinsic
+ tint_lang_core_ir
tint_lang_core_ir_transform
+ tint_lang_core_type
+ tint_lang_msl
+ tint_lang_msl_intrinsic
+ tint_lang_msl_ir
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
)
@@ -68,4 +80,48 @@
)
endif(TINT_BUILD_MSL_WRITER)
+endif(TINT_BUILD_MSL_WRITER)
+if(TINT_BUILD_MSL_WRITER)
+################################################################################
+# Target: tint_lang_msl_writer_raise_test
+# Kind: test
+# Condition: TINT_BUILD_MSL_WRITER
+################################################################################
+tint_add_target(tint_lang_msl_writer_raise_test test
+ lang/msl/writer/raise/builtin_polyfill_test.cc
+)
+
+tint_target_add_dependencies(tint_lang_msl_writer_raise_test test
+ tint_api_common
+ tint_lang_core
+ tint_lang_core_constant
+ tint_lang_core_intrinsic
+ tint_lang_core_ir
+ tint_lang_core_ir_transform_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_msl_writer_raise_test test
+ "gtest"
+)
+
+if(TINT_BUILD_MSL_WRITER)
+ tint_target_add_dependencies(tint_lang_msl_writer_raise_test test
+ tint_lang_msl_writer_raise
+ )
+endif(TINT_BUILD_MSL_WRITER)
+
endif(TINT_BUILD_MSL_WRITER)
\ No newline at end of file
diff --git a/src/tint/lang/msl/writer/raise/BUILD.gn b/src/tint/lang/msl/writer/raise/BUILD.gn
index 74cc477..d8e6ad5 100644
--- a/src/tint/lang/msl/writer/raise/BUILD.gn
+++ b/src/tint/lang/msl/writer/raise/BUILD.gn
@@ -37,25 +37,41 @@
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_msl_writer) {
libtint_source_set("raise") {
sources = [
+ "builtin_polyfill.cc",
+ "builtin_polyfill.h",
"raise.cc",
"raise.h",
]
deps = [
"${tint_src_dir}/api/common",
"${tint_src_dir}/api/options",
+ "${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/transform",
+ "${tint_src_dir}/lang/core/type",
+ "${tint_src_dir}/lang/msl",
+ "${tint_src_dir}/lang/msl/intrinsic",
+ "${tint_src_dir}/lang/msl/ir",
"${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",
]
@@ -65,3 +81,37 @@
}
}
}
+if (tint_build_unittests) {
+ if (tint_build_msl_writer) {
+ tint_unittests_source_set("unittests") {
+ sources = [ "builtin_polyfill_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/transform: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_msl_writer) {
+ deps += [ "${tint_src_dir}/lang/msl/writer/raise" ]
+ }
+ }
+ }
+}
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
new file mode 100644
index 0000000..1905246
--- /dev/null
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -0,0 +1,157 @@
+// 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/msl/writer/raise/builtin_polyfill.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/fluent_types.h"
+#include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/constant.h"
+#include "src/tint/lang/core/ir/core_builtin_call.h"
+#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/msl/barrier_type.h"
+#include "src/tint/lang/msl/ir/builtin_call.h"
+
+namespace tint::msl::writer::raise {
+namespace {
+
+using namespace tint::core::fluent_types; // NOLINT
+
+/// PIMPL state for the transform.
+struct State {
+ /// The IR module.
+ core::ir::Module& ir;
+
+ /// The IR builder.
+ core::ir::Builder b{ir};
+
+ /// The type manager.
+ core::type::Manager& ty{ir.Types()};
+
+ /// Process the module.
+ void Process() {
+ // Find the builtins that need replacing.
+ Vector<core::ir::CoreBuiltinCall*, 4> worklist;
+ for (auto* inst : ir.instructions.Objects()) {
+ if (!inst->Alive()) {
+ continue;
+ }
+ if (auto* builtin = inst->As<core::ir::CoreBuiltinCall>()) {
+ switch (builtin->Func()) {
+ case core::BuiltinFn::kStorageBarrier:
+ case core::BuiltinFn::kWorkgroupBarrier:
+ case core::BuiltinFn::kTextureBarrier:
+ worklist.Push(builtin);
+ break;
+ default:
+ break;
+ }
+ }
+ }
+
+ // Replace the builtins that we found.
+ for (auto* builtin : worklist) {
+ core::ir::Value* replacement = nullptr;
+ switch (builtin->Func()) {
+ case core::BuiltinFn::kStorageBarrier:
+ replacement = StorageBarrier(builtin);
+ break;
+ case core::BuiltinFn::kWorkgroupBarrier:
+ replacement = WorkgroupBarrier(builtin);
+ break;
+ case core::BuiltinFn::kTextureBarrier:
+ replacement = TextureBarrier(builtin);
+ break;
+ default:
+ break;
+ }
+ TINT_ASSERT_OR_RETURN(replacement);
+
+ // Replace the old builtin result with the new value.
+ if (auto name = ir.NameOf(builtin->Result(0))) {
+ ir.SetName(replacement, name);
+ }
+ builtin->Result(0)->ReplaceAllUsesWith(replacement);
+ builtin->Destroy();
+ }
+ }
+
+ /// Handle a `workgroupBarrier()` builtin.
+ /// @param builtin the builtin call instruction
+ /// @returns the replacement value
+ core::ir::Value* WorkgroupBarrier(core::ir::CoreBuiltinCall* builtin) {
+ // Replace the builtin call with a call to the msl.threadgroup_barrier intrinsic.
+ auto args = Vector<core::ir::Value*, 4>{b.Constant(u32(BarrierType::kThreadGroup))};
+
+ auto* call = b.Call<msl::ir::BuiltinCall>(
+ builtin->Result(0)->Type(), msl::BuiltinFn::kThreadgroupBarrier, std::move(args));
+ call->InsertBefore(builtin);
+ return call->Result(0);
+ }
+
+ /// Handle a `storageBarrier()` builtin.
+ /// @param builtin the builtin call instruction
+ /// @returns the replacement value
+ core::ir::Value* StorageBarrier(core::ir::CoreBuiltinCall* builtin) {
+ // Replace the builtin call with a call to the msl.threadgroup_barrier intrinsic.
+ auto args = Vector<core::ir::Value*, 4>{b.Constant(u32(BarrierType::kDevice))};
+
+ auto* call = b.Call<msl::ir::BuiltinCall>(
+ builtin->Result(0)->Type(), msl::BuiltinFn::kThreadgroupBarrier, std::move(args));
+ call->InsertBefore(builtin);
+ return call->Result(0);
+ }
+
+ /// Handle a `textureBarrier()` builtin.
+ /// @param builtin the builtin call instruction
+ /// @returns the replacement value
+ core::ir::Value* TextureBarrier(core::ir::CoreBuiltinCall* builtin) {
+ // Replace the builtin call with a call to the msl.threadgroup_barrier intrinsic.
+ auto args = Vector<core::ir::Value*, 4>{b.Constant(u32(BarrierType::kTexture))};
+
+ auto* call = b.Call<msl::ir::BuiltinCall>(
+ builtin->Result(0)->Type(), msl::BuiltinFn::kThreadgroupBarrier, std::move(args));
+ call->InsertBefore(builtin);
+ return call->Result(0);
+ }
+};
+
+} // namespace
+
+Result<SuccessType> BuiltinPolyfill(core::ir::Module& ir) {
+ auto result = ValidateAndDumpIfNeeded(ir, "BuiltinPolyfill transform");
+ if (!result) {
+ return result.Failure();
+ }
+
+ State{ir}.Process();
+
+ return Success;
+}
+
+} // namespace tint::msl::writer::raise
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.h b/src/tint/lang/msl/writer/raise/builtin_polyfill.h
new file mode 100644
index 0000000..344ed14
--- /dev/null
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.h
@@ -0,0 +1,54 @@
+// 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_MSL_WRITER_RAISE_BUILTIN_POLYFILL_H_
+#define SRC_TINT_LANG_MSL_WRITER_RAISE_BUILTIN_POLYFILL_H_
+
+#include <string>
+
+#include "src/tint/lang/core/ir/constant.h"
+#include "src/tint/lang/core/type/type.h"
+#include "src/tint/utils/diagnostic/diagnostic.h"
+#include "src/tint/utils/result/result.h"
+
+// Forward declarations.
+namespace tint::core::ir {
+class Module;
+class Texture;
+} // namespace tint::core::ir
+
+namespace tint::msl::writer::raise {
+
+/// BuiltinPolyfill is a transform that replaces calls to builtins with polyfills and calls to
+/// MSL backend intrinsic functions.
+/// @param module the module to transform
+/// @returns success or failure
+Result<SuccessType> BuiltinPolyfill(core::ir::Module& module);
+
+} // namespace tint::msl::writer::raise
+
+#endif // SRC_TINT_LANG_MSL_WRITER_RAISE_BUILTIN_POLYFILL_H_
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
new file mode 100644
index 0000000..b61a9a0
--- /dev/null
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -0,0 +1,143 @@
+// 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/msl/writer/raise/builtin_polyfill.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/ir/transform/helper_test.h"
+#include "src/tint/lang/core/type/array.h"
+#include "src/tint/lang/core/type/atomic.h"
+#include "src/tint/lang/core/type/builtin_structs.h"
+#include "src/tint/lang/core/type/depth_texture.h"
+#include "src/tint/lang/core/type/multisampled_texture.h"
+#include "src/tint/lang/core/type/sampled_texture.h"
+#include "src/tint/lang/core/type/storage_texture.h"
+
+namespace tint::msl::writer::raise {
+namespace {
+
+using MslWriter_BuiltinPolyfillTest = core::ir::transform::TransformTest;
+
+TEST_F(MslWriter_BuiltinPolyfillTest, WorkgroupBarrier) {
+ auto* func = b.Function("foo", ty.void_());
+ func->SetStage(core::ir::Function::PipelineStage::kCompute);
+ b.Append(func->Block(), [&] {
+ b.Call(ty.void_(), core::BuiltinFn::kWorkgroupBarrier);
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%foo = @compute func():void -> %b1 {
+ %b1 = block {
+ %2:void = workgroupBarrier
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = @compute func():void -> %b1 {
+ %b1 = block {
+ %2:void = msl.threadgroup_barrier 4u
+ ret
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, StorageBarrier) {
+ auto* func = b.Function("foo", ty.void_());
+ func->SetStage(core::ir::Function::PipelineStage::kCompute);
+ b.Append(func->Block(), [&] {
+ b.Call(ty.void_(), core::BuiltinFn::kStorageBarrier);
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%foo = @compute func():void -> %b1 {
+ %b1 = block {
+ %2:void = storageBarrier
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = @compute func():void -> %b1 {
+ %b1 = block {
+ %2:void = msl.threadgroup_barrier 1u
+ ret
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, TextureBarrier) {
+ auto* func = b.Function("foo", ty.void_());
+ func->SetStage(core::ir::Function::PipelineStage::kCompute);
+ b.Append(func->Block(), [&] {
+ b.Call(ty.void_(), core::BuiltinFn::kTextureBarrier);
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%foo = @compute func():void -> %b1 {
+ %b1 = block {
+ %2:void = textureBarrier
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = @compute func():void -> %b1 {
+ %b1 = block {
+ %2:void = msl.threadgroup_barrier 2u
+ ret
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+} // namespace
+} // namespace tint::msl::writer::raise
diff --git a/src/tint/lang/msl/writer/raise/raise.cc b/src/tint/lang/msl/writer/raise/raise.cc
index da70216..4140f46 100644
--- a/src/tint/lang/msl/writer/raise/raise.cc
+++ b/src/tint/lang/msl/writer/raise/raise.cc
@@ -41,6 +41,7 @@
#include "src/tint/lang/core/ir/transform/vectorize_scalar_matrix_constructors.h"
#include "src/tint/lang/core/ir/transform/zero_init_workgroup_memory.h"
#include "src/tint/lang/msl/writer/common/option_helpers.h"
+#include "src/tint/lang/msl/writer/raise/builtin_polyfill.h"
namespace tint::msl::writer::raise {
@@ -96,13 +97,13 @@
// PreservePadding must come before DirectVariableAccess.
RUN_TRANSFORM(core::ir::transform::PreservePadding);
-
RUN_TRANSFORM(core::ir::transform::VectorizeScalarMatrixConstructors);
// DemoteToHelper must come before any transform that introduces non-core instructions.
RUN_TRANSFORM(core::ir::transform::DemoteToHelper);
RUN_TRANSFORM(core::ir::transform::ValueToLet);
+ RUN_TRANSFORM(BuiltinPolyfill);
return Success;
}
