Tint/MSL: Fix translation of dot4{I|U}8Packed
This patch fixes the translation of `dot4I8Packed()` and
`dot4U8Packed()` on MSL backend by replacing `packed_char4` and
`packed_uchar4` with `char4` and `uchar4`.
Bug: tint:1497
Change-Id: I82122de00d4339651eacb3a103bcdb4787e8ea90
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/165382
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
index a662d59..75c793b 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -1383,8 +1383,8 @@
const sem::BuiltinFn* builtin) {
return CallBuiltinHelper(
out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
- Line(b) << "packed_char4 vec1 = as_type<packed_char4>(" << params[0] << ");";
- Line(b) << "packed_char4 vec2 = as_type<packed_char4>(" << params[1] << ");";
+ Line(b) << "char4 vec1 = as_type<char4>(" << params[0] << ");";
+ Line(b) << "char4 vec2 = as_type<char4>(" << params[1] << ");";
Line(b) << "return vec1[0] * vec2[0] + vec1[1] * vec2[1] + vec1[2] * vec2[2] + vec1[3] "
"* vec2[3];";
return true;
@@ -1396,8 +1396,8 @@
const sem::BuiltinFn* builtin) {
return CallBuiltinHelper(
out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
- Line(b) << "packed_uchar4 vec1 = as_type<packed_uchar4>(" << params[0] << ");";
- Line(b) << "packed_uchar4 vec2 = as_type<packed_uchar4>(" << params[1] << ");";
+ Line(b) << "uchar4 vec1 = as_type<uchar4>(" << params[0] << ");";
+ Line(b) << "uchar4 vec2 = as_type<uchar4>(" << params[1] << ");";
Line(b) << "return vec1[0] * vec2[0] + vec1[1] * vec2[1] + vec1[2] * vec2[2] + vec1[3] "
"* vec2[3];";
return true;
diff --git a/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.msl b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.msl
index 82bb3cb..4529cff 100644
--- a/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/dot4I8Packed/881e62.wgsl.expected.msl
@@ -1,29 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-enable chromium_experimental_dp4a;
-
-fn dot4I8Packed_881e62() {
- var res : i32 = dot4I8Packed(1u, 1u);
+int tint_dot4I8Packed(uint param_0, uint param_1) {
+ char4 vec1 = as_type<char4>(param_0);
+ char4 vec2 = as_type<char4>(param_1);
+ return vec1[0] * vec2[0] + vec1[1] * vec2[1] + vec1[2] * vec2[2] + vec1[3] * vec2[3];
}
-@vertex
-fn vertex_main() -> @builtin(position) vec4<f32> {
- dot4I8Packed_881e62();
- return vec4<f32>();
+void dot4I8Packed_881e62(device int* const tint_symbol_1) {
+ int res = tint_dot4I8Packed(1u, 1u);
+ *(tint_symbol_1) = res;
}
-@fragment
-fn fragment_main() {
- dot4I8Packed_881e62();
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+float4 vertex_main_inner(device int* const tint_symbol_2) {
+ dot4I8Packed_881e62(tint_symbol_2);
+ return float4(0.0f);
}
-@compute @workgroup_size(1)
-fn compute_main() {
- dot4I8Packed_881e62();
+vertex tint_symbol vertex_main(device int* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
}
-Failed to generate: builtins/gen/literal/dot4I8Packed/881e62.wgsl:24:8 error: MSL backend does not support extension 'chromium_experimental_dp4a'
-enable chromium_experimental_dp4a;
- ^^^^^^^^^^^^^^^^^^^^^^^^^^
+fragment void fragment_main(device int* tint_symbol_4 [[buffer(0)]]) {
+ dot4I8Packed_881e62(tint_symbol_4);
+ return;
+}
+
+kernel void compute_main(device int* tint_symbol_5 [[buffer(0)]]) {
+ dot4I8Packed_881e62(tint_symbol_5);
+ return;
+}
diff --git a/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.msl b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.msl
index e1bd1eb..fbf0fab 100644
--- a/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/dot4U8Packed/fbed7b.wgsl.expected.msl
@@ -1,29 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-enable chromium_experimental_dp4a;
-
-fn dot4U8Packed_fbed7b() {
- var res : u32 = dot4U8Packed(1u, 1u);
+uint tint_dot4U8Packed(uint param_0, uint param_1) {
+ uchar4 vec1 = as_type<uchar4>(param_0);
+ uchar4 vec2 = as_type<uchar4>(param_1);
+ return vec1[0] * vec2[0] + vec1[1] * vec2[1] + vec1[2] * vec2[2] + vec1[3] * vec2[3];
}
-@vertex
-fn vertex_main() -> @builtin(position) vec4<f32> {
- dot4U8Packed_fbed7b();
- return vec4<f32>();
+void dot4U8Packed_fbed7b(device uint* const tint_symbol_1) {
+ uint res = tint_dot4U8Packed(1u, 1u);
+ *(tint_symbol_1) = res;
}
-@fragment
-fn fragment_main() {
- dot4U8Packed_fbed7b();
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+float4 vertex_main_inner(device uint* const tint_symbol_2) {
+ dot4U8Packed_fbed7b(tint_symbol_2);
+ return float4(0.0f);
}
-@compute @workgroup_size(1)
-fn compute_main() {
- dot4U8Packed_fbed7b();
+vertex tint_symbol vertex_main(device uint* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
}
-Failed to generate: builtins/gen/literal/dot4U8Packed/fbed7b.wgsl:24:8 error: MSL backend does not support extension 'chromium_experimental_dp4a'
-enable chromium_experimental_dp4a;
- ^^^^^^^^^^^^^^^^^^^^^^^^^^
+fragment void fragment_main(device uint* tint_symbol_4 [[buffer(0)]]) {
+ dot4U8Packed_fbed7b(tint_symbol_4);
+ return;
+}
+
+kernel void compute_main(device uint* tint_symbol_5 [[buffer(0)]]) {
+ dot4U8Packed_fbed7b(tint_symbol_5);
+ return;
+}
diff --git a/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.msl b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.msl
index 78cb718..f178b9e 100644
--- a/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.msl
@@ -1,31 +1,43 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-enable chromium_experimental_dp4a;
-
-fn dot4I8Packed_881e62() {
- var arg_0 = 1u;
- var arg_1 = 1u;
- var res : i32 = dot4I8Packed(arg_0, arg_1);
+int tint_dot4I8Packed(uint param_0, uint param_1) {
+ char4 vec1 = as_type<char4>(param_0);
+ char4 vec2 = as_type<char4>(param_1);
+ return vec1[0] * vec2[0] + vec1[1] * vec2[1] + vec1[2] * vec2[2] + vec1[3] * vec2[3];
}
-@vertex
-fn vertex_main() -> @builtin(position) vec4<f32> {
- dot4I8Packed_881e62();
- return vec4<f32>();
+void dot4I8Packed_881e62(device int* const tint_symbol_1) {
+ uint arg_0 = 1u;
+ uint arg_1 = 1u;
+ int res = tint_dot4I8Packed(arg_0, arg_1);
+ *(tint_symbol_1) = res;
}
-@fragment
-fn fragment_main() {
- dot4I8Packed_881e62();
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+float4 vertex_main_inner(device int* const tint_symbol_2) {
+ dot4I8Packed_881e62(tint_symbol_2);
+ return float4(0.0f);
}
-@compute @workgroup_size(1)
-fn compute_main() {
- dot4I8Packed_881e62();
+vertex tint_symbol vertex_main(device int* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
}
-Failed to generate: builtins/gen/var/dot4I8Packed/881e62.wgsl:24:8 error: MSL backend does not support extension 'chromium_experimental_dp4a'
-enable chromium_experimental_dp4a;
- ^^^^^^^^^^^^^^^^^^^^^^^^^^
+fragment void fragment_main(device int* tint_symbol_4 [[buffer(0)]]) {
+ dot4I8Packed_881e62(tint_symbol_4);
+ return;
+}
+
+kernel void compute_main(device int* tint_symbol_5 [[buffer(0)]]) {
+ dot4I8Packed_881e62(tint_symbol_5);
+ return;
+}
diff --git a/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.msl b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.msl
index be47b5c..880de13 100644
--- a/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.msl
@@ -1,31 +1,43 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-enable chromium_experimental_dp4a;
-
-fn dot4U8Packed_fbed7b() {
- var arg_0 = 1u;
- var arg_1 = 1u;
- var res : u32 = dot4U8Packed(arg_0, arg_1);
+uint tint_dot4U8Packed(uint param_0, uint param_1) {
+ uchar4 vec1 = as_type<uchar4>(param_0);
+ uchar4 vec2 = as_type<uchar4>(param_1);
+ return vec1[0] * vec2[0] + vec1[1] * vec2[1] + vec1[2] * vec2[2] + vec1[3] * vec2[3];
}
-@vertex
-fn vertex_main() -> @builtin(position) vec4<f32> {
- dot4U8Packed_fbed7b();
- return vec4<f32>();
+void dot4U8Packed_fbed7b(device uint* const tint_symbol_1) {
+ uint arg_0 = 1u;
+ uint arg_1 = 1u;
+ uint res = tint_dot4U8Packed(arg_0, arg_1);
+ *(tint_symbol_1) = res;
}
-@fragment
-fn fragment_main() {
- dot4U8Packed_fbed7b();
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+float4 vertex_main_inner(device uint* const tint_symbol_2) {
+ dot4U8Packed_fbed7b(tint_symbol_2);
+ return float4(0.0f);
}
-@compute @workgroup_size(1)
-fn compute_main() {
- dot4U8Packed_fbed7b();
+vertex tint_symbol vertex_main(device uint* tint_symbol_3 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
}
-Failed to generate: builtins/gen/var/dot4U8Packed/fbed7b.wgsl:24:8 error: MSL backend does not support extension 'chromium_experimental_dp4a'
-enable chromium_experimental_dp4a;
- ^^^^^^^^^^^^^^^^^^^^^^^^^^
+fragment void fragment_main(device uint* tint_symbol_4 [[buffer(0)]]) {
+ dot4U8Packed_fbed7b(tint_symbol_4);
+ return;
+}
+
+kernel void compute_main(device uint* tint_symbol_5 [[buffer(0)]]) {
+ dot4U8Packed_fbed7b(tint_symbol_5);
+ return;
+}
