Skip to content
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

Resolve register spills in dispatch of __subgroup_radix_sort #1626

Merged
merged 10 commits into from
Jun 20, 2024
Merged

Resolve register spills in dispatch of __subgroup_radix_sort #1626

Changes from all commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
7e08e76
Resolve register spills in single-wg radix_sort
mmichel11 Jun 5, 2024
4dd6520
Switch to a runtime check for sub-group size in problematic case
mmichel11 Jun 11, 2024
b0fe6c3
Add missing include
mmichel11 Jun 11, 2024
dc9ec48
Correct <algorithm> include location
mmichel11 Jun 11, 2024
d10fa17
Clang format
mmichel11 Jun 11, 2024
bf9dc32
Add clarifying comment on sub-group size check
mmichel11 Jun 17, 2024
34e633a
Add second check to avoid register spills
mmichel11 Jun 17, 2024
c72cb6e
Move comment to before the first subgroup size check
mmichel11 Jun 17, 2024
cee1a50
Make variables const where appropriate
mmichel11 Jun 18, 2024
a489b24
Clarify comment on sub-group sizes and register spills
mmichel11 Jun 18, 2024
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
13 changes: 11 additions & 2 deletions include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_radix_sort.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,6 +20,7 @@
#include <type_traits>
#include <utility>
#include <cstdint>
#include <algorithm>

#include "sycl_defs.h"
#include "parallel_backend_sycl_utils.h"
Expand Down Expand Up @@ -778,6 +779,9 @@ __parallel_radix_sort(oneapi::dpl::__internal::__device_backend_tag, _ExecutionP

//TODO: 1.to reduce number of the kernels; 2.to define work group size in runtime, depending on number of elements
constexpr auto __wg_size = 64;
const auto __subgroup_sizes = __exec.queue().get_device().template get_info<sycl::info::device::sub_group_sizes>();
const bool __dev_has_sg16 = std::find(__subgroup_sizes.begin(), __subgroup_sizes.end(),
static_cast<std::size_t>(16)) != __subgroup_sizes.end();

//TODO: with _RadixSortKernel also the following a couple of compile time constants is used for unique kernel name
using _RadixSortKernel = oneapi::dpl::__internal::__policy_kernel_name<_ExecutionPolicy>;
Expand All @@ -803,10 +807,15 @@ __parallel_radix_sort(oneapi::dpl::__internal::__device_backend_tag, _ExecutionP
else if (__n <= 4096 && __wg_size * 4 <= __max_wg_size)
__event = __subgroup_radix_sort<_RadixSortKernel, __wg_size * 4, 16, __radix_bits, __is_ascending>{}(
__exec.queue(), ::std::forward<_Range>(__in_rng), __proj);
else if (__n <= 8192 && __wg_size * 8 <= __max_wg_size)
// In __subgroup_radix_sort, we request a sub-group size via _ONEDPL_SYCL_REQD_SUB_GROUP_SIZE_IF_SUPPORTED
// based upon the iters per item. For the below cases, register spills that result in runtime exceptions have
// been observed on accelerators that do not support the requested sub-group size of 16. For the above cases
// that request but may not receive a sub-group size of 16, inputs are small enough to avoid register
// spills on assessed hardware.
else if (__n <= 8192 && __wg_size * 8 <= __max_wg_size && __dev_has_sg16)
Copy link
Contributor

Choose a reason for hiding this comment

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

This avoids single workgroup radix sort in some of the cases which expect / request subgroup size 16, but not all.
To avoid all cases which want sg 16, it looks like we would need to check __dev_has_sg16 in all cases __n > 256. That is likely overkill, but do we have justification for this being the size cutoff of cases affected by this register overflow error?

In other words, on different hardware, might we see the same error for the __n <= 4096 case or smaller?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I have found that in the other cases that request subgroup sizes of 16, using a subgroup size of 32 is safe from register spills. I have looked through the different CUDA architectures and the register file size seems to have remained constant over time, so I believe on NvGPUs it will resolve the issue. I have also verified with sm_75.

In the case of a general device, I think this is a risk anywhere we use private memory. It is difficult to fully protect against since there is no SYCL check for maximum private memory per group. On some hardware platforms such as Intel GPUs, the registers will spill into global memory and only impact performance. On CUDA devices, it causes a runtime exception.

Copy link
Contributor

Choose a reason for hiding this comment

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

Fair enough, I just wanted to make sure that we had good justification for this choice, rather than this merely being where we have experienced errors.

It may be good to mention this in the comment, that while smaller cases would prefer subgroup size 16 and may end up as 32, they still fit within the register file for hardware we are aware of so that the intention is clear for future maintenance.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I have updated the comment

__event = __subgroup_radix_sort<_RadixSortKernel, __wg_size * 8, 16, __radix_bits, __is_ascending>{}(
__exec.queue(), ::std::forward<_Range>(__in_rng), __proj);
else if (__n <= 16384 && __wg_size * 8 <= __max_wg_size)
else if (__n <= 16384 && __wg_size * 8 <= __max_wg_size && __dev_has_sg16)
__event = __subgroup_radix_sort<_RadixSortKernel, __wg_size * 8, 32, __radix_bits, __is_ascending>{}(
__exec.queue(), ::std::forward<_Range>(__in_rng), __proj);
else
Expand Down
Loading