NVIDIA / NVIDIA/cccl

[FEA]: Not all copies in Thrust support `par_nosync`: only D->D for non-trivial types can be executed without synchronization.

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

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.