From 0db566cd543cc7bbd07ac4e2c0403807bf0f777f Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Wed, 10 Mar 2021 21:34:14 -0800 Subject: [PATCH 01/13] [WIP][SYCL] Implement SYCL 2020 specialization constants This is a draft patch to get early feedback. Spec is currently under review here - https://github.com/intel/llvm/pull/3331 Signed-off-by: Elizabeth Andrews --- clang/include/clang/Sema/Sema.h | 1 + clang/lib/Sema/SemaSYCL.cpp | 129 ++++++++++++++++++++++++- clang/test/SemaSYCL/kernel-handler.cpp | 51 ++++++++++ 3 files changed, 178 insertions(+), 3 deletions(-) create mode 100644 clang/test/SemaSYCL/kernel-handler.cpp diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index b3f2a14baf973..3e6bd3a83a4da 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -314,6 +314,7 @@ class SYCLIntegrationHeader { kind_accessor = kind_first, kind_std_layout, kind_sampler, + kind_specialization_constants_buffer, kind_pointer, kind_last = kind_pointer }; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d6a43df48e54b..c62f7e2c5adc8 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -58,6 +58,8 @@ enum KernelInvocationKind { const static std::string InitMethodName = "__init"; const static std::string InitESIMDMethodName = "__init_esimd"; +const static std::string InitSpecConstantsBuffer = + "__init_specialization_constants_buffer"; const static std::string FinalizeMethodName = "__finalize"; constexpr unsigned MaxKernelArgsSize = 2048; @@ -109,6 +111,10 @@ class Util { /// specialization constant class. static bool isSyclSpecConstantType(const QualType &Ty); + /// Checks whether given clang type is a full specialization of the SYCL + /// kernel_handler class. + static bool isSyclKernelHandlerType(const QualType &Ty); + // Checks declaration context hierarchy. /// \param DC the context of the item to be checked. /// \param Scopes the declaration scopes leading from the item context to the @@ -743,6 +749,13 @@ static ParamDesc makeParamDesc(ASTContext &Ctx, const CXXBaseSpecifier &Src, Ctx.getTrivialTypeSourceInfo(Ty)); } +// FIXME: Should we refactor makeParamDesc to just accept Name in all cases +// i.e. remove overloads. +static ParamDesc makeParamDesc(ASTContext &Ctx, StringRef Name, QualType Ty) { + return std::make_tuple(Ty, &Ctx.Idents.get(Name), + Ctx.getTrivialTypeSourceInfo(Ty)); +} + /// \return the target of given SYCL accessor type static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) { return static_cast( @@ -778,6 +791,19 @@ constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, KernelNameType)}; } +static bool hasSyclKernelHandlerArg(FunctionDecl *KernelCallerFunc) { + // Specialization constants in SYCL 2020 are not captured by lambda and + // accessed through new optional lambda argument kernel_handler + if (KernelCallerFunc->getNumParams() > 1) + return true; + // FIXME: Remember to correct this. Why does check not work? + // Are we replacing this using a special attribute? + // return + // Util::isSyclKernelHandlerType(KernelCallerFunc->getParamDecl(1)->getType()); + + return false; +} + // anonymous namespace so these don't get linkage. namespace { @@ -1647,6 +1673,12 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { makeParamDesc(SemaRef.getASTContext(), BS, FieldTy); addParam(newParamDesc, FieldTy); } + // Add a parameter with specified name and type + void addParam(StringRef Name, QualType ParamTy) { + ParamDesc newParamDesc = + makeParamDesc(SemaRef.getASTContext(), Name, ParamTy); + addParam(newParamDesc, ParamTy); + } void addParam(ParamDesc newParamDesc, QualType FieldTy) { // Create a new ParmVarDecl based on the new info. @@ -1947,6 +1979,13 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return true; } + void handleSyclKernelHandlerType() { + // Create parameters used to initialize spec constant + ASTContext &Context = SemaRef.getASTContext(); + StringRef Name = "specialization_constants_buffer"; + addParam(Name, Context.getPointerType(Context.CharTy)); + } + void setBody(CompoundStmt *KB) { KernelDecl->setBody(KB); } FunctionDecl *getKernelDecl() { return KernelDecl; } @@ -2092,6 +2131,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // pointer-struct-wrapping code to ensure that we don't try to wrap // non-top-level pointers. uint64_t StructDepth = 0; + VarDecl *KernelHandlerClone; // Using the statements/init expressions that we've created, this generates // the kernel body compound stmt. CompoundStmt needs to know its number of @@ -2114,6 +2154,17 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { SemaRef.PushFunctionScope(); KernelBodyTransform KBT(MappingPair, SemaRef); Stmt *NewBody = KBT.TransformStmt(FunctionBody).get(); + + if (hasSyclKernelHandlerArg(KernelCallerFunc)) { + // Factor this out. Repetitive code. + ParmVarDecl *KernelHandlerParam = KernelCallerFunc->getParamDecl(1); + KernelHandlerClone->setIsUsed(); + std::pair MappingPairKernelHandler = + std::make_pair(KernelHandlerParam, KernelHandlerClone); + KernelBodyTransform KBT(MappingPairKernelHandler, SemaRef); + NewBody = KBT.TransformStmt(NewBody).get(); + } + BodyStmts.push_back(NewBody); BodyStmts.insert(BodyStmts.end(), FinalizeStmts.begin(), @@ -2413,6 +2464,17 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return true; } + VarDecl *createKernelHandlerClone(ASTContext &Ctx, DeclContext *DC, + ParmVarDecl *KernelHandlerArg) { + QualType Ty = KernelHandlerArg->getType(); + TypeSourceInfo *TSInfo = Ctx.getTrivialTypeSourceInfo(Ty); + VarDecl *VD = + VarDecl::Create(Ctx, DC, KernelCallerSrcLoc, KernelCallerSrcLoc, + KernelHandlerArg->getIdentifier(), Ty, TSInfo, SC_None); + + return VD; + } + public: static constexpr const bool VisitInsideSimpleContainers = false; SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, @@ -2423,7 +2485,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { DC.getKernelDecl(), KernelObj)), VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)), KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc), - KernelCallerSrcLoc(KernelCallerFunc->getLocation()) { + KernelCallerSrcLoc(KernelCallerFunc->getLocation()), + KernelHandlerClone(nullptr) { CollectionInitExprs.push_back(createInitListExpr(KernelObj)); markParallelWorkItemCalls(); @@ -2517,6 +2580,40 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return true; } + // Default inits the type, then calls the init-method in the body + void handleSyclKernelHandlerType(ParmVarDecl *KernelHandlerArg) { + + // Create local clone of kernel handler + KernelHandlerClone = createKernelHandlerClone( + SemaRef.getASTContext(), DeclCreator.getKernelDecl(), KernelHandlerArg); + + // Default initialize clone + InitializedEntity VarEntity = + InitializedEntity::InitializeVariable(KernelHandlerClone); + InitializationKind InitKind = + InitializationKind::CreateDefault(KernelCallerSrcLoc); + InitializationSequence InitSeq(SemaRef, VarEntity, InitKind, None); + ExprResult Init = InitSeq.Perform(SemaRef, VarEntity, InitKind, None); + KernelHandlerClone->setInit( + SemaRef.MaybeCreateExprWithCleanups(Init.get())); + KernelHandlerClone->setInitStyle(VarDecl::CallInit); + + // Add declaration statement to openCL kernel body + Stmt *DS = + new (SemaRef.Context) DeclStmt(DeclGroupRef(KernelHandlerClone), + KernelCallerSrcLoc, KernelCallerSrcLoc); + BodyStmts.push_back(DS); + + // Generate init call + // FIXME: Should this be restricted to targets which do not have native + // support for specialization constants? + const auto *RecordDecl = + KernelHandlerClone->getType()->getAsCXXRecordDecl(); + // FIXME: This call generates __init function bound to kernel object clone. + // Fix this. + createSpecialMethodCall(RecordDecl, InitSpecConstantsBuffer, BodyStmts); + } + bool enterStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { ++StructDepth; // Add a dummy init expression to catch the accessor initializers. @@ -2670,11 +2767,13 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { addParam(ArgTy, Kind, offsetOf(FD, ArgTy)); } void addParam(QualType ArgTy, SYCLIntegrationHeader::kernel_param_kind_t Kind, - uint64_t OffsetAdj) { + uint64_t OffsetAdj, bool IsZeroOffset = false) { uint64_t Size; Size = SemaRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); Header.addParamDesc(Kind, static_cast(Size), - static_cast(CurOffset + OffsetAdj)); + ((IsZeroOffset) + ? static_cast(OffsetAdj) + : static_cast(CurOffset + OffsetAdj))); } // Returns 'true' if the thing we're visiting (Based on the FD/QualType pair) @@ -2871,6 +2970,14 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } + void handleSyclKernelHandlerType(QualType Ty) { + // Add corresponding entry in integration header. + // Offset is zero since kernel_handler argument is not part of + // kernel object (i.e. it is not captured) + addParam(Ty, SYCLIntegrationHeader::kind_specialization_constants_buffer, 0, + /*IsZeroOffset*/ true); + } + bool enterStream(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { ++StructDepth; CurOffset += offsetOf(FD, Ty); @@ -3202,6 +3309,13 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, KernelObjVisitor Visitor{*this}; Visitor.VisitRecordBases(KernelObj, kernel_decl, kernel_body, int_header); Visitor.VisitRecordFields(KernelObj, kernel_decl, kernel_body, int_header); + + if (hasSyclKernelHandlerArg(KernelCallerFunc)) { + ParmVarDecl *KernelHandlerArg = KernelCallerFunc->getParamDecl(1); + kernel_decl.handleSyclKernelHandlerType(); + kernel_body.handleSyclKernelHandlerType(KernelHandlerArg); + int_header.handleSyclKernelHandlerType(KernelHandlerArg->getType()); + } } void Sema::MarkDevice(void) { @@ -4034,6 +4148,15 @@ bool Util::isSyclSpecConstantType(const QualType &Ty) { return matchQualifiedTypeName(Ty, Scopes); } +bool Util::isSyclKernelHandlerType(const QualType &Ty) { + const StringRef &Name = "kernel_handler"; + std::array Scopes = { + Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"}, + Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"}, + Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}}; + return matchQualifiedTypeName(Ty, Scopes); +} + bool Util::isSyclBufferLocationType(const QualType &Ty) { const StringRef &PropertyName = "buffer_location"; const StringRef &InstanceName = "instance"; diff --git a/clang/test/SemaSYCL/kernel-handler.cpp b/clang/test/SemaSYCL/kernel-handler.cpp new file mode 100644 index 0000000000000..43c166a08ada2 --- /dev/null +++ b/clang/test/SemaSYCL/kernel-handler.cpp @@ -0,0 +1,51 @@ +// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s + +// This test checks that the compiler handles kernel_handler type (for +// SYCL 2020 specialization constants) correctly. + +//FIXME: Move to headers +namespace cl { +namespace sycl { +class kernel_handler { + void __init_specialization_constants_buffer(char *specialization_constants_buffer) {} +}; +} // namespace sycl +} // namespace cl + +template +__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc, cl::sycl::kernel_handler kh) { + kernelFunc(kh); +} + +int main() { + int a; + cl::sycl::kernel_handler kh; + + a_kernel( + [=](auto) { + int local = a; + }, + kh); +} + +// Check test_kernel_handler parameters +// CHECK: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int, char *)' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' +// CHECK-NEXT: ParmVarDecl {{.*}} used specialization_constants_buffer 'char *' + +// Check declaration and initialization of kernel object local clone +// CHECK-NEXT: CompoundStmt +// CHECK-NEXT: DeclStmt +// CHECK-NEXT: VarDecl {{.*}} cinit +// CHECK-NEXT: InitListExpr +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' + +// Check declaration and initialization of kernel object local clone using default constructor +// CHECK-NEXT: DeclStmt +// CHECK-NEXT: VarDecl {{.*}} callinit +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' 'void () noexcept' + +// Check call to __init_specialization_constants_buffer +// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (char *)' lvalue .__init_specialization_constants_buffer From ea11a193f16386b9b2e60287b119f1be77627396 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Thu, 18 Mar 2021 01:06:35 -0700 Subject: [PATCH 02/13] Review comments, refactor, corrections, tests Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 169 +++++++++++------- .../int_header_sycl2020_spec_const.cpp | 36 ++++ clang/test/SemaSYCL/kernel-handler.cpp | 76 ++++++-- 3 files changed, 196 insertions(+), 85 deletions(-) create mode 100644 clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c62f7e2c5adc8..7311c328f1311 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -791,17 +791,16 @@ constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, KernelNameType)}; } -static bool hasSyclKernelHandlerArg(FunctionDecl *KernelCallerFunc) { +static ParmVarDecl *getSyclKernelHandlerArg(FunctionDecl *KernelCallerFunc) { // Specialization constants in SYCL 2020 are not captured by lambda and // accessed through new optional lambda argument kernel_handler - if (KernelCallerFunc->getNumParams() > 1) - return true; - // FIXME: Remember to correct this. Why does check not work? - // Are we replacing this using a special attribute? - // return - // Util::isSyclKernelHandlerType(KernelCallerFunc->getParamDecl(1)->getType()); - - return false; + ParmVarDecl *PVD; + auto It = std::find_if(KernelCallerFunc->param_begin(), + KernelCallerFunc->param_end(), [](ParmVarDecl *PVD) { + return Util::isSyclKernelHandlerType(PVD->getType()); + }); + PVD = (It != KernelCallerFunc->param_end()) ? *It : nullptr; + return PVD; } // anonymous namespace so these don't get linkage. @@ -1979,9 +1978,16 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return true; } + // Generate kernel argument to intialize specialization constants. This + // argument is only generated when the target has no native support for + // specialization constants void handleSyclKernelHandlerType() { - // Create parameters used to initialize spec constant + ASTContext &Context = SemaRef.getASTContext(); + llvm::Triple T = Context.getTargetInfo().getTriple(); + if (T.isSPIR() && T.getSubArch() == llvm::Triple::NoSubArch) + return; + StringRef Name = "specialization_constants_buffer"; addParam(Name, Context.getPointerType(Context.CharTy)); } @@ -2131,40 +2137,46 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // pointer-struct-wrapping code to ensure that we don't try to wrap // non-top-level pointers. uint64_t StructDepth = 0; - VarDecl *KernelHandlerClone; + VarDecl *KernelHandlerClone = nullptr; + + Stmt *replaceWithLocalClone(ParmVarDecl *OriginalParam, VarDecl *LocalClone, + Stmt *FunctionBody) { + // DeclRefExpr with valid source location but with decl which is not marked + // as used is invalid. + LocalClone->setIsUsed(); + std::pair MappingPair = + std::make_pair(OriginalParam, LocalClone); + KernelBodyTransform KBT(MappingPair, SemaRef); + return KBT.TransformStmt(FunctionBody).get(); + } // Using the statements/init expressions that we've created, this generates // the kernel body compound stmt. CompoundStmt needs to know its number of // statements in advance to allocate it, so we cannot do this as we go along. CompoundStmt *createKernelBody() { + // Push the Kernel function scope to ensure the scope isn't empty + SemaRef.PushFunctionScope(); + + // Initialize kernel object local clone assert(CollectionInitExprs.size() == 1 && "Should have been popped down to just the first one"); KernelObjClone->setInit(CollectionInitExprs.back()); - Stmt *FunctionBody = KernelCallerFunc->getBody(); - - ParmVarDecl *KernelObjParam = *(KernelCallerFunc->param_begin()); - - // DeclRefExpr with valid source location but with decl which is not marked - // as used is invalid. - KernelObjClone->setIsUsed(); - std::pair MappingPair = - std::make_pair(KernelObjParam, KernelObjClone); - - // Push the Kernel function scope to ensure the scope isn't empty - SemaRef.PushFunctionScope(); - KernelBodyTransform KBT(MappingPair, SemaRef); - Stmt *NewBody = KBT.TransformStmt(FunctionBody).get(); - - if (hasSyclKernelHandlerArg(KernelCallerFunc)) { - // Factor this out. Repetitive code. - ParmVarDecl *KernelHandlerParam = KernelCallerFunc->getParamDecl(1); - KernelHandlerClone->setIsUsed(); - std::pair MappingPairKernelHandler = - std::make_pair(KernelHandlerParam, KernelHandlerClone); - KernelBodyTransform KBT(MappingPairKernelHandler, SemaRef); - NewBody = KBT.TransformStmt(NewBody).get(); - } + // Replace references to the kernel object in kernel body, to use the + // compiler generated local clone + Stmt *NewBody = + replaceWithLocalClone(KernelCallerFunc->getParamDecl(0), KernelObjClone, + KernelCallerFunc->getBody()); + + // If kernel_handler argument is passed by SYCL kernel, replace references + // to this argument in kernel body, to use the compiler generated local + // clone + ParmVarDecl *KernelHandlerParam = getSyclKernelHandlerArg(KernelCallerFunc); + if (KernelHandlerParam) + NewBody = replaceWithLocalClone(KernelHandlerParam, KernelHandlerClone, + NewBody); + + // Use transformed body (with clones) as kernel body BodyStmts.push_back(NewBody); BodyStmts.insert(BodyStmts.end(), FinalizeStmts.begin(), @@ -2464,15 +2476,37 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return true; } - VarDecl *createKernelHandlerClone(ASTContext &Ctx, DeclContext *DC, - ParmVarDecl *KernelHandlerArg) { + // Generate __init call for kernel handler argument + void handleSpecialType(QualType KernelHandlerTy) { + DeclRefExpr *KernelHandlerCloneRef = + DeclRefExpr::Create(SemaRef.Context, NestedNameSpecifierLoc(), + KernelCallerSrcLoc, KernelHandlerClone, false, + DeclarationNameInfo(), KernelHandlerTy, VK_LValue); + const auto *RecordDecl = + KernelHandlerClone->getType()->getAsCXXRecordDecl(); + MemberExprBases.push_back(KernelHandlerCloneRef); + createSpecialMethodCall(RecordDecl, InitSpecConstantsBuffer, BodyStmts); + MemberExprBases.pop_back(); + } + + void createKernelHandlerClone(ASTContext &Ctx, DeclContext *DC, + ParmVarDecl *KernelHandlerArg) { QualType Ty = KernelHandlerArg->getType(); TypeSourceInfo *TSInfo = Ctx.getTrivialTypeSourceInfo(Ty); - VarDecl *VD = + KernelHandlerClone = VarDecl::Create(Ctx, DC, KernelCallerSrcLoc, KernelCallerSrcLoc, KernelHandlerArg->getIdentifier(), Ty, TSInfo, SC_None); - return VD; + // Default initialize clone + InitializedEntity VarEntity = + InitializedEntity::InitializeVariable(KernelHandlerClone); + InitializationKind InitKind = + InitializationKind::CreateDefault(KernelCallerSrcLoc); + InitializationSequence InitSeq(SemaRef, VarEntity, InitKind, None); + ExprResult Init = InitSeq.Perform(SemaRef, VarEntity, InitKind, None); + KernelHandlerClone->setInit( + SemaRef.MaybeCreateExprWithCleanups(Init.get())); + KernelHandlerClone->setInitStyle(VarDecl::CallInit); } public: @@ -2485,8 +2519,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { DC.getKernelDecl(), KernelObj)), VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)), KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc), - KernelCallerSrcLoc(KernelCallerFunc->getLocation()), - KernelHandlerClone(nullptr) { + KernelCallerSrcLoc(KernelCallerFunc->getLocation()) { CollectionInitExprs.push_back(createInitListExpr(KernelObj)); markParallelWorkItemCalls(); @@ -2583,20 +2616,9 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // Default inits the type, then calls the init-method in the body void handleSyclKernelHandlerType(ParmVarDecl *KernelHandlerArg) { - // Create local clone of kernel handler - KernelHandlerClone = createKernelHandlerClone( - SemaRef.getASTContext(), DeclCreator.getKernelDecl(), KernelHandlerArg); - - // Default initialize clone - InitializedEntity VarEntity = - InitializedEntity::InitializeVariable(KernelHandlerClone); - InitializationKind InitKind = - InitializationKind::CreateDefault(KernelCallerSrcLoc); - InitializationSequence InitSeq(SemaRef, VarEntity, InitKind, None); - ExprResult Init = InitSeq.Perform(SemaRef, VarEntity, InitKind, None); - KernelHandlerClone->setInit( - SemaRef.MaybeCreateExprWithCleanups(Init.get())); - KernelHandlerClone->setInitStyle(VarDecl::CallInit); + // Create and default initialize local clone of kernel handler + createKernelHandlerClone(SemaRef.getASTContext(), + DeclCreator.getKernelDecl(), KernelHandlerArg); // Add declaration statement to openCL kernel body Stmt *DS = @@ -2604,14 +2626,14 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { KernelCallerSrcLoc, KernelCallerSrcLoc); BodyStmts.push_back(DS); - // Generate init call - // FIXME: Should this be restricted to targets which do not have native - // support for specialization constants? - const auto *RecordDecl = - KernelHandlerClone->getType()->getAsCXXRecordDecl(); - // FIXME: This call generates __init function bound to kernel object clone. - // Fix this. - createSpecialMethodCall(RecordDecl, InitSpecConstantsBuffer, BodyStmts); + // Generate + // KernelHandlerClone.__init_specialization_constants_buffer(specialization_constants_buffer) + // call if target does not have native support for specialization constants. + // Here, specialization_constants_buffer is the compiler generated kernel + // argument of type char*. + llvm::Triple T = SemaRef.Context.getTargetInfo().getTriple(); + if (!(T.isSPIR() && T.getSubArch() == llvm::Triple::NoSubArch)) + handleSpecialType(KernelHandlerArg->getType()); } bool enterStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { @@ -2971,10 +2993,20 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { } void handleSyclKernelHandlerType(QualType Ty) { - // Add corresponding entry in integration header. + // The compiler generated kernel argument used to initialize SYCL 2020 + // specialization constants, `specialization_constants_buffer`, should + // have corresponding entry in integration header. This argument is + // only generated when target has no native support for specialization + // constants. + ASTContext &Context = SemaRef.getASTContext(); + llvm::Triple T = Context.getTargetInfo().getTriple(); + if (T.isSPIR() && T.getSubArch() == llvm::Triple::NoSubArch) + return; + // Offset is zero since kernel_handler argument is not part of // kernel object (i.e. it is not captured) - addParam(Ty, SYCLIntegrationHeader::kind_specialization_constants_buffer, 0, + addParam(Context.getPointerType(Context.CharTy), + SYCLIntegrationHeader::kind_specialization_constants_buffer, 0, /*IsZeroOffset*/ true); } @@ -3310,8 +3342,8 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, Visitor.VisitRecordBases(KernelObj, kernel_decl, kernel_body, int_header); Visitor.VisitRecordFields(KernelObj, kernel_decl, kernel_body, int_header); - if (hasSyclKernelHandlerArg(KernelCallerFunc)) { - ParmVarDecl *KernelHandlerArg = KernelCallerFunc->getParamDecl(1); + ParmVarDecl *KernelHandlerArg = getSyclKernelHandlerArg(KernelCallerFunc); + if (KernelHandlerArg) { kernel_decl.handleSyclKernelHandlerType(); kernel_body.handleSyclKernelHandlerType(KernelHandlerArg); int_header.handleSyclKernelHandlerType(KernelHandlerArg->getType()); @@ -3563,6 +3595,7 @@ static const char *paramKind2Str(KernelParamKind K) { CASE(accessor); CASE(std_layout); CASE(sampler); + CASE(specialization_constants_buffer); CASE(pointer); } return ""; @@ -4153,7 +4186,7 @@ bool Util::isSyclKernelHandlerType(const QualType &Ty) { std::array Scopes = { Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"}, Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"}, - Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}}; + Util::DeclContextDesc{Decl::Kind::CXXRecord, Name}}; return matchQualifiedTypeName(Ty, Scopes); } diff --git a/clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp b/clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp new file mode 100644 index 0000000000000..c8627c15325a6 --- /dev/null +++ b/clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp @@ -0,0 +1,36 @@ +// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple nvptx64-unknown-unknown -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out +// RUN: FileCheck -input-file=%t.h %s --check-prefix=NONATIVESUPPORT --check-prefix=ALL +// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out +// RUN: FileCheck -input-file=%t.h %s --check-prefix=NATIVESUPPORT --check-prefix=ALL + +// This test checks that the compiler generates required information +// in integration header for kernel_handler type (SYCL 2020 specialization +// constants). + +//FIXME: Move to headers +namespace cl { +namespace sycl { +class kernel_handler { + void __init_specialization_constants_buffer(char *specialization_constants_buffer) {} +}; +} // namespace sycl +} // namespace cl + +template +__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc, cl::sycl::kernel_handler kh) { + kernelFunc(kh); +} + +int main() { + int a; + cl::sycl::kernel_handler kh; + + a_kernel( + [=](auto) { + int local = a; + }, + kh); +} +// ALL: const kernel_param_desc_t kernel_signatures[] = { +// NONATIVESUPPORT: { kernel_param_kind_t::kind_specialization_constants_buffer, 8, 0 } +// NATIVESUPPORT-NOT: { kernel_param_kind_t::kind_specialization_constants_buffer, 8, 0 } diff --git a/clang/test/SemaSYCL/kernel-handler.cpp b/clang/test/SemaSYCL/kernel-handler.cpp index 43c166a08ada2..bdd9bff65a1b3 100644 --- a/clang/test/SemaSYCL/kernel-handler.cpp +++ b/clang/test/SemaSYCL/kernel-handler.cpp @@ -1,4 +1,6 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple x86_64-pc-linux-gnu -ast-dump %s | FileCheck %s --check-prefix=NONATIVESUPPORT +// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple nvptx64-unknown-unknown -ast-dump %s | FileCheck %s --check-prefix=NONATIVESUPPORT +// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -ast-dump %s | FileCheck %s --check-prefix=NATIVESUPPORT // This test checks that the compiler handles kernel_handler type (for // SYCL 2020 specialization constants) correctly. @@ -29,23 +31,63 @@ int main() { } // Check test_kernel_handler parameters -// CHECK: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int, char *)' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used specialization_constants_buffer 'char *' +// NONATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int, char *)' +// NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' +// NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used specialization_constants_buffer 'char *' // Check declaration and initialization of kernel object local clone -// CHECK-NEXT: CompoundStmt -// CHECK-NEXT: DeclStmt -// CHECK-NEXT: VarDecl {{.*}} cinit -// CHECK-NEXT: InitListExpr -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' - -// Check declaration and initialization of kernel object local clone using default constructor -// CHECK-NEXT: DeclStmt -// CHECK-NEXT: VarDecl {{.*}} callinit -// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' 'void () noexcept' +// NONATIVESUPPORT-NEXT: CompoundStmt +// NONATIVESUPPORT-NEXT: DeclStmt +// NONATIVESUPPORT-NEXT: VarDecl {{.*}} cinit +// NONATIVESUPPORT-NEXT: InitListExpr +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' + +// Check declaration and initialization of kernel handler local clone using default constructor +// NONATIVESUPPORT-NEXT: DeclStmt +// NONATIVESUPPORT-NEXT: VarDecl {{.*}} callinit +// NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' 'void () noexcept' // Check call to __init_specialization_constants_buffer -// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (char *)' lvalue .__init_specialization_constants_buffer +// NONATIVESUPPORT-NEXT: CXXMemberCallExpr {{.*}} 'void' +// NONATIVESUPPORT-NEXT: MemberExpr {{.*}} 'void (char *)' lvalue .__init_specialization_constants_buffer +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' lvalue Var {{.*}} 'kh' +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'char *' +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'char *' lvalue ParmVar {{.*}} 'specialization_constants_buffer' 'char *' +// NONATIVESUPPORT-NEXT: CompoundStmt +// NONATIVESUPPORT-NEXT: CXXOperatorCallExpr +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'void (*)(cl::sycl::kernel_handler) const' +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'void (cl::sycl::kernel_handler) const' lvalue CXXMethod {{.*}} 'operator()' 'void (cl::sycl::kernel_handler) const' +// Kernel body with clones +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const (lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' +// NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' 'void (const cl::sycl::kernel_handler &) noexcept' +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::kernel_handler' lvalue +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' lvalue Var {{.*}} 'kh' 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' + +// Check test_kernel_handler parameters +// NATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int)' +// NATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' + +// Check declaration and initialization of kernel object local clone +// NATIVESUPPORT-NEXT: CompoundStmt +// NATIVESUPPORT-NEXT: DeclStmt +// NATIVESUPPORT-NEXT: VarDecl {{.*}} cinit +// NATIVESUPPORT-NEXT: InitListExpr +// NATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' +// NATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' + +// Check declaration and initialization of kernel handler local clone using default constructor +// NATIVESUPPORT-NEXT: DeclStmt +// NATIVESUPPORT-NEXT: VarDecl {{.*}} callinit +// NATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' 'void () noexcept' + +// Check no call to __init_specialization_constants_buffer +// NATIVESUPPORT-NOT: MemberExpr {{.*}} 'void (char *)' lvalue .__init_specialization_constants_buffer + +// Kernel body with clones +// NATIVESUPPORT: ImplicitCastExpr {{.*}} 'const (lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue +// NATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' +// NATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' 'void (const cl::sycl::kernel_handler &) noexcept' +// NATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::kernel_handler' lvalue +// NATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' lvalue Var {{.*}} 'kh' 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' From 83c0c65a784413da329bad4b69633c2e2792d5b9 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Thu, 18 Mar 2021 03:20:28 -0700 Subject: [PATCH 03/13] Review comment + fix issues after merge Signed-off-by: Elizabeth Andrews --- clang/include/clang/Sema/Sema.h | 4 ++-- clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp | 8 ++++---- clang/test/SemaSYCL/kernel-handler.cpp | 5 ++--- 3 files changed, 8 insertions(+), 9 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index bf33b1068fd9e..f61c83297c8ed 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -314,9 +314,9 @@ class SYCLIntegrationHeader { kind_accessor = kind_first, kind_std_layout, kind_sampler, - kind_specialization_constants_buffer, kind_pointer, - kind_last = kind_pointer + kind_specialization_constants_buffer, + kind_last = kind_specialization_constants_buffer }; public: diff --git a/clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp b/clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp index c8627c15325a6..78a6619e71747 100644 --- a/clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp +++ b/clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp @@ -1,10 +1,10 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple nvptx64-unknown-unknown -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out +// RUN: %clang_cc1 -fsycl-is-device -triple nvptx64-unknown-unknown -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out // RUN: FileCheck -input-file=%t.h %s --check-prefix=NONATIVESUPPORT --check-prefix=ALL -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out +// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out // RUN: FileCheck -input-file=%t.h %s --check-prefix=NATIVESUPPORT --check-prefix=ALL // This test checks that the compiler generates required information -// in integration header for kernel_handler type (SYCL 2020 specialization +// in integration header for kernel_handler type (SYCL 2020 specialization // constants). //FIXME: Move to headers @@ -32,5 +32,5 @@ int main() { kh); } // ALL: const kernel_param_desc_t kernel_signatures[] = { -// NONATIVESUPPORT: { kernel_param_kind_t::kind_specialization_constants_buffer, 8, 0 } +// NONATIVESUPPORT: { kernel_param_kind_t::kind_specialization_constants_buffer, 8, 0 } // NATIVESUPPORT-NOT: { kernel_param_kind_t::kind_specialization_constants_buffer, 8, 0 } diff --git a/clang/test/SemaSYCL/kernel-handler.cpp b/clang/test/SemaSYCL/kernel-handler.cpp index bdd9bff65a1b3..9d69720bffc6b 100644 --- a/clang/test/SemaSYCL/kernel-handler.cpp +++ b/clang/test/SemaSYCL/kernel-handler.cpp @@ -1,6 +1,5 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple x86_64-pc-linux-gnu -ast-dump %s | FileCheck %s --check-prefix=NONATIVESUPPORT -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple nvptx64-unknown-unknown -ast-dump %s | FileCheck %s --check-prefix=NONATIVESUPPORT -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -ast-dump %s | FileCheck %s --check-prefix=NATIVESUPPORT +// RUN: %clang_cc1 -fsycl-is-device -triple nvptx64-unknown-unknown -ast-dump %s | FileCheck %s --check-prefix=NONATIVESUPPORT +// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -ast-dump %s | FileCheck %s --check-prefix=NATIVESUPPORT // This test checks that the compiler handles kernel_handler type (for // SYCL 2020 specialization constants) correctly. From dbdb9060c45b694d34353b794d7150e2091443b5 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Thu, 18 Mar 2021 08:28:48 -0700 Subject: [PATCH 04/13] Add kernel_handler type to test headers and correct tests Signed-off-by: Elizabeth Andrews --- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 20 +++++++ .../int_header_sycl2020_spec_const.cpp | 35 +++++------- .../test/CodeGenSYCL/kernel-by-reference.cpp | 4 +- clang/test/SemaSYCL/Inputs/sycl.hpp | 18 ++++++ clang/test/SemaSYCL/args-size-overflow.cpp | 4 +- .../deferred-diagnostics-aux-builtin.cpp | 4 +- .../SemaSYCL/deferred-diagnostics-emit.cpp | 4 +- clang/test/SemaSYCL/float128.cpp | 6 +- clang/test/SemaSYCL/implicit_kernel_type.cpp | 16 +++--- clang/test/SemaSYCL/int128.cpp | 6 +- clang/test/SemaSYCL/kernel-handler.cpp | 57 ++++++++----------- clang/test/SemaSYCL/kernelname-enum.cpp | 8 +-- clang/test/SemaSYCL/stdtypes_kernel_type.cpp | 24 ++++---- clang/test/SemaSYCL/unnamed-kernel.cpp | 36 ++++++------ 14 files changed, 133 insertions(+), 109 deletions(-) diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 5462a2d13d1f6..c941a950f135d 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -291,12 +291,22 @@ class spec_constant { } // namespace experimental } // namespace ONEAPI +class kernel_handler { + void __init_specialization_constants_buffer(char *specialization_constants_buffer) {} +}; + #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) template ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } +#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) +template +ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc, kernel_handler kh) { + kernelFunc(kh); +} + template ATTR_SYCL_KERNEL void kernel_single_task_2017(KernelType kernelFunc) { kernelFunc(); @@ -347,6 +357,16 @@ class handler { #endif } + template + void single_task(const KernelType &kernelFunc, kernel_handler kh) { + using NameT = typename get_kernel_name_t::name; +#ifdef __SYCL_DEVICE_ONLY__ + kernel_single_task(kernelFunc, kh); +#else + kernelFunc(kh); +#endif + } + template void single_task_2017(KernelType kernelFunc) { using NameT = typename get_kernel_name_t::name; diff --git a/clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp b/clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp index 78a6619e71747..55733076e3efc 100644 --- a/clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp +++ b/clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp @@ -1,35 +1,28 @@ -// RUN: %clang_cc1 -fsycl-is-device -triple nvptx64-unknown-unknown -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx64-unknown-unknown -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out // RUN: FileCheck -input-file=%t.h %s --check-prefix=NONATIVESUPPORT --check-prefix=ALL -// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out // RUN: FileCheck -input-file=%t.h %s --check-prefix=NATIVESUPPORT --check-prefix=ALL // This test checks that the compiler generates required information // in integration header for kernel_handler type (SYCL 2020 specialization // constants). -//FIXME: Move to headers -namespace cl { -namespace sycl { -class kernel_handler { - void __init_specialization_constants_buffer(char *specialization_constants_buffer) {} -}; -} // namespace sycl -} // namespace cl +#include "sycl.hpp" -template -__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc, cl::sycl::kernel_handler kh) { - kernelFunc(kh); -} +using namespace cl::sycl; +queue q; int main() { - int a; - cl::sycl::kernel_handler kh; + q.submit([&](handler &h) { + int a; + kernel_handler kh; - a_kernel( - [=](auto) { - int local = a; - }, - kh); + h.single_task( + [=](auto) { + int local = a; + }, + kh); + }); } // ALL: const kernel_param_desc_t kernel_signatures[] = { // NONATIVESUPPORT: { kernel_param_kind_t::kind_specialization_constants_buffer, 8, 0 } diff --git a/clang/test/CodeGenSYCL/kernel-by-reference.cpp b/clang/test/CodeGenSYCL/kernel-by-reference.cpp index 056eea081cadf..33c057ca6be19 100644 --- a/clang/test/CodeGenSYCL/kernel-by-reference.cpp +++ b/clang/test/CodeGenSYCL/kernel-by-reference.cpp @@ -15,7 +15,7 @@ int simple_add(int i) { int main() { queue q; #if defined(SYCL2020) - // expected-warning@Inputs/sycl.hpp:301 {{Passing kernel functions by value is deprecated in SYCL 2020}} + // expected-warning@Inputs/sycl.hpp:311 {{Passing kernel functions by value is deprecated in SYCL 2020}} // expected-note@+3 {{in instantiation of function template specialization}} #endif q.submit([&](handler &h) { @@ -23,7 +23,7 @@ int main() { }); #if defined(SYCL2017) - // expected-warning@Inputs/sycl.hpp:296 {{Passing of kernel functions by reference is a SYCL 2020 extension}} + // expected-warning@Inputs/sycl.hpp:300 {{Passing of kernel functions by reference is a SYCL 2020 extension}} // expected-note@+3 {{in instantiation of function template specialization}} #endif q.submit([&](handler &h) { diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index a1a8d626b7c40..e0e7579d2a96d 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -207,6 +207,10 @@ struct get_kernel_name_t { using name = Type; }; +class kernel_handler { + void __init_specialization_constants_buffer(char *specialization_constants_buffer) {} +}; + // Used when parallel_for range is rounded-up. template class __pf_kernel_wrapper; @@ -220,6 +224,11 @@ template ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } +#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) +template +ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc, kernel_handler kh) { + kernelFunc(kh); +} template ATTR_SYCL_KERNEL void kernel_parallel_for(const KernelType &kernelFunc) { kernelFunc(); @@ -233,6 +242,15 @@ class handler { kernel_single_task(kernelFunc); #else kernelFunc(); +#endif + } + template + void single_task(const KernelType &kernelFunc, kernel_handler kh) { + using NameT = typename get_kernel_name_t::name; +#ifdef __SYCL_DEVICE_ONLY__ + kernel_single_task(kernelFunc, kh); +#else + kernelFunc(kh); #endif } template diff --git a/clang/test/SemaSYCL/args-size-overflow.cpp b/clang/test/SemaSYCL/args-size-overflow.cpp index fe28acde50d7e..5b4bd96a8ab87 100644 --- a/clang/test/SemaSYCL/args-size-overflow.cpp +++ b/clang/test/SemaSYCL/args-size-overflow.cpp @@ -11,9 +11,9 @@ queue q; using Accessor = accessor; #ifdef SPIR64 -// expected-warning@Inputs/sycl.hpp:233 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}} +// expected-warning@Inputs/sycl.hpp:242 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}} #elif SPIR32 -// expected-warning@Inputs/sycl.hpp:233 {{size of kernel arguments (7986 bytes) may exceed the supported maximum of 2048 bytes on some devices}} +// expected-warning@Inputs/sycl.hpp:242 {{size of kernel arguments (7986 bytes) may exceed the supported maximum of 2048 bytes on some devices}} #endif void use() { diff --git a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp index 6e0f8d8659f38..e056a2d136f05 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp @@ -12,7 +12,7 @@ int main(int argc, char **argv) { _mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}} deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task([]() { _mm_prefetch("test", 4); // expected-error {{builtin is not supported on this target}} _mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}} expected-error {{builtin is not supported on this target}} @@ -20,4 +20,4 @@ int main(int argc, char **argv) { }); return 0; -} \ No newline at end of file +} diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index 220dfde50e507..d5746953a28c1 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -64,7 +64,7 @@ template void setup_sycl_operation(const T VA[]) { deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task([]() { // ======= Zero Length Arrays Not Allowed in Kernel ========== // expected-error@+1 {{zero-length arrays are not permitted in C++}} @@ -156,7 +156,7 @@ int main(int argc, char **argv) { // --- direct lambda testing --- deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 2 {{called by 'kernel_single_task([]() { // expected-error@+1 {{zero-length arrays are not permitted in C++}} int BadArray[0]; diff --git a/clang/test/SemaSYCL/float128.cpp b/clang/test/SemaSYCL/float128.cpp index ff6b2f71b311d..ff26f814a2665 100644 --- a/clang/test/SemaSYCL/float128.cpp +++ b/clang/test/SemaSYCL/float128.cpp @@ -71,7 +71,7 @@ int main() { __float128 CapturedToDevice = 1; host_ok(); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task([=]() { // expected-error@+1 {{'__float128' is not supported on this target}} decltype(CapturedToDevice) D; @@ -88,7 +88,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 4{{called by 'kernel_single_task([=]() { // expected-note@+1 2{{called by 'operator()'}} usage(); @@ -104,7 +104,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task([=]() { // expected-note@+1 3{{used here}} Z<__float128> S; diff --git a/clang/test/SemaSYCL/implicit_kernel_type.cpp b/clang/test/SemaSYCL/implicit_kernel_type.cpp index d178a4c2d1256..a33624341f1d2 100644 --- a/clang/test/SemaSYCL/implicit_kernel_type.cpp +++ b/clang/test/SemaSYCL/implicit_kernel_type.cpp @@ -25,12 +25,12 @@ int main() { queue q; #if defined(WARN) - // expected-error@Inputs/sycl.hpp:233 {{'InvalidKernelName1' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:242 {{'InvalidKernelName1' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{'InvalidKernelName1' should be globally-visible}} // expected-note@+8 {{in instantiation of function template specialization}} #elif defined(ERROR) - // expected-error@Inputs/sycl.hpp:233 {{'InvalidKernelName1' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:242 {{'InvalidKernelName1' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{'InvalidKernelName1' should be globally-visible}} // expected-note@+4 {{in instantiation of function template specialization}} #endif class InvalidKernelName1 {}; @@ -39,9 +39,9 @@ int main() { }); #if defined(WARN) - // expected-warning@Inputs/sycl.hpp:233 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-warning@Inputs/sycl.hpp:242 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #elif defined(ERROR) - // expected-error@Inputs/sycl.hpp:233 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-error@Inputs/sycl.hpp:242 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #endif q.submit([&](handler &h) { @@ -53,9 +53,9 @@ int main() { }); #if defined(WARN) - // expected-warning@Inputs/sycl.hpp:233 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-warning@Inputs/sycl.hpp:242 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #elif defined(ERROR) - // expected-error@Inputs/sycl.hpp:233 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-error@Inputs/sycl.hpp:242 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #endif q.submit([&](handler &h) { diff --git a/clang/test/SemaSYCL/int128.cpp b/clang/test/SemaSYCL/int128.cpp index 196f2dfad66d4..4fd77fb61794d 100644 --- a/clang/test/SemaSYCL/int128.cpp +++ b/clang/test/SemaSYCL/int128.cpp @@ -80,7 +80,7 @@ int main() { __int128 CapturedToDevice = 1; host_ok(); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task([=]() { // expected-error@+1 {{'__int128' is not supported on this target}} decltype(CapturedToDevice) D; @@ -97,7 +97,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 5{{called by 'kernel_single_task([=]() { // expected-note@+1 2{{called by 'operator()'}} usage(); @@ -115,7 +115,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:221 {{called by 'kernel_single_task([=]() { // expected-note@+1 3{{used here}} Z<__int128> S; diff --git a/clang/test/SemaSYCL/kernel-handler.cpp b/clang/test/SemaSYCL/kernel-handler.cpp index 9d69720bffc6b..e9be5b2091f20 100644 --- a/clang/test/SemaSYCL/kernel-handler.cpp +++ b/clang/test/SemaSYCL/kernel-handler.cpp @@ -1,32 +1,25 @@ -// RUN: %clang_cc1 -fsycl-is-device -triple nvptx64-unknown-unknown -ast-dump %s | FileCheck %s --check-prefix=NONATIVESUPPORT -// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -ast-dump %s | FileCheck %s --check-prefix=NATIVESUPPORT +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx64-unknown-unknown -ast-dump %s | FileCheck %s --check-prefix=NONATIVESUPPORT +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -ast-dump %s | FileCheck %s --check-prefix=NATIVESUPPORT // This test checks that the compiler handles kernel_handler type (for // SYCL 2020 specialization constants) correctly. -//FIXME: Move to headers -namespace cl { -namespace sycl { -class kernel_handler { - void __init_specialization_constants_buffer(char *specialization_constants_buffer) {} -}; -} // namespace sycl -} // namespace cl +#include "sycl.hpp" -template -__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc, cl::sycl::kernel_handler kh) { - kernelFunc(kh); -} +using namespace cl::sycl; +queue q; int main() { - int a; - cl::sycl::kernel_handler kh; + q.submit([&](handler &h) { + int a; + kernel_handler kh; - a_kernel( - [=](auto) { - int local = a; - }, - kh); + h.single_task( + [=](auto) { + int local = a; + }, + kh); + }); } // Check test_kernel_handler parameters @@ -45,24 +38,24 @@ int main() { // Check declaration and initialization of kernel handler local clone using default constructor // NONATIVESUPPORT-NEXT: DeclStmt // NONATIVESUPPORT-NEXT: VarDecl {{.*}} callinit -// NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' 'void () noexcept' +// NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'sycl::kernel_handler' 'void () noexcept' // Check call to __init_specialization_constants_buffer // NONATIVESUPPORT-NEXT: CXXMemberCallExpr {{.*}} 'void' // NONATIVESUPPORT-NEXT: MemberExpr {{.*}} 'void (char *)' lvalue .__init_specialization_constants_buffer -// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' lvalue Var {{.*}} 'kh' +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh' // NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'char *' // NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'char *' lvalue ParmVar {{.*}} 'specialization_constants_buffer' 'char *' // NONATIVESUPPORT-NEXT: CompoundStmt // NONATIVESUPPORT-NEXT: CXXOperatorCallExpr -// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'void (*)(cl::sycl::kernel_handler) const' -// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'void (cl::sycl::kernel_handler) const' lvalue CXXMethod {{.*}} 'operator()' 'void (cl::sycl::kernel_handler) const' +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'void (*)(sycl::kernel_handler) const' +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'void (sycl::kernel_handler) const' lvalue CXXMethod {{.*}} 'operator()' 'void (sycl::kernel_handler) const' // Kernel body with clones // NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const (lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue // NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' -// NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' 'void (const cl::sycl::kernel_handler &) noexcept' -// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::kernel_handler' lvalue -// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' lvalue Var {{.*}} 'kh' 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' +// NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'sycl::kernel_handler':'sycl::kernel_handler' 'void (const sycl::kernel_handler &) noexcept' +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const sycl::kernel_handler' lvalue +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh' 'sycl::kernel_handler' // Check test_kernel_handler parameters // NATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int)' @@ -79,7 +72,7 @@ int main() { // Check declaration and initialization of kernel handler local clone using default constructor // NATIVESUPPORT-NEXT: DeclStmt // NATIVESUPPORT-NEXT: VarDecl {{.*}} callinit -// NATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' 'void () noexcept' +// NATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'sycl::kernel_handler' 'void () noexcept' // Check no call to __init_specialization_constants_buffer // NATIVESUPPORT-NOT: MemberExpr {{.*}} 'void (char *)' lvalue .__init_specialization_constants_buffer @@ -87,6 +80,6 @@ int main() { // Kernel body with clones // NATIVESUPPORT: ImplicitCastExpr {{.*}} 'const (lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue // NATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' -// NATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' 'void (const cl::sycl::kernel_handler &) noexcept' -// NATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::kernel_handler' lvalue -// NATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' lvalue Var {{.*}} 'kh' 'cl::sycl::kernel_handler':'cl::sycl::kernel_handler' +// NATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'sycl::kernel_handler':'sycl::kernel_handler' 'void (const sycl::kernel_handler &) noexcept' +// NATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const sycl::kernel_handler' lvalue +// NATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh' 'sycl::kernel_handler' diff --git a/clang/test/SemaSYCL/kernelname-enum.cpp b/clang/test/SemaSYCL/kernelname-enum.cpp index 62e54d61f51c8..6bc25873dcb6b 100644 --- a/clang/test/SemaSYCL/kernelname-enum.cpp +++ b/clang/test/SemaSYCL/kernelname-enum.cpp @@ -67,15 +67,15 @@ int main() { }); q.submit([&](cl::sycl::handler &cgh) { - // expected-error@Inputs/sycl.hpp:233 {{'dummy_functor_2' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} + // expected-error@Inputs/sycl.hpp:242 {{'dummy_functor_2' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task(f2); }); q.submit([&](cl::sycl::handler &cgh) { - // expected-error@Inputs/sycl.hpp:233 {{'templated_functor' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} + // expected-error@Inputs/sycl.hpp:242 {{'templated_functor' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task(f5); }); diff --git a/clang/test/SemaSYCL/stdtypes_kernel_type.cpp b/clang/test/SemaSYCL/stdtypes_kernel_type.cpp index 1209c4083ace0..fafa89dd176cb 100644 --- a/clang/test/SemaSYCL/stdtypes_kernel_type.cpp +++ b/clang/test/SemaSYCL/stdtypes_kernel_type.cpp @@ -25,38 +25,38 @@ queue q; int main() { #ifdef CHECK_ERROR q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:233 {{'nullptr_t' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{type 'nullptr_t' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:242 {{'nullptr_t' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{type 'nullptr_t' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task([=] {}); }); q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:233 {{'std::T' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{type 'std::T' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:242 {{'std::T' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{type 'std::T' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task([=] {}); }); q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:233 {{'Templated_kernel_name' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{type 'nullptr_t' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:242 {{'Templated_kernel_name' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{type 'nullptr_t' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task>([=] {}); }); q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:233 {{'Templated_kernel_name' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{type 'std::U' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:242 {{'Templated_kernel_name' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{type 'std::U' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task>([=] {}); }); q.submit([&](handler &cgh) { - // expected-error@Inputs/sycl.hpp:233 {{'Templated_kernel_name2>' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233{{type 'std::Foo' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:242 {{'Templated_kernel_name2>' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242{{type 'std::Foo' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task>>([]() {}); }); q.submit([&](handler &cgh) { - // expected-error@Inputs/sycl.hpp:233 {{'TemplParamPack' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{type 'nullptr_t' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:242 {{'TemplParamPack' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{type 'nullptr_t' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task>([]() {}); }); diff --git a/clang/test/SemaSYCL/unnamed-kernel.cpp b/clang/test/SemaSYCL/unnamed-kernel.cpp index bb3dc914ab345..c499ede958730 100644 --- a/clang/test/SemaSYCL/unnamed-kernel.cpp +++ b/clang/test/SemaSYCL/unnamed-kernel.cpp @@ -33,8 +33,8 @@ struct MyWrapper { void test() { cl::sycl::queue q; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'InvalidKernelName1' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:242 {{'InvalidKernelName1' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{'InvalidKernelName1' should be globally-visible}} // expected-note@+4{{in instantiation of function template specialization}} #endif class InvalidKernelName1 {}; @@ -43,8 +43,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'namespace1::KernelName' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName2' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:242 {{'namespace1::KernelName' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{'InvalidKernelName2' should be globally-visible}} // expected-note@+4{{in instantiation of function template specialization}} #endif class InvalidKernelName2 {}; @@ -53,8 +53,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName0' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName0' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:242 {{'MyWrapper::InvalidKernelName0' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{'MyWrapper::InvalidKernelName0' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -62,8 +62,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'namespace1::KernelName' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName3' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:242 {{'namespace1::KernelName' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{'MyWrapper::InvalidKernelName3' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -76,8 +76,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'std::max_align_t' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{type 'std::max_align_t' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:242 {{'std::max_align_t' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{type 'std::max_align_t' cannot be in the "std" namespace}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -86,8 +86,8 @@ struct MyWrapper { using InvalidAlias = InvalidKernelName4; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName4' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName4' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:242 {{'MyWrapper::InvalidKernelName4' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{'MyWrapper::InvalidKernelName4' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -96,16 +96,16 @@ struct MyWrapper { using InvalidAlias1 = InvalidKernelName5; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'namespace1::KernelName' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'MyWrapper::InvalidKernelName5' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:242 {{'namespace1::KernelName' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{'MyWrapper::InvalidKernelName5' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { h.single_task>([] {}); }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:233 {{'Templated_kernel_name2>' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:242 {{'Templated_kernel_name2>' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{'InvalidKernelName1' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -117,8 +117,8 @@ struct MyWrapper { int main() { cl::sycl::queue q; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error-re@Inputs/sycl.hpp:233 {{'(lambda at {{.*}}unnamed-kernel.cpp{{.*}}' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:233 {{unnamed type used in a SYCL kernel name}} + // expected-error-re@Inputs/sycl.hpp:242 {{'(lambda at {{.*}}unnamed-kernel.cpp{{.*}}' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:242 {{unnamed type used in a SYCL kernel name}} // expected-note@+2{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); }); From 7fdf3682e0704b24d568188fb2a64c73080d8071 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Thu, 18 Mar 2021 09:46:55 -0700 Subject: [PATCH 05/13] Append _arg__ to kernel argument to be consistent + Minor refactor Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 20 ++++++-------------- clang/test/SemaSYCL/kernel-handler.cpp | 4 ++-- 2 files changed, 8 insertions(+), 16 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 3961263925428..f66fac60569db 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -738,18 +738,6 @@ static ParamDesc makeParamDesc(const FieldDecl *Src, QualType Ty) { Ctx.getTrivialTypeSourceInfo(Ty)); } -static ParamDesc makeParamDesc(ASTContext &Ctx, const CXXBaseSpecifier &Src, - QualType Ty) { - // TODO: There is no name for the base available, but duplicate names are - // seemingly already possible, so we'll give them all the same name for now. - // This only happens with the accessor types. - std::string Name = "_arg__base"; - return std::make_tuple(Ty, &Ctx.Idents.get(Name), - Ctx.getTrivialTypeSourceInfo(Ty)); -} - -// FIXME: Should we refactor makeParamDesc to just accept Name in all cases -// i.e. remove overloads. static ParamDesc makeParamDesc(ASTContext &Ctx, StringRef Name, QualType Ty) { return std::make_tuple(Ty, &Ctx.Idents.get(Name), Ctx.getTrivialTypeSourceInfo(Ty)); @@ -1667,8 +1655,12 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { } void addParam(const CXXBaseSpecifier &BS, QualType FieldTy) { + // TODO: There is no name for the base available, but duplicate names are + // seemingly already possible, so we'll give them all the same name for now. + // This only happens with the accessor types. + StringRef Name = "_arg__base"; ParamDesc newParamDesc = - makeParamDesc(SemaRef.getASTContext(), BS, FieldTy); + makeParamDesc(SemaRef.getASTContext(), Name, FieldTy); addParam(newParamDesc, FieldTy); } // Add a parameter with specified name and type @@ -1987,7 +1979,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { if (T.isSPIR() && T.getSubArch() == llvm::Triple::NoSubArch) return; - StringRef Name = "specialization_constants_buffer"; + StringRef Name = "_arg__specialization_constants_buffer"; addParam(Name, Context.getPointerType(Context.CharTy)); } diff --git a/clang/test/SemaSYCL/kernel-handler.cpp b/clang/test/SemaSYCL/kernel-handler.cpp index e9be5b2091f20..b10c52557b559 100644 --- a/clang/test/SemaSYCL/kernel-handler.cpp +++ b/clang/test/SemaSYCL/kernel-handler.cpp @@ -25,7 +25,7 @@ int main() { // Check test_kernel_handler parameters // NONATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int, char *)' // NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' -// NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used specialization_constants_buffer 'char *' +// NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg__specialization_constants_buffer 'char *' // Check declaration and initialization of kernel object local clone // NONATIVESUPPORT-NEXT: CompoundStmt @@ -45,7 +45,7 @@ int main() { // NONATIVESUPPORT-NEXT: MemberExpr {{.*}} 'void (char *)' lvalue .__init_specialization_constants_buffer // NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh' // NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'char *' -// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'char *' lvalue ParmVar {{.*}} 'specialization_constants_buffer' 'char *' +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'char *' lvalue ParmVar {{.*}} '_arg__specialization_constants_buffer' 'char *' // NONATIVESUPPORT-NEXT: CompoundStmt // NONATIVESUPPORT-NEXT: CXXOperatorCallExpr // NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'void (*)(sycl::kernel_handler) const' From c1f7cdfe55c2d4a210f792998cc977d5c7d60d89 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Thu, 18 Mar 2021 13:01:01 -0700 Subject: [PATCH 06/13] Implement review comments Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 29 +++++++++++++++-------------- 1 file changed, 15 insertions(+), 14 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index f66fac60569db..a3dd4e889db67 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -778,16 +778,21 @@ constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, KernelNameType)}; } +static bool hasNativeSycl2020SpecConstantSupport(ASTContext &Context) { + llvm::Triple T = Context.getTargetInfo().getTriple(); + if (T.isSPIR() && T.getSubArch() == llvm::Triple::NoSubArch) + return true; + return false; +} + static ParmVarDecl *getSyclKernelHandlerArg(FunctionDecl *KernelCallerFunc) { // Specialization constants in SYCL 2020 are not captured by lambda and // accessed through new optional lambda argument kernel_handler - ParmVarDecl *PVD; auto It = std::find_if(KernelCallerFunc->param_begin(), KernelCallerFunc->param_end(), [](ParmVarDecl *PVD) { return Util::isSyclKernelHandlerType(PVD->getType()); }); - PVD = (It != KernelCallerFunc->param_end()) ? *It : nullptr; - return PVD; + return ((It != KernelCallerFunc->param_end()) ? *It : nullptr); } // anonymous namespace so these don't get linkage. @@ -1973,10 +1978,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // argument is only generated when the target has no native support for // specialization constants void handleSyclKernelHandlerType() { - ASTContext &Context = SemaRef.getASTContext(); - llvm::Triple T = Context.getTargetInfo().getTriple(); - if (T.isSPIR() && T.getSubArch() == llvm::Triple::NoSubArch) + if (hasNativeSycl2020SpecConstantSupport(Context)) return; StringRef Name = "_arg__specialization_constants_buffer"; @@ -2162,8 +2165,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // If kernel_handler argument is passed by SYCL kernel, replace references // to this argument in kernel body, to use the compiler generated local // clone - ParmVarDecl *KernelHandlerParam = getSyclKernelHandlerArg(KernelCallerFunc); - if (KernelHandlerParam) + if (ParmVarDecl *KernelHandlerParam = + getSyclKernelHandlerArg(KernelCallerFunc)) NewBody = replaceWithLocalClone(KernelHandlerParam, KernelHandlerClone, NewBody); @@ -2622,8 +2625,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // call if target does not have native support for specialization constants. // Here, specialization_constants_buffer is the compiler generated kernel // argument of type char*. - llvm::Triple T = SemaRef.Context.getTargetInfo().getTriple(); - if (!(T.isSPIR() && T.getSubArch() == llvm::Triple::NoSubArch)) + if (!hasNativeSycl2020SpecConstantSupport(SemaRef.Context)) handleSpecialType(KernelHandlerArg->getType()); } @@ -2990,8 +2992,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { // only generated when target has no native support for specialization // constants. ASTContext &Context = SemaRef.getASTContext(); - llvm::Triple T = Context.getTargetInfo().getTriple(); - if (T.isSPIR() && T.getSubArch() == llvm::Triple::NoSubArch) + if (hasNativeSycl2020SpecConstantSupport(Context)) return; // Offset is zero since kernel_handler argument is not part of @@ -3389,8 +3390,8 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, Visitor.VisitRecordBases(KernelObj, kernel_decl, kernel_body, int_header); Visitor.VisitRecordFields(KernelObj, kernel_decl, kernel_body, int_header); - ParmVarDecl *KernelHandlerArg = getSyclKernelHandlerArg(KernelCallerFunc); - if (KernelHandlerArg) { + if (ParmVarDecl *KernelHandlerArg = + getSyclKernelHandlerArg(KernelCallerFunc)) { kernel_decl.handleSyclKernelHandlerType(); kernel_body.handleSyclKernelHandlerType(KernelHandlerArg); int_header.handleSyclKernelHandlerType(KernelHandlerArg->getType()); From ca9f7288f00bd3506d43ba67a45e3dfefd8b9893 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Thu, 18 Mar 2021 13:58:31 -0700 Subject: [PATCH 07/13] IR test Signed-off-by: Elizabeth Andrews --- clang/test/CodeGenSYCL/kernel-handler.cpp | 38 +++++++++++++++++++++++ 1 file changed, 38 insertions(+) create mode 100644 clang/test/CodeGenSYCL/kernel-handler.cpp diff --git a/clang/test/CodeGenSYCL/kernel-handler.cpp b/clang/test/CodeGenSYCL/kernel-handler.cpp new file mode 100644 index 0000000000000..50eea67ddabeb --- /dev/null +++ b/clang/test/CodeGenSYCL/kernel-handler.cpp @@ -0,0 +1,38 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefix=NONATIVESUPPORT +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefix=NATIVESUPPORT + +// This test checks IR generated when kernel_handler argument +// (used to handle SYCL 2020 specialization constants) is passed +// by kernel + +#include "sycl.hpp" + +using namespace cl::sycl; + +void test(int val) { + queue q; + q.submit([&](handler &h) { + int a; + kernel_handler kh; + h.single_task( + [=](auto) { + int local = a; + }, + kh); + }); +} + +// NONATIVESUPPORT: define dso_local void @"_ZTSZZ4testiENK3$_0clERN2cl4sycl7handlerEE19test_kernel_handler" +// NONATIVESUPPORT-SAME: (i32 %_arg_, i8* %_arg__specialization_constants_buffer) +// NONATIVESUPPORT: %kh = alloca %"class._ZTSN2cl4sycl14kernel_handlerE.cl::sycl::kernel_handler", align 1 +// NONATIVESUPPORT: %[[KH:[0-9]+]] = load i8*, i8** %_arg__specialization_constants_buffer.addr, align 8 +// NONATIVESUPPORT: call void @_ZN2cl4sycl14kernel_handler38__init_specialization_constants_bufferEPc(%"class._ZTSN2cl4sycl14kernel_handlerE.cl::sycl::kernel_handler"* nonnull dereferenceable(1) %kh, i8* %[[KH]]) +// NONATIVESUPPORT: void @"_ZZZ4testiENK3$_0clERN2cl4sycl7handlerEENKUlT_E_clINS1_14kernel_handlerEEEDaS4_" +// NONATIVESUPPORT-SAME: byval(%"class._ZTSN2cl4sycl14kernel_handlerE.cl::sycl::kernel_handler") + +// NATIVESUPPORT: define dso_local spir_kernel void @"_ZTSZZ4testiENK3$_0clERN2cl4sycl7handlerEE19test_kernel_handler" +// NATIVESUPPORT-SAME: (i32 %_arg_) +// NATIVESUPPORT: %kh = alloca %"class._ZTSN2cl4sycl14kernel_handlerE.cl::sycl::kernel_handler" +// NATIVESUPPORT-NOT: __init_specialization_constants_buffer +// NATIVE-SUPPORT: call spir_func void @"_ZZZ4testiENK3$_0clERN2cl4sycl7handlerEENKUlT_E_clINS1_14kernel_handlerEEEDaS4_" +// NATIVE-SUPPORT-SAME: byval(%"class._ZTSN2cl4sycl14kernel_handlerE.cl::sycl::kernel_handler") From d1daf583fc2bdbdc1e13261bf4d16a0a3758f8ac Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Fri, 19 Mar 2021 12:47:24 -0700 Subject: [PATCH 08/13] Fix a crash on parallel_for_work_group kernel when kernel_handler is passed Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index a3dd4e889db67..9271c6b48cd51 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -622,11 +622,15 @@ class FindPFWGLambdaFnVisitor auto *M = dyn_cast(Call->getDirectCallee()); if (!M || (M->getOverloadedOperator() != OO_Call)) return true; - const int NumPFWGLambdaArgs = 2; // group and lambda obj + + unsigned int NumPFWGLambdaArgs = + M->getNumParams() + 1; // group, optional kernel_handler and lambda obj if (Call->getNumArgs() != NumPFWGLambdaArgs) return true; if (!Util::isSyclType(Call->getArg(1)->getType(), "group", true /*Tmpl*/)) return true; + if (!Util::isSyclKernelHandlerType(Call->getArg(2)->getType())) + return true; if (Call->getArg(0)->getType()->getAsCXXRecordDecl() != LambdaObjTy) return true; LambdaFn = M; // call to PFWG lambda found - record the lambda From 4ef981b673e1615bd602843981339d96be17e0ba Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Mon, 22 Mar 2021 03:30:18 -0700 Subject: [PATCH 09/13] Fix crash Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 9271c6b48cd51..437dc6273db66 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -629,7 +629,8 @@ class FindPFWGLambdaFnVisitor return true; if (!Util::isSyclType(Call->getArg(1)->getType(), "group", true /*Tmpl*/)) return true; - if (!Util::isSyclKernelHandlerType(Call->getArg(2)->getType())) + if ((Call->getNumArgs() > 2) && + !Util::isSyclKernelHandlerType(Call->getArg(2)->getType())) return true; if (Call->getArg(0)->getType()->getAsCXXRecordDecl() != LambdaObjTy) return true; From 6a436895bb0ec9cd7bf1e0a28527062a5c44bf7f Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Tue, 23 Mar 2021 13:41:46 -0700 Subject: [PATCH 10/13] Implement review comments, add test for PFWG and fix tests after header change Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 42 +++++++++++------ clang/test/SemaSYCL/Inputs/sycl.hpp | 22 ++++++++- clang/test/SemaSYCL/args-size-overflow.cpp | 4 +- .../deferred-diagnostics-aux-builtin.cpp | 2 +- .../SemaSYCL/deferred-diagnostics-emit.cpp | 4 +- clang/test/SemaSYCL/float128.cpp | 6 +-- clang/test/SemaSYCL/implicit_kernel_type.cpp | 16 +++---- clang/test/SemaSYCL/int128.cpp | 6 +-- clang/test/SemaSYCL/kernel-handler.cpp | 47 +++++++++++++++++++ clang/test/SemaSYCL/kernelname-enum.cpp | 8 ++-- clang/test/SemaSYCL/stdtypes_kernel_type.cpp | 24 +++++----- clang/test/SemaSYCL/unnamed-kernel.cpp | 36 +++++++------- 12 files changed, 148 insertions(+), 69 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 437dc6273db66..4b25015af7bbf 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -783,7 +783,7 @@ constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, KernelNameType)}; } -static bool hasNativeSycl2020SpecConstantSupport(ASTContext &Context) { +static bool isDefaultSPIRArch(ASTContext &Context) { llvm::Triple T = Context.getTargetInfo().getTriple(); if (T.isSPIR() && T.getSubArch() == llvm::Triple::NoSubArch) return true; @@ -793,11 +793,26 @@ static bool hasNativeSycl2020SpecConstantSupport(ASTContext &Context) { static ParmVarDecl *getSyclKernelHandlerArg(FunctionDecl *KernelCallerFunc) { // Specialization constants in SYCL 2020 are not captured by lambda and // accessed through new optional lambda argument kernel_handler - auto It = std::find_if(KernelCallerFunc->param_begin(), - KernelCallerFunc->param_end(), [](ParmVarDecl *PVD) { - return Util::isSyclKernelHandlerType(PVD->getType()); - }); - return ((It != KernelCallerFunc->param_end()) ? *It : nullptr); + auto KHArg = + std::find_if(KernelCallerFunc->param_begin(), + KernelCallerFunc->param_end(), [](ParmVarDecl *PVD) { + return Util::isSyclKernelHandlerType(PVD->getType()); + }); + + ParmVarDecl *KernelHandlerArg = + (KHArg != KernelCallerFunc->param_end()) ? *KHArg : nullptr; + + if (KernelHandlerArg) { + auto KHArgTooMany = std::find_if( + std::next(KHArg), KernelCallerFunc->param_end(), [](ParmVarDecl *PVD) { + return Util::isSyclKernelHandlerType(PVD->getType()); + }); + + assert(KHArgTooMany == KernelCallerFunc->param_end() && + "Too many kernel_handler arguments"); + } + + return KernelHandlerArg; } // anonymous namespace so these don't get linkage. @@ -1984,7 +1999,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // specialization constants void handleSyclKernelHandlerType() { ASTContext &Context = SemaRef.getASTContext(); - if (hasNativeSycl2020SpecConstantSupport(Context)) + if (isDefaultSPIRArch(Context)) return; StringRef Name = "_arg__specialization_constants_buffer"; @@ -2630,7 +2645,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // call if target does not have native support for specialization constants. // Here, specialization_constants_buffer is the compiler generated kernel // argument of type char*. - if (!hasNativeSycl2020SpecConstantSupport(SemaRef.Context)) + if (!isDefaultSPIRArch(SemaRef.Context)) handleSpecialType(KernelHandlerArg->getType()); } @@ -2787,13 +2802,11 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { addParam(ArgTy, Kind, offsetOf(FD, ArgTy)); } void addParam(QualType ArgTy, SYCLIntegrationHeader::kernel_param_kind_t Kind, - uint64_t OffsetAdj, bool IsZeroOffset = false) { + uint64_t OffsetAdj) { uint64_t Size; Size = SemaRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); Header.addParamDesc(Kind, static_cast(Size), - ((IsZeroOffset) - ? static_cast(OffsetAdj) - : static_cast(CurOffset + OffsetAdj))); + static_cast(CurOffset + OffsetAdj)); } // Returns 'true' if the thing we're visiting (Based on the FD/QualType pair) @@ -2997,14 +3010,13 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { // only generated when target has no native support for specialization // constants. ASTContext &Context = SemaRef.getASTContext(); - if (hasNativeSycl2020SpecConstantSupport(Context)) + if (isDefaultSPIRArch(Context)) return; // Offset is zero since kernel_handler argument is not part of // kernel object (i.e. it is not captured) addParam(Context.getPointerType(Context.CharTy), - SYCLIntegrationHeader::kind_specialization_constants_buffer, 0, - /*IsZeroOffset*/ true); + SYCLIntegrationHeader::kind_specialization_constants_buffer, 0); } bool enterStream(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index e0e7579d2a96d..fa6ce419662b5 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -207,6 +207,12 @@ struct get_kernel_name_t { using name = Type; }; +template +class group { +public: + group() = default; // fake constructor +}; + class kernel_handler { void __init_specialization_constants_buffer(char *specialization_constants_buffer) {} }; @@ -224,7 +230,6 @@ template ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } -#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) template ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc, kernel_handler kh) { kernelFunc(kh); @@ -233,6 +238,11 @@ template ATTR_SYCL_KERNEL void kernel_parallel_for(const KernelType &kernelFunc) { kernelFunc(); } +template +ATTR_SYCL_KERNEL void kernel_parallel_for_work_group(const KernelType &KernelFunc, kernel_handler kh) { + KernelFunc(group<1>(), kh); +} + class handler { public: template @@ -261,6 +271,16 @@ class handler { kernel_parallel_for(kernelObj); #else kernelObj(); +#endif + } + template + void parallel_for_work_group(const KernelType &kernelFunc, kernel_handler kh) { + using NameT = typename get_kernel_name_t::name; +#ifdef __SYCL_DEVICE_ONLY__ + kernel_parallel_for_work_group(kernelFunc, kh); +#else + group<1> G; + kernelFunc(G, kh); #endif } }; diff --git a/clang/test/SemaSYCL/args-size-overflow.cpp b/clang/test/SemaSYCL/args-size-overflow.cpp index 5b4bd96a8ab87..827fda68d6d35 100644 --- a/clang/test/SemaSYCL/args-size-overflow.cpp +++ b/clang/test/SemaSYCL/args-size-overflow.cpp @@ -11,9 +11,9 @@ queue q; using Accessor = accessor; #ifdef SPIR64 -// expected-warning@Inputs/sycl.hpp:242 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}} +// expected-warning@Inputs/sycl.hpp:252 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}} #elif SPIR32 -// expected-warning@Inputs/sycl.hpp:242 {{size of kernel arguments (7986 bytes) may exceed the supported maximum of 2048 bytes on some devices}} +// expected-warning@Inputs/sycl.hpp:252 {{size of kernel arguments (7986 bytes) may exceed the supported maximum of 2048 bytes on some devices}} #endif void use() { diff --git a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp index e056a2d136f05..e16cc96182c55 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp @@ -12,7 +12,7 @@ int main(int argc, char **argv) { _mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}} deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:225 {{called by 'kernel_single_task([]() { _mm_prefetch("test", 4); // expected-error {{builtin is not supported on this target}} _mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}} expected-error {{builtin is not supported on this target}} diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index d5746953a28c1..f69eb508fadcf 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -64,7 +64,7 @@ template void setup_sycl_operation(const T VA[]) { deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:225 {{called by 'kernel_single_task([]() { // ======= Zero Length Arrays Not Allowed in Kernel ========== // expected-error@+1 {{zero-length arrays are not permitted in C++}} @@ -156,7 +156,7 @@ int main(int argc, char **argv) { // --- direct lambda testing --- deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:225 2 {{called by 'kernel_single_task([]() { // expected-error@+1 {{zero-length arrays are not permitted in C++}} int BadArray[0]; diff --git a/clang/test/SemaSYCL/float128.cpp b/clang/test/SemaSYCL/float128.cpp index ff26f814a2665..4444a6c30548c 100644 --- a/clang/test/SemaSYCL/float128.cpp +++ b/clang/test/SemaSYCL/float128.cpp @@ -71,7 +71,7 @@ int main() { __float128 CapturedToDevice = 1; host_ok(); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:225 {{called by 'kernel_single_task([=]() { // expected-error@+1 {{'__float128' is not supported on this target}} decltype(CapturedToDevice) D; @@ -88,7 +88,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:225 4{{called by 'kernel_single_task([=]() { // expected-note@+1 2{{called by 'operator()'}} usage(); @@ -104,7 +104,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:225 {{called by 'kernel_single_task([=]() { // expected-note@+1 3{{used here}} Z<__float128> S; diff --git a/clang/test/SemaSYCL/implicit_kernel_type.cpp b/clang/test/SemaSYCL/implicit_kernel_type.cpp index a33624341f1d2..a7fdccb67cab1 100644 --- a/clang/test/SemaSYCL/implicit_kernel_type.cpp +++ b/clang/test/SemaSYCL/implicit_kernel_type.cpp @@ -25,12 +25,12 @@ int main() { queue q; #if defined(WARN) - // expected-error@Inputs/sycl.hpp:242 {{'InvalidKernelName1' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:252 {{'InvalidKernelName1' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{'InvalidKernelName1' should be globally-visible}} // expected-note@+8 {{in instantiation of function template specialization}} #elif defined(ERROR) - // expected-error@Inputs/sycl.hpp:242 {{'InvalidKernelName1' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:252 {{'InvalidKernelName1' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{'InvalidKernelName1' should be globally-visible}} // expected-note@+4 {{in instantiation of function template specialization}} #endif class InvalidKernelName1 {}; @@ -39,9 +39,9 @@ int main() { }); #if defined(WARN) - // expected-warning@Inputs/sycl.hpp:242 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-warning@Inputs/sycl.hpp:252 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #elif defined(ERROR) - // expected-error@Inputs/sycl.hpp:242 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-error@Inputs/sycl.hpp:252 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #endif q.submit([&](handler &h) { @@ -53,9 +53,9 @@ int main() { }); #if defined(WARN) - // expected-warning@Inputs/sycl.hpp:242 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-warning@Inputs/sycl.hpp:252 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #elif defined(ERROR) - // expected-error@Inputs/sycl.hpp:242 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-error@Inputs/sycl.hpp:252 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #endif q.submit([&](handler &h) { diff --git a/clang/test/SemaSYCL/int128.cpp b/clang/test/SemaSYCL/int128.cpp index 4fd77fb61794d..76e4668a3c502 100644 --- a/clang/test/SemaSYCL/int128.cpp +++ b/clang/test/SemaSYCL/int128.cpp @@ -80,7 +80,7 @@ int main() { __int128 CapturedToDevice = 1; host_ok(); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:225 {{called by 'kernel_single_task([=]() { // expected-error@+1 {{'__int128' is not supported on this target}} decltype(CapturedToDevice) D; @@ -97,7 +97,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:225 5{{called by 'kernel_single_task([=]() { // expected-note@+1 2{{called by 'operator()'}} usage(); @@ -115,7 +115,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:225 {{called by 'kernel_single_task([=]() { // expected-note@+1 3{{used here}} Z<__int128> S; diff --git a/clang/test/SemaSYCL/kernel-handler.cpp b/clang/test/SemaSYCL/kernel-handler.cpp index b10c52557b559..a5df0e186e10b 100644 --- a/clang/test/SemaSYCL/kernel-handler.cpp +++ b/clang/test/SemaSYCL/kernel-handler.cpp @@ -19,6 +19,11 @@ int main() { int local = a; }, kh); + h.parallel_for_work_group( + [=](group<1> G, kernel_handler kh) { + int local = a; + }, + kh); }); } @@ -57,6 +62,48 @@ int main() { // NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const sycl::kernel_handler' lvalue // NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh' 'sycl::kernel_handler' +// Check test_pfwg_kernel_handler parameters +// NONATIVESUPPORT: FunctionDecl {{.*}}test_pfwg_kernel_handler{{.*}} 'void (int, char *)' +// NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' +// NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg__specialization_constants_buffer 'char *' + +// Check declaration and initialization of kernel object local clone +// NONATIVESUPPORT-NEXT: CompoundStmt +// NONATIVESUPPORT-NEXT: DeclStmt +// NONATIVESUPPORT-NEXT: VarDecl {{.*}} cinit +// NONATIVESUPPORT-NEXT: InitListExpr +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'int' +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' + +// Check declaration and initialization of kernel handler local clone using default constructor +// NONATIVESUPPORT-NEXT: DeclStmt +// NONATIVESUPPORT-NEXT: VarDecl {{.*}} callinit +// NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'sycl::kernel_handler' 'void () noexcept' + +// Check call to __init_specialization_constants_buffer +// NONATIVESUPPORT-NEXT: CXXMemberCallExpr {{.*}} 'void' +// NONATIVESUPPORT-NEXT: MemberExpr {{.*}} 'void (char *)' lvalue .__init_specialization_constants_buffer +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh' +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'char *' +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'char *' lvalue ParmVar {{.*}} '_arg__specialization_constants_buffer' 'char *' +// NONATIVESUPPORT-NEXT: CompoundStmt +// NONATIVESUPPORT-NEXT: ExprWithCleanups +// NONATIVESUPPORT-NEXT: CXXOperatorCallExpr +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'void (*)(group<1>, sycl::kernel_handler) const' +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'void (group<1>, sycl::kernel_handler) const' lvalue CXXMethod {{.*}} 'operator()' 'void (group<1>, sycl::kernel_handler) const' + +// Kernel body with clones +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const (lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}kernel-handler.cpp{{.*}})' +// NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'group<1>':'sycl::group<1>' 'void (sycl::group<1> &&) noexcept' +// NONATIVESUPPORT-NEXT: MaterializeTemporaryExpr +// NONATIVESUPPORT-NEXT: CXXTemporaryObjectExpr +// NONATIVESUPPORT-NEXT: CXXConstructExpr {{.*}} 'sycl::kernel_handler' 'void (const sycl::kernel_handler &) noexcept' +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'const sycl::kernel_handler' lvalue +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh' 'sycl::kernel_handler' + +// Test AST for default SPIR architecture + // Check test_kernel_handler parameters // NATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int)' // NATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' diff --git a/clang/test/SemaSYCL/kernelname-enum.cpp b/clang/test/SemaSYCL/kernelname-enum.cpp index 6bc25873dcb6b..20912670b0dda 100644 --- a/clang/test/SemaSYCL/kernelname-enum.cpp +++ b/clang/test/SemaSYCL/kernelname-enum.cpp @@ -67,15 +67,15 @@ int main() { }); q.submit([&](cl::sycl::handler &cgh) { - // expected-error@Inputs/sycl.hpp:242 {{'dummy_functor_2' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} + // expected-error@Inputs/sycl.hpp:252 {{'dummy_functor_2' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task(f2); }); q.submit([&](cl::sycl::handler &cgh) { - // expected-error@Inputs/sycl.hpp:242 {{'templated_functor' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} + // expected-error@Inputs/sycl.hpp:252 {{'templated_functor' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task(f5); }); diff --git a/clang/test/SemaSYCL/stdtypes_kernel_type.cpp b/clang/test/SemaSYCL/stdtypes_kernel_type.cpp index fafa89dd176cb..b8ee24a1dbacb 100644 --- a/clang/test/SemaSYCL/stdtypes_kernel_type.cpp +++ b/clang/test/SemaSYCL/stdtypes_kernel_type.cpp @@ -25,38 +25,38 @@ queue q; int main() { #ifdef CHECK_ERROR q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:242 {{'nullptr_t' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{type 'nullptr_t' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:252 {{'nullptr_t' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{type 'nullptr_t' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task([=] {}); }); q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:242 {{'std::T' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{type 'std::T' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:252 {{'std::T' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{type 'std::T' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task([=] {}); }); q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:242 {{'Templated_kernel_name' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{type 'nullptr_t' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:252 {{'Templated_kernel_name' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{type 'nullptr_t' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task>([=] {}); }); q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:242 {{'Templated_kernel_name' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{type 'std::U' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:252 {{'Templated_kernel_name' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{type 'std::U' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task>([=] {}); }); q.submit([&](handler &cgh) { - // expected-error@Inputs/sycl.hpp:242 {{'Templated_kernel_name2>' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242{{type 'std::Foo' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:252 {{'Templated_kernel_name2>' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252{{type 'std::Foo' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task>>([]() {}); }); q.submit([&](handler &cgh) { - // expected-error@Inputs/sycl.hpp:242 {{'TemplParamPack' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{type 'nullptr_t' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:252 {{'TemplParamPack' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{type 'nullptr_t' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task>([]() {}); }); diff --git a/clang/test/SemaSYCL/unnamed-kernel.cpp b/clang/test/SemaSYCL/unnamed-kernel.cpp index c499ede958730..40434a69f1d2c 100644 --- a/clang/test/SemaSYCL/unnamed-kernel.cpp +++ b/clang/test/SemaSYCL/unnamed-kernel.cpp @@ -33,8 +33,8 @@ struct MyWrapper { void test() { cl::sycl::queue q; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:242 {{'InvalidKernelName1' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:252 {{'InvalidKernelName1' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{'InvalidKernelName1' should be globally-visible}} // expected-note@+4{{in instantiation of function template specialization}} #endif class InvalidKernelName1 {}; @@ -43,8 +43,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:242 {{'namespace1::KernelName' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{'InvalidKernelName2' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:252 {{'namespace1::KernelName' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{'InvalidKernelName2' should be globally-visible}} // expected-note@+4{{in instantiation of function template specialization}} #endif class InvalidKernelName2 {}; @@ -53,8 +53,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:242 {{'MyWrapper::InvalidKernelName0' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{'MyWrapper::InvalidKernelName0' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:252 {{'MyWrapper::InvalidKernelName0' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{'MyWrapper::InvalidKernelName0' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -62,8 +62,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:242 {{'namespace1::KernelName' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{'MyWrapper::InvalidKernelName3' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:252 {{'namespace1::KernelName' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{'MyWrapper::InvalidKernelName3' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -76,8 +76,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:242 {{'std::max_align_t' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{type 'std::max_align_t' cannot be in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:252 {{'std::max_align_t' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{type 'std::max_align_t' cannot be in the "std" namespace}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -86,8 +86,8 @@ struct MyWrapper { using InvalidAlias = InvalidKernelName4; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:242 {{'MyWrapper::InvalidKernelName4' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{'MyWrapper::InvalidKernelName4' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:252 {{'MyWrapper::InvalidKernelName4' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{'MyWrapper::InvalidKernelName4' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -96,16 +96,16 @@ struct MyWrapper { using InvalidAlias1 = InvalidKernelName5; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:242 {{'namespace1::KernelName' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{'MyWrapper::InvalidKernelName5' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:252 {{'namespace1::KernelName' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{'MyWrapper::InvalidKernelName5' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { h.single_task>([] {}); }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:242 {{'Templated_kernel_name2>' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@Inputs/sycl.hpp:252 {{'Templated_kernel_name2>' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{'InvalidKernelName1' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -117,8 +117,8 @@ struct MyWrapper { int main() { cl::sycl::queue q; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error-re@Inputs/sycl.hpp:242 {{'(lambda at {{.*}}unnamed-kernel.cpp{{.*}}' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:242 {{unnamed type used in a SYCL kernel name}} + // expected-error-re@Inputs/sycl.hpp:252 {{'(lambda at {{.*}}unnamed-kernel.cpp{{.*}}' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:252 {{unnamed type used in a SYCL kernel name}} // expected-note@+2{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); }); From 573a98be915b1557bba6e0eeded6bfa72c78b921 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Wed, 24 Mar 2021 11:01:11 -0700 Subject: [PATCH 11/13] Implement review comments Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 24 +++++++----------------- 1 file changed, 7 insertions(+), 17 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 4b25015af7bbf..9ffa0ce900698 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -793,26 +793,16 @@ static bool isDefaultSPIRArch(ASTContext &Context) { static ParmVarDecl *getSyclKernelHandlerArg(FunctionDecl *KernelCallerFunc) { // Specialization constants in SYCL 2020 are not captured by lambda and // accessed through new optional lambda argument kernel_handler - auto KHArg = - std::find_if(KernelCallerFunc->param_begin(), - KernelCallerFunc->param_end(), [](ParmVarDecl *PVD) { - return Util::isSyclKernelHandlerType(PVD->getType()); - }); - - ParmVarDecl *KernelHandlerArg = - (KHArg != KernelCallerFunc->param_end()) ? *KHArg : nullptr; + auto IsHandlerLambda = [](ParmVarDecl *PVD) { + return Util::isSyclKernelHandlerType(PVD->getType()); + }; - if (KernelHandlerArg) { - auto KHArgTooMany = std::find_if( - std::next(KHArg), KernelCallerFunc->param_end(), [](ParmVarDecl *PVD) { - return Util::isSyclKernelHandlerType(PVD->getType()); - }); + assert(llvm::count_if(KernelCallerFunc->parameters(), IsHandlerLambda) <= 1 && + "Multiple kernel_handler parameters"); - assert(KHArgTooMany == KernelCallerFunc->param_end() && - "Too many kernel_handler arguments"); - } + auto KHArg = llvm::find_if(KernelCallerFunc->parameters(), IsHandlerLambda); - return KernelHandlerArg; + return (KHArg != KernelCallerFunc->param_end()) ? *KHArg : nullptr; } // anonymous namespace so these don't get linkage. From d1e3854f6da4783d78b35f2b2bb04ee2b81665bf Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Wed, 24 Mar 2021 14:11:48 -0700 Subject: [PATCH 12/13] Remove unnecessary define Signed-off-by: Elizabeth Andrews --- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index c941a950f135d..1633728c29b92 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -301,7 +301,6 @@ ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { kernelFunc(); } -#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) template ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc, kernel_handler kh) { kernelFunc(kh); From 55510a900d2ae103593d6994d14cafacdc8af2dd Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Thu, 25 Mar 2021 23:53:30 -0700 Subject: [PATCH 13/13] Use markers in header. Update related tests. Make test mangling insensitive Signed-off-by: Elizabeth Andrews --- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 4 +-- .../test/CodeGenSYCL/kernel-by-reference.cpp | 4 +-- clang/test/CodeGenSYCL/kernel-handler.cpp | 18 +++++----- clang/test/SemaSYCL/Inputs/sycl.hpp | 4 +-- clang/test/SemaSYCL/args-size-overflow.cpp | 4 +-- .../deferred-diagnostics-aux-builtin.cpp | 2 +- .../SemaSYCL/deferred-diagnostics-emit.cpp | 4 +-- clang/test/SemaSYCL/float128.cpp | 6 ++-- clang/test/SemaSYCL/implicit_kernel_type.cpp | 16 ++++----- clang/test/SemaSYCL/int128.cpp | 6 ++-- clang/test/SemaSYCL/kernelname-enum.cpp | 8 ++--- clang/test/SemaSYCL/stdtypes_kernel_type.cpp | 24 ++++++------- clang/test/SemaSYCL/unnamed-kernel.cpp | 36 +++++++++---------- 13 files changed, 68 insertions(+), 68 deletions(-) diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 1633728c29b92..fc3ca2c146ad6 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -297,7 +297,7 @@ class kernel_handler { #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) template -ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { +ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { // #KernelSingleTask kernelFunc(); } @@ -307,7 +307,7 @@ ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc, kernel_ha } template -ATTR_SYCL_KERNEL void kernel_single_task_2017(KernelType kernelFunc) { +ATTR_SYCL_KERNEL void kernel_single_task_2017(KernelType kernelFunc) { // #KernelSingleTask2017 kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/kernel-by-reference.cpp b/clang/test/CodeGenSYCL/kernel-by-reference.cpp index 33c057ca6be19..720290c7b05cb 100644 --- a/clang/test/CodeGenSYCL/kernel-by-reference.cpp +++ b/clang/test/CodeGenSYCL/kernel-by-reference.cpp @@ -15,7 +15,7 @@ int simple_add(int i) { int main() { queue q; #if defined(SYCL2020) - // expected-warning@Inputs/sycl.hpp:311 {{Passing kernel functions by value is deprecated in SYCL 2020}} + // expected-warning@#KernelSingleTask2017 {{Passing kernel functions by value is deprecated in SYCL 2020}} // expected-note@+3 {{in instantiation of function template specialization}} #endif q.submit([&](handler &h) { @@ -23,7 +23,7 @@ int main() { }); #if defined(SYCL2017) - // expected-warning@Inputs/sycl.hpp:300 {{Passing of kernel functions by reference is a SYCL 2020 extension}} + // expected-warning@#KernelSingleTask {{Passing of kernel functions by reference is a SYCL 2020 extension}} // expected-note@+3 {{in instantiation of function template specialization}} #endif q.submit([&](handler &h) { diff --git a/clang/test/CodeGenSYCL/kernel-handler.cpp b/clang/test/CodeGenSYCL/kernel-handler.cpp index 50eea67ddabeb..9c24acc31a74c 100644 --- a/clang/test/CodeGenSYCL/kernel-handler.cpp +++ b/clang/test/CodeGenSYCL/kernel-handler.cpp @@ -22,17 +22,17 @@ void test(int val) { }); } -// NONATIVESUPPORT: define dso_local void @"_ZTSZZ4testiENK3$_0clERN2cl4sycl7handlerEE19test_kernel_handler" +// NONATIVESUPPORT: define dso_local void @"{{.*}}test_kernel_handler{{.*}}" // NONATIVESUPPORT-SAME: (i32 %_arg_, i8* %_arg__specialization_constants_buffer) -// NONATIVESUPPORT: %kh = alloca %"class._ZTSN2cl4sycl14kernel_handlerE.cl::sycl::kernel_handler", align 1 +// NONATIVESUPPORT: %kh = alloca %"class.[[MANGLEDCLASS:[a-zA-Z0-9_]+]].cl::sycl::kernel_handler", align 1 // NONATIVESUPPORT: %[[KH:[0-9]+]] = load i8*, i8** %_arg__specialization_constants_buffer.addr, align 8 -// NONATIVESUPPORT: call void @_ZN2cl4sycl14kernel_handler38__init_specialization_constants_bufferEPc(%"class._ZTSN2cl4sycl14kernel_handlerE.cl::sycl::kernel_handler"* nonnull dereferenceable(1) %kh, i8* %[[KH]]) -// NONATIVESUPPORT: void @"_ZZZ4testiENK3$_0clERN2cl4sycl7handlerEENKUlT_E_clINS1_14kernel_handlerEEEDaS4_" -// NONATIVESUPPORT-SAME: byval(%"class._ZTSN2cl4sycl14kernel_handlerE.cl::sycl::kernel_handler") +// NONATIVESUPPORT: call void @{{.*}}__init_specialization_constants_buffer{{.*}}(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler"* nonnull dereferenceable(1) %kh, i8* %[[KH]]) +// NONATIVESUPPORT: void @"[[MANGLEDKERNELCALL:[a-zA-Z0-9_$]+]]" +// NONATIVESUPPORT-SAME: byval(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler") -// NATIVESUPPORT: define dso_local spir_kernel void @"_ZTSZZ4testiENK3$_0clERN2cl4sycl7handlerEE19test_kernel_handler" +// NATIVESUPPORT: define dso_local spir_kernel void @"{{.*}}test_kernel_handler{{.*}}" // NATIVESUPPORT-SAME: (i32 %_arg_) -// NATIVESUPPORT: %kh = alloca %"class._ZTSN2cl4sycl14kernel_handlerE.cl::sycl::kernel_handler" +// NATIVESUPPORT: %kh = alloca %"class.[[MANGLEDCLASS:[a-zA-Z0-9_]+]].cl::sycl::kernel_handler" // NATIVESUPPORT-NOT: __init_specialization_constants_buffer -// NATIVE-SUPPORT: call spir_func void @"_ZZZ4testiENK3$_0clERN2cl4sycl7handlerEENKUlT_E_clINS1_14kernel_handlerEEEDaS4_" -// NATIVE-SUPPORT-SAME: byval(%"class._ZTSN2cl4sycl14kernel_handlerE.cl::sycl::kernel_handler") +// NATIVE-SUPPORT: call spir_func void @"[[MANGLEDKERNELCALL:[a-zA-Z0-9_$]+]]" +// NATIVE-SUPPORT-SAME: byval(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler") diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index fa6ce419662b5..4782bafa74ce5 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -228,7 +228,7 @@ template struct get_kernel_wrapper_name_t { #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) template ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { - kernelFunc(); + kernelFunc(); // #KernelSingleTaskKernelFuncCall } template ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc, kernel_handler kh) { @@ -249,7 +249,7 @@ class handler { void single_task(const KernelType &kernelFunc) { using NameT = typename get_kernel_name_t::name; #ifdef __SYCL_DEVICE_ONLY__ - kernel_single_task(kernelFunc); + kernel_single_task(kernelFunc); // #KernelSingleTask #else kernelFunc(); #endif diff --git a/clang/test/SemaSYCL/args-size-overflow.cpp b/clang/test/SemaSYCL/args-size-overflow.cpp index 827fda68d6d35..d3d88d39ca4e0 100644 --- a/clang/test/SemaSYCL/args-size-overflow.cpp +++ b/clang/test/SemaSYCL/args-size-overflow.cpp @@ -11,9 +11,9 @@ queue q; using Accessor = accessor; #ifdef SPIR64 -// expected-warning@Inputs/sycl.hpp:252 {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}} +// expected-warning@#KernelSingleTask {{size of kernel arguments (7994 bytes) may exceed the supported maximum of 2048 bytes on some devices}} #elif SPIR32 -// expected-warning@Inputs/sycl.hpp:252 {{size of kernel arguments (7986 bytes) may exceed the supported maximum of 2048 bytes on some devices}} +// expected-warning@#KernelSingleTask {{size of kernel arguments (7986 bytes) may exceed the supported maximum of 2048 bytes on some devices}} #endif void use() { diff --git a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp index e16cc96182c55..3caee03756cad 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp @@ -12,7 +12,7 @@ int main(int argc, char **argv) { _mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}} deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:231 {{called by 'kernel_single_task([]() { _mm_prefetch("test", 4); // expected-error {{builtin is not supported on this target}} _mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}} expected-error {{builtin is not supported on this target}} diff --git a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp index f69eb508fadcf..63bdc3b226e6e 100644 --- a/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp +++ b/clang/test/SemaSYCL/deferred-diagnostics-emit.cpp @@ -64,7 +64,7 @@ template void setup_sycl_operation(const T VA[]) { deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:231 {{called by 'kernel_single_task([]() { // ======= Zero Length Arrays Not Allowed in Kernel ========== // expected-error@+1 {{zero-length arrays are not permitted in C++}} @@ -156,7 +156,7 @@ int main(int argc, char **argv) { // --- direct lambda testing --- deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:231 2 {{called by 'kernel_single_task([]() { // expected-error@+1 {{zero-length arrays are not permitted in C++}} int BadArray[0]; diff --git a/clang/test/SemaSYCL/float128.cpp b/clang/test/SemaSYCL/float128.cpp index 4444a6c30548c..da7be1573151d 100644 --- a/clang/test/SemaSYCL/float128.cpp +++ b/clang/test/SemaSYCL/float128.cpp @@ -71,7 +71,7 @@ int main() { __float128 CapturedToDevice = 1; host_ok(); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:231 {{called by 'kernel_single_task([=]() { // expected-error@+1 {{'__float128' is not supported on this target}} decltype(CapturedToDevice) D; @@ -88,7 +88,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:231 4{{called by 'kernel_single_task([=]() { // expected-note@+1 2{{called by 'operator()'}} usage(); @@ -104,7 +104,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:231 {{called by 'kernel_single_task([=]() { // expected-note@+1 3{{used here}} Z<__float128> S; diff --git a/clang/test/SemaSYCL/implicit_kernel_type.cpp b/clang/test/SemaSYCL/implicit_kernel_type.cpp index a7fdccb67cab1..9f28a950affbe 100644 --- a/clang/test/SemaSYCL/implicit_kernel_type.cpp +++ b/clang/test/SemaSYCL/implicit_kernel_type.cpp @@ -25,12 +25,12 @@ int main() { queue q; #if defined(WARN) - // expected-error@Inputs/sycl.hpp:252 {{'InvalidKernelName1' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@#KernelSingleTask {{'InvalidKernelName1' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{'InvalidKernelName1' should be globally-visible}} // expected-note@+8 {{in instantiation of function template specialization}} #elif defined(ERROR) - // expected-error@Inputs/sycl.hpp:252 {{'InvalidKernelName1' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@#KernelSingleTask {{'InvalidKernelName1' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{'InvalidKernelName1' should be globally-visible}} // expected-note@+4 {{in instantiation of function template specialization}} #endif class InvalidKernelName1 {}; @@ -39,9 +39,9 @@ int main() { }); #if defined(WARN) - // expected-warning@Inputs/sycl.hpp:252 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-warning@#KernelSingleTask {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #elif defined(ERROR) - // expected-error@Inputs/sycl.hpp:252 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-error@#KernelSingleTask {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #endif q.submit([&](handler &h) { @@ -53,9 +53,9 @@ int main() { }); #if defined(WARN) - // expected-warning@Inputs/sycl.hpp:252 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-warning@#KernelSingleTask {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #elif defined(ERROR) - // expected-error@Inputs/sycl.hpp:252 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} + // expected-error@#KernelSingleTask {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}} #endif q.submit([&](handler &h) { diff --git a/clang/test/SemaSYCL/int128.cpp b/clang/test/SemaSYCL/int128.cpp index 76e4668a3c502..a833ccd6e448f 100644 --- a/clang/test/SemaSYCL/int128.cpp +++ b/clang/test/SemaSYCL/int128.cpp @@ -80,7 +80,7 @@ int main() { __int128 CapturedToDevice = 1; host_ok(); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:231 {{called by 'kernel_single_task([=]() { // expected-error@+1 {{'__int128' is not supported on this target}} decltype(CapturedToDevice) D; @@ -97,7 +97,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:231 5{{called by 'kernel_single_task([=]() { // expected-note@+1 2{{called by 'operator()'}} usage(); @@ -115,7 +115,7 @@ int main() { }); deviceQueue.submit([&](sycl::handler &h) { - // expected-note@Inputs/sycl.hpp:231 {{called by 'kernel_single_task([=]() { // expected-note@+1 3{{used here}} Z<__int128> S; diff --git a/clang/test/SemaSYCL/kernelname-enum.cpp b/clang/test/SemaSYCL/kernelname-enum.cpp index 20912670b0dda..b4f224729db8a 100644 --- a/clang/test/SemaSYCL/kernelname-enum.cpp +++ b/clang/test/SemaSYCL/kernelname-enum.cpp @@ -67,15 +67,15 @@ int main() { }); q.submit([&](cl::sycl::handler &cgh) { - // expected-error@Inputs/sycl.hpp:252 {{'dummy_functor_2' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} + // expected-error@#KernelSingleTask {{'dummy_functor_2' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task(f2); }); q.submit([&](cl::sycl::handler &cgh) { - // expected-error@Inputs/sycl.hpp:252 {{'templated_functor' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} + // expected-error@#KernelSingleTask {{'templated_functor' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task(f5); }); diff --git a/clang/test/SemaSYCL/stdtypes_kernel_type.cpp b/clang/test/SemaSYCL/stdtypes_kernel_type.cpp index b8ee24a1dbacb..aef5333dbc176 100644 --- a/clang/test/SemaSYCL/stdtypes_kernel_type.cpp +++ b/clang/test/SemaSYCL/stdtypes_kernel_type.cpp @@ -25,38 +25,38 @@ queue q; int main() { #ifdef CHECK_ERROR q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:252 {{'nullptr_t' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{type 'nullptr_t' cannot be in the "std" namespace}} + // expected-error@#KernelSingleTask {{'nullptr_t' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{type 'nullptr_t' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task([=] {}); }); q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:252 {{'std::T' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{type 'std::T' cannot be in the "std" namespace}} + // expected-error@#KernelSingleTask {{'std::T' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{type 'std::T' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task([=] {}); }); q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:252 {{'Templated_kernel_name' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{type 'nullptr_t' cannot be in the "std" namespace}} + // expected-error@#KernelSingleTask {{'Templated_kernel_name' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{type 'nullptr_t' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task>([=] {}); }); q.submit([&](handler &h) { - // expected-error@Inputs/sycl.hpp:252 {{'Templated_kernel_name' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{type 'std::U' cannot be in the "std" namespace}} + // expected-error@#KernelSingleTask {{'Templated_kernel_name' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{type 'std::U' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} h.single_task>([=] {}); }); q.submit([&](handler &cgh) { - // expected-error@Inputs/sycl.hpp:252 {{'Templated_kernel_name2>' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252{{type 'std::Foo' cannot be in the "std" namespace}} + // expected-error@#KernelSingleTask {{'Templated_kernel_name2>' is an invalid kernel name type}} + // expected-note@#KernelSingleTask{{type 'std::Foo' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task>>([]() {}); }); q.submit([&](handler &cgh) { - // expected-error@Inputs/sycl.hpp:252 {{'TemplParamPack' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{type 'nullptr_t' cannot be in the "std" namespace}} + // expected-error@#KernelSingleTask {{'TemplParamPack' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{type 'nullptr_t' cannot be in the "std" namespace}} // expected-note@+1{{in instantiation of function template specialization}} cgh.single_task>([]() {}); }); diff --git a/clang/test/SemaSYCL/unnamed-kernel.cpp b/clang/test/SemaSYCL/unnamed-kernel.cpp index 40434a69f1d2c..d728f60dbc72a 100644 --- a/clang/test/SemaSYCL/unnamed-kernel.cpp +++ b/clang/test/SemaSYCL/unnamed-kernel.cpp @@ -33,8 +33,8 @@ struct MyWrapper { void test() { cl::sycl::queue q; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:252 {{'InvalidKernelName1' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@#KernelSingleTask {{'InvalidKernelName1' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{'InvalidKernelName1' should be globally-visible}} // expected-note@+4{{in instantiation of function template specialization}} #endif class InvalidKernelName1 {}; @@ -43,8 +43,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:252 {{'namespace1::KernelName' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{'InvalidKernelName2' should be globally-visible}} + // expected-error@#KernelSingleTask {{'namespace1::KernelName' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{'InvalidKernelName2' should be globally-visible}} // expected-note@+4{{in instantiation of function template specialization}} #endif class InvalidKernelName2 {}; @@ -53,8 +53,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:252 {{'MyWrapper::InvalidKernelName0' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{'MyWrapper::InvalidKernelName0' should be globally-visible}} + // expected-error@#KernelSingleTask {{'MyWrapper::InvalidKernelName0' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{'MyWrapper::InvalidKernelName0' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -62,8 +62,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:252 {{'namespace1::KernelName' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{'MyWrapper::InvalidKernelName3' should be globally-visible}} + // expected-error@#KernelSingleTask {{'namespace1::KernelName' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{'MyWrapper::InvalidKernelName3' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -76,8 +76,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:252 {{'std::max_align_t' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{type 'std::max_align_t' cannot be in the "std" namespace}} + // expected-error@#KernelSingleTask {{'std::max_align_t' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{type 'std::max_align_t' cannot be in the "std" namespace}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -86,8 +86,8 @@ struct MyWrapper { using InvalidAlias = InvalidKernelName4; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:252 {{'MyWrapper::InvalidKernelName4' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{'MyWrapper::InvalidKernelName4' should be globally-visible}} + // expected-error@#KernelSingleTask {{'MyWrapper::InvalidKernelName4' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{'MyWrapper::InvalidKernelName4' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -96,16 +96,16 @@ struct MyWrapper { using InvalidAlias1 = InvalidKernelName5; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:252 {{'namespace1::KernelName' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{'MyWrapper::InvalidKernelName5' should be globally-visible}} + // expected-error@#KernelSingleTask {{'namespace1::KernelName' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{'MyWrapper::InvalidKernelName5' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { h.single_task>([] {}); }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:252 {{'Templated_kernel_name2>' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{'InvalidKernelName1' should be globally-visible}} + // expected-error@#KernelSingleTask {{'Templated_kernel_name2>' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{'InvalidKernelName1' should be globally-visible}} // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { @@ -117,8 +117,8 @@ struct MyWrapper { int main() { cl::sycl::queue q; #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error-re@Inputs/sycl.hpp:252 {{'(lambda at {{.*}}unnamed-kernel.cpp{{.*}}' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:252 {{unnamed type used in a SYCL kernel name}} + // expected-error-re@#KernelSingleTask {{'(lambda at {{.*}}unnamed-kernel.cpp{{.*}}' is an invalid kernel name type}} + // expected-note@#KernelSingleTask {{unnamed type used in a SYCL kernel name}} // expected-note@+2{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });