[SYCL] Made access::decorated::no default for get_multi_ptr.

[JackAKirk]

At the moment users always need to specify the decorated mode (on/off) for get_multi_ptr(), which is not ideal, since it is an unnecessary complication since in common usage SYCL programmers do not need to worry about access::decorated. From what I can tell it seems that it was intended to have a default value access::decorated::no.

[YuriPlyakhin]

What are the arguments to set default to "access::decorated::no" as opposed to "access::decorated::yes"?

[JackAKirk]

What are the arguments to set default to "access::decorated::no" as opposed to "access::decorated::yes"?

As I understand it in most cases decorated::yes is unused, and apparently decorated::yes can have undesirable secondary side effects (apparently some projects don't support it).
As I understand it the only difference is decorated::yes is passing the address space information to C++ code, which only has niche applications.

CC @ProGTX in case there is anything more to add to this.

[YuriPlyakhin]

@JackAKirk, do you have specific example of when decorated::yes can have undesirable secondary side effect and what is the effect?

[dkhaldi]

The joint matrix tests were changed in #8874 from:
accC.get_pointer()
to:
accC.template get_multi_ptraccess::decorated::no().

Should we change these to remove the template argument as access::decorated::no is the default now? This will make the syntax more concise.

[JackAKirk]

@JackAKirk, do you have specific example of when decorated::yes can have undesirable secondary side effect and what is the effect?

I don't, but I was told that there are some examples. Point is that there should be a default whether it is yes or no because it is an unnecessary complication that the vast majority of code doesn't care about: In fact, despite understanding the general purpose of it (exposing address space to C++), I still don't know what specific examples motivated its addition in the first place to the spec: it appears to be very niche! I was told that no was the obvious choice.

[JackAKirk]

The joint matrix tests were changed in #8874 from: accC.get_pointer() to: accC.template get_multi_ptraccess::decorated::no().

Should we change these to remove the template argument as access::decorated::no is the default now? This will make the syntax more concise.

Yes this is the point. I opened this PR to try to resolve this and raise a discussion to set the default. But when a default is chosen we should definitely update tests etc to make them more concise. If Da Vinci read these tests I don't think he would be happy: the quote is "simplicity is the ultimate sophistication".

[Naghasan]

@JackAKirk, do you have specific example of when decorated::yes can have undesirable secondary side effect and what is the effect?

break badly with template / type traits. This doesn't work for instance.

template<typename T>
void foo(T* p) {
  static_assert(std::numeric_limits<T>::is_integer);
}

[...]

auto* ptr = accessor_to_int.get_multi_ptr<access::decorated::yes>();
foo(ptr);

With this implementation that is, decorated pointers are entirely implementation defined so it is also bad for portability (they may not be decorated at all actually).

[MrSidims]

According to the SYCL 2020 (at least the one I'm reading) there is no default value for the method, neither ::yes nor ::no. So why do we even have it in our APIs?

[dkhaldi]

@JackAKirk, do you have specific example of when decorated::yes can have undesirable secondary side effect and what is the effect?

I don't, but I was told that there are some examples. Point is that there should be a default whether it is yes or no because it is an unnecessary complication that the vast majority of code doesn't care about: In fact, despite understanding the general purpose of it (exposing address space to C++), I still don't know what specific examples motivated its addition in the first place to the spec: it appears to be very niche! I was told that no was the obvious choice.

@JackAKirk if decorated is no, does it mean we are not passing address space information to SPIRV and underlying layers?
address space is important piece of information, so the default should be yes if we were going to have a default.

[Naghasan]

The decoration flag is only controlling the exposed interface, not what the multi_ptr holds. Its sole purpose is to portably carry that address space around.

[JackAKirk]

@JackAKirk, do you have specific example of when decorated::yes can have undesirable secondary side effect and what is the effect?

I don't, but I was told that there are some examples. Point is that there should be a default whether it is yes or no because it is an unnecessary complication that the vast majority of code doesn't care about: In fact, despite understanding the general purpose of it (exposing address space to C++), I still don't know what specific examples motivated its addition in the first place to the spec: it appears to be very niche! I was told that no was the obvious choice.

@JackAKirk if decorated is no, does it mean we are not passing address space information to SPIRV and underlying layers? address space is important piece of information, so the default should be yes if we were going to have a default.

I see here that the spirv functions take the decorated pointer type, but also that you seem to support the multi_ptr decorated::no case because you call sycl::detail::getDecorated : https://github.com/intel/llvm/pull/9244/files#diff-e5457c55cc4c6ae06dcc828220b499fd4e55a1e3f7755a27c34da20f636c86ceR83

Possibly I am misunderstanding things?

I was initially thrown by this: #9496 but then I realized that the cuda implementation doesn't need decorated pointer at all. Cuda implementation is loading the correct address space ptx without #9499.

[YuriPlyakhin]

I verified on a Joint Matrix test case that even if I set decorated::no, I still see correct address space in IGC LLVM IR, so default "no" should not impact Joint Matrix in IGC. But does this change (adding default) conform to SYCL 2020 standard? I second Dmitry's question about introducing non-standard API.

[dm-vodopyanov]

According to the SYCL 2020 (at least the one I'm reading) there is no default value for the method, neither ::yes nor ::no. So why do we even have it in our APIs?

But does this change (adding default) conform to SYCL 2020 standard? I second Dmitry's question about introducing non-standard API.

+1

This goes against the current SYCL 2020 specification (Table 57). You may ask specification writers (e.g., @jbrodman and @gmlueck) to make this change in the spec or create an issue against spec. If the spec change will be approved and merged, the current patch with implementation can be applied.

[JackAKirk]

According to the SYCL 2020 (at least the one I'm reading) there is no default value for the method, neither ::yes nor ::no. So why do we even have it in our APIs?

But does this change (adding default) conform to SYCL 2020 standard? I second Dmitry's question about introducing non-standard API.

+1

This goes against the current SYCL 2020 specification (Table 57). You may ask specification writers (e.g., @jbrodman and @gmlueck) to make this change in the spec or create an issue against spec. If the spec change will be approved and merged, the current patch with implementation can be applied.

OK I've made a sycl-docs issue.

[JackAKirk]

closed by #10174