From 819a973e10f39ed558c5d4e0883501a34b679c1d Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Sun, 19 Jul 2020 02:29:09 +0300 Subject: [PATCH 01/12] [SYCL][FPGA] Add clang support for buffer_location property This is a compiler-time known accessor property which serves as an optimization hint for a compiler on where exactly buffer was allocated. This is needed when a board has multiple disjoint global memories that must be managed explicitly by a programmer. When the property is added as a template parameter of an accessor - SemaSYCL will implicitly add ``intelfpga::kernel_arg_buffer_location`` attribute 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'. Signed-off-by: Dmitry Sidorov --- clang/include/clang/Basic/Attr.td | 22 +++++ clang/include/clang/Basic/AttrDocs.td | 18 ++++ .../clang/Basic/DiagnosticSemaKinds.td | 2 + clang/lib/CodeGen/CodeGenFunction.cpp | 10 +++ clang/lib/Sema/SemaDeclAttr.cpp | 9 ++ clang/lib/Sema/SemaSYCL.cpp | 84 ++++++++++++++++--- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 30 +++++-- .../test/CodeGenSYCL/accessor_inheritance.cpp | 6 +- clang/test/CodeGenSYCL/buffer_location.cpp | 18 ++++ clang/test/CodeGenSYCL/integration_header.cpp | 18 ++-- .../CodeGenSYCL/kernel-param-acc-array-ih.cpp | 2 +- .../kernel-param-member-acc-array-ih.cpp | 2 +- .../test/CodeGenSYCL/struct_kernel_param.cpp | 2 +- .../intel-kernel-arg-buffer-location.cpp | 34 ++++++++ 14 files changed, 223 insertions(+), 34 deletions(-) create mode 100644 clang/test/CodeGenSYCL/buffer_location.cpp create mode 100644 clang/test/SemaSYCL/intel-kernel-arg-buffer-location.cpp diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 94d337043e3a..e3474546cbda 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/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 ActualArgs; + + void setActualArgs(std::vector ArgVec) { + ActualArgs = ArgVec; + } + + std::vector 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>; diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 467b96793263..8a797b3820eb 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/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"; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index a071825f87ed..9e14ff04275d 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/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, DefaultError; +def warn_sycl_implicit_attr_usage : Warning < + "%0 attribute cannot be used explicitly">, InGroup; def err_ivdep_duplicate_arg : Error< "duplicate argument to 'ivdep'. attribute requires one or both of a safelen " "and array">; diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index ea9dee22378b..3e174ab2d017 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/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()) { + std::vector Args = A->getActualArgs(); + std::vector 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. diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 67f9b4b9ad9f..555a3ee13c66 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3082,6 +3082,12 @@ 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 +7795,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; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8cbd33b279e8..28a21324553e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/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 + /// property_list 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 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,41 @@ 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(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(AccTy->getTemplateArgs()[5]); + QualType PropTy = Prop.getAsType(); + if (!Util::isSyclBufferLocation(PropTy)) + return -1; + + const auto *PropDecl = cast( + PropTy->getAsRecordDecl()); + return static_cast( + 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 +1250,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 +1316,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); } @@ -1285,13 +1336,13 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // (if any). size_t ParamIndex = Params.size(); for (const ParmVarDecl *Param : InitMethod->parameters()) - addParam(BS, Param->getType().getCanonicalType()); + addParam(BS, Param->getType().getCanonicalType(), 42); 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 +2871,19 @@ bool Util::isSyclSpecConstantType(const QualType &Ty) { return matchQualifiedTypeName(Ty, Scopes); } +bool Util::isSyclBufferLocation(const QualType &Ty) { + const StringRef &Name = "buffer_location"; + std::array 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; diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 3184c58edcbf..bc71850146ed 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/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 +class buffer_location {}; + struct property_base { virtual prop_type type() const = 0; }; } // namespace property +template class property_list { public: - template - property_list(propertyTN... props) {} + property_list(properties... props) {} template bool has_property() const { return true; } @@ -127,7 +134,8 @@ struct _ImplT { template + 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 AccessRange, range MemRange, id Offset) {} + + propertyListT prop_list; }; template @@ -326,7 +336,8 @@ const stream& operator<<(const stream &S, T&&) { } template + 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 &bufferRange, - const property_list &propList = {}) {} + const property_list &propList = {}) {} buffer(T *hostData, const range &bufferRange, - const property_list &propList = {}) {} + const property_list &propList = {}) {} buffer(const T *hostData, const range &bufferRange, - const property_list &propList = {}) {} + const property_list &propList = {}) {} buffer(const buffer &rhs) = default; @@ -412,11 +423,12 @@ enum class image_channel_type : unsigned int { fp32 }; -template +template class image { public: image(image_channel_order Order, image_channel_type Type, - const range &Range, const property_list &PropList = {}) {} + const range &Range, + const property_list &PropList = {}) {} /* -- common interface members -- */ diff --git a/clang/test/CodeGenSYCL/accessor_inheritance.cpp b/clang/test/CodeGenSYCL/accessor_inheritance.cpp index e197c339c125..b3857806a2bc 100644 --- a/clang/test/CodeGenSYCL/accessor_inheritance.cpp +++ b/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 diff --git a/clang/test/CodeGenSYCL/buffer_location.cpp b/clang/test/CodeGenSYCL/buffer_location.cpp new file mode 100644 index 000000000000..c39c15f3c8ec --- /dev/null +++ b/clang/test/CodeGenSYCL/buffer_location.cpp @@ -0,0 +1,18 @@ +// 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> accessorA; + cl::sycl::kernel_single_task( + [=]() { + accessorA.use(); + }); + return 0; +} diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index c63e64a37f11..e285bfce8f53 100644 --- a/clang/test/CodeGenSYCL/integration_header.cpp +++ b/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, 24 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 40 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 28 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 48 }, // 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, 16 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 24 }, // 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, 16 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 24 }, // 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, 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-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-EMPTY: // CHECK-NEXT: }; // diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp index 8c2cfb2a1bd8..902ad7ddc339 100644 --- a/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp +++ b/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, 12 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 16 }, // CHECK-EMPTY: // CHECK-NEXT: }; diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp index f5f679f7d365..ddfaca966485 100644 --- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp +++ b/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, 12 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 16 }, // CHECK-EMPTY: // CHECK-NEXT: }; diff --git a/clang/test/CodeGenSYCL/struct_kernel_param.cpp b/clang/test/CodeGenSYCL/struct_kernel_param.cpp index a00f147b0dee..b1382ec6c2b9 100644 --- a/clang/test/CodeGenSYCL/struct_kernel_param.cpp +++ b/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:}; diff --git a/clang/test/SemaSYCL/intel-kernel-arg-buffer-location.cpp b/clang/test/SemaSYCL/intel-kernel-arg-buffer-location.cpp new file mode 100644 index 000000000000..23d201bf3f8b --- /dev/null +++ b/clang/test/SemaSYCL/intel-kernel-arg-buffer-location.cpp @@ -0,0 +1,34 @@ +// RUN: %clang %s -fsyntax-only -fsycl-device-only -DCHECKDIAG -Xclang -verify +// RUN: %clang %s -fsyntax-only -I %S/Inputs -Xclang -ast-dump -fsycl-device-only | FileCheck %s + +#ifndef CHECKDIAG +#include "sycl.hpp" +#endif // CHECKDIAG + +#ifdef CHECKDIAG +struct FuncObj { + [[intelfpga::kernel_arg_buffer_location]] // expected-warning{{'kernel_arg_buffer_location' attribute cannot be used explicitly}} + void operator()() {} +}; + +template +__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { + kernelFunc(); + [[intelfpga::kernel_arg_buffer_location]] int invalid = 42; // expected-error{{'kernel_arg_buffer_location' attribute only applies to functions}} +} +#endif // CHECKDIAG + +int main() { +#ifdef CHECKDIAG + kernel( + FuncObj()); +#else + // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ4mainE15kernel_function + // CHECK: SYCLIntelBufferLocationAttr + cl::sycl::accessor accessorA; + cl::sycl::kernel_single_task( + [=]() { + accessorA.use(); + }); +#endif // CHECKDIAG +} From a350e20910765f895496fad456af23a7d8b2b9d1 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Thu, 23 Jul 2020 16:56:40 +0300 Subject: [PATCH 02/12] Fix typos Signed-off-by: Dmitry Sidorov --- clang/lib/Sema/SemaSYCL.cpp | 14 +++++++++++--- 1 file changed, 11 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 28a21324553e..02e49c958870 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -81,7 +81,7 @@ class Util { static bool isSyclHalfType(const QualType &Ty); /// Checks whether given clang type is a full specialization of the SYCL - /// property_list class. + /// buffer_location class. static bool isSyclBufferLocation(const QualType &Ty); /// Checks whether given clang type is a standard SYCL API class with given @@ -1335,8 +1335,16 @@ 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(), 42); + 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; } From 49dce750c70d0663ef975e2d07f16723ed97e23a Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Fri, 24 Jul 2020 13:57:45 +0300 Subject: [PATCH 03/12] Apply clang-format + some fixes Signed-off-by: Dmitry Sidorov --- clang/lib/Sema/SemaDeclAttr.cpp | 3 +-- clang/lib/Sema/SemaSYCL.cpp | 7 +++---- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 4 ++-- clang/test/CodeGenSYCL/buffer_location.cpp | 3 ++- clang/test/SemaSYCL/intel-kernel-arg-buffer-location.cpp | 3 ++- 5 files changed, 10 insertions(+), 10 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 555a3ee13c66..593d794d883c 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3084,8 +3084,7 @@ static void handleMaxGlobalWorkDimAttr(Sema &S, Decl *D, // 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; + S.Diag(Attr.getLoc(), diag::warn_sycl_implicit_attr_usage) << Attr; } static void handleVecTypeHint(Sema &S, Decl *D, const ParsedAttr &AL) { diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 02e49c958870..db33856c0fcd 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1224,14 +1224,13 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // 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(AccTy->getTemplateArgs()[5]); + const auto Prop = cast(AccTy->getTemplateArgs()[5]); QualType PropTy = Prop.getAsType(); if (!Util::isSyclBufferLocation(PropTy)) return -1; - const auto *PropDecl = cast( - PropTy->getAsRecordDecl()); + const auto *PropDecl = + cast(PropTy->getAsRecordDecl()); return static_cast( PropDecl->getTemplateArgs()[0].getAsIntegral().getExtValue()); } diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index bc71850146ed..0863d3d4a0b7 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -75,7 +75,7 @@ enum prop_type { // 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 +template class buffer_location {}; struct property_base { @@ -83,7 +83,7 @@ struct property_base { }; } // namespace property -template +template class property_list { public: property_list(properties... props) {} diff --git a/clang/test/CodeGenSYCL/buffer_location.cpp b/clang/test/CodeGenSYCL/buffer_location.cpp index c39c15f3c8ec..c779c41ecb78 100644 --- a/clang/test/CodeGenSYCL/buffer_location.cpp +++ b/clang/test/CodeGenSYCL/buffer_location.cpp @@ -9,7 +9,8 @@ int main() { cl::sycl::accessor> accessorA; + cl::sycl::property::buffer_location<3>> + accessorA; cl::sycl::kernel_single_task( [=]() { accessorA.use(); diff --git a/clang/test/SemaSYCL/intel-kernel-arg-buffer-location.cpp b/clang/test/SemaSYCL/intel-kernel-arg-buffer-location.cpp index 23d201bf3f8b..3e555ecf7e3a 100644 --- a/clang/test/SemaSYCL/intel-kernel-arg-buffer-location.cpp +++ b/clang/test/SemaSYCL/intel-kernel-arg-buffer-location.cpp @@ -8,7 +8,8 @@ #ifdef CHECKDIAG struct FuncObj { [[intelfpga::kernel_arg_buffer_location]] // expected-warning{{'kernel_arg_buffer_location' attribute cannot be used explicitly}} - void operator()() {} + void + operator()() {} }; template From 9e86bffebf2de28e1c3588dc4ba7dfdaed092420 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Tue, 28 Jul 2020 14:02:21 +0300 Subject: [PATCH 04/12] Apply comments 1. Move attribute from FuncDecl to ParamVarDecl 2. Add several diagnostics 3. Reimplement property handler Signed-off-by: Dmitry Sidorov --- clang/include/clang/Basic/Attr.td | 17 +-- clang/include/clang/Basic/AttrDocs.td | 10 +- .../clang/Basic/DiagnosticSemaKinds.td | 8 +- clang/lib/CodeGen/CodeGenFunction.cpp | 10 -- clang/lib/CodeGen/CodeGenModule.cpp | 14 +++ clang/lib/Sema/SemaDeclAttr.cpp | 8 -- clang/lib/Sema/SemaSYCL.cpp | 109 ++++++++++++------ clang/test/CodeGenSYCL/Inputs/sycl.hpp | 3 +- clang/test/CodeGenSYCL/buffer_location.cpp | 3 +- clang/test/SemaSYCL/Inputs/sycl.hpp | 12 +- clang/test/SemaSYCL/accessor_inheritance.cpp | 4 +- clang/test/SemaSYCL/basic-kernel-wrapper.cpp | 2 +- clang/test/SemaSYCL/buffer_location.cpp | 61 ++++++++++ .../intel-kernel-arg-buffer-location.cpp | 35 ------ clang/test/SemaSYCL/wrapped-accessor.cpp | 8 +- 15 files changed, 187 insertions(+), 117 deletions(-) create mode 100644 clang/test/SemaSYCL/buffer_location.cpp delete mode 100644 clang/test/SemaSYCL/intel-kernel-arg-buffer-location.cpp diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index e3474546cbda..58c2b91e6b9f 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1189,22 +1189,11 @@ def SYCLDeviceIndirectlyCallable : InheritableAttr { } def SYCLIntelBufferLocation : InheritableAttr { - let Spellings = [CXX11<"intelfpga","kernel_arg_buffer_location">]; + // 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 AdditionalMembers = [{ - std::vector ActualArgs; - - void setActualArgs(std::vector ArgVec) { - ActualArgs = ArgVec; - } - - std::vector getActualArgs() const { - return ActualArgs; - } - }]; - let Documentation = [SYCLIntelBufferLocationAttrDocs]; let HasCustomParsing = 1; let PragmaAttributeSupport = 0; diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 8a797b3820eb..6884ab5db60b 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -1998,14 +1998,14 @@ 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. +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 pointer -kernel parameters. Metadata values are of an integer type and is being set +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'. diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 9e14ff04275d..68c55a355b02 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/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 " @@ -10982,8 +10988,6 @@ def warn_sycl_implicit_decl def warn_sycl_restrict_recursion : Warning<"SYCL kernel cannot call a recursive function">, InGroup, DefaultError; -def warn_sycl_implicit_attr_usage : Warning < - "%0 attribute cannot be used explicitly">, InGroup; def err_ivdep_duplicate_arg : Error< "duplicate argument to 'ivdep'. attribute requires one or both of a safelen " "and array">; diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 3e174ab2d017..ea9dee22378b 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -660,16 +660,6 @@ 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()) { - std::vector Args = A->getActualArgs(); - std::vector 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. diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index eb72b21e35c3..0f1db8af3d81 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1413,6 +1413,9 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, // MDNode for the kernel argument names. SmallVector argNames; + // MDNode for the intel_buffer_location attribute. + SmallVector 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(); + 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 diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 593d794d883c..67f9b4b9ad9f 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3082,11 +3082,6 @@ 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; @@ -7794,9 +7789,6 @@ 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; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index db33856c0fcd..d547f5f9c702 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -80,9 +80,13 @@ 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 isSyclBufferLocation(const QualType &Ty); + static bool isSyclBufferLocationType(const QualType &Ty); /// Checks whether given clang type is a standard SYCL API class with given /// name. @@ -1175,12 +1179,8 @@ 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 BufferLocationMD; - void addParam(const FieldDecl *FD, QualType FieldTy, size_t LocationID = -1) { + void addParam(const FieldDecl *FD, QualType FieldTy, int LocationID = -1) { const ConstantArrayType *CAT = SemaRef.getASTContext().getAsConstantArrayType(FieldTy); if (CAT) @@ -1190,16 +1190,17 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { } void addParam(const CXXBaseSpecifier &BS, QualType FieldTy, - size_t LocationID = -1) { + int LocationID = -1) { ParamDesc newParamDesc = makeParamDesc(SemaRef.getASTContext(), BS, FieldTy); addParam(newParamDesc, FieldTy, LocationID); } - void addParam(ParamDesc newParamDesc, QualType FieldTy, size_t LocationID) { + 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()); @@ -1207,32 +1208,72 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { LastParamIndex = Params.size(); Params.push_back(NewParam); - BufferLocationMD.push_back(LocationID); + 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 - size_t handleBufferLocationProperty(QualType FieldTy) { + int handleBufferLocationProperty(QualType FieldTy, SourceLocation Loc) { const auto *AccTy = cast(FieldTy->getAsRecordDecl()); - - // TODO: when SYCL headers' part is ready - replace this 'if' with an assert + // TODO: when SYCL headers' part is ready - replace this 'if' with an error 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(AccTy->getTemplateArgs()[5]); - QualType PropTy = Prop.getAsType(); - if (!Util::isSyclBufferLocation(PropTy)) + const auto PropList = cast(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; + } - const auto *PropDecl = - cast(PropTy->getAsRecordDecl()); - return static_cast( - PropDecl->getTemplateArgs()[0].getAsIntegral().getExtValue()); + int LocationID = -1; + const auto *PropListDecl = + cast(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(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(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 @@ -1254,7 +1295,9 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // Add meaningful argument (not '-1') to buffer_location attribute only // for an accessor pointer size_t BufferLocAttrArg = - isAccessorType ? handleBufferLocationProperty(FieldTy) : -1; + isAccessorType + ? handleBufferLocationProperty(FieldTy, FD->getLocation()) + : -1; addParam(FD, (*ParamIt)->getType().getCanonicalType(), BufferLocAttrArg); ++ParamIt; for (; ParamIt != InitMethod->parameters().end(); ++ParamIt) @@ -1315,11 +1358,6 @@ 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); } @@ -1338,7 +1376,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { if (*ParamIt) { // Add meaningful argument (not '-1') to buffer_location attribute only // for an accessor pointer - size_t BufferLocAttrArg = handleBufferLocationProperty(FieldTy); + size_t BufferLocAttrArg = + handleBufferLocationProperty(FieldTy, BS.getBeginLoc()); addParam(BS, (*ParamIt)->getType().getCanonicalType(), BufferLocAttrArg); ++ParamIt; for (; ParamIt != InitMethod->parameters().end(); ++ParamIt) @@ -2878,7 +2917,11 @@ bool Util::isSyclSpecConstantType(const QualType &Ty) { return matchQualifiedTypeName(Ty, Scopes); } -bool Util::isSyclBufferLocation(const QualType &Ty) { +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 Scopes = { Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"}, diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 0863d3d4a0b7..479679020789 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -86,7 +86,8 @@ struct property_base { template class property_list { public: - property_list(properties... props) {} + template + property_list(propertiesTN... props) {}; template bool has_property() const { return true; } diff --git a/clang/test/CodeGenSYCL/buffer_location.cpp b/clang/test/CodeGenSYCL/buffer_location.cpp index c779c41ecb78..babd6fb948bd 100644 --- a/clang/test/CodeGenSYCL/buffer_location.cpp +++ b/clang/test/CodeGenSYCL/buffer_location.cpp @@ -9,7 +9,8 @@ int main() { cl::sycl::accessor> + cl::sycl::property_list< + cl::sycl::property::buffer_location<3>>> accessorA; cl::sycl::kernel_single_task( [=]() { diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 9e3efc632109..4163eb15d616 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -37,6 +37,14 @@ enum class address_space : int { }; } // namespace access +namespace property { + template + class buffer_location {}; +} // namespace property + +template +class property_list {}; + namespace detail { namespace half_impl { struct half { @@ -86,7 +94,8 @@ struct DeviceValueType { template + access::placeholder isPlaceholder = access::placeholder::false_t, + typename propertyListT = property_list<>> class accessor { public: @@ -98,6 +107,7 @@ class accessor { using PtrType = typename DeviceValueType::type *; void __init(PtrType Ptr, range AccessRange, range MemRange, id Offset) {} + propertyListT prop_list; }; template diff --git a/clang/test/SemaSYCL/accessor_inheritance.cpp b/clang/test/SemaSYCL/accessor_inheritance.cpp index 17dafe7b4acd..db07ee25a545 100644 --- a/clang/test/SemaSYCL/accessor_inheritance.cpp +++ b/clang/test/SemaSYCL/accessor_inheritance.cpp @@ -42,8 +42,8 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_A' 'int' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_B' 'int' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor' 'void () noexcept' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor>' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor>' 'void () noexcept' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_C' 'int' diff --git a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp index 1f500eff0a88..54ae68ae5b8c 100644 --- a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp @@ -43,7 +43,7 @@ int main() { // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ({{.*}}PtrType, range<1>, range<1>, id<1>)' lvalue .__init -// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor' lvalue . +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor>' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})' lvalue Var // CHECK-NEXT: ImplicitCastExpr {{.*}} diff --git a/clang/test/SemaSYCL/buffer_location.cpp b/clang/test/SemaSYCL/buffer_location.cpp new file mode 100644 index 000000000000..b02653366287 --- /dev/null +++ b/clang/test/SemaSYCL/buffer_location.cpp @@ -0,0 +1,61 @@ +// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -verify -pedantic -DTRIGGER_ERROR %s + +#include "sycl.hpp" + +class another_property {}; + +template +class another_property_list { +}; + +int main() { +#ifndef TRIGGER_ERROR + cl::sycl::accessor>> +// CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 2 + accessorA; + cl::sycl::accessor>> +// CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 3 + accessorB; + cl::sycl::accessor> + accessorC; +#else + cl::sycl::accessor>> + accessorD; + cl::sycl::accessor> + accessorE; +#endif + cl::sycl::kernel_single_task( + [=]() { +#ifndef TRIGGER_ERROR + // expected-no-diagnostics + accessorA.use(); + accessorB.use(); + accessorC.use(); +#else + //expected-error@+1{{buffer_location template parameter must be a compiletime known non-negative integer}} + accessorD.use(); + //expected-error@+1{{accessor's 5th template parameter must be a property_list}} + accessorE.use(); +#endif + }); + return 0; +} diff --git a/clang/test/SemaSYCL/intel-kernel-arg-buffer-location.cpp b/clang/test/SemaSYCL/intel-kernel-arg-buffer-location.cpp deleted file mode 100644 index 3e555ecf7e3a..000000000000 --- a/clang/test/SemaSYCL/intel-kernel-arg-buffer-location.cpp +++ /dev/null @@ -1,35 +0,0 @@ -// RUN: %clang %s -fsyntax-only -fsycl-device-only -DCHECKDIAG -Xclang -verify -// RUN: %clang %s -fsyntax-only -I %S/Inputs -Xclang -ast-dump -fsycl-device-only | FileCheck %s - -#ifndef CHECKDIAG -#include "sycl.hpp" -#endif // CHECKDIAG - -#ifdef CHECKDIAG -struct FuncObj { - [[intelfpga::kernel_arg_buffer_location]] // expected-warning{{'kernel_arg_buffer_location' attribute cannot be used explicitly}} - void - operator()() {} -}; - -template -__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { - kernelFunc(); - [[intelfpga::kernel_arg_buffer_location]] int invalid = 42; // expected-error{{'kernel_arg_buffer_location' attribute only applies to functions}} -} -#endif // CHECKDIAG - -int main() { -#ifdef CHECKDIAG - kernel( - FuncObj()); -#else - // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ4mainE15kernel_function - // CHECK: SYCLIntelBufferLocationAttr - cl::sycl::accessor accessorA; - cl::sycl::kernel_single_task( - [=]() { - accessorA.use(); - }); -#endif // CHECKDIAG -} diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp index 1052b4ac24e0..9f9a3196024c 100644 --- a/clang/test/SemaSYCL/wrapped-accessor.cpp +++ b/clang/test/SemaSYCL/wrapped-accessor.cpp @@ -35,14 +35,14 @@ int main() { // argument // CHECK: VarDecl {{.*}}'(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' // CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' -// CHECK-NEXT: InitListExpr {{.*}}'AccWrapper>' -// CHECK-NEXT: CXXConstructExpr {{.*}}'cl::sycl::accessor':'cl::sycl::accessor' 'void () noexcept' +// CHECK-NEXT: InitListExpr {{.*}}'AccWrapper>>' +// CHECK-NEXT: CXXConstructExpr {{.*}}'cl::sycl::accessor>':'cl::sycl::accessor>' 'void () noexcept' // Check that accessor field of the wrapper object is initialized using __init method // CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ({{.*}}PtrType, range<1>, range<1>, id<1>)' lvalue .__init -// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor' lvalue .accessor {{.*}} -// CHECK-NEXT: MemberExpr {{.*}} 'AccWrapper>':'AccWrapper>' lvalue . +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor>':'cl::sycl::accessor>' lvalue .accessor {{.*}} +// CHECK-NEXT: MemberExpr {{.*}} 'AccWrapper>>':'AccWrapper>>' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' // Parameters of the _init method From f57d97d688bfcc9ec749301b6d4d19cd1491338f Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Tue, 28 Jul 2020 18:52:08 +0300 Subject: [PATCH 05/12] Apply clang format Signed-off-by: Dmitry Sidorov --- clang/lib/CodeGen/CodeGenModule.cpp | 6 +++--- clang/lib/Sema/SemaSYCL.cpp | 12 ++++++------ clang/test/CodeGenSYCL/Inputs/sycl.hpp | 4 ++-- clang/test/SemaSYCL/Inputs/sycl.hpp | 4 ++-- clang/test/SemaSYCL/buffer_location.cpp | 6 +++--- 5 files changed, 16 insertions(+), 16 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 0f1db8af3d81..38326d113dbe 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1414,7 +1414,7 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, SmallVector argNames; // MDNode for the intel_buffer_location attribute. - SmallVector argSYCLBufferLocationAttr; + SmallVector argSYCLBufferLocationAttr; if (FD && CGF) for (unsigned i = 0, e = FD->getNumParams(); i != e; ++i) { @@ -1543,9 +1543,9 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, auto *SYCLBufferLocationAttr = parm->getAttr(); argSYCLBufferLocationAttr.push_back( - (SYCLBufferLocationAttr) + (SYCLBufferLocationAttr) ? llvm::ConstantAsMetadata::get(CGF->Builder.getInt32( - SYCLBufferLocationAttr->getLocationID())) + SYCLBufferLocationAttr->getLocationID())) : llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(-1))); } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d547f5f9c702..633061301c44 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1224,13 +1224,13 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { const auto PropList = cast(AccTy->getTemplateArgs()[5]); if (PropList.getKind() != TemplateArgument::ArgKind::Type) { SemaRef.Diag(Loc, diag::err_sycl_invalid_property_template_param) - << "accessor's 5th" << /*type*/ 1; + << "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; + << "accessor's 5th" << /*property_list*/ 3; return -1; } @@ -1240,7 +1240,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { 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; + << "property_list" << /*parameter pack*/ 0; return -1; } // Move through TemplateArgs list of a property list and search for @@ -1254,7 +1254,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // accessor - emit an error if (LocationID != -1) { SemaRef.Diag(Loc, diag::err_sycl_compiletime_property_duplication) - << "buffer_location"; + << "buffer_location"; return -1; } const auto *PropDecl = @@ -1262,13 +1262,13 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { 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; + << "buffer_location" << /*non-negative integer*/ 2; return -1; } LocationID = static_cast(BufferLoc.getAsIntegral().getExtValue()); if (LocationID < 0) { SemaRef.Diag(Loc, diag::err_sycl_invalid_property_template_param) - << "buffer_location" << /*non-negative integer*/ 2; + << "buffer_location" << /*non-negative integer*/ 2; return -1; } } diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 479679020789..de3e8999961b 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -86,8 +86,8 @@ struct property_base { template class property_list { public: - template - property_list(propertiesTN... props) {}; + template + property_list(propertiesTN... props){}; template bool has_property() const { return true; } diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 4163eb15d616..1b3ab8b20763 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -38,8 +38,8 @@ enum class address_space : int { } // namespace access namespace property { - template - class buffer_location {}; +template +class buffer_location {}; } // namespace property template diff --git a/clang/test/SemaSYCL/buffer_location.cpp b/clang/test/SemaSYCL/buffer_location.cpp index b02653366287..2e9e3a95d714 100644 --- a/clang/test/SemaSYCL/buffer_location.cpp +++ b/clang/test/SemaSYCL/buffer_location.cpp @@ -16,14 +16,14 @@ int main() { cl::sycl::access::placeholder::false_t, cl::sycl::property_list< cl::sycl::property::buffer_location<2>>> -// CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 2 + // CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 2 accessorA; cl::sycl::accessor>> -// CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 3 + cl::sycl::property::buffer_location<3>>> + // CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 3 accessorB; cl::sycl::accessor Date: Tue, 28 Jul 2020 19:16:47 +0300 Subject: [PATCH 06/12] Add missing test Signed-off-by: Dmitry Sidorov --- clang/test/SemaSYCL/buffer_location.cpp | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/clang/test/SemaSYCL/buffer_location.cpp b/clang/test/SemaSYCL/buffer_location.cpp index 2e9e3a95d714..3ef0e4dc82af 100644 --- a/clang/test/SemaSYCL/buffer_location.cpp +++ b/clang/test/SemaSYCL/buffer_location.cpp @@ -42,6 +42,13 @@ int main() { cl::sycl::access::placeholder::false_t, another_property_list> accessorE; + cl::sycl::accessor, + cl::sycl::property::buffer_location<2>>> + accessorF; #endif cl::sycl::kernel_single_task( [=]() { @@ -55,6 +62,8 @@ int main() { accessorD.use(); //expected-error@+1{{accessor's 5th template parameter must be a property_list}} accessorE.use(); + //expected-error@+1{{Can't apply buffer_location property twice to the same accessor}} + accessorF.use(); #endif }); return 0; From c800cd65fd897aab0e77d4a1ad6d0f32a74b7511 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Thu, 30 Jul 2020 20:42:59 +0300 Subject: [PATCH 07/12] Apply comments and do a little refactoring Signed-off-by: Dmitry Sidorov --- clang/include/clang/Basic/Attr.td | 3 - clang/include/clang/Basic/AttrDocs.td | 21 ++-- .../clang/Basic/DiagnosticSemaKinds.td | 3 +- clang/lib/Sema/SemaSYCL.cpp | 119 +++++++++--------- clang/test/CodeGenSYCL/buffer_location.cpp | 23 +++- clang/test/SemaSYCL/buffer_location.cpp | 24 +++- 6 files changed, 114 insertions(+), 79 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 58c2b91e6b9f..e6005140a422 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1193,10 +1193,7 @@ def SYCLIntelBufferLocation : InheritableAttr { 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 { diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 6884ab5db60b..2ff3801062e9 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -1998,17 +1998,16 @@ 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'. +This attribute is implicitly added to OpenCL pointer kernel parameters generated +from a SYCL kernel object. It lacks a spelling, as it is not intended to be used +by the programmer. + +This attribute causes clang to generate metadata on the OpenCL kernel containing +the number of kernel parameters. The metadata contains an integer that is set +according to the values passed through the ``accessor`` property +``buffer_location``. These values are mapped to the actual locations of the +global buffers (such as DDR, QDR, etc) and applied to pointer kernel parameters. +The default value is -1. }]; } diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 68c55a355b02..775d84d13dec 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10973,8 +10973,7 @@ 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">; + "%select{parameter pack|type|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 " diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 633061301c44..89fa7618c81c 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1180,23 +1180,22 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // iterator as push_back invalidates iterators. size_t LastParamIndex = 0; - void addParam(const FieldDecl *FD, QualType FieldTy, int LocationID = -1) { + void addParam(const FieldDecl *FD, QualType FieldTy) { const ConstantArrayType *CAT = SemaRef.getASTContext().getAsConstantArrayType(FieldTy); if (CAT) FieldTy = CAT->getElementType(); ParamDesc newParamDesc = makeParamDesc(FD, FieldTy); - addParam(newParamDesc, FieldTy, LocationID); + addParam(newParamDesc, FieldTy); } - void addParam(const CXXBaseSpecifier &BS, QualType FieldTy, - int LocationID = -1) { + void addParam(const CXXBaseSpecifier &BS, QualType FieldTy) { ParamDesc newParamDesc = makeParamDesc(SemaRef.getASTContext(), BS, FieldTy); - addParam(newParamDesc, FieldTy, LocationID); + addParam(newParamDesc, FieldTy); } - void addParam(ParamDesc newParamDesc, QualType FieldTy, int LocationID) { + void addParam(ParamDesc newParamDesc, QualType FieldTy) { // Create a new ParmVarDecl based on the new info. ASTContext &Ctx = SemaRef.getASTContext(); auto *NewParam = ParmVarDecl::Create( @@ -1208,72 +1207,76 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { 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(FieldTy->getAsRecordDecl()); + // Handle accessor properties. If any properties were found in + // the property_list - add the appropriate attributes to ParmVarDecl. + void handleAccessorPropertyList(ParmVarDecl *Param, + const CXXRecordDecl *RecordDecl, + SourceLocation Loc) { + const auto *AccTy = cast(RecordDecl); // TODO: when SYCL headers' part is ready - replace this 'if' with an error if (AccTy->getTemplateArgs().size() < 6) - return -1; + return; const auto PropList = cast(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; + return; } 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; + return; } - int LocationID = -1; const auto *PropListDecl = cast(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; + return; } // 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. + // properties. If found - apply the appropriate attribute to ParmVarDecl. 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(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(BufferLoc.getAsIntegral().getExtValue()); - if (LocationID < 0) { - SemaRef.Diag(Loc, diag::err_sycl_invalid_property_template_param) - << "buffer_location" << /*non-negative integer*/ 2; - return -1; - } - } + if (Util::isSyclBufferLocationType(PropTy)) + handleBufferLocationProperty(Param, PropTy, Loc); + } + } + + // Obtain an integer value stored in a template parameter of buffer_location + // property to pass it to buffer_location kernel attribute + void handleBufferLocationProperty(ParmVarDecl *Param, QualType PropTy, + SourceLocation Loc) { + // If we have more than 1 buffer_location properties on a single + // accessor - emit an error + if (Param->hasAttr()) { + SemaRef.Diag(Loc, diag::err_sycl_compiletime_property_duplication) + << "buffer_location"; + return; + } + ASTContext &Ctx = SemaRef.getASTContext(); + const auto *PropDecl = + cast(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; + } + int LocationID = static_cast(BufferLoc.getAsIntegral().getExtValue()); + if (LocationID < 0) { + SemaRef.Diag(Loc, diag::err_sycl_invalid_property_template_param) + << "buffer_location" << /*non-negative integer*/ 2; + return; } - return LocationID; + Param->addAttr( + SYCLIntelBufferLocationAttr::CreateImplicit(Ctx, LocationID)); } // All special SYCL objects must have __init method. We extract types for @@ -1290,18 +1293,15 @@ 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(); - auto ParamIt = InitMethod->parameters().begin(); + ParmVarDecl **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); + addParam(FD, (*ParamIt)->getType().getCanonicalType()); + if (isAccessorType) + handleAccessorPropertyList(Params.back(), RecordDecl, + FD->getLocation()); ++ParamIt; for (; ParamIt != InitMethod->parameters().end(); ++ParamIt) - addParam(FD, (*ParamIt)->getType().getCanonicalType(), -1); + addParam(FD, (*ParamIt)->getType().getCanonicalType()); } LastParamIndex = ParamIndex; return true; @@ -1372,16 +1372,13 @@ 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(); - auto ParamIt = InitMethod->parameters().begin(); + ParmVarDecl **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); + addParam(BS, (*ParamIt)->getType().getCanonicalType()); + handleAccessorPropertyList(Params.back(), RecordDecl, BS.getBeginLoc()); ++ParamIt; for (; ParamIt != InitMethod->parameters().end(); ++ParamIt) - addParam(BS, (*ParamIt)->getType().getCanonicalType(), -1); + addParam(BS, (*ParamIt)->getType().getCanonicalType()); } LastParamIndex = ParamIndex; return true; diff --git a/clang/test/CodeGenSYCL/buffer_location.cpp b/clang/test/CodeGenSYCL/buffer_location.cpp index babd6fb948bd..0e661a34d738 100644 --- a/clang/test/CodeGenSYCL/buffer_location.cpp +++ b/clang/test/CodeGenSYCL/buffer_location.cpp @@ -1,11 +1,31 @@ // 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} +// CHECK: ![[MDBL]] = !{i32 3, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 2, i32 -1, i32 -1, i32 -1, i32 2, i32 -1, i32 -1, i32 -1, i32 -1} #include "sycl.hpp" +struct Base { + int A, B; + cl::sycl::accessor>> AccField; +}; + +struct Captured : Base, + cl::sycl::accessor> + > { + int C; +}; + int main() { + Captured Obj; cl::sycl::accessor( [=]() { accessorA.use(); + Obj.use(); }); return 0; } diff --git a/clang/test/SemaSYCL/buffer_location.cpp b/clang/test/SemaSYCL/buffer_location.cpp index 3ef0e4dc82af..b5c7aeb79385 100644 --- a/clang/test/SemaSYCL/buffer_location.cpp +++ b/clang/test/SemaSYCL/buffer_location.cpp @@ -9,8 +9,29 @@ template class another_property_list { }; +struct Base { + int A, B; + cl::sycl::accessor>> AccField; +}; + +struct Captured : Base, + cl::sycl::accessor> + > { + int C; +}; + int main() { #ifndef TRIGGER_ERROR + // CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 1 + Captured Obj; cl::sycl::accessor Date: Thu, 30 Jul 2020 21:25:33 +0300 Subject: [PATCH 08/12] Resolve clang-format Signed-off-by: Dmitry Sidorov --- clang/test/CodeGenSYCL/buffer_location.cpp | 12 ++++++------ clang/test/SemaSYCL/buffer_location.cpp | 10 +++++----- 2 files changed, 11 insertions(+), 11 deletions(-) diff --git a/clang/test/CodeGenSYCL/buffer_location.cpp b/clang/test/CodeGenSYCL/buffer_location.cpp index 0e661a34d738..ae1b9088cc27 100644 --- a/clang/test/CodeGenSYCL/buffer_location.cpp +++ b/clang/test/CodeGenSYCL/buffer_location.cpp @@ -8,10 +8,11 @@ struct Base { int A, B; cl::sycl::accessor>> AccField; + cl::sycl::access::target::global_buffer, + cl::sycl::access::placeholder::false_t, + cl::sycl::property_list< + cl::sycl::property::buffer_location<2>>> + AccField; }; struct Captured : Base, @@ -19,8 +20,7 @@ struct Captured : Base, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, cl::sycl::property_list< - cl::sycl::property::buffer_location<2>> - > { + cl::sycl::property::buffer_location<2>>> { int C; }; diff --git a/clang/test/SemaSYCL/buffer_location.cpp b/clang/test/SemaSYCL/buffer_location.cpp index b5c7aeb79385..37ae3ed35031 100644 --- a/clang/test/SemaSYCL/buffer_location.cpp +++ b/clang/test/SemaSYCL/buffer_location.cpp @@ -15,7 +15,8 @@ struct Base { cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, cl::sycl::property_list< - cl::sycl::property::buffer_location<1>>> AccField; + cl::sycl::property::buffer_location<1>>> + AccField; }; struct Captured : Base, @@ -23,8 +24,7 @@ struct Captured : Base, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, cl::sycl::property_list< - cl::sycl::property::buffer_location<1>> - > { + cl::sycl::property::buffer_location<1>>> { int C; }; @@ -67,8 +67,8 @@ int main() { cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, cl::sycl::property_list< - cl::sycl::property::buffer_location<1>, - cl::sycl::property::buffer_location<2>>> + cl::sycl::property::buffer_location<1>, + cl::sycl::property::buffer_location<2>>> accessorF; #endif cl::sycl::kernel_single_task( From 5c300937b88ac0cb12ed99fdb2db6f76f333495a Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Fri, 31 Jul 2020 18:56:23 +0300 Subject: [PATCH 09/12] Move diagnostics + don't rely on params order Signed-off-by: Dmitry Sidorov --- clang/include/clang/Basic/AttrDocs.td | 4 +- .../clang/Basic/DiagnosticSemaKinds.td | 6 +- clang/lib/Sema/SemaSYCL.cpp | 96 +++++++++++-------- clang/test/SemaSYCL/buffer_location.cpp | 2 +- 4 files changed, 64 insertions(+), 44 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 2ff3801062e9..56130e693434 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2007,7 +2007,9 @@ the number of kernel parameters. The metadata contains an integer that is set according to the values passed through the ``accessor`` property ``buffer_location``. These values are mapped to the actual locations of the global buffers (such as DDR, QDR, etc) and applied to pointer kernel parameters. -The default value is -1. +Number of metadata arguments is the same as a number of kernel parameters, so +any parameter that isn't an ``accessor`` with ``buffer_location`` property is +annotated by '-1' in the metadata node. }]; } diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 8b482154bfbc..b5ea32068277 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10976,9 +10976,11 @@ def err_sycl_attibute_cannot_be_applied_here "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< +def err_sycl_invalid_property_list_template_param : Error< "%0 template parameter must be a " - "%select{parameter pack|type|non-negative integer|property_list}1">; + "%select{parameter pack|type|non-negative integer}1">; +def err_sycl_invalid_accessor_property_template_param : Error< + "Fifth template parameter of the accessor must be of a property_list type">; 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 " diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 533494cc865b..54818cf448df 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1084,6 +1084,51 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { return false; } + void checkPropertyListType(TemplateArgument PropList, SourceLocation Loc) { + if (PropList.getKind() != TemplateArgument::ArgKind::Type) { + SemaRef.Diag(Loc, + diag::err_sycl_invalid_accessor_property_template_param); + return; + } + QualType PropListTy = PropList.getAsType(); + if (!Util::isPropertyListType(PropListTy)) { + SemaRef.Diag(Loc, + diag::err_sycl_invalid_accessor_property_template_param); + return; + } + const auto *PropListDecl = + cast(PropListTy->getAsRecordDecl()); + const auto TemplArg = PropListDecl->getTemplateArgs()[0]; + if (TemplArg.getKind() != TemplateArgument::ArgKind::Pack) { + SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param) + << "property_list" << /*parameter pack*/ 0; + return; + } + for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin(); + Prop != TemplArg.pack_end(); ++Prop) { + QualType PropTy = Prop->getAsType(); + if (Util::isSyclBufferLocationType(PropTy)) + checkBufferLocationType(PropTy, Loc); + } + } + + void checkBufferLocationType(QualType PropTy, SourceLocation Loc) { + const auto *PropDecl = + cast(PropTy->getAsRecordDecl()); + const auto BufferLoc = PropDecl->getTemplateArgs()[0]; + if (BufferLoc.getKind() != TemplateArgument::ArgKind::Integral) { + SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param) + << "buffer_location" << /*non-negative integer*/ 2; + return; + } + int LocationID = static_cast(BufferLoc.getAsIntegral().getExtValue()); + if (LocationID < 0) { + SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param) + << "buffer_location" << /*non-negative integer*/ 2; + return; + } + } + void checkAccessorType(QualType Ty, SourceRange Loc) { assert(Util::isSyclAccessorType(Ty) && "Should only be called on SYCL accessor types."); @@ -1095,6 +1140,8 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { TemplateArgument TA = TAL.get(0); const QualType TemplateArgTy = TA.getAsType(); + if (TAL.size() > 5) + checkPropertyListType(TAL.get(5), Loc.getBegin()); llvm::DenseSet Visited; checkSYCLType(SemaRef, TemplateArgTy, Loc, Visited); } @@ -1188,26 +1235,10 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { if (AccTy->getTemplateArgs().size() < 6) return; const auto PropList = cast(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; - } 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; - } - const auto *PropListDecl = cast(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; - } // 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(); @@ -1233,17 +1264,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { const auto *PropDecl = cast(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; - } int LocationID = static_cast(BufferLoc.getAsIntegral().getExtValue()); - if (LocationID < 0) { - SemaRef.Diag(Loc, diag::err_sycl_invalid_property_template_param) - << "buffer_location" << /*non-negative integer*/ 2; - return; - } Param->addAttr( SYCLIntelBufferLocationAttr::CreateImplicit(Ctx, LocationID)); } @@ -1262,15 +1283,12 @@ 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(); - ParmVarDecl **ParamIt = InitMethod->parameters().begin(); - if (*ParamIt) { - addParam(FD, (*ParamIt)->getType().getCanonicalType()); - if (isAccessorType) + for (const ParmVarDecl *Param : InitMethod->parameters()) { + QualType ParamTy = Param->getType(); + addParam(FD, ParamTy.getCanonicalType()); + if (ParamTy.getTypePtr()->isPointerType() && isAccessorType) handleAccessorPropertyList(Params.back(), RecordDecl, FD->getLocation()); - ++ParamIt; - for (; ParamIt != InitMethod->parameters().end(); ++ParamIt) - addParam(FD, (*ParamIt)->getType().getCanonicalType()); } LastParamIndex = ParamIndex; return true; @@ -1341,13 +1359,11 @@ 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(); - ParmVarDecl **ParamIt = InitMethod->parameters().begin(); - if (*ParamIt) { - addParam(BS, (*ParamIt)->getType().getCanonicalType()); - handleAccessorPropertyList(Params.back(), RecordDecl, BS.getBeginLoc()); - ++ParamIt; - for (; ParamIt != InitMethod->parameters().end(); ++ParamIt) - addParam(BS, (*ParamIt)->getType().getCanonicalType()); + for (const ParmVarDecl *Param : InitMethod->parameters()) { + QualType ParamTy = Param->getType(); + addParam(BS, ParamTy.getCanonicalType()); + if (ParamTy.getTypePtr()->isPointerType()) + handleAccessorPropertyList(Params.back(), RecordDecl, BS.getBeginLoc()); } LastParamIndex = ParamIndex; return true; diff --git a/clang/test/SemaSYCL/buffer_location.cpp b/clang/test/SemaSYCL/buffer_location.cpp index 37ae3ed35031..bd2daaa01ba4 100644 --- a/clang/test/SemaSYCL/buffer_location.cpp +++ b/clang/test/SemaSYCL/buffer_location.cpp @@ -82,7 +82,7 @@ int main() { #else //expected-error@+1{{buffer_location template parameter must be a non-negative integer}} accessorD.use(); - //expected-error@+1{{accessor's 5th template parameter must be a property_list}} + //expected-error@+1{{Fifth template parameter of the accessor must be of a property_list type}} accessorE.use(); //expected-error@+1{{Can't apply buffer_location property twice to the same accessor}} accessorF.use(); From ab5b6da096597e81415fd999ce5508c18260c428 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Sat, 1 Aug 2020 00:32:58 +0300 Subject: [PATCH 10/12] Revert "Add support for SPV_INTEL_fpga_buffer_location extension" This reverts commit f71cec58d14be2360a8be60f9fcfec1a4a10ea06. --- llvm-spirv/include/LLVMSPIRVExtensions.inc | 1 - llvm-spirv/lib/SPIRV/SPIRVReader.cpp | 30 --------- llvm-spirv/lib/SPIRV/SPIRVWriter.cpp | 16 ----- llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h | 4 -- llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h | 2 - .../lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h | 2 - llvm-spirv/lib/SPIRV/libSPIRV/spirv.hpp | 2 - .../test/transcoding/FPGABufferLocation.ll | 61 ------------------- 8 files changed, 118 deletions(-) delete mode 100644 llvm-spirv/test/transcoding/FPGABufferLocation.ll diff --git a/llvm-spirv/include/LLVMSPIRVExtensions.inc b/llvm-spirv/include/LLVMSPIRVExtensions.inc index adaa48efadc2..4a037aeddab2 100644 --- a/llvm-spirv/include/LLVMSPIRVExtensions.inc +++ b/llvm-spirv/include/LLVMSPIRVExtensions.inc @@ -23,4 +23,3 @@ EXT(SPV_INTEL_optimization_hints) EXT(SPV_INTEL_float_controls2) EXT(SPV_INTEL_vector_compute) EXT(SPV_INTEL_usm_storage_classes) -EXT(SPV_INTEL_fpga_buffer_location) diff --git a/llvm-spirv/lib/SPIRV/SPIRVReader.cpp b/llvm-spirv/lib/SPIRV/SPIRVReader.cpp index bc7386e26d76..d7b25c350a8d 100644 --- a/llvm-spirv/lib/SPIRV/SPIRVReader.cpp +++ b/llvm-spirv/lib/SPIRV/SPIRVReader.cpp @@ -168,27 +168,6 @@ static void addOCLKernelArgumentMetadata( Fn->setMetadata(MDName, MDNode::get(*Context, ValueVec)); } -static void addBufferLocationMetadata( - LLVMContext *Context, SPIRVFunction *BF, llvm::Function *Fn, - std::function Func) { - std::vector ValueVec; - bool DecorationFound = false; - BF->foreachArgument([&](SPIRVFunctionParameter *Arg) { - if (Arg->getType()->isTypePointer() && - Arg->hasDecorate(DecorationBufferLocationINTEL)) { - DecorationFound = true; - ValueVec.push_back(Func(Arg)); - } else { - llvm::Metadata *DefaultNode = ConstantAsMetadata::get( - ConstantInt::get(Type::getInt32Ty(*Context), -1)); - ValueVec.push_back(DefaultNode); - } - }); - if (DecorationFound) - Fn->setMetadata("kernel_arg_buffer_location", - MDNode::get(*Context, ValueVec)); -} - Value *SPIRVToLLVM::getTranslatedValue(SPIRVValue *BV) { auto Loc = ValueMap.find(BV); if (Loc != ValueMap.end()) @@ -3540,15 +3519,6 @@ bool SPIRVToLLVM::transOCLMetadata(SPIRVFunction *BF) { Arg->getName()); }); } - // Generate metadata for kernel_arg_buffer_location - addBufferLocationMetadata(Context, BF, F, [=](SPIRVFunctionParameter *Arg) { - auto Literals = Arg->getDecorationLiterals(DecorationBufferLocationINTEL); - assert(Literals.size() == 1 && - "BufferLocationINTEL decoration shall have 1 ID literal"); - - return ConstantAsMetadata::get( - ConstantInt::get(Type::getInt32Ty(*Context), Literals[0])); - }); return true; } diff --git a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp index 25988f58ea76..f187a47bf6cd 100644 --- a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp +++ b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp @@ -527,13 +527,6 @@ SPIRVFunction *LLVMToSPIRV::transFunctionDecl(Function *F) { BM->addEntryPoint(ExecutionModelKernel, BF->getId()); else if (F->getLinkage() != GlobalValue::InternalLinkage) BF->setLinkageType(transLinkageType(F)); - - // Translate OpenCL/SYCL buffer_location metadata if it's attached to the - // translated function declaration - MDNode *BufferLocation = nullptr; - if (BM->isAllowedToUseExtension(ExtensionID::SPV_INTEL_fpga_buffer_location)) - BufferLocation = ((*F).getMetadata("kernel_arg_buffer_location")); - auto Attrs = F->getAttributes(); for (Function::arg_iterator I = F->arg_begin(), E = F->arg_end(); I != E; @@ -559,15 +552,6 @@ SPIRVFunction *LLVMToSPIRV::transFunctionDecl(Function *F) { BA->addDecorate(DecorationMaxByteOffset, Attrs.getAttribute(ArgNo + 1, Attribute::Dereferenceable) .getDereferenceableBytes()); - if (BufferLocation && I->getType()->isPointerTy()) { - // Order of integer numbers in MD node follows the order of function - // parameters on which we shall attach the appropriate decoration. Add - // decoration only if MD value is not negative. - BM->addCapability(CapabilityFPGABufferLocationINTEL); - int LocID = getMDOperandAsInt(BufferLocation, ArgNo); - if (LocID >= 0) - BA->addDecorate(DecorationBufferLocationINTEL, LocID); - } } if (Attrs.hasAttribute(AttributeList::ReturnIndex, Attribute::ZExt)) BF->addDecorate(DecorationFuncParamAttr, FunctionParameterAttributeZext); diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h index ba03b181cf3c..31c60f0f2e24 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h @@ -164,8 +164,6 @@ class SPIRVDecorate : public SPIRVDecorateGeneric { return getSet(ExtensionID::SPV_INTEL_function_pointers); case DecorationIOPipeStorageINTEL: return getSet(ExtensionID::SPV_INTEL_io_pipes); - case DecorationBufferLocationINTEL: - return getSet(ExtensionID::SPV_INTEL_fpga_buffer_location); default: return SPIRVExtSet(); } @@ -269,8 +267,6 @@ class SPIRVMemberDecorate : public SPIRVDecorateGeneric { return getSet(ExtensionID::SPV_INTEL_fpga_memory_accesses); case DecorationIOPipeStorageINTEL: return getSet(ExtensionID::SPV_INTEL_io_pipes); - case DecorationBufferLocationINTEL: - return getSet(ExtensionID::SPV_INTEL_fpga_buffer_location); default: return SPIRVExtSet(); } diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h index c35b4f01f173..0ced7ce27f40 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h @@ -402,8 +402,6 @@ template <> inline void SPIRVMap::init() { ADD_VEC_INIT(DecorationDontStaticallyCoalesceINTEL, {CapabilityFPGAMemoryAccessesINTEL}); ADD_VEC_INIT(DecorationPrefetchINTEL, {CapabilityFPGAMemoryAccessesINTEL}); - ADD_VEC_INIT(DecorationBufferLocationINTEL, - {CapabilityFPGABufferLocationINTEL}); } template <> inline void SPIRVMap::init() { diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h index c9915382e4fc..e096040ed072 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h @@ -376,7 +376,6 @@ template <> inline void SPIRVMap::init() { add(DecorationGlobalVariableOffsetINTEL, "GlobalVariableOffsetINTEL"); add(DecorationFuncParamIOKind, "FuncParamIOKind"); add(DecorationSIMTCallINTEL, "SIMTCallINTEL"); - add(DecorationBufferLocationINTEL, "BufferLocationINTEL"); } SPIRV_DEF_NAMEMAP(Decoration, SPIRVDecorationNameMap) @@ -561,7 +560,6 @@ template <> inline void SPIRVMap::init() { "GroupNonUniformShuffleRelative"); add(CapabilityGroupNonUniformClustered, "GroupNonUniformClustered"); add(CapabilityUSMStorageClassesINTEL, "USMStorageClassesINTEL"); - add(CapabilityFPGABufferLocationINTEL, "FPGABufferLocationINTEL"); } SPIRV_DEF_NAMEMAP(Capability, SPIRVCapabilityNameMap) diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/spirv.hpp b/llvm-spirv/lib/SPIRV/libSPIRV/spirv.hpp index 53492791c7e7..22ccb3f8b17c 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/spirv.hpp +++ b/llvm-spirv/lib/SPIRV/libSPIRV/spirv.hpp @@ -508,7 +508,6 @@ enum Decoration { DecorationCacheSizeINTEL = 5900, DecorationDontStaticallyCoalesceINTEL = 5901, DecorationPrefetchINTEL = 5902, - DecorationBufferLocationINTEL = 5921, DecorationIOPipeStorageINTEL = 5944, DecorationMax = 0x7fffffff, }; @@ -963,7 +962,6 @@ enum Capability { CapabilityFPGARegINTEL = 5948, CapabilityKernelAttributesINTEL = 5892, CapabilityFPGAKernelAttributesINTEL = 5897, - CapabilityFPGABufferLocationINTEL = 5920, CapabilityUSMStorageClassesINTEL = 5935, CapabilityFPGAMemoryAccessesINTEL = 5898, CapabilityIOPipeINTEL = 5943, diff --git a/llvm-spirv/test/transcoding/FPGABufferLocation.ll b/llvm-spirv/test/transcoding/FPGABufferLocation.ll deleted file mode 100644 index e649a89625d5..000000000000 --- a/llvm-spirv/test/transcoding/FPGABufferLocation.ll +++ /dev/null @@ -1,61 +0,0 @@ -; RUN: llvm-as %s -o %t.bc -; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_fpga_buffer_location -o %t.spv -; RUN: llvm-spirv %t.spv -to-text -o %t.spt -; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV - -; RUN: llvm-spirv -r %t.spv -o %t.rev.bc -; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM - -; RUN: llvm-spirv -spirv-text -r %t.spt -o %t.rev.bc -; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM - -; CHECK-SPIRV: 2 Capability FPGABufferLocationINTEL -; CHECK-SPIRV: 9 Extension "SPV_INTEL_fpga_buffer_location" -; CHECK-SPIRV: 3 Name [[ARGA:[0-9]+]] "a" -; CHECK-SPIRV: 3 Name [[ARGB:[0-9]+]] "b" -; CHECK-SPIRV: 3 Name [[ARGC:[0-9]+]] "c" -; CHECK-SPIRV: 3 Name [[ARGD:[0-9]+]] "d" -; CHECK-SPIRV: 3 Name [[ARGE:[0-9]+]] "e" -; CHECK-SPIRV-NOT: 4 Decorate [[ARGC]] BufferLocationINTEL -1 -; CHECK-SPIRV-NOT: 4 Decorate [[ARGC]] BufferLocationINTEL -1 -; CHECK-SPIRV: 4 Decorate [[ARGA]] BufferLocationINTEL 1 -; CHECK-SPIRV: 4 Decorate [[ARGB]] BufferLocationINTEL 2 -; CHECK-SPIRV-NOT: 4 Decorate [[ARGD]] BufferLocationINTEL -1 -; CHECK-SPIRV-NOT: 4 Decorate [[ARGE]] BufferLocationINTEL 3 - -; CHECK-SPIRV: 5 Function -; CHECK-SPIRV: 3 FunctionParameter {{[0-9]+}} [[ARGA]] -; CHECK-SPIRV: 3 FunctionParameter {{[0-9]+}} [[ARGB]] -; CHECK-SPIRV: 3 FunctionParameter {{[0-9]+}} [[ARGC]] -; CHECK-SPIRV: 3 FunctionParameter {{[0-9]+}} [[ARGD]] -; CHECK-SPIRV: 3 FunctionParameter {{[0-9]+}} [[ARGE]] - -; CHECK-LLVM: define spir_kernel void @test{{.*}} !kernel_arg_buffer_location ![[BUFLOC_MD:[0-9]+]] {{.*}} -; CHECK-LLVM: ![[BUFLOC_MD]] = !{i32 1, i32 2, i32 -1, i32 -1, i32 -1} - -; ModuleID = 'buffer_location.cl' -source_filename = "buffer_location.cl" -target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" -target triple = "spir64-unknown-unknown" - -; Function Attrs: norecurse nounwind readnone -define spir_kernel void @test(i32 addrspace(1)* %a, float addrspace(1)* %b, i32 addrspace(1)* %c, i32 %d, i32 %e) local_unnamed_addr !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_buffer_location !6 { -entry: - ret void -} - -!opencl.enable.FP_CONTRACT = !{} -!opencl.ocl.version = !{!0} -!opencl.spir.version = !{!0} -!opencl.used.extensions = !{!1} -!opencl.used.optional.core.features = !{!1} -!opencl.compiler.options = !{!1} -!llvm.ident = !{!2} - -!0 = !{i32 2, i32 0} -!1 = !{} -!2 = !{!""} -!3 = !{i32 1, i32 1, i32 1} -!4 = !{!"none", !"none", !"none"} -!5 = !{!"int*", !"float*", !"int*"} -!6 = !{i32 1, i32 2, i32 -1, i32 -1, i32 3} From 677699cdaf61ff53f1d3c8e6f98a3fc81f45638a Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Sun, 2 Aug 2020 15:49:20 +0300 Subject: [PATCH 11/12] Revert "Revert "Add support for SPV_INTEL_fpga_buffer_location extension"" This reverts commit ab5b6da096597e81415fd999ce5508c18260c428. --- llvm-spirv/include/LLVMSPIRVExtensions.inc | 1 + llvm-spirv/lib/SPIRV/SPIRVReader.cpp | 30 +++++++++ llvm-spirv/lib/SPIRV/SPIRVWriter.cpp | 16 +++++ llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h | 4 ++ llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h | 2 + .../lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h | 2 + llvm-spirv/lib/SPIRV/libSPIRV/spirv.hpp | 2 + .../test/transcoding/FPGABufferLocation.ll | 61 +++++++++++++++++++ 8 files changed, 118 insertions(+) create mode 100644 llvm-spirv/test/transcoding/FPGABufferLocation.ll diff --git a/llvm-spirv/include/LLVMSPIRVExtensions.inc b/llvm-spirv/include/LLVMSPIRVExtensions.inc index 4a037aeddab2..adaa48efadc2 100644 --- a/llvm-spirv/include/LLVMSPIRVExtensions.inc +++ b/llvm-spirv/include/LLVMSPIRVExtensions.inc @@ -23,3 +23,4 @@ EXT(SPV_INTEL_optimization_hints) EXT(SPV_INTEL_float_controls2) EXT(SPV_INTEL_vector_compute) EXT(SPV_INTEL_usm_storage_classes) +EXT(SPV_INTEL_fpga_buffer_location) diff --git a/llvm-spirv/lib/SPIRV/SPIRVReader.cpp b/llvm-spirv/lib/SPIRV/SPIRVReader.cpp index d7b25c350a8d..bc7386e26d76 100644 --- a/llvm-spirv/lib/SPIRV/SPIRVReader.cpp +++ b/llvm-spirv/lib/SPIRV/SPIRVReader.cpp @@ -168,6 +168,27 @@ static void addOCLKernelArgumentMetadata( Fn->setMetadata(MDName, MDNode::get(*Context, ValueVec)); } +static void addBufferLocationMetadata( + LLVMContext *Context, SPIRVFunction *BF, llvm::Function *Fn, + std::function Func) { + std::vector ValueVec; + bool DecorationFound = false; + BF->foreachArgument([&](SPIRVFunctionParameter *Arg) { + if (Arg->getType()->isTypePointer() && + Arg->hasDecorate(DecorationBufferLocationINTEL)) { + DecorationFound = true; + ValueVec.push_back(Func(Arg)); + } else { + llvm::Metadata *DefaultNode = ConstantAsMetadata::get( + ConstantInt::get(Type::getInt32Ty(*Context), -1)); + ValueVec.push_back(DefaultNode); + } + }); + if (DecorationFound) + Fn->setMetadata("kernel_arg_buffer_location", + MDNode::get(*Context, ValueVec)); +} + Value *SPIRVToLLVM::getTranslatedValue(SPIRVValue *BV) { auto Loc = ValueMap.find(BV); if (Loc != ValueMap.end()) @@ -3519,6 +3540,15 @@ bool SPIRVToLLVM::transOCLMetadata(SPIRVFunction *BF) { Arg->getName()); }); } + // Generate metadata for kernel_arg_buffer_location + addBufferLocationMetadata(Context, BF, F, [=](SPIRVFunctionParameter *Arg) { + auto Literals = Arg->getDecorationLiterals(DecorationBufferLocationINTEL); + assert(Literals.size() == 1 && + "BufferLocationINTEL decoration shall have 1 ID literal"); + + return ConstantAsMetadata::get( + ConstantInt::get(Type::getInt32Ty(*Context), Literals[0])); + }); return true; } diff --git a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp index f187a47bf6cd..25988f58ea76 100644 --- a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp +++ b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp @@ -527,6 +527,13 @@ SPIRVFunction *LLVMToSPIRV::transFunctionDecl(Function *F) { BM->addEntryPoint(ExecutionModelKernel, BF->getId()); else if (F->getLinkage() != GlobalValue::InternalLinkage) BF->setLinkageType(transLinkageType(F)); + + // Translate OpenCL/SYCL buffer_location metadata if it's attached to the + // translated function declaration + MDNode *BufferLocation = nullptr; + if (BM->isAllowedToUseExtension(ExtensionID::SPV_INTEL_fpga_buffer_location)) + BufferLocation = ((*F).getMetadata("kernel_arg_buffer_location")); + auto Attrs = F->getAttributes(); for (Function::arg_iterator I = F->arg_begin(), E = F->arg_end(); I != E; @@ -552,6 +559,15 @@ SPIRVFunction *LLVMToSPIRV::transFunctionDecl(Function *F) { BA->addDecorate(DecorationMaxByteOffset, Attrs.getAttribute(ArgNo + 1, Attribute::Dereferenceable) .getDereferenceableBytes()); + if (BufferLocation && I->getType()->isPointerTy()) { + // Order of integer numbers in MD node follows the order of function + // parameters on which we shall attach the appropriate decoration. Add + // decoration only if MD value is not negative. + BM->addCapability(CapabilityFPGABufferLocationINTEL); + int LocID = getMDOperandAsInt(BufferLocation, ArgNo); + if (LocID >= 0) + BA->addDecorate(DecorationBufferLocationINTEL, LocID); + } } if (Attrs.hasAttribute(AttributeList::ReturnIndex, Attribute::ZExt)) BF->addDecorate(DecorationFuncParamAttr, FunctionParameterAttributeZext); diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h index 31c60f0f2e24..ba03b181cf3c 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h @@ -164,6 +164,8 @@ class SPIRVDecorate : public SPIRVDecorateGeneric { return getSet(ExtensionID::SPV_INTEL_function_pointers); case DecorationIOPipeStorageINTEL: return getSet(ExtensionID::SPV_INTEL_io_pipes); + case DecorationBufferLocationINTEL: + return getSet(ExtensionID::SPV_INTEL_fpga_buffer_location); default: return SPIRVExtSet(); } @@ -267,6 +269,8 @@ class SPIRVMemberDecorate : public SPIRVDecorateGeneric { return getSet(ExtensionID::SPV_INTEL_fpga_memory_accesses); case DecorationIOPipeStorageINTEL: return getSet(ExtensionID::SPV_INTEL_io_pipes); + case DecorationBufferLocationINTEL: + return getSet(ExtensionID::SPV_INTEL_fpga_buffer_location); default: return SPIRVExtSet(); } diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h index 0ced7ce27f40..c35b4f01f173 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h @@ -402,6 +402,8 @@ template <> inline void SPIRVMap::init() { ADD_VEC_INIT(DecorationDontStaticallyCoalesceINTEL, {CapabilityFPGAMemoryAccessesINTEL}); ADD_VEC_INIT(DecorationPrefetchINTEL, {CapabilityFPGAMemoryAccessesINTEL}); + ADD_VEC_INIT(DecorationBufferLocationINTEL, + {CapabilityFPGABufferLocationINTEL}); } template <> inline void SPIRVMap::init() { diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h index e096040ed072..c9915382e4fc 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h +++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h @@ -376,6 +376,7 @@ template <> inline void SPIRVMap::init() { add(DecorationGlobalVariableOffsetINTEL, "GlobalVariableOffsetINTEL"); add(DecorationFuncParamIOKind, "FuncParamIOKind"); add(DecorationSIMTCallINTEL, "SIMTCallINTEL"); + add(DecorationBufferLocationINTEL, "BufferLocationINTEL"); } SPIRV_DEF_NAMEMAP(Decoration, SPIRVDecorationNameMap) @@ -560,6 +561,7 @@ template <> inline void SPIRVMap::init() { "GroupNonUniformShuffleRelative"); add(CapabilityGroupNonUniformClustered, "GroupNonUniformClustered"); add(CapabilityUSMStorageClassesINTEL, "USMStorageClassesINTEL"); + add(CapabilityFPGABufferLocationINTEL, "FPGABufferLocationINTEL"); } SPIRV_DEF_NAMEMAP(Capability, SPIRVCapabilityNameMap) diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/spirv.hpp b/llvm-spirv/lib/SPIRV/libSPIRV/spirv.hpp index 22ccb3f8b17c..53492791c7e7 100644 --- a/llvm-spirv/lib/SPIRV/libSPIRV/spirv.hpp +++ b/llvm-spirv/lib/SPIRV/libSPIRV/spirv.hpp @@ -508,6 +508,7 @@ enum Decoration { DecorationCacheSizeINTEL = 5900, DecorationDontStaticallyCoalesceINTEL = 5901, DecorationPrefetchINTEL = 5902, + DecorationBufferLocationINTEL = 5921, DecorationIOPipeStorageINTEL = 5944, DecorationMax = 0x7fffffff, }; @@ -962,6 +963,7 @@ enum Capability { CapabilityFPGARegINTEL = 5948, CapabilityKernelAttributesINTEL = 5892, CapabilityFPGAKernelAttributesINTEL = 5897, + CapabilityFPGABufferLocationINTEL = 5920, CapabilityUSMStorageClassesINTEL = 5935, CapabilityFPGAMemoryAccessesINTEL = 5898, CapabilityIOPipeINTEL = 5943, diff --git a/llvm-spirv/test/transcoding/FPGABufferLocation.ll b/llvm-spirv/test/transcoding/FPGABufferLocation.ll new file mode 100644 index 000000000000..e649a89625d5 --- /dev/null +++ b/llvm-spirv/test/transcoding/FPGABufferLocation.ll @@ -0,0 +1,61 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_fpga_buffer_location -o %t.spv +; RUN: llvm-spirv %t.spv -to-text -o %t.spt +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV + +; RUN: llvm-spirv -r %t.spv -o %t.rev.bc +; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM + +; RUN: llvm-spirv -spirv-text -r %t.spt -o %t.rev.bc +; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM + +; CHECK-SPIRV: 2 Capability FPGABufferLocationINTEL +; CHECK-SPIRV: 9 Extension "SPV_INTEL_fpga_buffer_location" +; CHECK-SPIRV: 3 Name [[ARGA:[0-9]+]] "a" +; CHECK-SPIRV: 3 Name [[ARGB:[0-9]+]] "b" +; CHECK-SPIRV: 3 Name [[ARGC:[0-9]+]] "c" +; CHECK-SPIRV: 3 Name [[ARGD:[0-9]+]] "d" +; CHECK-SPIRV: 3 Name [[ARGE:[0-9]+]] "e" +; CHECK-SPIRV-NOT: 4 Decorate [[ARGC]] BufferLocationINTEL -1 +; CHECK-SPIRV-NOT: 4 Decorate [[ARGC]] BufferLocationINTEL -1 +; CHECK-SPIRV: 4 Decorate [[ARGA]] BufferLocationINTEL 1 +; CHECK-SPIRV: 4 Decorate [[ARGB]] BufferLocationINTEL 2 +; CHECK-SPIRV-NOT: 4 Decorate [[ARGD]] BufferLocationINTEL -1 +; CHECK-SPIRV-NOT: 4 Decorate [[ARGE]] BufferLocationINTEL 3 + +; CHECK-SPIRV: 5 Function +; CHECK-SPIRV: 3 FunctionParameter {{[0-9]+}} [[ARGA]] +; CHECK-SPIRV: 3 FunctionParameter {{[0-9]+}} [[ARGB]] +; CHECK-SPIRV: 3 FunctionParameter {{[0-9]+}} [[ARGC]] +; CHECK-SPIRV: 3 FunctionParameter {{[0-9]+}} [[ARGD]] +; CHECK-SPIRV: 3 FunctionParameter {{[0-9]+}} [[ARGE]] + +; CHECK-LLVM: define spir_kernel void @test{{.*}} !kernel_arg_buffer_location ![[BUFLOC_MD:[0-9]+]] {{.*}} +; CHECK-LLVM: ![[BUFLOC_MD]] = !{i32 1, i32 2, i32 -1, i32 -1, i32 -1} + +; ModuleID = 'buffer_location.cl' +source_filename = "buffer_location.cl" +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" +target triple = "spir64-unknown-unknown" + +; Function Attrs: norecurse nounwind readnone +define spir_kernel void @test(i32 addrspace(1)* %a, float addrspace(1)* %b, i32 addrspace(1)* %c, i32 %d, i32 %e) local_unnamed_addr !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_buffer_location !6 { +entry: + ret void +} + +!opencl.enable.FP_CONTRACT = !{} +!opencl.ocl.version = !{!0} +!opencl.spir.version = !{!0} +!opencl.used.extensions = !{!1} +!opencl.used.optional.core.features = !{!1} +!opencl.compiler.options = !{!1} +!llvm.ident = !{!2} + +!0 = !{i32 2, i32 0} +!1 = !{} +!2 = !{!""} +!3 = !{i32 1, i32 1, i32 1} +!4 = !{!"none", !"none", !"none"} +!5 = !{!"int*", !"float*", !"int*"} +!6 = !{i32 1, i32 2, i32 -1, i32 -1, i32 3} From 06e9d5c09f0d5ca2d65a17cd054b424c1893ea46 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Mon, 3 Aug 2020 22:47:28 +0300 Subject: [PATCH 12/12] Add few diag Signed-off-by: Dmitry Sidorov --- .../clang/Basic/DiagnosticSemaKinds.td | 9 +++++--- clang/lib/Sema/SemaSYCL.cpp | 23 +++++++++++++++---- 2 files changed, 25 insertions(+), 7 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index b5ea32068277..7df4222cc255 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10976,11 +10976,14 @@ def err_sycl_attibute_cannot_be_applied_here "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_list_template_param : Error< - "%0 template parameter must be a " - "%select{parameter pack|type|non-negative integer}1">; +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 " + "template parameter must be a " + "%select{parameter pack|type|non-negative integer}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 " diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 54818cf448df..dbbf1db2fadc 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1098,14 +1098,24 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { } const auto *PropListDecl = cast(PropListTy->getAsRecordDecl()); + if (PropListDecl->getTemplateArgs().size() != 1) { + SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_param_number) + << "property_list"; + return; + } const auto TemplArg = PropListDecl->getTemplateArgs()[0]; if (TemplArg.getKind() != TemplateArgument::ArgKind::Pack) { SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param) - << "property_list" << /*parameter pack*/ 0; + << /*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; + return; + } QualType PropTy = Prop->getAsType(); if (Util::isSyclBufferLocationType(PropTy)) checkBufferLocationType(PropTy, Loc); @@ -1114,17 +1124,22 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { void checkBufferLocationType(QualType PropTy, SourceLocation Loc) { const auto *PropDecl = - cast(PropTy->getAsRecordDecl()); + dyn_cast(PropTy->getAsRecordDecl()); + if (PropDecl->getTemplateArgs().size() != 1) { + SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_param_number) + << "buffer_location"; + return; + } const auto BufferLoc = PropDecl->getTemplateArgs()[0]; if (BufferLoc.getKind() != TemplateArgument::ArgKind::Integral) { SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param) - << "buffer_location" << /*non-negative integer*/ 2; + << /*buffer_location*/ 2 << /*non-negative integer*/ 2; return; } int LocationID = static_cast(BufferLoc.getAsIntegral().getExtValue()); if (LocationID < 0) { SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param) - << "buffer_location" << /*non-negative integer*/ 2; + << /*buffer_location*/ 2 << /*non-negative integer*/ 2; return; } }