commit ebda341c7b3c3bcc62ef959239a7e280ec6962da
author dan sinclair <dsinclair@chromium.org> Fri Jul 11 11:13:11 2025 -0700
committer Dawn LUCI CQ <dawn-scoped@luci-project-accounts.iam.gserviceaccount.com> Fri Jul 11 11:13:11 2025 -0700
tree b9d3ce596f9fc4044602f00723e2b2ea01b84f4e
parent 8b5299adc4daa7a06131bce2597ab9f852e97212

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

Add support to convert a `OpGroupNonUniformElect` into a `subgroupElect`
instruction.

Fixed: 431032556
Change-Id: I4413d2d0fec084283ad9448973eb37b764ae83fd
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/252076
Reviewed-by: David Neto <dneto@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 0b70ef4..86b560e 100644
--- a/src/tint/lang/spirv/reader/parser/builtin_test.cc
+++ b/src/tint/lang/spirv/reader/parser/builtin_test.cc
@@ -1759,5 +1759,36 @@
                   SPV_ENV_VULKAN_1_1);
 }
 
+TEST_F(SpirvParserTest, NonUniformElect) {
+    EXPECT_IR_SPV(R"(
+               OpCapability Shader
+               OpCapability GroupNonUniformVote
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %main "main"
+       %uint = OpTypeInt 32 0
+     %uint_3 = OpConstant %uint 3
+       %bool = OpTypeBool
+       %true = OpConstantTrue %bool
+       %void = OpTypeVoid
+         %23 = OpTypeFunction %void
+       %main = OpFunction %void None %23
+         %24 = OpLabel
+          %8 = OpGroupNonUniformElect %bool %uint_3
+               OpReturn
+               OpFunctionEnd
+)",
+                  R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:bool = subgroupElect
+    ret
+  }
+}
+)",
+                  SPV_ENV_VULKAN_1_1);
+}
+
 }  // 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 ac85dc6..9272c7e 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -2137,6 +2137,9 @@
                 case spv::Op::OpGroupNonUniformAny:
                     EmitSubgroupBuiltin(inst, core::BuiltinFn::kSubgroupAny);
                     break;
+                case spv::Op::OpGroupNonUniformElect:
+                    EmitSubgroupBuiltin(inst, core::BuiltinFn::kSubgroupElect);
+                    break;
                 default:
                     TINT_UNIMPLEMENTED()
                         << "unhandled SPIR-V instruction: " << static_cast<uint32_t>(inst.opcode());