NVIDIA / NVIDIA/cccl

[BUG]: extra unnecessary cudaStreamSynchronize

Open
#261 0 comments 0 reactions 0 assignees View on GitHub
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.