Add Java bindings for string literal support in AST

cpp/include/cudf/ast/detail/expression_evaluator.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -21,7 +21,6 @@
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/utilities/assert.cuh>
-#include <cudf/scalar/scalar_device_view.cuh>
 #include <cudf/table/table_device_view.cuh>
 #include <cudf/table/table_view.hpp>
 #include <cudf/types.hpp>

cpp/include/cudf/ast/detail/expression_parser.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -104,7 +104,7 @@ using IntermediateDataType = possibly_null_value_t<std::int64_t, has_nulls>;
  */
 struct expression_device_view {
   device_span<detail::device_data_reference const> data_references;
-  device_span<cudf::detail::fixed_width_scalar_device_view_base const> literals;
+  device_span<generic_scalar_device_view const> literals;
   device_span<ast_operator const> operators;
   device_span<cudf::size_type const> operator_source_indices;
   cudf::size_type num_intermediates;
@@ -281,11 +281,10 @@ class expression_parser {
       reinterpret_cast<detail::device_data_reference const*>(device_data_buffer_ptr +
                                                              buffer_offsets[0]),
       _data_references.size());
-    device_expression_data.literals =
-      device_span<cudf::detail::fixed_width_scalar_device_view_base const>(
-        reinterpret_cast<cudf::detail::fixed_width_scalar_device_view_base const*>(
-          device_data_buffer_ptr + buffer_offsets[1]),
-        _literals.size());
+    device_expression_data.literals = device_span<generic_scalar_device_view const>(
+      reinterpret_cast<generic_scalar_device_view const*>(device_data_buffer_ptr +
+                                                          buffer_offsets[1]),
+      _literals.size());
     device_expression_data.operators = device_span<ast_operator const>(
       reinterpret_cast<ast_operator const*>(device_data_buffer_ptr + buffer_offsets[2]),
       _operators.size());
@@ -335,7 +334,7 @@ class expression_parser {
   std::vector<detail::device_data_reference> _data_references;
   std::vector<ast_operator> _operators;
   std::vector<cudf::size_type> _operator_source_indices;
-  std::vector<cudf::detail::fixed_width_scalar_device_view_base> _literals;
+  std::vector<generic_scalar_device_view> _literals;
 };
 
 }  // namespace detail

cpp/include/cudf/ast/expressions.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -150,6 +150,96 @@ enum class table_reference {
   OUTPUT  ///< Column index in the output table
 };
 
+/**
+ * @brief A type-erased scalar_device_view where the value is a fixed width type or a string
+ */
+class generic_scalar_device_view : public cudf::detail::scalar_device_view_base {
+ public:
+  /**
+   * @brief Returns the stored value.
+   *
+   * @tparam T The desired type
+   * @returns The stored value
+   */
+  template <typename T>
+  __device__ T const value() const noexcept
+  {
+    if constexpr (std::is_same_v<T, cudf::string_view>) {
+      return string_view(static_cast<char const*>(_data), _size);
+    }
+    return *static_cast<T const*>(_data);
+  }
+
+  /** @brief Construct a new generic scalar device view object from a numeric scalar
+   *
+   * @param s The numeric scalar to construct from
+   */
+  template <typename T>
+  generic_scalar_device_view(numeric_scalar<T>& s)
+    : generic_scalar_device_view(s.type(), s.data(), s.validity_data())
+  {
+  }
+
+  /** @brief Construct a new generic scalar device view object from a timestamp scalar
+   *
+   * @param s The timestamp scalar to construct from
+   */
+  template <typename T>
+  generic_scalar_device_view(timestamp_scalar<T>& s)
+    : generic_scalar_device_view(s.type(), s.data(), s.validity_data())
+  {
+  }
+
+  /** @brief Construct a new generic scalar device view object from a duration scalar
+   *
+   * @param s The duration scalar to construct from
+   */
+  template <typename T>
+  generic_scalar_device_view(duration_scalar<T>& s)
+    : generic_scalar_device_view(s.type(), s.data(), s.validity_data())
+  {
+  }
+
+  /** @brief Construct a new generic scalar device view object from a string scalar
+   *
+   * @param s The string scalar to construct from
+   */
+  generic_scalar_device_view(string_scalar& s)
+    : generic_scalar_device_view(s.type(), s.data(), s.validity_data(), s.size())
+  {
+  }
+
+ protected:
+  void const* _data{};      ///< Pointer to device memory containing the value
+  size_type const _size{};  ///< Size of the string in bytes for string scalar
+
+  /**
+   * @brief Construct a new fixed width scalar device view object
+   *
+   * @param type The data type of the value
+   * @param data The pointer to the data in device memory
+   * @param is_valid The pointer to the bool in device memory that indicates the
+   * validity of the stored value
+   */
+  generic_scalar_device_view(data_type type, void const* data, bool* is_valid)
+    : cudf::detail::scalar_device_view_base(type, is_valid), _data(data)
+  {
+  }
+
+  /** @brief Construct a new string scalar device view object
+   *
+   * @param type The data type of the value
+   * @param data The pointer to the data in device memory
+   * @param is_valid The pointer to the bool in device memory that indicates the
+   * validity of the stored value
+   * @param size The size of the string in bytes
+   */
+  generic_scalar_device_view(data_type type, void const* data, bool* is_valid, size_type size)
+    : cudf::detail::scalar_device_view_base(type, is_valid), _data(data), _size(size)
+  {
+  }
+};
+
 /**
  * @brief A literal value used in an abstract syntax tree.
  */
@@ -162,8 +252,7 @@ class literal : public expression {
    * @param value A numeric scalar value
    */
   template <typename T>
-  literal(cudf::numeric_scalar<T>& value)
-    : scalar(value), value(cudf::get_scalar_device_view(value))
+  literal(cudf::numeric_scalar<T>& value) : scalar(value), value(value)
   {
   }
 
@@ -174,8 +263,7 @@ class literal : public expression {
    * @param value A timestamp scalar value
    */
   template <typename T>
-  literal(cudf::timestamp_scalar<T>& value)
-    : scalar(value), value(cudf::get_scalar_device_view(value))
+  literal(cudf::timestamp_scalar<T>& value) : scalar(value), value(value)
   {
   }
 
@@ -186,11 +274,17 @@ class literal : public expression {
    * @param value A duration scalar value
    */
   template <typename T>
-  literal(cudf::duration_scalar<T>& value)
-    : scalar(value), value(cudf::get_scalar_device_view(value))
+  literal(cudf::duration_scalar<T>& value) : scalar(value), value(value)
   {
   }
 
+  /**
+   * @brief Construct a new literal object.
+   *
+   * @param value A string scalar value
+   */
+  literal(cudf::string_scalar& value) : scalar(value), value(value) {}
+
   /**
    * @brief Get the data type.
    *
@@ -203,10 +297,7 @@ class literal : public expression {
    *
    * @return The device scalar object
    */
-  [[nodiscard]] cudf::detail::fixed_width_scalar_device_view_base get_value() const
-  {
-    return value;
-  }
+  [[nodiscard]] generic_scalar_device_view get_value() const { return value; }
 
   /**
    * @brief Accepts a visitor class.
@@ -236,7 +327,7 @@ class literal : public expression {
 
  private:
   cudf::scalar const& scalar;
-  cudf::detail::fixed_width_scalar_device_view_base const value;
+  generic_scalar_device_view const value;
 };
 
 /**

cpp/tests/ast/transform_tests.cpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -397,6 +397,40 @@ TEST_F(TransformTest, StringComparison)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view(), verbosity);
 }
 
+TEST_F(TransformTest, StringScalarComparison)
+{
+  auto c_0   = cudf::test::strings_column_wrapper({"1", "12", "123", "23"});
+  auto table = cudf::table_view{{c_0}};
+
+  auto literal_value = cudf::string_scalar("2");
+  auto literal       = cudf::ast::literal(literal_value);
+
+  auto col_ref_0  = cudf::ast::column_reference(0);
+  auto expression = cudf::ast::operation(cudf::ast::ast_operator::LESS, col_ref_0, literal);
+
+  auto expected = column_wrapper<bool>{true, true, true, false};
+  auto result   = cudf::compute_column(table, expression);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view(), verbosity);
+}
+
+TEST_F(TransformTest, NumericScalarComparison)
+{
+  auto c_0   = column_wrapper<int32_t>{1, 12, 123, 23};
+  auto table = cudf::table_view{{c_0}};
+
+  auto literal_value = cudf::numeric_scalar<int32_t>(2);
+  auto literal       = cudf::ast::literal(literal_value);
+
+  auto col_ref_0  = cudf::ast::column_reference(0);
+  auto expression = cudf::ast::operation(cudf::ast::ast_operator::LESS, col_ref_0, literal);
+
+  auto expected = column_wrapper<bool>{true, false, false, false};
+  auto result   = cudf::compute_column(table, expression);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view(), verbosity);
+}
+
 TEST_F(TransformTest, CopyColumn)
 {
   auto c_0   = column_wrapper<int32_t>{3, 0, 1, 50};

java/src/main/java/ai/rapids/cudf/ast/Literal.java

@@ -200,6 +200,15 @@ public static Literal ofDurationFromLong(DType type, Long value) {
     return ofDurationFromLong(type, value.longValue());
   }
 
+  /** Construct a string literal with the specified value or null. */
+  public static Literal ofString(String value) {
+    byte[] stringBytes = value.getBytes();
+    byte[] serializedValue = new byte[stringBytes.length + Integer.BYTES];
+    ByteBuffer.wrap(serializedValue).order(ByteOrder.nativeOrder()).putInt(stringBytes.length);
+    System.arraycopy(stringBytes, 0, serializedValue, Integer.BYTES, stringBytes.length);
+    return new Literal(DType.STRING, serializedValue);
+  }
+
   Literal(DType type, byte[] serializedValue) {
     this.type = type;
     this.serializedValue = serializedValue;

java/src/main/native/src/CompiledExpression.cpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -15,6 +15,7 @@
  */
 
 #include <cstdint>
+#include <memory>
 #include <stdexcept>
 #include <vector>
 
@@ -56,12 +57,20 @@ class jni_serialized_ast {
 
   /** Read a multi-byte value from the serialized AST data buffer */
   template <typename T> T read() {
-    check_for_eof(sizeof(T));
-    // use memcpy since data may be misaligned
-    T result;
-    memcpy(reinterpret_cast<jbyte *>(&result), data_ptr, sizeof(T));
-    data_ptr += sizeof(T);
-    return result;
+    if constexpr (std::is_same_v<T, std::string>) {
+      auto const size = read<cudf::size_type>();
+      check_for_eof(size);
+      auto const result = std::string(reinterpret_cast<char const *>(data_ptr), size);
+      data_ptr += size;
+      return result;
+    } else {
+      check_for_eof(sizeof(T));
+      // use memcpy since data may be misaligned
+      T result;
+      memcpy(reinterpret_cast<jbyte *>(&result), data_ptr, sizeof(T));
+      data_ptr += sizeof(T);
+      return result;
+    }
   }
 
   /** Decode a libcudf data type from the serialized AST data buffer */
@@ -254,9 +263,29 @@ struct make_literal {
                                      std::move(scalar_ptr));
   }
 
+  /** Construct an AST literal from a string value */
+  template <typename T, std::enable_if_t<std::is_same_v<T, cudf::string_view>> * = nullptr>
+  cudf::ast::literal &operator()(cudf::data_type dtype, bool is_valid,
+                                 cudf::jni::ast::compiled_expr &compiled_expr,
+                                 jni_serialized_ast &jni_ast) {
+    std::unique_ptr<cudf::scalar> scalar_ptr = [&]() {
+      if (is_valid) {
+        std::string val = jni_ast.read<std::string>();
+        return std::make_unique<cudf::string_scalar>(val, is_valid);
+      } else {
+        return std::make_unique<cudf::string_scalar>(rmm::device_buffer{}, is_valid);
+      }
+    }();
+
+    auto &str_scalar = static_cast<cudf::string_scalar &>(*scalar_ptr);
+    return compiled_expr.add_literal(std::make_unique<cudf::ast::literal>(str_scalar),
+                                     std::move(scalar_ptr));
+  }
+
   /** Default functor implementation to catch type dispatch errors */
   template <typename T, std::enable_if_t<!cudf::is_numeric<T>() && !cudf::is_timestamp<T>() &&
-                                         !cudf::is_duration<T>()> * = nullptr>
+                                         !cudf::is_duration<T>() &&
+                                         !std::is_same_v<T, cudf::string_view>> * = nullptr>
   cudf::ast::literal &operator()(cudf::data_type dtype, bool is_valid,
                                  cudf::jni::ast::compiled_expr &compiled_expr,
                                  jni_serialized_ast &jni_ast) {