[spirv-reader][ir] Support OpTranspose

Convert the SPIR-V `OpTranspose` instruction to a `transpose` builtin

Bug: 391485088
Change-Id: I4100f0d31a9f27f858d774594837897a65095e46
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/228014
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 a5ae877..f6af996 100644
--- a/src/tint/lang/spirv/reader/parser/builtin_test.cc
+++ b/src/tint/lang/spirv/reader/parser/builtin_test.cc
@@ -1410,5 +1410,122 @@
                                          BuiltinData{"OpDPdyCoarse", "dpdyCoarse"},
                                          BuiltinData{"OpFwidthCoarse", "fwidthCoarse"}));
 
+TEST_F(SpirvParserTest, Transpose_2x2) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+      %float = OpTypeFloat 32
+        %one = OpConstant %float 1
+        %two = OpConstant %float 2
+      %three = OpConstant %float 3
+    %v2float = OpTypeVector %float 2
+    %v3float = OpTypeVector %float 3
+  %m2v2float = OpTypeMatrix %v2float 2
+  %m2v3float = OpTypeMatrix %v3float 2
+  %m3v2float = OpTypeMatrix %v2float 3
+      %v2one = OpConstantComposite %v2float %one %one
+      %v2two = OpConstantComposite %v2float %two %two
+    %v3three = OpConstantComposite %v3float %three %three %three
+%m2v2_one_two = OpConstantComposite %m2v2float %v2one %v2two
+%m3v2_two_one_two = OpConstantComposite %m3v2float %v2two %v2one %v2two
+%m2v3_threee_three = OpConstantComposite %m2v3float %v3three %v3three
+    %void_fn = OpTypeFunction %void
+       %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+          %1 = OpTranspose %m2v2float %m2v2_one_two
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:mat2x2<f32> = transpose mat2x2<f32>(vec2<f32>(1.0f), vec2<f32>(2.0f))
+    ret
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Transpose_2x3) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+      %float = OpTypeFloat 32
+        %one = OpConstant %float 1
+        %two = OpConstant %float 2
+      %three = OpConstant %float 3
+    %v2float = OpTypeVector %float 2
+    %v3float = OpTypeVector %float 3
+  %m2v2float = OpTypeMatrix %v2float 2
+  %m2v3float = OpTypeMatrix %v3float 2
+  %m3v2float = OpTypeMatrix %v2float 3
+      %v2one = OpConstantComposite %v2float %one %one
+      %v2two = OpConstantComposite %v2float %two %two
+    %v3three = OpConstantComposite %v3float %three %three %three
+%m2v2_one_two = OpConstantComposite %m2v2float %v2one %v2two
+%m3v2_two_one_two = OpConstantComposite %m3v2float %v2two %v2one %v2two
+%m2v3_three_three = OpConstantComposite %m2v3float %v3three %v3three
+    %void_fn = OpTypeFunction %void
+       %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+          %1 = OpTranspose %m3v2float %m2v3_three_three
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:mat3x2<f32> = transpose mat2x3<f32>(vec3<f32>(3.0f))
+    ret
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Transpose_3x2) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+      %float = OpTypeFloat 32
+        %one = OpConstant %float 1
+        %two = OpConstant %float 2
+      %three = OpConstant %float 3
+    %v2float = OpTypeVector %float 2
+    %v3float = OpTypeVector %float 3
+  %m2v2float = OpTypeMatrix %v2float 2
+  %m2v3float = OpTypeMatrix %v3float 2
+  %m3v2float = OpTypeMatrix %v2float 3
+      %v2one = OpConstantComposite %v2float %one %one
+      %v2two = OpConstantComposite %v2float %two %two
+    %v3three = OpConstantComposite %v3float %three %three %three
+%m2v2_one_two = OpConstantComposite %m2v2float %v2one %v2two
+%m3v2_two_one_two = OpConstantComposite %m3v2float %v2two %v2one %v2two
+%m2v3_threee_three = OpConstantComposite %m2v3float %v3three %v3three
+    %void_fn = OpTypeFunction %void
+       %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+          %1 = OpTranspose %m2v3float %m3v2_two_one_two
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:mat2x3<f32> = transpose mat3x2<f32>(vec2<f32>(2.0f), vec2<f32>(1.0f), vec2<f32>(2.0f))
+    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 2bae38a..e06fd0e 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -976,6 +976,9 @@
                 case spv::Op::OpQuantizeToF16:
                     EmitBuiltinCall(inst, core::BuiltinFn::kQuantizeToF16);
                     break;
+                case spv::Op::OpTranspose:
+                    EmitBuiltinCall(inst, core::BuiltinFn::kTranspose);
+                    break;
                 default:
                     TINT_UNIMPLEMENTED()
                         << "unhandled SPIR-V instruction: " << static_cast<uint32_t>(inst.opcode());