[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); - ~~~~~~~~~~~~~~~~~~~~~ ^ ~~~ -