Add Glushkov regex fast-path to libcudf (#21936)

### Add bit-parallel Glushkov NFA regex engine with shared memory optimization
Implement Glushkov's NFA for regex string matching in libcudf: references (1) [hyperscan paper](https://www.usenix.org/system/files/nsdi19-wang-xiang.pdf) (2) [HybridSA paper](https://dl.acm.org/doi/10.1145/3689771) (3) [vectorscan repo](https://github.com/Vectorcamp/vectorscan). 
Basically, this is Glushkov's NFA compared with the other popular Thompson's NFA (also used in current libcudf regex). The Glushkov engine represents NFA state as a single uint64_t bitmask (max 64 positions), requiring no global memory per thread. Shared memory is used to hold the static instructions like in the current implementation.

### Key changes
  - Two-phase O(n) unanchored search algorithm (glushkov.inl): Phase 1 scans forward, injecting start states each character and recording provisional match ends. Phase 2 rescans only the match region to find the true leftmost start. Each character is processed at most twice.
  - Leftmost-first correctness via priority-kill (glushkov.cuh, glushkov_regcomp.cpp): A runtime glushkov_priority_kill clears lower-priority alternative paths at accept time. A compile-time conflict detector (frontier_has_priority_conflict) conservatively falls back to Thompson when bit-index ordering cannot guarantee Thompson-compatible leftmost-first semantics.
  - Automatic fallback: Patterns with anchors (^, $, \b, \B), >64 positions, match empty top-level expressions, capture group requirements (extract, backref_re), or priority conflicts transparently fall back to Thompson NFA — no user intervention needed.

### Limitations
- does not support capturing groups (e.g. extract, extract_all, findall, replace_with_backrefs)
- does not support zero-width assertions (empty-matchable) like BOL/EOL/BOW/NBOW
- max 64 character-consuming positions since we are using uint64_t as state data per row/thread
- does not support lazy quantifiers
- empty/degenerate patterns rejected
- does not support empty-matchable patterns as well as some ambiguous alternation patterns

When above condition is detected, it falls back to use the current Thompson's NFA.


### Unit tests + benchmark
  - Priority-kill parity tests: Verify Glushkov matches Thompson for overlapping-prefix alternations (foo|foobar, cat|catch, a|aa) across all 5 operations (contains, count, findall, replace, split)
  - Empty-matchable fallback parity: Confirm nullable patterns (a*, \d*, (ab)?) transparently fall back to Thompson and produce identical results
  - Spark-rapids compatibility: ~60 regex patterns from spark-rapids integration tests validated under both engines via parametrized Python tests
  - Benchmarks: 6–9 patterns per benchmark covering char classes, alternation, bounded repetition, dot wildcards, and late-failure stress patterns
  - Extended more complex regexes in the current split_re/contains/replace_re/count, it showed 1.01-6.62x speedup.

Authors:
  - Lingyan Yin (https://github.com/lingyany-nv)
  - David Wendt (https://github.com/davidwendt)

Approvers:
  - Basit Ayantunde (https://github.com/lamarrr)
  - Yunsong Wang (https://github.com/PointKernel)
  - Bradley Dice (https://github.com/bdice)

URL: #21936

Author: lingyany-nv

cpp/CMakeLists.txt

@@ -1045,6 +1045,8 @@ add_library(
   src/strings/merge/merge.cu
   src/strings/padding.cu
   src/strings/positions.cu
+  src/strings/regex/gkexec.cpp
+  src/strings/regex/glushkov_regcomp.cpp
   src/strings/regex/regcomp.cpp
   src/strings/regex/regexec.cpp
   src/strings/regex/regex_program.cpp

cpp/src/strings/char_types/char_flags.h

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2023, NVIDIA CORPORATION.  All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -23,7 +23,8 @@
 //    0 - decimal
 //
 
-uint8_t const g_character_codepoint_flags[] = {
+// clang-format off
+uint8_t const g_character_codepoint_flags[] = { // NOLINT
   0,   0,   0,   0,   0,   0,   0,   0,   0,   16,  16,  16,  16,  16,  0,   0,   0,   0,   0,
   0,   0,   0,   0,   0,   0,   0,   0,   0,   16,  16,  16,  16,  16,  0,   0,   0,   0,   0,
   0,   0,   0,   0,   0,   0,   0,   0,   0,   0,   7,   7,   7,   7,   7,   7,   7,   7,   7,
@@ -3475,3 +3476,4 @@ uint8_t const g_character_codepoint_flags[] = {
   0,   0,   0,   0,   0,   0,   0,   0,   0,   0,   0,   0,   0,   0,   0,   0,   0,   0,   0,
   0,   0,   0,   0,   0,
 };
+// clang-format on

cpp/src/strings/contains.cu

@@ -35,16 +35,16 @@ struct contains_fn {
   column_device_view const d_strings;
   bool const beginning_only;
 
-  __device__ bool operator()(size_type const idx,
-                             reprog_device const prog,
-                             int32_t const thread_idx)
+  template <typename ProgDevice>
+  __device__ bool operator()(size_type const idx, ProgDevice const prog, int32_t const thread_idx)
   {
     if (d_strings.is_null(idx)) return false;
     auto const d_str = d_strings.element<string_view>(idx);
 
     size_type end = beginning_only ? 1    // match only the beginning of the string;
                                    : -1;  // match anywhere in the string
-    return prog.find<positional::END_ONLY>(thread_idx, d_str, d_str.begin(), end).has_value();
+    return prog.template find<positional::END_ONLY>(thread_idx, d_str, d_str.begin(), end)
+      .has_value();
   }
 };
 
@@ -62,15 +62,18 @@ std::unique_ptr<column> contains_impl(strings_column_view const& input,
                                      mr);
   if (input.is_empty()) { return results; }
 
-  auto d_prog = regex_device_builder::create_prog_device(prog, stream);
-
   auto d_results       = results->mutable_view().data<bool>();
   auto const d_strings = column_device_view::create(input.parent(), stream);
 
-  launch_transform_kernel(
-    contains_fn{*d_strings, beginning_only}, *d_prog, d_results, input.size(), stream);
-
-  results->set_null_count(input.null_count());
+  if (regex_device_builder::glushkov_fast_path_supported(prog)) {
+    auto d_prog = regex_device_builder::create_gkprog_device(prog, stream);
+    launch_transform_kernel(
+      contains_fn{*d_strings, beginning_only}, *d_prog, d_results, input.size(), stream);
+  } else {
+    auto d_prog = regex_device_builder::create_prog_device(prog, stream);
+    launch_transform_kernel(
+      contains_fn{*d_strings, beginning_only}, *d_prog, d_results, input.size(), stream);
+  }
 
   return results;
 }
@@ -133,11 +136,9 @@ std::unique_ptr<column> count_re(strings_column_view const& input,
     return count(input, target, stream, mr);
   }
 
-  auto d_prog = regex_device_builder::create_prog_device(prog, stream);
-
   auto const d_strings = column_device_view::create(input.parent(), stream);
 
-  auto result = count_matches(*d_strings, *d_prog, stream, mr);
+  auto result = count_matches(*d_strings, prog, stream, mr);
   if (input.has_nulls()) {
     result->set_null_mask(cudf::detail::copy_bitmask(input.parent(), stream, mr),
                           input.null_count());

cpp/src/strings/count_matches.cu

@@ -1,16 +1,20 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
 #include "strings/count_matches.hpp"
+#include "strings/regex/regex_program_impl.h"
 #include "strings/regex/utilities.cuh"
 
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
+#include <cudf/strings/regex/regex_program.hpp>
 #include <cudf/strings/string_view.cuh>
 #include <cudf/utilities/memory_resource.hpp>
 
+#include <type_traits>
+
 namespace cudf {
 namespace strings {
 namespace detail {
@@ -23,8 +27,9 @@ template <positional P>
 struct count_fn {
   column_device_view const d_strings;
 
+  template <typename ProgDevice>
   __device__ int32_t operator()(size_type const idx,
-                                reprog_device const prog,
+                                ProgDevice const prog,
                                 int32_t const thread_idx)
   {
     if (d_strings.is_null(idx)) return 0;
@@ -34,7 +39,7 @@ struct count_fn {
 
     auto itr = d_str.begin();
     while (itr.position() <= nchars) {
-      auto result = prog.find<P>(thread_idx, d_str, itr);
+      auto result = prog.template find<P>(thread_idx, d_str, itr);
       if (!result) { break; }
       ++count;
       // increment the iterator is faster than creating a new one
@@ -47,29 +52,64 @@ struct count_fn {
 
 }  // namespace
 
+template <typename ProgDevice>
 std::unique_ptr<column> count_matches(column_device_view const& d_strings,
-                                      reprog_device& d_prog,
+                                      ProgDevice& d_prog,
+                                      size_type strings_count,
                                       rmm::cuda_stream_view stream,
                                       rmm::device_async_resource_ref mr)
 {
   auto results = make_numeric_column(
-    data_type{type_to_id<size_type>()}, d_strings.size(), mask_state::UNALLOCATED, stream, mr);
+    data_type{type_to_id<size_type>()}, strings_count, mask_state::UNALLOCATED, stream, mr);
 
-  if (d_strings.size() == 0) { return results; }
+  if (strings_count == 0) { return results; }
 
   auto d_results = results->mutable_view().data<cudf::size_type>();
 
-  if (d_prog.is_empty_match_possible()) {
+  // Glushkov's engine always requires the begin/end positional check; the Thompson
+  // engine can skip it (cheaper) when an empty match is not possible for this pattern.
+  if constexpr (std::is_same_v<ProgDevice, gkprog_device>) {
     launch_transform_kernel(
-      count_fn<positional::BEGIN_END>{d_strings}, d_prog, d_results, d_strings.size(), stream);
+      count_fn<positional::BEGIN_END>{d_strings}, d_prog, d_results, strings_count, stream);
   } else {
-    launch_transform_kernel(
-      count_fn<positional::END_ONLY>{d_strings}, d_prog, d_results, d_strings.size(), stream);
+    if (d_prog.is_empty_match_possible()) {
+      launch_transform_kernel(
+        count_fn<positional::BEGIN_END>{d_strings}, d_prog, d_results, strings_count, stream);
+    } else {
+      launch_transform_kernel(
+        count_fn<positional::END_ONLY>{d_strings}, d_prog, d_results, strings_count, stream);
+    }
   }
 
   return results;
 }
 
+template std::unique_ptr<column> count_matches<reprog_device>(column_device_view const&,
+                                                              reprog_device&,
+                                                              size_type,
+                                                              rmm::cuda_stream_view,
+                                                              rmm::device_async_resource_ref);
+
+template std::unique_ptr<column> count_matches<gkprog_device>(column_device_view const&,
+                                                              gkprog_device&,
+                                                              size_type,
+                                                              rmm::cuda_stream_view,
+                                                              rmm::device_async_resource_ref);
+
+std::unique_ptr<column> count_matches(column_device_view const& d_strings,
+                                      regex_program const& prog,
+                                      rmm::cuda_stream_view stream,
+                                      rmm::device_async_resource_ref mr)
+{
+  auto const strings_count = d_strings.size();
+  if (regex_device_builder::glushkov_fast_path_supported(prog)) {
+    auto d_prog = regex_device_builder::create_gkprog_device(prog, stream);
+    return count_matches(d_strings, *d_prog, strings_count, stream, mr);
+  }
+  auto d_prog = regex_device_builder::create_prog_device(prog, stream);
+  return count_matches(d_strings, *d_prog, strings_count, stream, mr);
+}
+
 }  // namespace detail
 }  // namespace strings
 }  // namespace cudf

cpp/src/strings/count_matches.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -15,23 +15,51 @@ namespace cudf {
 class column_device_view;
 
 namespace strings {
+
+class regex_program;
+
 namespace detail {
 
-class reprog_device;
+/**
+ * @brief Returns a column of regex match counts for each string in the given column.
+ *
+ * A null entry will result in a zero count for that output row.
+ *
+ * This overload evaluates against an already-built device regex program. Callers that
+ * also need the device program for other work (e.g. extraction) should build it once
+ * and pass it here to avoid a redundant device program build.
+ *
+ * @tparam The regex prog device instance used for this API
+ * @param d_strings Device view of the input strings column
+ * @param d_prog Device regex program to evaluate on each string
+ * @param strings_count Number of strings (and rows in the output column)
+ * @param stream CUDA stream used for device memory operations and kernel launches
+ * @param mr Device memory resource used to allocate the returned column's device memory
+ * @return Integer column of match counts
+ */
+template <typename ProgDevice>
+std::unique_ptr<column> count_matches(column_device_view const& d_strings,
+                                      ProgDevice& d_prog,
+                                      size_type strings_count,
+                                      rmm::cuda_stream_view stream,
+                                      rmm::device_async_resource_ref mr);
 
 /**
  * @brief Returns a column of regex match counts for each string in the given column.
  *
  * A null entry will result in a zero count for that output row.
  *
+ * This overload builds its own device regex program. Prefer the overload above when
+ * the device program is also needed for other work on the same call site.
+ *
  * @param d_strings Device view of the input strings column.
- * @param d_prog Regex instance to evaluate on each string.
+ * @param prog Regex program to evaluate on each string.
  * @param stream CUDA stream used for device memory operations and kernel launches.
  * @param mr Device memory resource used to allocate the returned column's device memory.
  * @return Integer column of match counts
  */
 std::unique_ptr<column> count_matches(column_device_view const& d_strings,
-                                      reprog_device& d_prog,
+                                      regex_program const& prog,
                                       rmm::cuda_stream_view stream,
                                       rmm::device_async_resource_ref mr);
 

cpp/src/strings/extract/extract_all.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -106,9 +106,10 @@ std::unique_ptr<column> extract_all_record(strings_column_view const& input,
   auto const groups = d_prog->group_counts();
   CUDF_EXPECTS(groups > 0, "extract_all requires group indicators in the regex pattern.");
 
-  // Get the match counts for each string.
+  // Get the match counts for each string, reusing the device program built above
+  // instead of building a second one.
   // This column will become the output lists child offsets column.
-  auto counts   = count_matches(*d_strings, *d_prog, stream, mr);
+  auto counts   = count_matches(*d_strings, *d_prog, strings_count, stream, mr);
   auto d_counts = counts->mutable_view().data<size_type>();
 
   // Compute null output rows

cpp/src/strings/regex/common.cuh

@@ -0,0 +1,90 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include "regcomp.h"
+
+#include <cudf/strings/detail/char_tables.hpp>
+#include <cudf/strings/detail/utf8.hpp>
+#include <cudf/types.hpp>
+
+#include <cuda/std/optional>
+#include <cuda/std/utility>
+#include <thrust/execution_policy.h>
+#include <thrust/logical.h>
+
+namespace cudf::strings::detail {
+
+/// Bitmask type: bit i is set when Glushkov position i is active.
+using glushkov_state_t = uint64_t;
+
+/// Maximum number of character-consuming positions (states) in the Glushkov NFA.
+/// Patterns with more positions fall back to Thompson NFA automatically.
+constexpr int32_t GLUSHKOV_MAX_STATES = sizeof(glushkov_state_t) * 8;
+
+/// Maximum shift amounts for the Hyperscan-style shift-and optimization.
+constexpr int32_t GLUSHKOV_MAX_SHIFTS = 8;
+
+/// Size of the precomputed ASCII reach table (characters 0–127).
+constexpr int32_t GLUSHKOV_ASCII_TABLE_SIZE = 128;
+
+/**
+ * @brief Regex class stored on the device and executed by reprog_device.
+ *
+ * This class holds the unique data for any regex CCLASS instruction.
+ */
+struct alignas(16) reclass_device {
+  int32_t builtins{};
+  int32_t count{};
+  reclass_range const* literals{};
+
+  __device__ inline bool is_match(char32_t const ch, uint8_t const* codepoint_flags) const
+  {
+    if (thrust::any_of(thrust::seq, literals, literals + count, [ch](auto literal) {
+          return ((ch >= literal.first) && (ch <= literal.last));
+        })) {
+      return true;
+    }
+
+    if (!builtins) { return false; }
+    auto const codept                = utf8_to_codepoint(ch);
+    constexpr uint32_t MAX_CODEPOINT = 0x00'FFFF;
+    if (codept > MAX_CODEPOINT) { return false; }
+    auto const fl = codepoint_flags[codept];
+    if ((builtins & CCLASS_W) && ((ch == '_') || IS_ALPHANUM(fl))) { return true; }     // \w
+    if ((builtins & CCLASS_S) && IS_SPACE(fl)) { return true; }                         // \s
+    if ((builtins & CCLASS_D) && IS_DIGIT(fl)) { return true; }                         // \d
+    if ((builtins & NCCLASS_W) && ((ch != '\n') && (ch != '_') && !IS_ALPHANUM(fl))) {  // \W
+      return true;
+    }
+    if ((builtins & NCCLASS_S) && !IS_SPACE(fl)) { return true; }                    // \S
+    if ((builtins & NCCLASS_D) && ((ch != '\n') && !IS_DIGIT(fl))) { return true; }  // \D
+
+    return false;
+  }
+};
+
+/**
+ * @brief Check for supported new-line characters
+ *
+ * '\n, \r, \u0085, \u2028, or \u2029'
+ */
+CUDF_HOST_DEVICE __forceinline__ constexpr bool is_newline(char32_t const ch)
+{
+  return (ch == '\n' || ch == '\r' || ch == 0x00c285 || ch == 0x00e280a8 || ch == 0x00e280a9);
+}
+
+/**
+ * @brief Template type used on `find` to specify desired position values in returned match_result
+ */
+enum class positional : int8_t {
+  BEGIN_END = 0,  /// both begin and end positions are returned
+  END_ONLY  = 1,  /// only the end position is returned
+};
+
+using match_pair   = cuda::std::pair<cudf::size_type, cudf::size_type>;
+using match_result = cuda::std::optional<match_pair>;
+
+}  // namespace cudf::strings::detail