[FEA] Support bool based on bit, Timestamp more range and StringView

[jinchengchenghh]

In Velox and Cudf integration, the bool type conversion is expensive, because the source Velox bool type is one bit for one value, but in cudf, it is one byte for one value, it is better to have a bool type align with Velox, then we only needs to copy a big buffer(data buffer and null buffer) instead of setting the value one by one, so as Timestamp.

With this approach, we can do the conversion in any time considering the memory/runtime statistic/operator fallback. E.g. The final round join is under memory pressure we cannot process in GPU without spill, use multiple threads CPU to do the following join instead of spilling to disk maybe a good trade off.

System goal	Timestamp type	Example systems / libraries	Notes
Maximum time range	int128	Some database engines, experimental extensions in Apache Arrow, specialized scientific time libraries	Used when timelines may span thousands or millions of years, or when avoiding overflow is critical
High-performance analytics	int64	RAPIDS cuDF, Pandas, Apache Arrow	Fast arithmetic, smaller memory footprint, and native support on CPUs/GPUs; covers ~1677–2262 for nanoseconds

For StringView, the arrow newest version also supports this format, inline the string with size <= 12 and store them in several continuous buffers, this will allow to not allocate buffer in one time, concat vectors will also benefit from it, as I know, copy from GPU to GPU is more expensive than CPU to GPU, please correct me if I'm wrong.

[GregoryKimball]

Thank you @jinchengchenghh for sharing these requests. Generally we avoid adding new data types to cuDF because it brings large compile time and performance penalties. However, we want to make interop from Velox and Arrow types as fast as possible with cuDF, and we aim to have GPU converters for all Velox data types.

• bit representation of boolean data. Arrow booleans are 1-bit so cuDF handles the convert to 1-byte booleans in to_arrow_* and from_arrow_*. If the CPU is doing the BOOL to BOOL8 conversion that would be expensive, but on the GPU this should be doable at memory bandwidth. Would you please share more about how you've seen the bool type conversion to be expensive?

• 128-bit timestamps. cuDF doesn't support this. My first approach would be to develop Velox operators that model 128-bit timestamp results using existing timestamp and integer types in cuDF. We could explore split word with two 64-bit int or using a seconds 64-bit timestamp with a nanosecond counter. We do have a 128-bit decimal type so perhaps a 128-bit timestamp type could make sense, but it would be a big project.

• stringview. cuDF has efficient interop for string view, so we can accept and emit string view data as needed by Velox. We have a benchmark suite that shows String representation to have overall better performance than String View in GPU execution.

[mhaseeb123]

You should be quickly able to convert your bit boolean data into byte boolean representation using thrust. Here's a snippet:

#include <cudf/types.hpp>

#include <thrust/for_each.h>
#include <cuda/iterator>

#include <rmm/exec_policy.hpp>

auto convert_boolean_representation(cudf::device_span<cuda::std::byte> bitmap, cudf::size_type num_rows, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr)
{
  auto output_buffer = rmm::device_buffer(num_rows * sizeof(bool), stream, mr);
  thrust::for_each(rmm::exec_policy_nosync(stream),
                   cuda::counting_iterator(0),
                   cuda::counting_iterator(num_rows),
                   [input = bitmap.data(), 
                    output = static_cast<bool*>(output_buffer.data())] 
                   __device__ (auto row_idx) {
                     constexpr auto bits_in_word = sizeof(cuda::std::byte) * 8;
                     auto const word_idx         = row_idx / bits_in_word;
                     auto const intra_word_idx   = row_idx % bits_in_word;
                     output[row_idx]             = cuda::std::to_integer<bool>(input[word_idx] & (cuda::std::byte{1} << intra_word_idx));
  });

  return std::make_unique<cudf::column>(
      cudf::data_type{cudf::type_id::BOOL8},
      static_cast<cudf::size_type>(num_rows),
      std::move(output_buffer),
      rmm::device_buffer{0, stream, mr},
      0);
}

[davidwendt]

Instead of thrust, you could use the public API cudf::mask_to_bools to convert a bitmask to a BOOL8 column:
https://docs.rapids.ai/api/libcudf/stable/group__transformation__transform.html#ga4ebc1bfb2a0117cd816890a694bcb136

[davidwendt]

This PR contains some results from the benchmarks mentioned by Greg concerning the ArrowStringView.

[jinchengchenghh]

StringView has benefit in expression evaluation, for example, substr(trim(string)), if we use String, we need to copy 2 times, if we use StringView, we don't need to copy them, the performance difference would make more sense for large string.

[jinchengchenghh]

For other type, we only need to copy the bulk buffer, but for bool type, we need to extract the bit one by one, and then set it one by one, so it is more expensive than other data types, as here https://github.com/facebookincubator/velox/pull/16760/changes#diff-b58f130ead38ac7ee5c9c2ade0bd797028aa1915d539024c1cabfbb290792876R656

[davidwendt]

StringView has benefit in expression evaluator, for example, substr(trim(string)), if we use String, we need to copy 2 times, if we use StringView, we don't need to copy them, the performance difference would make more sense for large string.

This will not work for libcudf since all strings are immutable and data is not shared so the output is always a copy regardless of the format. This is a fundamental part of the library and so would likely require an extensive amount of work to rebuild the library around sharing parts of common data with multiple columns.

I understand your substr example but there will always be examples where one format is slightly better than another in given situations. Please see the benchmark results mentioned in the above comment.

For these reasons, I would prefer to start with making the interop as fast as possible.