Add thrust::cuda::par_nosync execution policy

[jedbrown]

Thrust-1.9.4 made the breaking API change:

Synchronous Thrust algorithms now block until all of their operations have completed. Use the new asynchronous Thrust algorithms for non-blocking behavior.

The rationale seems to be (14f8a54):

* All Thrust synchronous algorithms for the CUDA backend now actually
  synchronize. Previously, any algorithm that did not allocate temporary
  storage (counterexample: `thrust::sort`) and did not have a
  computation-dependent result (counterexample: `thrust::reduce`) would actually
  be launched asynchronously.  Additionally, synchronous algorithms that
  allocated temporary storage would become asynchronous if a custom allocator
  was supplied that did not synchronize on allocation/deallocation, unlike
  `cudaMalloc`/`cudaFree`. So, now `thrust::for_each`, `thrust::transform`,
  `thrust::sort`, etc are truly synchronous. In some cases this may be a
  performance regression; if you need asynchrony, use the new asynchronous
  algorithms.

This is disruptive for libraries that wish to use Thrust internally without exposing its objects to the caller. The new thrust::async interfaces require holding the futures to be waited on eventually. In iterative linear algebra, this might be after an "essential" [^1] synchronization point, like a dot product reporting its result on the host. But the different operations can be implemented in different libraries (which may or may not call Thrust), so using async seems to imply that internal use of thrust::async requires all transitive callers to export Thrust futures or a suitable wrapper.

So my questions:

1. Is this transitive disruption avoidable? Is there a way to disown the futures without blocking?
2. What is your recommended migration path for legacy software that uses thrust::transform (for example) and desires stream-based nonblocking semantics?

[^1] I use scare quotes because it's possible to deliver such results on-device and thus avoid an extra round-trip latency.

[alliepiper]

There is an ongoing effort to expose the Thrust CUDA kernels as CUB device functions (https://github.com/NVIDIA/cub). These have the non-blocking, stream-ordered semantics that you're looking for.

Unfortunately, we don't have a suitable alternative for transform available in CUB yet, so there isn't a good option other than thrust::transform or thrust::async::transform at the moment. If you require this algorithm without blocking or using futures, a custom transform kernel would be needed. This is a large gap in our API that I'm hoping to have fixed soon.

I should point out that your usecase is a common request, and we're planning to provide better support for stream-ordered non-blocking "fire-and-forget"-type algorithms in future iterations of the Thrust and CUB APIs.

[jedbrown]

Thanks for your helpful reply. Should I file a cub::DeviceTransform issue so there's something to subscribe to?

BTW, I notice that the CUB website (linked from its GitHub) hasn't learned about any releases since early 2018.

[alliepiper]

Thanks for your helpful reply. Should I file a cub::DeviceTransform issue so there's something to subscribe to?

Sure, feel free to open an issue for this.

BTW, I notice that the CUB website (linked from its GitHub) hasn't learned about any releases since early 2018.

The Thrust and CUB docs are a mess currently. There's an ongoing effort to clean them up in #1475.

[brycelelbach]

We could add a `par_nosync` policy for folks who want the old behavior.

…

On Wed, Oct 6, 2021 at 4:38 PM Allison Vacanti ***@***.***> wrote: Thanks for your helpful reply. Should I file a cub::DeviceTransform issue so there's something to subscribe to? Sure, feel free to open an issue for this. BTW, I notice that the CUB website (linked from its GitHub) hasn't learned about any releases since early 2018. The Thrust and CUB docs are a mess currently. There's an ongoing effort to clean them up in #1475 <#1475>. — You are receiving this because you are subscribed to this thread. Reply to this email directly, view it on GitHub <#1515 (comment)>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/AADBG4VMJHXE4EFMPCGWRQDUFSXTHANCNFSM5DCVD37Q> .

-- Bryce Adelstein Lelbach aka wash (he/him/his) US Programming Language Standards (PL22) Chair ISO C++ Library Evolution Chair CppCon and C++Now Program Chair HPC Programming Models Architect @ NVIDIA --

[alliepiper]

We could add a par_nosync policy for folks who want the old behavior.

That could work, though it may be misleading since some thrust algorithms run multiple kernels that require synchronization between launches. But we could potentially drop the last sync.

[brycelelbach]

Yah, that's why we added the sync in the first place, for consistency. But par_nosync would be nice for folks just calling `transform`, which is a lot of people.

…

On Wed, Oct 6, 2021 at 6:10 PM Allison Vacanti ***@***.***> wrote: We could add a par_nosync policy for folks who want the old behavior. That could work, though it may be misleading since some thrust algorithms run multiple kernels that require synchronization between launches. But we could potentially drop the last sync. — You are receiving this because you commented. Reply to this email directly, view it on GitHub, or unsubscribe.

-- Bryce Adelstein Lelbach aka wash (he/him/his) US Programming Language Standards (PL22) Chair ISO C++ Library Evolution Chair CppCon and C++Now Program Chair HPC Programming Models Architect @ NVIDIA --

[jedbrown]

Indeed, this would be really convenient since our alternative right now is to replace PETSc's use of thrust::transform with raw CUDA to avoid some embarrassing latency costs.

[alliepiper]

I think it's doable. We can just document it as a hint that the implementation may ignore if needed.

I'll try to spend some time looking at this soon, since it's a common request.

[fkallen]

@allisonvacanti I created a pull request with a possible implementation of par_nosync