encounter/dawn-cmake
1
0
Fork 0
mirror of https://github.com/encounter/dawn-cmake.git synced 2025-12-17 00:47:13 +00:00
Code Issues Packages Projects Releases Wiki Activity

tint: Implement modf and frexp built-ins for f16 types

Browse Source 
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>
This commit is contained in:
Zhaoming Jiang
2022-08-05 15:11:44 +00:00
committed by Dawn LUCI CQ
parent 1cd4706f85
commit 20cddbf04c
130 changed files with 2519 additions and 1775 deletions
Download Patch File Download Diff File Expand all files Collapse all files

16
src/tint/intrinsics.def
View File

@@ -148,10 +148,10 @@ type texture_storage_2d_array<F: texel_format, A: access>
type texture_storage_3d<F: texel_format, A: access>
type texture_external
type __modf_result
@display("__modf_result_vec{N}") type __modf_result_vec<N: num>
type __frexp_result
@display("__frexp_result_vec{N}") type __frexp_result_vec<N: num>
@display("__modf_result_{T}") type __modf_result<T>
@display("__modf_result_vec{N}_{T}") type __modf_result_vec<N: num, T>
@display("__frexp_result_{T}") type __frexp_result<T>
@display("__frexp_result_vec{N}_{T}") type __frexp_result_vec<N: num, T>
type __atomic_compare_exchange_result<T>
@@ -463,8 +463,8 @@ fn fma<T: f32_f16>(T, T, T) -> T
fn fma<N: num, T: f32_f16>(vec<N, T>, vec<N, T>, vec<N, T>) -> vec<N, T>
fn fract<T: f32_f16>(T) -> T
fn fract<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
fn frexp(f32) -> __frexp_result
fn frexp<N: num>(vec<N, f32>) -> __frexp_result_vec<N>
fn frexp<T: f32_f16>(T) -> __frexp_result<T>
fn frexp<N: num, T: f32_f16>(vec<N, T>) -> __frexp_result_vec<N, T>
@stage("fragment") fn fwidth(f32) -> f32
@stage("fragment") fn fwidth<N: num>(vec<N, f32>) -> vec<N, f32>
@stage("fragment") fn fwidthCoarse(f32) -> f32
@@ -490,8 +490,8 @@ fn min<N: num, T: fiu32_f16>(vec<N, T>, vec<N, T>) -> vec<N, T>
fn mix<T: f32_f16>(T, T, T) -> T
fn mix<N: num, T: f32_f16>(vec<N, T>, vec<N, T>, vec<N, T>) -> vec<N, T>
fn mix<N: num, T: f32_f16>(vec<N, T>, vec<N, T>, T) -> vec<N, T>
fn modf(f32) -> __modf_result
fn modf<N: num>(vec<N, f32>) -> __modf_result_vec<N>
fn modf<T: f32_f16>(T) -> __modf_result<T>
fn modf<N: num, T: f32_f16>(vec<N, T>) -> __modf_result_vec<N, T>
fn normalize<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
fn pack2x16float(vec2<f32>) -> u32
fn pack2x16snorm(vec2<f32>) -> u32

178
src/tint/resolver/builtin_test.cc
View File

@@ -901,8 +901,9 @@ TEST_F(ResolverBuiltinFloatTest, Distance_NoParams) {
)");
}
// frexp: (f32) -> __frexp_result, (vecN<f32>) -> __frexp_result_vecN
TEST_F(ResolverBuiltinFloatTest, FrexpScalar) {
// frexp: (f32) -> __frexp_result, (vecN<f32>) -> __frexp_result_vecN, (f16) -> __frexp_result_16,
// (vecN<f16>) -> __frexp_result_vecN_f16
TEST_F(ResolverBuiltinFloatTest, FrexpScalar_f32) {
auto* call = Call("frexp", 1_f);
WrapInFunction(call);
@@ -931,7 +932,38 @@ TEST_F(ResolverBuiltinFloatTest, FrexpScalar) {
EXPECT_EQ(ty->SizeNoPadding(), 8u);
}
TEST_F(ResolverBuiltinFloatTest, FrexpVector) {
TEST_F(ResolverBuiltinFloatTest, FrexpScalar_f16) {
Enable(ast::Extension::kF16);
auto* call = Call("frexp", 1_h);
WrapInFunction(call);
EXPECT_TRUE(r()->Resolve()) << r()->error();
ASSERT_NE(TypeOf(call), nullptr);
auto* ty = TypeOf(call)->As<sem::Struct>();
ASSERT_NE(ty, nullptr);
ASSERT_EQ(ty->Members().size(), 2u);
auto* sig = ty->Members()[0];
EXPECT_TRUE(sig->Type()->Is<sem::F16>());
EXPECT_EQ(sig->Offset(), 0u);
EXPECT_EQ(sig->Size(), 2u);
EXPECT_EQ(sig->Align(), 2u);
EXPECT_EQ(sig->Name(), Sym("sig"));
auto* exp = ty->Members()[1];
EXPECT_TRUE(exp->Type()->Is<sem::I32>());
EXPECT_EQ(exp->Offset(), 4u);
EXPECT_EQ(exp->Size(), 4u);
EXPECT_EQ(exp->Align(), 4u);
EXPECT_EQ(exp->Name(), Sym("exp"));
EXPECT_EQ(ty->Size(), 8u);
EXPECT_EQ(ty->SizeNoPadding(), 8u);
}
TEST_F(ResolverBuiltinFloatTest, FrexpVector_f32) {
auto* call = Call("frexp", vec3<f32>());
WrapInFunction(call);
@@ -964,6 +996,41 @@ TEST_F(ResolverBuiltinFloatTest, FrexpVector) {
EXPECT_EQ(ty->SizeNoPadding(), 28u);
}
TEST_F(ResolverBuiltinFloatTest, FrexpVector_f16) {
Enable(ast::Extension::kF16);
auto* call = Call("frexp", vec3<f16>());
WrapInFunction(call);
EXPECT_TRUE(r()->Resolve()) << r()->error();
ASSERT_NE(TypeOf(call), nullptr);
auto* ty = TypeOf(call)->As<sem::Struct>();
ASSERT_NE(ty, nullptr);
ASSERT_EQ(ty->Members().size(), 2u);
auto* sig = ty->Members()[0];
ASSERT_TRUE(sig->Type()->Is<sem::Vector>());
EXPECT_EQ(sig->Type()->As<sem::Vector>()->Width(), 3u);
EXPECT_TRUE(sig->Type()->As<sem::Vector>()->type()->Is<sem::F16>());
EXPECT_EQ(sig->Offset(), 0u);
EXPECT_EQ(sig->Size(), 6u);
EXPECT_EQ(sig->Align(), 8u);
EXPECT_EQ(sig->Name(), Sym("sig"));
auto* exp = ty->Members()[1];
ASSERT_TRUE(exp->Type()->Is<sem::Vector>());
EXPECT_EQ(exp->Type()->As<sem::Vector>()->Width(), 3u);
EXPECT_TRUE(exp->Type()->As<sem::Vector>()->type()->Is<sem::I32>());
EXPECT_EQ(exp->Offset(), 16u);
EXPECT_EQ(exp->Size(), 12u);
EXPECT_EQ(exp->Align(), 16u);
EXPECT_EQ(exp->Name(), Sym("exp"));
EXPECT_EQ(ty->Size(), 32u);
EXPECT_EQ(ty->SizeNoPadding(), 28u);
}
TEST_F(ResolverBuiltinFloatTest, Frexp_Error_FirstParamInt) {
GlobalVar("v", ty.i32(), ast::StorageClass::kWorkgroup);
auto* call = Call("frexp", 1_i, AddressOf("v"));
@@ -975,8 +1042,8 @@ TEST_F(ResolverBuiltinFloatTest, Frexp_Error_FirstParamInt) {
R"(error: no matching call to frexp(i32, ptr<workgroup, i32, read_write>)
2 candidate functions:
frexp(f32) -> __frexp_result
frexp(vecN<f32>) -> __frexp_result_vecN
frexp(T) -> __frexp_result_T where: T is f32 or f16
frexp(vecN<T>) -> __frexp_result_vecN_T where: T is f32 or f16
)");
}
@@ -991,8 +1058,8 @@ TEST_F(ResolverBuiltinFloatTest, Frexp_Error_SecondParamFloatPtr) {
R"(error: no matching call to frexp(f32, ptr<workgroup, f32, read_write>)
2 candidate functions:
frexp(f32) -> __frexp_result
frexp(vecN<f32>) -> __frexp_result_vecN
frexp(T) -> __frexp_result_T where: T is f32 or f16
frexp(vecN<T>) -> __frexp_result_vecN_T where: T is f32 or f16
)");
}
@@ -1005,8 +1072,8 @@ TEST_F(ResolverBuiltinFloatTest, Frexp_Error_SecondParamNotAPointer) {
EXPECT_EQ(r()->error(), R"(error: no matching call to frexp(f32, i32)
2 candidate functions:
frexp(f32) -> __frexp_result
frexp(vecN<f32>) -> __frexp_result_vecN
frexp(T) -> __frexp_result_T where: T is f32 or f16
frexp(vecN<T>) -> __frexp_result_vecN_T where: T is f32 or f16
)");
}
@@ -1021,8 +1088,8 @@ TEST_F(ResolverBuiltinFloatTest, Frexp_Error_VectorSizesDontMatch) {
R"(error: no matching call to frexp(vec2<f32>, ptr<workgroup, vec4<i32>, read_write>)
2 candidate functions:
frexp(vecN<f32>) -> __frexp_result_vecN
frexp(f32) -> __frexp_result
frexp(T) -> __frexp_result_T where: T is f32 or f16
frexp(vecN<T>) -> __frexp_result_vecN_T where: T is f32 or f16
)");
}
@@ -1129,8 +1196,9 @@ TEST_F(ResolverBuiltinFloatTest, Mix_VectorScalar_f16) {
EXPECT_TRUE(TypeOf(call)->As<sem::Vector>()->type()->Is<sem::F16>());
}
// modf: (f32) -> __modf_result, (vecN<f32>) -> __modf_result_vecN
TEST_F(ResolverBuiltinFloatTest, ModfScalar) {
// modf: (f32) -> __modf_result, (vecN<f32>) -> __modf_result_vecN, (f16) -> __modf_result_f16,
// (vecN<f16>) -> __modf_result_vecN_f16
TEST_F(ResolverBuiltinFloatTest, ModfScalar_f32) {
auto* call = Call("modf", 1_f);
WrapInFunction(call);
@@ -1159,7 +1227,38 @@ TEST_F(ResolverBuiltinFloatTest, ModfScalar) {
EXPECT_EQ(ty->SizeNoPadding(), 8u);
}
TEST_F(ResolverBuiltinFloatTest, ModfVector) {
TEST_F(ResolverBuiltinFloatTest, ModfScalar_f16) {
Enable(ast::Extension::kF16);
auto* call = Call("modf", 1_h);
WrapInFunction(call);
EXPECT_TRUE(r()->Resolve()) << r()->error();
ASSERT_NE(TypeOf(call), nullptr);
auto* ty = TypeOf(call)->As<sem::Struct>();
ASSERT_NE(ty, nullptr);
ASSERT_EQ(ty->Members().size(), 2u);
auto* fract = ty->Members()[0];
EXPECT_TRUE(fract->Type()->Is<sem::F16>());
EXPECT_EQ(fract->Offset(), 0u);
EXPECT_EQ(fract->Size(), 2u);
EXPECT_EQ(fract->Align(), 2u);
EXPECT_EQ(fract->Name(), Sym("fract"));
auto* whole = ty->Members()[1];
EXPECT_TRUE(whole->Type()->Is<sem::F16>());
EXPECT_EQ(whole->Offset(), 2u);
EXPECT_EQ(whole->Size(), 2u);
EXPECT_EQ(whole->Align(), 2u);
EXPECT_EQ(whole->Name(), Sym("whole"));
EXPECT_EQ(ty->Size(), 4u);
EXPECT_EQ(ty->SizeNoPadding(), 4u);
}
TEST_F(ResolverBuiltinFloatTest, ModfVector_f32) {
auto* call = Call("modf", vec3<f32>());
WrapInFunction(call);
@@ -1192,6 +1291,41 @@ TEST_F(ResolverBuiltinFloatTest, ModfVector) {
EXPECT_EQ(ty->SizeNoPadding(), 28u);
}
TEST_F(ResolverBuiltinFloatTest, ModfVector_f16) {
Enable(ast::Extension::kF16);
auto* call = Call("modf", vec3<f16>());
WrapInFunction(call);
EXPECT_TRUE(r()->Resolve()) << r()->error();
ASSERT_NE(TypeOf(call), nullptr);
auto* ty = TypeOf(call)->As<sem::Struct>();
ASSERT_NE(ty, nullptr);
ASSERT_EQ(ty->Members().size(), 2u);
auto* fract = ty->Members()[0];
ASSERT_TRUE(fract->Type()->Is<sem::Vector>());
EXPECT_EQ(fract->Type()->As<sem::Vector>()->Width(), 3u);
EXPECT_TRUE(fract->Type()->As<sem::Vector>()->type()->Is<sem::F16>());
EXPECT_EQ(fract->Offset(), 0u);
EXPECT_EQ(fract->Size(), 6u);
EXPECT_EQ(fract->Align(), 8u);
EXPECT_EQ(fract->Name(), Sym("fract"));
auto* whole = ty->Members()[1];
ASSERT_TRUE(whole->Type()->Is<sem::Vector>());
EXPECT_EQ(whole->Type()->As<sem::Vector>()->Width(), 3u);
EXPECT_TRUE(whole->Type()->As<sem::Vector>()->type()->Is<sem::F16>());
EXPECT_EQ(whole->Offset(), 8u);
EXPECT_EQ(whole->Size(), 6u);
EXPECT_EQ(whole->Align(), 8u);
EXPECT_EQ(whole->Name(), Sym("whole"));
EXPECT_EQ(ty->Size(), 16u);
EXPECT_EQ(ty->SizeNoPadding(), 14u);
}
TEST_F(ResolverBuiltinFloatTest, Modf_Error_FirstParamInt) {
GlobalVar("whole", ty.f32(), ast::StorageClass::kWorkgroup);
auto* call = Call("modf", 1_i, AddressOf("whole"));
@@ -1203,8 +1337,8 @@ TEST_F(ResolverBuiltinFloatTest, Modf_Error_FirstParamInt) {
R"(error: no matching call to modf(i32, ptr<workgroup, f32, read_write>)
2 candidate functions:
modf(f32) -> __modf_result
modf(vecN<f32>) -> __modf_result_vecN
modf(T) -> __modf_result_T where: T is f32 or f16
modf(vecN<T>) -> __modf_result_vecN_T where: T is f32 or f16
)");
}
@@ -1219,8 +1353,8 @@ TEST_F(ResolverBuiltinFloatTest, Modf_Error_SecondParamIntPtr) {
R"(error: no matching call to modf(f32, ptr<workgroup, i32, read_write>)
2 candidate functions:
modf(f32) -> __modf_result
modf(vecN<f32>) -> __modf_result_vecN
modf(T) -> __modf_result_T where: T is f32 or f16
modf(vecN<T>) -> __modf_result_vecN_T where: T is f32 or f16
)");
}
@@ -1233,8 +1367,8 @@ TEST_F(ResolverBuiltinFloatTest, Modf_Error_SecondParamNotAPointer) {
EXPECT_EQ(r()->error(), R"(error: no matching call to modf(f32, f32)
2 candidate functions:
modf(f32) -> __modf_result
modf(vecN<f32>) -> __modf_result_vecN
modf(T) -> __modf_result_T where: T is f32 or f16
modf(vecN<T>) -> __modf_result_vecN_T where: T is f32 or f16
)");
}
@@ -1249,8 +1383,8 @@ TEST_F(ResolverBuiltinFloatTest, Modf_Error_VectorSizesDontMatch) {
R"(error: no matching call to modf(vec2<f32>, ptr<workgroup, vec4<f32>, read_write>)
2 candidate functions:
modf(vecN<f32>) -> __modf_result_vecN
modf(f32) -> __modf_result
modf(T) -> __modf_result_T where: T is f32 or f16
modf(vecN<T>) -> __modf_result_vecN_T where: T is f32 or f16
)");
}

79
src/tint/resolver/intrinsic_table.cc
View File

@@ -732,24 +732,34 @@ const sem::ExternalTexture* build_texture_external(MatchState& state) {
// Builtin types starting with a _ prefix cannot be declared in WGSL, so they
// can only be used as return types. Because of this, they must only match Any,
// which is used as the return type matcher.
bool match_modf_result(const sem::Type* ty) {
return ty->Is<Any>();
}
bool match_modf_result_vec(const sem::Type* ty, Number& N) {
bool match_modf_result(const sem::Type* ty, const sem::Type*& T) {
if (!ty->Is<Any>()) {
return false;
}
N = Number::any;
T = ty;
return true;
}
bool match_frexp_result(const sem::Type* ty) {
return ty->Is<Any>();
}
bool match_frexp_result_vec(const sem::Type* ty, Number& N) {
bool match_modf_result_vec(const sem::Type* ty, Number& N, const sem::Type*& T) {
if (!ty->Is<Any>()) {
return false;
}
N = Number::any;
T = ty;
return true;
}
bool match_frexp_result(const sem::Type* ty, const sem::Type*& T) {
if (!ty->Is<Any>()) {
return false;
}
T = ty;
return true;
}
bool match_frexp_result_vec(const sem::Type* ty, Number& N, const sem::Type*& T) {
if (!ty->Is<Any>()) {
return false;
}
N = Number::any;
T = ty;
return true;
}
@@ -763,7 +773,7 @@ bool match_atomic_compare_exchange_result(const sem::Type* ty, const sem::Type*&
struct NameAndType {
std::string name;
sem::Type* type;
const sem::Type* type;
};
const sem::Struct* build_struct(MatchState& state,
std::string name,
@@ -797,27 +807,46 @@ const sem::Struct* build_struct(MatchState& state,
/* size_no_padding */ size_without_padding);
}
const sem::Struct* build_modf_result(MatchState& state) {
auto* f32 = state.builder.create<sem::F32>();
return build_struct(state, "__modf_result", {{"fract", f32}, {"whole", f32}});
const sem::Struct* build_modf_result(MatchState& state, const sem::Type* el) {
std::string display_name;
if (el->Is<sem::F16>()) {
display_name = "__modf_result_f16";
} else {
display_name = "__modf_result";
}
return build_struct(state, display_name, {{"fract", el}, {"whole", el}});
}
const sem::Struct* build_modf_result_vec(MatchState& state, Number& n) {
auto* vec_f32 = state.builder.create<sem::Vector>(state.builder.create<sem::F32>(), n.Value());
return build_struct(state, "__modf_result_vec" + std::to_string(n.Value()),
{{"fract", vec_f32}, {"whole", vec_f32}});
const sem::Struct* build_modf_result_vec(MatchState& state, Number& n, const sem::Type* el) {
std::string display_name;
if (el->Is<sem::F16>()) {
display_name = "__modf_result_vec" + std::to_string(n.Value()) + "_f16";
} else {
display_name = "__modf_result_vec" + std::to_string(n.Value());
}
auto* vec = state.builder.create<sem::Vector>(el, n.Value());
return build_struct(state, display_name, {{"fract", vec}, {"whole", vec}});
}
const sem::Struct* build_frexp_result(MatchState& state) {
auto* f32 = state.builder.create<sem::F32>();
const sem::Struct* build_frexp_result(MatchState& state, const sem::Type* el) {
std::string display_name;
if (el->Is<sem::F16>()) {
display_name = "__frexp_result_f16";
} else {
display_name = "__frexp_result";
}
auto* i32 = state.builder.create<sem::I32>();
return build_struct(state, "__frexp_result", {{"sig", f32}, {"exp", i32}});
return build_struct(state, display_name, {{"sig", el}, {"exp", i32}});
}
const sem::Struct* build_frexp_result_vec(MatchState& state, Number& n) {
auto* vec_f32 = state.builder.create<sem::Vector>(state.builder.create<sem::F32>(), n.Value());
const sem::Struct* build_frexp_result_vec(MatchState& state, Number& n, const sem::Type* el) {
std::string display_name;
if (el->Is<sem::F16>()) {
display_name = "__frexp_result_vec" + std::to_string(n.Value()) + "_f16";
} else {
display_name = "__frexp_result_vec" + std::to_string(n.Value());
}
auto* vec = state.builder.create<sem::Vector>(el, n.Value());
auto* vec_i32 = state.builder.create<sem::Vector>(state.builder.create<sem::I32>(), n.Value());
return build_struct(state, "__frexp_result_vec" + std::to_string(n.Value()),
{{"sig", vec_f32}, {"exp", vec_i32}});
return build_struct(state, display_name, {{"sig", vec}, {"exp", vec_i32}});
}
const sem::Struct* build_atomic_compare_exchange_result(MatchState& state, const sem::Type* ty) {
return build_struct(
state, "__atomic_compare_exchange_result" + ty->FriendlyName(state.builder.Symbols()),

1980
src/tint/resolver/intrinsic_table.inl
View File

File diff suppressed because it is too large Load Diff

102
src/tint/writer/glsl/generator_impl.cc
View File

@@ -1205,97 +1205,51 @@ bool GeneratorImpl::EmitDotCall(std::ostream& out,
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;
});
}

151
src/tint/writer/glsl/generator_impl_builtin_test.cc
View File

@@ -412,7 +412,7 @@ TEST_F(GlslGeneratorImplTest_Builtin, FMA_f16) {
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 @@ void main() {
)");
}
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 @@ void main() {
)");
}
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 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
)"));
}
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 @@ void main() {
)"));
}
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);

16
src/tint/writer/hlsl/generator_impl.cc
View File

@@ -1855,16 +1855,15 @@ bool GeneratorImpl::EmitModfCall(std::ostream& out,
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 @@ bool GeneratorImpl::EmitFrexpCall(std::ostream& out,
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,

128
src/tint/writer/hlsl/generator_impl_builtin_test.cc
View File

@@ -377,7 +377,7 @@ TEST_F(HlslGeneratorImplTest_Builtin, Select_Vector) {
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 @@ TEST_F(HlslGeneratorImplTest_Builtin, Modf_Scalar) {
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 @@ void test_function() {
)");
}
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 @@ TEST_F(HlslGeneratorImplTest_Builtin, Modf_Vector) {
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 @@ void test_function() {
)");
}
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 @@ void test_function() {
)");
}
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 @@ void test_function() {
)");
}
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);

12
src/tint/writer/msl/generator_impl.cc
View File

@@ -1308,9 +1308,9 @@ bool GeneratorImpl::EmitModfCall(std::ostream& out,
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 @@ bool GeneratorImpl::EmitFrexpCall(std::ostream& out,
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;
});
}

240
src/tint/writer/msl/generator_impl_builtin_test.cc
View File

@@ -405,6 +405,246 @@ TEST_F(MslGeneratorImplTest, WorkgroupBarrier) {
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);

122
src/tint/writer/spirv/builder_builtin_test.cc
View File

@@ -1651,7 +1651,7 @@ OpFunctionEnd
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 @@ OpFunctionEnd
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 @@ OpFunctionEnd
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