[Review] replace.hpp strings functionality #3469

jwyles · 2019-11-26T23:12:35Z

This depends on PR 3187.
This is a work in progress PR and is not yet complete.
This is adding string support to the functions find_and_replace_all(), and replace_nulls() defined in replace.hpp.

…ll()

davidwendt · 2019-11-27T14:46:28Z

Also, the strings replace_nulls is in PR #3332

cpp/src/replace/replace.cu

cpp/tests/replace/replace_tests.cu

codecov · 2019-12-02T22:46:26Z

Codecov Report

❗ No coverage uploaded for pull request base (branch-0.11@53478c1). Click here to learn what that means.
The diff coverage is n/a.

@@              Coverage Diff               @@
##             branch-0.11    #3469   +/-   ##
==============================================
  Coverage               ?   87.38%           
==============================================
  Files                  ?       49           
  Lines                  ?     9376           
  Branches               ?        0           
==============================================
  Hits                   ?     8193           
  Misses                 ?     1183           
  Partials               ?        0

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 53478c1...cd3efd4. Read the comment docs.

davidwendt · 2019-12-03T13:39:51Z

cpp/src/replace/replace.cu

+                                           cudf::column_view const& values_to_replace,
+                                           cudf::column_view const& replacement_values,
+                                           rmm::mr::device_memory_resource* mr,
+                                           cudaStream_t stream = 0) {


I know you are still working on this but I think it would be good to refactor this into a strings::detail function and source/header file like we do for gather, scatter, copy, fill.

I am not sure that this would get much re-use if factored out that way, am I missing a use case for this?

Just trying to keep all the code that builds strings column in some central location in case we need to change the children types/layout. I suppose I can split the code out in a later PR.

cpp/src/replace/replace.cu

… added

davidwendt · 2019-12-04T02:31:14Z

cpp/src/replace/replace.cu

+    CUDA_TRY(cudaMemcpy(&size, offsets_data + offsets_view.size() - 1, sizeof(int32_t), cudaMemcpyDefault));
+
+    // Allocate chars array and output null mask
+    std::unique_ptr<cudf::column> output_chars = cudf::make_numeric_column(cudf::data_type(cudf::type_id::INT8), size);


This is not the correct way to create the chars column. There are some missing edge cases. Please use the create_chars_child_column method in cpp/include/cudf/strings/detail/utilities.hpp if possible.

harrism

Need to be more careful and consistent with stream semantics. The kernels are also unnecessarily complex and branchy.

cpp/src/replace/replace.cu

harrism · 2019-12-04T01:39:21Z

cpp/src/replace/replace.cu

+      output_is_valid = replacement.is_valid_nocheck(i);
+    }
+
+    if (phase == 0) {


Simplify lines 581-617 to something like this, it's a lot shorter and clearer.

cudf::string_view out; if (input_is_valid) { out = input.element<cudf::string_view>(i); } else if (output_is_valid) { out = replacement.element<cudf::string_view>(i); } bool nonzero_output = (input_is_valid || output_is_valid); if (phase == 0) { offsets[i] = nonzero_output ? out.size_bytes() : 0; } else if (phase == 1) { if (nonzero_output) std::memcpy(chars + offsets[i], out.data(), out.size_bytes()) uint32_t bitmask = __ballot_sync(active_mask, output_is_valid); if (0 == lane_id) { output_valid[cudf::word_index(i)] = bitmask; valid_sum += __popc(bitmask); } }

Note I think you can also replace that out.size_bytes() in phase 1 with offsets[i+1].

harrism · 2019-12-04T01:44:13Z

cpp/src/replace/replace.cu

+                                                              offsets_view.data<cudf::size_type>(),
+                                                              nullptr,
+                                                              valid_count);
+    cudaStreamSynchronize(stream);


No need to synchronize here.

Suggested change

cudaStreamSynchronize(stream);

harrism · 2019-12-04T01:44:43Z

cpp/src/replace/replace.cu

+                           offsets_data + offsets_view.size(),
+                           offsets_data);
+    int32_t size;
+    CUDA_TRY(cudaMemcpy(&size, offsets_data + offsets_view.size() - 1, sizeof(int32_t), cudaMemcpyDefault));


Suggested change

CUDA_TRY(cudaMemcpy(&size, offsets_data + offsets_view.size() - 1, sizeof(int32_t), cudaMemcpyDefault));

CUDA_TRY(cudaMemcpyAsync(&size, offsets_data + offsets_view.size() - 1, sizeof(int32_t), cudaMemcpyDefault, stream));

harrism · 2019-12-04T01:45:08Z

cpp/src/replace/replace.cu

+                                                               offsets_view.data<cudf::size_type>(),
+                                                               output_chars_view.data<char>(),
+                                                               valid_count);
+    cudaStreamSynchronize(stream);


Suggested change

cudaStreamSynchronize(stream);

harrism · 2019-12-04T02:36:23Z

cpp/src/replace/replace.cu

+      bool const input_is_valid = input.is_valid_nocheck(i);
+      output_is_valid = input_is_valid;
+      bitmask = __ballot_sync(active_mask, input_is_valid);
+      if (input_is_valid) {


There's a ton of duplication in this complex chain of nested if/else/elseifs. Factor out the calculation of output and memcpy to reduce duplication. See my comments on the other kernel for similar guidance.

cpp/src/replace/replace.cu

harrism · 2019-12-04T02:42:26Z

cpp/src/replace/replace.cu

+                                     cudf::column_device_view& input,
+                                     cudf::column_device_view& values_to_replace,
+                                     cudf::column_device_view& replacement_values) {
+  cudf::string_view inputString = input.element<cudf::string_view>(idx);


Suggested change

cudf::string_view inputString = input.element<cudf::string_view>(idx);

cudf::string_view input_string = input.element<cudf::string_view>(idx);

harrism · 2019-12-04T02:42:36Z

cpp/src/replace/replace.cu

+  cudf::string_view inputString = input.element<cudf::string_view>(idx);
+  int match = -1;
+  for (int i = 0; i < values_to_replace.size(); i++) {
+    cudf::string_view valueString = values_to_replace.element<cudf::string_view>(i);


Suggested change

cudf::string_view valueString = values_to_replace.element<cudf::string_view>(i);

cudf::string_view value_string = values_to_replace.element<cudf::string_view>(i);

davidwendt · 2019-12-04T21:32:42Z

cpp/src/replace/replace.cu

+
+    // Allocate chars array and output null mask
+    std::unique_ptr<cudf::column> output_chars =
+        cudf::strings::detail::create_chars_child_column(input_col.size(), 0, size, mr, stream);


The null-count parameter here must be accurate for the output strings column and not the chars column in order for this to work properly. This may mean you need to compute the nulls first.

harrism

Nearly there.

harrism · 2019-12-04T21:34:06Z

cpp/src/replace/replace.cu

+    bool input_is_valid = true;
+    if (input_has_nulls)
+      input_is_valid = input.is_valid_nocheck(i);
+    if (input_is_valid){


Suggestion, delete lines 118-120 and

Suggested change

if (input_is_valid){

if (input_has_nulls && input.is_valid_nocheck(i)){

That doesn't take care of the case where input_has_nulls is false, I don't think.

You are right. This would, I think.

Suggested change

if (input_is_valid){

if (!input_has_nulls || input.is_valid_nocheck(i)){

cpp/src/replace/replace.cu

Co-Authored-By: Mark Harris <[email protected]>

cpp/src/replace/replace.cu

harrism · 2019-12-05T10:17:13Z

rerun tests

cpp/src/replace/replace.cu

jrhemstad · 2019-12-05T14:53:14Z

cpp/src/replace/replace.cu

    }
    return output;
  }
+
  template<typename col_type,
      std::enable_if_t<not cudf::is_fixed_width<col_type>()>* = nullptr>


Strings may currently be the only non fixed-width type, but that won't be true in the future. This code path should only be enabled from string_view and not for all non fixed-width types.

I have updated all of these to have an explicit specialization for string_view and the general non-fixed-width type case will throw an exception.

jrhemstad · 2019-12-05T14:56:11Z

cpp/src/replace/replace.cu

+    auto device_indices = cudf::mutable_column_device_view::create(indices_view);
+
+    int valid_chars = cudf::bitmask_allocation_size_bytes(input_col.size());
+    rmm::device_buffer valid_bits(valid_chars, stream, mr);


Use create_null_mask instead https://github.com/rapidsai/cudf/blob/branch-0.11/cpp/include/cudf/null_mask.hpp#L78

jrhemstad · 2019-12-05T14:57:01Z

cpp/src/replace/replace.cu

+    // Compute the offsets from the sizes
+    int32_t * offsets_data = offsets_view.data<int32_t>();
+    thrust::exclusive_scan(rmm::exec_policy(stream)->on(stream),
+                           offsets_data,


Suggested change

offsets_data,

offsets_view.begin<int32_t>(),

jrhemstad · 2019-12-05T14:57:16Z

cpp/src/replace/replace.cu

+    int32_t * offsets_data = offsets_view.data<int32_t>();
+    thrust::exclusive_scan(rmm::exec_policy(stream)->on(stream),
+                           offsets_data,
+                           offsets_data + offsets_view.size(),


Suggested change

offsets_data + offsets_view.size(),

offsets_view.end<int32_t>(),

jrhemstad · 2019-12-05T14:57:31Z

cpp/src/replace/replace.cu

+    thrust::exclusive_scan(rmm::exec_policy(stream)->on(stream),
+                           offsets_data,
+                           offsets_data + offsets_view.size(),
+                           offsets_data);


Suggested change

offsets_data);

offsets_view.begin<int32_t>());

cpp/src/replace/replace.cu

jrhemstad · 2019-12-05T15:14:08Z

cpp/src/replace/replace.cu

@@ -87,6 +90,108 @@ __device__ auto get_new_value(cudf::size_type idx,
  return thrust::make_pair(new_value, output_is_valid);
 }

+__device__ int get_new_string_value(cudf::size_type idx,


get_new_value should be updated and then there is no need for two different functions. It should look like:

template <typename Element, bool replacement_has_nulls> __device__ auto get_new_value(size_type i, column_device_view input, cudf::column_device_view values_to_replace, column_device_view replacement_values){ auto found = thrust::find(thrust::seq, values_to_replace.begin<Element>(), values_to_replace.end<Element>(), input.element<Element>()); Element new_value; bool output_is_valid{true}; if(found != values_to_replace.end<Element>()){ auto d = thrust::distance(values_to_replace.begin<Element>(), found); new_value = replacement_values.element<Element>(d); if(replacement_has_nulls){ output_is_valid = replacement_values.is_valid(d); } } else { new_value = input.element<Element>(i); } return thrust::make_pair(new_value, output_is_valid); }

Otherwise separate function is just extra baggage that needs to be maintained.

The reason for having two functions is that the one for strings returns the index of the found value, rather than the value itself. This prevents the second pass kernel from needing to do a search again.

cpp/src/replace/replace.cu

harrism · 2019-12-06T00:04:39Z

@jrhemstad one more time...

harrism · 2019-12-06T03:16:51Z

Fixes #2944

jwyles added 4 commits November 21, 2019 13:12

Added initial implementation of strings version of find_and_replace_a…

873ec95

…ll()

WIP commit with test failure

c1ba632

Merge branch 'issue2944-2' into replace_strings

1bd4873

Bug fixes

6b6bddb

jwyles requested review from a team as code owners November 26, 2019 23:12

rgsl888prabhu reviewed Dec 2, 2019

View reviewed changes

cpp/src/replace/replace.cu Outdated Show resolved Hide resolved

cpp/src/replace/replace.cu Outdated Show resolved Hide resolved

cpp/src/replace/replace.cu Outdated Show resolved Hide resolved

cpp/tests/replace/replace_tests.cu Outdated Show resolved Hide resolved

rgsl888prabhu added 3 - Ready for Review Ready for review by team libcudf++ and removed 3 - Ready for Review Ready for review by team labels Dec 2, 2019

find_and_replace_all() now working for strings, added some tests

0aba59e

davidwendt reviewed Dec 3, 2019

View reviewed changes

jwyles added 3 commits December 3, 2019 13:53

Strings support for replace_nulls() column version implemented, tests…

fdd7db4

… added

Added missing file to commit

ced99e1

Merge remote-tracking branch 'upstream/branch-0.11' into replace_strings

ae73572

jwyles changed the title ~~[WIP] replace.hpp strings functionality~~ [Review] replace.hpp strings functionality Dec 4, 2019

davidwendt requested changes Dec 4, 2019

View reviewed changes

harrism requested changes Dec 4, 2019

View reviewed changes

harrism added 3 - Ready for Review Ready for review by team libcudf Affects libcudf (C++/CUDA) code. labels Dec 4, 2019

harrism assigned jwyles Dec 4, 2019

Updates per reviewer comments

599cab4

davidwendt requested changes Dec 4, 2019

View reviewed changes

harrism requested changes Dec 4, 2019

View reviewed changes

jwyles and others added 2 commits December 4, 2019 14:52

Update cpp/src/replace/replace.cu

e5987f9

Co-Authored-By: Mark Harris <[email protected]>

Update cpp/src/replace/replace.cu

d3570fa

Co-Authored-By: Mark Harris <[email protected]>

jwyles and others added 3 commits December 4, 2019 14:53

Update cpp/src/replace/replace.cu

c20f518

Co-Authored-By: Mark Harris <[email protected]>

Update cpp/src/replace/replace.cu

d37e518

Co-Authored-By: Mark Harris <[email protected]>

Update cpp/src/replace/replace.cu

7562059

Co-Authored-By: Mark Harris <[email protected]>

harrism approved these changes Dec 4, 2019

View reviewed changes

harrism requested a review from davidwendt December 4, 2019 22:27

davidwendt requested changes Dec 4, 2019

View reviewed changes

cpp/src/replace/replace.cu Outdated Show resolved Hide resolved

Updates per reviewer comments and updated change log

4ccdca4

davidwendt approved these changes Dec 4, 2019

View reviewed changes

davidwendt requested changes Dec 5, 2019

View reviewed changes

cpp/src/replace/replace.cu Outdated Show resolved Hide resolved

jrhemstad requested changes Dec 5, 2019

View reviewed changes

Updates per reviewer comments

cd3efd4

davidwendt approved these changes Dec 5, 2019

View reviewed changes

harrism added 5 - Ready to Merge Testing and reviews complete, ready to merge and removed 3 - Ready for Review Ready for review by team labels Dec 6, 2019

jrhemstad approved these changes Dec 6, 2019

View reviewed changes

jrhemstad merged commit 4f17e44 into rapidsai:branch-0.11 Dec 6, 2019

harrism mentioned this pull request Dec 6, 2019

[FEA] Port replace.hpp to cudf::column types #2944

Closed

	CUDA_TRY(cudaMemcpy(&size, offsets_data + offsets_view.size() - 1, sizeof(int32_t), cudaMemcpyDefault));
	CUDA_TRY(cudaMemcpyAsync(&size, offsets_data + offsets_view.size() - 1, sizeof(int32_t), cudaMemcpyDefault, stream));

	cudf::string_view inputString = input.element<cudf::string_view>(idx);
	cudf::string_view input_string = input.element<cudf::string_view>(idx);

	cudf::string_view valueString = values_to_replace.element<cudf::string_view>(i);
	cudf::string_view value_string = values_to_replace.element<cudf::string_view>(i);

	if (input_is_valid){
	if (input_has_nulls && input.is_valid_nocheck(i)){

	if (input_is_valid){
	if (!input_has_nulls \|\| input.is_valid_nocheck(i)){

	offsets_data + offsets_view.size(),
	offsets_view.end<int32_t>(),

[Review] replace.hpp strings functionality #3469

[Review] replace.hpp strings functionality #3469

Conversation

jwyles commented Nov 26, 2019

davidwendt commented Nov 27, 2019

codecov bot commented Dec 2, 2019 • edited Loading

Codecov Report

Choose a reason for hiding this comment

Choose a reason for hiding this comment

davidwendt Dec 4, 2019 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

harrism left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

harrism left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

harrism commented Dec 5, 2019

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

harrism commented Dec 6, 2019

harrism commented Dec 6, 2019

codecov bot commented Dec 2, 2019 •

edited

Loading

davidwendt Dec 4, 2019 •

edited

Loading