From a26dbef495d2c27615193219690413f920fcf267 Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Wed, 21 Oct 2026 13:39:01 -0400 Subject: [PATCH 1/2] [SYCL] Add diagnostic to work group scoped destructors. Non-trivial destructors are not currently supported in work group scopes. Depending on the backend, they either fail to link, are silently dropped, or cause other undefined behaviour. This change adds a warning message so users are made aware of this issue. Added a regression test to cover this new warning. Signed-off-by: Delisle, Thierry --- clang/include/clang/Basic/DiagnosticGroups.td | 1 + .../clang/Basic/DiagnosticSemaKinds.td | 5 ++ clang/lib/Sema/SemaSYCL.cpp | 41 ++++++++++-- clang/test/SemaSYCL/Inputs/sycl.hpp | 11 ++++ .../new-operator-unevaluated-context.cpp | 4 +- clang/test/SemaSYCL/sycl-wg-dtor-warning.cpp | 64 +++++++++++++++++++ 6 files changed, 117 insertions(+), 9 deletions(-) create mode 100644 clang/test/SemaSYCL/sycl-wg-dtor-warning.cpp 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..171425b8756ff 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, which is not supported">, + 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..a706ceab32262 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,24 @@ 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(); + } + } + } + } assert(CallOperator && "non callable object is passed as kernel obj"); // Mark the function that it "works" in a work group scope: @@ -3650,7 +3677,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, which is not supported}} + volatile NonTrivial vN; // expected-warning {{variable 'vN' declared at work group scope has a non-trivial destructor, which is not supported}} + const NonTrivial cN; // expected-warning {{variable 'cN' declared at work group scope has a non-trivial destructor, which is not supported}} + NonTrivial aN[2]; // expected-warning {{variable 'aN' declared at work group scope has a non-trivial destructor, which is not supported}} + Inherited I; // expected-warning {{variable 'I' declared at work group scope has a non-trivial destructor, which is not supported}} + Composed C; // expected-warning {{variable 'C' declared at work group scope has a non-trivial destructor, which is not supported}} + auto [i, sbN] = Composed{}; // expected-warning {{variable '[i, sbN]' declared at work group scope has a non-trivial destructor, which is not supported}} + + 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 Captured; // expected-warning {{captured variable 'Captured' copied into the kernel object at work group scope has a non-trivial destructor, which is not supported}} + CGH.parallel_for_work_group([=](sycl::group<1> G, sycl::kernel_handler) { // #PFWGCall + 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, which is not supported}} + // 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, which is not supported}} + const NonTrivial cN; // expected-warning {{variable 'cN' declared at work group scope has a non-trivial destructor, which is not supported}} + NonTrivial aN[2]; // expected-warning {{variable 'aN' declared at work group scope has a non-trivial destructor, which is not supported}} + Inherited I; // expected-warning {{variable 'I' declared at work group scope has a non-trivial destructor, which is not supported}} + Composed C; // expected-warning {{variable 'C' declared at work group scope has a non-trivial destructor, which is not supported}} + auto [i, sbN] = Composed{}; // expected-warning {{variable '[i, sbN]' declared at work group scope has a non-trivial destructor, which is not supported}} + (void)Captured; + + 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 From db02676107dbfcb51a359aa3cd334526e0db6f42 Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Fri, 23 Oct 2026 09:15:39 -0400 Subject: [PATCH 2/2] Implement code review suggestions and add extra test case. Signed-off-by: Delisle, Thierry --- .../clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/lib/Sema/SemaSYCL.cpp | 8 +++++ clang/test/SemaSYCL/sycl-wg-dtor-warning.cpp | 35 ++++++++++--------- 3 files changed, 27 insertions(+), 18 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 171425b8756ff..ec1472c7180a5 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -13967,7 +13967,7 @@ def err_attribute_argument_is_not_valid : Error< 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, which is not supported">, + "has a non-trivial destructor">, InGroup; // SPIR-V builtins diagnostics diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index a706ceab32262..7b00a780e9c9b 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3654,6 +3654,14 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { 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; } } } diff --git a/clang/test/SemaSYCL/sycl-wg-dtor-warning.cpp b/clang/test/SemaSYCL/sycl-wg-dtor-warning.cpp index a1d6aef1b4f95..aee7f3a66a312 100644 --- a/clang/test/SemaSYCL/sycl-wg-dtor-warning.cpp +++ b/clang/test/SemaSYCL/sycl-wg-dtor-warning.cpp @@ -19,13 +19,13 @@ struct NonTrivialRange { void wg_func(sycl::group<1> G) { Trivial T; // no warning - NonTrivial N; // expected-warning {{variable 'N' declared at work group scope has a non-trivial destructor, which is not supported}} - volatile NonTrivial vN; // expected-warning {{variable 'vN' declared at work group scope has a non-trivial destructor, which is not supported}} - const NonTrivial cN; // expected-warning {{variable 'cN' declared at work group scope has a non-trivial destructor, which is not supported}} - NonTrivial aN[2]; // expected-warning {{variable 'aN' declared at work group scope has a non-trivial destructor, which is not supported}} - Inherited I; // expected-warning {{variable 'I' declared at work group scope has a non-trivial destructor, which is not supported}} - Composed C; // expected-warning {{variable 'C' declared at work group scope has a non-trivial destructor, which is not supported}} - auto [i, sbN] = Composed{}; // expected-warning {{variable '[i, sbN]' declared at work group scope has a non-trivial destructor, which is not supported}} + 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 @@ -39,19 +39,20 @@ int main() { sycl::queue Q; Q.submit([&](sycl::handler &CGH) { sycl::kernel_handler KH; - NonTrivial Captured; // expected-warning {{captured variable 'Captured' copied into the kernel object at work group scope has a non-trivial destructor, which is not supported}} - CGH.parallel_for_work_group([=](sycl::group<1> G, sycl::kernel_handler) { // #PFWGCall + 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, which is not supported}} + 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, which is not supported}} - const NonTrivial cN; // expected-warning {{variable 'cN' declared at work group scope has a non-trivial destructor, which is not supported}} - NonTrivial aN[2]; // expected-warning {{variable 'aN' declared at work group scope has a non-trivial destructor, which is not supported}} - Inherited I; // expected-warning {{variable 'I' declared at work group scope has a non-trivial destructor, which is not supported}} - Composed C; // expected-warning {{variable 'C' declared at work group scope has a non-trivial destructor, which is not supported}} - auto [i, sbN] = Composed{}; // expected-warning {{variable '[i, sbN]' declared at work group scope has a non-trivial destructor, which is not supported}} - (void)Captured; + 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