Implement all core_float_math functions using core::intrinsics

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

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

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);
 }

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,
         },

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<u32>) {
     let tid = global_id.x;
@@ -18,9 +67,9 @@ fn main_cs(@builtin(global_invocation_id) global_id: vec3<u32>) {
     }
 
     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<u32>) {
     // Special values and conversions
     let int_val = i32(x);
     output[base_offset + 20u] = f32(int_val);
-}
+
+    // 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));
+}

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,
         },