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

Fix FXC compile errors on modulo by zero 

Just like for divide, FXC fails with the exact same error when
performing a modulo on a value that FXC determines to be zero. We
address it in the same way as we do for divide.

This also fixes a couple of the vk-gl-cts tests for which I manually
generated expectation files for.

Bug: tint:1083
Change-Id: Ia388bf002112afded542adb791d37e88e35a77ff
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/74220
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
This commit is contained in:
Antonio Maiorano 2021-12-22 15:02:09 +00:00 committed by Tint LUCI CQ
parent cc4d97b6e3
commit 9943de6813
159 changed files with 1965 additions and 69 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
src/program_builder.h
View File

@ -1665,6 +1665,16 @@ class ProgramBuilder {
Expr(std::forward<RHS>(rhs)));
}
/// @param lhs the left hand argument to the modulo operation
/// @param rhs the right hand argument to the modulo operation
/// @returns a `ast::BinaryExpression` applying modulo of `lhs` by `rhs`
template <typename LHS, typename RHS>
const ast::Expression* Mod(LHS&& lhs, RHS&& rhs) {
return create<ast::BinaryExpression>(ast::BinaryOp::kModulo,
Expr(std::forward<LHS>(lhs)),
Expr(std::forward<RHS>(rhs)));
}
/// @param lhs the left hand argument to the bit shift right operation
/// @param rhs the right hand argument to the bit shift right operation
/// @returns a `ast::BinaryExpression` bit shifting right `lhs` by `rhs`

10
src/writer/hlsl/generator_impl.cc
View File

@ -863,8 +863,8 @@ bool GeneratorImpl::EmitBinary(std::ostream& out,
break;
case ast::BinaryOp::kDivide:
out << "/";
// BUG(crbug.com/tint/1083): Integer divide by zero is a FXC compile
// error, and undefined behavior in WGSL.
// BUG(crbug.com/tint/1083): Integer divide/modulo by zero is a FXC
// compile error, and undefined behavior in WGSL.
if (TypeOf(expr->rhs)->UnwrapRef()->is_integer_scalar_or_vector()) {
out << " ";
return EmitExpressionOrOneIfZero(out, expr->rhs);
@ -872,6 +872,12 @@ bool GeneratorImpl::EmitBinary(std::ostream& out,
break;
case ast::BinaryOp::kModulo:
out << "%";
// BUG(crbug.com/tint/1083): Integer divide/modulo by zero is a FXC
// compile error, and undefined behavior in WGSL.
if (TypeOf(expr->rhs)->UnwrapRef()->is_integer_scalar_or_vector()) {
out << " ";
return EmitExpressionOrOneIfZero(out, expr->rhs);
}
break;
case ast::BinaryOp::kNone:
diagnostics_.add_error(diag::System::Writer,

126
src/writer/hlsl/generator_impl_binary_test.cc
View File

@ -137,10 +137,12 @@ INSTANTIATE_TEST_SUITE_P(
BinaryData{"(left + right)", ast::BinaryOp::kAdd},
BinaryData{"(left - right)", ast::BinaryOp::kSubtract},
BinaryData{"(left * right)", ast::BinaryOp::kMultiply},
// NOTE: Integer divide covered by DivideBy* tests below
// NOTE: Integer divide covered by DivOrModBy* tests below
BinaryData{"(left / right)", ast::BinaryOp::kDivide,
BinaryData::Types::Float},
BinaryData{"(left % right)", ast::BinaryOp::kModulo}));
// NOTE: Integer modulo covered by DivOrModBy* tests below
BinaryData{"(left % right)", ast::BinaryOp::kModulo,
BinaryData::Types::Float}));
TEST_F(HlslGeneratorImplTest_Binary, Multiply_VectorScalar) {
auto* lhs = vec3<f32>(1.f, 1.f, 1.f);
@ -526,11 +528,36 @@ foo((tint_tmp), (tint_tmp_1), (tint_tmp_2));
)");
}
TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_i32) {
namespace HlslGeneratorDivMod {
struct Params {
enum class Type { Div, Mod };
Type type;
};
struct HlslGeneratorDivModTest : TestParamHelper<Params> {
std::string Token() {
return GetParam().type == Params::Type::Div ? "/" : "%";
}
template <typename... Args>
auto Op(Args... args) {
return GetParam().type == Params::Type::Div
? Div(std::forward<Args>(args)...)
: Mod(std::forward<Args>(args)...);
}
};
INSTANTIATE_TEST_SUITE_P(HlslGeneratorImplTest,
HlslGeneratorDivModTest,
testing::Values(Params{Params::Type::Div},
Params{Params::Type::Mod}));
TEST_P(HlslGeneratorDivModTest, DivOrModByLiteralZero_i32) {
Func("fn", {}, ty.void_(),
{
Decl(Var("a", ty.i32())),
Decl(Const("r", nullptr, Div("a", 0))),
Decl(Const("r", nullptr, Op("a", 0))),
});
GeneratorImpl& gen = Build();
@ -538,16 +565,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_i32) {
ASSERT_TRUE(gen.Generate());
EXPECT_EQ(gen.result(), R"(void fn() {
int a = 0;
const int r = (a / 1);
const int r = (a )" + Token() +
R"( 1);
}
)");
}
TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_u32) {
TEST_P(HlslGeneratorDivModTest, DivOrModByLiteralZero_u32) {
Func("fn", {}, ty.void_(),
{
Decl(Var("a", ty.u32())),
Decl(Const("r", nullptr, Div("a", 0u))),
Decl(Const("r", nullptr, Op("a", 0u))),
});
GeneratorImpl& gen = Build();
@ -555,16 +583,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_u32) {
ASSERT_TRUE(gen.Generate());
EXPECT_EQ(gen.result(), R"(void fn() {
uint a = 0u;
const uint r = (a / 1u);
const uint r = (a )" + Token() +
R"( 1u);
}
)");
}
} // namespace HlslGeneratorDivMod
TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_vec_by_vec_i32) {
TEST_P(HlslGeneratorDivModTest, DivOrModByLiteralZero_vec_by_vec_i32) {
Func("fn", {}, ty.void_(),
{
Decl(Var("a", nullptr, vec4<i32>(100, 100, 100, 100))),
Decl(Const("r", nullptr, Div("a", vec4<i32>(50, 0, 25, 0)))),
Decl(Const("r", nullptr, Op("a", vec4<i32>(50, 0, 25, 0)))),
});
GeneratorImpl& gen = Build();
@ -572,16 +601,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_vec_by_vec_i32) {
ASSERT_TRUE(gen.Generate());
EXPECT_EQ(gen.result(), R"(void fn() {
int4 a = int4(100, 100, 100, 100);
const int4 r = (a / int4(50, 1, 25, 1));
const int4 r = (a )" + Token() +
R"( int4(50, 1, 25, 1));
}
)");
}
} // namespace
TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_vec_by_scalar_i32) {
TEST_P(HlslGeneratorDivModTest, DivOrModByLiteralZero_vec_by_scalar_i32) {
Func("fn", {}, ty.void_(),
{
Decl(Var("a", nullptr, vec4<i32>(100, 100, 100, 100))),
Decl(Const("r", nullptr, Div("a", 0))),
Decl(Const("r", nullptr, Op("a", 0))),
});
GeneratorImpl& gen = Build();
@ -589,16 +619,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByLiteralZero_vec_by_scalar_i32) {
ASSERT_TRUE(gen.Generate());
EXPECT_EQ(gen.result(), R"(void fn() {
int4 a = int4(100, 100, 100, 100);
const int4 r = (a / 1);
const int4 r = (a )" + Token() +
R"( 1);
}
)");
}
} // namespace hlsl
TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_i32) {
TEST_P(HlslGeneratorDivModTest, DivOrModByIdentifier_i32) {
Func("fn", {Param("b", ty.i32())}, ty.void_(),
{
Decl(Var("a", ty.i32())),
Decl(Const("r", nullptr, Div("a", "b"))),
Decl(Const("r", nullptr, Op("a", "b"))),
});
GeneratorImpl& gen = Build();
@ -606,16 +637,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_i32) {
ASSERT_TRUE(gen.Generate());
EXPECT_EQ(gen.result(), R"(void fn(int b) {
int a = 0;
const int r = (a / (b == 0 ? 1 : b));
const int r = (a )" + Token() +
R"( (b == 0 ? 1 : b));
}
)");
}
} // namespace writer
TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_u32) {
TEST_P(HlslGeneratorDivModTest, DivOrModByIdentifier_u32) {
Func("fn", {Param("b", ty.u32())}, ty.void_(),
{
Decl(Var("a", ty.u32())),
Decl(Const("r", nullptr, Div("a", "b"))),
Decl(Const("r", nullptr, Op("a", "b"))),
});
GeneratorImpl& gen = Build();
@ -623,16 +655,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_u32) {
ASSERT_TRUE(gen.Generate());
EXPECT_EQ(gen.result(), R"(void fn(uint b) {
uint a = 0u;
const uint r = (a / (b == 0u ? 1u : b));
const uint r = (a )" + Token() +
R"( (b == 0u ? 1u : b));
}
)");
}
} // namespace tint
TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_vec_by_vec_i32) {
TEST_P(HlslGeneratorDivModTest, DivOrModByIdentifier_vec_by_vec_i32) {
Func("fn", {Param("b", ty.vec3<i32>())}, ty.void_(),
{
Decl(Var("a", ty.vec3<i32>())),
Decl(Const("r", nullptr, Div("a", "b"))),
Decl(Const("r", nullptr, Op("a", "b"))),
});
GeneratorImpl& gen = Build();
@ -640,16 +673,17 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_vec_by_vec_i32) {
ASSERT_TRUE(gen.Generate());
EXPECT_EQ(gen.result(), R"(void fn(int3 b) {
int3 a = int3(0, 0, 0);
const int3 r = (a / (b == int3(0, 0, 0) ? int3(1, 1, 1) : b));
const int3 r = (a )" + Token() +
R"( (b == int3(0, 0, 0) ? int3(1, 1, 1) : b));
}
)");
}
TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_vec_by_scalar_i32) {
TEST_P(HlslGeneratorDivModTest, DivOrModByIdentifier_vec_by_scalar_i32) {
Func("fn", {Param("b", ty.i32())}, ty.void_(),
{
Decl(Var("a", ty.vec3<i32>())),
Decl(Const("r", nullptr, Div("a", "b"))),
Decl(Const("r", nullptr, Op("a", "b"))),
});
GeneratorImpl& gen = Build();
@ -657,12 +691,13 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByIdentifier_vec_by_scalar_i32) {
ASSERT_TRUE(gen.Generate());
EXPECT_EQ(gen.result(), R"(void fn(int b) {
int3 a = int3(0, 0, 0);
const int3 r = (a / (b == 0 ? 1 : b));
const int3 r = (a )" + Token() +
R"( (b == 0 ? 1 : b));
}
)");
}
TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_i32) {
TEST_P(HlslGeneratorDivModTest, DivOrModByExpression_i32) {
Func("zero", {}, ty.i32(),
{
Return(Expr(0)),
@ -671,7 +706,7 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_i32) {
Func("fn", {}, ty.void_(),
{
Decl(Var("a", ty.i32())),
Decl(Const("r", nullptr, Div("a", Call("zero")))),
Decl(Const("r", nullptr, Op("a", Call("zero")))),
});
GeneratorImpl& gen = Build();
@ -687,12 +722,13 @@ int zero() {
void fn() {
int a = 0;
const int r = (a / value_or_one_if_zero_int(zero()));
const int r = (a )" + Token() +
R"( value_or_one_if_zero_int(zero()));
}
)");
}
TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_u32) {
TEST_P(HlslGeneratorDivModTest, DivOrModByExpression_u32) {
Func("zero", {}, ty.u32(),
{
Return(Expr(0u)),
@ -701,7 +737,7 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_u32) {
Func("fn", {}, ty.void_(),
{
Decl(Var("a", ty.u32())),
Decl(Const("r", nullptr, Div("a", Call("zero")))),
Decl(Const("r", nullptr, Op("a", Call("zero")))),
});
GeneratorImpl& gen = Build();
@ -717,12 +753,13 @@ uint zero() {
void fn() {
uint a = 0u;
const uint r = (a / value_or_one_if_zero_uint(zero()));
const uint r = (a )" + Token() +
R"( value_or_one_if_zero_uint(zero()));
}
)");
}
TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_vec_by_vec_i32) {
TEST_P(HlslGeneratorDivModTest, DivOrModByExpression_vec_by_vec_i32) {
Func("zero", {}, ty.vec3<i32>(),
{
Return(vec3<i32>(0, 0, 0)),
@ -731,7 +768,7 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_vec_by_vec_i32) {
Func("fn", {}, ty.void_(),
{
Decl(Var("a", ty.vec3<i32>())),
Decl(Const("r", nullptr, Div("a", Call("zero")))),
Decl(Const("r", nullptr, Op("a", Call("zero")))),
});
GeneratorImpl& gen = Build();
@ -747,12 +784,13 @@ int3 zero() {
void fn() {
int3 a = int3(0, 0, 0);
const int3 r = (a / value_or_one_if_zero_int3(zero()));
const int3 r = (a )" + Token() +
R"( value_or_one_if_zero_int3(zero()));
}
)");
}
TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_vec_by_scalar_i32) {
TEST_P(HlslGeneratorDivModTest, DivOrModByExpression_vec_by_scalar_i32) {
Func("zero", {}, ty.i32(),
{
Return(0),
@ -761,7 +799,7 @@ TEST_F(HlslGeneratorImplTest_Binary, DivideByExpression_vec_by_scalar_i32) {
Func("fn", {}, ty.void_(),
{
Decl(Var("a", ty.vec3<i32>())),
Decl(Const("r", nullptr, Div("a", Call("zero")))),
Decl(Const("r", nullptr, Op("a", Call("zero")))),
});
GeneratorImpl& gen = Build();
@ -777,10 +815,12 @@ int zero() {
void fn() {
int3 a = int3(0, 0, 0);
const int3 r = (a / value_or_one_if_zero_int(zero()));
const int3 r = (a )" + Token() +
R"( value_or_one_if_zero_int(zero()));
}
)");
}
} // namespace HlslGeneratorDivMod
} // namespace
} // namespace hlsl

2
test/bug/chromium/1273230.wgsl.expected.hlsl
View File

@ -89,7 +89,7 @@ uint toIndex1D(uint gridSize, float3 voxelPos) {
uint3 toIndex4D(uint gridSize, uint index) {
uint z_1 = (gridSize / value_or_one_if_zero_uint((index * index)));
uint y_1 = ((gridSize - ((gridSize * gridSize) * z_1)) / (gridSize == 0u ? 1u : gridSize));
uint x_1 = (index % gridSize);
uint x_1 = (index % (gridSize == 0u ? 1u : gridSize));
return uint3(z_1, y_1, y_1);
}

2
test/bug/tint/1113.wgsl.expected.hlsl
View File

@ -60,7 +60,7 @@ uint toIndex1D(uint gridSize, float3 voxelPos) {
uint3 toIndex3D(uint gridSize, uint index) {
uint z_1 = (index / value_or_one_if_zero_uint((gridSize * gridSize)));
uint y_1 = ((index - ((gridSize * gridSize) * z_1)) / (gridSize == 0u ? 1u : gridSize));
uint x_1 = (index % gridSize);
uint x_1 = (index % (gridSize == 0u ? 1u : gridSize));
return uint3(x_1, y_1, z_1);
}

6
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 1.;
let b = 0.;
let r : f32 = a % b;
}

5
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,5 @@
[numthreads(1, 1, 1)]
void f() {
const float r = (1.0f % 0.0f);
return;
}

10
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
float const a = 1.0f;
float const b = 0.0f;
float const r = fmod(a, b);
return;
}

20
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,20 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 9
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
%float_1 = OpConstant %float 1
%float_0 = OpConstant %float 0
%f = OpFunction %void None %1
%4 = OpLabel
%8 = OpFRem %float %float_1 %float_0
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 1.0;
let b = 0.0;
let r : f32 = (a % b);
}

6
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 1;
let b = 0;
let r : i32 = a % b;
}

5
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,5 @@
[numthreads(1, 1, 1)]
void f() {
const int r = (1 % 1);
return;
}

10
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
int const a = 1;
int const b = 0;
int const r = (a % b);
return;
}

20
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,20 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 9
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%int = OpTypeInt 32 1
%int_1 = OpConstant %int 1
%int_0 = OpConstant %int 0
%f = OpFunction %void None %1
%4 = OpLabel
%8 = OpSMod %int %int_1 %int_0
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 1;
let b = 0;
let r : i32 = (a % b);
}

6
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 1u;
let b = 0u;
let r : u32 = a % b;
}

5
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,5 @@
[numthreads(1, 1, 1)]
void f() {
const uint r = (1u % 1u);
return;
}

10
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
uint const a = 1u;
uint const b = 0u;
uint const r = (a % b);
return;
}

20
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,20 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 9
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1
%uint_0 = OpConstant %uint 0
%f = OpFunction %void None %1
%4 = OpLabel
%8 = OpUMod %uint %uint_1 %uint_0
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 1u;
let b = 0u;
let r : u32 = (a % b);
}

6
test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 4;
let b = vec3<i32>(0, 2, 0);
let r : vec3<i32> = a % b;
}

7
test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,7 @@
[numthreads(1, 1, 1)]
void f() {
const int a = 4;
const int3 b = int3(0, 2, 0);
const int3 r = (a % int3(1, 2, 1));
return;
}

10
test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
int const a = 4;
int3 const b = int3(0, 2, 0);
int3 const r = (a % b);
return;
}

27
test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,27 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 16
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%int = OpTypeInt 32 1
%int_4 = OpConstant %int 4
%v3int = OpTypeVector %int 3
%int_0 = OpConstant %int 0
%int_2 = OpConstant %int 2
%10 = OpConstantComposite %v3int %int_0 %int_2 %int_0
%_ptr_Function_v3int = OpTypePointer Function %v3int
%14 = OpConstantNull %v3int
%f = OpFunction %void None %1
%4 = OpLabel
%12 = OpVariable %_ptr_Function_v3int Function %14
%15 = OpCompositeConstruct %v3int %int_4 %int_4 %int_4
%11 = OpSMod %v3int %15 %10
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 4;
let b = vec3<i32>(0, 2, 0);
let r : vec3<i32> = (a % b);
}

6
test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 4u;
let b = vec3<u32>(0u, 2u, 0u);
let r : vec3<u32> = a % b;
}

7
test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,7 @@
[numthreads(1, 1, 1)]
void f() {
const uint a = 4u;
const uint3 b = uint3(0u, 2u, 0u);
const uint3 r = (a % uint3(1u, 2u, 1u));
return;
}

10
test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
uint const a = 4u;
uint3 const b = uint3(0u, 2u, 0u);
uint3 const r = (a % b);
return;
}

27
test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,27 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 16
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%v3uint = OpTypeVector %uint 3
%uint_0 = OpConstant %uint 0
%uint_2 = OpConstant %uint 2
%10 = OpConstantComposite %v3uint %uint_0 %uint_2 %uint_0
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
%14 = OpConstantNull %v3uint
%f = OpFunction %void None %1
%4 = OpLabel
%12 = OpVariable %_ptr_Function_v3uint Function %14
%15 = OpCompositeConstruct %v3uint %uint_4 %uint_4 %uint_4
%11 = OpUMod %v3uint %15 %10
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = 4u;
let b = vec3<u32>(0u, 2u, 0u);
let r : vec3<u32> = (a % b);
}

6
test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<i32>(1, 2, 3);
let b = 0;
let r : vec3<i32> = a % b;
}

6
test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,6 @@
[numthreads(1, 1, 1)]
void f() {
const int3 a = int3(1, 2, 3);
const int3 r = (a % 1);
return;
}

10
test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
int3 const a = int3(1, 2, 3);
int const b = 0;
int3 const r = (a % b);
return;
}

28
test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,28 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 17
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%int = OpTypeInt 32 1
%v3int = OpTypeVector %int 3
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%10 = OpConstantComposite %v3int %int_1 %int_2 %int_3
%int_0 = OpConstant %int 0
%_ptr_Function_v3int = OpTypePointer Function %v3int
%15 = OpConstantNull %v3int
%f = OpFunction %void None %1
%4 = OpLabel
%13 = OpVariable %_ptr_Function_v3int Function %15
%16 = OpCompositeConstruct %v3int %int_0 %int_0 %int_0
%12 = OpSMod %v3int %10 %16
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<i32>(1, 2, 3);
let b = 0;
let r : vec3<i32> = (a % b);
}

6
test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<u32>(1u, 2u, 3u);
let b = 0u;
let r : vec3<u32> = a % b;
}

6
test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,6 @@
[numthreads(1, 1, 1)]
void f() {
const uint3 a = uint3(1u, 2u, 3u);
const uint3 r = (a % 1u);
return;
}

10
test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
uint3 const a = uint3(1u, 2u, 3u);
uint const b = 0u;
uint3 const r = (a % b);
return;
}

28
test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,28 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 17
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_3 = OpConstant %uint 3
%10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
%uint_0 = OpConstant %uint 0
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
%15 = OpConstantNull %v3uint
%f = OpFunction %void None %1
%4 = OpLabel
%13 = OpVariable %_ptr_Function_v3uint Function %15
%16 = OpCompositeConstruct %v3uint %uint_0 %uint_0 %uint_0
%12 = OpUMod %v3uint %10 %16
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<u32>(1u, 2u, 3u);
let b = 0u;
let r : vec3<u32> = (a % b);
}

6
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<f32>(1., 2., 3.);
let b = vec3<f32>(0., 5., 0.);
let r : vec3<f32> = a % b;
}

7
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,7 @@
[numthreads(1, 1, 1)]
void f() {
const float3 a = float3(1.0f, 2.0f, 3.0f);
const float3 b = float3(0.0f, 5.0f, 0.0f);
const float3 r = (a % b);
return;
}

10
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
float3 const a = float3(1.0f, 2.0f, 3.0f);
float3 const b = float3(0.0f, 5.0f, 0.0f);
float3 const r = fmod(a, b);
return;
}

26
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,26 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 15
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
%v3float = OpTypeVector %float 3
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%float_3 = OpConstant %float 3
%10 = OpConstantComposite %v3float %float_1 %float_2 %float_3
%float_0 = OpConstant %float 0
%float_5 = OpConstant %float 5
%13 = OpConstantComposite %v3float %float_0 %float_5 %float_0
%f = OpFunction %void None %1
%4 = OpLabel
%14 = OpFRem %v3float %10 %13
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<f32>(1.0, 2.0, 3.0);
let b = vec3<f32>(0.0, 5.0, 0.0);
let r : vec3<f32> = (a % b);
}

6
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<i32>(1, 2, 3);
let b = vec3<i32>(0, 5, 0);
let r : vec3<i32> = a % b;
}

7
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,7 @@
[numthreads(1, 1, 1)]
void f() {
const int3 a = int3(1, 2, 3);
const int3 b = int3(0, 5, 0);
const int3 r = (a % int3(1, 5, 1));
return;
}

10
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
int3 const a = int3(1, 2, 3);
int3 const b = int3(0, 5, 0);
int3 const r = (a % b);
return;
}

26
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,26 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 15
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%int = OpTypeInt 32 1
%v3int = OpTypeVector %int 3
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%10 = OpConstantComposite %v3int %int_1 %int_2 %int_3
%int_0 = OpConstant %int 0
%int_5 = OpConstant %int 5
%13 = OpConstantComposite %v3int %int_0 %int_5 %int_0
%f = OpFunction %void None %1
%4 = OpLabel
%14 = OpSMod %v3int %10 %13
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<i32>(1, 2, 3);
let b = vec3<i32>(0, 5, 0);
let r : vec3<i32> = (a % b);
}

6
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<u32>(1u, 2u, 3u);
let b = vec3<u32>(0u, 5u, 0u);
let r : vec3<u32> = a % b;
}

7
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,7 @@
[numthreads(1, 1, 1)]
void f() {
const uint3 a = uint3(1u, 2u, 3u);
const uint3 b = uint3(0u, 5u, 0u);
const uint3 r = (a % uint3(1u, 5u, 1u));
return;
}

10
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
uint3 const a = uint3(1u, 2u, 3u);
uint3 const b = uint3(0u, 5u, 0u);
uint3 const r = (a % b);
return;
}

26
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,26 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 15
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_3 = OpConstant %uint 3
%10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
%uint_0 = OpConstant %uint 0
%uint_5 = OpConstant %uint 5
%13 = OpConstantComposite %v3uint %uint_0 %uint_5 %uint_0
%f = OpFunction %void None %1
%4 = OpLabel
%14 = OpUMod %v3uint %10 %13
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
let a = vec3<u32>(1u, 2u, 3u);
let b = vec3<u32>(0u, 5u, 0u);
let r : vec3<u32> = (a % b);
}

6
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = 1.;
var b = 0.;
let r : f32 = a % (b + b);
}

7
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,7 @@
[numthreads(1, 1, 1)]
void f() {
float a = 1.0f;
float b = 0.0f;
const float r = (a % (b + b));
return;
}

10
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
float a = 1.0f;
float b = 0.0f;
float const r = fmod(a, (b + b));
return;
}

32
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,32 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 17
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
OpName %a "a"
OpName %b "b"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
%float_1 = OpConstant %float 1
%_ptr_Function_float = OpTypePointer Function %float
%9 = OpConstantNull %float
%float_0 = OpConstant %float 0
%f = OpFunction %void None %1
%4 = OpLabel
%a = OpVariable %_ptr_Function_float Function %9
%b = OpVariable %_ptr_Function_float Function %9
OpStore %a %float_1
OpStore %b %float_0
%12 = OpLoad %float %a
%13 = OpLoad %float %b
%14 = OpLoad %float %b
%15 = OpFAdd %float %13 %14
%16 = OpFRem %float %12 %15
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = 1.0;
var b = 0.0;
let r : f32 = (a % (b + b));
}

6
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = 1;
var b = 0;
let r : i32 = a % (b + b);
}

11
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,11 @@
int value_or_one_if_zero_int(int value) {
return value == 0 ? 1 : value;
}
[numthreads(1, 1, 1)]
void f() {
int a = 1;
int b = 0;
const int r = (a % value_or_one_if_zero_int((b + b)));
return;
}

10
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
int a = 1;
int b = 0;
int const r = (a % as_type<int>((as_type<uint>(b) + as_type<uint>(b))));
return;
}

32
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,32 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 17
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
OpName %a "a"
OpName %b "b"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%int = OpTypeInt 32 1
%int_1 = OpConstant %int 1
%_ptr_Function_int = OpTypePointer Function %int
%9 = OpConstantNull %int
%int_0 = OpConstant %int 0
%f = OpFunction %void None %1
%4 = OpLabel
%a = OpVariable %_ptr_Function_int Function %9
%b = OpVariable %_ptr_Function_int Function %9
OpStore %a %int_1
OpStore %b %int_0
%12 = OpLoad %int %a
%13 = OpLoad %int %b
%14 = OpLoad %int %b
%15 = OpIAdd %int %13 %14
%16 = OpSMod %int %12 %15
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = 1;
var b = 0;
let r : i32 = (a % (b + b));
}

6
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = 1u;
var b = 0u;
let r : u32 = a % (b + b);
}

11
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,11 @@
uint value_or_one_if_zero_uint(uint value) {
return value == 0u ? 1u : value;
}
[numthreads(1, 1, 1)]
void f() {
uint a = 1u;
uint b = 0u;
const uint r = (a % value_or_one_if_zero_uint((b + b)));
return;
}

10
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
uint a = 1u;
uint b = 0u;
uint const r = (a % (b + b));
return;
}

32
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,32 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 17
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
OpName %a "a"
OpName %b "b"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1
%_ptr_Function_uint = OpTypePointer Function %uint
%9 = OpConstantNull %uint
%uint_0 = OpConstant %uint 0
%f = OpFunction %void None %1
%4 = OpLabel
%a = OpVariable %_ptr_Function_uint Function %9
%b = OpVariable %_ptr_Function_uint Function %9
OpStore %a %uint_1
OpStore %b %uint_0
%12 = OpLoad %uint %a
%13 = OpLoad %uint %b
%14 = OpLoad %uint %b
%15 = OpIAdd %uint %13 %14
%16 = OpUMod %uint %12 %15
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = 1u;
var b = 0u;
let r : u32 = (a % (b + b));
}

6
test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = 4;
var b = vec3<i32>(0, 2, 0);
let r : vec3<i32> = a % (b + b);
}

11
test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,11 @@
int3 value_or_one_if_zero_int3(int3 value) {
return value == int3(0, 0, 0) ? int3(1, 1, 1) : value;
}
[numthreads(1, 1, 1)]
void f() {
int a = 4;
int3 b = int3(0, 2, 0);
const int3 r = (a % value_or_one_if_zero_int3((b + b)));
return;
}

10
test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
int a = 4;
int3 b = int3(0, 2, 0);
int3 const r = (a % as_type<int3>((as_type<uint3>(b) + as_type<uint3>(b))));
return;
}

39
test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,39 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 24
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
OpName %a "a"
OpName %b "b"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%int = OpTypeInt 32 1
%int_4 = OpConstant %int 4
%_ptr_Function_int = OpTypePointer Function %int
%9 = OpConstantNull %int
%v3int = OpTypeVector %int 3
%int_0 = OpConstant %int 0
%int_2 = OpConstant %int 2
%13 = OpConstantComposite %v3int %int_0 %int_2 %int_0
%_ptr_Function_v3int = OpTypePointer Function %v3int
%16 = OpConstantNull %v3int
%f = OpFunction %void None %1
%4 = OpLabel
%a = OpVariable %_ptr_Function_int Function %9
%b = OpVariable %_ptr_Function_v3int Function %16
%22 = OpVariable %_ptr_Function_v3int Function %16
OpStore %a %int_4
OpStore %b %13
%17 = OpLoad %int %a
%18 = OpLoad %v3int %b
%19 = OpLoad %v3int %b
%20 = OpIAdd %v3int %18 %19
%23 = OpCompositeConstruct %v3int %17 %17 %17
%21 = OpSMod %v3int %23 %20
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = 4;
var b = vec3<i32>(0, 2, 0);
let r : vec3<i32> = (a % (b + b));
}

6
test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = 4u;
var b = vec3<u32>(0u, 2u, 0u);
let r : vec3<u32> = a % (b + b);
}

11
test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,11 @@
uint3 value_or_one_if_zero_uint3(uint3 value) {
return value == uint3(0u, 0u, 0u) ? uint3(1u, 1u, 1u) : value;
}
[numthreads(1, 1, 1)]
void f() {
uint a = 4u;
uint3 b = uint3(0u, 2u, 0u);
const uint3 r = (a % value_or_one_if_zero_uint3((b + b)));
return;
}

10
test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
uint a = 4u;
uint3 b = uint3(0u, 2u, 0u);
uint3 const r = (a % (b + b));
return;
}

39
test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,39 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 24
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
OpName %a "a"
OpName %b "b"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4
%_ptr_Function_uint = OpTypePointer Function %uint
%9 = OpConstantNull %uint
%v3uint = OpTypeVector %uint 3
%uint_0 = OpConstant %uint 0
%uint_2 = OpConstant %uint 2
%13 = OpConstantComposite %v3uint %uint_0 %uint_2 %uint_0
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
%16 = OpConstantNull %v3uint
%f = OpFunction %void None %1
%4 = OpLabel
%a = OpVariable %_ptr_Function_uint Function %9
%b = OpVariable %_ptr_Function_v3uint Function %16
%22 = OpVariable %_ptr_Function_v3uint Function %16
OpStore %a %uint_4
OpStore %b %13
%17 = OpLoad %uint %a
%18 = OpLoad %v3uint %b
%19 = OpLoad %v3uint %b
%20 = OpIAdd %v3uint %18 %19
%23 = OpCompositeConstruct %v3uint %17 %17 %17
%21 = OpUMod %v3uint %23 %20
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = 4u;
var b = vec3<u32>(0u, 2u, 0u);
let r : vec3<u32> = (a % (b + b));
}

6
test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = vec3<i32>(1, 2, 3);
var b = 0;
let r : vec3<i32> = a % (b + b);
}

11
test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,11 @@
int value_or_one_if_zero_int(int value) {
return value == 0 ? 1 : value;
}
[numthreads(1, 1, 1)]
void f() {
int3 a = int3(1, 2, 3);
int b = 0;
const int3 r = (a % value_or_one_if_zero_int((b + b)));
return;
}

10
test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
int3 a = int3(1, 2, 3);
int b = 0;
int3 const r = (a % as_type<int>((as_type<uint>(b) + as_type<uint>(b))));
return;
}

40
test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,40 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 25
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
OpName %a "a"
OpName %b "b"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%int = OpTypeInt 32 1
%v3int = OpTypeVector %int 3
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%10 = OpConstantComposite %v3int %int_1 %int_2 %int_3
%_ptr_Function_v3int = OpTypePointer Function %v3int
%13 = OpConstantNull %v3int
%int_0 = OpConstant %int 0
%_ptr_Function_int = OpTypePointer Function %int
%17 = OpConstantNull %int
%f = OpFunction %void None %1
%4 = OpLabel
%a = OpVariable %_ptr_Function_v3int Function %13
%b = OpVariable %_ptr_Function_int Function %17
%23 = OpVariable %_ptr_Function_v3int Function %13
OpStore %a %10
OpStore %b %int_0
%18 = OpLoad %v3int %a
%19 = OpLoad %int %b
%20 = OpLoad %int %b
%21 = OpIAdd %int %19 %20
%24 = OpCompositeConstruct %v3int %21 %21 %21
%22 = OpSMod %v3int %18 %24
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = vec3<i32>(1, 2, 3);
var b = 0;
let r : vec3<i32> = (a % (b + b));
}

6
test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = vec3<u32>(1u, 2u, 3u);
var b = 0u;
let r : vec3<u32> = a % (b + b);
}

11
test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,11 @@
uint value_or_one_if_zero_uint(uint value) {
return value == 0u ? 1u : value;
}
[numthreads(1, 1, 1)]
void f() {
uint3 a = uint3(1u, 2u, 3u);
uint b = 0u;
const uint3 r = (a % value_or_one_if_zero_uint((b + b)));
return;
}

10
test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
uint3 a = uint3(1u, 2u, 3u);
uint b = 0u;
uint3 const r = (a % (b + b));
return;
}

40
test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,40 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 25
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
OpName %a "a"
OpName %b "b"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%uint_1 = OpConstant %uint 1
%uint_2 = OpConstant %uint 2
%uint_3 = OpConstant %uint 3
%10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
%13 = OpConstantNull %v3uint
%uint_0 = OpConstant %uint 0
%_ptr_Function_uint = OpTypePointer Function %uint
%17 = OpConstantNull %uint
%f = OpFunction %void None %1
%4 = OpLabel
%a = OpVariable %_ptr_Function_v3uint Function %13
%b = OpVariable %_ptr_Function_uint Function %17
%23 = OpVariable %_ptr_Function_v3uint Function %13
OpStore %a %10
OpStore %b %uint_0
%18 = OpLoad %v3uint %a
%19 = OpLoad %uint %b
%20 = OpLoad %uint %b
%21 = OpIAdd %uint %19 %20
%24 = OpCompositeConstruct %v3uint %21 %21 %21
%22 = OpUMod %v3uint %18 %24
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = vec3<u32>(1u, 2u, 3u);
var b = 0u;
let r : vec3<u32> = (a % (b + b));
}

6
test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = vec3<f32>(1., 2., 3.);
var b = vec3<f32>(0., 5., 0.);
let r : vec3<f32> = a % (b + b);
}

7
test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,7 @@
[numthreads(1, 1, 1)]
void f() {
float3 a = float3(1.0f, 2.0f, 3.0f);
float3 b = float3(0.0f, 5.0f, 0.0f);
const float3 r = (a % (b + b));
return;
}

10
test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
float3 a = float3(1.0f, 2.0f, 3.0f);
float3 b = float3(0.0f, 5.0f, 0.0f);
float3 const r = fmod(a, (b + b));
return;
}

38
test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,38 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 23
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
OpName %a "a"
OpName %b "b"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
%v3float = OpTypeVector %float 3
%float_1 = OpConstant %float 1
%float_2 = OpConstant %float 2
%float_3 = OpConstant %float 3
%10 = OpConstantComposite %v3float %float_1 %float_2 %float_3
%_ptr_Function_v3float = OpTypePointer Function %v3float
%13 = OpConstantNull %v3float
%float_0 = OpConstant %float 0
%float_5 = OpConstant %float 5
%16 = OpConstantComposite %v3float %float_0 %float_5 %float_0
%f = OpFunction %void None %1
%4 = OpLabel
%a = OpVariable %_ptr_Function_v3float Function %13
%b = OpVariable %_ptr_Function_v3float Function %13
OpStore %a %10
OpStore %b %16
%18 = OpLoad %v3float %a
%19 = OpLoad %v3float %b
%20 = OpLoad %v3float %b
%21 = OpFAdd %v3float %19 %20
%22 = OpFRem %v3float %18 %21
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = vec3<f32>(1.0, 2.0, 3.0);
var b = vec3<f32>(0.0, 5.0, 0.0);
let r : vec3<f32> = (a % (b + b));
}

6
test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = vec3<i32>(1, 2, 3);
var b = vec3<i32>(0, 5, 0);
let r : vec3<i32> = a % (b + b);
}

11
test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.hlsl Normal file
View File

@ -0,0 +1,11 @@
int3 value_or_one_if_zero_int3(int3 value) {
return value == int3(0, 0, 0) ? int3(1, 1, 1) : value;
}
[numthreads(1, 1, 1)]
void f() {
int3 a = int3(1, 2, 3);
int3 b = int3(0, 5, 0);
const int3 r = (a % value_or_one_if_zero_int3((b + b)));
return;
}

10
test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl Normal file
View File

@ -0,0 +1,10 @@
#include <metal_stdlib>
using namespace metal;
kernel void f() {
int3 a = int3(1, 2, 3);
int3 b = int3(0, 5, 0);
int3 const r = (a % as_type<int3>((as_type<uint3>(b) + as_type<uint3>(b))));
return;
}

38
test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,38 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 23
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %f "f"
OpExecutionMode %f LocalSize 1 1 1
OpName %f "f"
OpName %a "a"
OpName %b "b"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%int = OpTypeInt 32 1
%v3int = OpTypeVector %int 3
%int_1 = OpConstant %int 1
%int_2 = OpConstant %int 2
%int_3 = OpConstant %int 3
%10 = OpConstantComposite %v3int %int_1 %int_2 %int_3
%_ptr_Function_v3int = OpTypePointer Function %v3int
%13 = OpConstantNull %v3int
%int_0 = OpConstant %int 0
%int_5 = OpConstant %int 5
%16 = OpConstantComposite %v3int %int_0 %int_5 %int_0
%f = OpFunction %void None %1
%4 = OpLabel
%a = OpVariable %_ptr_Function_v3int Function %13
%b = OpVariable %_ptr_Function_v3int Function %13
OpStore %a %10
OpStore %b %16
%18 = OpLoad %v3int %a
%19 = OpLoad %v3int %b
%20 = OpLoad %v3int %b
%21 = OpIAdd %v3int %19 %20
%22 = OpSMod %v3int %18 %21
OpReturn
OpFunctionEnd

6
test/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,6 @@
[[stage(compute), workgroup_size(1)]]
fn f() {
var a = vec3<i32>(1, 2, 3);
var b = vec3<i32>(0, 5, 0);
let r : vec3<i32> = (a % (b + b));
}

Some files were not shown because too many files have changed in this diff Show More