[REVIEW] impl and refactor device_scalar::set_value and value, respectively. #167
Conversation
…(T)` to `void set_value(T)`, adjust docs
cwharris
changed the title
|
This change breaks cudf, so I've creating a sister pr here: rapidsai/cudf#3211 |
|
I added async versions of |
cwharris
changed the title
include/rmm/device_scalar.hpp
Outdated
| @@ -92,6 +92,14 @@ class device_scalar { | |||
|
|
|||
| private: | |||
| rmm::device_buffer buff{sizeof(T)}; | |||
|
|
|||
| inline void _memcpy(void *dst, const void *src, cudaStream_t stream) const { | |||
| auto status = cudaMemcpyAsync(dst, src, sizeof(T), cudaMemcpyDefault, | |||
There was a problem hiding this comment.
Please either use cudaMemcpy or add a cudaStreamSynchronize() after this. Otherwise it appears to be asynchronous and is therefore ambiguous with the documentation.
I know I said that cudaMemcpyAsync is not async when either src or dst is un-pinned host memory, but even we should make the synchronization explicit.
There was a problem hiding this comment.
We cannot specify the stream with cudaMemcpy, but Jake mentioned specifying the stream may still be beneficial. I'm not sure how to proceed.
There was a problem hiding this comment.
Right, it's still better to use cudaMemcpyAsync because it let's us sync a stream that's NOT the null stream (synchronizing the NULL stream is the same as cudaDeviceSynchronize(), which is overkill). To proceed, just add in cudaStreamSynchronize(stream) after the cudaMemcpyAsync like you had before.
Addresses #166
Depends on #169