Fix string to double conversion and row equivalent comparison

cpp/benchmarks/CMakeLists.txt

@@ -308,6 +308,7 @@ set(STRINGS_BENCH_SRC
   "${CMAKE_CURRENT_SOURCE_DIR}/string/case_benchmark.cpp"
   "${CMAKE_CURRENT_SOURCE_DIR}/string/contains_benchmark.cpp"
   "${CMAKE_CURRENT_SOURCE_DIR}/string/convert_durations_benchmark.cpp"
+  "${CMAKE_CURRENT_SOURCE_DIR}/string/convert_floats_benchmark.cpp"
   "${CMAKE_CURRENT_SOURCE_DIR}/string/copy_benchmark.cpp"
   "${CMAKE_CURRENT_SOURCE_DIR}/string/find_benchmark.cpp"
   "${CMAKE_CURRENT_SOURCE_DIR}/string/replace_benchmark.cpp"

cpp/benchmarks/string/convert_floats_benchmark.cpp

@@ -0,0 +1,131 @@
+/*
+ * Copyright (c) 2021, 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 <fixture/benchmark_fixture.hpp>
+#include <synchronization/synchronization.hpp>
+
+#include <benchmark/benchmark.h>
+
+#include <cudf/column/column_view.hpp>
+#include <cudf/strings/convert/convert_floats.hpp>
+#include <cudf/types.hpp>
+
+#include <cudf_test/base_fixture.hpp>
+#include <cudf_test/column_wrapper.hpp>
+
+#include <numeric>
+#include <random>
+#include <unordered_map>
+
+namespace {
+
+// For each array_size, this function is called twice from both StringToFloatNumber and
+// StringFromFloatNumber classes. Thus, the results are cached for reuse.
+template <class FloatType>
+static const std::vector<FloatType>& get_float_numbers(int64_t array_size)
+{
+  static std::unordered_map<int64_t, std::vector<FloatType>> number_arrays;
+  auto& numbers = number_arrays[array_size];
+  if (numbers.size() == 0) {
+    numbers.reserve(array_size);
+    cudf::test::UniformRandomGenerator<FloatType> rand_gen(std::numeric_limits<FloatType>::min(),
+                                                           std::numeric_limits<FloatType>::max());
+    std::generate_n(
+      std::back_inserter(numbers), array_size, [&rand_gen]() { return rand_gen.generate(); });
+  }
+  return numbers;
+}
+
+template <class FloatType>
+static std::vector<std::string> get_floats_numbers_as_string(int64_t array_size)
+{
+  std::vector<std::string> numbers_str(array_size);
+  const auto& numbers = get_float_numbers<FloatType>(array_size);
+  std::transform(
+    numbers.begin(), numbers.end(), numbers_str.begin(), [](auto x) { return std::to_string(x); });
+  return numbers_str;
+}
+
+}  // anonymous namespace
+
+class StringToFloatNumber : public cudf::benchmark {
+};
+template <cudf::type_id float_type>
+void convert_to_float_number(benchmark::State& state)
+{
+  const auto& h_strings   = get_floats_numbers_as_string<double>(state.range(0));
+  const auto strings_size = std::accumulate(
+    h_strings.begin(), h_strings.end(), std::size_t{0}, [](std::size_t size, const auto& str) {
+      return size + str.length();
+    });
+
+  cudf::test::strings_column_wrapper strings(h_strings.begin(), h_strings.end());
+  const auto strings_view = cudf::strings_column_view(strings);
+
+  for (auto _ : state) {
+    cuda_event_timer raii(state, true);  // flush_l2_cache = true, stream = 0
+    volatile auto results = cudf::strings::to_floats(strings_view, cudf::data_type{float_type});
+  }
+
+  state.SetBytesProcessed(state.iterations() * strings_size);
+}
+
+class StringFromFloatNumber : public cudf::benchmark {
+};
+template <class FloatType>
+void convert_from_float_number(benchmark::State& state)
+{
+  const auto& h_floats   = get_float_numbers<FloatType>(state.range(0));
+  const auto floats_size = h_floats.size() * sizeof(FloatType);
+
+  cudf::test::fixed_width_column_wrapper<FloatType> floats(h_floats.begin(), h_floats.end());
+  const auto floats_view = cudf::column_view(floats);
+
+  for (auto _ : state) {
+    cuda_event_timer raii(state, true);  // flush_l2_cache = true, stream = 0
+    volatile auto results = cudf::strings::from_floats(floats_view);
+  }
+
+  state.SetBytesProcessed(state.iterations() * floats_size);
+}
+
+#define CV_TO_FLOATS_BENCHMARK_DEFINE(name, float_type_id)                  \
+  BENCHMARK_DEFINE_F(StringToFloatNumber, name)(::benchmark::State & state) \
+  {                                                                         \
+    convert_to_float_number<float_type_id>(state);                          \
+  }                                                                         \
+  BENCHMARK_REGISTER_F(StringToFloatNumber, name)                           \
+    ->RangeMultiplier(1 << 2)                                               \
+    ->Range(1 << 10, 1 << 17)                                               \
+    ->UseManualTime()                                                       \
+    ->Unit(benchmark::kMicrosecond);
+
+#define CV_FROM_FLOATS_BENCHMARK_DEFINE(name, float_type)                     \
+  BENCHMARK_DEFINE_F(StringFromFloatNumber, name)(::benchmark::State & state) \
+  {                                                                           \
+    convert_from_float_number<float_type>(state);                             \
+  }                                                                           \
+  BENCHMARK_REGISTER_F(StringFromFloatNumber, name)                           \
+    ->RangeMultiplier(1 << 2)                                                 \
+    ->Range(1 << 10, 1 << 17)                                                 \
+    ->UseManualTime()                                                         \
+    ->Unit(benchmark::kMicrosecond);
+
+CV_TO_FLOATS_BENCHMARK_DEFINE(string_to_float32, cudf::type_id::FLOAT32);
+CV_TO_FLOATS_BENCHMARK_DEFINE(string_to_float64, cudf::type_id::FLOAT64);
+
+CV_FROM_FLOATS_BENCHMARK_DEFINE(string_from_float32, float);
+CV_FROM_FLOATS_BENCHMARK_DEFINE(string_from_float64, double);

cpp/src/strings/convert/convert_floats.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -58,15 +58,15 @@ __device__ inline double stod(string_view const& d_str)
   if (d_str.compare("NaN", 3) == 0) return std::numeric_limits<double>::quiet_NaN();
   if (d_str.compare("Inf", 3) == 0) return std::numeric_limits<double>::infinity();
   if (d_str.compare("-Inf", 4) == 0) return -std::numeric_limits<double>::infinity();
-  double sign = 1.0;
+  double sign{1.0};
   if (*in_ptr == '-' || *in_ptr == '+') {
     sign = (*in_ptr == '-' ? -1 : 1);
     ++in_ptr;
   }
-  unsigned long max_mantissa = 0x0FFFFFFFFFFFFF;
-  unsigned long digits       = 0;
-  int exp_off                = 0;
-  bool decimal               = false;
+  constexpr uint64_t max_holding = (std::numeric_limits<uint64_t>::max() - 9L) / 10L;
+  uint64_t digits                = 0;
+  int exp_off                    = 0;
+  bool decimal                   = false;
   while (in_ptr < end) {
     char ch = *in_ptr;
     if (ch == '.') {
@@ -75,18 +75,20 @@ __device__ inline double stod(string_view const& d_str)
       continue;
     }
     if (ch < '0' || ch > '9') break;
-    if (digits > max_mantissa)
+    if (digits > max_holding)
       exp_off += (int)!decimal;
     else {
-      digits = (digits * 10L) + (unsigned long)(ch - '0');
-      if (digits > max_mantissa) {
+      digits = (digits * 10L) + static_cast<uint64_t>(ch - '0');
+      if (digits > max_holding) {
         digits = digits / 10L;
         exp_off += (int)!decimal;
       } else
         exp_off -= (int)decimal;
     }
     ++in_ptr;
   }
+  if (digits == 0) return sign * static_cast<double>(0);
+
   // check for exponent char
   int exp_ten  = 0;
   int exp_sign = 1;
@@ -107,17 +109,23 @@ __device__ inline double stod(string_view const& d_str)
       }
     }
   }
+
+  int const num_digits = static_cast<int>(log10(digits)) + 1;
   exp_ten *= exp_sign;
   exp_ten += exp_off;
-  if (exp_ten > 308)
+  exp_ten += num_digits - 1;
+  if (exp_ten > std::numeric_limits<double>::max_exponent10)
     return sign > 0 ? std::numeric_limits<double>::infinity()
                     : -std::numeric_limits<double>::infinity();
-  else if (exp_ten < -308)
-    return 0.0;
+  else if (exp_ten < std::numeric_limits<double>::min_exponent10)
+    return double{0};
+
   // using exp10() since the pow(10.0,exp_ten) function is
   // very inaccurate in 10.2: http://nvbugs/2971187
-  double value = static_cast<double>(digits) * exp10(static_cast<double>(exp_ten));
-  return (value * sign);
+  double const base =
+    sign * static_cast<double>(digits) * exp10(static_cast<double>(1 - num_digits));
+  double const exponent = exp10(static_cast<double>(exp_ten));
+  return base * exponent;
 }
 
 /**
@@ -132,8 +140,8 @@ struct string_to_float_fn {
   __device__ FloatType operator()(size_type idx)
   {
     if (strings_column.is_null(idx)) return static_cast<FloatType>(0);
-    // the cast to FloatType will create predictable results
-    // for floats that are larger than the FloatType can hold
+    // The cast to FloatType will create predictable results for floats that are larger than the
+    // FloatType can hold
     return static_cast<FloatType>(stod(strings_column.element<string_view>(idx)));
   }
 };

cpp/tests/strings/floats_tests.cpp

@@ -45,31 +45,17 @@ TEST_F(StringsConvertTest, ToFloats32)
                                      "456e",
                                      "-1.78e+5",
                                      "-122.33644782123456789",
-                                     "12e+309"};
+                                     "12e+309",
+                                     "3.4028236E38"};
   cudf::test::strings_column_wrapper strings(
     h_strings.begin(),
     h_strings.end(),
     thrust::make_transform_iterator(h_strings.begin(), [](auto str) { return str != nullptr; }));
 
-  float nanval = std::numeric_limits<float>::quiet_NaN();
-  float infval = std::numeric_limits<float>::infinity();
-  std::vector<float> h_expected{1234.0,
-                                0,
-                                -876.0,
-                                543.2,
-                                -0.12,
-                                0.25,
-                                -0.002,
-                                0,
-                                -0.0,
-                                12000,
-                                nanval,
-                                0,
-                                123.0,
-                                456.0,
-                                -178000.0,
-                                -122.3364486694336,
-                                infval};
+  std::vector<float> h_expected;
+  std::for_each(h_strings.begin(), h_strings.end(), [&](const char* str) {
+    h_expected.push_back(str ? std::atof(str) : 0);
+  });
 
   auto strings_view = cudf::strings_column_view(strings);
   auto results = cudf::strings::to_floats(strings_view, cudf::data_type{cudf::type_id::FLOAT32});
@@ -78,7 +64,7 @@ TEST_F(StringsConvertTest, ToFloats32)
     h_expected.begin(),
     h_expected.end(),
     thrust::make_transform_iterator(h_strings.begin(), [](auto str) { return str != nullptr; }));
-  CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*results, expected);
+  CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*results, expected, true);
 }
 
 TEST_F(StringsConvertTest, FromFloats32)
@@ -107,7 +93,7 @@ TEST_F(StringsConvertTest, FromFloats32)
     h_expected.end(),
     thrust::make_transform_iterator(h_expected.begin(), [](auto str) { return str != nullptr; }));
 
-  CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*results, expected);
+  CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*results, expected, true);
 }
 
 TEST_F(StringsConvertTest, ToFloats64)
@@ -128,31 +114,17 @@ TEST_F(StringsConvertTest, ToFloats64)
                                      "456e",
                                      "-1.78e+5",
                                      "-122.33644782",
-                                     "12e+309"};
+                                     "12e+309",
+                                     "1.7976931348623159E308"};
   cudf::test::strings_column_wrapper strings(
     h_strings.begin(),
     h_strings.end(),
     thrust::make_transform_iterator(h_strings.begin(), [](auto str) { return str != nullptr; }));
 
-  double nanval = std::numeric_limits<double>::quiet_NaN();
-  double infval = std::numeric_limits<double>::infinity();
-  std::vector<double> h_expected{1234.0,
-                                 0,
-                                 -876.0,
-                                 543.2,
-                                 -0.12,
-                                 0.25,
-                                 -0.002,
-                                 0,
-                                 -0.0,
-                                 1.28e256,
-                                 nanval,
-                                 0,
-                                 123.0,
-                                 456.0,
-                                 -178000.0,
-                                 -122.33644781999999,
-                                 infval};
+  std::vector<double> h_expected;
+  std::for_each(h_strings.begin(), h_strings.end(), [&](const char* str) {
+    h_expected.push_back(str ? std::atof(str) : 0);
+  });
 
   auto strings_view = cudf::strings_column_view(strings);
   auto results = cudf::strings::to_floats(strings_view, cudf::data_type{cudf::type_id::FLOAT64});
@@ -161,7 +133,7 @@ TEST_F(StringsConvertTest, ToFloats64)
     h_expected.begin(),
     h_expected.end(),
     thrust::make_transform_iterator(h_strings.begin(), [](auto str) { return str != nullptr; }));
-  CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*results, expected);
+  CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*results, expected, true);
 }
 
 TEST_F(StringsConvertTest, FromFloats64)
@@ -190,7 +162,7 @@ TEST_F(StringsConvertTest, FromFloats64)
     h_expected.end(),
     thrust::make_transform_iterator(h_expected.begin(), [](auto str) { return str != nullptr; }));
 
-  CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*results, expected);
+  CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*results, expected, true);
 }
 
 TEST_F(StringsConvertTest, ZeroSizeStringsColumnFloat)

cpp/tests/utilities/column_utilities.cu

@@ -127,11 +127,20 @@ class corresponding_rows_not_equivalent {
       column_device_view const& lhs, column_device_view const& rhs, size_type index)
     {
       if (lhs.is_valid(index) and rhs.is_valid(index)) {
-        int ulp = 4;  // value taken from google test
-        T x     = lhs.element<T>(index);
-        T y     = rhs.element<T>(index);
-        return std::abs(x - y) > std::numeric_limits<T>::epsilon() * std::abs(x + y) * ulp &&
-               std::abs(x - y) >= std::numeric_limits<T>::min();
+        T const x = lhs.element<T>(index);
+        T const y = rhs.element<T>(index);
+
+        // Must handle inf and nan separately
+        if (std::isinf(x) || std::isinf(y)) {
+          return x != y;  // comparison of (inf==inf) returns true
+        } else if (std::isnan(x) || std::isnan(y)) {
+          return std::isnan(x) != std::isnan(y);  // comparison of (nan==nan) returns false
+        } else {
+          constexpr int ulp     = 4;  // ulp = unit of least precision, value taken from google test
+          T const abs_x_minus_y = std::abs(x - y);
+          return abs_x_minus_y >= std::numeric_limits<T>::min() &&
+                 abs_x_minus_y > std::numeric_limits<T>::epsilon() * std::abs(x + y) * ulp;
+        }
       } else {
         // if either is null, then the inequality was checked already
         return true;