Skip to content

Add nonstable env DeviceSegmentedSort::Pairs 2/2 - #8286

Merged
gonidelis merged 2 commits into
NVIDIA:mainfrom
gonidelis:seg_sort_pairs_env_nonstable
Apr 14, 2026
Merged

gonidelis merged 2 commits into
NVIDIA:mainfrom
gonidelis:seg_sort_pairs_env_nonstable

Conversation

@gonidelis

Copy link
Copy Markdown
Member

Split 2/2 from from #8003 for easier review

@gonidelis
gonidelis requested a review from a team as a code owner April 3, 2026 10:47
@github-project-automation github-project-automation Bot moved this to Todo in CCCL Apr 3, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Apr 3, 2026
@gonidelis gonidelis changed the title Add nonstable env DeviceSegmentedSort::Pairs Add nonstable env DeviceSegmentedSort::Pairs Apr 3, 2026
@github-actions

This comment has been minimized.

@gonidelis gonidelis changed the title Add nonstable env DeviceSegmentedSort::Pairs Add nonstable env DeviceSegmentedSort::Pairs 2/2 Apr 3, 2026
@gonidelis gonidelis mentioned this pull request Apr 8, 2026
5 of 20 tasks
@github-actions

This comment has been minimized.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Important: please add a custom-stream env test in this file, similar to the sibling env suites for segmented radix sort / merge sort (see

TEST_CASE("DeviceMergeSort::SortPairs uses custom stream", "[merge_sort][device]")
for example). The new cases here only pass expected_allocation_size(...), so they verify temp-storage plumbing but not that the launch-helper path / segmented sort dispatch actually
honors cuda::get_stream_t. We still have stream-bearing examples in _env_api.cu, but those don’t give the same coverage as the helper-based env

Add SortPairs and SortPairsDescending env overloads for both pointer
and DoubleBuffer variants. Each delegates to SortPairsNoNVTX<Order>
via dispatch_with_env, matching the pattern from the stable overloads.

Tests cover default env, custom stream propagation through the
launch-helper path, allocation size verification, and env-based API
examples with stream.sync() before assertions.
@gonidelis
gonidelis force-pushed the seg_sort_pairs_env_nonstable branch from bcfbd9e to 1568074 Compare April 9, 2026 10:34
@gonidelis

Copy link
Copy Markdown
Member Author

I repushed all changes clean after refactoring that was done in #8003 with the current reviews incorporated.

@github-actions

This comment has been minimized.

Comment on lines +207 to +208
cudaStream_t custom_stream;
REQUIRE(cudaSuccess == cudaStreamCreate(&custom_stream));

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Important: please use cuda::stream here for all the tests instead of the C APIs.

@gonidelis
gonidelis enabled auto-merge (squash) April 14, 2026 00:56
@github-actions

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 1h 18m: Pass: 100%/269 | Total: 2d 17h | Max: 57m 03s | Hits: 98%/177321

See results here.

@gonidelis
gonidelis merged commit faf620e into NVIDIA:main Apr 14, 2026
285 of 288 checks passed
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Apr 14, 2026
jainishmehta pushed a commit to jainishmehta/cccl that referenced this pull request Apr 19, 2026
* Add nonstable env overloads for DeviceSegmentedSort::SortPairs

Add SortPairs and SortPairsDescending env overloads for both pointer
and DoubleBuffer variants. Each delegates to SortPairsNoNVTX<Order>
via dispatch_with_env, matching the pattern from the stable overloads.

Tests cover default env, custom stream propagation through the
launch-helper path, allocation size verification, and env-based API
examples with stream.sync() before assertions.

* Pass stream_ref directly and replace cudaStream_t with cuda::stream
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Archived in project

Development

Successfully merging this pull request may close these issues.

4 participants