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>
diff --git a/src/program_builder.h b/src/program_builder.h
index 53929af..b724328 100644
--- a/src/program_builder.h
+++ b/src/program_builder.h
@@ -1665,6 +1665,16 @@
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`
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 53188fb..9108f9f 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -863,8 +863,8 @@
break;
case ast::BinaryOp::kDivide:
out << "/";
- // BUG(crbug.com/tint/1083): Integer divide by zero is a FXC compile
- // error, and undefined behavior in WGSL.
+ // 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);
@@ -872,6 +872,12 @@
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,
diff --git a/src/writer/hlsl/generator_impl_binary_test.cc b/src/writer/hlsl/generator_impl_binary_test.cc
index 6443c48..199f9b1 100644
--- a/src/writer/hlsl/generator_impl_binary_test.cc
+++ b/src/writer/hlsl/generator_impl_binary_test.cc
@@ -137,10 +137,12 @@
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 @@
)");
}
-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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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
diff --git a/test/bug/chromium/1273230.wgsl.expected.hlsl b/test/bug/chromium/1273230.wgsl.expected.hlsl
index aef846b..eb84f1a 100644
--- a/test/bug/chromium/1273230.wgsl.expected.hlsl
+++ b/test/bug/chromium/1273230.wgsl.expected.hlsl
@@ -89,7 +89,7 @@
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);
}
diff --git a/test/bug/tint/1113.wgsl.expected.hlsl b/test/bug/tint/1113.wgsl.expected.hlsl
index 656837e..b2da18f 100644
--- a/test/bug/tint/1113.wgsl.expected.hlsl
+++ b/test/bug/tint/1113.wgsl.expected.hlsl
@@ -60,7 +60,7 @@
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);
}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl
new file mode 100644
index 0000000..b34099c
--- /dev/null
+++ b/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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..d3de7db
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.msl
new file mode 100644
index 0000000..25c2f18
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..71bf9df
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..235be93
--- /dev/null
+++ b/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl
new file mode 100644
index 0000000..9900e05
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..b79aad9
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.msl
new file mode 100644
index 0000000..98f1698
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..37536a1
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..4c6fdf9
--- /dev/null
+++ b/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl
new file mode 100644
index 0000000..4fb0f9f
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..241c5f5
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.msl
new file mode 100644
index 0000000..456e3b9
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..e59fd6e
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..62a525b
--- /dev/null
+++ b/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl
new file mode 100644
index 0000000..9b27df2
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..fa4482f
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.msl
new file mode 100644
index 0000000..95a93f7
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..dc1cd05
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..dca7442
--- /dev/null
+++ b/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl
new file mode 100644
index 0000000..494a712
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..12aa8c0
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.msl
new file mode 100644
index 0000000..16fa025
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..019876b
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..c061d41
--- /dev/null
+++ b/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl
new file mode 100644
index 0000000..2fe5f2d
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..6906266
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.msl
new file mode 100644
index 0000000..cbac4c7
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..ac1f1d7
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..0aac1fa
--- /dev/null
+++ b/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl
new file mode 100644
index 0000000..25b884f
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..4d8b0b4
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.msl
new file mode 100644
index 0000000..78398ec
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..b8ce74f
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..7665b70
--- /dev/null
+++ b/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl
new file mode 100644
index 0000000..61df86d
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..5ea2fbd
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.msl
new file mode 100644
index 0000000..9506b9c
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..6f5c427
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..180e7d9
--- /dev/null
+++ b/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl
new file mode 100644
index 0000000..23aa96c
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..ee60674
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.msl
new file mode 100644
index 0000000..7619c9c
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..5205a19
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..259840d
--- /dev/null
+++ b/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl
new file mode 100644
index 0000000..ee9eecc
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..2389839
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.msl
new file mode 100644
index 0000000..028424a
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..f1fa325
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..08943bb
--- /dev/null
+++ b/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl
new file mode 100644
index 0000000..a4e2202
--- /dev/null
+++ 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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..468c33d
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.msl
new file mode 100644
index 0000000..c3d64df
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..33268b0
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..b6c7b67
--- /dev/null
+++ b/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));
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl
new file mode 100644
index 0000000..7f8d8d0
--- /dev/null
+++ 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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..fd4f932
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.msl
new file mode 100644
index 0000000..ead2dce
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..a37441f
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..581f6df
--- /dev/null
+++ b/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));
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl
new file mode 100644
index 0000000..edd134b
--- /dev/null
+++ 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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..f8f7511
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.msl
new file mode 100644
index 0000000..33bf9ab
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..ae27d46
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..f894e6d
--- /dev/null
+++ b/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));
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl
new file mode 100644
index 0000000..b90a404
--- /dev/null
+++ 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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..8e8a7d0
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
new file mode 100644
index 0000000..12f1a72
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..b293226
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..724d89b
--- /dev/null
+++ b/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));
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl
new file mode 100644
index 0000000..e6626f1
--- /dev/null
+++ 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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..2330363
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.msl
new file mode 100644
index 0000000..a7511f2
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..ee0be7a
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..d820f98
--- /dev/null
+++ b/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));
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl
new file mode 100644
index 0000000..e5dd983
--- /dev/null
+++ 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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..ea6b4cb
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.msl
new file mode 100644
index 0000000..ba6626b
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..d9dd3f1
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..c924762
--- /dev/null
+++ b/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));
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl
new file mode 100644
index 0000000..b775a66
--- /dev/null
+++ 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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..dca1662
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.msl
new file mode 100644
index 0000000..96dba81
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..78e1a7e
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..d73c1a9
--- /dev/null
+++ b/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));
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl
new file mode 100644
index 0000000..92e19cf
--- /dev/null
+++ 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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..8d21550
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.msl
new file mode 100644
index 0000000..6ef4b08
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..fb8ce28
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..06e70ae
--- /dev/null
+++ b/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));
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl
new file mode 100644
index 0000000..e843464
--- /dev/null
+++ 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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..f33a43d
--- /dev/null
+++ 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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
new file mode 100644
index 0000000..61356dc
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..b96d65a
--- /dev/null
+++ b/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
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..78f7040
--- /dev/null
+++ b/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));
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl
new file mode 100644
index 0000000..ad50323
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ var a = vec3<u32>(1u, 2u, 3u);
+ var b = vec3<u32>(0u, 5u, 0u);
+ let r : vec3<u32> = a % (b + b);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..f112cf8
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-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() {
+ uint3 a = uint3(1u, 2u, 3u);
+ uint3 b = uint3(0u, 5u, 0u);
+ const uint3 r = (a % value_or_one_if_zero_uint3((b + b)));
+ return;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.msl
new file mode 100644
index 0000000..c6d0db5
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint3 a = uint3(1u, 2u, 3u);
+ uint3 b = uint3(0u, 5u, 0u);
+ uint3 const r = (a % (b + b));
+ return;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..c4b1f90
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.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
+ %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
+ %uint_5 = OpConstant %uint 5
+ %16 = OpConstantComposite %v3uint %uint_0 %uint_5 %uint_0
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_v3uint Function %13
+ %b = OpVariable %_ptr_Function_v3uint Function %13
+ OpStore %a %10
+ OpStore %b %16
+ %18 = OpLoad %v3uint %a
+ %19 = OpLoad %v3uint %b
+ %20 = OpLoad %v3uint %b
+ %21 = OpIAdd %v3uint %19 %20
+ %22 = OpUMod %v3uint %18 %21
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..e7f2ce8
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ var a = vec3<u32>(1u, 2u, 3u);
+ var b = vec3<u32>(0u, 5u, 0u);
+ let r : vec3<u32> = (a % (b + b));
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl
new file mode 100644
index 0000000..b80c201
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..3d43aff
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+ return;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.msl
new file mode 100644
index 0000000..4a82529
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+ return;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..14b6b44
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.spvasm
@@ -0,0 +1,30 @@
+; 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"
+ 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 = OpFRem %float %12 %13
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..1b3af01
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl
new file mode 100644
index 0000000..38874c4
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..9f27720
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ int a = 1;
+ int b = 0;
+ const int r = (a % (b == 0 ? 1 : b));
+ return;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.msl
new file mode 100644
index 0000000..4a51c54
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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 % b);
+ return;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..f942f0b
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.spvasm
@@ -0,0 +1,30 @@
+; 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"
+ 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 = OpSMod %int %12 %13
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..1b19397
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl
new file mode 100644
index 0000000..a8b6555
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..cc3d3c1
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ uint a = 1u;
+ uint b = 0u;
+ const uint r = (a % (b == 0u ? 1u : b));
+ return;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.msl
new file mode 100644
index 0000000..8849d0d
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+ return;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..8d29bb6
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.spvasm
@@ -0,0 +1,30 @@
+; 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"
+ 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 = OpUMod %uint %12 %13
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..1a5f726
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl
new file mode 100644
index 0000000..0f9c970
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..0b14c77
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ int a = 4;
+ int3 b = int3(0, 2, 0);
+ const int3 r = (a % (b == int3(0, 0, 0) ? int3(1, 1, 1) : b));
+ return;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.msl
new file mode 100644
index 0000000..ae32a86
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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 % b);
+ return;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..1098fb2
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.spvasm
@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 22
+; 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
+ %20 = OpVariable %_ptr_Function_v3int Function %16
+ OpStore %a %int_4
+ OpStore %b %13
+ %17 = OpLoad %int %a
+ %18 = OpLoad %v3int %b
+ %21 = OpCompositeConstruct %v3int %17 %17 %17
+ %19 = OpSMod %v3int %21 %18
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..30f19f4
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl
new file mode 100644
index 0000000..9afcb3c
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..fa88822
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ uint a = 4u;
+ uint3 b = uint3(0u, 2u, 0u);
+ const uint3 r = (a % (b == uint3(0u, 0u, 0u) ? uint3(1u, 1u, 1u) : b));
+ return;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.msl
new file mode 100644
index 0000000..bc4bada
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+ return;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..45bce24
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.spvasm
@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 22
+; 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
+ %20 = OpVariable %_ptr_Function_v3uint Function %16
+ OpStore %a %uint_4
+ OpStore %b %13
+ %17 = OpLoad %uint %a
+ %18 = OpLoad %v3uint %b
+ %21 = OpCompositeConstruct %v3uint %17 %17 %17
+ %19 = OpUMod %v3uint %21 %18
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..0e406f7
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl
new file mode 100644
index 0000000..41f8c7c
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..208637e
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ int3 a = int3(1, 2, 3);
+ int b = 0;
+ const int3 r = (a % (b == 0 ? 1 : b));
+ return;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.msl
new file mode 100644
index 0000000..39d00f8
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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 % b);
+ return;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..f25b333
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/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
+%_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
+ %21 = OpVariable %_ptr_Function_v3int Function %13
+ OpStore %a %10
+ OpStore %b %int_0
+ %18 = OpLoad %v3int %a
+ %19 = OpLoad %int %b
+ %22 = OpCompositeConstruct %v3int %19 %19 %19
+ %20 = OpSMod %v3int %18 %22
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..dcb6939
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl
new file mode 100644
index 0000000..5ddae5b
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..0cb79cf
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ uint3 a = uint3(1u, 2u, 3u);
+ uint b = 0u;
+ const uint3 r = (a % (b == 0u ? 1u : b));
+ return;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.msl
new file mode 100644
index 0000000..3c8b16e
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+ return;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..0965b4a
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.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
+ %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
+ %21 = OpVariable %_ptr_Function_v3uint Function %13
+ OpStore %a %10
+ OpStore %b %uint_0
+ %18 = OpLoad %v3uint %a
+ %19 = OpLoad %uint %b
+ %22 = OpCompositeConstruct %v3uint %19 %19 %19
+ %20 = OpUMod %v3uint %18 %22
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..eabfbcd
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl
new file mode 100644
index 0000000..7a91bd0
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..636719b
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+ return;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.msl
new file mode 100644
index 0000000..4ebdd54
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+ return;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..1cdd95a
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.spvasm
@@ -0,0 +1,36 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 21
+; 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 = OpFRem %v3float %18 %19
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..6d96f89
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl
new file mode 100644
index 0000000..e9364db
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..97e59e6
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ int3 a = int3(1, 2, 3);
+ int3 b = int3(0, 5, 0);
+ const int3 r = (a % (b == int3(0, 0, 0) ? int3(1, 1, 1) : b));
+ return;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.msl
new file mode 100644
index 0000000..e6b07ad
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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 % b);
+ return;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..29c04a6
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.spvasm
@@ -0,0 +1,36 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 21
+; 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 = OpSMod %v3int %18 %19
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..ebd6b91
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/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);
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl
new file mode 100644
index 0000000..cf77b25
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ var a = vec3<u32>(1u, 2u, 3u);
+ var b = vec3<u32>(0u, 5u, 0u);
+ let r : vec3<u32> = a % b;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..b041731
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ uint3 a = uint3(1u, 2u, 3u);
+ uint3 b = uint3(0u, 5u, 0u);
+ const uint3 r = (a % (b == uint3(0u, 0u, 0u) ? uint3(1u, 1u, 1u) : b));
+ return;
+}
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.msl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.msl
new file mode 100644
index 0000000..46e36b0
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ uint3 a = uint3(1u, 2u, 3u);
+ uint3 b = uint3(0u, 5u, 0u);
+ uint3 const r = (a % b);
+ return;
+}
+
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..3fbdc22
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.spvasm
@@ -0,0 +1,36 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 21
+; 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
+ %uint_5 = OpConstant %uint 5
+ %16 = OpConstantComposite %v3uint %uint_0 %uint_5 %uint_0
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_v3uint Function %13
+ %b = OpVariable %_ptr_Function_v3uint Function %13
+ OpStore %a %10
+ OpStore %b %16
+ %18 = OpLoad %v3uint %a
+ %19 = OpLoad %v3uint %b
+ %20 = OpUMod %v3uint %18 %19
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..615b56c
--- /dev/null
+++ b/test/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ var a = vec3<u32>(1u, 2u, 3u);
+ var b = vec3<u32>(0u, 5u, 0u);
+ let r : vec3<u32> = (a % b);
+}
diff --git a/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.spvasm.expected.hlsl
index 535c006..0d695f6 100644
--- a/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.spvasm.expected.hlsl
+++ b/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.spvasm.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
cbuffer cbuffer_x_6 : register(b0, space0) {
uint4 x_6[4];
};
@@ -18,7 +16,7 @@
if ((x_35 == x_37)) {
a = (a + 1);
} else {
- a = (a / i);
+ a = (a / (i == 0 ? 1 : i));
}
}
}
@@ -59,6 +57,3 @@
wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1;
return wrapper_result;
}
-C:\src\tint\test\Shader@0x00000150DB0FB4A0(19,14-18): error X4010: Unsigned integer divide by zero
-C:\src\tint\test\Shader@0x00000150DB0FB4A0(19,14-18): warning X3556: integer divides may be much slower, try using uints if possible.
-
diff --git a/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.wgsl.expected.hlsl
index c8f9a07..0d695f6 100644
--- a/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.wgsl.expected.hlsl
+++ b/test/vk-gl-cts/graphicsfuzz/cov-loop-increment-or-divide-by-loop-index/0-opt.wgsl.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
cbuffer cbuffer_x_6 : register(b0, space0) {
uint4 x_6[4];
};
@@ -18,7 +16,7 @@
if ((x_35 == x_37)) {
a = (a + 1);
} else {
- a = (a / i);
+ a = (a / (i == 0 ? 1 : i));
}
}
}
@@ -59,6 +57,3 @@
wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1;
return wrapper_result;
}
-C:\src\tint\test\Shader@0x000001A53C9E4360(19,14-18): error X4010: Unsigned integer divide by zero
-C:\src\tint\test\Shader@0x000001A53C9E4360(19,14-18): warning X3556: integer divides may be much slower, try using uints if possible.
-
diff --git a/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.spvasm.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.spvasm.expected.hlsl
index 1589466..8455574 100644
--- a/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.spvasm.expected.hlsl
+++ b/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.spvasm.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
cbuffer cbuffer_x_8 : register(b2, space0) {
uint4 x_8[2];
};
@@ -24,7 +22,7 @@
if ((x_43 < x_45)) {
const uint scalar_offset_1 = ((16u * uint(0))) / 4;
const uint x_50 = x_12[scalar_offset_1 / 4][scalar_offset_1 % 4];
- b = asint((x_50 % a));
+ b = asint((x_50 % (a == 0u ? 1u : a)));
}
const int x_54 = b;
const int x_56 = asint(x_8[1].x);
@@ -68,5 +66,3 @@
wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1;
return wrapper_result;
}
-C:\src\tint\test\Shader@0x000002533EADD8C0(25,16-23): error X4010: Unsigned integer divide by zero
-
diff --git a/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.wgsl.expected.hlsl b/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.wgsl.expected.hlsl
index 4cecdf3..8455574 100644
--- a/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.wgsl.expected.hlsl
+++ b/test/vk-gl-cts/graphicsfuzz/cov-modulo-zero-never-executed/0-opt.wgsl.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
cbuffer cbuffer_x_8 : register(b2, space0) {
uint4 x_8[2];
};
@@ -24,7 +22,7 @@
if ((x_43 < x_45)) {
const uint scalar_offset_1 = ((16u * uint(0))) / 4;
const uint x_50 = x_12[scalar_offset_1 / 4][scalar_offset_1 % 4];
- b = asint((x_50 % a));
+ b = asint((x_50 % (a == 0u ? 1u : a)));
}
const int x_54 = b;
const int x_56 = asint(x_8[1].x);
@@ -68,5 +66,3 @@
wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1;
return wrapper_result;
}
-C:\src\tint\test\Shader@0x0000026AFDC08170(25,16-23): error X4010: Unsigned integer divide by zero
-