From 7698ff8f49e8929675ee3b2a5b363f5732937923 Mon Sep 17 00:00:00 2001
From: Hugh Delaney <hugh.delaney@codeplay.com>
Date: Tue, 14 May 2024 11:13:05 +0100
Subject: [PATCH] [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                    | 10 ++++++++++
 .../accessor/LocalAccessorDefaultCtor.cpp         | 15 ---------------
 2 files changed, 10 insertions(+), 15 deletions(-)

diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 83a05dd1e9d9d..af128aa312410 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/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());
   }
 
diff --git a/sycl/unittests/accessor/LocalAccessorDefaultCtor.cpp b/sycl/unittests/accessor/LocalAccessorDefaultCtor.cpp
index 912b5e6a82d53..7e1b6a9e22f74 100644
--- a/sycl/unittests/accessor/LocalAccessorDefaultCtor.cpp
+++ b/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;
-}