Skip to content

[hipCUB] Fix oob_default bug for block_merge_sort. #124

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
stanleytsang-amd merged 4 commits into from
May 30, 2025
Merged

[hipCUB] Fix oob_default bug for block_merge_sort. #124

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
1c5165c
fixed block mergesort for case with oob_defaul
NguyenNhuDi May 2, 2025
372a085
updated changelog
NguyenNhuDi May 2, 2025
69b931a
fixed warp merge sort error caused by changes
NguyenNhuDi May 15, 2025
1d91e70
Merge commit '69b931af63ac3249083f79dae0c3973abb73e685' into import/r…
assistant-librarian[bot] May 29, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 3 additions & 0 deletions projects/hipcub/CHANGELOG.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,6 +38,9 @@ Full documentation for hipCUB is available at [https://rocm.docs.amd.com/project
* The `hipcub::detail::accumulator_t` in rocPRIM backend has been changed to utilise `rocprim::accumulator_t`.
* The usage of `rocprim::invoke_result_binary_op_t` has been replaced with `rocprim::accumulator_t`.

### Resolved issues
* Fixed an issue where `Sort(keys, compare_op, valid_items, oob_default)` in `block_merge_sort.hpp` would not fill in elements that are out of range (items after `valid_items`) with `oob_default`.

### Known issues

* `BlockAdjacentDifference::FlagHeads`, `BlockAdjacentDifference::FlagTails` and `BlockAdjacentDifference::FlagHeadsAndTails` have been removed from hipCUB's CUB backend. They were already deprecated as of version 2.12.0 of hipCUB and they were removed from CCCL (CUB) as of CCCL's 2.6.0 release.
Expand Down
15 changes: 9 additions & 6 deletions projects/hipcub/hipcub/include/hipcub/backend/rocprim/block/block_merge_sort.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -172,7 +172,8 @@ template <typename KeyT,
typename ValueT,
int NUM_THREADS,
int ITEMS_PER_THREAD,
typename SynchronizationPolicy>
typename SynchronizationPolicy,
bool WARP_SORT = false>
class BlockMergeSortStrategy
{
static_assert(PowerOfTwo<NUM_THREADS>::VALUE,
Expand Down Expand Up @@ -387,7 +388,7 @@ class BlockMergeSortStrategy
KeyT max_key = oob_default;

#pragma unroll
for (int item = 1; item < ITEMS_PER_THREAD; ++item)
for (int item = WARP_SORT ? 1 : 0; item < ITEMS_PER_THREAD; ++item)
{
if (ITEMS_PER_THREAD * static_cast<int>(linear_tid) + item < valid_items)
{
Expand Down Expand Up @@ -439,14 +440,16 @@ class BlockMergeSortStrategy

int thread_idx_in_thread_group_being_merged = mask & linear_tid;

const int COMPARE_NUM = WARP_SORT ? valid_items : ITEMS_PER_THREAD * blockDim.x;

int diag =
(::rocprim::min)(valid_items,
(::rocprim::min)(COMPARE_NUM,
ITEMS_PER_THREAD * thread_idx_in_thread_group_being_merged);

int keys1_beg = (::rocprim::min)(valid_items, start);
int keys1_end = (::rocprim::min)(valid_items, keys1_beg + size);
int keys1_beg = (::rocprim::min)(COMPARE_NUM, start);
int keys1_end = (::rocprim::min)(COMPARE_NUM, keys1_beg + size);
int keys2_beg = keys1_end;
int keys2_end = (::rocprim::min)(valid_items, keys2_beg + size);
int keys2_end = (::rocprim::min)(COMPARE_NUM, keys2_beg + size);

int keys1_count = keys1_end - keys1_beg;
int keys2_count = keys2_end - keys2_beg;
Expand Down
6 changes: 4 additions & 2 deletions projects/hipcub/hipcub/include/hipcub/backend/rocprim/warp/warp_merge_sort.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -129,7 +129,8 @@ class WarpMergeSort
ValueT,
LOGICAL_WARP_THREADS,
ITEMS_PER_THREAD,
WarpMergeSort<KeyT, ITEMS_PER_THREAD, LOGICAL_WARP_THREADS, ValueT, PTX_ARCH>>
WarpMergeSort<KeyT, ITEMS_PER_THREAD, LOGICAL_WARP_THREADS, ValueT, PTX_ARCH>,
true>
{
private:
constexpr static bool IS_ARCH_WARP = LOGICAL_WARP_THREADS == HIPCUB_DEVICE_WARP_THREADS;
Expand All @@ -140,7 +141,8 @@ class WarpMergeSort
ValueT,
LOGICAL_WARP_THREADS,
ITEMS_PER_THREAD,
WarpMergeSort>;
WarpMergeSort,
true>;

const unsigned int warp_id;
const uint64_t member_mask;
Expand Down
Loading