Import Tint changes from Dawn
Changes:
- 26e8eb157779689ecbf762a26f1e52740b7dadc9 [tint] Add language features for status testing. by Corentin Wallez <cwallez@chromium.org>
- cf34b18f8f0359a43c7ce7fa6093205d405ee3fe [tint] Add lang/wgsl/features/status by Corentin Wallez <cwallez@chromium.org>
- 0ebbc5c2fe37c20591d2ca54ea3728d622dd7991 [ir][msl] Add msl def file by dan sinclair <dsinclair@chromium.org>
- 7c0a801574a89ad06cf99dc523c0eab722df8998 D3D11: Implement polyfill of DP4A by Jiawei Shao <jiawei.shao@intel.com>
GitOrigin-RevId: 26e8eb157779689ecbf762a26f1e52740b7dadc9
Change-Id: I2d4253d475fbcc5ecb446216b1f6e79c8da0f61f
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/163640
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/cmd/test/BUILD.bazel b/src/tint/cmd/test/BUILD.bazel
index 6f1e558..080bcb9 100644
--- a/src/tint/cmd/test/BUILD.bazel
+++ b/src/tint/cmd/test/BUILD.bazel
@@ -50,9 +50,9 @@
"//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",
"//src/tint/lang/wgsl/helpers:test",
"//src/tint/lang/wgsl/program:test",
"//src/tint/lang/wgsl/reader/lower:test",
@@ -107,6 +107,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..3709145 100644
--- a/src/tint/cmd/test/BUILD.cmake
+++ b/src/tint/cmd/test/BUILD.cmake
@@ -51,9 +51,9 @@
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
tint_lang_wgsl_helpers_test
tint_lang_wgsl_program_test
tint_lang_wgsl_reader_lower_test
@@ -117,6 +117,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..e307f93 100644
--- a/src/tint/cmd/test/BUILD.gn
+++ b/src/tint/cmd/test/BUILD.gn
@@ -56,10 +56,10 @@
"${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",
- "${tint_src_dir}/lang/wgsl/features:unittests",
"${tint_src_dir}/lang/wgsl/helpers:unittests",
"${tint_src_dir}/lang/wgsl/program:unittests",
"${tint_src_dir}/lang/wgsl/reader/lower:unittests",
@@ -114,6 +114,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/hlsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
index d66be8d..c668179 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
@@ -255,6 +255,7 @@
polyfills.reflect_vec2_f32 = options.polyfill_reflect_vec2_f32;
polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
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/hlsl/writer/common/options.h b/src/tint/lang/hlsl/writer/common/options.h
index 1ff17bb..44cb403 100644
--- a/src/tint/lang/hlsl/writer/common/options.h
+++ b/src/tint/lang/hlsl/writer/common/options.h
@@ -71,6 +71,9 @@
/// Set to `true` to generate polyfill for `reflect` builtin for vec2<f32>
bool polyfill_reflect_vec2_f32 = false;
+ /// Set to `true` to generate polyfill for `dot4I8Packed` and `dot4U8Packed` builtins
+ bool polyfill_dot_4x8_packed = false;
+
/// Options used to specify a mapping of binding points to indices into a UBO
/// from which to load buffer sizes.
ArrayLengthFromUniformOptions array_length_from_uniform = {};
@@ -102,6 +105,7 @@
disable_workgroup_init,
truncate_interstage_variables,
polyfill_reflect_vec2_f32,
+ polyfill_dot_4x8_packed,
array_length_from_uniform,
interstage_locations,
root_constant_binding_point,
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;
}
diff --git a/src/tint/lang/wgsl/BUILD.bazel b/src/tint/lang/wgsl/BUILD.bazel
index 6541b22..f30a202 100644
--- a/src/tint/lang/wgsl/BUILD.bazel
+++ b/src/tint/lang/wgsl/BUILD.bazel
@@ -71,6 +71,8 @@
"diagnostic_rule_test.cc",
"diagnostic_severity_test.cc",
"extension_test.cc",
+ "language_feature_status_test.cc",
+ "language_feature_test.cc",
"wgsl_test.cc",
] + select({
"//conditions:default": [],
diff --git a/src/tint/lang/wgsl/BUILD.cmake b/src/tint/lang/wgsl/BUILD.cmake
index 4e0a08b..17f4d73 100644
--- a/src/tint/lang/wgsl/BUILD.cmake
+++ b/src/tint/lang/wgsl/BUILD.cmake
@@ -82,6 +82,8 @@
lang/wgsl/diagnostic_rule_test.cc
lang/wgsl/diagnostic_severity_test.cc
lang/wgsl/extension_test.cc
+ lang/wgsl/language_feature_status_test.cc
+ lang/wgsl/language_feature_test.cc
lang/wgsl/wgsl_test.cc
)
diff --git a/src/tint/lang/wgsl/BUILD.gn b/src/tint/lang/wgsl/BUILD.gn
index 0b9801f..1b7590c 100644
--- a/src/tint/lang/wgsl/BUILD.gn
+++ b/src/tint/lang/wgsl/BUILD.gn
@@ -71,6 +71,8 @@
"diagnostic_rule_test.cc",
"diagnostic_severity_test.cc",
"extension_test.cc",
+ "language_feature_status_test.cc",
+ "language_feature_test.cc",
"wgsl_test.cc",
]
deps = [
diff --git a/src/tint/lang/wgsl/ast/transform/builtin_polyfill.cc b/src/tint/lang/wgsl/ast/transform/builtin_polyfill.cc
index 6e38d10..a222570 100644
--- a/src/tint/lang/wgsl/ast/transform/builtin_polyfill.cc
+++ b/src/tint/lang/wgsl/ast/transform/builtin_polyfill.cc
@@ -888,6 +888,58 @@
return name;
}
+ Symbol Dot4I8Packed() {
+ using vec4i = vec4<i32>;
+ using vec4u = vec4<u32>;
+
+ auto name = b.Symbols().New("tint_dot4_i8_packed");
+
+ auto body = tint::Vector{
+ // const n = vec4u(24, 16, 8, 0);
+ // let a_i8 = bitcast<vec4i>(vec4u(a) << n) >> vec4u(24);
+ // let b_i8 = bitcast<vec4i>(vec4u(b) << n) >> vec4u(24);
+ // return dot(a_i8, b_i8);
+ b.Decl(b.Const("n", b.Call<vec4u>(24_a, 16_a, 8_a, 0_a))),
+ b.Decl(b.Let("a_i8", b.Shr(b.Bitcast<vec4i>(b.Shl(b.Call<vec4u>("a"), "n")),
+ b.Call<vec4u>(24_a)))),
+ b.Decl(b.Let("b_i8", b.Shr(b.Bitcast<vec4i>(b.Shl(b.Call<vec4u>("b"), "n")),
+ b.Call<vec4u>(24_a)))),
+ b.Return(b.Call("dot", "a_i8", "b_i8")),
+ };
+ b.Func(name,
+ tint::Vector{
+ b.Param("a", b.ty.u32()),
+ b.Param("b", b.ty.u32()),
+ },
+ b.ty.i32(), body);
+
+ return name;
+ }
+
+ Symbol Dot4U8Packed() {
+ using vec4u = vec4<u32>;
+ auto name = b.Symbols().New("tint_dot4_u8_packed");
+
+ auto body = tint::Vector{
+ // const n = vec4u(24, 16, 8, 0);
+ // let a_u8 = (vec4u(a) >> n) & vec4u(0xff);
+ // let b_u8 = (vec4u(b) >> n) & vec4u(0xff);
+ // return dot(a_u8, b_u8);
+ b.Decl(b.Const("n", b.Call<vec4u>(24_a, 16_a, 8_a, 0_a))),
+ b.Decl(b.Let("a_u8", b.And(b.Shr(b.Call<vec4u>("a"), "n"), b.Call<vec4u>(0xff_a)))),
+ b.Decl(b.Let("b_u8", b.And(b.Shr(b.Call<vec4u>("b"), "n"), b.Call<vec4u>(0xff_a)))),
+ b.Return(b.Call("dot", "a_u8", "b_u8")),
+ };
+ b.Func(name,
+ tint::Vector{
+ b.Param("a", b.ty.u32()),
+ b.Param("b", b.ty.u32()),
+ },
+ b.ty.u32(), body);
+
+ return name;
+ }
+
////////////////////////////////////////////////////////////////////////////
// Inline polyfills
////////////////////////////////////////////////////////////////////////////
@@ -1270,6 +1322,22 @@
}
return Symbol{};
+ case wgsl::BuiltinFn::kDot4I8Packed: {
+ if (cfg.builtins.dot_4x8_packed) {
+ return builtin_polyfills.GetOrCreate(builtin,
+ [&] { return Dot4I8Packed(); });
+ }
+ return Symbol{};
+ }
+
+ case wgsl::BuiltinFn::kDot4U8Packed: {
+ if (cfg.builtins.dot_4x8_packed) {
+ return builtin_polyfills.GetOrCreate(builtin,
+ [&] { return Dot4U8Packed(); });
+ }
+ return Symbol{};
+ }
+
default:
return Symbol{};
}
diff --git a/src/tint/lang/wgsl/ast/transform/builtin_polyfill.h b/src/tint/lang/wgsl/ast/transform/builtin_polyfill.h
index 0252fce..86059b6 100644
--- a/src/tint/lang/wgsl/ast/transform/builtin_polyfill.h
+++ b/src/tint/lang/wgsl/ast/transform/builtin_polyfill.h
@@ -98,6 +98,8 @@
bool quantize_to_vec_f16 = false;
/// Should `workgroupUniformLoad()` be polyfilled?
bool workgroup_uniform_load = false;
+ /// Should `dot4I8Packed()` and `dot4U8Packed()` be polyfilled?
+ bool dot_4x8_packed = false;
};
/// Config is consumed by the BuiltinPolyfill transform.
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 e2d0dde..7045762 100644
--- a/src/tint/lang/wgsl/ast/transform/builtin_polyfill_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/builtin_polyfill_test.cc
@@ -4055,6 +4055,83 @@
}
////////////////////////////////////////////////////////////////////////////////
+// Built-in functions in packed_4x8_integer_dot_product
+////////////////////////////////////////////////////////////////////////////////
+DataMap polyfillPacked4x8IntegerDotProduct() {
+ BuiltinPolyfill::Builtins builtins;
+ builtins.dot_4x8_packed = true;
+ DataMap data;
+ data.Add<BuiltinPolyfill::Config>(builtins);
+ return data;
+}
+
+TEST_F(BuiltinPolyfillTest, Dot4I8Packed) {
+ auto* src = R"(
+enable chromium_experimental_dp4a;
+
+fn f() {
+ let v1 = 0x01020304u;
+ let v2 = 0xF1F2F3F4u;
+ _ = dot4I8Packed(v1, v2);
+}
+)";
+
+ 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));
+ let b_i8 = (bitcast<vec4<i32>>((vec4<u32>(b) << n)) >> vec4<u32>(24));
+ return dot(a_i8, b_i8);
+}
+
+fn f() {
+ let v1 = 16909060u;
+ let v2 = 4059231220u;
+ _ = tint_dot4_i8_packed(v1, v2);
+}
+)";
+
+ auto got = Run<BuiltinPolyfill>(src, polyfillPacked4x8IntegerDotProduct());
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, Dot4U8Packed) {
+ auto* src = R"(
+enable chromium_experimental_dp4a;
+
+fn f() {
+ let v1 = 0x01020304u;
+ let v2 = 0xF1F2F3F4u;
+ _ = dot4U8Packed(v1, v2);
+}
+)";
+
+ 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));
+ let b_u8 = ((vec4<u32>(b) >> n) & vec4<u32>(255));
+ return dot(a_u8, b_u8);
+}
+
+fn f() {
+ let v1 = 16909060u;
+ let v2 = 4059231220u;
+ _ = tint_dot4_u8_packed(v1, v2);
+}
+)";
+
+ auto got = Run<BuiltinPolyfill>(src, polyfillPacked4x8IntegerDotProduct());
+
+ EXPECT_EQ(expect, str(got));
+}
+
+////////////////////////////////////////////////////////////////////////////////
// Polyfill combinations
////////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/lang/wgsl/features/BUILD.bazel b/src/tint/lang/wgsl/features/BUILD.bazel
index cec9088..879fbfd 100644
--- a/src/tint/lang/wgsl/features/BUILD.bazel
+++ b/src/tint/lang/wgsl/features/BUILD.bazel
@@ -40,34 +40,15 @@
name = "features",
srcs = [
"language_feature.cc",
+ "status.cc",
],
hdrs = [
"language_feature.h",
+ "status.h",
],
deps = [
],
copts = COPTS,
visibility = ["//visibility:public"],
)
-cc_library(
- name = "test",
- alwayslink = True,
- srcs = [
- "language_feature_test.cc",
- ],
- deps = [
- "//src/tint/lang/wgsl/features",
- "//src/tint/utils/containers",
- "//src/tint/utils/ice",
- "//src/tint/utils/macros",
- "//src/tint/utils/math",
- "//src/tint/utils/memory",
- "//src/tint/utils/rtti",
- "//src/tint/utils/text",
- "//src/tint/utils/traits",
- "@gtest",
- ],
- copts = COPTS,
- visibility = ["//visibility:public"],
-)
diff --git a/src/tint/lang/wgsl/features/BUILD.cmake b/src/tint/lang/wgsl/features/BUILD.cmake
index 9bbfe60..7746c88 100644
--- a/src/tint/lang/wgsl/features/BUILD.cmake
+++ b/src/tint/lang/wgsl/features/BUILD.cmake
@@ -41,28 +41,6 @@
tint_add_target(tint_lang_wgsl_features lib
lang/wgsl/features/language_feature.cc
lang/wgsl/features/language_feature.h
-)
-
-################################################################################
-# Target: tint_lang_wgsl_features_test
-# Kind: test
-################################################################################
-tint_add_target(tint_lang_wgsl_features_test test
- lang/wgsl/features/language_feature_test.cc
-)
-
-tint_target_add_dependencies(tint_lang_wgsl_features_test test
- tint_lang_wgsl_features
- tint_utils_containers
- tint_utils_ice
- tint_utils_macros
- tint_utils_math
- tint_utils_memory
- tint_utils_rtti
- tint_utils_text
- tint_utils_traits
-)
-
-tint_target_add_external_dependencies(tint_lang_wgsl_features_test test
- "gtest"
+ lang/wgsl/features/status.cc
+ lang/wgsl/features/status.h
)
diff --git a/src/tint/lang/wgsl/features/BUILD.gn b/src/tint/lang/wgsl/features/BUILD.gn
index db3557a..95dea3f 100644
--- a/src/tint/lang/wgsl/features/BUILD.gn
+++ b/src/tint/lang/wgsl/features/BUILD.gn
@@ -38,6 +38,8 @@
sources = [
"language_feature.cc",
"language_feature.h",
+ "status.cc",
+ "status.h",
]
deps = []
diff --git a/src/tint/lang/wgsl/features/language_feature.cc b/src/tint/lang/wgsl/features/language_feature.cc
index df6a24c..708697b 100644
--- a/src/tint/lang/wgsl/features/language_feature.cc
+++ b/src/tint/lang/wgsl/features/language_feature.cc
@@ -42,6 +42,21 @@
/// @param str the string to parse
/// @returns the parsed enum, or LanguageFeature::kUndefined if the string could not be parsed.
LanguageFeature ParseLanguageFeature(std::string_view str) {
+ if (str == "chromium_testing_experimental") {
+ return LanguageFeature::kChromiumTestingExperimental;
+ }
+ if (str == "chromium_testing_shipped") {
+ return LanguageFeature::kChromiumTestingShipped;
+ }
+ if (str == "chromium_testing_shipped_with_killswitch") {
+ return LanguageFeature::kChromiumTestingShippedWithKillswitch;
+ }
+ if (str == "chromium_testing_unimplemented") {
+ return LanguageFeature::kChromiumTestingUnimplemented;
+ }
+ if (str == "chromium_testing_unsafe_experimental") {
+ return LanguageFeature::kChromiumTestingUnsafeExperimental;
+ }
if (str == "readonly_and_readwrite_storage_textures") {
return LanguageFeature::kReadonlyAndReadwriteStorageTextures;
}
@@ -52,6 +67,16 @@
switch (value) {
case LanguageFeature::kUndefined:
return "undefined";
+ case LanguageFeature::kChromiumTestingExperimental:
+ return "chromium_testing_experimental";
+ case LanguageFeature::kChromiumTestingShipped:
+ return "chromium_testing_shipped";
+ case LanguageFeature::kChromiumTestingShippedWithKillswitch:
+ return "chromium_testing_shipped_with_killswitch";
+ case LanguageFeature::kChromiumTestingUnimplemented:
+ return "chromium_testing_unimplemented";
+ case LanguageFeature::kChromiumTestingUnsafeExperimental:
+ return "chromium_testing_unsafe_experimental";
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 1187bfa..4788b2d 100644
--- a/src/tint/lang/wgsl/features/language_feature.h
+++ b/src/tint/lang/wgsl/features/language_feature.h
@@ -46,6 +46,11 @@
/// @see src/tint/lang/wgsl/intrinsics.def for language feature descriptions
enum class LanguageFeature : uint8_t {
kUndefined,
+ kChromiumTestingExperimental,
+ kChromiumTestingShipped,
+ kChromiumTestingShippedWithKillswitch,
+ kChromiumTestingUnimplemented,
+ kChromiumTestingUnsafeExperimental,
kReadonlyAndReadwriteStorageTextures,
};
@@ -59,11 +64,21 @@
LanguageFeature ParseLanguageFeature(std::string_view str);
constexpr std::string_view kLanguageFeatureStrings[] = {
+ "chromium_testing_experimental",
+ "chromium_testing_shipped",
+ "chromium_testing_shipped_with_killswitch",
+ "chromium_testing_unimplemented",
+ "chromium_testing_unsafe_experimental",
"readonly_and_readwrite_storage_textures",
};
/// All features
static constexpr LanguageFeature kAllLanguageFeatures[] = {
+ LanguageFeature::kChromiumTestingExperimental,
+ LanguageFeature::kChromiumTestingShipped,
+ LanguageFeature::kChromiumTestingShippedWithKillswitch,
+ LanguageFeature::kChromiumTestingUnimplemented,
+ LanguageFeature::kChromiumTestingUnsafeExperimental,
LanguageFeature::kReadonlyAndReadwriteStorageTextures,
};
diff --git a/src/tint/lang/wgsl/features/status.cc b/src/tint/lang/wgsl/features/status.cc
new file mode 100644
index 0000000..1520bbf
--- /dev/null
+++ b/src/tint/lang/wgsl/features/status.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/wgsl/features/status.h"
+
+#include "src/tint/lang/wgsl/features/language_feature.h"
+
+namespace tint::wgsl {
+
+FeatureStatus GetLanguageFeatureStatus(LanguageFeature f) {
+ switch (f) {
+ case LanguageFeature::kReadonlyAndReadwriteStorageTextures:
+ return FeatureStatus::kExperimental;
+ case LanguageFeature::kUndefined:
+ return FeatureStatus::kUnknown;
+
+ case LanguageFeature::kChromiumTestingUnimplemented:
+ return FeatureStatus::kUnimplemented;
+ case LanguageFeature::kChromiumTestingUnsafeExperimental:
+ return FeatureStatus::kUnsafeExperimental;
+ case LanguageFeature::kChromiumTestingExperimental:
+ return FeatureStatus::kExperimental;
+ case LanguageFeature::kChromiumTestingShippedWithKillswitch:
+ return FeatureStatus::kShippedWithKillswitch;
+ case LanguageFeature::kChromiumTestingShipped:
+ return FeatureStatus::kShipped;
+ }
+
+ return FeatureStatus::kUnknown;
+}
+
+} // namespace tint::wgsl
diff --git a/src/tint/lang/wgsl/features/status.h b/src/tint/lang/wgsl/features/status.h
new file mode 100644
index 0000000..3cb578a
--- /dev/null
+++ b/src/tint/lang/wgsl/features/status.h
@@ -0,0 +1,63 @@
+// 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_WGSL_FEATURES_STATUS_H_
+#define SRC_TINT_LANG_WGSL_FEATURES_STATUS_H_
+
+#include <cstdint>
+
+namespace tint::wgsl {
+
+enum class LanguageFeature : uint8_t;
+
+/// The status of the implementation of a WGSL language feature so that other components (like Dawn)
+/// can query it. The enum values are in the order of least implemented to most implemented.
+enum class FeatureStatus : uint8_t {
+ // The feature is not known.
+ kUnknown,
+ // The feature is known in wgsl.def but not implemented at all.
+ kUnimplemented,
+ // The feature is at least partially implemented but might contain big security of correctness
+ // issues.
+ kUnsafeExperimental,
+ // The feature is implemented and should be safe from a security standpoint, but shouldn't be
+ // exposed by default.
+ kExperimental,
+ // The feature is implemented and can be exposed by default, but is only turned on if the
+ // feature is explicitly enabled in the wgsl reader options.
+ kShippedWithKillswitch,
+ // The feature is exposed by default and cannot be turned off.
+ kShipped,
+};
+
+/// @param f the feature to get the status of.
+/// @returns the status, or kUnknown if the feature is not known.
+FeatureStatus GetLanguageFeatureStatus(LanguageFeature f);
+
+} // namespace tint::wgsl
+
+#endif // SRC_TINT_LANG_WGSL_FEATURES_STATUS_H_
diff --git a/src/tint/lang/wgsl/language_feature_status_test.cc b/src/tint/lang/wgsl/language_feature_status_test.cc
new file mode 100644
index 0000000..8ebab9b
--- /dev/null
+++ b/src/tint/lang/wgsl/language_feature_status_test.cc
@@ -0,0 +1,58 @@
+// 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 <gtest/gtest.h>
+
+#include "src/tint/lang/wgsl/features/language_feature.h"
+#include "src/tint/lang/wgsl/features/status.h"
+
+namespace tint::wgsl {
+namespace {
+
+TEST(LanguageFeatureStatusTest, AllFeaturesAreKnown) {
+ for (auto f : kAllLanguageFeatures) {
+ EXPECT_NE(FeatureStatus::kUnknown, GetLanguageFeatureStatus(f));
+ }
+}
+
+TEST(LanguageFeatureStatusTest, ChromiumTestingValues) {
+ EXPECT_EQ(FeatureStatus::kUnimplemented,
+ GetLanguageFeatureStatus(tint::wgsl::LanguageFeature::kChromiumTestingUnimplemented));
+ EXPECT_EQ(
+ FeatureStatus::kUnsafeExperimental,
+ GetLanguageFeatureStatus(tint::wgsl::LanguageFeature::kChromiumTestingUnsafeExperimental));
+ EXPECT_EQ(FeatureStatus::kExperimental,
+ GetLanguageFeatureStatus(tint::wgsl::LanguageFeature::kChromiumTestingExperimental));
+ EXPECT_EQ(FeatureStatus::kShippedWithKillswitch,
+ GetLanguageFeatureStatus(
+ tint::wgsl::LanguageFeature::kChromiumTestingShippedWithKillswitch));
+ EXPECT_EQ(FeatureStatus::kShipped,
+ GetLanguageFeatureStatus(tint::wgsl::LanguageFeature::kChromiumTestingShipped));
+}
+
+} // namespace
+} // namespace tint::wgsl
diff --git a/src/tint/lang/wgsl/features/language_feature_test.cc b/src/tint/lang/wgsl/language_feature_test.cc
similarity index 65%
rename from src/tint/lang/wgsl/features/language_feature_test.cc
rename to src/tint/lang/wgsl/language_feature_test.cc
index 295e3a5..045f462 100644
--- a/src/tint/lang/wgsl/features/language_feature_test.cc
+++ b/src/tint/lang/wgsl/language_feature_test.cc
@@ -27,7 +27,7 @@
////////////////////////////////////////////////////////////////////////////////
// File generated by 'tools/src/cmd/gen' using the template:
-// src/tint/lang/wgsl/features/language_feature_test.cc.tmpl
+// src/tint/lang/wgsl/language_feature_test.cc.tmpl
//
// To regenerate run: './tools/run gen'
//
@@ -57,14 +57,35 @@
}
static constexpr Case kValidCases[] = {
+ {"chromium_testing_experimental", LanguageFeature::kChromiumTestingExperimental},
+ {"chromium_testing_shipped", LanguageFeature::kChromiumTestingShipped},
+ {"chromium_testing_shipped_with_killswitch",
+ LanguageFeature::kChromiumTestingShippedWithKillswitch},
+ {"chromium_testing_unimplemented", LanguageFeature::kChromiumTestingUnimplemented},
+ {"chromium_testing_unsafe_experimental", LanguageFeature::kChromiumTestingUnsafeExperimental},
{"readonly_and_readwrite_storage_textures",
LanguageFeature::kReadonlyAndReadwriteStorageTextures},
};
static constexpr Case kInvalidCases[] = {
- {"eadonly_and_readwrite_stccrage_textures", LanguageFeature::kUndefined},
- {"rladonly_a3readrite_storage_textures", LanguageFeature::kUndefined},
- {"readonly_and_readwriVe_storage_textures", LanguageFeature::kUndefined},
+ {"chccomium_esting_experimental", LanguageFeature::kUndefined},
+ {"cr3mium_testlg_experimentl", LanguageFeature::kUndefined},
+ {"cVromium_testing_experimental", LanguageFeature::kUndefined},
+ {"1hromium_testing_shipped", LanguageFeature::kUndefined},
+ {"chromium_testing_hJpqqed", LanguageFeature::kUndefined},
+ {"cllromium_te77ting_shipped", LanguageFeature::kUndefined},
+ {"chromippm_testing_sHHipped_with_kqqlswitch", LanguageFeature::kUndefined},
+ {"chromicm_esting_shippvd_with_kilsitch", LanguageFeature::kUndefined},
+ {"chrbmium_testing_shGpped_wih_killswitch", LanguageFeature::kUndefined},
+ {"chromium_testing_iinimplemevted", LanguageFeature::kUndefined},
+ {"chromiumWWtesting_unimp8emented", LanguageFeature::kUndefined},
+ {"chxxoium_tMsting_unimplemented", LanguageFeature::kUndefined},
+ {"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},
};
using LanguageFeatureParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/lang/wgsl/features/language_feature_test.cc.tmpl b/src/tint/lang/wgsl/language_feature_test.cc.tmpl
similarity index 100%
rename from src/tint/lang/wgsl/features/language_feature_test.cc.tmpl
rename to src/tint/lang/wgsl/language_feature_test.cc.tmpl
diff --git a/src/tint/lang/wgsl/wgsl.def b/src/tint/lang/wgsl/wgsl.def
index 23ba104..7969b9f 100644
--- a/src/tint/lang/wgsl/wgsl.def
+++ b/src/tint/lang/wgsl/wgsl.def
@@ -95,6 +95,13 @@
// https://gpuweb.github.io/gpuweb/wgsl/#language-extensions-sec
enum language_feature {
readonly_and_readwrite_storage_textures
+
+ // Language features used only for testing whose status will never change.
+ chromium_testing_unimplemented
+ chromium_testing_unsafe_experimental
+ chromium_testing_experimental
+ chromium_testing_shipped_with_killswitch
+ chromium_testing_shipped
}
////////////////////////////////////////////////////////////////////////////////
