[SYCL] Implement SYCL part of sycl_ext_oneapi_prefetch

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 (auto &Property : Properties) {
+  for (const auto &[propName, propVal] : Properties) {
     // sycl-alignment is converted to align on
     // previous parseAlignmentAndApply(), dropping here
-    if (*Property.first == "sycl-alignment")
+    if (propName == "sycl-alignment")
       continue;
 
-    auto DecorIt = SpirvDecorMap.find(*Property.first);
+    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 (Property.second)
-      NewAnnotString += ":\"" + Property.second->str() + "\"";
+    if (propVal)
+      NewAnnotString += ":\"" + propVal->str() + "\"";
     NewAnnotString += "}";
   }
 

sycl/include/sycl/ext/oneapi/prefetch.hpp

@@ -0,0 +1,263 @@
+//==--------------- 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 = 0, L2 = 1, L3 = 2, L4 = 3 };
+
+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;
+
+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 <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 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) {
+#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);
+#else
+  std::ignore = ptr;
+  std::ignore = bytes;
+  std::ignore = properties;
+#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(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);
+}
+
+template <access::address_space AddressSpace, access::decorated IsDecorated,
+          typename Properties = empty_properties_t>
+std::enable_if_t<detail::check_prefetch_AS<AddressSpace>>
+prefetch(multi_ptr<void, AddressSpace, IsDecorated> ptr,
+         Properties properties = {}) {
+  prefetch(ptr.get(), properties);
+}
+
+template <access::address_space AddressSpace, access::decorated IsDecorated,
+          typename Properties = empty_properties_t>
+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);
+}
+
+template <typename T, access::address_space AddressSpace,
+          access::decorated IsDecorated,
+          typename Properties = empty_properties_t>
+std::enable_if_t<detail::check_prefetch_AS<AddressSpace>>
+prefetch(multi_ptr<T, AddressSpace, IsDecorated> ptr,
+         Properties properties = {}) {
+  prefetch(ptr.get(), properties);
+}
+
+template <typename T, access::address_space AddressSpace,
+          access::decorated IsDecorated,
+          typename Properties = empty_properties_t>
+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);
+}
+
+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)>
+prefetch(
+    accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
+    id<Dimensions> offset, Properties properties = {}) {
+  prefetch(&acc[offset], sizeof(DataT), properties);
+}
+
+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)>
+prefetch(
+    accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
+    size_t offset, size_t count, Properties properties = {}) {
+  prefetch(&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);
+}
+
+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);
+}
+
+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);
+}
+
+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);
+}
+
+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>>>
+joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
+               Properties properties = {}) {
+  joint_prefetch(g, ptr.get(), properties);
+}
+
+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>>>
+joint_prefetch(Group g, multi_ptr<void, AddressSpace, IsDecorated> ptr,
+               size_t bytes, Properties properties = {}) {
+  joint_prefetch(g, ptr.get(), bytes, properties);
+}
+
+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>>>
+joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
+               Properties properties = {}) {
+  joint_prefetch(g, ptr.get(), properties);
+}
+
+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>>>
+joint_prefetch(Group g, multi_ptr<T, AddressSpace, IsDecorated> ptr,
+               size_t count, Properties properties = {}) {
+  joint_prefetch(g, ptr.get(), count, properties);
+}
+
+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>>>
+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);
+}
+
+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>>>
+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);
+}
+
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/sycl.hpp

@@ -89,6 +89,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>

sycl/source/feature_test.hpp.in

@@ -90,6 +90,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