[SYCL] Get pointer should not be called on host for local accessors (…

…#13747)

Local accessors only have their memory allocated on device when a kernel
starts executing. Any pointer that refers to local memory on host before
kernel execution has begun is therefore invalid.

> This function may only be called from within a [SYCL kernel
function](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sycl-kernel-function).

Table 78 in

https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_interface_for_local_accessors

sycl/include/sycl/accessor.hpp

@@ -2589,11 +2589,21 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor
   __SYCL2020_DEPRECATED(
       "local_accessor::get_pointer() is deprecated, please use get_multi_ptr()")
   local_ptr<DataT> get_pointer() const noexcept {
+#ifndef __SYCL_DEVICE_ONLY__
+    throw sycl::exception(
+        make_error_code(errc::invalid),
+        "get_pointer must not be called on the host for a local accessor");
+#endif
     return local_ptr<DataT>(local_acc::getQualifiedPtr());
   }
 
   template <access::decorated IsDecorated>
   accessor_ptr<IsDecorated> get_multi_ptr() const noexcept {
+#ifndef __SYCL_DEVICE_ONLY__
+    throw sycl::exception(
+        make_error_code(errc::invalid),
+        "get_multi_ptr must not be called on the host for a local accessor");
+#endif
     return accessor_ptr<IsDecorated>(local_acc::getQualifiedPtr());
   }
 

sycl/unittests/accessor/LocalAccessorDefaultCtor.cpp

@@ -29,18 +29,3 @@ TEST(LocalAccessorDefaultCtorTest, LocalAcessorDefaultCtorSizeQueries) {
   EXPECT_TRUE(size == 0);
   EXPECT_TRUE(max_size == 0);
 }
-
-TEST(LocalAccessorDefaultCtorTest, LocalAcessorDefaultCtorPtrQueries) {
-  AccT acc;
-
-  // The return values of get_pointer() and get_multi_ptr() are
-  // unspecified. Just check they can run without any issue.
-  auto ptr = acc.get_pointer();
-  (void)ptr;
-  auto multi_ptr = acc.get_multi_ptr<access::decorated::yes>();
-  (void)multi_ptr;
-  auto multi_ptr_no_decorated = acc.get_multi_ptr<access::decorated::no>();
-  (void)multi_ptr_no_decorated;
-  auto multi_ptr_legacy = acc.get_multi_ptr<access::decorated::legacy>();
-  (void)multi_ptr_legacy;
-}