uxlfoundation / oneAPI-spec

oneAPI Specification source files
https://spec.oneapi.com
Other
189 stars 109 forks source link

Subgroup size for subgroup primitives #118

Open sravanikonda opened 4 years ago

sravanikonda commented 4 years ago

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 commented 4 years ago

@pennycook any thoughts on direction here?

Pennycook commented 4 years ago

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: https://github.com/intel/llvm/pull/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.