diff --git a/src/tint/transform/builtin_polyfill.cc b/src/tint/transform/builtin_polyfill.cc index 4dfc1495f0..15aa233069 100644 --- a/src/tint/transform/builtin_polyfill.cc +++ b/src/tint/transform/builtin_polyfill.cc @@ -41,6 +41,10 @@ struct BuiltinPolyfill::State { /// @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 @@ struct BuiltinPolyfill::State { 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 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 @@ struct BuiltinPolyfill::State { 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()) { + 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(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 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 @@ struct BuiltinPolyfill::State { /// The source clone context const sem::Info& sem = ctx.src->Sem(); + // Polyfill functions for binary operators. + utils::Hashmap 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 @@ Transform::ApplyResult BuiltinPolyfill::Apply(const Program* src, auto& polyfill = cfg->builtins; utils::Hashmap builtin_polyfills; - utils::Hashmap binary_op_polyfills; ProgramBuilder b; CloneContext ctx{&b, src, /* auto_clone_symbols */ true}; @@ -849,15 +881,7 @@ Transform::ApplyResult BuiltinPolyfill::Apply(const Program* src, 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()) { - 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 @@ Transform::ApplyResult BuiltinPolyfill::Apply(const Program* src, 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 fbbd90e4c5..87fd9d5cf7 100644 --- a/src/tint/transform/builtin_polyfill_test.cc +++ b/src/tint/transform/builtin_polyfill_test.cc @@ -3000,5 +3000,37 @@ fn f() { 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(builtins); + + auto got = Run(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 7f89b55301..8b61a3e3c9 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 @@ uint toIndex1D(uint gridSize, float3 voxelPos) { 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 7f89b55301..8b61a3e3c9 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 @@ uint toIndex1D(uint gridSize, float3 voxelPos) { 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 15000e6755..aa346060bb 100644 --- a/test/tint/bug/chromium/1273230.wgsl.expected.msl +++ b/test/tint/bug/chromium/1273230.wgsl.expected.msl @@ -14,14 +14,6 @@ struct tint_array { 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 @@ uint toIndex1D(uint gridSize, float3 voxelPos) { 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 2e6650348d..96d229ceb1 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 0000000000..b0745fe437 --- /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) { + 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 0000000000..424a72f575 --- /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 0000000000..424a72f575 --- /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 0000000000..ab0cb6e517 --- /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 0000000000..694a88a216 --- /dev/null +++ b/test/tint/bug/chromium/1386647.wgsl.expected.msl @@ -0,0 +1,18 @@ +#include + +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 0000000000..c54e606b7c --- /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 0000000000..b8c5bd328d --- /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) { + 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 6a0326e702..68248ca64c 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 @@ uint toIndex1D(uint gridSize, float3 voxelPos) { 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 6a0326e702..68248ca64c 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 @@ uint toIndex1D(uint gridSize, float3 voxelPos) { 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 0ab0884b2f..26605009bf 100644 --- a/test/tint/bug/tint/1113.wgsl.expected.msl +++ b/test/tint/bug/tint/1113.wgsl.expected.msl @@ -14,14 +14,6 @@ struct tint_array { 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 @@ uint toIndex1D(uint gridSize, float3 voxelPos) { 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 f1eee0ccab..9544d463c5 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 e07e2fa318..06255d05cb 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 float4 sk_FragColor = float4(0.0f, 0.0f, 0.0f, 0.0f); 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 e07e2fa318..06255d05cb 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 float4 sk_FragColor = float4(0.0f, 0.0f, 0.0f, 0.0f); 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 fbf2852b99..8cd026fe7c 100644 --- a/test/tint/bug/tint/1520.spvasm.expected.glsl +++ b/test/tint/bug/tint/1520.spvasm.expected.glsl @@ -3,10 +3,6 @@ precision mediump float; 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 @@ layout(binding = 0, std140) uniform x_4_block_ubo { 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 b997052bfe..335a96b1af 100644 --- a/test/tint/bug/tint/1520.spvasm.expected.msl +++ b/test/tint/bug/tint/1520.spvasm.expected.msl @@ -14,10 +14,6 @@ struct tint_array { 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 tint_pad; /* 0x0010 */ float unknownInput_S1_c0; @@ -27,6 +23,10 @@ struct UniformBuffer { /* 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 708904c0ea..40385661d0 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 708904c0ea..40385661d0 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 5f6890fed9..4bcee5e013 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 @@ layout(binding = 0, std430) buffer b_block_ssbo { 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 7e30a1864e..8a82c4348c 100644 --- a/test/tint/bug/tint/221.wgsl.expected.msl +++ b/test/tint/bug/tint/221.wgsl.expected.msl @@ -14,15 +14,15 @@ struct tint_array { 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 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 a893add710..6b741d66d6 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 @@ void mm_write(uint row, uint col, float value) { 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 a893add710..6b741d66d6 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 @@ void mm_write(uint row, uint col, float value) { 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 2797131645..fe3d6bcce6 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 @@ void mm_write(uint row, uint col, float value) { 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 b7dec3396e..5ea78f26fd 100644 --- a/test/tint/bug/tint/914.wgsl.expected.msl +++ b/test/tint/bug/tint/914.wgsl.expected.msl @@ -14,10 +14,6 @@ struct tint_array { 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 @@ void mm_write(uint row, uint col, float value, const constant Uniforms* const ti } } +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, 64>* const tint_symbol_10, threadgroup tint_array, 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 f7719e037f..829ad7d69f 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 e445348236..8b20a87eb4 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 @@ cbuffer cbuffer_flip : register(b3, space1) { }; 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 e445348236..8b20a87eb4 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 @@ cbuffer cbuffer_flip : register(b3, space1) { }; 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 77b463ae82..17703b1aa6 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 @@ layout(binding = 3, std140) uniform flip_block_ubo { } 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 507e1bcefa..991c4cf8b2 100644 --- a/test/tint/bug/tint/942.wgsl.expected.msl +++ b/test/tint/bug/tint/942.wgsl.expected.msl @@ -14,10 +14,6 @@ struct tint_array { 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 @@ struct Flip { /* 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, 4>* const tint_symbol_1, const constant Params* const tint_symbol_2, texture2d tint_symbol_3, const constant Flip* const tint_symbol_4, sampler tint_symbol_5, texture2d 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 b3bbd36e50..f3fc89828a 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 @@ uint tint_mod(uint lhs, uint rhs) { 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 b3bbd36e50..f3fc89828a 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 @@ uint tint_mod(uint lhs, uint rhs) { 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 d21534a2b5..0912c08785 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 @@ uint tint_mod(uint lhs, uint rhs) { 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 b3bbd36e50..f3fc89828a 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 @@ uint tint_mod(uint lhs, uint rhs) { 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 b3bbd36e50..f3fc89828a 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 @@ uint tint_mod(uint lhs, uint rhs) { 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 d21534a2b5..0912c08785 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 @@ uint tint_mod(uint lhs, uint rhs) { 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 35dd3113ef..e1f2af52c9 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 @@ void unused_entry_point() { 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 @@ int tint_mod(int lhs, int rhs) { 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 35dd3113ef..e1f2af52c9 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 @@ void unused_entry_point() { 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 @@ int tint_mod(int lhs, int rhs) { 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 3af799e954..61509234e0 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 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; 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 @@ int tint_mod(int lhs, int rhs) { 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 2de2195033..4efd805e06 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 @@ void unused_entry_point() { 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 2de2195033..4efd805e06 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 @@ void unused_entry_point() { 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 4efe337dd4..bbfe95eb20 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 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; 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 e84dd9abe5..961678047d 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 @@ void unused_entry_point() { 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 e84dd9abe5..961678047d 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 @@ void unused_entry_point() { 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 43823fcc6c..822f89b0b9 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 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; 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 @@ layout(binding = 0, std430) buffer v_block_ssbo { 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 855cb61e70..c839ac0133 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 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 92912b2e72..f7bac0aedd 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 @@ void unused_entry_point() { 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 92912b2e72..f7bac0aedd 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 @@ void unused_entry_point() { 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 f3cf594097..f84aaf67a9 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 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; 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 @@ layout(binding = 0, std430) buffer v_block_ssbo { 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 978e63014d..e8cc90c3c4 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 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 d519f35f4e..c2ffedcfcd 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 @@ void unused_entry_point() { 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 d519f35f4e..c2ffedcfcd 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 @@ void unused_entry_point() { 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 aa704fec0b..daea9b2647 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 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; 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 @@ layout(binding = 0, std430) buffer v_block_ssbo { 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 782112cf4c..f0dd095353 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 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 e0195aae3b..a6095f4453 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 @@ void unused_entry_point() { 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 e0195aae3b..a6095f4453 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 @@ void unused_entry_point() { 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 ed2d419ad1..5609858b2b 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 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; 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 @@ layout(binding = 0, std430) buffer v_block_ssbo { 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 db6580245d..7a065fd5cf 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 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 f413ff635f..8a5eb1a9c5 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 @@ void unused_entry_point() { 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 f413ff635f..8a5eb1a9c5 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 @@ void unused_entry_point() { 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 1f5c0c42f0..192c38d11a 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 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; 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 @@ layout(binding = 0, std430) buffer v_block_ssbo { 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 2435db2633..cb90ef8c19 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 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 d6246e401b..85be9f74d7 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 @@ void unused_entry_point() { 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 d6246e401b..85be9f74d7 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 @@ void unused_entry_point() { 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 3b920a20d8..797f37d03b 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 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; 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)));