commit 323a821f6bd7ca7a567a01bb942327ae290b34f7
author dan sinclair <dsinclair@chromium.org> Tue Feb 11 05:49:27 2025 -0800
committer Dawn LUCI CQ <dawn-scoped@luci-project-accounts.iam.gserviceaccount.com> Tue Feb 11 05:49:27 2025 -0800
tree 5ffe37ba28449c9f9266a5f2e49f1d73bfea4534
parent 10486f7b64a3e01c5ae798653bb627021ea8d238

[spirv-reader][ir] Add support for `OpAll`

Converts an `OpAll` into a WGSL `all` instruction.

Bug: 391485778
Change-Id: I95eb7c2794489377f5aeefe140652678a8cb57ae
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/225856
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>

src/tint/lang/spirv/reader/parser/builtin_test.cc

diff --git a/src/tint/lang/spirv/reader/parser/builtin_test.cc b/src/tint/lang/spirv/reader/parser/builtin_test.cc
index 2574cb0..f3f38cd 100644
--- a/src/tint/lang/spirv/reader/parser/builtin_test.cc
+++ b/src/tint/lang/spirv/reader/parser/builtin_test.cc
@@ -1264,5 +1264,34 @@
 )");
 }
 
+TEST_F(SpirvParserTest, All) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+       %bool = OpTypeBool
+     %v2bool = OpTypeVector %bool 2
+       %true = OpConstantTrue %bool
+      %false = OpConstantFalse %bool
+ %v2bool_true_false = OpConstantComposite %v2bool %true %false
+    %ep_type = OpTypeFunction %void
+       %main = OpFunction %void None %ep_type
+      %entry = OpLabel
+          %1 = OpAll %bool %v2bool_true_false
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:bool = all vec2<bool>(true, false)
+    ret
+  }
+}
+)");
+}
+
 }  // namespace
 }  // namespace tint::spirv::reader

src/tint/lang/spirv/reader/parser/parser.cc

diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index ebf8b56..13b4813 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -640,7 +640,9 @@
                 case spv::Op::OpBitReverse:
                     EmitBuiltinCall(inst, core::BuiltinFn::kReverseBits);
                     break;
-
+                case spv::Op::OpAll:
+                    EmitBuiltinCall(inst, core::BuiltinFn::kAll);
+                    break;
                 default:
                     TINT_UNIMPLEMENTED()
                         << "unhandled SPIR-V instruction: " << static_cast<uint32_t>(inst.opcode());