diff --git a/sycl/include/sycl/multi_ptr.hpp b/sycl/include/sycl/multi_ptr.hpp index 824a1470b2c9..686d0a629a4a 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::type const *; + using ptr_t = + typename detail::DecoratedType::type const *; __spirv_ocl_prefetch(reinterpret_cast(get_decorated()), NumBytes); +#endif } // Arithmetic operators @@ -1087,10 +1090,13 @@ class 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::type const *; + using ptr_t = + typename detail::DecoratedType::type const *; __spirv_ocl_prefetch(reinterpret_cast(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 000000000000..281a2d2db09d --- /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_EXTERNAL void +foo(sycl::multi_ptr mptr) { + mptr.prefetch(0); +} + +SYCL_EXTERNAL void +bar(sycl::multi_ptr + mptr) { + mptr.prefetch(0); +}