[spirv-reader][ir] Add `OpDot` support.

Adds support for the SPIR-V `OpDot` instruction.

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

src/tint/lang/spirv/reader/parser/BUILD.bazel

diff --git a/src/tint/lang/spirv/reader/parser/BUILD.bazel b/src/tint/lang/spirv/reader/parser/BUILD.bazel
index 74fa26b..165aba1 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/parser/BUILD.bazel
@@ -82,6 +82,7 @@
   alwayslink = True,
   srcs = [
     "binary_test.cc",
+    "builtin_test.cc",
     "composite_test.cc",
     "constant_test.cc",
     "function_test.cc",

src/tint/lang/spirv/reader/parser/BUILD.cmake

diff --git a/src/tint/lang/spirv/reader/parser/BUILD.cmake b/src/tint/lang/spirv/reader/parser/BUILD.cmake
index a87f1eb..c9d1414 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/parser/BUILD.cmake
@@ -91,6 +91,7 @@
 ################################################################################
 tint_add_target(tint_lang_spirv_reader_parser_test test
   lang/spirv/reader/parser/binary_test.cc
+  lang/spirv/reader/parser/builtin_test.cc
   lang/spirv/reader/parser/composite_test.cc
   lang/spirv/reader/parser/constant_test.cc
   lang/spirv/reader/parser/function_test.cc

src/tint/lang/spirv/reader/parser/BUILD.gn

diff --git a/src/tint/lang/spirv/reader/parser/BUILD.gn b/src/tint/lang/spirv/reader/parser/BUILD.gn
index 63ce831..55175db 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.gn
+++ b/src/tint/lang/spirv/reader/parser/BUILD.gn
@@ -90,6 +90,7 @@
     tint_unittests_source_set("unittests") {
       sources = [
         "binary_test.cc",
+        "builtin_test.cc",
         "composite_test.cc",
         "constant_test.cc",
         "function_test.cc",

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
new file mode 100644
index 0000000..fb8cf46
--- /dev/null
+++ b/src/tint/lang/spirv/reader/parser/builtin_test.cc
@@ -0,0 +1,63 @@
+// Copyright 2025 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+//    list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+//    this list of conditions and the following disclaimer in the documentation
+//    and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+//    contributors may be used to endorse or promote products derived from
+//    this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/reader/parser/helper_test.h"
+
+namespace tint::spirv::reader {
+namespace {
+
+TEST_F(SpirvParserTest, Dot) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+    %ep_type = OpTypeFunction %void
+   %float_50 = OpConstant %float 50
+   %float_60 = OpConstant %float 60
+%v2float_50_60 = OpConstantComposite %v2float %float_50 %float_60
+%v2float_60_50 = OpConstantComposite %v2float %float_60 %float_50
+       %main = OpFunction %void None %ep_type
+      %entry = OpLabel
+          %1 = OpDot %float %v2float_50_60 %v2float_60_50
+               OpReturn
+               OpFunctionEnd)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:f32 = dot vec2<f32>(50.0f, 60.0f), vec2<f32>(60.0f, 50.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 b73f9e6..3ab1ffc 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -619,6 +619,12 @@
                 case spv::Op::OpKill:
                     EmitKill(inst);
                     break;
+                case spv::Op::OpDot:
+                    Emit(b_.Call(Type(inst.type_id()), core::BuiltinFn::kDot,
+                                 Value(inst.GetSingleWordOperand(2)),
+                                 Value(inst.GetSingleWordOperand(3))),
+                         inst.result_id());
+                    break;
                 default:
                     TINT_UNIMPLEMENTED()
                         << "unhandled SPIR-V instruction: " << static_cast<uint32_t>(inst.opcode());