[FEA]: Thrust algorithms invoked with `par_nosync` should asynchronously allocate temporary storage
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
```
cudax::stream stream;
auto policy = thrust::cuda::par_nosync.on(stream.get());
thrust::device_vector A(...);
auto it = thrust::inclusive_scan(policy, A.begin(), A.end(), A.begin());
// Should allocate temporary storage asynchronously and not block.
```
Proposed semantics:
- If an algorithm that allocates temporary storage is invoked with a `par_nosync` policy and the allocator attached to the policy supports asynchronous allocation, the temporary storage should be allocated asynchronously.
- The default allocator for `par_nosync` should be an allocator that supports asynchronous allocation, so that you get asynchronous allocation with `par_nosync` by default.
Why: There's a lot of existing Thrust code out there today that uses algorithms that allocate temporary storage but don't have a compute-dependent result (like `inclusive_scan` and `sort`). This existing code could benefit from `par_nosync`, but doesn't today because the synchronous free of the temporary storage will block.
How:
- The Scary Way: Temporary allocations always go through `temporary_array` (which is implemented with `contiguous_buffer`), and uses Thrust allocator traits. I believe `temporary_array` takes a policy on construction, not an allocator, so we have the information we need. We could extend allocator traits to support async allocation.
- The Easy Way: We could just check for the async allocator/MR concepts and then use more modern libcu++ interfaces directly in `temporary_array`/`contiguous_buffer`.
- I believe at some point in Thrust's history we had a `get_temporary_buffer`/`return_temporary_buffer` interface that was intended to be customizable by users. I don't think it's still around, but if it is, that may be worth exploring. I recall it being incredibly hazardous though.
Concerns and considerations:
- Pooling of allocations.
- Exploding memory footprints.
- Touching Thrust temporary_array/contiguous_buffer
- Touching Thrust allocator traits is terrifying. Be careful
Contributor guide
Assessment
This issue has not been assessed yet.