[SYCL][FPGA] Align clang with new spec of accessor_property_list

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -11033,9 +11033,10 @@ def err_sycl_compiletime_property_duplication : Error<
 def err_sycl_invalid_property_list_param_number : Error<
   "%0 must have exactly one template parameter">;
 def err_sycl_invalid_accessor_property_template_param : Error<
-  "Fifth template parameter of the accessor must be of a property_list type">;
-def err_sycl_invalid_property_list_template_param : Error<
-  "%select{property_list|property_list pack argument|buffer_location}0 "
+  "sixth template parameter of the accessor must be of accessor_property_list "
+  "type">;
+def err_sycl_invalid_accessor_property_list_template_param : Error<
+  "%select{accessor_property_list|accessor_property_list pack argument|buffer_location}0 "
   "template parameter must be a "
   "%select{parameter pack|type|non-negative integer}1">;
 def warn_sycl_pass_by_value_deprecated 

clang/lib/Sema/SemaSYCL.cpp

@@ -82,8 +82,8 @@ class Util {
   static bool isSyclHalfType(const QualType &Ty);
 
   /// Checks whether given clang type is a full specialization of the SYCL
-  /// property_list class.
-  static bool isPropertyListType(const QualType &Ty);
+  /// accessor_property_list class.
+  static bool isAccessorPropertyListType(const QualType &Ty);
 
   /// Checks whether given clang type is a full specialization of the SYCL
   /// buffer_location class.
@@ -1194,29 +1194,31 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
       return;
     }
     QualType PropListTy = PropList.getAsType();
-    if (!Util::isPropertyListType(PropListTy)) {
+    if (!Util::isAccessorPropertyListType(PropListTy)) {
       SemaRef.Diag(Loc,
                    diag::err_sycl_invalid_accessor_property_template_param);
       return;
     }
-    const auto *PropListDecl =
+    const auto *AccPropListDecl =
         cast<ClassTemplateSpecializationDecl>(PropListTy->getAsRecordDecl());
-    if (PropListDecl->getTemplateArgs().size() != 1) {
+    if (AccPropListDecl->getTemplateArgs().size() != 1) {
       SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_param_number)
-          << "property_list";
+          << "accessor_property_list";
       return;
     }
-    const auto TemplArg = PropListDecl->getTemplateArgs()[0];
+    const auto TemplArg = AccPropListDecl->getTemplateArgs()[0];
     if (TemplArg.getKind() != TemplateArgument::ArgKind::Pack) {
-      SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param)
-          << /*property_list*/ 0 << /*parameter pack*/ 0;
+      SemaRef.Diag(Loc,
+                   diag::err_sycl_invalid_accessor_property_list_template_param)
+          << /*accessor_property_list*/ 0 << /*parameter pack*/ 0;
       return;
     }
     for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin();
          Prop != TemplArg.pack_end(); ++Prop) {
       if (Prop->getKind() != TemplateArgument::ArgKind::Type) {
-        SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param)
-            << /*property_list pack argument*/ 1 << /*type*/ 1;
+        SemaRef.Diag(
+            Loc, diag::err_sycl_invalid_accessor_property_list_template_param)
+            << /*accessor_property_list pack argument*/ 1 << /*type*/ 1;
         return;
       }
       QualType PropTy = Prop->getAsType();
@@ -1235,13 +1237,15 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
     }
     const auto BufferLoc = PropDecl->getTemplateArgs()[0];
     if (BufferLoc.getKind() != TemplateArgument::ArgKind::Integral) {
-      SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param)
+      SemaRef.Diag(Loc,
+                   diag::err_sycl_invalid_accessor_property_list_template_param)
           << /*buffer_location*/ 2 << /*non-negative integer*/ 2;
       return;
     }
     int LocationID = static_cast<int>(BufferLoc.getAsIntegral().getExtValue());
     if (LocationID < 0) {
-      SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param)
+      SemaRef.Diag(Loc,
+                   diag::err_sycl_invalid_accessor_property_list_template_param)
           << /*buffer_location*/ 2 << /*non-negative integer*/ 2;
       return;
     }
@@ -1402,19 +1406,18 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
   }
 
   // Handle accessor properties. If any properties were found in
-  // the property_list - add the appropriate attributes to ParmVarDecl.
+  // the accessor_property_list - add the appropriate attributes to ParmVarDecl.
   void handleAccessorPropertyList(ParmVarDecl *Param,
                                   const CXXRecordDecl *RecordDecl,
                                   SourceLocation Loc) {
     const auto *AccTy = cast<ClassTemplateSpecializationDecl>(RecordDecl);
-    // TODO: when SYCL headers' part is ready - replace this 'if' with an error
     if (AccTy->getTemplateArgs().size() < 6)
       return;
     const auto PropList = cast<TemplateArgument>(AccTy->getTemplateArgs()[5]);
     QualType PropListTy = PropList.getAsType();
-    const auto *PropListDecl =
+    const auto *AccPropListDecl =
         cast<ClassTemplateSpecializationDecl>(PropListTy->getAsRecordDecl());
-    const auto TemplArg = PropListDecl->getTemplateArgs()[0];
+    const auto TemplArg = AccPropListDecl->getTemplateArgs()[0];
     // Move through TemplateArgs list of a property list and search for
     // properties. If found - apply the appropriate attribute to ParmVarDecl.
     for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin();
@@ -3443,20 +3446,17 @@ bool Util::isSyclSpecConstantType(const QualType &Ty) {
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
-bool Util::isPropertyListType(const QualType &Ty) {
-  return isSyclType(Ty, "property_list", true /*Tmpl*/);
-}
-
 bool Util::isSyclBufferLocationType(const QualType &Ty) {
-  const StringRef &Name = "buffer_location";
-  std::array<DeclContextDesc, 4> Scopes = {
-      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
-      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
-      // TODO: this doesn't belong to property namespace, instead it shall be
-      // in its own namespace. Change it, when the actual implementation in SYCL
-      // headers is ready
-      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "property"},
-      Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}};
+  const StringRef &PropertyName = "buffer_location";
+  const StringRef &InstanceName = "instance";
+  std::array<DeclContextDesc, 6> Scopes = {
+      Util::DeclContextDesc{Decl::Kind::Namespace, "cl"},
+      Util::DeclContextDesc{Decl::Kind::Namespace, "sycl"},
+      Util::DeclContextDesc{Decl::Kind::Namespace, "INTEL"},
+      Util::DeclContextDesc{Decl::Kind::Namespace, "property"},
+      Util::DeclContextDesc{Decl::Kind::CXXRecord, PropertyName},
+      Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization,
+                            InstanceName}};
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
@@ -3470,6 +3470,16 @@ bool Util::isSyclType(const QualType &Ty, StringRef Name, bool Tmpl) {
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
+bool Util::isAccessorPropertyListType(const QualType &Ty) {
+  const StringRef &Name = "accessor_property_list";
+  std::array<DeclContextDesc, 4> Scopes = {
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "ONEAPI"},
+      Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}};
+  return matchQualifiedTypeName(Ty, Scopes);
+}
+
 bool Util::matchQualifiedTypeName(const QualType &Ty,
                                   ArrayRef<Util::DeclContextDesc> Scopes) {
   // The idea: check the declaration context chain starting from the type

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -71,19 +71,11 @@ enum prop_type {
   base_prop
 };
 
-// Compile time known accessor property
-// TODO: this doesn't belong to property namespace, instead it shall be in its
-// own namespace. Change it, when the actual implementation in SYCL headers is
-// ready
-template <int>
-class buffer_location {};
-
 struct property_base {
   virtual prop_type type() const = 0;
 };
 } // namespace property
 
-template <typename... properties>
 class property_list {
 public:
   template <typename... propertiesTN>
@@ -102,6 +94,20 @@ class property_list {
   bool operator!=(const property_list &rhs) const { return false; }
 };
 
+namespace INTEL {
+namespace property {
+// Compile time known accessor property
+struct buffer_location {
+  template <int> class instance {};
+};
+} // namespace property
+} // namespace INTEL
+
+namespace ONEAPI {
+template <typename... properties>
+class accessor_property_list {};
+} // namespace ONEAPI
+
 template <int dim>
 struct id {
   template <typename... T>
@@ -136,7 +142,7 @@ struct _ImplT {
 template <typename dataT, int dimensions, access::mode accessmode,
           access::target accessTarget = access::target::global_buffer,
           access::placeholder isPlaceholder = access::placeholder::false_t,
-          typename propertyListT = property_list<>>
+          typename propertyListT = ONEAPI::accessor_property_list<>>
 class accessor {
 
 public:
@@ -150,8 +156,6 @@ class accessor {
 private:
   void __init(__attribute__((opencl_global)) dataT *Ptr, range<dimensions> AccessRange,
               range<dimensions> MemRange, id<dimensions> Offset) {}
-
-  propertyListT prop_list;
 };
 
 template <int dimensions, access::mode accessmode, access::target accesstarget>
@@ -339,8 +343,7 @@ const stream& operator<<(const stream &S, T&&) {
 }
 
 template <typename T, int dimensions = 1,
-          typename AllocatorT = int /*fake type as AllocatorT is not used*/,
-          typename... properties>
+          typename AllocatorT = int /*fake type as AllocatorT is not used*/>
 class buffer {
 public:
   using value_type = T;
@@ -352,13 +355,13 @@ class buffer {
   buffer(ParamTypes... args) {} // fake constructor
 
   buffer(const range<dimensions> &bufferRange,
-         const property_list<properties...> &propList = {}) {}
+         const property_list &propList = {}) {}
 
   buffer(T *hostData, const range<dimensions> &bufferRange,
-         const property_list<properties...> &propList = {}) {}
+         const property_list &propList = {}) {}
 
   buffer(const T *hostData, const range<dimensions> &bufferRange,
-         const property_list<properties...> &propList = {}) {}
+         const property_list &propList = {}) {}
 
   buffer(const buffer &rhs) = default;
 
@@ -426,12 +429,12 @@ enum class image_channel_type : unsigned int {
   fp32
 };
 
-template <int dimensions = 1, typename AllocatorT = int, typename... properties>
+template <int dimensions = 1, typename AllocatorT = int>
 class image {
 public:
   image(image_channel_order Order, image_channel_type Type,
         const range<dimensions> &Range,
-        const property_list<properties...> &PropList = {}) {}
+        const property_list &PropList = {}) {}
 
   /* -- common interface members -- */
 

clang/test/CodeGenSYCL/accessor_inheritance.cpp

@@ -67,13 +67,13 @@ int main() {
 // CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 2
 // CHECK: [[ACC1_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
 // Default constructor call
-// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_13property_listIJEEEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]])
+// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]])
 // CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to i8*
-// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 24
+// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 20
 // CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8* [[GEP1]] to %"class{{.*}}cl::sycl::accessor"*
 // CHECK: [[ACC2_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[BITCAST2]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
 // Default constructor call
-// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_13property_listIJEEEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]])
+// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]])
 
 // CHECK C field initialization
 // CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured* [[GEP]], i32 0, i32 2

clang/test/CodeGenSYCL/buffer_location.cpp

@@ -10,17 +10,17 @@ struct Base {
   cl::sycl::accessor<char, 1, cl::sycl::access::mode::read,
                      cl::sycl::access::target::global_buffer,
                      cl::sycl::access::placeholder::false_t,
-                     cl::sycl::property_list<
-                         cl::sycl::property::buffer_location<2>>>
+                     cl::sycl::ONEAPI::accessor_property_list<
+                         cl::sycl::INTEL::property::buffer_location::instance<2>>>
       AccField;
 };
 
 struct Captured : Base,
                   cl::sycl::accessor<char, 1, cl::sycl::access::mode::read,
                                      cl::sycl::access::target::global_buffer,
                                      cl::sycl::access::placeholder::false_t,
-                                     cl::sycl::property_list<
-                                         cl::sycl::property::buffer_location<2>>> {
+                                     cl::sycl::ONEAPI::accessor_property_list<
+                                         cl::sycl::INTEL::property::buffer_location::instance<2>>> {
   int C;
 };
 
@@ -29,8 +29,8 @@ int main() {
   cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
                      cl::sycl::access::target::global_buffer,
                      cl::sycl::access::placeholder::false_t,
-                     cl::sycl::property_list<
-                         cl::sycl::property::buffer_location<3>>>
+                     cl::sycl::ONEAPI::accessor_property_list<
+                         cl::sycl::INTEL::property::buffer_location::instance<3>>>
       accessorA;
   cl::sycl::kernel_single_task<class kernel_function>(
       [=]() {

clang/test/CodeGenSYCL/integration_header.cpp

@@ -31,18 +31,18 @@
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 1, 4 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 8 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 12 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 28 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 48 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 24 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 40 },
 // CHECK-EMPTY:
 // CHECK-NEXT:   //--- _ZTSN16second_namespace13second_kernelIcEE
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 4 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 24 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 16 },
 // CHECK-EMPTY:
 // CHECK-NEXT:   //--- _ZTS12third_kernelILi1Ei5pointIZ4mainE1XEE
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 4 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 24 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 16 },
 // CHECK-EMPTY:
 // CHECK-NEXT:   //--- _ZTS13fourth_kernelIJN15template_arg_ns14namespaced_argILi1EEEEE
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
@@ -52,11 +52,11 @@
 // CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 0 },
 // CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 4 },
 // CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 8 },
-// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 24 },
-// CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 28 },
-// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 44 },
-// CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 48 },
-// CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 64 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 20 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 24 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 36 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 40 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 52 },
 // CHECK-EMPTY:
 // CHECK-NEXT: };
 //

clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp

@@ -21,7 +21,7 @@
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT:   //--- _ZTSZ4mainE8kernel_A
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 0 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 16 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 12 },
 // CHECK-EMPTY:
 // CHECK-NEXT: };
 

clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp

@@ -21,7 +21,7 @@
 // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT: //--- _ZTSZ4mainE8kernel_C
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 0 },
-// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 16 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 12 },
 // CHECK-EMPTY:
 // CHECK-NEXT: };
 

clang/test/CodeGenSYCL/struct_kernel_param.cpp

@@ -4,12 +4,12 @@
 // CHECK:     const kernel_param_desc_t kernel_signatures[] = {
 // CHECK-NEXT:  //--- _ZTSZZ5test0vENK3$_0clERN2cl4sycl7handlerEE8MyKernel
 // CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 0 },
+// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 12 },
 // CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 16 },
 // CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 20 },
 // CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 24 },
 // CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 28 },
 // CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 32 },
-// CHECK-NEXT:  { kernel_param_kind_t::kind_std_layout, 4, 36 },
 // CHECK-EMPTY:
 // CHECK-NEXT:};