Merge branch 'branch-22.04' into fix-reduce-pytest-time

.pre-commit-config.yaml

@@ -88,6 +88,13 @@ repos:
                 # of dependencies, so we'll have to update this manually.
                 additional_dependencies:
                   - cmake-format==0.6.11
+              - id: copyright-check
+                name: copyright-check
+                # This hook's use of Git tools appears to conflict with
+                # existing CI invocations so we don't invoke it during CI runs.
+                stages: [commit]
+                entry: python ./ci/checks/copyright.py --git-modified-only
+                language: python
 
 default_language_version:
       python: python3

cpp/benchmarks/binaryop/compiled_binaryop.cpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-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.
@@ -27,8 +27,8 @@
 class COMPILED_BINARYOP : public cudf::benchmark {
 };
 
-template <typename TypeLhs, typename TypeRhs, typename TypeOut, cudf::binary_operator binop>
-void BM_compiled_binaryop(benchmark::State& state)
+template <typename TypeLhs, typename TypeRhs, typename TypeOut>
+void BM_compiled_binaryop(benchmark::State& state, cudf::binary_operator binop)
 {
   const cudf::size_type column_size{(cudf::size_type)state.range(0)};
 
@@ -50,21 +50,26 @@ void BM_compiled_binaryop(benchmark::State& state)
 }
 
 // TODO tparam boolean for null.
-#define BINARYOP_BENCHMARK_DEFINE(TypeLhs, TypeRhs, binop, TypeOut) \
-  TEMPLATED_BENCHMARK_F(COMPILED_BINARYOP,                          \
-                        BM_compiled_binaryop,                       \
-                        TypeLhs,                                    \
-                        TypeRhs,                                    \
-                        TypeOut,                                    \
-                        cudf::binary_operator::binop)               \
-    ->Unit(benchmark::kMicrosecond)                                 \
-    ->UseManualTime()                                               \
-    ->Arg(10000)      /* 10k */                                     \
-    ->Arg(100000)     /* 100k */                                    \
-    ->Arg(1000000)    /* 1M */                                      \
-    ->Arg(10000000)   /* 10M */                                     \
+#define BM_BINARYOP_BENCHMARK_DEFINE(name, lhs, rhs, bop, tout)           \
+  BENCHMARK_DEFINE_F(COMPILED_BINARYOP, name)                             \
+  (::benchmark::State & st)                                               \
+  {                                                                       \
+    BM_compiled_binaryop<lhs, rhs, tout>(st, cudf::binary_operator::bop); \
+  }                                                                       \
+  BENCHMARK_REGISTER_F(COMPILED_BINARYOP, name)                           \
+    ->Unit(benchmark::kMicrosecond)                                       \
+    ->UseManualTime()                                                     \
+    ->Arg(10000)      /* 10k */                                           \
+    ->Arg(100000)     /* 100k */                                          \
+    ->Arg(1000000)    /* 1M */                                            \
+    ->Arg(10000000)   /* 10M */                                           \
     ->Arg(100000000); /* 100M */
 
+#define build_name(a, b, c, d) a##_##b##_##c##_##d
+
+#define BINARYOP_BENCHMARK_DEFINE(lhs, rhs, bop, tout) \
+  BM_BINARYOP_BENCHMARK_DEFINE(build_name(bop, lhs, rhs, tout), lhs, rhs, bop, tout)
+
 using namespace cudf;
 using namespace numeric;
 

cpp/benchmarks/string/url_decode.cpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-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.
@@ -70,8 +70,7 @@ cudf::test::strings_column_wrapper generate_column(cudf::size_type num_rows,
 class UrlDecode : public cudf::benchmark {
 };
 
-template <int esc_seq_pct>
-void BM_url_decode(benchmark::State& state)
+void BM_url_decode(benchmark::State& state, int esc_seq_pct)
 {
   cudf::size_type const num_rows      = state.range(0);
   cudf::size_type const chars_per_row = state.range(1);
@@ -88,12 +87,14 @@ void BM_url_decode(benchmark::State& state)
                           (chars_per_row + sizeof(cudf::size_type)));
 }
 
-#define URLD_BENCHMARK_DEFINE(esc_seq_pct)                     \
-  TEMPLATED_BENCHMARK_F(UrlDecode, BM_url_decode, esc_seq_pct) \
-    ->Args({100000000, 10})                                    \
-    ->Args({10000000, 100})                                    \
-    ->Args({1000000, 1000})                                    \
-    ->Unit(benchmark::kMillisecond)                            \
+#define URLD_BENCHMARK_DEFINE(esc_seq_pct)                      \
+  BENCHMARK_DEFINE_F(UrlDecode, esc_seq_pct)                    \
+  (::benchmark::State & st) { BM_url_decode(st, esc_seq_pct); } \
+  BENCHMARK_REGISTER_F(UrlDecode, esc_seq_pct)                  \
+    ->Args({100000000, 10})                                     \
+    ->Args({10000000, 100})                                     \
+    ->Args({1000000, 1000})                                     \
+    ->Unit(benchmark::kMillisecond)                             \
     ->UseManualTime();
 
 URLD_BENCHMARK_DEFINE(10)

cpp/src/groupby/sort/group_single_pass_reduction_util.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2019-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.
@@ -74,12 +74,16 @@ struct element_arg_minmax_fn {
 /**
  * @brief Value accessor for column which supports dictionary column too.
  *
+ * This is similar to `value_accessor` in `column_device_view.cuh` but with support of dictionary
+ * type.
+ *
  * @tparam T Type of the underlying column. For dictionary column, type of the key column.
  */
 template <typename T>
 struct value_accessor {
   column_device_view const col;
   bool const is_dict;
+
   value_accessor(column_device_view const& col) : col(col), is_dict(cudf::is_dictionary(col.type()))
   {
   }
@@ -93,27 +97,36 @@ struct value_accessor {
       return col.element<T>(i);
     }
   }
+
   __device__ auto operator()(size_type i) const { return value(i); }
 };
 
 /**
  * @brief Null replaced value accessor for column which supports dictionary column too.
  * For null value, returns null `init` value
  *
- * @tparam T Type of the underlying column. For dictionary column, type of the key column.
+ * @tparam SourceType Type of the underlying column. For dictionary column, type of the key column.
+ * @tparam TargetType Type that is used for computation.
  */
-template <typename T>
-struct null_replaced_value_accessor : value_accessor<T> {
-  using super_t = value_accessor<T>;
+template <typename SourceType, typename TargetType>
+struct null_replaced_value_accessor : value_accessor<SourceType> {
+  using super_t = value_accessor<SourceType>;
+
+  TargetType const init;
   bool const has_nulls;
-  T const init;
-  null_replaced_value_accessor(column_device_view const& col, T const& init, bool const has_nulls)
+
+  null_replaced_value_accessor(column_device_view const& col,
+                               TargetType const& init,
+                               bool const has_nulls)
     : super_t(col), init(init), has_nulls(has_nulls)
   {
   }
-  __device__ T operator()(size_type i) const
+
+  __device__ TargetType operator()(size_type i) const
   {
-    return has_nulls && super_t::col.is_null_nocheck(i) ? init : super_t::value(i);
+    return has_nulls && super_t::col.is_null_nocheck(i)
+             ? init
+             : static_cast<TargetType>(super_t::value(i));
   }
 };
 
@@ -168,7 +181,7 @@ struct group_reduction_functor<K, T, std::enable_if_t<is_group_reduction_support
                                         rmm::mr::device_memory_resource* mr)
 
   {
-    using DeviceType  = device_storage_type_t<T>;
+    using SourceDType = device_storage_type_t<T>;
     using ResultType  = cudf::detail::target_type_t<T, K>;
     using ResultDType = device_storage_type_t<ResultType>;
 
@@ -203,9 +216,11 @@ struct group_reduction_functor<K, T, std::enable_if_t<is_group_reduction_support
       do_reduction(count_iter, result_begin, binop);
     } else {
       using OpType    = cudf::detail::corresponding_operator_t<K>;
-      auto init       = OpType::template identity<DeviceType>();
+      auto init       = OpType::template identity<ResultDType>();
       auto inp_values = cudf::detail::make_counting_transform_iterator(
-        0, null_replaced_value_accessor{*d_values_ptr, init, values.has_nulls()});
+        0,
+        null_replaced_value_accessor<SourceDType, ResultDType>{
+          *d_values_ptr, init, values.has_nulls()});
       do_reduction(inp_values, result_begin, OpType{});
     }
 

cpp/src/io/utilities/file_io_utilities.cpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-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.
@@ -194,20 +194,13 @@ template <typename DataT,
 std::vector<std::future<ResultT>> make_sliced_tasks(
   F function, DataT* ptr, size_t offset, size_t size, cudf::detail::thread_pool& pool)
 {
+  constexpr size_t default_max_slice_size = 4 * 1024 * 1024;
+  static auto const max_slice_size = getenv_or("LIBCUDF_CUFILE_SLICE_SIZE", default_max_slice_size);
+  auto const slices                = make_file_io_slices(size, max_slice_size);
   std::vector<std::future<ResultT>> slice_tasks;
-  constexpr size_t default_max_slice_bytes = 4 * 1024 * 1024;
-  static auto const max_slice_bytes =
-    getenv_or("LIBCUDF_CUFILE_SLICE_SIZE", default_max_slice_bytes);
-  size_t const n_slices = util::div_rounding_up_safe(size, max_slice_bytes);
-  size_t slice_offset   = 0;
-  for (size_t t = 0; t < n_slices; ++t) {
-    DataT* ptr_slice = ptr + slice_offset;
-
-    size_t const slice_size = (t == n_slices - 1) ? size % max_slice_bytes : max_slice_bytes;
-    slice_tasks.push_back(pool.submit(function, ptr_slice, slice_size, offset + slice_offset));
-
-    slice_offset += slice_size;
-  }
+  std::transform(slices.cbegin(), slices.cend(), std::back_inserter(slice_tasks), [&](auto& slice) {
+    return pool.submit(function, ptr + slice.offset, slice.size, offset + slice.offset);
+  });
   return slice_tasks;
 }
 
@@ -318,6 +311,21 @@ std::unique_ptr<cufile_output_impl> make_cufile_output(std::string const& filepa
   return nullptr;
 }
 
+std::vector<file_io_slice> make_file_io_slices(size_t size, size_t max_slice_size)
+{
+  max_slice_size      = std::max(1024ul, max_slice_size);
+  auto const n_slices = util::div_rounding_up_safe(size, max_slice_size);
+  std::vector<file_io_slice> slices;
+  slices.reserve(n_slices);
+  std::generate_n(std::back_inserter(slices), n_slices, [&, idx = 0]() mutable {
+    auto const slice_offset = idx++ * max_slice_size;
+    auto const slice_size   = std::min(size - slice_offset, max_slice_size);
+    return file_io_slice{slice_offset, slice_size};
+  });
+
+  return slices;
+}
+
 }  // namespace detail
 }  // namespace io
 }  // namespace cudf

cpp/src/io/utilities/file_io_utilities.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-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.
@@ -291,6 +291,21 @@ std::unique_ptr<cufile_input_impl> make_cufile_input(std::string const& filepath
  */
 std::unique_ptr<cufile_output_impl> make_cufile_output(std::string const& filepath);
 
+/**
+ * @brief Byte range to be read/written in a single operation.
+ */
+struct file_io_slice {
+  size_t offset;
+  size_t size;
+};
+
+/**
+ * @brief Split the total number of bytes to read/write into slices to enable parallel IO.
+ *
+ * If `max_slice_size` is below 1024, 1024 will be used instead to prevent potential misuse.
+ */
+std::vector<file_io_slice> make_file_io_slices(size_t size, size_t max_slice_size);
+
 }  // namespace detail
 }  // namespace io
 }  // namespace cudf

cpp/tests/CMakeLists.txt

@@ -199,6 +199,7 @@ ConfigureTest(
 ConfigureTest(DECOMPRESSION_TEST io/comp/decomp_test.cpp)
 
 ConfigureTest(CSV_TEST io/csv_test.cpp)
+ConfigureTest(FILE_IO_TEST io/file_io_test.cpp)
 ConfigureTest(ORC_TEST io/orc_test.cpp)
 ConfigureTest(PARQUET_TEST io/parquet_test.cpp)
 ConfigureTest(JSON_TEST io/json_test.cpp)

cpp/tests/groupby/sum_tests.cpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2019-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.
@@ -156,6 +156,27 @@ TYPED_TEST(groupby_sum_test, dictionary)
                   force_use_sort_impl::YES);
 }
 
+struct overflow_test : public cudf::test::BaseFixture {
+};
+TEST_F(overflow_test, overflow_integer)
+{
+  using int32_col = fixed_width_column_wrapper<int32_t>;
+  using int64_col = fixed_width_column_wrapper<int64_t>;
+
+  auto const keys        = int32_col{0, 0};
+  auto const vals        = int32_col{-2147483648, -2147483648};
+  auto const expect_keys = int32_col{0};
+  auto const expect_vals = int64_col{-4294967296L};
+
+  auto test_sum = [&](auto const use_sort) {
+    auto agg = make_sum_aggregation<groupby_aggregation>();
+    test_single_agg(keys, vals, expect_keys, expect_vals, std::move(agg), use_sort);
+  };
+
+  test_sum(force_use_sort_impl::NO);
+  test_sum(force_use_sort_impl::YES);
+}
+
 template <typename T>
 struct FixedPointTestAllReps : public cudf::test::BaseFixture {
 };

cpp/tests/io/file_io_test.cpp

@@ -0,0 +1,46 @@
+/*
+ * 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 <cudf_test/base_fixture.hpp>
+#include <cudf_test/cudf_gtest.hpp>
+
+#include <src/io/utilities/file_io_utilities.hpp>
+
+#include <type_traits>
+
+// Base test fixture for tests
+struct CuFileIOTest : public cudf::test::BaseFixture {
+};
+
+TEST_F(CuFileIOTest, SliceSize)
+{
+  std::vector<std::pair<size_t, size_t>> test_cases{
+    {1 << 20, 1 << 18}, {1 << 18, 1 << 20}, {1 << 20, 3333}, {0, 1 << 18}, {0, 0}, {1 << 20, 0}};
+  for (auto const& test_case : test_cases) {
+    auto const slices = cudf::io::detail::make_file_io_slices(test_case.first, test_case.second);
+    if (slices.empty()) {
+      ASSERT_EQ(test_case.first, 0);
+    } else {
+      ASSERT_EQ(slices.front().offset, 0);
+      ASSERT_EQ(slices.back().offset + slices.back().size, test_case.first);
+      for (auto i = 1u; i < slices.size(); ++i) {
+        ASSERT_EQ(slices[i].offset, slices[i - 1].offset + slices[i - 1].size);
+      }
+    }
+  }
+}
+
+CUDF_TEST_PROGRAM_MAIN()