[SYCL][ESIMD] Support get_multi_ptr in ESIMD context (#10638)

accessor::get_pointer() is now deprecated in favor of get_multi_ptr.
This also fixes a number of tests that use deprecated get_pointer:
accessor::get_pointer() has been changed to conform to a recent change
to the SYCL specification, where it now returns multi_ptr rather than a
raw pointer, where the result of get_pointer is being cast to uintptr_t,
and there is no conversion operator from multi_ptr, so it's necessary to
retrieve the raw pointer first. get_multi_ptr is used instead.

Author: turinevgeny

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -39,6 +39,7 @@ static const char *LegalSYCLFunctions[] = {
     "^sycl::_V1::local_accessor<.+>::local_accessor",
     "^sycl::_V1::local_accessor<.+>::__init_esimd",
     "^sycl::_V1::local_accessor<.+>::get_pointer",
+    "^sycl::_V1::local_accessor<.+>::get_multi_ptr",
     "^sycl::_V1::local_accessor_base<.+>::local_accessor_base",
     "^sycl::_V1::local_accessor_base<.+>::__init_esimd",
     "^sycl::_V1::local_accessor_base<.+>::getQualifiedPtr",
@@ -67,6 +68,7 @@ static const char *LegalSYCLFunctions[] = {
 
 static const char *LegalSYCLFunctionsInStatelessMode[] = {
     "^sycl::_V1::accessor<.+>::get_pointer.+",
+    "^sycl::_V1::accessor<.+>::get_multi_ptr.+",
     "^sycl::_V1::accessor<.+>::getPointerAdjusted",
     "^sycl::_V1::accessor<.+>::getTotalOffset",
     "^sycl::_V1::accessor<.+>::getLinearIndex",

sycl/test-e2e/InvokeSimd/Regression/nbarrier_exec_in_order.cpp

@@ -53,7 +53,9 @@ ESIMD_INLINE void ESIMD_CALLEE_nbarrier(local_accessor<int, 1> local_acc,
 
   esimd::simd<int, VL> val(local_id);
 
-  unsigned int slm_base = static_cast<uint32_t>(reinterpret_cast<std::uintptr_t>(local_acc.get_pointer()));
+  unsigned int slm_base =
+      static_cast<uint32_t>(reinterpret_cast<std::uintptr_t>(
+          local_acc.get_multi_ptr<access::decorated::no>().get_raw()));
 
   /* Each thread operates on a region of memory (global or slm) that overlaps
    * with that of the previous and next threads.

sycl/test-e2e/InvokeSimd/Regression/nbarrier_loop.cpp

@@ -60,8 +60,9 @@ ESIMD_INLINE void ESIMD_CALLEE_nbarrier(local_accessor<int, 1> local_acc,
   unsigned int prods[2] = {2, 1};                // number of producers
 
   // Producer writes to SLM, consumer reads what producer wrote.
-  unsigned int slm_base = static_cast<uint32_t>(
-      reinterpret_cast<std::uintptr_t>(local_acc.get_pointer()));
+  unsigned int slm_base =
+      static_cast<uint32_t>(reinterpret_cast<std::uintptr_t>(
+          local_acc.get_multi_ptr<access::decorated::no>().get_raw()));
 
   esimd::barrier();
 

sycl/test-e2e/InvokeSimd/Regression/nbarrier_multiple_wg.cpp

@@ -64,8 +64,8 @@ ESIMD_INLINE void ESIMD_CALLEE_nbarrier(local_accessor<int, 1> local_acc,
   unsigned int global_off = VL * (is_producer ? global_id : (global_id - 1));
 
   unsigned int slm_base =
-      static_cast<uint32_t>(
-          reinterpret_cast<std::uintptr_t>(local_acc.get_pointer()));
+      static_cast<uint32_t>(reinterpret_cast<std::uintptr_t>(
+          local_acc.get_multi_ptr<access::decorated::no>().get_raw()));
   unsigned int slm_off = slm_base + global_off * sizeof(int);
 
   esimd::barrier();

sycl/test/esimd/esimd_verify.cpp

@@ -10,6 +10,7 @@ using namespace sycl;
 using namespace sycl::ext::intel::esimd;
 
 // CHECK-NEGATIVE-DAG: error: function 'sycl::_V1::multi_ptr<{{.+}}> sycl::_V1::accessor<{{.+}}>::get_pointer<{{.+}}>() const' is not supported in ESIMD context
+// CHECK-NEGATIVE-DAG: error: function 'std::conditional<true, sycl::_V1::multi_ptr{{.+}}>::type sycl::_V1::accessor<{{.+}}>::get_multi_ptr<{{.+}}>() const' is not supported in ESIMD context
 // CHECK-NEGATIVE-DAG: error: function '{{.+}} sycl::_V1::accessor<{{.+}}>::operator[]<{{.+}}>({{.+}}) const' is not supported in ESIMD context
 // CHECK-NEGATIVE-DAG: error: function '{{.+}}combine(int const&)' is not supported in ESIMD context
 
@@ -21,13 +22,19 @@ test0(accessor<int, 1, access::mode::read_write, access::target::device> &acc)
   return acc.get_pointer();
 }
 
-SYCL_EXTERNAL void
+SYCL_EXTERNAL auto
 test1(accessor<int, 1, access::mode::read_write, access::target::device> &acc)
     SYCL_ESIMD_FUNCTION {
+  return acc.get_multi_ptr<access::decorated::no>();
+}
+
+SYCL_EXTERNAL void
+test2(accessor<int, 1, access::mode::read_write, access::target::device> &acc)
+    SYCL_ESIMD_FUNCTION {
   acc[0] = 0;
 }
 
-void test2(sycl::handler &cgh, int *buf) {
+void test3(sycl::handler &cgh, int *buf) {
   auto reduction = sycl::reduction(buf, sycl::plus<int>());
   cgh.parallel_for<class Test2>(sycl::range<1>(1), reduction,
                                 [=](sycl::id<1>, auto &reducer)