This repository has been archived by the owner on Mar 21, 2024. It is now read-only.
Add DeviceMergeSort::StableSortKeysCopy API #565
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 changes
Sep 8, 2022
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.
Suggested change
| 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); | |
| } |
This is not required. The debug_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.
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 request
Sep 9, 2022
miscco
approved these changes
Sep 9, 2022
gevtushenko
added
testing: gpuCI passed
Sign up for free
to subscribe to this conversation on GitHub.
Already have an account?
Sign in.
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.