commit 2d70c1d55f54f209c25e8a6aed3cf55eb9d255ad
author James Price <jrprice@google.com> Tue Dec 19 00:57:45 2023 +0000
committer James Price <jrprice@google.com> Tue Dec 19 00:57:45 2023 +0000
tree 16d7ba2b376153e1fac3f208d67546af3ed87ccd
parent 884af6b070cbee6e32731ad6bf3746b30231aa32

[spirv-reader] Implement OpReturnValue

Bug: tint:1907
Change-Id: I839dd9cbea9d4b2c0a1c3e3a33c2e7a9c047b00e
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/165562
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

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

diff --git a/src/tint/lang/spirv/reader/parser/function_test.cc b/src/tint/lang/spirv/reader/parser/function_test.cc
index 15ef150..39e58b2 100644
--- a/src/tint/lang/spirv/reader/parser/function_test.cc
+++ b/src/tint/lang/spirv/reader/parser/function_test.cc
@@ -317,5 +317,82 @@
 )");
 }
 
+TEST_F(SpirvParserTest, FunctionCall_ReturnValue) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+       %bool = OpTypeBool
+       %true = OpConstantTrue %bool
+   %foo_type = OpTypeFunction %bool
+  %main_type = OpTypeFunction %void
+
+        %foo = OpFunction %bool None %foo_type
+  %foo_start = OpLabel
+               OpReturnValue %true
+               OpFunctionEnd
+
+       %main = OpFunction %void None %main_type
+ %main_start = OpLabel
+          %1 = OpFunctionCall %bool %foo
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+%1 = func():bool -> %b1 {
+  %b1 = block {
+    ret true
+  }
+}
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
+  %b2 = block {
+    %3:bool = call %1
+    ret
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, FunctionCall_ParamAndReturnValue) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+       %bool = OpTypeBool
+       %true = OpConstantTrue %bool
+   %foo_type = OpTypeFunction %bool %bool
+  %main_type = OpTypeFunction %void
+
+        %foo = OpFunction %bool None %foo_type
+      %param = OpFunctionParameter %bool
+  %foo_start = OpLabel
+               OpReturnValue %param
+               OpFunctionEnd
+
+       %main = OpFunction %void None %main_type
+ %main_start = OpLabel
+          %1 = OpFunctionCall %bool %foo %true
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+%1 = func(%2:bool):bool -> %b1 {
+  %b1 = block {
+    ret %2
+  }
+}
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
+  %b2 = block {
+    %4:bool = call %1, true
+    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 e5e04c7..b72a042 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -205,6 +205,9 @@
                 case spv::Op::OpReturn:
                     dst->Append(b_.Return(current_function_));
                     break;
+                case spv::Op::OpReturnValue:
+                    dst->Append(b_.Return(current_function_, Value(inst.GetSingleWordOperand(0))));
+                    break;
                 default:
                     TINT_UNIMPLEMENTED()
                         << "unhandled SPIR-V instruction: " << static_cast<uint32_t>(inst.opcode());