[spirv-reader][ir] Handle GLSL 450 Normalize
The SPIR-V `Normalize` method allows scalar values. When converting to
WGSL, convert a scalar normalize with a WGSL `sign` call.
Bug: 42250952
Change-Id: Id41ffacaf726c64faad94d183e18bc85a38f16c2
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/220854
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/builtin_fn.cc b/src/tint/lang/spirv/builtin_fn.cc
index c2380c36..4ea7ab7 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -108,6 +108,8 @@
return "vector_times_matrix";
case BuiltinFn::kVectorTimesScalar:
return "vector_times_scalar";
+ case BuiltinFn::kNormalize:
+ return "normalize";
case BuiltinFn::kSdot:
return "sdot";
case BuiltinFn::kUdot:
@@ -161,6 +163,7 @@
case BuiltinFn::kSdot:
case BuiltinFn::kUdot:
case BuiltinFn::kNone:
+ case BuiltinFn::kNormalize:
break;
}
return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/spirv/builtin_fn.cc.tmpl b/src/tint/lang/spirv/builtin_fn.cc.tmpl
index a539fa5..cba1b82 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -72,6 +72,7 @@
case BuiltinFn::kSdot:
case BuiltinFn::kUdot:
case BuiltinFn::kNone:
+ case BuiltinFn::kNormalize:
break;
}
return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/spirv/builtin_fn.h b/src/tint/lang/spirv/builtin_fn.h
index 8e582ab..0e49a9b 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -81,6 +81,7 @@
kSelect,
kVectorTimesMatrix,
kVectorTimesScalar,
+ kNormalize,
kSdot,
kUdot,
kNone,
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index 5b40302..8e9b6d0 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -4881,6 +4881,28 @@
{
/* [149] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 1,
+ /* templates */ TemplateIndex(21),
+ /* parameters */ ParameterIndex(4),
+ /* return_matcher_indices */ MatcherIndicesIndex(3),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [150] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 2,
+ /* templates */ TemplateIndex(42),
+ /* parameters */ ParameterIndex(279),
+ /* return_matcher_indices */ MatcherIndicesIndex(25),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [151] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
@@ -4890,7 +4912,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [150] */
+ /* [152] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 4,
/* num_explicit_templates */ 0,
@@ -4901,7 +4923,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [151] */
+ /* [153] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 6,
/* num_explicit_templates */ 0,
@@ -4912,7 +4934,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [152] */
+ /* [154] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -4923,7 +4945,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [153] */
+ /* [155] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 4,
/* num_explicit_templates */ 0,
@@ -4934,7 +4956,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [154] */
+ /* [156] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -4945,7 +4967,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [155] */
+ /* [157] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -4956,7 +4978,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [156] */
+ /* [158] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -4967,7 +4989,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [157] */
+ /* [159] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -4978,7 +5000,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [158] */
+ /* [160] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -4989,7 +5011,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [159] */
+ /* [161] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5000,7 +5022,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [160] */
+ /* [162] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -5011,7 +5033,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [161] */
+ /* [163] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -5031,91 +5053,91 @@
/* [0] */
/* fn array_length[I : u32, A : access](ptr<storage, struct_with_runtime_array, A>, I) -> u32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(149),
+ /* overloads */ OverloadIndex(151),
},
{
/* [1] */
/* fn atomic_and[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(150),
+ /* overloads */ OverloadIndex(152),
},
{
/* [2] */
/* fn atomic_compare_exchange[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, U, T, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(151),
+ /* overloads */ OverloadIndex(153),
},
{
/* [3] */
/* fn atomic_exchange[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(150),
+ /* overloads */ OverloadIndex(152),
},
{
/* [4] */
/* fn atomic_iadd[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(150),
+ /* overloads */ OverloadIndex(152),
},
{
/* [5] */
/* fn atomic_isub[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(150),
+ /* overloads */ OverloadIndex(152),
},
{
/* [6] */
/* fn atomic_load[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(152),
+ /* overloads */ OverloadIndex(154),
},
{
/* [7] */
/* fn atomic_or[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(150),
+ /* overloads */ OverloadIndex(152),
},
{
/* [8] */
/* fn atomic_smax[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(150),
+ /* overloads */ OverloadIndex(152),
},
{
/* [9] */
/* fn atomic_smin[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(150),
+ /* overloads */ OverloadIndex(152),
},
{
/* [10] */
/* fn atomic_store[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(153),
+ /* overloads */ OverloadIndex(155),
},
{
/* [11] */
/* fn atomic_umax[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(150),
+ /* overloads */ OverloadIndex(152),
},
{
/* [12] */
/* fn atomic_umin[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(150),
+ /* overloads */ OverloadIndex(152),
},
{
/* [13] */
/* fn atomic_xor[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(150),
+ /* overloads */ OverloadIndex(152),
},
{
/* [14] */
/* fn dot[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(154),
+ /* overloads */ OverloadIndex(156),
},
{
/* [15] */
@@ -5313,19 +5335,19 @@
/* [26] */
/* fn matrix_times_matrix[T : f32_f16, K : num, C : num, R : num](mat<K, R, T>, mat<C, K, T>) -> mat<C, R, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(155),
+ /* overloads */ OverloadIndex(157),
},
{
/* [27] */
/* fn matrix_times_scalar[T : f32_f16, N : num, M : num](mat<N, M, T>, T) -> mat<N, M, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(156),
+ /* overloads */ OverloadIndex(158),
},
{
/* [28] */
/* fn matrix_times_vector[T : f32_f16, N : num, M : num](mat<N, M, T>, vec<N, T>) -> vec<M, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(157),
+ /* overloads */ OverloadIndex(159),
},
{
/* [29] */
@@ -5353,25 +5375,32 @@
/* [31] */
/* fn vector_times_matrix[T : f32_f16, N : num, M : num](vec<N, T>, mat<M, N, T>) -> vec<M, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(158),
+ /* overloads */ OverloadIndex(160),
},
{
/* [32] */
/* fn vector_times_scalar[T : f32_f16, N : num](vec<N, T>, T) -> vec<N, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(159),
+ /* overloads */ OverloadIndex(161),
},
{
/* [33] */
- /* fn sdot(u32, u32, u32) -> i32 */
- /* num overloads */ 1,
- /* overloads */ OverloadIndex(160),
+ /* fn normalize[T : f32_f16](T) -> T */
+ /* fn normalize[N : num, T : f32_f16](vec<N, T>) -> vec<N, T> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(149),
},
{
/* [34] */
+ /* fn sdot(u32, u32, u32) -> i32 */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(162),
+ },
+ {
+ /* [35] */
/* fn udot(u32, u32, u32) -> u32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(161),
+ /* overloads */ OverloadIndex(163),
},
};
diff --git a/src/tint/lang/spirv/reader/helper_test.h b/src/tint/lang/spirv/reader/helper_test.h
index 7bb17bf..0d14aca 100644
--- a/src/tint/lang/spirv/reader/helper_test.h
+++ b/src/tint/lang/spirv/reader/helper_test.h
@@ -28,14 +28,12 @@
#ifndef SRC_TINT_LANG_SPIRV_READER_HELPER_TEST_H_
#define SRC_TINT_LANG_SPIRV_READER_HELPER_TEST_H_
-#include <iostream>
#include <string>
-#include <vector>
#include "gmock/gmock.h"
#include "gtest/gtest.h"
+
#include "src/tint/lang/core/ir/disassembler.h"
-#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/validator.h"
#include "src/tint/lang/spirv/reader/common/helper_test.h"
#include "src/tint/lang/spirv/reader/reader.h"
diff --git a/src/tint/lang/spirv/reader/import_glsl_std450_test.cc b/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
index 0434bc1..47b240a 100644
--- a/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
+++ b/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
@@ -832,7 +832,7 @@
// So we have to test it separately, as it does not fit the patterns tested
// above.
-TEST_F(SpirvReaderTest, DISABLED_Normalize_Scalar) {
+TEST_F(SpirvReaderTest, Normalize_Scalar) {
// Scalar normalize maps to sign.
EXPECT_IR(Preamble() + R"(
%1 = OpExtInst %float %glsl Normalize %float_50
@@ -851,7 +851,7 @@
)");
}
-TEST_F(SpirvReaderTest, DISABLED_Normalize_Vector2) {
+TEST_F(SpirvReaderTest, Normalize_Vector2) {
EXPECT_IR(Preamble() + R"(
%1 = OpExtInst %v2float %glsl Normalize %v2float_50_60
%2 = OpCopyObject %v2float %1
@@ -869,7 +869,7 @@
)");
}
-TEST_F(SpirvReaderTest, DISABLED_Normalize_Vector3) {
+TEST_F(SpirvReaderTest, Normalize_Vector3) {
EXPECT_IR(Preamble() + R"(
%1 = OpExtInst %v3float %glsl Normalize %v3float_50_60_70
%2 = OpCopyObject %v3float %1
@@ -887,7 +887,7 @@
)");
}
-TEST_F(SpirvReaderTest, DISABLED_Normalize_Vector4) {
+TEST_F(SpirvReaderTest, Normalize_Vector4) {
EXPECT_IR(Preamble() + R"(
%1 = OpExtInst %v4float %glsl Normalize %v4float_50_50_50_50
%2 = OpCopyObject %v4float %1
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.bazel b/src/tint/lang/spirv/reader/lower/BUILD.bazel
index b29d625..eac8130 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/lower/BUILD.bazel
@@ -39,11 +39,13 @@
cc_library(
name = "lower",
srcs = [
+ "builtins.cc",
"lower.cc",
"shader_io.cc",
"vector_element_pointer.cc",
],
hdrs = [
+ "builtins.h",
"lower.h",
"shader_io.h",
"vector_element_pointer.h",
@@ -55,6 +57,9 @@
"//src/tint/lang/core/intrinsic",
"//src/tint/lang/core/ir",
"//src/tint/lang/core/type",
+ "//src/tint/lang/spirv",
+ "//src/tint/lang/spirv/intrinsic",
+ "//src/tint/lang/spirv/ir",
"//src/tint/utils",
"//src/tint/utils/containers",
"//src/tint/utils/diagnostic",
@@ -75,6 +80,7 @@
name = "test",
alwayslink = True,
srcs = [
+ "builtins_test.cc",
"shader_io_test.cc",
"vector_element_pointer_test.cc",
],
@@ -86,6 +92,9 @@
"//src/tint/lang/core/ir",
"//src/tint/lang/core/ir/transform:test",
"//src/tint/lang/core/type",
+ "//src/tint/lang/spirv",
+ "//src/tint/lang/spirv/intrinsic",
+ "//src/tint/lang/spirv/ir",
"//src/tint/lang/spirv/reader/lower",
"//src/tint/utils",
"//src/tint/utils/containers",
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.cmake b/src/tint/lang/spirv/reader/lower/BUILD.cmake
index 6869be8..e39b8bb 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/lower/BUILD.cmake
@@ -39,6 +39,8 @@
# Kind: lib
################################################################################
tint_add_target(tint_lang_spirv_reader_lower lib
+ lang/spirv/reader/lower/builtins.cc
+ lang/spirv/reader/lower/builtins.h
lang/spirv/reader/lower/lower.cc
lang/spirv/reader/lower/lower.h
lang/spirv/reader/lower/shader_io.cc
@@ -54,6 +56,9 @@
tint_lang_core_intrinsic
tint_lang_core_ir
tint_lang_core_type
+ tint_lang_spirv
+ tint_lang_spirv_intrinsic
+ tint_lang_spirv_ir
tint_utils
tint_utils_containers
tint_utils_diagnostic
@@ -76,6 +81,7 @@
# Kind: test
################################################################################
tint_add_target(tint_lang_spirv_reader_lower_test test
+ lang/spirv/reader/lower/builtins_test.cc
lang/spirv/reader/lower/shader_io_test.cc
lang/spirv/reader/lower/vector_element_pointer_test.cc
)
@@ -88,6 +94,9 @@
tint_lang_core_ir
tint_lang_core_ir_transform_test
tint_lang_core_type
+ tint_lang_spirv
+ tint_lang_spirv_intrinsic
+ tint_lang_spirv_ir
tint_lang_spirv_reader_lower
tint_utils
tint_utils_containers
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.gn b/src/tint/lang/spirv/reader/lower/BUILD.gn
index 6012672..37bc08d 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.gn
+++ b/src/tint/lang/spirv/reader/lower/BUILD.gn
@@ -45,6 +45,8 @@
libtint_source_set("lower") {
sources = [
+ "builtins.cc",
+ "builtins.h",
"lower.cc",
"lower.h",
"shader_io.cc",
@@ -60,6 +62,9 @@
"${tint_src_dir}/lang/core/intrinsic",
"${tint_src_dir}/lang/core/ir",
"${tint_src_dir}/lang/core/type",
+ "${tint_src_dir}/lang/spirv",
+ "${tint_src_dir}/lang/spirv/intrinsic",
+ "${tint_src_dir}/lang/spirv/ir",
"${tint_src_dir}/utils",
"${tint_src_dir}/utils/containers",
"${tint_src_dir}/utils/diagnostic",
@@ -76,6 +81,7 @@
if (tint_build_unittests) {
tint_unittests_source_set("unittests") {
sources = [
+ "builtins_test.cc",
"shader_io_test.cc",
"vector_element_pointer_test.cc",
]
@@ -89,6 +95,9 @@
"${tint_src_dir}/lang/core/ir",
"${tint_src_dir}/lang/core/ir/transform:unittests",
"${tint_src_dir}/lang/core/type",
+ "${tint_src_dir}/lang/spirv",
+ "${tint_src_dir}/lang/spirv/intrinsic",
+ "${tint_src_dir}/lang/spirv/ir",
"${tint_src_dir}/lang/spirv/reader/lower",
"${tint_src_dir}/utils",
"${tint_src_dir}/utils/containers",
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
new file mode 100644
index 0000000..c9edc17
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -0,0 +1,100 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/reader/lower/builtins.h"
+
+#include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/spirv/ir/builtin_call.h"
+
+namespace tint::spirv::reader::lower {
+
+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() {
+ Vector<spirv::ir::BuiltinCall*, 4> builtin_worklist;
+ for (auto* inst : ir.Instructions()) {
+ if (auto* builtin = inst->As<spirv::ir::BuiltinCall>()) {
+ builtin_worklist.Push(builtin);
+ }
+ }
+
+ // Replace the builtins that we found.
+ for (auto* builtin : builtin_worklist) {
+ switch (builtin->Func()) {
+ case spirv::BuiltinFn::kNormalize:
+ Normalize(builtin);
+ break;
+ default:
+ TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
+ }
+ }
+ }
+
+ void Normalize(spirv::ir::BuiltinCall* call) {
+ auto* arg = call->Args()[0];
+
+ b.InsertBefore(call, [&] {
+ core::BuiltinFn fn = core::BuiltinFn::kNormalize;
+ if (arg->Type()->IsScalar()) {
+ fn = core::BuiltinFn::kSign;
+ }
+ b.CallWithResult(call->DetachResult(), fn, Vector<core::ir::Value*, 1>{arg});
+ });
+ call->Destroy();
+ }
+};
+
+} // namespace
+
+Result<SuccessType> Builtins(core::ir::Module& ir) {
+ auto result = ValidateAndDumpIfNeeded(ir, "spirv.Builtins");
+ if (result != Success) {
+ return result.Failure();
+ }
+
+ State{ir}.Process();
+
+ return Success;
+}
+
+} // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/lower/builtins.h b/src/tint/lang/spirv/reader/lower/builtins.h
new file mode 100644
index 0000000..f3b8009
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/builtins.h
@@ -0,0 +1,48 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#ifndef SRC_TINT_LANG_SPIRV_READER_LOWER_BUILTINS_H_
+#define SRC_TINT_LANG_SPIRV_READER_LOWER_BUILTINS_H_
+
+#include "src/tint/utils/result/result.h"
+
+// Forward declarations.
+namespace tint::core::ir {
+class Module;
+}
+
+namespace tint::spirv::reader::lower {
+
+/// Builtins is a transform that converts SPIR-V builtin methods into the form expected by Tint's
+/// core IR.
+/// @param module the module to transform
+/// @returns success or failure
+Result<SuccessType> Builtins(core::ir::Module& module);
+
+} // namespace tint::spirv::reader::lower
+
+#endif // SRC_TINT_LANG_SPIRV_READER_LOWER_BUILTINS_H_
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
new file mode 100644
index 0000000..95c814e
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -0,0 +1,103 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/reader/lower/builtins.h"
+
+#include "src/tint/lang/core/ir/transform/helper_test.h"
+#include "src/tint/lang/spirv/ir/builtin_call.h"
+
+namespace tint::spirv::reader::lower {
+namespace {
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+using SpirvParser_BuiltinsTest = core::ir::transform::TransformTest;
+
+TEST_F(SpirvParser_BuiltinsTest, Normalize_Scalar) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.f32(), spirv::BuiltinFn::kNormalize, 10_f);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:f32 = spirv.normalize 10.0f
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:f32 = sign 10.0f
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, Normalize_Vector) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.vec2<f32>(), spirv::BuiltinFn::kNormalize,
+ b.Splat(ty.vec2<f32>(), 10_f));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<f32> = spirv.normalize vec2<f32>(10.0f)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<f32> = normalize vec2<f32>(10.0f)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+} // namespace
+} // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/lower/lower.cc b/src/tint/lang/spirv/reader/lower/lower.cc
index f1861d4..2eecf38 100644
--- a/src/tint/lang/spirv/reader/lower/lower.cc
+++ b/src/tint/lang/spirv/reader/lower/lower.cc
@@ -28,6 +28,7 @@
#include "src/tint/lang/spirv/reader/lower/lower.h"
#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/spirv/reader/lower/builtins.h"
#include "src/tint/lang/spirv/reader/lower/shader_io.h"
#include "src/tint/lang/spirv/reader/lower/vector_element_pointer.h"
@@ -44,6 +45,7 @@
RUN_TRANSFORM(lower::VectorElementPointer, mod);
RUN_TRANSFORM(lower::ShaderIO, mod);
+ RUN_TRANSFORM(lower::Builtins, mod);
if (auto res = core::ir::ValidateAndDumpIfNeeded(mod, "spirv.Lower"); res != Success) {
return res.Failure();
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.bazel b/src/tint/lang/spirv/reader/parser/BUILD.bazel
index 1fcad15..e99e6d0 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/parser/BUILD.bazel
@@ -51,6 +51,9 @@
"//src/tint/lang/core/intrinsic",
"//src/tint/lang/core/ir",
"//src/tint/lang/core/type",
+ "//src/tint/lang/spirv",
+ "//src/tint/lang/spirv/intrinsic",
+ "//src/tint/lang/spirv/ir",
"//src/tint/utils",
"//src/tint/utils/containers",
"//src/tint/utils/diagnostic",
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.cmake b/src/tint/lang/spirv/reader/parser/BUILD.cmake
index e8ddda8..d4c22cb 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/parser/BUILD.cmake
@@ -52,6 +52,9 @@
tint_lang_core_intrinsic
tint_lang_core_ir
tint_lang_core_type
+ tint_lang_spirv
+ tint_lang_spirv_intrinsic
+ tint_lang_spirv_ir
tint_utils
tint_utils_containers
tint_utils_diagnostic
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.gn b/src/tint/lang/spirv/reader/parser/BUILD.gn
index c54f228..9f4c271 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.gn
+++ b/src/tint/lang/spirv/reader/parser/BUILD.gn
@@ -56,6 +56,9 @@
"${tint_src_dir}/lang/core/intrinsic",
"${tint_src_dir}/lang/core/ir",
"${tint_src_dir}/lang/core/type",
+ "${tint_src_dir}/lang/spirv",
+ "${tint_src_dir}/lang/spirv/intrinsic",
+ "${tint_src_dir}/lang/spirv/ir",
"${tint_src_dir}/utils",
"${tint_src_dir}/utils/containers",
"${tint_src_dir}/utils/diagnostic",
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 3a70f3b..a32a8fb 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -46,6 +46,8 @@
#include "src/tint/lang/core/ir/builder.h"
#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/spirv/builtin_fn.h"
+#include "src/tint/lang/spirv/ir/builtin_call.h"
#include "src/tint/lang/spirv/validate/validate.h"
using namespace tint::core::fluent_types; // NOLINT
@@ -633,7 +635,7 @@
// Returns the WGSL standard library function for the given GLSL.std.450 extended instruction
// operation code. This handles GLSL functions which directly translate to the WGSL equivalent.
// Any non-direct translation is returned as `kNone`.
- core::BuiltinFn GetGlslStd450FuncName(uint32_t ext_opcode) {
+ core::BuiltinFn GetGlslStd450WgslEquivalentFuncName(uint32_t ext_opcode) {
switch (ext_opcode) {
case GLSLstd450Acos:
return core::BuiltinFn::kAcos;
@@ -753,22 +755,40 @@
return core::BuiltinFn::kNone;
}
+ spirv::BuiltinFn GetGlslStd450SpirvEquivalentFuncName(uint32_t ext_opcode) {
+ switch (ext_opcode) {
+ case GLSLstd450Normalize:
+ return spirv::BuiltinFn::kNormalize;
+ default:
+ break;
+ }
+ return spirv::BuiltinFn::kNone;
+ }
+
/// @param inst the SPIR-V instruction for OpAccessChain
void EmitGlslStd450ExtInst(const spvtools::opt::Instruction& inst) {
const auto ext_opcode = inst.GetSingleWordInOperand(1);
auto* result_ty = Type(inst.type_id());
- const auto fn = GetGlslStd450FuncName(ext_opcode);
- if (fn == core::BuiltinFn::kNone) {
- TINT_UNIMPLEMENTED() << "unhandled GLSL.std.450 instruction " << ext_opcode;
- }
-
Vector<core::ir::Value*, 4> operands;
// All parameters to GLSL.std.450 extended instructions are IDs.
for (uint32_t idx = 2; idx < inst.NumInOperands(); ++idx) {
operands.Push(Value(inst.GetSingleWordInOperand(idx)));
}
- Emit(b_.Call(result_ty, fn, operands), inst.result_id());
+
+ const auto wgsl_fn = GetGlslStd450WgslEquivalentFuncName(ext_opcode);
+ if (wgsl_fn != core::BuiltinFn::kNone) {
+ Emit(b_.Call(result_ty, wgsl_fn, operands), inst.result_id());
+ return;
+ }
+
+ const auto spv_fn = GetGlslStd450SpirvEquivalentFuncName(ext_opcode);
+ if (spv_fn != spirv::BuiltinFn::kNone) {
+ Emit(b_.Call<spirv::ir::BuiltinCall>(result_ty, spv_fn, operands), inst.result_id());
+ return;
+ }
+
+ TINT_UNIMPLEMENTED() << "unhandled GLSL.std.450 instruction " << ext_opcode;
}
/// @param inst the SPIR-V instruction for OpAccessChain
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index 0208789..cf7b5fc 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -318,6 +318,9 @@
implicit(T: f32_f16, N: num, M: num) fn vector_times_matrix(vec<N, T>, mat<M, N, T>) -> vec<M, T>
implicit(T: f32_f16, N: num) fn vector_times_scalar(vec<N, T>, T) -> vec<N, T>
+implicit(T: f32_f16) fn normalize(T) -> T
+implicit(N: num, T: f32_f16) fn normalize(vec<N, T>) -> vec<N, T>
+
////////////////////////////////////////////////////////////////////////////////
// SPV_KHR_integer_dot_product instructions
////////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index 7dbdd6c..b69b23c 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -116,6 +116,7 @@
namespace {
constexpr uint32_t kWriterVersion = 1;
+constexpr const char* kGLSLstd450 = "GLSL.std.450";
SpvStorageClass StorageClass(core::AddressSpace addrspace) {
switch (addrspace) {
@@ -1271,6 +1272,11 @@
void EmitSpirvBuiltinCall(spirv::ir::BuiltinCall* builtin) {
auto id = Value(builtin);
+ OperandList operands;
+ if (!builtin->Result(0)->Type()->Is<core::type::Void>()) {
+ operands = {Type(builtin->Result(0)->Type()), id};
+ }
+
spv::Op op = spv::Op::Max;
switch (builtin->Func()) {
case spirv::BuiltinFn::kArrayLength:
@@ -1362,6 +1368,11 @@
case spirv::BuiltinFn::kMatrixTimesVector:
op = spv::Op::OpMatrixTimesVector;
break;
+ case spirv::BuiltinFn::kNormalize:
+ op = spv::Op::OpExtInst;
+ operands.push_back(ImportGlslStd450());
+ operands.push_back(U32Operand(GLSLstd450Normalize));
+ break;
case spirv::BuiltinFn::kSampledImage:
op = spv::Op::OpSampledImage;
break;
@@ -1390,16 +1401,21 @@
TINT_ICE() << "undefined spirv ir function";
}
- OperandList operands;
- if (!builtin->Result(0)->Type()->Is<core::type::Void>()) {
- operands = {Type(builtin->Result(0)->Type()), id};
- }
for (auto* arg : builtin->Args()) {
operands.push_back(Value(arg));
}
current_function_.PushInst(op, operands);
}
+ uint32_t ImportGlslStd450() {
+ return imports_.GetOrAdd(kGLSLstd450, [&] {
+ // Import the instruction set the first time it is requested.
+ auto import = module_.NextId();
+ module_.PushExtImport(spv::Op::OpExtInstImport, {import, Operand(kGLSLstd450)});
+ return import;
+ });
+ }
+
/// Emit a builtin function call instruction.
/// @param builtin the builtin call instruction to emit
void EmitCoreBuiltinCall(core::ir::CoreBuiltinCall* builtin) {
@@ -1426,14 +1442,8 @@
// Helper to set up the opcode and operand list for a GLSL extended instruction.
auto glsl_ext_inst = [&](enum GLSLstd450 inst) {
- constexpr const char* kGLSLstd450 = "GLSL.std.450";
op = spv::Op::OpExtInst;
- operands.push_back(imports_.GetOrAdd(kGLSLstd450, [&] {
- // Import the instruction set the first time it is requested.
- auto import = module_.NextId();
- module_.PushExtImport(spv::Op::OpExtInstImport, {import, Operand(kGLSLstd450)});
- return import;
- }));
+ operands.push_back(ImportGlslStd450());
operands.push_back(U32Operand(inst));
};
