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

Support user-defined types in radix sort and introduce initial sphinx docs #671

Merged
merged 7 commits into from
May 21, 2023

Conversation

gevtushenko
Copy link
Collaborator

@gevtushenko gevtushenko commented Apr 6, 2023

This PR introduces support for user-defined types in the block / device radix sort (addresses #52):

struct custom_t
{
  std::uint16_t i;
  int unused;
  float f;
};

struct decomposer_t
{
  __host__ __device__ //
  ::cuda::std::tuple<std::uint16_t&, float&> operator()(custom_t &key) const
  {
    return {key.i, key.f};
  }
};

cub::DeviceRadixSort::SortKeys(d_temp_storage,
                               temp_storage_bytes,
                               d_in,
                               d_out,
                               num_items,
                               decomposer_t{});

There's no difference in performance between sorting a key that's a pair of int32 and sorting a single int64 key. The PR also doesn't affect SASS for fundamental types.

Documentation for the new API is written using sphinx, whose initial support is also introduced by this PR. Doxygen is temporarily broken, and a follow-up PR will provide the complete transition to sphinx. To build sphinx docs:

:cd docs 
:./gen_docs.sh 
:open _build/docs/CUB/latest/index.html

gevtushenko added a commit to gevtushenko/thrust that referenced this pull request Apr 6, 2023
gevtushenko added a commit to gevtushenko/thrust that referenced this pull request Apr 7, 2023
@gevtushenko gevtushenko added the testing: gpuCI in progress Started gpuCI testing. label Apr 7, 2023
gevtushenko added a commit to gevtushenko/thrust that referenced this pull request Apr 7, 2023
@gevtushenko gevtushenko added testing: gpuCI passed Passed gpuCI testing. and removed testing: gpuCI in progress Started gpuCI testing. labels Apr 8, 2023
@jrhemstad jrhemstad linked an issue May 10, 2023 that may be closed by this pull request
cub/agent/agent_radix_sort_downsweep.cuh Show resolved Hide resolved
cub/block/block_radix_sort.cuh Show resolved Hide resolved
cub/block/radix_rank_sort_operations.cuh Outdated Show resolved Hide resolved
cub/block/radix_rank_sort_operations.cuh Show resolved Hide resolved
cub/block/radix_rank_sort_operations.cuh Show resolved Hide resolved
cub/block/radix_rank_sort_operations.cuh Show resolved Hide resolved
cub/block/radix_rank_sort_operations.cuh Show resolved Hide resolved
cub/block/radix_rank_sort_operations.cuh Show resolved Hide resolved
Copy link
Collaborator

@elstehle elstehle left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Really solid work! Like how we can keep most of the implementation details unchanged.

Partial review; review of BlockRadixSort and tests remaining. Also want to push radix_rank_sort_operations through its paces a bit.

cub/device/device_radix_sort.cuh Show resolved Hide resolved
cub/agent/agent_radix_sort_upsweep.cuh Show resolved Hide resolved
cub/device/device_radix_sort.cuh Outdated Show resolved Hide resolved
Copy link
Collaborator

@elstehle elstehle left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Only a few, mostly minor, comments remaining 🙂

test/catch2_test_device_radix_sort_custom.cu Outdated Show resolved Hide resolved
test/catch2_test_device_radix_sort_custom.cu Outdated Show resolved Hide resolved
test/catch2_test_device_radix_sort_custom.cu Show resolved Hide resolved
test/catch2_test_device_radix_sort_custom.cu Show resolved Hide resolved
test/test_device_radix_sort.cu Outdated Show resolved Hide resolved
test/catch2_test_block_radix_sort.cuh Outdated Show resolved Hide resolved
test/catch2_test_block_radix_sort_custom.cu Outdated Show resolved Hide resolved
test/catch2_test_block_radix_sort.cuh Show resolved Hide resolved
test/catch2_test_radix_operations.cu Show resolved Hide resolved
test/catch2_test_block_radix_sort_custom.cu Show resolved Hide resolved
@gevtushenko gevtushenko force-pushed the enh-main/github/custom_radix branch from 803e61a to d428ad3 Compare May 20, 2023 11:33
@gevtushenko gevtushenko force-pushed the enh-main/github/custom_radix branch from e1ea3c9 to d1dcb77 Compare May 20, 2023 13:44
gevtushenko added a commit to gevtushenko/thrust that referenced this pull request May 20, 2023
Sign up for free to subscribe to this conversation on GitHub. Already have an account? Sign in.
Labels
testing: gpuCI passed Passed gpuCI testing.
Projects
Archived in project
Development

Successfully merging this pull request may close these issues.

Sphinx docs proof of concept
5 participants