[SYCL] Fix multi_ptr::prefetch regression (#19013)

The regression was introduced in #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).

Author: AlexeySachkov

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:

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);
+}