[FEA] Add fixed-width list support to JIT transforms

cpp/CMakeLists.txt

@@ -513,7 +513,7 @@ function(precompile_fixed_width_kernel_fragments)
     )
     set(VARIANT_NAME transform_kernel_${FILE_INDEX})
     set(INSTANCE
-        "cudf::jit::transform_kernel<false, false, cudf::jit::type_list<cudf::jit::column_accessor<0, cudf::column_device_view_core, ${TYPE}, false, 0>>, cudf::jit::type_list<cudf::jit::column_accessor<0, cudf::mutable_column_device_view_core, ${TYPE}, false, 0>>>"
+        "cudf::jit::transform_kernel<false, false, cudf::jit::type_list<cudf::jit::column_accessor<0, ${TYPE}, false, 0>>, cudf::jit::type_list<cudf::jit::column_accessor<0, ${TYPE}, false, 0>>>"
     )
     add_fragment(
       cudf_fragments
@@ -548,7 +548,7 @@ function(precompile_fixed_width_kernel_fragments)
       )
       set(VARIANT_NAME transform_kernel_${FILE_INDEX})
       set(INSTANCE
-          "cudf::jit::transform_kernel<false, false, cudf::jit::type_list<cudf::jit::column_accessor<0, cudf::column_device_view_core, ${TYPE}, false, 0>, cudf::jit::column_accessor<1, cudf::column_device_view_core, ${TYPE}, ${RHS_IS_SCALAR}, 0>>, cudf::jit::type_list<cudf::jit::column_accessor<0, cudf::mutable_column_device_view_core, ${TYPE}, false, 0>>>"
+          "cudf::jit::transform_kernel<false, false, cudf::jit::type_list<cudf::jit::column_accessor<0, ${TYPE}, false, 0>, cudf::jit::column_accessor<1, ${TYPE}, ${RHS_IS_SCALAR}, 0>>, cudf::jit::type_list<cudf::jit::column_accessor<0, ${TYPE}, false, 0>>>"
       )
       add_fragment(
         cudf_fragments
@@ -578,20 +578,9 @@ endfunction()
 # amortize the cost of JIT compilation for kernels that have matching signatures
 function(precompile_string_kernel_fragments)
   list(APPEND CUDF_PRECOMPILE_STRING_OUTPUT_PHYSICAL_TYPES uint8_t uint16_t uint32_t uint64_t
-       cuda::std::span<char>
-  )
-  list(
-    APPEND
-    CUDF_PRECOMPILE_STRING_OUTPUT_COLUMN_TYPES
-    cudf::mutable_column_device_view_core
-    cudf::mutable_column_device_view_core
-    cudf::mutable_column_device_view_core
-    cudf::mutable_column_device_view_core
-    cudf::jit::mutable_strings_column_device_view
-  )
-  foreach(OUTPUT_TYPE OUTPUT_COLUMN_TYPE IN ZIP_LISTS CUDF_PRECOMPILE_STRING_OUTPUT_PHYSICAL_TYPES
-                      CUDF_PRECOMPILE_STRING_OUTPUT_COLUMN_TYPES
+       cudf::mutable_string_view
   )
+  foreach(OUTPUT_TYPE IN LISTS CUDF_PRECOMPILE_STRING_OUTPUT_PHYSICAL_TYPES)
     # Pre-compile for up to 8 output columns
     foreach(OUTPUT_INDICES IN ITEMS "0" "0;1" "0;1;2" "0;1;2;3" "0;1;2;3;4" "0;1;2;3;4;5"
                                     "0;1;2;3;4;5;6" "0;1;2;3;4;5;6;7"
@@ -606,17 +595,15 @@ function(precompile_string_kernel_fragments)
       set(OUTPUT_ACCESSORS "")
       set(OUTPUT_POINTERS "")
       foreach(OUTPUT_INDEX IN LISTS OUTPUT_INDICES)
-        list(
-          APPEND
-          OUTPUT_ACCESSORS
-          "cudf::jit::column_accessor<${OUTPUT_INDEX}, ${OUTPUT_COLUMN_TYPE}, ${OUTPUT_TYPE}, false, 0>"
+        list(APPEND OUTPUT_ACCESSORS
+             "cudf::jit::column_accessor<${OUTPUT_INDEX}, ${OUTPUT_TYPE}, false, 0>"
         )
         list(APPEND OUTPUT_POINTERS "${OUTPUT_TYPE} *")
       endforeach()
       list(JOIN OUTPUT_ACCESSORS " ," OUTPUT_ACCESSORS_STR)
       list(JOIN OUTPUT_POINTERS " ," OUTPUT_POINTERS_STR)
       set(INSTANCE
-          "cudf::jit::transform_kernel<false, false, cudf::jit::type_list<cudf::jit::column_accessor<0, cudf::column_device_view_core, cudf::string_view, false, 0>>, cudf::jit::type_list<${OUTPUT_ACCESSORS_STR}>>"
+          "cudf::jit::transform_kernel<false, false, cudf::jit::type_list<cudf::jit::column_accessor<0, cudf::string_view, false, 0>>, cudf::jit::type_list<${OUTPUT_ACCESSORS_STR}>>"
       )
       add_fragment(
         cudf_fragments

cpp/examples/string_transforms/url_logs/fragments.cu

@@ -3,11 +3,11 @@
  * SPDX-License-Identifier: Apache-2.0
  */
 
+#include <cudf/strings/mutable_string_view.hpp>
 #include <cudf/strings/string_view.cuh>
 #include <cudf/types.hpp>
 
 #include <cuda/std/cstdint>
-#include <cuda/std/span>
 
 struct range32 {
   int32_t begin{};
@@ -179,18 +179,18 @@ __device__ int compute_url_component_sizes(int32_t* protocol_size,
 }
 
 // Copies the six parsed URL components into their preallocated string buffers.
-__device__ int write_url_components(cuda::std::span<char>* protocol,
-                                    cuda::std::span<char>* host,
-                                    cuda::std::span<char>* port,
-                                    cuda::std::span<char>* path,
-                                    cuda::std::span<char>* query,
-                                    cuda::std::span<char>* fragment,
+__device__ int write_url_components(cudf::mutable_string_view* protocol,
+                                    cudf::mutable_string_view* host,
+                                    cudf::mutable_string_view* port,
+                                    cudf::mutable_string_view* path,
+                                    cudf::mutable_string_view* query,
+                                    cudf::mutable_string_view* fragment,
                                     cudf::string_view input)
 {
   url_ranges ranges;
   if (!parse_url(input, &ranges)) { return 0; }
-  cuda::std::span<char>* outputs[] = {protocol, host, port, path, query, fragment};
-  range32 components[]             = {
+  cudf::mutable_string_view* outputs[] = {protocol, host, port, path, query, fragment};
+  range32 components[]                 = {
     ranges.protocol, ranges.host, ranges.port, ranges.path, ranges.query, ranges.fragment};
   for (auto component = 0; component < 6; ++component) {
     auto range = components[component];
@@ -216,12 +216,12 @@ extern "C" __device__ int transform(int32_t* protocol_size,
 #else
 #ifdef UDF_WRITE_OUTPUT
 // Exposes the component-writing pass through the transform LTO ABI.
-extern "C" __device__ int transform(cuda::std::span<char>* protocol,
-                                    cuda::std::span<char>* host,
-                                    cuda::std::span<char>* port,
-                                    cuda::std::span<char>* path,
-                                    cuda::std::span<char>* query,
-                                    cuda::std::span<char>* fragment,
+extern "C" __device__ int transform(cudf::mutable_string_view* protocol,
+                                    cudf::mutable_string_view* host,
+                                    cudf::mutable_string_view* port,
+                                    cudf::mutable_string_view* path,
+                                    cudf::mutable_string_view* query,
+                                    cudf::mutable_string_view* fragment,
                                     cudf::string_view input)
 {
   return write_url_components(protocol, host, port, path, query, fragment, input);

cpp/examples/string_transforms/url_logs/transforms.cpp

@@ -205,17 +205,17 @@ __device__ int compute_url_component_sizes(int32_t* protocol_size,
 // Builds the output UDF from the same parser so both CUDA passes use identical ranges.
 std::string const url_component_output_udf = std::string{R"***(
 // Copies the six parsed URL components into their preallocated string buffers.
-__device__ int write_url_components(cuda::std::span<char>* protocol,
-                                    cuda::std::span<char>* host,
-                                    cuda::std::span<char>* port,
-                                    cuda::std::span<char>* path,
-                                    cuda::std::span<char>* query,
-                                    cuda::std::span<char>* fragment,
+__device__ int write_url_components(cudf::mutable_string_view* protocol,
+                                    cudf::mutable_string_view* host,
+                                    cudf::mutable_string_view* port,
+                                    cudf::mutable_string_view* path,
+                                    cudf::mutable_string_view* query,
+                                    cudf::mutable_string_view* fragment,
                                     cudf::string_view input) {
 )***"} + parse_url_udf + R"***(
   url_ranges ranges;
   if (!parse_url(&ranges)) { return 0; }
-  cuda::std::span<char>* outputs[] = {protocol, host, port, path, query, fragment};
+  cudf::mutable_string_view* outputs[] = {protocol, host, port, path, query, fragment};
   range32 components[]     = {
     ranges.protocol, ranges.host, ranges.port, ranges.path, ranges.query, ranges.fragment};
   for (auto component = 0; component < 6; ++component) {

cpp/include/cudf/column/column_child_offsets.hpp

@@ -14,12 +14,28 @@
 
 namespace CUDF_EXPORT cudf {
 
-static constexpr size_type offsets_column_index = 0;  ///< Child index of the offsets column
+/**
+ * @brief Child index of the offsets column.
+ * @ingroup column_classes
+ */
+constexpr size_type offsets_column_index = 0;
 
 static constexpr size_type dictionary_indices_column_index =
   0;  ///< Child index of the dictionary indices column
 
 static constexpr size_type dictionary_keys_column_index =
   1;  ///< Child index of the dictionary key column
 
+/**
+ * @brief Child index of the list offsets column.
+ * @ingroup lists_classes
+ */
+constexpr size_type lists_offsets_column_index = offsets_column_index;
+
+/**
+ * @brief Child index of the list values column.
+ * @ingroup lists_classes
+ */
+constexpr size_type lists_child_column_index = 1;
+
 }  // namespace CUDF_EXPORT cudf