[msl] Add polyfill for modf
MSL's modf returns `fract` and outputs `whole` as an output
parameter. Polyfill it by declaring the result struct and then setting
the values.
Bug: 42251016
Change-Id: I20b1b44e5384948b4236cca8f563320d7e0f5b0d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/196359
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/msl/builtin_fn.cc b/src/tint/lang/msl/builtin_fn.cc
index e23f008..c522933 100644
--- a/src/tint/lang/msl/builtin_fn.cc
+++ b/src/tint/lang/msl/builtin_fn.cc
@@ -98,6 +98,8 @@
return "fmod";
case BuiltinFn::kLength:
return "length";
+ case BuiltinFn::kModf:
+ return "modf";
case BuiltinFn::kSign:
return "sign";
case BuiltinFn::kThreadgroupBarrier:
diff --git a/src/tint/lang/msl/builtin_fn.h b/src/tint/lang/msl/builtin_fn.h
index d28064f..8da3f08 100644
--- a/src/tint/lang/msl/builtin_fn.h
+++ b/src/tint/lang/msl/builtin_fn.h
@@ -75,6 +75,7 @@
kDot,
kFmod,
kLength,
+ kModf,
kSign,
kThreadgroupBarrier,
kNone,
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
index b73cce94..b60151f 100644
--- a/src/tint/lang/msl/intrinsic/data.cc
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -4870,13 +4870,20 @@
},
{
/* [28] */
+ /* fn modf[T : f32_f16](T, T) -> T */
+ /* fn modf[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> vec<N, T> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(156),
+ },
+ {
+ /* [29] */
/* fn sign[T : f32_f16](T) -> T */
/* fn sign[N : num, T : f32_f16](vec<N, T>) -> vec<N, T> */
/* num overloads */ 2,
/* overloads */ OverloadIndex(158),
},
{
- /* [29] */
+ /* [30] */
/* fn threadgroup_barrier(u32) */
/* num overloads */ 1,
/* overloads */ OverloadIndex(166),
diff --git a/src/tint/lang/msl/msl.def b/src/tint/lang/msl/msl.def
index 01ec90f..d2e8264 100644
--- a/src/tint/lang/msl/msl.def
+++ b/src/tint/lang/msl/msl.def
@@ -326,6 +326,8 @@
fn fmod[T: f32_f16](T, T) -> T
fn fmod[N: num, T: f32_f16](vec<N, T>, vec<N, T>) -> vec<N, T>
fn length[N: num, T: f32_f16](vec<N, T>) -> T
+fn modf[T: f32_f16](T, T) -> T
+fn modf[N: num, T: f32_f16](vec<N, T>, vec<N, T>) -> vec<N, T>
fn sign[T: f32_f16](T) -> T
fn sign[N: num, T: f32_f16](vec<N, T>) -> vec<N, T>
@stage("compute") fn threadgroup_barrier(u32)
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index a2426b8..0b7ce2b 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -100,6 +100,7 @@
case core::BuiltinFn::kDistance:
case core::BuiltinFn::kDot:
case core::BuiltinFn::kLength:
+ case core::BuiltinFn::kModf:
case core::BuiltinFn::kQuantizeToF16:
case core::BuiltinFn::kSign:
case core::BuiltinFn::kTextureDimensions:
@@ -181,6 +182,9 @@
case core::BuiltinFn::kLength:
Length(builtin);
break;
+ case core::BuiltinFn::kModf:
+ Modf(builtin);
+ break;
case core::BuiltinFn::kQuantizeToF16:
QuantizeToF16(builtin);
break;
@@ -386,6 +390,33 @@
builtin->Destroy();
}
+ /// Polyfill a modf call.
+ /// @param builtin the builtin call instruction
+ void Modf(core::ir::CoreBuiltinCall* builtin) {
+ b.InsertBefore(builtin, [&] {
+ // MSL's modf returns `fract` and outputs `whole` as an output parameter.
+ // Polyfill it by declaring the result struct and then setting the values:
+ // __modf_result result = {};
+ // result.fract = modf(arg, result.whole);
+ //
+ // Note: We need to use a `load` instruction to pass `result.whole`, as the intrinsic
+ // definition expects a value type (as we do not have reference types in the IR). The
+ // printer will just fold away the load, which achieves the pass-by-reference semantics
+ // that we want.
+ //
+ auto* result_type = builtin->Result(0)->Type();
+ auto* element_type = result_type->Element(0);
+ auto* result = b.Var(ty.ptr(function, result_type));
+ auto* whole = b.Access(ty.ptr(function, element_type), result, u32(1));
+ auto args = Vector<core::ir::Value*, 2>{builtin->Args()[0], b.Load(whole)->Result(0)};
+ auto* call =
+ b.Call<msl::ir::BuiltinCall>(element_type, msl::BuiltinFn::kModf, std::move(args));
+ b.Store(b.Access(ty.ptr(function, element_type), result, u32(0)), call);
+ builtin->Result(0)->ReplaceAllUsesWith(b.Load(result)->Result(0));
+ });
+ builtin->Destroy();
+ }
+
/// Polyfill a quantizeToF16 call.
/// @param builtin the builtin call instruction
void QuantizeToF16(core::ir::CoreBuiltinCall* builtin) {
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
index 41ca478..6537676 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -1088,6 +1088,122 @@
EXPECT_EQ(expect, str());
}
+TEST_F(MslWriter_BuiltinPolyfillTest, Modf_Scalar) {
+ auto* value = b.FunctionParam<f32>("value");
+ auto* func = b.Function("foo", ty.f32());
+ func->SetParams({value});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call(core::type::CreateModfResult(ty, mod.symbols, ty.f32()),
+ core::BuiltinFn::kModf, value);
+ auto* fract = b.Access<f32>(result, 0_u);
+ auto* whole = b.Access<f32>(result, 1_u);
+ b.Return(func, b.Add<f32>(fract, whole));
+ });
+
+ auto* src = R"(
+__modf_result_f32 = struct @align(4) {
+ fract:f32 @offset(0)
+ whole:f32 @offset(4)
+}
+
+%foo = func(%value:f32):f32 {
+ $B1: {
+ %3:__modf_result_f32 = modf %value
+ %4:f32 = access %3, 0u
+ %5:f32 = access %3, 1u
+ %6:f32 = add %4, %5
+ ret %6
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+__modf_result_f32 = struct @align(4) {
+ fract:f32 @offset(0)
+ whole:f32 @offset(4)
+}
+
+%foo = func(%value:f32):f32 {
+ $B1: {
+ %3:ptr<function, __modf_result_f32, read_write> = var
+ %4:ptr<function, f32, read_write> = access %3, 1u
+ %5:f32 = load %4
+ %6:f32 = msl.modf %value, %5
+ %7:ptr<function, f32, read_write> = access %3, 0u
+ store %7, %6
+ %8:__modf_result_f32 = load %3
+ %9:f32 = access %8, 0u
+ %10:f32 = access %8, 1u
+ %11:f32 = add %9, %10
+ ret %11
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, Modf_Vector) {
+ auto* value = b.FunctionParam<vec4<f32>>("value");
+ auto* func = b.Function("foo", ty.vec4<f32>());
+ func->SetParams({value});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call(core::type::CreateModfResult(ty, mod.symbols, ty.vec4<f32>()),
+ core::BuiltinFn::kModf, value);
+ auto* fract = b.Access<vec4<f32>>(result, 0_u);
+ auto* whole = b.Access<vec4<f32>>(result, 1_u);
+ b.Return(func, b.Add<vec4<f32>>(fract, whole));
+ });
+
+ auto* src = R"(
+__modf_result_vec4_f32 = struct @align(16) {
+ fract:vec4<f32> @offset(0)
+ whole:vec4<f32> @offset(16)
+}
+
+%foo = func(%value:vec4<f32>):vec4<f32> {
+ $B1: {
+ %3:__modf_result_vec4_f32 = modf %value
+ %4:vec4<f32> = access %3, 0u
+ %5:vec4<f32> = access %3, 1u
+ %6:vec4<f32> = add %4, %5
+ ret %6
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+__modf_result_vec4_f32 = struct @align(16) {
+ fract:vec4<f32> @offset(0)
+ whole:vec4<f32> @offset(16)
+}
+
+%foo = func(%value:vec4<f32>):vec4<f32> {
+ $B1: {
+ %3:ptr<function, __modf_result_vec4_f32, read_write> = var
+ %4:ptr<function, vec4<f32>, read_write> = access %3, 1u
+ %5:vec4<f32> = load %4
+ %6:vec4<f32> = msl.modf %value, %5
+ %7:ptr<function, vec4<f32>, read_write> = access %3, 0u
+ store %7, %6
+ %8:__modf_result_vec4_f32 = load %3
+ %9:vec4<f32> = access %8, 0u
+ %10:vec4<f32> = access %8, 1u
+ %11:vec4<f32> = add %9, %10
+ ret %11
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
TEST_F(MslWriter_BuiltinPolyfillTest, QuantizeToF16_Scalar) {
auto* value = b.FunctionParam<f32>("value");
auto* func = b.Function("foo", ty.f32());
diff --git a/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.ir.msl
index 9a4469c..6e9881b 100644
--- a/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.ir.msl
@@ -1,9 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: modf
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct modf_result_vec2_f32 {
+ float2 fract;
+ float2 whole;
+};
+
+struct VertexOutput {
+ float4 pos;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+};
+
+void modf_2d50da() {
+ float2 arg_0 = float2(-1.5f);
+ modf_result_vec2_f32 v = {};
+ v.fract = modf(arg_0, v.whole);
+ modf_result_vec2_f32 res = v;
+}
+
+fragment void fragment_main() {
+ modf_2d50da();
+}
+
+kernel void compute_main() {
+ modf_2d50da();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ modf_2d50da();
+ return out;
+}
+
+vertex vertex_main_outputs vertex_main() {
+ return vertex_main_outputs{.VertexOutput_pos=vertex_main_inner().pos};
+}
diff --git a/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.ir.msl
index 9a4469c..e8d3a98 100644
--- a/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.ir.msl
@@ -1,9 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: modf
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct modf_result_vec3_f16 {
+ half3 fract;
+ half3 whole;
+};
+
+struct VertexOutput {
+ float4 pos;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+};
+
+void modf_45005f() {
+ half3 arg_0 = half3(-1.5h);
+ modf_result_vec3_f16 v = {};
+ v.fract = modf(arg_0, v.whole);
+ modf_result_vec3_f16 res = v;
+}
+
+fragment void fragment_main() {
+ modf_45005f();
+}
+
+kernel void compute_main() {
+ modf_45005f();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ modf_45005f();
+ return out;
+}
+
+vertex vertex_main_outputs vertex_main() {
+ return vertex_main_outputs{.VertexOutput_pos=vertex_main_inner().pos};
+}
diff --git a/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.ir.msl
index 9a4469c..dfa275f 100644
--- a/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.ir.msl
@@ -1,9 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: modf
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct modf_result_vec4_f32 {
+ float4 fract;
+ float4 whole;
+};
+
+struct VertexOutput {
+ float4 pos;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+};
+
+void modf_4bfced() {
+ float4 arg_0 = float4(-1.5f);
+ modf_result_vec4_f32 v = {};
+ v.fract = modf(arg_0, v.whole);
+ modf_result_vec4_f32 res = v;
+}
+
+fragment void fragment_main() {
+ modf_4bfced();
+}
+
+kernel void compute_main() {
+ modf_4bfced();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ modf_4bfced();
+ return out;
+}
+
+vertex vertex_main_outputs vertex_main() {
+ return vertex_main_outputs{.VertexOutput_pos=vertex_main_inner().pos};
+}
diff --git a/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.ir.msl
index 9a4469c..4e3a867 100644
--- a/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.ir.msl
@@ -1,9 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: modf
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct modf_result_vec3_f32 {
+ float3 fract;
+ float3 whole;
+};
+
+struct VertexOutput {
+ float4 pos;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+};
+
+void modf_5ea256() {
+ float3 arg_0 = float3(-1.5f);
+ modf_result_vec3_f32 v = {};
+ v.fract = modf(arg_0, v.whole);
+ modf_result_vec3_f32 res = v;
+}
+
+fragment void fragment_main() {
+ modf_5ea256();
+}
+
+kernel void compute_main() {
+ modf_5ea256();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ modf_5ea256();
+ return out;
+}
+
+vertex vertex_main_outputs vertex_main() {
+ return vertex_main_outputs{.VertexOutput_pos=vertex_main_inner().pos};
+}
diff --git a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.ir.msl
index 9a4469c..bb92f98 100644
--- a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.ir.msl
@@ -1,9 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: modf
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct modf_result_f16 {
+ half fract;
+ half whole;
+};
+
+struct VertexOutput {
+ float4 pos;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+};
+
+void modf_8dbbbf() {
+ half arg_0 = -1.5h;
+ modf_result_f16 v = {};
+ v.fract = modf(arg_0, v.whole);
+ modf_result_f16 res = v;
+}
+
+fragment void fragment_main() {
+ modf_8dbbbf();
+}
+
+kernel void compute_main() {
+ modf_8dbbbf();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ modf_8dbbbf();
+ return out;
+}
+
+vertex vertex_main_outputs vertex_main() {
+ return vertex_main_outputs{.VertexOutput_pos=vertex_main_inner().pos};
+}
diff --git a/test/tint/builtins/gen/var/modf/995934.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/modf/995934.wgsl.expected.ir.msl
index 9a4469c..3f1f8a3 100644
--- a/test/tint/builtins/gen/var/modf/995934.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/modf/995934.wgsl.expected.ir.msl
@@ -1,9 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: modf
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct modf_result_vec4_f16 {
+ half4 fract;
+ half4 whole;
+};
+
+struct VertexOutput {
+ float4 pos;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+};
+
+void modf_995934() {
+ half4 arg_0 = half4(-1.5h);
+ modf_result_vec4_f16 v = {};
+ v.fract = modf(arg_0, v.whole);
+ modf_result_vec4_f16 res = v;
+}
+
+fragment void fragment_main() {
+ modf_995934();
+}
+
+kernel void compute_main() {
+ modf_995934();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ modf_995934();
+ return out;
+}
+
+vertex vertex_main_outputs vertex_main() {
+ return vertex_main_outputs{.VertexOutput_pos=vertex_main_inner().pos};
+}
diff --git a/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.ir.msl
index 9a4469c..8095cbc 100644
--- a/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.ir.msl
@@ -1,9 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: modf
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct modf_result_vec2_f16 {
+ half2 fract;
+ half2 whole;
+};
+
+struct VertexOutput {
+ float4 pos;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+};
+
+void modf_a545b9() {
+ half2 arg_0 = half2(-1.5h);
+ modf_result_vec2_f16 v = {};
+ v.fract = modf(arg_0, v.whole);
+ modf_result_vec2_f16 res = v;
+}
+
+fragment void fragment_main() {
+ modf_a545b9();
+}
+
+kernel void compute_main() {
+ modf_a545b9();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ modf_a545b9();
+ return out;
+}
+
+vertex vertex_main_outputs vertex_main() {
+ return vertex_main_outputs{.VertexOutput_pos=vertex_main_inner().pos};
+}
diff --git a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.ir.msl
index 9a4469c..68035c9 100644
--- a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.ir.msl
@@ -1,9 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: modf
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct modf_result_f32 {
+ float fract;
+ float whole;
+};
+
+struct VertexOutput {
+ float4 pos;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+};
+
+void modf_bbf7f7() {
+ float arg_0 = -1.5f;
+ modf_result_f32 v = {};
+ v.fract = modf(arg_0, v.whole);
+ modf_result_f32 res = v;
+}
+
+fragment void fragment_main() {
+ modf_bbf7f7();
+}
+
+kernel void compute_main() {
+ modf_bbf7f7();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ modf_bbf7f7();
+ return out;
+}
+
+vertex vertex_main_outputs vertex_main() {
+ return vertex_main_outputs{.VertexOutput_pos=vertex_main_inner().pos};
+}
diff --git a/test/tint/builtins/modf/scalar/mixed.wgsl.expected.ir.msl b/test/tint/builtins/modf/scalar/mixed.wgsl.expected.ir.msl
index 9a4469c..e51943a 100644
--- a/test/tint/builtins/modf/scalar/mixed.wgsl.expected.ir.msl
+++ b/test/tint/builtins/modf/scalar/mixed.wgsl.expected.ir.msl
@@ -1,9 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: modf
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct modf_result_f32 {
+ float fract;
+ float whole;
+};
+
+kernel void tint_symbol() {
+ float const runtime_in = 1.25f;
+ modf_result_f32 res = modf_result_f32{.fract=0.25f, .whole=1.0f};
+ modf_result_f32 v = {};
+ v.fract = modf(runtime_in, v.whole);
+ res = v;
+ res = modf_result_f32{.fract=0.25f, .whole=1.0f};
+ float const fract = res.fract;
+ float const whole = res.whole;
+}
diff --git a/test/tint/builtins/modf/scalar/runtime.wgsl.expected.ir.msl b/test/tint/builtins/modf/scalar/runtime.wgsl.expected.ir.msl
index 9a4469c..fb2b667 100644
--- a/test/tint/builtins/modf/scalar/runtime.wgsl.expected.ir.msl
+++ b/test/tint/builtins/modf/scalar/runtime.wgsl.expected.ir.msl
@@ -1,9 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: modf
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct modf_result_f32 {
+ float fract;
+ float whole;
+};
+
+kernel void tint_symbol() {
+ float const in = 1.25f;
+ modf_result_f32 v = {};
+ v.fract = modf(in, v.whole);
+ modf_result_f32 const res = v;
+ float const fract = res.fract;
+ float const whole = res.whole;
+}
diff --git a/test/tint/builtins/modf/vector/mixed.wgsl.expected.ir.msl b/test/tint/builtins/modf/vector/mixed.wgsl.expected.ir.msl
index 9a4469c..33408f9 100644
--- a/test/tint/builtins/modf/vector/mixed.wgsl.expected.ir.msl
+++ b/test/tint/builtins/modf/vector/mixed.wgsl.expected.ir.msl
@@ -1,9 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: modf
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct modf_result_vec2_f32 {
+ float2 fract;
+ float2 whole;
+};
+
+kernel void tint_symbol() {
+ float2 const runtime_in = float2(1.25f, 3.75f);
+ modf_result_vec2_f32 res = modf_result_vec2_f32{.fract=float2(0.25f, 0.75f), .whole=float2(1.0f, 3.0f)};
+ modf_result_vec2_f32 v = {};
+ v.fract = modf(runtime_in, v.whole);
+ res = v;
+ res = modf_result_vec2_f32{.fract=float2(0.25f, 0.75f), .whole=float2(1.0f, 3.0f)};
+ float2 const fract = res.fract;
+ float2 const whole = res.whole;
+}
diff --git a/test/tint/builtins/modf/vector/runtime.wgsl.expected.ir.msl b/test/tint/builtins/modf/vector/runtime.wgsl.expected.ir.msl
index 9a4469c..9cb2a79 100644
--- a/test/tint/builtins/modf/vector/runtime.wgsl.expected.ir.msl
+++ b/test/tint/builtins/modf/vector/runtime.wgsl.expected.ir.msl
@@ -1,9 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: modf
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct modf_result_vec2_f32 {
+ float2 fract;
+ float2 whole;
+};
+
+kernel void tint_symbol() {
+ float2 const in = float2(1.25f, 3.75f);
+ modf_result_vec2_f32 v = {};
+ v.fract = modf(in, v.whole);
+ modf_result_vec2_f32 const res = v;
+ float2 const fract = res.fract;
+ float2 const whole = res.whole;
+}
