commit 0767dc6ae9a188e088cbf6b2dc973e5f727bec18
author dan sinclair <dsinclair@chromium.org> Thu Feb 06 08:23:55 2025 -0800
committer Dawn LUCI CQ <dawn-scoped@luci-project-accounts.iam.gserviceaccount.com> Thu Feb 06 08:23:55 2025 -0800
tree 8a5357536e16871c44220168e79f05fb6ac7cf7e
parent 9d2d2611c727360d6f3679d174b41441c4d8087f

[spirv-reader][ir] Support `OpUnreachable`

Translate an `OpUnreachable` SPIR-V instruction into an `unreachable` IR
instruction.

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

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

diff --git a/src/tint/lang/spirv/reader/parser/misc_test.cc b/src/tint/lang/spirv/reader/parser/misc_test.cc
index 71ce054..d814c5e 100644
--- a/src/tint/lang/spirv/reader/parser/misc_test.cc
+++ b/src/tint/lang/spirv/reader/parser/misc_test.cc
@@ -253,5 +253,130 @@
 )");
 }
 
+TEST_F(SpirvParserTest, OpUnreachable_TopLevel) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+    %ep_type = OpTypeFunction %void
+       %main = OpFunction %void None %ep_type
+         %10 = OpLabel
+               OpUnreachable
+               OpFunctionEnd
+  )",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    unreachable
+  }
+}
+)");
+}
+
+// TODO(dsinclair): Requires OpSelectionMerge and OpBranchConditional
+TEST_F(SpirvParserTest, DISABLED_OpUnreachable_InsideIf) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+       %bool = OpTypeBool
+    %ep_type = OpTypeFunction %void
+       %true = OpConstantTrue %bool
+       %main = OpFunction %void None %ep_type
+         %10 = OpLabel
+               OpSelectionMerge %99 None
+               OpBranchConditional %true %20 %99
+         %20 = OpLabel
+               OpUnreachable
+         %99 = OpLabel
+               OpReturn
+               OpFunctionEnd
+  )",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:void = if true [t: $B2] {
+      unreachable
+    }
+  }
+}
+)");
+}
+
+// TODO(dsinclair): Requires OpBranch
+TEST_F(SpirvParserTest, DISABLED_OpUnreachable_InsideLoop) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+       %bool = OpTypeBool
+    %ep_type = OpTypeFunction %void
+       %true = OpConstantTrue %bool
+       %main = OpFunction %void None %ep_type
+         %10 = OpLabel
+               OpBranch %20
+         %20 = OpLabel
+               OpLoopMerge %99 %80 None
+               OpBranchConditional %true %30 %30
+         %30 = OpLabel
+               OpUnreachable
+         %80 = OpLabel
+               OpBranch %20
+         %99 = OpLabel
+               OpReturn
+               OpFunctionEnd
+  )",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:void = loop %true []
+      unreachable
+    }
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, OpUnreachable_InNonVoidFunction) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+       %bool = OpTypeBool
+    %ep_type = OpTypeFunction %void
+     %boolfn = OpTypeFunction %bool
+        %200 = OpFunction %bool None %boolfn
+        %210 = OpLabel
+               OpUnreachable
+               OpFunctionEnd
+       %main = OpFunction %void None %ep_type
+         %10 = OpLabel
+         %11 = OpFunctionCall %bool %200
+               OpReturn
+               OpFunctionEnd
+  )",
+              R"(
+%1 = func():bool {
+  $B1: {
+    unreachable
+  }
+}
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B2: {
+    %3:bool = call %1
+    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 cd0684c..86edda6 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -614,6 +614,9 @@
                 case spv::Op::OpVariable:
                     EmitVar(inst);
                     break;
+                case spv::Op::OpUnreachable:
+                    EmitWithoutResult(b_.Unreachable());
+                    break;
                 default:
                     TINT_UNIMPLEMENTED()
                         << "unhandled SPIR-V instruction: " << static_cast<uint32_t>(inst.opcode());