instrinsics: Implement dot() for integer vector types
Fixed: tint:1263
Change-Id: I642ea0b6c9be7f04930cf6ea1a8059825661e326
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68520
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/docs/origin-trial-changes.md b/docs/origin-trial-changes.md
index 062a297..38f9279 100644
--- a/docs/origin-trial-changes.md
+++ b/docs/origin-trial-changes.md
@@ -6,6 +6,10 @@
* Taking the address of a vector component is no longer allowed.
+### New Features
+
+* The `dot()` builtin now supports integer vector types.
+
## Changes for M97
### Breaking Changes
diff --git a/src/intrinsic_table.inl b/src/intrinsic_table.inl
index 861fe97..22fa903 100644
--- a/src/intrinsic_table.inl
+++ b/src/intrinsic_table.inl
@@ -3780,12 +3780,12 @@
{
/* [398] */
/* usage */ ParameterUsage::kNone,
- /* matcher indices */ &kMatcherIndices[21],
+ /* matcher indices */ &kMatcherIndices[39],
},
{
/* [399] */
/* usage */ ParameterUsage::kNone,
- /* matcher indices */ &kMatcherIndices[21],
+ /* matcher indices */ &kMatcherIndices[39],
},
{
/* [400] */
@@ -7791,12 +7791,12 @@
{
/* [251] */
/* num parameters */ 2,
- /* num open types */ 0,
+ /* num open types */ 1,
/* num open numbers */ 1,
- /* open types */ &kOpenTypes[4],
+ /* open types */ &kOpenTypes[1],
/* open numbers */ &kOpenNumbers[3],
/* parameters */ &kParameters[398],
- /* return matcher indices */ &kMatcherIndices[12],
+ /* return matcher indices */ &kMatcherIndices[1],
/* supported_stages */ PipelineStageSet(PipelineStage::kVertex, PipelineStage::kFragment, PipelineStage::kCompute),
/* is_deprecated */ false,
},
@@ -8082,7 +8082,7 @@
},
{
/* [16] */
- /* fn dot<N : num>(vec<N, f32>, vec<N, f32>) -> f32 */
+ /* fn dot<N : num, T : fiu32>(vec<N, T>, vec<N, T>) -> T */
/* num overloads */ 1,
/* overloads */ &kOverloads[251],
},
diff --git a/src/intrinsics.def b/src/intrinsics.def
index 50fea3b..9f14296 100644
--- a/src/intrinsics.def
+++ b/src/intrinsics.def
@@ -292,7 +292,7 @@
fn determinant<N: num>(mat<N, N, f32>) -> f32
fn distance(f32, f32) -> f32
fn distance<N: num>(vec<N, f32>, vec<N, f32>) -> f32
-fn dot<N: num>(vec<N, f32>, vec<N, f32>) -> f32
+fn dot<N: num, T: fiu32>(vec<N, T>, vec<N, T>) -> T
[[stage("fragment")]] fn dpdx(f32) -> f32
[[stage("fragment")]] fn dpdx<N: num>(vec<N, f32>) -> vec<N, f32>
[[stage("fragment")]] fn dpdxCoarse(f32) -> f32
diff --git a/src/resolver/intrinsic_test.cc b/src/resolver/intrinsic_test.cc
index 305f001..8581765 100644
--- a/src/resolver/intrinsic_test.cc
+++ b/src/resolver/intrinsic_test.cc
@@ -339,7 +339,7 @@
}
TEST_F(ResolverIntrinsicTest, Dot_Vec3) {
- Global("my_var", ty.vec3<f32>(), ast::StorageClass::kPrivate);
+ Global("my_var", ty.vec3<i32>(), ast::StorageClass::kPrivate);
auto* expr = Call("dot", "my_var", "my_var");
WrapInFunction(expr);
@@ -347,11 +347,11 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
ASSERT_NE(TypeOf(expr), nullptr);
- EXPECT_TRUE(TypeOf(expr)->Is<sem::F32>());
+ EXPECT_TRUE(TypeOf(expr)->Is<sem::I32>());
}
TEST_F(ResolverIntrinsicTest, Dot_Vec4) {
- Global("my_var", ty.vec4<f32>(), ast::StorageClass::kPrivate);
+ Global("my_var", ty.vec4<u32>(), ast::StorageClass::kPrivate);
auto* expr = Call("dot", "my_var", "my_var");
WrapInFunction(expr);
@@ -359,7 +359,7 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
ASSERT_NE(TypeOf(expr), nullptr);
- EXPECT_TRUE(TypeOf(expr)->Is<sem::F32>());
+ EXPECT_TRUE(TypeOf(expr)->Is<sem::U32>());
}
TEST_F(ResolverIntrinsicTest, Dot_Error_Scalar) {
@@ -372,23 +372,7 @@
R"(error: no matching call to dot(f32, f32)
1 candidate function:
- dot(vecN<f32>, vecN<f32>) -> f32
-)");
-}
-
-TEST_F(ResolverIntrinsicTest, Dot_Error_VectorInt) {
- Global("my_var", ty.vec4<i32>(), ast::StorageClass::kPrivate);
-
- auto* expr = Call("dot", "my_var", "my_var");
- WrapInFunction(expr);
-
- EXPECT_FALSE(r()->Resolve());
-
- EXPECT_EQ(r()->error(),
- R"(error: no matching call to dot(vec4<i32>, vec4<i32>)
-
-1 candidate function:
- dot(vecN<f32>, vecN<f32>) -> f32
+ dot(vecN<T>, vecN<T>) -> T where: T is f32, i32 or u32
)");
}
diff --git a/src/writer/glsl/generator_impl.cc b/src/writer/glsl/generator_impl.cc
index 5a13b11..6e78944 100644
--- a/src/writer/glsl/generator_impl.cc
+++ b/src/writer/glsl/generator_impl.cc
@@ -107,7 +107,11 @@
size_t last_padding_line = 0;
line() << "#version 310 es";
- line() << "precision mediump float;" << std::endl;
+ line() << "precision mediump float;";
+
+ auto helpers_insertion_point = current_buffer_->lines.size();
+
+ line();
for (auto* decl : builder_.AST().GlobalDeclarations()) {
if (decl->Is<ast::Alias>()) {
@@ -153,7 +157,8 @@
}
if (!helpers_.lines.empty()) {
- current_buffer_->Insert(helpers_, 0, 0);
+ current_buffer_->Insert("", helpers_insertion_point++, 0);
+ current_buffer_->Insert(helpers_, helpers_insertion_point++, 0);
}
return true;
@@ -407,6 +412,8 @@
return EmitTextureCall(out, expr, intrinsic);
} else if (intrinsic->Type() == sem::IntrinsicType::kSelect) {
return EmitSelectCall(out, expr);
+ } else if (intrinsic->Type() == sem::IntrinsicType::kDot) {
+ return EmitDotCall(out, expr, intrinsic);
} else if (intrinsic->Type() == sem::IntrinsicType::kModf) {
return EmitModfCall(out, expr, intrinsic);
} else if (intrinsic->Type() == sem::IntrinsicType::kFrexp) {
@@ -671,6 +678,78 @@
return true;
}
+bool GeneratorImpl::EmitDotCall(std::ostream& out,
+ const ast::CallExpression* expr,
+ const sem::Intrinsic* intrinsic) {
+ auto* vec_ty = intrinsic->Parameters()[0]->Type()->As<sem::Vector>();
+ std::string fn = "dot";
+ if (vec_ty->type()->is_integer_scalar()) {
+ // GLSL does not have a builtin for dot() with integer vector types.
+ // Generate the helper function if it hasn't been created already
+ fn = utils::GetOrCreate(int_dot_funcs_, vec_ty, [&]() -> std::string {
+ TextBuffer b;
+ TINT_DEFER(helpers_.Append(b));
+
+ auto fn_name = UniqueIdentifier("tint_int_dot");
+
+ std::string v;
+ {
+ std::stringstream s;
+ if (!EmitType(s, vec_ty->type(), ast::StorageClass::kNone,
+ ast::Access::kRead, "")) {
+ return "";
+ }
+ v = s.str();
+ }
+ { // (u)int tint_int_dot([i|u]vecN a, [i|u]vecN b) {
+ auto l = line(&b);
+ if (!EmitType(l, vec_ty->type(), ast::StorageClass::kNone,
+ ast::Access::kRead, "")) {
+ return "";
+ }
+ l << " " << fn_name << "(";
+ if (!EmitType(l, vec_ty, ast::StorageClass::kNone, ast::Access::kRead,
+ "")) {
+ return "";
+ }
+ l << " a, ";
+ if (!EmitType(l, vec_ty, ast::StorageClass::kNone, ast::Access::kRead,
+ "")) {
+ return "";
+ }
+ l << " b) {";
+ }
+ {
+ auto l = line(&b);
+ l << " return ";
+ for (uint32_t i = 0; i < vec_ty->Width(); i++) {
+ if (i > 0) {
+ l << " + ";
+ }
+ l << "a[" << i << "]*b[" << i << "]";
+ }
+ l << ";";
+ }
+ line(&b) << "}";
+ return fn_name;
+ });
+ if (fn.empty()) {
+ return false;
+ }
+ }
+
+ out << fn << "(";
+ if (!EmitExpression(out, expr->args[0])) {
+ return false;
+ }
+ out << ", ";
+ if (!EmitExpression(out, expr->args[1])) {
+ return false;
+ }
+ out << ")";
+ return true;
+}
+
bool GeneratorImpl::EmitModfCall(std::ostream& out,
const ast::CallExpression* expr,
const sem::Intrinsic* intrinsic) {
@@ -2216,9 +2295,6 @@
}
if (auto* c = stmt->As<ast::CallStatement>()) {
auto out = line();
- if (!TypeOf(c->expr)->Is<sem::Void>()) {
- out << "(void) ";
- }
if (!EmitCall(out, c->expr)) {
return false;
}
diff --git a/src/writer/glsl/generator_impl.h b/src/writer/glsl/generator_impl.h
index ba82fa8..7539da5 100644
--- a/src/writer/glsl/generator_impl.h
+++ b/src/writer/glsl/generator_impl.h
@@ -136,6 +136,14 @@
/// @param expr the call expression
/// @returns true if the call expression is emitted
bool EmitSelectCall(std::ostream& out, const ast::CallExpression* expr);
+ /// Handles generating a call to the `dot()` intrinsic
+ /// @param out the output of the expression stream
+ /// @param expr the call expression
+ /// @param intrinsic the semantic information for the intrinsic
+ /// @returns true if the call expression is emitted
+ bool EmitDotCall(std::ostream& out,
+ const ast::CallExpression* expr,
+ const sem::Intrinsic* intrinsic);
/// Handles generating a call to the `modf()` intrinsic
/// @param out the output of the expression stream
/// @param expr the call expression
@@ -416,6 +424,7 @@
std::unordered_map<const sem::Intrinsic*, std::string> intrinsics_;
std::unordered_map<const sem::Struct*, std::string> structure_builders_;
std::unordered_map<const sem::Vector*, std::string> dynamic_vector_write_;
+ std::unordered_map<const sem::Vector*, std::string> int_dot_funcs_;
};
} // namespace glsl
diff --git a/src/writer/glsl/generator_impl_intrinsic_test.cc b/src/writer/glsl/generator_impl_intrinsic_test.cc
index 4d47310..7891beb 100644
--- a/src/writer/glsl/generator_impl_intrinsic_test.cc
+++ b/src/writer/glsl/generator_impl_intrinsic_test.cc
@@ -599,6 +599,64 @@
}
#endif
+TEST_F(GlslGeneratorImplTest_Intrinsic, DotI32) {
+ Global("v", ty.vec3<i32>(), ast::StorageClass::kPrivate);
+ WrapInFunction(Call("dot", "v", "v"));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#version 310 es
+precision mediump float;
+
+int tint_int_dot(ivec3 a, ivec3 b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2];
+}
+
+ivec3 v = ivec3(0, 0, 0);
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void test_function() {
+ tint_int_dot(v, v);
+ return;
+}
+void main() {
+ test_function();
+}
+
+
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Intrinsic, DotU32) {
+ Global("v", ty.vec3<u32>(), ast::StorageClass::kPrivate);
+ WrapInFunction(Call("dot", "v", "v"));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#version 310 es
+precision mediump float;
+
+uint tint_int_dot(uvec3 a, uvec3 b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2];
+}
+
+uvec3 v = uvec3(0u, 0u, 0u);
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void test_function() {
+ tint_int_dot(v, v);
+ return;
+}
+void main() {
+ test_function();
+}
+
+
+)");
+}
+
} // namespace
} // namespace glsl
} // namespace writer
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index 83071cc..2d7ad03 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -543,6 +543,8 @@
auto name = generate_builtin_name(intrinsic);
switch (intrinsic->Type()) {
+ case sem::IntrinsicType::kDot:
+ return EmitDotCall(out, expr, intrinsic);
case sem::IntrinsicType::kModf:
return EmitModfCall(out, expr, intrinsic);
case sem::IntrinsicType::kFrexp:
@@ -1005,6 +1007,53 @@
return true;
}
+bool GeneratorImpl::EmitDotCall(std::ostream& out,
+ const ast::CallExpression* expr,
+ const sem::Intrinsic* intrinsic) {
+ auto* vec_ty = intrinsic->Parameters()[0]->Type()->As<sem::Vector>();
+ std::string fn = "dot";
+ if (vec_ty->type()->is_integer_scalar()) {
+ // MSL does not have a builtin for dot() with integer vector types.
+ // Generate the helper function if it hasn't been created already
+ fn = utils::GetOrCreate(
+ int_dot_funcs_, vec_ty->Width(), [&]() -> std::string {
+ TextBuffer b;
+ TINT_DEFER(helpers_.Append(b));
+
+ auto fn_name =
+ UniqueIdentifier("tint_dot" + std::to_string(vec_ty->Width()));
+ auto v = "vec<T," + std::to_string(vec_ty->Width()) + ">";
+
+ line(&b) << "template<typename T>";
+ line(&b) << "T " << fn_name << "(" << v << " a, " << v << " b) {";
+ {
+ auto l = line(&b);
+ l << " return ";
+ for (uint32_t i = 0; i < vec_ty->Width(); i++) {
+ if (i > 0) {
+ l << " + ";
+ }
+ l << "a[" << i << "]*b[" << i << "]";
+ }
+ l << ";";
+ }
+ line(&b) << "}";
+ return fn_name;
+ });
+ }
+
+ out << fn << "(";
+ if (!EmitExpression(out, expr->args[0])) {
+ return false;
+ }
+ out << ", ";
+ if (!EmitExpression(out, expr->args[1])) {
+ return false;
+ }
+ out << ")";
+ return true;
+}
+
bool GeneratorImpl::EmitModfCall(std::ostream& out,
const ast::CallExpression* expr,
const sem::Intrinsic* intrinsic) {
diff --git a/src/writer/msl/generator_impl.h b/src/writer/msl/generator_impl.h
index 6009c92..5c68d31 100644
--- a/src/writer/msl/generator_impl.h
+++ b/src/writer/msl/generator_impl.h
@@ -154,6 +154,14 @@
bool EmitTextureCall(std::ostream& out,
const ast::CallExpression* expr,
const sem::Intrinsic* intrinsic);
+ /// Handles generating a call to the `dot()` intrinsic
+ /// @param out the output of the expression stream
+ /// @param expr the call expression
+ /// @param intrinsic the semantic information for the intrinsic
+ /// @returns true if the call expression is emitted
+ bool EmitDotCall(std::ostream& out,
+ const ast::CallExpression* expr,
+ const sem::Intrinsic* intrinsic);
/// Handles generating a call to the `modf()` intrinsic
/// @param out the output of the expression stream
/// @param expr the call expression
@@ -394,6 +402,7 @@
std::unordered_map<const sem::Intrinsic*, std::string> intrinsics_;
std::unordered_map<const sem::Type*, std::string> unary_minus_funcs_;
+ std::unordered_map<uint32_t, std::string> int_dot_funcs_;
};
} // namespace msl
diff --git a/src/writer/msl/generator_impl_intrinsic_test.cc b/src/writer/msl/generator_impl_intrinsic_test.cc
index 89bb44a..1f5c78e 100644
--- a/src/writer/msl/generator_impl_intrinsic_test.cc
+++ b/src/writer/msl/generator_impl_intrinsic_test.cc
@@ -343,6 +343,30 @@
EXPECT_EQ(out.str(), "float2(as_type<half2>(p1))");
}
+TEST_F(MslGeneratorImplTest, DotI32) {
+ Global("v", ty.vec3<i32>(), ast::StorageClass::kPrivate);
+ WrapInFunction(Call("dot", "v", "v"));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T>
+T tint_dot3(vec<T,3> a, vec<T,3> b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2];
+}
+kernel void test_function() {
+ thread int3 tint_symbol = 0;
+ tint_dot3(tint_symbol, tint_symbol);
+ return;
+}
+
+)");
+}
+
TEST_F(MslGeneratorImplTest, Ignore) {
Func("f", {Param("a", ty.i32()), Param("b", ty.i32()), Param("c", ty.i32())},
ty.i32(), {Return(Mul(Add("a", "b"), "c"))});
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc
index 621838a..d0340e9 100644
--- a/src/writer/spirv/builder.cc
+++ b/src/writer/spirv/builder.cc
@@ -2423,9 +2423,48 @@
case IntrinsicType::kCountOneBits:
op = spv::Op::OpBitCount;
break;
- case IntrinsicType::kDot:
+ case IntrinsicType::kDot: {
op = spv::Op::OpDot;
+ auto* vec_ty = intrinsic->Parameters()[0]->Type()->As<sem::Vector>();
+ if (vec_ty->type()->is_integer_scalar()) {
+ // TODO(crbug.com/tint/1267): OpDot requires floating-point types, but
+ // WGSL also supports integer types. SPV_KHR_integer_dot_product adds
+ // support for integer vectors. Use it if it is available.
+ auto el_ty = Operand::Int(GenerateTypeIfNeeded(vec_ty->type()));
+ auto vec_a = Operand::Int(get_arg_as_value_id(0));
+ auto vec_b = Operand::Int(get_arg_as_value_id(1));
+ if (vec_a.to_i() == 0 || vec_b.to_i() == 0) {
+ return 0;
+ }
+
+ auto sum = Operand::Int(0);
+ for (uint32_t i = 0; i < vec_ty->Width(); i++) {
+ auto a = result_op();
+ auto b = result_op();
+ auto mul = result_op();
+ if (!push_function_inst(spv::Op::OpCompositeExtract,
+ {el_ty, a, vec_a, Operand::Int(i)}) ||
+ !push_function_inst(spv::Op::OpCompositeExtract,
+ {el_ty, b, vec_b, Operand::Int(i)}) ||
+ !push_function_inst(spv::Op::OpIMul, {el_ty, mul, a, b})) {
+ return 0;
+ }
+ if (i == 0) {
+ sum = mul;
+ } else {
+ auto prev_sum = sum;
+ auto is_last_el = i == (vec_ty->Width() - 1);
+ sum = is_last_el ? Operand::Int(result_id) : result_op();
+ if (!push_function_inst(spv::Op::OpIAdd,
+ {el_ty, sum, prev_sum, mul})) {
+ return 0;
+ }
+ }
+ }
+ return result_id;
+ }
break;
+ }
case IntrinsicType::kDpdx:
op = spv::Op::OpDPdx;
break;
diff --git a/src/writer/spirv/builder_intrinsic_test.cc b/src/writer/spirv/builder_intrinsic_test.cc
index b69404e..34ac6af 100644
--- a/src/writer/spirv/builder_intrinsic_test.cc
+++ b/src/writer/spirv/builder_intrinsic_test.cc
@@ -407,7 +407,7 @@
testing::Values(IntrinsicData{"countOneBits", "OpBitCount"},
IntrinsicData{"reverseBits", "OpBitReverse"}));
-TEST_F(IntrinsicBuilderTest, Call_Dot) {
+TEST_F(IntrinsicBuilderTest, Call_Dot_F32) {
auto* var = Global("v", ty.vec3<f32>(), ast::StorageClass::kPrivate);
auto* expr = Call("dot", "v", "v");
@@ -432,6 +432,76 @@
)");
}
+TEST_F(IntrinsicBuilderTest, Call_Dot_U32) {
+ auto* var = Global("v", ty.vec3<u32>(), ast::StorageClass::kPrivate);
+
+ auto* expr = Call("dot", "v", "v");
+ WrapInFunction(expr);
+
+ spirv::Builder& b = Build();
+
+ b.push_function(Function{});
+ ASSERT_TRUE(b.GenerateGlobalVariable(var)) << b.error();
+
+ EXPECT_EQ(b.GenerateCallExpression(expr), 6u) << b.error();
+ EXPECT_EQ(DumpInstructions(b.types()), R"(%4 = OpTypeInt 32 0
+%3 = OpTypeVector %4 3
+%2 = OpTypePointer Private %3
+%5 = OpConstantNull %3
+%1 = OpVariable %2 Private %5
+)");
+ EXPECT_EQ(DumpInstructions(b.functions()[0].instructions()),
+ R"(%7 = OpLoad %3 %1
+%8 = OpLoad %3 %1
+%9 = OpCompositeExtract %4 %7 0
+%10 = OpCompositeExtract %4 %8 0
+%11 = OpIMul %4 %9 %10
+%12 = OpCompositeExtract %4 %7 1
+%13 = OpCompositeExtract %4 %8 1
+%14 = OpIMul %4 %12 %13
+%15 = OpIAdd %4 %11 %14
+%16 = OpCompositeExtract %4 %7 2
+%17 = OpCompositeExtract %4 %8 2
+%18 = OpIMul %4 %16 %17
+%6 = OpIAdd %4 %15 %18
+)");
+}
+
+TEST_F(IntrinsicBuilderTest, Call_Dot_I32) {
+ auto* var = Global("v", ty.vec3<i32>(), ast::StorageClass::kPrivate);
+
+ auto* expr = Call("dot", "v", "v");
+ WrapInFunction(expr);
+
+ spirv::Builder& b = Build();
+
+ b.push_function(Function{});
+ ASSERT_TRUE(b.GenerateGlobalVariable(var)) << b.error();
+
+ EXPECT_EQ(b.GenerateCallExpression(expr), 6u) << b.error();
+ EXPECT_EQ(DumpInstructions(b.types()), R"(%4 = OpTypeInt 32 1
+%3 = OpTypeVector %4 3
+%2 = OpTypePointer Private %3
+%5 = OpConstantNull %3
+%1 = OpVariable %2 Private %5
+)");
+ EXPECT_EQ(DumpInstructions(b.functions()[0].instructions()),
+ R"(%7 = OpLoad %3 %1
+%8 = OpLoad %3 %1
+%9 = OpCompositeExtract %4 %7 0
+%10 = OpCompositeExtract %4 %8 0
+%11 = OpIMul %4 %9 %10
+%12 = OpCompositeExtract %4 %7 1
+%13 = OpCompositeExtract %4 %8 1
+%14 = OpIMul %4 %12 %13
+%15 = OpIAdd %4 %11 %14
+%16 = OpCompositeExtract %4 %7 2
+%17 = OpCompositeExtract %4 %8 2
+%18 = OpIMul %4 %16 %17
+%6 = OpIAdd %4 %15 %18
+)");
+}
+
using IntrinsicDeriveTest = IntrinsicBuilderTestWithParam<IntrinsicData>;
TEST_P(IntrinsicDeriveTest, Call_Derivative_Scalar) {
auto param = GetParam();
diff --git a/test/intrinsics/gen/dot/7548a0.wgsl b/test/intrinsics/gen/dot/7548a0.wgsl
new file mode 100644
index 0000000..62d39d6
--- /dev/null
+++ b/test/intrinsics/gen/dot/7548a0.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+// test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+// src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn dot(vec<3, u32>, vec<3, u32>) -> u32
+fn dot_7548a0() {
+ var res: u32 = dot(vec3<u32>(), vec3<u32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+ dot_7548a0();
+ return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+ dot_7548a0();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+ dot_7548a0();
+}
diff --git a/test/intrinsics/gen/dot/7548a0.wgsl.expected.hlsl b/test/intrinsics/gen/dot/7548a0.wgsl.expected.hlsl
new file mode 100644
index 0000000..87fc5ae
--- /dev/null
+++ b/test/intrinsics/gen/dot/7548a0.wgsl.expected.hlsl
@@ -0,0 +1,30 @@
+void dot_7548a0() {
+ uint res = dot(uint3(0u, 0u, 0u), uint3(0u, 0u, 0u));
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ dot_7548a0();
+ return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ dot_7548a0();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ dot_7548a0();
+ return;
+}
diff --git a/test/intrinsics/gen/dot/7548a0.wgsl.expected.msl b/test/intrinsics/gen/dot/7548a0.wgsl.expected.msl
new file mode 100644
index 0000000..fb3daab
--- /dev/null
+++ b/test/intrinsics/gen/dot/7548a0.wgsl.expected.msl
@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T>
+T tint_dot3(vec<T,3> a, vec<T,3> b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2];
+}
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+void dot_7548a0() {
+ uint res = tint_dot3(uint3(), uint3());
+}
+
+float4 vertex_main_inner() {
+ dot_7548a0();
+ return float4();
+}
+
+vertex tint_symbol vertex_main() {
+ float4 const inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+fragment void fragment_main() {
+ dot_7548a0();
+ return;
+}
+
+kernel void compute_main() {
+ dot_7548a0();
+ return;
+}
+
diff --git a/test/intrinsics/gen/dot/7548a0.wgsl.expected.spvasm b/test/intrinsics/gen/dot/7548a0.wgsl.expected.spvasm
new file mode 100644
index 0000000..7d10efb
--- /dev/null
+++ b/test/intrinsics/gen/dot/7548a0.wgsl.expected.spvasm
@@ -0,0 +1,78 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 44
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %dot_7548a0 "dot_7548a0"
+ OpName %res "res"
+ OpName %vertex_main_inner "vertex_main_inner"
+ OpName %vertex_main "vertex_main"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+ %16 = OpConstantNull %v3uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %29 = OpConstantNull %uint
+ %30 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
+ %dot_7548a0 = OpFunction %void None %9
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_uint Function %29
+ %17 = OpCompositeExtract %uint %16 0
+ %18 = OpCompositeExtract %uint %16 0
+ %19 = OpIMul %uint %17 %18
+ %20 = OpCompositeExtract %uint %16 1
+ %21 = OpCompositeExtract %uint %16 1
+ %22 = OpIMul %uint %20 %21
+ %23 = OpIAdd %uint %19 %22
+ %24 = OpCompositeExtract %uint %16 2
+ %25 = OpCompositeExtract %uint %16 2
+ %26 = OpIMul %uint %24 %25
+ %13 = OpIAdd %uint %23 %26
+ OpStore %res %13
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %30
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %dot_7548a0
+ OpReturnValue %5
+ OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+ %35 = OpLabel
+ %36 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %36
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+ %39 = OpLabel
+ %40 = OpFunctionCall %void %dot_7548a0
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %9
+ %42 = OpLabel
+ %43 = OpFunctionCall %void %dot_7548a0
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/gen/dot/7548a0.wgsl.expected.wgsl b/test/intrinsics/gen/dot/7548a0.wgsl.expected.wgsl
new file mode 100644
index 0000000..27b8265
--- /dev/null
+++ b/test/intrinsics/gen/dot/7548a0.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn dot_7548a0() {
+ var res : u32 = dot(vec3<u32>(), vec3<u32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+ dot_7548a0();
+ return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+ dot_7548a0();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+ dot_7548a0();
+}
diff --git a/test/intrinsics/gen/dot/97c7ee.wgsl b/test/intrinsics/gen/dot/97c7ee.wgsl
new file mode 100644
index 0000000..d2e7e5d
--- /dev/null
+++ b/test/intrinsics/gen/dot/97c7ee.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+// test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+// src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn dot(vec<2, u32>, vec<2, u32>) -> u32
+fn dot_97c7ee() {
+ var res: u32 = dot(vec2<u32>(), vec2<u32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+ dot_97c7ee();
+ return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+ dot_97c7ee();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+ dot_97c7ee();
+}
diff --git a/test/intrinsics/gen/dot/97c7ee.wgsl.expected.hlsl b/test/intrinsics/gen/dot/97c7ee.wgsl.expected.hlsl
new file mode 100644
index 0000000..b262398
--- /dev/null
+++ b/test/intrinsics/gen/dot/97c7ee.wgsl.expected.hlsl
@@ -0,0 +1,30 @@
+void dot_97c7ee() {
+ uint res = dot(uint2(0u, 0u), uint2(0u, 0u));
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ dot_97c7ee();
+ return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ dot_97c7ee();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ dot_97c7ee();
+ return;
+}
diff --git a/test/intrinsics/gen/dot/97c7ee.wgsl.expected.msl b/test/intrinsics/gen/dot/97c7ee.wgsl.expected.msl
new file mode 100644
index 0000000..d288445
--- /dev/null
+++ b/test/intrinsics/gen/dot/97c7ee.wgsl.expected.msl
@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T>
+T tint_dot2(vec<T,2> a, vec<T,2> b) {
+ return a[0]*b[0] + a[1]*b[1];
+}
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+void dot_97c7ee() {
+ uint res = tint_dot2(uint2(), uint2());
+}
+
+float4 vertex_main_inner() {
+ dot_97c7ee();
+ return float4();
+}
+
+vertex tint_symbol vertex_main() {
+ float4 const inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+fragment void fragment_main() {
+ dot_97c7ee();
+ return;
+}
+
+kernel void compute_main() {
+ dot_97c7ee();
+ return;
+}
+
diff --git a/test/intrinsics/gen/dot/97c7ee.wgsl.expected.spvasm b/test/intrinsics/gen/dot/97c7ee.wgsl.expected.spvasm
new file mode 100644
index 0000000..6cc49f91
--- /dev/null
+++ b/test/intrinsics/gen/dot/97c7ee.wgsl.expected.spvasm
@@ -0,0 +1,74 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 40
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %dot_97c7ee "dot_97c7ee"
+ OpName %res "res"
+ OpName %vertex_main_inner "vertex_main_inner"
+ OpName %vertex_main "vertex_main"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %v2uint = OpTypeVector %uint 2
+ %16 = OpConstantNull %v2uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %25 = OpConstantNull %uint
+ %26 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
+ %dot_97c7ee = OpFunction %void None %9
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_uint Function %25
+ %17 = OpCompositeExtract %uint %16 0
+ %18 = OpCompositeExtract %uint %16 0
+ %19 = OpIMul %uint %17 %18
+ %20 = OpCompositeExtract %uint %16 1
+ %21 = OpCompositeExtract %uint %16 1
+ %22 = OpIMul %uint %20 %21
+ %13 = OpIAdd %uint %19 %22
+ OpStore %res %13
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %26
+ %28 = OpLabel
+ %29 = OpFunctionCall %void %dot_97c7ee
+ OpReturnValue %5
+ OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+ %31 = OpLabel
+ %32 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %32
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+ %35 = OpLabel
+ %36 = OpFunctionCall %void %dot_97c7ee
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %9
+ %38 = OpLabel
+ %39 = OpFunctionCall %void %dot_97c7ee
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/gen/dot/97c7ee.wgsl.expected.wgsl b/test/intrinsics/gen/dot/97c7ee.wgsl.expected.wgsl
new file mode 100644
index 0000000..7741f7d
--- /dev/null
+++ b/test/intrinsics/gen/dot/97c7ee.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn dot_97c7ee() {
+ var res : u32 = dot(vec2<u32>(), vec2<u32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+ dot_97c7ee();
+ return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+ dot_97c7ee();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+ dot_97c7ee();
+}
diff --git a/test/intrinsics/gen/dot/e994c7.wgsl b/test/intrinsics/gen/dot/e994c7.wgsl
new file mode 100644
index 0000000..8ffe117
--- /dev/null
+++ b/test/intrinsics/gen/dot/e994c7.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+// test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+// src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn dot(vec<4, u32>, vec<4, u32>) -> u32
+fn dot_e994c7() {
+ var res: u32 = dot(vec4<u32>(), vec4<u32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+ dot_e994c7();
+ return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+ dot_e994c7();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+ dot_e994c7();
+}
diff --git a/test/intrinsics/gen/dot/e994c7.wgsl.expected.hlsl b/test/intrinsics/gen/dot/e994c7.wgsl.expected.hlsl
new file mode 100644
index 0000000..66f8a0d
--- /dev/null
+++ b/test/intrinsics/gen/dot/e994c7.wgsl.expected.hlsl
@@ -0,0 +1,30 @@
+void dot_e994c7() {
+ uint res = dot(uint4(0u, 0u, 0u, 0u), uint4(0u, 0u, 0u, 0u));
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ dot_e994c7();
+ return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ dot_e994c7();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ dot_e994c7();
+ return;
+}
diff --git a/test/intrinsics/gen/dot/e994c7.wgsl.expected.msl b/test/intrinsics/gen/dot/e994c7.wgsl.expected.msl
new file mode 100644
index 0000000..70fa0e3
--- /dev/null
+++ b/test/intrinsics/gen/dot/e994c7.wgsl.expected.msl
@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T>
+T tint_dot4(vec<T,4> a, vec<T,4> b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
+}
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+void dot_e994c7() {
+ uint res = tint_dot4(uint4(), uint4());
+}
+
+float4 vertex_main_inner() {
+ dot_e994c7();
+ return float4();
+}
+
+vertex tint_symbol vertex_main() {
+ float4 const inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+fragment void fragment_main() {
+ dot_e994c7();
+ return;
+}
+
+kernel void compute_main() {
+ dot_e994c7();
+ return;
+}
+
diff --git a/test/intrinsics/gen/dot/e994c7.wgsl.expected.spvasm b/test/intrinsics/gen/dot/e994c7.wgsl.expected.spvasm
new file mode 100644
index 0000000..d852f99
--- /dev/null
+++ b/test/intrinsics/gen/dot/e994c7.wgsl.expected.spvasm
@@ -0,0 +1,82 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %dot_e994c7 "dot_e994c7"
+ OpName %res "res"
+ OpName %vertex_main_inner "vertex_main_inner"
+ OpName %vertex_main "vertex_main"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %v4uint = OpTypeVector %uint 4
+ %16 = OpConstantNull %v4uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %33 = OpConstantNull %uint
+ %34 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
+ %dot_e994c7 = OpFunction %void None %9
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_uint Function %33
+ %17 = OpCompositeExtract %uint %16 0
+ %18 = OpCompositeExtract %uint %16 0
+ %19 = OpIMul %uint %17 %18
+ %20 = OpCompositeExtract %uint %16 1
+ %21 = OpCompositeExtract %uint %16 1
+ %22 = OpIMul %uint %20 %21
+ %23 = OpIAdd %uint %19 %22
+ %24 = OpCompositeExtract %uint %16 2
+ %25 = OpCompositeExtract %uint %16 2
+ %26 = OpIMul %uint %24 %25
+ %27 = OpIAdd %uint %23 %26
+ %28 = OpCompositeExtract %uint %16 3
+ %29 = OpCompositeExtract %uint %16 3
+ %30 = OpIMul %uint %28 %29
+ %13 = OpIAdd %uint %27 %30
+ OpStore %res %13
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+ %36 = OpLabel
+ %37 = OpFunctionCall %void %dot_e994c7
+ OpReturnValue %5
+ OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+ %39 = OpLabel
+ %40 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %40
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+ %43 = OpLabel
+ %44 = OpFunctionCall %void %dot_e994c7
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %9
+ %46 = OpLabel
+ %47 = OpFunctionCall %void %dot_e994c7
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/gen/dot/e994c7.wgsl.expected.wgsl b/test/intrinsics/gen/dot/e994c7.wgsl.expected.wgsl
new file mode 100644
index 0000000..365e0ce
--- /dev/null
+++ b/test/intrinsics/gen/dot/e994c7.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn dot_e994c7() {
+ var res : u32 = dot(vec4<u32>(), vec4<u32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+ dot_e994c7();
+ return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+ dot_e994c7();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+ dot_e994c7();
+}
diff --git a/test/intrinsics/gen/dot/ef6b1d.wgsl b/test/intrinsics/gen/dot/ef6b1d.wgsl
new file mode 100644
index 0000000..b6e5eac
--- /dev/null
+++ b/test/intrinsics/gen/dot/ef6b1d.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+// test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+// src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn dot(vec<4, i32>, vec<4, i32>) -> i32
+fn dot_ef6b1d() {
+ var res: i32 = dot(vec4<i32>(), vec4<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+ dot_ef6b1d();
+ return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+ dot_ef6b1d();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+ dot_ef6b1d();
+}
diff --git a/test/intrinsics/gen/dot/ef6b1d.wgsl.expected.hlsl b/test/intrinsics/gen/dot/ef6b1d.wgsl.expected.hlsl
new file mode 100644
index 0000000..90b3adb
--- /dev/null
+++ b/test/intrinsics/gen/dot/ef6b1d.wgsl.expected.hlsl
@@ -0,0 +1,30 @@
+void dot_ef6b1d() {
+ int res = dot(int4(0, 0, 0, 0), int4(0, 0, 0, 0));
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ dot_ef6b1d();
+ return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ dot_ef6b1d();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ dot_ef6b1d();
+ return;
+}
diff --git a/test/intrinsics/gen/dot/ef6b1d.wgsl.expected.msl b/test/intrinsics/gen/dot/ef6b1d.wgsl.expected.msl
new file mode 100644
index 0000000..f7a85b8
--- /dev/null
+++ b/test/intrinsics/gen/dot/ef6b1d.wgsl.expected.msl
@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T>
+T tint_dot4(vec<T,4> a, vec<T,4> b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2] + a[3]*b[3];
+}
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+void dot_ef6b1d() {
+ int res = tint_dot4(int4(), int4());
+}
+
+float4 vertex_main_inner() {
+ dot_ef6b1d();
+ return float4();
+}
+
+vertex tint_symbol vertex_main() {
+ float4 const inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+fragment void fragment_main() {
+ dot_ef6b1d();
+ return;
+}
+
+kernel void compute_main() {
+ dot_ef6b1d();
+ return;
+}
+
diff --git a/test/intrinsics/gen/dot/ef6b1d.wgsl.expected.spvasm b/test/intrinsics/gen/dot/ef6b1d.wgsl.expected.spvasm
new file mode 100644
index 0000000..db22cad
--- /dev/null
+++ b/test/intrinsics/gen/dot/ef6b1d.wgsl.expected.spvasm
@@ -0,0 +1,82 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %dot_ef6b1d "dot_ef6b1d"
+ OpName %res "res"
+ OpName %vertex_main_inner "vertex_main_inner"
+ OpName %vertex_main "vertex_main"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %v4int = OpTypeVector %int 4
+ %16 = OpConstantNull %v4int
+%_ptr_Function_int = OpTypePointer Function %int
+ %33 = OpConstantNull %int
+ %34 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
+ %dot_ef6b1d = OpFunction %void None %9
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_int Function %33
+ %17 = OpCompositeExtract %int %16 0
+ %18 = OpCompositeExtract %int %16 0
+ %19 = OpIMul %int %17 %18
+ %20 = OpCompositeExtract %int %16 1
+ %21 = OpCompositeExtract %int %16 1
+ %22 = OpIMul %int %20 %21
+ %23 = OpIAdd %int %19 %22
+ %24 = OpCompositeExtract %int %16 2
+ %25 = OpCompositeExtract %int %16 2
+ %26 = OpIMul %int %24 %25
+ %27 = OpIAdd %int %23 %26
+ %28 = OpCompositeExtract %int %16 3
+ %29 = OpCompositeExtract %int %16 3
+ %30 = OpIMul %int %28 %29
+ %13 = OpIAdd %int %27 %30
+ OpStore %res %13
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+ %36 = OpLabel
+ %37 = OpFunctionCall %void %dot_ef6b1d
+ OpReturnValue %5
+ OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+ %39 = OpLabel
+ %40 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %40
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+ %43 = OpLabel
+ %44 = OpFunctionCall %void %dot_ef6b1d
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %9
+ %46 = OpLabel
+ %47 = OpFunctionCall %void %dot_ef6b1d
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/gen/dot/ef6b1d.wgsl.expected.wgsl b/test/intrinsics/gen/dot/ef6b1d.wgsl.expected.wgsl
new file mode 100644
index 0000000..4c35395
--- /dev/null
+++ b/test/intrinsics/gen/dot/ef6b1d.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn dot_ef6b1d() {
+ var res : i32 = dot(vec4<i32>(), vec4<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+ dot_ef6b1d();
+ return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+ dot_ef6b1d();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+ dot_ef6b1d();
+}
diff --git a/test/intrinsics/gen/dot/f1312c.wgsl b/test/intrinsics/gen/dot/f1312c.wgsl
new file mode 100644
index 0000000..f08b6f2
--- /dev/null
+++ b/test/intrinsics/gen/dot/f1312c.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+// test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+// src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn dot(vec<3, i32>, vec<3, i32>) -> i32
+fn dot_f1312c() {
+ var res: i32 = dot(vec3<i32>(), vec3<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+ dot_f1312c();
+ return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+ dot_f1312c();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+ dot_f1312c();
+}
diff --git a/test/intrinsics/gen/dot/f1312c.wgsl.expected.hlsl b/test/intrinsics/gen/dot/f1312c.wgsl.expected.hlsl
new file mode 100644
index 0000000..2c67fcc
--- /dev/null
+++ b/test/intrinsics/gen/dot/f1312c.wgsl.expected.hlsl
@@ -0,0 +1,30 @@
+void dot_f1312c() {
+ int res = dot(int3(0, 0, 0), int3(0, 0, 0));
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ dot_f1312c();
+ return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ dot_f1312c();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ dot_f1312c();
+ return;
+}
diff --git a/test/intrinsics/gen/dot/f1312c.wgsl.expected.msl b/test/intrinsics/gen/dot/f1312c.wgsl.expected.msl
new file mode 100644
index 0000000..aab9ef1
--- /dev/null
+++ b/test/intrinsics/gen/dot/f1312c.wgsl.expected.msl
@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T>
+T tint_dot3(vec<T,3> a, vec<T,3> b) {
+ return a[0]*b[0] + a[1]*b[1] + a[2]*b[2];
+}
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+void dot_f1312c() {
+ int res = tint_dot3(int3(), int3());
+}
+
+float4 vertex_main_inner() {
+ dot_f1312c();
+ return float4();
+}
+
+vertex tint_symbol vertex_main() {
+ float4 const inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+fragment void fragment_main() {
+ dot_f1312c();
+ return;
+}
+
+kernel void compute_main() {
+ dot_f1312c();
+ return;
+}
+
diff --git a/test/intrinsics/gen/dot/f1312c.wgsl.expected.spvasm b/test/intrinsics/gen/dot/f1312c.wgsl.expected.spvasm
new file mode 100644
index 0000000..45222d1
--- /dev/null
+++ b/test/intrinsics/gen/dot/f1312c.wgsl.expected.spvasm
@@ -0,0 +1,78 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 44
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %dot_f1312c "dot_f1312c"
+ OpName %res "res"
+ OpName %vertex_main_inner "vertex_main_inner"
+ OpName %vertex_main "vertex_main"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+ %16 = OpConstantNull %v3int
+%_ptr_Function_int = OpTypePointer Function %int
+ %29 = OpConstantNull %int
+ %30 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
+ %dot_f1312c = OpFunction %void None %9
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_int Function %29
+ %17 = OpCompositeExtract %int %16 0
+ %18 = OpCompositeExtract %int %16 0
+ %19 = OpIMul %int %17 %18
+ %20 = OpCompositeExtract %int %16 1
+ %21 = OpCompositeExtract %int %16 1
+ %22 = OpIMul %int %20 %21
+ %23 = OpIAdd %int %19 %22
+ %24 = OpCompositeExtract %int %16 2
+ %25 = OpCompositeExtract %int %16 2
+ %26 = OpIMul %int %24 %25
+ %13 = OpIAdd %int %23 %26
+ OpStore %res %13
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %30
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %dot_f1312c
+ OpReturnValue %5
+ OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+ %35 = OpLabel
+ %36 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %36
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+ %39 = OpLabel
+ %40 = OpFunctionCall %void %dot_f1312c
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %9
+ %42 = OpLabel
+ %43 = OpFunctionCall %void %dot_f1312c
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/gen/dot/f1312c.wgsl.expected.wgsl b/test/intrinsics/gen/dot/f1312c.wgsl.expected.wgsl
new file mode 100644
index 0000000..1e9ffe1
--- /dev/null
+++ b/test/intrinsics/gen/dot/f1312c.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn dot_f1312c() {
+ var res : i32 = dot(vec3<i32>(), vec3<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+ dot_f1312c();
+ return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+ dot_f1312c();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+ dot_f1312c();
+}
diff --git a/test/intrinsics/gen/dot/fc5f7c.wgsl b/test/intrinsics/gen/dot/fc5f7c.wgsl
new file mode 100644
index 0000000..7b0ae29
--- /dev/null
+++ b/test/intrinsics/gen/dot/fc5f7c.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+// test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+// src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn dot(vec<2, i32>, vec<2, i32>) -> i32
+fn dot_fc5f7c() {
+ var res: i32 = dot(vec2<i32>(), vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+ dot_fc5f7c();
+ return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+ dot_fc5f7c();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+ dot_fc5f7c();
+}
diff --git a/test/intrinsics/gen/dot/fc5f7c.wgsl.expected.hlsl b/test/intrinsics/gen/dot/fc5f7c.wgsl.expected.hlsl
new file mode 100644
index 0000000..22ff835
--- /dev/null
+++ b/test/intrinsics/gen/dot/fc5f7c.wgsl.expected.hlsl
@@ -0,0 +1,30 @@
+void dot_fc5f7c() {
+ int res = dot(int2(0, 0), int2(0, 0));
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ dot_fc5f7c();
+ return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ dot_fc5f7c();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ dot_fc5f7c();
+ return;
+}
diff --git a/test/intrinsics/gen/dot/fc5f7c.wgsl.expected.msl b/test/intrinsics/gen/dot/fc5f7c.wgsl.expected.msl
new file mode 100644
index 0000000..138e822
--- /dev/null
+++ b/test/intrinsics/gen/dot/fc5f7c.wgsl.expected.msl
@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T>
+T tint_dot2(vec<T,2> a, vec<T,2> b) {
+ return a[0]*b[0] + a[1]*b[1];
+}
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+void dot_fc5f7c() {
+ int res = tint_dot2(int2(), int2());
+}
+
+float4 vertex_main_inner() {
+ dot_fc5f7c();
+ return float4();
+}
+
+vertex tint_symbol vertex_main() {
+ float4 const inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+fragment void fragment_main() {
+ dot_fc5f7c();
+ return;
+}
+
+kernel void compute_main() {
+ dot_fc5f7c();
+ return;
+}
+
diff --git a/test/intrinsics/gen/dot/fc5f7c.wgsl.expected.spvasm b/test/intrinsics/gen/dot/fc5f7c.wgsl.expected.spvasm
new file mode 100644
index 0000000..6f05cef
--- /dev/null
+++ b/test/intrinsics/gen/dot/fc5f7c.wgsl.expected.spvasm
@@ -0,0 +1,74 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 40
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %dot_fc5f7c "dot_fc5f7c"
+ OpName %res "res"
+ OpName %vertex_main_inner "vertex_main_inner"
+ OpName %vertex_main "vertex_main"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %v2int = OpTypeVector %int 2
+ %16 = OpConstantNull %v2int
+%_ptr_Function_int = OpTypePointer Function %int
+ %25 = OpConstantNull %int
+ %26 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
+ %dot_fc5f7c = OpFunction %void None %9
+ %12 = OpLabel
+ %res = OpVariable %_ptr_Function_int Function %25
+ %17 = OpCompositeExtract %int %16 0
+ %18 = OpCompositeExtract %int %16 0
+ %19 = OpIMul %int %17 %18
+ %20 = OpCompositeExtract %int %16 1
+ %21 = OpCompositeExtract %int %16 1
+ %22 = OpIMul %int %20 %21
+ %13 = OpIAdd %int %19 %22
+ OpStore %res %13
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %26
+ %28 = OpLabel
+ %29 = OpFunctionCall %void %dot_fc5f7c
+ OpReturnValue %5
+ OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+ %31 = OpLabel
+ %32 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %32
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+ %35 = OpLabel
+ %36 = OpFunctionCall %void %dot_fc5f7c
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %9
+ %38 = OpLabel
+ %39 = OpFunctionCall %void %dot_fc5f7c
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/gen/dot/fc5f7c.wgsl.expected.wgsl b/test/intrinsics/gen/dot/fc5f7c.wgsl.expected.wgsl
new file mode 100644
index 0000000..4aaabfc
--- /dev/null
+++ b/test/intrinsics/gen/dot/fc5f7c.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn dot_fc5f7c() {
+ var res : i32 = dot(vec2<i32>(), vec2<i32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+ dot_fc5f7c();
+ return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+ dot_fc5f7c();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+ dot_fc5f7c();
+}
