Conversation
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
…tests Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Contributor
There was a problem hiding this comment.
Pull request overview
This PR adds an optimization to the SYCL radix-sort “onesweep” kernel to early-exit the SLM reorder/scatter path when an entire work-group’s tile maps into a single radix bin, and augments KT tests with constrained-range inputs intended to trigger that fast path.
Changes:
- Add single-bin detection in
__rank_globaland a direct-copy output path to skip SLM reorder when applicable. - Add
generate_constrained_range_datatest utility and new constrained-range test cases for in-place, out-of-place, and by-key variants. - Extend existing radix-sort KT test executables to run the new constrained-range scenarios for both ascending/descending.
Reviewed changes
Copilot reviewed 6 out of 6 changed files in this pull request and generated 1 comment.
Show a summary per file
| File | Description |
|---|---|
| include/oneapi/dpl/experimental/kt/internal/sycl_radix_sort_kernels.h | Implements single-bin detection and direct-copy fast path in the onesweep kernel. |
| test/kt/radix_sort_utils.h | Adds constrained-range data generator used to trigger single-bin optimization in tests. |
| test/kt/radix_sort.cpp | Adds constrained-range in-place radix-sort test coverage. |
| test/kt/radix_sort_out_of_place.cpp | Adds constrained-range out-of-place radix-sort test coverage. |
| test/kt/radix_sort_by_key.cpp | Adds constrained-range radix-sort-by-key (in-place) test coverage. |
| test/kt/radix_sort_by_key_out_of_place.cpp | Adds constrained-range radix-sort-by-key (out-of-place) test coverage. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Comment on lines
+424
to
428
| // Detect single-bin tiles: if all elements land in one bin, we can skip SLM reorder. | ||
| bool __my_bin_is_full = (__item_bin_count == __data_per_work_group); | ||
| bool __is_single_bin = sycl::any_of_group(__group, __my_bin_is_full); | ||
|
|
||
| // 1.4. Finalize the partial scans from step 1.3 by connecting the independent sub-group segments |
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
Signed-off-by: Matthew Michel <matthew.michel@intel.com>
This branch has not been deployed
This file contains hidden or 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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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.
In the SYCL radix sort KT, keys are reordered into SLM before being scattered into memory. In the case where all keys
collapse to the same bin in a work-group, each work-group is able to copy its already loaded keys and values into global memory in sequential order.
When scanning over keys prior to lookback, we can detect this case and entirely skip SLM reorder after the lookback.