Add rmm::exec_policy_nosync

[fkallen]

Add rmm:exec_policy_nosync.

For Thrust 1.16 and newer, this policy is based on execute_on_stream_nosync_base instead of execute_on_stream_base to allow Thrust algorithms to skip synchronizations.

For older Thrust versions, this is an alias to rmm::exec_policy

[GPUtester]

Can one of the admins verify this patch?

[GPUtester]

Can one of the admins verify this patch?

[jrhemstad]

tbh, I'd been thinking about just changing the rmm::exec_policy to use the nosync version.

Technically this would be a breaking change, but I'd be surprised if anything in RAPIDS was depending on the synchronous behavior of Thrust algorithms. I suspect most would appreciate the elimination of the extra syncs.

I'll see what @harrism thinks

[harrism]

but I'd be surprised if anything in RAPIDS was depending on the synchronous behavior of Thrust algorithms.

Really? I won't be surprised. Not saying I oppose it, just that we should do an investigation before flipping a switch like this.

[jrhemstad]

Really? I won't be surprised.

What can I say, people always tell me I'm such an optimist ¯_(ツ)_/¯

To be clear, algorithms like thrust::reduce would still synchronize with the nosync policy because they have to.

The nosync policy controls if an algorithm like thrust::transform or thrust::for_each synchronizes or not. I'd hope we didn't have code relying on that behavior.

[harrism]

Just saying we should test all RAPIDS libraries before merging such a change (not just RMM tests). Also test Spark since they rely on PTDS.

[jrhemstad]

Just saying we should test all RAPIDS libraries before merging such a change (not just RMM tests). Also test Spark since they rely on PTDS.

I think I've changed my mind. I don't think the hassle of having to verify every usage of rmm::exec_policy is worth it. It's easy enough to find/replace rmm::exec_policy with rmm::exec_policy_async in consumer libraries.

[jrhemstad]

Thanks @fkallen! Could you also add some tests?

[fkallen]

Do you have a test in mind? I don't know how to test this other than taking a look at the profiler to check if there is a pool allocation and a missing sync.

[jrhemstad]

Do you have a test in mind? I don't know how to test this other than taking a look at the profiler to check if there is a pool allocation and a missing sync.

I was just thinking making sure that it works as a valid execution policy. So simply launch a Thrust algorithm with device data and make sure the result is sound.

I was going to say "just do what we do for our current tests", but apparently that's nothing!

[harrism]

ok to test

[harrism]

rerun tests

[cwharris]

nosync isn't a term we use in rmm, or any rapids libraries as far as I know. What if we use async, which is common?

Nevermind, I see now that this is based on thrust policy.

[harrism]

@gpucibot merge