Resolve register spills in dispatch of __subgroup_radix_sort

[mmichel11]

On certain accelerators, particularly NvGPUs with sm_80 and sm_90, I have observed the following runtime exceptions in our single group radix sort:

terminate called after throwing an instance of 'sycl::_V1::nd_range_error'
  what():  Exceeded the number of registers available on the hardware.
    The number registers per work-group cannot exceed 65536 for this kernel on this device.
    The kernel uses 144 registers per work-item for a total of 512 work-items per work-group.
 -54 (PI_ERROR_INVALID_WORK_GROUP_SIZE) 

In the cases causing this crash, we attempt to use a sub-group size of 16 and pass an empty attribute to the kernel for compilation targets that we know do not have this support. The sub-group size of 16 gives a greater portion of register space to each work-item for use. The largest two single work-group case results in register spills on common hardware that does not support size 16 subgroups. I have added a check to query the device's subgroup sizes and skip this case when it occurs.

[danhoeflinger]

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?

[mmichel11]

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.

[danhoeflinger]

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.

[mmichel11]

I have updated the comment

[danhoeflinger]

LGTM

[SergeyKopienko]

LGTM