Address Brad's review: stream-safe table init, docs, C++ stream tests

- strings/utilities.cu: synchronize the stream before returning the cached
  flags/cases/special-case tables in each find_or_initialize lambda, so the
  first-use async copy completes before the shared pointer is handed to any
  caller/stream (stream safety).
- char_types.hpp: document the special-case table's void* entry layout
  (16-byte special_case_mapping; counts at 0/8, code points at 2/10) and the
  cp % 499 hash index as part of the public contract.
- tests/streams/strings/filter_test.cpp: add a CharacterTables stream test
  exercising the three table accessors on the test stream.
- replace docstring/comments + empty-old test doc: state the empty-target
  behavior matches Python's str.replace (libcudf rejects an empty target),
  per Brad's clarification.

Author: brandon-b-miller

cpp/include/cudf/strings/char_types/char_types.hpp

@@ -148,6 +148,18 @@ std::uint16_t const* get_character_cases_table(
  * call the table is copied into device memory and cached; subsequent calls return the
  * cached pointer. This call is thread-safe.
  *
+ * The pointer is returned as `void const*` because the entry type lives in a detail
+ * header. The layout is part of this public contract: it is an array of fixed
+ * 16-byte entries matching `cudf::strings::detail::special_case_mapping`, i.e. for a
+ * little-endian `uint16_t` view of an entry `e`:
+ *   - `e[0]`      : number of upper-case code points (0..3), at byte offset 0
+ *   - `e[1..3]`   : up to 3 upper-case code points,         at byte offset 2
+ *   - `e[4]`      : number of lower-case code points (0..3), at byte offset 8
+ *   - `e[5..7]`   : up to 3 lower-case code points,         at byte offset 10
+ * A code point is looked up at entry index `code_point % 499`
+ * (`cudf::strings::detail::get_special_case_hash_index`). Changing either the
+ * `special_case_mapping` layout or the hash prime is a change to this public API.
+ *
  * @param stream CUDA stream used for device memory operations and kernel launches
  * @return Device memory pointer to the special-case mapping table
  */

cpp/src/strings/utilities.cu

@@ -107,6 +107,10 @@ character_flags_table_type const* get_character_flags_table(cuda::stream_ref str
                                           cudaMemcpyDefault,
                                           stream.get()));
     CUDF_CUDA_TRY(cudaGetSymbolAddress((void**)&table, character_codepoint_flags));
+    // The copy above is asynchronous and the cached pointer is shared across
+    // callers/streams; synchronize so the table is fully populated before it is
+    // handed out (first-use only).
+    CUDF_CUDA_TRY(cudaStreamSynchronize(stream.get()));
     return table;
   });
 }
@@ -125,6 +129,10 @@ character_cases_table_type const* get_character_cases_table(cuda::stream_ref str
                                           cudaMemcpyDefault,
                                           stream.get()));
     CUDF_CUDA_TRY(cudaGetSymbolAddress((void**)&table, character_cases_table));
+    // The copy above is asynchronous and the cached pointer is shared across
+    // callers/streams; synchronize so the table is fully populated before it is
+    // handed out (first-use only).
+    CUDF_CUDA_TRY(cudaStreamSynchronize(stream.get()));
     return table;
   });
 }
@@ -143,6 +151,10 @@ special_case_mapping const* get_special_case_mapping_table(cuda::stream_ref stre
                                           cudaMemcpyDefault,
                                           stream.get()));
     CUDF_CUDA_TRY(cudaGetSymbolAddress((void**)&table, character_special_case_mappings));
+    // The copy above is asynchronous and the cached pointer is shared across
+    // callers/streams; synchronize so the table is fully populated before it is
+    // handed out (first-use only).
+    CUDF_CUDA_TRY(cudaStreamSynchronize(stream.get()));
     return table;
   });
 }

cpp/tests/streams/strings/filter_test.cpp

@@ -64,3 +64,12 @@ TEST_F(StringsFilterTest, FilterTypes)
   cudf::strings::filter_characters_of_type(
     view, all_types, repl, space_types, cudf::test::get_default_stream());
 }
+
+TEST_F(StringsFilterTest, CharacterTables)
+{
+  // The character flags/cases/special-case tables are lazily copied to the
+  // device on first use using the supplied stream.
+  cudf::strings::get_character_flags_table(cudf::test::get_default_stream());
+  cudf::strings::get_character_cases_table(cudf::test::get_default_stream());
+  cudf::strings::get_special_case_mapping_table(cudf::test::get_default_stream());
+}

python/cudf/cudf/core/udf/mlir_backend/string_lowering_impl.py

@@ -394,7 +394,8 @@ def lower_replace(
     Byte-level substring matching (equivalent to character matching for
     well-formed UTF-8). Two passes: count occurrences to size the result, then
     scan-and-copy emitting replacement or source bytes. An empty ``old`` inserts
-    ``new`` around every byte, matching libcudf's replace.
+    ``new`` before every character and at the end, matching Python's
+    ``str.replace`` (libcudf's ``replace`` rejects an empty target).
 
     Parameters
     ----------
@@ -456,7 +457,7 @@ def lower_replace(
         c_new_skip = arith.select(c_valid, arith.addi(ci, o_nb), c_skip)
         scf.YieldOp([c_new_count, c_new_skip])
 
-    # For an empty target, libcudf inserts ``new`` before each *character*
+    # For an empty target, Python inserts ``new`` before each *character*
     # (UTF-8 start byte) and once more at the end; counting bytes would split
     # multibyte characters and produce invalid UTF-8. Use the character count
     # (+1 for the trailing insertion) as the occurrence count in that case.
@@ -504,7 +505,7 @@ def lower_replace(
             b_match_valid, arith.cmpi(eq, target_empty, false_val)
         )
         # Empty target: insert ``new`` only before UTF-8 start bytes so
-        # multibyte characters are not split (matches libcudf).
+        # multibyte characters are not split (matches Python's str.replace).
         is_start_byte = _is_begin_utf8_char(_byte_at(s_data, si))
         write_repl = arith.ori(
             arith.andi(arith.andi(target_empty, active), is_start_byte),

python/cudf/cudf/tests/private_objects/mlir_backend/test_strings_lowering.py

@@ -1368,7 +1368,7 @@ def test_replace_empty_old():
     """``s.replace("", new)`` inserts ``new`` around each character.
 
     For multibyte sources the insertion must land on UTF-8 character
-    boundaries (start bytes) so the result stays valid UTF-8, matching libcudf.
+    boundaries (start bytes) so the result stays valid UTF-8, matching Python.
     """
     vals = ["abc", "h\u00e9llo", "\U0001f600x", "", "z"]
     expected = [v.replace("", "-") for v in vals]