Update to official libcu++ on Github(#6275)

Update to libcu++ on Github. Authors: - ptaylor <[email protected]> - Paul Taylor <[email protected]> Approvers: - Mark Harris - Keith Kraus - Christopher Harris - Mark Harris URL: #6275
rapidsai · Dec 8, 2020 · 78f9789 · 78f9789
1 parent 9120992
commit 78f9789
Show file tree

Hide file tree

Showing 38 changed files with 399 additions and 332 deletions.
diff --git a/CHANGELOG.md b/CHANGELOG.md
@@ -9,6 +9,7 @@
 
 ## Improvements
 
+- PR #6275 Update to official libcu++ on Github
 - PR #6838 Fix `columns` & `index` handling in dataframe constructor
 
 ## Bug Fixes

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
diff --git a/cpp/benchmarks/common/generate_benchmark_input.cpp b/cpp/benchmarks/common/generate_benchmark_input.cpp
@@ -121,11 +121,11 @@ struct random_value_fn<T, typename std::enable_if_t<cudf::is_chrono<T>()>> {
 
   random_value_fn(distribution_params<T> params)
   {
-    using simt::std::chrono::duration_cast;
+    using cuda::std::chrono::duration_cast;
 
     std::pair<cudf::duration_s, cudf::duration_s> const range_s = {
-      duration_cast<simt::std::chrono::seconds>(typename T::duration{params.lower_bound}),
-      duration_cast<simt::std::chrono::seconds>(typename T::duration{params.upper_bound})};
+      duration_cast<cuda::std::chrono::seconds>(typename T::duration{params.lower_bound}),
+      duration_cast<cuda::std::chrono::seconds>(typename T::duration{params.upper_bound})};
     if (range_s.first != range_s.second) {
       seconds_gen =
         make_distribution<int64_t>(params.id, range_s.first.count(), range_s.second.count());
@@ -149,7 +149,7 @@ struct random_value_fn<T, typename std::enable_if_t<cudf::is_chrono<T>()>> {
     auto const timestamp_ns =
       cudf::duration_s{seconds_gen(engine)} + cudf::duration_ns{nanoseconds_gen(engine)};
     // Return value in the type's precision
-    return T(simt::std::chrono::duration_cast<typename T::duration>(timestamp_ns));
+    return T(cuda::std::chrono::duration_cast<typename T::duration>(timestamp_ns));
   }
 };
 

diff --git a/cpp/benchmarks/common/generate_benchmark_input.hpp b/cpp/benchmarks/common/generate_benchmark_input.hpp
@@ -93,7 +93,7 @@ distribution_id default_distribution_id()
 template <typename T, std::enable_if_t<cudf::is_timestamp<T>()>* = nullptr>
 std::pair<int64_t, int64_t> default_range()
 {
-  using simt::std::chrono::duration_cast;
+  using cuda::std::chrono::duration_cast;
   auto const year = duration_cast<typename T::duration>(cudf::duration_D{365l});
   return {50 * year.count(), 0};
 }
@@ -106,7 +106,7 @@ std::pair<int64_t, int64_t> default_range()
 template <typename T, std::enable_if_t<cudf::is_duration<T>()>* = nullptr>
 std::pair<int64_t, int64_t> default_range()
 {
-  using simt::std::chrono::duration_cast;
+  using cuda::std::chrono::duration_cast;
   auto const year = duration_cast<typename T::duration>(cudf::duration_D{365l});
   return {0, 2 * year.count()};
 }

diff --git a/cpp/include/cudf/ast/detail/operators.hpp b/cpp/include/cudf/ast/detail/operators.hpp
@@ -15,12 +15,14 @@
  */
 #pragma once
 
-#include <cmath>
 #include <cudf/ast/operators.hpp>
 #include <cudf/types.hpp>
 #include <cudf/utilities/error.hpp>
 #include <cudf/utilities/type_dispatcher.hpp>
-#include <simt/type_traits>
+
+#include <cuda/std/type_traits>
+
+#include <cmath>
 #include <type_traits>
 #include <utility>
 #include <vector>
@@ -33,10 +35,10 @@ namespace detail {
 
 // Traits for valid operator / type combinations
 template <typename Op, typename LHS, typename RHS>
-constexpr bool is_valid_binary_op = simt::std::is_invocable<Op, LHS, RHS>::value;
+constexpr bool is_valid_binary_op = cuda::std::is_invocable<Op, LHS, RHS>::value;
 
 template <typename Op, typename T>
-constexpr bool is_valid_unary_op = simt::std::is_invocable<Op, T>::value;
+constexpr bool is_valid_unary_op = cuda::std::is_invocable<Op, T>::value;
 
 /**
  * @brief Operator dispatcher
@@ -983,7 +985,7 @@ struct return_type_functor {
             std::enable_if_t<is_valid_binary_op<OperatorFunctor, LHS, RHS>>* = nullptr>
   CUDA_HOST_DEVICE_CALLABLE void operator()(cudf::data_type& result)
   {
-    using Out = simt::std::invoke_result_t<OperatorFunctor, LHS, RHS>;
+    using Out = cuda::std::invoke_result_t<OperatorFunctor, LHS, RHS>;
     result    = cudf::data_type(cudf::type_to_id<Out>());
   }
 
@@ -1012,7 +1014,7 @@ struct return_type_functor {
             std::enable_if_t<is_valid_unary_op<OperatorFunctor, T>>* = nullptr>
   CUDA_HOST_DEVICE_CALLABLE void operator()(cudf::data_type& result)
   {
-    using Out = simt::std::invoke_result_t<OperatorFunctor, T>;
+    using Out = cuda::std::invoke_result_t<OperatorFunctor, T>;
     result    = cudf::data_type(cudf::type_to_id<Out>());
   }
 

diff --git a/cpp/include/cudf/ast/detail/transform.cuh b/cpp/include/cudf/ast/detail/transform.cuh
@@ -76,7 +76,7 @@ struct unary_row_output : public row_output {
                              detail::device_data_reference output) const
   {
     using OperatorFunctor = detail::operator_functor<op>;
-    using Out             = simt::std::invoke_result_t<OperatorFunctor, Input>;
+    using Out             = cuda::std::invoke_result_t<OperatorFunctor, Input>;
     resolve_output<Out>(output, row_index, OperatorFunctor{}(input));
   }
 
@@ -104,7 +104,7 @@ struct binary_row_output : public row_output {
                              detail::device_data_reference output) const
   {
     using OperatorFunctor = detail::operator_functor<op>;
-    using Out             = simt::std::invoke_result_t<OperatorFunctor, LHS, RHS>;
+    using Out             = cuda::std::invoke_result_t<OperatorFunctor, LHS, RHS>;
     resolve_output<Out>(output, row_index, OperatorFunctor{}(lhs, rhs));
   }
 

diff --git a/cpp/include/cudf/detail/copy_if.cuh b/cpp/include/cudf/detail/copy_if.cuh
@@ -200,6 +200,26 @@ __launch_bounds__(block_size) __global__
   }
 }
 
+template <typename T, typename Enable = void>
+struct DeviceType {
+  using type = T;
+};
+
+template <typename T>
+struct DeviceType<T, std::enable_if_t<cudf::is_timestamp<T>()>> {
+  using type = typename T::rep;
+};
+
+template <typename T>
+struct DeviceType<T, std::enable_if_t<std::is_same<numeric::decimal32, T>::value>> {
+  using type = typename cudf::device_storage_type_t<T>;
+};
+
+template <typename T>
+struct DeviceType<T, std::enable_if_t<std::is_same<numeric::decimal64, T>::value>> {
+  using type = typename cudf::device_storage_type_t<T>;
+};
+
 // Dispatch functor which performs the scatter for fixed column types and gather for other
 template <typename Filter, int block_size>
 struct scatter_gather_functor {
@@ -219,7 +239,8 @@ struct scatter_gather_functor {
 
     bool has_valid = input.nullable();
 
-    using Type   = cudf::device_storage_type_t<T>;
+    using Type = typename DeviceType<T>::type;
+
     auto scatter = (has_valid) ? scatter_kernel<Type, Filter, block_size, true>
                                : scatter_kernel<Type, Filter, block_size, false>;
 

diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp
@@ -16,17 +16,16 @@
 
 #pragma once
 
-#define _LIBCUDACXX_USE_CXX17_TYPE_TRAITS
+#include <cudf/types.hpp>
 
-// Note: The <simt/*> versions are used in order for Jitify to work with our fixed_point type.
+// Note: The <cuda/std/*> versions are used in order for Jitify to work with our fixed_point type.
 //       Jitify is needed for several algorithms (binaryop, rolling, etc)
-#include <simt/limits>
-#include <simt/type_traits>  // add simt namespace
+#include <cuda/std/limits>
+#include <cuda/std/type_traits>  // add cuda namespace
 
 #include <algorithm>
 #include <cassert>
 #include <cmath>
-
 #include <string>
 
 //! `fixed_point` and supporting types
@@ -60,13 +59,13 @@ enum class Radix : int32_t { BASE_2 = 2, BASE_10 = 10 };
 template <typename T>
 constexpr inline auto is_supported_representation_type()
 {
-  return simt::std::is_same<T, int32_t>::value || simt::std::is_same<T, int64_t>::value;
+  return cuda::std::is_same<T, int32_t>::value || cuda::std::is_same<T, int64_t>::value;
 }
 
 template <typename T>
 constexpr inline auto is_supported_construction_value_type()
 {
-  return simt::std::is_integral<T>::value || simt::std::is_floating_point<T>::value;
+  return cuda::std::is_integral<T>::value || cuda::std::is_floating_point<T>::value;
 }
 
 // Helper functions for `fixed_point` type
@@ -87,7 +86,7 @@ namespace detail {
 template <typename Rep,
           Radix Base,
           typename T,
-          typename simt::std::enable_if_t<(simt::std::is_same<int32_t, T>::value &&
+          typename cuda::std::enable_if_t<(cuda::std::is_same<int32_t, T>::value &&
                                            is_supported_representation_type<Rep>())>* = nullptr>
 CUDA_HOST_DEVICE_CALLABLE Rep ipow(T exponent)
 {
@@ -192,7 +191,7 @@ CUDA_HOST_DEVICE_CALLABLE constexpr T shift(T const& val, scale_type const& scal
  * @tparam Rep The representation type (either `int32_t` or `int64_t`)
  */
 template <typename Rep,
-          typename simt::std::enable_if_t<is_supported_representation_type<Rep>()>* = nullptr>
+          typename cuda::std::enable_if_t<is_supported_representation_type<Rep>()>* = nullptr>
 struct scaled_integer {
   Rep value;
   scale_type scale;
@@ -224,7 +223,7 @@ class fixed_point {
    * @param scale The exponent that is applied to Rad to perform shifting
    */
   template <typename T,
-            typename simt::std::enable_if_t<is_supported_construction_value_type<T>() &&
+            typename cuda::std::enable_if_t<is_supported_construction_value_type<T>() &&
                                             is_supported_representation_type<Rep>()>* = nullptr>
   CUDA_HOST_DEVICE_CALLABLE explicit fixed_point(T const& value, scale_type const& scale)
     : _value{static_cast<Rep>(detail::shift<Rep, Rad>(value, scale))}, _scale{scale}
@@ -245,7 +244,7 @@ class fixed_point {
    * value and scale of zero
    */
   template <typename T,
-            typename simt::std::enable_if_t<is_supported_construction_value_type<T>()>* = nullptr>
+            typename cuda::std::enable_if_t<is_supported_construction_value_type<T>()>* = nullptr>
   CUDA_HOST_DEVICE_CALLABLE fixed_point(T const& value)
     : _value{static_cast<Rep>(value)}, _scale{scale_type{0}}
   {
@@ -265,7 +264,7 @@ class fixed_point {
    * @return The `fixed_point` number in base 10 (aka human readable format)
    */
   template <typename U,
-            typename simt::std::enable_if_t<is_supported_construction_value_type<U>()>* = nullptr>
+            typename cuda::std::enable_if_t<is_supported_construction_value_type<U>()>* = nullptr>
   CUDA_HOST_DEVICE_CALLABLE explicit constexpr operator U() const
   {
     return detail::shift<Rep, Rad>(static_cast<U>(_value), detail::negate(_scale));
@@ -538,9 +537,9 @@ class fixed_point {
 template <typename Rep>
 std::string print_rep()
 {
-  if (simt::std::is_same<Rep, int32_t>::value)
+  if (cuda::std::is_same<Rep, int32_t>::value)
     return "int32_t";
-  else if (simt::std::is_same<Rep, int64_t>::value)
+  else if (cuda::std::is_same<Rep, int64_t>::value)
     return "int64_t";
   else
     return "unknown type";
@@ -557,8 +556,8 @@ std::string print_rep()
 template <typename Rep, typename T>
 CUDA_HOST_DEVICE_CALLABLE auto addition_overflow(T lhs, T rhs)
 {
-  return rhs > 0 ? lhs > simt::std::numeric_limits<Rep>::max() - rhs
-                 : lhs < simt::std::numeric_limits<Rep>::min() - rhs;
+  return rhs > 0 ? lhs > cuda::std::numeric_limits<Rep>::max() - rhs
+                 : lhs < cuda::std::numeric_limits<Rep>::min() - rhs;
 }
 
 /** @brief Function for identifying integer overflow when subtracting
@@ -572,8 +571,8 @@ CUDA_HOST_DEVICE_CALLABLE auto addition_overflow(T lhs, T rhs)
 template <typename Rep, typename T>
 CUDA_HOST_DEVICE_CALLABLE auto subtraction_overflow(T lhs, T rhs)
 {
-  return rhs > 0 ? lhs < simt::std::numeric_limits<Rep>::min() + rhs
-                 : lhs > simt::std::numeric_limits<Rep>::max() + rhs;
+  return rhs > 0 ? lhs < cuda::std::numeric_limits<Rep>::min() + rhs
+                 : lhs > cuda::std::numeric_limits<Rep>::max() + rhs;
 }
 
 /** @brief Function for identifying integer overflow when dividing
@@ -587,7 +586,7 @@ CUDA_HOST_DEVICE_CALLABLE auto subtraction_overflow(T lhs, T rhs)
 template <typename Rep, typename T>
 CUDA_HOST_DEVICE_CALLABLE auto division_overflow(T lhs, T rhs)
 {
-  return lhs == simt::std::numeric_limits<Rep>::min() && rhs == -1;
+  return lhs == cuda::std::numeric_limits<Rep>::min() && rhs == -1;
 }
 
 /** @brief Function for identifying integer overflow when multiplying
@@ -601,8 +600,8 @@ CUDA_HOST_DEVICE_CALLABLE auto division_overflow(T lhs, T rhs)
 template <typename Rep, typename T>
 CUDA_HOST_DEVICE_CALLABLE auto multiplication_overflow(T lhs, T rhs)
 {
-  auto const min = simt::std::numeric_limits<Rep>::min();
-  auto const max = simt::std::numeric_limits<Rep>::max();
+  auto const min = cuda::std::numeric_limits<Rep>::min();
+  auto const max = cuda::std::numeric_limits<Rep>::max();
   if (rhs > 0)
     return lhs > max / rhs || lhs < min / rhs;
   else if (rhs < -1)

diff --git a/cpp/include/cudf/utilities/traits.hpp b/cpp/include/cudf/utilities/traits.hpp
@@ -24,7 +24,7 @@
 #include <cudf/wrappers/durations.hpp>
 #include <cudf/wrappers/timestamps.hpp>
 
-#include <simt/type_traits>
+#include <cuda/std/type_traits>
 
 namespace cudf {
 
@@ -66,14 +66,14 @@ struct is_equality_comparable_impl<L, R, void_t<equality_comparable<L, R>>> : st
 };
 
 template <typename T>
-using is_timestamp_t = simt::std::disjunction<std::is_same<cudf::timestamp_D, T>,
+using is_timestamp_t = cuda::std::disjunction<std::is_same<cudf::timestamp_D, T>,
                                               std::is_same<cudf::timestamp_s, T>,
                                               std::is_same<cudf::timestamp_ms, T>,
                                               std::is_same<cudf::timestamp_us, T>,
                                               std::is_same<cudf::timestamp_ns, T>>;
 
 template <typename T>
-using is_duration_t = simt::std::disjunction<std::is_same<cudf::duration_D, T>,
+using is_duration_t = cuda::std::disjunction<std::is_same<cudf::duration_D, T>,
                                              std::is_same<cudf::duration_s, T>,
                                              std::is_same<cudf::duration_ms, T>,
                                              std::is_same<cudf::duration_us, T>,

diff --git a/cpp/include/cudf/utilities/type_dispatcher.hpp b/cpp/include/cudf/utilities/type_dispatcher.hpp
@@ -17,14 +17,14 @@
 #pragma once
 
 #include <cudf/detail/utilities/release_assert.cuh>
+#include <cudf/fixed_point/fixed_point.hpp>
 #include <cudf/types.hpp>
 #include <cudf/utilities/error.hpp>
 #include <cudf/wrappers/dictionary.hpp>
 #include <cudf/wrappers/durations.hpp>
 #include <cudf/wrappers/timestamps.hpp>
-#include <string>
 
-#include <cudf/fixed_point/fixed_point.hpp>
+#include <string>
 
 /**
  * @file

diff --git a/cpp/include/cudf/wrappers/durations.hpp b/cpp/include/cudf/wrappers/durations.hpp
@@ -18,10 +18,7 @@
 
 #include <limits>
 
-#define _LIBCUDACXX_USE_CXX20_CHRONO
-#define _LIBCUDACXX_USE_CXX17_TYPE_TRAITS
-
-#include <simt/chrono>
+#include <cuda/std/chrono>
 
 namespace cudf {
 
@@ -35,23 +32,23 @@ namespace cudf {
 /**
  * @brief Type alias representing an int32_t duration of days.
  **/
-using duration_D = simt::std::chrono::duration<int32_t, simt::std::chrono::days::period>;
+using duration_D = cuda::std::chrono::duration<int32_t, cuda::std::chrono::days::period>;
 /**
  * @brief Type alias representing an int64_t duration of seconds.
  **/
-using duration_s = simt::std::chrono::duration<int64_t, simt::std::chrono::seconds::period>;
+using duration_s = cuda::std::chrono::duration<int64_t, cuda::std::chrono::seconds::period>;
 /**
  * @brief Type alias representing an int64_t duration of milliseconds.
  **/
-using duration_ms = simt::std::chrono::duration<int64_t, simt::std::chrono::milliseconds::period>;
+using duration_ms = cuda::std::chrono::duration<int64_t, cuda::std::chrono::milliseconds::period>;
 /**
  * @brief Type alias representing an int64_t duration of microseconds.
  **/
-using duration_us = simt::std::chrono::duration<int64_t, simt::std::chrono::microseconds::period>;
+using duration_us = cuda::std::chrono::duration<int64_t, cuda::std::chrono::microseconds::period>;
 /**
  * @brief Type alias representing an int64_t duration of nanoseconds.
  **/
-using duration_ns = simt::std::chrono::duration<int64_t, simt::std::chrono::nanoseconds::period>;
+using duration_ns = cuda::std::chrono::duration<int64_t, cuda::std::chrono::nanoseconds::period>;
 
 static_assert(sizeof(duration_D) == sizeof(typename duration_D::rep), "");
 static_assert(sizeof(duration_s) == sizeof(typename duration_s::rep), "");