diff --git a/crates/spirv-std/src/float.rs b/crates/spirv-std/src/float.rs index 0f957c28416..71d4a03eb64 100644 --- a/crates/spirv-std/src/float.rs +++ b/crates/spirv-std/src/float.rs @@ -8,22 +8,234 @@ use core::intrinsics; #[cfg(target_arch = "spirv")] impl f32 { + /// Returns the largest integer that is less than or equal to `self`. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn floor(self) -> f32 { + intrinsics::floorf32(self) + } + + /// Returns the smallest integer that is greater than or equal to `self`. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn ceil(self) -> f32 { + intrinsics::ceilf32(self) + } + + /// Returns the nearest integer to `self`. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn round(self) -> f32 { + intrinsics::roundf32(self) + } + + /// Returns the nearest integer to a number. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn round_ties_even(self) -> f32 { + intrinsics::round_ties_even_f32(self) + } + + /// Returns the integer part of `self`. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn trunc(self) -> f32 { + intrinsics::truncf32(self) + } + + /// Returns the fractional part of `self`. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn fract(self) -> f32 { + self - self.trunc() + } + + /// Fused multiply-add. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn mul_add(self, a: f32, b: f32) -> f32 { + intrinsics::fmaf32(self, a, b) + } + + /// Calculates Euclidean division, the matching method for `rem_euclid`. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn div_euclid(self, rhs: f32) -> f32 { + let q = (self / rhs).trunc(); + if self % rhs < 0.0 { + return if rhs > 0.0 { q - 1.0 } else { q + 1.0 }; + } + q + } + + /// Calculates the least nonnegative remainder of `self` when divided by `rhs`. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn rem_euclid(self, rhs: f32) -> f32 { + let r = self % rhs; + if r < 0.0 { r + rhs.abs() } else { r } + } + /// Raises a number to an integer power. #[inline] #[rustc_allow_incoherent_impl] pub fn powi(self, n: i32) -> f32 { intrinsics::powif32(self, n) } + + /// Returns the square root of a number. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn sqrt(self) -> f32 { + intrinsics::sqrtf32(self) + } + + /// Returns the cube root of a number. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn cbrt(self) -> f32 { + intrinsics::powf32(self, 1.0 / 3.0) + } + + /// The positive difference of two numbers. + #[inline] + #[rustc_allow_incoherent_impl] + #[deprecated( + since = "0.10.0", + note = "you probably meant `(self - other).abs()`: \ + this operation is `(self - other).max(0.0)` \ + except that `abs_sub` also propagates NaNs (also \ + known as `fdimf` in C). If you truly need the positive \ + difference, consider using that expression or the C function \ + `fdimf`, depending on how you wish to handle NaN (please consider \ + filing an issue describing your use-case too)." + )] + pub fn abs_sub(self, other: f32) -> f32 { + let r = self - other; + if r.is_nan() { + r + } else if r > 0.0 { + r + } else { + 0.0 + } + } } #[cfg(target_arch = "spirv")] impl f64 { + /// Returns the largest integer that is less than or equal to `self`. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn floor(self) -> f64 { + intrinsics::floorf64(self) + } + + /// Returns the smallest integer that is greater than or equal to `self`. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn ceil(self) -> f64 { + intrinsics::ceilf64(self) + } + + /// Returns the nearest integer to `self`. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn round(self) -> f64 { + intrinsics::roundf64(self) + } + + /// Returns the nearest integer to a number. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn round_ties_even(self) -> f64 { + intrinsics::round_ties_even_f64(self) + } + + /// Returns the integer part of `self`. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn trunc(self) -> f64 { + intrinsics::truncf64(self) + } + + /// Returns the fractional part of `self`. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn fract(self) -> f64 { + self - self.trunc() + } + + /// Fused multiply-add. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn mul_add(self, a: f64, b: f64) -> f64 { + intrinsics::fmaf64(self, a, b) + } + + /// Calculates Euclidean division, the matching method for `rem_euclid`. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn div_euclid(self, rhs: f64) -> f64 { + let q = (self / rhs).trunc(); + if self % rhs < 0.0 { + return if rhs > 0.0 { q - 1.0 } else { q + 1.0 }; + } + q + } + + /// Calculates the least nonnegative remainder of `self` when divided by `rhs`. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn rem_euclid(self, rhs: f64) -> f64 { + let r = self % rhs; + if r < 0.0 { r + rhs.abs() } else { r } + } + /// Raises a number to an integer power. #[inline] #[rustc_allow_incoherent_impl] pub fn powi(self, n: i32) -> f64 { intrinsics::powif64(self, n) } + + /// Returns the square root of a number. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn sqrt(self) -> f64 { + intrinsics::sqrtf64(self) + } + + /// Returns the cube root of a number. + #[inline] + #[rustc_allow_incoherent_impl] + pub fn cbrt(self) -> f64 { + intrinsics::powf64(self, 1.0 / 3.0) + } + + /// The positive difference of two numbers. + #[inline] + #[rustc_allow_incoherent_impl] + #[deprecated( + since = "0.10.0", + note = "you probably meant `(self - other).abs()`: \ + this operation is `(self - other).max(0.0)` \ + except that `abs_sub` also propagates NaNs (also \ + known as `fdim` in C). If you truly need the positive \ + difference, consider using that expression or the C function \ + `fdim`, depending on how you wish to handle NaN (please consider \ + filing an issue describing your use-case too)." + )] + pub fn abs_sub(self, other: f64) -> f64 { + let r = self - other; + if r.is_nan() { + r + } else if r > 0.0 { + r + } else { + 0.0 + } + } } /// Converts two f32 values (floats) into two f16 values (halfs). The result is a u32, with the low diff --git a/tests/compiletests/ui/dis/powi.stderr b/tests/compiletests/ui/dis/powi.stderr index 4d189108626..30aeb48af6a 100644 --- a/tests/compiletests/ui/dis/powi.stderr +++ b/tests/compiletests/ui/dis/powi.stderr @@ -2,7 +2,7 @@ %4 = OpLabel OpLine %5 10 12 %6 = OpLoad %7 %8 - OpLine %9 15 8 + OpLine %9 83 8 %10 = OpConvertSToF %7 %11 %12 = OpExtInst %7 %13 26 %6 %10 OpLine %5 11 4 diff --git a/tests/difftests/tests/lang/core/ops/math_ops/math_ops-rust/src/lib.rs b/tests/difftests/tests/lang/core/ops/math_ops/math_ops-rust/src/lib.rs index 1703e2a0362..7cd2f149a8e 100644 --- a/tests/difftests/tests/lang/core/ops/math_ops/math_ops-rust/src/lib.rs +++ b/tests/difftests/tests/lang/core/ops/math_ops/math_ops-rust/src/lib.rs @@ -18,9 +18,9 @@ pub fn main_cs( } let x = input[tid]; - let base_offset = tid * 21; + let base_offset = tid * 30; - if base_offset + 20 >= output.len() { + if base_offset + 29 >= output.len() { return; } @@ -53,4 +53,17 @@ pub fn main_cs( // Special values and conversions let int_val = x as i32; output[base_offset + 20] = int_val as f32; + + // core_float_math methods not covered above + output[base_offset + 21] = x.round_ties_even(); + output[base_offset + 22] = x.trunc(); + output[base_offset + 23] = compat_round!(x.fract()); + output[base_offset + 24] = compat_round!(x.mul_add(2.0, 1.0)); + output[base_offset + 25] = compat_round!(x.div_euclid(3.0)); + output[base_offset + 26] = compat_round!(x.rem_euclid(3.0)); + output[base_offset + 27] = compat_round!(x.powi(2)); + output[base_offset + 28] = compat_round!(x.abs().cbrt()); + #[allow(deprecated)] + let abs_sub = x.abs_sub(1.0); + output[base_offset + 29] = compat_round!(abs_sub); } diff --git a/tests/difftests/tests/lang/core/ops/math_ops/math_ops-rust/src/main.rs b/tests/difftests/tests/lang/core/ops/math_ops/math_ops-rust/src/main.rs index 9358fb37fc7..68cc2db01f4 100644 --- a/tests/difftests/tests/lang/core/ops/math_ops/math_ops-rust/src/main.rs +++ b/tests/difftests/tests/lang/core/ops/math_ops/math_ops-rust/src/main.rs @@ -39,7 +39,7 @@ fn main() { initial_data: Some(input_bytes), }, BufferConfig { - size: 2688, // 672 f32 values (32 threads * 21 outputs each) + size: 3840, // 960 f32 values (32 threads * 30 outputs each) usage: BufferUsage::Storage, initial_data: None, }, diff --git a/tests/difftests/tests/lang/core/ops/math_ops/math_ops-wgsl/shader.wgsl b/tests/difftests/tests/lang/core/ops/math_ops/math_ops-wgsl/shader.wgsl index c143e8a9df9..88f42173f51 100644 --- a/tests/difftests/tests/lang/core/ops/math_ops/math_ops-wgsl/shader.wgsl +++ b/tests/difftests/tests/lang/core/ops/math_ops/math_ops-wgsl/shader.wgsl @@ -9,6 +9,55 @@ fn compat_round(v: f32) -> f32 { return round(v * 100000.0) / 100000.0; } +fn is_even_int(x: f32) -> bool { + return (x * 0.5) == floor(x * 0.5); +} + +fn round_ties_even(x: f32) -> f32 { + let floored = floor(x); + let frac = x - floored; + + if (frac == 0.5) { + return select(floored + 1.0, floored, is_even_int(floored)); + } else if (frac > 0.5) { + return floored + 1.0; + } else { + return floored; + } +} + +fn div_euclid(x: f32, rhs: f32) -> f32 { + let q = trunc(x / rhs); + + if (x % rhs < 0.0) { + return select(q + 1.0, q - 1.0, rhs > 0.0); + } + + return q; +} + +fn rem_euclid(x: f32, rhs: f32) -> f32 { + let r = x % rhs; + + if (r < 0.0) { + return r + abs(rhs); + } + + return r; +} + +fn abs_sub(x: f32, other: f32) -> f32 { + let r = x - other; + + if (r != r) { + return r; + } else if (r > 0.0) { + return r; + } else { + return 0.0; + } +} + @compute @workgroup_size(32, 1, 1) fn main_cs(@builtin(global_invocation_id) global_id: vec3) { let tid = global_id.x; @@ -18,9 +67,9 @@ fn main_cs(@builtin(global_invocation_id) global_id: vec3) { } let x = input[tid]; - let base_offset = tid * 21u; + let base_offset = tid * 30u; - if (base_offset + 20u >= arrayLength(&output)) { + if (base_offset + 29u >= arrayLength(&output)) { return; } @@ -53,4 +102,15 @@ fn main_cs(@builtin(global_invocation_id) global_id: vec3) { // Special values and conversions let int_val = i32(x); output[base_offset + 20u] = f32(int_val); -} \ No newline at end of file + + // core_float_math methods not covered above + output[base_offset + 21u] = round_ties_even(x); + output[base_offset + 22u] = trunc(x); + output[base_offset + 23u] = compat_round(x - trunc(x)); + output[base_offset + 24u] = compat_round(fma(x, 2.0, 1.0)); + output[base_offset + 25u] = compat_round(div_euclid(x, 3.0)); + output[base_offset + 26u] = compat_round(rem_euclid(x, 3.0)); + output[base_offset + 27u] = compat_round(pow(x, 2.0)); + output[base_offset + 28u] = compat_round(pow(abs(x), 1.0 / 3.0)); + output[base_offset + 29u] = compat_round(abs_sub(x, 1.0)); +} diff --git a/tests/difftests/tests/lang/core/ops/math_ops/math_ops-wgsl/src/main.rs b/tests/difftests/tests/lang/core/ops/math_ops/math_ops-wgsl/src/main.rs index 50c6a693531..689d5b824e0 100644 --- a/tests/difftests/tests/lang/core/ops/math_ops/math_ops-wgsl/src/main.rs +++ b/tests/difftests/tests/lang/core/ops/math_ops/math_ops-wgsl/src/main.rs @@ -36,7 +36,7 @@ fn main() { initial_data: Some(input_bytes), }, BufferConfig { - size: 2688, // 672 f32 values (32 threads * 21 outputs each) + size: 3840, // 960 f32 values (32 threads * 30 outputs each) usage: BufferUsage::Storage, initial_data: None, },