[SYCL] Refine annotated_ptr address space based on compile-time USM kind (#11623)

Right now the annotated_ptr extension applies compile-time properties on `decorated_global_ptr` for SYCL device code. As the SYCL extension [sycl_ext_oneapi_usm_malloc_properties](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_usm_malloc_properties.asciidoc#sycl_ext_oneapi_usm_malloc_properties) introduces compile-time property "usm_kind"(implemented in #10235 ), this PR further refines the address space of this pointer (using `decorated_device_ptr` or `decorated_host_ptr`) when 1. `usm_kind` property is specified in annotated_ptr type, and 2. both `__SYCL_DEVICE_ONLY__` and `__ENABLE_USM_ADDR_SPACE__` are turned on (i.e. requiring compiler flags "-fsycl-device-only -fsycl-targets=spir64_fpga") Other changes: - move `annotated_ptr.hpp` out from `annotated_arg` file into a separate folder - add new header file `annotated_ptr_properties.hpp` for annotated_ptr specific properties

Author: wangdi4

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

@@ -177,10 +177,20 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr<T, detail::properties_t<Props...>> {
  T, typename unpack<filtered_properties>::type>;
 
 #ifdef __SYCL_DEVICE_ONLY__
+#ifdef __ENABLE_USM_ADDR_SPACE__
+ using global_pointer_t = std::conditional_t<
+ detail::IsUsmKindDevice<property_list_t>::value,
+ typename sycl::ext::intel::decorated_device_ptr<T>::pointer,
+ std::conditional_t<
+ detail::IsUsmKindHost<property_list_t>::value,
+ typename sycl::ext::intel::decorated_host_ptr<T>::pointer,
+ typename decorated_global_ptr<T>::pointer>>;
+#else
  using global_pointer_t = typename decorated_global_ptr<T>::pointer;
+#endif // __ENABLE_USM_ADDR_SPACE__
 #else
  using global_pointer_t = T *;
-#endif
+#endif // __SYCL_DEVICE_ONLY__
 
  global_pointer_t m_Ptr;
 

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

@@ -62,6 +62,24 @@ struct PropertyMetaInfo<usm_kind_key::value_t<Kind>> {
  static constexpr sycl::usm::alloc value = Kind;
 };
 
+template <typename PropertyListT> struct IsUsmKindDevice : std::false_type {};
+template <typename... Props>
+struct IsUsmKindDevice<detail::properties_t<Props...>>
+ : detail::ContainsProperty<std::remove_const_t<decltype(usm_kind_device)>,
+ std::tuple<Props...>> {};
+
+template <typename PropertyListT> struct IsUsmKindHost : std::false_type {};
+template <typename... Props>
+struct IsUsmKindHost<detail::properties_t<Props...>>
+ : detail::ContainsProperty<std::remove_const_t<decltype(usm_kind_host)>,
+ std::tuple<Props...>> {};
+
+template <typename PropertyListT> struct IsUsmKindShared : std::false_type {};
+template <typename... Props>
+struct IsUsmKindShared<detail::properties_t<Props...>>
+ : detail::ContainsProperty<std::remove_const_t<decltype(usm_kind_shared)>,
+ std::tuple<Props...>> {};
+
 } // namespace detail
 
 } // namespace experimental

sycl/test/extensions/annotated_arg/annotated_ptr.cpp renamed to sycl/test/extensions/annotated_ptr/annotated_ptr.cpp

@@ -0,0 +1,56 @@
+// RUN: %clangxx -fsycl-device-only -fsycl-targets=spir64_fpga -S -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR
+
+// Tests that the address space of `annotated_ptr` kernel argument is refined
+// when:
+// 1. `usm_kind` property is specified in annotated_ptr type, and
+// 2. both `__SYCL_DEVICE_ONLY__` and `__ENABLE_USM_ADDR_SPACE__` are turned on
+// (i.e. equiv to flags "-fsycl-device-only -fsycl-targets=spir64_fpga")
+
+#include "sycl/sycl.hpp"
+
+#include <sycl/ext/intel/fpga_extensions.hpp>
+
+#include <iostream>
+
+using namespace sycl;
+using namespace ext::oneapi::experimental;
+using namespace ext::intel::experimental;
+
+using annotated_ptr_t1 =
+ annotated_ptr<int,
+ decltype(properties(buffer_location<0>, usm_kind_device))>;
+using annotated_ptr_t2 =
+ annotated_ptr<int, decltype(properties(buffer_location<1>, usm_kind_host))>;
+using annotated_ptr_t3 =
+ annotated_ptr<int,
+ decltype(properties(buffer_location<2>, usm_kind_shared))>;
+
+struct MyIP {
+
+ // CHECK-IR: spir_kernel void @_ZTS4MyIP(ptr addrspace(5) {{.*}} %_arg_a, ptr addrspace(6) {{.*}} %_arg_b, ptr addrspace(1) {{.*}} %_arg_c)
+ annotated_ptr_t1 a;
+ annotated_ptr_t2 b;
+ annotated_ptr_t3 c;
+
+ MyIP(int *a_, int *b_, int *c_) : a(a_), b(b_), c(c_) {}
+
+ void operator()() const { *a = *b + *c; }
+};
+
+void TestVectorAddWithAnnotatedMMHosts() {
+ queue q;
+ auto p1 = malloc_device<int>(5, q);
+ auto p2 = malloc_host<int>(5, q);
+ auto p3 = malloc_shared<int>(5, q);
+
+ q.submit([&](handler &h) { h.single_task(MyIP{p1, p2, p3}); }).wait();
+
+ free(p1, q);
+ free(p2, q);
+ free(p3, q);
+}
+
+int main() {
+ TestVectorAddWithAnnotatedMMHosts();
+ return 0;
+}