[REVIEW] Port scatter to libcudf++

CHANGELOG.md

@@ -77,7 +77,7 @@
 - PR #2971 Added initial gather and scatter methods for strings_column_view
 - PR #3133 Port NVStrings to cudf column: count_characters and count_bytes
 - PR #2991 Added strings column functions concatenate and join_strings
-- PR #3028 Port gather and scatter to libcudf++
+- PR #3028 Port gather to libcudf++
 - PR #3135 Add nvtx utilities to cudf::nvtx namespace
 - PR #3021 Java host side concat of serialized buffers
 - PR #3138 Move unary files to legacy
@@ -122,6 +122,7 @@
 - PR #3294 Update to arrow-cpp and pyarrow 0.15.1
 - PR #3310 Add `row_hasher` and `element_hasher` utilities
 - PR #3286 Clean up the starter code on README
+- PR #3354 Port gather and scatter to libcudf++
 - PR #3322 Port NVStrings pad operations to cudf strings column
 - PR #3345 Add cache member for number of characters in string_view class
 - PR #3299 Define and implement new `is_sorted` APIs

conda/recipes/libcudf/meta.yaml

@@ -46,6 +46,7 @@ test:
   commands:
     - test -f $PREFIX/lib/libcudf.so
     - test -f $PREFIX/lib/libcudftestutil.a
+    - test -f $PREFIX/include/cudf/detail/scatter.hpp
     - test -f $PREFIX/include/cudf/legacy/bitmask.hpp
     - test -f $PREFIX/include/cudf/legacy/column.hpp
     - test -f $PREFIX/include/cudf/legacy/reshape.hpp

cpp/CMakeLists.txt

@@ -459,11 +459,12 @@ add_library(cudf
             src/io/utilities/legacy/parsing_utils.cu
             src/utilities/legacy/cuda_utils.cu
             src/utilities/column_utils.cpp
-	    src/copying/gather.cu
             src/utilities/legacy/error_utils.cpp
             src/utilities/nvtx/nvtx_utils.cpp
             src/utilities/nvtx/legacy/nvtx_utils.cpp
             src/copying/copy.cpp
+            src/copying/gather.cu
+            src/copying/scatter.cu
             src/copying/slice.cpp
             src/copying/split.cpp
             src/copying/legacy/copy.cpp

cpp/include/cudf/copying.hpp

@@ -20,6 +20,7 @@
 #include "types.hpp"
 #include <cudf/cudf.h>
 #include <cudf/types.hpp>
+#include <cudf/scalar/scalar.hpp>
 #include <cudf/column/column_view.hpp>
 #include <cudf/table/table.hpp>
 
@@ -48,14 +49,89 @@ namespace experimental {
  * rows in the source columns to rows in the destination columns.
  * @param[in] check_bounds Optionally perform bounds checking on the values
  * of `gather_map` and throw an error if any of its values are out of bounds.
- * @params[in] mr The resource to use for all allocations
- * @return cudf::table Result of the gather
+ * @param[in] mr The resource to use for all allocations
+ * @return std::unique_ptr<table> Result of the gather
  */
 std::unique_ptr<table> gather(table_view const& source_table, column_view const& gather_map,
-			      bool check_bounds = false,
-			      rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource());
+                              bool check_bounds = false,
+                              rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource());
+
+/**
+ * @brief Scatters the rows of the source table into a copy of the target table
+ * according to a scatter map.
+ *
+ * Scatters values from the source table into the target table out-of-place,
+ * returning a "destination table". The scatter is performed according to a
+ * scatter map such that row `scatter_map[i]` of the destination table gets row
+ * `i` of the source table. All other rows of the destination table equal
+ * corresponding rows of the target table.
+ *
+ * The number of columns in source must match the number of columns in target
+ * and their corresponding datatypes must be the same.
+ * 
+ * If the same index appears more than once in the scatter map, the result is
+ * undefined.
+ *
+ * A negative value `i` in the `scatter_map` is interpreted as `i+n`, where `n`
+ * is the number of rows in the `target` table.
+ *
+ * @throws `cudf::logic_error` if `check_bounds == true` and an index exists in
+ * `scatter_map` outside the range `[-n, n)`, where `n` is the number of rows in
+ * the target table. If `check_bounds == false`, the behavior is undefined.
+ *
+ * @param source The input columns containing values to be scattered into the
+ * target columns
+ * @param scatter_map A non-nullable column of integral indices that maps the
+ * rows in the source table to rows in the target table. The size must be equal
+ * to or less than the number of elements in the source columns.
+ * @param target The set of columns into which values from the source_table
+ * are to be scattered
+ * @param check_bounds Optionally perform bounds checking on the values of
+ * `scatter_map` and throw an error if any of its values are out of bounds.
+ * @param mr The resource to use for all allocations
+ * @return Result of scattering values from source to target
+ *---------------------------------------------------------------------------**/
+std::unique_ptr<table> scatter(
+    table_view const& source, column_view const& scatter_map,
+    table_view const& target, bool check_bounds = false,
+    rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource());
+
+/**
+ * @brief Scatters a row of scalar values into a copy of the target table
+ * according to a scatter map.
+ *
+ * Scatters values from the source row into the target table out-of-place,
+ * returning a "destination table". The scatter is performed according to a
+ * scatter map such that row `scatter_map[i]` of the destination table is
+ * replaced by the source row. All other rows of the destination table equal
+ * corresponding rows of the target table.
+ *
+ * The number of elements in source must match the number of columns in target
+ * and their corresponding datatypes must be the same.
+ * 
+ * If the same index appears more than once in the scatter map, the result is
+ * undefined.
+ *
+ * @throws `cudf::logic_error` if `check_bounds == true` and an index exists in
+ * `scatter_map` outside the range `[-n, n)`, where `n` is the number of rows in
+ * the target table. If `check_bounds == false`, the behavior is undefined.
+ *
+ * @param source The input scalars containing values to be scattered into the
+ * target columns
+ * @param indices A non-nullable column of integral indices that indicate
+ * the rows in the target table to be replaced by source.
+ * @param target The set of columns into which values from the source_table
+ * are to be scattered
+ * @param check_bounds Optionally perform bounds checking on the values of
+ * `scatter_map` and throw an error if any of its values are out of bounds.
+ * @param mr The resource to use for all allocations
+ * @return Result of scattering values from source to target
+ *---------------------------------------------------------------------------**/
+std::unique_ptr<table> scatter(
+    std::vector<std::unique_ptr<scalar>> const& source, column_view const& indices,
+    table_view const& target, bool check_bounds = false,
+    rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource());
 
-
 /** ---------------------------------------------------------------------------*
 * @brief Indicates when to allocate a mask, based on an existing mask.
 * ---------------------------------------------------------------------------**/
@@ -65,7 +141,6 @@ enum class  mask_allocation_policy {
     ALWAYS ///< Allocate a null mask, regardless of input
 };
 
-
 /*
  * Initializes and returns an empty column of the same type as the `input`.
  *

cpp/include/cudf/detail/gather.cuh

@@ -55,12 +55,10 @@ __global__ void gather_bitmask_kernel(table_device_view source_table,
 
   for (size_type i = 0; i < source_table.num_columns(); i++) {
 
-    constexpr int warp_size = 32;
-
     column_device_view source_col = source_table.column(i);
     mutable_column_device_view destination_col = destination_table.column(i);
 
-    if (source_col.has_nulls()) {
+    if (destination_col.nullable()) {
       size_type destination_row_base = blockIdx.x * blockDim.x;
       cudf::size_type valid_count_accumulate = 0;
 
@@ -71,14 +69,17 @@ __global__ void gather_bitmask_kernel(table_device_view source_table,
         size_type source_row =
           thread_active ? gather_map[destination_row] : 0;
 
-        bool source_bit_is_valid = source_col.has_nulls()
-          ? source_col.is_valid_nocheck(source_row)
-          : true;
+        bool bit_is_valid;
+        if (ignore_out_of_bounds && (source_row < 0 || source_row >= source_col.size())) {
+          bit_is_valid = thread_active && destination_col.is_valid_nocheck(destination_row);
+        } else {
+          bit_is_valid = source_col.is_valid(source_row);
+        }
 
         // Use ballot to find all valid bits in this warp and create the output
         // bitmask element
         const uint32_t valid_warp =
-          __ballot_sync(0xffffffff, thread_active && source_bit_is_valid);
+          __ballot_sync(0xffffffff, thread_active && bit_is_valid);
 
         const size_type valid_index = word_index(destination_row);
 

cpp/include/cudf/detail/scatter.hpp

@@ -0,0 +1,110 @@
+/*
+ * Copyright (c) 2019, 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/table/table_view.hpp>
+#include <cudf/table/table.hpp>
+
+#include <memory>
+
+namespace cudf {
+namespace experimental {
+namespace detail {
+
+/**
+ * @brief Scatters the rows of the source table into a copy of the target table
+ * according to a scatter map.
+ *
+ * Scatters values from the source table into the target table out-of-place,
+ * returning a "destination table". The scatter is performed according to a
+ * scatter map such that row `scatter_map[i]` of the destination table gets row
+ * `i` of the source table. All other rows of the destination table equal
+ * corresponding rows of the target table.
+ *
+ * The number of columns in source must match the number of columns in target
+ * and their corresponding datatypes must be the same.
+ *
+ * A negative value `i` in the `scatter_map` is interpreted as `i+n`, where `n`
+ * is the number of rows in the `target` table.
+ * 
+ * If the same index appears more than once in the scatter map, the result is
+ * undefined.
+ *
+ * @throws `cudf::logic_error` if `check_bounds == true` and an index exists in
+ * `scatter_map` outside the range `[-n, n)`, where `n` is the number of rows in
+ * the target table. If `check_bounds == false`, the behavior is undefined.
+ *
+ * @param source The input columns containing values to be scattered into the
+ * target columns
+ * @param scatter_map A non-nullable column of integral indices that maps the
+ * rows in the source table to rows in the target table. The size must be equal
+ * to or less than the number of elements in the source columns.
+ * @param target The set of columns into which values from the source_table
+ * are to be scattered
+ * @param check_bounds Optionally perform bounds checking on the values of
+ * `scatter_map` and throw an error if any of its values are out of bounds.
+ * @param mr The resource to use for all allocations
+ * @param stream The stream to use for CUDA operations
+ * @return Result of scattering values from source to target
+ *---------------------------------------------------------------------------**/
+std::unique_ptr<table> scatter(
+    table_view const& source, column_view const& scatter_map,
+    table_view const& target, bool check_bounds = false,
+    rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource(),
+    cudaStream_t stream = 0);
+
+/**
+ * @brief Scatters a row of scalar values into a copy of the target table
+ * according to a scatter map.
+ *
+ * Scatters values from the source row into the target table out-of-place,
+ * returning a "destination table". The scatter is performed according to a
+ * scatter map such that row `scatter_map[i]` of the destination table is
+ * replaced by the source row. All other rows of the destination table equal
+ * corresponding rows of the target table.
+ *
+ * The number of elements in source must match the number of columns in target
+ * and their corresponding datatypes must be the same.
+ * 
+ * If the same index appears more than once in the scatter map, the result is
+ * undefined.
+ *
+ * @throws `cudf::logic_error` if `check_bounds == true` and an index exists in
+ * `scatter_map` outside the range `[-n, n)`, where `n` is the number of rows in
+ * the target table. If `check_bounds == false`, the behavior is undefined.
+ *
+ * @param source The input scalars containing values to be scattered into the
+ * target columns
+ * @param indices A non-nullable column of integral indices that indicate
+ * the rows in the target table to be replaced by source.
+ * @param target The set of columns into which values from the source_table
+ * are to be scattered
+ * @param check_bounds Optionally perform bounds checking on the values of
+ * `scatter_map` and throw an error if any of its values are out of bounds.
+ * @param mr The resource to use for all allocations
+ * @param stream The stream to use for CUDA operations
+ * @return Result of scattering values from source to target
+ *---------------------------------------------------------------------------**/
+std::unique_ptr<table> scatter(
+    std::vector<std::unique_ptr<scalar>> const& source, column_view const& indices,
+    table_view const& target, bool check_bounds = false,
+    rmm::mr::device_memory_resource* mr = rmm::mr::get_default_resource(),
+    cudaStream_t stream = 0);
+
+}  // namespace detail
+}  // namespace experimental
+}  // namespace cudf