Add env overloads for DeviceSegmentedRadixSort by gonidelis · Pull Request #7999 · NVIDIA/cccl

gonidelis · 2026-03-12T02:47:39Z

No deterministic guarantees. It's sorting, if it's not deterministic it's just wrong.

NaderAlAwar

Important: we are missing the DoubleBuffer overloads

NaderAlAwar · 2026-03-12T13:32:49Z

cub/test/catch2_test_device_segmented_radix_sort_env_api.cu

+    offsets.begin() + 1,
+    0,
+    sizeof(int) * 8,
+    env);


Question: do we want to add a test where we do not explicitly pass an env? Other env API tests seem to have both kinds, explicit and implicit env objects.

i think it's fine. unit tests cover that case and users can see it throughout docs as soon as they start using our env overloads. don't wanna over bloat our test base for

We have that test in the other file added in this PR.

gonidelis · 2026-03-12T23:08:07Z

Important: we are missing the DoubleBuffer overloads

@NaderAlAwar will be added in subsequent pr for economy of reviewing effort per pr

cub/cub/device/device_segmented_radix_sort.cuh

bernhardmgruber · 2026-03-13T19:04:51Z

cub/test/catch2_test_device_segmented_radix_sort_env_api.cu

+    offsets.begin() + 1,
+    0,
+    sizeof(int) * 8,
+    env);


We have that test in the other file added in this PR.

github-actions · 2026-03-19T01:02:57Z

🥳 CI Workflow Results

🟩 Finished in 1h 10m: Pass: 100%/249 | Total: 3d 19h | Max: 55m 34s | Hits: 94%/157462

See results here.

pauleonix · 2026-03-19T23:35:55Z

No deterministic guarantees. It's sorting, if it's not deterministic it's just wrong.

Only because our radix-sort is stable by design. Unstable sort could be non-deterministic 😉

pauleonix

Still some nits

pauleonix · 2026-03-20T00:01:30Z

cub/cub/device/device_segmented_radix_sort.cuh

+  //!
+  //! Snippet
+  //!
+  //! .. literalinclude:: ../../../cub/test/catch2_test_device_segmented_radix_sort_env_api.cu


Suggested change

//! .. literalinclude:: ../../../cub/test/catch2_test_device_segmented_radix_sort_env_api.cu

//! .. literalinclude:: ../../test/catch2_test_device_segmented_radix_sort_env_api.cu

pauleonix · 2026-03-20T00:08:37Z

cub/cub/device/device_segmented_radix_sort.cuh

+  //! This is an environment-based API that allows customization of:
+  //!
+  //! - Stream: Query via ``cuda::get_stream``
+  //! - Memory resource: Query via ``cuda::mr::get_memory_resource``


Not sure I like this better than the - Can use a specific stream or cuda memory resource through the ``env`` parameter I saw in the previous PR. I'm also not sure this list wont be rendered/read as part as the following list.

pauleonix · 2026-03-20T00:12:10Z

cub/cub/device/device_segmented_radix_sort.cuh

+  //!   yield a corresponding performance improvement.
+  //! - Note, the size of any segment may not exceed ``INT_MAX``. Please consider using ``DeviceSegmentedSort`` instead,
+  //!   if the size of at least one of your segments could exceed ``INT_MAX``.
+  //!


Maybe we should have a @env like the @devicestorage on the old APIs to avoid the repetition. Even if it needs multiple like @env-with-guarantees.

I considered suggesting the same! No need to do it in this PR though.

pauleonix · 2026-03-20T00:21:25Z

cub/cub/device/device_segmented_radix_sort.cuh

+  //! - An optional bit subrange ``[begin_bit, end_bit)`` of differentiating key
+  //!   bits can be specified. This can reduce overall sorting overhead and
+  //!   yield a corresponding performance improvement.
+  //! - Note, the size of any segment may not exceed ``INT_MAX``. Please consider using ``DeviceSegmentedSort`` instead,


Missing

//! - Let ``in`` be one of ``{d_keys_in, d_values_in}`` and `out` be any of //! ``{d_keys_out, d_values_out}``. The range ``[out, out + num_items)`` shall //! not overlap ``[in, in + num_items)``, //! ``[d_begin_offsets, d_begin_offsets + num_segments)`` nor //! ``[d_end_offsets, d_end_offsets + num_segments)`` in any way.

and

//! - Segments are not required to be contiguous. For all index values ``i`` //! outside the specified segments ``d_keys_in[i]``, ``d_values_in[i]``, //! ``d_keys_out[i]``, ``d_values_out[i]`` will not be accessed nor modified.

Maybe also

//! - @devicestorageNP For sorting using only ``O(P)`` temporary storage, see //! the sorting interface using DoubleBuffer wrappers below.

pauleonix · 2026-03-20T00:26:25Z

cub/cub/device/device_segmented_radix_sort.cuh

  }

+  //! @rst
+  //! Overview


This "title" is new in comparison to the existing docs. If it stays, it looks like it should be underlined (rendered as a sub-title)

Suggested change

//! Overview

//! Overview

//! +++++++++++++++++++++++++++++++++++++++++++++

Same for Snippet below (which is not new but has this underline in existing docs)

pauleonix · 2026-03-20T00:35:41Z

cub/cub/device/device_segmented_radix_sort.cuh

+            typename BeginOffsetIteratorT,
+            typename EndOffsetIteratorT,
+            typename EnvT = ::cuda::std::execution::env<>>
+  [[nodiscard]] CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t SortKeys(


Is _CCCL_FORCEINLINE necessary here? Existing overload does not have it.

Not necessary, please remove!

pauleonix · 2026-03-20T00:37:57Z

cub/cub/device/device_segmented_radix_sort.cuh

+  //! - An optional bit subrange ``[begin_bit, end_bit)`` of differentiating key
+  //!   bits can be specified. This can reduce overall sorting overhead and
+  //!   yield a corresponding performance improvement.
+  //! - Note, the size of any segment may not exceed ``INT_MAX``. Please consider using ``DeviceSegmentedSort`` instead,


Missing entries again.

pauleonix · 2026-03-20T00:43:27Z

cub/test/catch2_test_device_segmented_radix_sort_env.cu

+    thrust::raw_pointer_cast(keys_out.data()),
+    thrust::raw_pointer_cast(values_in.data()),
+    thrust::raw_pointer_cast(values_out.data()),
+    static_cast<int>(keys_in.size()),


Suggested change

static_cast<int>(keys_in.size()),

static_cast<cuda::std::int64_t>(keys_in.size()),

pauleonix · 2026-03-20T00:47:34Z

cub/test/catch2_test_device_segmented_radix_sort_env.cu

+  REQUIRE(keys_out == expected_keys);
+}
+
+TEST_CASE("DeviceSegmentedRadixSort::SortPairs uses custom stream", "[segmented_radix_sort][device]")


Should this be tested for all overloads?

pauleonix · 2026-03-20T00:49:45Z

cub/test/catch2_test_device_segmented_radix_sort_env_api.cu

+    thrust::raw_pointer_cast(keys_out.data()),
+    thrust::raw_pointer_cast(values_in.data()),
+    thrust::raw_pointer_cast(values_out.data()),
+    static_cast<int>(keys_in.size()),


Suggested change

static_cast<int>(keys_in.size()),

static_cast<cuda::std::int64_t>(keys_in.size()),

Do we need a cast at all here? To suppress the unsigned -> signed warning, right?

bernhardmgruber

Please apply the remaining feedback from other reviewers. Otherwise LGTM

gonidelis requested a review from a team as a code owner March 12, 2026 02:47

gonidelis requested a review from pauleonix March 12, 2026 02:47

github-project-automation bot added this to CCCL Mar 12, 2026

github-project-automation bot moved this to Todo in CCCL Mar 12, 2026

gonidelis requested a review from bernhardmgruber March 12, 2026 02:47

cccl-authenticator-app bot moved this from Todo to In Review in CCCL Mar 12, 2026

gonidelis requested review from NaderAlAwar, elstehle and srinivasyadav18 March 12, 2026 02:47

Add env overloads for DeviceSegmentedRadixSort

55ca027

gonidelis force-pushed the seg-radix-sort-env branch from 2c8d80a to 55ca027 Compare March 12, 2026 04:13

This comment has been minimized.

Sign in to view

NaderAlAwar reviewed Mar 12, 2026

View reviewed changes

bernhardmgruber reviewed Mar 13, 2026

View reviewed changes

cub/cub/device/device_segmented_radix_sort.cuh Outdated Show resolved Hide resolved

bernhardmgruber reviewed Mar 13, 2026

View reviewed changes

gonidelis mentioned this pull request Mar 13, 2026

Revamp enable_if constraints in all env overloads #8033

Open

Simplify template params in env overloads (remove sfinae, NumItemsT)

c577169

gonidelis requested review from NaderAlAwar and bernhardmgruber March 18, 2026 23:50

pauleonix reviewed Mar 20, 2026

View reviewed changes

bernhardmgruber approved these changes Mar 20, 2026

View reviewed changes

	//! .. literalinclude:: ../../../cub/test/catch2_test_device_segmented_radix_sort_env_api.cu
	//! .. literalinclude:: ../../test/catch2_test_device_segmented_radix_sort_env_api.cu

	//! Overview
	//! Overview
	//! +++++++++++++++++++++++++++++++++++++++++++++

	static_cast<int>(keys_in.size()),
	static_cast<cuda::std::int64_t>(keys_in.size()),

Conversation

gonidelis commented Mar 12, 2026

Uh oh!

This comment has been minimized.

NaderAlAwar left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

gonidelis commented Mar 12, 2026

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

github-actions bot commented Mar 19, 2026

🥳 CI Workflow Results

🟩 Finished in 1h 10m: Pass: 100%/249 | Total: 3d 19h | Max: 55m 34s | Hits: 94%/157462

Uh oh!

pauleonix commented Mar 19, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

pauleonix left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

bernhardmgruber Mar 20, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

pauleonix Mar 20, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

bernhardmgruber left a comment

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

pauleonix commented Mar 19, 2026 •

edited

Loading

bernhardmgruber Mar 20, 2026 •

edited

Loading

pauleonix Mar 20, 2026 •

edited

Loading