[Review] Add stream support for static_map execution #113
Conversation
include/cuco/detail/static_map.inl
Outdated
| slots_ = std::allocator_traits<slot_allocator_type>::allocate(slot_allocator_, capacity_, | ||
| exec_stream_); | ||
| num_successes_ = std::allocator_traits<counter_allocator_type>::allocate(counter_allocator_, 1, | ||
| exec_stream_); |
There was a problem hiding this comment.
| slots_ = std::allocator_traits<slot_allocator_type>::allocate(slot_allocator_, capacity_, | |
| exec_stream_); | |
| num_successes_ = std::allocator_traits<counter_allocator_type>::allocate(counter_allocator_, 1, | |
| exec_stream_); | |
| slots_ = slot_allocator_.allocate(capacity_, exec_stream_); | |
| num_successes_ = counter_allocator_.allocate(1, exec_stream_); |
I could be wrong but passing CUDA stream as a hint to std::allocator_traits::allocate (see here) may not be a proper use case.
There was a problem hiding this comment.
I assumed it can given this function taking stream argument. Please let me know what is the correct use here.
There was a problem hiding this comment.
Passing the stream to allocator_traits is not going to work. You need to call allocate directly like Yunsong showed.
There was a problem hiding this comment.
It will indeed invoke cuco::allocator::allocate but the stream argument won't be passed properly in this way. It doesn't matter much with cuco::allocator since it takes stream but does not really use it. A real case would be using e.g. rmm allocator allocate. My suggestions will work just fine.
include/cuco/detail/static_map.inl
Outdated
| std::allocator_traits<slot_allocator_type>::deallocate(slot_allocator_, slots_, capacity_); | ||
| std::allocator_traits<counter_allocator_type>::deallocate(counter_allocator_, num_successes_, 1); |
There was a problem hiding this comment.
| std::allocator_traits<slot_allocator_type>::deallocate(slot_allocator_, slots_, capacity_); | |
| std::allocator_traits<counter_allocator_type>::deallocate(counter_allocator_, num_successes_, 1); | |
| counter_allocator_.deallocate(num_successes_, 1, exec_stream_); | |
| slot_allocator_.deallocate(slots_, capacity_, exec_stream_); |
include/cuco/detail/static_map.inl
Outdated
| stream)); | ||
| // stream'd execution assumes sync not required | ||
| if (stream == NULL) | ||
| CUCO_CUDA_TRY(cudaDeviceSynchronize()); // ensures legacy behavior |
There was a problem hiding this comment.
| CUCO_CUDA_TRY(cudaDeviceSynchronize()); // ensures legacy behavior | |
| CUCO_CUDA_TRY(cudaStreamSynchronize(stream)); // ensures legacy behavior |
We should replace all device sync with stream sync
There was a problem hiding this comment.
I am fine with replacing deviceSync with streamSync, however this change may reveal unintentional data race in applications that unknowingly benefitted from the device level synchronization provided by cudaDeviceSynchronize()
There was a problem hiding this comment.
I don't think we need to worry about people inadvertently relying on the device sync here. They should not have been doing that in the first place.
|
Can you please update the docs also since we are adding a new argument to bulk functions? |
chirayuG-nvidia
changed the title
|
Added unit test and addressed review comments. |
| @@ -42,6 +44,7 @@ int main(void) | |||
|
|
|||
| // Inserts all pairs into the map | |||
| map.insert(pairs.begin(), pairs.end()); | |||
| cudaStreamSynchronize(str); | |||
There was a problem hiding this comment.
this sync isn't really necessary
PointKernel
added
type: feature request
cuCollection added stream support for static_map in NVIDIA/cuCollections#113 This allows us to avoid unnecessary stream/device synchronizations and also allows to enable multi-stream execution without unnecessary serialization. This PR updates libcugraph code to pass `handle.get_stream()` to `cuco::static_map`. Authors: - Seunghwa Kang (https://github.com/seunghwak) Approvers: - Chuck Hastings (https://github.com/ChuckHastings) - Kumar Aatish (https://github.com/kaatish) URL: #1984
Adding plumbing for execution static_map functions on user supplied cuda stream.