[BUG]: extra unnecessary cudaStreamSynchronize
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 487
- Avg merge
- 2d 7h
- Merged PRs (30d)
- 296
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
`thrust::transform_reduce(thrust::device, ...` flow is the following:
1. DeviceReduce kernel
2. DeviceReduceSingleTile kernel
3. cudaStreamSynchronize
4. `get_value`: cudaMemcpyAsync + cudaStreamSynchronize
The cudaStreamSynchronize in step 3 is unnecessary when `get_value` is going to perform a stream-ordered operation and synchronize.
This was discovered in TeaLeaf and impacts `std::transform_reduce` performance.
### How to Reproduce
Use the transform reduce algorithm on trivial types on device memory or unified memory.
### Expected behavior
No unnecessary cudaStreamSynchronize.
### 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.