[SYCL] Implement SYCL2020 sub_group class

[alexbatashev]

Link to tests: intel/llvm-test-suite#283

[cperkinsintel]

I was planning on moving the sub_group stuff out of ONEAPI namespace for the SYCL2020 sub-group support where I'm renaming the routines. Does the existing ONEAPI/sub_group API have to be left intact?

[cperkinsintel]

Apparently yes. No problem.

[alexbatashev]

That's right. nd_item returns ONEAPI::sub_group. To comply with the spec we'd need to either change the return type (which is a breaking change), or leave it as is and just add new members.

[gmlueck]

Apologies for commenting on a merged PR, but @Pennycook and I have a concern about adding an alias from ONEAPI::sub_group into the sycl:: namespace. There are two problems we see:

1. We will probably want to deprecate the ONEAPI version of the this API at some point by adding the [[deprecated]] attribute. However, it's not clear how we can do this if sycl::sub_group is an alias. By adding [[deprecated]] to the ONEAPI version we would also deprecate the sycl version.

2. The ONEAPI::sub_group contains many member functions that are not part of the standard sycl::sub_group type. By adding an alias, we silently add all of these non-standard member functions into the standard type. This is a violation of the extension policy.

Point (2) is a big concern to us because we do not want applications to mistakenly use non-standard APIs. The extension convention was created, in part, to make it obvious when an application uses an extended API. For example, consider the following code:

sycl::sub_group sg = it.get_sub_group(); // create a standard SYCL 2020 sub-group
auto x = sg.load(ptr); // This uses a non-standard member function from sycl::ONEAPI::sub_group, but it's not obvious

Application code like that above may not realize that sg.load() is non-standard, and they will be surprised when it doesn't work on another SYCL implementation. We had customer feedback during the design of the extension policy that customers want extensions to be named in way that makes it clear that they are extensions.

I think the right way to add support for SYCL 2020 sub_group would be to add a new class sycl::sub_group that is not an alias to the ONEAPI version. The new class would either:

• Implement only the standard SYCL 2020 member functions, or
• Implement the standard member functions and add extended ones using the extension naming convention (which means they start with the prefix ext_oneapi_).

Can we reconsider the way this is implemented?

[alexbatashev]

@gmlueck I see one big problem with that code. Consider I'm a big fan of always auto. If I write code like this:

auto sg = it.get_sub_group();

What kind of behavior should I expect? If I previously used ONEAPI extension, and we update the return type, my code is broken. If I write new code and we do not update the return type, it will also be broken (although in a slightly less painful way). Either way, the fact that nd_item::get_sub_group() must return different types breaks some use cases.

Application code like that above may not realize that sg.load() is non-standard

We can add more warnings for this case.

[Pennycook]

I acknowledge that whatever we do here we'll probably break somebody's code. Unfortunately I think it's unavoidable, and it's a side-effect of the sub_group class predating the extension mechanism.

If we'd had the forethought to use it.ext_oneapi_get_sub_group() we'd have no problems. Instead, we're in a situation where SYCL 1.2.1 + SYCL_INTEL_sub_group defines a function with the same name as SYCL 2020.

Can we guard things based on SYCL language version, maybe? If somebody compiles with SYCL 1.2.1 + DPC++ extensions, it.get_sub_group() could continue to return sycl::ONEAPI::sub_group as before. But if somebody opts-in to compiling with SYCL 2020, the only reasonable expectation is that it.get_sub_group() returns a sycl::sub_group.

[gmlueck]

Yes, I agree that things would be better if we had named this it.ext_oneapi_get_sub_group(). Actually, this is a good illustration of the problems that can occur when we don't follow the extension mechanism!

I like @Pennycook's suggestion about defining it.get_sub_group() differently depending on whether the compiler is in SYCL 1.2.1 mode or SYCL 2020 mode. In addition, I'd suggest adding a new member function it.ext_intel_get_sub_group(). That member function would always return the ONEAPI::sub_group type. This provides a quick fix for any applications that use the old sub-group extension; they can just change their call from it.get_sub_group() to it.ext_intel_get_sub_group().

[cperkinsintel]

Looks good to me.

[dm-vodopyanov]

LGTM

[cperkinsintel]

LGTM

[alexbatashev]

@AlexeySachkov @Pennycook could you please review? This PR needs your approval