spirv-reader: Polyfill GLSLStd450 Degrees and Radians
Fixed: tint:1044
Change-Id: I5b2f3820b35c47bdc589ef41fb7a8735a7c6dff1
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59660
Auto-Submit: Corentin Wallez <cwallez@chromium.org>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/reader/spirv/function.cc b/src/reader/spirv/function.cc
index 90e63e7..fcf87f5 100644
--- a/src/reader/spirv/function.cc
+++ b/src/reader/spirv/function.cc
@@ -4065,6 +4065,41 @@
}
}
+ // Some GLSLStd450 builtins don't have a WGSL equivalent. Polyfill them.
+ switch (ext_opcode) {
+ case GLSLstd450Radians: {
+ auto degrees = MakeOperand(inst, 2);
+ TINT_ASSERT(Reader, degrees.type->IsFloatScalarOrVector());
+
+ constexpr auto kPiOver180 = static_cast<float>(3.141592653589793 / 180.0);
+ auto* factor = builder_.Expr(kPiOver180);
+ if (degrees.type->Is<F32>()) {
+ return {degrees.type, builder_.Mul(degrees.expr, factor)};
+ } else {
+ uint32_t size = degrees.type->As<Vector>()->size;
+ return {degrees.type,
+ builder_.Mul(degrees.expr,
+ builder_.vec(builder_.ty.f32(), size, factor))};
+ }
+ }
+
+ case GLSLstd450Degrees: {
+ auto radians = MakeOperand(inst, 2);
+ TINT_ASSERT(Reader, radians.type->IsFloatScalarOrVector());
+
+ constexpr auto k180OverPi = static_cast<float>(180.0 / 3.141592653589793);
+ auto* factor = builder_.Expr(k180OverPi);
+ if (radians.type->Is<F32>()) {
+ return {radians.type, builder_.Mul(radians.expr, factor)};
+ } else {
+ uint32_t size = radians.type->As<Vector>()->size;
+ return {radians.type,
+ builder_.Mul(radians.expr,
+ builder_.vec(builder_.ty.f32(), size, factor))};
+ }
+ }
+ }
+
const auto name = GetGlslStd450FuncName(ext_opcode);
if (name.empty()) {
Fail() << "unhandled GLSL.std.450 instruction " << ext_opcode;
diff --git a/src/reader/spirv/function_glsl_std_450_test.cc b/src/reader/spirv/function_glsl_std_450_test.cc
index be6d5a4..05ad700 100644
--- a/src/reader/spirv/function_glsl_std_450_test.cc
+++ b/src/reader/spirv/function_glsl_std_450_test.cc
@@ -2051,6 +2051,142 @@
EXPECT_THAT(body, HasSubstr(expected)) << body;
}
+TEST_F(SpvParserTest, GlslStd450_Degrees_Scalar) {
+ const auto assembly = Preamble() + R"(
+ %1 = OpExtInst %float %glsl Degrees %float_50
+ OpReturn
+ OpFunctionEnd
+ )";
+ auto p = parser(test::Assemble(assembly));
+ ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
+ auto fe = p->function_emitter(100);
+ EXPECT_TRUE(fe.EmitBody()) << p->error();
+ const auto body = ToString(p->builder(), fe.ast_body());
+ const auto* expected = R"(VariableDeclStatement{
+ VariableConst{
+ x_1
+ none
+ undefined
+ __f32
+ {
+ Binary[not set]{
+ ScalarConstructor[not set]{50.000000}
+ multiply
+ ScalarConstructor[not set]{57.295780}
+ }
+ }
+ }
+})";
+
+ EXPECT_THAT(body, HasSubstr(expected)) << body;
+}
+
+TEST_F(SpvParserTest, GlslStd450_Degrees_Vector) {
+ const auto assembly = Preamble() + R"(
+ %1 = OpExtInst %v3float %glsl Degrees %v3float_60_70_50
+ OpReturn
+ OpFunctionEnd
+ )";
+ auto p = parser(test::Assemble(assembly));
+ ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
+ auto fe = p->function_emitter(100);
+ EXPECT_TRUE(fe.EmitBody()) << p->error();
+ const auto body = ToString(p->builder(), fe.ast_body());
+ const auto* expected = R"(VariableDeclStatement{
+ VariableConst{
+ x_1
+ none
+ undefined
+ __vec_3__f32
+ {
+ Binary[not set]{
+ TypeConstructor[not set]{
+ __vec_3__f32
+ ScalarConstructor[not set]{60.000000}
+ ScalarConstructor[not set]{70.000000}
+ ScalarConstructor[not set]{50.000000}
+ }
+ multiply
+ TypeConstructor[not set]{
+ __vec_3__f32
+ ScalarConstructor[not set]{57.295780}
+ }
+ }
+ }
+ }
+})";
+
+ EXPECT_THAT(body, HasSubstr(expected)) << body;
+}
+
+TEST_F(SpvParserTest, GlslStd450_Radians_Scalar) {
+ const auto assembly = Preamble() + R"(
+ %1 = OpExtInst %float %glsl Radians %float_50
+ OpReturn
+ OpFunctionEnd
+ )";
+ auto p = parser(test::Assemble(assembly));
+ ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
+ auto fe = p->function_emitter(100);
+ EXPECT_TRUE(fe.EmitBody()) << p->error();
+ const auto body = ToString(p->builder(), fe.ast_body());
+ const auto* expected = R"(VariableDeclStatement{
+ VariableConst{
+ x_1
+ none
+ undefined
+ __f32
+ {
+ Binary[not set]{
+ ScalarConstructor[not set]{50.000000}
+ multiply
+ ScalarConstructor[not set]{0.017453}
+ }
+ }
+ }
+})";
+
+ EXPECT_THAT(body, HasSubstr(expected)) << body;
+}
+
+TEST_F(SpvParserTest, GlslStd450_Radians_Vector) {
+ const auto assembly = Preamble() + R"(
+ %1 = OpExtInst %v3float %glsl Radians %v3float_60_70_50
+ OpReturn
+ OpFunctionEnd
+ )";
+ auto p = parser(test::Assemble(assembly));
+ ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions());
+ auto fe = p->function_emitter(100);
+ EXPECT_TRUE(fe.EmitBody()) << p->error();
+ const auto body = ToString(p->builder(), fe.ast_body());
+ const auto* expected = R"(VariableDeclStatement{
+ VariableConst{
+ x_1
+ none
+ undefined
+ __vec_3__f32
+ {
+ Binary[not set]{
+ TypeConstructor[not set]{
+ __vec_3__f32
+ ScalarConstructor[not set]{60.000000}
+ ScalarConstructor[not set]{70.000000}
+ ScalarConstructor[not set]{50.000000}
+ }
+ multiply
+ TypeConstructor[not set]{
+ __vec_3__f32
+ ScalarConstructor[not set]{0.017453}
+ }
+ }
+ }
+ }
+})";
+
+ EXPECT_THAT(body, HasSubstr(expected)) << body;
+}
+
} // namespace
} // namespace spirv
} // namespace reader
diff --git a/test/intrinsics/degrees.spvasm b/test/intrinsics/degrees.spvasm
new file mode 100644
index 0000000..ddd7b86
--- /dev/null
+++ b/test/intrinsics/degrees.spvasm
@@ -0,0 +1,29 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 10
+; Bound: 13
+; Schema: 0
+ OpCapability Shader
+ %1 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpSource GLSL 460
+ OpName %main "main"
+ OpName %a "a"
+ OpName %b "b"
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+%_ptr_Function_float = OpTypePointer Function %float
+ %float_42 = OpConstant %float 42
+ %main = OpFunction %void None %3
+ %5 = OpLabel
+ %a = OpVariable %_ptr_Function_float Function
+ %b = OpVariable %_ptr_Function_float Function
+ OpStore %a %float_42
+ %11 = OpLoad %float %a
+ %12 = OpExtInst %float %1 Degrees %11
+ OpStore %b %12
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/degrees.spvasm.expected.hlsl b/test/intrinsics/degrees.spvasm.expected.hlsl
new file mode 100644
index 0000000..223f68e
--- /dev/null
+++ b/test/intrinsics/degrees.spvasm.expected.hlsl
@@ -0,0 +1,13 @@
+void main_1() {
+ float a = 0.0f;
+ float b = 0.0f;
+ a = 42.0f;
+ b = (a * 57.295780182f);
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ main_1();
+ return;
+}
diff --git a/test/intrinsics/degrees.spvasm.expected.msl b/test/intrinsics/degrees.spvasm.expected.msl
new file mode 100644
index 0000000..96eb526
--- /dev/null
+++ b/test/intrinsics/degrees.spvasm.expected.msl
@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void main_1() {
+ float a = 0.0f;
+ float b = 0.0f;
+ a = 42.0f;
+ float const x_11 = a;
+ b = (x_11 * 57.295780182f);
+ return;
+}
+
+kernel void tint_symbol() {
+ main_1();
+ return;
+}
+
diff --git a/test/intrinsics/degrees.spvasm.expected.spvasm b/test/intrinsics/degrees.spvasm.expected.spvasm
new file mode 100644
index 0000000..70266bc
--- /dev/null
+++ b/test/intrinsics/degrees.spvasm.expected.spvasm
@@ -0,0 +1,35 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main_1 "main_1"
+ OpName %a "a"
+ OpName %b "b"
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+%_ptr_Function_float = OpTypePointer Function %float
+ %8 = OpConstantNull %float
+ %float_42 = OpConstant %float 42
+%float_57_2957802 = OpConstant %float 57.2957802
+ %main_1 = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_float Function %8
+ %b = OpVariable %_ptr_Function_float Function %8
+ OpStore %a %float_42
+ %11 = OpLoad %float %a
+ %13 = OpFMul %float %11 %float_57_2957802
+ OpStore %b %13
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %15 = OpLabel
+ %16 = OpFunctionCall %void %main_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/degrees.spvasm.expected.wgsl b/test/intrinsics/degrees.spvasm.expected.wgsl
new file mode 100644
index 0000000..f1c58e0
--- /dev/null
+++ b/test/intrinsics/degrees.spvasm.expected.wgsl
@@ -0,0 +1,13 @@
+fn main_1() {
+ var a : f32;
+ var b : f32;
+ a = 42.0;
+ let x_11 : f32 = a;
+ b = (x_11 * 57.295780182);
+ return;
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+ main_1();
+}
diff --git a/test/intrinsics/radians.spvasm b/test/intrinsics/radians.spvasm
new file mode 100644
index 0000000..93445ff
--- /dev/null
+++ b/test/intrinsics/radians.spvasm
@@ -0,0 +1,29 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 10
+; Bound: 13
+; Schema: 0
+ OpCapability Shader
+ %1 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpSource GLSL 460
+ OpName %main "main"
+ OpName %a "a"
+ OpName %b "b"
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+%_ptr_Function_float = OpTypePointer Function %float
+ %float_42 = OpConstant %float 42
+ %main = OpFunction %void None %3
+ %5 = OpLabel
+ %a = OpVariable %_ptr_Function_float Function
+ %b = OpVariable %_ptr_Function_float Function
+ OpStore %a %float_42
+ %11 = OpLoad %float %a
+ %12 = OpExtInst %float %1 Radians %11
+ OpStore %b %12
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/radians.spvasm.expected.hlsl b/test/intrinsics/radians.spvasm.expected.hlsl
new file mode 100644
index 0000000..ec281fc
--- /dev/null
+++ b/test/intrinsics/radians.spvasm.expected.hlsl
@@ -0,0 +1,13 @@
+void main_1() {
+ float a = 0.0f;
+ float b = 0.0f;
+ a = 42.0f;
+ b = (a * 0.017453292f);
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ main_1();
+ return;
+}
diff --git a/test/intrinsics/radians.spvasm.expected.msl b/test/intrinsics/radians.spvasm.expected.msl
new file mode 100644
index 0000000..e7a33e4
--- /dev/null
+++ b/test/intrinsics/radians.spvasm.expected.msl
@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void main_1() {
+ float a = 0.0f;
+ float b = 0.0f;
+ a = 42.0f;
+ float const x_11 = a;
+ b = (x_11 * 0.017453292f);
+ return;
+}
+
+kernel void tint_symbol() {
+ main_1();
+ return;
+}
+
diff --git a/test/intrinsics/radians.spvasm.expected.spvasm b/test/intrinsics/radians.spvasm.expected.spvasm
new file mode 100644
index 0000000..bc75a2f
--- /dev/null
+++ b/test/intrinsics/radians.spvasm.expected.spvasm
@@ -0,0 +1,35 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main_1 "main_1"
+ OpName %a "a"
+ OpName %b "b"
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+%_ptr_Function_float = OpTypePointer Function %float
+ %8 = OpConstantNull %float
+ %float_42 = OpConstant %float 42
+%float_0_0174532924 = OpConstant %float 0.0174532924
+ %main_1 = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_float Function %8
+ %b = OpVariable %_ptr_Function_float Function %8
+ OpStore %a %float_42
+ %11 = OpLoad %float %a
+ %13 = OpFMul %float %11 %float_0_0174532924
+ OpStore %b %13
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %15 = OpLabel
+ %16 = OpFunctionCall %void %main_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/radians.spvasm.expected.wgsl b/test/intrinsics/radians.spvasm.expected.wgsl
new file mode 100644
index 0000000..b9900af
--- /dev/null
+++ b/test/intrinsics/radians.spvasm.expected.wgsl
@@ -0,0 +1,13 @@
+fn main_1() {
+ var a : f32;
+ var b : f32;
+ a = 42.0;
+ let x_11 : f32 = a;
+ b = (x_11 * 0.017453292);
+ return;
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn main() {
+ main_1();
+}
