[FEA] libcudf should not rely on the ABI stability of Thrust types

[jrhemstad]

Is your feature request related to a problem? Please describe.

libcudf's public API currently includes use of Thrust types like thrust::device_vector. This means libcudf implicitly depends on the ABI of those Thrust types to be stable not only across Thrust versions, but even within the same version (e.g., host TUs vs device TUs).

This is problematic because Thrust and CUB symbols make no ABI stability guarantees.

Describe the solution you'd like
Thrust types in libcudf's public API should be replaced with types that have ABI guarantees. This likely means replacing with equivalent cuda::std:: types.

Additional context

This is related to #14734 in a roundabout way that isn't worth going into the details.

See also:
NVIDIA/cccl#1246
NVIDIA/cccl#1262

[robertmaynard]

public symbols that use thrust types:

cudf::dictionary::detail::merge(cudf::dictionary_column_view const&, cudf::dictionary_column_view const&, rmm::device_uvector<thrust::pair<cudf::detail::side, int> > const&, rmm::cuda_stream_view, rmm::mr::device_memory_resource*)

cudf::make_strings_column(cudf::device_span<thrust::pair<char const*, int> const, 18446744073709551615ul>, rmm::cuda_stream_view, rmm::mr::device_memory_resource*)

cudf::io::avro::gpu::DecodeAvroColumnData(cudf::device_span<cudf::io::avro::block_desc_s const, 18446744073709551615ul>, cudf::io::avro::gpu::schemadesc_s*, cudf::device_span<thrust::pair<char const*, int> const, 18446744073709551615ul>, unsigned char const*, unsigned int, unsigned int, rmm::cuda_stream_view)

cudf::io::json::detail::parse_data(char const*, thrust::zip_iterator<thrust::tuple<int const*, int const*, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, int, cudf::data_type, rmm::device_buffer&&, int, cudf::io::parse_options_view const&, rmm::cuda_stream_view, rmm::mr::device_memory_resource*)

cudf::io::json::detail::legacy::get_data_types(cudf::io::json_reader_options const&, cudf::io::parse_options_view const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > > const&, concurrent_unordered_map<unsigned int, int, cudf::hashing::detail::MurmurHash3_x86_32<unsigned int>, equal_to<unsigned int>, default_allocator<thrust::pair<unsigned int, int> > >*, cudf::device_span<unsigned long const, 18446744073709551615ul>, cudf::device_span<char const, 18446744073709551615ul>, rmm::cuda_stream_view)

cudf::io::json::detail::legacy::collect_keys_info(cudf::io::parse_options_view const&, cudf::device_span<char const, 18446744073709551615ul>, cudf::device_span<unsigned long const, 18446744073709551615ul>, unsigned long long*, thrust::optional<cudf::mutable_table_device_view>, rmm::cuda_stream_view)

cudf::io::json::detail::legacy::detect_data_types(cudf::io::parse_options_view const&, cudf::device_span<char const, 18446744073709551615ul>, cudf::device_span<unsigned long const, 18446744073709551615ul>, bool, int, concurrent_unordered_map<unsigned int, int, cudf::hashing::detail::MurmurHash3_x86_32<unsigned int>, equal_to<unsigned int>, default_allocator<thrust::pair<unsigned int, int> > >*, rmm::cuda_stream_view)

cudf::io::json::detail::legacy::convert_data_to_table(cudf::io::parse_options_view const&, std::vector<cudf::data_type, std::allocator<cudf::data_type> > const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > >&&, concurrent_unordered_map<unsigned int, int, cudf::hashing::detail::MurmurHash3_x86_32<unsigned int>, equal_to<unsigned int>, default_allocator<thrust::pair<unsigned int, int> > >*, cudf::device_span<unsigned long const, 18446744073709551615ul>, cudf::device_span<char const, 18446744073709551615ul>, rmm::cuda_stream_view, rmm::mr::device_memory_resource*)

cudf::io::json::detail::legacy::convert_json_to_columns(cudf::io::parse_options_view const&, cudf::device_span<char const, 18446744073709551615ul>, cudf::device_span<unsigned long const, 18446744073709551615ul>, cudf::device_span<cudf::data_type const, 18446744073709551615ul>, concurrent_unordered_map<unsigned int, int, cudf::hashing::detail::MurmurHash3_x86_32<unsigned int>, equal_to<unsigned int>, default_allocator<thrust::pair<unsigned int, int> > >*, cudf::device_span<void* const, 18446744073709551615ul>, cudf::device_span<unsigned int* const, 18446744073709551615ul>, cudf::device_span<int, 18446744073709551615ul>, rmm::cuda_stream_view)

cudf::io::detail::infer_data_type(cudf::io::json_inference_options_view const&, cudf::device_span<char const, 18446744073709551615ul>, thrust::zip_iterator<thrust::tuple<int const*, int const*, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> >, unsigned long, rmm::cuda_stream_view)

cudf::io::detail::avro::decode_data(cudf::io::detail::avro::metadata&, rmm::device_buffer const&, std::vector<std::pair<unsigned int, unsigned int>, std::allocator<std::pair<unsigned int, unsigned int> > > const&, cudf::device_span<thrust::pair<char const*, int> const, 18446744073709551615ul>, unsigned long, std::vector<std::pair<int, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > >, std::allocator<std::pair<int, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > > > const&, std::vector<cudf::data_type, std::allocator<cudf::data_type> > const&, rmm::cuda_stream_view, rmm::mr::device_memory_resource*)

[robertmaynard]

That is an incomplete list of symbols since the restriction of no ABI support between C++ and CUDA also means we can't pass internally Thrust types between C++ and CUDA sources. This means that components within our detail:: namespace also have to not have thrust types if we expect internal uses of that function across C++ and CUDA sources ( which we should )