[REVIEW] Cross join feature.

CHANGELOG.md

@@ -3,6 +3,7 @@
 ## New Features
 
 - PR #5222 Adding clip feature support to DataFrame and Series
+- PR #5327 Add `cudf::cross_join` feature
 
 ## Improvements
 - PR #5245 Add column reduction benchmark

conda/recipes/libcudf/meta.yaml

@@ -69,6 +69,7 @@ test:
     - test -f $PREFIX/include/cudf/detail/reduction_functions.hpp
     - test -f $PREFIX/include/cudf/detail/repeat.hpp
     - test -f $PREFIX/include/cudf/detail/replace.hpp
+    - test -f $PREFIX/include/cudf/detail/reshape.hpp
     - test -f $PREFIX/include/cudf/detail/scatter.hpp
     - test -f $PREFIX/include/cudf/detail/search.hpp
     - test -f $PREFIX/include/cudf/detail/sequence.hpp

cpp/CMakeLists.txt

@@ -415,6 +415,7 @@ add_library(cudf
             src/merge/merge.cu
             src/partitioning/round_robin.cu
             src/join/join.cu
+            src/join/cross_join.cu
             src/join/semi_join.cu
             src/sort/is_sorted.cu
             src/binaryop/binaryop.cpp

cpp/include/cudf/detail/reshape.hpp

@@ -0,0 +1,35 @@
+/*
+ * Copyright (c) 2020, 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/types.hpp>
+
+#include <memory>
+
+namespace cudf {
+namespace detail {
+/**
+ * @copydoc cudf::tile
+ *
+ * @param stream CUDA stream used for device memory operations and kernel launches.
+ */
+std::unique_ptr<table> tile(table_view const& input,
+                            size_type count,
+                            rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource(),
+                            cudaStream_t stream                 = 0);
+}  // namespace detail
+}  // namespace cudf

cpp/include/cudf/join.hpp

@@ -327,5 +327,26 @@ std::unique_ptr<cudf::table> left_anti_join(
   std::vector<cudf::size_type> const& return_columns,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource());
 
+/**
+ * @brief  Performs a cross join on two tables (left, right)
+ *
+ * The cross join returns the cartesian product of rows from each table.
+ *
+ * The approach is to repeat the left table by the number of rows in the right table
+ * and tile the right table by the number of rows in the left table.
+ *
+ * @throws cudf::logic_error if number of columns in either `left` or `right` table is 0
+ *
+ * @param[in] left                  The left table
+ * @param[in] right                 The right table
+ * @param[in] mr                    Device memory resource to use for device memory allocation
+ *
+ * @returns                         Result of cross joining `left` and `right` tables
+ */
+std::unique_ptr<cudf::table> cross_join(
+  cudf::table_view const& left,
+  cudf::table_view const& right,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource());
+
 /** @} */  // end of group
 }  // namespace cudf

cpp/src/join/cross_join.cu

@@ -0,0 +1,94 @@
+/*
+ * Copyright (c) 2020, 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/column/column.hpp>
+#include <cudf/copying.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/detail/repeat.hpp>
+#include <cudf/detail/reshape.hpp>
+#include <cudf/filling.hpp>
+#include <cudf/join.hpp>
+#include <cudf/reshape.hpp>
+#include <cudf/table/table.hpp>
+#include <cudf/table/table_view.hpp>
+#include <cudf/types.hpp>
+#include <cudf/utilities/error.hpp>
+
+namespace cudf {
+namespace detail {
+/**
+ * @brief  Performs a cross join on two tables (left, right)
+ *
+ * The cross join returns the cartesian product of rows from each table.
+ *
+ * The approach is to repeat the left table by the number of rows in the right table
+ * and tile the right table by the number of rows in the left table.
+ *
+ * @throws cudf::logic_error if number of columns in either `left` or `right` table is 0
+ *
+ * @param[in] left                  The left table
+ * @param[in] right                 The right table
+ * @param[in] mr                    Device memory resource to use for device memory allocation
+ * @param[in] stream                Cuda stream
+ *
+ * @returns                         Result of cross joining `left` and `right` tables
+ */
+std::unique_ptr<cudf::table> cross_join(
+  cudf::table_view const& left,
+  cudf::table_view const& right,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource(),
+  cudaStream_t stream                 = 0)
+{
+  CUDF_EXPECTS(0 != left.num_columns(), "Left table is empty");
+  CUDF_EXPECTS(0 != right.num_columns(), "Right table is empty");
+
+  // If left or right table has no rows, return an empty table with all columns
+  if ((0 == left.num_rows()) || (0 == right.num_rows())) {
+    auto empty_left_columns  = empty_like(left)->release();
+    auto empty_right_columns = empty_like(right)->release();
+    std::move(empty_right_columns.begin(),
+              empty_right_columns.end(),
+              std::back_inserter(empty_left_columns));
+    return std::make_unique<table>(std::move(empty_left_columns));
+  }
+
+  // Repeat left table
+  numeric_scalar<size_type> num_repeats = right.num_rows();
+  auto left_repeated                    = detail::repeat(left, num_repeats, mr, stream);
+
+  // Tile right table
+  auto right_tiled = detail::tile(right, left.num_rows(), mr, stream);
+
+  // Concatenate all repeated/tiled columns into one table
+  auto left_repeated_columns = left_repeated->release();
+  auto right_tiled_columns   = right_tiled->release();
+  std::move(right_tiled_columns.begin(),
+            right_tiled_columns.end(),
+            std::back_inserter(left_repeated_columns));
+
+  return std::make_unique<table>(std::move(left_repeated_columns));
+}
+}  // namespace detail
+
+std::unique_ptr<cudf::table> cross_join(cudf::table_view const& left,
+                                        cudf::table_view const& right,
+                                        rmm::mr::device_memory_resource* mr)
+{
+  CUDF_FUNC_RANGE();
+  return detail::cross_join(left, right, mr, 0);
+}
+
+}  // namespace cudf

cpp/src/join/semi_join.cu

@@ -1,3 +1,19 @@
+/*
+ * Copyright (c) 2020, 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/column/column_factories.hpp>
 #include <cudf/copying.hpp>
 #include <cudf/detail/gather.hpp>

cpp/src/reshape/tile.cu

@@ -16,6 +16,7 @@
 
 #include <cudf/copying.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/detail/reshape.hpp>
 #include <cudf/table/table.hpp>
 #include <cudf/types.hpp>
 #include <cudf/utilities/error.hpp>
@@ -35,11 +36,12 @@ struct tile_functor {
 
 }  // anonymous namespace
 
+namespace detail {
 std::unique_ptr<table> tile(const table_view &in,
                             size_type count,
-                            rmm::mr::device_memory_resource *mr)
+                            rmm::mr::device_memory_resource *mr,
+                            cudaStream_t stream)
 {
-  CUDF_FUNC_RANGE();
   CUDF_EXPECTS(count >= 0, "Count cannot be negative");
 
   auto in_num_rows = in.num_rows();
@@ -50,7 +52,16 @@ std::unique_ptr<table> tile(const table_view &in,
   auto counting_it  = thrust::make_counting_iterator<size_type>(0);
   auto tiled_it     = thrust::make_transform_iterator(counting_it, tile_functor{in_num_rows});
 
-  return detail::gather(in, tiled_it, tiled_it + out_num_rows, false, mr);
+  return detail::gather(in, tiled_it, tiled_it + out_num_rows, false, mr, stream);
+}
+}  // namespace detail
+
+std::unique_ptr<table> tile(const table_view &in,
+                            size_type count,
+                            rmm::mr::device_memory_resource *mr)
+{
+  CUDF_FUNC_RANGE();
+  return detail::tile(in, count, mr, 0);
 }
 
 }  // namespace cudf

cpp/tests/CMakeLists.txt

@@ -158,6 +158,7 @@ ConfigureTest(GROUPBY_TEST "${GROUPBY_TEST_SRC}")
 
 set(JOIN_TEST_SRC
     "${CMAKE_CURRENT_SOURCE_DIR}/join/join_tests.cu"
+    "${CMAKE_CURRENT_SOURCE_DIR}/join/cross_join_tests.cpp"
     "${CMAKE_CURRENT_SOURCE_DIR}/join/semi_join_tests.cpp")
 
 ConfigureTest(JOIN_TEST "${JOIN_TEST_SRC}")