Merge branch 'branch-21.06' of https://github.com/rapidsai/cudf into …

…groupshift-python

README.md

@@ -65,15 +65,15 @@ Please see the [Demo Docker Repository](https://hub.docker.com/r/rapidsai/rapids
 
 cuDF can be installed with conda ([miniconda](https://conda.io/miniconda.html), or the full [Anaconda distribution](https://www.anaconda.com/download)) from the `rapidsai` channel:
 
-For `cudf version == 21.06` :
+For `cudf version == 0.19.2` :
 ```bash
 # for CUDA 10.1
 conda install -c rapidsai -c nvidia -c numba -c conda-forge \
-    cudf=21.06 python=3.7 cudatoolkit=10.1
+    cudf=0.19 python=3.7 cudatoolkit=10.1
 
 # or, for CUDA 10.2
 conda install -c rapidsai -c nvidia -c numba -c conda-forge \
-    cudf=21.06 python=3.7 cudatoolkit=10.2
+    cudf=0.19 python=3.7 cudatoolkit=10.2
 
 ```
 

conda/recipes/cudf/meta.yaml

@@ -3,7 +3,7 @@
 {% set version = environ.get('GIT_DESCRIBE_TAG', '0.0.0.dev').lstrip('v') + environ.get('VERSION_SUFFIX', '') %}
 {% set minor_version =  version.split('.')[0] + '.' + version.split('.')[1] %}
 {% set py_version=environ.get('CONDA_PY', 36) %}
-{% set cuda_version='.'.join(environ.get('CUDA_VERSION', '10.1').split('.')[:2]) %}
+{% set cuda_version='.'.join(environ.get('CUDA', '10.1').split('.')[:2]) %}
 
 package:
   name: cudf

conda/recipes/cudf_kafka/meta.yaml

@@ -3,7 +3,7 @@
 {% set version = environ.get('GIT_DESCRIBE_TAG', '0.0.0.dev').lstrip('v') + environ.get('VERSION_SUFFIX', '') %}
 {% set minor_version =  version.split('.')[0] + '.' + version.split('.')[1] %}
 {% set py_version=environ.get('CONDA_PY', 36) %}
-{% set cuda_version='.'.join(environ.get('CUDA_VERSION', '10.1').split('.')[:2]) %}
+{% set cuda_version='.'.join(environ.get('CUDA', '10.1').split('.')[:2]) %}
 
 package:
   name: cudf_kafka

conda/recipes/custreamz/meta.yaml

@@ -3,7 +3,7 @@
 {% set version = environ.get('GIT_DESCRIBE_TAG', '0.0.0.dev').lstrip('v') + environ.get('VERSION_SUFFIX', '') %}
 {% set minor_version =  version.split('.')[0] + '.' + version.split('.')[1] %}
 {% set py_version=environ.get('CONDA_PY', 36) %}
-{% set cuda_version='.'.join(environ.get('CUDA_VERSION', '10.1').split('.')[:2]) %}
+{% set cuda_version='.'.join(environ.get('CUDA', '10.1').split('.')[:2]) %}
 
 package:
   name: custreamz

conda/recipes/dask-cudf/meta.yaml

@@ -3,7 +3,7 @@
 {% set version = environ.get('GIT_DESCRIBE_TAG', '0.0.0.dev').lstrip('v') + environ.get('VERSION_SUFFIX', '') %}
 {% set minor_version =  version.split('.')[0] + '.' + version.split('.')[1] %}
 {% set py_version=environ.get('CONDA_PY', 36) %}
-{% set cuda_version='.'.join(environ.get('CUDA_VERSION', '10.1').split('.')[:2]) %}
+{% set cuda_version='.'.join(environ.get('CUDA', '10.1').split('.')[:2]) %}
 
 package:
   name: dask-cudf

conda/recipes/libcudf/meta.yaml

@@ -2,7 +2,7 @@
 
 {% set version = environ.get('GIT_DESCRIBE_TAG', '0.0.0.dev').lstrip('v') + environ.get('VERSION_SUFFIX', '') %}
 {% set minor_version =  version.split('.')[0] + '.' + version.split('.')[1] %}
-{% set cuda_version='.'.join(environ.get('CUDA_VERSION', '10.1').split('.')[:2]) %}
+{% set cuda_version='.'.join(environ.get('CUDA', '10.1').split('.')[:2]) %}
 
 package:
   name: libcudf
@@ -74,6 +74,7 @@ test:
     - test -f $PREFIX/include/cudf/detail/gather.hpp
     - test -f $PREFIX/include/cudf/detail/groupby.hpp
     - test -f $PREFIX/include/cudf/detail/groupby/sort_helper.hpp
+    - test -f $PREFIX/include/cudf/detail/groupby/group_replace_nulls.hpp
     - test -f $PREFIX/include/cudf/detail/hashing.hpp
     - test -f $PREFIX/include/cudf/detail/interop.hpp
     - test -f $PREFIX/include/cudf/detail/is_element_valid.hpp

cpp/CMakeLists.txt

@@ -209,6 +209,7 @@ add_library(cudf
     src/groupby/sort/group_max_scan.cu
     src/groupby/sort/group_min_scan.cu
     src/groupby/sort/group_sum_scan.cu
+    src/groupby/sort/group_replace_nulls.cu
     src/groupby/sort/sort_helper.cu
     src/hash/hashing.cu
     src/hash/md5_hash.cu
@@ -333,8 +334,8 @@ add_library(cudf
     src/strings/char_types/char_cases.cu
     src/strings/char_types/char_types.cu
     src/strings/combine/concatenate.cu
-    src/strings/combine/concatenate_list_elements.cu
     src/strings/combine/join.cu
+    src/strings/combine/join_list_elements.cu
     src/strings/contains.cu
     src/strings/convert/convert_booleans.cu
     src/strings/convert/convert_datetime.cu

cpp/benchmarks/join/join_benchmark.cu

@@ -37,8 +37,8 @@ template <typename key_type, typename payload_type>
 class Join : public cudf::benchmark {
 };
 
-template <typename key_type, typename payload_type, bool Nullable>
-static void BM_join(benchmark::State &state)
+template <typename key_type, typename payload_type, bool Nullable, typename Join>
+static void BM_join(benchmark::State& state, Join JoinFunc)
 {
   const cudf::size_type build_table_size{(cudf::size_type)state.range(0)};
   const cudf::size_type probe_table_size{(cudf::size_type)state.range(1)};
@@ -105,20 +105,69 @@ static void BM_join(benchmark::State &state)
   for (auto _ : state) {
     cuda_event_timer raii(state, true, rmm::cuda_stream_default);
 
-    auto result = cudf::inner_join(
+    auto result = JoinFunc(
       probe_table, build_table, columns_to_join, columns_to_join, cudf::null_equality::UNEQUAL);
   }
 }
 
-#define JOIN_BENCHMARK_DEFINE(name, key_type, payload_type, nullable) \
-  BENCHMARK_TEMPLATE_DEFINE_F(Join, name, key_type, payload_type)     \
-  (::benchmark::State & st) { BM_join<key_type, payload_type, nullable>(st); }
+#define JOIN_BENCHMARK_DEFINE(name, key_type, payload_type, nullable)         \
+  BENCHMARK_TEMPLATE_DEFINE_F(Join, name, key_type, payload_type)             \
+  (::benchmark::State & st)                                                   \
+  {                                                                           \
+    auto join = [](cudf::table_view const& left,                              \
+                   cudf::table_view const& right,                             \
+                   std::vector<cudf::size_type> const& left_on,               \
+                   std::vector<cudf::size_type> const& right_on,              \
+                   cudf::null_equality compare_nulls) {                       \
+      return cudf::inner_join(left, right, left_on, right_on, compare_nulls); \
+    };                                                                        \
+    BM_join<key_type, payload_type, nullable>(st, join);                      \
+  }
 
 JOIN_BENCHMARK_DEFINE(join_32bit, int32_t, int32_t, false);
 JOIN_BENCHMARK_DEFINE(join_64bit, int64_t, int64_t, false);
 JOIN_BENCHMARK_DEFINE(join_32bit_nulls, int32_t, int32_t, true);
 JOIN_BENCHMARK_DEFINE(join_64bit_nulls, int64_t, int64_t, true);
 
+#define LEFT_ANTI_JOIN_BENCHMARK_DEFINE(name, key_type, payload_type, nullable)   \
+  BENCHMARK_TEMPLATE_DEFINE_F(Join, name, key_type, payload_type)                 \
+  (::benchmark::State & st)                                                       \
+  {                                                                               \
+    auto join = [](cudf::table_view const& left,                                  \
+                   cudf::table_view const& right,                                 \
+                   std::vector<cudf::size_type> const& left_on,                   \
+                   std::vector<cudf::size_type> const& right_on,                  \
+                   cudf::null_equality compare_nulls) {                           \
+      return cudf::left_anti_join(left, right, left_on, right_on, compare_nulls); \
+    };                                                                            \
+    BM_join<key_type, payload_type, nullable>(st, join);                          \
+  }
+
+LEFT_ANTI_JOIN_BENCHMARK_DEFINE(left_anti_join_32bit, int32_t, int32_t, false);
+LEFT_ANTI_JOIN_BENCHMARK_DEFINE(left_anti_join_64bit, int64_t, int64_t, false);
+LEFT_ANTI_JOIN_BENCHMARK_DEFINE(left_anti_join_32bit_nulls, int32_t, int32_t, true);
+LEFT_ANTI_JOIN_BENCHMARK_DEFINE(left_anti_join_64bit_nulls, int64_t, int64_t, true);
+
+#define LEFT_SEMI_JOIN_BENCHMARK_DEFINE(name, key_type, payload_type, nullable)   \
+  BENCHMARK_TEMPLATE_DEFINE_F(Join, name, key_type, payload_type)                 \
+  (::benchmark::State & st)                                                       \
+  {                                                                               \
+    auto join = [](cudf::table_view const& left,                                  \
+                   cudf::table_view const& right,                                 \
+                   std::vector<cudf::size_type> const& left_on,                   \
+                   std::vector<cudf::size_type> const& right_on,                  \
+                   cudf::null_equality compare_nulls) {                           \
+      return cudf::left_semi_join(left, right, left_on, right_on, compare_nulls); \
+    };                                                                            \
+    BM_join<key_type, payload_type, nullable>(st, join);                          \
+  }
+
+LEFT_SEMI_JOIN_BENCHMARK_DEFINE(left_semi_join_32bit, int32_t, int32_t, false);
+LEFT_SEMI_JOIN_BENCHMARK_DEFINE(left_semi_join_64bit, int64_t, int64_t, false);
+LEFT_SEMI_JOIN_BENCHMARK_DEFINE(left_semi_join_32bit_nulls, int32_t, int32_t, true);
+LEFT_SEMI_JOIN_BENCHMARK_DEFINE(left_semi_join_64bit_nulls, int64_t, int64_t, true);
+
+// join -----------------------------------------------------------------------
 BENCHMARK_REGISTER_F(Join, join_32bit)
   ->Unit(benchmark::kMillisecond)
   ->Args({100'000, 100'000})
@@ -154,3 +203,77 @@ BENCHMARK_REGISTER_F(Join, join_64bit_nulls)
   ->Args({50'000'000, 50'000'000})
   ->Args({40'000'000, 120'000'000})
   ->UseManualTime();
+
+// left anti-join -------------------------------------------------------------
+BENCHMARK_REGISTER_F(Join, left_anti_join_32bit)
+  ->Unit(benchmark::kMillisecond)
+  ->Args({100'000, 100'000})
+  ->Args({100'000, 400'000})
+  ->Args({100'000, 1'000'000})
+  ->Args({10'000'000, 10'000'000})
+  ->Args({10'000'000, 40'000'000})
+  ->Args({10'000'000, 100'000'000})
+  ->Args({100'000'000, 100'000'000})
+  ->Args({80'000'000, 240'000'000})
+  ->UseManualTime();
+
+BENCHMARK_REGISTER_F(Join, left_anti_join_64bit)
+  ->Unit(benchmark::kMillisecond)
+  ->Args({50'000'000, 50'000'000})
+  ->Args({40'000'000, 120'000'000})
+  ->UseManualTime();
+
+BENCHMARK_REGISTER_F(Join, left_anti_join_32bit_nulls)
+  ->Unit(benchmark::kMillisecond)
+  ->Args({100'000, 100'000})
+  ->Args({100'000, 400'000})
+  ->Args({100'000, 1'000'000})
+  ->Args({10'000'000, 10'000'000})
+  ->Args({10'000'000, 40'000'000})
+  ->Args({10'000'000, 100'000'000})
+  ->Args({100'000'000, 100'000'000})
+  ->Args({80'000'000, 240'000'000})
+  ->UseManualTime();
+
+BENCHMARK_REGISTER_F(Join, left_anti_join_64bit_nulls)
+  ->Unit(benchmark::kMillisecond)
+  ->Args({50'000'000, 50'000'000})
+  ->Args({40'000'000, 120'000'000})
+  ->UseManualTime();
+
+// left semi-join -------------------------------------------------------------
+BENCHMARK_REGISTER_F(Join, left_semi_join_32bit)
+  ->Unit(benchmark::kMillisecond)
+  ->Args({100'000, 100'000})
+  ->Args({100'000, 400'000})
+  ->Args({100'000, 1'000'000})
+  ->Args({10'000'000, 10'000'000})
+  ->Args({10'000'000, 40'000'000})
+  ->Args({10'000'000, 100'000'000})
+  ->Args({100'000'000, 100'000'000})
+  ->Args({80'000'000, 240'000'000})
+  ->UseManualTime();
+
+BENCHMARK_REGISTER_F(Join, left_semi_join_64bit)
+  ->Unit(benchmark::kMillisecond)
+  ->Args({50'000'000, 50'000'000})
+  ->Args({40'000'000, 120'000'000})
+  ->UseManualTime();
+
+BENCHMARK_REGISTER_F(Join, left_semi_join_32bit_nulls)
+  ->Unit(benchmark::kMillisecond)
+  ->Args({100'000, 100'000})
+  ->Args({100'000, 400'000})
+  ->Args({100'000, 1'000'000})
+  ->Args({10'000'000, 10'000'000})
+  ->Args({10'000'000, 40'000'000})
+  ->Args({10'000'000, 100'000'000})
+  ->Args({100'000'000, 100'000'000})
+  ->Args({80'000'000, 240'000'000})
+  ->UseManualTime();
+
+BENCHMARK_REGISTER_F(Join, left_semi_join_64bit_nulls)
+  ->Unit(benchmark::kMillisecond)
+  ->Args({50'000'000, 50'000'000})
+  ->Args({40'000'000, 120'000'000})
+  ->UseManualTime();

cpp/cmake/thrust.patch

@@ -81,25 +81,3 @@ index c0c6d59..937ee31 100644
      {
          typedef AgentScanPolicy<
                  128, 15,                                        ///< Threads per block, items per thread
-diff --git a/thrust/system/cuda/detail/scan_by_key.h b/thrust/system/cuda/detail/scan_by_key.h
-index fe4b321c..b3974c69 100644
---- a/thrust/system/cuda/detail/scan_by_key.h
-+++ b/thrust/system/cuda/detail/scan_by_key.h
-@@ -513,7 +513,7 @@ namespace __scan_by_key {
-             scan_op(scan_op_)
-       {
-         int  tile_idx      = blockIdx.x;
--        Size tile_base     = ITEMS_PER_TILE * tile_idx;
-+        Size tile_base     = ITEMS_PER_TILE * static_cast<Size>(tile_idx);
-         Size num_remaining = num_items - tile_base;
-
-         if (num_remaining > ITEMS_PER_TILE)
-@@ -734,7 +734,7 @@ namespace __scan_by_key {
-                              ScanOp                     scan_op,
-                              AddInitToScan              add_init_to_scan)
-   {
--    int          num_items    = static_cast<int>(thrust::distance(keys_first, keys_last));
-+    size_t       num_items    = static_cast<size_t>(thrust::distance(keys_first, keys_last));
-     size_t       storage_size = 0;
-     cudaStream_t stream       = cuda_cub::stream(policy);
-     bool         debug_sync   = THRUST_DEBUG_SYNC_FLAG;

cpp/include/cudf/detail/aggregation/aggregation.cuh

@@ -53,6 +53,14 @@ struct corresponding_operator<aggregation::MAX> {
   using type = DeviceMax;
 };
 template <>
+struct corresponding_operator<aggregation::ARGMIN> {
+  using type = DeviceMin;
+};
+template <>
+struct corresponding_operator<aggregation::ARGMAX> {
+  using type = DeviceMax;
+};
+template <>
 struct corresponding_operator<aggregation::ANY> {
   using type = DeviceMax;
 };
@@ -81,6 +89,10 @@ struct corresponding_operator<aggregation::VARIANCE> {
   using type = DeviceSum;
 };
 template <>
+struct corresponding_operator<aggregation::MEAN> {
+  using type = DeviceSum;
+};
+template <>
 struct corresponding_operator<aggregation::COUNT_VALID> {
   using type = DeviceCount;
 };
@@ -92,6 +104,12 @@ struct corresponding_operator<aggregation::COUNT_ALL> {
 template <aggregation::Kind k>
 using corresponding_operator_t = typename corresponding_operator<k>::type;
 
+template <aggregation::Kind k>
+constexpr bool has_corresponding_operator()
+{
+  return !std::is_same<typename corresponding_operator<k>::type, void>::value;
+}
+
 template <typename Source,
           aggregation::Kind k,
           bool target_has_nulls,

cpp/include/cudf/detail/groupby/group_replace_nulls.hpp

@@ -0,0 +1,47 @@
+/*
+ * 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.
+ */
+
+#pragma once
+
+#include <cudf/column/column_view.hpp>
+#include <cudf/replace.hpp>
+#include <cudf/types.hpp>
+#include <cudf/utilities/span.hpp>
+
+#include <rmm/exec_policy.hpp>
+namespace cudf {
+namespace groupby {
+namespace detail {
+
+/**
+ * @brief Internal API to replace nulls with preceding/following non-null values in @p value
+ *
+ * @param[in] grouped_value A column whose null values will be replaced.
+ * @param[in] group_labels Group labels for @p grouped_value, corresponding to group keys.
+ * @param[in] replace_policy Specify the position of replacement values relative to null values.
+ * @param stream CUDA stream used for device memory operations and kernel launches.
+ * @param[in] mr Device memory resource used to allocate device memory of the returned column.
+ */
+std::unique_ptr<column> group_replace_nulls(
+  cudf::column_view const& grouped_value,
+  device_span<size_type const> group_labels,
+  cudf::replace_policy replace_policy,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+}  // namespace detail
+}  // namespace groupby
+}  // namespace cudf