[SYCL][Joint Matrix] Pass on address space to Load/Store

[YuriPlyakhin]

Pass on address space information to SPIR-V Joint Matrix Load/Store intrinsics.
TODO:

• Implement the same for Legacy Matrix API
• One test fails, need to investigate
• add test to sycl/test/check_device_code - check for correct address spaces for SLM and global accessor 

[MrSidims]

LGTM
one nit: I believe we don't support constant AS for joint matrix as well (also AFAIK this AS is being deprecated in SYCL 2020)

[MrSidims]

LGTM one nit: I believe we don't support constant AS for joint matrix as well (also AFAIK this AS is being deprecated in SYCL 2020)

Discussed offline: no issues here

[YuriPlyakhin]

LGTM one nit: I believe we don't support constant AS for joint matrix as well (also AFAIK this AS is being deprecated in SYCL 2020)

As we discussed in theory constant AS to load from should be supported for Joint Matrix.

[YuriPlyakhin]

Failed Tests (1): SYCL :: Basic/vec_bool.cpp - The failed test is unrelated to my patch. Also mentioned in #9263, #8344 as failing.

[steffenlarsen]

LGTM!

[steffenlarsen]

Failed Tests (1):
SYCL :: Basic/vec_bool.cpp - This test was recently changed (#9259) and is not failing in post-commit. Failures are likely due to a desync in testing.

[JackAKirk]

@YuriPlyakhin @dkhaldi
I think there is a problem here. If I adjust your test sycl/test/check_device_code/matrix/matrix_load_store_as.cpp so that joint_matrix_load takes the tileA.get_pointer() or tileA.get_multi_ptr() from tileA local_accessor directly instead of tileA.template get_multi_ptr<sycl::access::decorated::yes>() then compiler errors:

matrixas.cpp:43:11: error: no matching function for call to 'joint_matrix_load'
          joint_matrix_load(
          ^~~~~~~~~~~~~~~~~
llvm/build/bin/../include/sycl/ext/oneapi/matrix/matrix-unified.hpp:261:1: note: candidate template ignored: could not match 'multi_ptr<T, Space, IsDecorated>' against 'std::add_pointer_t<value_type>' (aka 'unsigned short *')
joint_matrix_load(Group sg,

whereas before this patch it would have passed. The same happens in the nvptx64 backend with this patch #9499. Maybe I'm doing something wrong but this seems to give no option to just pass in a multi_ptr from an accessor where we don't specify whether or not it is decorated? Do we really want the requirement that passed pointers must specify whether they are decorated or not? Can there be a default constructor that is not decorated perhaps?

Since the implementations do not make any use of decorated pointer it seems bad that a programmer has to now pass in a decorated pointer for the local_accessor case. Note that this is not the case for a standard accessor. Local memory usage is very important for joint_matrix so it should not be penalised in this way.
In general I don't understand why the concept of pointer decoration has to be exposed to the user at all, although I am not aware of the circumstances surrounding its introduction to the spec, but this is another question.
I don't think it makes sense to merge #9499 until this issue is resolved since decorating the pointer has no advantage in the implementation since there is no specialization for different address spaces and in the existing implementation the local/global address spaces are correctly supported.

Note that this issue can be fixed with:

--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -2148,7 +2148,7 @@ public:
     return constant_ptr<DataT>(getPointerAdjusted());
   }
 
-  template <access::decorated IsDecorated>
+  template <access::decorated IsDecorated=access::decorated::no>
   accessor_ptr<IsDecorated> get_multi_ptr() const noexcept {
     return accessor_ptr<IsDecorated>(getPointerAdjusted());
   }
@@ -2916,7 +2916,7 @@ public:
     return std::add_pointer_t<value_type>(local_acc::getQualifiedPtr());
   }
 
-  template <access::decorated IsDecorated>
+  template <access::decorated IsDecorated=access::decorated::no>
   accessor_ptr<IsDecorated> get_multi_ptr() const noexcept {
     return accessor_ptr<IsDecorated>(local_acc::getQualifiedPtr());
   }