[msl] Polyfill pack2x16float
Also add unit test for unpack2x16float which was missed when that
polyfill was added.
Bug: 42251016
Change-Id: Ifea9f38185d629395b11ade214625e661cda9c5d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/199897
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Natalie Chouinard <chouinard@google.com>
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index 5fda95e..b4c0759 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -101,6 +101,7 @@
case core::BuiltinFn::kFrexp:
case core::BuiltinFn::kLength:
case core::BuiltinFn::kModf:
+ case core::BuiltinFn::kPack2X16Float:
case core::BuiltinFn::kQuantizeToF16:
case core::BuiltinFn::kSign:
case core::BuiltinFn::kTextureDimensions:
@@ -246,6 +247,9 @@
break;
// Pack/unpack builtins.
+ case core::BuiltinFn::kPack2X16Float:
+ Pack2x16Float(builtin);
+ break;
case core::BuiltinFn::kUnpack2X16Float:
Unpack2x16Float(builtin);
break;
@@ -438,6 +442,18 @@
builtin->Destroy();
}
+ /// Polyfill an Pack2x16Float call.
+ /// @param builtin the builtin call instruction
+ void Pack2x16Float(core::ir::CoreBuiltinCall* builtin) {
+ // Replace the call with `as_type<uint>(half2(value))`.
+ b.InsertBefore(builtin, [&] {
+ auto* convert = b.Convert<vec2<f16>>(builtin->Args()[0]);
+ auto* bitcast = b.Bitcast(ty.u32(), convert);
+ bitcast->SetResults(Vector{builtin->DetachResult()});
+ });
+ 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 496f342..5c1a409 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -3213,5 +3213,73 @@
EXPECT_EQ(expect, str());
}
+TEST_F(MslWriter_BuiltinPolyfillTest, Pack2x16Float) {
+ auto* func = b.Function("foo", ty.u32());
+ auto* input = b.FunctionParam("input", ty.vec2<f32>());
+ func->SetParams(Vector{input});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call<u32>(core::BuiltinFn::kPack2X16Float, input);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%input:vec2<f32>):u32 {
+ $B1: {
+ %3:u32 = pack2x16float %input
+ ret %3
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%input:vec2<f32>):u32 {
+ $B1: {
+ %3:vec2<f16> = convert %input
+ %4:u32 = bitcast %3
+ ret %4
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, Unpack2x16Float) {
+ auto* func = b.Function("foo", ty.vec2<f32>());
+ auto* input = b.FunctionParam("input", ty.u32());
+ func->SetParams(Vector{input});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call<vec2<f32>>(core::BuiltinFn::kUnpack2X16Float, input);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%input:u32):vec2<f32> {
+ $B1: {
+ %3:vec2<f32> = unpack2x16float %input
+ ret %3
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%input:u32):vec2<f32> {
+ $B1: {
+ %3:vec2<f16> = bitcast %input
+ %4:vec2<f32> = convert %3
+ ret %4
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::msl::writer::raise
diff --git a/test/tint/builtins/gen/var/pack2x16float/0e97b3.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/pack2x16float/0e97b3.wgsl.expected.ir.msl
index 88b650d..b99130b0 100644
--- a/test/tint/builtins/gen/var/pack2x16float/0e97b3.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/pack2x16float/0e97b3.wgsl.expected.ir.msl
@@ -1,9 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: pack2x16float
-********************************************************************
-* 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 tint_module_vars_struct {
+ device uint* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ uint prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ uint VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+uint pack2x16float_0e97b3() {
+ float2 arg_0 = float2(1.0f);
+ uint res = as_type<uint>(half2(arg_0));
+ return res;
+}
+
+fragment void fragment_main(device uint* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = pack2x16float_0e97b3();
+}
+
+kernel void compute_main(device uint* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = pack2x16float_0e97b3();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = pack2x16float_0e97b3();
+ return out;
+}
+
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v.pos, .VertexOutput_prevent_dce=v.prevent_dce};
+}
