[SYCL-E2E][group local memory] question about access::decorated::legacy

[jinz2014]

Need tests that do not specify access::decorated::legacy be added ?
Is access::decorated::legacy deprecated ?
Both tests contain "XFAIL: hip_nvidia", but I didn't see any compile/run failure for the tests.

[AlexeySachkov]

Hi @jinz2014,

Is access::decorated::legacy deprecated ?

Yes, access::decorated::legacy is deprecated in SYCL-2020, see 4.7.7.1. Multi-pointer class

Need tests that do not specify access::decorated::legacy be added ?

Right. Even though the the extension spec spec does not allow to specify whether or not you want a decorated multi_ptr and relies on default legacy (tagging @Pennycook to see if we want to change that), it should still be possible to convert legacy-decorated multi_ptr into yes-/no-decorated multi_ptr if I read the SYCL specification correctly.

In particular, the following methods should be available:

  // Available only when:
  //   (Space == access::address_space::generic_space &&
  //    AS != access::address_space::constant_space)
  template <access::address_space AS, access::decorated IsDecorated>
  multi_ptr& operator=(const multi_ptr<value_type, AS, IsDecorated>&);

  // Available only when:
  //   (Space == access::address_space::generic_space &&
  //    AS != access::address_space::constant_space)
  template <access::address_space AS, access::decorated IsDecorated>
  multi_ptr& operator=(multi_ptr<value_type, AS, IsDecorated>&&);

Example usage:

void foo(sycl::multi_ptr<int, sycl::access::address_space::local_space, sycl::access::decorated::legacy> mptr) {
  sycl::multi_ptr<int, sycl::access::address_space::generic_space, sycl::access::decorated::no> a(mptr);
  sycl::multi_ptr<int, sycl::access::address_space::generic_space, sycl::access::decorated::yes> b(mptr);
}

Note that it converts local_space pointer into generic_space pointer.

  // Implicit conversion to multi_ptr<const value_type, Space>.
  template <access::decorated IsDecorated>
  operator multi_ptr<const value_type, Space, IsDecorated>() const;

This implicit conversion doesn't make much sense with this particular extension, because you won't be able to write to such pointer, but this conversion seems to allow changing the decoration of multi_ptr. However, I can't get it work with our implementation, perhaps we have some bugs, or perhaps this is a spec bug and this conversion should now allow changing decoration type of multi_ptr (tagging @gmlueck here).

Both tests contain "XFAIL: hip_nvidia", but I didn't see any compile/run failure for the tests.

We have not been diligent enough when disabling some tests and did not create any issues for us to track such disabled tests. Therefore, some of them have been disabled for years, even though everything could be ok now. We are aware of this problem and hopefully we will do a review of XFAILed tests soon to have every one of them covered with a tracker for us to analyze.

Also, due to the way of how our infrastructure works at the moment, we treat such test as UNSUPPORTED and don't run them at all. Otherwise, we would have seen that they pass and removed XFAIL.

llvm/sycl/test-e2e/format.py

Lines 129 to 130 in 9d22e38

	# Treat XFAIL as UNSUPPORTED if the test is to be executed on multiple
	# devices.

[gmlueck]

This implicit conversion doesn't make much sense with this particular extension, because you won't be able to write to such pointer, but this conversion seems to allow changing the decoration of multi_ptr. However, I can't get it work with our implementation, perhaps we have some bugs, or perhaps this is a spec bug and this conversion should now allow changing decoration type of multi_ptr (tagging @gmlueck here).

I think there is a spec bug here. If I understand correctly, you expect to have an implicit conversion from a legacy multi-pointer to either a yes or no decorated multi-ptr. I see that we have implicit conversion from yes to no and vice-versa:

177 // Implicit conversion to the non-decorated version of multi_ptr.
178 // Available only when: (is_decorated == true)
179 operator multi_ptr<value_type, Space, access::decorated::no>() const;
180
181 // Implicit conversion to the decorated version of multi_ptr.
182 // Available only when: (is_decorated == false)
183 operator multi_ptr<value_type, Space, access::decorated::yes>() const;

However, I don't see any similar conversion from legacy to yes or no. It seems to me that we should have such implicit conversions. Tagging @AerialMantis for his thoughts.