Replace direct cudaMemcpyAsync calls with utility functions (limited to cudf::io)
#17132
Conversation
…fea-remove-cudamemcpy-io
…fea-remove-cudamemcpy-io
vuule
added
Performance
| @@ -218,7 +218,7 @@ void generate_depth_remappings( | |||
| */ | |||
| [[nodiscard]] std::future<void> read_column_chunks_async( | |||
| std::vector<std::unique_ptr<datasource>> const& sources, | |||
| std::vector<std::unique_ptr<datasource::buffer>>& page_data, | |||
| cudf::host_span<rmm::device_buffer> page_data, | |||
There was a problem hiding this comment.
simplified outdated complexity
vuule
changed the title
cudaMemcpyAsync calls with utility functions
vuule
changed the title
cudaMemcpyAsync calls with utility functionscudaMemcpyAsync calls with utility functions (limited to cudf::io)
| std::pair(source_ptr->device_read_async( | ||
| read_info.offset, read_info.length, dst_base + read_info.dst_pos, _stream), | ||
| read_info.length)); | ||
| device_read_tasks.emplace_back( |
There was a problem hiding this comment.
Unrelated change; noticed clang-tidy complaining that we used to make an unnecessary move here :)
|
Do you have (run) any benchmark to make sure there is no regression? |
I haven't because we currently don't do anything differently - we end up calling |
| @@ -87,8 +87,10 @@ class datasource_chunk_reader : public data_chunk_reader { | |||
| _source->host_read(_offset, read_size, reinterpret_cast<uint8_t*>(h_ticket.buffer.data())); | |||
|
|
|||
| // copy the host-pinned data on to device | |||
| CUDF_CUDA_TRY(cudaMemcpyAsync( | |||
| chunk.data(), h_ticket.buffer.data(), read_size, cudaMemcpyDefault, stream.value())); | |||
| cudf::detail::cuda_memcpy_async<char>( | |||
There was a problem hiding this comment.
There are a number of places where the template argument (char) is given explicitly ... is the compiler really not able to deduce it from the inputs?
There was a problem hiding this comment.
hm, maybe it wasn't required here
In general, compiler can't do template type deduction + implicit conversion. So passing a container that get implicitly converted to a span requires the template type for cuda_memcpy_async.
|
/merge |
Description
Issue #15620
Replaced the calls to
cudaMemcpyAsyncwith the newcuda_memcpy/cuda_memcpy_asyncutility, which optionally avoids using the copy engine. Changes are limited to cuIO to make the PR easier to review (repetitive enough as-is!).Also took the opportunity to use
cudf::detail::host_vectorand its factories to enable wider pinned memory use.Skipped a few instances of
cudaMemcpyAsync; few are underio::comp, which we don't want to invest in further (if possible). The othercudaMemcpyAsyncinstances are D2D copies, whichcuda_memcpy/cuda_memcpy_asyncdon't support. Perhaps they should, just to make the use ubiquitous.Checklist