commit 45b4e0ab199d1fc95ca523bf48e81b39755dc1c7
author dan sinclair <dsinclair@chromium.org> Mon Jan 13 10:13:40 2025 -0800
committer Dawn LUCI CQ <dawn-scoped@luci-project-accounts.iam.gserviceaccount.com> Mon Jan 13 10:13:40 2025 -0800
tree fc513938491180933c5a32010d3076a9c1a86da6
parent e31cabd0ee0dfd26ccc06993137688c9dea5e0b4

[spirv-reader][ir] Support OpUndef in a function

Handle the SPIR-V `OpUndef` instruction inside a function by converting
it onto a zero `Value` object.

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

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 7d55b5a..6ef1a61 100644
--- a/src/tint/lang/spirv/reader/parser/misc_test.cc
+++ b/src/tint/lang/spirv/reader/parser/misc_test.cc
@@ -52,5 +52,67 @@
 )");
 }
 
+TEST_F(SpirvParserTest, Misc_OpUndefInFunction) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+       %bool = OpTypeBool
+       %uint = OpTypeInt 32 0
+        %int = OpTypeInt 32 1
+      %float = OpTypeFloat 32
+     %uint_2 = OpConstant %uint 2
+%_arr_uint_uint_2 = OpTypeArray %uint %uint_2
+     %v2uint = OpTypeVector %uint 2
+      %v2int = OpTypeVector %int 2
+    %v2float = OpTypeVector %float 2
+%mat2v2float = OpTypeMatrix %v2float 2
+ %_struct_13 = OpTypeStruct %bool %uint %int %float
+    %ep_type = OpTypeFunction %void
+         %11 = OpTypeFunction %int
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+          %1 = OpUndef %bool
+          %2 = OpUndef %uint
+          %3 = OpUndef %int
+          %4 = OpUndef %float
+          %5 = OpUndef %_arr_uint_uint_2
+          %6 = OpUndef %mat2v2float
+          %7 = OpUndef %v2uint
+          %8 = OpUndef %v2int
+          %9 = OpUndef %v2float
+         %10 = OpUndef %_struct_13
+               OpReturn
+               OpFunctionEnd
+          %m = OpFunction %int None %11
+         %12 = OpLabel
+         %13 = OpUndef %int
+               OpReturnValue %13
+               OpFunctionEnd
+
+)",
+              R"(
+tint_symbol_4 = struct @align(4) {
+  tint_symbol:bool @offset(0)
+  tint_symbol_1:u32 @offset(4)
+  tint_symbol_2:i32 @offset(8)
+  tint_symbol_3:f32 @offset(12)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    ret
+  }
+}
+%2 = func():i32 {
+  $B2: {
+    ret 0i
+  }
+}
+)");
+}
+
 }  // 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 7f5606b..418f7fe 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -504,6 +504,9 @@
             switch (inst.opcode()) {
                 case spv::Op::OpNop:
                     break;
+                case spv::Op::OpUndef:
+                    AddValue(inst.result_id(), b_.Zero(Type(inst.type_id())));
+                    break;
                 case spv::Op::OpAccessChain:
                 case spv::Op::OpInBoundsAccessChain:
                     EmitAccess(inst);