From 33aaa289961ece67049a59e4fd33e51263966987 Mon Sep 17 00:00:00 2001 From: Fredrik Fornwall Date: Fri, 1 Sep 2023 14:20:14 +0200 Subject: [PATCH 1/4] [wgsl-in] Allow sign() to take int argument --- src/back/msl/keywords.rs | 1 + src/back/msl/writer.rs | 31 ++- src/valid/expression.rs | 17 +- tests/in/math-functions.wgsl | 1 + .../glsl/math-functions.main.Fragment.glsl | 5 +- tests/out/hlsl/math-functions.hlsl | 5 +- tests/out/msl/access.msl | 4 + tests/out/msl/array-in-ctor.msl | 4 + .../out/msl/array-in-function-return-type.msl | 4 + tests/out/msl/atomicOps.msl | 4 + tests/out/msl/binding-arrays.msl | 4 + tests/out/msl/bitcast.msl | 4 + tests/out/msl/bits.msl | 4 + tests/out/msl/boids.msl | 4 + tests/out/msl/bounds-check-image-restrict.msl | 4 + tests/out/msl/bounds-check-image-rzsw.msl | 4 + tests/out/msl/bounds-check-restrict.msl | 4 + tests/out/msl/bounds-check-zero-atomic.msl | 4 + tests/out/msl/bounds-check-zero.msl | 4 + tests/out/msl/break-if.msl | 4 + tests/out/msl/collatz.msl | 4 + tests/out/msl/constructors.msl | 4 + tests/out/msl/control-flow.msl | 4 + tests/out/msl/do-while.msl | 4 + tests/out/msl/dualsource.msl | 4 + tests/out/msl/empty-global-name.msl | 4 + tests/out/msl/empty.msl | 4 + tests/out/msl/extra.msl | 4 + tests/out/msl/fragment-output.msl | 4 + tests/out/msl/functions.msl | 4 + tests/out/msl/globals.msl | 4 + tests/out/msl/image.msl | 4 + tests/out/msl/interface.msl | 4 + tests/out/msl/interpolate.msl | 4 + tests/out/msl/math-functions.msl | 17 +- tests/out/msl/operators.msl | 4 + tests/out/msl/padding.msl | 4 + tests/out/msl/policy-mix.msl | 4 + tests/out/msl/quad-vert.msl | 4 + tests/out/msl/quad.msl | 4 + tests/out/msl/ray-query.msl | 4 + tests/out/msl/resource-binding-map.msl | 4 + tests/out/msl/shadow.msl | 4 + tests/out/msl/skybox.msl | 4 + tests/out/msl/standard.msl | 4 + tests/out/msl/texture-arg.msl | 4 + tests/out/msl/workgroup-uniform-load.msl | 4 + tests/out/msl/workgroup-var-init.msl | 4 + tests/out/spv/math-functions.spvasm | 181 +++++++++--------- tests/out/wgsl/math-functions.wgsl | 1 + 50 files changed, 318 insertions(+), 105 deletions(-) diff --git a/src/back/msl/keywords.rs b/src/back/msl/keywords.rs index f34b618db8..670e2f85ed 100644 --- a/src/back/msl/keywords.rs +++ b/src/back/msl/keywords.rs @@ -216,4 +216,5 @@ pub const RESERVED: &[&str] = &[ "clamped_lod_e", super::writer::FREXP_FUNCTION, super::writer::MODF_FUNCTION, + super::writer::ISIGN_FUNCTION, ]; diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index 24c54ff8ab..f60d9f7b3c 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -34,6 +34,7 @@ const RAY_QUERY_FUN_MAP_INTERSECTION: &str = "_map_intersection_type"; pub(crate) const MODF_FUNCTION: &str = "naga_modf"; pub(crate) const FREXP_FUNCTION: &str = "naga_frexp"; +pub(crate) const ISIGN_FUNCTION: &str = "naga_isign"; /// Write the Metal name for a Naga numeric type: scalar, vector, or matrix. /// @@ -1647,8 +1648,9 @@ impl Writer { } => { use crate::MathFunction as Mf; - let scalar_argument = match *context.resolve_type(arg) { - crate::TypeInner::Scalar { .. } => true, + let arg_type = context.resolve_type(arg); + let scalar_argument = match arg_type { + &crate::TypeInner::Scalar { .. } => true, _ => false, }; @@ -1713,7 +1715,17 @@ impl Writer { Mf::Reflect => "reflect", Mf::Refract => "refract", // computational - Mf::Sign => "sign", + Mf::Sign => match arg_type { + &crate::TypeInner::Scalar { + kind: crate::ScalarKind::Sint, + .. + } + | &crate::TypeInner::Vector { + kind: crate::ScalarKind::Sint, + .. + } => ISIGN_FUNCTION, + _ => "sign", + }, Mf::Fma => "fma", Mf::Mix => "mix", Mf::Step => "step", @@ -1816,7 +1828,7 @@ impl Writer { write!(self.out, "((")?; self.put_expression(arg, context, false)?; write!(self.out, ") * 57.295779513082322865)")?; - } else if fun == Mf::Modf || fun == Mf::Frexp { + } else if fun == Mf::Modf || fun == Mf::Frexp || fun_name == ISIGN_FUNCTION { write!(self.out, "{fun_name}")?; self.put_call_parameters(iter::once(arg), context)?; } else { @@ -3091,6 +3103,7 @@ impl Writer { self.write_type_defs(module)?; self.write_global_constants(module)?; + self.write_polyfills()?; self.write_functions(module, info, options, pipeline_options) } @@ -4103,6 +4116,16 @@ impl Writer { Ok(info) } + fn write_polyfills(&mut self) -> Result<(), Error> { + writeln!( + self.out, + "\ntemplate inline T {ISIGN_FUNCTION}(T arg) {{ + return {NAMESPACE}::select({NAMESPACE}::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +}}" + )?; + Ok(()) + } + fn write_barrier(&mut self, flags: crate::Barrier, level: back::Level) -> BackendResult { // Note: OR-ring bitflags requires `__HAVE_MEMFLAG_OPERATORS__`, // so we try to avoid it here. diff --git a/src/valid/expression.rs b/src/valid/expression.rs index 548b6c1451..ddef1aaea0 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -976,7 +976,6 @@ impl super::Validator { | Mf::Log | Mf::Log2 | Mf::Length - | Mf::Sign | Mf::Sqrt | Mf::InverseSqrt => { if arg1_ty.is_some() | arg2_ty.is_some() | arg3_ty.is_some() { @@ -992,6 +991,22 @@ impl super::Validator { _ => return Err(ExpressionError::InvalidArgumentType(fun, 0, arg)), } } + Mf::Sign => { + if arg1_ty.is_some() | arg2_ty.is_some() | arg3_ty.is_some() { + return Err(ExpressionError::WrongArgumentCount(fun)); + } + match *arg_ty { + Ti::Scalar { + kind: Sk::Float | Sk::Sint, + .. + } + | Ti::Vector { + kind: Sk::Float | Sk::Sint, + .. + } => {} + _ => return Err(ExpressionError::InvalidArgumentType(fun, 0, arg)), + } + } Mf::Atan2 | Mf::Pow | Mf::Distance | Mf::Step => { let arg1_ty = match (arg1_ty, arg2_ty, arg3_ty) { (Some(ty1), None, None) => ty1, diff --git a/tests/in/math-functions.wgsl b/tests/in/math-functions.wgsl index 408f8a74f8..6e7bccf7f6 100644 --- a/tests/in/math-functions.wgsl +++ b/tests/in/math-functions.wgsl @@ -8,6 +8,7 @@ fn main() { let d = radians(v); let e = saturate(v); let g = refract(v, v, f); + let h = sign(-1); let const_dot = dot(vec2(), vec2()); let first_leading_bit_abs = firstLeadingBit(abs(0u)); let flb_a = firstLeadingBit(-1); diff --git a/tests/out/glsl/math-functions.main.Fragment.glsl b/tests/out/glsl/math-functions.main.Fragment.glsl index be81715ce1..595235bdd1 100644 --- a/tests/out/glsl/math-functions.main.Fragment.glsl +++ b/tests/out/glsl/math-functions.main.Fragment.glsl @@ -62,6 +62,7 @@ void main() { vec4 d = radians(v); vec4 e = clamp(v, vec4(0.0), vec4(1.0)); vec4 g = refract(v, v, 1.0); + int h = sign(-1); int const_dot = ( + ivec2(0).x * ivec2(0).x + ivec2(0).y * ivec2(0).y); uint first_leading_bit_abs = uint(findMSB(uint(abs(int(0u))))); int flb_a = findMSB(-1); @@ -81,8 +82,8 @@ void main() { ivec2 ctz_h = ivec2(min(uvec2(findLSB(ivec2(1))), uvec2(32u))); int clz_a = (-1 < 0 ? 0 : 31 - findMSB(-1)); uint clz_b = uint(31 - findMSB(1u)); - ivec2 _e58 = ivec2(-1); - ivec2 clz_c = mix(ivec2(31) - findMSB(_e58), ivec2(0), lessThan(_e58, ivec2(0))); + ivec2 _e60 = ivec2(-1); + ivec2 clz_c = mix(ivec2(31) - findMSB(_e60), ivec2(0), lessThan(_e60, ivec2(0))); uvec2 clz_d = uvec2(ivec2(31) - findMSB(uvec2(1u))); float lde_a = ldexp(1.0, 2); vec2 lde_b = ldexp(vec2(1.0, 2.0), ivec2(3, 4)); diff --git a/tests/out/hlsl/math-functions.hlsl b/tests/out/hlsl/math-functions.hlsl index afdc6f4671..6a2e4e6787 100644 --- a/tests/out/hlsl/math-functions.hlsl +++ b/tests/out/hlsl/math-functions.hlsl @@ -72,6 +72,7 @@ void main() float4 d = radians(v); float4 e = saturate(v); float4 g = refract(v, v, 1.0); + int h = sign(-1); int const_dot = dot((int2)0, (int2)0); uint first_leading_bit_abs = firstbithigh(abs(0u)); int flb_a = asint(firstbithigh(-1)); @@ -91,8 +92,8 @@ void main() int2 ctz_h = asint(min((32u).xx, firstbitlow((1).xx))); int clz_a = (-1 < 0 ? 0 : 31 - asint(firstbithigh(-1))); uint clz_b = (31u - firstbithigh(1u)); - int2 _expr58 = (-1).xx; - int2 clz_c = (_expr58 < (0).xx ? (0).xx : (31).xx - asint(firstbithigh(_expr58))); + int2 _expr60 = (-1).xx; + int2 clz_c = (_expr60 < (0).xx ? (0).xx : (31).xx - asint(firstbithigh(_expr60))); uint2 clz_d = ((31u).xx - firstbithigh((1u).xx)); float lde_a = ldexp(1.0, 2); float2 lde_b = ldexp(float2(1.0, 2.0), int2(3, 4)); diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index e5d875dd19..d1fe75ee5d 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -58,6 +58,10 @@ struct type_26 { metal::float4 inner[2]; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + void test_matrix_within_struct_accesses( constant Baz& baz ) { diff --git a/tests/out/msl/array-in-ctor.msl b/tests/out/msl/array-in-ctor.msl index 9428cb1e74..e322333d61 100644 --- a/tests/out/msl/array-in-ctor.msl +++ b/tests/out/msl/array-in-ctor.msl @@ -11,6 +11,10 @@ struct Ah { type_1 inner; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + kernel void cs_main( device Ah const& ah [[user(fake0)]] ) { diff --git a/tests/out/msl/array-in-function-return-type.msl b/tests/out/msl/array-in-function-return-type.msl index c2c2379cce..7475f01ff0 100644 --- a/tests/out/msl/array-in-function-return-type.msl +++ b/tests/out/msl/array-in-function-return-type.msl @@ -8,6 +8,10 @@ struct type_1 { float inner[2]; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + type_1 ret_array( ) { return type_1 {1.0, 2.0}; diff --git a/tests/out/msl/atomicOps.msl b/tests/out/msl/atomicOps.msl index b7264e883d..7bff77dff7 100644 --- a/tests/out/msl/atomicOps.msl +++ b/tests/out/msl/atomicOps.msl @@ -12,6 +12,10 @@ struct Struct { type_2 atomic_arr; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + struct cs_mainInput { }; kernel void cs_main( diff --git a/tests/out/msl/binding-arrays.msl b/tests/out/msl/binding-arrays.msl index 10a5397910..b826460e01 100644 --- a/tests/out/msl/binding-arrays.msl +++ b/tests/out/msl/binding-arrays.msl @@ -17,6 +17,10 @@ struct FragmentIn { uint index; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + struct main_Input { uint index [[user(loc0), flat]]; }; diff --git a/tests/out/msl/bitcast.msl b/tests/out/msl/bitcast.msl index a0cf093b7f..75d6228a6c 100644 --- a/tests/out/msl/bitcast.msl +++ b/tests/out/msl/bitcast.msl @@ -5,6 +5,10 @@ using metal::uint; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + kernel void main_( ) { metal::int2 i2_ = {}; diff --git a/tests/out/msl/bits.msl b/tests/out/msl/bits.msl index 2a00b6b843..785f0adf60 100644 --- a/tests/out/msl/bits.msl +++ b/tests/out/msl/bits.msl @@ -5,6 +5,10 @@ using metal::uint; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + kernel void main_( ) { int i = {}; diff --git a/tests/out/msl/boids.msl b/tests/out/msl/boids.msl index 1a81aaf684..dcb92f51a5 100644 --- a/tests/out/msl/boids.msl +++ b/tests/out/msl/boids.msl @@ -28,6 +28,10 @@ struct Particles { }; constant uint NUM_PARTICLES = 1500u; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + struct main_Input { }; kernel void main_( diff --git a/tests/out/msl/bounds-check-image-restrict.msl b/tests/out/msl/bounds-check-image-restrict.msl index 9f94ef0a6e..4b631593d1 100644 --- a/tests/out/msl/bounds-check-image-restrict.msl +++ b/tests/out/msl/bounds-check-image-restrict.msl @@ -5,6 +5,10 @@ using metal::uint; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + metal::float4 test_textureLoad_1d( int coords, int level, diff --git a/tests/out/msl/bounds-check-image-rzsw.msl b/tests/out/msl/bounds-check-image-rzsw.msl index a93014cb27..8d26f222a2 100644 --- a/tests/out/msl/bounds-check-image-rzsw.msl +++ b/tests/out/msl/bounds-check-image-rzsw.msl @@ -11,6 +11,10 @@ struct DefaultConstructible { }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + metal::float4 test_textureLoad_1d( int coords, int level, diff --git a/tests/out/msl/bounds-check-restrict.msl b/tests/out/msl/bounds-check-restrict.msl index 49a232e706..7c3c5ca6ee 100644 --- a/tests/out/msl/bounds-check-restrict.msl +++ b/tests/out/msl/bounds-check-restrict.msl @@ -20,6 +20,10 @@ struct Globals { type_4 d; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + float index_array( int i, device Globals const& globals, diff --git a/tests/out/msl/bounds-check-zero-atomic.msl b/tests/out/msl/bounds-check-zero-atomic.msl index daaa079233..a191700235 100644 --- a/tests/out/msl/bounds-check-zero-atomic.msl +++ b/tests/out/msl/bounds-check-zero-atomic.msl @@ -24,6 +24,10 @@ struct Globals { type_2 c; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + uint fetch_add_atomic( device Globals& globals, constant _mslBufferSizes& _buffer_sizes diff --git a/tests/out/msl/bounds-check-zero.msl b/tests/out/msl/bounds-check-zero.msl index 816983d98b..9dc96a6237 100644 --- a/tests/out/msl/bounds-check-zero.msl +++ b/tests/out/msl/bounds-check-zero.msl @@ -26,6 +26,10 @@ struct Globals { type_4 d; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + float index_array( int i, device Globals const& globals, diff --git a/tests/out/msl/break-if.msl b/tests/out/msl/break-if.msl index 3a6f2e9bff..18b8a8ff98 100644 --- a/tests/out/msl/break-if.msl +++ b/tests/out/msl/break-if.msl @@ -5,6 +5,10 @@ using metal::uint; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + void breakIfEmpty( ) { bool loop_init = true; diff --git a/tests/out/msl/collatz.msl b/tests/out/msl/collatz.msl index 88f9521a27..04b902177a 100644 --- a/tests/out/msl/collatz.msl +++ b/tests/out/msl/collatz.msl @@ -13,6 +13,10 @@ struct PrimeIndices { type_1 data; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + uint collatz_iterations( uint n_base ) { diff --git a/tests/out/msl/constructors.msl b/tests/out/msl/constructors.msl index b3d3b9dd43..68face5ebb 100644 --- a/tests/out/msl/constructors.msl +++ b/tests/out/msl/constructors.msl @@ -30,6 +30,10 @@ constant type_10 cz6_ = type_10 {}; constant Foo cz7_ = Foo {}; constant type_11 cp3_ = type_11 {0, 1, 2, 3}; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + kernel void main_( ) { Foo foo = {}; diff --git a/tests/out/msl/control-flow.msl b/tests/out/msl/control-flow.msl index be396e23a8..75312986e7 100644 --- a/tests/out/msl/control-flow.msl +++ b/tests/out/msl/control-flow.msl @@ -5,6 +5,10 @@ using metal::uint; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + void switch_default_break( int i ) { diff --git a/tests/out/msl/do-while.msl b/tests/out/msl/do-while.msl index 4f928d1f3c..8f6c56388e 100644 --- a/tests/out/msl/do-while.msl +++ b/tests/out/msl/do-while.msl @@ -5,6 +5,10 @@ using metal::uint; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + void fb1_( thread bool& cond ) { diff --git a/tests/out/msl/dualsource.msl b/tests/out/msl/dualsource.msl index 92b1909f8b..c5a49ce558 100644 --- a/tests/out/msl/dualsource.msl +++ b/tests/out/msl/dualsource.msl @@ -9,6 +9,10 @@ struct FragmentOutput { metal::float4 mask; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + struct main_Input { }; struct main_Output { diff --git a/tests/out/msl/empty-global-name.msl b/tests/out/msl/empty-global-name.msl index f456649bc6..4c64c79aae 100644 --- a/tests/out/msl/empty-global-name.msl +++ b/tests/out/msl/empty-global-name.msl @@ -8,6 +8,10 @@ struct type_1 { int member; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + void function( device type_1& unnamed ) { diff --git a/tests/out/msl/empty.msl b/tests/out/msl/empty.msl index 4f8bf9f5e9..724e63aa0a 100644 --- a/tests/out/msl/empty.msl +++ b/tests/out/msl/empty.msl @@ -5,6 +5,10 @@ using metal::uint; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + kernel void main_( ) { return; diff --git a/tests/out/msl/extra.msl b/tests/out/msl/extra.msl index 8288dfad92..40272282cd 100644 --- a/tests/out/msl/extra.msl +++ b/tests/out/msl/extra.msl @@ -14,6 +14,10 @@ struct FragmentIn { uint primitive_index; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + struct main_Input { metal::float4 color [[user(loc0), center_perspective]]; }; diff --git a/tests/out/msl/fragment-output.msl b/tests/out/msl/fragment-output.msl index 4d25809e4f..6ed78d6bb1 100644 --- a/tests/out/msl/fragment-output.msl +++ b/tests/out/msl/fragment-output.msl @@ -21,6 +21,10 @@ struct FragmentOutputVec2Scalar { uint scalaru; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + struct main_vec4vec3_Output { metal::float4 vec4f [[color(0)]]; metal::int4 vec4i [[color(1)]]; diff --git a/tests/out/msl/functions.msl b/tests/out/msl/functions.msl index 31f57c656e..abd1ff5718 100644 --- a/tests/out/msl/functions.msl +++ b/tests/out/msl/functions.msl @@ -5,6 +5,10 @@ using metal::uint; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + metal::float2 test_fma( ) { metal::float2 a = metal::float2(2.0, 2.0); diff --git a/tests/out/msl/globals.msl b/tests/out/msl/globals.msl index d9142c1990..24ccdb73fa 100644 --- a/tests/out/msl/globals.msl +++ b/tests/out/msl/globals.msl @@ -33,6 +33,10 @@ struct type_15 { }; constant bool Foo_1 = true; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + void test_msl_packed_vec3_as_arg( metal::float3 arg ) { diff --git a/tests/out/msl/image.msl b/tests/out/msl/image.msl index e390c2e0fc..d59862f9b8 100644 --- a/tests/out/msl/image.msl +++ b/tests/out/msl/image.msl @@ -5,6 +5,10 @@ using metal::uint; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + struct main_Input { }; kernel void main_( diff --git a/tests/out/msl/interface.msl b/tests/out/msl/interface.msl index c9b93d86a6..711e021743 100644 --- a/tests/out/msl/interface.msl +++ b/tests/out/msl/interface.msl @@ -23,6 +23,10 @@ struct Input2_ { uint index; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + struct vertex_Input { uint color [[attribute(10)]]; }; diff --git a/tests/out/msl/interpolate.msl b/tests/out/msl/interpolate.msl index 5d8b67111e..999edbd9e2 100644 --- a/tests/out/msl/interpolate.msl +++ b/tests/out/msl/interpolate.msl @@ -15,6 +15,10 @@ struct FragmentInput { float perspective_sample; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + struct vert_mainOutput { metal::float4 position [[position]]; uint _flat [[user(loc0), flat]]; diff --git a/tests/out/msl/math-functions.msl b/tests/out/msl/math-functions.msl index 14824ec30f..e2e148e6c6 100644 --- a/tests/out/msl/math-functions.msl +++ b/tests/out/msl/math-functions.msl @@ -55,6 +55,10 @@ __frexp_result_vec4_f32_ naga_frexp(metal::float4 arg) { return __frexp_result_vec4_f32_{ fract, other }; } +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + fragment void main_( ) { metal::float4 v = metal::float4(0.0); @@ -64,14 +68,15 @@ fragment void main_( metal::float4 d = ((v) * 0.017453292519943295474); metal::float4 e = metal::saturate(v); metal::float4 g = metal::refract(v, v, 1.0); + int h = naga_isign(-1); int const_dot = ( + metal::int2 {}.x * metal::int2 {}.x + metal::int2 {}.y * metal::int2 {}.y); - uint _e13 = metal::abs(0u); - uint first_leading_bit_abs = metal::select(31 - metal::clz(_e13), uint(-1), _e13 == 0 || _e13 == -1); + uint _e15 = metal::abs(0u); + uint first_leading_bit_abs = metal::select(31 - metal::clz(_e15), uint(-1), _e15 == 0 || _e15 == -1); int flb_a = metal::select(31 - metal::clz(metal::select(-1, ~-1, -1 < 0)), int(-1), -1 == 0 || -1 == -1); - metal::int2 _e18 = metal::int2(-1); - metal::int2 flb_b = metal::select(31 - metal::clz(metal::select(_e18, ~_e18, _e18 < 0)), int2(-1), _e18 == 0 || _e18 == -1); - metal::uint2 _e21 = metal::uint2(1u); - metal::uint2 flb_c = metal::select(31 - metal::clz(_e21), uint2(-1), _e21 == 0 || _e21 == -1); + metal::int2 _e20 = metal::int2(-1); + metal::int2 flb_b = metal::select(31 - metal::clz(metal::select(_e20, ~_e20, _e20 < 0)), int2(-1), _e20 == 0 || _e20 == -1); + metal::uint2 _e23 = metal::uint2(1u); + metal::uint2 flb_c = metal::select(31 - metal::clz(_e23), uint2(-1), _e23 == 0 || _e23 == -1); int ftb_a = (((metal::ctz(-1) + 1) % 33) - 1); uint ftb_b = (((metal::ctz(1u) + 1) % 33) - 1); metal::int2 ftb_c = (((metal::ctz(metal::int2(-1)) + 1) % 33) - 1); diff --git a/tests/out/msl/operators.msl b/tests/out/msl/operators.msl index de3e990db2..76ef5ebbe2 100644 --- a/tests/out/msl/operators.msl +++ b/tests/out/msl/operators.msl @@ -9,6 +9,10 @@ constant metal::float4 v_f32_zero = metal::float4(0.0, 0.0, 0.0, 0.0); constant metal::float4 v_f32_half = metal::float4(0.5, 0.5, 0.5, 0.5); constant metal::int4 v_i32_one = metal::int4(1, 1, 1, 1); +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + metal::float4 builtins( ) { int s1_ = true ? 1 : 0; diff --git a/tests/out/msl/padding.msl b/tests/out/msl/padding.msl index 4d99bb4c4c..0aa7659414 100644 --- a/tests/out/msl/padding.msl +++ b/tests/out/msl/padding.msl @@ -23,6 +23,10 @@ struct Test3_ { float b; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + struct vertex_Output { metal::float4 member [[position]]; }; diff --git a/tests/out/msl/policy-mix.msl b/tests/out/msl/policy-mix.msl index f6a4fe5d6d..2b05111c59 100644 --- a/tests/out/msl/policy-mix.msl +++ b/tests/out/msl/policy-mix.msl @@ -32,6 +32,10 @@ struct type_9 { metal::float4 inner[2]; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + metal::float4 mock_function( metal::int2 c, int i, diff --git a/tests/out/msl/quad-vert.msl b/tests/out/msl/quad-vert.msl index d322cbcb9b..370bfc8036 100644 --- a/tests/out/msl/quad-vert.msl +++ b/tests/out/msl/quad-vert.msl @@ -18,6 +18,10 @@ struct type_9 { metal::float4 gl_Position; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + void main_1( thread metal::float2& v_uv, thread metal::float2& a_uv_1, diff --git a/tests/out/msl/quad.msl b/tests/out/msl/quad.msl index 5fa5788aac..368c89d179 100644 --- a/tests/out/msl/quad.msl +++ b/tests/out/msl/quad.msl @@ -10,6 +10,10 @@ struct VertexOutput { }; constant float c_scale = 1.2; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + struct vert_mainInput { metal::float2 pos [[attribute(0)]]; metal::float2 uv [[attribute(1)]]; diff --git a/tests/out/msl/ray-query.msl b/tests/out/msl/ray-query.msl index 0d4560f313..f773cade66 100644 --- a/tests/out/msl/ray-query.msl +++ b/tests/out/msl/ray-query.msl @@ -41,6 +41,10 @@ struct RayDesc { metal::float3 dir; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + metal::float3 get_torus_normal( metal::float3 world_point, RayIntersection intersection diff --git a/tests/out/msl/resource-binding-map.msl b/tests/out/msl/resource-binding-map.msl index b4a53d97b5..f53f6c8423 100644 --- a/tests/out/msl/resource-binding-map.msl +++ b/tests/out/msl/resource-binding-map.msl @@ -11,6 +11,10 @@ struct DefaultConstructible { }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + struct entry_point_oneInput { }; struct entry_point_oneOutput { diff --git a/tests/out/msl/shadow.msl b/tests/out/msl/shadow.msl index 53f320344a..0c06512c03 100644 --- a/tests/out/msl/shadow.msl +++ b/tests/out/msl/shadow.msl @@ -33,6 +33,10 @@ struct type_7 { constant metal::float3 c_ambient = metal::float3(0.05, 0.05, 0.05); constant uint c_max_lights = 10u; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + float fetch_shadow( uint light_id, metal::float4 homogeneous_coords, diff --git a/tests/out/msl/skybox.msl b/tests/out/msl/skybox.msl index 7b10ea23e7..cacc687a9a 100644 --- a/tests/out/msl/skybox.msl +++ b/tests/out/msl/skybox.msl @@ -13,6 +13,10 @@ struct Data { metal::float4x4 view; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + struct vs_mainInput { }; struct vs_mainOutput { diff --git a/tests/out/msl/standard.msl b/tests/out/msl/standard.msl index f02243eaac..f1d5e10d0b 100644 --- a/tests/out/msl/standard.msl +++ b/tests/out/msl/standard.msl @@ -5,6 +5,10 @@ using metal::uint; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + bool test_any_and_all_for_bool( ) { return true; diff --git a/tests/out/msl/texture-arg.msl b/tests/out/msl/texture-arg.msl index 5fb9b25649..b11d2f9caa 100644 --- a/tests/out/msl/texture-arg.msl +++ b/tests/out/msl/texture-arg.msl @@ -5,6 +5,10 @@ using metal::uint; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + metal::float4 test( metal::texture2d Passed_Texture, metal::sampler Passed_Sampler diff --git a/tests/out/msl/workgroup-uniform-load.msl b/tests/out/msl/workgroup-uniform-load.msl index 37a8781739..28a0df7ac5 100644 --- a/tests/out/msl/workgroup-uniform-load.msl +++ b/tests/out/msl/workgroup-uniform-load.msl @@ -9,6 +9,10 @@ struct type_2 { }; constant uint SIZE = 128u; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + struct test_workgroupUniformLoadInput { }; kernel void test_workgroupUniformLoad( diff --git a/tests/out/msl/workgroup-var-init.msl b/tests/out/msl/workgroup-var-init.msl index ac300d4337..ba495894fd 100644 --- a/tests/out/msl/workgroup-var-init.msl +++ b/tests/out/msl/workgroup-var-init.msl @@ -19,6 +19,10 @@ struct WStruct { type_4 atom_arr; }; +template inline T naga_isign(T arg) { + return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); +} + kernel void main_( metal::uint3 __local_invocation_id [[thread_position_in_threadgroup]] , threadgroup WStruct& w_mem diff --git a/tests/out/spv/math-functions.spvasm b/tests/out/spv/math-functions.spvasm index 260c3b4bd4..50fcc08a93 100644 --- a/tests/out/spv/math-functions.spvasm +++ b/tests/out/spv/math-functions.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 127 +; Bound: 128 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 @@ -32,13 +32,13 @@ OpMemberDecorate %13 1 Offset 16 %16 = OpTypeFunction %2 %17 = OpConstant %4 1.0 %18 = OpConstant %4 0.0 -%19 = OpConstantNull %5 -%20 = OpTypeInt 32 0 -%21 = OpConstant %20 0 -%22 = OpConstant %6 -1 -%23 = OpConstant %20 1 +%19 = OpConstant %6 -1 +%20 = OpConstantNull %5 +%21 = OpTypeInt 32 0 +%22 = OpConstant %21 0 +%23 = OpConstant %21 1 %24 = OpConstant %6 0 -%25 = OpConstant %20 4294967295 +%25 = OpConstant %21 4294967295 %26 = OpConstant %6 1 %27 = OpConstant %6 2 %28 = OpConstant %4 2.0 @@ -47,12 +47,12 @@ OpMemberDecorate %13 1 Offset 16 %31 = OpConstant %4 1.5 %39 = OpConstantComposite %3 %18 %18 %18 %18 %40 = OpConstantComposite %3 %17 %17 %17 %17 -%43 = OpConstantNull %6 -%56 = OpTypeVector %20 2 -%66 = OpConstant %20 32 -%76 = OpConstantComposite %56 %66 %66 -%88 = OpConstant %6 31 -%94 = OpConstantComposite %5 %88 %88 +%44 = OpConstantNull %6 +%57 = OpTypeVector %21 2 +%67 = OpConstant %21 32 +%77 = OpConstantComposite %57 %67 %67 +%89 = OpConstant %6 31 +%95 = OpConstantComposite %5 %89 %89 %15 = OpFunction %2 None %16 %14 = OpLabel OpBranch %32 @@ -64,84 +64,85 @@ OpBranch %32 %37 = OpExtInst %3 %1 Radians %33 %38 = OpExtInst %3 %1 FClamp %33 %39 %40 %41 = OpExtInst %3 %1 Refract %33 %33 %17 -%44 = OpCompositeExtract %6 %19 0 -%45 = OpCompositeExtract %6 %19 0 -%46 = OpIMul %6 %44 %45 -%47 = OpIAdd %6 %43 %46 -%48 = OpCompositeExtract %6 %19 1 -%49 = OpCompositeExtract %6 %19 1 -%50 = OpIMul %6 %48 %49 -%42 = OpIAdd %6 %47 %50 -%51 = OpCopyObject %20 %21 -%52 = OpExtInst %20 %1 FindUMsb %51 -%53 = OpExtInst %6 %1 FindSMsb %22 -%54 = OpCompositeConstruct %5 %22 %22 -%55 = OpExtInst %5 %1 FindSMsb %54 -%57 = OpCompositeConstruct %56 %23 %23 -%58 = OpExtInst %56 %1 FindUMsb %57 -%59 = OpExtInst %6 %1 FindILsb %22 -%60 = OpExtInst %20 %1 FindILsb %23 -%61 = OpCompositeConstruct %5 %22 %22 -%62 = OpExtInst %5 %1 FindILsb %61 -%63 = OpCompositeConstruct %56 %23 %23 -%64 = OpExtInst %56 %1 FindILsb %63 -%67 = OpExtInst %20 %1 FindILsb %21 -%65 = OpExtInst %20 %1 UMin %66 %67 -%69 = OpExtInst %6 %1 FindILsb %24 -%68 = OpExtInst %6 %1 UMin %66 %69 -%71 = OpExtInst %20 %1 FindILsb %25 -%70 = OpExtInst %20 %1 UMin %66 %71 -%73 = OpExtInst %6 %1 FindILsb %22 -%72 = OpExtInst %6 %1 UMin %66 %73 -%74 = OpCompositeConstruct %56 %21 %21 -%77 = OpExtInst %56 %1 FindILsb %74 -%75 = OpExtInst %56 %1 UMin %76 %77 -%78 = OpCompositeConstruct %5 %24 %24 -%80 = OpExtInst %5 %1 FindILsb %78 -%79 = OpExtInst %5 %1 UMin %76 %80 -%81 = OpCompositeConstruct %56 %23 %23 -%83 = OpExtInst %56 %1 FindILsb %81 -%82 = OpExtInst %56 %1 UMin %76 %83 -%84 = OpCompositeConstruct %5 %26 %26 -%86 = OpExtInst %5 %1 FindILsb %84 -%85 = OpExtInst %5 %1 UMin %76 %86 -%89 = OpExtInst %6 %1 FindUMsb %22 -%87 = OpISub %6 %88 %89 -%91 = OpExtInst %6 %1 FindUMsb %23 -%90 = OpISub %20 %88 %91 -%92 = OpCompositeConstruct %5 %22 %22 -%95 = OpExtInst %5 %1 FindUMsb %92 -%93 = OpISub %5 %94 %95 -%96 = OpCompositeConstruct %56 %23 %23 -%98 = OpExtInst %5 %1 FindUMsb %96 -%97 = OpISub %56 %94 %98 -%99 = OpExtInst %4 %1 Ldexp %17 %27 -%100 = OpCompositeConstruct %7 %17 %28 -%101 = OpCompositeConstruct %5 %29 %30 -%102 = OpExtInst %7 %1 Ldexp %100 %101 -%103 = OpExtInst %8 %1 ModfStruct %31 +%42 = OpExtInst %6 %1 SSign %19 +%45 = OpCompositeExtract %6 %20 0 +%46 = OpCompositeExtract %6 %20 0 +%47 = OpIMul %6 %45 %46 +%48 = OpIAdd %6 %44 %47 +%49 = OpCompositeExtract %6 %20 1 +%50 = OpCompositeExtract %6 %20 1 +%51 = OpIMul %6 %49 %50 +%43 = OpIAdd %6 %48 %51 +%52 = OpCopyObject %21 %22 +%53 = OpExtInst %21 %1 FindUMsb %52 +%54 = OpExtInst %6 %1 FindSMsb %19 +%55 = OpCompositeConstruct %5 %19 %19 +%56 = OpExtInst %5 %1 FindSMsb %55 +%58 = OpCompositeConstruct %57 %23 %23 +%59 = OpExtInst %57 %1 FindUMsb %58 +%60 = OpExtInst %6 %1 FindILsb %19 +%61 = OpExtInst %21 %1 FindILsb %23 +%62 = OpCompositeConstruct %5 %19 %19 +%63 = OpExtInst %5 %1 FindILsb %62 +%64 = OpCompositeConstruct %57 %23 %23 +%65 = OpExtInst %57 %1 FindILsb %64 +%68 = OpExtInst %21 %1 FindILsb %22 +%66 = OpExtInst %21 %1 UMin %67 %68 +%70 = OpExtInst %6 %1 FindILsb %24 +%69 = OpExtInst %6 %1 UMin %67 %70 +%72 = OpExtInst %21 %1 FindILsb %25 +%71 = OpExtInst %21 %1 UMin %67 %72 +%74 = OpExtInst %6 %1 FindILsb %19 +%73 = OpExtInst %6 %1 UMin %67 %74 +%75 = OpCompositeConstruct %57 %22 %22 +%78 = OpExtInst %57 %1 FindILsb %75 +%76 = OpExtInst %57 %1 UMin %77 %78 +%79 = OpCompositeConstruct %5 %24 %24 +%81 = OpExtInst %5 %1 FindILsb %79 +%80 = OpExtInst %5 %1 UMin %77 %81 +%82 = OpCompositeConstruct %57 %23 %23 +%84 = OpExtInst %57 %1 FindILsb %82 +%83 = OpExtInst %57 %1 UMin %77 %84 +%85 = OpCompositeConstruct %5 %26 %26 +%87 = OpExtInst %5 %1 FindILsb %85 +%86 = OpExtInst %5 %1 UMin %77 %87 +%90 = OpExtInst %6 %1 FindUMsb %19 +%88 = OpISub %6 %89 %90 +%92 = OpExtInst %6 %1 FindUMsb %23 +%91 = OpISub %21 %89 %92 +%93 = OpCompositeConstruct %5 %19 %19 +%96 = OpExtInst %5 %1 FindUMsb %93 +%94 = OpISub %5 %95 %96 +%97 = OpCompositeConstruct %57 %23 %23 +%99 = OpExtInst %5 %1 FindUMsb %97 +%98 = OpISub %57 %95 %99 +%100 = OpExtInst %4 %1 Ldexp %17 %27 +%101 = OpCompositeConstruct %7 %17 %28 +%102 = OpCompositeConstruct %5 %29 %30 +%103 = OpExtInst %7 %1 Ldexp %101 %102 %104 = OpExtInst %8 %1 ModfStruct %31 -%105 = OpCompositeExtract %4 %104 0 -%106 = OpExtInst %8 %1 ModfStruct %31 -%107 = OpCompositeExtract %4 %106 1 -%108 = OpCompositeConstruct %7 %31 %31 -%109 = OpExtInst %9 %1 ModfStruct %108 -%110 = OpCompositeConstruct %3 %31 %31 %31 %31 -%111 = OpExtInst %10 %1 ModfStruct %110 -%112 = OpCompositeExtract %3 %111 1 -%113 = OpCompositeExtract %4 %112 0 -%114 = OpCompositeConstruct %7 %31 %31 -%115 = OpExtInst %9 %1 ModfStruct %114 -%116 = OpCompositeExtract %7 %115 0 -%117 = OpCompositeExtract %4 %116 1 -%118 = OpExtInst %11 %1 FrexpStruct %31 +%105 = OpExtInst %8 %1 ModfStruct %31 +%106 = OpCompositeExtract %4 %105 0 +%107 = OpExtInst %8 %1 ModfStruct %31 +%108 = OpCompositeExtract %4 %107 1 +%109 = OpCompositeConstruct %7 %31 %31 +%110 = OpExtInst %9 %1 ModfStruct %109 +%111 = OpCompositeConstruct %3 %31 %31 %31 %31 +%112 = OpExtInst %10 %1 ModfStruct %111 +%113 = OpCompositeExtract %3 %112 1 +%114 = OpCompositeExtract %4 %113 0 +%115 = OpCompositeConstruct %7 %31 %31 +%116 = OpExtInst %9 %1 ModfStruct %115 +%117 = OpCompositeExtract %7 %116 0 +%118 = OpCompositeExtract %4 %117 1 %119 = OpExtInst %11 %1 FrexpStruct %31 -%120 = OpCompositeExtract %4 %119 0 -%121 = OpExtInst %11 %1 FrexpStruct %31 -%122 = OpCompositeExtract %6 %121 1 -%123 = OpCompositeConstruct %3 %31 %31 %31 %31 -%124 = OpExtInst %13 %1 FrexpStruct %123 -%125 = OpCompositeExtract %12 %124 1 -%126 = OpCompositeExtract %6 %125 0 +%120 = OpExtInst %11 %1 FrexpStruct %31 +%121 = OpCompositeExtract %4 %120 0 +%122 = OpExtInst %11 %1 FrexpStruct %31 +%123 = OpCompositeExtract %6 %122 1 +%124 = OpCompositeConstruct %3 %31 %31 %31 %31 +%125 = OpExtInst %13 %1 FrexpStruct %124 +%126 = OpCompositeExtract %12 %125 1 +%127 = OpCompositeExtract %6 %126 0 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/math-functions.wgsl b/tests/out/wgsl/math-functions.wgsl index 149ebff8e0..488eadf3e9 100644 --- a/tests/out/wgsl/math-functions.wgsl +++ b/tests/out/wgsl/math-functions.wgsl @@ -7,6 +7,7 @@ fn main() { let d = radians(v); let e = saturate(v); let g = refract(v, v, 1.0); + let h = sign(-1); let const_dot = dot(vec2(), vec2()); let first_leading_bit_abs = firstLeadingBit(abs(0u)); let flb_a = firstLeadingBit(-1); From d16e605aa6f58043cef24b6cb8204bcf0dfb3cab Mon Sep 17 00:00:00 2001 From: Fredrik Fornwall Date: Wed, 6 Sep 2023 08:09:19 +0200 Subject: [PATCH 2/4] Use TypeInner::scalar_kind() --- src/back/msl/writer.rs | 11 ++--------- 1 file changed, 2 insertions(+), 9 deletions(-) diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index f60d9f7b3c..de0532f400 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -1715,15 +1715,8 @@ impl Writer { Mf::Reflect => "reflect", Mf::Refract => "refract", // computational - Mf::Sign => match arg_type { - &crate::TypeInner::Scalar { - kind: crate::ScalarKind::Sint, - .. - } - | &crate::TypeInner::Vector { - kind: crate::ScalarKind::Sint, - .. - } => ISIGN_FUNCTION, + Mf::Sign => match arg_type.scalar_kind() { + Some(crate::ScalarKind::Sint) => ISIGN_FUNCTION, _ => "sign", }, Mf::Fma => "fma", From f6a2a918d48176cf255fd52b2ed85686950b07c3 Mon Sep 17 00:00:00 2001 From: Fredrik Fornwall Date: Wed, 6 Sep 2023 08:13:35 +0200 Subject: [PATCH 3/4] Test more sign() variants --- tests/in/math-functions.wgsl | 5 +- .../glsl/math-functions.main.Fragment.glsl | 9 +- tests/out/hlsl/math-functions.hlsl | 9 +- tests/out/msl/math-functions.msl | 17 +- tests/out/spv/math-functions.spvasm | 222 +++++++++--------- tests/out/wgsl/math-functions.wgsl | 5 +- 6 files changed, 144 insertions(+), 123 deletions(-) diff --git a/tests/in/math-functions.wgsl b/tests/in/math-functions.wgsl index 6e7bccf7f6..d08e76e4f2 100644 --- a/tests/in/math-functions.wgsl +++ b/tests/in/math-functions.wgsl @@ -8,7 +8,10 @@ fn main() { let d = radians(v); let e = saturate(v); let g = refract(v, v, f); - let h = sign(-1); + let sign_a = sign(-1); + let sign_b = sign(vec4(-1)); + let sign_c = sign(-1.0); + let sign_d = sign(vec4(-1.0)); let const_dot = dot(vec2(), vec2()); let first_leading_bit_abs = firstLeadingBit(abs(0u)); let flb_a = firstLeadingBit(-1); diff --git a/tests/out/glsl/math-functions.main.Fragment.glsl b/tests/out/glsl/math-functions.main.Fragment.glsl index 595235bdd1..37c072a6fa 100644 --- a/tests/out/glsl/math-functions.main.Fragment.glsl +++ b/tests/out/glsl/math-functions.main.Fragment.glsl @@ -62,7 +62,10 @@ void main() { vec4 d = radians(v); vec4 e = clamp(v, vec4(0.0), vec4(1.0)); vec4 g = refract(v, v, 1.0); - int h = sign(-1); + int sign_a = sign(-1); + ivec4 sign_b = sign(ivec4(-1)); + float sign_c = sign(-1.0); + vec4 sign_d = sign(vec4(-1.0)); int const_dot = ( + ivec2(0).x * ivec2(0).x + ivec2(0).y * ivec2(0).y); uint first_leading_bit_abs = uint(findMSB(uint(abs(int(0u))))); int flb_a = findMSB(-1); @@ -82,8 +85,8 @@ void main() { ivec2 ctz_h = ivec2(min(uvec2(findLSB(ivec2(1))), uvec2(32u))); int clz_a = (-1 < 0 ? 0 : 31 - findMSB(-1)); uint clz_b = uint(31 - findMSB(1u)); - ivec2 _e60 = ivec2(-1); - ivec2 clz_c = mix(ivec2(31) - findMSB(_e60), ivec2(0), lessThan(_e60, ivec2(0))); + ivec2 _e68 = ivec2(-1); + ivec2 clz_c = mix(ivec2(31) - findMSB(_e68), ivec2(0), lessThan(_e68, ivec2(0))); uvec2 clz_d = uvec2(ivec2(31) - findMSB(uvec2(1u))); float lde_a = ldexp(1.0, 2); vec2 lde_b = ldexp(vec2(1.0, 2.0), ivec2(3, 4)); diff --git a/tests/out/hlsl/math-functions.hlsl b/tests/out/hlsl/math-functions.hlsl index 6a2e4e6787..fc5cadb65e 100644 --- a/tests/out/hlsl/math-functions.hlsl +++ b/tests/out/hlsl/math-functions.hlsl @@ -72,7 +72,10 @@ void main() float4 d = radians(v); float4 e = saturate(v); float4 g = refract(v, v, 1.0); - int h = sign(-1); + int sign_a = sign(-1); + int4 sign_b = sign((-1).xxxx); + float sign_c = sign(-1.0); + float4 sign_d = sign((-1.0).xxxx); int const_dot = dot((int2)0, (int2)0); uint first_leading_bit_abs = firstbithigh(abs(0u)); int flb_a = asint(firstbithigh(-1)); @@ -92,8 +95,8 @@ void main() int2 ctz_h = asint(min((32u).xx, firstbitlow((1).xx))); int clz_a = (-1 < 0 ? 0 : 31 - asint(firstbithigh(-1))); uint clz_b = (31u - firstbithigh(1u)); - int2 _expr60 = (-1).xx; - int2 clz_c = (_expr60 < (0).xx ? (0).xx : (31).xx - asint(firstbithigh(_expr60))); + int2 _expr68 = (-1).xx; + int2 clz_c = (_expr68 < (0).xx ? (0).xx : (31).xx - asint(firstbithigh(_expr68))); uint2 clz_d = ((31u).xx - firstbithigh((1u).xx)); float lde_a = ldexp(1.0, 2); float2 lde_b = ldexp(float2(1.0, 2.0), int2(3, 4)); diff --git a/tests/out/msl/math-functions.msl b/tests/out/msl/math-functions.msl index e2e148e6c6..7476c51f93 100644 --- a/tests/out/msl/math-functions.msl +++ b/tests/out/msl/math-functions.msl @@ -68,15 +68,18 @@ fragment void main_( metal::float4 d = ((v) * 0.017453292519943295474); metal::float4 e = metal::saturate(v); metal::float4 g = metal::refract(v, v, 1.0); - int h = naga_isign(-1); + int sign_a = naga_isign(-1); + metal::int4 sign_b = naga_isign(metal::int4(-1)); + float sign_c = metal::sign(-1.0); + metal::float4 sign_d = metal::sign(metal::float4(-1.0)); int const_dot = ( + metal::int2 {}.x * metal::int2 {}.x + metal::int2 {}.y * metal::int2 {}.y); - uint _e15 = metal::abs(0u); - uint first_leading_bit_abs = metal::select(31 - metal::clz(_e15), uint(-1), _e15 == 0 || _e15 == -1); + uint _e23 = metal::abs(0u); + uint first_leading_bit_abs = metal::select(31 - metal::clz(_e23), uint(-1), _e23 == 0 || _e23 == -1); int flb_a = metal::select(31 - metal::clz(metal::select(-1, ~-1, -1 < 0)), int(-1), -1 == 0 || -1 == -1); - metal::int2 _e20 = metal::int2(-1); - metal::int2 flb_b = metal::select(31 - metal::clz(metal::select(_e20, ~_e20, _e20 < 0)), int2(-1), _e20 == 0 || _e20 == -1); - metal::uint2 _e23 = metal::uint2(1u); - metal::uint2 flb_c = metal::select(31 - metal::clz(_e23), uint2(-1), _e23 == 0 || _e23 == -1); + metal::int2 _e28 = metal::int2(-1); + metal::int2 flb_b = metal::select(31 - metal::clz(metal::select(_e28, ~_e28, _e28 < 0)), int2(-1), _e28 == 0 || _e28 == -1); + metal::uint2 _e31 = metal::uint2(1u); + metal::uint2 flb_c = metal::select(31 - metal::clz(_e31), uint2(-1), _e31 == 0 || _e31 == -1); int ftb_a = (((metal::ctz(-1) + 1) % 33) - 1); uint ftb_b = (((metal::ctz(1u) + 1) % 33) - 1); metal::int2 ftb_c = (((metal::ctz(metal::int2(-1)) + 1) % 33) - 1); diff --git a/tests/out/spv/math-functions.spvasm b/tests/out/spv/math-functions.spvasm index 50fcc08a93..7edb0e26b4 100644 --- a/tests/out/spv/math-functions.spvasm +++ b/tests/out/spv/math-functions.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 128 +; Bound: 134 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 @@ -33,116 +33,122 @@ OpMemberDecorate %13 1 Offset 16 %17 = OpConstant %4 1.0 %18 = OpConstant %4 0.0 %19 = OpConstant %6 -1 -%20 = OpConstantNull %5 -%21 = OpTypeInt 32 0 -%22 = OpConstant %21 0 -%23 = OpConstant %21 1 -%24 = OpConstant %6 0 -%25 = OpConstant %21 4294967295 -%26 = OpConstant %6 1 -%27 = OpConstant %6 2 -%28 = OpConstant %4 2.0 -%29 = OpConstant %6 3 -%30 = OpConstant %6 4 -%31 = OpConstant %4 1.5 -%39 = OpConstantComposite %3 %18 %18 %18 %18 -%40 = OpConstantComposite %3 %17 %17 %17 %17 -%44 = OpConstantNull %6 -%57 = OpTypeVector %21 2 -%67 = OpConstant %21 32 -%77 = OpConstantComposite %57 %67 %67 -%89 = OpConstant %6 31 -%95 = OpConstantComposite %5 %89 %89 +%20 = OpConstant %4 -1.0 +%21 = OpConstantNull %5 +%22 = OpTypeInt 32 0 +%23 = OpConstant %22 0 +%24 = OpConstant %22 1 +%25 = OpConstant %6 0 +%26 = OpConstant %22 4294967295 +%27 = OpConstant %6 1 +%28 = OpConstant %6 2 +%29 = OpConstant %4 2.0 +%30 = OpConstant %6 3 +%31 = OpConstant %6 4 +%32 = OpConstant %4 1.5 +%40 = OpConstantComposite %3 %18 %18 %18 %18 +%41 = OpConstantComposite %3 %17 %17 %17 %17 +%50 = OpConstantNull %6 +%63 = OpTypeVector %22 2 +%73 = OpConstant %22 32 +%83 = OpConstantComposite %63 %73 %73 +%95 = OpConstant %6 31 +%101 = OpConstantComposite %5 %95 %95 %15 = OpFunction %2 None %16 %14 = OpLabel -OpBranch %32 -%32 = OpLabel -%33 = OpCompositeConstruct %3 %18 %18 %18 %18 -%34 = OpExtInst %4 %1 Degrees %17 -%35 = OpExtInst %4 %1 Radians %17 -%36 = OpExtInst %3 %1 Degrees %33 -%37 = OpExtInst %3 %1 Radians %33 -%38 = OpExtInst %3 %1 FClamp %33 %39 %40 -%41 = OpExtInst %3 %1 Refract %33 %33 %17 -%42 = OpExtInst %6 %1 SSign %19 -%45 = OpCompositeExtract %6 %20 0 -%46 = OpCompositeExtract %6 %20 0 -%47 = OpIMul %6 %45 %46 -%48 = OpIAdd %6 %44 %47 -%49 = OpCompositeExtract %6 %20 1 -%50 = OpCompositeExtract %6 %20 1 -%51 = OpIMul %6 %49 %50 -%43 = OpIAdd %6 %48 %51 -%52 = OpCopyObject %21 %22 -%53 = OpExtInst %21 %1 FindUMsb %52 -%54 = OpExtInst %6 %1 FindSMsb %19 -%55 = OpCompositeConstruct %5 %19 %19 -%56 = OpExtInst %5 %1 FindSMsb %55 -%58 = OpCompositeConstruct %57 %23 %23 -%59 = OpExtInst %57 %1 FindUMsb %58 -%60 = OpExtInst %6 %1 FindILsb %19 -%61 = OpExtInst %21 %1 FindILsb %23 -%62 = OpCompositeConstruct %5 %19 %19 -%63 = OpExtInst %5 %1 FindILsb %62 -%64 = OpCompositeConstruct %57 %23 %23 -%65 = OpExtInst %57 %1 FindILsb %64 -%68 = OpExtInst %21 %1 FindILsb %22 -%66 = OpExtInst %21 %1 UMin %67 %68 -%70 = OpExtInst %6 %1 FindILsb %24 -%69 = OpExtInst %6 %1 UMin %67 %70 -%72 = OpExtInst %21 %1 FindILsb %25 -%71 = OpExtInst %21 %1 UMin %67 %72 -%74 = OpExtInst %6 %1 FindILsb %19 -%73 = OpExtInst %6 %1 UMin %67 %74 -%75 = OpCompositeConstruct %57 %22 %22 -%78 = OpExtInst %57 %1 FindILsb %75 -%76 = OpExtInst %57 %1 UMin %77 %78 -%79 = OpCompositeConstruct %5 %24 %24 -%81 = OpExtInst %5 %1 FindILsb %79 -%80 = OpExtInst %5 %1 UMin %77 %81 -%82 = OpCompositeConstruct %57 %23 %23 -%84 = OpExtInst %57 %1 FindILsb %82 -%83 = OpExtInst %57 %1 UMin %77 %84 -%85 = OpCompositeConstruct %5 %26 %26 +OpBranch %33 +%33 = OpLabel +%34 = OpCompositeConstruct %3 %18 %18 %18 %18 +%35 = OpExtInst %4 %1 Degrees %17 +%36 = OpExtInst %4 %1 Radians %17 +%37 = OpExtInst %3 %1 Degrees %34 +%38 = OpExtInst %3 %1 Radians %34 +%39 = OpExtInst %3 %1 FClamp %34 %40 %41 +%42 = OpExtInst %3 %1 Refract %34 %34 %17 +%43 = OpExtInst %6 %1 SSign %19 +%44 = OpCompositeConstruct %12 %19 %19 %19 %19 +%45 = OpExtInst %12 %1 SSign %44 +%46 = OpExtInst %4 %1 FSign %20 +%47 = OpCompositeConstruct %3 %20 %20 %20 %20 +%48 = OpExtInst %3 %1 FSign %47 +%51 = OpCompositeExtract %6 %21 0 +%52 = OpCompositeExtract %6 %21 0 +%53 = OpIMul %6 %51 %52 +%54 = OpIAdd %6 %50 %53 +%55 = OpCompositeExtract %6 %21 1 +%56 = OpCompositeExtract %6 %21 1 +%57 = OpIMul %6 %55 %56 +%49 = OpIAdd %6 %54 %57 +%58 = OpCopyObject %22 %23 +%59 = OpExtInst %22 %1 FindUMsb %58 +%60 = OpExtInst %6 %1 FindSMsb %19 +%61 = OpCompositeConstruct %5 %19 %19 +%62 = OpExtInst %5 %1 FindSMsb %61 +%64 = OpCompositeConstruct %63 %24 %24 +%65 = OpExtInst %63 %1 FindUMsb %64 +%66 = OpExtInst %6 %1 FindILsb %19 +%67 = OpExtInst %22 %1 FindILsb %24 +%68 = OpCompositeConstruct %5 %19 %19 +%69 = OpExtInst %5 %1 FindILsb %68 +%70 = OpCompositeConstruct %63 %24 %24 +%71 = OpExtInst %63 %1 FindILsb %70 +%74 = OpExtInst %22 %1 FindILsb %23 +%72 = OpExtInst %22 %1 UMin %73 %74 +%76 = OpExtInst %6 %1 FindILsb %25 +%75 = OpExtInst %6 %1 UMin %73 %76 +%78 = OpExtInst %22 %1 FindILsb %26 +%77 = OpExtInst %22 %1 UMin %73 %78 +%80 = OpExtInst %6 %1 FindILsb %19 +%79 = OpExtInst %6 %1 UMin %73 %80 +%81 = OpCompositeConstruct %63 %23 %23 +%84 = OpExtInst %63 %1 FindILsb %81 +%82 = OpExtInst %63 %1 UMin %83 %84 +%85 = OpCompositeConstruct %5 %25 %25 %87 = OpExtInst %5 %1 FindILsb %85 -%86 = OpExtInst %5 %1 UMin %77 %87 -%90 = OpExtInst %6 %1 FindUMsb %19 -%88 = OpISub %6 %89 %90 -%92 = OpExtInst %6 %1 FindUMsb %23 -%91 = OpISub %21 %89 %92 -%93 = OpCompositeConstruct %5 %19 %19 -%96 = OpExtInst %5 %1 FindUMsb %93 -%94 = OpISub %5 %95 %96 -%97 = OpCompositeConstruct %57 %23 %23 -%99 = OpExtInst %5 %1 FindUMsb %97 -%98 = OpISub %57 %95 %99 -%100 = OpExtInst %4 %1 Ldexp %17 %27 -%101 = OpCompositeConstruct %7 %17 %28 -%102 = OpCompositeConstruct %5 %29 %30 -%103 = OpExtInst %7 %1 Ldexp %101 %102 -%104 = OpExtInst %8 %1 ModfStruct %31 -%105 = OpExtInst %8 %1 ModfStruct %31 -%106 = OpCompositeExtract %4 %105 0 -%107 = OpExtInst %8 %1 ModfStruct %31 -%108 = OpCompositeExtract %4 %107 1 -%109 = OpCompositeConstruct %7 %31 %31 -%110 = OpExtInst %9 %1 ModfStruct %109 -%111 = OpCompositeConstruct %3 %31 %31 %31 %31 -%112 = OpExtInst %10 %1 ModfStruct %111 -%113 = OpCompositeExtract %3 %112 1 -%114 = OpCompositeExtract %4 %113 0 -%115 = OpCompositeConstruct %7 %31 %31 +%86 = OpExtInst %5 %1 UMin %83 %87 +%88 = OpCompositeConstruct %63 %24 %24 +%90 = OpExtInst %63 %1 FindILsb %88 +%89 = OpExtInst %63 %1 UMin %83 %90 +%91 = OpCompositeConstruct %5 %27 %27 +%93 = OpExtInst %5 %1 FindILsb %91 +%92 = OpExtInst %5 %1 UMin %83 %93 +%96 = OpExtInst %6 %1 FindUMsb %19 +%94 = OpISub %6 %95 %96 +%98 = OpExtInst %6 %1 FindUMsb %24 +%97 = OpISub %22 %95 %98 +%99 = OpCompositeConstruct %5 %19 %19 +%102 = OpExtInst %5 %1 FindUMsb %99 +%100 = OpISub %5 %101 %102 +%103 = OpCompositeConstruct %63 %24 %24 +%105 = OpExtInst %5 %1 FindUMsb %103 +%104 = OpISub %63 %101 %105 +%106 = OpExtInst %4 %1 Ldexp %17 %28 +%107 = OpCompositeConstruct %7 %17 %29 +%108 = OpCompositeConstruct %5 %30 %31 +%109 = OpExtInst %7 %1 Ldexp %107 %108 +%110 = OpExtInst %8 %1 ModfStruct %32 +%111 = OpExtInst %8 %1 ModfStruct %32 +%112 = OpCompositeExtract %4 %111 0 +%113 = OpExtInst %8 %1 ModfStruct %32 +%114 = OpCompositeExtract %4 %113 1 +%115 = OpCompositeConstruct %7 %32 %32 %116 = OpExtInst %9 %1 ModfStruct %115 -%117 = OpCompositeExtract %7 %116 0 -%118 = OpCompositeExtract %4 %117 1 -%119 = OpExtInst %11 %1 FrexpStruct %31 -%120 = OpExtInst %11 %1 FrexpStruct %31 -%121 = OpCompositeExtract %4 %120 0 -%122 = OpExtInst %11 %1 FrexpStruct %31 -%123 = OpCompositeExtract %6 %122 1 -%124 = OpCompositeConstruct %3 %31 %31 %31 %31 -%125 = OpExtInst %13 %1 FrexpStruct %124 -%126 = OpCompositeExtract %12 %125 1 -%127 = OpCompositeExtract %6 %126 0 +%117 = OpCompositeConstruct %3 %32 %32 %32 %32 +%118 = OpExtInst %10 %1 ModfStruct %117 +%119 = OpCompositeExtract %3 %118 1 +%120 = OpCompositeExtract %4 %119 0 +%121 = OpCompositeConstruct %7 %32 %32 +%122 = OpExtInst %9 %1 ModfStruct %121 +%123 = OpCompositeExtract %7 %122 0 +%124 = OpCompositeExtract %4 %123 1 +%125 = OpExtInst %11 %1 FrexpStruct %32 +%126 = OpExtInst %11 %1 FrexpStruct %32 +%127 = OpCompositeExtract %4 %126 0 +%128 = OpExtInst %11 %1 FrexpStruct %32 +%129 = OpCompositeExtract %6 %128 1 +%130 = OpCompositeConstruct %3 %32 %32 %32 %32 +%131 = OpExtInst %13 %1 FrexpStruct %130 +%132 = OpCompositeExtract %12 %131 1 +%133 = OpCompositeExtract %6 %132 0 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/math-functions.wgsl b/tests/out/wgsl/math-functions.wgsl index 488eadf3e9..0b20291f57 100644 --- a/tests/out/wgsl/math-functions.wgsl +++ b/tests/out/wgsl/math-functions.wgsl @@ -7,7 +7,10 @@ fn main() { let d = radians(v); let e = saturate(v); let g = refract(v, v, 1.0); - let h = sign(-1); + let sign_a = sign(-1); + let sign_b = sign(vec4(-1)); + let sign_c = sign(-1.0); + let sign_d = sign(vec4(-1.0)); let const_dot = dot(vec2(), vec2()); let first_leading_bit_abs = firstLeadingBit(abs(0u)); let flb_a = firstLeadingBit(-1); From 32e96a560a39f8a90b18e8daa28d7ff7fc480d5e Mon Sep 17 00:00:00 2001 From: Fredrik Fornwall Date: Wed, 6 Sep 2023 08:38:11 +0200 Subject: [PATCH 4/4] [msl-out] Emit sign(i32) code inline --- src/back/msl/keywords.rs | 1 - src/back/msl/writer.rs | 53 ++++++++++++++----- src/valid/expression.rs | 2 +- tests/out/msl/access.msl | 4 -- tests/out/msl/array-in-ctor.msl | 4 -- .../out/msl/array-in-function-return-type.msl | 4 -- tests/out/msl/atomicOps.msl | 4 -- tests/out/msl/binding-arrays.msl | 4 -- tests/out/msl/bitcast.msl | 4 -- tests/out/msl/bits.msl | 4 -- tests/out/msl/boids.msl | 4 -- tests/out/msl/bounds-check-image-restrict.msl | 4 -- tests/out/msl/bounds-check-image-rzsw.msl | 4 -- tests/out/msl/bounds-check-restrict.msl | 4 -- tests/out/msl/bounds-check-zero-atomic.msl | 4 -- tests/out/msl/bounds-check-zero.msl | 4 -- tests/out/msl/break-if.msl | 4 -- tests/out/msl/collatz.msl | 4 -- tests/out/msl/constructors.msl | 4 -- tests/out/msl/control-flow.msl | 4 -- tests/out/msl/do-while.msl | 4 -- tests/out/msl/dualsource.msl | 4 -- tests/out/msl/empty-global-name.msl | 4 -- tests/out/msl/empty.msl | 4 -- tests/out/msl/extra.msl | 4 -- tests/out/msl/fragment-output.msl | 4 -- tests/out/msl/functions.msl | 4 -- tests/out/msl/globals.msl | 4 -- tests/out/msl/image.msl | 4 -- tests/out/msl/interface.msl | 4 -- tests/out/msl/interpolate.msl | 4 -- tests/out/msl/math-functions.msl | 9 ++-- tests/out/msl/operators.msl | 4 -- tests/out/msl/padding.msl | 4 -- tests/out/msl/policy-mix.msl | 4 -- tests/out/msl/quad-vert.msl | 4 -- tests/out/msl/quad.msl | 4 -- tests/out/msl/ray-query.msl | 4 -- tests/out/msl/resource-binding-map.msl | 4 -- tests/out/msl/shadow.msl | 4 -- tests/out/msl/skybox.msl | 4 -- tests/out/msl/standard.msl | 4 -- tests/out/msl/texture-arg.msl | 4 -- tests/out/msl/workgroup-uniform-load.msl | 4 -- tests/out/msl/workgroup-var-init.msl | 4 -- 45 files changed, 43 insertions(+), 186 deletions(-) diff --git a/src/back/msl/keywords.rs b/src/back/msl/keywords.rs index 670e2f85ed..f34b618db8 100644 --- a/src/back/msl/keywords.rs +++ b/src/back/msl/keywords.rs @@ -216,5 +216,4 @@ pub const RESERVED: &[&str] = &[ "clamped_lod_e", super::writer::FREXP_FUNCTION, super::writer::MODF_FUNCTION, - super::writer::ISIGN_FUNCTION, ]; diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index de0532f400..67ab887285 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -34,7 +34,6 @@ const RAY_QUERY_FUN_MAP_INTERSECTION: &str = "_map_intersection_type"; pub(crate) const MODF_FUNCTION: &str = "naga_modf"; pub(crate) const FREXP_FUNCTION: &str = "naga_frexp"; -pub(crate) const ISIGN_FUNCTION: &str = "naga_isign"; /// Write the Metal name for a Naga numeric type: scalar, vector, or matrix. /// @@ -1185,6 +1184,31 @@ impl Writer { Ok(()) } + /// Emit code for the sign(i32) expression. + /// + fn put_isign( + &mut self, + arg: Handle, + context: &ExpressionContext, + ) -> BackendResult { + write!(self.out, "{NAMESPACE}::select({NAMESPACE}::select(")?; + match context.resolve_type(arg) { + &crate::TypeInner::Vector { size, .. } => { + let size = back::vector_size_str(size); + write!(self.out, "int{size}(-1), int{size}(1)")?; + } + _ => { + write!(self.out, "-1, 1")?; + } + } + write!(self.out, ", (")?; + self.put_expression(arg, context, true)?; + write!(self.out, " > 0)), 0, (")?; + self.put_expression(arg, context, true)?; + write!(self.out, " == 0))")?; + Ok(()) + } + fn put_const_expression( &mut self, expr_handle: Handle, @@ -1716,7 +1740,9 @@ impl Writer { Mf::Refract => "refract", // computational Mf::Sign => match arg_type.scalar_kind() { - Some(crate::ScalarKind::Sint) => ISIGN_FUNCTION, + Some(crate::ScalarKind::Sint) => { + return self.put_isign(arg, context); + } _ => "sign", }, Mf::Fma => "fma", @@ -1821,7 +1847,7 @@ impl Writer { write!(self.out, "((")?; self.put_expression(arg, context, false)?; write!(self.out, ") * 57.295779513082322865)")?; - } else if fun == Mf::Modf || fun == Mf::Frexp || fun_name == ISIGN_FUNCTION { + } else if fun == Mf::Modf || fun == Mf::Frexp { write!(self.out, "{fun_name}")?; self.put_call_parameters(iter::once(arg), context)?; } else { @@ -2428,6 +2454,16 @@ impl Writer { crate::MathFunction::FindMsb => { self.need_bake_expressions.insert(arg); } + crate::MathFunction::Sign => { + // WGSL's `sign` function works also on signed ints, but Metal's only + // works on floating points, so we emit inline code for integer `sign` + // calls. But that code uses each argument 2 times (see `put_isign`), + // so to avoid duplicated evaluation, we must bake the argument. + let inner = context.resolve_type(expr_handle); + if inner.scalar_kind() == Some(crate::ScalarKind::Sint) { + self.need_bake_expressions.insert(arg); + } + } _ => {} } } @@ -3096,7 +3132,6 @@ impl Writer { self.write_type_defs(module)?; self.write_global_constants(module)?; - self.write_polyfills()?; self.write_functions(module, info, options, pipeline_options) } @@ -4109,16 +4144,6 @@ impl Writer { Ok(info) } - fn write_polyfills(&mut self) -> Result<(), Error> { - writeln!( - self.out, - "\ntemplate inline T {ISIGN_FUNCTION}(T arg) {{ - return {NAMESPACE}::select({NAMESPACE}::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -}}" - )?; - Ok(()) - } - fn write_barrier(&mut self, flags: crate::Barrier, level: back::Level) -> BackendResult { // Note: OR-ring bitflags requires `__HAVE_MEMFLAG_OPERATORS__`, // so we try to avoid it here. diff --git a/src/valid/expression.rs b/src/valid/expression.rs index ddef1aaea0..502ce7420a 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -992,7 +992,7 @@ impl super::Validator { } } Mf::Sign => { - if arg1_ty.is_some() | arg2_ty.is_some() | arg3_ty.is_some() { + if arg1_ty.is_some() || arg2_ty.is_some() || arg3_ty.is_some() { return Err(ExpressionError::WrongArgumentCount(fun)); } match *arg_ty { diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index d1fe75ee5d..e5d875dd19 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -58,10 +58,6 @@ struct type_26 { metal::float4 inner[2]; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - void test_matrix_within_struct_accesses( constant Baz& baz ) { diff --git a/tests/out/msl/array-in-ctor.msl b/tests/out/msl/array-in-ctor.msl index e322333d61..9428cb1e74 100644 --- a/tests/out/msl/array-in-ctor.msl +++ b/tests/out/msl/array-in-ctor.msl @@ -11,10 +11,6 @@ struct Ah { type_1 inner; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - kernel void cs_main( device Ah const& ah [[user(fake0)]] ) { diff --git a/tests/out/msl/array-in-function-return-type.msl b/tests/out/msl/array-in-function-return-type.msl index 7475f01ff0..c2c2379cce 100644 --- a/tests/out/msl/array-in-function-return-type.msl +++ b/tests/out/msl/array-in-function-return-type.msl @@ -8,10 +8,6 @@ struct type_1 { float inner[2]; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - type_1 ret_array( ) { return type_1 {1.0, 2.0}; diff --git a/tests/out/msl/atomicOps.msl b/tests/out/msl/atomicOps.msl index 7bff77dff7..b7264e883d 100644 --- a/tests/out/msl/atomicOps.msl +++ b/tests/out/msl/atomicOps.msl @@ -12,10 +12,6 @@ struct Struct { type_2 atomic_arr; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - struct cs_mainInput { }; kernel void cs_main( diff --git a/tests/out/msl/binding-arrays.msl b/tests/out/msl/binding-arrays.msl index b826460e01..10a5397910 100644 --- a/tests/out/msl/binding-arrays.msl +++ b/tests/out/msl/binding-arrays.msl @@ -17,10 +17,6 @@ struct FragmentIn { uint index; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - struct main_Input { uint index [[user(loc0), flat]]; }; diff --git a/tests/out/msl/bitcast.msl b/tests/out/msl/bitcast.msl index 75d6228a6c..a0cf093b7f 100644 --- a/tests/out/msl/bitcast.msl +++ b/tests/out/msl/bitcast.msl @@ -5,10 +5,6 @@ using metal::uint; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - kernel void main_( ) { metal::int2 i2_ = {}; diff --git a/tests/out/msl/bits.msl b/tests/out/msl/bits.msl index 785f0adf60..2a00b6b843 100644 --- a/tests/out/msl/bits.msl +++ b/tests/out/msl/bits.msl @@ -5,10 +5,6 @@ using metal::uint; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - kernel void main_( ) { int i = {}; diff --git a/tests/out/msl/boids.msl b/tests/out/msl/boids.msl index dcb92f51a5..1a81aaf684 100644 --- a/tests/out/msl/boids.msl +++ b/tests/out/msl/boids.msl @@ -28,10 +28,6 @@ struct Particles { }; constant uint NUM_PARTICLES = 1500u; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - struct main_Input { }; kernel void main_( diff --git a/tests/out/msl/bounds-check-image-restrict.msl b/tests/out/msl/bounds-check-image-restrict.msl index 4b631593d1..9f94ef0a6e 100644 --- a/tests/out/msl/bounds-check-image-restrict.msl +++ b/tests/out/msl/bounds-check-image-restrict.msl @@ -5,10 +5,6 @@ using metal::uint; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - metal::float4 test_textureLoad_1d( int coords, int level, diff --git a/tests/out/msl/bounds-check-image-rzsw.msl b/tests/out/msl/bounds-check-image-rzsw.msl index 8d26f222a2..a93014cb27 100644 --- a/tests/out/msl/bounds-check-image-rzsw.msl +++ b/tests/out/msl/bounds-check-image-rzsw.msl @@ -11,10 +11,6 @@ struct DefaultConstructible { }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - metal::float4 test_textureLoad_1d( int coords, int level, diff --git a/tests/out/msl/bounds-check-restrict.msl b/tests/out/msl/bounds-check-restrict.msl index 7c3c5ca6ee..49a232e706 100644 --- a/tests/out/msl/bounds-check-restrict.msl +++ b/tests/out/msl/bounds-check-restrict.msl @@ -20,10 +20,6 @@ struct Globals { type_4 d; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - float index_array( int i, device Globals const& globals, diff --git a/tests/out/msl/bounds-check-zero-atomic.msl b/tests/out/msl/bounds-check-zero-atomic.msl index a191700235..daaa079233 100644 --- a/tests/out/msl/bounds-check-zero-atomic.msl +++ b/tests/out/msl/bounds-check-zero-atomic.msl @@ -24,10 +24,6 @@ struct Globals { type_2 c; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - uint fetch_add_atomic( device Globals& globals, constant _mslBufferSizes& _buffer_sizes diff --git a/tests/out/msl/bounds-check-zero.msl b/tests/out/msl/bounds-check-zero.msl index 9dc96a6237..816983d98b 100644 --- a/tests/out/msl/bounds-check-zero.msl +++ b/tests/out/msl/bounds-check-zero.msl @@ -26,10 +26,6 @@ struct Globals { type_4 d; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - float index_array( int i, device Globals const& globals, diff --git a/tests/out/msl/break-if.msl b/tests/out/msl/break-if.msl index 18b8a8ff98..3a6f2e9bff 100644 --- a/tests/out/msl/break-if.msl +++ b/tests/out/msl/break-if.msl @@ -5,10 +5,6 @@ using metal::uint; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - void breakIfEmpty( ) { bool loop_init = true; diff --git a/tests/out/msl/collatz.msl b/tests/out/msl/collatz.msl index 04b902177a..88f9521a27 100644 --- a/tests/out/msl/collatz.msl +++ b/tests/out/msl/collatz.msl @@ -13,10 +13,6 @@ struct PrimeIndices { type_1 data; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - uint collatz_iterations( uint n_base ) { diff --git a/tests/out/msl/constructors.msl b/tests/out/msl/constructors.msl index 68face5ebb..b3d3b9dd43 100644 --- a/tests/out/msl/constructors.msl +++ b/tests/out/msl/constructors.msl @@ -30,10 +30,6 @@ constant type_10 cz6_ = type_10 {}; constant Foo cz7_ = Foo {}; constant type_11 cp3_ = type_11 {0, 1, 2, 3}; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - kernel void main_( ) { Foo foo = {}; diff --git a/tests/out/msl/control-flow.msl b/tests/out/msl/control-flow.msl index 75312986e7..be396e23a8 100644 --- a/tests/out/msl/control-flow.msl +++ b/tests/out/msl/control-flow.msl @@ -5,10 +5,6 @@ using metal::uint; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - void switch_default_break( int i ) { diff --git a/tests/out/msl/do-while.msl b/tests/out/msl/do-while.msl index 8f6c56388e..4f928d1f3c 100644 --- a/tests/out/msl/do-while.msl +++ b/tests/out/msl/do-while.msl @@ -5,10 +5,6 @@ using metal::uint; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - void fb1_( thread bool& cond ) { diff --git a/tests/out/msl/dualsource.msl b/tests/out/msl/dualsource.msl index c5a49ce558..92b1909f8b 100644 --- a/tests/out/msl/dualsource.msl +++ b/tests/out/msl/dualsource.msl @@ -9,10 +9,6 @@ struct FragmentOutput { metal::float4 mask; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - struct main_Input { }; struct main_Output { diff --git a/tests/out/msl/empty-global-name.msl b/tests/out/msl/empty-global-name.msl index 4c64c79aae..f456649bc6 100644 --- a/tests/out/msl/empty-global-name.msl +++ b/tests/out/msl/empty-global-name.msl @@ -8,10 +8,6 @@ struct type_1 { int member; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - void function( device type_1& unnamed ) { diff --git a/tests/out/msl/empty.msl b/tests/out/msl/empty.msl index 724e63aa0a..4f8bf9f5e9 100644 --- a/tests/out/msl/empty.msl +++ b/tests/out/msl/empty.msl @@ -5,10 +5,6 @@ using metal::uint; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - kernel void main_( ) { return; diff --git a/tests/out/msl/extra.msl b/tests/out/msl/extra.msl index 40272282cd..8288dfad92 100644 --- a/tests/out/msl/extra.msl +++ b/tests/out/msl/extra.msl @@ -14,10 +14,6 @@ struct FragmentIn { uint primitive_index; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - struct main_Input { metal::float4 color [[user(loc0), center_perspective]]; }; diff --git a/tests/out/msl/fragment-output.msl b/tests/out/msl/fragment-output.msl index 6ed78d6bb1..4d25809e4f 100644 --- a/tests/out/msl/fragment-output.msl +++ b/tests/out/msl/fragment-output.msl @@ -21,10 +21,6 @@ struct FragmentOutputVec2Scalar { uint scalaru; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - struct main_vec4vec3_Output { metal::float4 vec4f [[color(0)]]; metal::int4 vec4i [[color(1)]]; diff --git a/tests/out/msl/functions.msl b/tests/out/msl/functions.msl index abd1ff5718..31f57c656e 100644 --- a/tests/out/msl/functions.msl +++ b/tests/out/msl/functions.msl @@ -5,10 +5,6 @@ using metal::uint; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - metal::float2 test_fma( ) { metal::float2 a = metal::float2(2.0, 2.0); diff --git a/tests/out/msl/globals.msl b/tests/out/msl/globals.msl index 24ccdb73fa..d9142c1990 100644 --- a/tests/out/msl/globals.msl +++ b/tests/out/msl/globals.msl @@ -33,10 +33,6 @@ struct type_15 { }; constant bool Foo_1 = true; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - void test_msl_packed_vec3_as_arg( metal::float3 arg ) { diff --git a/tests/out/msl/image.msl b/tests/out/msl/image.msl index d59862f9b8..e390c2e0fc 100644 --- a/tests/out/msl/image.msl +++ b/tests/out/msl/image.msl @@ -5,10 +5,6 @@ using metal::uint; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - struct main_Input { }; kernel void main_( diff --git a/tests/out/msl/interface.msl b/tests/out/msl/interface.msl index 711e021743..c9b93d86a6 100644 --- a/tests/out/msl/interface.msl +++ b/tests/out/msl/interface.msl @@ -23,10 +23,6 @@ struct Input2_ { uint index; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - struct vertex_Input { uint color [[attribute(10)]]; }; diff --git a/tests/out/msl/interpolate.msl b/tests/out/msl/interpolate.msl index 999edbd9e2..5d8b67111e 100644 --- a/tests/out/msl/interpolate.msl +++ b/tests/out/msl/interpolate.msl @@ -15,10 +15,6 @@ struct FragmentInput { float perspective_sample; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - struct vert_mainOutput { metal::float4 position [[position]]; uint _flat [[user(loc0), flat]]; diff --git a/tests/out/msl/math-functions.msl b/tests/out/msl/math-functions.msl index 7476c51f93..dccb90ad6c 100644 --- a/tests/out/msl/math-functions.msl +++ b/tests/out/msl/math-functions.msl @@ -55,10 +55,6 @@ __frexp_result_vec4_f32_ naga_frexp(metal::float4 arg) { return __frexp_result_vec4_f32_{ fract, other }; } -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - fragment void main_( ) { metal::float4 v = metal::float4(0.0); @@ -68,8 +64,9 @@ fragment void main_( metal::float4 d = ((v) * 0.017453292519943295474); metal::float4 e = metal::saturate(v); metal::float4 g = metal::refract(v, v, 1.0); - int sign_a = naga_isign(-1); - metal::int4 sign_b = naga_isign(metal::int4(-1)); + int sign_a = metal::select(metal::select(-1, 1, (-1 > 0)), 0, (-1 == 0)); + metal::int4 _e12 = metal::int4(-1); + metal::int4 sign_b = metal::select(metal::select(int4(-1), int4(1), (_e12 > 0)), 0, (_e12 == 0)); float sign_c = metal::sign(-1.0); metal::float4 sign_d = metal::sign(metal::float4(-1.0)); int const_dot = ( + metal::int2 {}.x * metal::int2 {}.x + metal::int2 {}.y * metal::int2 {}.y); diff --git a/tests/out/msl/operators.msl b/tests/out/msl/operators.msl index 76ef5ebbe2..de3e990db2 100644 --- a/tests/out/msl/operators.msl +++ b/tests/out/msl/operators.msl @@ -9,10 +9,6 @@ constant metal::float4 v_f32_zero = metal::float4(0.0, 0.0, 0.0, 0.0); constant metal::float4 v_f32_half = metal::float4(0.5, 0.5, 0.5, 0.5); constant metal::int4 v_i32_one = metal::int4(1, 1, 1, 1); -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - metal::float4 builtins( ) { int s1_ = true ? 1 : 0; diff --git a/tests/out/msl/padding.msl b/tests/out/msl/padding.msl index 0aa7659414..4d99bb4c4c 100644 --- a/tests/out/msl/padding.msl +++ b/tests/out/msl/padding.msl @@ -23,10 +23,6 @@ struct Test3_ { float b; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - struct vertex_Output { metal::float4 member [[position]]; }; diff --git a/tests/out/msl/policy-mix.msl b/tests/out/msl/policy-mix.msl index 2b05111c59..f6a4fe5d6d 100644 --- a/tests/out/msl/policy-mix.msl +++ b/tests/out/msl/policy-mix.msl @@ -32,10 +32,6 @@ struct type_9 { metal::float4 inner[2]; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - metal::float4 mock_function( metal::int2 c, int i, diff --git a/tests/out/msl/quad-vert.msl b/tests/out/msl/quad-vert.msl index 370bfc8036..d322cbcb9b 100644 --- a/tests/out/msl/quad-vert.msl +++ b/tests/out/msl/quad-vert.msl @@ -18,10 +18,6 @@ struct type_9 { metal::float4 gl_Position; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - void main_1( thread metal::float2& v_uv, thread metal::float2& a_uv_1, diff --git a/tests/out/msl/quad.msl b/tests/out/msl/quad.msl index 368c89d179..5fa5788aac 100644 --- a/tests/out/msl/quad.msl +++ b/tests/out/msl/quad.msl @@ -10,10 +10,6 @@ struct VertexOutput { }; constant float c_scale = 1.2; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - struct vert_mainInput { metal::float2 pos [[attribute(0)]]; metal::float2 uv [[attribute(1)]]; diff --git a/tests/out/msl/ray-query.msl b/tests/out/msl/ray-query.msl index f773cade66..0d4560f313 100644 --- a/tests/out/msl/ray-query.msl +++ b/tests/out/msl/ray-query.msl @@ -41,10 +41,6 @@ struct RayDesc { metal::float3 dir; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - metal::float3 get_torus_normal( metal::float3 world_point, RayIntersection intersection diff --git a/tests/out/msl/resource-binding-map.msl b/tests/out/msl/resource-binding-map.msl index f53f6c8423..b4a53d97b5 100644 --- a/tests/out/msl/resource-binding-map.msl +++ b/tests/out/msl/resource-binding-map.msl @@ -11,10 +11,6 @@ struct DefaultConstructible { }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - struct entry_point_oneInput { }; struct entry_point_oneOutput { diff --git a/tests/out/msl/shadow.msl b/tests/out/msl/shadow.msl index 0c06512c03..53f320344a 100644 --- a/tests/out/msl/shadow.msl +++ b/tests/out/msl/shadow.msl @@ -33,10 +33,6 @@ struct type_7 { constant metal::float3 c_ambient = metal::float3(0.05, 0.05, 0.05); constant uint c_max_lights = 10u; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - float fetch_shadow( uint light_id, metal::float4 homogeneous_coords, diff --git a/tests/out/msl/skybox.msl b/tests/out/msl/skybox.msl index cacc687a9a..7b10ea23e7 100644 --- a/tests/out/msl/skybox.msl +++ b/tests/out/msl/skybox.msl @@ -13,10 +13,6 @@ struct Data { metal::float4x4 view; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - struct vs_mainInput { }; struct vs_mainOutput { diff --git a/tests/out/msl/standard.msl b/tests/out/msl/standard.msl index f1d5e10d0b..f02243eaac 100644 --- a/tests/out/msl/standard.msl +++ b/tests/out/msl/standard.msl @@ -5,10 +5,6 @@ using metal::uint; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - bool test_any_and_all_for_bool( ) { return true; diff --git a/tests/out/msl/texture-arg.msl b/tests/out/msl/texture-arg.msl index b11d2f9caa..5fb9b25649 100644 --- a/tests/out/msl/texture-arg.msl +++ b/tests/out/msl/texture-arg.msl @@ -5,10 +5,6 @@ using metal::uint; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - metal::float4 test( metal::texture2d Passed_Texture, metal::sampler Passed_Sampler diff --git a/tests/out/msl/workgroup-uniform-load.msl b/tests/out/msl/workgroup-uniform-load.msl index 28a0df7ac5..37a8781739 100644 --- a/tests/out/msl/workgroup-uniform-load.msl +++ b/tests/out/msl/workgroup-uniform-load.msl @@ -9,10 +9,6 @@ struct type_2 { }; constant uint SIZE = 128u; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - struct test_workgroupUniformLoadInput { }; kernel void test_workgroupUniformLoad( diff --git a/tests/out/msl/workgroup-var-init.msl b/tests/out/msl/workgroup-var-init.msl index ba495894fd..ac300d4337 100644 --- a/tests/out/msl/workgroup-var-init.msl +++ b/tests/out/msl/workgroup-var-init.msl @@ -19,10 +19,6 @@ struct WStruct { type_4 atom_arr; }; -template inline T naga_isign(T arg) { - return metal::select(metal::select(T(-1), T(1), (arg > 0)), 0, (arg == 0)); -} - kernel void main_( metal::uint3 __local_invocation_id [[thread_position_in_threadgroup]] , threadgroup WStruct& w_mem