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>
This commit is contained in:
Corentin Wallez
Tint LUCI CQ
parent
13081d4e46
commit
b32c22cead
|
|
@ -4065,6 +4065,41 @@ TypedExpression FunctionEmitter::EmitGlslStd450ExtInst(
|
|||
}
|
||||
}
|
||||
|
||||
// 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;
|
||||
|
|
|
|||
|
|
@ -2051,6 +2051,142 @@ VariableDeclStatement{
|
|||
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
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
||||
|
|
@ -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
|
||||
|
|
@ -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();
|
||||
}
|
||||
|
|
@ -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
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
||||
|
|
@ -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
|
||||
|
|
@ -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();
|
||||
}
|
||||
Loading…
Reference in New Issue