add lit & update spec

intel · Jan 31, 2024 · dfbaad2 · dfbaad2
1 parent 3ebd695
commit dfbaad2
Show file tree

Hide file tree

Showing 6 changed files with 55 additions and 25 deletions.
diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
@@ -780,11 +780,6 @@ 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")
-      continue;
-
     auto DecorIt = SpirvDecorMap.find(*PropName);
     if (DecorIt == SpirvDecorMap.end())
       continue;

diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_annotated_ptr.asciidoc b/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|

diff --git a/sycl/include/sycl/ext/intel/experimental/fpga_annotated_properties.hpp b/sycl/include/sycl/ext/intel/experimental/fpga_annotated_properties.hpp
@@ -356,6 +356,9 @@ struct is_valid_property<T, stable_key::value_t> : std::true_type {};
 // buffer_location is applied on PtrAnnotation
 template <>
 struct propagateToPtrAnnotation<buffer_location_key> : std::true_type {};
+template <int K>
+struct propagateToPtrAnnotation<buffer_location_key::value_t<K>>
+    : std::true_type {};
 
 //===----------------------------------------------------------------------===//
 //   Utility for FPGA properties

diff --git a/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp b/sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp
@@ -75,14 +75,6 @@ struct annotationHelper<I, detail::properties_t<P...>> {
         detail::PropertyMetaInfo<P>::value...);
   }
 
-  // static I load(I *ptr) {
-  //   return *annotate(ptr);
-  // }
-
-  // template <class O> static I store(I *ptr, O &&Obj) {
-  //   return *annotate(ptr) = std::forward<O>(Obj);
-  // }
-
   static I load(I *ptr) {
     return *__builtin_intel_sycl_ptr_annotation(
         ptr, detail::PropertyMetaInfo<P>::name...,
@@ -138,16 +130,6 @@ class annotated_ref<T, detail::properties_t<Props...>> {
     return *this = t2;
   }
 
-  // address-of operator
-  T *operator&() const {
-#ifdef __SYCL_DEVICE_ONLY__
-    return annotationHelper<T, detail::annotation_filter<Props...>>::annotate(
-        m_Ptr);
-#else
-    return *m_Ptr;
-#endif
-  }
-
   // propagate compound operators
 #define PROPAGATE_OP(op)                                                       \
   template <class O, typename = std::enable_if_t<!detail::is_ann_ref_v<O>>>    \

diff --git a/sycl/include/sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp b/sycl/include/sycl/ext/oneapi/experimental/common_annotated_properties/properties.hpp
@@ -83,6 +83,8 @@ struct is_property_key_of<alignment_key, annotated_arg<T, PropertyListT>>
     : std::true_type {};
 
 template <> struct propagateToPtrAnnotation<alignment_key> : std::true_type {};
+template <int K>
+struct propagateToPtrAnnotation<alignment_key::value_t<K>> : std::true_type {};
 
 namespace detail {
 

diff --git a/sycl/test/extensions/annotated_ptr/annotation_insertion.cpp b/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;
+}