msl-writer: support subgroupBroadcast
Update end-to-end subgroupBroadcast tests for MSL and WGSL
Bug: tint:2041
Change-Id: I9363bbce7d4f43ca3271d5d02e5cf980c54dbf98
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/152582
Auto-Submit: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: David Neto <dneto@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 8dced2c..e2fd37e 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -762,6 +762,20 @@
break;
}
+ case core::Function::kSubgroupBroadcast: {
+ // The lane argument is ushort.
+ out << "simd_broadcast(";
+ if (!EmitExpression(out, expr->args[0])) {
+ return false;
+ }
+ out << ",ushort(";
+ if (!EmitExpression(out, expr->args[1])) {
+ return false;
+ }
+ out << "))";
+ return true;
+ }
+
default:
break;
}
diff --git a/test/tint/builtins/gen/literal/subgroupBroadcast/08beca.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupBroadcast/08beca.wgsl.expected.msl
index e959257..d27a852 100644
--- a/test/tint/builtins/gen/literal/subgroupBroadcast/08beca.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/subgroupBroadcast/08beca.wgsl.expected.msl
@@ -1,18 +1,13 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-enable chromium_experimental_subgroups;
-
-fn subgroupBroadcast_08beca() {
- var res : f32 = subgroupBroadcast(1.0f, 1u);
- prevent_dce = res;
+using namespace metal;
+void subgroupBroadcast_08beca(device float* const tint_symbol) {
+ float res = simd_broadcast(1.0f,ushort(1u));
+ *(tint_symbol) = res;
}
-@group(2) @binding(0) var<storage, read_write> prevent_dce : f32;
-
-@compute @workgroup_size(1)
-fn compute_main() {
- subgroupBroadcast_08beca();
+kernel void compute_main(device float* tint_symbol_1 [[buffer(0)]]) {
+ subgroupBroadcast_08beca(tint_symbol_1);
+ return;
}
-Failed to generate: error: Unknown import method: subgroupBroadcast
diff --git a/test/tint/builtins/gen/literal/subgroupBroadcast/08beca.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupBroadcast/08beca.wgsl.expected.wgsl
index 665982c..6fb0ab3 100644
--- a/test/tint/builtins/gen/literal/subgroupBroadcast/08beca.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/subgroupBroadcast/08beca.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
enable chromium_experimental_subgroups;
fn subgroupBroadcast_08beca() {
diff --git a/test/tint/builtins/gen/literal/subgroupBroadcast/1d79c7.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupBroadcast/1d79c7.wgsl.expected.msl
index bd8e5d7..78f4e90 100644
--- a/test/tint/builtins/gen/literal/subgroupBroadcast/1d79c7.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/subgroupBroadcast/1d79c7.wgsl.expected.msl
@@ -1,18 +1,13 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-enable chromium_experimental_subgroups;
-
-fn subgroupBroadcast_1d79c7() {
- var res : i32 = subgroupBroadcast(1i, 1u);
- prevent_dce = res;
+using namespace metal;
+void subgroupBroadcast_1d79c7(device int* const tint_symbol) {
+ int res = simd_broadcast(1,ushort(1u));
+ *(tint_symbol) = res;
}
-@group(2) @binding(0) var<storage, read_write> prevent_dce : i32;
-
-@compute @workgroup_size(1)
-fn compute_main() {
- subgroupBroadcast_1d79c7();
+kernel void compute_main(device int* tint_symbol_1 [[buffer(0)]]) {
+ subgroupBroadcast_1d79c7(tint_symbol_1);
+ return;
}
-Failed to generate: error: Unknown import method: subgroupBroadcast
diff --git a/test/tint/builtins/gen/literal/subgroupBroadcast/1d79c7.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupBroadcast/1d79c7.wgsl.expected.wgsl
index e350734..9962b7f 100644
--- a/test/tint/builtins/gen/literal/subgroupBroadcast/1d79c7.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/subgroupBroadcast/1d79c7.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
enable chromium_experimental_subgroups;
fn subgroupBroadcast_1d79c7() {
diff --git a/test/tint/builtins/gen/literal/subgroupBroadcast/c36fe1.wgsl.expected.msl b/test/tint/builtins/gen/literal/subgroupBroadcast/c36fe1.wgsl.expected.msl
index 1b1de1d..4554b31 100644
--- a/test/tint/builtins/gen/literal/subgroupBroadcast/c36fe1.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/subgroupBroadcast/c36fe1.wgsl.expected.msl
@@ -1,18 +1,13 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-enable chromium_experimental_subgroups;
-
-fn subgroupBroadcast_c36fe1() {
- var res : u32 = subgroupBroadcast(1u, 1u);
- prevent_dce = res;
+using namespace metal;
+void subgroupBroadcast_c36fe1(device uint* const tint_symbol) {
+ uint res = simd_broadcast(1u,ushort(1u));
+ *(tint_symbol) = res;
}
-@group(2) @binding(0) var<storage, read_write> prevent_dce : u32;
-
-@compute @workgroup_size(1)
-fn compute_main() {
- subgroupBroadcast_c36fe1();
+kernel void compute_main(device uint* tint_symbol_1 [[buffer(0)]]) {
+ subgroupBroadcast_c36fe1(tint_symbol_1);
+ return;
}
-Failed to generate: error: Unknown import method: subgroupBroadcast
diff --git a/test/tint/builtins/gen/literal/subgroupBroadcast/c36fe1.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/subgroupBroadcast/c36fe1.wgsl.expected.wgsl
index a448f05..284ac16 100644
--- a/test/tint/builtins/gen/literal/subgroupBroadcast/c36fe1.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/subgroupBroadcast/c36fe1.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
enable chromium_experimental_subgroups;
fn subgroupBroadcast_c36fe1() {
diff --git a/test/tint/builtins/gen/var/subgroupBroadcast/08beca.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupBroadcast/08beca.wgsl.expected.msl
index 3bb9248..0ce7f74 100644
--- a/test/tint/builtins/gen/var/subgroupBroadcast/08beca.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/subgroupBroadcast/08beca.wgsl.expected.msl
@@ -1,20 +1,14 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-enable chromium_experimental_subgroups;
-
-fn subgroupBroadcast_08beca() {
- var arg_0 = 1.0f;
- const arg_1 = 1u;
- var res : f32 = subgroupBroadcast(arg_0, arg_1);
- prevent_dce = res;
+using namespace metal;
+void subgroupBroadcast_08beca(device float* const tint_symbol) {
+ float arg_0 = 1.0f;
+ float res = simd_broadcast(arg_0,ushort(1u));
+ *(tint_symbol) = res;
}
-@group(2) @binding(0) var<storage, read_write> prevent_dce : f32;
-
-@compute @workgroup_size(1)
-fn compute_main() {
- subgroupBroadcast_08beca();
+kernel void compute_main(device float* tint_symbol_1 [[buffer(0)]]) {
+ subgroupBroadcast_08beca(tint_symbol_1);
+ return;
}
-Failed to generate: error: Unknown import method: subgroupBroadcast
diff --git a/test/tint/builtins/gen/var/subgroupBroadcast/08beca.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupBroadcast/08beca.wgsl.expected.wgsl
new file mode 100644
index 0000000..0cf02e4
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBroadcast/08beca.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_experimental_subgroups;
+
+fn subgroupBroadcast_08beca() {
+ var arg_0 = 1.0f;
+ const arg_1 = 1u;
+ var res : f32 = subgroupBroadcast(arg_0, arg_1);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : f32;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBroadcast_08beca();
+}
diff --git a/test/tint/builtins/gen/var/subgroupBroadcast/1d79c7.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupBroadcast/1d79c7.wgsl.expected.msl
index c022db4..ea2106e 100644
--- a/test/tint/builtins/gen/var/subgroupBroadcast/1d79c7.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/subgroupBroadcast/1d79c7.wgsl.expected.msl
@@ -1,20 +1,14 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-enable chromium_experimental_subgroups;
-
-fn subgroupBroadcast_1d79c7() {
- var arg_0 = 1i;
- const arg_1 = 1u;
- var res : i32 = subgroupBroadcast(arg_0, arg_1);
- prevent_dce = res;
+using namespace metal;
+void subgroupBroadcast_1d79c7(device int* const tint_symbol) {
+ int arg_0 = 1;
+ int res = simd_broadcast(arg_0,ushort(1u));
+ *(tint_symbol) = res;
}
-@group(2) @binding(0) var<storage, read_write> prevent_dce : i32;
-
-@compute @workgroup_size(1)
-fn compute_main() {
- subgroupBroadcast_1d79c7();
+kernel void compute_main(device int* tint_symbol_1 [[buffer(0)]]) {
+ subgroupBroadcast_1d79c7(tint_symbol_1);
+ return;
}
-Failed to generate: error: Unknown import method: subgroupBroadcast
diff --git a/test/tint/builtins/gen/var/subgroupBroadcast/1d79c7.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupBroadcast/1d79c7.wgsl.expected.wgsl
new file mode 100644
index 0000000..fc70347
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBroadcast/1d79c7.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_experimental_subgroups;
+
+fn subgroupBroadcast_1d79c7() {
+ var arg_0 = 1i;
+ const arg_1 = 1u;
+ var res : i32 = subgroupBroadcast(arg_0, arg_1);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : i32;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBroadcast_1d79c7();
+}
diff --git a/test/tint/builtins/gen/var/subgroupBroadcast/c36fe1.wgsl.expected.msl b/test/tint/builtins/gen/var/subgroupBroadcast/c36fe1.wgsl.expected.msl
index c4a291c..c9a0fc3 100644
--- a/test/tint/builtins/gen/var/subgroupBroadcast/c36fe1.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/subgroupBroadcast/c36fe1.wgsl.expected.msl
@@ -1,20 +1,14 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-enable chromium_experimental_subgroups;
-
-fn subgroupBroadcast_c36fe1() {
- var arg_0 = 1u;
- const arg_1 = 1u;
- var res : u32 = subgroupBroadcast(arg_0, arg_1);
- prevent_dce = res;
+using namespace metal;
+void subgroupBroadcast_c36fe1(device uint* const tint_symbol) {
+ uint arg_0 = 1u;
+ uint res = simd_broadcast(arg_0,ushort(1u));
+ *(tint_symbol) = res;
}
-@group(2) @binding(0) var<storage, read_write> prevent_dce : u32;
-
-@compute @workgroup_size(1)
-fn compute_main() {
- subgroupBroadcast_c36fe1();
+kernel void compute_main(device uint* tint_symbol_1 [[buffer(0)]]) {
+ subgroupBroadcast_c36fe1(tint_symbol_1);
+ return;
}
-Failed to generate: error: Unknown import method: subgroupBroadcast
diff --git a/test/tint/builtins/gen/var/subgroupBroadcast/c36fe1.wgsl.expected.wgsl b/test/tint/builtins/gen/var/subgroupBroadcast/c36fe1.wgsl.expected.wgsl
new file mode 100644
index 0000000..6dd5c5e
--- /dev/null
+++ b/test/tint/builtins/gen/var/subgroupBroadcast/c36fe1.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+enable chromium_experimental_subgroups;
+
+fn subgroupBroadcast_c36fe1() {
+ var arg_0 = 1u;
+ const arg_1 = 1u;
+ var res : u32 = subgroupBroadcast(arg_0, arg_1);
+ prevent_dce = res;
+}
+
+@group(2) @binding(0) var<storage, read_write> prevent_dce : u32;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ subgroupBroadcast_c36fe1();
+}
