From 4592e773db53185567e85775fc6186c383b8fc5f Mon Sep 17 00:00:00 2001
From: Alexey Sachkov <alexey.sachkov@intel.com>
Date: Tue, 17 Jun 2025 11:07:33 +0200
Subject: [PATCH] [SYCL] Fix `multi_ptr::prefetch` regression

The regression was introduced in intel/llvm#18839 and the error looks
like:

```
/include/sycl/multi_ptr.hpp:397:5: error: no matching function for call to '__spirv_ocl_prefetch'
  397 |     __spirv_ocl_prefetch(reinterpret_cast<ptr_t>(get_decorated()), NumBytes);
/include/sycl/multi_ptr.hpp:397:5: note: candidate function not viable: no known conversion from 'ptr_t' (aka 'const __global char *') to 'const __global signed char *' for 1st argument
  397 |     __spirv_ocl_prefetch(reinterpret_cast<ptr_t>(get_decorated()), NumBytes);
```

We previously switched to use `__spirv_ocl_prefetch` that is
automatically defined by the compiler instead of having its
forward-declaration in our headers. However, compiler-provided
declaration does not define an overload for plain `char` for some
reason.

Changing that may trigger some unknown and unwanted side effects, so for
the meantime the problem is worked around by using `unsigned char`
overload instead which should be safe (from strict aliasing rules point
of view).
---
 sycl/include/sycl/multi_ptr.hpp            | 14 ++++++++++----
 sycl/test/regression/multiptr-prefetch.cpp | 16 ++++++++++++++++
 2 files changed, 26 insertions(+), 4 deletions(-)
 create mode 100644 sycl/test/regression/multiptr-prefetch.cpp

diff --git a/sycl/include/sycl/multi_ptr.hpp b/sycl/include/sycl/multi_ptr.hpp
index 824a1470b2c9b..686d0a629a4a9 100644
--- a/sycl/include/sycl/multi_ptr.hpp
+++ b/sycl/include/sycl/multi_ptr.hpp
@@ -391,10 +391,13 @@ class __SYCL_TYPE(multi_ptr) multi_ptr {
       access::address_space _Space = Space,
       typename = typename std::enable_if_t<
           _Space == Space && Space == access::address_space::global_space>>
-  void prefetch(size_t NumElements) const {
+  void prefetch([[maybe_unused]] size_t NumElements) const {
+#ifdef __SYCL_DEVICE_ONLY__
     size_t NumBytes = NumElements * sizeof(ElementType);
-    using ptr_t = typename detail::DecoratedType<char, Space>::type const *;
+    using ptr_t =
+        typename detail::DecoratedType<unsigned char, Space>::type const *;
     __spirv_ocl_prefetch(reinterpret_cast<ptr_t>(get_decorated()), NumBytes);
+#endif
   }
 
   // Arithmetic operators
@@ -1087,10 +1090,13 @@ class multi_ptr<ElementType, Space, access::decorated::legacy> {
       access::address_space _Space = Space,
       typename = typename std::enable_if_t<
           _Space == Space && Space == access::address_space::global_space>>
-  void prefetch(size_t NumElements) const {
+  void prefetch([[maybe_unused]] size_t NumElements) const {
+#ifdef __SYCL_DEVICE_ONLY__
     size_t NumBytes = NumElements * sizeof(ElementType);
-    using ptr_t = typename detail::DecoratedType<char, Space>::type const *;
+    using ptr_t =
+        typename detail::DecoratedType<unsigned char, Space>::type const *;
     __spirv_ocl_prefetch(reinterpret_cast<ptr_t>(m_Pointer), NumBytes);
+#endif
   }
 
 private:
diff --git a/sycl/test/regression/multiptr-prefetch.cpp b/sycl/test/regression/multiptr-prefetch.cpp
new file mode 100644
index 0000000000000..281a2d2db09d1
--- /dev/null
+++ b/sycl/test/regression/multiptr-prefetch.cpp
@@ -0,0 +1,16 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s
+// expected-no-diagnostics
+
+#include <sycl/multi_ptr.hpp>
+
+SYCL_EXTERNAL void
+foo(sycl::multi_ptr<int, sycl::access::address_space::global_space> mptr) {
+  mptr.prefetch(0);
+}
+
+SYCL_EXTERNAL void
+bar(sycl::multi_ptr<int, sycl::access::address_space::global_space,
+                    sycl::access::decorated::legacy>
+        mptr) {
+  mptr.prefetch(0);
+}