Subgroup size for subgroup primitives

[sravanikonda]

Querying sycl::info::device::sub_group_size gives several numbers. For example, we get 8, 16 and 32 for Gen9. We would like to specify the sub group size and this feature is supported. All three sizes seem to work except that subgroup primitives such as shuffle_down do not work for all sizes. By try and error, we have found that shuffle_down works for 16. Could oneAPI provide a query function for returning the "primitive" subgroup size?

[mkinsner]

@Pennycook any thoughts on direction here?

[Pennycook]

Sorry for the late response here. There was a lot of discussion about how to expose this functionality, and whether additional guarantees were required to make it useful.

Proposed changes to the sub-group class are being reviewed here: intel/llvm#2452. @sravanikonda , could you please take a look and confirm that the changes are sufficient?

I think that the "core" sub-group size is equivalent to the requested "primitive" sub-group size. The other changes clarify that not all features are supported by all sub-group sizes, and make it easy for developers to use a sub-group size that is guaranteed to work for all of their kernels.