[REVIEW] Cross join feature.

CHANGELOG.md

@@ -4,6 +4,7 @@
 
 - PR #5287 Add `index.join` support
 - PR #5222 Adding clip feature support to DataFrame and Series
+- PR #5327 Add `cudf::cross_join` feature
 - PR #5204 Concatenate strings columns using row separator as strings column
 - PR #5342 Add support for `StringMethods.__getitem__`
 

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,
+                            cudaStream_t stream                 = 0,
+                            rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource());
+}  // namespace detail
+}  // namespace cudf

cpp/include/cudf/join.hpp

@@ -327,5 +327,32 @@ 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.
+ *
+ * @note Warning: This function can easily cause out-of-memory errors. The size of the output is
+ * equal to `left.num_rows() * right.num_rows()`. Use with caution.
+ *
+ * @code{.pseudo}
+ *          Left a: {0, 1, 2}
+ *          Right b: {3, 4, 5}
+ * Result: { a: {0, 0, 0, 1, 1, 1, 2, 2, 2}, b: {3, 4, 5, 3, 4, 5, 3, 4, 5} }
+ * @endcode
+
+ * @throw cudf::logic_error if number of columns in either `left` or `right` table is 0
+ *
+ * @param left  The left table
+ * @param right The right table
+ * @param mr    Device memory resource used to allocate the returned table's device memory
+ *
+ * @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,82 @@
+/*
+ * 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 {
+/**
+ * @copydoc cudf::cross_join
+ *
+ * @param stream CUDA stream used for device memory operations and kernel launches
+ */
+std::unique_ptr<cudf::table> cross_join(
+  cudf::table_view const& left,
+  cudf::table_view const& right,
+  cudaStream_t stream,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource())
+{
+  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(), stream, mr);
+
+  // 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, 0, mr);
+}
+
+}  // 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,
+                            cudaStream_t stream,
                             rmm::mr::device_memory_resource *mr)
 {
-  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, 0, mr);
 }
 
 }  // 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}")

cpp/tests/join/cross_join_tests.cpp

@@ -0,0 +1,144 @@
+/*
+ * 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/column/column_view.hpp>
+#include <cudf/join.hpp>
+#include <cudf/table/table.hpp>
+#include <cudf/table/table_view.hpp>
+
+#include <tests/utilities/base_fixture.hpp>
+#include <tests/utilities/column_utilities.hpp>
+#include <tests/utilities/column_wrapper.hpp>
+#include <tests/utilities/table_utilities.hpp>
+#include <tests/utilities/type_lists.hpp>
+
+template <typename T>
+using column_wrapper = cudf::test::fixed_width_column_wrapper<T>;
+
+template <typename T>
+class CrossJoinTypeTests : public cudf::test::BaseFixture {
+};
+
+TYPED_TEST_CASE(CrossJoinTypeTests, cudf::test::FixedWidthTypes);
+
+TYPED_TEST(CrossJoinTypeTests, CrossJoin)
+{
+  auto a_0 = column_wrapper<int32_t>{10, 20, 20, 50};
+  auto a_1 = column_wrapper<float>{5.0, .5, .5, .7};
+  auto a_2 = column_wrapper<TypeParam>{0, 0, 0, 0};
+  auto a_3 = cudf::test::strings_column_wrapper({"quick", "accénted", "turtlé", "composéd"});
+
+  auto b_0 = column_wrapper<int32_t>{10, 20, 20};
+  auto b_1 = column_wrapper<float>{5.0, .7, .7};
+  auto b_2 = column_wrapper<TypeParam>{0, 0, 0};
+  auto b_3 = cudf::test::strings_column_wrapper({"result", "", "words"});
+
+  auto expect_0 = column_wrapper<int32_t>{10, 10, 10, 20, 20, 20, 20, 20, 20, 50, 50, 50};
+  auto expect_1 = column_wrapper<float>{5.0, 5.0, 5.0, .5, .5, .5, .5, .5, .5, .7, .7, .7};
+  auto expect_2 = column_wrapper<TypeParam>{0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  auto expect_3 = cudf::test::strings_column_wrapper({"quick",
+                                                      "quick",
+                                                      "quick",
+                                                      "accénted",
+                                                      "accénted",
+                                                      "accénted",
+                                                      "turtlé",
+                                                      "turtlé",
+                                                      "turtlé",
+                                                      "composéd",
+                                                      "composéd",
+                                                      "composéd"});
+  auto expect_4 = column_wrapper<int32_t>{10, 20, 20, 10, 20, 20, 10, 20, 20, 10, 20, 20};
+  auto expect_5 = column_wrapper<float>{5.0, .7, .7, 5.0, .7, .7, 5.0, .7, .7, 5.0, .7, .7};
+  auto expect_6 = column_wrapper<TypeParam>{0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  auto expect_7 = cudf::test::strings_column_wrapper(
+    {"result", "", "words", "result", "", "words", "result", "", "words", "result", "", "words"});
+
+  auto table_a      = cudf::table_view{{a_0, a_1, a_2, a_3}};
+  auto table_b      = cudf::table_view{{b_0, b_1, b_2, b_3}};
+  auto table_expect = cudf::table_view{
+    {expect_0, expect_1, expect_2, expect_3, expect_4, expect_5, expect_6, expect_7}};
+
+  auto join_table = cudf::cross_join(table_a, table_b);
+
+  EXPECT_EQ(join_table->num_columns(), table_a.num_columns() + table_b.num_columns());
+  EXPECT_EQ(join_table->num_rows(), table_a.num_rows() * table_b.num_rows());
+  cudf::test::expect_tables_equal(join_table->view(), table_expect);
+}
+
+class CrossJoinInvalidInputs : public cudf::test::BaseFixture {
+};
+
+TEST_F(CrossJoinInvalidInputs, EmptyTable)
+{
+  auto b_0 = column_wrapper<int32_t>{10, 20, 20, 50};
+  auto b_1 = column_wrapper<float>{5.0, .7, .7, .7};
+  auto b_2 = column_wrapper<int8_t>{90, 75, 62, 41};
+  auto b_3 = cudf::test::strings_column_wrapper({"quick", "words", "result", ""});
+
+  auto column_a = std::vector<std::unique_ptr<cudf::column>>{};
+  auto table_a  = cudf::table(std::move(column_a));
+  auto table_b  = cudf::table_view{{b_0, b_1, b_2, b_3}};
+
+  //
+  //  table_a has no columns, table_b has columns
+  //  Let's check different permutations of passing table
+  //  with no columns to verify that exceptions are thrown
+  //
+  EXPECT_THROW(cudf::cross_join(table_a, table_b), cudf::logic_error);
+  EXPECT_THROW(cudf::cross_join(table_b, table_a), cudf::logic_error);
+}
+
+class CrossJoinEmptyResult : public cudf::test::BaseFixture {
+};
+
+TEST_F(CrossJoinEmptyResult, NoRows)
+{
+  auto a_0           = column_wrapper<int32_t>{};
+  auto a_1           = column_wrapper<float>{};
+  auto a_2           = column_wrapper<int8_t>{};
+  auto empty_strings = std::vector<std::string>();
+  auto a_3 = cudf::test::strings_column_wrapper(empty_strings.begin(), empty_strings.end());
+
+  auto b_0 = column_wrapper<int32_t>{10, 20, 20, 50};
+  auto b_1 = column_wrapper<float>{5.0, .7, .7, .7};
+  auto b_2 = column_wrapper<int8_t>{90, 75, 62, 41};
+  auto b_3 = cudf::test::strings_column_wrapper({"quick", "words", "result", ""});
+
+  auto expect_0 = column_wrapper<int32_t>{};
+  auto expect_1 = column_wrapper<float>{};
+  auto expect_2 = column_wrapper<int8_t>{};
+  auto expect_3 = cudf::test::strings_column_wrapper(empty_strings.begin(), empty_strings.end());
+  auto expect_4 = column_wrapper<int32_t>{};
+  auto expect_5 = column_wrapper<float>{};
+  auto expect_6 = column_wrapper<int8_t>{};
+  auto expect_7 = cudf::test::strings_column_wrapper(empty_strings.begin(), empty_strings.end());
+
+  auto table_a      = cudf::table_view{{a_0, a_1, a_2, a_3}};
+  auto table_b      = cudf::table_view{{b_0, b_1, b_2, b_3}};
+  auto table_expect = cudf::table_view{
+    {expect_0, expect_1, expect_2, expect_3, expect_4, expect_5, expect_6, expect_7}};
+
+  auto join_table         = cudf::cross_join(table_a, table_b);
+  auto join_table_reverse = cudf::cross_join(table_b, table_a);
+
+  EXPECT_EQ(join_table->num_columns(), table_a.num_columns() + table_b.num_columns());
+  EXPECT_EQ(join_table->num_rows(), 0);
+  cudf::test::expect_tables_equal(join_table->view(), table_expect);
+  EXPECT_EQ(join_table_reverse->num_columns(), table_a.num_columns() + table_b.num_columns());
+  EXPECT_EQ(join_table_reverse->num_rows(), 0);
+}