Implement parallel cuda::std::reduce
#6777
Conversation
| _CCCL_TRY_CUDA_API( | ||
| ::cub::DeviceReduce::Reduce, | ||
| "__pstl_cuda_reduce: cub::DeviceReduce::Reduce failed", | ||
| ::cuda::std::move(__first), | ||
| __device_ret_ptr, | ||
| __count, | ||
| ::cuda::std::move(__func), | ||
| ::cuda::std::move(__init), | ||
| ::cuda::std::move(__policy)); |
There was a problem hiding this comment.
Important: In the for_each_n implementation we create a ::cuda::stream_ref __stream{cudaStreamPerThread}; and pass the stream instead of the policy here. I think we need to add __stream to __policy here.
There was a problem hiding this comment.
Agreed, we should pass a stream. Regarding adding stream to policy, it's a complicated subject that we are deferring until after reduction is merged.
| _CCCL_TRY_CUDA_API( | ||
| ::cudaMalloc, "__pstl_cuda_reduce: allocation failed", reinterpret_cast<void**>(&__device_ret_ptr), sizeof(_Tp)); |
There was a problem hiding this comment.
Important: This must be cudaMallocAsync.
There was a problem hiding this comment.
I don't think cudaMallocAsync is enough. Many calls below throw without ever freeing allocated memory. We need a RAII abstraction, like async_buffer, at least internally.
There was a problem hiding this comment.
How about a unique_ptr?
| _CCCL_TRY_CUDA_API( | ||
| ::cudaMemcpy, | ||
| "__pstl_cuda_reduce: copy of result from device to host failed", | ||
| ::cuda::std::addressof(__ret), | ||
| __device_ret_ptr, | ||
| sizeof(_Tp), | ||
| ::cudaMemcpyDeviceToHost); | ||
|
|
||
| _CCCL_TRY_CUDA_API(::cudaFree, "__pstl_cuda_reduce: deallocate failed", __device_ret_ptr); |
There was a problem hiding this comment.
Important: this should be cudaMemcpyAsync and cudaFreeAsync, followed by a sync of the stream.
🥳 CI Workflow Results🟩 Finished in 3h 12m: Pass: 100%/90 | Total: 3d 05h | Max: 3h 02m | Hits: 49%/212260See results here. |
| _Tp* __device_ret_ptr = nullptr; | ||
|
|
||
| _CCCL_TRY_CUDA_API( | ||
| ::cudaMalloc, "__pstl_cuda_reduce: allocation failed", reinterpret_cast<void**>(&__device_ret_ptr), sizeof(_Tp)); | ||
|
|
There was a problem hiding this comment.
important: this might lead to an issue if user type has a non-trivial constructor. Thrust hanles that by invoking a kernel only when a constructor is needed. We might do a bit better. Consider having a fancy iterator that does in-place new for non-trivial types and a raw pointer otherwise.
There was a problem hiding this comment.
Yep, I think the iterator handling the output should do placement new.
| // Allocate memory for result | ||
| _Tp* __device_ret_ptr = nullptr; | ||
|
|
||
| _CCCL_TRY_CUDA_API( |
There was a problem hiding this comment.
important: given that we throw below, this looks like a memory leak. Consider a RAII abstraction
| thrust::sequence(data.begin(), data.end(), 1); | ||
|
|
||
| const auto policy = cuda::execution::__cub_par_unseq; | ||
| decltype(auto) res = cuda::std::reduce(policy, data.begin(), data.end(), 42, plus_two{}); |
There was a problem hiding this comment.
suggestion: consider adding a check in the binary operator to see if it's actually invoked on GPU
3 participants
This implements
cuda::std::reduceutilizing a CUB backend