Fix FXC compile errors on divide by zero

FXC fails to compile when it determines that the rhs of an integral
division is zero with "error X4010: Unsigned integer divide by zero".

bclayton's fix (https://dawn-review.googlesource.com/c/tint/+/60500)
addressed cases for division by an integer constant 0. This CL adds the
missing support for division by integral vectors with 0 components.

FXC also fails on division by integral expressions that it can fold to
0. To handle these cases, we now emit a runtime check for 0 and replace
by 1. In the cases I've tested, FXC seems able to optimize these checks
away.

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

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() = default;
 
 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

src/sem/constant.h

@@ -97,6 +97,9 @@ class Constant {
   /// @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

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 @@ bool GeneratorImpl::EmitAssign(const ast::AssignmentStatement* stmt) {
   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 @@ bool GeneratorImpl::EmitBinary(std::ostream& out,
       break;
     case ast::BinaryOp::kDivide:
       out << "/";
-
+      // BUG(crbug.com/tint/1083): Integer divide by zero is a FXC compile
-      if (auto val = builder_.Sem().Get(expr->rhs)->ConstantValue()) {
+      // error, and undefined behavior in WGSL.
-        // Integer divide by zero is a DXC compile error, and undefined behavior
+      if (TypeOf(expr->rhs)->UnwrapRef()->is_integer_scalar_or_vector()) {
-        // in WGSL. Replace the 0 with 1.
+        out << " ";
-        if (val.Type()->Is<sem::I32>() && val.Elements()[0].i32 == 0) {
+        return EmitExpressionOrOneIfZero(out, expr->rhs);
-          out << " 1";
-          return true;
-        }
-        if (val.Type()->Is<sem::U32>() && val.Elements()[0].u32 == 0u) {
-          out << " 1u";
-          return true;
-        }
       }
       break;
     case ast::BinaryOp::kModulo:
@@ -3006,15 +3121,17 @@ bool GeneratorImpl::EmitLiteral(std::ostream& out,
   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 @@ bool GeneratorImpl::EmitZeroValue(std::ostream& out, const sem::Type* type) {
       if (i != 0) {
         out << ", ";
       }
-      if (!EmitZeroValue(out, vec->type())) {
+      if (!EmitValue(out, vec->type(), value)) {
         return false;
       }
     }
@@ -3039,7 +3156,7 @@ bool GeneratorImpl::EmitZeroValue(std::ostream& out, const sem::Type* type) {
       if (i != 0) {
         out << ", ";
       }
-      if (!EmitZeroValue(out, mat->type())) {
+      if (!EmitValue(out, mat->type(), value)) {
         return false;
       }
     }
@@ -3049,16 +3166,20 @@ bool GeneratorImpl::EmitZeroValue(std::ostream& out, const sem::Type* type) {
                   "")) {
       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()) {

src/writer/hlsl/generator_impl.h

@@ -100,6 +100,12 @@ class GeneratorImpl : public TextGenerator {
   /// @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 @@ class GeneratorImpl : public TextGenerator {
   /// @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 @@ class GeneratorImpl : public TextGenerator {
       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

src/writer/hlsl/generator_impl_binary_test.cc

@@ -26,6 +26,9 @@ using HlslGeneratorImplTest_Binary = TestHelper;
 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 @@ using HlslBinaryTest = TestParamHelper<BinaryData>;
 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_f32) {
 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_u32) {
 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 @@ INSTANTIATE_TEST_SUITE_P(
         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 @@ foo((tint_tmp), (tint_tmp_1), (tint_tmp_2));
 }
 
 TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_i32) {
-  Global("a", ty.i32(), ast::StorageClass::kPrivate);
+  Func("fn", {}, ty.void_(),
-
+       {
-  auto* expr = Div("a", 0);
+           Decl(Var("a", ty.i32())),
-  WrapInFunction(expr);
+           Decl(Const("r", nullptr, Div("a", 0))),
+       });
 
   GeneratorImpl& gen = Build();
 
-  std::stringstream out;
+  ASSERT_TRUE(gen.Generate());
-  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(void fn() {
-  EXPECT_EQ(out.str(), R"((a / 1))");
+  int a = 0;
+  const int r = (a / 1);
+}
+)");
 }
 
 TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_u32) {
-  Global("a", ty.u32(), ast::StorageClass::kPrivate);
+  Func("fn", {}, ty.void_(),
-
+       {
-  auto* expr = Div("a", 0u);
+           Decl(Var("a", ty.u32())),
-  WrapInFunction(expr);
+           Decl(Const("r", nullptr, Div("a", 0u))),
+       });
 
   GeneratorImpl& gen = Build();
 
-  std::stringstream out;
+  ASSERT_TRUE(gen.Generate());
-  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(void fn() {
-  EXPECT_EQ(out.str(), R"((a / 1u))");
+  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

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

@@ -30,6 +30,10 @@ bool4 tint_isNormal_1(float4 param_0) {
   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 @@ uint toIndex1D(uint gridSize, float3 voxelPos) {
 }
 
 uint3 toIndex4D(uint gridSize, uint index) {
-  uint z_1 = (gridSize / (index * index));
+  uint z_1 = (gridSize / value_or_one_if_zero_uint((index * index)));
-  uint y_1 = ((gridSize - ((gridSize * gridSize) * z_1)) / gridSize);
+  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);
 }

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 @@ uint toIndex1D(uint gridSize, float3 voxelPos) {
 }
 
 uint3 toIndex3D(uint gridSize, uint index) {
-  uint z_1 = (index / (gridSize * gridSize));
+  uint z_1 = (index / value_or_one_if_zero_uint((gridSize * gridSize)));
-  uint y_1 = ((index - ((gridSize * gridSize) * z_1)) / 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);
 }

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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