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>

src/tint/fuzzers/tint_ast_fuzzer/mutations/change_binary_operator.cc

@@ -59,16 +59,10 @@ bool CanReplaceAddSubtractWith(const sem::Type* lhs_type,
       // 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 @@ bool CanReplaceMultiplyWith(const sem::Type* lhs_type,
       // 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 CanReplaceMultiplyWith(const sem::Type* lhs_type,
   }
 }
 
-bool CanReplaceDivideWith(const sem::Type* lhs_type,
+bool CanReplaceDivideOrModuloWith(const sem::Type* lhs_type,
-                          const sem::Type* rhs_type,
+                                  const sem::Type* rhs_type,
-                          ast::BinaryOp new_operator) {
+                                  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 @@ bool CanReplaceDivideWith(const sem::Type* lhs_type,
     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 @@ bool CanReplaceDivideWith(const sem::Type* lhs_type,
   }
 }
 
-// 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 @@ bool MutationChangeBinaryOperator::CanReplaceBinaryOperator(
       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);

src/tint/fuzzers/tint_ast_fuzzer/mutations/change_binary_operator_test.cc

@@ -281,18 +281,12 @@ TEST(ChangeBinaryOperatorTest, AddSubtract) {
     }
     for (std::string vector_type : {"vec2<f32>", "vec3<f32>", "vec4<f32>"}) {
       std::string scalar_type = "f32";
-      CheckMutations(
+      CheckMutations(vector_type, scalar_type, vector_type, op,
-          vector_type, scalar_type, vector_type, op,
+                     {other_op, ast::BinaryOp::kMultiply,
-          {
+                      ast::BinaryOp::kDivide, ast::BinaryOp::kModulo});
-              other_op, ast::BinaryOp::kMultiply, ast::BinaryOp::kDivide
+      CheckMutations(scalar_type, vector_type, vector_type, op,
-              // TODO(https://crbug.com/tint/1370): once fixed, add kModulo
+                     {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
-              // TODO(https://crbug.com/tint/1370): once fixed, add kModulo
-          });
     }
     for (std::string square_matrix_type :
          {"mat2x2<f32>", "mat3x3<f32>", "mat4x4<f32>"}) {
@@ -353,20 +347,14 @@ TEST(ChangeBinaryOperatorTest, Mul) {
   }
   for (std::string vector_type : {"vec2<f32>", "vec3<f32>", "vec4<f32>"}) {
     std::string scalar_type = "f32";
-    CheckMutations(
+    CheckMutations(vector_type, scalar_type, vector_type,
-        vector_type, scalar_type, vector_type, ast::BinaryOp::kMultiply,
+                   ast::BinaryOp::kMultiply,
-        {
+                   {ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
-            ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
+                    ast::BinaryOp::kDivide, ast::BinaryOp::kModulo});
-            ast::BinaryOp::kDivide
+    CheckMutations(scalar_type, vector_type, vector_type,
-            // TODO(https://crbug.com/tint/1370): once fixed, add kModulo
+                   ast::BinaryOp::kMultiply,
-        });
+                   {ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
-    CheckMutations(
+                    ast::BinaryOp::kDivide, ast::BinaryOp::kModulo});
-        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
-        });
   }
   for (std::string square_matrix_type :
        {"mat2x2<f32>", "mat3x3<f32>", "mat4x4<f32>"}) {
@@ -472,7 +460,7 @@ TEST(ChangeBinaryOperatorTest, Mul) {
                  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 @@ TEST(ChangeBinaryOperatorTest, Divide) {
   }
   for (std::string vector_type : {"vec2<f32>", "vec3<f32>", "vec4<f32>"}) {
     std::string scalar_type = "f32";
-    CheckMutations(
+    CheckMutations(vector_type, scalar_type, vector_type,
-        vector_type, scalar_type, vector_type, ast::BinaryOp::kDivide,
+                   ast::BinaryOp::kDivide,
-        {
+                   {ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
-            ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
+                    ast::BinaryOp::kMultiply, ast::BinaryOp::kModulo});
-            ast::BinaryOp::kMultiply
+    CheckMutations(scalar_type, vector_type, vector_type,
-            // TODO(https://crbug.com/tint/1370): once fixed, add kModulo
+                   ast::BinaryOp::kDivide,
-        });
+                   {ast::BinaryOp::kAdd, ast::BinaryOp::kSubtract,
-    CheckMutations(
+                    ast::BinaryOp::kMultiply, ast::BinaryOp::kModulo});
-        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
-        });
   }
-}
-
-// 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 @@ TEST(ChangeBinaryOperatorTest, Modulo) {
                    {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) {

src/tint/resolver/resolver.cc

@@ -1901,23 +1901,13 @@ sem::Expression* Resolver::Binary(const ast::BinaryExpression* expr) {
     }
 
     // Binary arithmetic expressions with mixed scalar and vector operands
-    if (lhs_vec_elem_type && (lhs_vec_elem_type == rhs_ty)) {
+    if (lhs_vec_elem_type && (lhs_vec_elem_type == rhs_ty) &&
-      if (expr->IsModulo()) {
+        rhs_ty->is_numeric_scalar()) {
-        if (rhs_ty->is_integer_scalar()) {
+      return build(lhs_ty);
-          return build(lhs_ty);
-        }
-      } else if (rhs_ty->is_numeric_scalar()) {
-        return build(lhs_ty);
-      }
     }
-    if (rhs_vec_elem_type && (rhs_vec_elem_type == lhs_ty)) {
+    if (rhs_vec_elem_type && (rhs_vec_elem_type == lhs_ty) &&
-      if (expr->IsModulo()) {
+        lhs_ty->is_numeric_scalar()) {
-        if (lhs_ty->is_integer_scalar()) {
+      return build(rhs_ty);
-          return build(rhs_ty);
-        }
-      } else if (lhs_ty->is_numeric_scalar()) {
-        return build(rhs_ty);
-      }
     }
   }
 

src/tint/resolver/resolver_test.cc

@@ -1367,15 +1367,13 @@ static constexpr Params all_valid_cases[] = {
     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),

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;
+}

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;
+}

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;
+}

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;
+}
+

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

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);
+}

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;
+}

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;
+}

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;
+}

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;
+}
+

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

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);
+}

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;
+}

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;
+}

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;
+}

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;
+}
+

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

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);
+}

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;
+}

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;
+}

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;
+}

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;
+}
+

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

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);
+}

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;
+}

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;
+}

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;
+}

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;
+}
+

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

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);
+}

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;
+}

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;
+}

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;
+}

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;
+}
+

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

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);
+}