From 0b069cc8ff0f072f2a5b49d530ecd97b7ca4145d Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Mon, 24 Aug 2020 07:37:08 +0300 Subject: [PATCH 01/10] [SYCL] Warn when number of kernel args exceeds maximum available on GPU Emit a warning when number of resulting kernel arguments exceeds 2k - maximum available number of kernel arguments on GPU device. Emit a warning only in GPU AOT mode since other devices don't have such limitation. --- .../include/clang/Basic/DiagnosticSemaKinds.td | 3 +++ clang/lib/Sema/SemaSYCL.cpp | 9 +++++++++ clang/test/SemaSYCL/num-args-overflow.cpp | 17 +++++++++++++++++ 3 files changed, 29 insertions(+) create mode 100644 clang/test/SemaSYCL/num-args-overflow.cpp diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index a571fd4080bcb..041d5c85dbfd6 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10990,6 +10990,9 @@ def err_sycl_restrict : Error< "|use a const static or global variable that is neither zero-initialized " "nor constant-initialized" "}0">; +def warn_sycl_kernel_too_many_args : Warning< + "resulting number of kernel arguments %0 is greater than maximum supported " + "on GPU device - %1">; def err_sycl_virtual_types : Error< "No class with a vtable can be used in a SYCL kernel or any code included in the kernel">; def note_sycl_recursive_function_declared_here: Note<"function implemented using recursion declared here">; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 6d62f035a9674..88720da6a3374 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -56,6 +56,7 @@ enum KernelInvocationKind { const static std::string InitMethodName = "__init"; const static std::string FinalizeMethodName = "__finalize"; +const static unsigned GPUMaxKernelArgsNum= 2000; namespace { @@ -1487,6 +1488,14 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { std::back_inserter(ArgTys), [](const ParmVarDecl *PVD) { return PVD->getType(); }); + // TODO: enable template instantiation tree for this diagnostic + if (SemaRef.Context.getTargetInfo().getTriple().getSubArch() == + llvm::Triple::SPIRSubArch_gen) + if (Params.size() > GPUMaxKernelArgsNum) + SemaRef.Diag(KernelDecl->getLocation(), + diag::warn_sycl_kernel_too_many_args) + << static_cast(Params.size()) << GPUMaxKernelArgsNum; + QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, ArgTys, Info); KernelDecl->setType(FuncType); KernelDecl->setParams(Params); diff --git a/clang/test/SemaSYCL/num-args-overflow.cpp b/clang/test/SemaSYCL/num-args-overflow.cpp new file mode 100644 index 0000000000000..a23bf06b98baa --- /dev/null +++ b/clang/test/SemaSYCL/num-args-overflow.cpp @@ -0,0 +1,17 @@ +// RUN: %clang_cc1 -fsycl -triple spir64_gen -DGPU -fsycl-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl -triple spir64 -fsycl-is-device -fsyntax-only -verify %s + +template +__attribute__((sycl_kernel)) void kernel(F kernelFunc) { + kernelFunc(); +} + +void use() { + int Arr[2001]; +#ifdef GPU + // expected-warning@+4 {{resulting number of kernel arguments 2001 is greater than maximum supported on GPU device - 2000}} +#else + // expected-no-diagnostics +#endif + kernel([=]() { (void)Arr[0]; }); +} From ef6a8e0466d7b1d436750e9017b093721663d60e Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Mon, 24 Aug 2020 08:16:58 +0300 Subject: [PATCH 02/10] Fix formatting --- clang/lib/Sema/SemaSYCL.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 88720da6a3374..0e96e92c752ca 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -56,7 +56,7 @@ enum KernelInvocationKind { const static std::string InitMethodName = "__init"; const static std::string FinalizeMethodName = "__finalize"; -const static unsigned GPUMaxKernelArgsNum= 2000; +const static unsigned GPUMaxKernelArgsNum = 2000; namespace { From 05d5cdb5b31fe04e10383061df0e761d5259c594 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 25 Aug 2020 17:50:22 +0300 Subject: [PATCH 03/10] Add diagnostic to SyclStrict group --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/test/SemaSYCL/num-args-overflow.cpp | 6 +++++- 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 041d5c85dbfd6..26e74f6a4c2ea 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10992,7 +10992,7 @@ def err_sycl_restrict : Error< "}0">; def warn_sycl_kernel_too_many_args : Warning< "resulting number of kernel arguments %0 is greater than maximum supported " - "on GPU device - %1">; + "on GPU device - %1">, InGroup; def err_sycl_virtual_types : Error< "No class with a vtable can be used in a SYCL kernel or any code included in the kernel">; def note_sycl_recursive_function_declared_here: Note<"function implemented using recursion declared here">; diff --git a/clang/test/SemaSYCL/num-args-overflow.cpp b/clang/test/SemaSYCL/num-args-overflow.cpp index a23bf06b98baa..3d942bb1c09a4 100644 --- a/clang/test/SemaSYCL/num-args-overflow.cpp +++ b/clang/test/SemaSYCL/num-args-overflow.cpp @@ -1,5 +1,7 @@ // RUN: %clang_cc1 -fsycl -triple spir64_gen -DGPU -fsycl-is-device -fsyntax-only -verify %s // RUN: %clang_cc1 -fsycl -triple spir64 -fsycl-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl -triple spir64_gen -Wno-sycl-strict -fsycl-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl -triple spir64_gen -Werror=sycl-strict -DERROR -fsycl-is-device -fsyntax-only -verify %s template __attribute__((sycl_kernel)) void kernel(F kernelFunc) { @@ -9,7 +11,9 @@ __attribute__((sycl_kernel)) void kernel(F kernelFunc) { void use() { int Arr[2001]; #ifdef GPU - // expected-warning@+4 {{resulting number of kernel arguments 2001 is greater than maximum supported on GPU device - 2000}} + // expected-warning@+6 {{resulting number of kernel arguments 2001 is greater than maximum supported on GPU device - 2000}} +#elif ERROR + // expected-error@+4 {{resulting number of kernel arguments 2001 is greater than maximum supported on GPU device - 2000}} #else // expected-no-diagnostics #endif From 25b7cccfdb713d02c81838e92628ccee86d6711a Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 25 Aug 2020 20:32:54 +0300 Subject: [PATCH 04/10] Update diagnostic message --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 4 ++-- clang/test/SemaSYCL/num-args-overflow.cpp | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 26e74f6a4c2ea..f7791cca29a14 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10991,8 +10991,8 @@ def err_sycl_restrict : Error< "nor constant-initialized" "}0">; def warn_sycl_kernel_too_many_args : Warning< - "resulting number of kernel arguments %0 is greater than maximum supported " - "on GPU device - %1">, InGroup; + "resulting number of kernel arguments (%0) is greater than maximum supported " + "on GPU device - (%1)">, InGroup; def err_sycl_virtual_types : Error< "No class with a vtable can be used in a SYCL kernel or any code included in the kernel">; def note_sycl_recursive_function_declared_here: Note<"function implemented using recursion declared here">; diff --git a/clang/test/SemaSYCL/num-args-overflow.cpp b/clang/test/SemaSYCL/num-args-overflow.cpp index 3d942bb1c09a4..1e27f3ff8e8cd 100644 --- a/clang/test/SemaSYCL/num-args-overflow.cpp +++ b/clang/test/SemaSYCL/num-args-overflow.cpp @@ -11,9 +11,9 @@ __attribute__((sycl_kernel)) void kernel(F kernelFunc) { void use() { int Arr[2001]; #ifdef GPU - // expected-warning@+6 {{resulting number of kernel arguments 2001 is greater than maximum supported on GPU device - 2000}} + // expected-warning@+6 {{resulting number of kernel arguments (2001) is greater than maximum supported on GPU device - (2000)}} #elif ERROR - // expected-error@+4 {{resulting number of kernel arguments 2001 is greater than maximum supported on GPU device - 2000}} + // expected-error@+4 {{resulting number of kernel arguments (2001) is greater than maximum supported on GPU device - (2000)}} #else // expected-no-diagnostics #endif From cfe01094bdbbafa1b4ec058feb1c17d9be1c251a Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Wed, 26 Aug 2020 11:00:39 +0300 Subject: [PATCH 05/10] Add clarifying note --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 6 ++++-- clang/lib/Sema/SemaSYCL.cpp | 8 ++++++-- clang/test/SemaSYCL/num-args-overflow.cpp | 6 ++++-- 3 files changed, 14 insertions(+), 6 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index f7791cca29a14..96695fa88728d 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10991,8 +10991,10 @@ def err_sycl_restrict : Error< "nor constant-initialized" "}0">; def warn_sycl_kernel_too_many_args : Warning< - "resulting number of kernel arguments (%0) is greater than maximum supported " - "on GPU device - (%1)">, InGroup; + "kernel argument count (%0) exceeds supported maximum of %1 on GPU">, + InGroup; +def note_sycl_kernel_args_count : Note<"array elements and fields of a " + "class/struct may be counted separately">; def err_sycl_virtual_types : Error< "No class with a vtable can be used in a SYCL kernel or any code included in the kernel">; def note_sycl_recursive_function_declared_here: Note<"function implemented using recursion declared here">; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 0e96e92c752ca..d1fd9e4b5cff0 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1490,11 +1490,15 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // TODO: enable template instantiation tree for this diagnostic if (SemaRef.Context.getTargetInfo().getTriple().getSubArch() == - llvm::Triple::SPIRSubArch_gen) - if (Params.size() > GPUMaxKernelArgsNum) + llvm::Triple::SPIRSubArch_gen) { + if (Params.size() > GPUMaxKernelArgsNum) { SemaRef.Diag(KernelDecl->getLocation(), diag::warn_sycl_kernel_too_many_args) << static_cast(Params.size()) << GPUMaxKernelArgsNum; + SemaRef.Diag(KernelDecl->getLocation(), + diag::note_sycl_kernel_args_count); + } + } QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, ArgTys, Info); KernelDecl->setType(FuncType); diff --git a/clang/test/SemaSYCL/num-args-overflow.cpp b/clang/test/SemaSYCL/num-args-overflow.cpp index 1e27f3ff8e8cd..0ca0aae66f041 100644 --- a/clang/test/SemaSYCL/num-args-overflow.cpp +++ b/clang/test/SemaSYCL/num-args-overflow.cpp @@ -11,9 +11,11 @@ __attribute__((sycl_kernel)) void kernel(F kernelFunc) { void use() { int Arr[2001]; #ifdef GPU - // expected-warning@+6 {{resulting number of kernel arguments (2001) is greater than maximum supported on GPU device - (2000)}} + // expected-warning@+8 {{kernel argument count (2001) exceeds supported maximum of 2000 on GPU}} + // expected-note@+7 {{array elements and fields of a class/struct may be counted separately}} #elif ERROR - // expected-error@+4 {{resulting number of kernel arguments (2001) is greater than maximum supported on GPU device - (2000)}} + // expected-error@+5 {{kernel argument count (2001) exceeds supported maximum of 2000 on GPU}} + // expected-note@+4 {{array elements and fields of a class/struct may be counted separately}} #else // expected-no-diagnostics #endif From 822367c88dba0363553c6009c11dc28135f9138c Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Mon, 31 Aug 2020 16:47:47 +0300 Subject: [PATCH 06/10] WIP on template thing --- clang/lib/Sema/SemaSYCL.cpp | 92 ++++++++++++++++++++++++++++++------- 1 file changed, 75 insertions(+), 17 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d1fd9e4b5cff0..8604d9b71127d 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1334,6 +1334,7 @@ class SyclKernelUnionChecker : public SyclKernelFieldHandler { }; // A type to Create and own the FunctionDecl for the kernel. +template class SyclKernelDeclCreator : public SyclKernelFieldHandler { FunctionDecl *KernelDecl; llvm::SmallVector Params; @@ -1344,6 +1345,11 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // Keeps track of whether we are currently handling fields inside a struct. int StructDepth = 0; + // Keeps track of number of kernel arguments, used only if this is an arg + // counter. + unsigned NumOfParams = 0; + SourceLocation KernelLoc; + void addParam(const FieldDecl *FD, QualType FieldTy) { const ConstantArrayType *CAT = SemaRef.getASTContext().getAsConstantArrayType(FieldTy); @@ -1488,18 +1494,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { std::back_inserter(ArgTys), [](const ParmVarDecl *PVD) { return PVD->getType(); }); - // TODO: enable template instantiation tree for this diagnostic - if (SemaRef.Context.getTargetInfo().getTriple().getSubArch() == - llvm::Triple::SPIRSubArch_gen) { - if (Params.size() > GPUMaxKernelArgsNum) { - SemaRef.Diag(KernelDecl->getLocation(), - diag::warn_sycl_kernel_too_many_args) - << static_cast(Params.size()) << GPUMaxKernelArgsNum; - SemaRef.Diag(KernelDecl->getLocation(), - diag::note_sycl_kernel_args_count); - } - } - QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, ArgTys, Info); KernelDecl->setType(FuncType); KernelDecl->setParams(Params); @@ -1631,8 +1625,69 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { using SyclKernelFieldHandler::leaveStruct; }; +template<> +void SyclKernelDeclCreator::addParam(const FieldDecl *FD, + QualType FieldTy) { + NumOfParams++; +} + +template <> +bool SyclKernelDeclCreator::handleSpecialType( + FieldDecl *FD, QualType FieldTy, bool isAccessorType) { + const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); + assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); + CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); + assert(InitMethod && "The accessor/sampler must have the __init method"); + for (const ParmVarDecl *Param : InitMethod->parameters()) { + QualType ParamTy = Param->getType(); + addParam(FD, ParamTy.getCanonicalType()); + } + return true; +} + +template<> +bool SyclKernelDeclCreator::handleSyclAccessorType( + const CXXBaseSpecifier &BS, QualType FieldTy) { + const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); + assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); + CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); + assert(InitMethod && "The accessor/sampler must have the __init method"); + for (const ParmVarDecl *Param : InitMethod->parameters()) { + QualType ParamTy = Param->getType(); + addParam(BS, ParamTy.getCanonicalType()); + } + return true; +} + +template<> +bool SyclKernelDeclCreator::handlePointerType(FieldDecl *FD, + QualType FieldTy) { + addParam(FD, FieldTy); + return true; +} + +// template <> +// SyclKernelDeclCreator::SyclKernelDeclCreator(Sema &S, StringRef Name, +// SourceLocation Loc, +// bool IsInline, +// bool IsSIMDKernel) +// : SyclKernelFieldHandler(S), KernelLoc(Loc) {} + +template <> SyclKernelDeclCreator::~SyclKernelDeclCreator() { + if (SemaRef.Context.getTargetInfo().getTriple().getSubArch() == + llvm::Triple::SPIRSubArch_gen) { + if (NumOfParams > GPUMaxKernelArgsNum) { + SemaRef.Diag(KernelDecl->getLocation(), + diag::warn_sycl_kernel_too_many_args) + << NumOfParams << GPUMaxKernelArgsNum; + SemaRef.Diag(KernelDecl->getLocation(), + diag::note_sycl_kernel_args_count); + } + } +} + class SyclKernelBodyCreator : public SyclKernelFieldHandler { - SyclKernelDeclCreator &DeclCreator; + SyclKernelDeclCreator &DeclCreator; llvm::SmallVector BodyStmts; llvm::SmallVector FinalizeStmts; llvm::SmallVector InitExprs; @@ -1867,7 +1922,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } public: - SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, + SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, const CXXRecordDecl *KernelObj, FunctionDecl *KernelCallerFunc) : SyclKernelFieldHandler(S), DeclCreator(DC), @@ -2262,6 +2317,9 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, SyclKernelFieldChecker FieldChecker(*this); SyclKernelUnionChecker UnionChecker(*this); + SyclKernelDeclCreator DeclCreator( + *this, StringRef(), KernelObj->getLocation(), + KernelFunc->isInlined(), KernelFunc->hasAttr()); // check that calling kernel conforms to spec QualType KernelParamTy = KernelFunc->getParamDecl(0)->getType(); if (KernelParamTy->isReferenceType()) { @@ -2276,8 +2334,8 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, KernelObjVisitor Visitor{*this}; DiagnosingSYCLKernel = true; - Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker); - Visitor.VisitRecordFields(KernelObj, FieldChecker, UnionChecker); + Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker, DeclCreator); + Visitor.VisitRecordFields(KernelObj, FieldChecker, UnionChecker, DeclCreator); DiagnosingSYCLKernel = false; if (!FieldChecker.isValid() || !UnionChecker.isValid()) KernelFunc->setInvalidDecl(); @@ -2317,7 +2375,7 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, constructKernelName(*this, KernelCallerFunc, MC); StringRef KernelName(getLangOpts().SYCLUnnamedLambda ? StableName : CalculatedName); - SyclKernelDeclCreator kernel_decl(*this, KernelName, KernelObj->getLocation(), + SyclKernelDeclCreator kernel_decl(*this, KernelName, KernelObj->getLocation(), KernelCallerFunc->isInlined(), KernelCallerFunc->hasAttr()); SyclKernelBodyCreator kernel_body(*this, kernel_decl, KernelObj, From f96d73f7bf4b646804da895c7482daa1a51aba56 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 1 Sep 2020 14:03:28 +0300 Subject: [PATCH 07/10] Revert "WIP on template thing" This reverts commit 822367c88dba0363553c6009c11dc28135f9138c. --- clang/lib/Sema/SemaSYCL.cpp | 92 +++++++------------------------------ 1 file changed, 17 insertions(+), 75 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d4bbec4e786df..6bf6ce9b900da 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1372,7 +1372,6 @@ class SyclKernelUnionChecker : public SyclKernelFieldHandler { }; // A type to Create and own the FunctionDecl for the kernel. -template class SyclKernelDeclCreator : public SyclKernelFieldHandler { FunctionDecl *KernelDecl; llvm::SmallVector Params; @@ -1383,11 +1382,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // Keeps track of whether we are currently handling fields inside a struct. int StructDepth = 0; - // Keeps track of number of kernel arguments, used only if this is an arg - // counter. - unsigned NumOfParams = 0; - SourceLocation KernelLoc; - void addParam(const FieldDecl *FD, QualType FieldTy) { const ConstantArrayType *CAT = SemaRef.getASTContext().getAsConstantArrayType(FieldTy); @@ -1532,6 +1526,18 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { std::back_inserter(ArgTys), [](const ParmVarDecl *PVD) { return PVD->getType(); }); + // TODO: enable template instantiation tree for this diagnostic + if (SemaRef.Context.getTargetInfo().getTriple().getSubArch() == + llvm::Triple::SPIRSubArch_gen) { + if (Params.size() > GPUMaxKernelArgsNum) { + SemaRef.Diag(KernelDecl->getLocation(), + diag::warn_sycl_kernel_too_many_args) + << static_cast(Params.size()) << GPUMaxKernelArgsNum; + SemaRef.Diag(KernelDecl->getLocation(), + diag::note_sycl_kernel_args_count); + } + } + QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, ArgTys, Info); KernelDecl->setType(FuncType); KernelDecl->setParams(Params); @@ -1664,69 +1670,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { using SyclKernelFieldHandler::leaveStruct; }; -template<> -void SyclKernelDeclCreator::addParam(const FieldDecl *FD, - QualType FieldTy) { - NumOfParams++; -} - -template <> -bool SyclKernelDeclCreator::handleSpecialType( - FieldDecl *FD, QualType FieldTy, bool isAccessorType) { - const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); - assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); - CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); - assert(InitMethod && "The accessor/sampler must have the __init method"); - for (const ParmVarDecl *Param : InitMethod->parameters()) { - QualType ParamTy = Param->getType(); - addParam(FD, ParamTy.getCanonicalType()); - } - return true; -} - -template<> -bool SyclKernelDeclCreator::handleSyclAccessorType( - const CXXBaseSpecifier &BS, QualType FieldTy) { - const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); - assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); - CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); - assert(InitMethod && "The accessor/sampler must have the __init method"); - for (const ParmVarDecl *Param : InitMethod->parameters()) { - QualType ParamTy = Param->getType(); - addParam(BS, ParamTy.getCanonicalType()); - } - return true; -} - -template<> -bool SyclKernelDeclCreator::handlePointerType(FieldDecl *FD, - QualType FieldTy) { - addParam(FD, FieldTy); - return true; -} - -// template <> -// SyclKernelDeclCreator::SyclKernelDeclCreator(Sema &S, StringRef Name, -// SourceLocation Loc, -// bool IsInline, -// bool IsSIMDKernel) -// : SyclKernelFieldHandler(S), KernelLoc(Loc) {} - -template <> SyclKernelDeclCreator::~SyclKernelDeclCreator() { - if (SemaRef.Context.getTargetInfo().getTriple().getSubArch() == - llvm::Triple::SPIRSubArch_gen) { - if (NumOfParams > GPUMaxKernelArgsNum) { - SemaRef.Diag(KernelDecl->getLocation(), - diag::warn_sycl_kernel_too_many_args) - << NumOfParams << GPUMaxKernelArgsNum; - SemaRef.Diag(KernelDecl->getLocation(), - diag::note_sycl_kernel_args_count); - } - } -} - class SyclKernelBodyCreator : public SyclKernelFieldHandler { - SyclKernelDeclCreator &DeclCreator; + SyclKernelDeclCreator &DeclCreator; llvm::SmallVector BodyStmts; llvm::SmallVector FinalizeStmts; llvm::SmallVector InitExprs; @@ -2020,7 +1965,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } public: - SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, + SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, const CXXRecordDecl *KernelObj, FunctionDecl *KernelCallerFunc) : SyclKernelFieldHandler(S), DeclCreator(DC), @@ -2419,9 +2364,6 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, SyclKernelFieldChecker FieldChecker(*this); SyclKernelUnionChecker UnionChecker(*this); - SyclKernelDeclCreator DeclCreator( - *this, StringRef(), KernelObj->getLocation(), - KernelFunc->isInlined(), KernelFunc->hasAttr()); // check that calling kernel conforms to spec QualType KernelParamTy = KernelFunc->getParamDecl(0)->getType(); if (KernelParamTy->isReferenceType()) { @@ -2436,8 +2378,8 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, KernelObjVisitor Visitor{*this}; DiagnosingSYCLKernel = true; - Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker, DeclCreator); - Visitor.VisitRecordFields(KernelObj, FieldChecker, UnionChecker, DeclCreator); + Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker); + Visitor.VisitRecordFields(KernelObj, FieldChecker, UnionChecker); DiagnosingSYCLKernel = false; if (!FieldChecker.isValid() || !UnionChecker.isValid()) KernelFunc->setInvalidDecl(); @@ -2477,7 +2419,7 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, constructKernelName(*this, KernelCallerFunc, MC); StringRef KernelName(getLangOpts().SYCLUnnamedLambda ? StableName : CalculatedName); - SyclKernelDeclCreator kernel_decl(*this, KernelName, KernelObj->getLocation(), + SyclKernelDeclCreator kernel_decl(*this, KernelName, KernelObj->getLocation(), KernelCallerFunc->isInlined(), KernelCallerFunc->hasAttr()); SyclKernelBodyCreator kernel_body(*this, kernel_decl, KernelObj, From 3650e918d36f0dbb97a403a0285f5b012b663c32 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 1 Sep 2020 14:48:25 +0300 Subject: [PATCH 08/10] Add NumArgsChecker, extend the test --- clang/lib/Sema/SemaSYCL.cpp | 100 +++++++++++++++++++--- clang/test/SemaSYCL/num-args-overflow.cpp | 38 ++++++-- 2 files changed, 115 insertions(+), 23 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 6bf6ce9b900da..5dd3b1fbb740a 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1526,18 +1526,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { std::back_inserter(ArgTys), [](const ParmVarDecl *PVD) { return PVD->getType(); }); - // TODO: enable template instantiation tree for this diagnostic - if (SemaRef.Context.getTargetInfo().getTriple().getSubArch() == - llvm::Triple::SPIRSubArch_gen) { - if (Params.size() > GPUMaxKernelArgsNum) { - SemaRef.Diag(KernelDecl->getLocation(), - diag::warn_sycl_kernel_too_many_args) - << static_cast(Params.size()) << GPUMaxKernelArgsNum; - SemaRef.Diag(KernelDecl->getLocation(), - diag::note_sycl_kernel_args_count); - } - } - QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, ArgTys, Info); KernelDecl->setType(FuncType); KernelDecl->setParams(Params); @@ -1670,6 +1658,87 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { using SyclKernelFieldHandler::leaveStruct; }; +class SyclKernelNumArgsChecker : public SyclKernelFieldHandler { + SourceLocation KernelLoc; + unsigned NumOfParams = 0; + + void addParam() { NumOfParams++; } + + bool handleSpecialType(QualType FieldTy) { + const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); + assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); + CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); + assert(InitMethod && "The accessor/sampler must have the __init method"); + unsigned NumOfParams = InitMethod->getNumParams(); + for (unsigned I = 0; I < NumOfParams; ++I) + addParam(); + return true; + } + +public: + SyclKernelNumArgsChecker(Sema &S, SourceLocation Loc) + : SyclKernelFieldHandler(S), KernelLoc(Loc) {} + + ~SyclKernelNumArgsChecker() { + if (SemaRef.Context.getTargetInfo().getTriple().getSubArch() == + llvm::Triple::SPIRSubArch_gen) { + if (NumOfParams > GPUMaxKernelArgsNum) { + SemaRef.Diag(KernelLoc, diag::warn_sycl_kernel_too_many_args) + << NumOfParams << GPUMaxKernelArgsNum; + SemaRef.Diag(KernelLoc, diag::note_sycl_kernel_args_count); + } + } + } + + bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final { + return handleSpecialType(FieldTy); + } + + bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &, + QualType FieldTy) final { + return handleSpecialType(FieldTy); + } + + bool handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final { + return handleSpecialType(FieldTy); + } + + bool handleSyclSamplerType(const CXXRecordDecl *, const CXXBaseSpecifier &BS, + QualType FieldTy) final { + return handleSpecialType(FieldTy); + } + + bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { + addParam(); + return true; + } + + bool handleScalarType(FieldDecl *FD, QualType FieldTy) final { + addParam(); + return true; + } + + bool handleUnionType(FieldDecl *FD, QualType FieldTy) final { + return handleScalarType(FD, FieldTy); + } + + bool handleSyclHalfType(FieldDecl *FD, QualType FieldTy) final { + addParam(); + return true; + } + + bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final { + addParam(); + return true; + } + bool handleSyclStreamType(const CXXRecordDecl *, const CXXBaseSpecifier &, + QualType FieldTy) final { + addParam(); + return true; + } + using SyclKernelFieldHandler::handleSyclHalfType; +}; + class SyclKernelBodyCreator : public SyclKernelFieldHandler { SyclKernelDeclCreator &DeclCreator; llvm::SmallVector BodyStmts; @@ -2364,6 +2433,7 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, SyclKernelFieldChecker FieldChecker(*this); SyclKernelUnionChecker UnionChecker(*this); + SyclKernelNumArgsChecker NumArgsChecker(*this, Args[0]->getExprLoc()); // check that calling kernel conforms to spec QualType KernelParamTy = KernelFunc->getParamDecl(0)->getType(); if (KernelParamTy->isReferenceType()) { @@ -2378,8 +2448,10 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, KernelObjVisitor Visitor{*this}; DiagnosingSYCLKernel = true; - Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker); - Visitor.VisitRecordFields(KernelObj, FieldChecker, UnionChecker); + Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker, + NumArgsChecker); + Visitor.VisitRecordFields(KernelObj, FieldChecker, UnionChecker, + NumArgsChecker); DiagnosingSYCLKernel = false; if (!FieldChecker.isValid() || !UnionChecker.isValid()) KernelFunc->setInvalidDecl(); diff --git a/clang/test/SemaSYCL/num-args-overflow.cpp b/clang/test/SemaSYCL/num-args-overflow.cpp index 0ca0aae66f041..e4fcc1c2f2d71 100644 --- a/clang/test/SemaSYCL/num-args-overflow.cpp +++ b/clang/test/SemaSYCL/num-args-overflow.cpp @@ -1,15 +1,17 @@ -// RUN: %clang_cc1 -fsycl -triple spir64_gen -DGPU -fsycl-is-device -fsyntax-only -verify %s -// RUN: %clang_cc1 -fsycl -triple spir64 -fsycl-is-device -fsyntax-only -verify %s -// RUN: %clang_cc1 -fsycl -triple spir64_gen -Wno-sycl-strict -fsycl-is-device -fsyntax-only -verify %s -// RUN: %clang_cc1 -fsycl -triple spir64_gen -Werror=sycl-strict -DERROR -fsycl-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -I %S/Inputs -fsycl -triple spir64_gen -DGPU -fsycl-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -I %S/Inputs -fsycl -triple spir64 -fsycl-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -I %S/Inputs -fsycl -triple spir64_gen -Wno-sycl-strict -fsycl-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -I %S/Inputs -fsycl -triple spir64_gen -Werror=sycl-strict -DERROR -fsycl-is-device -fsyntax-only -verify %s + +#include template -__attribute__((sycl_kernel)) void kernel(F kernelFunc) { - kernelFunc(); +__attribute__((sycl_kernel)) void kernel(F KernelFunc) { + KernelFunc(); } -void use() { - int Arr[2001]; +template +void parallel_for(F KernelFunc) { #ifdef GPU // expected-warning@+8 {{kernel argument count (2001) exceeds supported maximum of 2000 on GPU}} // expected-note@+7 {{array elements and fields of a class/struct may be counted separately}} @@ -19,5 +21,23 @@ void use() { #else // expected-no-diagnostics #endif - kernel([=]() { (void)Arr[0]; }); + kernel(KernelFunc); +} + +using Accessor = + cl::sycl::accessor; + +void use() { + struct S { + int A; + int B; + Accessor AAcc; + Accessor BAcc; + int Array[1991]; + } Args; + auto L = [=]() { (void)Args; }; +#if defined(GPU) || defined(ERROR) + // expected-note@+2 {{in instantiation of function template specialization 'parallel_for(L); } From 8aa2b27a2087e6785b012e87c8870d0c666359b3 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 1 Sep 2020 18:08:17 +0300 Subject: [PATCH 09/10] Apply comments --- clang/lib/Sema/SemaSYCL.cpp | 18 +++++++----------- 1 file changed, 7 insertions(+), 11 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 5dd3b1fbb740a..9d34b5bb02442 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -56,7 +56,7 @@ enum KernelInvocationKind { const static std::string InitMethodName = "__init"; const static std::string FinalizeMethodName = "__finalize"; -const static unsigned GPUMaxKernelArgsNum = 2000; +constexpr unsigned GPUMaxKernelArgsNum = 2000; namespace { @@ -1662,16 +1662,12 @@ class SyclKernelNumArgsChecker : public SyclKernelFieldHandler { SourceLocation KernelLoc; unsigned NumOfParams = 0; - void addParam() { NumOfParams++; } - bool handleSpecialType(QualType FieldTy) { const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); assert(InitMethod && "The accessor/sampler must have the __init method"); - unsigned NumOfParams = InitMethod->getNumParams(); - for (unsigned I = 0; I < NumOfParams; ++I) - addParam(); + NumOfParams += InitMethod->getNumParams(); return true; } @@ -1709,12 +1705,12 @@ class SyclKernelNumArgsChecker : public SyclKernelFieldHandler { } bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { - addParam(); + NumOfParams++; return true; } bool handleScalarType(FieldDecl *FD, QualType FieldTy) final { - addParam(); + NumOfParams++; return true; } @@ -1723,17 +1719,17 @@ class SyclKernelNumArgsChecker : public SyclKernelFieldHandler { } bool handleSyclHalfType(FieldDecl *FD, QualType FieldTy) final { - addParam(); + NumOfParams++; return true; } bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final { - addParam(); + NumOfParams++; return true; } bool handleSyclStreamType(const CXXRecordDecl *, const CXXBaseSpecifier &, QualType FieldTy) final { - addParam(); + NumOfParams++; return true; } using SyclKernelFieldHandler::handleSyclHalfType; From 493e64be1ecac80fbd9cbffb521748f18d972eda Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 1 Sep 2020 18:33:29 +0300 Subject: [PATCH 10/10] Apply comment --- clang/lib/Sema/SemaSYCL.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 9d34b5bb02442..3fadf06a12a13 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1663,7 +1663,7 @@ class SyclKernelNumArgsChecker : public SyclKernelFieldHandler { unsigned NumOfParams = 0; bool handleSpecialType(QualType FieldTy) { - const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); + const CXXRecordDecl *RecordDecl = FieldTy->getAsCXXRecordDecl(); assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); assert(InitMethod && "The accessor/sampler must have the __init method");