Implement mixed vector-scalar float % operator
W3C consensus on https://github.com/gpuweb/gpuweb/issues/2450
Spec change: https://github.com/gpuweb/gpuweb/pull/2495
Bug: tint:1370
Change-Id: I85bb9c802b0355bc53aa8dbacca8427fb7be1ff6
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/84880
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
diff --git a/src/tint/fuzzers/tint_ast_fuzzer/mutations/change_binary_operator.cc b/src/tint/fuzzers/tint_ast_fuzzer/mutations/change_binary_operator.cc
index 174424a..60a2a1c 100644
--- a/src/tint/fuzzers/tint_ast_fuzzer/mutations/change_binary_operator.cc
+++ b/src/tint/fuzzers/tint_ast_fuzzer/mutations/change_binary_operator.cc
@@ -59,16 +59,10 @@
// type-compatible if the matrices are square.
return !lhs_type->is_float_matrix() || lhs_type->is_square_float_matrix();
case ast::BinaryOp::kDivide:
+ case ast::BinaryOp::kModulo:
// '/' is not defined for matrices.
return lhs_type->is_numeric_scalar_or_vector() &&
rhs_type->is_numeric_scalar_or_vector();
- case ast::BinaryOp::kModulo:
- // TODO(https://crbug.com/tint/1370): once fixed, the rules should be the
- // same as for divide.
- if (lhs_type->is_float_vector() || rhs_type->is_float_vector()) {
- return lhs_type == rhs_type;
- }
- return !lhs_type->is_float_matrix() && !rhs_type->is_float_matrix();
case ast::BinaryOp::kShiftLeft:
case ast::BinaryOp::kShiftRight:
return IsSuitableForShift(lhs_type, rhs_type);
@@ -102,16 +96,10 @@
// These operators require homogeneous integer types.
return lhs_type == rhs_type && lhs_type->is_integer_scalar_or_vector();
case ast::BinaryOp::kDivide:
+ case ast::BinaryOp::kModulo:
// '/' is not defined for matrices.
return lhs_type->is_numeric_scalar_or_vector() &&
rhs_type->is_numeric_scalar_or_vector();
- case ast::BinaryOp::kModulo:
- // TODO(https://crbug.com/tint/1370): once fixed, this should be the same
- // as for divide
- if (lhs_type->is_float_vector() || rhs_type->is_float_vector()) {
- return lhs_type == rhs_type;
- }
- return !lhs_type->is_float_matrix() && !rhs_type->is_float_matrix();
case ast::BinaryOp::kShiftLeft:
case ast::BinaryOp::kShiftRight:
return IsSuitableForShift(lhs_type, rhs_type);
@@ -120,9 +108,9 @@
}
}
-bool CanReplaceDivideWith(const sem::Type* lhs_type,
- const sem::Type* rhs_type,
- ast::BinaryOp new_operator) {
+bool CanReplaceDivideOrModuloWith(const sem::Type* lhs_type,
+ const sem::Type* rhs_type,
+ ast::BinaryOp new_operator) {
// The program is assumed to be well-typed, so this method determines when
// 'new_operator' can be used as a type-preserving replacement in a '/'
// expression.
@@ -131,12 +119,9 @@
case ast::BinaryOp::kSubtract:
case ast::BinaryOp::kMultiply:
case ast::BinaryOp::kDivide:
+ case ast::BinaryOp::kModulo:
// These operators work in all contexts where '/' works.
return true;
- case ast::BinaryOp::kModulo:
- // TODO(https://crbug.com/tint/1370): this special case should not be
- // required; modulo and divide should work in the same contexts.
- return lhs_type->is_integer_scalar_or_vector() || lhs_type == rhs_type;
case ast::BinaryOp::kAnd:
case ast::BinaryOp::kOr:
case ast::BinaryOp::kXor:
@@ -150,30 +135,6 @@
}
}
-// TODO(https://crbug.com/tint/1370): once fixed, this method will be removed
-// and the same method will be used to check Divide and Modulo.
-bool CanReplaceModuloWith(const sem::Type* lhs_type,
- const sem::Type* rhs_type,
- ast::BinaryOp new_operator) {
- switch (new_operator) {
- case ast::BinaryOp::kAdd:
- case ast::BinaryOp::kSubtract:
- case ast::BinaryOp::kMultiply:
- case ast::BinaryOp::kDivide:
- case ast::BinaryOp::kModulo:
- return true;
- case ast::BinaryOp::kAnd:
- case ast::BinaryOp::kOr:
- case ast::BinaryOp::kXor:
- return lhs_type == rhs_type && lhs_type->is_integer_scalar_or_vector();
- case ast::BinaryOp::kShiftLeft:
- case ast::BinaryOp::kShiftRight:
- return IsSuitableForShift(lhs_type, rhs_type);
- default:
- return false;
- }
-}
-
bool CanReplaceLogicalAndLogicalOrWith(ast::BinaryOp new_operator) {
switch (new_operator) {
case ast::BinaryOp::kLogicalAnd:
@@ -362,9 +323,9 @@
return CanReplaceMultiplyWith(lhs_basic_type, rhs_basic_type,
new_operator);
case ast::BinaryOp::kDivide:
- return CanReplaceDivideWith(lhs_basic_type, rhs_basic_type, new_operator);
case ast::BinaryOp::kModulo:
- return CanReplaceModuloWith(lhs_basic_type, rhs_basic_type, new_operator);
+ return CanReplaceDivideOrModuloWith(lhs_basic_type, rhs_basic_type,
+ new_operator);
case ast::BinaryOp::kAnd:
case ast::BinaryOp::kOr:
return CanReplaceAndOrWith(lhs_basic_type, rhs_basic_type, new_operator);
diff --git a/src/tint/fuzzers/tint_ast_fuzzer/mutations/change_binary_operator_test.cc b/src/tint/fuzzers/tint_ast_fuzzer/mutations/change_binary_operator_test.cc
index 1120d6f..d9994e3 100644
--- a/src/tint/fuzzers/tint_ast_fuzzer/mutations/change_binary_operator_test.cc
+++ b/src/tint/fuzzers/tint_ast_fuzzer/mutations/change_binary_operator_test.cc
@@ -281,18 +281,12 @@
}
for (std::string vector_type : {"vec2<f32>", "vec3<f32>", "vec4<f32>"}) {
std::string scalar_type = "f32";
- CheckMutations(
- vector_type, scalar_type, vector_type, op,
- {
- other_op, ast::BinaryOp::kMultiply, ast::BinaryOp::kDivide
- // TODO(https://crbug.com/tint/1370): once fixed, add kModulo
- });
- CheckMutations(
- scalar_type, vector_type, vector_type, op,
- {
- other_op, ast::BinaryOp::kMultiply, ast::BinaryOp::kDivide
- // TODO(https://crbug.com/tint/1370): once fixed, add kModulo
- });
+ CheckMutations(vector_type, scalar_type, vector_type, op,
+ {other_op, ast::BinaryOp::kMultiply,
+ ast::BinaryOp::kDivide, ast::BinaryOp::kModulo});
+ CheckMutations(scalar_type, vector_type, vector_type, op,
+ {other_op, ast::BinaryOp::kMultiply,
+ ast::BinaryOp::kDivide, ast::BinaryOp::kModulo});
}
for (std::string square_matrix_type :
{"mat2x2<f32>", "mat3x3<f32>", "mat4x4<f32>"}) {
@@ -353,20 +347,14 @@
}
for (std::string vector_type : {"vec2<f32>", "vec3<f32>", "vec4<f32>"}) {
std::string scalar_type = "f32";
- CheckMutations(
- vector_type, scalar_type, vector_type, ast::BinaryOp::kMultiply,
- {
- ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
- ast::BinaryOp::kDivide
- // TODO(https://crbug.com/tint/1370): once fixed, add kModulo
- });
- CheckMutations(
- scalar_type, vector_type, vector_type, ast::BinaryOp::kMultiply,
- {
- ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
- ast::BinaryOp::kDivide
- // TODO(https://crbug.com/tint/1370): once fixed, add kModulo
- });
+ CheckMutations(vector_type, scalar_type, vector_type,
+ ast::BinaryOp::kMultiply,
+ {ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
+ ast::BinaryOp::kDivide, ast::BinaryOp::kModulo});
+ CheckMutations(scalar_type, vector_type, vector_type,
+ ast::BinaryOp::kMultiply,
+ {ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
+ ast::BinaryOp::kDivide, ast::BinaryOp::kModulo});
}
for (std::string square_matrix_type :
{"mat2x2<f32>", "mat3x3<f32>", "mat4x4<f32>"}) {
@@ -472,7 +460,7 @@
ast::BinaryOp::kMultiply, {});
}
-TEST(ChangeBinaryOperatorTest, Divide) {
+TEST(ChangeBinaryOperatorTest, DivideAndModulo) {
for (std::string type : {"i32", "vec2<i32>", "vec3<i32>", "vec4<i32>"}) {
CheckMutations(
type, type, type, ast::BinaryOp::kDivide,
@@ -517,26 +505,15 @@
}
for (std::string vector_type : {"vec2<f32>", "vec3<f32>", "vec4<f32>"}) {
std::string scalar_type = "f32";
- CheckMutations(
- vector_type, scalar_type, vector_type, ast::BinaryOp::kDivide,
- {
- ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
- ast::BinaryOp::kMultiply
- // TODO(https://crbug.com/tint/1370): once fixed, add kModulo
- });
- CheckMutations(
- scalar_type, vector_type, vector_type, ast::BinaryOp::kDivide,
- {
- ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
- ast::BinaryOp::kMultiply
- // TODO(https://crbug.com/tint/1370): once fixed, add kModulo
- });
+ CheckMutations(vector_type, scalar_type, vector_type,
+ ast::BinaryOp::kDivide,
+ {ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
+ ast::BinaryOp::kMultiply, ast::BinaryOp::kModulo});
+ CheckMutations(scalar_type, vector_type, vector_type,
+ ast::BinaryOp::kDivide,
+ {ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
+ ast::BinaryOp::kMultiply, ast::BinaryOp::kModulo});
}
-}
-
-// TODO(https://crbug.com/tint/1370): once fixed, combine this with the Divide
-// test
-TEST(ChangeBinaryOperatorTest, Modulo) {
for (std::string type : {"i32", "vec2<i32>", "vec3<i32>", "vec4<i32>"}) {
CheckMutations(
type, type, type, ast::BinaryOp::kModulo,
@@ -579,8 +556,6 @@
{ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
ast::BinaryOp::kMultiply, ast::BinaryOp::kDivide});
}
- // TODO(https://crbug.com/tint/1370): mixed float scalars/vectors will be
- // added when this test is combined with the Divide test
}
TEST(ChangeBinaryOperatorTest, AndOrXor) {
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index bb05c9d..6068a73 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -1901,23 +1901,13 @@
}
// Binary arithmetic expressions with mixed scalar and vector operands
- if (lhs_vec_elem_type && (lhs_vec_elem_type == rhs_ty)) {
- if (expr->IsModulo()) {
- if (rhs_ty->is_integer_scalar()) {
- return build(lhs_ty);
- }
- } else if (rhs_ty->is_numeric_scalar()) {
- return build(lhs_ty);
- }
+ if (lhs_vec_elem_type && (lhs_vec_elem_type == rhs_ty) &&
+ rhs_ty->is_numeric_scalar()) {
+ return build(lhs_ty);
}
- if (rhs_vec_elem_type && (rhs_vec_elem_type == lhs_ty)) {
- if (expr->IsModulo()) {
- if (lhs_ty->is_integer_scalar()) {
- return build(rhs_ty);
- }
- } else if (lhs_ty->is_numeric_scalar()) {
- return build(rhs_ty);
- }
+ if (rhs_vec_elem_type && (rhs_vec_elem_type == lhs_ty) &&
+ lhs_ty->is_numeric_scalar()) {
+ return build(rhs_ty);
}
}
diff --git a/src/tint/resolver/resolver_test.cc b/src/tint/resolver/resolver_test.cc
index 3d12b75..e9b55f8 100644
--- a/src/tint/resolver/resolver_test.cc
+++ b/src/tint/resolver/resolver_test.cc
@@ -1367,15 +1367,13 @@
ParamsFor<vec3<f32>, f32, vec3<f32>>(Op::kSubtract),
ParamsFor<vec3<f32>, f32, vec3<f32>>(Op::kMultiply),
ParamsFor<vec3<f32>, f32, vec3<f32>>(Op::kDivide),
- // NOTE: no kModulo for vec3<f32>, f32
- // ParamsFor<vec3<f32>, f32, vec3<f32>>(Op::kModulo),
+ ParamsFor<vec3<f32>, f32, vec3<f32>>(Op::kModulo),
ParamsFor<f32, vec3<f32>, vec3<f32>>(Op::kAdd),
ParamsFor<f32, vec3<f32>, vec3<f32>>(Op::kSubtract),
ParamsFor<f32, vec3<f32>, vec3<f32>>(Op::kMultiply),
ParamsFor<f32, vec3<f32>, vec3<f32>>(Op::kDivide),
- // NOTE: no kModulo for f32, vec3<f32>
- // ParamsFor<f32, vec3<f32>, vec3<f32>>(Op::kModulo),
+ ParamsFor<f32, vec3<f32>, vec3<f32>>(Op::kModulo),
// Matrix arithmetic
ParamsFor<mat2x3<f32>, f32, mat2x3<f32>>(Op::kMultiply),
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl
new file mode 100644
index 0000000..55a3acc
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl
@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a = 4.;
+ let b = vec3<f32>(1., 2., 3.);
+ let r : vec3<f32> = a % b;
+}
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.glsl b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.glsl
new file mode 100644
index 0000000..bccb68c
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.glsl
@@ -0,0 +1,18 @@
+#version 310 es
+
+vec3 tint_float_modulo(float lhs, vec3 rhs) {
+ return (lhs - rhs * trunc(lhs / rhs));
+}
+
+
+void f() {
+ float a = 4.0f;
+ vec3 b = vec3(1.0f, 2.0f, 3.0f);
+ vec3 r = tint_float_modulo(a, b);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.hlsl b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..b2653df
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float a = 4.0f;
+ const float3 b = float3(1.0f, 2.0f, 3.0f);
+ const float3 r = (a % b);
+ return;
+}
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.msl b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.msl
new file mode 100644
index 0000000..bbcdd8e
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float const a = 4.0f;
+ float3 const b = float3(1.0f, 2.0f, 3.0f);
+ float3 const r = fmod(a, b);
+ return;
+}
+
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.spvasm b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..2264cc4
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.spvasm
@@ -0,0 +1,28 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %f "f"
+ OpExecutionMode %f LocalSize 1 1 1
+ OpName %f "f"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+ %float_4 = OpConstant %float 4
+ %v3float = OpTypeVector %float 3
+ %float_1 = OpConstant %float 1
+ %float_2 = OpConstant %float 2
+ %float_3 = OpConstant %float 3
+ %11 = OpConstantComposite %v3float %float_1 %float_2 %float_3
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+ %15 = OpConstantNull %v3float
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %13 = OpVariable %_ptr_Function_v3float Function %15
+ %16 = OpCompositeConstruct %v3float %float_4 %float_4 %float_4
+ %12 = OpFRem %v3float %16 %11
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.wgsl b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..7bfc341
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a = 4.0;
+ let b = vec3<f32>(1.0, 2.0, 3.0);
+ let r : vec3<f32> = (a % b);
+}
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl b/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl
new file mode 100644
index 0000000..9e874a5
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl
@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a = 4;
+ let b = vec3<i32>(1, 2, 3);
+ let r : vec3<i32> = a % b;
+}
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.glsl b/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.glsl
new file mode 100644
index 0000000..0167016
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.glsl
@@ -0,0 +1,13 @@
+#version 310 es
+
+void f() {
+ int a = 4;
+ ivec3 b = ivec3(1, 2, 3);
+ ivec3 r = (a % b);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.hlsl b/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..5b0bfca
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const int a = 4;
+ const int3 b = int3(1, 2, 3);
+ const int3 r = (a % b);
+ return;
+}
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.msl
new file mode 100644
index 0000000..533d477
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ int const a = 4;
+ int3 const b = int3(1, 2, 3);
+ int3 const r = (a % b);
+ return;
+}
+
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.spvasm b/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..6dd950c
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.spvasm
@@ -0,0 +1,28 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %f "f"
+ OpExecutionMode %f LocalSize 1 1 1
+ OpName %f "f"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %int_4 = OpConstant %int 4
+ %v3int = OpTypeVector %int 3
+ %int_1 = OpConstant %int 1
+ %int_2 = OpConstant %int 2
+ %int_3 = OpConstant %int 3
+ %11 = OpConstantComposite %v3int %int_1 %int_2 %int_3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+ %15 = OpConstantNull %v3int
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %13 = OpVariable %_ptr_Function_v3int Function %15
+ %16 = OpCompositeConstruct %v3int %int_4 %int_4 %int_4
+ %12 = OpSMod %v3int %16 %11
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.wgsl b/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..6b76bbb
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a = 4;
+ let b = vec3<i32>(1, 2, 3);
+ let r : vec3<i32> = (a % b);
+}
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl b/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl
new file mode 100644
index 0000000..e4d45ca
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl
@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a = 4u;
+ let b = vec3<u32>(1u, 2u, 3u);
+ let r : vec3<u32> = a % b;
+}
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.glsl b/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.glsl
new file mode 100644
index 0000000..f150a92
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.glsl
@@ -0,0 +1,13 @@
+#version 310 es
+
+void f() {
+ uint a = 4u;
+ uvec3 b = uvec3(1u, 2u, 3u);
+ uvec3 r = (a % b);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.hlsl b/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..97533fb
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const uint a = 4u;
+ const uint3 b = uint3(1u, 2u, 3u);
+ const uint3 r = (a % b);
+ return;
+}
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.msl b/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.msl
new file mode 100644
index 0000000..a5e85ce
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint const a = 4u;
+ uint3 const b = uint3(1u, 2u, 3u);
+ uint3 const r = (a % b);
+ return;
+}
+
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.spvasm b/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..99afb76
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.spvasm
@@ -0,0 +1,28 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %f "f"
+ OpExecutionMode %f LocalSize 1 1 1
+ OpName %f "f"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_4 = OpConstant %uint 4
+ %v3uint = OpTypeVector %uint 3
+ %uint_1 = OpConstant %uint 1
+ %uint_2 = OpConstant %uint 2
+ %uint_3 = OpConstant %uint 3
+ %11 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+ %15 = OpConstantNull %v3uint
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %13 = OpVariable %_ptr_Function_v3uint Function %15
+ %16 = OpCompositeConstruct %v3uint %uint_4 %uint_4 %uint_4
+ %12 = OpUMod %v3uint %16 %11
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.wgsl b/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..309c295
--- /dev/null
+++ b/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a = 4u;
+ let b = vec3<u32>(1u, 2u, 3u);
+ let r : vec3<u32> = (a % b);
+}
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl
new file mode 100644
index 0000000..7a9cd13
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl
@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a = vec3<f32>(1., 2., 3.);
+ let b = 4.;
+ let r : vec3<f32> = a % b;
+}
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.glsl b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.glsl
new file mode 100644
index 0000000..8aef558
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.glsl
@@ -0,0 +1,17 @@
+#version 310 es
+
+vec3 tint_float_modulo(vec3 lhs, float rhs) {
+ return (lhs - rhs * trunc(lhs / rhs));
+}
+
+
+void f() {
+ vec3 a = vec3(1.0f, 2.0f, 3.0f);
+ vec3 r = tint_float_modulo(a, 4.0f);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.hlsl b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..3aca50d
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float3 a = float3(1.0f, 2.0f, 3.0f);
+ const float3 r = (a % 4.0f);
+ return;
+}
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.msl b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.msl
new file mode 100644
index 0000000..ab3f908
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float3 const a = float3(1.0f, 2.0f, 3.0f);
+ float const b = 4.0f;
+ float3 const r = fmod(a, b);
+ return;
+}
+
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.spvasm b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..46b895a
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.spvasm
@@ -0,0 +1,28 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %f "f"
+ OpExecutionMode %f LocalSize 1 1 1
+ OpName %f "f"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+ %float_1 = OpConstant %float 1
+ %float_2 = OpConstant %float 2
+ %float_3 = OpConstant %float 3
+ %10 = OpConstantComposite %v3float %float_1 %float_2 %float_3
+ %float_4 = OpConstant %float 4
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+ %15 = OpConstantNull %v3float
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %13 = OpVariable %_ptr_Function_v3float Function %15
+ %16 = OpCompositeConstruct %v3float %float_4 %float_4 %float_4
+ %12 = OpFRem %v3float %10 %16
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.wgsl b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..31c0bbd
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a = vec3<f32>(1.0, 2.0, 3.0);
+ let b = 4.0;
+ let r : vec3<f32> = (a % b);
+}
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl b/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl
new file mode 100644
index 0000000..e06a539
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl
@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a = vec3<i32>(1, 2, 3);
+ let b = 4;
+ let r : vec3<i32> = a % b;
+}
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.glsl b/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.glsl
new file mode 100644
index 0000000..4ffb72d
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.glsl
@@ -0,0 +1,12 @@
+#version 310 es
+
+void f() {
+ ivec3 a = ivec3(1, 2, 3);
+ ivec3 r = (a % 4);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.hlsl b/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..1abee48
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const int3 a = int3(1, 2, 3);
+ const int3 r = (a % 4);
+ return;
+}
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.msl b/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.msl
new file mode 100644
index 0000000..8df9cbe
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ int3 const a = int3(1, 2, 3);
+ int const b = 4;
+ int3 const r = (a % b);
+ return;
+}
+
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.spvasm b/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..70a8d1e
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.spvasm
@@ -0,0 +1,28 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %f "f"
+ OpExecutionMode %f LocalSize 1 1 1
+ OpName %f "f"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+ %int_1 = OpConstant %int 1
+ %int_2 = OpConstant %int 2
+ %int_3 = OpConstant %int 3
+ %10 = OpConstantComposite %v3int %int_1 %int_2 %int_3
+ %int_4 = OpConstant %int 4
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+ %15 = OpConstantNull %v3int
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %13 = OpVariable %_ptr_Function_v3int Function %15
+ %16 = OpCompositeConstruct %v3int %int_4 %int_4 %int_4
+ %12 = OpSMod %v3int %10 %16
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.wgsl b/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..afa5ad7
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a = vec3<i32>(1, 2, 3);
+ let b = 4;
+ let r : vec3<i32> = (a % b);
+}
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl b/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl
new file mode 100644
index 0000000..af3a204
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl
@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a = vec3<u32>(1u, 2u, 3u);
+ let b = 4u;
+ let r : vec3<u32> = a % b;
+}
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.glsl b/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.glsl
new file mode 100644
index 0000000..e4ac530
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.glsl
@@ -0,0 +1,12 @@
+#version 310 es
+
+void f() {
+ uvec3 a = uvec3(1u, 2u, 3u);
+ uvec3 r = (a % 4u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.hlsl b/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..c2d051d
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const uint3 a = uint3(1u, 2u, 3u);
+ const uint3 r = (a % 4u);
+ return;
+}
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.msl b/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.msl
new file mode 100644
index 0000000..3adb345
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint3 const a = uint3(1u, 2u, 3u);
+ uint const b = 4u;
+ uint3 const r = (a % b);
+ return;
+}
+
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.spvasm b/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..3fdde36
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.spvasm
@@ -0,0 +1,28 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %f "f"
+ OpExecutionMode %f LocalSize 1 1 1
+ OpName %f "f"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+ %uint_1 = OpConstant %uint 1
+ %uint_2 = OpConstant %uint 2
+ %uint_3 = OpConstant %uint 3
+ %10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
+ %uint_4 = OpConstant %uint 4
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+ %15 = OpConstantNull %v3uint
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %13 = OpVariable %_ptr_Function_v3uint Function %15
+ %16 = OpCompositeConstruct %v3uint %uint_4 %uint_4 %uint_4
+ %12 = OpUMod %v3uint %10 %16
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.wgsl b/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..3c97f77
--- /dev/null
+++ b/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a = vec3<u32>(1u, 2u, 3u);
+ let b = 4u;
+ let r : vec3<u32> = (a % b);
+}
