[SYCL][FPGA] Add clang support for buffer_location property

clang/include/clang/Basic/Attr.td

@@ -1188,6 +1188,17 @@ def SYCLDeviceIndirectlyCallable : InheritableAttr {
   let PragmaAttributeSupport = 0;
 }
 
+def SYCLIntelBufferLocation : InheritableAttr {
+  // No spelling, as this attribute can't be created in the source code.
+  let Spellings = [];
+  let Args = [UnsignedArgument<"LocationID">];
+  let LangOpts = [SYCLIsDevice, SYCLIsHost];
+  let Subjects = SubjectList<[Function], ErrorDiag>;
+  let Documentation = [SYCLIntelBufferLocationAttrDocs];
+  let HasCustomParsing = 1;
+  let PragmaAttributeSupport = 0;
+}
+
 def SYCLIntelKernelArgsRestrict : InheritableAttr {
   let Spellings = [ CXX11<"intel", "kernel_args_restrict"> ];
   let Subjects = SubjectList<[Function], ErrorDiag>;

clang/include/clang/Basic/AttrDocs.td

@@ -1994,6 +1994,24 @@ can be lowered.
   }];
 }
 
+def SYCLIntelBufferLocationAttrDocs : Documentation {
+  let Category = DocCatFunction;
+  let Heading = "kernel_args_buffer_location";
+  let Content = [{
+The attribute is being implicitly applied to an OpenCL kernel parameters
+generated from SYCL kernel object. It accepts a non-negative compiletime known
+integer. It is not allowed to use the attribute explicitly in SYCL code.
+
+When the attribute is applied, clang generates metadata attached to OpenCL
+kernel. Number of values stored in the metadata is the same as number of kernel
+parameters. Order of metadata values is following the order of kernel
+parameters. Metadata values are of an integer type and is being set
+accordingly values passed through accessor property ``buffer_location``. This
+values are mapped in hardware backend to the actual locations of buffers
+(DDR, QDR etc). Default value passed in the metadata is '-1'.
+  }];
+}
+
 def SYCLIntelKernelArgsRestrictDocs : Documentation {
   let Category = DocCatVariable;
   let Heading = "kernel_args_restrict";

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10969,6 +10969,12 @@ def warn_boolean_attribute_argument_is_not_valid: Warning<
 def err_sycl_attibute_cannot_be_applied_here
     : Error<"%0 attribute cannot be applied to a "
             "static function or function in an anonymous namespace">;
+def err_sycl_compiletime_property_duplication : Error<
+  "Can't apply %0 property twice to the same accessor">;
+def err_sycl_invalid_property_template_param : Error<
+  "%0 template parameter must be a "
+  "%select{parameter pack|type|compiletime known non-negative integer|"
+  "property_list}1">;
 def warn_sycl_attibute_function_raw_ptr
     : Warning<"SYCL 1.2.1 specification does not allow %0 attribute applied "
               "to a function with a raw pointer "

clang/lib/CodeGen/CodeGenModule.cpp

@@ -1413,6 +1413,9 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
   // MDNode for the kernel argument names.
   SmallVector<llvm::Metadata *, 8> argNames;
 
+  // MDNode for the intel_buffer_location attribute.
+  SmallVector<llvm::Metadata *, 8> argSYCLBufferLocationAttr;
+
   if (FD && CGF)
     for (unsigned i = 0, e = FD->getNumParams(); i != e; ++i) {
       const ParmVarDecl *parm = FD->getParamDecl(i);
@@ -1536,6 +1539,14 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
 
       // Get argument name.
       argNames.push_back(llvm::MDString::get(VMContext, parm->getName()));
+
+      auto *SYCLBufferLocationAttr =
+          parm->getAttr<SYCLIntelBufferLocationAttr>();
+      argSYCLBufferLocationAttr.push_back(
+          (SYCLBufferLocationAttr)
+              ? llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(
+                    SYCLBufferLocationAttr->getLocationID()))
+              : llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(-1)));
     }
 
   Fn->setMetadata("kernel_arg_addr_space",
@@ -1551,6 +1562,9 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
   if (getCodeGenOpts().EmitOpenCLArgMetadata)
     Fn->setMetadata("kernel_arg_name",
                     llvm::MDNode::get(VMContext, argNames));
+  if (LangOpts.SYCLIsDevice)
+    Fn->setMetadata("kernel_arg_buffer_location",
+                    llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr));
 }
 
 /// Determines whether the language options require us to model

clang/lib/Sema/SemaSYCL.cpp

@@ -80,6 +80,14 @@ class Util {
   /// half class.
   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);
+
+  /// Checks whether given clang type is a full specialization of the SYCL
+  /// buffer_location class.
+  static bool isSyclBufferLocationType(const QualType &Ty);
+
   /// Checks whether given clang type is a standard SYCL API class with given
   /// name.
   /// \param Ty    the clang type being checked
@@ -1172,39 +1180,108 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
   // iterator as push_back invalidates iterators.
   size_t LastParamIndex = 0;
 
-  void addParam(const FieldDecl *FD, QualType FieldTy) {
+  void addParam(const FieldDecl *FD, QualType FieldTy, int LocationID = -1) {
     const ConstantArrayType *CAT =
         SemaRef.getASTContext().getAsConstantArrayType(FieldTy);
     if (CAT)
       FieldTy = CAT->getElementType();
     ParamDesc newParamDesc = makeParamDesc(FD, FieldTy);
-    addParam(newParamDesc, FieldTy);
+    addParam(newParamDesc, FieldTy, LocationID);
   }
 
-  void addParam(const CXXBaseSpecifier &BS, QualType FieldTy) {
+  void addParam(const CXXBaseSpecifier &BS, QualType FieldTy,
+                int LocationID = -1) {
     ParamDesc newParamDesc =
         makeParamDesc(SemaRef.getASTContext(), BS, FieldTy);
-    addParam(newParamDesc, FieldTy);
+    addParam(newParamDesc, FieldTy, LocationID);
   }
 
-  void addParam(ParamDesc newParamDesc, QualType FieldTy) {
+  void addParam(ParamDesc newParamDesc, QualType FieldTy, int LocationID) {
     // Create a new ParmVarDecl based on the new info.
+    ASTContext &Ctx = SemaRef.getASTContext();
     auto *NewParam = ParmVarDecl::Create(
-        SemaRef.getASTContext(), KernelDecl, SourceLocation(), SourceLocation(),
+        Ctx, KernelDecl, SourceLocation(), SourceLocation(),
         std::get<1>(newParamDesc), std::get<0>(newParamDesc),
         std::get<2>(newParamDesc), SC_None, /*DefArg*/ nullptr);
     NewParam->setScopeInfo(0, Params.size());
     NewParam->setIsUsed();
 
     LastParamIndex = Params.size();
     Params.push_back(NewParam);
+    if (LocationID != -1)
+      NewParam->addAttr(
+          SYCLIntelBufferLocationAttr::CreateImplicit(Ctx, LocationID));
+  }
+
+  // Obtain an integer value stored in a template parameter of buffer_location
+  // property to pass it to buffer_location kernel attribute
+  int handleBufferLocationProperty(QualType FieldTy, SourceLocation Loc) {
+    const auto *AccTy =
+        cast<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl());
+    // TODO: when SYCL headers' part is ready - replace this 'if' with an error
+    if (AccTy->getTemplateArgs().size() < 6)
+      return -1;
+    const auto PropList = cast<TemplateArgument>(AccTy->getTemplateArgs()[5]);
+    if (PropList.getKind() != TemplateArgument::ArgKind::Type) {
+      SemaRef.Diag(Loc, diag::err_sycl_invalid_property_template_param)
+          << "accessor's 5th" << /*type*/ 1;
+      return -1;
+    }
+    QualType PropListTy = PropList.getAsType();
+    if (!Util::isPropertyListType(PropListTy)) {
+      SemaRef.Diag(Loc, diag::err_sycl_invalid_property_template_param)
+          << "accessor's 5th" << /*property_list*/ 3;
+      return -1;
+    }
+
+    int LocationID = -1;
+    const auto *PropListDecl =
+        cast<ClassTemplateSpecializationDecl>(PropListTy->getAsRecordDecl());
+    const auto TemplArg = PropListDecl->getTemplateArgs()[0];
+    if (TemplArg.getKind() != TemplateArgument::ArgKind::Pack) {
+      SemaRef.Diag(Loc, diag::err_sycl_invalid_property_template_param)
+          << "property_list" << /*parameter pack*/ 0;
+      return -1;
+    }
+    // Move through TemplateArgs list of a property list and search for
+    // buffer_location property. If found - return the stored integer value in
+    // its template parameter, if not - return -1.
+    for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin();
+         Prop != TemplArg.pack_end(); ++Prop) {
+      QualType PropTy = Prop->getAsType();
+      if (Util::isSyclBufferLocationType(PropTy)) {
+        // If we have more than 1 buffer_location properties on a single
+        // accessor - emit an error
+        if (LocationID != -1) {
+          SemaRef.Diag(Loc, diag::err_sycl_compiletime_property_duplication)
+              << "buffer_location";
+          return -1;
+        }
+        const auto *PropDecl =
+            cast<ClassTemplateSpecializationDecl>(PropTy->getAsRecordDecl());
+        const auto BufferLoc = PropDecl->getTemplateArgs()[0];
+        if (BufferLoc.getKind() != TemplateArgument::ArgKind::Integral) {
+          SemaRef.Diag(Loc, diag::err_sycl_invalid_property_template_param)
+              << "buffer_location" << /*non-negative integer*/ 2;
+          return -1;
+        }
+        LocationID = static_cast<int>(BufferLoc.getAsIntegral().getExtValue());
+        if (LocationID < 0) {
+          SemaRef.Diag(Loc, diag::err_sycl_invalid_property_template_param)
+              << "buffer_location" << /*non-negative integer*/ 2;
+          return -1;
+        }
+      }
+    }
+    return LocationID;
   }
 
   // All special SYCL objects must have __init method. We extract types for
   // kernel parameters from __init method parameters. We will use __init method
   // and kernel parameters which we build here to initialize special objects in
   // the kernel body.
-  bool handleSpecialType(FieldDecl *FD, QualType FieldTy) {
+  bool handleSpecialType(FieldDecl *FD, QualType FieldTy,
+                         bool isAccessorType = false) {
     const auto *RecordDecl = FieldTy->getAsCXXRecordDecl();
     assert(RecordDecl && "The accessor/sampler must be a RecordDecl");
     CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName);
@@ -1213,8 +1290,19 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     // Don't do -1 here because we count on this to be the first parameter added
     // (if any).
     size_t ParamIndex = Params.size();
-    for (const ParmVarDecl *Param : InitMethod->parameters())
-      addParam(FD, Param->getType().getCanonicalType());
+    auto ParamIt = InitMethod->parameters().begin();
+    if (*ParamIt) {
+      // Add meaningful argument (not '-1') to buffer_location attribute only
+      // for an accessor pointer
+      size_t BufferLocAttrArg =
+          isAccessorType
+              ? handleBufferLocationProperty(FieldTy, FD->getLocation())
+              : -1;
+      addParam(FD, (*ParamIt)->getType().getCanonicalType(), BufferLocAttrArg);
+      ++ParamIt;
+      for (; ParamIt != InitMethod->parameters().end(); ++ParamIt)
+        addParam(FD, (*ParamIt)->getType().getCanonicalType(), -1);
+    }
     LastParamIndex = ParamIndex;
     return true;
   }
@@ -1284,14 +1372,23 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     // Don't do -1 here because we count on this to be the first parameter added
     // (if any).
     size_t ParamIndex = Params.size();
-    for (const ParmVarDecl *Param : InitMethod->parameters())
-      addParam(BS, Param->getType().getCanonicalType());
+    auto ParamIt = InitMethod->parameters().begin();
+    if (*ParamIt) {
+      // Add meaningful argument (not '-1') to buffer_location attribute only
+      // for an accessor pointer
+      size_t BufferLocAttrArg =
+          handleBufferLocationProperty(FieldTy, BS.getBeginLoc());
+      addParam(BS, (*ParamIt)->getType().getCanonicalType(), BufferLocAttrArg);
+      ++ParamIt;
+      for (; ParamIt != InitMethod->parameters().end(); ++ParamIt)
+        addParam(BS, (*ParamIt)->getType().getCanonicalType(), -1);
+    }
     LastParamIndex = ParamIndex;
     return true;
   }
 
   bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final {
-    return handleSpecialType(FD, FieldTy);
+    return handleSpecialType(FD, FieldTy, /*isAccessorType*/ true);
   }
 
   bool handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final {
@@ -2820,6 +2917,23 @@ 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}};
+  return matchQualifiedTypeName(Ty, Scopes);
+}
+
 bool Util::isSyclType(const QualType &Ty, StringRef Name, bool Tmpl) {
   Decl::Kind ClassDeclKind =
       Tmpl ? Decl::Kind::ClassTemplateSpecialization : Decl::Kind::CXXRecord;

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -71,15 +71,23 @@ 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... propertyTN>
-  property_list(propertyTN... props) {}
+  template <typename... propertiesTN>
+  property_list(propertiesTN... props){};
 
   template <typename propertyT>
   bool has_property() const { return true; }
@@ -127,7 +135,8 @@ struct _ImplT {
 
 template <typename dataT, int dimensions, access::mode accessmode,
           access::target accessTarget = access::target::global_buffer,
-          access::placeholder isPlaceholder = access::placeholder::false_t>
+          access::placeholder isPlaceholder = access::placeholder::false_t,
+          typename propertyListT = property_list<>>
 class accessor {
 
 public:
@@ -141,6 +150,8 @@ 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>
@@ -326,7 +337,8 @@ const stream& operator<<(const stream &S, T&&) {
 }
 
 template <typename T, int dimensions = 1,
-          typename AllocatorT = int /*fake type as AllocatorT is not used*/>
+          typename AllocatorT = int /*fake type as AllocatorT is not used*/,
+          typename... properties>
 class buffer {
 public:
   using value_type = T;
@@ -338,13 +350,13 @@ class buffer {
   buffer(ParamTypes... args) {} // fake constructor
 
   buffer(const range<dimensions> &bufferRange,
-         const property_list &propList = {}) {}
+         const property_list<properties...> &propList = {}) {}
 
   buffer(T *hostData, const range<dimensions> &bufferRange,
-         const property_list &propList = {}) {}
+         const property_list<properties...> &propList = {}) {}
 
   buffer(const T *hostData, const range<dimensions> &bufferRange,
-         const property_list &propList = {}) {}
+         const property_list<properties...> &propList = {}) {}
 
   buffer(const buffer &rhs) = default;
 
@@ -412,11 +424,12 @@ enum class image_channel_type : unsigned int {
   fp32
 };
 
-template <int dimensions = 1, typename AllocatorT = int>
+template <int dimensions = 1, typename AllocatorT = int, typename... properties>
 class image {
 public:
   image(image_channel_order Order, image_channel_type Type,
-        const range<dimensions> &Range, const property_list &PropList = {}) {}
+        const range<dimensions> &Range,
+        const property_list<properties...> &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_11placeholderE0EEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]])
+// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_13property_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 20
+// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 24
 // 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_11placeholderE0EEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]])
+// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_13property_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