[REVIEW] find_and_replace function

cpp/CMakeLists.txt

@@ -146,6 +146,7 @@ add_library(cudf SHARED
             src/hash/hash_ops.cu
             src/quantiles/quantiles.cu
             src/reductions/reductions.cu
+            src/replace/replace.cu
             src/reductions/scan.cu
             src/unary/unary_ops.cu
             # src/windowed/windowed_ops.cu ... this is broken

cpp/include/cudf/functions.h

@@ -852,3 +852,18 @@ gdf_error gdf_quantile_aprrox(	gdf_column*  col_in,       //input column with 0
                                 double       q,            //requested quantile in [0,1]
                                 void*        t_erased_res, //type-erased result of same type as column;
                                 gdf_context* ctxt);        //context info
+
+/* --------------------------------------------------------------------------*/
+/** 
+ * @brief Replace elements from col according to the mapping old_values to new_values
+ * 
+ * @Param[in] col gdf_column with the data to be modified
+ * @Param[in] old_values gdf_column with the old values to be replaced
+ * @Param[in] new_values gdf_column with the new values
+ * 
+ * @Returns GDF_SUCCESS upon successful completion
+ */
+/* ----------------------------------------------------------------------------*/
+gdf_error gdf_find_and_replace_all(gdf_column*       col,
+                                   const gdf_column* old_values,
+                                   const gdf_column* new_values);

cpp/src/dataframe/type_dispatcher.hpp

@@ -83,6 +83,9 @@
 /* ----------------------------------------------------------------------------*/
 namespace cudf{
 
+// This pragma disables a compiler warning that complains about the valid usage	
+// of calling a __host__ functor from this function which is __host__ __device__
+#pragma hd_warning_disable
 template < class functor_t, 
            typename... Ts>
 CUDA_HOST_DEVICE_CALLABLE
@@ -104,12 +107,8 @@ decltype(auto) type_dispatcher(gdf_dtype dtype,
     case GDF_DATE64:    { return f.template operator()< date64 >(std::forward<Ts>(args)...); }
     case GDF_TIMESTAMP: { return f.template operator()< timestamp >(std::forward<Ts>(args)...); }
     case GDF_CATEGORY:  { return f.template operator()< category >(std::forward<Ts>(args)...); }
-    //case GDF_STRING:    { return f.template operator()< enum_map<GDF_STRING> >(std::forward<Ts>(args)...); }
+    default: { assert(false && "type_dispatcher: invalid gdf_dtype"); }// this will only fire on a Debug build
   }
-
-  // This will only fire with a DEBUG build
-  assert(0 && "type_dispatcher: invalid gdf_dtype");
-
   // Need to find out what the return type is in order to have a default return value
   // and solve the compiler warning for lack of a default return
   using return_type = decltype(f.template operator()<int8_t>(std::forward<Ts>(args)...));

cpp/src/replace/replace.cu

@@ -0,0 +1,126 @@
+/*
+ * Copyright 2018 BlazingDB, Inc.
+ *     Copyright 2018 Cristhian Alberto Gonzales Castillo <[email protected]>
+ *
+ * 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 <thrust/device_ptr.h>
+#include <thrust/find.h>
+#include <thrust/execution_policy.h>
+
+#include "cudf.h"
+#include "utilities/error_utils.h"
+#include "dataframe/type_dispatcher.hpp"
+
+namespace{ //anonymous
+
+  constexpr int BLOCK_SIZE = 256;
+
+  template <class T>
+  struct elements_are_equal{
+    __device__
+    elements_are_equal(T el) : element{el}
+    {
+    }
+
+    __device__
+    bool operator()(T other)
+    { 
+      return (cudf::detail::unwrap(other) == cudf::detail::unwrap(element));
+    }
+
+  private:
+    T element;
+  };
+
+  template <class T>
+  __global__
+  void replace_kernel(T*                          d_col_data,
+                      size_t                      nrows,
+                      thrust::device_ptr<const T> old_values_begin,
+                      thrust::device_ptr<const T> old_values_end,
+                      const T*                    d_new_values)
+  {
+    size_t i = blockIdx.x * blockDim.x + threadIdx.x;
+    while(i < nrows)
+    {
+      auto found_ptr = thrust::find_if(thrust::seq, old_values_begin, old_values_end, elements_are_equal<T>(d_col_data[i]));
+
+      if (found_ptr != old_values_end) {
+          auto d = thrust::distance(old_values_begin, found_ptr);
+          d_col_data[i] = d_new_values[d];
+      }
+
+      i += blockDim.x * gridDim.x;
+    }
+  }
+
+  struct replace_kernel_forwarder {
+    template <typename col_type>
+    void operator()(void*       d_col_data,
+                    size_t      nrows,
+                    const void* d_old_values,
+                    const void* d_new_values,
+                    size_t      nvalues)
+    {
+      thrust::device_ptr<const col_type> old_values_begin(static_cast<const col_type*>(d_old_values));
+      thrust::device_ptr<const col_type> old_values_end(static_cast<const col_type*>(d_old_values) + nvalues);
+
+      const size_t grid_size = nrows / BLOCK_SIZE + (nrows % BLOCK_SIZE != 0);
+      replace_kernel<<<grid_size, BLOCK_SIZE>>>(static_cast<col_type*>(d_col_data),
+                                             nrows,
+                                             old_values_begin,
+                                             old_values_end,
+                                             static_cast<const col_type*>(d_new_values));
+    }
+  };
+
+  gdf_error find_and_replace_all(gdf_column*       col,
+                                 const gdf_column* old_values,
+                                 const gdf_column* new_values)
+  {
+    GDF_REQUIRE(col != nullptr && old_values != nullptr && new_values != nullptr, GDF_DATASET_EMPTY);
+    GDF_REQUIRE(old_values->size == new_values->size, GDF_COLUMN_SIZE_MISMATCH);
+    GDF_REQUIRE(col->dtype == old_values->dtype && col->dtype == new_values->dtype, GDF_DTYPE_MISMATCH);
+    GDF_REQUIRE(col->valid == nullptr && col->null_count == 0, GDF_VALIDITY_UNSUPPORTED);
+
+    cudf::type_dispatcher(col->dtype, replace_kernel_forwarder{},
+                          col->data,
+                          col->size,
+                          old_values->data,
+                          new_values->data,
+                          new_values->size); 
+
+    return GDF_SUCCESS;
+  }
+
+} //end anonymous namespace
+
+/* --------------------------------------------------------------------------*/
+/** 
+ * @brief Replace elements from col according to the mapping old_values to new_values
+ * 
+ * @Param[in] col gdf_column with the data to be modified
+ * @Param[in] old_values gdf_column with the old values to be replaced
+ * @Param[in] new_values gdf_column with the new values
+ * 
+ * @Returns GDF_SUCCESS upon successful completion
+ */
+/* ----------------------------------------------------------------------------*/
+gdf_error gdf_find_and_replace_all(gdf_column*       col,
+                                   const gdf_column* old_values,
+                                   const gdf_column* new_values)
+{
+  return find_and_replace_all(col, old_values, new_values);
+}

cpp/tests/CMakeLists.txt

@@ -135,6 +135,14 @@ set(QUANTILES_TEST_SRC
 
 ConfigureTest(QUANTILES_TEST "${QUANTILES_TEST_SRC}")
 
+###################################################################################################
+# - replace tests ---------------------------------------------------------------------------------
+
+set(REPLACE_TEST_SRC 
+    "${CMAKE_CURRENT_SOURCE_DIR}/replace/replace_tests.cu")
+
+ConfigureTest(REPLACE_TEST "${REPLACE_TEST_SRC}")
+
 ###################################################################################################
 # - unary tests -----------------------------------------------------------------------------------