Enable use of cudaMemcpyAsync for thrust::copy
#211
Merged
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
miscco
added
thrust
|
@senior-zero does this resolve NVIDIA/thrust#1536 ? |
In case of contigous ranges of trivially relocatable types we can directly utilize `cudaMemcpyAsync` instead of going through transform. Fixes NVIDIA#210
miscco
force-pushed
the
utilize_cudaMemcpyAsync
branch
from
3f8dedd
to
c659362
Compare
gevtushenko
approved these changes
Jul 14, 2023
There was a problem hiding this comment.
This PR is potentially breaking on systems with concurrentManagedAccess=false (windows?) when used with unified memory (details). Ideally, we should check if concurrentManagedAccess=1 before dispatching to memcpy. Since this issue already exists in the case of cross-system copy, I think we can proceed with this PR and prioritize #210 (@jrhemstad @gonzalobg).
7 tasks
rapids-bot bot
pushed a commit
to rapidsai/rapids-cmake
that referenced
this pull request
Dec 19, 2023
Something is going wrong in `thrust::copy`, so I am adding a patch to revert NVIDIA/cccl#211. This will fix the failing tests for rapidsai/cuspatial#1315. This is where it's failing, with `thrust::copy`: https://github.com/rapidsai/cuspatial/blob/8e8fa0689339c0288746b01a8af66fb599f5c770/cpp/tests/range/multipolygon_range_test.cu#L800 CI log: https://github.com/rapidsai/cuspatial/actions/runs/7229405991/job/19700477218?pr=1315#step:7:3280 ``` [ RUN ] MultipolygonRangeOneTest/1.OneMultipolygonRange unknown file: Failure C++ exception with description "__copy:: D->D: failed: cudaErrorInvalidValue: invalid argument" thrown in the test body. [ FAILED ] MultipolygonRangeOneTest/1.OneMultipolygonRange, where TypeParam = double (0 ms) ``` I traced it down to this CCCL PR - note that `thrust::transform(..., thrust::identity{})` works, but the new `thrust::copy` code path in NVIDIA/cccl#211 that calls `cudaMemcpyAsync` doesn't work for some reason. I was not able to identify the exact reason why that PR causes the failure. Patching this out will unblock RAPIDS and let us migrate forward with CCCL 2.2.0. This helps with #502. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Michael Schellenberger Costa (https://github.com/miscco) - Robert Maynard (https://github.com/robertmaynard) URL: #511
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
In case of contigous ranges of trivially relocatable types we can directly utilize
cudaMemcpyAsyncinstead of going through transform.Fixes NVIDIA/thrust#1536