From 6ab79b1b4e6d5f3d0c86855bbd832dc409c1f84c Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Mon, 25 Sep 2023 10:10:45 -0700
Subject: [PATCH 01/25] Initial

---
 .../SYCLLowerIR/CompileTimePropertiesPass.cpp |   6 +
 sycl/include/sycl/ext/oneapi/prefetch.hpp     | 186 ++++++++++++++++++
 .../sycl/ext/oneapi/properties/properties.hpp |   2 +
 sycl/include/sycl/sycl.hpp                    |   1 +
 sycl/source/feature_test.hpp.in               |   1 +
 sycl/test/extensions/prefetch.cpp             |  38 ++++
 6 files changed, 234 insertions(+)
 create mode 100644 sycl/include/sycl/ext/oneapi/prefetch.hpp
 create mode 100644 sycl/test/extensions/prefetch.cpp

diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
index 7d845d9212e4..b00503879f70 100644
--- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
+++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
@@ -619,6 +619,12 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
     if (*Property.first == "sycl-alignment")
       continue;
 
+    // leave these annotations as is. They will be processed by SPIRVWriter.
+    if (*Property.first == "sycl-prefetch-hint" ||
+        *Property.first == "sycl-prefetch-hint-nt") {
+      return false;
+    }
+
     auto DecorIt = SpirvDecorMap.find(*Property.first);
     if (DecorIt == SpirvDecorMap.end())
       continue;
diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/prefetch.hpp
new file mode 100644
index 000000000000..57981f3c2a45
--- /dev/null
+++ b/sycl/include/sycl/ext/oneapi/prefetch.hpp
@@ -0,0 +1,186 @@
+//==--------------- prefetch.hpp --- SYCL prefetch extension ---------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/__spirv/spirv_ops.hpp>
+#include <sycl/ext/oneapi/properties/properties.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+
+enum class cache_level {
+  L1,
+  L2,
+  L3,
+  L4,
+};
+
+struct nontemporal;
+
+struct prefetch_hint_key {
+  template <cache_level Level, typename Hint>
+  using value_t =
+      property_value<prefetch_hint_key,
+                     std::integral_constant<cache_level, Level>, Hint>;
+};
+
+template <cache_level Level, typename Hint>
+inline constexpr prefetch_hint_key::value_t<Level, Hint> prefetch_hint;
+
+inline constexpr prefetch_hint_key::value_t<cache_level::L1, void>
+    prefetch_hint_L1;
+inline constexpr prefetch_hint_key::value_t<cache_level::L2, void>
+    prefetch_hint_L2;
+inline constexpr prefetch_hint_key::value_t<cache_level::L3, void>
+    prefetch_hint_L3;
+inline constexpr prefetch_hint_key::value_t<cache_level::L4, void>
+    prefetch_hint_L4;
+
+inline constexpr prefetch_hint_key::value_t<cache_level::L1, nontemporal>
+    prefetch_hint_L1_nt;
+inline constexpr prefetch_hint_key::value_t<cache_level::L2, nontemporal>
+    prefetch_hint_L2_nt;
+inline constexpr prefetch_hint_key::value_t<cache_level::L3, nontemporal>
+    prefetch_hint_L3_nt;
+inline constexpr prefetch_hint_key::value_t<cache_level::L4, nontemporal>
+    prefetch_hint_L4_nt;
+
+template <typename T, cache_level Level, typename Hint>
+struct is_valid_property<T, prefetch_hint_key::value_t<Level, Hint>>
+    : std::bool_constant<std::is_pointer<T>::value> {};
+
+namespace detail {
+template <> struct IsCompileTimeProperty<prefetch_hint_key> : std::true_type {};
+
+template <cache_level Level, typename Hint>
+struct PropertyMetaInfo<prefetch_hint_key::value_t<Level, Hint>> {
+  static constexpr const char *name = std::is_same_v<Hint, nontemporal>
+                                          ? "sycl-prefetch-hint-nt"
+                                          : "sycl-prefetch-hint";
+  static constexpr int value = static_cast<int>(Level);
+};
+
+template <typename Properties>
+void prefetch_impl(void *ptr, size_t bytes, Properties properties) {
+#ifdef __SYCL_DEVICE_ONLY__
+  auto *ptrGlobalAS = __SYCL_GenericCastToPtrExplicit_ToGlobal<char>(ptr);
+  __attribute__((opencl_global)) char *ptrAnnotated = nullptr;
+  if constexpr (!properties.template has_property<prefetch_hint_key>()) {
+    ptrAnnotated = __builtin_intel_sycl_ptr_annotation(
+        ptrGlobalAS, "sycl-prefetch-hint", static_cast<int>(cache_level::L1));
+  } else {
+    auto prop = properties.template get_property<prefetch_hint_key>();
+    ptrAnnotated = __builtin_intel_sycl_ptr_annotation(
+        ptrGlobalAS, PropertyMetaInfo<decltype(prop)>::name,
+        PropertyMetaInfo<decltype(prop)>::value);
+  }
+  __spirv_ocl_prefetch(ptrAnnotated, bytes);
+#endif
+}
+} // namespace detail
+
+template <typename Properties = empty_properties_t>
+void prefetch(void *ptr, Properties properties = {}) {
+  detail::prefetch_impl(ptr, 1, properties);
+}
+
+template <typename Properties = empty_properties_t>
+void prefetch(void *ptr, size_t bytes, Properties properties = {}) {
+  detail::prefetch_impl(ptr, bytes, properties);
+}
+
+template <typename T, typename Properties = empty_properties_t>
+void prefetch(T *ptr, Properties properties = {}) {
+  prefetch((void *)ptr, sizeof(T), properties);
+}
+
+template <typename T, typename Properties = empty_properties_t>
+void prefetch(T *ptr, size_t count, Properties properties = {}) {
+  prefetch((void *)ptr, count * sizeof(T), properties);
+}
+
+// Only available if AddressSpace == global_space || AddressSpace ==
+// generic_space
+template <
+    access::address_space AddressSpace, access::decorated IsDecorated,
+    typename Properties = empty_properties_t,
+    std::enable_if_t<AddressSpace == access::address_space::global_space ||
+                     AddressSpace == access::address_space::generic_space>>
+void prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr,
+              Properties properties = {}) {
+  prefetch(ptr.get(), properties);
+}
+
+// Only available if AddressSpace == global_space || AddressSpace ==
+// generic_space
+template <
+    access::address_space AddressSpace, access::decorated IsDecorated,
+    typename Properties = empty_properties_t,
+    std::enable_if_t<AddressSpace == access::address_space::global_space ||
+                     AddressSpace == access::address_space::generic_space>>
+void prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, size_t bytes,
+              Properties properties = {}) {
+  prefetch(ptr.get(), bytes, properties);
+}
+
+// Only available if AddressSpace == global_space || AddressSpace ==
+// generic_space
+template <
+    typename T, access::address_space AddressSpace,
+    access::decorated IsDecorated, typename Properties = empty_properties_t,
+    std::enable_if_t<AddressSpace == access::address_space::global_space ||
+                     AddressSpace == access::address_space::generic_space>>
+void prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr,
+              Properties properties = {}) {
+  prefetch(ptr.get(), properties);
+}
+
+// Only available if AddressSpace == global_space || AddressSpace ==
+// generic_space
+template <
+    typename T, access::address_space AddressSpace,
+    access::decorated IsDecorated, typename Properties = empty_properties_t,
+    std::enable_if_t<AddressSpace == access::address_space::global_space ||
+                     AddressSpace == access::address_space::generic_space>>
+void prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count,
+              Properties properties = {}) {
+  prefetch(ptr.get(), count, properties);
+}
+
+// Only available if Dimensions > 0 && (AccessMode == read || AccessMode ==
+// read_write)
+template <typename DataT, int Dimensions, access_mode AccessMode,
+          access::placeholder IsPlaceholder,
+          typename Properties = empty_properties_t,
+          std::enable_if_t<(Dimensions > 0) &&
+                           (AccessMode == access_mode::read ||
+                            AccessMode == access_mode::read_write)>>
+void prefetch(
+    accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
+    id<Dimensions> offset, Properties properties = {}) {
+  prefetch((void *)&acc[offset], sizeof(DataT), properties);
+}
+
+// Only available if Dimensions > 0 && (AccessMode == read || AccessMode ==
+// read_write)
+template <typename DataT, int Dimensions, access_mode AccessMode,
+          access::placeholder IsPlaceholder,
+          typename Properties = empty_properties_t,
+          std::enable_if_t<(Dimensions > 0) &&
+                           (AccessMode == access_mode::read ||
+                            AccessMode == access_mode::read_write)>>
+void prefetch(
+    accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
+    size_t offset, size_t count, Properties properties = {}) {
+  prefetch((void *)&acc[offset], count * sizeof(DataT), properties);
+}
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/include/sycl/ext/oneapi/properties/properties.hpp b/sycl/include/sycl/ext/oneapi/properties/properties.hpp
index bb89db3e9335..9c361cd63b05 100644
--- a/sycl/include/sycl/ext/oneapi/properties/properties.hpp
+++ b/sycl/include/sycl/ext/oneapi/properties/properties.hpp
@@ -187,6 +187,8 @@ properties(PropertyValueTs... props)
     -> properties<typename detail::Sorted<PropertyValueTs...>::type>;
 #endif
 
+using empty_properties_t = decltype(properties{});
+
 // Property list traits
 template <typename propertiesT> struct is_property_list : std::false_type {};
 template <typename... PropertyValueTs>
diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp
index 21d0d9015de0..0ea5f5c1b240 100644
--- a/sycl/include/sycl/sycl.hpp
+++ b/sycl/include/sycl/sycl.hpp
@@ -86,6 +86,7 @@
 #include <sycl/ext/oneapi/kernel_properties/properties.hpp>
 #include <sycl/ext/oneapi/matrix/matrix.hpp>
 #include <sycl/ext/oneapi/owner_less.hpp>
+#include <sycl/ext/oneapi/prefetch.hpp>
 #include <sycl/ext/oneapi/properties/properties.hpp>
 #include <sycl/ext/oneapi/properties/property_value.hpp>
 #include <sycl/ext/oneapi/sub_group.hpp>
diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in
index 30a76930ca36..7ae92e675c31 100644
--- a/sycl/source/feature_test.hpp.in
+++ b/sycl/source/feature_test.hpp.in
@@ -88,6 +88,7 @@ inline namespace _V1 {
 #define SYCL_EXT_CODEPLAY_MAX_REGISTERS_PER_WORK_GROUP_QUERY 1
 #define SYCL_EXT_ONEAPI_DEVICE_GLOBAL 1
 #define SYCL_EXT_INTEL_QUEUE_IMMEDIATE_COMMAND_LIST 1
+#define SYCL_EXT_ONEAPI_PREFETCH 1
 
 #ifndef __has_include
 #define __has_include(x) 0
diff --git a/sycl/test/extensions/prefetch.cpp b/sycl/test/extensions/prefetch.cpp
new file mode 100644
index 000000000000..9f9620db1512
--- /dev/null
+++ b/sycl/test/extensions/prefetch.cpp
@@ -0,0 +1,38 @@
+// RUN: %clangxx -fsycl-device-only -S %s -o - | FileCheck %s
+
+#include <sycl/sycl.hpp>
+
+char data[] = {0, 1, 2, 3};
+
+// CHECK: [[PREFETCH_STR:@.*]] = private unnamed_addr addrspace(1) constant [19 x i8] c"sycl-prefetch-hint\00", section "llvm.metadata"
+// CHECK: [[PREFETCH_LVL0:@.*]] = private unnamed_addr addrspace(1) constant [2 x i8] c"0\00", section "llvm.metadata"
+// CHECK: [[ANNOTATION1:@.*]] = private unnamed_addr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) [[PREFETCH_STR]], ptr addrspace(1) [[PREFETCH_LVL0]] }, section "llvm.metadata"
+// CHECK: [[PREFETCH_LVL1:@.*]] = private unnamed_addr addrspace(1) constant [2 x i8] c"1\00", section "llvm.metadata"
+// CHECK: [[ANNOTATION2:@.*]] = private unnamed_addr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) [[PREFETCH_STR]], ptr addrspace(1) [[PREFETCH_LVL1]] }, section "llvm.metadata"
+// CHECK: [[PREFETCH_STR_NT:@.*]] = private unnamed_addr addrspace(1) constant [22 x i8] c"sycl-prefetch-hint-nt\00", section "llvm.metadata"
+// CHECK: [[PREFETCH_LVL2:@.*]] = private unnamed_addr addrspace(1) constant [2 x i8] c"2\00", section "llvm.metadata"
+// CHECK: [[ANNOTATION3:@.*]] = private unnamed_addr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) [[PREFETCH_STR_NT]], ptr addrspace(1) [[PREFETCH_LVL2]] }, section "llvm.metadata"
+
+int main() {
+  namespace syclex = sycl::ext::oneapi::experimental;
+  sycl::queue q;
+  void *dataPtr = &data;
+  q.parallel_for(1, [=](sycl::id<1> idx) {
+    // CHECK: [[CASTED:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}
+
+    // CHECK: [[ANNOTATED1:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], ptr addrspace(1) {{.*}}, ptr addrspace(1) {{.*}}, i32 76, ptr addrspace(1) [[ANNOTATION1]])
+    // CHECK: tail call spir_func void @_Z20__spirv_ocl_prefetch{{.*}}(ptr addrspace(1) noundef [[ANNOTATED1]], i64 noundef 1)
+    syclex::prefetch(dataPtr);
+
+    // CHECK: [[ANNOTATED2:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], ptr addrspace(1) {{.*}}, ptr addrspace(1) {{.*}}, i32 80, ptr addrspace(1) [[ANNOTATION2]])
+    // CHECK: tail call spir_func void @_Z20__spirv_ocl_prefetch{{.*}}(ptr addrspace(1) noundef [[ANNOTATED2]], i64 noundef 1)
+    syclex::prefetch(dataPtr, syclex::properties{syclex::prefetch_hint_L2});
+
+    // CHECK: [[ANNOTATED3:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], ptr addrspace(1){{.*}}, ptr addrspace(1) {{.*}}, i32 80, ptr addrspace(1) [[ANNOTATION3]])
+    // CHECK: tail call spir_func void @_Z20__spirv_ocl_prefetch{{.*}}(ptr addrspace(1) noundef [[ANNOTATED3]], i64 noundef 4)
+    syclex::prefetch(dataPtr, 4, syclex::properties{syclex::prefetch_hint_L3_nt});
+  });
+  q.wait();
+
+  return 0;
+}

From c4123e9e42274f8967de09bf0179e6d9ab0a5091 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Mon, 9 Oct 2023 01:43:13 -0700
Subject: [PATCH 02/25] clang format

---
 sycl/test/extensions/prefetch.cpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/sycl/test/extensions/prefetch.cpp b/sycl/test/extensions/prefetch.cpp
index 9f9620db1512..64969d85c342 100644
--- a/sycl/test/extensions/prefetch.cpp
+++ b/sycl/test/extensions/prefetch.cpp
@@ -30,7 +30,8 @@ int main() {
 
     // CHECK: [[ANNOTATED3:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], ptr addrspace(1){{.*}}, ptr addrspace(1) {{.*}}, i32 80, ptr addrspace(1) [[ANNOTATION3]])
     // CHECK: tail call spir_func void @_Z20__spirv_ocl_prefetch{{.*}}(ptr addrspace(1) noundef [[ANNOTATED3]], i64 noundef 4)
-    syclex::prefetch(dataPtr, 4, syclex::properties{syclex::prefetch_hint_L3_nt});
+    syclex::prefetch(dataPtr, 4,
+                     syclex::properties{syclex::prefetch_hint_L3_nt});
   });
   q.wait();
 

From 549f9e885b198a7ebf7a27af21a3d31be16405b8 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Mon, 9 Oct 2023 04:43:27 -0700
Subject: [PATCH 03/25] Add joint_prefetch

---
 sycl/include/sycl/ext/oneapi/prefetch.hpp | 119 ++++++++++++++++++++++
 1 file changed, 119 insertions(+)

diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/prefetch.hpp
index 57981f3c2a45..0dc081c9f711 100644
--- a/sycl/include/sycl/ext/oneapi/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/prefetch.hpp
@@ -181,6 +181,125 @@ void prefetch(
     size_t offset, size_t count, Properties properties = {}) {
   prefetch((void *)&acc[offset], count * sizeof(DataT), properties);
 }
+
+template <typename Group, typename Properties = empty_properties_t>
+typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
+joint_prefetch(Group g, void *ptr, Properties properties = {}) {
+  detail::prefetch_impl(ptr, 1, properties);
+}
+
+template <typename Group, typename Properties = empty_properties_t>
+typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
+joint_prefetch(Group g, void *ptr, size_t bytes, Properties properties = {}) {
+  detail::prefetch_impl(ptr, bytes, properties);
+}
+
+template <typename Group, typename T, typename Properties = empty_properties_t>
+typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
+joint_prefetch(Group g, T *ptr, Properties properties = {}) {
+  joint_prefetch((void *)ptr, sizeof(T), properties);
+}
+
+template <typename Group, typename T, typename Properties = empty_properties_t>
+typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
+joint_prefetch(Group g, T *ptr, size_t count, Properties properties = {}) {
+  joint_prefetch((void *)ptr, count * sizeof(T), properties);
+}
+
+// Only available if AddressSpace == global_space || AddressSpace ==
+// generic_space
+template <typename Group, access::address_space AddressSpace,
+          access::decorated IsDecorated,
+          typename Properties = ext::oneapi::experimental::empty_properties_t>
+typename std::enable_if_t<
+    sycl::is_group_v<std::decay_t<Group>> &&
+        (AddressSpace == access::address_space::global_space ||
+         AddressSpace == access::address_space::generic_space),
+    void>
+joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
+               Properties properties = {}) {
+  joint_prefetch(g, ptr.get(), properties);
+}
+
+// Only available if AddressSpace == global_space || AddressSpace ==
+// generic_space
+template <typename Group, access::address_space AddressSpace,
+          access::decorated IsDecorated,
+          typename Properties = ext::oneapi::experimental::empty_properties_t>
+typename std::enable_if_t<
+    sycl::is_group_v<std::decay_t<Group>> &&
+        (AddressSpace == access::address_space::global_space ||
+         AddressSpace == access::address_space::generic_space),
+    void>
+joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
+               size_t bytes, Properties properties = {}) {
+  joint_prefetch(g, ptr.get(), bytes, properties);
+}
+
+// Only available if AddressSpace == global_space || AddressSpace ==
+// generic_space
+template <typename Group, typename T, access::address_space AddressSpace,
+          access::decorated IsDecorated,
+          typename Properties = ext::oneapi::experimental::empty_properties_t>
+typename std::enable_if_t<
+    sycl::is_group_v<std::decay_t<Group>> &&
+        (AddressSpace == access::address_space::global_space ||
+         AddressSpace == access::address_space::generic_space),
+    void>
+joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
+               Properties properties = {}) {
+  joint_prefetch(g, ptr.get(), properties);
+}
+
+// Only available if AddressSpace == global_space || AddressSpace ==
+// generic_space
+template <typename Group, typename T, access::address_space AddressSpace,
+          access::decorated IsDecorated,
+          typename Properties = ext::oneapi::experimental::empty_properties_t>
+typename std::enable_if_t<
+    sycl::is_group_v<std::decay_t<Group>> &&
+        (AddressSpace == access::address_space::global_space ||
+         AddressSpace == access::address_space::generic_space),
+    void>
+joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
+               size_t count, Properties properties = {}) {
+  joint_prefetch(g, ptr.get(), count, properties);
+}
+
+// Only available if Dimensions > 0 && (AccessMode == read || AccessMode ==
+// read_write)
+template <typename Group, typename DataT, int Dimensions,
+          access_mode AccessMode, access::placeholder IsPlaceholder,
+          typename Properties = ext::oneapi::experimental::empty_properties_t>
+typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
+                              (Dimensions > 0) &&
+                              (AccessMode == access_mode::read ||
+                               AccessMode == access_mode::read_write),
+                          void>
+joint_prefetch(
+    Group g,
+    accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
+    size_t offset, Properties properties = {}) {
+  joint_prefetch(g, (void *)&acc[offset], sizeof(DataT), properties);
+}
+
+// Only available if Dimensions > 0 && (AccessMode == read || AccessMode ==
+// read_write)
+template <typename Group, typename DataT, int Dimensions,
+          access_mode AccessMode, access::placeholder IsPlaceholder,
+          typename Properties = ext::oneapi::experimental::empty_properties_t>
+typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
+                              (Dimensions > 0) &&
+                              (AccessMode == access_mode::read ||
+                               AccessMode == access_mode::read_write),
+                          void>
+joint_prefetch(
+    Group g,
+    accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
+    size_t offset, size_t count, Properties properties = {}) {
+  joint_prefetch(g, (void *)&acc[offset], count * sizeof(DataT), properties);
+}
+
 } // namespace ext::oneapi::experimental
 } // namespace _V1
 } // namespace sycl

From 19e19a8d79a821da89a90a3a127bfdd1a5a975fd Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Mon, 9 Oct 2023 05:18:43 -0700
Subject: [PATCH 04/25] Fix comment

---
 llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
index b00503879f70..af5578613b19 100644
--- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
+++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
@@ -619,7 +619,7 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
     if (*Property.first == "sycl-alignment")
       continue;
 
-    // leave these annotations as is. They will be processed by SPIRVWriter.
+    // Leave these annotations as is. They will be processed by SPIRVWriter.
     if (*Property.first == "sycl-prefetch-hint" ||
         *Property.first == "sycl-prefetch-hint-nt") {
       return false;

From 30e531a17d6d59e337a85a36c22276be7d81ca10 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Mon, 9 Oct 2023 05:58:11 -0700
Subject: [PATCH 05/25] Use structured bindings

---
 llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp | 13 ++++++-------
 1 file changed, 6 insertions(+), 7 deletions(-)

diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
index af5578613b19..dd018b11e026 100644
--- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
+++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
@@ -613,19 +613,18 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
   // Read the annotation values and create the new annotation string.
   std::string NewAnnotString = "";
   auto Properties = parseSYCLPropertiesString(M, IntrInst);
-  for (auto &Property : Properties) {
+  for (const auto &[first, second] : Properties) {
     // sycl-alignment is converted to align on
     // previous parseAlignmentAndApply(), dropping here
-    if (*Property.first == "sycl-alignment")
+    if (first == "sycl-alignment")
       continue;
 
     // Leave these annotations as is. They will be processed by SPIRVWriter.
-    if (*Property.first == "sycl-prefetch-hint" ||
-        *Property.first == "sycl-prefetch-hint-nt") {
+    if (first == "sycl-prefetch-hint" || first == "sycl-prefetch-hint-nt") {
       return false;
     }
 
-    auto DecorIt = SpirvDecorMap.find(*Property.first);
+    auto DecorIt = SpirvDecorMap.find(*first);
     if (DecorIt == SpirvDecorMap.end())
       continue;
     uint32_t DecorCode = DecorIt->second.Code;
@@ -635,8 +634,8 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
     // string values are handled correctly. Note that " around values are
     // always valid, even if the decoration parameters are not strings.
     NewAnnotString += "{" + std::to_string(DecorCode);
-    if (Property.second)
-      NewAnnotString += ":\"" + Property.second->str() + "\"";
+    if (second)
+      NewAnnotString += ":\"" + second->str() + "\"";
     NewAnnotString += "}";
   }
 

From 42de79f1183b750d905557b77f53011bd7669abe Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Mon, 9 Oct 2023 06:55:58 -0700
Subject: [PATCH 06/25] Remove redundant is_valid_property

---
 sycl/include/sycl/ext/oneapi/prefetch.hpp | 4 ----
 1 file changed, 4 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/prefetch.hpp
index 0dc081c9f711..8394b45c9930 100644
--- a/sycl/include/sycl/ext/oneapi/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/prefetch.hpp
@@ -52,10 +52,6 @@ inline constexpr prefetch_hint_key::value_t<cache_level::L3, nontemporal>
 inline constexpr prefetch_hint_key::value_t<cache_level::L4, nontemporal>
     prefetch_hint_L4_nt;
 
-template <typename T, cache_level Level, typename Hint>
-struct is_valid_property<T, prefetch_hint_key::value_t<Level, Hint>>
-    : std::bool_constant<std::is_pointer<T>::value> {};
-
 namespace detail {
 template <> struct IsCompileTimeProperty<prefetch_hint_key> : std::true_type {};
 

From 2d5a3607fdaab946b69b632dab85e3f085a29477 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Mon, 9 Oct 2023 06:56:22 -0700
Subject: [PATCH 07/25] Add explicit enum values

---
 sycl/include/sycl/ext/oneapi/prefetch.hpp | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/prefetch.hpp
index 8394b45c9930..6497f7057f20 100644
--- a/sycl/include/sycl/ext/oneapi/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/prefetch.hpp
@@ -16,10 +16,10 @@ inline namespace _V1 {
 namespace ext::oneapi::experimental {
 
 enum class cache_level {
-  L1,
-  L2,
-  L3,
-  L4,
+  L1 = 0,
+  L2 = 1,
+  L3 = 2,
+  L4 = 3
 };
 
 struct nontemporal;

From c40e4bd3cba3b809272518fbd0922d475900ad92 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Mon, 9 Oct 2023 07:05:49 -0700
Subject: [PATCH 08/25] Ignore unused vars on host

---
 sycl/include/sycl/ext/oneapi/prefetch.hpp | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/prefetch.hpp
index 6497f7057f20..ddbdc517d011 100644
--- a/sycl/include/sycl/ext/oneapi/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/prefetch.hpp
@@ -79,6 +79,9 @@ void prefetch_impl(void *ptr, size_t bytes, Properties properties) {
   }
   __spirv_ocl_prefetch(ptrAnnotated, bytes);
 #endif
+  std::ignore = ptr;
+  std::ignore = bytes;
+  std::ignore = properties;
 }
 } // namespace detail
 

From 5de267d25f7b3453a670d778906eac8986466ba7 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Mon, 9 Oct 2023 07:14:44 -0700
Subject: [PATCH 09/25] Use enable_if as ret type

---
 sycl/include/sycl/ext/oneapi/prefetch.hpp | 78 ++++++++++++-----------
 1 file changed, 40 insertions(+), 38 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/prefetch.hpp
index ddbdc517d011..2ba0e54b6f00 100644
--- a/sycl/include/sycl/ext/oneapi/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/prefetch.hpp
@@ -107,49 +107,51 @@ void prefetch(T *ptr, size_t count, Properties properties = {}) {
 
 // Only available if AddressSpace == global_space || AddressSpace ==
 // generic_space
-template <
-    access::address_space AddressSpace, access::decorated IsDecorated,
-    typename Properties = empty_properties_t,
-    std::enable_if_t<AddressSpace == access::address_space::global_space ||
-                     AddressSpace == access::address_space::generic_space>>
-void prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr,
-              Properties properties = {}) {
+template <access::address_space AddressSpace, access::decorated IsDecorated,
+          typename Properties = empty_properties_t>
+std::enable_if_t<AddressSpace == access::address_space::global_space ||
+                     AddressSpace == access::address_space::generic_space,
+                 void>
+prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr,
+         Properties properties = {}) {
   prefetch(ptr.get(), properties);
 }
 
 // Only available if AddressSpace == global_space || AddressSpace ==
 // generic_space
-template <
-    access::address_space AddressSpace, access::decorated IsDecorated,
-    typename Properties = empty_properties_t,
-    std::enable_if_t<AddressSpace == access::address_space::global_space ||
-                     AddressSpace == access::address_space::generic_space>>
-void prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, size_t bytes,
-              Properties properties = {}) {
+template <access::address_space AddressSpace, access::decorated IsDecorated,
+          typename Properties = empty_properties_t>
+std::enable_if_t<AddressSpace == access::address_space::global_space ||
+                     AddressSpace == access::address_space::generic_space,
+                 void>
+prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, size_t bytes,
+         Properties properties = {}) {
   prefetch(ptr.get(), bytes, properties);
 }
 
 // Only available if AddressSpace == global_space || AddressSpace ==
 // generic_space
-template <
-    typename T, access::address_space AddressSpace,
-    access::decorated IsDecorated, typename Properties = empty_properties_t,
-    std::enable_if_t<AddressSpace == access::address_space::global_space ||
-                     AddressSpace == access::address_space::generic_space>>
-void prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr,
-              Properties properties = {}) {
+template <typename T, access::address_space AddressSpace,
+          access::decorated IsDecorated,
+          typename Properties = empty_properties_t>
+std::enable_if_t<AddressSpace == access::address_space::global_space ||
+                     AddressSpace == access::address_space::generic_space,
+                 void>
+prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr,
+         Properties properties = {}) {
   prefetch(ptr.get(), properties);
 }
 
 // Only available if AddressSpace == global_space || AddressSpace ==
 // generic_space
-template <
-    typename T, access::address_space AddressSpace,
-    access::decorated IsDecorated, typename Properties = empty_properties_t,
-    std::enable_if_t<AddressSpace == access::address_space::global_space ||
-                     AddressSpace == access::address_space::generic_space>>
-void prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count,
-              Properties properties = {}) {
+template <typename T, access::address_space AddressSpace,
+          access::decorated IsDecorated,
+          typename Properties = empty_properties_t>
+std::enable_if_t<AddressSpace == access::address_space::global_space ||
+                     AddressSpace == access::address_space::generic_space,
+                 void>
+prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count,
+         Properties properties = {}) {
   prefetch(ptr.get(), count, properties);
 }
 
@@ -157,11 +159,11 @@ void prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count,
 // read_write)
 template <typename DataT, int Dimensions, access_mode AccessMode,
           access::placeholder IsPlaceholder,
-          typename Properties = empty_properties_t,
-          std::enable_if_t<(Dimensions > 0) &&
-                           (AccessMode == access_mode::read ||
-                            AccessMode == access_mode::read_write)>>
-void prefetch(
+          typename Properties = empty_properties_t>
+std::enable_if_t<(Dimensions > 0) && (AccessMode == access_mode::read ||
+                                      AccessMode == access_mode::read_write),
+                 void>
+prefetch(
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     id<Dimensions> offset, Properties properties = {}) {
   prefetch((void *)&acc[offset], sizeof(DataT), properties);
@@ -171,11 +173,11 @@ void prefetch(
 // read_write)
 template <typename DataT, int Dimensions, access_mode AccessMode,
           access::placeholder IsPlaceholder,
-          typename Properties = empty_properties_t,
-          std::enable_if_t<(Dimensions > 0) &&
-                           (AccessMode == access_mode::read ||
-                            AccessMode == access_mode::read_write)>>
-void prefetch(
+          typename Properties = empty_properties_t>
+std::enable_if_t<(Dimensions > 0) && (AccessMode == access_mode::read ||
+                                      AccessMode == access_mode::read_write),
+                 void>
+prefetch(
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     size_t offset, size_t count, Properties properties = {}) {
   prefetch((void *)&acc[offset], count * sizeof(DataT), properties);

From 7afe1c961241cce2b09d4fc2ac2054f0accf3ccd Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Mon, 9 Oct 2023 07:31:03 -0700
Subject: [PATCH 10/25] Use single_task in test

---
 sycl/test/extensions/prefetch.cpp | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/sycl/test/extensions/prefetch.cpp b/sycl/test/extensions/prefetch.cpp
index 64969d85c342..9397d3d7eb30 100644
--- a/sycl/test/extensions/prefetch.cpp
+++ b/sycl/test/extensions/prefetch.cpp
@@ -17,18 +17,18 @@ int main() {
   namespace syclex = sycl::ext::oneapi::experimental;
   sycl::queue q;
   void *dataPtr = &data;
-  q.parallel_for(1, [=](sycl::id<1> idx) {
+  q.single_task([=]() {
     // CHECK: [[CASTED:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}
 
-    // CHECK: [[ANNOTATED1:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], ptr addrspace(1) {{.*}}, ptr addrspace(1) {{.*}}, i32 76, ptr addrspace(1) [[ANNOTATION1]])
+    // CHECK: [[ANNOTATED1:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], {{.*}}, ptr addrspace(1) [[ANNOTATION1]])
     // CHECK: tail call spir_func void @_Z20__spirv_ocl_prefetch{{.*}}(ptr addrspace(1) noundef [[ANNOTATED1]], i64 noundef 1)
     syclex::prefetch(dataPtr);
 
-    // CHECK: [[ANNOTATED2:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], ptr addrspace(1) {{.*}}, ptr addrspace(1) {{.*}}, i32 80, ptr addrspace(1) [[ANNOTATION2]])
+    // CHECK: [[ANNOTATED2:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], {{.*}}, ptr addrspace(1) [[ANNOTATION2]])
     // CHECK: tail call spir_func void @_Z20__spirv_ocl_prefetch{{.*}}(ptr addrspace(1) noundef [[ANNOTATED2]], i64 noundef 1)
     syclex::prefetch(dataPtr, syclex::properties{syclex::prefetch_hint_L2});
 
-    // CHECK: [[ANNOTATED3:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], ptr addrspace(1){{.*}}, ptr addrspace(1) {{.*}}, i32 80, ptr addrspace(1) [[ANNOTATION3]])
+    // CHECK: [[ANNOTATED3:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], {{.*}}, ptr addrspace(1) [[ANNOTATION3]])
     // CHECK: tail call spir_func void @_Z20__spirv_ocl_prefetch{{.*}}(ptr addrspace(1) noundef [[ANNOTATED3]], i64 noundef 4)
     syclex::prefetch(dataPtr, 4,
                      syclex::properties{syclex::prefetch_hint_L3_nt});

From 886f2abc816bb92c08cc468b901a0bfd86ec4f7d Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Mon, 9 Oct 2023 07:36:29 -0700
Subject: [PATCH 11/25] clang-format

---
 sycl/include/sycl/ext/oneapi/prefetch.hpp | 7 +------
 1 file changed, 1 insertion(+), 6 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/prefetch.hpp
index 2ba0e54b6f00..3e67defa29fa 100644
--- a/sycl/include/sycl/ext/oneapi/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/prefetch.hpp
@@ -15,12 +15,7 @@ namespace sycl {
 inline namespace _V1 {
 namespace ext::oneapi::experimental {
 
-enum class cache_level {
-  L1 = 0,
-  L2 = 1,
-  L3 = 2,
-  L4 = 3
-};
+enum class cache_level { L1 = 0, L2 = 1, L3 = 2, L4 = 3 };
 
 struct nontemporal;
 

From a8e907bf97d98bbd79501fdc1074540ea832340c Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Tue, 10 Oct 2023 02:01:24 -0700
Subject: [PATCH 12/25] Rename first and second

---
 llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp | 13 +++++++------
 1 file changed, 7 insertions(+), 6 deletions(-)

diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
index dd018b11e026..9273a4c574e1 100644
--- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
+++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
@@ -613,18 +613,19 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
   // Read the annotation values and create the new annotation string.
   std::string NewAnnotString = "";
   auto Properties = parseSYCLPropertiesString(M, IntrInst);
-  for (const auto &[first, second] : Properties) {
+  for (const auto &[propName, propVal] : Properties) {
     // sycl-alignment is converted to align on
     // previous parseAlignmentAndApply(), dropping here
-    if (first == "sycl-alignment")
+    if (propName == "sycl-alignment")
       continue;
 
     // Leave these annotations as is. They will be processed by SPIRVWriter.
-    if (first == "sycl-prefetch-hint" || first == "sycl-prefetch-hint-nt") {
+    if (propName == "sycl-prefetch-hint" ||
+        propName == "sycl-prefetch-hint-nt") {
       return false;
     }
 
-    auto DecorIt = SpirvDecorMap.find(*first);
+    auto DecorIt = SpirvDecorMap.find(*propName);
     if (DecorIt == SpirvDecorMap.end())
       continue;
     uint32_t DecorCode = DecorIt->second.Code;
@@ -634,8 +635,8 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
     // string values are handled correctly. Note that " around values are
     // always valid, even if the decoration parameters are not strings.
     NewAnnotString += "{" + std::to_string(DecorCode);
-    if (second)
-      NewAnnotString += ":\"" + second->str() + "\"";
+    if (propVal)
+      NewAnnotString += ":\"" + propVal->str() + "\"";
     NewAnnotString += "}";
   }
 

From 01a8b5ab0cf65faf86a79a0ec2acfbd07e66a8d2 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Tue, 10 Oct 2023 02:02:14 -0700
Subject: [PATCH 13/25] Fix warning

---
 sycl/include/sycl/ext/oneapi/prefetch.hpp | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/prefetch.hpp
index 3e67defa29fa..41ece6391217 100644
--- a/sycl/include/sycl/ext/oneapi/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/prefetch.hpp
@@ -181,24 +181,28 @@ prefetch(
 template <typename Group, typename Properties = empty_properties_t>
 typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
 joint_prefetch(Group g, void *ptr, Properties properties = {}) {
+  std::ignore = g;
   detail::prefetch_impl(ptr, 1, properties);
 }
 
 template <typename Group, typename Properties = empty_properties_t>
 typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
 joint_prefetch(Group g, void *ptr, size_t bytes, Properties properties = {}) {
+  std::ignore = g;
   detail::prefetch_impl(ptr, bytes, properties);
 }
 
 template <typename Group, typename T, typename Properties = empty_properties_t>
 typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
 joint_prefetch(Group g, T *ptr, Properties properties = {}) {
+  std::ignore = g;
   joint_prefetch((void *)ptr, sizeof(T), properties);
 }
 
 template <typename Group, typename T, typename Properties = empty_properties_t>
 typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
 joint_prefetch(Group g, T *ptr, size_t count, Properties properties = {}) {
+  std::ignore = g;
   joint_prefetch((void *)ptr, count * sizeof(T), properties);
 }
 

From 154660a893d44e868104238157f67e5aa5edd23d Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Tue, 10 Oct 2023 02:06:31 -0700
Subject: [PATCH 14/25] Delete comments

---
 sycl/include/sycl/ext/oneapi/prefetch.hpp | 24 -----------------------
 1 file changed, 24 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/prefetch.hpp
index 41ece6391217..afabe25ab95d 100644
--- a/sycl/include/sycl/ext/oneapi/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/prefetch.hpp
@@ -100,8 +100,6 @@ void prefetch(T *ptr, size_t count, Properties properties = {}) {
   prefetch((void *)ptr, count * sizeof(T), properties);
 }
 
-// Only available if AddressSpace == global_space || AddressSpace ==
-// generic_space
 template <access::address_space AddressSpace, access::decorated IsDecorated,
           typename Properties = empty_properties_t>
 std::enable_if_t<AddressSpace == access::address_space::global_space ||
@@ -112,8 +110,6 @@ prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr,
   prefetch(ptr.get(), properties);
 }
 
-// Only available if AddressSpace == global_space || AddressSpace ==
-// generic_space
 template <access::address_space AddressSpace, access::decorated IsDecorated,
           typename Properties = empty_properties_t>
 std::enable_if_t<AddressSpace == access::address_space::global_space ||
@@ -124,8 +120,6 @@ prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, size_t bytes,
   prefetch(ptr.get(), bytes, properties);
 }
 
-// Only available if AddressSpace == global_space || AddressSpace ==
-// generic_space
 template <typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = empty_properties_t>
@@ -137,8 +131,6 @@ prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr,
   prefetch(ptr.get(), properties);
 }
 
-// Only available if AddressSpace == global_space || AddressSpace ==
-// generic_space
 template <typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = empty_properties_t>
@@ -150,8 +142,6 @@ prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count,
   prefetch(ptr.get(), count, properties);
 }
 
-// Only available if Dimensions > 0 && (AccessMode == read || AccessMode ==
-// read_write)
 template <typename DataT, int Dimensions, access_mode AccessMode,
           access::placeholder IsPlaceholder,
           typename Properties = empty_properties_t>
@@ -164,8 +154,6 @@ prefetch(
   prefetch((void *)&acc[offset], sizeof(DataT), properties);
 }
 
-// Only available if Dimensions > 0 && (AccessMode == read || AccessMode ==
-// read_write)
 template <typename DataT, int Dimensions, access_mode AccessMode,
           access::placeholder IsPlaceholder,
           typename Properties = empty_properties_t>
@@ -206,8 +194,6 @@ joint_prefetch(Group g, T *ptr, size_t count, Properties properties = {}) {
   joint_prefetch((void *)ptr, count * sizeof(T), properties);
 }
 
-// Only available if AddressSpace == global_space || AddressSpace ==
-// generic_space
 template <typename Group, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = ext::oneapi::experimental::empty_properties_t>
@@ -221,8 +207,6 @@ joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
   joint_prefetch(g, ptr.get(), properties);
 }
 
-// Only available if AddressSpace == global_space || AddressSpace ==
-// generic_space
 template <typename Group, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = ext::oneapi::experimental::empty_properties_t>
@@ -236,8 +220,6 @@ joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
   joint_prefetch(g, ptr.get(), bytes, properties);
 }
 
-// Only available if AddressSpace == global_space || AddressSpace ==
-// generic_space
 template <typename Group, typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = ext::oneapi::experimental::empty_properties_t>
@@ -251,8 +233,6 @@ joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
   joint_prefetch(g, ptr.get(), properties);
 }
 
-// Only available if AddressSpace == global_space || AddressSpace ==
-// generic_space
 template <typename Group, typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = ext::oneapi::experimental::empty_properties_t>
@@ -266,8 +246,6 @@ joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
   joint_prefetch(g, ptr.get(), count, properties);
 }
 
-// Only available if Dimensions > 0 && (AccessMode == read || AccessMode ==
-// read_write)
 template <typename Group, typename DataT, int Dimensions,
           access_mode AccessMode, access::placeholder IsPlaceholder,
           typename Properties = ext::oneapi::experimental::empty_properties_t>
@@ -283,8 +261,6 @@ joint_prefetch(
   joint_prefetch(g, (void *)&acc[offset], sizeof(DataT), properties);
 }
 
-// Only available if Dimensions > 0 && (AccessMode == read || AccessMode ==
-// read_write)
 template <typename Group, typename DataT, int Dimensions,
           access_mode AccessMode, access::placeholder IsPlaceholder,
           typename Properties = ext::oneapi::experimental::empty_properties_t>

From 82badcc34a2c0353f6fa95da22c95d7cb5ba44b9 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Tue, 10 Oct 2023 02:08:13 -0700
Subject: [PATCH 15/25] Fix

---
 sycl/include/sycl/ext/oneapi/prefetch.hpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/prefetch.hpp
index afabe25ab95d..9068b100b7c2 100644
--- a/sycl/include/sycl/ext/oneapi/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/prefetch.hpp
@@ -73,10 +73,11 @@ void prefetch_impl(void *ptr, size_t bytes, Properties properties) {
         PropertyMetaInfo<decltype(prop)>::value);
   }
   __spirv_ocl_prefetch(ptrAnnotated, bytes);
-#endif
+#else
   std::ignore = ptr;
   std::ignore = bytes;
   std::ignore = properties;
+#endif
 }
 } // namespace detail
 

From ebc3b23a418d43b68ba2cc7fb9ca4f1590574a8e Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Tue, 10 Oct 2023 02:45:03 -0700
Subject: [PATCH 16/25] Add helper

---
 sycl/include/sycl/ext/oneapi/prefetch.hpp | 93 +++++++++++------------
 1 file changed, 44 insertions(+), 49 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/prefetch.hpp
index 9068b100b7c2..4dd9cf05d199 100644
--- a/sycl/include/sycl/ext/oneapi/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/prefetch.hpp
@@ -58,6 +58,27 @@ struct PropertyMetaInfo<prefetch_hint_key::value_t<Level, Hint>> {
   static constexpr int value = static_cast<int>(Level);
 };
 
+template <typename T, T Parameter, bool AdditionalCondition = true>
+struct prefetch_helper {};
+
+template <access::address_space AddressSpace, bool AdditionalCondition>
+struct prefetch_helper<access::address_space, AddressSpace,
+                       AdditionalCondition> {
+  using type =
+      std::enable_if_t<(AddressSpace == access::address_space::global_space ||
+                        AddressSpace == access::address_space::generic_space) &&
+                           AdditionalCondition,
+                       void>;
+};
+
+template <access_mode AccessMode, bool AdditionalCondition>
+struct prefetch_helper<access_mode, AccessMode, AdditionalCondition> {
+  using type = std::enable_if_t<(AccessMode == access_mode::read ||
+                                 AccessMode == access_mode::write) &&
+                                    AdditionalCondition,
+                                void>;
+};
+
 template <typename Properties>
 void prefetch_impl(void *ptr, size_t bytes, Properties properties) {
 #ifdef __SYCL_DEVICE_ONLY__
@@ -103,9 +124,7 @@ void prefetch(T *ptr, size_t count, Properties properties = {}) {
 
 template <access::address_space AddressSpace, access::decorated IsDecorated,
           typename Properties = empty_properties_t>
-std::enable_if_t<AddressSpace == access::address_space::global_space ||
-                     AddressSpace == access::address_space::generic_space,
-                 void>
+typename detail::prefetch_helper<access::address_space, AddressSpace>::type
 prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr,
          Properties properties = {}) {
   prefetch(ptr.get(), properties);
@@ -113,9 +132,7 @@ prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr,
 
 template <access::address_space AddressSpace, access::decorated IsDecorated,
           typename Properties = empty_properties_t>
-std::enable_if_t<AddressSpace == access::address_space::global_space ||
-                     AddressSpace == access::address_space::generic_space,
-                 void>
+typename detail::prefetch_helper<access::address_space, AddressSpace>::type
 prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, size_t bytes,
          Properties properties = {}) {
   prefetch(ptr.get(), bytes, properties);
@@ -124,9 +141,7 @@ prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, size_t bytes,
 template <typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = empty_properties_t>
-std::enable_if_t<AddressSpace == access::address_space::global_space ||
-                     AddressSpace == access::address_space::generic_space,
-                 void>
+typename detail::prefetch_helper<access::address_space, AddressSpace>::type
 prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr,
          Properties properties = {}) {
   prefetch(ptr.get(), properties);
@@ -135,9 +150,7 @@ prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr,
 template <typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = empty_properties_t>
-std::enable_if_t<AddressSpace == access::address_space::global_space ||
-                     AddressSpace == access::address_space::generic_space,
-                 void>
+typename detail::prefetch_helper<access::address_space, AddressSpace>::type
 prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count,
          Properties properties = {}) {
   prefetch(ptr.get(), count, properties);
@@ -146,9 +159,8 @@ prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count,
 template <typename DataT, int Dimensions, access_mode AccessMode,
           access::placeholder IsPlaceholder,
           typename Properties = empty_properties_t>
-std::enable_if_t<(Dimensions > 0) && (AccessMode == access_mode::read ||
-                                      AccessMode == access_mode::read_write),
-                 void>
+typename detail::prefetch_helper<access_mode, AccessMode,
+                                 (Dimensions > 0)>::type
 prefetch(
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     id<Dimensions> offset, Properties properties = {}) {
@@ -158,9 +170,8 @@ prefetch(
 template <typename DataT, int Dimensions, access_mode AccessMode,
           access::placeholder IsPlaceholder,
           typename Properties = empty_properties_t>
-std::enable_if_t<(Dimensions > 0) && (AccessMode == access_mode::read ||
-                                      AccessMode == access_mode::read_write),
-                 void>
+typename detail::prefetch_helper<access_mode, AccessMode,
+                                 (Dimensions > 0)>::type
 prefetch(
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     size_t offset, size_t count, Properties properties = {}) {
@@ -197,12 +208,9 @@ joint_prefetch(Group g, T *ptr, size_t count, Properties properties = {}) {
 
 template <typename Group, access::address_space AddressSpace,
           access::decorated IsDecorated,
-          typename Properties = ext::oneapi::experimental::empty_properties_t>
-typename std::enable_if_t<
-    sycl::is_group_v<std::decay_t<Group>> &&
-        (AddressSpace == access::address_space::global_space ||
-         AddressSpace == access::address_space::generic_space),
-    void>
+          typename Properties = empty_properties_t>
+typename detail::prefetch_helper<access::address_space, AddressSpace,
+                                 sycl::is_group_v<std::decay_t<Group>>>::type
 joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
                Properties properties = {}) {
   joint_prefetch(g, ptr.get(), properties);
@@ -211,11 +219,8 @@ joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
 template <typename Group, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = ext::oneapi::experimental::empty_properties_t>
-typename std::enable_if_t<
-    sycl::is_group_v<std::decay_t<Group>> &&
-        (AddressSpace == access::address_space::global_space ||
-         AddressSpace == access::address_space::generic_space),
-    void>
+typename detail::prefetch_helper<access::address_space, AddressSpace,
+                                 sycl::is_group_v<std::decay_t<Group>>>::type
 joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
                size_t bytes, Properties properties = {}) {
   joint_prefetch(g, ptr.get(), bytes, properties);
@@ -224,11 +229,8 @@ joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
 template <typename Group, typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = ext::oneapi::experimental::empty_properties_t>
-typename std::enable_if_t<
-    sycl::is_group_v<std::decay_t<Group>> &&
-        (AddressSpace == access::address_space::global_space ||
-         AddressSpace == access::address_space::generic_space),
-    void>
+typename detail::prefetch_helper<access::address_space, AddressSpace,
+                                 sycl::is_group_v<std::decay_t<Group>>>::type
 joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
                Properties properties = {}) {
   joint_prefetch(g, ptr.get(), properties);
@@ -237,11 +239,8 @@ joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
 template <typename Group, typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = ext::oneapi::experimental::empty_properties_t>
-typename std::enable_if_t<
-    sycl::is_group_v<std::decay_t<Group>> &&
-        (AddressSpace == access::address_space::global_space ||
-         AddressSpace == access::address_space::generic_space),
-    void>
+typename detail::prefetch_helper<access::address_space, AddressSpace,
+                                 sycl::is_group_v<std::decay_t<Group>>>::type
 joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
                size_t count, Properties properties = {}) {
   joint_prefetch(g, ptr.get(), count, properties);
@@ -250,11 +249,9 @@ joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
 template <typename Group, typename DataT, int Dimensions,
           access_mode AccessMode, access::placeholder IsPlaceholder,
           typename Properties = ext::oneapi::experimental::empty_properties_t>
-typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
-                              (Dimensions > 0) &&
-                              (AccessMode == access_mode::read ||
-                               AccessMode == access_mode::read_write),
-                          void>
+typename detail::prefetch_helper<
+    access_mode, AccessMode,
+    (Dimensions > 0) && sycl::is_group_v<std::decay_t<Group>>>::type
 joint_prefetch(
     Group g,
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
@@ -265,11 +262,9 @@ joint_prefetch(
 template <typename Group, typename DataT, int Dimensions,
           access_mode AccessMode, access::placeholder IsPlaceholder,
           typename Properties = ext::oneapi::experimental::empty_properties_t>
-typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
-                              (Dimensions > 0) &&
-                              (AccessMode == access_mode::read ||
-                               AccessMode == access_mode::read_write),
-                          void>
+typename detail::prefetch_helper<
+    access_mode, AccessMode,
+    (Dimensions > 0) && sycl::is_group_v<std::decay_t<Group>>>::type
 joint_prefetch(
     Group g,
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,

From a3dbe9608822613874122ca472efdd13884bbe57 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Tue, 10 Oct 2023 09:42:56 -0700
Subject: [PATCH 17/25] Remove redundant C-style casts

---
 sycl/include/sycl/ext/oneapi/prefetch.hpp | 16 ++++++++--------
 1 file changed, 8 insertions(+), 8 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/prefetch.hpp
index 4dd9cf05d199..9868b3b9583b 100644
--- a/sycl/include/sycl/ext/oneapi/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/prefetch.hpp
@@ -114,12 +114,12 @@ void prefetch(void *ptr, size_t bytes, Properties properties = {}) {
 
 template <typename T, typename Properties = empty_properties_t>
 void prefetch(T *ptr, Properties properties = {}) {
-  prefetch((void *)ptr, sizeof(T), properties);
+  prefetch(ptr, sizeof(T), properties);
 }
 
 template <typename T, typename Properties = empty_properties_t>
 void prefetch(T *ptr, size_t count, Properties properties = {}) {
-  prefetch((void *)ptr, count * sizeof(T), properties);
+  prefetch(ptr, count * sizeof(T), properties);
 }
 
 template <access::address_space AddressSpace, access::decorated IsDecorated,
@@ -164,7 +164,7 @@ typename detail::prefetch_helper<access_mode, AccessMode,
 prefetch(
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     id<Dimensions> offset, Properties properties = {}) {
-  prefetch((void *)&acc[offset], sizeof(DataT), properties);
+  prefetch(&acc[offset], sizeof(DataT), properties);
 }
 
 template <typename DataT, int Dimensions, access_mode AccessMode,
@@ -175,7 +175,7 @@ typename detail::prefetch_helper<access_mode, AccessMode,
 prefetch(
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     size_t offset, size_t count, Properties properties = {}) {
-  prefetch((void *)&acc[offset], count * sizeof(DataT), properties);
+  prefetch(&acc[offset], count * sizeof(DataT), properties);
 }
 
 template <typename Group, typename Properties = empty_properties_t>
@@ -196,14 +196,14 @@ template <typename Group, typename T, typename Properties = empty_properties_t>
 typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
 joint_prefetch(Group g, T *ptr, Properties properties = {}) {
   std::ignore = g;
-  joint_prefetch((void *)ptr, sizeof(T), properties);
+  joint_prefetch(ptr, sizeof(T), properties);
 }
 
 template <typename Group, typename T, typename Properties = empty_properties_t>
 typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
 joint_prefetch(Group g, T *ptr, size_t count, Properties properties = {}) {
   std::ignore = g;
-  joint_prefetch((void *)ptr, count * sizeof(T), properties);
+  joint_prefetch(ptr, count * sizeof(T), properties);
 }
 
 template <typename Group, access::address_space AddressSpace,
@@ -256,7 +256,7 @@ joint_prefetch(
     Group g,
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     size_t offset, Properties properties = {}) {
-  joint_prefetch(g, (void *)&acc[offset], sizeof(DataT), properties);
+  joint_prefetch(g, &acc[offset], sizeof(DataT), properties);
 }
 
 template <typename Group, typename DataT, int Dimensions,
@@ -269,7 +269,7 @@ joint_prefetch(
     Group g,
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     size_t offset, size_t count, Properties properties = {}) {
-  joint_prefetch(g, (void *)&acc[offset], count * sizeof(DataT), properties);
+  joint_prefetch(g, &acc[offset], count * sizeof(DataT), properties);
 }
 
 } // namespace ext::oneapi::experimental

From 9e81d89e11539826a51efffa0bce91811cda2e47 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Tue, 10 Oct 2023 10:26:49 -0700
Subject: [PATCH 18/25] Change helper

---
 sycl/include/sycl/ext/oneapi/prefetch.hpp | 92 +++++++++++------------
 1 file changed, 42 insertions(+), 50 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/prefetch.hpp
index 9868b3b9583b..9d47b0044456 100644
--- a/sycl/include/sycl/ext/oneapi/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/prefetch.hpp
@@ -58,26 +58,14 @@ struct PropertyMetaInfo<prefetch_hint_key::value_t<Level, Hint>> {
   static constexpr int value = static_cast<int>(Level);
 };
 
-template <typename T, T Parameter, bool AdditionalCondition = true>
-struct prefetch_helper {};
-
-template <access::address_space AddressSpace, bool AdditionalCondition>
-struct prefetch_helper<access::address_space, AddressSpace,
-                       AdditionalCondition> {
-  using type =
-      std::enable_if_t<(AddressSpace == access::address_space::global_space ||
-                        AddressSpace == access::address_space::generic_space) &&
-                           AdditionalCondition,
-                       void>;
-};
+template <access::address_space AS>
+inline constexpr bool check_prefetch_AS =
+    AS == access::address_space::global_space ||
+    AS == access::address_space::generic_space;
 
-template <access_mode AccessMode, bool AdditionalCondition>
-struct prefetch_helper<access_mode, AccessMode, AdditionalCondition> {
-  using type = std::enable_if_t<(AccessMode == access_mode::read ||
-                                 AccessMode == access_mode::write) &&
-                                    AdditionalCondition,
-                                void>;
-};
+template <access_mode mode>
+inline constexpr bool check_prefetch_acc_mode =
+    mode == access_mode::read || mode == access_mode::read_write;
 
 template <typename Properties>
 void prefetch_impl(void *ptr, size_t bytes, Properties properties) {
@@ -124,7 +112,7 @@ void prefetch(T *ptr, size_t count, Properties properties = {}) {
 
 template <access::address_space AddressSpace, access::decorated IsDecorated,
           typename Properties = empty_properties_t>
-typename detail::prefetch_helper<access::address_space, AddressSpace>::type
+std::enable_if_t<detail::check_prefetch_AS<AddressSpace>, void>
 prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr,
          Properties properties = {}) {
   prefetch(ptr.get(), properties);
@@ -132,7 +120,7 @@ prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr,
 
 template <access::address_space AddressSpace, access::decorated IsDecorated,
           typename Properties = empty_properties_t>
-typename detail::prefetch_helper<access::address_space, AddressSpace>::type
+std::enable_if_t<detail::check_prefetch_AS<AddressSpace>, void>
 prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, size_t bytes,
          Properties properties = {}) {
   prefetch(ptr.get(), bytes, properties);
@@ -141,7 +129,7 @@ prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, size_t bytes,
 template <typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = empty_properties_t>
-typename detail::prefetch_helper<access::address_space, AddressSpace>::type
+std::enable_if_t<detail::check_prefetch_AS<AddressSpace>, void>
 prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr,
          Properties properties = {}) {
   prefetch(ptr.get(), properties);
@@ -150,7 +138,7 @@ prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr,
 template <typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = empty_properties_t>
-typename detail::prefetch_helper<access::address_space, AddressSpace>::type
+std::enable_if_t<detail::check_prefetch_AS<AddressSpace>, void>
 prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count,
          Properties properties = {}) {
   prefetch(ptr.get(), count, properties);
@@ -159,8 +147,8 @@ prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count,
 template <typename DataT, int Dimensions, access_mode AccessMode,
           access::placeholder IsPlaceholder,
           typename Properties = empty_properties_t>
-typename detail::prefetch_helper<access_mode, AccessMode,
-                                 (Dimensions > 0)>::type
+std::enable_if_t<
+    detail::check_prefetch_acc_mode<AccessMode> && (Dimensions > 0), void>
 prefetch(
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     id<Dimensions> offset, Properties properties = {}) {
@@ -170,8 +158,8 @@ prefetch(
 template <typename DataT, int Dimensions, access_mode AccessMode,
           access::placeholder IsPlaceholder,
           typename Properties = empty_properties_t>
-typename detail::prefetch_helper<access_mode, AccessMode,
-                                 (Dimensions > 0)>::type
+std::enable_if_t<
+    detail::check_prefetch_acc_mode<AccessMode> && (Dimensions > 0), void>
 prefetch(
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     size_t offset, size_t count, Properties properties = {}) {
@@ -179,28 +167,28 @@ prefetch(
 }
 
 template <typename Group, typename Properties = empty_properties_t>
-typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
+std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
 joint_prefetch(Group g, void *ptr, Properties properties = {}) {
   std::ignore = g;
   detail::prefetch_impl(ptr, 1, properties);
 }
 
 template <typename Group, typename Properties = empty_properties_t>
-typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
+std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
 joint_prefetch(Group g, void *ptr, size_t bytes, Properties properties = {}) {
   std::ignore = g;
   detail::prefetch_impl(ptr, bytes, properties);
 }
 
 template <typename Group, typename T, typename Properties = empty_properties_t>
-typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
+std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
 joint_prefetch(Group g, T *ptr, Properties properties = {}) {
   std::ignore = g;
   joint_prefetch(ptr, sizeof(T), properties);
 }
 
 template <typename Group, typename T, typename Properties = empty_properties_t>
-typename std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
+std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
 joint_prefetch(Group g, T *ptr, size_t count, Properties properties = {}) {
   std::ignore = g;
   joint_prefetch(ptr, count * sizeof(T), properties);
@@ -209,8 +197,9 @@ joint_prefetch(Group g, T *ptr, size_t count, Properties properties = {}) {
 template <typename Group, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = empty_properties_t>
-typename detail::prefetch_helper<access::address_space, AddressSpace,
-                                 sycl::is_group_v<std::decay_t<Group>>>::type
+std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
+                     sycl::is_group_v<std::decay_t<Group>>,
+                 void>
 joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
                Properties properties = {}) {
   joint_prefetch(g, ptr.get(), properties);
@@ -218,9 +207,10 @@ joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
 
 template <typename Group, access::address_space AddressSpace,
           access::decorated IsDecorated,
-          typename Properties = ext::oneapi::experimental::empty_properties_t>
-typename detail::prefetch_helper<access::address_space, AddressSpace,
-                                 sycl::is_group_v<std::decay_t<Group>>>::type
+          typename Properties = empty_properties_t>
+std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
+                     sycl::is_group_v<std::decay_t<Group>>,
+                 void>
 joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
                size_t bytes, Properties properties = {}) {
   joint_prefetch(g, ptr.get(), bytes, properties);
@@ -228,9 +218,10 @@ joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
 
 template <typename Group, typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
-          typename Properties = ext::oneapi::experimental::empty_properties_t>
-typename detail::prefetch_helper<access::address_space, AddressSpace,
-                                 sycl::is_group_v<std::decay_t<Group>>>::type
+          typename Properties = empty_properties_t>
+std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
+                     sycl::is_group_v<std::decay_t<Group>>,
+                 void>
 joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
                Properties properties = {}) {
   joint_prefetch(g, ptr.get(), properties);
@@ -238,9 +229,10 @@ joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
 
 template <typename Group, typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
-          typename Properties = ext::oneapi::experimental::empty_properties_t>
-typename detail::prefetch_helper<access::address_space, AddressSpace,
-                                 sycl::is_group_v<std::decay_t<Group>>>::type
+          typename Properties = empty_properties_t>
+std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
+                     sycl::is_group_v<std::decay_t<Group>>,
+                 void>
 joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
                size_t count, Properties properties = {}) {
   joint_prefetch(g, ptr.get(), count, properties);
@@ -248,10 +240,10 @@ joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
 
 template <typename Group, typename DataT, int Dimensions,
           access_mode AccessMode, access::placeholder IsPlaceholder,
-          typename Properties = ext::oneapi::experimental::empty_properties_t>
-typename detail::prefetch_helper<
-    access_mode, AccessMode,
-    (Dimensions > 0) && sycl::is_group_v<std::decay_t<Group>>>::type
+          typename Properties = empty_properties_t>
+std::enable_if_t<detail::check_prefetch_acc_mode<AccessMode> &&
+                     (Dimensions > 0) && sycl::is_group_v<std::decay_t<Group>>,
+                 void>
 joint_prefetch(
     Group g,
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
@@ -261,10 +253,10 @@ joint_prefetch(
 
 template <typename Group, typename DataT, int Dimensions,
           access_mode AccessMode, access::placeholder IsPlaceholder,
-          typename Properties = ext::oneapi::experimental::empty_properties_t>
-typename detail::prefetch_helper<
-    access_mode, AccessMode,
-    (Dimensions > 0) && sycl::is_group_v<std::decay_t<Group>>>::type
+          typename Properties = empty_properties_t>
+std::enable_if_t<detail::check_prefetch_acc_mode<AccessMode> &&
+                     (Dimensions > 0) && sycl::is_group_v<std::decay_t<Group>>,
+                 void>
 joint_prefetch(
     Group g,
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,

From 3e554acf96d408cd59aded226d46d096e1e7ea0a Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Wed, 11 Oct 2023 01:40:27 -0700
Subject: [PATCH 19/25] Drop void in enable_if

---
 sycl/include/sycl/ext/oneapi/prefetch.hpp | 42 ++++++++++-------------
 1 file changed, 18 insertions(+), 24 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/prefetch.hpp
index 9d47b0044456..92814a2d9f32 100644
--- a/sycl/include/sycl/ext/oneapi/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/prefetch.hpp
@@ -112,7 +112,7 @@ void prefetch(T *ptr, size_t count, Properties properties = {}) {
 
 template <access::address_space AddressSpace, access::decorated IsDecorated,
           typename Properties = empty_properties_t>
-std::enable_if_t<detail::check_prefetch_AS<AddressSpace>, void>
+std::enable_if_t<detail::check_prefetch_AS<AddressSpace>>
 prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr,
          Properties properties = {}) {
   prefetch(ptr.get(), properties);
@@ -120,7 +120,7 @@ prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr,
 
 template <access::address_space AddressSpace, access::decorated IsDecorated,
           typename Properties = empty_properties_t>
-std::enable_if_t<detail::check_prefetch_AS<AddressSpace>, void>
+std::enable_if_t<detail::check_prefetch_AS<AddressSpace>>
 prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, size_t bytes,
          Properties properties = {}) {
   prefetch(ptr.get(), bytes, properties);
@@ -129,7 +129,7 @@ prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, size_t bytes,
 template <typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = empty_properties_t>
-std::enable_if_t<detail::check_prefetch_AS<AddressSpace>, void>
+std::enable_if_t<detail::check_prefetch_AS<AddressSpace>>
 prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr,
          Properties properties = {}) {
   prefetch(ptr.get(), properties);
@@ -138,7 +138,7 @@ prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr,
 template <typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = empty_properties_t>
-std::enable_if_t<detail::check_prefetch_AS<AddressSpace>, void>
+std::enable_if_t<detail::check_prefetch_AS<AddressSpace>>
 prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count,
          Properties properties = {}) {
   prefetch(ptr.get(), count, properties);
@@ -147,8 +147,8 @@ prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count,
 template <typename DataT, int Dimensions, access_mode AccessMode,
           access::placeholder IsPlaceholder,
           typename Properties = empty_properties_t>
-std::enable_if_t<
-    detail::check_prefetch_acc_mode<AccessMode> && (Dimensions > 0), void>
+std::enable_if_t<detail::check_prefetch_acc_mode<AccessMode> &&
+                 (Dimensions > 0)>
 prefetch(
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     id<Dimensions> offset, Properties properties = {}) {
@@ -158,8 +158,8 @@ prefetch(
 template <typename DataT, int Dimensions, access_mode AccessMode,
           access::placeholder IsPlaceholder,
           typename Properties = empty_properties_t>
-std::enable_if_t<
-    detail::check_prefetch_acc_mode<AccessMode> && (Dimensions > 0), void>
+std::enable_if_t<detail::check_prefetch_acc_mode<AccessMode> &&
+                 (Dimensions > 0)>
 prefetch(
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     size_t offset, size_t count, Properties properties = {}) {
@@ -167,28 +167,28 @@ prefetch(
 }
 
 template <typename Group, typename Properties = empty_properties_t>
-std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
+std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, void *ptr, Properties properties = {}) {
   std::ignore = g;
   detail::prefetch_impl(ptr, 1, properties);
 }
 
 template <typename Group, typename Properties = empty_properties_t>
-std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
+std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, void *ptr, size_t bytes, Properties properties = {}) {
   std::ignore = g;
   detail::prefetch_impl(ptr, bytes, properties);
 }
 
 template <typename Group, typename T, typename Properties = empty_properties_t>
-std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
+std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, T *ptr, Properties properties = {}) {
   std::ignore = g;
   joint_prefetch(ptr, sizeof(T), properties);
 }
 
 template <typename Group, typename T, typename Properties = empty_properties_t>
-std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>, void>
+std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, T *ptr, size_t count, Properties properties = {}) {
   std::ignore = g;
   joint_prefetch(ptr, count * sizeof(T), properties);
@@ -198,8 +198,7 @@ template <typename Group, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = empty_properties_t>
 std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
-                     sycl::is_group_v<std::decay_t<Group>>,
-                 void>
+                 sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
                Properties properties = {}) {
   joint_prefetch(g, ptr.get(), properties);
@@ -209,8 +208,7 @@ template <typename Group, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = empty_properties_t>
 std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
-                     sycl::is_group_v<std::decay_t<Group>>,
-                 void>
+                 sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
                size_t bytes, Properties properties = {}) {
   joint_prefetch(g, ptr.get(), bytes, properties);
@@ -220,8 +218,7 @@ template <typename Group, typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = empty_properties_t>
 std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
-                     sycl::is_group_v<std::decay_t<Group>>,
-                 void>
+                 sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
                Properties properties = {}) {
   joint_prefetch(g, ptr.get(), properties);
@@ -231,8 +228,7 @@ template <typename Group, typename T, access::address_space AddressSpace,
           access::decorated IsDecorated,
           typename Properties = empty_properties_t>
 std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
-                     sycl::is_group_v<std::decay_t<Group>>,
-                 void>
+                 sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
                size_t count, Properties properties = {}) {
   joint_prefetch(g, ptr.get(), count, properties);
@@ -242,8 +238,7 @@ template <typename Group, typename DataT, int Dimensions,
           access_mode AccessMode, access::placeholder IsPlaceholder,
           typename Properties = empty_properties_t>
 std::enable_if_t<detail::check_prefetch_acc_mode<AccessMode> &&
-                     (Dimensions > 0) && sycl::is_group_v<std::decay_t<Group>>,
-                 void>
+                 (Dimensions > 0) && sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(
     Group g,
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
@@ -255,8 +250,7 @@ template <typename Group, typename DataT, int Dimensions,
           access_mode AccessMode, access::placeholder IsPlaceholder,
           typename Properties = empty_properties_t>
 std::enable_if_t<detail::check_prefetch_acc_mode<AccessMode> &&
-                     (Dimensions > 0) && sycl::is_group_v<std::decay_t<Group>>,
-                 void>
+                 (Dimensions > 0) && sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(
     Group g,
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,

From 024c901e19660c5695d89e44dd63d600264f06e8 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Thu, 12 Oct 2023 05:16:21 -0700
Subject: [PATCH 20/25] Remove sycl-post-link & test in favor on a follow-up
 patch

---
 .../SYCLLowerIR/CompileTimePropertiesPass.cpp |  6 ---
 sycl/test/extensions/prefetch.cpp             | 39 -------------------
 2 files changed, 45 deletions(-)
 delete mode 100644 sycl/test/extensions/prefetch.cpp

diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
index fd42d847ff95..8e6276455607 100644
--- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
+++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
@@ -634,12 +634,6 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
     if (propName == "sycl-alignment")
       continue;
 
-    // Leave these annotations as is. They will be processed by SPIRVWriter.
-    if (propName == "sycl-prefetch-hint" ||
-        propName == "sycl-prefetch-hint-nt") {
-      return false;
-    }
-
     auto DecorIt = SpirvDecorMap.find(*propName);
     if (DecorIt == SpirvDecorMap.end())
       continue;
diff --git a/sycl/test/extensions/prefetch.cpp b/sycl/test/extensions/prefetch.cpp
deleted file mode 100644
index 9397d3d7eb30..000000000000
--- a/sycl/test/extensions/prefetch.cpp
+++ /dev/null
@@ -1,39 +0,0 @@
-// RUN: %clangxx -fsycl-device-only -S %s -o - | FileCheck %s
-
-#include <sycl/sycl.hpp>
-
-char data[] = {0, 1, 2, 3};
-
-// CHECK: [[PREFETCH_STR:@.*]] = private unnamed_addr addrspace(1) constant [19 x i8] c"sycl-prefetch-hint\00", section "llvm.metadata"
-// CHECK: [[PREFETCH_LVL0:@.*]] = private unnamed_addr addrspace(1) constant [2 x i8] c"0\00", section "llvm.metadata"
-// CHECK: [[ANNOTATION1:@.*]] = private unnamed_addr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) [[PREFETCH_STR]], ptr addrspace(1) [[PREFETCH_LVL0]] }, section "llvm.metadata"
-// CHECK: [[PREFETCH_LVL1:@.*]] = private unnamed_addr addrspace(1) constant [2 x i8] c"1\00", section "llvm.metadata"
-// CHECK: [[ANNOTATION2:@.*]] = private unnamed_addr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) [[PREFETCH_STR]], ptr addrspace(1) [[PREFETCH_LVL1]] }, section "llvm.metadata"
-// CHECK: [[PREFETCH_STR_NT:@.*]] = private unnamed_addr addrspace(1) constant [22 x i8] c"sycl-prefetch-hint-nt\00", section "llvm.metadata"
-// CHECK: [[PREFETCH_LVL2:@.*]] = private unnamed_addr addrspace(1) constant [2 x i8] c"2\00", section "llvm.metadata"
-// CHECK: [[ANNOTATION3:@.*]] = private unnamed_addr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) [[PREFETCH_STR_NT]], ptr addrspace(1) [[PREFETCH_LVL2]] }, section "llvm.metadata"
-
-int main() {
-  namespace syclex = sycl::ext::oneapi::experimental;
-  sycl::queue q;
-  void *dataPtr = &data;
-  q.single_task([=]() {
-    // CHECK: [[CASTED:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobal{{.*}}
-
-    // CHECK: [[ANNOTATED1:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], {{.*}}, ptr addrspace(1) [[ANNOTATION1]])
-    // CHECK: tail call spir_func void @_Z20__spirv_ocl_prefetch{{.*}}(ptr addrspace(1) noundef [[ANNOTATED1]], i64 noundef 1)
-    syclex::prefetch(dataPtr);
-
-    // CHECK: [[ANNOTATED2:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], {{.*}}, ptr addrspace(1) [[ANNOTATION2]])
-    // CHECK: tail call spir_func void @_Z20__spirv_ocl_prefetch{{.*}}(ptr addrspace(1) noundef [[ANNOTATED2]], i64 noundef 1)
-    syclex::prefetch(dataPtr, syclex::properties{syclex::prefetch_hint_L2});
-
-    // CHECK: [[ANNOTATED3:%.*]] = tail call ptr addrspace(1) @llvm.ptr.annotation.p1.p1(ptr addrspace(1) [[CASTED]], {{.*}}, ptr addrspace(1) [[ANNOTATION3]])
-    // CHECK: tail call spir_func void @_Z20__spirv_ocl_prefetch{{.*}}(ptr addrspace(1) noundef [[ANNOTATED3]], i64 noundef 4)
-    syclex::prefetch(dataPtr, 4,
-                     syclex::properties{syclex::prefetch_hint_L3_nt});
-  });
-  q.wait();
-
-  return 0;
-}

From c796b9e063313c04dcda28b0feeb46f08b0fbcff Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Fri, 13 Oct 2023 05:56:23 -0700
Subject: [PATCH 21/25] use capital initial letters

---
 llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp | 10 +++++-----
 1 file changed, 5 insertions(+), 5 deletions(-)

diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
index 8e6276455607..b16b6b515d04 100644
--- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
+++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
@@ -628,13 +628,13 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
   // Read the annotation values and create the new annotation string.
   std::string NewAnnotString = "";
   auto Properties = parseSYCLPropertiesString(M, IntrInst);
-  for (const auto &[propName, propVal] : Properties) {
+  for (const auto &[PropName, PropVal] : Properties) {
     // sycl-alignment is converted to align on
     // previous parseAlignmentAndApply(), dropping here
-    if (propName == "sycl-alignment")
+    if (PropName == "sycl-alignment")
       continue;
 
-    auto DecorIt = SpirvDecorMap.find(*propName);
+    auto DecorIt = SpirvDecorMap.find(*PropName);
     if (DecorIt == SpirvDecorMap.end())
       continue;
     uint32_t DecorCode = DecorIt->second.Code;
@@ -644,8 +644,8 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
     // string values are handled correctly. Note that " around values are
     // always valid, even if the decoration parameters are not strings.
     NewAnnotString += "{" + std::to_string(DecorCode);
-    if (propVal)
-      NewAnnotString += ":\"" + propVal->str() + "\"";
+    if (PropVal)
+      NewAnnotString += ":\"" + PropVal->str() + "\"";
     NewAnnotString += "}";
   }
 

From bf3d6c937f643cd7bb59f3e7b4d4045b70b7c8ed Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Fri, 13 Oct 2023 06:01:40 -0700
Subject: [PATCH 22/25] Move prefetch.hpp to experimental/ dir

---
 sycl/include/sycl/ext/oneapi/{ => experimental}/prefetch.hpp | 0
 sycl/include/sycl/sycl.hpp                                   | 2 +-
 2 files changed, 1 insertion(+), 1 deletion(-)
 rename sycl/include/sycl/ext/oneapi/{ => experimental}/prefetch.hpp (100%)

diff --git a/sycl/include/sycl/ext/oneapi/prefetch.hpp b/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp
similarity index 100%
rename from sycl/include/sycl/ext/oneapi/prefetch.hpp
rename to sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp
diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp
index 67f24f11c740..ca663c981c5e 100644
--- a/sycl/include/sycl/sycl.hpp
+++ b/sycl/include/sycl/sycl.hpp
@@ -82,6 +82,7 @@
 #include <sycl/ext/oneapi/experimental/cuda/barrier.hpp>
 #include <sycl/ext/oneapi/experimental/fixed_size_group.hpp>
 #include <sycl/ext/oneapi/experimental/opportunistic_group.hpp>
+#include <sycl/ext/oneapi/experimental/prefetch.hpp>
 #include <sycl/ext/oneapi/experimental/tangle_group.hpp>
 #include <sycl/ext/oneapi/filter_selector.hpp>
 #include <sycl/ext/oneapi/functional.hpp>
@@ -89,7 +90,6 @@
 #include <sycl/ext/oneapi/kernel_properties/properties.hpp>
 #include <sycl/ext/oneapi/matrix/matrix.hpp>
 #include <sycl/ext/oneapi/owner_less.hpp>
-#include <sycl/ext/oneapi/prefetch.hpp>
 #include <sycl/ext/oneapi/properties/properties.hpp>
 #include <sycl/ext/oneapi/properties/property_value.hpp>
 #include <sycl/ext/oneapi/sub_group.hpp>

From 6d9de5aaa5b4f2816476a3d229f8e1f92642bc3d Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Fri, 13 Oct 2023 08:02:57 -0700
Subject: [PATCH 23/25] Improve impl

---
 .../sycl/ext/oneapi/experimental/prefetch.hpp | 58 ++++++++++---------
 1 file changed, 32 insertions(+), 26 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp b/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp
index 92814a2d9f32..1a5b6eae4c44 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp
@@ -67,11 +67,11 @@ template <access_mode mode>
 inline constexpr bool check_prefetch_acc_mode =
     mode == access_mode::read || mode == access_mode::read_write;
 
-template <typename Properties>
-void prefetch_impl(void *ptr, size_t bytes, Properties properties) {
+template <typename T, typename Properties>
+void prefetch_impl(T *ptr, size_t bytes, Properties properties) {
 #ifdef __SYCL_DEVICE_ONLY__
-  auto *ptrGlobalAS = __SYCL_GenericCastToPtrExplicit_ToGlobal<char>(ptr);
-  __attribute__((opencl_global)) char *ptrAnnotated = nullptr;
+  auto *ptrGlobalAS = __SYCL_GenericCastToPtrExplicit_ToGlobal<const char>(ptr);
+  const __attribute__((opencl_global)) char *ptrAnnotated = nullptr;
   if constexpr (!properties.template has_property<prefetch_hint_key>()) {
     ptrAnnotated = __builtin_intel_sycl_ptr_annotation(
         ptrGlobalAS, "sycl-prefetch-hint", static_cast<int>(cache_level::L1));
@@ -88,6 +88,15 @@ void prefetch_impl(void *ptr, size_t bytes, Properties properties) {
   std::ignore = properties;
 #endif
 }
+
+template <typename Group, typename T, typename Properties>
+void joint_prefetch_impl(Group g, T *ptr, size_t bytes, Properties properties) {
+  // Although calling joint_prefetch is functionally equivalent to calling
+  // prefetch from every work-item in a group, native suppurt may be added to to
+  // issue cooperative prefetches more efficiently on some hardware.
+  std::ignore = g;
+  prefetch_impl(ptr, bytes, properties);
+}
 } // namespace detail
 
 template <typename Properties = empty_properties_t>
@@ -102,12 +111,12 @@ void prefetch(void *ptr, size_t bytes, Properties properties = {}) {
 
 template <typename T, typename Properties = empty_properties_t>
 void prefetch(T *ptr, Properties properties = {}) {
-  prefetch(ptr, sizeof(T), properties);
+  detail::prefetch_impl(ptr, sizeof(T), properties);
 }
 
 template <typename T, typename Properties = empty_properties_t>
 void prefetch(T *ptr, size_t count, Properties properties = {}) {
-  prefetch(ptr, count * sizeof(T), properties);
+  detail::prefetch_impl(ptr, count * sizeof(T), properties);
 }
 
 template <access::address_space AddressSpace, access::decorated IsDecorated,
@@ -115,7 +124,7 @@ template <access::address_space AddressSpace, access::decorated IsDecorated,
 std::enable_if_t<detail::check_prefetch_AS<AddressSpace>>
 prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr,
          Properties properties = {}) {
-  prefetch(ptr.get(), properties);
+  detail::prefetch_impl(ptr.get(), 1, properties);
 }
 
 template <access::address_space AddressSpace, access::decorated IsDecorated,
@@ -123,7 +132,7 @@ template <access::address_space AddressSpace, access::decorated IsDecorated,
 std::enable_if_t<detail::check_prefetch_AS<AddressSpace>>
 prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr, size_t bytes,
          Properties properties = {}) {
-  prefetch(ptr.get(), bytes, properties);
+  detail::prefetch_impl(ptr.get(), bytes, properties);
 }
 
 template <typename T, access::address_space AddressSpace,
@@ -132,7 +141,7 @@ template <typename T, access::address_space AddressSpace,
 std::enable_if_t<detail::check_prefetch_AS<AddressSpace>>
 prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr,
          Properties properties = {}) {
-  prefetch(ptr.get(), properties);
+  detail::prefetch_impl(ptr.get(), sizeof(T), properties);
 }
 
 template <typename T, access::address_space AddressSpace,
@@ -141,7 +150,7 @@ template <typename T, access::address_space AddressSpace,
 std::enable_if_t<detail::check_prefetch_AS<AddressSpace>>
 prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr, size_t count,
          Properties properties = {}) {
-  prefetch(ptr.get(), count, properties);
+  detail::prefetch_impl(ptr.get(), count * sizeof(T), properties);
 }
 
 template <typename DataT, int Dimensions, access_mode AccessMode,
@@ -152,7 +161,7 @@ std::enable_if_t<detail::check_prefetch_acc_mode<AccessMode> &&
 prefetch(
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     id<Dimensions> offset, Properties properties = {}) {
-  prefetch(&acc[offset], sizeof(DataT), properties);
+  detail::prefetch_impl(&acc[offset], sizeof(DataT), properties);
 }
 
 template <typename DataT, int Dimensions, access_mode AccessMode,
@@ -163,35 +172,31 @@ std::enable_if_t<detail::check_prefetch_acc_mode<AccessMode> &&
 prefetch(
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     size_t offset, size_t count, Properties properties = {}) {
-  prefetch(&acc[offset], count * sizeof(DataT), properties);
+  detail::prefetch_impl(&acc[offset], count * sizeof(DataT), properties);
 }
 
 template <typename Group, typename Properties = empty_properties_t>
 std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, void *ptr, Properties properties = {}) {
-  std::ignore = g;
-  detail::prefetch_impl(ptr, 1, properties);
+  detail::joint_prefetch_impl(g, ptr, 1, properties);
 }
 
 template <typename Group, typename Properties = empty_properties_t>
 std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, void *ptr, size_t bytes, Properties properties = {}) {
-  std::ignore = g;
-  detail::prefetch_impl(ptr, bytes, properties);
+  detail::joint_prefetch_impl(g, ptr, bytes, properties);
 }
 
 template <typename Group, typename T, typename Properties = empty_properties_t>
 std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, T *ptr, Properties properties = {}) {
-  std::ignore = g;
-  joint_prefetch(ptr, sizeof(T), properties);
+  detail::joint_prefetch_impl(g, ptr, sizeof(T), properties);
 }
 
 template <typename Group, typename T, typename Properties = empty_properties_t>
 std::enable_if_t<sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, T *ptr, size_t count, Properties properties = {}) {
-  std::ignore = g;
-  joint_prefetch(ptr, count * sizeof(T), properties);
+  detail::joint_prefetch_impl(g, ptr, count * sizeof(T), properties);
 }
 
 template <typename Group, access::address_space AddressSpace,
@@ -201,7 +206,7 @@ std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
                  sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
                Properties properties = {}) {
-  joint_prefetch(g, ptr.get(), properties);
+  detail::joint_prefetch_impl(g, ptr.get(), 1, properties);
 }
 
 template <typename Group, access::address_space AddressSpace,
@@ -211,7 +216,7 @@ std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
                  sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
                size_t bytes, Properties properties = {}) {
-  joint_prefetch(g, ptr.get(), bytes, properties);
+  detail::joint_prefetch_impl(g, ptr.get(), bytes, properties);
 }
 
 template <typename Group, typename T, access::address_space AddressSpace,
@@ -221,7 +226,7 @@ std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
                  sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
                Properties properties = {}) {
-  joint_prefetch(g, ptr.get(), properties);
+  detail::joint_prefetch_impl(g, ptr.get(), properties);
 }
 
 template <typename Group, typename T, access::address_space AddressSpace,
@@ -231,7 +236,7 @@ std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
                  sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
                size_t count, Properties properties = {}) {
-  joint_prefetch(g, ptr.get(), count, properties);
+  detail::joint_prefetch_impl(g, ptr.get(), count, properties);
 }
 
 template <typename Group, typename DataT, int Dimensions,
@@ -243,7 +248,7 @@ joint_prefetch(
     Group g,
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     size_t offset, Properties properties = {}) {
-  joint_prefetch(g, &acc[offset], sizeof(DataT), properties);
+  detail::joint_prefetch_impl(g, &acc[offset], sizeof(DataT), properties);
 }
 
 template <typename Group, typename DataT, int Dimensions,
@@ -255,7 +260,8 @@ joint_prefetch(
     Group g,
     accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
     size_t offset, size_t count, Properties properties = {}) {
-  joint_prefetch(g, &acc[offset], count * sizeof(DataT), properties);
+  detail::joint_prefetch_impl(g, &acc[offset], count * sizeof(DataT),
+                              properties);
 }
 
 } // namespace ext::oneapi::experimental

From 9ff283ee17d6d5c64a2bef8c2c5dd8b3d466fc51 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Fri, 13 Oct 2023 08:05:53 -0700
Subject: [PATCH 24/25] Add new test

---
 sycl/test/extensions/prefetch.cpp | 62 +++++++++++++++++++++++++++++++
 1 file changed, 62 insertions(+)
 create mode 100644 sycl/test/extensions/prefetch.cpp

diff --git a/sycl/test/extensions/prefetch.cpp b/sycl/test/extensions/prefetch.cpp
new file mode 100644
index 000000000000..56fd678f8333
--- /dev/null
+++ b/sycl/test/extensions/prefetch.cpp
@@ -0,0 +1,62 @@
+// RUN: %clangxx -fsycl -fsyntax-only %s
+
+#include <sycl/sycl.hpp>
+
+int data[] = {0, 1, 2, 3};
+
+int main() {
+  namespace syclex = sycl::ext::oneapi::experimental;
+  void *dataPtrVoid = data;
+  int *dataPtrInt = data;
+  auto prop = syclex::properties{syclex::prefetch_hint_L1};
+
+  {
+    sycl::buffer<int, 1> buf(data, 4);
+    sycl::queue q;
+    q.submit([&](sycl::handler &h) {
+      auto acc = buf.get_access<sycl::access_mode::read>(h);
+      h.parallel_for<class Kernel>(
+          sycl::nd_range<1>(1, 1), ([=](sycl::nd_item<1> index) {
+            syclex::prefetch(dataPtrVoid, prop);
+            syclex::prefetch(dataPtrVoid, 16, prop);
+
+            syclex::prefetch(dataPtrInt, prop);
+            syclex::prefetch(dataPtrInt, 4, prop);
+
+            auto mPtrVoid = sycl::address_space_cast<
+                sycl::access::address_space::global_space,
+                sycl::access::decorated::yes>(dataPtrVoid);
+            syclex::prefetch(mPtrVoid, prop);
+            syclex::prefetch(mPtrVoid, 16, prop);
+
+            auto mPtrInt = sycl::address_space_cast<
+                sycl::access::address_space::global_space,
+                sycl::access::decorated::yes>(dataPtrInt);
+            syclex::prefetch(mPtrInt, prop);
+            syclex::prefetch(mPtrInt, 8, prop);
+
+            syclex::prefetch(acc, sycl::id(0), prop);
+            syclex::prefetch(acc, sycl::id(0), 4, prop);
+
+            auto g = index.get_group();
+            syclex::joint_prefetch(g, dataPtrVoid, prop);
+            syclex::joint_prefetch(g, dataPtrVoid, 16, prop);
+
+            syclex::joint_prefetch(g, dataPtrInt, prop);
+            syclex::joint_prefetch(g, dataPtrInt, 4, prop);
+
+            syclex::joint_prefetch(g, mPtrVoid, prop);
+            syclex::joint_prefetch(g, mPtrVoid, 16, prop);
+
+            syclex::joint_prefetch(g, mPtrInt, prop);
+            syclex::joint_prefetch(g, mPtrInt, 8, prop);
+
+            syclex::joint_prefetch(g, acc, sycl::id(0), prop);
+            syclex::joint_prefetch(g, acc, sycl::id(0), 4, prop);
+          }));
+    });
+    q.wait();
+  }
+
+  return 0;
+}

From 9b2ecc83bffee57ed6bcd4cf29f517f95c6622c5 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Fri, 13 Oct 2023 09:22:29 -0700
Subject: [PATCH 25/25] small fix

---
 sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp b/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp
index 1a5b6eae4c44..9271af505940 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/prefetch.hpp
@@ -226,7 +226,7 @@ std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
                  sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
                Properties properties = {}) {
-  detail::joint_prefetch_impl(g, ptr.get(), properties);
+  detail::joint_prefetch_impl(g, ptr.get(), sizeof(T), properties);
 }
 
 template <typename Group, typename T, access::address_space AddressSpace,
@@ -236,7 +236,7 @@ std::enable_if_t<detail::check_prefetch_AS<AddressSpace> &&
                  sycl::is_group_v<std::decay_t<Group>>>
 joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
                size_t count, Properties properties = {}) {
-  detail::joint_prefetch_impl(g, ptr.get(), count, properties);
+  detail::joint_prefetch_impl(g, ptr.get(), count * sizeof(T), properties);
 }
 
 template <typename Group, typename DataT, int Dimensions,