This repository was archived by the owner on Mar 21, 2024. It is now read-only.
- Notifications
You must be signed in to change notification settings - Fork449
Add DeviceMergeSort::StableSortKeysCopy API#565
Merged
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
gevtushenko approved these changesSep 8, 2022
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
Thank you for the contribution! A few minor suggestions below.
cub/device/device_merge_sort.cuh Outdated
Comment on lines 1117 to 1142
| template <typename KeyInputIteratorT, | ||
| typename KeyIteratorT, | ||
| typename OffsetT, | ||
| typename CompareOpT> | ||
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED | ||
| CUB_RUNTIME_FUNCTION static cudaError_t | ||
| StableSortKeysCopy(void *d_temp_storage, | ||
| std::size_t &temp_storage_bytes, | ||
| KeyInputIteratorT d_input_keys, | ||
| KeyIteratorT d_output_keys, | ||
| OffsetT num_items, | ||
| CompareOpT compare_op, | ||
| cudaStream_t stream, | ||
| bool debug_synchronous) | ||
| { | ||
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG | ||
| return StableSortKeysCopy<KeyInputIteratorT, KeyIteratorT, OffsetT, CompareOpT>(d_temp_storage, | ||
| temp_storage_bytes, | ||
| d_input_keys, | ||
| d_output_keys, | ||
| num_items, | ||
| compare_op, | ||
| stream); | ||
| } | ||
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
Suggested change
| template<typename KeyInputIteratorT, | |
| typename KeyIteratorT, | |
| typename OffsetT, | |
| typename CompareOpT> | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_IS_NOT_SUPPORTED | |
| CUB_RUNTIME_FUNCTIONstatic cudaError_t | |
| StableSortKeysCopy(void *d_temp_storage, | |
| std::size_t &temp_storage_bytes, | |
| KeyInputIteratorT d_input_keys, | |
| KeyIteratorT d_output_keys, | |
| OffsetT num_items, | |
| CompareOpT compare_op, | |
| cudaStream_t stream, | |
| bool debug_synchronous) | |
| { | |
| CUB_DETAIL_RUNTIME_DEBUG_SYNC_USAGE_LOG | |
| return StableSortKeysCopy<KeyInputIteratorT, KeyIteratorT, OffsetT, CompareOpT>(d_temp_storage, | |
| temp_storage_bytes, | |
| d_input_keys, | |
| d_output_keys, | |
| num_items, | |
| compare_op, | |
| stream); | |
| } |
This is not required. Thedebug_synchronous overloads exist so as not to break API. For the new functions, we don't have anything to break, so let's remove this overload.
| @@ -234,6 +234,18 @@ void TestKeys(std::int64_t num_items, | |||
| CustomLess())); | |||
| AssertTrue(CheckResult(d_keys)); | |||
| CubDebugExit(cub::DeviceMergeSort::StableSortKeysCopy( | |||
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
Just realized that there might be a value in something like:
Suggested change
| CubDebugExit(cub::DeviceMergeSort::StableSortKeysCopy( | |
| thrust::fill(d_keys.begin(), d_keys.end(), KeyType{}); | |
| CubDebugExit(cub::DeviceMergeSort::StableSortKeysCopy( |
gevtushenko added a commit to gevtushenko/thrust that referenced this pull requestSep 9, 2022
miscco approved these changesSep 9, 2022
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others.Learn more.
Thanks for the contribution and the excellent documentation
Sign up for freeto subscribe to this conversation on GitHub. Already have an account?Sign in.
Labels
3 participants
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.
CUB has a
SortKeysCopyand aStableSortKeys. This PR addsStableSortKeysCopyso a sort in-place is not required. The in-placeStableSortKeysrequires a temporary vector if a copy is needed.Since
StableSortKeyscallsSortKeysandSortKeysCopyuses the same kernels asSortKeys(which appear to already be stable-sort enabled), the newStableSortKeysCopysimply callsSortKeysCopy.