[SYCL] Insert annotation in annotated_ptr::get() (#12343)

When properties like alignment is specified in a `annotated_ptr` type,
certain operators (like `[]`, `+=`, `++`) are disabled. This results in
loop code to be written as follows:
```
annotated_ptr<int, decltype(properties{...alignment<8>...})> ann_ptr;
...
int *p = ann_ptr.get();     // ann_ptr cannot be used in the for loop directly
for (int i = 0; i < n; i++) {
    p[i] = i;
}
```
When getting the underlying pointer, the annotation gets lost, so does
the the possible optimization on the for-loop brought by the
annotated_ptr properties.

This PR includes changes on spec, header and clang compiler:
1. In `annotated_ptr` spec, update the spec for the `get()` function
2. In the `annotated_ptr` header, update the `get()` function by
inserting `llvm.ptr.annotation`, so that on the target machines like
FPGA for which clang FE only performs O0 optimization, the annotation
inserted can be preserved for the corresponding backends to perform
platform-specific optimizations. For the example above, the `alignment`
information can help the FPGA compiler to build aligned loads/stores.
3. In the clang compiler, the pass `CompileTimePropertiesPass` used to
always drop `alignment` from the annotation string. This PR changes this
behavior to dropping `alignment` only when the compiler finds
load/store/MemIntrinsics in the users of `llvm.ptr.annotation` and
applies the alignment to these instructions.

Author: wangdi4

llvm/include/llvm/SYCLLowerIR/CompileTimePropertiesPass.h

@@ -40,7 +40,7 @@ class CompileTimePropertiesPass
       Module &M, IntrinsicInst *IntrInst,
       SmallVectorImpl<IntrinsicInst *> &RemovableAnnotations);
 
-  void parseAlignmentAndApply(Module &M, IntrinsicInst *IntrInst);
+  bool parseAlignmentAndApply(Module &M, IntrinsicInst *IntrInst);
 
   // Map for keeping track of global variables generated for annotation strings.
   // This allows reuse for annotations with the same generated annotation

llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

@@ -685,7 +685,7 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M,
                                   : PreservedAnalyses::all();
 }
 
-void CompileTimePropertiesPass::parseAlignmentAndApply(
+bool CompileTimePropertiesPass::parseAlignmentAndApply(
     Module &M, IntrinsicInst *IntrInst) {
   // Get the global variable with the annotation string.
   const GlobalVariable *AnnotStrArgGV = nullptr;
@@ -695,11 +695,11 @@ void CompileTimePropertiesPass::parseAlignmentAndApply(
   else if (auto *GEP = dyn_cast<GEPOperator>(IntrAnnotStringArg))
     AnnotStrArgGV = dyn_cast<GlobalVariable>(GEP->getOperand(0));
   if (!AnnotStrArgGV)
-    return;
+    return false;
 
   std::optional<StringRef> AnnotStr = getGlobalVariableString(AnnotStrArgGV);
   if (!AnnotStr)
-    return;
+    return false;
 
   // parse properties string to decoration-value pairs
   auto Properties = parseSYCLPropertiesString(M, IntrInst);
@@ -710,6 +710,7 @@ void CompileTimePropertiesPass::parseAlignmentAndApply(
   getUserListIgnoringCast<StoreInst>(IntrInst, TargetedInstList);
   getUserListIgnoringCast<MemTransferInst>(IntrInst, TargetedInstList);
 
+  bool AlignApplied = false;
   for (auto &Property : Properties) {
     auto DecorStr = Property.first->str();
     auto DecorValue = Property.second;
@@ -733,18 +734,26 @@ void CompileTimePropertiesPass::parseAlignmentAndApply(
         auto Op_num = Pair.second;
         if (auto *LInst = dyn_cast<LoadInst>(Inst)) {
           LInst->setAlignment(Align_val);
+          AlignApplied = true;
         } else if (auto *SInst = dyn_cast<StoreInst>(Inst)) {
-          if (Op_num == 1)
+          if (Op_num == 1) {
             SInst->setAlignment(Align_val);
+            AlignApplied = true;
+          }
         } else if (auto *MI = dyn_cast<MemTransferInst>(Inst)) {
-          if (Op_num == 0)
+          if (Op_num == 0) {
             MI->setDestAlignment(Align_val);
-          else if (Op_num == 1)
+            AlignApplied = true;
+          } else if (Op_num == 1) {
             MI->setSourceAlignment(Align_val);
+            AlignApplied = true;
+          }
         }
       }
     }
   }
+
+  return AlignApplied;
 }
 
 // Returns true if the transformation changed IntrInst.
@@ -773,7 +782,7 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
     return false;
 
   // check alignment annotation and apply it to load/store
-  parseAlignmentAndApply(M, IntrInst);
+  bool AlignApplied = parseAlignmentAndApply(M, IntrInst);
 
   // Read the annotation values and create new annotation strings.
   std::string NewAnnotString = "";
@@ -782,9 +791,9 @@ bool CompileTimePropertiesPass::transformSYCLPropertiesAnnotation(
   bool CacheProp = false;
   bool FPGAProp = false;
   for (const auto &[PropName, PropVal] : Properties) {
-    // sycl-alignment is converted to align on
-    // previous parseAlignmentAndApply(), dropping here
-    if (PropName == "sycl-alignment")
+    // if sycl-alignment is converted to align on IR constructs
+    // during parseAlignmentAndApply(), dropping here
+    if (PropName == "sycl-alignment" && AlignApplied)
       continue;
 
     auto DecorIt = SpirvDecorMap.find(*PropName);

llvm/test/SYCLLowerIR/CompileTimePropertiesPass/sycl-properties-alignment-loadstore.ll renamed to llvm/test/SYCLLowerIR/CompileTimePropertiesPass/sycl-properties-alignment.ll

@@ -1,7 +1,7 @@
 ; RUN: opt -passes=compile-time-properties -S %s -o %t.ll
 ; RUN: FileCheck %s -input-file=%t.ll
 ;
-; Tests the translation of "sycl-alignment" to alignment attributes on load/store
+; Tests the translation of "sycl-alignment" to alignment attributes on load/store/non-memory instructions
 
 target triple = "spir64_fpga-unknown-unknown"
 
@@ -11,13 +11,14 @@ target triple = "spir64_fpga-unknown-unknown"
 $_ZN7ann_refIiEC2EPi = comdat any
 $_ZN7ann_refIiEcvRiEv = comdat any
 $_ZN7ann_refIiEC2EPi1= comdat any
+$no_load_store = comdat any
 
 @.str = private unnamed_addr addrspace(1) constant [16 x i8] c"sycl-properties\00", section "llvm.metadata"
 @.str.1 = private unnamed_addr addrspace(1) constant [9 x i8] c"main.cpp\00", section "llvm.metadata"
 @.str.2 = private unnamed_addr addrspace(1) constant [15 x i8] c"sycl-alignment\00", section "llvm.metadata"
 @.str.3 = private unnamed_addr addrspace(1) constant [3 x i8] c"64\00", section "llvm.metadata"
-@.args = private unnamed_addr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) @.str.2, ptr addrspace(1) @.str.3 }, section "llvm.met
-adata"
+@.args = private unnamed_addr addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) @.str.2, ptr addrspace(1) @.str.3 }, section "llvm.metadata"
+; CHECK: @[[AnnoStr:.*]] = private unnamed_addr addrspace(1) constant [10 x i8] c"{44:\2264\22}\00"
 
 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite)
 declare ptr addrspace(4) @llvm.ptr.annotation.p4.p1(ptr addrspace(4), ptr addrspace(1), ptr addrspace(1), i32, ptr addrspace(1)) #5
@@ -77,4 +78,19 @@ entry:
   ret void
 }
 
+; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
+define linkonce_odr dso_local spir_func noundef ptr addrspace(4) @no_load_store(ptr addrspace(4) noundef %ptr) comdat align 2 {
+entry:
+  %retval = alloca ptr addrspace(4), align 8
+  %ptr.addr = alloca ptr addrspace(4), align 8
+  %retval.ascast = addrspacecast ptr %retval to ptr addrspace(4)
+  %ptr.addr.ascast = addrspacecast ptr %ptr.addr to ptr addrspace(4)
+  store ptr addrspace(4) %ptr, ptr addrspace(4) %ptr.addr.ascast, align 8
+  %0 = load ptr addrspace(4), ptr addrspace(4) %ptr.addr.ascast, align 8
+  ; CHECK: %[[AnnoPtr:.*]] = call ptr addrspace(4) @llvm.ptr.annotation.p4.p1(ptr addrspace(4) %0, ptr addrspace(1) @[[AnnoStr]]
+  ; CHECK: ret ptr addrspace(4) %[[AnnoPtr]]
+  %1 = call ptr addrspace(4) @llvm.ptr.annotation.p4.p1(ptr addrspace(4) %0, ptr addrspace(1) @.str, ptr addrspace(1) @.str.1, i32 73, ptr addrspace(1) @.args)
+  ret ptr addrspace(4) %1
+}
+
 declare void @llvm.memcpy.p4.p4.i32(ptr addrspace(4), ptr addrspace(4), i32, i1)

sycl/doc/extensions/experimental/sycl_ext_oneapi_annotated_ptr.asciidoc

@@ -476,8 +476,8 @@ a|
 T* get() const noexcept;
 ----
 |
-Returns the underlying raw pointer. The raw pointer will not retain the
-annotations.
+Returns the underlying raw pointer. Implementations are free to propagate information from properties of
+an annotated_ptr to the raw pointer.
 
 // --- ROW BREAK ---
 a|

sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp

@@ -53,8 +53,41 @@ template <typename... Ts>
 using contains_alignment =
     detail::ContainsProperty<alignment_key, std::tuple<Ts...>>;
 
+// properties filter
+template <typename property_list, template <class...> typename filter>
+using PropertiesFilter =
+    sycl::detail::boost::mp11::mp_copy_if<property_list, filter>;
+
+// filter properties that are applied on annotations
+template <typename... Props>
+using annotation_filter = properties<
+    PropertiesFilter<std::tuple<Props...>, propagateToPtrAnnotation>>;
 } // namespace detail
 
+template <typename I, typename P> struct annotationHelper {};
+
+// unpack properties to varadic template
+template <typename I, typename... P>
+struct annotationHelper<I, detail::properties_t<P...>> {
+  static I *annotate(I *ptr) {
+    return __builtin_intel_sycl_ptr_annotation(
+        ptr, detail::PropertyMetaInfo<P>::name...,
+        detail::PropertyMetaInfo<P>::value...);
+  }
+
+  static I load(I *ptr) {
+    return *__builtin_intel_sycl_ptr_annotation(
+        ptr, detail::PropertyMetaInfo<P>::name...,
+        detail::PropertyMetaInfo<P>::value...);
+  }
+
+  template <class O> static I store(I *ptr, O &&Obj) {
+    return *__builtin_intel_sycl_ptr_annotation(
+               ptr, detail::PropertyMetaInfo<P>::name...,
+               detail::PropertyMetaInfo<P>::value...) = std::forward<O>(Obj);
+  }
+};
+
 template <typename T, typename... Props>
 class annotated_ref<T, detail::properties_t<Props...>> {
   using property_list_t = detail::properties_t<Props...>;
@@ -67,44 +100,14 @@ class annotated_ref<T, detail::properties_t<Props...>> {
   T *m_Ptr;
   explicit annotated_ref(T *Ptr) : m_Ptr(Ptr) {}
 
-  // properties filter
-  template <typename property_list, template <class...> typename filter>
-  using PropertiesFilter =
-      sycl::detail::boost::mp11::mp_copy_if<property_list, filter>;
-
-  template <typename p>
-  using annotation_filter = propagateToPtrAnnotation<typename p::key_t>;
-
-  // filter properties that are applied on annotations
-  using property_tuple_t = std::tuple<Props...>;
-  using annotation_props =
-      properties<PropertiesFilter<property_tuple_t, annotation_filter>>;
-
-  template <typename I, typename P> struct annotationHelper {};
-
-  // unpack properties to varadic template
-  template <typename I, typename... P>
-  struct annotationHelper<I, detail::properties_t<P...>> {
-    static I load(I *ptr) {
-      return *__builtin_intel_sycl_ptr_annotation(
-          ptr, detail::PropertyMetaInfo<P>::name...,
-          detail::PropertyMetaInfo<P>::value...);
-    }
-
-    template <class O> static I store(I *ptr, O &&Obj) {
-      return *__builtin_intel_sycl_ptr_annotation(
-                 ptr, detail::PropertyMetaInfo<P>::name...,
-                 detail::PropertyMetaInfo<P>::value...) = std::forward<O>(Obj);
-    }
-  };
-
 public:
   annotated_ref(const annotated_ref &) = delete;
 
   // implicit conversion with annotaion
   operator T() const {
 #ifdef __SYCL_DEVICE_ONLY__
-    return annotationHelper<T, annotation_props>::load(m_Ptr);
+    return annotationHelper<T, detail::annotation_filter<Props...>>::load(
+        m_Ptr);
 #else
     return *m_Ptr;
 #endif
@@ -114,7 +117,8 @@ class annotated_ref<T, detail::properties_t<Props...>> {
   template <class O, typename = std::enable_if_t<!detail::is_ann_ref_v<O>>>
   T operator=(O &&Obj) const {
 #ifdef __SYCL_DEVICE_ONLY__
-    return annotationHelper<T, annotation_props>::store(m_Ptr, Obj);
+    return annotationHelper<T, detail::annotation_filter<Props...>>::store(
+        m_Ptr, Obj);
 #else
     return *m_Ptr = std::forward<O>(Obj);
 #endif
@@ -376,7 +380,14 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr<T, detail::properties_t<Props...>> {
 
   operator T *() const noexcept = delete;
 
-  T *get() const noexcept { return m_Ptr; }
+  T *get() const noexcept {
+#ifdef __SYCL_DEVICE_ONLY__
+    return annotationHelper<T, detail::annotation_filter<Props...>>::annotate(
+        m_Ptr);
+#else
+    return m_Ptr;
+#endif
+  }
 
   // When the properties contain alignment, operator '[]', '+', '++' and '--'
   // (both post- and prefix) are disabled. Calling these operators when

sycl/include/sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp

@@ -58,6 +58,11 @@ struct check_property_list<T, Prop, Props...>
 
 template <typename PropTy> struct propagateToPtrAnnotation : std::false_type {};
 
+// Partial specilization for property_value
+template <typename PropKeyT, typename... PropValuesTs>
+struct propagateToPtrAnnotation<property_value<PropKeyT, PropValuesTs...>>
+    : propagateToPtrAnnotation<PropKeyT> {};
+
 //===----------------------------------------------------------------------===//
 //        Common properties of annotated_arg/annotated_ptr
 //===----------------------------------------------------------------------===//

sycl/test/extensions/annotated_ptr/annotation_insertion.cpp

@@ -0,0 +1,48 @@
+// RUN: %clangxx -fsycl-device-only -fsycl-targets=spir64_fpga -S -emit-llvm %s -o - | FileCheck %s
+
+// Tests that `@llvm.ptr.annotation` is inserted when calling
+// `annotated_ptr::get()`
+
+#include "sycl/sycl.hpp"
+#include <sycl/ext/intel/fpga_extensions.hpp>
+
+#include <iostream>
+
+// clang-format on
+
+using namespace sycl;
+using namespace ext::oneapi::experimental;
+using namespace ext::intel::experimental;
+
+// CHECK: @[[AnnStr:.*]] = private unnamed_addr addrspace(1) constant [19 x i8] c"{5921:\220\22}{44:\228\22}\00"
+
+using ann_ptr_t1 =
+    annotated_ptr<int, decltype(properties(buffer_location<0>, alignment<8>))>;
+
+struct MyIP {
+  ann_ptr_t1 a;
+
+  MyIP(int *a_) : a(a_) {}
+
+  void operator()() const {
+    // CHECK: %ptr.addr = alloca ptr addrspace(4), align 8
+    // CHECK: store ptr addrspace(4) %ptr, ptr %ptr.addr, align 8
+    // CHECK: %[[LoadPtr:.*]] = load ptr addrspace(4), ptr %ptr.addr, align 8
+    // CHECK: %[[AnnPtr:.*]] = call ptr addrspace(4) @llvm.ptr.annotation.p4.p1(ptr addrspace(4) %[[LoadPtr]], ptr addrspace(1) @[[AnnStr]]
+    // CHECK: ret ptr addrspace(4) %[[AnnPtr]]
+    int *ptr = a.get(); // llvm.ptr.annotation is inserted
+    *ptr = 15;
+  }
+};
+
+void TestVectorAddWithAnnotatedMMHosts() {
+  sycl::queue q;
+  auto raw = malloc_shared<int>(5, q);
+  q.submit([&](handler &h) { h.single_task(MyIP{raw}); }).wait();
+  free(raw, q);
+}
+
+int main() {
+  TestVectorAddWithAnnotatedMMHosts();
+  return 0;
+}

sycl/test/extensions/properties/properties_cache_control.cpp

@@ -9,14 +9,12 @@ using namespace ext::intel::experimental;
 
 using load_hint = annotated_ptr<
     float, decltype(properties(
-               alignment<8>,
                read_hint<cache_control<cache_mode::cached, cache_level::L1>,
                          cache_control<cache_mode::uncached, cache_level::L2,
                                        cache_level::L3>>))>;
 using load_assertion = annotated_ptr<
     int,
     decltype(properties(
-        alignment<8>,
         read_assertion<cache_control<cache_mode::constant, cache_level::L1>,
                        cache_control<cache_mode::invalidate, cache_level::L2,
                                      cache_level::L3>>))>;