[SYCL] Implement SYCL part of sycl_ext_oneapi_prefetch

llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

@@ -628,13 +628,19 @@ 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);
+    // 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;
     uint32_t DecorCode = DecorIt->second.Code;
@@ -644,8 +650,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,277 @@
+//==--------------- 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 <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__
+  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((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);
+}
+
+template <access::address_space AddressSpace, access::decorated IsDecorated,
+          typename Properties = empty_properties_t>
+typename detail::prefetch_helper<access::address_space, AddressSpace>::type
+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>
+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);
+}
+
+template <typename T, access::address_space AddressSpace,
+          access::decorated IsDecorated,
+          typename Properties = empty_properties_t>
+typename detail::prefetch_helper<access::address_space, AddressSpace>::type
+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>
+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);
+}
+
+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
+prefetch(
+    accessor<DataT, Dimensions, AccessMode, target::device, IsPlaceholder> acc,
+    id<Dimensions> offset, Properties properties = {}) {
+  prefetch((void *)&acc[offset], sizeof(DataT), properties);
+}
+
+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
+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);
+}
+
+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);
+}
+
+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
+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 = ext::oneapi::experimental::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,
+               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 = ext::oneapi::experimental::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<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 = ext::oneapi::experimental::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<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 = ext::oneapi::experimental::empty_properties_t>
+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,
+    size_t offset, Properties properties = {}) {
+  joint_prefetch(g, (void *)&acc[offset], sizeof(DataT), properties);
+}
+
+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
+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

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

sycl/test/extensions/prefetch.cpp

@@ -0,0 +1,39 @@
+// 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;
+}