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