[SYCL] Add diagnostic to work group scoped destructors.

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">;

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -13964,6 +13964,11 @@ def warn_sycl_old_and_new_kernel_attributes : Warning<
   InGroup<IgnoredAttributes>;
 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<SyclWorkGroupDtor>;
 
 // SPIR-V builtins diagnostics
 def err_spirv_invalid_target : Error<

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<VarDecl>(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);
     }
   }
 

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -272,6 +272,8 @@ template <int dimensions = 1>
 class __SYCL_TYPE(group) group {
 public:
   group() = default; // fake constructor
+  template <typename WorkItemFunctionT>
+  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<DataT> LocalMem;
 };
 
+template <typename T, int Dimensions = 1>
+class __SYCL_TYPE(private_memory) private_memory {
+public:
+  private_memory(const group<Dimensions> &G) { (void)G; }
+
+private:
+  T Val; // device layout: one instance per work-item
+};
+
 namespace ext {
 namespace oneapi {
 namespace experimental {

clang/test/SemaSYCL/new-operator-unevaluated-context.cpp

@@ -108,14 +108,14 @@ int main() {
 
   q.submit([&](sycl::handler &h) {
     h.single_task<class TestActualAlloc>([&]() {
-      // expected-note@Inputs/sycl.hpp:312 {{called by 'kernel_single_task<TestActualAlloc}}
+      // expected-note@Inputs/sycl.hpp:314 {{called by 'kernel_single_task<TestActualAlloc}}
       actual_allocation(); // expected-note {{called by 'operator()'}}
     });
   });
 
   q.submit([&](sycl::handler &h) {
     h.single_task<class TestActualArrayAlloc>([&]() {
-      // expected-note@Inputs/sycl.hpp:312 {{called by 'kernel_single_task<TestActualArrayAlloc}}
+      // expected-note@Inputs/sycl.hpp:314 {{called by 'kernel_single_task<TestActualArrayAlloc}}
       actual_array_allocation(); // expected-note {{called by 'operator()'}}
     });
   });

clang/test/SemaSYCL/sycl-wg-dtor-warning.cpp

@@ -0,0 +1,65 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -verify %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-work-group-dtor -verify=quiet %s
+// RUN: %clang_cc1 -fsycl-is-host -internal-isystem %S/Inputs -verify=quiet %s
+// quiet-no-diagnostics
+#include "sycl.hpp"
+
+struct Trivial { int I; };
+struct NonTrivial { ~NonTrivial() {} };
+struct Inherited : NonTrivial { ~Inherited() = default; };
+struct Composed { ~Composed() = default; int i; NonTrivial N; };
+
+NonTrivial MakeNonTrivial() { return NonTrivial{}; }
+
+struct NonTrivialRange {
+  ~NonTrivialRange() {}
+  int *begin() const { return nullptr; }
+  int *end() const { return nullptr; }
+};
+
+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}}
+    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<NonTrivial> 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<NonTrivial> 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);
+  });
+}