[FEA]: Not all copies in Thrust support `par_nosync`: only D->D for non-trivial types can be executed without synchronization.
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
### Is this a duplicate?
- [X] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Type of Bug
Performance
### Component
Thrust
### Describe the bug
For device-to-device copies using `thrust::copy`, non-trivial types are routed through `thrust::transform`, which is implemented using `cuda_cub::parallel_for` and will not synchronized when used with CUDA's `par_nosync` execution policy. This is good.
Trivially relocatable types, however, use `cudaMemcpyAsync` followed by a non-optional synchronization.
We should change the `cuda_cub::synchronize(policy)` call to `cuda_cub::synchronize_optional(policy)`, after ensuring that this won't break any internal copies (e.g. to/from temp_storage) that may rely on non-optional synchronization.
Trivially relocatable H->D and D->H copies also use the same pattern of `cudaMemcpyAsync` + sync, but would take additional refactoring to pass the policy into the implementation, instead of just the stream. This will also require some changes to the implementation of cross-system non-trivial copies, which rely on the trivial copy implementations internally and need to always synchronize (D->H in particular).
Copy currently takes a single execution policy that defines the memory space for both input and output ranges, and this poses an additional barrier to using the CUDA `par_nosync` policy for cross-system copies.
These implementations are located [here](https://github.com/NVIDIA/cccl/blob/edda6228347fb411e199c1de886176e78c341665/thrust/thrust/system/cuda/detail/util.h#L151-L208).
### How to Reproduce
Use `thrust::copy(thrust::cuda::par_nosync, ...)` for anything but a device-to-device copy of non-trivially-relocatable types. The call will synchronize.
### Expected behavior
The call should not synchronize.
### Reproduction link
_No response_
### Operating System
_No response_
### nvidia-smi output
_No response_
### NVCC version
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.