diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 73dec5e69c61a..8e8d16f9e5995 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1188,6 +1188,14 @@ def SYCLDeviceIndirectlyCallable : InheritableAttr { let PragmaAttributeSupport = 0; } +def SYCLIntelBufferLocation : InheritableAttr { + // No spelling, as this attribute can't be created in the source code. + let Spellings = []; + let Args = [UnsignedArgument<"LocationID">]; + let LangOpts = [SYCLIsDevice, SYCLIsHost]; + let Documentation = [SYCLIntelBufferLocationAttrDocs]; +} + 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 467b96793263f..56130e6934348 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -1994,6 +1994,25 @@ can be lowered. }]; } +def SYCLIntelBufferLocationAttrDocs : Documentation { + let Category = DocCatFunction; + let Heading = "kernel_args_buffer_location"; + let Content = [{ +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. +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. + }]; +} + 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 62926c9edaf16..7df4222cc255e 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10974,6 +10974,16 @@ 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_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/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 43f0db10f49b9..c4ef7d42ecda1 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/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index ad5a3768846c7..dbbf1db2fadca 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -80,6 +80,14 @@ class Util { /// half class. static bool isSyclHalfType(const QualType &Ty); + /// Checks whether given clang type is a full specialization of the SYCL + /// property_list class. + static bool isPropertyListType(const QualType &Ty); + + /// Checks whether given clang type is a full specialization of the SYCL + /// buffer_location class. + static bool isSyclBufferLocationType(const QualType &Ty); + /// Checks whether given clang type is a standard SYCL API class with given /// name. /// \param Ty the clang type being checked @@ -1076,6 +1084,66 @@ 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()); + 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*/ 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); + } + } + + void checkBufferLocationType(QualType PropTy, SourceLocation Loc) { + const auto *PropDecl = + 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*/ 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*/ 2 << /*non-negative integer*/ 2; + return; + } + } + void checkAccessorType(QualType Ty, SourceRange Loc) { assert(Util::isSyclAccessorType(Ty) && "Should only be called on SYCL accessor types."); @@ -1087,6 +1155,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); } @@ -1158,8 +1228,9 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { void addParam(ParamDesc newParamDesc, QualType FieldTy) { // 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()); @@ -1169,11 +1240,56 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { Params.push_back(NewParam); } + // 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; + const auto PropList = cast(AccTy->getTemplateArgs()[5]); + QualType PropListTy = PropList.getAsType(); + const auto *PropListDecl = + cast(PropListTy->getAsRecordDecl()); + const auto TemplArg = PropListDecl->getTemplateArgs()[0]; + // Move through TemplateArgs list of a property list and search for + // properties. If found - apply the appropriate attribute to ParmVarDecl. + for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin(); + Prop != TemplArg.pack_end(); ++Prop) { + QualType PropTy = Prop->getAsType(); + 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]; + int LocationID = static_cast(BufferLoc.getAsIntegral().getExtValue()); + Param->addAttr( + SYCLIntelBufferLocationAttr::CreateImplicit(Ctx, LocationID)); + } + // All special SYCL objects must have __init method. We extract types for // kernel parameters from __init method parameters. We will use __init method // and kernel parameters which we build here to initialize special objects in // the kernel body. - bool handleSpecialType(FieldDecl *FD, QualType FieldTy) { + bool handleSpecialType(FieldDecl *FD, QualType FieldTy, + bool isAccessorType = false) { const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); @@ -1182,8 +1298,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(); - for (const ParmVarDecl *Param : InitMethod->parameters()) - addParam(FD, Param->getType().getCanonicalType()); + 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()); + } LastParamIndex = ParamIndex; return true; } @@ -1253,14 +1374,18 @@ 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()); + 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; } bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final { - return handleSpecialType(FD, FieldTy); + return handleSpecialType(FD, FieldTy, /*isAccessorType*/ true); } bool handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final { @@ -2821,6 +2946,23 @@ bool Util::isSyclSpecConstantType(const QualType &Ty) { return matchQualifiedTypeName(Ty, Scopes); } +bool Util::isPropertyListType(const QualType &Ty) { + return isSyclType(Ty, "property_list", true /*Tmpl*/); +} + +bool Util::isSyclBufferLocationType(const QualType &Ty) { + const StringRef &Name = "buffer_location"; + std::array 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 3184c58edcbfc..de3e8999961b0 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -71,15 +71,23 @@ enum prop_type { base_prop }; +// Compile time known accessor property +// TODO: this doesn't belong to property namespace, instead it shall be in its +// own namespace. Change it, when the actual implementation in SYCL headers is +// ready +template +class buffer_location {}; + struct property_base { virtual prop_type type() const = 0; }; } // namespace property +template class property_list { public: - template - property_list(propertyTN... props) {} + template + property_list(propertiesTN... props){}; template bool has_property() const { return true; } @@ -127,7 +135,8 @@ struct _ImplT { template + access::placeholder isPlaceholder = access::placeholder::false_t, + typename propertyListT = property_list<>> class accessor { public: @@ -141,6 +150,8 @@ class accessor { private: void __init(__attribute__((opencl_global)) dataT *Ptr, range AccessRange, range MemRange, id Offset) {} + + propertyListT prop_list; }; template @@ -326,7 +337,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 +350,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 +424,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 e197c339c1251..b3857806a2bcd 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 0000000000000..ae1b9088cc271 --- /dev/null +++ b/clang/test/CodeGenSYCL/buffer_location.cpp @@ -0,0 +1,41 @@ +// 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, 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; + cl::sycl::kernel_single_task( + [=]() { + accessorA.use(); + Obj.use(); + }); + return 0; +} diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index c63e64a37f117..e285bfce8f536 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 8c2cfb2a1bd8b..902ad7ddc339c 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 f5f679f7d3650..ddfaca9664851 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 a00f147b0dee0..b1382ec6c2b92 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/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 9e3efc6321096..1b3ab8b20763f 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 17dafe7b4acdd..db07ee25a5457 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 1f500eff0a888..54ae68ae5b8ce 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 0000000000000..bd2daaa01ba41 --- /dev/null +++ b/clang/test/SemaSYCL/buffer_location.cpp @@ -0,0 +1,92 @@ +// 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 { +}; + +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>> + // 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; + cl::sycl::accessor, + cl::sycl::property::buffer_location<2>>> + accessorF; +#endif + cl::sycl::kernel_single_task( + [=]() { +#ifndef TRIGGER_ERROR + // expected-no-diagnostics + Obj.use(); + accessorA.use(); + accessorB.use(); + accessorC.use(); +#else + //expected-error@+1{{buffer_location template parameter must be a non-negative integer}} + accessorD.use(); + //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(); +#endif + }); + return 0; +} diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp index 1052b4ac24e0f..9f9a3196024c1 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