Add regex rewrite kernel to find `literal[a,b]{x,y}` in a string #2041

thirtiseven · 2024-05-15T10:51:24Z

Related plugin pr: NVIDIA/spark-rapids#10822

This pr adds a custom kernel to rewrite regex patterns like literal[a-b]{x,y} to get performance gain for rlike in spark-rapids.

It checks if each string contains a substring that starts with a given prefix and follows at least x characters whose codepoints are in the range of [a,b].

Signed-off-by: Haoyang Li <[email protected]>

ttnghia · 2024-05-17T04:00:48Z

src/main/cpp/src/RegexRewriteUtilsJni.cpp

+extern "C" {
+
+JNIEXPORT jlong JNICALL Java_com_nvidia_spark_rapids_jni_RegexRewriteUtils_literalRangePattern(
+  JNIEnv* env, jclass, jlong column_view, jlong target, jint d, jint start, jint end)


Do not use the existing class name as variable name to avoid potential name clashing.

Suggested change

JNIEnv* env, jclass, jlong column_view, jlong target, jint d, jint start, jint end)

JNIEnv* env, jclass, jlong input, jlong target, jint d, jint start, jint end)

ttnghia · 2024-05-17T04:02:19Z

src/main/cpp/src/regex_rewrite_utils.cu

+template <typename BoolFunction>
+std::unique_ptr<cudf::column> literal_range_pattern_fn(cudf::strings_column_view const& strings,


Functions/classes used only within the current file need to be wrapped in anonymous namespace to avoid symbol conflict.

ttnghia · 2024-05-17T04:19:04Z

src/main/cpp/src/regex_rewrite_utils.hpp

+  int const start,
+  int const end,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());


Please start using device_async_resource_ref.

Suggested change

rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());

rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());

src/main/cpp/src/regex_rewrite_utils.hpp

ttnghia · 2024-05-17T04:20:36Z

src/main/java/com/nvidia/spark/rapids/jni/RegexRewriteUtils.java

+    NativeDepsLoader.loadNativeDeps();
+  }
+
+  public static ColumnVector literalRangePattern(ColumnVector input, Scalar pattern, int d, int start, int end) {


Similarly, please add docs for this function.

ttnghia · 2024-05-17T04:21:07Z

src/main/cpp/src/regex_rewrite_utils.hpp

+#include <cudf/column/column_device_view.cuh>
+#include <cudf/column/column_factories.hpp>
+#include <cudf/detail/iterator.cuh>
+#include <cudf/detail/null_mask.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/scalar/scalar_factories.hpp>
+#include <cudf/strings/detail/utilities.hpp>
+#include <cudf/strings/find.hpp>
+#include <cudf/strings/string_view.cuh>
+#include <cudf/strings/strings_column_view.hpp>
+#include <cudf/utilities/default_stream.hpp>
+#include <cudf/utilities/error.hpp>


Not all of these headers are used in this file. Please remove the unused ones.

ttnghia · 2024-05-17T04:22:45Z

src/test/java/com/nvidia/spark/rapids/jni/RegexRewriteUtilsTest.java

+import ai.rapids.cudf.Scalar;
+import org.junit.jupiter.api.Test;
+
+import com.nvidia.spark.rapids.jni.JSONUtils;


It is used here?

ttnghia · 2024-05-17T04:24:32Z

src/main/cpp/src/regex_rewrite_utils.cu

+#include <cudf/column/column_device_view.cuh>
+#include <cudf/column/column_factories.hpp>
+#include <cudf/detail/iterator.cuh>
+#include <cudf/detail/null_mask.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/detail/utilities/cuda.cuh>
+#include <cudf/scalar/scalar_factories.hpp>
+#include <cudf/strings/detail/utf8.hpp>
+#include <cudf/strings/detail/utilities.hpp>
+#include <cudf/strings/find.hpp>
+#include <cudf/strings/string_view.cuh>
+#include <cudf/strings/strings_column_view.hpp>
+#include <cudf/utilities/default_stream.hpp>
+#include <cudf/utilities/error.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
+
+#include <cuda/atomic>
+#include <thrust/binary_search.h>
+#include <thrust/fill.h>
+#include <thrust/for_each.h>
+#include <thrust/iterator/constant_iterator.h>
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/transform.h>


I also see a lot of unused headers here. Please remove them.

ttnghia · 2024-05-17T04:25:05Z

src/main/cpp/src/regex_rewrite_utils.cu

+ * @param start Minimum code point value to check for in the range.
+ * @param end Maximum code point value to check for in the range.


What is code point?

Plugin will pass utf-8 codepoints to jni of start and end in regex pattern literal[start-end]{len,}

ttnghia · 2024-05-17T04:25:58Z

src/main/cpp/src/regex_rewrite_utils.cu

+                                                       rmm::cuda_stream_view stream,
+                                                       rmm::mr::device_memory_resource* mr)
+{
+  auto strings_count = strings.size();


Suggested change

auto strings_count = strings.size();

auto const strings_count = strings.size();

ttnghia · 2024-05-17T04:26:34Z

src/main/cpp/src/regex_rewrite_utils.cu

+  auto d_prefix       = cudf::string_view(prefix.data(), prefix.size());
+  auto strings_column = cudf::column_device_view::create(strings.parent(), stream);
+  auto d_strings      = *strings_column;


d_strings should be a reference to avoid copying.

Suggested change

auto d_prefix = cudf::string_view(prefix.data(), prefix.size());

auto strings_column = cudf::column_device_view::create(strings.parent(), stream);

auto d_strings = *strings_column;

auto const d_prefix = cudf::string_view(prefix.data(), prefix.size());

auto const strings_column = cudf::column_device_view::create(strings.parent(), stream);

auto const& d_strings = *strings_column;

ttnghia · 2024-05-17T04:26:57Z

src/main/cpp/src/regex_rewrite_utils.cu

+  auto d_prefix       = cudf::string_view(prefix.data(), prefix.size());
+  auto strings_column = cudf::column_device_view::create(strings.parent(), stream);
+  auto d_strings      = *strings_column;
+  // create output column


It is very trivial 😄

Suggested change

// create output column

ttnghia · 2024-05-17T04:27:37Z

src/main/cpp/src/regex_rewrite_utils.cu

+  auto results_view = results->mutable_view();
+  auto d_results    = results_view.data<bool>();


Suggested change

auto results_view = results->mutable_view();

auto d_results = results_view.data<bool>();

auto const d_results = results->mutable_view().data<bool>();

ttnghia · 2024-05-17T04:30:13Z

src/main/cpp/src/regex_rewrite_utils.cu

+                                                    rmm::cuda_stream_view stream,
+                                                    rmm::mr::device_memory_resource* mr)
+{
+  auto pfn =


Suggested change

auto pfn =

auto const pfn =

In addition, it is better to define this lamba as a separate functor in an anonymous namespace:

Suggested change

auto pfn =

namespace {

struct literal_range_pattern_fn {

__device__ bool operator()(...) const {

....

}

};

}

The template function literal_range_pattern_fn above can be renamed into something like:

namespace { std::unique_ptr<cudf::column> find_literal_range_pattern(....) { ... } }

Signed-off-by: Haoyang Li <[email protected]>

thirtiseven · 2024-05-17T07:01:01Z

@ttnghia Thanks for the review, I think all addressed now. Will check those comments like using const and removing useless headers locally first next time 👀

ttnghia

Sorry I have a couple of more requests.

src/main/cpp/src/regex_rewrite_utils.cu

ttnghia · 2024-05-21T04:01:50Z

src/main/cpp/src/regex_rewrite_utils.cu

+  auto results         = make_numeric_column(cudf::data_type{cudf::type_id::BOOL8},
+                                     strings_count,
+                                     cudf::detail::copy_bitmask(strings.parent(), stream, mr),
+                                     strings.null_count(),
+                                     stream,
+                                     mr);


I suspect that the file was not properly formated (using clang-format).

It passed clang-format in the pre-commit hook

src/main/cpp/src/regex_rewrite_utils.cu

Co-authored-by: Nghia Truong <[email protected]>

Signed-off-by: Haoyang Li <[email protected]>

thirtiseven · 2024-05-21T05:50:42Z

build

Signed-off-by: Haoyang Li <[email protected]>

thirtiseven · 2024-05-21T07:26:30Z

build

thirtiseven added 12 commits April 22, 2024 18:18

wip

f54ea85

Signed-off-by: Haoyang Li <[email protected]>

wip

6fffca8

Signed-off-by: Haoyang Li <[email protected]>

wip

64d3784

Signed-off-by: Haoyang Li <[email protected]>

Merge branch 'branch-24.06' into str_dig

5e8fe37

wip

7b8d375

Signed-off-by: Haoyang Li <[email protected]>

Merge branch 'branch-24.06' into str_dig

a92beea

Merge branch 'branch-24.06' into str_dig

21c19be

Merge branch 'branch-24.06' into str_dig

28f0ad5

support range filter

d30e025

Signed-off-by: Haoyang Li <[email protected]>

clean up

8293324

Signed-off-by: Haoyang Li <[email protected]>

Merge branch 'branch-24.06' into str_dig

03ba4fa

wip

ac6575c

Signed-off-by: Haoyang Li <[email protected]>

thirtiseven mentioned this pull request May 15, 2024

Rewrite regex pattern literal[a-b]{x} to custom kernel in rlike NVIDIA/spark-rapids#10822

Merged

thirtiseven added 2 commits May 16, 2024 11:39

change some names

c335a8d

Signed-off-by: Haoyang Li <[email protected]>

clean up

cf9fde7

Signed-off-by: Haoyang Li <[email protected]>

thirtiseven marked this pull request as ready for review May 16, 2024 09:55

thirtiseven requested review from revans2, mythrocks and NVnavkumar May 16, 2024 10:00

ttnghia reviewed May 17, 2024

View reviewed changes

src/main/cpp/src/regex_rewrite_utils.hpp Show resolved Hide resolved

ttnghia reviewed May 17, 2024

View reviewed changes

address comments

9e7b06a

Signed-off-by: Haoyang Li <[email protected]>

thirtiseven self-assigned this May 21, 2024

thirtiseven requested a review from ttnghia May 21, 2024 01:54

thirtiseven mentioned this pull request May 21, 2024

Add regex rewrite kernel to find literal[a,b]{x,y} in a string wjxiz1992/spark-rapids-jni#1

Merged

ttnghia requested changes May 21, 2024

View reviewed changes

thirtiseven and others added 2 commits May 21, 2024 13:24

Apply suggestions from code review

ee96c44

Co-authored-by: Nghia Truong <[email protected]>

format

0be1c23

Signed-off-by: Haoyang Li <[email protected]>

ttnghia previously approved these changes May 21, 2024

View reviewed changes

fix build

7e46ab2

Signed-off-by: Haoyang Li <[email protected]>

thirtiseven dismissed ttnghia’s stale review via 7e46ab2 May 21, 2024 07:25

thirtiseven requested a review from ttnghia May 21, 2024 16:43

ttnghia approved these changes May 21, 2024

View reviewed changes

thirtiseven merged commit 956fb53 into NVIDIA:branch-24.06 May 21, 2024
3 checks passed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Add regex rewrite kernel to find `literal[a,b]{x,y}` in a string #2041

Add regex rewrite kernel to find `literal[a,b]{x,y}` in a string #2041

thirtiseven commented May 15, 2024 •

edited

Loading

ttnghia May 17, 2024

ttnghia May 17, 2024

ttnghia May 17, 2024

ttnghia May 17, 2024

ttnghia May 17, 2024 •

edited

Loading

ttnghia May 17, 2024

ttnghia May 17, 2024

ttnghia May 17, 2024

thirtiseven May 17, 2024

ttnghia May 17, 2024

ttnghia May 17, 2024

ttnghia May 17, 2024

ttnghia May 17, 2024

ttnghia May 17, 2024

ttnghia May 17, 2024

thirtiseven commented May 17, 2024

ttnghia left a comment

ttnghia May 21, 2024

thirtiseven May 21, 2024

thirtiseven commented May 21, 2024

thirtiseven commented May 21, 2024

	JNIEnv* env, jclass, jlong column_view, jlong target, jint d, jint start, jint end)
	JNIEnv* env, jclass, jlong input, jlong target, jint d, jint start, jint end)

		template <typename BoolFunction>
		std::unique_ptr<cudf::column> literal_range_pattern_fn(cudf::strings_column_view const& strings,

	rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
	rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());

		* @param start Minimum code point value to check for in the range.
		* @param end Maximum code point value to check for in the range.

	auto strings_count = strings.size();
	auto const strings_count = strings.size();

		auto results_view = results->mutable_view();
		auto d_results = results_view.data<bool>();

	auto results_view = results->mutable_view();
	auto d_results = results_view.data<bool>();
	auto const d_results = results->mutable_view().data<bool>();

-  auto pfn =
+namespace {
+  struct literal_range_pattern_fn {
+    __device__ bool operator()(...) const {
+    ....
+    }
+    };
+  }

Add regex rewrite kernel to find literal[a,b]{x,y} in a string #2041

Add regex rewrite kernel to find literal[a,b]{x,y} in a string #2041

Conversation

thirtiseven commented May 15, 2024 • edited Loading

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

ttnghia May 17, 2024 • edited Loading

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

thirtiseven commented May 17, 2024

ttnghia left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

thirtiseven commented May 21, 2024

thirtiseven commented May 21, 2024

Add regex rewrite kernel to find `literal[a,b]{x,y}` in a string #2041

Add regex rewrite kernel to find `literal[a,b]{x,y}` in a string #2041

thirtiseven commented May 15, 2024 •

edited

Loading

ttnghia May 17, 2024 •

edited

Loading