encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

tint/transform: Fix ICE when combining polyfills 

There's a reason the overload of `ctx.Replace()` that takes a pointer to the replacement is deprecated - it doesn't play well when used as part of another replacement.
Switch to using the callback overload of Replace() to fix bad transform output.

Bug: tint:1386647
Change-Id: I94292eeb65d24d7b2446b16b8b4ad13bdd27965a
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/111000
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
Ben Clayton 2022-11-21 17:11:05 +00:00 committed by Dawn LUCI CQ
parent f5eec817de
commit 619f9bd639
69 changed files with 941 additions and 756 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

166
src/tint/transform/builtin_polyfill.cc
View File

@ -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<const ast::Statement*, 4> body;
if (lhs_width < width) {
// lhs is scalar, rhs is vector. Convert lhs to vector.
body.Push(b.Decl(b.Let("l", b.vec(T(lhs_el_ty), width, b.Expr(lhs)))));
lhs = "l";
}
if (rhs_width < width) {
// lhs is vector, rhs is scalar. Convert rhs to vector.
body.Push(b.Decl(b.Let("r", b.vec(T(rhs_el_ty), width, b.Expr(rhs)))));
rhs = "r";
}
auto name = b.Symbols().New(is_div ? "tint_div" : "tint_mod");
auto* use_one = b.Equal(rhs, ScalarOrVector(width, 0_a));
if (lhs_ty->is_signed_scalar_or_vector()) {
const auto bits = lhs_el_ty->Size() * 8;
auto min_int = AInt(AInt::kLowestValue >> (AInt::kNumBits - bits));
const ast::Expression* lhs_is_min = b.Equal(lhs, ScalarOrVector(width, min_int));
const ast::Expression* rhs_is_minus_one = b.Equal(rhs, ScalarOrVector(width, -1_a));
// use_one = use_one | ((lhs == MIN_INT) & (rhs == -1))
use_one = b.Or(use_one, b.And(lhs_is_min, rhs_is_minus_one));
}
auto* select = b.Call("select", rhs, ScalarOrVector(width, 1_a), use_one);
body.Push(b.Return(is_div ? b.Div(lhs, select) : b.Mod(lhs, select)));
b.Func(name,
utils::Vector{
b.Param("lhs", T(lhs_ty)),
b.Param("rhs", T(rhs_ty)),
},
width == 1 ? T(lhs_ty) : b.ty.vec(T(lhs_el_ty), width), // return type
std::move(body));
return name;
}
/// Builds the polyfill function for the `saturate` builtin
/// @param ty the parameter and return type for the function
/// @return the polyfill function name
@ -677,6 +624,89 @@ 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<sem::Vector>()) {
mask = b.Construct(CreateASTTypeFor(ctx, rhs_ty), mask);
}
auto* lhs = ctx.Clone(bin_op->lhs);
auto* rhs = b.And(ctx.Clone(bin_op->rhs), mask);
return b.create<ast::BinaryExpression>(ctx.Clone(bin_op->source), bin_op->op, lhs, rhs);
}
/// Builds the polyfill inline expression for a integer divide or modulo, preventing DBZs and
/// integer overflows.
/// @param bin_op the original BinaryExpression
/// @return the polyfill divide or modulo
const ast::Expression* IntDivMod(const ast::BinaryExpression* bin_op) {
auto* lhs_ty = ctx.src->TypeOf(bin_op->lhs)->UnwrapRef();
auto* rhs_ty = ctx.src->TypeOf(bin_op->rhs)->UnwrapRef();
BinaryOpSignature sig{bin_op->op, lhs_ty, rhs_ty};
auto fn = binary_op_polyfills.GetOrCreate(sig, [&] {
const bool is_div = bin_op->op == ast::BinaryOp::kDivide;
uint32_t lhs_width = 1;
uint32_t rhs_width = 1;
const auto* lhs_el_ty = sem::Type::ElementOf(lhs_ty, &lhs_width);
const auto* rhs_el_ty = sem::Type::ElementOf(rhs_ty, &rhs_width);
const uint32_t width = std::max(lhs_width, rhs_width);
const char* lhs = "lhs";
const char* rhs = "rhs";
utils::Vector<const ast::Statement*, 4> body;
if (lhs_width < width) {
// lhs is scalar, rhs is vector. Convert lhs to vector.
body.Push(b.Decl(b.Let("l", b.vec(T(lhs_el_ty), width, b.Expr(lhs)))));
lhs = "l";
}
if (rhs_width < width) {
// lhs is vector, rhs is scalar. Convert rhs to vector.
body.Push(b.Decl(b.Let("r", b.vec(T(rhs_el_ty), width, b.Expr(rhs)))));
rhs = "r";
}
auto name = b.Symbols().New(is_div ? "tint_div" : "tint_mod");
auto* use_one = b.Equal(rhs, ScalarOrVector(width, 0_a));
if (lhs_ty->is_signed_scalar_or_vector()) {
const auto bits = lhs_el_ty->Size() * 8;
auto min_int = AInt(AInt::kLowestValue >> (AInt::kNumBits - bits));
const ast::Expression* lhs_is_min = b.Equal(lhs, ScalarOrVector(width, min_int));
const ast::Expression* rhs_is_minus_one = b.Equal(rhs, ScalarOrVector(width, -1_a));
// use_one = use_one | ((lhs == MIN_INT) & (rhs == -1))
use_one = b.Or(use_one, b.And(lhs_is_min, rhs_is_minus_one));
}
auto* select = b.Call("select", rhs, ScalarOrVector(width, 1_a), use_one);
body.Push(b.Return(is_div ? b.Div(lhs, select) : b.Mod(lhs, select)));
b.Func(name,
utils::Vector{
b.Param("lhs", T(lhs_ty)),
b.Param("rhs", T(rhs_ty)),
},
width == 1 ? T(lhs_ty) : b.ty.vec(T(lhs_el_ty), width), // return type
std::move(body));
return name;
});
auto* lhs = ctx.Clone(bin_op->lhs);
auto* rhs = ctx.Clone(bin_op->rhs);
return b.Call(fn, lhs, rhs);
}
private:
/// The clone context
CloneContext& ctx;
@ -687,6 +717,9 @@ struct BuiltinPolyfill::State {
/// The source clone context
const sem::Info& sem = ctx.src->Sem();
// Polyfill functions for binary operators.
utils::Hashmap<BinaryOpSignature, Symbol, 8> binary_op_polyfills;
/// @returns the AST type for the given sem type
const ast::Type* T(const sem::Type* ty) const { return CreateASTTypeFor(ctx, ty); }
@ -724,7 +757,6 @@ Transform::ApplyResult BuiltinPolyfill::Apply(const Program* src,
auto& polyfill = cfg->builtins;
utils::Hashmap<const sem::Builtin*, Symbol, 8> builtin_polyfills;
utils::Hashmap<BinaryOpSignature, Symbol, 8> binary_op_polyfills;
ProgramBuilder b;
CloneContext ctx{&b, src, /* auto_clone_symbols */ true};
@ -849,15 +881,7 @@ 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<sem::Vector>()) {
mask = b.Construct(CreateASTTypeFor(ctx, rhs_ty), mask);
}
auto* mod = b.And(ctx.Clone(bin_op->rhs), mask);
ctx.Replace(bin_op->rhs, mod);
ctx.Replace(bin_op, [bin_op, &s] { return s.BitshiftModulo(bin_op); });
made_changes = true;
}
break;
@ -867,13 +891,7 @@ 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;
}
}

32
src/tint/transform/builtin_polyfill_test.cc
View File

@ -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<BuiltinPolyfill::Config>(builtins);
auto got = Run<BuiltinPolyfill>(src, std::move(data));
EXPECT_EQ(expect, str(got));
}
} // namespace
} // namespace tint::transform

16
test/tint/bug/chromium/1273230.wgsl.expected.dxc.hlsl
View File

@ -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);

16
test/tint/bug/chromium/1273230.wgsl.expected.fxc.hlsl
View File

@ -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);

16
test/tint/bug/chromium/1273230.wgsl.expected.msl
View File

@ -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);

304
test/tint/bug/chromium/1273230.wgsl.expected.spvasm
View File

@ -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

4
test/tint/bug/chromium/1386647.wgsl Normal file
View File

@ -0,0 +1,4 @@
@compute @workgroup_size(1)
fn f(@builtin(global_invocation_id) v : vec3<u32>) {
let l = v.x << (v.y % 1);
}

19
test/tint/bug/chromium/1386647.wgsl.expected.dxc.hlsl Normal file
View File

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

19
test/tint/bug/chromium/1386647.wgsl.expected.fxc.hlsl Normal file
View File

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

17
test/tint/bug/chromium/1386647.wgsl.expected.glsl Normal file
View File

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

18
test/tint/bug/chromium/1386647.wgsl.expected.msl Normal file
View File

@ -0,0 +1,18 @@
#include <metal_stdlib>
using namespace metal;
uint tint_mod(uint lhs, uint rhs) {
return (lhs % select(rhs, 1u, (rhs == 0u)));
}
void f_inner(uint3 v) {
uint const tint_symbol = v[0];
uint const tint_symbol_1 = tint_mod(v[1], 1u);
uint const l = (tint_symbol << (tint_symbol_1 & 31u));
}
kernel void f(uint3 v [[thread_position_in_grid]]) {
f_inner(v);
return;
}

54
test/tint/bug/chromium/1386647.wgsl.expected.spvasm Normal file
View File

@ -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

4
test/tint/bug/chromium/1386647.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,4 @@
@compute @workgroup_size(1)
fn f(@builtin(global_invocation_id) v : vec3<u32>) {
let l = (v.x << (v.y % 1));
}

16
test/tint/bug/tint/1113.wgsl.expected.dxc.hlsl
View File

@ -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);

16
test/tint/bug/tint/1113.wgsl.expected.fxc.hlsl
View File

@ -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);

16
test/tint/bug/tint/1113.wgsl.expected.msl
View File

@ -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);

330
test/tint/bug/tint/1113.wgsl.expected.spvasm
View File

@ -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

8
test/tint/bug/tint/1520.spvasm.expected.dxc.hlsl
View File

@ -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;

8
test/tint/bug/tint/1520.spvasm.expected.fxc.hlsl
View File

@ -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;

8
test/tint/bug/tint/1520.spvasm.expected.glsl
View File

@ -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;

8
test/tint/bug/tint/1520.spvasm.expected.msl
View File

@ -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<int8_t, 16> 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;

4
test/tint/bug/tint/221.wgsl.expected.dxc.hlsl
View File

@ -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;

4
test/tint/bug/tint/221.wgsl.expected.fxc.hlsl
View File

@ -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;

8
test/tint/bug/tint/221.wgsl.expected.glsl
View File

@ -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) {

8
test/tint/bug/tint/221.wgsl.expected.msl
View File

@ -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<uint, 50> data;
};
uint tint_mod(uint lhs, uint rhs) {
return (lhs % select(rhs, 1u, (rhs == 0u)));
}
kernel void tint_symbol(device Buf* tint_symbol_2 [[buffer(0)]]) {
uint i = 0u;
while (true) {

8
test/tint/bug/tint/914.wgsl.expected.dxc.hlsl
View File

@ -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;

8
test/tint/bug/tint/914.wgsl.expected.fxc.hlsl
View File

@ -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;

8
test/tint/bug/tint/914.wgsl.expected.glsl
View File

@ -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)) {

8
test/tint/bug/tint/914.wgsl.expected.msl
View File

@ -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<tint_array<float, 64>, 64>* const tint_symbol_10, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_11, const constant Uniforms* const tint_symbol_12, const device Matrix* const tint_symbol_13, const device Matrix* const tint_symbol_14, device Matrix* const tint_symbol_15) {
for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 256u)) {
uint const i = (idx / 64u);

300
test/tint/bug/tint/914.wgsl.expected.spvasm
View File

@ -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

8
test/tint/bug/tint/942.wgsl.expected.dxc.hlsl
View File

@ -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;

8
test/tint/bug/tint/942.wgsl.expected.fxc.hlsl
View File

@ -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;

8
test/tint/bug/tint/942.wgsl.expected.glsl
View File

@ -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;

8
test/tint/bug/tint/942.wgsl.expected.msl
View File

@ -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<tint_array<float3, 256>, 4>* const tint_symbol_1, const constant Params* const tint_symbol_2, texture2d<float, access::sample> tint_symbol_3, const constant Flip* const tint_symbol_4, sampler tint_symbol_5, texture2d<float, access::write> tint_symbol_6) {
for(uint idx = local_invocation_index; (idx < 1024u); idx = (idx + 64u)) {
uint const i_1 = (idx / 256u);

6
test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.dxc.hlsl
View File

@ -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;

6
test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.fxc.hlsl
View File

@ -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;

4
test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.glsl
View File

@ -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;

6
test/tint/builtins/atomicStore/array/arrays.spvasm.expected.dxc.hlsl
View File

@ -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;

6
test/tint/builtins/atomicStore/array/arrays.spvasm.expected.fxc.hlsl
View File

@ -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;

4
test/tint/builtins/atomicStore/array/arrays.spvasm.expected.glsl
View File

@ -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;

6
test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.dxc.hlsl
View File

@ -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);

6
test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.fxc.hlsl
View File

@ -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);

4
test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.glsl
View File

@ -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);

8
test/tint/statements/compound_assign/private.wgsl.expected.dxc.hlsl
View File

@ -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);

8
test/tint/statements/compound_assign/private.wgsl.expected.fxc.hlsl
View File

@ -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);

6
test/tint/statements/compound_assign/private.wgsl.expected.glsl
View File

@ -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)));

4
test/tint/statements/compound_assign/scalar/divide.wgsl.expected.dxc.hlsl
View File

@ -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));

4
test/tint/statements/compound_assign/scalar/divide.wgsl.expected.fxc.hlsl
View File

@ -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));

8
test/tint/statements/compound_assign/scalar/divide.wgsl.expected.glsl
View File

@ -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;

8
test/tint/statements/compound_assign/scalar/divide.wgsl.expected.msl
View File

@ -1,14 +1,14 @@
#include <metal_stdlib>
using namespace metal;
int tint_div(int lhs, int rhs) {
return (lhs / select(rhs, 1, bool((rhs == 0) | bool((lhs == (-2147483647 - 1)) & (rhs == -1)))));
}
struct S {
/* 0x0000 */ int a;
};
int tint_div(int lhs, int rhs) {
return (lhs / select(rhs, 1, bool((rhs == 0) | bool((lhs == (-2147483647 - 1)) & (rhs == -1)))));
}
void foo(device S* const tint_symbol_1) {
int const tint_symbol = tint_div((*(tint_symbol_1)).a, 2);
(*(tint_symbol_1)).a = tint_symbol;

4
test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.dxc.hlsl
View File

@ -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));

4
test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.fxc.hlsl
View File

@ -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));

8
test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.glsl
View File

@ -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;

8
test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.msl
View File

@ -1,14 +1,14 @@
#include <metal_stdlib>
using namespace metal;
int tint_mod(int lhs, int rhs) {
return (lhs % select(rhs, 1, bool((rhs == 0) | bool((lhs == (-2147483647 - 1)) & (rhs == -1)))));
}
struct S {
/* 0x0000 */ int a;
};
int tint_mod(int lhs, int rhs) {
return (lhs % select(rhs, 1, bool((rhs == 0) | bool((lhs == (-2147483647 - 1)) & (rhs == -1)))));
}
void foo(device S* const tint_symbol_1) {
int const tint_symbol = tint_mod((*(tint_symbol_1)).a, 2);
(*(tint_symbol_1)).a = tint_symbol;

4
test/tint/statements/compound_assign/vector/divide.wgsl.expected.dxc.hlsl
View File

@ -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));

4
test/tint/statements/compound_assign/vector/divide.wgsl.expected.fxc.hlsl
View File

@ -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));

8
test/tint/statements/compound_assign/vector/divide.wgsl.expected.glsl
View File

@ -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;

8
test/tint/statements/compound_assign/vector/divide.wgsl.expected.msl
View File

@ -1,14 +1,14 @@
#include <metal_stdlib>
using namespace metal;
int4 tint_div(int4 lhs, int4 rhs) {
return (lhs / select(rhs, int4(1), ((rhs == int4(0)) | ((lhs == int4((-2147483647 - 1))) & (rhs == int4(-1))))));
}
struct S {
/* 0x0000 */ int4 a;
};
int4 tint_div(int4 lhs, int4 rhs) {
return (lhs / select(rhs, int4(1), ((rhs == int4(0)) | ((lhs == int4((-2147483647 - 1))) & (rhs == int4(-1))))));
}
void foo(device S* const tint_symbol_1) {
int4 const tint_symbol = tint_div((*(tint_symbol_1)).a, int4(2));
(*(tint_symbol_1)).a = tint_symbol;

4
test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.dxc.hlsl
View File

@ -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));

4
test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.fxc.hlsl
View File

@ -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));

10
test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.glsl
View File

@ -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;

8
test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.msl
View File

@ -1,15 +1,15 @@
#include <metal_stdlib>
using namespace metal;
struct S {
/* 0x0000 */ int4 a;
};
int4 tint_mod(int4 lhs, int rhs) {
int4 const r = int4(rhs);
return (lhs % select(r, int4(1), ((r == int4(0)) | ((lhs == int4((-2147483647 - 1))) & (r == int4(-1))))));
}
struct S {
/* 0x0000 */ int4 a;
};
void foo(device S* const tint_symbol_1) {
int4 const tint_symbol = tint_mod((*(tint_symbol_1)).a, 2);
(*(tint_symbol_1)).a = tint_symbol;

4
test/tint/statements/compound_assign/vector/modulo.wgsl.expected.dxc.hlsl
View File

@ -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));

4
test/tint/statements/compound_assign/vector/modulo.wgsl.expected.fxc.hlsl
View File

@ -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));

8
test/tint/statements/compound_assign/vector/modulo.wgsl.expected.glsl
View File

@ -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;

8
test/tint/statements/compound_assign/vector/modulo.wgsl.expected.msl
View File

@ -1,14 +1,14 @@
#include <metal_stdlib>
using namespace metal;
int4 tint_mod(int4 lhs, int4 rhs) {
return (lhs % select(rhs, int4(1), ((rhs == int4(0)) | ((lhs == int4((-2147483647 - 1))) & (rhs == int4(-1))))));
}
struct S {
/* 0x0000 */ int4 a;
};
int4 tint_mod(int4 lhs, int4 rhs) {
return (lhs % select(rhs, int4(1), ((rhs == int4(0)) | ((lhs == int4((-2147483647 - 1))) & (rhs == int4(-1))))));
}
void foo(device S* const tint_symbol_1) {
int4 const tint_symbol = tint_mod((*(tint_symbol_1)).a, int4(2));
(*(tint_symbol_1)).a = tint_symbol;

8
test/tint/statements/compound_assign/workgroup.wgsl.expected.dxc.hlsl
View File

@ -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);

8
test/tint/statements/compound_assign/workgroup.wgsl.expected.fxc.hlsl
View File

@ -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);

6
test/tint/statements/compound_assign/workgroup.wgsl.expected.glsl
View File

@ -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)));