diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td index 080cf39c8c0d6..19766a84dd74c 100644 --- a/clang/include/clang/Basic/DiagnosticGroups.td +++ b/clang/include/clang/Basic/DiagnosticGroups.td @@ -1798,6 +1798,7 @@ def SyclAspectMismatch : DiagGroup<"sycl-aspect-mismatch">; def SyclNativeCPUTargets: DiagGroup<"sycl-native-cpu-targets">; def SyclPrivateAllocaPositiveSize : DiagGroup<"sycl-private-alloca-positive-size">; def SyclUndefinedFuncInImage : DiagGroup<"sycl-undefined-func-in-image">; +def SyclWorkGroupDtor : DiagGroup<"sycl-work-group-dtor">; // Common warnings for SYCL and OpenMP offloading def OffloadTarget : DiagGroup<"offload-target">; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 32db360ef4508..ec1472c7180a5 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -13964,6 +13964,11 @@ def warn_sycl_old_and_new_kernel_attributes : Warning< InGroup; def err_attribute_argument_is_not_valid : Error< "%0 attribute requires integer constant value 0 or 1">; +def warn_sycl_work_group_scope_dtor : Warning< + "%select{variable|captured variable}0 %1 " + "%select{declared at|copied into the kernel object at}0 work group scope " + "has a non-trivial destructor">, + InGroup; // SPIR-V builtins diagnostics def err_spirv_invalid_target : Error< diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 6cfe09546f600..7b00a780e9c9b 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -863,7 +863,7 @@ static bool isSYCLPrivateMemoryVar(VarDecl *VD) { return SemaSYCL::isSyclType(VD->getType(), SYCLTypeAttr::private_memory); } -static void addScopeAttrToLocalVars(FunctionDecl &F) { +static void addScopeAttrToLocalVars(SemaSYCL &SemaSYCLRef, FunctionDecl &F) { for (Decl *D : F.decls()) { VarDecl *VD = dyn_cast(D); @@ -876,10 +876,19 @@ static void addScopeAttrToLocalVars(FunctionDecl &F) { // locals not declared in user code (lambda object parameter, byval // arguments) which will result in alloca w/o any attribute, so need WI // scope too. - SYCLScopeAttr::Level L = isSYCLPrivateMemoryVar(VD) - ? SYCLScopeAttr::Level::WorkItem - : SYCLScopeAttr::Level::WorkGroup; - VD->addAttr(SYCLScopeAttr::CreateImplicit(F.getASTContext(), L)); + if (isSYCLPrivateMemoryVar(VD)) { + VD->addAttr(SYCLScopeAttr::CreateImplicit( + F.getASTContext(), SYCLScopeAttr::Level::WorkItem)); + } else { + if (VD->needsDestruction(SemaSYCLRef.getASTContext()) == + QualType::DK_cxx_destructor) { + SemaSYCLRef.Diag(VD->getLocation(), + diag::warn_sycl_work_group_scope_dtor) + << 0 << VD << VD->getSourceRange(); + } + VD->addAttr(SYCLScopeAttr::CreateImplicit( + F.getASTContext(), SYCLScopeAttr::Level::WorkGroup)); + } } } @@ -948,7 +957,7 @@ class SingleDeviceFunctionTracker { Caller->addAttr( SYCLScopeAttr::CreateImplicit(Parent.SemaSYCLRef.getASTContext(), SYCLScopeAttr::Level::WorkGroup)); - addScopeAttrToLocalVars(*Caller); + addScopeAttrToLocalVars(Parent.SemaSYCLRef, *Caller); } } @@ -3631,6 +3640,32 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // scope memory allocation. KernelObjClone->addAttr(SYCLScopeAttr::CreateImplicit( SemaSYCLRef.getASTContext(), SYCLScopeAttr::Level::WorkGroup)); + if (!SemaSYCLRef.getDiagnostics().isIgnored( + diag::warn_sycl_work_group_scope_dtor, + KernelObjClone->getBeginLoc())) { + if (const CXXRecordDecl *RD = + KernelObjClone->getType()->getAsCXXRecordDecl(); + RD && RD->isLambda()) { + for (const LambdaCapture &C : RD->captures()) { + if (C.capturesVariable() && C.getCaptureKind() == LCK_ByCopy && + C.getCapturedVar()->getType().isDestructedType() == + QualType::DK_cxx_destructor) { + SemaSYCLRef.Diag(C.getCapturedVar()->getLocation(), + diag::warn_sycl_work_group_scope_dtor) + << 1 << C.getCapturedVar() + << C.getCapturedVar()->getSourceRange(); + // If this variable is implicitly captured it's worth pointing out + // where it is captured. Note: For explicit captures, the warning + // already indicates where the variable is captured, so we only add + // a note for implicit captures. + if (!C.isExplicit()) + SemaSYCLRef.Diag(C.getLocation(), + diag::note_var_explicitly_captured_here) + << C.getCapturedVar() << /*explicitly*/ false; + } + } + } + } assert(CallOperator && "non callable object is passed as kernel obj"); // Mark the function that it "works" in a work group scope: @@ -3650,7 +3685,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { MarkWIScope.TraverseDecl(CallOperator); // Now mark local variables declared in the PFWG lambda with work group // scope attribute - addScopeAttrToLocalVars(*CallOperator); + addScopeAttrToLocalVars(SemaSYCLRef, *CallOperator); } } diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 086a68e099adc..d331e719d5de3 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -272,6 +272,8 @@ template class __SYCL_TYPE(group) group { public: group() = default; // fake constructor + template + void parallel_for_work_item(const WorkItemFunctionT &Func) const {} }; class __SYCL_TYPE(kernel_handler) kernel_handler { @@ -475,6 +477,15 @@ __SYCL_TYPE(dynamic_work_group_memory) dynamic_work_group_memory { work_group_memory LocalMem; }; +template +class __SYCL_TYPE(private_memory) private_memory { +public: + private_memory(const group &G) { (void)G; } + +private: + T Val; // device layout: one instance per work-item +}; + namespace ext { namespace oneapi { namespace experimental { diff --git a/clang/test/SemaSYCL/new-operator-unevaluated-context.cpp b/clang/test/SemaSYCL/new-operator-unevaluated-context.cpp index 722e6952ba63b..7703a15239b3c 100644 --- a/clang/test/SemaSYCL/new-operator-unevaluated-context.cpp +++ b/clang/test/SemaSYCL/new-operator-unevaluated-context.cpp @@ -108,14 +108,14 @@ int main() { q.submit([&](sycl::handler &h) { h.single_task([&]() { - // expected-note@Inputs/sycl.hpp:312 {{called by 'kernel_single_task([&]() { - // expected-note@Inputs/sycl.hpp:312 {{called by 'kernel_single_task G) { + Trivial T; // no warning + NonTrivial N; // expected-warning {{variable 'N' declared at work group scope has a non-trivial destructor}} + volatile NonTrivial vN; // expected-warning {{variable 'vN' declared at work group scope has a non-trivial destructor}} + const NonTrivial cN; // expected-warning {{variable 'cN' declared at work group scope has a non-trivial destructor}} + NonTrivial aN[2]; // expected-warning {{variable 'aN' declared at work group scope has a non-trivial destructor}} + Inherited I; // expected-warning {{variable 'I' declared at work group scope has a non-trivial destructor}} + Composed C; // expected-warning {{variable 'C' declared at work group scope has a non-trivial destructor}} + auto [i, sbN] = Composed{}; // expected-warning {{variable '[i, sbN]' declared at work group scope has a non-trivial destructor}} + + sycl::private_memory PM(G); // no warning: work-item scope + const NonTrivial& rN = MakeNonTrivial(); // no warning: temporary is a scoped local, destructor runs + for (int i : NonTrivialRange{}) + (void)i; + + G.parallel_for_work_item([&]() {}); +} + +int main() { + sycl::queue Q; + Q.submit([&](sycl::handler &CGH) { + sycl::kernel_handler KH; + NonTrivial Implicit; // expected-warning {{captured variable 'Implicit' copied into the kernel object at work group scope has a non-trivial destructor}} + NonTrivial Explicit; + CGH.parallel_for_work_group([=, Explicit = Explicit](sycl::group<1> G, sycl::kernel_handler) { // #PFWGCall expected-warning {{captured variable 'Explicit' copied into the kernel object at work group scope has a non-trivial destructor}} + Trivial T; // no warning + static const NonTrivial S; // no warning: not automatic storage + NonTrivial N; // expected-warning {{variable 'N' declared at work group scope has a non-trivial destructor}} + // expected-note@#PFWGCall {{in instantiation of function template specialization}} + volatile NonTrivial vN; // expected-warning {{variable 'vN' declared at work group scope has a non-trivial destructor}} + const NonTrivial cN; // expected-warning {{variable 'cN' declared at work group scope has a non-trivial destructor}} + NonTrivial aN[2]; // expected-warning {{variable 'aN' declared at work group scope has a non-trivial destructor}} + Inherited I; // expected-warning {{variable 'I' declared at work group scope has a non-trivial destructor}} + Composed C; // expected-warning {{variable 'C' declared at work group scope has a non-trivial destructor}} + auto [i, sbN] = Composed{}; // expected-warning {{variable '[i, sbN]' declared at work group scope has a non-trivial destructor}} + (void)Implicit; // expected-note {{variable 'Implicit' is captured here}} + + sycl::private_memory PM(G); // no warning: work-item scope + const NonTrivial& rN = MakeNonTrivial(); // no warning: temporary is a scoped local, destructor runs + for (int i : NonTrivialRange{}) + (void)i; + + wg_func(G); + }, KH); + }); +} \ No newline at end of file