diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 138edfd63f550..f18f0409e7337 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -14,13 +14,13 @@ #include "clang/AST/QualTypeNames.h" #include "clang/AST/RecordLayout.h" #include "clang/AST/RecursiveASTVisitor.h" +#include "clang/Analysis/CallGraph.h" #include "clang/Sema/Sema.h" #include "llvm/ADT/SmallPtrSet.h" #include "llvm/ADT/SmallVector.h" #include "llvm/Support/FileSystem.h" #include "llvm/Support/Path.h" #include "llvm/Support/raw_ostream.h" -#include "clang/Analysis/CallGraph.h" #include @@ -95,17 +95,17 @@ class MarkDeviceFunction : public RecursiveASTVisitor { // all functions used by kernel have already been parsed and have // definitions. if (RecursiveSet.count(Callee)) { - SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) << - KernelCallRecursiveFunction; + SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) + << KernelCallRecursiveFunction; SemaRef.Diag(Callee->getSourceRange().getBegin(), - diag::note_sycl_recursive_function_declared_here) - << KernelCallRecursiveFunction; + diag::note_sycl_recursive_function_declared_here) + << KernelCallRecursiveFunction; } if (const CXXMethodDecl *Method = dyn_cast(Callee)) if (Method->isVirtual()) - SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) << - KernelCallVirtualFunction; + SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) + << KernelCallVirtualFunction; CheckSYCLType(Callee->getReturnType(), Callee->getSourceRange()); @@ -116,8 +116,8 @@ class MarkDeviceFunction : public RecursiveASTVisitor { } } } else if (!SemaRef.getLangOpts().SYCLAllowFuncPtr) - SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) << - KernelCallFunctionPointer; + SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) + << KernelCallFunctionPointer; return true; } @@ -178,8 +178,8 @@ class MarkDeviceFunction : public RecursiveASTVisitor { if (VarDecl *VD = dyn_cast(E->getMemberDecl())) { bool IsConst = VD->getType().getNonReferenceType().isConstQualified(); if (VD->isStaticDataMember() && !IsConst) - SemaRef.Diag(E->getExprLoc(), diag::err_sycl_restrict) << - KernelNonConstStaticDataVariable; + SemaRef.Diag(E->getExprLoc(), diag::err_sycl_restrict) + << KernelNonConstStaticDataVariable; } return true; } @@ -190,24 +190,24 @@ class MarkDeviceFunction : public RecursiveASTVisitor { bool IsConst = VD->getType().getNonReferenceType().isConstQualified(); if (!IsConst && VD->hasGlobalStorage() && !VD->isStaticLocal() && !VD->isStaticDataMember() && !isa(VD)) - SemaRef.Diag(E->getLocation(), diag::err_sycl_restrict) << - KernelGlobalVariable; + SemaRef.Diag(E->getLocation(), diag::err_sycl_restrict) + << KernelGlobalVariable; } return true; } bool VisitCXXNewExpr(CXXNewExpr *E) { - // Memory storage allocation is not allowed in kernels. - // All memory allocation for the device is done on - // the host using accessor classes. Consequently, the default - // allocation operator new overloads that allocate - // storage are disallowed in a SYCL kernel. The placement - // new operator and any user-defined overloads that - // do not allocate storage are permitted. + // Memory storage allocation is not allowed in kernels. + // All memory allocation for the device is done on + // the host using accessor classes. Consequently, the default + // allocation operator new overloads that allocate + // storage are disallowed in a SYCL kernel. The placement + // new operator and any user-defined overloads that + // do not allocate storage are permitted. if (FunctionDecl *FD = E->getOperatorNew()) { if (FD->isReplaceableGlobalAllocationFunction()) { - SemaRef.Diag(E->getExprLoc(), diag::err_sycl_restrict) << - KernelAllocateStorage; + SemaRef.Diag(E->getExprLoc(), diag::err_sycl_restrict) + << KernelAllocateStorage; } else if (FunctionDecl *Def = FD->getDefinition()) { if (!Def->hasAttr()) { Def->addAttr(SYCLDeviceAttr::CreateImplicit(SemaRef.Context)); @@ -219,26 +219,26 @@ class MarkDeviceFunction : public RecursiveASTVisitor { } bool VisitCXXThrowExpr(CXXThrowExpr *E) { - SemaRef.Diag(E->getExprLoc(), diag::err_sycl_restrict) << - KernelUseExceptions; + SemaRef.Diag(E->getExprLoc(), diag::err_sycl_restrict) + << KernelUseExceptions; return true; } bool VisitCXXCatchStmt(CXXCatchStmt *S) { - SemaRef.Diag(S->getBeginLoc(), diag::err_sycl_restrict) << - KernelUseExceptions; + SemaRef.Diag(S->getBeginLoc(), diag::err_sycl_restrict) + << KernelUseExceptions; return true; } bool VisitCXXTryStmt(CXXTryStmt *S) { - SemaRef.Diag(S->getBeginLoc(), diag::err_sycl_restrict) << - KernelUseExceptions; + SemaRef.Diag(S->getBeginLoc(), diag::err_sycl_restrict) + << KernelUseExceptions; return true; } bool VisitSEHTryStmt(SEHTryStmt *S) { - SemaRef.Diag(S->getBeginLoc(), diag::err_sycl_restrict) << - KernelUseExceptions; + SemaRef.Diag(S->getBeginLoc(), diag::err_sycl_restrict) + << KernelUseExceptions; return true; } @@ -291,8 +291,8 @@ class MarkDeviceFunction : public RecursiveASTVisitor { } } } -private: +private: bool CheckSYCLType(QualType Ty, SourceRange Loc) { if (Ty->isVariableArrayType()) { SemaRef.Diag(Loc.getBegin(), diag::err_vla_unsupported); @@ -306,7 +306,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor { // FIXME: this seems like a temporary fix for SYCL programs // that pre-declare, use, but not define OclCXX classes, // which are later translated into SPIRV types. - if(!CRD->hasDefinition()) + if (!CRD->hasDefinition()) return true; if (CRD->isPolymorphic()) { @@ -450,9 +450,10 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { QualType FieldType = Field->getType(); CXXRecordDecl *CRD = FieldType->getAsCXXRecordDecl(); if (CRD && Util::isSyclAccessorType(FieldType)) { - // Since this is an accessor next 3 TargetFuncParams including current - // should be set in __init method: _ValueType*, range, id - const size_t NumParams = 3; + // Since this is an accessor next 4 TargetFuncParams including current + // should be set in __init method: _ValueType*, range, range, + // id + const size_t NumParams = 4; llvm::SmallVector ParamDREs(NumParams); auto TFP = TargetFuncParam; for (size_t I = 0; I < NumParams; ++TFP, ++I) { @@ -495,7 +496,7 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { ExprValueKind VK = Expr::getValueKindForType(ResultTy); ResultTy = ResultTy.getNonLValueExprType(S.Context); - // __init needs three parameter + // __init needs four parameter auto ParamItr = InitMethod->param_begin(); // kernel_parameters llvm::SmallVector ParamStmts; @@ -505,7 +506,10 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { S, ((*ParamItr++))->getOriginalType(), ParamDREs[1])); ParamStmts.push_back(getExprForRangeOrOffset( S, ((*ParamItr++))->getOriginalType(), ParamDREs[2])); - // kernel_obj.accessor.__init(_ValueType*, range, id) + ParamStmts.push_back(getExprForRangeOrOffset( + S, ((*ParamItr++))->getOriginalType(), ParamDREs[3])); + // kernel_obj.accessor.__init(_ValueType*, range, range, + // id) CXXMemberCallExpr *Call = CXXMemberCallExpr::Create( S.Context, ME, ParamStmts, ResultTy, VK, SourceLocation()); BodyStmts.push_back(Call); @@ -643,10 +647,17 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, CreateAndAddPrmDsc(Fld, PointerType); - FieldDecl *RangeFld = getFieldDeclByName(RecordDecl, {"__impl", "Range"}); - assert(RangeFld && - "The accessor must contain the Range from the __impl field"); - CreateAndAddPrmDsc(RangeFld, RangeFld->getType()); + FieldDecl *AccessRangeFld = + getFieldDeclByName(RecordDecl, {"__impl", "AccessRange"}); + assert(AccessRangeFld && + "The accessor must contain the AccessRange from the __impl field"); + CreateAndAddPrmDsc(AccessRangeFld, AccessRangeFld->getType()); + + FieldDecl *MemRangeFld = + getFieldDeclByName(RecordDecl, {"__impl", "MemRange"}); + assert(MemRangeFld && + "The accessor must contain the MemRange from the __impl field"); + CreateAndAddPrmDsc(MemRangeFld, MemRangeFld->getType()); FieldDecl *OffsetFld = getFieldDeclByName(RecordDecl, {"__impl", "Offset"}); @@ -705,13 +716,22 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name, const auto *AccTmplTy = cast(AccTy); H.addParamDesc(SYCLIntegrationHeader::kind_accessor, getAccessTarget(AccTmplTy), Offset); - // ... second descriptor (translated to range kernel parameter): - FieldDecl *RngFld = - getFieldDeclByName(AccTy, {"__impl", "Range"}, &Offset); - uint64_t Sz = Ctx.getTypeSizeInChars(RngFld->getType()).getQuantity(); + // ... second descriptor (translated to access range kernel parameter): + FieldDecl *AccessRngFld = + getFieldDeclByName(AccTy, {"__impl", "AccessRange"}, &Offset); + uint64_t Sz = + Ctx.getTypeSizeInChars(AccessRngFld->getType()).getQuantity(); + H.addParamDesc(SYCLIntegrationHeader::kind_std_layout, + static_cast(Sz), static_cast(Offset)); + // ... third descriptor (translated to mem range kernel parameter): + // Get offset in bytes + Offset = Layout.getFieldOffset(Fld->getFieldIndex()) / 8; + FieldDecl *MemRngFld = + getFieldDeclByName(AccTy, {"__impl", "MemRange"}, &Offset); + Sz = Ctx.getTypeSizeInChars(MemRngFld->getType()).getQuantity(); H.addParamDesc(SYCLIntegrationHeader::kind_std_layout, static_cast(Sz), static_cast(Offset)); - // ... third descriptor (translated to id kernel parameter): + // ... fourth descriptor (translated to id kernel parameter): // Get offset in bytes Offset = Layout.getFieldOffset(Fld->getFieldIndex()) / 8; FieldDecl *OffstFld = diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index def0bb0b8d4a9..cffc9804eb507 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -64,9 +64,11 @@ class property_list { template property_list(propertyTN... props) {} - template bool has_property() const { return true; } + template + bool has_property() const { return true; } - template propertyT get_property() const { + template + propertyT get_property() const { return propertyT{}; } @@ -77,12 +79,14 @@ class property_list { template struct id { - template id(T...args) {} // fake constructor + template + id(T... args) {} // fake constructor }; template struct range { - template range(T...args) {} // fake constructor + template + range(T... args) {} // fake constructor }; template @@ -91,35 +95,36 @@ struct nd_range { template struct _ImplT { - range Range; - id Offset; + range AccessRange; + range MemRange; + id Offset; }; template - class accessor { - - public: - - void __init(__global dataT *Ptr, range Range, - id Offset) { - } - void use(void) const {} - template void use(T...args) { } - template void use(T...args) const { } - _ImplT __impl; + access::target accessTarget = access::target::global_buffer, + access::placeholder isPlaceholder = access::placeholder::false_t> +class accessor { +public: + void __init(__global dataT *Ptr, range AccessRange, + range MemRange, id Offset) {} + void use(void) const {} + template + void use(T... args) {} + template + void use(T... args) const {} + _ImplT __impl; }; class kernel {}; class context {}; class device {}; -class event{}; +class event {}; class queue { public: - template event submit(T cgf) { return event{}; } + template + event submit(T cgf) { return event{}; } void wait() {} void wait_and_throw() {} @@ -177,7 +182,8 @@ class buffer { using const_reference = const value_type &; using allocator_type = AllocatorT; - template buffer(ParamTypes...args) {} // fake constructor + template + buffer(ParamTypes... args) {} // fake constructor buffer(const range &bufferRange, const property_list &propList = {}) {} @@ -221,4 +227,3 @@ class buffer { } // namespace sycl } // namespace cl - diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index dfa256ebde142..729861e544cac 100644 --- a/clang/test/CodeGenSYCL/integration_header.cpp +++ b/clang/test/CodeGenSYCL/integration_header.cpp @@ -25,21 +25,25 @@ // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2014, 4 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 5 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 6 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 6 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 7 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 7 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 8 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 9 }, // 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, 2016, 4 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 5 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 6 }, // 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, 2016, 4 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 5 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 6 }, // CHECK-EMPTY: // CHECK-NEXT: }; // diff --git a/clang/test/CodeGenSYCL/struct_kernel_param.cpp b/clang/test/CodeGenSYCL/struct_kernel_param.cpp index c66ec4f2b16d7..089c3ab1d987e 100644 --- a/clang/test/CodeGenSYCL/struct_kernel_param.cpp +++ b/clang/test/CodeGenSYCL/struct_kernel_param.cpp @@ -6,6 +6,7 @@ // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2014, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 1 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 2 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 24, 4 }, // CHECK-EMPTY: // CHECK-NEXT:}; diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index edd6ebbc41061..8015e3ecd9210 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -51,7 +51,8 @@ struct id { template struct _ImplT { - range Range; + range AccessRange; + range MemRange; id Offset; }; @@ -63,10 +64,8 @@ class accessor { public: void use(void) const {} void use(void*) const {} - void __init(__global dataT *Ptr, range Range, - id Offset) { - } - + void __init(__global dataT *Ptr, range AccessRange, + range MemRange, id Offset) {} _ImplT __impl; }; diff --git a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp index 16ce5b8c0b4fe..ba3b080a66760 100644 --- a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp +++ b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp @@ -21,5 +21,5 @@ int main() { }); return 0; } -// CHECK: kernel_function 'void (__global int *, range<1>, id<1>) -// CHECK: kernel_local_acc 'void (__local int *, range<1>, id<1>) +// CHECK: kernel_function 'void (__global int *, range<1>, range<1>, id<1>) +// CHECK: kernel_local_acc 'void (__local int *, range<1>, range<1>, id<1>) diff --git a/clang/test/SemaSYCL/fake-accessors.cpp b/clang/test/SemaSYCL/fake-accessors.cpp index 280dc850167e1..112f986d20083 100644 --- a/clang/test/SemaSYCL/fake-accessors.cpp +++ b/clang/test/SemaSYCL/fake-accessors.cpp @@ -51,6 +51,6 @@ int main() { }); return 0; } -// CHECK: fake_accessors 'void (__global int *, range<1>, id<1>, foo::cl::sycl::accessor, accessor) -// CHECK: accessor_typedef 'void (__global int *, range<1>, id<1>, foo::cl::sycl::accessor, accessor) -// CHECK: accessor_alias 'void (__global int *, range<1>, id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: fake_accessors 'void (__global int *, range<1>, range<1>, id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_typedef 'void (__global int *, range<1>, range<1>, id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_alias 'void (__global int *, range<1>, range<1>, id<1>, foo::cl::sycl::accessor, accessor)