Skip to content
This repository has been archived by the owner on Mar 21, 2024. It is now read-only.

Add async scan algorithms #1251

Merged
merged 2 commits into from
Jan 21, 2021
Merged

Add async scan algorithms #1251

merged 2 commits into from
Jan 21, 2021

Conversation

alliepiper
Copy link
Collaborator

@alliepiper alliepiper commented Aug 6, 2020 •
edited
Loading

Prerequisite: NVIDIA/cub#210

@alliepiper alliepiper marked this pull request as draft August 6, 2020 01:24
alliepiper added a commit to alliepiper/thrust that referenced this pull request Aug 6, 2020
@alliepiper
Restore some THRUST_DECLTYPE_RETURNS macros in async test implementat…
582252c 
…ions.

This partially reverts 7ff227a and
fixes NVIDIA#1250.

I'm not sure why changing these broke the tests, but since these
usages are just testing details that are being refactored by NVIDIA#1251
let's just revert the change for now.

The test failures were only happening on GCC, MSVC was fine with both
versions of these functions, so it may be a compiler issue.
@alliepiper alliepiper mentioned this pull request Aug 6, 2020
Merged
@alliepiper alliepiper force-pushed the feature/async_scan_algos branch from 45c8380 to 87ae049 Compare August 7, 2020 01:12
@brycelelbach brycelelbach added this to the 1.10.0 milestone Aug 7, 2020
alliepiper added a commit that referenced this pull request Aug 12, 2020
@alliepiper
Restore some THRUST_DECLTYPE_RETURNS macros in async test implementat…
a713f93 
…ions.

This partially reverts 7ff227a and
fixes #1250.

I'm not sure why changing these broke the tests, but since these
usages are just testing details that are being refactored by #1251
let's just revert the change for now.

The test failures were only happening on GCC, MSVC was fine with both
versions of these functions, so it may be a compiler issue.
@alliepiper alliepiper force-pushed the feature/async_scan_algos branch 2 times, most recently from d6ae713 to 7caa692 Compare August 27, 2020 20:58
@alliepiper alliepiper self-assigned this Sep 3, 2020
@alliepiper alliepiper force-pushed the feature/async_scan_algos branch 4 times, most recently from ab507ce to 9e7dfe2 Compare September 10, 2020 18:12
@brycelelbach brycelelbach modified the milestones: 1.10.0, 1.11.0 Sep 16, 2020
@alliepiper alliepiper force-pushed the feature/async_scan_algos branch from 9e7dfe2 to c953231 Compare October 1, 2020 23:04
@alliepiper alliepiper mentioned this pull request Oct 2, 2020
Closed
@alliepiper alliepiper force-pushed the feature/async_scan_algos branch from c953231 to c236ed8 Compare October 2, 2020 21:26
@alliepiper alliepiper marked this pull request as ready for review October 5, 2020 18:49
@alliepiper
Copy link
Collaborator Author

The "large indices" tests are currently failing because temporary memory allocation fails for the larger problem sizes. This doesn't happen for the synchronous scan (which uses a custom implementation instead of just calling cub::DeviceScan). I'll pull those tests out before merging if I can't find a solution/workaround.

@alliepiper alliepiper removed their assignment Oct 5, 2020
@alliepiper alliepiper force-pushed the feature/async_scan_algos branch from c236ed8 to 3ff9d84 Compare October 14, 2020 20:57
@alliepiper alliepiper added the blocked Cannot make progress. label Oct 14, 2020
@alliepiper
Copy link
Collaborator Author

mixed_types test should still be failing, but running DVS to try to repro build timeouts.

DVS CL: 29371528

@alliepiper alliepiper removed the testing: gpuCI passed Passed gpuCI testing. label Nov 30, 2020
@alliepiper
Copy link
Collaborator Author

Latest round of DVS runs show the following issues on Clang:

thrust/system/cuda/detail/async/exclusive_scan.h:160:8: error: moving a local object in a return statement prevents copy elision [-Werror,-Wpessimizing-move]
thrust/system/cuda/detail/async/inclusive_scan.h:158:8: error: moving a local object in a return statement prevents copy elision [-Werror,-Wpessimizing-move]

@alliepiper alliepiper force-pushed the feature/async_scan_algos branch from 56c0eec to a071c17 Compare December 8, 2020 23:03
@alliepiper
Copy link
Collaborator Author

run tests

1 similar comment
@brycelelbach
Copy link
Collaborator

run tests

@alliepiper
Add iterator_*_t helpers for iterator traits.
1469009 
- iterator_value_t
- iterator_pointer_t
- iterator_reference_t
- iterator_difference_t
- iterator_system_t
@alliepiper alliepiper force-pushed the feature/async_scan_algos branch from a071c17 to 888512b Compare January 6, 2021 00:18
@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper alliepiper force-pushed the feature/async_scan_algos branch from 888512b to e1b3caa Compare January 6, 2021 02:20
@alliepiper
Copy link
Collaborator Author

run tests

@alliepiper
Copy link
Collaborator Author

DVS CL: 29473661

@alliepiper alliepiper added testing: internal ci in progress Currently testing on internal NVIDIA CI (DVS). testing: gpuCI in progress Started gpuCI testing. labels Jan 6, 2021
@alliepiper alliepiper added testing: gpuCI passed Passed gpuCI testing. testing: internal ci passed Passed internal NVIDIA CI (DVS). and removed testing: gpuCI in progress Started gpuCI testing. testing: internal ci in progress Currently testing on internal NVIDIA CI (DVS). labels Jan 7, 2021
@alliepiper alliepiper merged commit 0d51e20 into NVIDIA:main Jan 21, 2021
@alliepiper alliepiper deleted the feature/async_scan_algos branch January 21, 2021 20:45
Sign up for free to subscribe to this conversation on GitHub. Already have an account? Sign in.
Reviewers

@trivialfis trivialfis trivialfis left review comments

@brycelelbach brycelelbach brycelelbach approved these changes

Assignees

@alliepiper alliepiper

Labels
testing: gpuCI passed Passed gpuCI testing. testing: internal ci passed Passed internal NVIDIA CI (DVS).
Projects
None yet
Milestone
1.12.0
Development

Successfully merging this pull request may close these issues.
4 participants
@alliepiper @brycelelbach @hcedwar @trivialfis