tint: Implement modf and frexp built-ins for f16 types
This patch implement modf and frexp built-ins for f16 types, and also
simplify their implementation for f32 in MSL and HLSL, and clean up
deprecated code in GLSL writer. Corresponding unittests are also
implemented, but end-to-end tests for f16 are not implemented yet.
Bug: tint:1473, tint:1502
Change-Id: I12887ae5303c6dc032a51f619e1afeb19b4603b6
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/98102
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Zhaoming Jiang <zhaoming.jiang@intel.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index b7a3be6..c66e174 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -1205,97 +1205,51 @@
bool GeneratorImpl::EmitModfCall(std::ostream& out,
const ast::CallExpression* expr,
const sem::Builtin* builtin) {
- if (expr->args.Length() == 1) {
- return CallBuiltinHelper(
- out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
- // Emit the builtin return type unique to this overload. This does not
- // exist in the AST, so it will not be generated in Generate().
- if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
+ TINT_ASSERT(Writer, expr->args.Length() == 1);
+ return CallBuiltinHelper(
+ out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
+ // Emit the builtin return type unique to this overload. This does not
+ // exist in the AST, so it will not be generated in Generate().
+ if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
+ return false;
+ }
+
+ {
+ auto l = line(b);
+ if (!EmitType(l, builtin->ReturnType(), ast::StorageClass::kNone,
+ ast::Access::kUndefined, "")) {
return false;
}
-
- {
- auto l = line(b);
- if (!EmitType(l, builtin->ReturnType(), ast::StorageClass::kNone,
- ast::Access::kUndefined, "")) {
- return false;
- }
- l << " result;";
- }
- line(b) << "result.fract = modf(" << params[0] << ", result.whole);";
- line(b) << "return result;";
- return true;
- });
- }
-
- // DEPRECATED
- out << "modf";
- ScopedParen sp(out);
- if (!EmitExpression(out, expr->args[0])) {
- return false;
- }
- out << ", ";
- if (!EmitExpression(out, expr->args[1])) {
- return false;
- }
- return true;
+ l << " result;";
+ }
+ line(b) << "result.fract = modf(" << params[0] << ", result.whole);";
+ line(b) << "return result;";
+ return true;
+ });
}
bool GeneratorImpl::EmitFrexpCall(std::ostream& out,
const ast::CallExpression* expr,
const sem::Builtin* builtin) {
- if (expr->args.Length() == 1) {
- return CallBuiltinHelper(
- out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
- // Emit the builtin return type unique to this overload. This does not
- // exist in the AST, so it will not be generated in Generate().
- if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
- return false;
- }
-
- {
- auto l = line(b);
- if (!EmitType(l, builtin->ReturnType(), ast::StorageClass::kNone,
- ast::Access::kUndefined, "")) {
- return false;
- }
- l << " result;";
- }
- line(b) << "result.sig = frexp(" << params[0] << ", result.exp);";
- line(b) << "return result;";
- return true;
- });
- }
- // DEPRECATED
- // Exponent is an integer in WGSL, but HLSL wants a float.
- // We need to make the call with a temporary float, and then cast.
+ TINT_ASSERT(Writer, expr->args.Length() == 1);
return CallBuiltinHelper(
out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
- auto* significand_ty = builtin->Parameters()[0]->Type();
- auto significand = params[0];
- auto* exponent_ty = builtin->Parameters()[1]->Type();
- auto exponent = params[1];
-
- std::string width;
- if (auto* vec = significand_ty->As<sem::Vector>()) {
- width = std::to_string(vec->Width());
+ // Emit the builtin return type unique to this overload. This does not
+ // exist in the AST, so it will not be generated in Generate().
+ if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
+ return false;
}
- // Exponent is an integer, which HLSL does not have an overload for.
- // We need to cast from a float.
- line(b) << "float" << width << " float_exp;";
- line(b) << "float" << width << " significand = frexp(" << significand
- << ", float_exp);";
{
auto l = line(b);
- l << exponent << " = ";
- if (!EmitType(l, exponent_ty->UnwrapPtr(), ast::StorageClass::kNone,
+ if (!EmitType(l, builtin->ReturnType(), ast::StorageClass::kNone,
ast::Access::kUndefined, "")) {
return false;
}
- l << "(float_exp);";
+ l << " result;";
}
- line(b) << "return significand;";
+ line(b) << "result.sig = frexp(" << params[0] << ", result.exp);";
+ line(b) << "return result;";
return true;
});
}
diff --git a/src/tint/writer/glsl/generator_impl_builtin_test.cc b/src/tint/writer/glsl/generator_impl_builtin_test.cc
index c7c8d53..b8161e2 100644
--- a/src/tint/writer/glsl/generator_impl_builtin_test.cc
+++ b/src/tint/writer/glsl/generator_impl_builtin_test.cc
@@ -412,7 +412,7 @@
EXPECT_EQ(out.str(), "((a) * (b) + (c))");
}
-TEST_F(GlslGeneratorImplTest_Builtin, Modf_Scalar) {
+TEST_F(GlslGeneratorImplTest_Builtin, Modf_Scalar_f32) {
auto* call = Call("modf", 1_f);
WrapInFunction(CallStmt(call));
@@ -445,7 +445,43 @@
)");
}
-TEST_F(GlslGeneratorImplTest_Builtin, Modf_Vector) {
+TEST_F(GlslGeneratorImplTest_Builtin, Modf_Scalar_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* call = Call("modf", 1_h);
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+struct modf_result_f16 {
+ float16_t fract;
+ float16_t whole;
+};
+
+modf_result_f16 tint_modf(float16_t param_0) {
+ modf_result_f16 result;
+ result.fract = modf(param_0, result.whole);
+ return result;
+}
+
+
+void test_function() {
+ tint_modf(1.0hf);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ test_function();
+ return;
+}
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Builtin, Modf_Vector_f32) {
auto* call = Call("modf", vec3<f32>());
WrapInFunction(CallStmt(call));
@@ -478,7 +514,43 @@
)");
}
-TEST_F(GlslGeneratorImplTest_Builtin, Frexp_Scalar_i32) {
+TEST_F(GlslGeneratorImplTest_Builtin, Modf_Vector_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* call = Call("modf", vec3<f16>());
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+struct modf_result_vec3_f16 {
+ f16vec3 fract;
+ f16vec3 whole;
+};
+
+modf_result_vec3_f16 tint_modf(f16vec3 param_0) {
+ modf_result_vec3_f16 result;
+ result.fract = modf(param_0, result.whole);
+ return result;
+}
+
+
+void test_function() {
+ tint_modf(f16vec3(0.0hf));
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ test_function();
+ return;
+}
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Builtin, Frexp_Scalar_f32) {
auto* call = Call("frexp", 1_f);
WrapInFunction(CallStmt(call));
@@ -505,7 +577,42 @@
)"));
}
-TEST_F(GlslGeneratorImplTest_Builtin, Frexp_Vector_i32) {
+TEST_F(GlslGeneratorImplTest_Builtin, Frexp_Scalar_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* call = Call("frexp", 1_h);
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(R"(#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+struct frexp_result_f16 {
+ float16_t sig;
+ int exp;
+};
+
+frexp_result_f16 tint_frexp(float16_t param_0) {
+ frexp_result_f16 result;
+ result.sig = frexp(param_0, result.exp);
+ return result;
+}
+
+
+void test_function() {
+ tint_frexp(1.0hf);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ test_function();
+ return;
+)"));
+}
+
+TEST_F(GlslGeneratorImplTest_Builtin, Frexp_Vector_f32) {
auto* call = Call("frexp", vec3<f32>());
WrapInFunction(CallStmt(call));
@@ -537,6 +644,42 @@
)"));
}
+TEST_F(GlslGeneratorImplTest_Builtin, Frexp_Vector_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* call = Call("frexp", vec3<f16>());
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_THAT(gen.result(), HasSubstr(R"(#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+struct frexp_result_vec3_f16 {
+ f16vec3 sig;
+ ivec3 exp;
+};
+
+frexp_result_vec3_f16 tint_frexp(f16vec3 param_0) {
+ frexp_result_vec3_f16 result;
+ result.sig = frexp(param_0, result.exp);
+ return result;
+}
+
+
+void test_function() {
+ tint_frexp(f16vec3(0.0hf));
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ test_function();
+ return;
+}
+)"));
+}
+
TEST_F(GlslGeneratorImplTest_Builtin, Degrees_Scalar_f32) {
auto* val = Var("val", ty.f32());
auto* call = Call("degrees", val);
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index 726a459..444746c 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -1855,16 +1855,15 @@
return false;
}
- line(b) << "float" << width << " whole;";
- line(b) << "float" << width << " fract = modf(" << in << ", whole);";
{
auto l = line(b);
if (!EmitType(l, builtin->ReturnType(), ast::StorageClass::kNone,
ast::Access::kUndefined, "")) {
return false;
}
- l << " result = {fract, whole};";
+ l << " result;";
}
+ line(b) << "result.fract = modf(" << params[0] << ", result.whole);";
line(b) << "return result;";
return true;
});
@@ -1889,8 +1888,15 @@
return false;
}
- line(b) << "float" << width << " exp;";
- line(b) << "float" << width << " sig = frexp(" << in << ", exp);";
+ std::string member_type;
+ if (Is<sem::F16>(sem::Type::DeepestElementOf(ty))) {
+ member_type = width.empty() ? "float16_t" : ("vector<float16_t, " + width + ">");
+ } else {
+ member_type = "float" + width;
+ }
+
+ line(b) << member_type << " exp;";
+ line(b) << member_type << " sig = frexp(" << in << ", exp);";
{
auto l = line(b);
if (!EmitType(l, builtin->ReturnType(), ast::StorageClass::kNone,
diff --git a/src/tint/writer/hlsl/generator_impl_builtin_test.cc b/src/tint/writer/hlsl/generator_impl_builtin_test.cc
index 82c7e7e..3f9cfa7 100644
--- a/src/tint/writer/hlsl/generator_impl_builtin_test.cc
+++ b/src/tint/writer/hlsl/generator_impl_builtin_test.cc
@@ -377,7 +377,7 @@
EXPECT_EQ(out.str(), "(bool2(true, false) ? int2(3, 4) : int2(1, 2))");
}
-TEST_F(HlslGeneratorImplTest_Builtin, Modf_Scalar) {
+TEST_F(HlslGeneratorImplTest_Builtin, Modf_Scalar_f32) {
auto* call = Call("modf", 1_f);
WrapInFunction(CallStmt(call));
@@ -389,9 +389,8 @@
float whole;
};
modf_result tint_modf(float param_0) {
- float whole;
- float fract = modf(param_0, whole);
- modf_result result = {fract, whole};
+ modf_result result;
+ result.fract = modf(param_0, result.whole);
return result;
}
@@ -403,7 +402,34 @@
)");
}
-TEST_F(HlslGeneratorImplTest_Builtin, Modf_Vector) {
+TEST_F(HlslGeneratorImplTest_Builtin, Modf_Scalar_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* call = Call("modf", 1_h);
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct modf_result_f16 {
+ float16_t fract;
+ float16_t whole;
+};
+modf_result_f16 tint_modf(float16_t param_0) {
+ modf_result_f16 result;
+ result.fract = modf(param_0, result.whole);
+ return result;
+}
+
+[numthreads(1, 1, 1)]
+void test_function() {
+ tint_modf(float16_t(1.0h));
+ return;
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Builtin, Modf_Vector_f32) {
auto* call = Call("modf", vec3<f32>());
WrapInFunction(CallStmt(call));
@@ -415,9 +441,8 @@
float3 whole;
};
modf_result_vec3 tint_modf(float3 param_0) {
- float3 whole;
- float3 fract = modf(param_0, whole);
- modf_result_vec3 result = {fract, whole};
+ modf_result_vec3 result;
+ result.fract = modf(param_0, result.whole);
return result;
}
@@ -429,7 +454,34 @@
)");
}
-TEST_F(HlslGeneratorImplTest_Builtin, Frexp_Scalar_i32) {
+TEST_F(HlslGeneratorImplTest_Builtin, Modf_Vector_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* call = Call("modf", vec3<f16>());
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct modf_result_vec3_f16 {
+ vector<float16_t, 3> fract;
+ vector<float16_t, 3> whole;
+};
+modf_result_vec3_f16 tint_modf(vector<float16_t, 3> param_0) {
+ modf_result_vec3_f16 result;
+ result.fract = modf(param_0, result.whole);
+ return result;
+}
+
+[numthreads(1, 1, 1)]
+void test_function() {
+ tint_modf((float16_t(0.0h)).xxx);
+ return;
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Builtin, Frexp_Scalar_f32) {
auto* call = Call("frexp", 1_f);
WrapInFunction(CallStmt(call));
@@ -455,7 +507,35 @@
)");
}
-TEST_F(HlslGeneratorImplTest_Builtin, Frexp_Vector_i32) {
+TEST_F(HlslGeneratorImplTest_Builtin, Frexp_Scalar_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* call = Call("frexp", 1_h);
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct frexp_result_f16 {
+ float16_t sig;
+ int exp;
+};
+frexp_result_f16 tint_frexp(float16_t param_0) {
+ float16_t exp;
+ float16_t sig = frexp(param_0, exp);
+ frexp_result_f16 result = {sig, int(exp)};
+ return result;
+}
+
+[numthreads(1, 1, 1)]
+void test_function() {
+ tint_frexp(float16_t(1.0h));
+ return;
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Builtin, Frexp_Vector_f32) {
auto* call = Call("frexp", vec3<f32>());
WrapInFunction(CallStmt(call));
@@ -481,6 +561,34 @@
)");
}
+TEST_F(HlslGeneratorImplTest_Builtin, Frexp_Vector_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* call = Call("frexp", vec3<f16>());
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct frexp_result_vec3_f16 {
+ vector<float16_t, 3> sig;
+ int3 exp;
+};
+frexp_result_vec3_f16 tint_frexp(vector<float16_t, 3> param_0) {
+ vector<float16_t, 3> exp;
+ vector<float16_t, 3> sig = frexp(param_0, exp);
+ frexp_result_vec3_f16 result = {sig, int3(exp)};
+ return result;
+}
+
+[numthreads(1, 1, 1)]
+void test_function() {
+ tint_frexp((float16_t(0.0h)).xxx);
+ return;
+}
+)");
+}
+
TEST_F(HlslGeneratorImplTest_Builtin, Degrees_Scalar_f32) {
auto* val = Var("val", ty.f32());
auto* call = Call("degrees", val);
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index 577ecd7..28d2c10 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -1308,9 +1308,9 @@
return false;
}
- line(b) << "float" << width << " whole;";
- line(b) << "float" << width << " fract = modf(" << in << ", whole);";
- line(b) << "return {fract, whole};";
+ line(b) << StructName(builtin->ReturnType()->As<sem::Struct>()) << " result;";
+ line(b) << "result.fract = modf(" << in << ", result.whole);";
+ line(b) << "return result;";
return true;
});
}
@@ -1334,9 +1334,9 @@
return false;
}
- line(b) << "int" << width << " exp;";
- line(b) << "float" << width << " sig = frexp(" << in << ", exp);";
- line(b) << "return {sig, exp};";
+ line(b) << StructName(builtin->ReturnType()->As<sem::Struct>()) << " result;";
+ line(b) << "result.sig = frexp(" << in << ", result.exp);";
+ line(b) << "return result;";
return true;
});
}
diff --git a/src/tint/writer/msl/generator_impl_builtin_test.cc b/src/tint/writer/msl/generator_impl_builtin_test.cc
index b97a108..9f75691 100644
--- a/src/tint/writer/msl/generator_impl_builtin_test.cc
+++ b/src/tint/writer/msl/generator_impl_builtin_test.cc
@@ -405,6 +405,246 @@
EXPECT_EQ(out.str(), "threadgroup_barrier(mem_flags::mem_threadgroup)");
}
+TEST_F(MslGeneratorImplTest, Modf_Scalar_f32) {
+ auto* call = Call("modf", 1_f);
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result {
+ float fract;
+ float whole;
+};
+modf_result tint_modf(float param_0) {
+ modf_result result;
+ result.fract = modf(param_0, result.whole);
+ return result;
+}
+
+kernel void test_function() {
+ tint_modf(1.0f);
+ return;
+}
+
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Modf_Scalar_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* call = Call("modf", 1_h);
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result_f16 {
+ half fract;
+ half whole;
+};
+modf_result_f16 tint_modf(half param_0) {
+ modf_result_f16 result;
+ result.fract = modf(param_0, result.whole);
+ return result;
+}
+
+kernel void test_function() {
+ tint_modf(1.0h);
+ return;
+}
+
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Modf_Vector_f32) {
+ auto* call = Call("modf", vec3<f32>());
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result_vec3 {
+ float3 fract;
+ float3 whole;
+};
+modf_result_vec3 tint_modf(float3 param_0) {
+ modf_result_vec3 result;
+ result.fract = modf(param_0, result.whole);
+ return result;
+}
+
+kernel void test_function() {
+ tint_modf(float3(0.0f));
+ return;
+}
+
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Modf_Vector_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* call = Call("modf", vec3<f16>());
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result_vec3_f16 {
+ half3 fract;
+ half3 whole;
+};
+modf_result_vec3_f16 tint_modf(half3 param_0) {
+ modf_result_vec3_f16 result;
+ result.fract = modf(param_0, result.whole);
+ return result;
+}
+
+kernel void test_function() {
+ tint_modf(half3(0.0h));
+ return;
+}
+
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Frexp_Scalar_f32) {
+ auto* call = Call("frexp", 1_f);
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+
+struct frexp_result {
+ float sig;
+ int exp;
+};
+frexp_result tint_frexp(float param_0) {
+ frexp_result result;
+ result.sig = frexp(param_0, result.exp);
+ return result;
+}
+
+kernel void test_function() {
+ tint_frexp(1.0f);
+ return;
+}
+
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Frexp_Scalar_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* call = Call("frexp", 1_h);
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+
+struct frexp_result_f16 {
+ half sig;
+ int exp;
+};
+frexp_result_f16 tint_frexp(half param_0) {
+ frexp_result_f16 result;
+ result.sig = frexp(param_0, result.exp);
+ return result;
+}
+
+kernel void test_function() {
+ tint_frexp(1.0h);
+ return;
+}
+
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Frexp_Vector_f32) {
+ auto* call = Call("frexp", vec3<f32>());
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+
+struct frexp_result_vec3 {
+ float3 sig;
+ int3 exp;
+};
+frexp_result_vec3 tint_frexp(float3 param_0) {
+ frexp_result_vec3 result;
+ result.sig = frexp(param_0, result.exp);
+ return result;
+}
+
+kernel void test_function() {
+ tint_frexp(float3(0.0f));
+ return;
+}
+
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Frexp_Vector_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* call = Call("frexp", vec3<f16>());
+ WrapInFunction(CallStmt(call));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+
+struct frexp_result_vec3_f16 {
+ half3 sig;
+ int3 exp;
+};
+frexp_result_vec3_f16 tint_frexp(half3 param_0) {
+ frexp_result_vec3_f16 result;
+ result.sig = frexp(param_0, result.exp);
+ return result;
+}
+
+kernel void test_function() {
+ tint_frexp(half3(0.0h));
+ return;
+}
+
+)");
+}
+
TEST_F(MslGeneratorImplTest, Degrees_Scalar_f32) {
auto* val = Var("val", ty.f32());
auto* call = Call("degrees", val);
diff --git a/src/tint/writer/spirv/builder_builtin_test.cc b/src/tint/writer/spirv/builder_builtin_test.cc
index 1faf474..52254d2 100644
--- a/src/tint/writer/spirv/builder_builtin_test.cc
+++ b/src/tint/writer/spirv/builder_builtin_test.cc
@@ -1651,7 +1651,7 @@
EXPECT_EQ(got, expect);
}
-TEST_F(BuiltinBuilderTest, Call_Modf) {
+TEST_F(BuiltinBuilderTest, Call_Modf_f32) {
auto* vec = Var("vec", nullptr, vec2<f32>(1_f, 2_f));
auto* expr = Call("modf", vec);
Func("a_func", utils::Empty, ty.void_(),
@@ -1703,7 +1703,65 @@
Validate(b);
}
-TEST_F(BuiltinBuilderTest, Call_Frexp) {
+TEST_F(BuiltinBuilderTest, Call_Modf_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* vec = Var("vec", nullptr, vec2<f16>(1_h, 2_h));
+ auto* expr = Call("modf", vec);
+ Func("a_func", utils::Empty, ty.void_(),
+ utils::Vector{
+ Decl(vec),
+ CallStmt(expr),
+ },
+ utils::Vector{
+ Stage(ast::PipelineStage::kFragment),
+ });
+
+ spirv::Builder& b = Build();
+
+ ASSERT_TRUE(b.Build()) << b.error();
+ auto got = DumpBuilder(b);
+ auto* expect = R"(OpCapability Shader
+OpCapability Float16
+OpCapability UniformAndStorageBuffer16BitAccess
+OpCapability StorageBuffer16BitAccess
+OpCapability StorageInputOutput16
+%15 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint Fragment %3 "a_func"
+OpExecutionMode %3 OriginUpperLeft
+OpName %3 "a_func"
+OpName %10 "vec"
+OpName %14 "__modf_result_vec2_f16"
+OpMemberName %14 0 "fract"
+OpMemberName %14 1 "whole"
+OpMemberDecorate %14 0 Offset 0
+OpMemberDecorate %14 1 Offset 4
+%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%6 = OpTypeFloat 16
+%5 = OpTypeVector %6 2
+%7 = OpConstant %6 0x1p+0
+%8 = OpConstant %6 0x1p+1
+%9 = OpConstantComposite %5 %7 %8
+%11 = OpTypePointer Function %5
+%12 = OpConstantNull %5
+%14 = OpTypeStruct %5 %5
+%3 = OpFunction %2 None %1
+%4 = OpLabel
+%10 = OpVariable %11 Function %12
+OpStore %10 %9
+%16 = OpLoad %5 %10
+%13 = OpExtInst %14 %15 ModfStruct %16
+OpReturn
+OpFunctionEnd
+)";
+ EXPECT_EQ(expect, got);
+
+ Validate(b);
+}
+
+TEST_F(BuiltinBuilderTest, Call_Frexp_f32) {
auto* vec = Var("vec", nullptr, vec2<f32>(1_f, 2_f));
auto* expr = Call("frexp", vec);
Func("a_func", utils::Empty, ty.void_(),
@@ -1757,6 +1815,66 @@
Validate(b);
}
+TEST_F(BuiltinBuilderTest, Call_Frexp_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* vec = Var("vec", nullptr, vec2<f16>(1_h, 2_h));
+ auto* expr = Call("frexp", vec);
+ Func("a_func", utils::Empty, ty.void_(),
+ utils::Vector{
+ Decl(vec),
+ CallStmt(expr),
+ },
+ utils::Vector{
+ Stage(ast::PipelineStage::kFragment),
+ });
+
+ spirv::Builder& b = Build();
+
+ ASSERT_TRUE(b.Build()) << b.error();
+ auto got = DumpBuilder(b);
+ auto* expect = R"(OpCapability Shader
+OpCapability Float16
+OpCapability UniformAndStorageBuffer16BitAccess
+OpCapability StorageBuffer16BitAccess
+OpCapability StorageInputOutput16
+%17 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint Fragment %3 "a_func"
+OpExecutionMode %3 OriginUpperLeft
+OpName %3 "a_func"
+OpName %10 "vec"
+OpName %14 "__frexp_result_vec2_f16"
+OpMemberName %14 0 "sig"
+OpMemberName %14 1 "exp"
+OpMemberDecorate %14 0 Offset 0
+OpMemberDecorate %14 1 Offset 8
+%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%6 = OpTypeFloat 16
+%5 = OpTypeVector %6 2
+%7 = OpConstant %6 0x1p+0
+%8 = OpConstant %6 0x1p+1
+%9 = OpConstantComposite %5 %7 %8
+%11 = OpTypePointer Function %5
+%12 = OpConstantNull %5
+%16 = OpTypeInt 32 1
+%15 = OpTypeVector %16 2
+%14 = OpTypeStruct %5 %15
+%3 = OpFunction %2 None %1
+%4 = OpLabel
+%10 = OpVariable %11 Function %12
+OpStore %10 %9
+%18 = OpLoad %5 %10
+%13 = OpExtInst %14 %17 FrexpStruct %18
+OpReturn
+OpFunctionEnd
+)";
+ EXPECT_EQ(expect, got);
+
+ Validate(b);
+}
+
} // namespace float_builtin_tests
// Tests for Numeric builtins with all integer parameter
