diff --git a/src/back/msl/keywords.rs b/src/back/msl/keywords.rs index f34b618db8..c7da1e332c 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::NAGA_ISIGN_FUNCTION, ]; diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index 24c54ff8ab..459c71569f 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 NAGA_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, + .. + } => NAGA_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 == NAGA_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,17 @@ impl Writer { Ok(info) } + fn write_polyfills(&mut self) -> Result<(), Error> { + writeln!(self.out)?; + writeln!( + self.out, + "template inline T {NAGA_ISIGN_FUNCTION}(T arg) {{ + return {NAMESPACE}::select({NAMESPACE}::select(-1, 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..2387f5cbaf 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(-1, 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..940768628a 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(-1, 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..7180504600 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(-1, 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..788c3b7f28 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(-1, 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..dc99248cb7 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(-1, 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..cb25dda2e8 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(-1, 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..0d90adf9a2 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(-1, 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..1f5c69227a 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(-1, 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..ee0fa34c03 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(-1, 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..f7223a2a7b 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(-1, 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..6adf899670 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(-1, 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..9eaa27a858 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(-1, 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..35211d2777 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(-1, 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..28abeac7a2 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(-1, 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..9d58b9c9d5 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(-1, 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..7015fb390f 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(-1, 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..be8a29a131 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(-1, 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..95761b4cba 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(-1, 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..9de32cc576 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(-1, 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..1638a2f9d8 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(-1, 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..dbac56a1f2 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(-1, 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..e5c10865aa 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(-1, 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..d2d00f4714 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(-1, 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..ad0bcf7414 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(-1, 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..d0b0349e24 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(-1, 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..3260b92361 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(-1, 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..49d89836ec 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(-1, 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..dc9f33041f 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(-1, 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..020a77b861 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(-1, 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..ae651c6303 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(-1, 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..511e7fbac7 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(-1, 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..943eb5db0f 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(-1, 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..6771c4cf37 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(-1, 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..d42ce427c5 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(-1, 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..a27e989a61 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(-1, 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..40b3e40e46 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(-1, 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..8a1181c234 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(-1, 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..ffcba770e5 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(-1, 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..2c0313bd59 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(-1, 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..de5b4708c3 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(-1, 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..c8db8041d5 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(-1, 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..adc3f58631 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(-1, 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);