From dfbaad274de63b8759d0b4fc81dec04991ac5798 Mon Sep 17 00:00:00 2001 From: "Wang, Di5" Date: Wed, 31 Jan 2024 05:42:32 -0800 Subject: [PATCH] add lit & update spec --- .../SYCLLowerIR/CompileTimePropertiesPass.cpp | 5 -- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 4 +- .../fpga_annotated_properties.hpp | 3 ++ .../annotated_ptr/annotated_ptr.hpp | 18 ------- .../properties.hpp | 2 + .../annotated_ptr/annotation_insertion.cpp | 48 +++++++++++++++++++ 6 files changed, 55 insertions(+), 25 deletions(-) create mode 100644 sycl/test/extensions/annotated_ptr/annotation_insertion.cpp diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index 18b3ca253c0f2..5135846ea0f25 100644 --- 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 index b85904e975587..162cf9a8bbe08 100644 --- 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 index 7bc04324d6b0e..fcf4419e35052 100644 --- 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 : std::true_type {}; // buffer_location is applied on PtrAnnotation template <> struct propagateToPtrAnnotation : std::true_type {}; +template +struct propagateToPtrAnnotation> + : 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 index d40611c54e136..4f32802ae548b 100644 --- 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> { detail::PropertyMetaInfo

::value...); } - // static I load(I *ptr) { - // return *annotate(ptr); - // } - - // template static I store(I *ptr, O &&Obj) { - // return *annotate(ptr) = std::forward(Obj); - // } - static I load(I *ptr) { return *__builtin_intel_sycl_ptr_annotation( ptr, detail::PropertyMetaInfo

::name..., @@ -138,16 +130,6 @@ class annotated_ref> { return *this = t2; } - // address-of operator - T *operator&() const { -#ifdef __SYCL_DEVICE_ONLY__ - return annotationHelper>::annotate( - m_Ptr); -#else - return *m_Ptr; -#endif - } - // propagate compound operators #define PROPAGATE_OP(op) \ template >> \ 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 index 9864e916bc475..fa1b65d43d3c0 100644 --- 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> : std::true_type {}; template <> struct propagateToPtrAnnotation : std::true_type {}; +template +struct propagateToPtrAnnotation> : 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 new file mode 100644 index 0000000000000..78d20a0f7f5d7 --- /dev/null +++ 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 + +#include + +// 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, 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(5, q); + q.submit([&](handler &h) { h.single_task(MyIP{raw}); }).wait(); + free(raw, q); +} + +int main() { + TestVectorAddWithAnnotatedMMHosts(); + return 0; +}