[msl] Add polyfill for quantizeToF16
Convert the argument to f16 and then back to f32.
Bug: 42251016
Change-Id: I7006e0ce91e34a8a2a372b5a1f5d6b24bc5b0c5a
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/195214
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index 0b75610..0d2e61b 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -99,6 +99,7 @@
case core::BuiltinFn::kDistance:
case core::BuiltinFn::kDot:
case core::BuiltinFn::kLength:
+ case core::BuiltinFn::kQuantizeToF16:
case core::BuiltinFn::kTextureDimensions:
case core::BuiltinFn::kTextureGather:
case core::BuiltinFn::kTextureGatherCompare:
@@ -173,6 +174,9 @@
case core::BuiltinFn::kLength:
Length(builtin);
break;
+ case core::BuiltinFn::kQuantizeToF16:
+ QuantizeToF16(builtin);
+ break;
// Texture builtins.
case core::BuiltinFn::kTextureDimensions:
@@ -367,6 +371,23 @@
builtin->Destroy();
}
+ /// Polyfill a quantizeToF16 call.
+ /// @param builtin the builtin call instruction
+ void QuantizeToF16(core::ir::CoreBuiltinCall* builtin) {
+ auto* arg = builtin->Args()[0];
+ auto* type_f32 = arg->Type();
+ const core::type::Type* type_f16 = ty.f16();
+ if (auto* vec = type_f32->As<core::type::Vector>()) {
+ type_f16 = ty.vec(ty.f16(), vec->Width());
+ }
+
+ // Convert the argument to f16 and then back again.
+ b.InsertBefore(builtin, [&] {
+ b.ConvertWithResult(builtin->DetachResult(), b.Convert(type_f16, arg));
+ });
+ builtin->Destroy();
+ }
+
/// Replace a textureDimensions call with the equivalent MSL intrinsics.
/// @param builtin the builtin call instruction
void TextureDimensions(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 d9bc55b..467bb5e 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,74 @@
EXPECT_EQ(expect, str());
}
+TEST_F(MslWriter_BuiltinPolyfillTest, QuantizeToF16_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<f32>(core::BuiltinFn::kQuantizeToF16, value);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%value:f32):f32 {
+ $B1: {
+ %3:f32 = quantizeToF16 %value
+ ret %3
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%value:f32):f32 {
+ $B1: {
+ %3:f16 = convert %value
+ %4:f32 = convert %3
+ ret %4
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, QuantizeToF16_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<vec4<f32>>(core::BuiltinFn::kQuantizeToF16, value);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%value:vec4<f32>):vec4<f32> {
+ $B1: {
+ %3:vec4<f32> = quantizeToF16 %value
+ ret %3
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%value:vec4<f32>):vec4<f32> {
+ $B1: {
+ %3:vec4<f16> = convert %value
+ %4:vec4<f32> = convert %3
+ ret %4
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
TEST_F(MslWriter_BuiltinPolyfillTest, TextureDimensions_1d) {
auto* t = b.FunctionParam(
"t", ty.Get<core::type::SampledTexture>(core::type::TextureDimension::k1d, ty.f32()));
diff --git a/test/tint/builtins/gen/var/quantizeToF16/12e50e.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/quantizeToF16/12e50e.wgsl.expected.ir.msl
index ec7fec1..60aadc8 100644
--- a/test/tint/builtins/gen/var/quantizeToF16/12e50e.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/quantizeToF16/12e50e.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device float* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ float prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+float quantizeToF16_12e50e() {
+ float arg_0 = 1.0f;
+ float res = float(half(arg_0));
+ return res;
}
-%quantizeToF16_12e50e = func():void {
- $B2: {
- %arg_0:ptr<function, f32, read_write> = var, 1.0f
- %4:f32 = load %arg_0
- %5:f32 = quantizeToF16 %4
- %res:ptr<function, f32, read_write> = var, %5
- %7:f32 = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %quantizeToF16_12e50e
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %quantizeToF16_12e50e
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %quantizeToF16_12e50e
- ret
- }
+fragment void fragment_main(device float* 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) = quantizeToF16_12e50e();
}
-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. *
-********************************************************************
+kernel void compute_main(device float* 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) = quantizeToF16_12e50e();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = quantizeToF16_12e50e();
+ 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};
+}
diff --git a/test/tint/builtins/gen/var/quantizeToF16/2cddf3.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/quantizeToF16/2cddf3.wgsl.expected.ir.msl
index 97979d6..bb54664 100644
--- a/test/tint/builtins/gen/var/quantizeToF16/2cddf3.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/quantizeToF16/2cddf3.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../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)
+struct tint_module_vars_struct {
+ device float2* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ float2 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float2 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+float2 quantizeToF16_2cddf3() {
+ float2 arg_0 = float2(1.0f);
+ float2 res = float2(half2(arg_0));
+ return res;
}
-%quantizeToF16_2cddf3 = func():void {
- $B2: {
- %arg_0:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
- %4:vec2<f32> = load %arg_0
- %5:vec2<f32> = quantizeToF16 %4
- %res:ptr<function, vec2<f32>, read_write> = var, %5
- %7:vec2<f32> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %quantizeToF16_2cddf3
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %quantizeToF16_2cddf3
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %quantizeToF16_2cddf3
- ret
- }
+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};
+ (*tint_module_vars.prevent_dce) = quantizeToF16_2cddf3();
}
-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. *
-********************************************************************
+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};
+ (*tint_module_vars.prevent_dce) = quantizeToF16_2cddf3();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = quantizeToF16_2cddf3();
+ 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};
+}
diff --git a/test/tint/builtins/gen/var/quantizeToF16/cba294.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/quantizeToF16/cba294.wgsl.expected.ir.msl
index f4fe649..643d6f8 100644
--- a/test/tint/builtins/gen/var/quantizeToF16/cba294.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/quantizeToF16/cba294.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device float4* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ float4 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float4 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+float4 quantizeToF16_cba294() {
+ float4 arg_0 = float4(1.0f);
+ float4 res = float4(half4(arg_0));
+ return res;
}
-%quantizeToF16_cba294 = func():void {
- $B2: {
- %arg_0:ptr<function, vec4<f32>, read_write> = var, vec4<f32>(1.0f)
- %4:vec4<f32> = load %arg_0
- %5:vec4<f32> = quantizeToF16 %4
- %res:ptr<function, vec4<f32>, read_write> = var, %5
- %7:vec4<f32> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %quantizeToF16_cba294
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %quantizeToF16_cba294
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %quantizeToF16_cba294
- ret
- }
+fragment void fragment_main(device float4* 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) = quantizeToF16_cba294();
}
-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. *
-********************************************************************
+kernel void compute_main(device float4* 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) = quantizeToF16_cba294();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = quantizeToF16_cba294();
+ 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};
+}
diff --git a/test/tint/builtins/gen/var/quantizeToF16/e8fd14.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/quantizeToF16/e8fd14.wgsl.expected.ir.msl
index 1e11ec8..236518a 100644
--- a/test/tint/builtins/gen/var/quantizeToF16/e8fd14.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/quantizeToF16/e8fd14.wgsl.expected.ir.msl
@@ -1,43 +1,44 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: { # root
- %prevent_dce:ptr<storage, vec3<f32>, read_write> = var @binding_point(2, 0)
+struct tint_module_vars_struct {
+ device float3* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ float3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ float3 VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+float3 quantizeToF16_e8fd14() {
+ float3 arg_0 = float3(1.0f);
+ float3 res = float3(half3(arg_0));
+ return res;
}
-%quantizeToF16_e8fd14 = func():void {
- $B2: {
- %arg_0:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
- %4:vec3<f32> = load %arg_0
- %5:vec3<f32> = quantizeToF16 %4
- %res:ptr<function, vec3<f32>, read_write> = var, %5
- %7:vec3<f32> = load %res
- store %prevent_dce, %7
- ret
- }
-}
-%vertex_main = @vertex func():vec4<f32> [@position] {
- $B3: {
- %9:void = call %quantizeToF16_e8fd14
- ret vec4<f32>(0.0f)
- }
-}
-%fragment_main = @fragment func():void {
- $B4: {
- %11:void = call %quantizeToF16_e8fd14
- ret
- }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
- $B5: {
- %13:void = call %quantizeToF16_e8fd14
- ret
- }
+fragment void fragment_main(device float3* 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) = quantizeToF16_e8fd14();
}
-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. *
-********************************************************************
+kernel void compute_main(device float3* 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) = quantizeToF16_e8fd14();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = quantizeToF16_e8fd14();
+ 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};
+}
