Fix FXC compile errors on modulo by zero

Just like for divide, FXC fails with the exact same error when
performing a modulo on a value that FXC determines to be zero. We
address it in the same way as we do for divide.

This also fixes a couple of the vk-gl-cts tests for which I manually
generated expectation files for.

Bug: tint:1083
Change-Id: Ia388bf002112afded542adb791d37e88e35a77ff
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/74220
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>

src/program_builder.h

@@ -1665,6 +1665,16 @@ class ProgramBuilder {
                                          Expr(std::forward<RHS>(rhs)));
   }
 
+  /// @param lhs the left hand argument to the modulo operation
+  /// @param rhs the right hand argument to the modulo operation
+  /// @returns a `ast::BinaryExpression` applying modulo of `lhs` by `rhs`
+  template <typename LHS, typename RHS>
+  const ast::Expression* Mod(LHS&& lhs, RHS&& rhs) {
+    return create<ast::BinaryExpression>(ast::BinaryOp::kModulo,
+                                         Expr(std::forward<LHS>(lhs)),
+                                         Expr(std::forward<RHS>(rhs)));
+  }
+
   /// @param lhs the left hand argument to the bit shift right operation
   /// @param rhs the right hand argument to the bit shift right operation
   /// @returns a `ast::BinaryExpression` bit shifting right `lhs` by `rhs`

src/writer/hlsl/generator_impl.cc

@@ -863,8 +863,8 @@ bool GeneratorImpl::EmitBinary(std::ostream& out,
       break;
     case ast::BinaryOp::kDivide:
       out << "/";
-      // BUG(crbug.com/tint/1083): Integer divide by zero is a FXC compile
+      // BUG(crbug.com/tint/1083): Integer divide/modulo by zero is a FXC
-      // error, and undefined behavior in WGSL.
+      // compile error, and undefined behavior in WGSL.
       if (TypeOf(expr->rhs)->UnwrapRef()->is_integer_scalar_or_vector()) {
         out << " ";
         return EmitExpressionOrOneIfZero(out, expr->rhs);
@@ -872,6 +872,12 @@ bool GeneratorImpl::EmitBinary(std::ostream& out,
       break;
     case ast::BinaryOp::kModulo:
       out << "%";
+      // BUG(crbug.com/tint/1083): Integer divide/modulo by zero is a FXC
+      // compile error, and undefined behavior in WGSL.
+      if (TypeOf(expr->rhs)->UnwrapRef()->is_integer_scalar_or_vector()) {
+        out << " ";
+        return EmitExpressionOrOneIfZero(out, expr->rhs);
+      }
       break;
     case ast::BinaryOp::kNone:
       diagnostics_.add_error(diag::System::Writer,

src/writer/hlsl/generator_impl_binary_test.cc

@@ -137,10 +137,12 @@ INSTANTIATE_TEST_SUITE_P(
         BinaryData{"(left + right)", ast::BinaryOp::kAdd},
         BinaryData{"(left - right)", ast::BinaryOp::kSubtract},
         BinaryData{"(left * right)", ast::BinaryOp::kMultiply},
-        // NOTE: Integer divide covered by DivideBy* tests below
+        // NOTE: Integer divide covered by DivOrModBy* tests below
         BinaryData{"(left / right)", ast::BinaryOp::kDivide,
                    BinaryData::Types::Float},
-        BinaryData{"(left % right)", ast::BinaryOp::kModulo}));
+        // NOTE: Integer modulo covered by DivOrModBy* tests below
+        BinaryData{"(left % right)", ast::BinaryOp::kModulo,
+                   BinaryData::Types::Float}));
 
 TEST_F(HlslGeneratorImplTest_Binary, Multiply_VectorScalar) {
   auto* lhs = vec3<f32>(1.f, 1.f, 1.f);
@@ -526,11 +528,36 @@ foo((tint_tmp), (tint_tmp_1), (tint_tmp_2));
 )");
 }
 
-TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_i32) {
+namespace HlslGeneratorDivMod {
+
+struct Params {
+  enum class Type { Div, Mod };
+  Type type;
+};
+
+struct HlslGeneratorDivModTest : TestParamHelper<Params> {
+  std::string Token() {
+    return GetParam().type == Params::Type::Div ? "/" : "%";
+  }
+
+  template <typename... Args>
+  auto Op(Args... args) {
+    return GetParam().type == Params::Type::Div
+               ? Div(std::forward<Args>(args)...)
+               : Mod(std::forward<Args>(args)...);
+  }
+};
+
+INSTANTIATE_TEST_SUITE_P(HlslGeneratorImplTest,
+                         HlslGeneratorDivModTest,
+                         testing::Values(Params{Params::Type::Div},
+                                         Params{Params::Type::Mod}));
+
+TEST_P(HlslGeneratorDivModTest, DivOrModByLiteralZero_i32) {
   Func("fn", {}, ty.void_(),
        {
            Decl(Var("a", ty.i32())),
-           Decl(Const("r", nullptr, Div("a", 0))),
+           Decl(Const("r", nullptr, Op("a", 0))),
        });
 
   GeneratorImpl& gen = Build();
@@ -538,16 +565,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_i32) {
   ASSERT_TRUE(gen.Generate());
   EXPECT_EQ(gen.result(), R"(void fn() {
   int a = 0;
-  const int r = (a / 1);
+  const int r = (a )" + Token() +
+                              R"( 1);
 }
 )");
 }
 
-TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_u32) {
+TEST_P(HlslGeneratorDivModTest, DivOrModByLiteralZero_u32) {
   Func("fn", {}, ty.void_(),
        {
            Decl(Var("a", ty.u32())),
-           Decl(Const("r", nullptr, Div("a", 0u))),
+           Decl(Const("r", nullptr, Op("a", 0u))),
        });
 
   GeneratorImpl& gen = Build();
@@ -555,16 +583,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_u32) {
   ASSERT_TRUE(gen.Generate());
   EXPECT_EQ(gen.result(), R"(void fn() {
   uint a = 0u;
-  const uint r = (a / 1u);
+  const uint r = (a )" + Token() +
+                              R"( 1u);
 }
 )");
-}
+}  // namespace HlslGeneratorDivMod
 
-TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_vec_by_vec_i32) {
+TEST_P(HlslGeneratorDivModTest, DivOrModByLiteralZero_vec_by_vec_i32) {
   Func("fn", {}, ty.void_(),
        {
            Decl(Var("a", nullptr, vec4<i32>(100, 100, 100, 100))),
-           Decl(Const("r", nullptr, Div("a", vec4<i32>(50, 0, 25, 0)))),
+           Decl(Const("r", nullptr, Op("a", vec4<i32>(50, 0, 25, 0)))),
        });
 
   GeneratorImpl& gen = Build();
@@ -572,16 +601,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_vec_by_vec_i32) {
   ASSERT_TRUE(gen.Generate());
   EXPECT_EQ(gen.result(), R"(void fn() {
   int4 a = int4(100, 100, 100, 100);
-  const int4 r = (a / int4(50, 1, 25, 1));
+  const int4 r = (a )" + Token() +
+                              R"( int4(50, 1, 25, 1));
 }
 )");
-}
+}  // namespace
 
-TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_vec_by_scalar_i32) {
+TEST_P(HlslGeneratorDivModTest, DivOrModByLiteralZero_vec_by_scalar_i32) {
   Func("fn", {}, ty.void_(),
        {
            Decl(Var("a", nullptr, vec4<i32>(100, 100, 100, 100))),
-           Decl(Const("r", nullptr, Div("a", 0))),
+           Decl(Const("r", nullptr, Op("a", 0))),
        });
 
   GeneratorImpl& gen = Build();
@@ -589,16 +619,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_vec_by_scalar_i32) {
   ASSERT_TRUE(gen.Generate());
   EXPECT_EQ(gen.result(), R"(void fn() {
   int4 a = int4(100, 100, 100, 100);
-  const int4 r = (a / 1);
+  const int4 r = (a )" + Token() +
+                              R"( 1);
 }
 )");
-}
+}  // namespace hlsl
 
-TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_i32) {
+TEST_P(HlslGeneratorDivModTest, DivOrModByIdentifier_i32) {
   Func("fn", {Param("b", ty.i32())}, ty.void_(),
        {
            Decl(Var("a", ty.i32())),
-           Decl(Const("r", nullptr, Div("a", "b"))),
+           Decl(Const("r", nullptr, Op("a", "b"))),
        });
 
   GeneratorImpl& gen = Build();
@@ -606,16 +637,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_i32) {
   ASSERT_TRUE(gen.Generate());
   EXPECT_EQ(gen.result(), R"(void fn(int b) {
   int a = 0;
-  const int r = (a / (b == 0 ? 1 : b));
+  const int r = (a )" + Token() +
+                              R"( (b == 0 ? 1 : b));
 }
 )");
-}
+}  // namespace writer
 
-TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_u32) {
+TEST_P(HlslGeneratorDivModTest, DivOrModByIdentifier_u32) {
   Func("fn", {Param("b", ty.u32())}, ty.void_(),
        {
            Decl(Var("a", ty.u32())),
-           Decl(Const("r", nullptr, Div("a", "b"))),
+           Decl(Const("r", nullptr, Op("a", "b"))),
        });
 
   GeneratorImpl& gen = Build();
@@ -623,16 +655,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_u32) {
   ASSERT_TRUE(gen.Generate());
   EXPECT_EQ(gen.result(), R"(void fn(uint b) {
   uint a = 0u;
-  const uint r = (a / (b == 0u ? 1u : b));
+  const uint r = (a )" + Token() +
+                              R"( (b == 0u ? 1u : b));
 }
 )");
-}
+}  // namespace tint
 
-TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_vec_by_vec_i32) {
+TEST_P(HlslGeneratorDivModTest, DivOrModByIdentifier_vec_by_vec_i32) {
   Func("fn", {Param("b", ty.vec3<i32>())}, ty.void_(),
        {
            Decl(Var("a", ty.vec3<i32>())),
-           Decl(Const("r", nullptr, Div("a", "b"))),
+           Decl(Const("r", nullptr, Op("a", "b"))),
        });
 
   GeneratorImpl& gen = Build();
@@ -640,16 +673,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_vec_by_vec_i32) {
   ASSERT_TRUE(gen.Generate());
   EXPECT_EQ(gen.result(), R"(void fn(int3 b) {
   int3 a = int3(0, 0, 0);
-  const int3 r = (a / (b == int3(0, 0, 0) ? int3(1, 1, 1) : b));
+  const int3 r = (a )" + Token() +
+                              R"( (b == int3(0, 0, 0) ? int3(1, 1, 1) : b));
 }
 )");
 }
 
-TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_vec_by_scalar_i32) {
+TEST_P(HlslGeneratorDivModTest, DivOrModByIdentifier_vec_by_scalar_i32) {
   Func("fn", {Param("b", ty.i32())}, ty.void_(),
        {
            Decl(Var("a", ty.vec3<i32>())),
-           Decl(Const("r", nullptr, Div("a", "b"))),
+           Decl(Const("r", nullptr, Op("a", "b"))),
        });
 
   GeneratorImpl& gen = Build();
@@ -657,12 +691,13 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_vec_by_scalar_i32) {
   ASSERT_TRUE(gen.Generate());
   EXPECT_EQ(gen.result(), R"(void fn(int b) {
   int3 a = int3(0, 0, 0);
-  const int3 r = (a / (b == 0 ? 1 : b));
+  const int3 r = (a )" + Token() +
+                              R"( (b == 0 ? 1 : b));
 }
 )");
 }
 
-TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_i32) {
+TEST_P(HlslGeneratorDivModTest, DivOrModByExpression_i32) {
   Func("zero", {}, ty.i32(),
        {
            Return(Expr(0)),
@@ -671,7 +706,7 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_i32) {
   Func("fn", {}, ty.void_(),
        {
            Decl(Var("a", ty.i32())),
-           Decl(Const("r", nullptr, Div("a", Call("zero")))),
+           Decl(Const("r", nullptr, Op("a", Call("zero")))),
        });
 
   GeneratorImpl& gen = Build();
@@ -687,12 +722,13 @@ int zero() {
 
 void fn() {
   int a = 0;
-  const int r = (a / value_or_one_if_zero_int(zero()));
+  const int r = (a )" + Token() +
+                              R"( value_or_one_if_zero_int(zero()));
 }
 )");
 }
 
-TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_u32) {
+TEST_P(HlslGeneratorDivModTest, DivOrModByExpression_u32) {
   Func("zero", {}, ty.u32(),
        {
            Return(Expr(0u)),
@@ -701,7 +737,7 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_u32) {
   Func("fn", {}, ty.void_(),
        {
            Decl(Var("a", ty.u32())),
-           Decl(Const("r", nullptr, Div("a", Call("zero")))),
+           Decl(Const("r", nullptr, Op("a", Call("zero")))),
        });
 
   GeneratorImpl& gen = Build();
@@ -717,12 +753,13 @@ uint zero() {
 
 void fn() {
   uint a = 0u;
-  const uint r = (a / value_or_one_if_zero_uint(zero()));
+  const uint r = (a )" + Token() +
+                              R"( value_or_one_if_zero_uint(zero()));
 }
 )");
 }
 
-TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_vec_by_vec_i32) {
+TEST_P(HlslGeneratorDivModTest, DivOrModByExpression_vec_by_vec_i32) {
   Func("zero", {}, ty.vec3<i32>(),
        {
            Return(vec3<i32>(0, 0, 0)),
@@ -731,7 +768,7 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_vec_by_vec_i32) {
   Func("fn", {}, ty.void_(),
        {
            Decl(Var("a", ty.vec3<i32>())),
-           Decl(Const("r", nullptr, Div("a", Call("zero")))),
+           Decl(Const("r", nullptr, Op("a", Call("zero")))),
        });
 
   GeneratorImpl& gen = Build();
@@ -747,12 +784,13 @@ int3 zero() {
 
 void fn() {
   int3 a = int3(0, 0, 0);
-  const int3 r = (a / value_or_one_if_zero_int3(zero()));
+  const int3 r = (a )" + Token() +
+                              R"( value_or_one_if_zero_int3(zero()));
 }
 )");
 }
 
-TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_vec_by_scalar_i32) {
+TEST_P(HlslGeneratorDivModTest, DivOrModByExpression_vec_by_scalar_i32) {
   Func("zero", {}, ty.i32(),
        {
            Return(0),
@@ -761,7 +799,7 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_vec_by_scalar_i32) {
   Func("fn", {}, ty.void_(),
        {
            Decl(Var("a", ty.vec3<i32>())),
-           Decl(Const("r", nullptr, Div("a", Call("zero")))),
+           Decl(Const("r", nullptr, Op("a", Call("zero")))),
        });
 
   GeneratorImpl& gen = Build();
@@ -777,10 +815,12 @@ int zero() {
 
 void fn() {
   int3 a = int3(0, 0, 0);
-  const int3 r = (a / value_or_one_if_zero_int(zero()));
+  const int3 r = (a )" + Token() +
+                              R"( value_or_one_if_zero_int(zero()));
 }
 )");
 }
+}  // namespace HlslGeneratorDivMod
 
 }  // namespace
 }  // namespace hlsl

test/bug/chromium/1273230.wgsl.expected.hlsl

@@ -89,7 +89,7 @@ uint toIndex1D(uint gridSize, float3 voxelPos) {
 uint3 toIndex4D(uint gridSize, uint index) {
   uint z_1 = (gridSize / value_or_one_if_zero_uint((index * index)));
   uint y_1 = ((gridSize - ((gridSize * gridSize) * z_1)) / (gridSize == 0u ? 1u : gridSize));
-  uint x_1 = (index % gridSize);
+  uint x_1 = (index % (gridSize == 0u ? 1u : gridSize));
   return uint3(z_1, y_1, y_1);
 }
 

test/bug/tint/1113.wgsl.expected.hlsl

@@ -60,7 +60,7 @@ uint toIndex1D(uint gridSize, float3 voxelPos) {
 uint3 toIndex3D(uint gridSize, uint index) {
   uint z_1 = (index / value_or_one_if_zero_uint((gridSize * gridSize)));
   uint y_1 = ((index - ((gridSize * gridSize) * z_1)) / (gridSize == 0u ? 1u : gridSize));
-  uint x_1 = (index % gridSize);
+  uint x_1 = (index % (gridSize == 0u ? 1u : gridSize));
   return uint3(x_1, y_1, z_1);
 }
 

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = 1.;
+    let b = 0.;
+    let r : f32 = a % b;
+}

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.hlsl

@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const float r = (1.0f % 0.0f);
+  return;
+}

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  float const a = 1.0f;
+  float const b = 0.0f;
+  float const r = fmod(a, b);
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.spvasm

@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 9
+; 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_1 = OpConstant %float 1
+    %float_0 = OpConstant %float 0
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %8 = OpFRem %float %float_1 %float_0
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = 1.0;
+  let b = 0.0;
+  let r : f32 = (a % b);
+}

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = 1;
+    let b = 0;
+    let r : i32 = a % b;
+}

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.hlsl

@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const int r = (1 % 1);
+  return;
+}

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  int const a = 1;
+  int const b = 0;
+  int const r = (a % b);
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.spvasm

@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 9
+; 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_1 = OpConstant %int 1
+      %int_0 = OpConstant %int 0
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %8 = OpSMod %int %int_1 %int_0
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = 1;
+  let b = 0;
+  let r : i32 = (a % b);
+}

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = 1u;
+    let b = 0u;
+    let r : u32 = a % b;
+}

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.hlsl

@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const uint r = (1u % 1u);
+  return;
+}

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  uint const a = 1u;
+  uint const b = 0u;
+  uint const r = (a % b);
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.spvasm

@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 9
+; 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_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %8 = OpUMod %uint %uint_1 %uint_0
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = 1u;
+  let b = 0u;
+  let r : u32 = (a % b);
+}

test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = 4;
+    let b = vec3<i32>(0, 2, 0);
+    let r : vec3<i32> = a % b;
+}

test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.hlsl

@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const int a = 4;
+  const int3 b = int3(0, 2, 0);
+  const int3 r = (a % int3(1, 2, 1));
+  return;
+}

test/expressions/binary/mod_by_zero/by_constant/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(0, 2, 0);
+  int3 const r = (a % b);
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.spvasm

@@ -0,0 +1,27 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; 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_0 = OpConstant %int 0
+      %int_2 = OpConstant %int 2
+         %10 = OpConstantComposite %v3int %int_0 %int_2 %int_0
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+         %14 = OpConstantNull %v3int
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+         %12 = OpVariable %_ptr_Function_v3int Function %14
+         %15 = OpCompositeConstruct %v3int %int_4 %int_4 %int_4
+         %11 = OpSMod %v3int %15 %10
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = 4;
+  let b = vec3<i32>(0, 2, 0);
+  let r : vec3<i32> = (a % b);
+}

test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = 4u;
+    let b = vec3<u32>(0u, 2u, 0u);
+    let r : vec3<u32> = a % b;
+}

test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.hlsl

@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const uint a = 4u;
+  const uint3 b = uint3(0u, 2u, 0u);
+  const uint3 r = (a % uint3(1u, 2u, 1u));
+  return;
+}

test/expressions/binary/mod_by_zero/by_constant/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(0u, 2u, 0u);
+  uint3 const r = (a % b);
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.spvasm

@@ -0,0 +1,27 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; 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_0 = OpConstant %uint 0
+     %uint_2 = OpConstant %uint 2
+         %10 = OpConstantComposite %v3uint %uint_0 %uint_2 %uint_0
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+         %14 = OpConstantNull %v3uint
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+         %12 = OpVariable %_ptr_Function_v3uint Function %14
+         %15 = OpCompositeConstruct %v3uint %uint_4 %uint_4 %uint_4
+         %11 = OpUMod %v3uint %15 %10
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = 4u;
+  let b = vec3<u32>(0u, 2u, 0u);
+  let r : vec3<u32> = (a % b);
+}

test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = vec3<i32>(1, 2, 3);
+    let b = 0;
+    let r : vec3<i32> = a % b;
+}

test/expressions/binary/mod_by_zero/by_constant/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 % 1);
+  return;
+}

test/expressions/binary/mod_by_zero/by_constant/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 = 0;
+  int3 const r = (a % b);
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_constant/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_0 = OpConstant %int 0
+%_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_0 %int_0 %int_0
+         %12 = OpSMod %v3int %10 %16
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_constant/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 = 0;
+  let r : vec3<i32> = (a % b);
+}

test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = vec3<u32>(1u, 2u, 3u);
+    let b = 0u;
+    let r : vec3<u32> = a % b;
+}

test/expressions/binary/mod_by_zero/by_constant/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 % 1u);
+  return;
+}

test/expressions/binary/mod_by_zero/by_constant/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 = 0u;
+  uint3 const r = (a % b);
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_constant/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_0 = OpConstant %uint 0
+%_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_0 %uint_0 %uint_0
+         %12 = OpUMod %v3uint %10 %16
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_constant/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 = 0u;
+  let r : vec3<u32> = (a % b);
+}

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = vec3<f32>(1., 2., 3.);
+    let b = vec3<f32>(0., 5., 0.);
+    let r : vec3<f32> = a % b;
+}

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.hlsl

@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const float3 a = float3(1.0f, 2.0f, 3.0f);
+  const float3 b = float3(0.0f, 5.0f, 0.0f);
+  const float3 r = (a % b);
+  return;
+}

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/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);
+  float3 const b = float3(0.0f, 5.0f, 0.0f);
+  float3 const r = fmod(a, b);
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.spvasm

@@ -0,0 +1,26 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 15
+; 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_0 = OpConstant %float 0
+    %float_5 = OpConstant %float 5
+         %13 = OpConstantComposite %v3float %float_0 %float_5 %float_0
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+         %14 = OpFRem %v3float %10 %13
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/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 = vec3<f32>(0.0, 5.0, 0.0);
+  let r : vec3<f32> = (a % b);
+}

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = vec3<i32>(1, 2, 3);
+    let b = vec3<i32>(0, 5, 0);
+    let r : vec3<i32> = a % b;
+}

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.hlsl

@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const int3 a = int3(1, 2, 3);
+  const int3 b = int3(0, 5, 0);
+  const int3 r = (a % int3(1, 5, 1));
+  return;
+}

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  int3 const a = int3(1, 2, 3);
+  int3 const b = int3(0, 5, 0);
+  int3 const r = (a % b);
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.spvasm

@@ -0,0 +1,26 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 15
+; 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_0 = OpConstant %int 0
+      %int_5 = OpConstant %int 5
+         %13 = OpConstantComposite %v3int %int_0 %int_5 %int_0
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+         %14 = OpSMod %v3int %10 %13
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = vec3<i32>(1, 2, 3);
+  let b = vec3<i32>(0, 5, 0);
+  let r : vec3<i32> = (a % b);
+}

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = vec3<u32>(1u, 2u, 3u);
+    let b = vec3<u32>(0u, 5u, 0u);
+    let r : vec3<u32> = a % b;
+}

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.hlsl

@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const uint3 a = uint3(1u, 2u, 3u);
+  const uint3 b = uint3(0u, 5u, 0u);
+  const uint3 r = (a % uint3(1u, 5u, 1u));
+  return;
+}

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  uint3 const a = uint3(1u, 2u, 3u);
+  uint3 const b = uint3(0u, 5u, 0u);
+  uint3 const r = (a % b);
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.spvasm

@@ -0,0 +1,26 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 15
+; 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_0 = OpConstant %uint 0
+     %uint_5 = OpConstant %uint 5
+         %13 = OpConstantComposite %v3uint %uint_0 %uint_5 %uint_0
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+         %14 = OpUMod %v3uint %10 %13
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = vec3<u32>(1u, 2u, 3u);
+  let b = vec3<u32>(0u, 5u, 0u);
+  let r : vec3<u32> = (a % b);
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    var a = 1.;
+    var b = 0.;
+    let r : f32 = a % (b + b);
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.hlsl

@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+  float a = 1.0f;
+  float b = 0.0f;
+  const float r = (a % (b + b));
+  return;
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  float a = 1.0f;
+  float b = 0.0f;
+  float const r = fmod(a, (b + b));
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; 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"
+               OpName %a "a"
+               OpName %b "b"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+    %float_1 = OpConstant %float 1
+%_ptr_Function_float = OpTypePointer Function %float
+          %9 = OpConstantNull %float
+    %float_0 = OpConstant %float 0
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %a = OpVariable %_ptr_Function_float Function %9
+          %b = OpVariable %_ptr_Function_float Function %9
+               OpStore %a %float_1
+               OpStore %b %float_0
+         %12 = OpLoad %float %a
+         %13 = OpLoad %float %b
+         %14 = OpLoad %float %b
+         %15 = OpFAdd %float %13 %14
+         %16 = OpFRem %float %12 %15
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var a = 1.0;
+  var b = 0.0;
+  let r : f32 = (a % (b + b));
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    var a = 1;
+    var b = 0;
+    let r : i32 = a % (b + b);
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.hlsl

@@ -0,0 +1,11 @@
+int value_or_one_if_zero_int(int value) {
+  return value == 0 ? 1 : value;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  int a = 1;
+  int b = 0;
+  const int r = (a % value_or_one_if_zero_int((b + b)));
+  return;
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  int a = 1;
+  int b = 0;
+  int const r = (a % as_type<int>((as_type<uint>(b) + as_type<uint>(b))));
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; 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"
+               OpName %a "a"
+               OpName %b "b"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+%_ptr_Function_int = OpTypePointer Function %int
+          %9 = OpConstantNull %int
+      %int_0 = OpConstant %int 0
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %a = OpVariable %_ptr_Function_int Function %9
+          %b = OpVariable %_ptr_Function_int Function %9
+               OpStore %a %int_1
+               OpStore %b %int_0
+         %12 = OpLoad %int %a
+         %13 = OpLoad %int %b
+         %14 = OpLoad %int %b
+         %15 = OpIAdd %int %13 %14
+         %16 = OpSMod %int %12 %15
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var a = 1;
+  var b = 0;
+  let r : i32 = (a % (b + b));
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    var a = 1u;
+    var b = 0u;
+    let r : u32 = a % (b + b);
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.hlsl

@@ -0,0 +1,11 @@
+uint value_or_one_if_zero_uint(uint value) {
+  return value == 0u ? 1u : value;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  uint a = 1u;
+  uint b = 0u;
+  const uint r = (a % value_or_one_if_zero_uint((b + b)));
+  return;
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  uint a = 1u;
+  uint b = 0u;
+  uint const r = (a % (b + b));
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; 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"
+               OpName %a "a"
+               OpName %b "b"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_uint = OpTypePointer Function %uint
+          %9 = OpConstantNull %uint
+     %uint_0 = OpConstant %uint 0
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %a = OpVariable %_ptr_Function_uint Function %9
+          %b = OpVariable %_ptr_Function_uint Function %9
+               OpStore %a %uint_1
+               OpStore %b %uint_0
+         %12 = OpLoad %uint %a
+         %13 = OpLoad %uint %b
+         %14 = OpLoad %uint %b
+         %15 = OpIAdd %uint %13 %14
+         %16 = OpUMod %uint %12 %15
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var a = 1u;
+  var b = 0u;
+  let r : u32 = (a % (b + b));
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    var a = 4;
+    var b = vec3<i32>(0, 2, 0);
+    let r : vec3<i32> = a % (b + b);
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.hlsl

@@ -0,0 +1,11 @@
+int3 value_or_one_if_zero_int3(int3 value) {
+  return value == int3(0, 0, 0) ? int3(1, 1, 1) : value;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  int a = 4;
+  int3 b = int3(0, 2, 0);
+  const int3 r = (a % value_or_one_if_zero_int3((b + b)));
+  return;
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  int a = 4;
+  int3 b = int3(0, 2, 0);
+  int3 const r = (a % as_type<int3>((as_type<uint3>(b) + as_type<uint3>(b))));
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.spvasm

@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+               OpName %a "a"
+               OpName %b "b"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %int_4 = OpConstant %int 4
+%_ptr_Function_int = OpTypePointer Function %int
+          %9 = OpConstantNull %int
+      %v3int = OpTypeVector %int 3
+      %int_0 = OpConstant %int 0
+      %int_2 = OpConstant %int 2
+         %13 = OpConstantComposite %v3int %int_0 %int_2 %int_0
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+         %16 = OpConstantNull %v3int
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %a = OpVariable %_ptr_Function_int Function %9
+          %b = OpVariable %_ptr_Function_v3int Function %16
+         %22 = OpVariable %_ptr_Function_v3int Function %16
+               OpStore %a %int_4
+               OpStore %b %13
+         %17 = OpLoad %int %a
+         %18 = OpLoad %v3int %b
+         %19 = OpLoad %v3int %b
+         %20 = OpIAdd %v3int %18 %19
+         %23 = OpCompositeConstruct %v3int %17 %17 %17
+         %21 = OpSMod %v3int %23 %20
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var a = 4;
+  var b = vec3<i32>(0, 2, 0);
+  let r : vec3<i32> = (a % (b + b));
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    var a = 4u;
+    var b = vec3<u32>(0u, 2u, 0u);
+    let r : vec3<u32> = a % (b + b);
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.hlsl

@@ -0,0 +1,11 @@
+uint3 value_or_one_if_zero_uint3(uint3 value) {
+  return value == uint3(0u, 0u, 0u) ? uint3(1u, 1u, 1u) : value;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  uint a = 4u;
+  uint3 b = uint3(0u, 2u, 0u);
+  const uint3 r = (a % value_or_one_if_zero_uint3((b + b)));
+  return;
+}

test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  uint a = 4u;
+  uint3 b = uint3(0u, 2u, 0u);
+  uint3 const r = (a % (b + b));
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.spvasm

@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+               OpName %a "a"
+               OpName %b "b"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_4 = OpConstant %uint 4
+%_ptr_Function_uint = OpTypePointer Function %uint
+          %9 = OpConstantNull %uint
+     %v3uint = OpTypeVector %uint 3
+     %uint_0 = OpConstant %uint 0
+     %uint_2 = OpConstant %uint 2
+         %13 = OpConstantComposite %v3uint %uint_0 %uint_2 %uint_0
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+         %16 = OpConstantNull %v3uint
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %a = OpVariable %_ptr_Function_uint Function %9
+          %b = OpVariable %_ptr_Function_v3uint Function %16
+         %22 = OpVariable %_ptr_Function_v3uint Function %16
+               OpStore %a %uint_4
+               OpStore %b %13
+         %17 = OpLoad %uint %a
+         %18 = OpLoad %v3uint %b
+         %19 = OpLoad %v3uint %b
+         %20 = OpIAdd %v3uint %18 %19
+         %23 = OpCompositeConstruct %v3uint %17 %17 %17
+         %21 = OpUMod %v3uint %23 %20
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var a = 4u;
+  var b = vec3<u32>(0u, 2u, 0u);
+  let r : vec3<u32> = (a % (b + b));
+}

test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    var a = vec3<i32>(1, 2, 3);
+    var b = 0;
+    let r : vec3<i32> = a % (b + b);
+}

test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.hlsl

@@ -0,0 +1,11 @@
+int value_or_one_if_zero_int(int value) {
+  return value == 0 ? 1 : value;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  int3 a = int3(1, 2, 3);
+  int b = 0;
+  const int3 r = (a % value_or_one_if_zero_int((b + b)));
+  return;
+}

test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  int3 a = int3(1, 2, 3);
+  int b = 0;
+  int3 const r = (a % as_type<int>((as_type<uint>(b) + as_type<uint>(b))));
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.spvasm

@@ -0,0 +1,40 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+               OpName %a "a"
+               OpName %b "b"
+       %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
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+         %13 = OpConstantNull %v3int
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %17 = OpConstantNull %int
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %a = OpVariable %_ptr_Function_v3int Function %13
+          %b = OpVariable %_ptr_Function_int Function %17
+         %23 = OpVariable %_ptr_Function_v3int Function %13
+               OpStore %a %10
+               OpStore %b %int_0
+         %18 = OpLoad %v3int %a
+         %19 = OpLoad %int %b
+         %20 = OpLoad %int %b
+         %21 = OpIAdd %int %19 %20
+         %24 = OpCompositeConstruct %v3int %21 %21 %21
+         %22 = OpSMod %v3int %18 %24
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var a = vec3<i32>(1, 2, 3);
+  var b = 0;
+  let r : vec3<i32> = (a % (b + b));
+}

test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    var a = vec3<u32>(1u, 2u, 3u);
+    var b = 0u;
+    let r : vec3<u32> = a % (b + b);
+}

test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.hlsl

@@ -0,0 +1,11 @@
+uint value_or_one_if_zero_uint(uint value) {
+  return value == 0u ? 1u : value;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  uint3 a = uint3(1u, 2u, 3u);
+  uint b = 0u;
+  const uint3 r = (a % value_or_one_if_zero_uint((b + b)));
+  return;
+}

test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  uint3 a = uint3(1u, 2u, 3u);
+  uint b = 0u;
+  uint3 const r = (a % (b + b));
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.spvasm

@@ -0,0 +1,40 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+               OpName %a "a"
+               OpName %b "b"
+       %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
+%_ptr_Function_v3uint = OpTypePointer Function %v3uint
+         %13 = OpConstantNull %v3uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %17 = OpConstantNull %uint
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %a = OpVariable %_ptr_Function_v3uint Function %13
+          %b = OpVariable %_ptr_Function_uint Function %17
+         %23 = OpVariable %_ptr_Function_v3uint Function %13
+               OpStore %a %10
+               OpStore %b %uint_0
+         %18 = OpLoad %v3uint %a
+         %19 = OpLoad %uint %b
+         %20 = OpLoad %uint %b
+         %21 = OpIAdd %uint %19 %20
+         %24 = OpCompositeConstruct %v3uint %21 %21 %21
+         %22 = OpUMod %v3uint %18 %24
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var a = vec3<u32>(1u, 2u, 3u);
+  var b = 0u;
+  let r : vec3<u32> = (a % (b + b));
+}

test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    var a = vec3<f32>(1., 2., 3.);
+    var b = vec3<f32>(0., 5., 0.);
+    let r : vec3<f32> = a % (b + b);
+}

test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.hlsl

@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+  float3 a = float3(1.0f, 2.0f, 3.0f);
+  float3 b = float3(0.0f, 5.0f, 0.0f);
+  const float3 r = (a % (b + b));
+  return;
+}

test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  float3 a = float3(1.0f, 2.0f, 3.0f);
+  float3 b = float3(0.0f, 5.0f, 0.0f);
+  float3 const r = fmod(a, (b + b));
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.spvasm

@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 23
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+               OpName %a "a"
+               OpName %b "b"
+       %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
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %13 = OpConstantNull %v3float
+    %float_0 = OpConstant %float 0
+    %float_5 = OpConstant %float 5
+         %16 = OpConstantComposite %v3float %float_0 %float_5 %float_0
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %a = OpVariable %_ptr_Function_v3float Function %13
+          %b = OpVariable %_ptr_Function_v3float Function %13
+               OpStore %a %10
+               OpStore %b %16
+         %18 = OpLoad %v3float %a
+         %19 = OpLoad %v3float %b
+         %20 = OpLoad %v3float %b
+         %21 = OpFAdd %v3float %19 %20
+         %22 = OpFRem %v3float %18 %21
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var a = vec3<f32>(1.0, 2.0, 3.0);
+  var b = vec3<f32>(0.0, 5.0, 0.0);
+  let r : vec3<f32> = (a % (b + b));
+}

test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    var a = vec3<i32>(1, 2, 3);
+    var b = vec3<i32>(0, 5, 0);
+    let r : vec3<i32> = a % (b + b);
+}

test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.hlsl

@@ -0,0 +1,11 @@
+int3 value_or_one_if_zero_int3(int3 value) {
+  return value == int3(0, 0, 0) ? int3(1, 1, 1) : value;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  int3 a = int3(1, 2, 3);
+  int3 b = int3(0, 5, 0);
+  const int3 r = (a % value_or_one_if_zero_int3((b + b)));
+  return;
+}

test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  int3 a = int3(1, 2, 3);
+  int3 b = int3(0, 5, 0);
+  int3 const r = (a % as_type<int3>((as_type<uint3>(b) + as_type<uint3>(b))));
+  return;
+}
+

test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.spvasm

@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 23
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+               OpName %a "a"
+               OpName %b "b"
+       %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
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+         %13 = OpConstantNull %v3int
+      %int_0 = OpConstant %int 0
+      %int_5 = OpConstant %int 5
+         %16 = OpConstantComposite %v3int %int_0 %int_5 %int_0
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %a = OpVariable %_ptr_Function_v3int Function %13
+          %b = OpVariable %_ptr_Function_v3int Function %13
+               OpStore %a %10
+               OpStore %b %16
+         %18 = OpLoad %v3int %a
+         %19 = OpLoad %v3int %b
+         %20 = OpLoad %v3int %b
+         %21 = OpIAdd %v3int %19 %20
+         %22 = OpSMod %v3int %18 %21
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  var a = vec3<i32>(1, 2, 3);
+  var b = vec3<i32>(0, 5, 0);
+  let r : vec3<i32> = (a % (b + b));
+}