Use nosync execution policy everywhere - #20807
Conversation
bdice
added
improvement
|
AFAIK compute-sanitizer won't help us find the issues with host object going out of scope because we removed a sync point. I'm not sure we should replace all without examining each instance for this kind of issue. |
|
It may help to investigate these errors with https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/index.html#stream-ordered-race-detection: |
| auto const end = begin + input.num_rows(); | ||
| auto const indices_end = | ||
| thrust::copy_if(rmm::exec_policy(stream), begin, end, indices.begin(), filter); | ||
| thrust::copy_if(rmm::exec_policy_nosync(stream), begin, end, indices.begin(), filter); |
There was a problem hiding this comment.
I guess what worries me is that I can foresee where the nosync policy allows the copy_if to assume the return value is not needed and implement an optimization that no longer performs a synchronize on the stream.
This comment only applies to those few places where we know a sync is needed.
I feel being purposeful also more clearly communicates intent to the reader/reviewer.
There was a problem hiding this comment.
Thrust cannot implement this algorithm without a sync. Thrust will always uphold the contract of returning a value to host that is ready to use. In other words, “nosync” is best-effort.
Thought about this a bit more. The scenario I was concerned with is when we use a |
|
I'm putting this on pause until after I get #20800 and some other projects finished up first. |
|
I can take this over in January if you want. |
|
@davidwendt That would be great! Thank you for offering. |
|
I'd like to merge this once CI passes since it's a large PR. There are a few follow-up tasks:
|
|
/merge |
… add docs (#20978) Follow-up to #20807. This PR documents the expectations for `nosync` execution policies and applies the changes to remaining tests, benchmarks, Java, and Python code. It also adds pre-commit validation that `rmm::exec_policy` is not used. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Matthew Roeschke (https://github.com/mroeschke) - James Lamb (https://github.com/jameslamb) - Tianyu Liu (https://github.com/kingcrimsontianyu) - Nghia Truong (https://github.com/ttnghia) - Muhammad Haseeb (https://github.com/mhaseeb123) URL: #20978
This is a follow-up to NVIDIA#11577 and NVIDIA#12086. We discussed the `exec_policy_nosync` and we would like to experiment with enabling it everywhere in libcudf. Since last time we investigated this, we have refined a lot of the library to be more stream-friendly. We also have weekly compute-sanitizer runs now, to help us identify any issues. We would like to see if this changes provides any performance improvements or reduces synchronization in the multi-thread, multi-stream workflow engines that use libcudf. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Nghia Truong (https://github.com/ttnghia) - David Wendt (https://github.com/davidwendt) URL: NVIDIA#20807
… add docs (NVIDIA#20978) Follow-up to NVIDIA#20807. This PR documents the expectations for `nosync` execution policies and applies the changes to remaining tests, benchmarks, Java, and Python code. It also adds pre-commit validation that `rmm::exec_policy` is not used. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Matthew Roeschke (https://github.com/mroeschke) - James Lamb (https://github.com/jameslamb) - Tianyu Liu (https://github.com/kingcrimsontianyu) - Nghia Truong (https://github.com/ttnghia) - Muhammad Haseeb (https://github.com/mhaseeb123) URL: NVIDIA#20978
4 participants
Description
This is a follow-up to #11577 and #12086. We discussed the
exec_policy_nosyncand we would like to experiment with enabling it everywhere in libcudf.Since last time we investigated this, we have refined a lot of the library to be more stream-friendly. We also have weekly compute-sanitizer runs now, to help us identify any issues. We would like to see if this changes provides any performance improvements or reduces synchronization in the multi-thread, multi-stream workflow engines that use libcudf.
Checklist