[msl] Polyfill unpack2x16float builtin
Bug: 42251016
Change-Id: I1143abc5ccc245674d349e46ae8fb9614f68bf67
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/193141
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/core/ir/builder.h b/src/tint/lang/core/ir/builder.h
index f85ec01..716a5a4 100644
--- a/src/tint/lang/core/ir/builder.h
+++ b/src/tint/lang/core/ir/builder.h
@@ -1091,6 +1091,16 @@
Values(std::forward<ARGS>(args)...));
}
+ /// Creates a value conversion instruction with an existing instruction result.
+ /// @param result the instruction result to use
+ /// @param val the value to be converted
+ /// @returns the instruction
+ template <typename VAL>
+ ir::Convert* ConvertWithResult(ir::InstructionResult* result, VAL&& val) {
+ return Append(
+ ir.allocators.instructions.Create<ir::Convert>(result, Value(std::forward<VAL>(val))));
+ }
+
/// Creates a value conversion instruction to the template type T
/// @param val the value to be converted
/// @returns the instruction
@@ -1106,8 +1116,7 @@
/// @returns the instruction
template <typename VAL>
ir::Convert* Convert(const core::type::Type* to, VAL&& val) {
- return Append(ir.allocators.instructions.Create<ir::Convert>(
- InstructionResult(to), Value(std::forward<VAL>(val))));
+ return ConvertWithResult(InstructionResult(to), Value(std::forward<VAL>(val)));
}
/// Creates a value constructor instruction with an existing instruction result
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index 89e38bc..09bd680 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -94,6 +94,7 @@
case core::BuiltinFn::kStorageBarrier:
case core::BuiltinFn::kWorkgroupBarrier:
case core::BuiltinFn::kTextureBarrier:
+ case core::BuiltinFn::kUnpack2X16Float:
worklist.Push(builtin);
break;
default:
@@ -165,6 +166,11 @@
ThreadgroupBarrier(builtin, BarrierType::kTexture);
break;
+ // Pack/unpack builtins.
+ case core::BuiltinFn::kUnpack2X16Float:
+ Unpack2x16Float(builtin);
+ break;
+
default:
break;
}
@@ -381,6 +387,17 @@
call->InsertBefore(builtin);
builtin->Destroy();
}
+
+ /// Polyfill an Unpack2x16Float call.
+ /// @param builtin the builtin call instruction
+ void Unpack2x16Float(core::ir::CoreBuiltinCall* builtin) {
+ // Replace the call with `float2(as_type<half2>(value))`.
+ b.InsertBefore(builtin, [&] {
+ auto* bitcast = b.Bitcast<vec2<f16>>(builtin->Args()[0]);
+ b.ConvertWithResult(builtin->DetachResult(), bitcast);
+ });
+ builtin->Destroy();
+ }
};
} // namespace
diff --git a/test/tint/builtins/gen/var/unpack2x16float/32a5cf.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/unpack2x16float/32a5cf.wgsl.expected.ir.msl
index 03355e9..e409ee9 100644
--- a/test/tint/builtins/gen/var/unpack2x16float/32a5cf.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/unpack2x16float/32a5cf.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+ device float2* prevent_dce;
+};
+struct vertex_main_outputs {
+ float4 tint_symbol [[position]];
+};
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec2<f32>, read_write> = var @binding_point(2, 0)
+void unpack2x16float_32a5cf(tint_module_vars_struct tint_module_vars) {
+ uint arg_0 = 1u;
+ float2 res = float2(as_type<half2>(arg_0));
+ (*tint_module_vars.prevent_dce) = res;
}
-
-%unpack2x16float_32a5cf = func():void {
- $B2: {
- %arg_0:ptr<function, u32, read_write> = var, 1u
- %4:u32 = load %arg_0
- %5:vec2<f32> = unpack2x16float %4
- %res:ptr<function, vec2<f32>, read_write> = var, %5
- %7:vec2<f32> = load %res
- store %prevent_dce, %7
- ret
- }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+ unpack2x16float_32a5cf(tint_module_vars);
+ return float4(0.0f);
}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %unpack2x16float_32a5cf
- ret vec4<f32>(0.0f)
- }
+fragment void fragment_main(device float2* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ unpack2x16float_32a5cf(tint_module_vars);
}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %unpack2x16float_32a5cf
- ret
- }
+kernel void compute_main(device float2* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ unpack2x16float_32a5cf(tint_module_vars);
}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %unpack2x16float_32a5cf
- ret
- }
+vertex vertex_main_outputs vertex_main(device float2* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
}
-
-unhandled variable address space
-********************************************************************
-* 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. *
-********************************************************************
