Add Glushkov regex fast-path to libcudf

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

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -33,16 +33,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();
   }
 };
 
@@ -60,15 +60,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;
 }
@@ -96,12 +99,9 @@ std::unique_ptr<column> count_re(strings_column_view const& input,
                                  rmm::cuda_stream_view stream,
                                  rmm::device_async_resource_ref mr)
 {
-  // create device object from regex_program
-  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