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

clang/include/clang/Basic/Attr.td

@@ -1188,6 +1188,28 @@ def SYCLDeviceIndirectlyCallable : InheritableAttr {
   let PragmaAttributeSupport = 0;
 }
 
+def SYCLIntelBufferLocation : InheritableAttr {
+  let Spellings = [CXX11<"intelfpga","kernel_arg_buffer_location">];
+  let LangOpts = [SYCLIsDevice, SYCLIsHost];
+  let Subjects = SubjectList<[Function], ErrorDiag>;
+
+  let AdditionalMembers = [{
+    std::vector<size_t> ActualArgs;
+
+    void setActualArgs(std::vector<size_t> ArgVec) {
+      ActualArgs = ArgVec;
+    }
+
+    std::vector<size_t> getActualArgs() const {
+      return ActualArgs;
+    }
+  }];
+
+  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 ``intelfpga::kernel_arg_buffer_location`` is being implicitly
+applied to an OpenCL kernel generated from SYCL kernel object. 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 pointer
+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

@@ -10982,6 +10982,8 @@ def warn_sycl_implicit_decl
 def warn_sycl_restrict_recursion
     : Warning<"SYCL kernel cannot call a recursive function">,
       InGroup<SyclStrict>, DefaultError;
+def warn_sycl_implicit_attr_usage : Warning <
+  "%0 attribute cannot be used explicitly">, InGroup<IgnoredAttributes>;
 def err_ivdep_duplicate_arg : Error<
   "duplicate argument to 'ivdep'. attribute requires one or both of a safelen "
   "and array">;

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -660,6 +660,16 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
     if (A->getEnabled())
       Fn->setMetadata("no_global_work_offset", llvm::MDNode::get(Context, {}));
   }
+
+  if (const SYCLIntelBufferLocationAttr *A =
+          FD->getAttr<SYCLIntelBufferLocationAttr>()) {
+    std::vector<size_t> Args = A->getActualArgs();
+    std::vector<llvm::Metadata *> AttrMDArgs;
+    for (auto A : Args)
+      AttrMDArgs.push_back(llvm::ConstantAsMetadata::get(Builder.getInt32(A)));
+    Fn->setMetadata("kernel_arg_buffer_location",
+                    llvm::MDNode::get(Context, AttrMDArgs));
+  }
 }
 
 /// Determine whether the function F ends with a return stmt.

clang/lib/Sema/SemaDeclAttr.cpp

@@ -3082,6 +3082,11 @@ static void handleMaxGlobalWorkDimAttr(Sema &S, Decl *D,
         S.Context, Attr, MaxGlobalWorkDim));
 }
 
+// Handles kernel_arg_buffer_location attr.
+static void handleBufferLocationAttr(Sema &S, Decl *D, const ParsedAttr &Attr) {
+  S.Diag(Attr.getLoc(), diag::warn_sycl_implicit_attr_usage) << Attr;
+}
+
 static void handleVecTypeHint(Sema &S, Decl *D, const ParsedAttr &AL) {
   if (!AL.hasParsedType()) {
     S.Diag(AL.getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
@@ -7789,6 +7794,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset:
     handleNoGlobalWorkOffsetAttr(S, D, AL);
     break;
+  case ParsedAttr::AT_SYCLIntelBufferLocation:
+    handleBufferLocationAttr(S, D, AL);
+    break;
   case ParsedAttr::AT_VecTypeHint:
     handleVecTypeHint(S, D, AL);
     break;

clang/lib/Sema/SemaSYCL.cpp

@@ -80,6 +80,10 @@ class Util {
   /// half class.
   static bool isSyclHalfType(const QualType &Ty);
 
+  /// Checks whether given clang type is a full specialization of the SYCL
+  /// buffer_location class.
+  static bool isSyclBufferLocation(const QualType &Ty);
+
   /// Checks whether given clang type is a standard SYCL API class with given
   /// name.
   /// \param Ty    the clang type being checked
@@ -1171,23 +1175,28 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
   // Holds the last handled field's first parameter. This doesn't store an
   // iterator as push_back invalidates iterators.
   size_t LastParamIndex = 0;
+  // This vector stores information about buffer location. If no buffer_location
+  // property of an accessor is set - the appropriate value stored in the
+  // vector = -1.
+  std::vector<size_t> BufferLocationMD;
 
-  void addParam(const FieldDecl *FD, QualType FieldTy) {
+  void addParam(const FieldDecl *FD, QualType FieldTy, size_t 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,
+                size_t 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, size_t LocationID) {
     // Create a new ParmVarDecl based on the new info.
     auto *NewParam = ParmVarDecl::Create(
         SemaRef.getASTContext(), KernelDecl, SourceLocation(), SourceLocation(),
@@ -1198,13 +1207,40 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
 
     LastParamIndex = Params.size();
     Params.push_back(NewParam);
+    BufferLocationMD.push_back(LocationID);
+  }
+
+  // Obtain an integer value stored in a template parameter of buffer_location
+  // property to pass it to buffer_location kernel attribute
+  size_t handleBufferLocationProperty(QualType FieldTy) {
+    const auto *AccTy =
+        cast<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl());
+
+    // TODO: when SYCL headers' part is ready - replace this 'if' with an assert
+    if (AccTy->getTemplateArgs().size() < 6)
+      return -1;
+
+    // TODO: at this point of time it's unclear, what representation in LLVM IR
+    // is going to be for other compile time known accessor properties, hence
+    // it's not clear, how handle them in SemaSYCL. But in general property_list
+    // is a parameter pack and shall be handled appropriately.
+    const auto Prop = cast<TemplateArgument>(AccTy->getTemplateArgs()[5]);
+    QualType PropTy = Prop.getAsType();
+    if (!Util::isSyclBufferLocation(PropTy))
+      return -1;
+
+    const auto *PropDecl =
+        cast<ClassTemplateSpecializationDecl>(PropTy->getAsRecordDecl());
+    return static_cast<int>(
+        PropDecl->getTemplateArgs()[0].getAsIntegral().getExtValue());
   }
 
   // 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 +1249,17 @@ 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) : -1;
+      addParam(FD, (*ParamIt)->getType().getCanonicalType(), BufferLocAttrArg);
+      ++ParamIt;
+      for (; ParamIt != InitMethod->parameters().end(); ++ParamIt)
+        addParam(FD, (*ParamIt)->getType().getCanonicalType(), -1);
+    }
     LastParamIndex = ParamIndex;
     return true;
   }
@@ -1270,6 +1315,11 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     KernelDecl->setType(FuncType);
     KernelDecl->setParams(Params);
 
+    // Add SYCLIntelBufferLocationAttr to the kernel declaration
+    auto *BufferLocAttr = SYCLIntelBufferLocationAttr::CreateImplicit(Ctx);
+    BufferLocAttr->setActualArgs(BufferLocationMD);
+    KernelDecl->addAttr(BufferLocAttr);
+
     if (ArgChecker.isValid())
       SemaRef.addSyclDeviceDecl(KernelDecl);
   }
@@ -1284,14 +1334,22 @@ 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);
+      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 +2878,19 @@ bool Util::isSyclSpecConstantType(const QualType &Ty) {
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
+bool Util::isSyclBufferLocation(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,22 @@ 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) {}
+  property_list(properties... props) {}
 
   template <typename propertyT>
   bool has_property() const { return true; }
@@ -127,7 +134,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 +149,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 +336,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 +349,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 +423,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

clang/test/CodeGenSYCL/buffer_location.cpp

@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
+
+// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function{{.*}} !kernel_arg_buffer_location ![[MDBL:[0-9]+]]
+// CHECK: ![[MDBL]] = !{i32 3, i32 -1, i32 -1, i32 -1}
+
+#include "sycl.hpp"
+
+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::buffer_location<3>>
+      accessorA;
+  cl::sycl::kernel_single_task<class kernel_function>(
+      [=]() {
+        accessorA.use();
+      });
+  return 0;
+}