tint/transform: Fix ICE when combining polyfills
There's a reason the overload of `ctx.Replace()` that takes a pointer to the replacement is deprecated - it doesn't play well when used as part of another replacement.
Switch to using the callback overload of Replace() to fix bad transform output.
Bug: tint:1386647
Change-Id: I94292eeb65d24d7b2446b16b8b4ad13bdd27965a
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/111000
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/transform/builtin_polyfill.cc b/src/tint/transform/builtin_polyfill.cc
index 4dfc149..15aa233 100644
--- a/src/tint/transform/builtin_polyfill.cc
+++ b/src/tint/transform/builtin_polyfill.cc
@@ -41,6 +41,10 @@
/// @param p the builtins to polyfill
State(CloneContext& c, Builtins p) : ctx(c), polyfill(p) {}
+ ////////////////////////////////////////////////////////////////////////////
+ // Function polyfills
+ ////////////////////////////////////////////////////////////////////////////
+
/// Builds the polyfill function for the `acosh` builtin
/// @param ty the parameter and return type for the function
/// @return the polyfill function name
@@ -559,63 +563,6 @@
return name;
}
- /// Builds the polyfill function for a divide or modulo operator with integer scalar or vector
- /// operands.
- /// @param sig the signature of the binary operator
- /// @return the polyfill function name
- Symbol int_div_mod(const BinaryOpSignature& sig) {
- const auto op = std::get<0>(sig);
- const auto* lhs_ty = std::get<1>(sig);
- const auto* rhs_ty = std::get<2>(sig);
- const bool is_div = op == ast::BinaryOp::kDivide;
-
- uint32_t lhs_width = 1;
- uint32_t rhs_width = 1;
- const auto* lhs_el_ty = sem::Type::ElementOf(lhs_ty, &lhs_width);
- const auto* rhs_el_ty = sem::Type::ElementOf(rhs_ty, &rhs_width);
-
- const uint32_t width = std::max(lhs_width, rhs_width);
-
- const char* lhs = "lhs";
- const char* rhs = "rhs";
-
- utils::Vector<const ast::Statement*, 4> body;
-
- if (lhs_width < width) {
- // lhs is scalar, rhs is vector. Convert lhs to vector.
- body.Push(b.Decl(b.Let("l", b.vec(T(lhs_el_ty), width, b.Expr(lhs)))));
- lhs = "l";
- }
- if (rhs_width < width) {
- // lhs is vector, rhs is scalar. Convert rhs to vector.
- body.Push(b.Decl(b.Let("r", b.vec(T(rhs_el_ty), width, b.Expr(rhs)))));
- rhs = "r";
- }
-
- auto name = b.Symbols().New(is_div ? "tint_div" : "tint_mod");
- auto* use_one = b.Equal(rhs, ScalarOrVector(width, 0_a));
- if (lhs_ty->is_signed_scalar_or_vector()) {
- const auto bits = lhs_el_ty->Size() * 8;
- auto min_int = AInt(AInt::kLowestValue >> (AInt::kNumBits - bits));
- const ast::Expression* lhs_is_min = b.Equal(lhs, ScalarOrVector(width, min_int));
- const ast::Expression* rhs_is_minus_one = b.Equal(rhs, ScalarOrVector(width, -1_a));
- // use_one = use_one | ((lhs == MIN_INT) & (rhs == -1))
- use_one = b.Or(use_one, b.And(lhs_is_min, rhs_is_minus_one));
- }
- auto* select = b.Call("select", rhs, ScalarOrVector(width, 1_a), use_one);
-
- body.Push(b.Return(is_div ? b.Div(lhs, select) : b.Mod(lhs, select)));
- b.Func(name,
- utils::Vector{
- b.Param("lhs", T(lhs_ty)),
- b.Param("rhs", T(rhs_ty)),
- },
- width == 1 ? T(lhs_ty) : b.ty.vec(T(lhs_el_ty), width), // return type
- std::move(body));
-
- return name;
- }
-
/// Builds the polyfill function for the `saturate` builtin
/// @param ty the parameter and return type for the function
/// @return the polyfill function name
@@ -677,6 +624,89 @@
return name;
}
+ ////////////////////////////////////////////////////////////////////////////
+ // Inline polyfills
+ ////////////////////////////////////////////////////////////////////////////
+
+ /// Builds the polyfill inline expression for a bitshift left or bitshift right, ensuring that
+ /// the RHS is modulo the bit-width of the LHS.
+ /// @param bin_op the original BinaryExpression
+ /// @return the polyfill value for bitshift operation
+ const ast::Expression* BitshiftModulo(const ast::BinaryExpression* bin_op) {
+ auto* lhs_ty = ctx.src->TypeOf(bin_op->lhs)->UnwrapRef();
+ auto* rhs_ty = ctx.src->TypeOf(bin_op->rhs)->UnwrapRef();
+ auto* lhs_el_ty = sem::Type::DeepestElementOf(lhs_ty);
+ const ast::Expression* mask = b.Expr(AInt(lhs_el_ty->Size() * 8 - 1));
+ if (rhs_ty->Is<sem::Vector>()) {
+ mask = b.Construct(CreateASTTypeFor(ctx, rhs_ty), mask);
+ }
+ auto* lhs = ctx.Clone(bin_op->lhs);
+ auto* rhs = b.And(ctx.Clone(bin_op->rhs), mask);
+ return b.create<ast::BinaryExpression>(ctx.Clone(bin_op->source), bin_op->op, lhs, rhs);
+ }
+
+ /// Builds the polyfill inline expression for a integer divide or modulo, preventing DBZs and
+ /// integer overflows.
+ /// @param bin_op the original BinaryExpression
+ /// @return the polyfill divide or modulo
+ const ast::Expression* IntDivMod(const ast::BinaryExpression* bin_op) {
+ auto* lhs_ty = ctx.src->TypeOf(bin_op->lhs)->UnwrapRef();
+ auto* rhs_ty = ctx.src->TypeOf(bin_op->rhs)->UnwrapRef();
+ BinaryOpSignature sig{bin_op->op, lhs_ty, rhs_ty};
+ auto fn = binary_op_polyfills.GetOrCreate(sig, [&] {
+ const bool is_div = bin_op->op == ast::BinaryOp::kDivide;
+
+ uint32_t lhs_width = 1;
+ uint32_t rhs_width = 1;
+ const auto* lhs_el_ty = sem::Type::ElementOf(lhs_ty, &lhs_width);
+ const auto* rhs_el_ty = sem::Type::ElementOf(rhs_ty, &rhs_width);
+
+ const uint32_t width = std::max(lhs_width, rhs_width);
+
+ const char* lhs = "lhs";
+ const char* rhs = "rhs";
+
+ utils::Vector<const ast::Statement*, 4> body;
+
+ if (lhs_width < width) {
+ // lhs is scalar, rhs is vector. Convert lhs to vector.
+ body.Push(b.Decl(b.Let("l", b.vec(T(lhs_el_ty), width, b.Expr(lhs)))));
+ lhs = "l";
+ }
+ if (rhs_width < width) {
+ // lhs is vector, rhs is scalar. Convert rhs to vector.
+ body.Push(b.Decl(b.Let("r", b.vec(T(rhs_el_ty), width, b.Expr(rhs)))));
+ rhs = "r";
+ }
+
+ auto name = b.Symbols().New(is_div ? "tint_div" : "tint_mod");
+ auto* use_one = b.Equal(rhs, ScalarOrVector(width, 0_a));
+ if (lhs_ty->is_signed_scalar_or_vector()) {
+ const auto bits = lhs_el_ty->Size() * 8;
+ auto min_int = AInt(AInt::kLowestValue >> (AInt::kNumBits - bits));
+ const ast::Expression* lhs_is_min = b.Equal(lhs, ScalarOrVector(width, min_int));
+ const ast::Expression* rhs_is_minus_one = b.Equal(rhs, ScalarOrVector(width, -1_a));
+ // use_one = use_one | ((lhs == MIN_INT) & (rhs == -1))
+ use_one = b.Or(use_one, b.And(lhs_is_min, rhs_is_minus_one));
+ }
+ auto* select = b.Call("select", rhs, ScalarOrVector(width, 1_a), use_one);
+
+ body.Push(b.Return(is_div ? b.Div(lhs, select) : b.Mod(lhs, select)));
+ b.Func(name,
+ utils::Vector{
+ b.Param("lhs", T(lhs_ty)),
+ b.Param("rhs", T(rhs_ty)),
+ },
+ width == 1 ? T(lhs_ty) : b.ty.vec(T(lhs_el_ty), width), // return type
+ std::move(body));
+
+ return name;
+ });
+ auto* lhs = ctx.Clone(bin_op->lhs);
+ auto* rhs = ctx.Clone(bin_op->rhs);
+ return b.Call(fn, lhs, rhs);
+ }
+
private:
/// The clone context
CloneContext& ctx;
@@ -687,6 +717,9 @@
/// The source clone context
const sem::Info& sem = ctx.src->Sem();
+ // Polyfill functions for binary operators.
+ utils::Hashmap<BinaryOpSignature, Symbol, 8> binary_op_polyfills;
+
/// @returns the AST type for the given sem type
const ast::Type* T(const sem::Type* ty) const { return CreateASTTypeFor(ctx, ty); }
@@ -724,7 +757,6 @@
auto& polyfill = cfg->builtins;
utils::Hashmap<const sem::Builtin*, Symbol, 8> builtin_polyfills;
- utils::Hashmap<BinaryOpSignature, Symbol, 8> binary_op_polyfills;
ProgramBuilder b;
CloneContext ctx{&b, src, /* auto_clone_symbols */ true};
@@ -849,15 +881,7 @@
case ast::BinaryOp::kShiftLeft:
case ast::BinaryOp::kShiftRight: {
if (polyfill.bitshift_modulo) {
- auto* lhs_ty = src->TypeOf(bin_op->lhs)->UnwrapRef();
- auto* rhs_ty = src->TypeOf(bin_op->rhs)->UnwrapRef();
- auto* lhs_el_ty = sem::Type::DeepestElementOf(lhs_ty);
- const ast::Expression* mask = b.Expr(AInt(lhs_el_ty->Size() * 8 - 1));
- if (rhs_ty->Is<sem::Vector>()) {
- mask = b.Construct(CreateASTTypeFor(ctx, rhs_ty), mask);
- }
- auto* mod = b.And(ctx.Clone(bin_op->rhs), mask);
- ctx.Replace(bin_op->rhs, mod);
+ ctx.Replace(bin_op, [bin_op, &s] { return s.BitshiftModulo(bin_op); });
made_changes = true;
}
break;
@@ -867,13 +891,7 @@
if (polyfill.int_div_mod) {
auto* lhs_ty = src->TypeOf(bin_op->lhs)->UnwrapRef();
if (lhs_ty->is_integer_scalar_or_vector()) {
- auto* rhs_ty = src->TypeOf(bin_op->rhs)->UnwrapRef();
- BinaryOpSignature sig{bin_op->op, lhs_ty, rhs_ty};
- auto fn = binary_op_polyfills.GetOrCreate(
- sig, [&] { return s.int_div_mod(sig); });
- auto* lhs = ctx.Clone(bin_op->lhs);
- auto* rhs = ctx.Clone(bin_op->rhs);
- ctx.Replace(bin_op, b.Call(fn, lhs, rhs));
+ ctx.Replace(bin_op, [bin_op, &s] { return s.IntDivMod(bin_op); });
made_changes = true;
}
}
diff --git a/src/tint/transform/builtin_polyfill_test.cc b/src/tint/transform/builtin_polyfill_test.cc
index fbbd90e..87fd9d5 100644
--- a/src/tint/transform/builtin_polyfill_test.cc
+++ b/src/tint/transform/builtin_polyfill_test.cc
@@ -3000,5 +3000,37 @@
EXPECT_EQ(expect, str(got));
}
+////////////////////////////////////////////////////////////////////////////////
+// Polyfill combinations
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_F(BuiltinPolyfillTest, BitshiftAndModulo) {
+ auto* src = R"(
+fn f(x : i32, y : u32, z : u32) {
+ let l = x << (y % z);
+}
+)";
+
+ auto* expect = R"(
+fn tint_mod(lhs : u32, rhs : u32) -> u32 {
+ return (lhs % select(rhs, 1, (rhs == 0)));
+}
+
+fn f(x : i32, y : u32, z : u32) {
+ let l = (x << (tint_mod(y, z) & 31));
+}
+)";
+
+ BuiltinPolyfill::Builtins builtins;
+ builtins.bitshift_modulo = true;
+ builtins.int_div_mod = true;
+ DataMap data;
+ data.Add<BuiltinPolyfill::Config>(builtins);
+
+ auto got = Run<BuiltinPolyfill>(src, std::move(data));
+
+ EXPECT_EQ(expect, str(got));
+}
+
} // namespace
} // namespace tint::transform
diff --git a/test/tint/bug/chromium/1273230.wgsl.expected.dxc.hlsl b/test/tint/bug/chromium/1273230.wgsl.expected.dxc.hlsl
index 7f89b55..8b61a3e 100644
--- a/test/tint/bug/chromium/1273230.wgsl.expected.dxc.hlsl
+++ b/test/tint/bug/chromium/1273230.wgsl.expected.dxc.hlsl
@@ -1,11 +1,3 @@
-uint tint_div(uint lhs, uint rhs) {
- return (lhs / ((rhs == 0u) ? 1u : rhs));
-}
-
-uint tint_mod(uint lhs, uint rhs) {
- return (lhs % ((rhs == 0u) ? 1u : rhs));
-}
-
void marg8uintin() {
}
@@ -35,6 +27,14 @@
return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z));
}
+uint tint_div(uint lhs, uint rhs) {
+ return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
+uint tint_mod(uint lhs, uint rhs) {
+ return (lhs % ((rhs == 0u) ? 1u : rhs));
+}
+
uint3 toIndex4D(uint gridSize, uint index) {
uint z_1 = tint_div(gridSize, (index * index));
uint y_1 = tint_div((gridSize - ((gridSize * gridSize) * z_1)), gridSize);
diff --git a/test/tint/bug/chromium/1273230.wgsl.expected.fxc.hlsl b/test/tint/bug/chromium/1273230.wgsl.expected.fxc.hlsl
index 7f89b55..8b61a3e 100644
--- a/test/tint/bug/chromium/1273230.wgsl.expected.fxc.hlsl
+++ b/test/tint/bug/chromium/1273230.wgsl.expected.fxc.hlsl
@@ -1,11 +1,3 @@
-uint tint_div(uint lhs, uint rhs) {
- return (lhs / ((rhs == 0u) ? 1u : rhs));
-}
-
-uint tint_mod(uint lhs, uint rhs) {
- return (lhs % ((rhs == 0u) ? 1u : rhs));
-}
-
void marg8uintin() {
}
@@ -35,6 +27,14 @@
return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z));
}
+uint tint_div(uint lhs, uint rhs) {
+ return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
+uint tint_mod(uint lhs, uint rhs) {
+ return (lhs % ((rhs == 0u) ? 1u : rhs));
+}
+
uint3 toIndex4D(uint gridSize, uint index) {
uint z_1 = tint_div(gridSize, (index * index));
uint y_1 = tint_div((gridSize - ((gridSize * gridSize) * z_1)), gridSize);
diff --git a/test/tint/bug/chromium/1273230.wgsl.expected.msl b/test/tint/bug/chromium/1273230.wgsl.expected.msl
index 15000e6..aa34606 100644
--- a/test/tint/bug/chromium/1273230.wgsl.expected.msl
+++ b/test/tint/bug/chromium/1273230.wgsl.expected.msl
@@ -14,14 +14,6 @@
T elements[N];
};
-uint tint_div(uint lhs, uint rhs) {
- return (lhs / select(rhs, 1u, (rhs == 0u)));
-}
-
-uint tint_mod(uint lhs, uint rhs) {
- return (lhs % select(rhs, 1u, (rhs == 0u)));
-}
-
void marg8uintin() {
}
@@ -88,6 +80,14 @@
return ((icoord[0] + (gridSize * icoord[1])) + ((gridSize * gridSize) * icoord[2]));
}
+uint tint_div(uint lhs, uint rhs) {
+ return (lhs / select(rhs, 1u, (rhs == 0u)));
+}
+
+uint tint_mod(uint lhs, uint rhs) {
+ return (lhs % select(rhs, 1u, (rhs == 0u)));
+}
+
uint3 toIndex4D(uint gridSize, uint index) {
uint z_1 = tint_div(gridSize, (index * index));
uint y_1 = tint_div((gridSize - ((gridSize * gridSize) * z_1)), gridSize);
diff --git a/test/tint/bug/chromium/1273230.wgsl.expected.spvasm b/test/tint/bug/chromium/1273230.wgsl.expected.spvasm
index 2e66503..96d229c 100644
--- a/test/tint/bug/chromium/1273230.wgsl.expected.spvasm
+++ b/test/tint/bug/chromium/1273230.wgsl.expected.spvasm
@@ -4,7 +4,7 @@
; Bound: 290
; Schema: 0
OpCapability Shader
- %86 = OpExtInstImport "GLSL.std.450"
+ %69 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main_count "main_count" %GlobalInvocationID_1
OpExecutionMode %main_count LocalSize 128 1 1
@@ -47,12 +47,6 @@
OpMemberName %Dbg 10 "value_f32_2"
OpMemberName %Dbg 11 "value_f32_3"
OpName %dbg "dbg"
- OpName %tint_div "tint_div"
- OpName %lhs "lhs"
- OpName %rhs "rhs"
- OpName %tint_mod "tint_mod"
- OpName %lhs_0 "lhs"
- OpName %rhs_0 "rhs"
OpName %marg8uintin "marg8uintin"
OpName %toVoxelPos "toVoxelPos"
OpName %position "position"
@@ -68,6 +62,12 @@
OpName %gridSize_0 "gridSize"
OpName %voxelPos "voxelPos"
OpName %icoord "icoord"
+ OpName %tint_div "tint_div"
+ OpName %lhs "lhs"
+ OpName %rhs "rhs"
+ OpName %tint_mod "tint_mod"
+ OpName %lhs_0 "lhs"
+ OpName %rhs_0 "rhs"
OpName %toIndex4D "toIndex4D"
OpName %gridSize_1 "gridSize"
OpName %index "index"
@@ -177,27 +177,27 @@
%dbg_block = OpTypeStruct %Dbg
%_ptr_StorageBuffer_dbg_block = OpTypePointer StorageBuffer %dbg_block
%dbg = OpVariable %_ptr_StorageBuffer_dbg_block StorageBuffer
- %32 = OpTypeFunction %uint %uint %uint
- %38 = OpConstantNull %uint
- %bool = OpTypeBool
- %uint_1 = OpConstant %uint 1
%void = OpTypeVoid
- %50 = OpTypeFunction %void
- %54 = OpTypeFunction %v3float %v3float
+ %32 = OpTypeFunction %void
+ %36 = OpTypeFunction %v3float %v3float
%uint_0 = OpConstant %uint 0
%uint_4 = OpConstant %uint 4
%_ptr_Uniform_float = OpTypePointer Uniform %float
+ %uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%_ptr_Function_v3float = OpTypePointer Function %v3float
- %71 = OpConstantNull %v3float
+ %54 = OpConstantNull %v3float
%uint_5 = OpConstant %uint 5
%_ptr_Function_float = OpTypePointer Function %float
- %96 = OpConstantNull %float
+ %79 = OpConstantNull %float
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
- %133 = OpTypeFunction %uint %uint %v3float
+ %116 = OpTypeFunction %uint %uint %v3float
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
- %141 = OpConstantNull %v3uint
+ %124 = OpConstantNull %v3uint
%_ptr_Function_uint = OpTypePointer Function %uint
+ %137 = OpTypeFunction %uint %uint %uint
+ %143 = OpConstantNull %uint
+ %bool = OpTypeBool
%154 = OpTypeFunction %v3uint %uint %uint
%174 = OpTypeFunction %v3float %uint
%uint_3 = OpConstant %uint 3
@@ -210,132 +210,132 @@
%222 = OpTypeFunction %void %v3uint
%float_3 = OpConstant %float 3
%int_1 = OpConstant %int 1
- %tint_div = OpFunction %uint None %32
- %lhs = OpFunctionParameter %uint
- %rhs = OpFunctionParameter %uint
- %36 = OpLabel
- %39 = OpIEqual %bool %rhs %38
- %37 = OpSelect %uint %39 %uint_1 %rhs
- %42 = OpUDiv %uint %lhs %37
- OpReturnValue %42
- OpFunctionEnd
- %tint_mod = OpFunction %uint None %32
- %lhs_0 = OpFunctionParameter %uint
- %rhs_0 = OpFunctionParameter %uint
- %46 = OpLabel
- %48 = OpIEqual %bool %rhs_0 %38
- %47 = OpSelect %uint %48 %uint_1 %rhs_0
- %49 = OpUMod %uint %lhs_0 %47
- OpReturnValue %49
- OpFunctionEnd
-%marg8uintin = OpFunction %void None %50
- %53 = OpLabel
+%marg8uintin = OpFunction %void None %32
+ %35 = OpLabel
OpReturn
OpFunctionEnd
- %toVoxelPos = OpFunction %v3float None %54
+ %toVoxelPos = OpFunction %v3float None %36
%position = OpFunctionParameter %v3float
- %57 = OpLabel
- %bbMin = OpVariable %_ptr_Function_v3float Function %71
- %bbMax = OpVariable %_ptr_Function_v3float Function %71
- %bbSize = OpVariable %_ptr_Function_v3float Function %71
- %cubeSize = OpVariable %_ptr_Function_float Function %96
- %gridSize = OpVariable %_ptr_Function_float Function %96
- %gx = OpVariable %_ptr_Function_float Function %96
- %gy = OpVariable %_ptr_Function_float Function %96
- %gz = OpVariable %_ptr_Function_float Function %96
- %61 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_0
- %62 = OpLoad %float %61
- %63 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_1
- %64 = OpLoad %float %63
- %66 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_2
- %67 = OpLoad %float %66
- %68 = OpCompositeConstruct %v3float %62 %64 %67
- OpStore %bbMin %68
- %73 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_5 %uint_0
- %74 = OpLoad %float %73
- %75 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_5 %uint_1
- %76 = OpLoad %float %75
- %77 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_5 %uint_2
- %78 = OpLoad %float %77
- %79 = OpCompositeConstruct %v3float %74 %76 %78
- OpStore %bbMax %79
- %81 = OpLoad %v3float %bbMin
- %82 = OpLoad %v3float %bbMin
- %83 = OpFSub %v3float %81 %82
- OpStore %bbSize %83
- %89 = OpAccessChain %_ptr_Function_float %bbMax %uint_0
- %90 = OpLoad %float %89
- %91 = OpAccessChain %_ptr_Function_float %bbMax %uint_1
- %92 = OpLoad %float %91
- %87 = OpExtInst %float %86 NMax %90 %92
- %93 = OpAccessChain %_ptr_Function_float %bbSize %uint_2
- %94 = OpLoad %float %93
- %85 = OpExtInst %float %86 NMax %87 %94
- OpStore %cubeSize %85
- %99 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_1
- %100 = OpLoad %uint %99
- %97 = OpConvertUToF %float %100
- OpStore %gridSize %97
- %102 = OpLoad %float %cubeSize
- %103 = OpCompositeExtract %float %position 0
- %104 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_0
- %105 = OpLoad %float %104
- %106 = OpFSub %float %103 %105
- %107 = OpFMul %float %102 %106
- %108 = OpLoad %float %cubeSize
- %109 = OpFDiv %float %107 %108
- OpStore %gx %109
- %111 = OpLoad %float %gx
- %112 = OpCompositeExtract %float %position 1
- %113 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_1
- %114 = OpLoad %float %113
- %115 = OpFSub %float %112 %114
- %116 = OpFMul %float %111 %115
- %117 = OpLoad %float %gridSize
- %118 = OpFDiv %float %116 %117
- OpStore %gy %118
- %120 = OpLoad %float %gridSize
- %121 = OpCompositeExtract %float %position 2
- %122 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_2
- %123 = OpLoad %float %122
- %124 = OpFSub %float %121 %123
- %125 = OpFMul %float %120 %124
- %126 = OpLoad %float %gridSize
- %127 = OpFDiv %float %125 %126
- OpStore %gz %127
- %129 = OpLoad %float %gz
- %130 = OpLoad %float %gz
- %131 = OpLoad %float %gz
- %132 = OpCompositeConstruct %v3float %129 %130 %131
- OpReturnValue %132
+ %39 = OpLabel
+ %bbMin = OpVariable %_ptr_Function_v3float Function %54
+ %bbMax = OpVariable %_ptr_Function_v3float Function %54
+ %bbSize = OpVariable %_ptr_Function_v3float Function %54
+ %cubeSize = OpVariable %_ptr_Function_float Function %79
+ %gridSize = OpVariable %_ptr_Function_float Function %79
+ %gx = OpVariable %_ptr_Function_float Function %79
+ %gy = OpVariable %_ptr_Function_float Function %79
+ %gz = OpVariable %_ptr_Function_float Function %79
+ %43 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_0
+ %44 = OpLoad %float %43
+ %46 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_1
+ %47 = OpLoad %float %46
+ %49 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_2
+ %50 = OpLoad %float %49
+ %51 = OpCompositeConstruct %v3float %44 %47 %50
+ OpStore %bbMin %51
+ %56 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_5 %uint_0
+ %57 = OpLoad %float %56
+ %58 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_5 %uint_1
+ %59 = OpLoad %float %58
+ %60 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_5 %uint_2
+ %61 = OpLoad %float %60
+ %62 = OpCompositeConstruct %v3float %57 %59 %61
+ OpStore %bbMax %62
+ %64 = OpLoad %v3float %bbMin
+ %65 = OpLoad %v3float %bbMin
+ %66 = OpFSub %v3float %64 %65
+ OpStore %bbSize %66
+ %72 = OpAccessChain %_ptr_Function_float %bbMax %uint_0
+ %73 = OpLoad %float %72
+ %74 = OpAccessChain %_ptr_Function_float %bbMax %uint_1
+ %75 = OpLoad %float %74
+ %70 = OpExtInst %float %69 NMax %73 %75
+ %76 = OpAccessChain %_ptr_Function_float %bbSize %uint_2
+ %77 = OpLoad %float %76
+ %68 = OpExtInst %float %69 NMax %70 %77
+ OpStore %cubeSize %68
+ %82 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_1
+ %83 = OpLoad %uint %82
+ %80 = OpConvertUToF %float %83
+ OpStore %gridSize %80
+ %85 = OpLoad %float %cubeSize
+ %86 = OpCompositeExtract %float %position 0
+ %87 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_0
+ %88 = OpLoad %float %87
+ %89 = OpFSub %float %86 %88
+ %90 = OpFMul %float %85 %89
+ %91 = OpLoad %float %cubeSize
+ %92 = OpFDiv %float %90 %91
+ OpStore %gx %92
+ %94 = OpLoad %float %gx
+ %95 = OpCompositeExtract %float %position 1
+ %96 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_1
+ %97 = OpLoad %float %96
+ %98 = OpFSub %float %95 %97
+ %99 = OpFMul %float %94 %98
+ %100 = OpLoad %float %gridSize
+ %101 = OpFDiv %float %99 %100
+ OpStore %gy %101
+ %103 = OpLoad %float %gridSize
+ %104 = OpCompositeExtract %float %position 2
+ %105 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_2
+ %106 = OpLoad %float %105
+ %107 = OpFSub %float %104 %106
+ %108 = OpFMul %float %103 %107
+ %109 = OpLoad %float %gridSize
+ %110 = OpFDiv %float %108 %109
+ OpStore %gz %110
+ %112 = OpLoad %float %gz
+ %113 = OpLoad %float %gz
+ %114 = OpLoad %float %gz
+ %115 = OpCompositeConstruct %v3float %112 %113 %114
+ OpReturnValue %115
OpFunctionEnd
- %toIndex1D = OpFunction %uint None %133
+ %toIndex1D = OpFunction %uint None %116
%gridSize_0 = OpFunctionParameter %uint
%voxelPos = OpFunctionParameter %v3float
- %137 = OpLabel
- %icoord = OpVariable %_ptr_Function_v3uint Function %141
- %138 = OpConvertFToU %v3uint %voxelPos
- OpStore %icoord %138
- %143 = OpAccessChain %_ptr_Function_uint %icoord %uint_0
- %144 = OpLoad %uint %143
- %145 = OpAccessChain %_ptr_Function_uint %icoord %uint_1
- %146 = OpLoad %uint %145
- %147 = OpIMul %uint %gridSize_0 %146
- %148 = OpIAdd %uint %144 %147
- %149 = OpIMul %uint %gridSize_0 %gridSize_0
- %150 = OpAccessChain %_ptr_Function_uint %icoord %uint_2
- %151 = OpLoad %uint %150
- %152 = OpIMul %uint %149 %151
- %153 = OpIAdd %uint %148 %152
+ %120 = OpLabel
+ %icoord = OpVariable %_ptr_Function_v3uint Function %124
+ %121 = OpConvertFToU %v3uint %voxelPos
+ OpStore %icoord %121
+ %126 = OpAccessChain %_ptr_Function_uint %icoord %uint_0
+ %127 = OpLoad %uint %126
+ %128 = OpAccessChain %_ptr_Function_uint %icoord %uint_1
+ %129 = OpLoad %uint %128
+ %130 = OpIMul %uint %gridSize_0 %129
+ %131 = OpIAdd %uint %127 %130
+ %132 = OpIMul %uint %gridSize_0 %gridSize_0
+ %133 = OpAccessChain %_ptr_Function_uint %icoord %uint_2
+ %134 = OpLoad %uint %133
+ %135 = OpIMul %uint %132 %134
+ %136 = OpIAdd %uint %131 %135
+ OpReturnValue %136
+ OpFunctionEnd
+ %tint_div = OpFunction %uint None %137
+ %lhs = OpFunctionParameter %uint
+ %rhs = OpFunctionParameter %uint
+ %141 = OpLabel
+ %144 = OpIEqual %bool %rhs %143
+ %142 = OpSelect %uint %144 %uint_1 %rhs
+ %146 = OpUDiv %uint %lhs %142
+ OpReturnValue %146
+ OpFunctionEnd
+ %tint_mod = OpFunction %uint None %137
+ %lhs_0 = OpFunctionParameter %uint
+ %rhs_0 = OpFunctionParameter %uint
+ %150 = OpLabel
+ %152 = OpIEqual %bool %rhs_0 %143
+ %151 = OpSelect %uint %152 %uint_1 %rhs_0
+ %153 = OpUMod %uint %lhs_0 %151
OpReturnValue %153
OpFunctionEnd
%toIndex4D = OpFunction %v3uint None %154
%gridSize_1 = OpFunctionParameter %uint
%index = OpFunctionParameter %uint
%158 = OpLabel
- %z = OpVariable %_ptr_Function_uint Function %38
- %y = OpVariable %_ptr_Function_uint Function %38
- %x = OpVariable %_ptr_Function_uint Function %38
+ %z = OpVariable %_ptr_Function_uint Function %143
+ %y = OpVariable %_ptr_Function_uint Function %143
+ %x = OpVariable %_ptr_Function_uint Function %143
%160 = OpIMul %uint %index %index
%159 = OpFunctionCall %uint %tint_div %gridSize_1 %160
OpStore %z %159
@@ -356,9 +356,9 @@
%loadPosition = OpFunction %v3float None %174
%vertexIndex = OpFunctionParameter %uint
%177 = OpLabel
- %position_0 = OpVariable %_ptr_Function_v3float Function %71
+ %position_0 = OpVariable %_ptr_Function_v3float Function %54
%179 = OpIMul %uint %uint_3 %vertexIndex
- %180 = OpIAdd %uint %179 %38
+ %180 = OpIAdd %uint %179 %143
%182 = OpAccessChain %_ptr_StorageBuffer_float %positions %uint_0 %180
%183 = OpLoad %float %182
%184 = OpIMul %uint %uint_3 %vertexIndex
@@ -374,13 +374,13 @@
%194 = OpLoad %v3float %position_0
OpReturnValue %194
OpFunctionEnd
- %doIgnore = OpFunction %void None %50
+ %doIgnore = OpFunction %void None %32
%196 = OpLabel
- %g43 = OpVariable %_ptr_Function_uint Function %38
- %kj6 = OpVariable %_ptr_Function_uint Function %38
- %b53 = OpVariable %_ptr_Function_uint Function %38
- %rwg = OpVariable %_ptr_Function_uint Function %38
- %rb5 = OpVariable %_ptr_Function_float Function %96
+ %g43 = OpVariable %_ptr_Function_uint Function %143
+ %kj6 = OpVariable %_ptr_Function_uint Function %143
+ %b53 = OpVariable %_ptr_Function_uint Function %143
+ %rwg = OpVariable %_ptr_Function_uint Function %143
+ %rb5 = OpVariable %_ptr_Function_float Function %79
%g55 = OpVariable %_ptr_Function_int Function %206
%197 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_0
%198 = OpLoad %uint %197
@@ -405,17 +405,17 @@
%main_count_inner = OpFunction %void None %222
%GlobalInvocationID = OpFunctionParameter %v3uint
%225 = OpLabel
-%triangleIndex = OpVariable %_ptr_Function_uint Function %38
- %i0 = OpVariable %_ptr_Function_uint Function %38
- %i1 = OpVariable %_ptr_Function_uint Function %38
- %i2 = OpVariable %_ptr_Function_uint Function %38
- %p0 = OpVariable %_ptr_Function_v3float Function %71
- %p1 = OpVariable %_ptr_Function_v3float Function %71
- %p2 = OpVariable %_ptr_Function_v3float Function %71
- %269 = OpVariable %_ptr_Function_v3float Function %71
- %center = OpVariable %_ptr_Function_v3float Function %71
- %voxelPos_0 = OpVariable %_ptr_Function_v3float Function %71
- %lIndex = OpVariable %_ptr_Function_uint Function %38
+%triangleIndex = OpVariable %_ptr_Function_uint Function %143
+ %i0 = OpVariable %_ptr_Function_uint Function %143
+ %i1 = OpVariable %_ptr_Function_uint Function %143
+ %i2 = OpVariable %_ptr_Function_uint Function %143
+ %p0 = OpVariable %_ptr_Function_v3float Function %54
+ %p1 = OpVariable %_ptr_Function_v3float Function %54
+ %p2 = OpVariable %_ptr_Function_v3float Function %54
+ %269 = OpVariable %_ptr_Function_v3float Function %54
+ %center = OpVariable %_ptr_Function_v3float Function %54
+ %voxelPos_0 = OpVariable %_ptr_Function_v3float Function %54
+ %lIndex = OpVariable %_ptr_Function_uint Function %143
%triangleOffset = OpVariable %_ptr_Function_int Function %206
%226 = OpCompositeExtract %uint %GlobalInvocationID 0
OpStore %triangleIndex %226
@@ -431,7 +431,7 @@
%234 = OpFunctionCall %void %doIgnore
%235 = OpLoad %uint %triangleIndex
%236 = OpIMul %uint %uint_3 %235
- %237 = OpIAdd %uint %236 %38
+ %237 = OpIAdd %uint %236 %143
%238 = OpAccessChain %_ptr_StorageBuffer_uint %indices %uint_0 %237
%239 = OpLoad %uint %238
OpStore %i0 %239
@@ -478,7 +478,7 @@
OpStore %triangleOffset %280
OpReturn
OpFunctionEnd
- %main_count = OpFunction %void None %50
+ %main_count = OpFunction %void None %32
%287 = OpLabel
%289 = OpLoad %v3uint %GlobalInvocationID_1
%288 = OpFunctionCall %void %main_count_inner %289
diff --git a/test/tint/bug/chromium/1386647.wgsl b/test/tint/bug/chromium/1386647.wgsl
new file mode 100644
index 0000000..b0745fe
--- /dev/null
+++ b/test/tint/bug/chromium/1386647.wgsl
@@ -0,0 +1,4 @@
+@compute @workgroup_size(1)
+fn f(@builtin(global_invocation_id) v : vec3<u32>) {
+ let l = v.x << (v.y % 1);
+}
diff --git a/test/tint/bug/chromium/1386647.wgsl.expected.dxc.hlsl b/test/tint/bug/chromium/1386647.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..424a72f
--- /dev/null
+++ b/test/tint/bug/chromium/1386647.wgsl.expected.dxc.hlsl
@@ -0,0 +1,19 @@
+uint tint_mod(uint lhs, uint rhs) {
+ return (lhs % ((rhs == 0u) ? 1u : rhs));
+}
+
+struct tint_symbol_1 {
+ uint3 v : SV_DispatchThreadID;
+};
+
+void f_inner(uint3 v) {
+ const uint tint_symbol_2 = v.x;
+ const uint tint_symbol_3 = tint_mod(v.y, 1u);
+ const uint l = (tint_symbol_2 << (tint_symbol_3 & 31u));
+}
+
+[numthreads(1, 1, 1)]
+void f(tint_symbol_1 tint_symbol) {
+ f_inner(tint_symbol.v);
+ return;
+}
diff --git a/test/tint/bug/chromium/1386647.wgsl.expected.fxc.hlsl b/test/tint/bug/chromium/1386647.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..424a72f
--- /dev/null
+++ b/test/tint/bug/chromium/1386647.wgsl.expected.fxc.hlsl
@@ -0,0 +1,19 @@
+uint tint_mod(uint lhs, uint rhs) {
+ return (lhs % ((rhs == 0u) ? 1u : rhs));
+}
+
+struct tint_symbol_1 {
+ uint3 v : SV_DispatchThreadID;
+};
+
+void f_inner(uint3 v) {
+ const uint tint_symbol_2 = v.x;
+ const uint tint_symbol_3 = tint_mod(v.y, 1u);
+ const uint l = (tint_symbol_2 << (tint_symbol_3 & 31u));
+}
+
+[numthreads(1, 1, 1)]
+void f(tint_symbol_1 tint_symbol) {
+ f_inner(tint_symbol.v);
+ return;
+}
diff --git a/test/tint/bug/chromium/1386647.wgsl.expected.glsl b/test/tint/bug/chromium/1386647.wgsl.expected.glsl
new file mode 100644
index 0000000..ab0cb6e
--- /dev/null
+++ b/test/tint/bug/chromium/1386647.wgsl.expected.glsl
@@ -0,0 +1,17 @@
+#version 310 es
+
+uint tint_mod(uint lhs, uint rhs) {
+ return (lhs % ((rhs == 0u) ? 1u : rhs));
+}
+
+void f(uvec3 v) {
+ uint tint_symbol = v.x;
+ uint tint_symbol_1 = tint_mod(v.y, 1u);
+ uint l = (tint_symbol << (tint_symbol_1 & 31u));
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f(gl_GlobalInvocationID);
+ return;
+}
diff --git a/test/tint/bug/chromium/1386647.wgsl.expected.msl b/test/tint/bug/chromium/1386647.wgsl.expected.msl
new file mode 100644
index 0000000..694a88a
--- /dev/null
+++ b/test/tint/bug/chromium/1386647.wgsl.expected.msl
@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+
+using namespace metal;
+uint tint_mod(uint lhs, uint rhs) {
+ return (lhs % select(rhs, 1u, (rhs == 0u)));
+}
+
+void f_inner(uint3 v) {
+ uint const tint_symbol = v[0];
+ uint const tint_symbol_1 = tint_mod(v[1], 1u);
+ uint const l = (tint_symbol << (tint_symbol_1 & 31u));
+}
+
+kernel void f(uint3 v [[thread_position_in_grid]]) {
+ f_inner(v);
+ return;
+}
+
diff --git a/test/tint/bug/chromium/1386647.wgsl.expected.spvasm b/test/tint/bug/chromium/1386647.wgsl.expected.spvasm
new file mode 100644
index 0000000..c54e606
--- /dev/null
+++ b/test/tint/bug/chromium/1386647.wgsl.expected.spvasm
@@ -0,0 +1,54 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 32
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %f "f" %v_1
+ OpExecutionMode %f LocalSize 1 1 1
+ OpName %v_1 "v_1"
+ OpName %tint_mod "tint_mod"
+ OpName %lhs "lhs"
+ OpName %rhs "rhs"
+ OpName %f_inner "f_inner"
+ OpName %v "v"
+ OpName %f "f"
+ OpDecorate %v_1 BuiltIn GlobalInvocationId
+ %uint = OpTypeInt 32 0
+ %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+ %v_1 = OpVariable %_ptr_Input_v3uint Input
+ %5 = OpTypeFunction %uint %uint %uint
+ %11 = OpConstantNull %uint
+ %bool = OpTypeBool
+ %uint_1 = OpConstant %uint 1
+ %void = OpTypeVoid
+ %16 = OpTypeFunction %void %v3uint
+ %uint_31 = OpConstant %uint 31
+ %27 = OpTypeFunction %void
+ %tint_mod = OpFunction %uint None %5
+ %lhs = OpFunctionParameter %uint
+ %rhs = OpFunctionParameter %uint
+ %9 = OpLabel
+ %12 = OpIEqual %bool %rhs %11
+ %10 = OpSelect %uint %12 %uint_1 %rhs
+ %15 = OpUMod %uint %lhs %10
+ OpReturnValue %15
+ OpFunctionEnd
+ %f_inner = OpFunction %void None %16
+ %v = OpFunctionParameter %v3uint
+ %20 = OpLabel
+ %21 = OpCompositeExtract %uint %v 0
+ %23 = OpCompositeExtract %uint %v 1
+ %22 = OpFunctionCall %uint %tint_mod %23 %uint_1
+ %25 = OpBitwiseAnd %uint %22 %uint_31
+ %26 = OpShiftLeftLogical %uint %21 %25
+ OpReturn
+ OpFunctionEnd
+ %f = OpFunction %void None %27
+ %29 = OpLabel
+ %31 = OpLoad %v3uint %v_1
+ %30 = OpFunctionCall %void %f_inner %31
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/chromium/1386647.wgsl.expected.wgsl b/test/tint/bug/chromium/1386647.wgsl.expected.wgsl
new file mode 100644
index 0000000..b8c5bd3
--- /dev/null
+++ b/test/tint/bug/chromium/1386647.wgsl.expected.wgsl
@@ -0,0 +1,4 @@
+@compute @workgroup_size(1)
+fn f(@builtin(global_invocation_id) v : vec3<u32>) {
+ let l = (v.x << (v.y % 1));
+}
diff --git a/test/tint/bug/tint/1113.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1113.wgsl.expected.dxc.hlsl
index 6a0326e..68248ca 100644
--- a/test/tint/bug/tint/1113.wgsl.expected.dxc.hlsl
+++ b/test/tint/bug/tint/1113.wgsl.expected.dxc.hlsl
@@ -1,11 +1,3 @@
-uint tint_div(uint lhs, uint rhs) {
- return (lhs / ((rhs == 0u) ? 1u : rhs));
-}
-
-uint tint_mod(uint lhs, uint rhs) {
- return (lhs % ((rhs == 0u) ? 1u : rhs));
-}
-
cbuffer cbuffer_uniforms : register(b0, space0) {
uint4 uniforms[3];
};
@@ -32,6 +24,14 @@
return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z));
}
+uint tint_div(uint lhs, uint rhs) {
+ return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
+uint tint_mod(uint lhs, uint rhs) {
+ return (lhs % ((rhs == 0u) ? 1u : rhs));
+}
+
uint3 toIndex3D(uint gridSize, uint index) {
uint z_1 = tint_div(index, (gridSize * gridSize));
uint y_1 = tint_div((index - ((gridSize * gridSize) * z_1)), gridSize);
diff --git a/test/tint/bug/tint/1113.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1113.wgsl.expected.fxc.hlsl
index 6a0326e..68248ca 100644
--- a/test/tint/bug/tint/1113.wgsl.expected.fxc.hlsl
+++ b/test/tint/bug/tint/1113.wgsl.expected.fxc.hlsl
@@ -1,11 +1,3 @@
-uint tint_div(uint lhs, uint rhs) {
- return (lhs / ((rhs == 0u) ? 1u : rhs));
-}
-
-uint tint_mod(uint lhs, uint rhs) {
- return (lhs % ((rhs == 0u) ? 1u : rhs));
-}
-
cbuffer cbuffer_uniforms : register(b0, space0) {
uint4 uniforms[3];
};
@@ -32,6 +24,14 @@
return ((icoord.x + (gridSize * icoord.y)) + ((gridSize * gridSize) * icoord.z));
}
+uint tint_div(uint lhs, uint rhs) {
+ return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
+uint tint_mod(uint lhs, uint rhs) {
+ return (lhs % ((rhs == 0u) ? 1u : rhs));
+}
+
uint3 toIndex3D(uint gridSize, uint index) {
uint z_1 = tint_div(index, (gridSize * gridSize));
uint y_1 = tint_div((index - ((gridSize * gridSize) * z_1)), gridSize);
diff --git a/test/tint/bug/tint/1113.wgsl.expected.msl b/test/tint/bug/tint/1113.wgsl.expected.msl
index 0ab0884..2660500 100644
--- a/test/tint/bug/tint/1113.wgsl.expected.msl
+++ b/test/tint/bug/tint/1113.wgsl.expected.msl
@@ -14,14 +14,6 @@
T elements[N];
};
-uint tint_div(uint lhs, uint rhs) {
- return (lhs / select(rhs, 1u, (rhs == 0u)));
-}
-
-uint tint_mod(uint lhs, uint rhs) {
- return (lhs % select(rhs, 1u, (rhs == 0u)));
-}
-
struct Uniforms {
/* 0x0000 */ uint numTriangles;
/* 0x0004 */ uint gridSize;
@@ -85,6 +77,14 @@
return ((icoord[0] + (gridSize * icoord[1])) + ((gridSize * gridSize) * icoord[2]));
}
+uint tint_div(uint lhs, uint rhs) {
+ return (lhs / select(rhs, 1u, (rhs == 0u)));
+}
+
+uint tint_mod(uint lhs, uint rhs) {
+ return (lhs % select(rhs, 1u, (rhs == 0u)));
+}
+
uint3 toIndex3D(uint gridSize, uint index) {
uint z_1 = tint_div(index, (gridSize * gridSize));
uint y_1 = tint_div((index - ((gridSize * gridSize) * z_1)), gridSize);
diff --git a/test/tint/bug/tint/1113.wgsl.expected.spvasm b/test/tint/bug/tint/1113.wgsl.expected.spvasm
index f1eee0c..9544d46 100644
--- a/test/tint/bug/tint/1113.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/1113.wgsl.expected.spvasm
@@ -4,7 +4,7 @@
; Bound: 419
; Schema: 0
OpCapability Shader
- %84 = OpExtInstImport "GLSL.std.450"
+ %67 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main_count "main_count" %GlobalInvocationID_1
OpEntryPoint GLCompute %main_create_lut "main_create_lut" %GlobalInvocationID_2
@@ -53,12 +53,6 @@
OpMemberName %Dbg 10 "value_f32_2"
OpMemberName %Dbg 11 "value_f32_3"
OpName %dbg "dbg"
- OpName %tint_div "tint_div"
- OpName %lhs "lhs"
- OpName %rhs "rhs"
- OpName %tint_mod "tint_mod"
- OpName %lhs_0 "lhs"
- OpName %rhs_0 "rhs"
OpName %toVoxelPos "toVoxelPos"
OpName %position "position"
OpName %bbMin "bbMin"
@@ -73,6 +67,12 @@
OpName %gridSize_0 "gridSize"
OpName %voxelPos "voxelPos"
OpName %icoord "icoord"
+ OpName %tint_div "tint_div"
+ OpName %lhs "lhs"
+ OpName %rhs "rhs"
+ OpName %tint_mod "tint_mod"
+ OpName %lhs_0 "lhs"
+ OpName %rhs_0 "rhs"
OpName %toIndex3D "toIndex3D"
OpName %gridSize_1 "gridSize"
OpName %index "index"
@@ -207,25 +207,25 @@
%dbg_block = OpTypeStruct %Dbg
%_ptr_StorageBuffer_dbg_block = OpTypePointer StorageBuffer %dbg_block
%dbg = OpVariable %_ptr_StorageBuffer_dbg_block StorageBuffer
- %34 = OpTypeFunction %uint %uint %uint
- %40 = OpConstantNull %uint
- %bool = OpTypeBool
- %uint_1 = OpConstant %uint 1
- %52 = OpTypeFunction %v3float %v3float
+ %34 = OpTypeFunction %v3float %v3float
%uint_0 = OpConstant %uint 0
%uint_4 = OpConstant %uint 4
%_ptr_Uniform_float = OpTypePointer Uniform %float
+ %uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%_ptr_Function_v3float = OpTypePointer Function %v3float
- %69 = OpConstantNull %v3float
+ %52 = OpConstantNull %v3float
%uint_5 = OpConstant %uint 5
%_ptr_Function_float = OpTypePointer Function %float
- %94 = OpConstantNull %float
+ %77 = OpConstantNull %float
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
- %131 = OpTypeFunction %uint %uint %v3float
+ %114 = OpTypeFunction %uint %uint %v3float
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
- %139 = OpConstantNull %v3uint
+ %122 = OpConstantNull %v3uint
%_ptr_Function_uint = OpTypePointer Function %uint
+ %135 = OpTypeFunction %uint %uint %uint
+ %141 = OpConstantNull %uint
+ %bool = OpTypeBool
%152 = OpTypeFunction %v3uint %uint %uint
%172 = OpTypeFunction %v3float %uint
%uint_3 = OpConstant %uint 3
@@ -244,128 +244,128 @@
%uint_10 = OpConstant %uint 10
%int_n1 = OpConstant %int -1
%int_1 = OpConstant %int 1
- %tint_div = OpFunction %uint None %34
- %lhs = OpFunctionParameter %uint
- %rhs = OpFunctionParameter %uint
- %38 = OpLabel
- %41 = OpIEqual %bool %rhs %40
- %39 = OpSelect %uint %41 %uint_1 %rhs
- %44 = OpUDiv %uint %lhs %39
- OpReturnValue %44
- OpFunctionEnd
- %tint_mod = OpFunction %uint None %34
- %lhs_0 = OpFunctionParameter %uint
- %rhs_0 = OpFunctionParameter %uint
- %48 = OpLabel
- %50 = OpIEqual %bool %rhs_0 %40
- %49 = OpSelect %uint %50 %uint_1 %rhs_0
- %51 = OpUMod %uint %lhs_0 %49
- OpReturnValue %51
- OpFunctionEnd
- %toVoxelPos = OpFunction %v3float None %52
+ %toVoxelPos = OpFunction %v3float None %34
%position = OpFunctionParameter %v3float
- %55 = OpLabel
- %bbMin = OpVariable %_ptr_Function_v3float Function %69
- %bbMax = OpVariable %_ptr_Function_v3float Function %69
- %bbSize = OpVariable %_ptr_Function_v3float Function %69
- %cubeSize = OpVariable %_ptr_Function_float Function %94
- %gridSize = OpVariable %_ptr_Function_float Function %94
- %gx = OpVariable %_ptr_Function_float Function %94
- %gy = OpVariable %_ptr_Function_float Function %94
- %gz = OpVariable %_ptr_Function_float Function %94
- %59 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_0
- %60 = OpLoad %float %59
- %61 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_1
- %62 = OpLoad %float %61
- %64 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_2
- %65 = OpLoad %float %64
- %66 = OpCompositeConstruct %v3float %60 %62 %65
- OpStore %bbMin %66
- %71 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_5 %uint_0
- %72 = OpLoad %float %71
- %73 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_5 %uint_1
- %74 = OpLoad %float %73
- %75 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_5 %uint_2
- %76 = OpLoad %float %75
- %77 = OpCompositeConstruct %v3float %72 %74 %76
- OpStore %bbMax %77
- %79 = OpLoad %v3float %bbMax
- %80 = OpLoad %v3float %bbMin
- %81 = OpFSub %v3float %79 %80
- OpStore %bbSize %81
- %87 = OpAccessChain %_ptr_Function_float %bbSize %uint_0
- %88 = OpLoad %float %87
- %89 = OpAccessChain %_ptr_Function_float %bbSize %uint_1
- %90 = OpLoad %float %89
- %85 = OpExtInst %float %84 NMax %88 %90
- %91 = OpAccessChain %_ptr_Function_float %bbSize %uint_2
- %92 = OpLoad %float %91
- %83 = OpExtInst %float %84 NMax %85 %92
- OpStore %cubeSize %83
- %97 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_1
- %98 = OpLoad %uint %97
- %95 = OpConvertUToF %float %98
- OpStore %gridSize %95
- %100 = OpLoad %float %gridSize
- %101 = OpCompositeExtract %float %position 0
- %102 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_0
- %103 = OpLoad %float %102
- %104 = OpFSub %float %101 %103
- %105 = OpFMul %float %100 %104
- %106 = OpLoad %float %cubeSize
- %107 = OpFDiv %float %105 %106
- OpStore %gx %107
- %109 = OpLoad %float %gridSize
- %110 = OpCompositeExtract %float %position 1
- %111 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_1
- %112 = OpLoad %float %111
- %113 = OpFSub %float %110 %112
- %114 = OpFMul %float %109 %113
- %115 = OpLoad %float %cubeSize
- %116 = OpFDiv %float %114 %115
- OpStore %gy %116
- %118 = OpLoad %float %gridSize
- %119 = OpCompositeExtract %float %position 2
- %120 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_2
- %121 = OpLoad %float %120
- %122 = OpFSub %float %119 %121
- %123 = OpFMul %float %118 %122
- %124 = OpLoad %float %cubeSize
- %125 = OpFDiv %float %123 %124
- OpStore %gz %125
- %127 = OpLoad %float %gx
- %128 = OpLoad %float %gy
- %129 = OpLoad %float %gz
- %130 = OpCompositeConstruct %v3float %127 %128 %129
- OpReturnValue %130
+ %37 = OpLabel
+ %bbMin = OpVariable %_ptr_Function_v3float Function %52
+ %bbMax = OpVariable %_ptr_Function_v3float Function %52
+ %bbSize = OpVariable %_ptr_Function_v3float Function %52
+ %cubeSize = OpVariable %_ptr_Function_float Function %77
+ %gridSize = OpVariable %_ptr_Function_float Function %77
+ %gx = OpVariable %_ptr_Function_float Function %77
+ %gy = OpVariable %_ptr_Function_float Function %77
+ %gz = OpVariable %_ptr_Function_float Function %77
+ %41 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_0
+ %42 = OpLoad %float %41
+ %44 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_1
+ %45 = OpLoad %float %44
+ %47 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_2
+ %48 = OpLoad %float %47
+ %49 = OpCompositeConstruct %v3float %42 %45 %48
+ OpStore %bbMin %49
+ %54 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_5 %uint_0
+ %55 = OpLoad %float %54
+ %56 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_5 %uint_1
+ %57 = OpLoad %float %56
+ %58 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_5 %uint_2
+ %59 = OpLoad %float %58
+ %60 = OpCompositeConstruct %v3float %55 %57 %59
+ OpStore %bbMax %60
+ %62 = OpLoad %v3float %bbMax
+ %63 = OpLoad %v3float %bbMin
+ %64 = OpFSub %v3float %62 %63
+ OpStore %bbSize %64
+ %70 = OpAccessChain %_ptr_Function_float %bbSize %uint_0
+ %71 = OpLoad %float %70
+ %72 = OpAccessChain %_ptr_Function_float %bbSize %uint_1
+ %73 = OpLoad %float %72
+ %68 = OpExtInst %float %67 NMax %71 %73
+ %74 = OpAccessChain %_ptr_Function_float %bbSize %uint_2
+ %75 = OpLoad %float %74
+ %66 = OpExtInst %float %67 NMax %68 %75
+ OpStore %cubeSize %66
+ %80 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_1
+ %81 = OpLoad %uint %80
+ %78 = OpConvertUToF %float %81
+ OpStore %gridSize %78
+ %83 = OpLoad %float %gridSize
+ %84 = OpCompositeExtract %float %position 0
+ %85 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_0
+ %86 = OpLoad %float %85
+ %87 = OpFSub %float %84 %86
+ %88 = OpFMul %float %83 %87
+ %89 = OpLoad %float %cubeSize
+ %90 = OpFDiv %float %88 %89
+ OpStore %gx %90
+ %92 = OpLoad %float %gridSize
+ %93 = OpCompositeExtract %float %position 1
+ %94 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_1
+ %95 = OpLoad %float %94
+ %96 = OpFSub %float %93 %95
+ %97 = OpFMul %float %92 %96
+ %98 = OpLoad %float %cubeSize
+ %99 = OpFDiv %float %97 %98
+ OpStore %gy %99
+ %101 = OpLoad %float %gridSize
+ %102 = OpCompositeExtract %float %position 2
+ %103 = OpAccessChain %_ptr_Uniform_float %uniforms %uint_0 %uint_4 %uint_2
+ %104 = OpLoad %float %103
+ %105 = OpFSub %float %102 %104
+ %106 = OpFMul %float %101 %105
+ %107 = OpLoad %float %cubeSize
+ %108 = OpFDiv %float %106 %107
+ OpStore %gz %108
+ %110 = OpLoad %float %gx
+ %111 = OpLoad %float %gy
+ %112 = OpLoad %float %gz
+ %113 = OpCompositeConstruct %v3float %110 %111 %112
+ OpReturnValue %113
OpFunctionEnd
- %toIndex1D = OpFunction %uint None %131
+ %toIndex1D = OpFunction %uint None %114
%gridSize_0 = OpFunctionParameter %uint
%voxelPos = OpFunctionParameter %v3float
- %135 = OpLabel
- %icoord = OpVariable %_ptr_Function_v3uint Function %139
- %136 = OpConvertFToU %v3uint %voxelPos
- OpStore %icoord %136
- %141 = OpAccessChain %_ptr_Function_uint %icoord %uint_0
- %142 = OpLoad %uint %141
- %143 = OpAccessChain %_ptr_Function_uint %icoord %uint_1
- %144 = OpLoad %uint %143
- %145 = OpIMul %uint %gridSize_0 %144
- %146 = OpIAdd %uint %142 %145
- %147 = OpIMul %uint %gridSize_0 %gridSize_0
- %148 = OpAccessChain %_ptr_Function_uint %icoord %uint_2
- %149 = OpLoad %uint %148
- %150 = OpIMul %uint %147 %149
- %151 = OpIAdd %uint %146 %150
+ %118 = OpLabel
+ %icoord = OpVariable %_ptr_Function_v3uint Function %122
+ %119 = OpConvertFToU %v3uint %voxelPos
+ OpStore %icoord %119
+ %124 = OpAccessChain %_ptr_Function_uint %icoord %uint_0
+ %125 = OpLoad %uint %124
+ %126 = OpAccessChain %_ptr_Function_uint %icoord %uint_1
+ %127 = OpLoad %uint %126
+ %128 = OpIMul %uint %gridSize_0 %127
+ %129 = OpIAdd %uint %125 %128
+ %130 = OpIMul %uint %gridSize_0 %gridSize_0
+ %131 = OpAccessChain %_ptr_Function_uint %icoord %uint_2
+ %132 = OpLoad %uint %131
+ %133 = OpIMul %uint %130 %132
+ %134 = OpIAdd %uint %129 %133
+ OpReturnValue %134
+ OpFunctionEnd
+ %tint_div = OpFunction %uint None %135
+ %lhs = OpFunctionParameter %uint
+ %rhs = OpFunctionParameter %uint
+ %139 = OpLabel
+ %142 = OpIEqual %bool %rhs %141
+ %140 = OpSelect %uint %142 %uint_1 %rhs
+ %144 = OpUDiv %uint %lhs %140
+ OpReturnValue %144
+ OpFunctionEnd
+ %tint_mod = OpFunction %uint None %135
+ %lhs_0 = OpFunctionParameter %uint
+ %rhs_0 = OpFunctionParameter %uint
+ %148 = OpLabel
+ %150 = OpIEqual %bool %rhs_0 %141
+ %149 = OpSelect %uint %150 %uint_1 %rhs_0
+ %151 = OpUMod %uint %lhs_0 %149
OpReturnValue %151
OpFunctionEnd
%toIndex3D = OpFunction %v3uint None %152
%gridSize_1 = OpFunctionParameter %uint
%index = OpFunctionParameter %uint
%156 = OpLabel
- %z = OpVariable %_ptr_Function_uint Function %40
- %y = OpVariable %_ptr_Function_uint Function %40
- %x = OpVariable %_ptr_Function_uint Function %40
+ %z = OpVariable %_ptr_Function_uint Function %141
+ %y = OpVariable %_ptr_Function_uint Function %141
+ %x = OpVariable %_ptr_Function_uint Function %141
%158 = OpIMul %uint %gridSize_1 %gridSize_1
%157 = OpFunctionCall %uint %tint_div %index %158
OpStore %z %157
@@ -386,9 +386,9 @@
%loadPosition = OpFunction %v3float None %172
%vertexIndex = OpFunctionParameter %uint
%175 = OpLabel
- %position_0 = OpVariable %_ptr_Function_v3float Function %69
+ %position_0 = OpVariable %_ptr_Function_v3float Function %52
%177 = OpIMul %uint %uint_3 %vertexIndex
- %178 = OpIAdd %uint %177 %40
+ %178 = OpIAdd %uint %177 %141
%180 = OpAccessChain %_ptr_StorageBuffer_float %positions %uint_0 %178
%181 = OpLoad %float %180
%182 = OpIMul %uint %uint_3 %vertexIndex
@@ -406,11 +406,11 @@
OpFunctionEnd
%doIgnore = OpFunction %void None %193
%196 = OpLabel
- %g42 = OpVariable %_ptr_Function_uint Function %40
- %kj6 = OpVariable %_ptr_Function_uint Function %40
- %b53 = OpVariable %_ptr_Function_uint Function %40
- %rwg = OpVariable %_ptr_Function_uint Function %40
- %rb5 = OpVariable %_ptr_Function_float Function %94
+ %g42 = OpVariable %_ptr_Function_uint Function %141
+ %kj6 = OpVariable %_ptr_Function_uint Function %141
+ %b53 = OpVariable %_ptr_Function_uint Function %141
+ %rwg = OpVariable %_ptr_Function_uint Function %141
+ %rb5 = OpVariable %_ptr_Function_float Function %77
%g55 = OpVariable %_ptr_Function_int Function %206
%197 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_0
%198 = OpLoad %uint %197
@@ -435,18 +435,18 @@
%main_count_inner = OpFunction %void None %222
%GlobalInvocationID = OpFunctionParameter %v3uint
%225 = OpLabel
-%triangleIndex = OpVariable %_ptr_Function_uint Function %40
- %i0 = OpVariable %_ptr_Function_uint Function %40
- %i1 = OpVariable %_ptr_Function_uint Function %40
- %i2 = OpVariable %_ptr_Function_uint Function %40
- %p0 = OpVariable %_ptr_Function_v3float Function %69
- %p1 = OpVariable %_ptr_Function_v3float Function %69
- %p2 = OpVariable %_ptr_Function_v3float Function %69
- %269 = OpVariable %_ptr_Function_v3float Function %69
- %center = OpVariable %_ptr_Function_v3float Function %69
- %voxelPos_0 = OpVariable %_ptr_Function_v3float Function %69
- %voxelIndex = OpVariable %_ptr_Function_uint Function %40
- %acefg = OpVariable %_ptr_Function_uint Function %40
+%triangleIndex = OpVariable %_ptr_Function_uint Function %141
+ %i0 = OpVariable %_ptr_Function_uint Function %141
+ %i1 = OpVariable %_ptr_Function_uint Function %141
+ %i2 = OpVariable %_ptr_Function_uint Function %141
+ %p0 = OpVariable %_ptr_Function_v3float Function %52
+ %p1 = OpVariable %_ptr_Function_v3float Function %52
+ %p2 = OpVariable %_ptr_Function_v3float Function %52
+ %269 = OpVariable %_ptr_Function_v3float Function %52
+ %center = OpVariable %_ptr_Function_v3float Function %52
+ %voxelPos_0 = OpVariable %_ptr_Function_v3float Function %52
+ %voxelIndex = OpVariable %_ptr_Function_uint Function %141
+ %acefg = OpVariable %_ptr_Function_uint Function %141
%226 = OpCompositeExtract %uint %GlobalInvocationID 0
OpStore %triangleIndex %226
%228 = OpLoad %uint %triangleIndex
@@ -461,7 +461,7 @@
%234 = OpFunctionCall %void %doIgnore
%235 = OpLoad %uint %triangleIndex
%236 = OpIMul %uint %uint_3 %235
- %237 = OpIAdd %uint %236 %40
+ %237 = OpIAdd %uint %236 %141
%238 = OpAccessChain %_ptr_StorageBuffer_uint %indices %uint_0 %237
%239 = OpLoad %uint %238
OpStore %i0 %239
@@ -507,7 +507,7 @@
%280 = OpAtomicIAdd %uint %283 %uint_1 %uint_0 %uint_1
OpStore %acefg %280
%285 = OpLoad %uint %triangleIndex
- %286 = OpIEqual %bool %285 %40
+ %286 = OpIEqual %bool %285 %141
OpSelectionMerge %287 None
OpBranchConditional %286 %288 %287
%288 = OpLabel
@@ -540,9 +540,9 @@
%main_create_lut_inner = OpFunction %void None %222
%GlobalInvocationID_0 = OpFunctionParameter %v3uint
%310 = OpLabel
-%voxelIndex_0 = OpVariable %_ptr_Function_uint Function %40
- %maxVoxels = OpVariable %_ptr_Function_uint Function %40
-%numTriangles = OpVariable %_ptr_Function_uint Function %40
+%voxelIndex_0 = OpVariable %_ptr_Function_uint Function %141
+ %maxVoxels = OpVariable %_ptr_Function_uint Function %141
+%numTriangles = OpVariable %_ptr_Function_uint Function %141
%offset = OpVariable %_ptr_Function_int Function %206
%311 = OpCompositeExtract %uint %GlobalInvocationID_0 0
OpStore %voxelIndex_0 %311
@@ -570,7 +570,7 @@
OpStore %numTriangles %328
OpStore %offset %int_n1
%335 = OpLoad %uint %numTriangles
- %336 = OpUGreaterThan %bool %335 %40
+ %336 = OpUGreaterThan %bool %335 %141
OpSelectionMerge %337 None
OpBranchConditional %336 %338 %337
%338 = OpLabel
@@ -596,17 +596,17 @@
%main_sort_triangles_inner = OpFunction %void None %222
%GlobalInvocationID_4 = OpFunctionParameter %v3uint
%355 = OpLabel
-%triangleIndex_0 = OpVariable %_ptr_Function_uint Function %40
- %i0_0 = OpVariable %_ptr_Function_uint Function %40
- %i1_0 = OpVariable %_ptr_Function_uint Function %40
- %i2_0 = OpVariable %_ptr_Function_uint Function %40
- %p0_0 = OpVariable %_ptr_Function_v3float Function %69
- %p1_0 = OpVariable %_ptr_Function_v3float Function %69
- %p2_0 = OpVariable %_ptr_Function_v3float Function %69
- %398 = OpVariable %_ptr_Function_v3float Function %69
- %center_0 = OpVariable %_ptr_Function_v3float Function %69
- %voxelPos_1 = OpVariable %_ptr_Function_v3float Function %69
-%voxelIndex_1 = OpVariable %_ptr_Function_uint Function %40
+%triangleIndex_0 = OpVariable %_ptr_Function_uint Function %141
+ %i0_0 = OpVariable %_ptr_Function_uint Function %141
+ %i1_0 = OpVariable %_ptr_Function_uint Function %141
+ %i2_0 = OpVariable %_ptr_Function_uint Function %141
+ %p0_0 = OpVariable %_ptr_Function_v3float Function %52
+ %p1_0 = OpVariable %_ptr_Function_v3float Function %52
+ %p2_0 = OpVariable %_ptr_Function_v3float Function %52
+ %398 = OpVariable %_ptr_Function_v3float Function %52
+ %center_0 = OpVariable %_ptr_Function_v3float Function %52
+ %voxelPos_1 = OpVariable %_ptr_Function_v3float Function %52
+%voxelIndex_1 = OpVariable %_ptr_Function_uint Function %141
%triangleOffset = OpVariable %_ptr_Function_int Function %206
%356 = OpCompositeExtract %uint %GlobalInvocationID_4 0
OpStore %triangleIndex_0 %356
@@ -622,7 +622,7 @@
%363 = OpLabel
%365 = OpLoad %uint %triangleIndex_0
%366 = OpIMul %uint %uint_3 %365
- %367 = OpIAdd %uint %366 %40
+ %367 = OpIAdd %uint %366 %141
%368 = OpAccessChain %_ptr_StorageBuffer_uint %indices %uint_0 %367
%369 = OpLoad %uint %368
OpStore %i0_0 %369
diff --git a/test/tint/bug/tint/1520.spvasm.expected.dxc.hlsl b/test/tint/bug/tint/1520.spvasm.expected.dxc.hlsl
index e07e2fa..06255d0 100644
--- a/test/tint/bug/tint/1520.spvasm.expected.dxc.hlsl
+++ b/test/tint/bug/tint/1520.spvasm.expected.dxc.hlsl
@@ -1,7 +1,3 @@
-int4 tint_div(int4 lhs, int4 rhs) {
- return (lhs / (((rhs == (0).xxxx) | ((lhs == (-2147483648).xxxx) & (rhs == (-1).xxxx))) ? (1).xxxx : rhs));
-}
-
cbuffer cbuffer_x_4 : register(b0, space0) {
uint4 x_4[7];
};
@@ -9,6 +5,10 @@
static bool sk_Clockwise = false;
static float4 vcolor_S0 = float4(0.0f, 0.0f, 0.0f, 0.0f);
+int4 tint_div(int4 lhs, int4 rhs) {
+ return (lhs / (((rhs == (0).xxxx) | ((lhs == (-2147483648).xxxx) & (rhs == (-1).xxxx))) ? (1).xxxx : rhs));
+}
+
bool test_int_S1_c0_b() {
int unknown = 0;
bool ok = false;
diff --git a/test/tint/bug/tint/1520.spvasm.expected.fxc.hlsl b/test/tint/bug/tint/1520.spvasm.expected.fxc.hlsl
index e07e2fa..06255d0 100644
--- a/test/tint/bug/tint/1520.spvasm.expected.fxc.hlsl
+++ b/test/tint/bug/tint/1520.spvasm.expected.fxc.hlsl
@@ -1,7 +1,3 @@
-int4 tint_div(int4 lhs, int4 rhs) {
- return (lhs / (((rhs == (0).xxxx) | ((lhs == (-2147483648).xxxx) & (rhs == (-1).xxxx))) ? (1).xxxx : rhs));
-}
-
cbuffer cbuffer_x_4 : register(b0, space0) {
uint4 x_4[7];
};
@@ -9,6 +5,10 @@
static bool sk_Clockwise = false;
static float4 vcolor_S0 = float4(0.0f, 0.0f, 0.0f, 0.0f);
+int4 tint_div(int4 lhs, int4 rhs) {
+ return (lhs / (((rhs == (0).xxxx) | ((lhs == (-2147483648).xxxx) & (rhs == (-1).xxxx))) ? (1).xxxx : rhs));
+}
+
bool test_int_S1_c0_b() {
int unknown = 0;
bool ok = false;
diff --git a/test/tint/bug/tint/1520.spvasm.expected.glsl b/test/tint/bug/tint/1520.spvasm.expected.glsl
index fbf2852..8cd026f 100644
--- a/test/tint/bug/tint/1520.spvasm.expected.glsl
+++ b/test/tint/bug/tint/1520.spvasm.expected.glsl
@@ -3,10 +3,6 @@
layout(location = 0) in vec4 vcolor_S0_param_1;
layout(location = 0) out vec4 sk_FragColor_1_1;
-ivec4 tint_div(ivec4 lhs, ivec4 rhs) {
- return (lhs / mix(rhs, ivec4(1), bvec4(uvec4(equal(rhs, ivec4(0))) | uvec4(bvec4(uvec4(equal(lhs, ivec4(-2147483648))) & uvec4(equal(rhs, ivec4(-1))))))));
-}
-
struct UniformBuffer {
uint pad;
uint pad_1;
@@ -28,6 +24,10 @@
vec4 sk_FragColor = vec4(0.0f, 0.0f, 0.0f, 0.0f);
bool sk_Clockwise = false;
vec4 vcolor_S0 = vec4(0.0f, 0.0f, 0.0f, 0.0f);
+ivec4 tint_div(ivec4 lhs, ivec4 rhs) {
+ return (lhs / mix(rhs, ivec4(1), bvec4(uvec4(equal(rhs, ivec4(0))) | uvec4(bvec4(uvec4(equal(lhs, ivec4(-2147483648))) & uvec4(equal(rhs, ivec4(-1))))))));
+}
+
bool test_int_S1_c0_b() {
int unknown = 0;
bool ok = false;
diff --git a/test/tint/bug/tint/1520.spvasm.expected.msl b/test/tint/bug/tint/1520.spvasm.expected.msl
index b997052..335a96b 100644
--- a/test/tint/bug/tint/1520.spvasm.expected.msl
+++ b/test/tint/bug/tint/1520.spvasm.expected.msl
@@ -14,10 +14,6 @@
T elements[N];
};
-int4 tint_div(int4 lhs, int4 rhs) {
- return (lhs / select(rhs, int4(1), ((rhs == int4(0)) | ((lhs == int4((-2147483647 - 1))) & (rhs == int4(-1))))));
-}
-
struct UniformBuffer {
/* 0x0000 */ tint_array<int8_t, 16> tint_pad;
/* 0x0010 */ float unknownInput_S1_c0;
@@ -27,6 +23,10 @@
/* 0x0040 */ float3x3 umatrix_S1;
};
+int4 tint_div(int4 lhs, int4 rhs) {
+ return (lhs / select(rhs, int4(1), ((rhs == int4(0)) | ((lhs == int4((-2147483647 - 1))) & (rhs == int4(-1))))));
+}
+
bool test_int_S1_c0_b(const constant UniformBuffer* const tint_symbol_6) {
int unknown = 0;
bool ok = false;
diff --git a/test/tint/bug/tint/221.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/221.wgsl.expected.dxc.hlsl
index 708904c..4038566 100644
--- a/test/tint/bug/tint/221.wgsl.expected.dxc.hlsl
+++ b/test/tint/bug/tint/221.wgsl.expected.dxc.hlsl
@@ -1,9 +1,9 @@
+RWByteAddressBuffer b : register(u0, space0);
+
uint tint_mod(uint lhs, uint rhs) {
return (lhs % ((rhs == 0u) ? 1u : rhs));
}
-RWByteAddressBuffer b : register(u0, space0);
-
[numthreads(1, 1, 1)]
void main() {
uint i = 0u;
diff --git a/test/tint/bug/tint/221.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/221.wgsl.expected.fxc.hlsl
index 708904c..4038566 100644
--- a/test/tint/bug/tint/221.wgsl.expected.fxc.hlsl
+++ b/test/tint/bug/tint/221.wgsl.expected.fxc.hlsl
@@ -1,9 +1,9 @@
+RWByteAddressBuffer b : register(u0, space0);
+
uint tint_mod(uint lhs, uint rhs) {
return (lhs % ((rhs == 0u) ? 1u : rhs));
}
-RWByteAddressBuffer b : register(u0, space0);
-
[numthreads(1, 1, 1)]
void main() {
uint i = 0u;
diff --git a/test/tint/bug/tint/221.wgsl.expected.glsl b/test/tint/bug/tint/221.wgsl.expected.glsl
index 5f6890f..4bcee5e 100644
--- a/test/tint/bug/tint/221.wgsl.expected.glsl
+++ b/test/tint/bug/tint/221.wgsl.expected.glsl
@@ -1,9 +1,5 @@
#version 310 es
-uint tint_mod(uint lhs, uint rhs) {
- return (lhs % ((rhs == 0u) ? 1u : rhs));
-}
-
struct Buf {
uint count;
uint data[50];
@@ -13,6 +9,10 @@
Buf inner;
} b;
+uint tint_mod(uint lhs, uint rhs) {
+ return (lhs % ((rhs == 0u) ? 1u : rhs));
+}
+
void tint_symbol() {
uint i = 0u;
while (true) {
diff --git a/test/tint/bug/tint/221.wgsl.expected.msl b/test/tint/bug/tint/221.wgsl.expected.msl
index 7e30a18..8a82c43 100644
--- a/test/tint/bug/tint/221.wgsl.expected.msl
+++ b/test/tint/bug/tint/221.wgsl.expected.msl
@@ -14,15 +14,15 @@
T elements[N];
};
-uint tint_mod(uint lhs, uint rhs) {
- return (lhs % select(rhs, 1u, (rhs == 0u)));
-}
-
struct Buf {
/* 0x0000 */ uint count;
/* 0x0004 */ tint_array<uint, 50> data;
};
+uint tint_mod(uint lhs, uint rhs) {
+ return (lhs % select(rhs, 1u, (rhs == 0u)));
+}
+
kernel void tint_symbol(device Buf* tint_symbol_2 [[buffer(0)]]) {
uint i = 0u;
while (true) {
diff --git a/test/tint/bug/tint/914.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/914.wgsl.expected.dxc.hlsl
index a893add..6b741d6 100644
--- a/test/tint/bug/tint/914.wgsl.expected.dxc.hlsl
+++ b/test/tint/bug/tint/914.wgsl.expected.dxc.hlsl
@@ -1,7 +1,3 @@
-uint tint_div(uint lhs, uint rhs) {
- return (lhs / ((rhs == 0u) ? 1u : rhs));
-}
-
ByteAddressBuffer firstMatrix : register(t0, space0);
ByteAddressBuffer secondMatrix : register(t1, space0);
RWByteAddressBuffer resultMatrix : register(u2, space0);
@@ -47,6 +43,10 @@
groupshared float mm_Asub[64][64];
groupshared float mm_Bsub[64][64];
+uint tint_div(uint lhs, uint rhs) {
+ return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
struct tint_symbol_1 {
uint3 local_id : SV_GroupThreadID;
uint local_invocation_index : SV_GroupIndex;
diff --git a/test/tint/bug/tint/914.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/914.wgsl.expected.fxc.hlsl
index a893add..6b741d6 100644
--- a/test/tint/bug/tint/914.wgsl.expected.fxc.hlsl
+++ b/test/tint/bug/tint/914.wgsl.expected.fxc.hlsl
@@ -1,7 +1,3 @@
-uint tint_div(uint lhs, uint rhs) {
- return (lhs / ((rhs == 0u) ? 1u : rhs));
-}
-
ByteAddressBuffer firstMatrix : register(t0, space0);
ByteAddressBuffer secondMatrix : register(t1, space0);
RWByteAddressBuffer resultMatrix : register(u2, space0);
@@ -47,6 +43,10 @@
groupshared float mm_Asub[64][64];
groupshared float mm_Bsub[64][64];
+uint tint_div(uint lhs, uint rhs) {
+ return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
struct tint_symbol_1 {
uint3 local_id : SV_GroupThreadID;
uint local_invocation_index : SV_GroupIndex;
diff --git a/test/tint/bug/tint/914.wgsl.expected.glsl b/test/tint/bug/tint/914.wgsl.expected.glsl
index 2797131..fe3d6bc 100644
--- a/test/tint/bug/tint/914.wgsl.expected.glsl
+++ b/test/tint/bug/tint/914.wgsl.expected.glsl
@@ -1,9 +1,5 @@
#version 310 es
-uint tint_div(uint lhs, uint rhs) {
- return (lhs / ((rhs == 0u) ? 1u : rhs));
-}
-
struct Uniforms {
uint dimAOuter;
uint dimInner;
@@ -64,6 +60,10 @@
shared float mm_Asub[64][64];
shared float mm_Bsub[64][64];
+uint tint_div(uint lhs, uint rhs) {
+ return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
void tint_symbol(uvec3 local_id, uvec3 global_id, uint local_invocation_index) {
{
for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 256u)) {
diff --git a/test/tint/bug/tint/914.wgsl.expected.msl b/test/tint/bug/tint/914.wgsl.expected.msl
index b7dec33..5ea78f2 100644
--- a/test/tint/bug/tint/914.wgsl.expected.msl
+++ b/test/tint/bug/tint/914.wgsl.expected.msl
@@ -14,10 +14,6 @@
T elements[N];
};
-uint tint_div(uint lhs, uint rhs) {
- return (lhs / select(rhs, 1u, (rhs == 0u)));
-}
-
struct Uniforms {
/* 0x0000 */ uint dimAOuter;
/* 0x0004 */ uint dimInner;
@@ -51,6 +47,10 @@
}
}
+uint tint_div(uint lhs, uint rhs) {
+ return (lhs / select(rhs, 1u, (rhs == 0u)));
+}
+
void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_index, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_10, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_11, const constant Uniforms* const tint_symbol_12, const device Matrix* const tint_symbol_13, const device Matrix* const tint_symbol_14, device Matrix* const tint_symbol_15) {
for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 256u)) {
uint const i = (idx / 64u);
diff --git a/test/tint/bug/tint/914.wgsl.expected.spvasm b/test/tint/bug/tint/914.wgsl.expected.spvasm
index f7719e0..829ad7d 100644
--- a/test/tint/bug/tint/914.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/914.wgsl.expected.spvasm
@@ -24,9 +24,6 @@
OpName %uniforms "uniforms"
OpName %mm_Asub "mm_Asub"
OpName %mm_Bsub "mm_Bsub"
- OpName %tint_div "tint_div"
- OpName %lhs "lhs"
- OpName %rhs "rhs"
OpName %mm_readA "mm_readA"
OpName %row "row"
OpName %col "col"
@@ -41,6 +38,9 @@
OpName %row_1 "row"
OpName %col_1 "col"
OpName %value "value"
+ OpName %tint_div "tint_div"
+ OpName %lhs "lhs"
+ OpName %rhs "rhs"
OpName %main_inner "main_inner"
OpName %local_id "local_id"
OpName %global_id "global_id"
@@ -112,22 +112,22 @@
%_ptr_Workgroup__arr__arr_float_uint_64_uint_64 = OpTypePointer Workgroup %_arr__arr_float_uint_64_uint_64
%mm_Asub = OpVariable %_ptr_Workgroup__arr__arr_float_uint_64_uint_64 Workgroup
%mm_Bsub = OpVariable %_ptr_Workgroup__arr__arr_float_uint_64_uint_64 Workgroup
- %25 = OpTypeFunction %uint %uint %uint
- %31 = OpConstantNull %uint
+ %25 = OpTypeFunction %float %uint %uint
%bool = OpTypeBool
- %uint_1 = OpConstant %uint 1
- %36 = OpTypeFunction %float %uint %uint
%_ptr_Function_bool = OpTypePointer Function %bool
- %43 = OpConstantNull %bool
+ %33 = OpConstantNull %bool
%_ptr_Function_float = OpTypePointer Function %float
- %46 = OpConstantNull %float
+ %36 = OpConstantNull %float
%uint_0 = OpConstant %uint 0
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+ %uint_1 = OpConstant %uint 1
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
%true = OpConstantTrue %bool
%uint_2 = OpConstant %uint 2
%void = OpTypeVoid
- %102 = OpTypeFunction %void %uint %uint %float
+ %93 = OpTypeFunction %void %uint %uint %float
+ %116 = OpTypeFunction %uint %uint %uint
+ %122 = OpConstantNull %uint
%125 = OpTypeFunction %void %v3uint %v3uint %uint
%_ptr_Function_uint = OpTypePointer Function %uint
%uint_4096 = OpConstant %uint 4096
@@ -143,152 +143,152 @@
%_ptr_Function__arr_float_uint_4 = OpTypePointer Function %_arr_float_uint_4
%178 = OpConstantNull %_arr_float_uint_4
%390 = OpTypeFunction %void
- %tint_div = OpFunction %uint None %25
- %lhs = OpFunctionParameter %uint
- %rhs = OpFunctionParameter %uint
- %29 = OpLabel
- %32 = OpIEqual %bool %rhs %31
- %30 = OpSelect %uint %32 %uint_1 %rhs
- %35 = OpUDiv %uint %lhs %30
- OpReturnValue %35
- OpFunctionEnd
- %mm_readA = OpFunction %float None %36
+ %mm_readA = OpFunction %float None %25
%row = OpFunctionParameter %uint
%col = OpFunctionParameter %uint
- %40 = OpLabel
-%tint_return_flag = OpVariable %_ptr_Function_bool Function %43
-%tint_return_value = OpVariable %_ptr_Function_float Function %46
- %49 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_0
- %50 = OpLoad %uint %49
- %51 = OpULessThan %bool %row %50
- OpSelectionMerge %52 None
- OpBranchConditional %51 %53 %52
- %53 = OpLabel
- %54 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_1
- %55 = OpLoad %uint %54
- %56 = OpULessThan %bool %col %55
- OpBranch %52
- %52 = OpLabel
- %57 = OpPhi %bool %51 %40 %56 %53
- OpSelectionMerge %58 None
- OpBranchConditional %57 %59 %58
- %59 = OpLabel
- %60 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_1
- %61 = OpLoad %uint %60
- %62 = OpIMul %uint %row %61
- %63 = OpIAdd %uint %62 %col
- %65 = OpAccessChain %_ptr_StorageBuffer_float %firstMatrix %uint_0 %63
- %66 = OpLoad %float %65
+ %29 = OpLabel
+%tint_return_flag = OpVariable %_ptr_Function_bool Function %33
+%tint_return_value = OpVariable %_ptr_Function_float Function %36
+ %39 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_0
+ %40 = OpLoad %uint %39
+ %41 = OpULessThan %bool %row %40
+ OpSelectionMerge %42 None
+ OpBranchConditional %41 %43 %42
+ %43 = OpLabel
+ %45 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_1
+ %46 = OpLoad %uint %45
+ %47 = OpULessThan %bool %col %46
+ OpBranch %42
+ %42 = OpLabel
+ %48 = OpPhi %bool %41 %29 %47 %43
+ OpSelectionMerge %49 None
+ OpBranchConditional %48 %50 %49
+ %50 = OpLabel
+ %51 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_1
+ %52 = OpLoad %uint %51
+ %53 = OpIMul %uint %row %52
+ %54 = OpIAdd %uint %53 %col
+ %56 = OpAccessChain %_ptr_StorageBuffer_float %firstMatrix %uint_0 %54
+ %57 = OpLoad %float %56
OpStore %tint_return_flag %true
- OpStore %tint_return_value %66
- OpBranch %58
- %58 = OpLabel
- %69 = OpLoad %bool %tint_return_flag
- %68 = OpLogicalNot %bool %69
- OpSelectionMerge %70 None
- OpBranchConditional %68 %71 %70
- %71 = OpLabel
+ OpStore %tint_return_value %57
+ OpBranch %49
+ %49 = OpLabel
+ %60 = OpLoad %bool %tint_return_flag
+ %59 = OpLogicalNot %bool %60
+ OpSelectionMerge %61 None
+ OpBranchConditional %59 %62 %61
+ %62 = OpLabel
OpStore %tint_return_flag %true
- OpStore %tint_return_value %46
- OpBranch %70
- %70 = OpLabel
- %72 = OpLoad %float %tint_return_value
- OpReturnValue %72
+ OpStore %tint_return_value %36
+ OpBranch %61
+ %61 = OpLabel
+ %63 = OpLoad %float %tint_return_value
+ OpReturnValue %63
OpFunctionEnd
- %mm_readB = OpFunction %float None %36
+ %mm_readB = OpFunction %float None %25
%row_0 = OpFunctionParameter %uint
%col_0 = OpFunctionParameter %uint
- %76 = OpLabel
-%tint_return_flag_1 = OpVariable %_ptr_Function_bool Function %43
-%tint_return_value_1 = OpVariable %_ptr_Function_float Function %46
- %79 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_1
- %80 = OpLoad %uint %79
- %81 = OpULessThan %bool %row_0 %80
- OpSelectionMerge %82 None
- OpBranchConditional %81 %83 %82
- %83 = OpLabel
- %85 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_2
- %86 = OpLoad %uint %85
- %87 = OpULessThan %bool %col_0 %86
- OpBranch %82
- %82 = OpLabel
- %88 = OpPhi %bool %81 %76 %87 %83
- OpSelectionMerge %89 None
- OpBranchConditional %88 %90 %89
+ %67 = OpLabel
+%tint_return_flag_1 = OpVariable %_ptr_Function_bool Function %33
+%tint_return_value_1 = OpVariable %_ptr_Function_float Function %36
+ %70 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_1
+ %71 = OpLoad %uint %70
+ %72 = OpULessThan %bool %row_0 %71
+ OpSelectionMerge %73 None
+ OpBranchConditional %72 %74 %73
+ %74 = OpLabel
+ %76 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_2
+ %77 = OpLoad %uint %76
+ %78 = OpULessThan %bool %col_0 %77
+ OpBranch %73
+ %73 = OpLabel
+ %79 = OpPhi %bool %72 %67 %78 %74
+ OpSelectionMerge %80 None
+ OpBranchConditional %79 %81 %80
+ %81 = OpLabel
+ %82 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_2
+ %83 = OpLoad %uint %82
+ %84 = OpIMul %uint %row_0 %83
+ %85 = OpIAdd %uint %84 %col_0
+ %86 = OpAccessChain %_ptr_StorageBuffer_float %secondMatrix %uint_0 %85
+ %87 = OpLoad %float %86
+ OpStore %tint_return_flag_1 %true
+ OpStore %tint_return_value_1 %87
+ OpBranch %80
+ %80 = OpLabel
+ %89 = OpLoad %bool %tint_return_flag_1
+ %88 = OpLogicalNot %bool %89
+ OpSelectionMerge %90 None
+ OpBranchConditional %88 %91 %90
+ %91 = OpLabel
+ OpStore %tint_return_flag_1 %true
+ OpStore %tint_return_value_1 %36
+ OpBranch %90
%90 = OpLabel
- %91 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_2
- %92 = OpLoad %uint %91
- %93 = OpIMul %uint %row_0 %92
- %94 = OpIAdd %uint %93 %col_0
- %95 = OpAccessChain %_ptr_StorageBuffer_float %secondMatrix %uint_0 %94
- %96 = OpLoad %float %95
- OpStore %tint_return_flag_1 %true
- OpStore %tint_return_value_1 %96
- OpBranch %89
- %89 = OpLabel
- %98 = OpLoad %bool %tint_return_flag_1
- %97 = OpLogicalNot %bool %98
- OpSelectionMerge %99 None
- OpBranchConditional %97 %100 %99
- %100 = OpLabel
- OpStore %tint_return_flag_1 %true
- OpStore %tint_return_value_1 %46
- OpBranch %99
- %99 = OpLabel
- %101 = OpLoad %float %tint_return_value_1
- OpReturnValue %101
+ %92 = OpLoad %float %tint_return_value_1
+ OpReturnValue %92
OpFunctionEnd
- %mm_write = OpFunction %void None %102
+ %mm_write = OpFunction %void None %93
%row_1 = OpFunctionParameter %uint
%col_1 = OpFunctionParameter %uint
%value = OpFunctionParameter %float
- %108 = OpLabel
- %109 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_0
- %110 = OpLoad %uint %109
- %111 = OpULessThan %bool %row_1 %110
- OpSelectionMerge %112 None
- OpBranchConditional %111 %113 %112
- %113 = OpLabel
- %114 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_2
- %115 = OpLoad %uint %114
- %116 = OpULessThan %bool %col_1 %115
- OpBranch %112
- %112 = OpLabel
- %117 = OpPhi %bool %111 %108 %116 %113
- OpSelectionMerge %118 None
- OpBranchConditional %117 %119 %118
- %119 = OpLabel
- %120 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_2
- %121 = OpLoad %uint %120
- %122 = OpIMul %uint %row_1 %121
- %123 = OpIAdd %uint %col_1 %122
- %124 = OpAccessChain %_ptr_StorageBuffer_float %resultMatrix %uint_0 %123
- OpStore %124 %value
- OpBranch %118
- %118 = OpLabel
+ %99 = OpLabel
+ %100 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_0
+ %101 = OpLoad %uint %100
+ %102 = OpULessThan %bool %row_1 %101
+ OpSelectionMerge %103 None
+ OpBranchConditional %102 %104 %103
+ %104 = OpLabel
+ %105 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_2
+ %106 = OpLoad %uint %105
+ %107 = OpULessThan %bool %col_1 %106
+ OpBranch %103
+ %103 = OpLabel
+ %108 = OpPhi %bool %102 %99 %107 %104
+ OpSelectionMerge %109 None
+ OpBranchConditional %108 %110 %109
+ %110 = OpLabel
+ %111 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_2
+ %112 = OpLoad %uint %111
+ %113 = OpIMul %uint %row_1 %112
+ %114 = OpIAdd %uint %col_1 %113
+ %115 = OpAccessChain %_ptr_StorageBuffer_float %resultMatrix %uint_0 %114
+ OpStore %115 %value
+ OpBranch %109
+ %109 = OpLabel
OpReturn
OpFunctionEnd
+ %tint_div = OpFunction %uint None %116
+ %lhs = OpFunctionParameter %uint
+ %rhs = OpFunctionParameter %uint
+ %120 = OpLabel
+ %123 = OpIEqual %bool %rhs %122
+ %121 = OpSelect %uint %123 %uint_1 %rhs
+ %124 = OpUDiv %uint %lhs %121
+ OpReturnValue %124
+ OpFunctionEnd
%main_inner = OpFunction %void None %125
%local_id = OpFunctionParameter %v3uint
%global_id = OpFunctionParameter %v3uint
%local_invocation_index = OpFunctionParameter %uint
%130 = OpLabel
- %idx = OpVariable %_ptr_Function_uint Function %31
+ %idx = OpVariable %_ptr_Function_uint Function %122
%acc = OpVariable %_ptr_Function__arr_float_uint_16 Function %173
- %ACached = OpVariable %_ptr_Function_float Function %46
+ %ACached = OpVariable %_ptr_Function_float Function %36
%BCached = OpVariable %_ptr_Function__arr_float_uint_4 Function %178
- %index = OpVariable %_ptr_Function_uint Function %31
- %t = OpVariable %_ptr_Function_uint Function %31
- %innerRow = OpVariable %_ptr_Function_uint Function %31
- %innerCol = OpVariable %_ptr_Function_uint Function %31
- %innerRow_0 = OpVariable %_ptr_Function_uint Function %31
- %innerCol_0 = OpVariable %_ptr_Function_uint Function %31
- %k = OpVariable %_ptr_Function_uint Function %31
- %inner = OpVariable %_ptr_Function_uint Function %31
- %innerRow_1 = OpVariable %_ptr_Function_uint Function %31
- %innerCol_1 = OpVariable %_ptr_Function_uint Function %31
- %innerRow_2 = OpVariable %_ptr_Function_uint Function %31
- %innerCol_2 = OpVariable %_ptr_Function_uint Function %31
+ %index = OpVariable %_ptr_Function_uint Function %122
+ %t = OpVariable %_ptr_Function_uint Function %122
+ %innerRow = OpVariable %_ptr_Function_uint Function %122
+ %innerCol = OpVariable %_ptr_Function_uint Function %122
+ %innerRow_0 = OpVariable %_ptr_Function_uint Function %122
+ %innerCol_0 = OpVariable %_ptr_Function_uint Function %122
+ %k = OpVariable %_ptr_Function_uint Function %122
+ %inner = OpVariable %_ptr_Function_uint Function %122
+ %innerRow_1 = OpVariable %_ptr_Function_uint Function %122
+ %innerCol_1 = OpVariable %_ptr_Function_uint Function %122
+ %innerRow_2 = OpVariable %_ptr_Function_uint Function %122
+ %innerCol_2 = OpVariable %_ptr_Function_uint Function %122
OpStore %idx %local_invocation_index
OpBranch %133
%133 = OpLabel
@@ -308,9 +308,9 @@
%145 = OpLoad %uint %idx
%146 = OpUMod %uint %145 %uint_64
%148 = OpAccessChain %_ptr_Workgroup_float %mm_Asub %144 %146
- OpStore %148 %46
+ OpStore %148 %36
%149 = OpAccessChain %_ptr_Workgroup_float %mm_Bsub %144 %146
- OpStore %149 %46
+ OpStore %149 %36
OpBranch %135
%135 = OpLabel
%150 = OpLoad %uint %idx
@@ -332,7 +332,7 @@
%167 = OpISub %uint %166 %uint_1
%164 = OpFunctionCall %uint %tint_div %167 %uint_64
%168 = OpIAdd %uint %164 %uint_1
- OpStore %index %31
+ OpStore %index %122
OpBranch %180
%180 = OpLabel
OpLoopMerge %181 %182 None
@@ -348,7 +348,7 @@
%187 = OpLabel
%189 = OpLoad %uint %index
%190 = OpAccessChain %_ptr_Function_float %acc %189
- OpStore %190 %46
+ OpStore %190 %36
OpBranch %182
%182 = OpLabel
%191 = OpLoad %uint %index
@@ -360,7 +360,7 @@
%194 = OpIMul %uint %193 %uint_4
%195 = OpCompositeExtract %uint %local_id 1
%196 = OpIMul %uint %195 %uint_4
- OpStore %t %31
+ OpStore %t %122
OpBranch %198
%198 = OpLabel
OpLoopMerge %199 %200 None
@@ -374,7 +374,7 @@
%206 = OpLabel
OpBranch %199
%205 = OpLabel
- OpStore %innerRow %31
+ OpStore %innerRow %122
OpBranch %208
%208 = OpLabel
OpLoopMerge %209 %210 None
@@ -388,7 +388,7 @@
%216 = OpLabel
OpBranch %209
%215 = OpLabel
- OpStore %innerCol %31
+ OpStore %innerCol %122
OpBranch %218
%218 = OpLabel
OpLoopMerge %219 %220 None
@@ -428,7 +428,7 @@
OpStore %innerRow %241
OpBranch %208
%209 = OpLabel
- OpStore %innerRow_0 %31
+ OpStore %innerRow_0 %122
OpBranch %243
%243 = OpLabel
OpLoopMerge %244 %245 None
@@ -442,7 +442,7 @@
%251 = OpLabel
OpBranch %244
%250 = OpLabel
- OpStore %innerCol_0 %31
+ OpStore %innerCol_0 %122
OpBranch %253
%253 = OpLabel
OpLoopMerge %254 %255 None
@@ -484,7 +484,7 @@
OpBranch %243
%244 = OpLabel
OpControlBarrier %uint_2 %uint_2 %uint_264
- OpStore %k %31
+ OpStore %k %122
OpBranch %280
%280 = OpLabel
OpLoopMerge %281 %282 None
@@ -498,7 +498,7 @@
%288 = OpLabel
OpBranch %281
%287 = OpLabel
- OpStore %inner %31
+ OpStore %inner %122
OpBranch %290
%290 = OpLabel
OpLoopMerge %291 %292 None
@@ -527,7 +527,7 @@
OpStore %inner %307
OpBranch %290
%291 = OpLabel
- OpStore %innerRow_1 %31
+ OpStore %innerRow_1 %122
OpBranch %309
%309 = OpLabel
OpLoopMerge %310 %311 None
@@ -547,7 +547,7 @@
%321 = OpAccessChain %_ptr_Workgroup_float %mm_Asub %319 %320
%322 = OpLoad %float %321
OpStore %ACached %322
- OpStore %innerCol_1 %31
+ OpStore %innerCol_1 %122
OpBranch %324
%324 = OpLabel
OpLoopMerge %325 %326 None
@@ -604,7 +604,7 @@
OpStore %t %354
OpBranch %198
%199 = OpLabel
- OpStore %innerRow_2 %31
+ OpStore %innerRow_2 %122
OpBranch %356
%356 = OpLabel
OpLoopMerge %357 %358 None
@@ -618,7 +618,7 @@
%364 = OpLabel
OpBranch %357
%363 = OpLabel
- OpStore %innerCol_2 %31
+ OpStore %innerCol_2 %122
OpBranch %366
%366 = OpLabel
OpLoopMerge %367 %368 None
diff --git a/test/tint/bug/tint/942.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/942.wgsl.expected.dxc.hlsl
index e445348..8b20a87 100644
--- a/test/tint/bug/tint/942.wgsl.expected.dxc.hlsl
+++ b/test/tint/bug/tint/942.wgsl.expected.dxc.hlsl
@@ -1,7 +1,3 @@
-uint tint_div(uint lhs, uint rhs) {
- return (lhs / ((rhs == 0u) ? 1u : rhs));
-}
-
SamplerState samp : register(s0, space0);
cbuffer cbuffer_params : register(b1, space0) {
uint4 params[1];
@@ -14,6 +10,10 @@
};
groupshared float3 tile[4][256];
+uint tint_div(uint lhs, uint rhs) {
+ return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
struct tint_symbol_1 {
uint3 LocalInvocationID : SV_GroupThreadID;
uint local_invocation_index : SV_GroupIndex;
diff --git a/test/tint/bug/tint/942.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/942.wgsl.expected.fxc.hlsl
index e445348..8b20a87 100644
--- a/test/tint/bug/tint/942.wgsl.expected.fxc.hlsl
+++ b/test/tint/bug/tint/942.wgsl.expected.fxc.hlsl
@@ -1,7 +1,3 @@
-uint tint_div(uint lhs, uint rhs) {
- return (lhs / ((rhs == 0u) ? 1u : rhs));
-}
-
SamplerState samp : register(s0, space0);
cbuffer cbuffer_params : register(b1, space0) {
uint4 params[1];
@@ -14,6 +10,10 @@
};
groupshared float3 tile[4][256];
+uint tint_div(uint lhs, uint rhs) {
+ return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
struct tint_symbol_1 {
uint3 LocalInvocationID : SV_GroupThreadID;
uint local_invocation_index : SV_GroupIndex;
diff --git a/test/tint/bug/tint/942.wgsl.expected.glsl b/test/tint/bug/tint/942.wgsl.expected.glsl
index 77b463a..17703b1 100644
--- a/test/tint/bug/tint/942.wgsl.expected.glsl
+++ b/test/tint/bug/tint/942.wgsl.expected.glsl
@@ -1,9 +1,5 @@
#version 310 es
-uint tint_div(uint lhs, uint rhs) {
- return (lhs / ((rhs == 0u) ? 1u : rhs));
-}
-
struct Params {
uint filterDim;
uint blockDim;
@@ -28,6 +24,10 @@
} flip;
shared vec3 tile[4][256];
+uint tint_div(uint lhs, uint rhs) {
+ return (lhs / ((rhs == 0u) ? 1u : rhs));
+}
+
uniform highp sampler2D inputTex_1;
uniform highp sampler2D inputTex_samp;
diff --git a/test/tint/bug/tint/942.wgsl.expected.msl b/test/tint/bug/tint/942.wgsl.expected.msl
index 507e1bc..991c4cf 100644
--- a/test/tint/bug/tint/942.wgsl.expected.msl
+++ b/test/tint/bug/tint/942.wgsl.expected.msl
@@ -14,10 +14,6 @@
T elements[N];
};
-uint tint_div(uint lhs, uint rhs) {
- return (lhs / select(rhs, 1u, (rhs == 0u)));
-}
-
struct Params {
/* 0x0000 */ uint filterDim;
/* 0x0004 */ uint blockDim;
@@ -27,6 +23,10 @@
/* 0x0000 */ uint value;
};
+uint tint_div(uint lhs, uint rhs) {
+ return (lhs / select(rhs, 1u, (rhs == 0u)));
+}
+
void tint_symbol_inner(uint3 WorkGroupID, uint3 LocalInvocationID, uint local_invocation_index, threadgroup tint_array<tint_array<float3, 256>, 4>* const tint_symbol_1, const constant Params* const tint_symbol_2, texture2d<float, access::sample> tint_symbol_3, const constant Flip* const tint_symbol_4, sampler tint_symbol_5, texture2d<float, access::write> tint_symbol_6) {
for(uint idx = local_invocation_index; (idx < 1024u); idx = (idx + 64u)) {
uint const i_1 = (idx / 256u);
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.dxc.hlsl b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.dxc.hlsl
index b3bbd36..f3fc898 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.dxc.hlsl
@@ -1,3 +1,6 @@
+static uint local_invocation_index_1 = 0u;
+groupshared uint wg[3][2][1];
+
uint tint_div(uint lhs, uint rhs) {
return (lhs / ((rhs == 0u) ? 1u : rhs));
}
@@ -6,9 +9,6 @@
return (lhs % ((rhs == 0u) ? 1u : rhs));
}
-static uint local_invocation_index_1 = 0u;
-groupshared uint wg[3][2][1];
-
void compute_main_inner(uint local_invocation_index) {
uint idx = 0u;
idx = local_invocation_index;
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.fxc.hlsl b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.fxc.hlsl
index b3bbd36..f3fc898 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.fxc.hlsl
@@ -1,3 +1,6 @@
+static uint local_invocation_index_1 = 0u;
+groupshared uint wg[3][2][1];
+
uint tint_div(uint lhs, uint rhs) {
return (lhs / ((rhs == 0u) ? 1u : rhs));
}
@@ -6,9 +9,6 @@
return (lhs % ((rhs == 0u) ? 1u : rhs));
}
-static uint local_invocation_index_1 = 0u;
-groupshared uint wg[3][2][1];
-
void compute_main_inner(uint local_invocation_index) {
uint idx = 0u;
idx = local_invocation_index;
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.glsl b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.glsl
index d21534a..0912c08 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.glsl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.glsl
@@ -1,5 +1,7 @@
#version 310 es
+uint local_invocation_index_1 = 0u;
+shared uint wg[3][2][1];
uint tint_div(uint lhs, uint rhs) {
return (lhs / ((rhs == 0u) ? 1u : rhs));
}
@@ -8,8 +10,6 @@
return (lhs % ((rhs == 0u) ? 1u : rhs));
}
-uint local_invocation_index_1 = 0u;
-shared uint wg[3][2][1];
void compute_main_inner(uint local_invocation_index) {
uint idx = 0u;
idx = local_invocation_index;
diff --git a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.dxc.hlsl b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.dxc.hlsl
index b3bbd36..f3fc898 100644
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.dxc.hlsl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.dxc.hlsl
@@ -1,3 +1,6 @@
+static uint local_invocation_index_1 = 0u;
+groupshared uint wg[3][2][1];
+
uint tint_div(uint lhs, uint rhs) {
return (lhs / ((rhs == 0u) ? 1u : rhs));
}
@@ -6,9 +9,6 @@
return (lhs % ((rhs == 0u) ? 1u : rhs));
}
-static uint local_invocation_index_1 = 0u;
-groupshared uint wg[3][2][1];
-
void compute_main_inner(uint local_invocation_index) {
uint idx = 0u;
idx = local_invocation_index;
diff --git a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.fxc.hlsl b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.fxc.hlsl
index b3bbd36..f3fc898 100644
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.fxc.hlsl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.fxc.hlsl
@@ -1,3 +1,6 @@
+static uint local_invocation_index_1 = 0u;
+groupshared uint wg[3][2][1];
+
uint tint_div(uint lhs, uint rhs) {
return (lhs / ((rhs == 0u) ? 1u : rhs));
}
@@ -6,9 +9,6 @@
return (lhs % ((rhs == 0u) ? 1u : rhs));
}
-static uint local_invocation_index_1 = 0u;
-groupshared uint wg[3][2][1];
-
void compute_main_inner(uint local_invocation_index) {
uint idx = 0u;
idx = local_invocation_index;
diff --git a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.glsl b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.glsl
index d21534a..0912c08 100644
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.glsl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.glsl
@@ -1,5 +1,7 @@
#version 310 es
+uint local_invocation_index_1 = 0u;
+shared uint wg[3][2][1];
uint tint_div(uint lhs, uint rhs) {
return (lhs / ((rhs == 0u) ? 1u : rhs));
}
@@ -8,8 +10,6 @@
return (lhs % ((rhs == 0u) ? 1u : rhs));
}
-uint local_invocation_index_1 = 0u;
-shared uint wg[3][2][1];
void compute_main_inner(uint local_invocation_index) {
uint idx = 0u;
idx = local_invocation_index;
diff --git a/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.dxc.hlsl b/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.dxc.hlsl
index 35dd311..e1f2af5 100644
--- a/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.dxc.hlsl
+++ b/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.dxc.hlsl
@@ -3,6 +3,9 @@
return;
}
+static int a = 0;
+static float b = 0.0f;
+
int tint_div(int lhs, int rhs) {
return (lhs / (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
}
@@ -11,9 +14,6 @@
return (lhs % (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
}
-static int a = 0;
-static float b = 0.0f;
-
void foo(int maybe_zero) {
a = tint_div(a, 0);
a = tint_mod(a, 0);
diff --git a/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.fxc.hlsl b/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.fxc.hlsl
index 35dd311..e1f2af5 100644
--- a/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.fxc.hlsl
+++ b/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.fxc.hlsl
@@ -3,6 +3,9 @@
return;
}
+static int a = 0;
+static float b = 0.0f;
+
int tint_div(int lhs, int rhs) {
return (lhs / (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
}
@@ -11,9 +14,6 @@
return (lhs % (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
}
-static int a = 0;
-static float b = 0.0f;
-
void foo(int maybe_zero) {
a = tint_div(a, 0);
a = tint_mod(a, 0);
diff --git a/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.glsl b/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.glsl
index 3af799e..6150923 100644
--- a/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.glsl
+++ b/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.glsl
@@ -9,6 +9,8 @@
void unused_entry_point() {
return;
}
+int a = 0;
+float b = 0.0f;
int tint_div(int lhs, int rhs) {
return (lhs / (bool(uint((rhs == 0)) | uint(bool(uint((lhs == -2147483648)) & uint((rhs == -1))))) ? 1 : rhs));
}
@@ -17,8 +19,6 @@
return (lhs % (bool(uint((rhs == 0)) | uint(bool(uint((lhs == -2147483648)) & uint((rhs == -1))))) ? 1 : rhs));
}
-int a = 0;
-float b = 0.0f;
void foo(int maybe_zero) {
a = tint_div(a, 0);
a = tint_mod(a, 0);
diff --git a/test/tint/statements/compound_assign/private.wgsl.expected.dxc.hlsl b/test/tint/statements/compound_assign/private.wgsl.expected.dxc.hlsl
index 2de2195..4efd805 100644
--- a/test/tint/statements/compound_assign/private.wgsl.expected.dxc.hlsl
+++ b/test/tint/statements/compound_assign/private.wgsl.expected.dxc.hlsl
@@ -3,14 +3,14 @@
return;
}
-int tint_div(int lhs, int rhs) {
- return (lhs / (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
-}
-
static int a = 0;
static float4 b = float4(0.0f, 0.0f, 0.0f, 0.0f);
static float2x2 c = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+int tint_div(int lhs, int rhs) {
+ return (lhs / (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
+}
+
void foo() {
a = tint_div(a, 2);
b = mul(float4x4((0.0f).xxxx, (0.0f).xxxx, (0.0f).xxxx, (0.0f).xxxx), b);
diff --git a/test/tint/statements/compound_assign/private.wgsl.expected.fxc.hlsl b/test/tint/statements/compound_assign/private.wgsl.expected.fxc.hlsl
index 2de2195..4efd805 100644
--- a/test/tint/statements/compound_assign/private.wgsl.expected.fxc.hlsl
+++ b/test/tint/statements/compound_assign/private.wgsl.expected.fxc.hlsl
@@ -3,14 +3,14 @@
return;
}
-int tint_div(int lhs, int rhs) {
- return (lhs / (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
-}
-
static int a = 0;
static float4 b = float4(0.0f, 0.0f, 0.0f, 0.0f);
static float2x2 c = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+int tint_div(int lhs, int rhs) {
+ return (lhs / (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
+}
+
void foo() {
a = tint_div(a, 2);
b = mul(float4x4((0.0f).xxxx, (0.0f).xxxx, (0.0f).xxxx, (0.0f).xxxx), b);
diff --git a/test/tint/statements/compound_assign/private.wgsl.expected.glsl b/test/tint/statements/compound_assign/private.wgsl.expected.glsl
index 4efe337..bbfe95e 100644
--- a/test/tint/statements/compound_assign/private.wgsl.expected.glsl
+++ b/test/tint/statements/compound_assign/private.wgsl.expected.glsl
@@ -4,13 +4,13 @@
void unused_entry_point() {
return;
}
+int a = 0;
+vec4 b = vec4(0.0f, 0.0f, 0.0f, 0.0f);
+mat2 c = mat2(0.0f, 0.0f, 0.0f, 0.0f);
int tint_div(int lhs, int rhs) {
return (lhs / (bool(uint((rhs == 0)) | uint(bool(uint((lhs == -2147483648)) & uint((rhs == -1))))) ? 1 : rhs));
}
-int a = 0;
-vec4 b = vec4(0.0f, 0.0f, 0.0f, 0.0f);
-mat2 c = mat2(0.0f, 0.0f, 0.0f, 0.0f);
void foo() {
a = tint_div(a, 2);
b = (b * mat4(vec4(0.0f), vec4(0.0f), vec4(0.0f), vec4(0.0f)));
diff --git a/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.dxc.hlsl b/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.dxc.hlsl
index e84dd9a..9616780 100644
--- a/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.dxc.hlsl
+++ b/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.dxc.hlsl
@@ -3,12 +3,12 @@
return;
}
+RWByteAddressBuffer v : register(u0, space0);
+
int tint_div(int lhs, int rhs) {
return (lhs / (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
}
-RWByteAddressBuffer v : register(u0, space0);
-
void foo() {
const int tint_symbol = tint_div(asint(v.Load(0u)), 2);
v.Store(0u, asuint(tint_symbol));
diff --git a/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.fxc.hlsl b/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.fxc.hlsl
index e84dd9a..9616780 100644
--- a/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.fxc.hlsl
+++ b/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.fxc.hlsl
@@ -3,12 +3,12 @@
return;
}
+RWByteAddressBuffer v : register(u0, space0);
+
int tint_div(int lhs, int rhs) {
return (lhs / (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
}
-RWByteAddressBuffer v : register(u0, space0);
-
void foo() {
const int tint_symbol = tint_div(asint(v.Load(0u)), 2);
v.Store(0u, asuint(tint_symbol));
diff --git a/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.glsl b/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.glsl
index 43823fc..822f89b 100644
--- a/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.glsl
+++ b/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.glsl
@@ -4,10 +4,6 @@
void unused_entry_point() {
return;
}
-int tint_div(int lhs, int rhs) {
- return (lhs / (bool(uint((rhs == 0)) | uint(bool(uint((lhs == -2147483648)) & uint((rhs == -1))))) ? 1 : rhs));
-}
-
struct S {
int a;
};
@@ -16,6 +12,10 @@
S inner;
} v;
+int tint_div(int lhs, int rhs) {
+ return (lhs / (bool(uint((rhs == 0)) | uint(bool(uint((lhs == -2147483648)) & uint((rhs == -1))))) ? 1 : rhs));
+}
+
void foo() {
int tint_symbol = tint_div(v.inner.a, 2);
v.inner.a = tint_symbol;
diff --git a/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.msl b/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.msl
index 855cb61..c839ac0 100644
--- a/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/scalar/divide.wgsl.expected.msl
@@ -1,14 +1,14 @@
#include <metal_stdlib>
using namespace metal;
-int tint_div(int lhs, int rhs) {
- return (lhs / select(rhs, 1, bool((rhs == 0) | bool((lhs == (-2147483647 - 1)) & (rhs == -1)))));
-}
-
struct S {
/* 0x0000 */ int a;
};
+int tint_div(int lhs, int rhs) {
+ return (lhs / select(rhs, 1, bool((rhs == 0) | bool((lhs == (-2147483647 - 1)) & (rhs == -1)))));
+}
+
void foo(device S* const tint_symbol_1) {
int const tint_symbol = tint_div((*(tint_symbol_1)).a, 2);
(*(tint_symbol_1)).a = tint_symbol;
diff --git a/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.dxc.hlsl b/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.dxc.hlsl
index 92912b2..f7bac0a 100644
--- a/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.dxc.hlsl
+++ b/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.dxc.hlsl
@@ -3,12 +3,12 @@
return;
}
+RWByteAddressBuffer v : register(u0, space0);
+
int tint_mod(int lhs, int rhs) {
return (lhs % (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
}
-RWByteAddressBuffer v : register(u0, space0);
-
void foo() {
const int tint_symbol = tint_mod(asint(v.Load(0u)), 2);
v.Store(0u, asuint(tint_symbol));
diff --git a/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.fxc.hlsl b/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.fxc.hlsl
index 92912b2..f7bac0a 100644
--- a/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.fxc.hlsl
+++ b/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.fxc.hlsl
@@ -3,12 +3,12 @@
return;
}
+RWByteAddressBuffer v : register(u0, space0);
+
int tint_mod(int lhs, int rhs) {
return (lhs % (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
}
-RWByteAddressBuffer v : register(u0, space0);
-
void foo() {
const int tint_symbol = tint_mod(asint(v.Load(0u)), 2);
v.Store(0u, asuint(tint_symbol));
diff --git a/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.glsl b/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.glsl
index f3cf594..f84aaf6 100644
--- a/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.glsl
+++ b/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.glsl
@@ -4,10 +4,6 @@
void unused_entry_point() {
return;
}
-int tint_mod(int lhs, int rhs) {
- return (lhs % (bool(uint((rhs == 0)) | uint(bool(uint((lhs == -2147483648)) & uint((rhs == -1))))) ? 1 : rhs));
-}
-
struct S {
int a;
};
@@ -16,6 +12,10 @@
S inner;
} v;
+int tint_mod(int lhs, int rhs) {
+ return (lhs % (bool(uint((rhs == 0)) | uint(bool(uint((lhs == -2147483648)) & uint((rhs == -1))))) ? 1 : rhs));
+}
+
void foo() {
int tint_symbol = tint_mod(v.inner.a, 2);
v.inner.a = tint_symbol;
diff --git a/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.msl b/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.msl
index 978e630..e8cc90c 100644
--- a/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.msl
@@ -1,14 +1,14 @@
#include <metal_stdlib>
using namespace metal;
-int tint_mod(int lhs, int rhs) {
- return (lhs % select(rhs, 1, bool((rhs == 0) | bool((lhs == (-2147483647 - 1)) & (rhs == -1)))));
-}
-
struct S {
/* 0x0000 */ int a;
};
+int tint_mod(int lhs, int rhs) {
+ return (lhs % select(rhs, 1, bool((rhs == 0) | bool((lhs == (-2147483647 - 1)) & (rhs == -1)))));
+}
+
void foo(device S* const tint_symbol_1) {
int const tint_symbol = tint_mod((*(tint_symbol_1)).a, 2);
(*(tint_symbol_1)).a = tint_symbol;
diff --git a/test/tint/statements/compound_assign/vector/divide.wgsl.expected.dxc.hlsl b/test/tint/statements/compound_assign/vector/divide.wgsl.expected.dxc.hlsl
index d519f35..c2ffedc 100644
--- a/test/tint/statements/compound_assign/vector/divide.wgsl.expected.dxc.hlsl
+++ b/test/tint/statements/compound_assign/vector/divide.wgsl.expected.dxc.hlsl
@@ -3,12 +3,12 @@
return;
}
+RWByteAddressBuffer v : register(u0, space0);
+
int4 tint_div(int4 lhs, int4 rhs) {
return (lhs / (((rhs == (0).xxxx) | ((lhs == (-2147483648).xxxx) & (rhs == (-1).xxxx))) ? (1).xxxx : rhs));
}
-RWByteAddressBuffer v : register(u0, space0);
-
void foo() {
const int4 tint_symbol = tint_div(asint(v.Load4(0u)), (2).xxxx);
v.Store4(0u, asuint(tint_symbol));
diff --git a/test/tint/statements/compound_assign/vector/divide.wgsl.expected.fxc.hlsl b/test/tint/statements/compound_assign/vector/divide.wgsl.expected.fxc.hlsl
index d519f35..c2ffedc 100644
--- a/test/tint/statements/compound_assign/vector/divide.wgsl.expected.fxc.hlsl
+++ b/test/tint/statements/compound_assign/vector/divide.wgsl.expected.fxc.hlsl
@@ -3,12 +3,12 @@
return;
}
+RWByteAddressBuffer v : register(u0, space0);
+
int4 tint_div(int4 lhs, int4 rhs) {
return (lhs / (((rhs == (0).xxxx) | ((lhs == (-2147483648).xxxx) & (rhs == (-1).xxxx))) ? (1).xxxx : rhs));
}
-RWByteAddressBuffer v : register(u0, space0);
-
void foo() {
const int4 tint_symbol = tint_div(asint(v.Load4(0u)), (2).xxxx);
v.Store4(0u, asuint(tint_symbol));
diff --git a/test/tint/statements/compound_assign/vector/divide.wgsl.expected.glsl b/test/tint/statements/compound_assign/vector/divide.wgsl.expected.glsl
index aa704fe..daea9b2 100644
--- a/test/tint/statements/compound_assign/vector/divide.wgsl.expected.glsl
+++ b/test/tint/statements/compound_assign/vector/divide.wgsl.expected.glsl
@@ -4,10 +4,6 @@
void unused_entry_point() {
return;
}
-ivec4 tint_div(ivec4 lhs, ivec4 rhs) {
- return (lhs / mix(rhs, ivec4(1), bvec4(uvec4(equal(rhs, ivec4(0))) | uvec4(bvec4(uvec4(equal(lhs, ivec4(-2147483648))) & uvec4(equal(rhs, ivec4(-1))))))));
-}
-
struct S {
ivec4 a;
};
@@ -16,6 +12,10 @@
S inner;
} v;
+ivec4 tint_div(ivec4 lhs, ivec4 rhs) {
+ return (lhs / mix(rhs, ivec4(1), bvec4(uvec4(equal(rhs, ivec4(0))) | uvec4(bvec4(uvec4(equal(lhs, ivec4(-2147483648))) & uvec4(equal(rhs, ivec4(-1))))))));
+}
+
void foo() {
ivec4 tint_symbol = tint_div(v.inner.a, ivec4(2));
v.inner.a = tint_symbol;
diff --git a/test/tint/statements/compound_assign/vector/divide.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/divide.wgsl.expected.msl
index 782112c..f0dd095 100644
--- a/test/tint/statements/compound_assign/vector/divide.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/vector/divide.wgsl.expected.msl
@@ -1,14 +1,14 @@
#include <metal_stdlib>
using namespace metal;
-int4 tint_div(int4 lhs, int4 rhs) {
- return (lhs / select(rhs, int4(1), ((rhs == int4(0)) | ((lhs == int4((-2147483647 - 1))) & (rhs == int4(-1))))));
-}
-
struct S {
/* 0x0000 */ int4 a;
};
+int4 tint_div(int4 lhs, int4 rhs) {
+ return (lhs / select(rhs, int4(1), ((rhs == int4(0)) | ((lhs == int4((-2147483647 - 1))) & (rhs == int4(-1))))));
+}
+
void foo(device S* const tint_symbol_1) {
int4 const tint_symbol = tint_div((*(tint_symbol_1)).a, int4(2));
(*(tint_symbol_1)).a = tint_symbol;
diff --git a/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.dxc.hlsl b/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.dxc.hlsl
index e0195aa..a6095f4 100644
--- a/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.dxc.hlsl
+++ b/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.dxc.hlsl
@@ -3,13 +3,13 @@
return;
}
+RWByteAddressBuffer v : register(u0, space0);
+
int4 tint_mod(int4 lhs, int rhs) {
const int4 r = int4((rhs).xxxx);
return (lhs % (((r == (0).xxxx) | ((lhs == (-2147483648).xxxx) & (r == (-1).xxxx))) ? (1).xxxx : r));
}
-RWByteAddressBuffer v : register(u0, space0);
-
void foo() {
const int4 tint_symbol = tint_mod(asint(v.Load4(0u)), 2);
v.Store4(0u, asuint(tint_symbol));
diff --git a/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.fxc.hlsl b/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.fxc.hlsl
index e0195aa..a6095f4 100644
--- a/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.fxc.hlsl
+++ b/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.fxc.hlsl
@@ -3,13 +3,13 @@
return;
}
+RWByteAddressBuffer v : register(u0, space0);
+
int4 tint_mod(int4 lhs, int rhs) {
const int4 r = int4((rhs).xxxx);
return (lhs % (((r == (0).xxxx) | ((lhs == (-2147483648).xxxx) & (r == (-1).xxxx))) ? (1).xxxx : r));
}
-RWByteAddressBuffer v : register(u0, space0);
-
void foo() {
const int4 tint_symbol = tint_mod(asint(v.Load4(0u)), 2);
v.Store4(0u, asuint(tint_symbol));
diff --git a/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.glsl b/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.glsl
index ed2d419..5609858 100644
--- a/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.glsl
+++ b/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.glsl
@@ -4,11 +4,6 @@
void unused_entry_point() {
return;
}
-ivec4 tint_mod(ivec4 lhs, int rhs) {
- ivec4 r = ivec4(rhs);
- return (lhs % mix(r, ivec4(1), bvec4(uvec4(equal(r, ivec4(0))) | uvec4(bvec4(uvec4(equal(lhs, ivec4(-2147483648))) & uvec4(equal(r, ivec4(-1))))))));
-}
-
struct S {
ivec4 a;
};
@@ -17,6 +12,11 @@
S inner;
} v;
+ivec4 tint_mod(ivec4 lhs, int rhs) {
+ ivec4 r = ivec4(rhs);
+ return (lhs % mix(r, ivec4(1), bvec4(uvec4(equal(r, ivec4(0))) | uvec4(bvec4(uvec4(equal(lhs, ivec4(-2147483648))) & uvec4(equal(r, ivec4(-1))))))));
+}
+
void foo() {
ivec4 tint_symbol = tint_mod(v.inner.a, 2);
v.inner.a = tint_symbol;
diff --git a/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.msl
index db65802..7a065fd 100644
--- a/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.msl
@@ -1,15 +1,15 @@
#include <metal_stdlib>
using namespace metal;
+struct S {
+ /* 0x0000 */ int4 a;
+};
+
int4 tint_mod(int4 lhs, int rhs) {
int4 const r = int4(rhs);
return (lhs % select(r, int4(1), ((r == int4(0)) | ((lhs == int4((-2147483647 - 1))) & (r == int4(-1))))));
}
-struct S {
- /* 0x0000 */ int4 a;
-};
-
void foo(device S* const tint_symbol_1) {
int4 const tint_symbol = tint_mod((*(tint_symbol_1)).a, 2);
(*(tint_symbol_1)).a = tint_symbol;
diff --git a/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.dxc.hlsl b/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.dxc.hlsl
index f413ff6..8a5eb1a 100644
--- a/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.dxc.hlsl
+++ b/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.dxc.hlsl
@@ -3,12 +3,12 @@
return;
}
+RWByteAddressBuffer v : register(u0, space0);
+
int4 tint_mod(int4 lhs, int4 rhs) {
return (lhs % (((rhs == (0).xxxx) | ((lhs == (-2147483648).xxxx) & (rhs == (-1).xxxx))) ? (1).xxxx : rhs));
}
-RWByteAddressBuffer v : register(u0, space0);
-
void foo() {
const int4 tint_symbol = tint_mod(asint(v.Load4(0u)), (2).xxxx);
v.Store4(0u, asuint(tint_symbol));
diff --git a/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.fxc.hlsl b/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.fxc.hlsl
index f413ff6..8a5eb1a 100644
--- a/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.fxc.hlsl
+++ b/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.fxc.hlsl
@@ -3,12 +3,12 @@
return;
}
+RWByteAddressBuffer v : register(u0, space0);
+
int4 tint_mod(int4 lhs, int4 rhs) {
return (lhs % (((rhs == (0).xxxx) | ((lhs == (-2147483648).xxxx) & (rhs == (-1).xxxx))) ? (1).xxxx : rhs));
}
-RWByteAddressBuffer v : register(u0, space0);
-
void foo() {
const int4 tint_symbol = tint_mod(asint(v.Load4(0u)), (2).xxxx);
v.Store4(0u, asuint(tint_symbol));
diff --git a/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.glsl b/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.glsl
index 1f5c0c4..192c38d 100644
--- a/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.glsl
+++ b/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.glsl
@@ -4,10 +4,6 @@
void unused_entry_point() {
return;
}
-ivec4 tint_mod(ivec4 lhs, ivec4 rhs) {
- return (lhs % mix(rhs, ivec4(1), bvec4(uvec4(equal(rhs, ivec4(0))) | uvec4(bvec4(uvec4(equal(lhs, ivec4(-2147483648))) & uvec4(equal(rhs, ivec4(-1))))))));
-}
-
struct S {
ivec4 a;
};
@@ -16,6 +12,10 @@
S inner;
} v;
+ivec4 tint_mod(ivec4 lhs, ivec4 rhs) {
+ return (lhs % mix(rhs, ivec4(1), bvec4(uvec4(equal(rhs, ivec4(0))) | uvec4(bvec4(uvec4(equal(lhs, ivec4(-2147483648))) & uvec4(equal(rhs, ivec4(-1))))))));
+}
+
void foo() {
ivec4 tint_symbol = tint_mod(v.inner.a, ivec4(2));
v.inner.a = tint_symbol;
diff --git a/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.msl
index 2435db2..cb90ef8 100644
--- a/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.msl
@@ -1,14 +1,14 @@
#include <metal_stdlib>
using namespace metal;
-int4 tint_mod(int4 lhs, int4 rhs) {
- return (lhs % select(rhs, int4(1), ((rhs == int4(0)) | ((lhs == int4((-2147483647 - 1))) & (rhs == int4(-1))))));
-}
-
struct S {
/* 0x0000 */ int4 a;
};
+int4 tint_mod(int4 lhs, int4 rhs) {
+ return (lhs % select(rhs, int4(1), ((rhs == int4(0)) | ((lhs == int4((-2147483647 - 1))) & (rhs == int4(-1))))));
+}
+
void foo(device S* const tint_symbol_1) {
int4 const tint_symbol = tint_mod((*(tint_symbol_1)).a, int4(2));
(*(tint_symbol_1)).a = tint_symbol;
diff --git a/test/tint/statements/compound_assign/workgroup.wgsl.expected.dxc.hlsl b/test/tint/statements/compound_assign/workgroup.wgsl.expected.dxc.hlsl
index d6246e4..85be9f7 100644
--- a/test/tint/statements/compound_assign/workgroup.wgsl.expected.dxc.hlsl
+++ b/test/tint/statements/compound_assign/workgroup.wgsl.expected.dxc.hlsl
@@ -3,14 +3,14 @@
return;
}
-int tint_div(int lhs, int rhs) {
- return (lhs / (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
-}
-
groupshared int a;
groupshared float4 b;
groupshared float2x2 c;
+int tint_div(int lhs, int rhs) {
+ return (lhs / (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
+}
+
void foo() {
a = tint_div(a, 2);
b = mul(float4x4((0.0f).xxxx, (0.0f).xxxx, (0.0f).xxxx, (0.0f).xxxx), b);
diff --git a/test/tint/statements/compound_assign/workgroup.wgsl.expected.fxc.hlsl b/test/tint/statements/compound_assign/workgroup.wgsl.expected.fxc.hlsl
index d6246e4..85be9f7 100644
--- a/test/tint/statements/compound_assign/workgroup.wgsl.expected.fxc.hlsl
+++ b/test/tint/statements/compound_assign/workgroup.wgsl.expected.fxc.hlsl
@@ -3,14 +3,14 @@
return;
}
-int tint_div(int lhs, int rhs) {
- return (lhs / (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
-}
-
groupshared int a;
groupshared float4 b;
groupshared float2x2 c;
+int tint_div(int lhs, int rhs) {
+ return (lhs / (((rhs == 0) | ((lhs == -2147483648) & (rhs == -1))) ? 1 : rhs));
+}
+
void foo() {
a = tint_div(a, 2);
b = mul(float4x4((0.0f).xxxx, (0.0f).xxxx, (0.0f).xxxx, (0.0f).xxxx), b);
diff --git a/test/tint/statements/compound_assign/workgroup.wgsl.expected.glsl b/test/tint/statements/compound_assign/workgroup.wgsl.expected.glsl
index 3b920a2..797f37d 100644
--- a/test/tint/statements/compound_assign/workgroup.wgsl.expected.glsl
+++ b/test/tint/statements/compound_assign/workgroup.wgsl.expected.glsl
@@ -4,13 +4,13 @@
void unused_entry_point() {
return;
}
+shared int a;
+shared vec4 b;
+shared mat2 c;
int tint_div(int lhs, int rhs) {
return (lhs / (bool(uint((rhs == 0)) | uint(bool(uint((lhs == -2147483648)) & uint((rhs == -1))))) ? 1 : rhs));
}
-shared int a;
-shared vec4 b;
-shared mat2 c;
void foo() {
a = tint_div(a, 2);
b = (b * mat4(vec4(0.0f), vec4(0.0f), vec4(0.0f), vec4(0.0f)));
