Introduce ULP Based Floating Point Equality Test to Device Function

cpp/benchmarks/CMakeLists.txt

@@ -93,3 +93,6 @@ ConfigureNVBench(POINT_IN_POLYGON_BENCH
 
 ConfigureNVBench(QUADTREE_ON_POINTS_BENCH
     quadtree_on_points.cu)
+
+ConfigureNVBench(FLOATING_POINT_EQUALITY_BENCH
+    floating_point_equality.cu)

cpp/benchmarks/floating_point_equality.cu

@@ -0,0 +1,96 @@
+/*
+ * Copyright (c) 2022, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <benchmarks/fixture/rmm_pool_raii.hpp>
+#include <cuspatial_test/random.cuh>
+
+#include <cuspatial/detail/utility/floating_point.cuh>
+#include <cuspatial/error.hpp>
+
+#include <rmm/device_vector.hpp>
+#include <rmm/exec_policy.hpp>
+
+#include <nvbench/nvbench.cuh>
+
+#include <thrust/tabulate.h>
+
+#include <memory>
+#include <type_traits>
+
+using namespace cuspatial;
+
+/**
+ * @brief Helper to generate floats
+ *
+ * @p begin and @p end must be iterators to device-accessible memory
+ *
+ * @tparam FloatsIter The type of the iterator to the output floats container
+ * @param begin The start of the sequence of floats to generate
+ * @param end The end of the sequence of floats to generate
+ */
+template <class FloatsIter>
+void generate_floats(FloatsIter begin, FloatsIter end)
+{
+  using T       = typename std::iterator_traits<FloatsIter>::value_type;
+  auto engine_x = deterministic_engine(std::distance(begin, end));
+
+  auto lo = std::numeric_limits<T>::min();
+  auto hi = std::numeric_limits<T>::max();
+
+  auto x_dist = make_uniform_dist(lo, hi);
+
+  auto x_gen = value_generator{lo, hi, engine_x, x_dist};
+
+  thrust::tabulate(
+    rmm::exec_policy(), begin, end, [x_gen] __device__(size_t n) mutable { return x_gen(n); });
+}
+
+template <typename Float>
+struct eq_comp {
+  using element_t = Float;
+  bool __device__ operator()(Float lhs, Float rhs)
+  {
+    // return lhs == rhs;
+    return detail::float_equal(lhs, rhs);
+  }
+};
+
+template <typename T>
+void floating_point_equivalence_benchmark(nvbench::state& state, nvbench::type_list<T>)
+{
+  // TODO: to be replaced by nvbench fixture once it's ready
+  cuspatial::rmm_pool_raii rmm_pool;
+
+  int64_t const num_floats{state.get_int64("NumFloats")};
+  rmm::device_vector<T> floats(num_floats);
+  rmm::device_vector<bool> results(num_floats);
+
+  generate_floats(floats.begin(), floats.end());
+
+  CUSPATIAL_CUDA_TRY(cudaDeviceSynchronize());
+
+  state.add_element_count(num_floats);
+
+  state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) {
+    auto stream = rmm::cuda_stream_view(launch.get_stream());
+    thrust::transform(floats.begin(), floats.begin(), floats.end(), results.begin(), eq_comp<T>{});
+  });
+}
+
+using floating_point_type = nvbench::type_list<float, double>;
+NVBENCH_BENCH_TYPES(floating_point_equivalence_benchmark, NVBENCH_TYPE_AXES(floating_point_type))
+  .set_type_axes_names({"FloatingPointType"})
+  .add_int64_axis("NumFloats", {100'000, 1'000'000, 10'000'000, 100'000'000});

cpp/include/cuspatial/detail/utility/floating_point.cuh

@@ -0,0 +1,142 @@
+/*
+ * Copyright (c) 2022, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <cuspatial/cuda_utils.hpp>
+
+#include <cmath>
+#include <cstdint>
+#include <type_traits>
+
+namespace cuspatial {
+namespace detail {
+
+constexpr unsigned default_max_ulp = 4;
+
+template <int size, typename = void>
+struct uint_selector;
+
+template <int size>
+struct uint_selector<size, std::enable_if_t<size == 2>> {
+  using type = uint16_t;
+};
+
+template <int size>
+struct uint_selector<size, std::enable_if_t<size == 4>> {
+  using type = uint32_t;
+};
+
+template <int size>
+struct uint_selector<size, std::enable_if_t<size == 8>> {
+  using type = uint64_t;
+};
+
+template <typename Bits>
+Bits constexpr sign_bit_mask()
+{
+  return Bits{1} << 8 * sizeof(Bits) - 1;
+}
+
+template <typename T>
+union FloatingPointBits {
+  using Bits = typename uint_selector<sizeof(T)>::type;
+  CUSPATIAL_HOST_DEVICE FloatingPointBits(T float_number) : _f(float_number) {}
+  T _f;
+  Bits _b;
+};
+
+/**
+ * @internal
+ * @brief Converts integer of sign-magnitude representation to biased representation.
+ *
+ * Biased representation has 1 representation of zero while sign-magnitude has 2.
+ * This conversion will collapse the two representations into 1. This is in line with
+ * our expectation that a positive number 1 differ from a negative number -1 by 2 hops
+ * instead of 3 in biased representation.
+ *
+ * Example:
+ * Assume `N` bits in the type `Bits`. In total 2^(N-1) representable numbers.
+ * (N=4):
+ *              |--------------|  |-----------------|
+ * decimal    -2^3+1          -0 +0                2^3-1
+ * SaM         1111          1000 0000             0111
+ *
+ * In SaM, 0 is represented twice. In biased representation we need to collapse
+ * them to single representation, resulting in 1 more representable number in
+ * biased form.
+ *
+ * Naturally, lowest bit should map to the smallest number representable in the range.
+ * With 1 more representable number in biased form, we discard the lowest bit and start
+ * at the next lowest bit.
+ *              |--------------|-----------------|
+ * decimal    -2^3+1           0                2^3-1
+ * biased      0001           0111              1110
+ *
+ * The following implements the mapping independently in negative and positive range.
+ *
+ * Read http://en.wikipedia.org/wiki/Signed_number_representations for more
+ * details on signed number representations.
+ *
+ * @tparam Bits Unsigned type to store the bits
+ * @param sam Sign and magnitude representation
+ * @return Biased representation
+ */
+template <typename Bits>
+std::enable_if_t<std::is_unsigned_v<Bits>, Bits> CUSPATIAL_HOST_DEVICE
+signmagnitude_to_biased(Bits const& sam)
+{
+  return sam & sign_bit_mask<Bits>() ? ~sam + 1 : sam | sign_bit_mask<Bits>();
+}
+
+/**
+ * @brief Floating-point equivalence comparator based on ULP (Unit in the last place).
+ *
+ * @tparam T Type of floating point
+ * @tparam max_ulp Maximum tolerable unit in the last place
+ * @param lhs First floating point to compare
+ * @param rhs Second floating point to compare
+ * @return `true` if two floating points differ by less or equal to `ulp`.
+ */
+template <typename T, unsigned max_ulp = default_max_ulp>
+bool CUSPATIAL_HOST_DEVICE float_equal(T const& flhs, T const& frhs)
+{
+  FloatingPointBits<T> lhs{flhs};
+  FloatingPointBits<T> rhs{frhs};
+  if (std::isnan(lhs._f) || std::isnan(rhs._f)) return false;
+  auto lhsbiased = signmagnitude_to_biased(lhs._b);
+  auto rhsbiased = signmagnitude_to_biased(rhs._b);
+
+  return lhsbiased >= rhsbiased ? (lhsbiased - rhsbiased) <= max_ulp
+                                : (rhsbiased - lhsbiased) <= max_ulp;
+}
+
+/**
+ * @brief Floating-point non equivalence comparator based on ULP (Unit in the last place).
+ *
+ * @tparam T Type of floating point
+ * @tparam max_ulp Maximum tolerable unit in the last place
+ * @param lhs First floating point to compare
+ * @param rhs Second floating point to compare
+ * @return `true` if two floating points differ by greater `ulp`.
+ */
+template <typename T, unsigned max_ulp = default_max_ulp>
+bool CUSPATIAL_HOST_DEVICE not_float_equal(FloatingPointBits<T> const& lhs,
+                                           FloatingPointBits<T> const& rhs)
+{
+  return !float_equal(lhs, rhs);
+}
+
+}  // namespace detail
+}  // namespace cuspatial

cpp/tests/CMakeLists.txt

@@ -110,6 +110,9 @@ ConfigureTest(TRAJECTORY_BOUNDING_BOXES_TEST
 ConfigureTest(SPATIAL_WINDOW_POINT_TEST
     spatial_window/spatial_window_test.cpp)
 
+ConfigureTest(FLOAT_EQUIVALENT_UTILITY_TEST
+    utility_test/test_float_equivalent.cu)
+
 # Experimental API
 ConfigureTest(HAVERSINE_TEST_EXP
     experimental/spatial/haversine_test.cu)

cpp/tests/utility_test/test_float_equivalent.cu

@@ -0,0 +1,122 @@
+#include <cuspatial/detail/utility/floating_point.cuh>
+
+#include <limits>
+#include <rmm/device_vector.hpp>
+
+#include <gtest/gtest.h>
+
+using namespace cuspatial;
+
+template <typename T>
+struct ULPFloatingPointEquivalenceTest : public ::testing::Test {
+};
+
+using TestTypes = ::testing::Types<float, double>;
+
+TYPED_TEST_CASE(ULPFloatingPointEquivalenceTest, TestTypes);
+
+template <typename Float>
+struct float_eq_comp {
+  bool __device__ operator()(Float lhs, Float rhs) { return detail::float_equal(lhs, rhs); }
+};
+
+template <typename T>
+T increment(T f, unsigned step)
+{
+  if (!step) return f;
+  return increment(std::nextafter(f, std::numeric_limits<T>::max()), step - 1);
+}
+
+template <typename T>
+T decrement(T f, unsigned step)
+{
+  if (!step) return f;
+  return decrement(std::nextafter(f, std::numeric_limits<T>::min()), step - 1);
+}
+
+template <typename T>
+void run_test(T base)
+{
+  T FourULPGreater = increment(base, 4);
+  T FiveULPGreater = increment(base, 5);
+  T FourULPLess    = decrement(base, 4);
+  T FiveULPLess    = decrement(base, 5);
+
+  std::vector<T> first{base, base, base, base};
+  std::vector<T> second{FourULPGreater, FiveULPGreater, FourULPLess, FiveULPLess};
+
+  rmm::device_vector<T> d_first(first);
+  rmm::device_vector<T> d_second(second);
+
+  std::vector<bool> expected{true, false, true, false};
+  rmm::device_vector<bool> got(4);
+
+  thrust::transform(
+    d_first.begin(), d_first.end(), d_second.begin(), got.begin(), float_eq_comp<T>{});
+
+  EXPECT_EQ(expected, got);
+}
+
+TYPED_TEST(ULPFloatingPointEquivalenceTest, BiasedFromPositiveZero)
+{
+  using T = TypeParam;
+  run_test(T{0.0});
+}
+
+TYPED_TEST(ULPFloatingPointEquivalenceTest, BiasedFromNegativeZero)
+{
+  using T = TypeParam;
+  run_test(T{-0.0});
+}
+
+TYPED_TEST(ULPFloatingPointEquivalenceTest, TestVeryNearZeroPositive)
+{
+  using T                     = TypeParam;
+  T very_small_positive_float = increment(T{0.0}, 1);
+  run_test(very_small_positive_float);
+}
+
+TYPED_TEST(ULPFloatingPointEquivalenceTest, TestVeryNearZeroNegative)
+{
+  using T                     = TypeParam;
+  T very_small_negative_float = decrement(T{0.0}, 1);
+  run_test(very_small_negative_float);
+}
+
+TYPED_TEST(ULPFloatingPointEquivalenceTest, BiasedFromSmallPostiveFloat)
+{
+  using T = TypeParam;
+  run_test(T{0.1});
+}
+
+TYPED_TEST(ULPFloatingPointEquivalenceTest, BiasedFromSmallNegativeFloat)
+{
+  using T = TypeParam;
+  run_test(T{-0.1});
+}
+
+TYPED_TEST(ULPFloatingPointEquivalenceTest, BiasedFromPostiveFloat)
+{
+  using T = TypeParam;
+  run_test(T{1234.0});
+}
+
+TYPED_TEST(ULPFloatingPointEquivalenceTest, BiasedFromNegativeFloat)
+{
+  using T = TypeParam;
+  run_test(T{-5678.0});
+}
+
+TYPED_TEST(ULPFloatingPointEquivalenceTest, BiasedFromVeryLargePositiveFloat)
+{
+  using T                     = TypeParam;
+  T very_large_positive_float = decrement(std::numeric_limits<T>::max(), 10);
+  run_test(very_large_positive_float);
+}
+
+TYPED_TEST(ULPFloatingPointEquivalenceTest, BiasedFromVeryLargeNegativeFloat)
+{
+  using T                     = TypeParam;
+  T very_large_negative_float = increment(std::numeric_limits<T>::min(), 10);
+  run_test(very_large_negative_float);
+}