[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>
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",
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
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",
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
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());
