Best practices for handling variable width sub_group sizes?

allanmac1 — Wed, 24 Aug 2016 17:35:25 GMT

I wrote the following back in May 2015:

It would be great to get some sort of compile-time indication or guarantee that a certain size subgroup was selected in order to close off certain code paths.

If only the compiler can determine the subgroup size and your codebase is dependent on a certain subgroup size and variations in the code result in changes to the subgroup size then ... it becomes a circular mess.

There is always the option of probing the subgroup size at runtime with get_sub_group_size() but I would caution against that approach. It may be portable but it's probably not performant.

But if get_sub_group_size() was a compile time constant resolved early enough to snip dead code then it would provide portability and performance.

It would be great to hear how Intel actually implements this super useful extension.

Is there any new info on how and when sub_group sizes are determined in Intel Processor Graphics OpenCL and whether or not a sub_group size can ever be an optimizable compile time constant (maybe it can't if get_sub_group_size() and get_max_sub_group_size() are not always equal)?

As implied above, Intel GEN appears to be unique in allowing the compiler to determine the "warp" (SIMD) width at compile time.

A small change to a kernel might result in switching from SIMD8 to SIMD16 sub groups.

My experience so far is that this flexibility forces a defensive assumption that SIMD8 is the widest sub_group that can be relied upon and that various work item id calculations have to be performed explicitly: e.g. "get_local_id(0) & 7" instead of "get_sub_group_local_id()".

Some tips and tricks on this subject would be appreciated!

Indeed.

Timothy_B_Intel — Mon, 29 Aug 2016 21:31:46 GMT

Indeed.

https://www.khronos.org/registry/cl/extensions/intel/cl_intel_required_subgroup_size.txt

There is a minimum HW requirement, so I am not sure which platform it becomes available.

Aww yiss! That looks perfect

allanmac1 — Mon, 29 Aug 2016 21:47:32 GMT

Aww yiss! That looks perfect.

topic Best practices for handling variable width sub_group sizes? in OpenCL* for CPU

Best practices for handling variable width sub_group sizes?

Indeed.

Aww yiss! That looks perfect