[msl] Add polyfill for floating point modulo
Use the `fmod` builtin for this case.
Bug: 42251016
Change-Id: Id7f34869391ee61768eca7f2c48d3c73a8c226db
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/196355
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/msl/builtin_fn.cc b/src/tint/lang/msl/builtin_fn.cc
index bda3140..9bc5c80 100644
--- a/src/tint/lang/msl/builtin_fn.cc
+++ b/src/tint/lang/msl/builtin_fn.cc
@@ -94,6 +94,8 @@
return "distance";
case BuiltinFn::kDot:
return "dot";
+ case BuiltinFn::kFmod:
+ return "fmod";
case BuiltinFn::kLength:
return "length";
case BuiltinFn::kThreadgroupBarrier:
diff --git a/src/tint/lang/msl/builtin_fn.h b/src/tint/lang/msl/builtin_fn.h
index b0cac68..2ad0d28 100644
--- a/src/tint/lang/msl/builtin_fn.h
+++ b/src/tint/lang/msl/builtin_fn.h
@@ -73,6 +73,7 @@
kWrite,
kDistance,
kDot,
+ kFmod,
kLength,
kThreadgroupBarrier,
kNone,
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
index 62e30c9..1272c5b 100644
--- a/src/tint/lang/msl/intrinsic/data.cc
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -2610,6 +2610,16 @@
},
{
/* [302] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(3),
+ },
+ {
+ /* [303] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(3),
+ },
+ {
+ /* [304] */
/* usage */ core::ParameterUsage::kTexture,
/* matcher_indices */ MatcherIndicesIndex(9),
},
@@ -3557,7 +3567,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(302),
+ /* parameters */ ParameterIndex(304),
/* return_matcher_indices */ MatcherIndicesIndex(59),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4338,7 +4348,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(302),
+ /* parameters */ ParameterIndex(304),
/* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4421,6 +4431,28 @@
},
{
/* [156] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 2,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 1,
+ /* templates */ TemplateIndex(12),
+ /* parameters */ ParameterIndex(302),
+ /* return_matcher_indices */ MatcherIndicesIndex(3),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [157] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 2,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 2,
+ /* templates */ TemplateIndex(11),
+ /* parameters */ ParameterIndex(300),
+ /* return_matcher_indices */ MatcherIndicesIndex(93),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [158] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 5,
/* num_explicit_templates */ 0,
@@ -4431,7 +4463,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [157] */
+ /* [159] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -4442,7 +4474,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [158] */
+ /* [160] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -4453,7 +4485,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [159] */
+ /* [161] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -4464,7 +4496,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [160] */
+ /* [162] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -4475,7 +4507,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [161] */
+ /* [163] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4486,7 +4518,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [162] */
+ /* [164] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
@@ -4506,67 +4538,67 @@
/* [0] */
/* fn atomic_compare_exchange_weak_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, ptr<function, T, read_write>, T, u32, u32) -> bool */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(156),
+ /* overloads */ OverloadIndex(158),
},
{
/* [1] */
/* fn atomic_exchange_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(157),
+ /* overloads */ OverloadIndex(159),
},
{
/* [2] */
/* fn atomic_fetch_add_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(157),
+ /* overloads */ OverloadIndex(159),
},
{
/* [3] */
/* fn atomic_fetch_and_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(157),
+ /* overloads */ OverloadIndex(159),
},
{
/* [4] */
/* fn atomic_fetch_max_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(157),
+ /* overloads */ OverloadIndex(159),
},
{
/* [5] */
/* fn atomic_fetch_min_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(157),
+ /* overloads */ OverloadIndex(159),
},
{
/* [6] */
/* fn atomic_fetch_or_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(157),
+ /* overloads */ OverloadIndex(159),
},
{
/* [7] */
/* fn atomic_fetch_sub_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(157),
+ /* overloads */ OverloadIndex(159),
},
{
/* [8] */
/* fn atomic_fetch_xor_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(157),
+ /* overloads */ OverloadIndex(159),
},
{
/* [9] */
/* fn atomic_load_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, u32) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(158),
+ /* overloads */ OverloadIndex(160),
},
{
/* [10] */
/* fn atomic_store_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(159),
+ /* overloads */ OverloadIndex(161),
},
{
/* [11] */
@@ -4793,25 +4825,32 @@
/* [24] */
/* fn distance[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(160),
+ /* overloads */ OverloadIndex(162),
},
{
/* [25] */
/* fn dot[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(160),
+ /* overloads */ OverloadIndex(162),
},
{
/* [26] */
- /* fn length[N : num, T : f32_f16](vec<N, T>) -> T */
- /* num overloads */ 1,
- /* overloads */ OverloadIndex(161),
+ /* fn fmod[T : f32_f16](T, T) -> T */
+ /* fn fmod[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> vec<N, T> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(156),
},
{
/* [27] */
+ /* fn length[N : num, T : f32_f16](vec<N, T>) -> T */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(163),
+ },
+ {
+ /* [28] */
/* fn threadgroup_barrier(u32) */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(162),
+ /* overloads */ OverloadIndex(164),
},
};
diff --git a/src/tint/lang/msl/msl.def b/src/tint/lang/msl/msl.def
index 3894e23..4c887be 100644
--- a/src/tint/lang/msl/msl.def
+++ b/src/tint/lang/msl/msl.def
@@ -323,6 +323,8 @@
fn distance[N: num, T: f32_f16](vec<N, T>, vec<N, T>) -> T
fn dot[N: num, T: f32_f16](vec<N, T>, vec<N, T>) -> T
+fn fmod[T: f32_f16](T, T) -> T
+fn fmod[N: num, T: f32_f16](vec<N, T>, vec<N, T>) -> vec<N, T>
fn length[N: num, T: f32_f16](vec<N, T>) -> T
@stage("compute") fn threadgroup_barrier(u32)
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index 0d2e61b..c4ea261 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -80,8 +80,9 @@
/// Process the module.
void Process() {
- // Find the builtins that need replacing.
- Vector<core::ir::CoreBuiltinCall*, 4> worklist;
+ // Find the builtins and binary operators that need replacing.
+ Vector<core::ir::CoreBinary*, 4> fmod_worklist;
+ Vector<core::ir::CoreBuiltinCall*, 4> builtin_worklist;
for (auto* inst : ir.Instructions()) {
if (auto* builtin = inst->As<core::ir::CoreBuiltinCall>()) {
switch (builtin->Func()) {
@@ -118,16 +119,21 @@
case core::BuiltinFn::kWorkgroupBarrier:
case core::BuiltinFn::kTextureBarrier:
case core::BuiltinFn::kUnpack2X16Float:
- worklist.Push(builtin);
+ builtin_worklist.Push(builtin);
break;
default:
break;
}
+ } else if (auto* binary = inst->As<core::ir::CoreBinary>()) {
+ if (binary->Op() == core::BinaryOp::kModulo &&
+ binary->LHS()->Type()->is_float_scalar_or_vector()) {
+ fmod_worklist.Push(binary);
+ }
}
}
// Replace the builtins that we found.
- for (auto* builtin : worklist) {
+ for (auto* builtin : builtin_worklist) {
switch (builtin->Func()) {
// Atomics.
case core::BuiltinFn::kAtomicAdd:
@@ -242,6 +248,11 @@
break;
}
}
+
+ // Replace the fmod instructions that we found.
+ for (auto* fmod : fmod_worklist) {
+ FMod(fmod);
+ }
}
/// Replace an atomic builtin call with an equivalent MSL intrinsic.
@@ -778,6 +789,15 @@
});
builtin->Destroy();
}
+
+ /// Replace a floating point modulo binary instruction with the equivalent MSL intrinsic.
+ /// @param binary the float point modulo binary instruction
+ void FMod(core::ir::CoreBinary* binary) {
+ auto* call = b.CallWithResult<msl::ir::BuiltinCall>(
+ binary->DetachResult(), msl::BuiltinFn::kFmod, binary->Operands());
+ call->InsertBefore(binary);
+ binary->Destroy();
+ }
};
} // namespace
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
index 467bb5e..6280415 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -2872,5 +2872,73 @@
EXPECT_EQ(expect, str());
}
+TEST_F(MslWriter_BuiltinPolyfillTest, FMod_Scalar) {
+ auto* lhs = b.FunctionParam<f32>("lhs");
+ auto* rhs = b.FunctionParam<f32>("rhs");
+ auto* func = b.Function("foo", ty.f32());
+ func->SetParams({lhs, rhs});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Modulo<f32>(lhs, rhs);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%lhs:f32, %rhs:f32):f32 {
+ $B1: {
+ %4:f32 = mod %lhs, %rhs
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%lhs:f32, %rhs:f32):f32 {
+ $B1: {
+ %4:f32 = msl.fmod %lhs, %rhs
+ ret %4
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, FMod_Vector) {
+ auto* lhs = b.FunctionParam<vec4<f16>>("lhs");
+ auto* rhs = b.FunctionParam<vec4<f16>>("rhs");
+ auto* func = b.Function("foo", ty.vec4<f16>());
+ func->SetParams({lhs, rhs});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Modulo<vec4<f16>>(lhs, rhs);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%lhs:vec4<f16>, %rhs:vec4<f16>):vec4<f16> {
+ $B1: {
+ %4:vec4<f16> = mod %lhs, %rhs
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%lhs:vec4<f16>, %rhs:vec4<f16>):vec4<f16> {
+ $B1: {
+ %4:vec4<f16> = msl.fmod %lhs, %rhs
+ ret %4
+ }
+}
+)";
+
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::msl::writer::raise
diff --git a/test/tint/bug/chromium/1434271.wgsl.expected.ir.msl b/test/tint/bug/chromium/1434271.wgsl.expected.ir.msl
index 028ef20..e15baf4 100644
--- a/test/tint/bug/chromium/1434271.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/1434271.wgsl.expected.ir.msl
@@ -141,7 +141,7 @@
float const c = (*tint_module_vars.buf_in).weights[((src_offset + 1u) + (*tint_module_vars.ubo).width)];
float const d = (*tint_module_vars.buf_in).weights[((src_offset + 1u) + (*tint_module_vars.ubo).width)];
float const sum = dot(float4(a, b, c, d), float4(1.0f));
- (*tint_module_vars.buf_out).weights[dst_offset] = (sum % 4.0f);
+ (*tint_module_vars.buf_out).weights[dst_offset] = fmod(sum, 4.0f);
float4 const v_4 = float4(a, (a * b), ((a / b) + c), sum);
float4 const probabilities = (v_4 + max(sum, 0.0f));
tint_module_vars.tex_out.write(probabilities, uint2(int2(coord.xy)));
@@ -171,9 +171,6 @@
program_source:89:8: warning: unused variable 'res' [-Wunused-variable]
half res = asinh(arg_0);
^
-program_source:142:60: error: invalid operands to binary expression ('const float' and 'float')
- (*tint_module_vars.buf_out).weights[dst_offset] = (sum % 4.0f);
- ~~~ ^ ~~~~
program_source:154:104: error: call to deleted constructor of 'texture1d<float, access::sample>'
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.render_params=render_params};
^
diff --git a/test/tint/bug/tint/948.wgsl.expected.ir.msl b/test/tint/bug/tint/948.wgsl.expected.ir.msl
index 65f1035..2d0fd9c 100644
--- a/test/tint/bug/tint/948.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/948.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
@@ -144,7 +142,7 @@
if ((x_174 > 0.0f)) {
float const x_181 = (*tint_module_vars.x_20).time;
float const x_184 = animationData[2u];
- (*tint_module_vars.mt) = ((x_181 * x_184) % 1.0f);
+ (*tint_module_vars.mt) = fmod((x_181 * x_184), 1.0f);
f = 0.0f;
{
while(true) {
@@ -257,16 +255,3 @@
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.x_20=x_20, .frameMapTexture=frameMapTexture, .frameMapSampler=frameMapSampler, .tUV=(&tUV), .tileMapsTexture0=tileMapsTexture0, .tileMapsSampler=tileMapsSampler, .tileMapsTexture1=tileMapsTexture1, .animationMapTexture=animationMapTexture, .animationMapSampler=animationMapSampler, .mt=(&mt), .spriteSheetTexture=spriteSheetTexture, .spriteSheetSampler=spriteSheetSampler, .glFragColor=(&glFragColor), .tileID_1=(&tileID_1), .levelUnits=(&levelUnits), .stageUnits_1=(&stageUnits_1), .vPosition=(&vPosition), .vUV=(&vUV)};
return tint_symbol_outputs{.main_out_glFragColor_1=tint_symbol_inner(inputs.tUV_param, inputs.tileID_1_param, inputs.levelUnits_param, inputs.stageUnits_1_param, inputs.vPosition_param, inputs.vUV_param, tint_module_vars).glFragColor_1};
}
-program_source:145:51: error: invalid operands to binary expression ('float' and 'float')
- (*tint_module_vars.mt) = ((x_181 * x_184) % 1.0f);
- ~~~~~~~~~~~~~~~ ^ ~~~~
-program_source:161:25: warning: unused variable 'x_208' [-Wunused-variable]
- float const x_208 = frameID_1;
- ^
-program_source:162:25: warning: unused variable 'x_211' [-Wunused-variable]
- float const x_211 = (*tint_module_vars.x_20).spriteCount;
- ^
-program_source:163:25: warning: unused variable 'x_214' [-Wunused-variable]
- float const x_214 = f;
- ^
-
diff --git a/test/tint/expressions/binary/mod/scalar-scalar/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/scalar-scalar/f16.wgsl.expected.ir.msl
index 812cbe6..fa2c3ee 100644
--- a/test/tint/expressions/binary/mod/scalar-scalar/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/scalar-scalar/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
half const a = 1.0h;
half const b = 2.0h;
- half const r = (a % b);
+ half const r = fmod(a, b);
}
-program_source:7:21: error: invalid operands to binary expression ('const half' and 'const half')
- half const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod/scalar-scalar/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/scalar-scalar/f32.wgsl.expected.ir.msl
index c55eb34..7360105 100644
--- a/test/tint/expressions/binary/mod/scalar-scalar/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/scalar-scalar/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
float const a = 1.0f;
float const b = 2.0f;
- float const r = (a % b);
+ float const r = fmod(a, b);
}
-program_source:7:22: error: invalid operands to binary expression ('const float' and 'const float')
- float const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/scalar-vec3/f16.wgsl.expected.ir.msl
index 4e7e3dd..9d36d1e 100644
--- a/test/tint/expressions/binary/mod/scalar-vec3/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/scalar-vec3/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
half const a = 4.0h;
half3 const b = half3(1.0h, 2.0h, 3.0h);
- half3 const r = (a % b);
+ half3 const r = fmod(a, b);
}
-program_source:7:22: error: invalid operands to binary expression ('const half' and 'const half3' (vector of 3 'half' values))
- half3 const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.ir.msl
index c20de57..00748aa 100644
--- a/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
float const a = 4.0f;
float3 const b = float3(1.0f, 2.0f, 3.0f);
- float3 const r = (a % b);
+ float3 const r = fmod(a, b);
}
-program_source:7:23: error: invalid operands to binary expression ('const float' and 'const float3' (vector of 3 'float' values))
- float3 const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/vec3-scalar/f16.wgsl.expected.ir.msl
index 339b67d..eb99c54 100644
--- a/test/tint/expressions/binary/mod/vec3-scalar/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/vec3-scalar/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
half3 const a = half3(1.0h, 2.0h, 3.0h);
half const b = 4.0h;
- half3 const r = (a % b);
+ half3 const r = fmod(a, b);
}
-program_source:7:22: error: invalid operands to binary expression ('const half3' (vector of 3 'half' values) and 'const half')
- half3 const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.ir.msl
index fcd961f..3e9ccce 100644
--- a/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
float3 const a = float3(1.0f, 2.0f, 3.0f);
float const b = 4.0f;
- float3 const r = (a % b);
+ float3 const r = fmod(a, b);
}
-program_source:7:23: error: invalid operands to binary expression ('const float3' (vector of 3 'float' values) and 'const float')
- float3 const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod/vec3-vec3/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/vec3-vec3/f16.wgsl.expected.ir.msl
index 5eb95c9..c331b93 100644
--- a/test/tint/expressions/binary/mod/vec3-vec3/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/vec3-vec3/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
half3 const a = half3(1.0h, 2.0h, 3.0h);
half3 const b = half3(4.0h, 5.0h, 6.0h);
- half3 const r = (a % b);
+ half3 const r = fmod(a, b);
}
-program_source:7:22: error: invalid operands to binary expression ('const half3' (vector of 3 'half' values) and 'const half3')
- half3 const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod/vec3-vec3/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/vec3-vec3/f32.wgsl.expected.ir.msl
index 7ad70c2..c1fcab3 100644
--- a/test/tint/expressions/binary/mod/vec3-vec3/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/vec3-vec3/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
float3 const a = float3(1.0f, 2.0f, 3.0f);
float3 const b = float3(4.0f, 5.0f, 6.0f);
- float3 const r = (a % b);
+ float3 const r = fmod(a, b);
}
-program_source:7:23: error: invalid operands to binary expression ('const float3' (vector of 3 'float' values) and 'const float3')
- float3 const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f16.wgsl.expected.ir.msl
index 3883afa..d169baa 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
half const a = 1.0h;
half const b = 0.0h;
- half const r = (a % b);
+ half const r = fmod(a, b);
}
-program_source:7:21: error: invalid operands to binary expression ('const half' and 'const half')
- half const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.ir.msl
index f4e0d01..d46406f 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
float const a = 1.0f;
float const b = 0.0f;
- float const r = (a % b);
+ float const r = fmod(a, b);
}
-program_source:7:22: error: invalid operands to binary expression ('const float' and 'const float')
- float const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f16.wgsl.expected.ir.msl
index 500fa80..8826b2a 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
half3 const a = half3(1.0h, 2.0h, 3.0h);
half3 const b = half3(0.0h, 5.0h, 0.0h);
- half3 const r = (a % b);
+ half3 const r = fmod(a, b);
}
-program_source:7:22: error: invalid operands to binary expression ('const half3' (vector of 3 'half' values) and 'const half3')
- half3 const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.ir.msl
index b771d78..6496f21 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
float3 const a = float3(1.0f, 2.0f, 3.0f);
float3 const b = float3(0.0f, 5.0f, 0.0f);
- float3 const r = (a % b);
+ float3 const r = fmod(a, b);
}
-program_source:7:23: error: invalid operands to binary expression ('const float3' (vector of 3 'float' values) and 'const float3')
- float3 const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f16.wgsl.expected.ir.msl
index 347f215..cfb7230 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
half a = 1.0h;
half b = 0.0h;
- half const r = (a % (b + b));
+ half const r = fmod(a, (b + b));
}
-program_source:7:21: error: invalid operands to binary expression ('half' and 'half')
- half const r = (a % (b + b));
- ~ ^ ~~~~~~~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.ir.msl
index 5b44ae7..2a021b0 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
float a = 1.0f;
float b = 0.0f;
- float const r = (a % (b + b));
+ float const r = fmod(a, (b + b));
}
-program_source:7:22: error: invalid operands to binary expression ('float' and 'float')
- float const r = (a % (b + b));
- ~ ^ ~~~~~~~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f16.wgsl.expected.ir.msl
index cfcd9fa..5cb5c8d 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
half3 a = half3(1.0h, 2.0h, 3.0h);
half3 b = half3(0.0h, 5.0h, 0.0h);
- half3 const r = (a % (b + b));
+ half3 const r = fmod(a, (b + b));
}
-program_source:7:22: error: invalid operands to binary expression ('half3' (vector of 3 'half' values) and 'half3')
- half3 const r = (a % (b + b));
- ~ ^ ~~~~~~~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.ir.msl
index 3238577..b6f792c 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#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 = (a % (b + b));
+ float3 const r = fmod(a, (b + b));
}
-program_source:7:23: error: invalid operands to binary expression ('float3' (vector of 3 'float' values) and 'float3')
- float3 const r = (a % (b + b));
- ~ ^ ~~~~~~~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f16.wgsl.expected.ir.msl
index b925c09..eddd331 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
half a = 1.0h;
half b = 0.0h;
- half const r = (a % b);
+ half const r = fmod(a, b);
}
-program_source:7:21: error: invalid operands to binary expression ('half' and 'half')
- half const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.ir.msl
index 730fb96..bf6392d 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
float a = 1.0f;
float b = 0.0f;
- float const r = (a % b);
+ float const r = fmod(a, b);
}
-program_source:7:22: error: invalid operands to binary expression ('float' and 'float')
- float const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f16.wgsl.expected.ir.msl
index a81f54b..b04f111 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
kernel void f() {
half3 a = half3(1.0h, 2.0h, 3.0h);
half3 b = half3(0.0h, 5.0h, 0.0h);
- half3 const r = (a % b);
+ half3 const r = fmod(a, b);
}
-program_source:7:22: error: invalid operands to binary expression ('half3' (vector of 3 'half' values) and 'half3')
- half3 const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.ir.msl
index 0dbb18d..29578fd 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
#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 = (a % b);
+ float3 const r = fmod(a, b);
}
-program_source:7:23: error: invalid operands to binary expression ('float3' (vector of 3 'float' values) and 'float3')
- float3 const r = (a % b);
- ~ ^ ~
-
diff --git a/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.ir.msl
index 4ab3cb3..8494843 100644
--- a/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
#include <metal_stdlib>
using namespace metal;
@@ -23,16 +21,9 @@
(*tint_module_vars.a) = tint_div_i32((*tint_module_vars.a), maybe_zero);
(*tint_module_vars.a) = tint_mod_i32((*tint_module_vars.a), maybe_zero);
(*tint_module_vars.b) = ((*tint_module_vars.b) / 0.0f);
- (*tint_module_vars.b) = ((*tint_module_vars.b) % 0.0f);
+ (*tint_module_vars.b) = fmod((*tint_module_vars.b), 0.0f);
float const v_1 = float(maybe_zero);
(*tint_module_vars.b) = ((*tint_module_vars.b) / v_1);
float const v_2 = float(maybe_zero);
- (*tint_module_vars.b) = ((*tint_module_vars.b) % v_2);
+ (*tint_module_vars.b) = fmod((*tint_module_vars.b), v_2);
}
-program_source:24:50: error: invalid operands to binary expression ('float' and 'float')
- (*tint_module_vars.b) = ((*tint_module_vars.b) % 0.0f);
- ~~~~~~~~~~~~~~~~~~~~~ ^ ~~~~
-program_source:28:50: error: invalid operands to binary expression ('float' and 'const float')
- (*tint_module_vars.b) = ((*tint_module_vars.b) % v_2);
- ~~~~~~~~~~~~~~~~~~~~~ ^ ~~~
-
