Fix FXC compile errors on divide by zero
FXC fails to compile when it determines that the rhs of an integral
division is zero with "error X4010: Unsigned integer divide by zero".
bclayton's fix (https://dawn-review.googlesource.com/c/tint/+/60500)
addressed cases for division by an integer constant 0. This CL adds the
missing support for division by integral vectors with 0 components.
FXC also fails on division by integral expressions that it can fold to
0. To handle these cases, we now emit a runtime check for 0 and replace
by 1. In the cases I've tested, FXC seems able to optimize these checks
away.
Bug: tint:1083
Change-Id: I02f08e9077882f03c1e42b62dacb742a48fa48ba
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/73580
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
diff --git a/src/sem/constant.cc b/src/sem/constant.cc
index b39a115..df2e67e 100644
--- a/src/sem/constant.cc
+++ b/src/sem/constant.cc
@@ -14,6 +14,7 @@
#include "src/sem/constant.h"
+#include <functional>
#include <utility>
#include "src/debug.h"
@@ -61,5 +62,23 @@
Constant& Constant::operator=(const Constant& rhs) = default;
+bool Constant::AnyZero() const {
+ for (size_t i = 0; i < Elements().size(); ++i) {
+ if (WithScalarAt(i, [&](auto&& s) {
+ // Use std::equal_to to work around -Wfloat-equal warnings
+ auto equals_to =
+ std::equal_to<std::remove_reference_t<decltype(s)>>{};
+
+ if (equals_to(s, 0)) {
+ return true;
+ }
+ return false;
+ })) {
+ return true;
+ }
+ }
+ return false;
+}
+
} // namespace sem
} // namespace tint
diff --git a/src/sem/constant.h b/src/sem/constant.h
index d28ee7b..4f5e198 100644
--- a/src/sem/constant.h
+++ b/src/sem/constant.h
@@ -97,6 +97,9 @@
/// @returns the constant's scalar elements
const Scalars& Elements() const { return elems_; }
+ /// @returns true if any scalar element is zero
+ bool AnyZero() const;
+
/// Calls `func(s)` with s being the current scalar value at `index`.
/// `func` is typically a lambda of the form '[](auto&& s)'.
/// @param index the index of the scalar value
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 82c36cb..53188fb 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -16,6 +16,7 @@
#include <algorithm>
#include <cmath>
+#include <functional>
#include <iomanip>
#include <set>
#include <utility>
@@ -618,6 +619,127 @@
return true;
}
+bool GeneratorImpl::EmitExpressionOrOneIfZero(std::ostream& out,
+ const ast::Expression* expr) {
+ // For constants, replace literal 0 with 1.
+ sem::Constant::Scalars elems;
+ if (const auto& val = builder_.Sem().Get(expr)->ConstantValue()) {
+ if (!val.AnyZero()) {
+ return EmitExpression(out, expr);
+ }
+
+ if (val.Type()->IsAnyOf<sem::I32, sem::U32>()) {
+ return EmitValue(out, val.Type(), 1);
+ }
+
+ if (auto* vec = val.Type()->As<sem::Vector>()) {
+ auto* elem_ty = vec->type();
+
+ if (!EmitType(out, val.Type(), ast::StorageClass::kNone,
+ ast::Access::kUndefined, "")) {
+ return false;
+ }
+
+ out << "(";
+ for (size_t i = 0; i < val.Elements().size(); ++i) {
+ if (i != 0) {
+ out << ", ";
+ }
+ if (!val.WithScalarAt(i, [&](auto&& s) -> bool {
+ // Use std::equal_to to work around -Wfloat-equal warnings
+ auto equals_to =
+ std::equal_to<std::remove_reference_t<decltype(s)>>{};
+
+ bool is_zero = equals_to(s, 0);
+ return EmitValue(out, elem_ty, is_zero ? 1 : static_cast<int>(s));
+ })) {
+ return false;
+ }
+ }
+ out << ")";
+ return true;
+ }
+
+ TINT_ICE(Writer, diagnostics_)
+ << "EmitExpressionOrOneIfZero expects integer scalar or vector";
+ return false;
+ }
+
+ auto* ty = TypeOf(expr)->UnwrapRef();
+
+ // For non-constants, we need to emit runtime code to check if the value is 0,
+ // and return 1 in that case.
+ std::string zero;
+ {
+ std::ostringstream ss;
+ EmitValue(ss, ty, 0);
+ zero = ss.str();
+ }
+ std::string one;
+ {
+ std::ostringstream ss;
+ EmitValue(ss, ty, 1);
+ one = ss.str();
+ }
+
+ // For identifiers, no need for a function call as it's fine to evaluate
+ // `expr` more than once.
+ if (expr->Is<ast::IdentifierExpression>()) {
+ out << "(";
+ if (!EmitExpression(out, expr)) {
+ return false;
+ }
+ out << " == " << zero << " ? " << one << " : ";
+ if (!EmitExpression(out, expr)) {
+ return false;
+ }
+ out << ")";
+ return true;
+ }
+
+ // For non-identifier expressions, call a function to make sure `expr` is only
+ // evaluated once.
+ auto name =
+ utils::GetOrCreate(value_or_one_if_zero_, ty, [&]() -> std::string {
+ // Example:
+ // int4 tint_value_or_one_if_zero_int4(int4 value) {
+ // return value == 0 ? 0 : value;
+ // }
+ std::string ty_name;
+ {
+ std::ostringstream ss;
+ if (!EmitType(ss, ty, tint::ast::StorageClass::kInvalid,
+ ast::Access::kUndefined, "")) {
+ return "";
+ }
+ ty_name = ss.str();
+ }
+
+ std::string fn = UniqueIdentifier("value_or_one_if_zero_" + ty_name);
+ line(&helpers_) << ty_name << " " << fn << "(" << ty_name
+ << " value) {";
+ {
+ ScopedIndent si(&helpers_);
+ line(&helpers_) << "return value == " << zero << " ? " << one
+ << " : value;";
+ }
+ line(&helpers_) << "}";
+ line(&helpers_);
+ return fn;
+ });
+
+ if (name.empty()) {
+ return false;
+ }
+
+ out << name << "(";
+ if (!EmitExpression(out, expr)) {
+ return false;
+ }
+ out << ")";
+ return true;
+}
+
bool GeneratorImpl::EmitBinary(std::ostream& out,
const ast::BinaryExpression* expr) {
if (expr->op == ast::BinaryOp::kLogicalAnd ||
@@ -741,18 +863,11 @@
break;
case ast::BinaryOp::kDivide:
out << "/";
-
- if (auto val = builder_.Sem().Get(expr->rhs)->ConstantValue()) {
- // Integer divide by zero is a DXC compile error, and undefined behavior
- // in WGSL. Replace the 0 with 1.
- if (val.Type()->Is<sem::I32>() && val.Elements()[0].i32 == 0) {
- out << " 1";
- return true;
- }
- if (val.Type()->Is<sem::U32>() && val.Elements()[0].u32 == 0u) {
- out << " 1u";
- return true;
- }
+ // BUG(crbug.com/tint/1083): Integer divide 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::kModulo:
@@ -3006,15 +3121,17 @@
return true;
}
-bool GeneratorImpl::EmitZeroValue(std::ostream& out, const sem::Type* type) {
+bool GeneratorImpl::EmitValue(std::ostream& out,
+ const sem::Type* type,
+ int value) {
if (type->Is<sem::Bool>()) {
- out << "false";
+ out << (value == 0 ? "false" : "true");
} else if (type->Is<sem::F32>()) {
- out << "0.0f";
+ out << value << ".0f";
} else if (type->Is<sem::I32>()) {
- out << "0";
+ out << value;
} else if (type->Is<sem::U32>()) {
- out << "0u";
+ out << value << "u";
} else if (auto* vec = type->As<sem::Vector>()) {
if (!EmitType(out, type, ast::StorageClass::kNone, ast::Access::kReadWrite,
"")) {
@@ -3025,7 +3142,7 @@
if (i != 0) {
out << ", ";
}
- if (!EmitZeroValue(out, vec->type())) {
+ if (!EmitValue(out, vec->type(), value)) {
return false;
}
}
@@ -3039,7 +3156,7 @@
if (i != 0) {
out << ", ";
}
- if (!EmitZeroValue(out, mat->type())) {
+ if (!EmitValue(out, mat->type(), value)) {
return false;
}
}
@@ -3049,16 +3166,20 @@
"")) {
return false;
}
- out << ")0";
+ out << ")" << value;
} else {
diagnostics_.add_error(
diag::System::Writer,
- "Invalid type for zero emission: " + type->type_name());
+ "Invalid type for value emission: " + type->type_name());
return false;
}
return true;
}
+bool GeneratorImpl::EmitZeroValue(std::ostream& out, const sem::Type* type) {
+ return EmitValue(out, type, 0);
+}
+
bool GeneratorImpl::EmitLoop(const ast::LoopStatement* stmt) {
auto emit_continuing = [this, stmt]() {
if (stmt->continuing && !stmt->continuing->Empty()) {
diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h
index 12124d2..97cb5d7 100644
--- a/src/writer/hlsl/generator_impl.h
+++ b/src/writer/hlsl/generator_impl.h
@@ -100,6 +100,12 @@
/// @param stmt the statement to emit
/// @returns true if the statement was emitted successfully
bool EmitAssign(const ast::AssignmentStatement* stmt);
+ /// Emits code such that if `expr` is zero, it emits one, else `expr`
+ /// @param out the output of the expression stream
+ /// @param expr the expression
+ /// @returns true if the expression was emitted, false otherwise
+ bool EmitExpressionOrOneIfZero(std::ostream& out,
+ const ast::Expression* expr);
/// Handles generating a binary expression
/// @param out the output of the expression stream
/// @param expr the binary expression
@@ -401,6 +407,12 @@
/// @param expr the expression to emit
/// @returns true if the expression was emitted
bool EmitUnaryOp(std::ostream& out, const ast::UnaryOpExpression* expr);
+ /// Emits `value` for the given type
+ /// @param out the output stream
+ /// @param type the type to emit the value for
+ /// @param value the value to emit
+ /// @returns true if the value was successfully emitted.
+ bool EmitValue(std::ostream& out, const sem::Type* type, int value);
/// Emits the zero value for the given type
/// @param out the output stream
/// @param type the type to emit the value for
@@ -515,6 +527,7 @@
dynamic_matrix_vector_write_;
std::unordered_map<const sem::Matrix*, std::string>
dynamic_matrix_scalar_write_;
+ std::unordered_map<const sem::Type*, std::string> value_or_one_if_zero_;
};
} // namespace hlsl
diff --git a/src/writer/hlsl/generator_impl_binary_test.cc b/src/writer/hlsl/generator_impl_binary_test.cc
index c14a89a..6443c48 100644
--- a/src/writer/hlsl/generator_impl_binary_test.cc
+++ b/src/writer/hlsl/generator_impl_binary_test.cc
@@ -26,6 +26,9 @@
struct BinaryData {
const char* result;
ast::BinaryOp op;
+
+ enum Types { All = 0b11, Integer = 0b10, Float = 0b01 };
+ Types valid_for = Types::All;
};
inline std::ostream& operator<<(std::ostream& out, BinaryData data) {
out << data.op;
@@ -36,6 +39,10 @@
TEST_P(HlslBinaryTest, Emit_f32) {
auto params = GetParam();
+ if ((params.valid_for & BinaryData::Types::Float) == 0) {
+ return;
+ }
+
// Skip ops that are illegal for this type
if (params.op == ast::BinaryOp::kAnd || params.op == ast::BinaryOp::kOr ||
params.op == ast::BinaryOp::kXor ||
@@ -63,6 +70,10 @@
TEST_P(HlslBinaryTest, Emit_u32) {
auto params = GetParam();
+ if ((params.valid_for & BinaryData::Types::Integer) == 0) {
+ return;
+ }
+
Global("left", ty.u32(), ast::StorageClass::kPrivate);
Global("right", ty.u32(), ast::StorageClass::kPrivate);
@@ -82,6 +93,10 @@
TEST_P(HlslBinaryTest, Emit_i32) {
auto params = GetParam();
+ if ((params.valid_for & BinaryData::Types::Integer) == 0) {
+ return;
+ }
+
// Skip ops that are illegal for this type
if (params.op == ast::BinaryOp::kShiftLeft ||
params.op == ast::BinaryOp::kShiftRight) {
@@ -122,7 +137,9 @@
BinaryData{"(left + right)", ast::BinaryOp::kAdd},
BinaryData{"(left - right)", ast::BinaryOp::kSubtract},
BinaryData{"(left * right)", ast::BinaryOp::kMultiply},
- BinaryData{"(left / right)", ast::BinaryOp::kDivide},
+ // NOTE: Integer divide covered by DivideBy* tests below
+ BinaryData{"(left / right)", ast::BinaryOp::kDivide,
+ BinaryData::Types::Float},
BinaryData{"(left % right)", ast::BinaryOp::kModulo}));
TEST_F(HlslGeneratorImplTest_Binary, Multiply_VectorScalar) {
@@ -510,29 +527,259 @@
}
TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_i32) {
- Global("a", ty.i32(), ast::StorageClass::kPrivate);
-
- auto* expr = Div("a", 0);
- WrapInFunction(expr);
+ Func("fn", {}, ty.void_(),
+ {
+ Decl(Var("a", ty.i32())),
+ Decl(Const("r", nullptr, Div("a", 0))),
+ });
GeneratorImpl& gen = Build();
- std::stringstream out;
- ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
- EXPECT_EQ(out.str(), R"((a / 1))");
+ ASSERT_TRUE(gen.Generate());
+ EXPECT_EQ(gen.result(), R"(void fn() {
+ int a = 0;
+ const int r = (a / 1);
+}
+)");
}
TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_u32) {
- Global("a", ty.u32(), ast::StorageClass::kPrivate);
-
- auto* expr = Div("a", 0u);
- WrapInFunction(expr);
+ Func("fn", {}, ty.void_(),
+ {
+ Decl(Var("a", ty.u32())),
+ Decl(Const("r", nullptr, Div("a", 0u))),
+ });
GeneratorImpl& gen = Build();
- std::stringstream out;
- ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
- EXPECT_EQ(out.str(), R"((a / 1u))");
+ ASSERT_TRUE(gen.Generate());
+ EXPECT_EQ(gen.result(), R"(void fn() {
+ uint a = 0u;
+ const uint r = (a / 1u);
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_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)))),
+ });
+
+ GeneratorImpl& gen = Build();
+
+ 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));
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_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))),
+ });
+
+ GeneratorImpl& gen = Build();
+
+ ASSERT_TRUE(gen.Generate());
+ EXPECT_EQ(gen.result(), R"(void fn() {
+ int4 a = int4(100, 100, 100, 100);
+ const int4 r = (a / 1);
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_i32) {
+ Func("fn", {Param("b", ty.i32())}, ty.void_(),
+ {
+ Decl(Var("a", ty.i32())),
+ Decl(Const("r", nullptr, Div("a", "b"))),
+ });
+
+ GeneratorImpl& gen = Build();
+
+ ASSERT_TRUE(gen.Generate());
+ EXPECT_EQ(gen.result(), R"(void fn(int b) {
+ int a = 0;
+ const int r = (a / (b == 0 ? 1 : b));
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_u32) {
+ Func("fn", {Param("b", ty.u32())}, ty.void_(),
+ {
+ Decl(Var("a", ty.u32())),
+ Decl(Const("r", nullptr, Div("a", "b"))),
+ });
+
+ GeneratorImpl& gen = Build();
+
+ ASSERT_TRUE(gen.Generate());
+ EXPECT_EQ(gen.result(), R"(void fn(uint b) {
+ uint a = 0u;
+ const uint r = (a / (b == 0u ? 1u : b));
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_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"))),
+ });
+
+ GeneratorImpl& gen = Build();
+
+ 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));
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_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"))),
+ });
+
+ GeneratorImpl& gen = Build();
+
+ 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));
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_i32) {
+ Func("zero", {}, ty.i32(),
+ {
+ Return(Expr(0)),
+ });
+
+ Func("fn", {}, ty.void_(),
+ {
+ Decl(Var("a", ty.i32())),
+ Decl(Const("r", nullptr, Div("a", Call("zero")))),
+ });
+
+ GeneratorImpl& gen = Build();
+
+ ASSERT_TRUE(gen.Generate());
+ EXPECT_EQ(gen.result(), R"(int value_or_one_if_zero_int(int value) {
+ return value == 0 ? 1 : value;
+}
+
+int zero() {
+ return 0;
+}
+
+void fn() {
+ int a = 0;
+ const int r = (a / value_or_one_if_zero_int(zero()));
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_u32) {
+ Func("zero", {}, ty.u32(),
+ {
+ Return(Expr(0u)),
+ });
+
+ Func("fn", {}, ty.void_(),
+ {
+ Decl(Var("a", ty.u32())),
+ Decl(Const("r", nullptr, Div("a", Call("zero")))),
+ });
+
+ GeneratorImpl& gen = Build();
+
+ ASSERT_TRUE(gen.Generate());
+ EXPECT_EQ(gen.result(), R"(uint value_or_one_if_zero_uint(uint value) {
+ return value == 0u ? 1u : value;
+}
+
+uint zero() {
+ return 0u;
+}
+
+void fn() {
+ uint a = 0u;
+ const uint r = (a / value_or_one_if_zero_uint(zero()));
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_vec_by_vec_i32) {
+ Func("zero", {}, ty.vec3<i32>(),
+ {
+ Return(vec3<i32>(0, 0, 0)),
+ });
+
+ Func("fn", {}, ty.void_(),
+ {
+ Decl(Var("a", ty.vec3<i32>())),
+ Decl(Const("r", nullptr, Div("a", Call("zero")))),
+ });
+
+ GeneratorImpl& gen = Build();
+
+ ASSERT_TRUE(gen.Generate());
+ EXPECT_EQ(gen.result(), R"(int3 value_or_one_if_zero_int3(int3 value) {
+ return value == int3(0, 0, 0) ? int3(1, 1, 1) : value;
+}
+
+int3 zero() {
+ return int3(0, 0, 0);
+}
+
+void fn() {
+ int3 a = int3(0, 0, 0);
+ const int3 r = (a / value_or_one_if_zero_int3(zero()));
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_vec_by_scalar_i32) {
+ Func("zero", {}, ty.i32(),
+ {
+ Return(0),
+ });
+
+ Func("fn", {}, ty.void_(),
+ {
+ Decl(Var("a", ty.vec3<i32>())),
+ Decl(Const("r", nullptr, Div("a", Call("zero")))),
+ });
+
+ GeneratorImpl& gen = Build();
+
+ ASSERT_TRUE(gen.Generate());
+ EXPECT_EQ(gen.result(), R"(int value_or_one_if_zero_int(int value) {
+ return value == 0 ? 1 : value;
+}
+
+int zero() {
+ return 0;
+}
+
+void fn() {
+ int3 a = int3(0, 0, 0);
+ const int3 r = (a / value_or_one_if_zero_int(zero()));
+}
+)");
}
} // namespace
diff --git a/test/bug/chromium/1273230.wgsl.expected.hlsl b/test/bug/chromium/1273230.wgsl.expected.hlsl
index f721475..aef846b 100644
--- a/test/bug/chromium/1273230.wgsl.expected.hlsl
+++ b/test/bug/chromium/1273230.wgsl.expected.hlsl
@@ -30,6 +30,10 @@
return clamped == exponent;
}
+uint value_or_one_if_zero_uint(uint value) {
+ return value == 0u ? 1u : value;
+}
+
uint atomicLoad_1(RWByteAddressBuffer buffer, uint offset) {
uint value = 0;
buffer.InterlockedOr(offset, 0, value);
@@ -83,8 +87,8 @@
}
uint3 toIndex4D(uint gridSize, uint index) {
- uint z_1 = (gridSize / (index * index));
- uint y_1 = ((gridSize - ((gridSize * gridSize) * z_1)) / gridSize);
+ 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);
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 e630a4c..656837e 100644
--- a/test/bug/tint/1113.wgsl.expected.hlsl
+++ b/test/bug/tint/1113.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+uint value_or_one_if_zero_uint(uint value) {
+ return value == 0u ? 1u : value;
+}
+
uint atomicLoad_1(RWByteAddressBuffer buffer, uint offset) {
uint value = 0;
buffer.InterlockedOr(offset, 0, value);
@@ -54,8 +58,8 @@
}
uint3 toIndex3D(uint gridSize, uint index) {
- uint z_1 = (index / (gridSize * gridSize));
- uint y_1 = ((index - ((gridSize * gridSize) * z_1)) / gridSize);
+ 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);
return uint3(x_1, y_1, z_1);
}
diff --git a/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/f32.wgsl b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/f32.wgsl
new file mode 100644
index 0000000..7ce4276
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..a06ee90
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.msl
new file mode 100644
index 0000000..b0953f8
--- /dev/null
+++ b/test/expressions/binary/div_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 = (a / b);
+ return;
+}
+
diff --git a/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..aba66b3
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpFDiv %float %float_1 %float_0
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..da2c352
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-scalar/i32.wgsl b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/i32.wgsl
new file mode 100644
index 0000000..7c980aa
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..67f002e
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.msl
new file mode 100644
index 0000000..0923684
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..47589eb
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpSDiv %int %int_1 %int_0
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..63bdb5a
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-scalar/u32.wgsl b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/u32.wgsl
new file mode 100644
index 0000000..e31e484
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..60e18ed
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.msl
new file mode 100644
index 0000000..2c1b78a
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..96e2a30
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpUDiv %uint %uint_1 %uint_0
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..3b06f40
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-vec3/f32.wgsl b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/f32.wgsl
new file mode 100644
index 0000000..83e438c
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/f32.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a = 4.;
+ let b = vec3<f32>(0., 2., 0.);
+ let r : vec3<f32> = a / b;
+}
diff --git a/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/f32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..85db3bf
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/f32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float a = 4.0f;
+ const float3 b = float3(0.0f, 2.0f, 0.0f);
+ const float3 r = (a / b);
+ return;
+}
diff --git a/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/f32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/f32.wgsl.expected.msl
new file mode 100644
index 0000000..e251ea7
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/f32.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float const a = 4.0f;
+ float3 const b = float3(0.0f, 2.0f, 0.0f);
+ float3 const r = (a / b);
+ return;
+}
+
diff --git a/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/f32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..517826d
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/f32.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
+ %float = OpTypeFloat 32
+ %float_4 = OpConstant %float 4
+ %v3float = OpTypeVector %float 3
+ %float_0 = OpConstant %float 0
+ %float_2 = OpConstant %float 2
+ %10 = OpConstantComposite %v3float %float_0 %float_2 %float_0
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+ %14 = OpConstantNull %v3float
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %12 = OpVariable %_ptr_Function_v3float Function %14
+ %15 = OpCompositeConstruct %v3float %float_4 %float_4 %float_4
+ %11 = OpFDiv %v3float %15 %10
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/f32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..040875e
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/f32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a = 4.0;
+ let b = vec3<f32>(0.0, 2.0, 0.0);
+ let r : vec3<f32> = (a / b);
+}
diff --git a/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/i32.wgsl b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/i32.wgsl
new file mode 100644
index 0000000..b981e35
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..35cc532
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.msl
new file mode 100644
index 0000000..89df256
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..7cf0d58
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpSDiv %v3int %15 %10
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..34804c3
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-vec3/u32.wgsl b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/u32.wgsl
new file mode 100644
index 0000000..ae85d37
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..f57e7ed
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.msl
new file mode 100644
index 0000000..eb1c975
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..af6d858
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpUDiv %v3uint %15 %10
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..a20cb6a
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-scalar/f32.wgsl b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/f32.wgsl
new file mode 100644
index 0000000..0a15e5b
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/f32.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a = vec3<f32>(1., 2., 3.);
+ let b = 0.;
+ let r : vec3<f32> = a / b;
+}
diff --git a/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/f32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..c661ab1
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/f32.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void f() {
+ const float3 a = float3(1.0f, 2.0f, 3.0f);
+ const float3 r = (a / 0.0f);
+ return;
+}
diff --git a/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/f32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/f32.wgsl.expected.msl
new file mode 100644
index 0000000..91df91e
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/f32.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float3 const a = float3(1.0f, 2.0f, 3.0f);
+ float const b = 0.0f;
+ float3 const r = (a / b);
+ return;
+}
+
diff --git a/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/f32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..cc78cc9
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/f32.wgsl.expected.spvasm
@@ -0,0 +1,28 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %f "f"
+ OpExecutionMode %f LocalSize 1 1 1
+ OpName %f "f"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+ %float_1 = OpConstant %float 1
+ %float_2 = OpConstant %float 2
+ %float_3 = OpConstant %float 3
+ %10 = OpConstantComposite %v3float %float_1 %float_2 %float_3
+ %float_0 = OpConstant %float 0
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+ %15 = OpConstantNull %v3float
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %13 = OpVariable %_ptr_Function_v3float Function %15
+ %16 = OpCompositeConstruct %v3float %float_0 %float_0 %float_0
+ %12 = OpFDiv %v3float %10 %16
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/f32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..7396529
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/f32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ let a = vec3<f32>(1.0, 2.0, 3.0);
+ let b = 0.0;
+ let r : vec3<f32> = (a / b);
+}
diff --git a/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/i32.wgsl b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/i32.wgsl
new file mode 100644
index 0000000..93243bc
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..8a44ac5
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.msl
new file mode 100644
index 0000000..c187dd2
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..7f5c03f
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpSDiv %v3int %10 %16
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..c15fc0d
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-scalar/u32.wgsl b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/u32.wgsl
new file mode 100644
index 0000000..e4cc9f0
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..18d1da5
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.msl
new file mode 100644
index 0000000..2d342e7
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..a81a4c2
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpUDiv %v3uint %10 %16
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..a7aaf27
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-vec3/f32.wgsl b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/f32.wgsl
new file mode 100644
index 0000000..3800a7b
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..94b990b
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.msl
new file mode 100644
index 0000000..a4765b3
--- /dev/null
+++ b/test/expressions/binary/div_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 = (a / b);
+ return;
+}
+
diff --git a/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..5a8a4d8
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpFDiv %v3float %10 %13
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..b7be414
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-vec3/i32.wgsl b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/i32.wgsl
new file mode 100644
index 0000000..6f5bac7
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..3cd640e
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.msl
new file mode 100644
index 0000000..7cf2681
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..8176141
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpSDiv %v3int %10 %13
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..72c4004
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-vec3/u32.wgsl b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/u32.wgsl
new file mode 100644
index 0000000..4d73587
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..ea7ca06
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.msl
new file mode 100644
index 0000000..66ee158
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..276d688
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpUDiv %v3uint %10 %13
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..03c8910
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-scalar/f32.wgsl b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/f32.wgsl
new file mode 100644
index 0000000..12f1a15
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..bf3a16e
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.msl
new file mode 100644
index 0000000..d317b59
--- /dev/null
+++ b/test/expressions/binary/div_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 = (a / (b + b));
+ return;
+}
+
diff --git a/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..c4eaf69
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpFDiv %float %12 %15
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..6c35a6e
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-scalar/i32.wgsl b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/i32.wgsl
new file mode 100644
index 0000000..35b4427
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..6ed4430
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.msl
new file mode 100644
index 0000000..c319da3
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..e1a8969
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpSDiv %int %12 %15
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..6680c0b
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-scalar/u32.wgsl b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/u32.wgsl
new file mode 100644
index 0000000..1add65a
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..ec658be
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.msl
new file mode 100644
index 0000000..ad2302f
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..8d771cc
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpUDiv %uint %12 %15
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..ef947c3
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-vec3/f32.wgsl b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/f32.wgsl
new file mode 100644
index 0000000..bc06d3b
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/f32.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ var a = 4.;
+ var b = vec3<f32>(0., 2., 0.);
+ let r : vec3<f32> = a / (b + b);
+}
diff --git a/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/f32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..399b2b9
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/f32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ float a = 4.0f;
+ float3 b = float3(0.0f, 2.0f, 0.0f);
+ const float3 r = (a / (b + b));
+ return;
+}
diff --git a/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/f32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/f32.wgsl.expected.msl
new file mode 100644
index 0000000..c78571c
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/f32.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float a = 4.0f;
+ float3 b = float3(0.0f, 2.0f, 0.0f);
+ float3 const r = (a / (b + b));
+ return;
+}
+
diff --git a/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/f32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..ac14577
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/f32.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
+ %float = OpTypeFloat 32
+ %float_4 = OpConstant %float 4
+%_ptr_Function_float = OpTypePointer Function %float
+ %9 = OpConstantNull %float
+ %v3float = OpTypeVector %float 3
+ %float_0 = OpConstant %float 0
+ %float_2 = OpConstant %float 2
+ %13 = OpConstantComposite %v3float %float_0 %float_2 %float_0
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+ %16 = OpConstantNull %v3float
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_float Function %9
+ %b = OpVariable %_ptr_Function_v3float Function %16
+ %22 = OpVariable %_ptr_Function_v3float Function %16
+ OpStore %a %float_4
+ OpStore %b %13
+ %17 = OpLoad %float %a
+ %18 = OpLoad %v3float %b
+ %19 = OpLoad %v3float %b
+ %20 = OpFAdd %v3float %18 %19
+ %23 = OpCompositeConstruct %v3float %17 %17 %17
+ %21 = OpFDiv %v3float %23 %20
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/f32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..6c0a259
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/f32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ var a = 4.0;
+ var b = vec3<f32>(0.0, 2.0, 0.0);
+ let r : vec3<f32> = (a / (b + b));
+}
diff --git a/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl
new file mode 100644
index 0000000..7625e4c
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..f23bda3
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
new file mode 100644
index 0000000..72cb90c
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..cfc7c26
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpSDiv %v3int %23 %20
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..58c7456
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-vec3/u32.wgsl b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/u32.wgsl
new file mode 100644
index 0000000..86989de
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..ec3621c
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.msl
new file mode 100644
index 0000000..f742cd5
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..e0abc2f
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpUDiv %v3uint %23 %20
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..563beef
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-scalar/f32.wgsl b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/f32.wgsl
new file mode 100644
index 0000000..6e95ded
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/f32.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ var a = vec3<f32>(1., 2., 3.);
+ var b = 0.;
+ let r : vec3<f32> = a / (b + b);
+}
diff --git a/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/f32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..85f1063
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/f32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ float3 a = float3(1.0f, 2.0f, 3.0f);
+ float b = 0.0f;
+ const float3 r = (a / (b + b));
+ return;
+}
diff --git a/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/f32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/f32.wgsl.expected.msl
new file mode 100644
index 0000000..aadbf3b
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/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);
+ float b = 0.0f;
+ float3 const r = (a / (b + b));
+ return;
+}
+
diff --git a/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/f32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..6f77e08
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/f32.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
+ %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
+%_ptr_Function_float = OpTypePointer Function %float
+ %17 = OpConstantNull %float
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_v3float Function %13
+ %b = OpVariable %_ptr_Function_float Function %17
+ %23 = OpVariable %_ptr_Function_v3float Function %13
+ OpStore %a %10
+ OpStore %b %float_0
+ %18 = OpLoad %v3float %a
+ %19 = OpLoad %float %b
+ %20 = OpLoad %float %b
+ %21 = OpFAdd %float %19 %20
+ %24 = OpCompositeConstruct %v3float %21 %21 %21
+ %22 = OpFDiv %v3float %18 %24
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/f32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..27f709e
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/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 = 0.0;
+ let r : vec3<f32> = (a / (b + b));
+}
diff --git a/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/i32.wgsl b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/i32.wgsl
new file mode 100644
index 0000000..c04ada0
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..964d275
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.msl
new file mode 100644
index 0000000..b2a2e6a
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..755ab0e
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpSDiv %v3int %18 %24
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..6bd3b6d
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-scalar/u32.wgsl b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/u32.wgsl
new file mode 100644
index 0000000..eef2625
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..f4bd21b
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.msl
new file mode 100644
index 0000000..05e7e6e
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..d83d27b
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpUDiv %v3uint %18 %24
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..31867ef
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-vec3/f32.wgsl b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/f32.wgsl
new file mode 100644
index 0000000..351527d
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..4135478
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.msl
new file mode 100644
index 0000000..25e716a
--- /dev/null
+++ b/test/expressions/binary/div_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 = (a / (b + b));
+ return;
+}
+
diff --git a/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..eed89f7
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpFDiv %v3float %18 %21
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..f3eb992
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-vec3/i32.wgsl b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl
new file mode 100644
index 0000000..ab6ecd4
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..fefc614
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
new file mode 100644
index 0000000..e2a2a76
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..5552161
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpSDiv %v3int %18 %21
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..0e81db7
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-vec3/u32.wgsl b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/u32.wgsl
new file mode 100644
index 0000000..f55edd5
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..841bf62
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.msl
new file mode 100644
index 0000000..e055047
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..8f37304
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpUDiv %v3uint %18 %21
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..1d60a13
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-scalar/f32.wgsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/f32.wgsl
new file mode 100644
index 0000000..2dff3b7
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..47c924e
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.msl
new file mode 100644
index 0000000..80f3377
--- /dev/null
+++ b/test/expressions/binary/div_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 = (a / b);
+ return;
+}
+
diff --git a/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..b1ac052
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpFDiv %float %12 %13
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..332c2b7
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-scalar/i32.wgsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/i32.wgsl
new file mode 100644
index 0000000..375c773
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..111cfc9
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.msl
new file mode 100644
index 0000000..a3201dc
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..14dcaa0
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpSDiv %int %12 %13
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..90d99f4
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-scalar/u32.wgsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/u32.wgsl
new file mode 100644
index 0000000..b555e24
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..289c446
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.msl
new file mode 100644
index 0000000..bf7517d
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..0519b51
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpUDiv %uint %12 %13
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..478767e
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-vec3/f32.wgsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/f32.wgsl
new file mode 100644
index 0000000..16588a9
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/f32.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ var a = 4.;
+ var b = vec3<f32>(0., 2., 0.);
+ let r : vec3<f32> = a / b;
+}
diff --git a/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/f32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..6f9cc20
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/f32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ float a = 4.0f;
+ float3 b = float3(0.0f, 2.0f, 0.0f);
+ const float3 r = (a / b);
+ return;
+}
diff --git a/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/f32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/f32.wgsl.expected.msl
new file mode 100644
index 0000000..92c9809
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/f32.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+ float a = 4.0f;
+ float3 b = float3(0.0f, 2.0f, 0.0f);
+ float3 const r = (a / b);
+ return;
+}
+
diff --git a/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/f32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..a213c44
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/f32.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
+ %float = OpTypeFloat 32
+ %float_4 = OpConstant %float 4
+%_ptr_Function_float = OpTypePointer Function %float
+ %9 = OpConstantNull %float
+ %v3float = OpTypeVector %float 3
+ %float_0 = OpConstant %float 0
+ %float_2 = OpConstant %float 2
+ %13 = OpConstantComposite %v3float %float_0 %float_2 %float_0
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+ %16 = OpConstantNull %v3float
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_float Function %9
+ %b = OpVariable %_ptr_Function_v3float Function %16
+ %20 = OpVariable %_ptr_Function_v3float Function %16
+ OpStore %a %float_4
+ OpStore %b %13
+ %17 = OpLoad %float %a
+ %18 = OpLoad %v3float %b
+ %21 = OpCompositeConstruct %v3float %17 %17 %17
+ %19 = OpFDiv %v3float %21 %18
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/f32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..0d44b76
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/f32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ var a = 4.0;
+ var b = vec3<f32>(0.0, 2.0, 0.0);
+ let r : vec3<f32> = (a / b);
+}
diff --git a/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/i32.wgsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/i32.wgsl
new file mode 100644
index 0000000..8d25a71
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..24f0e8f
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.msl
new file mode 100644
index 0000000..03ac311
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..9164966
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpSDiv %v3int %21 %18
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..2bc4678
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-vec3/u32.wgsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/u32.wgsl
new file mode 100644
index 0000000..f417803
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..0692d21
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.msl
new file mode 100644
index 0000000..b7bf1d7
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..ce88d36
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpUDiv %v3uint %21 %18
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..147f4d2
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-scalar/f32.wgsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/f32.wgsl
new file mode 100644
index 0000000..2f23ba4
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/f32.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+ var a = vec3<f32>(1., 2., 3.);
+ var b = 0.;
+ let r : vec3<f32> = a / b;
+}
diff --git a/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/f32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..8d90669
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/f32.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+ float3 a = float3(1.0f, 2.0f, 3.0f);
+ float b = 0.0f;
+ const float3 r = (a / b);
+ return;
+}
diff --git a/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/f32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/f32.wgsl.expected.msl
new file mode 100644
index 0000000..86cae16
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/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);
+ float b = 0.0f;
+ float3 const r = (a / b);
+ return;
+}
+
diff --git a/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/f32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..292b631
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/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
+%_ptr_Function_float = OpTypePointer Function %float
+ %17 = OpConstantNull %float
+ %f = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_v3float Function %13
+ %b = OpVariable %_ptr_Function_float Function %17
+ %21 = OpVariable %_ptr_Function_v3float Function %13
+ OpStore %a %10
+ OpStore %b %float_0
+ %18 = OpLoad %v3float %a
+ %19 = OpLoad %float %b
+ %22 = OpCompositeConstruct %v3float %19 %19 %19
+ %20 = OpFDiv %v3float %18 %22
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/f32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..fed430a
--- /dev/null
+++ b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/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 = 0.0;
+ let r : vec3<f32> = (a / b);
+}
diff --git a/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/i32.wgsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/i32.wgsl
new file mode 100644
index 0000000..55aa4bc
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..4f87f70
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.msl
new file mode 100644
index 0000000..8e25431
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..35ceb4a
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpSDiv %v3int %18 %22
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..56d8e9d
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-scalar/u32.wgsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/u32.wgsl
new file mode 100644
index 0000000..4076d88
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..012cb29
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.msl
new file mode 100644
index 0000000..126ec2a
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..d04ce1e
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpUDiv %v3uint %18 %22
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..e55b987
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-vec3/f32.wgsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/f32.wgsl
new file mode 100644
index 0000000..f5092ad
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000..5ed2f85
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.msl
new file mode 100644
index 0000000..5894f5d
--- /dev/null
+++ b/test/expressions/binary/div_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 = (a / b);
+ return;
+}
+
diff --git a/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000..f5b1805
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpFDiv %v3float %18 %19
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000..534f264
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-vec3/i32.wgsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/i32.wgsl
new file mode 100644
index 0000000..07d8e62
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000..3e1ba7f
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.msl
new file mode 100644
index 0000000..298e7fa
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000..5647910
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpSDiv %v3int %18 %19
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000..5ce7a80
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-vec3/u32.wgsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/u32.wgsl
new file mode 100644
index 0000000..218e516
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.hlsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000..5f022e4
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.msl b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.msl
new file mode 100644
index 0000000..d300954
--- /dev/null
+++ b/test/expressions/binary/div_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/div_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.spvasm b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000..aba2a2b
--- /dev/null
+++ b/test/expressions/binary/div_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 = OpUDiv %v3uint %18 %19
+ OpReturn
+ OpFunctionEnd
diff --git a/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.wgsl b/test/expressions/binary/div_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000..35b85a1
--- /dev/null
+++ b/test/expressions/binary/div_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);
+}
