Merge remote-tracking branch 'intel_llvm/sycl' into public_vklochkov_syclmtd_dll

Author: v-klochkov

clang/lib/Sema/SemaSYCL.cpp

@@ -14,6 +14,8 @@
 #include "clang/AST/QualTypeNames.h"
 #include "clang/AST/RecordLayout.h"
 #include "clang/AST/RecursiveASTVisitor.h"
+#include "clang/AST/TemplateArgumentVisitor.h"
+#include "clang/AST/TypeVisitor.h"
 #include "clang/Analysis/CallGraph.h"
 #include "clang/Basic/Attributes.h"
 #include "clang/Basic/Builtins.h"
@@ -2473,9 +2475,111 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
 
 } // namespace
 
+class SYCLKernelNameTypeVisitor
+    : public TypeVisitor<SYCLKernelNameTypeVisitor>,
+      public ConstTemplateArgumentVisitor<SYCLKernelNameTypeVisitor> {
+  Sema &S;
+  SourceLocation KernelInvocationFuncLoc;
+  using InnerTypeVisitor = TypeVisitor<SYCLKernelNameTypeVisitor>;
+  using InnerTAVisitor =
+      ConstTemplateArgumentVisitor<SYCLKernelNameTypeVisitor>;
+
+public:
+  SYCLKernelNameTypeVisitor(Sema &S, SourceLocation KernelInvocationFuncLoc)
+      : S(S), KernelInvocationFuncLoc(KernelInvocationFuncLoc) {}
+
+  void Visit(QualType T) {
+    if (T.isNull())
+      return;
+    const CXXRecordDecl *RD = T->getAsCXXRecordDecl();
+    if (!RD)
+      return;
+    // If KernelNameType has template args visit each template arg via
+    // ConstTemplateArgumentVisitor
+    if (const auto *TSD = dyn_cast<ClassTemplateSpecializationDecl>(RD)) {
+      const TemplateArgumentList &Args = TSD->getTemplateArgs();
+      for (unsigned I = 0; I < Args.size(); I++) {
+        Visit(Args[I]);
+      }
+    } else {
+      InnerTypeVisitor::Visit(T.getTypePtr());
+    }
+  }
+
+  void Visit(const TemplateArgument &TA) {
+    if (TA.isNull())
+      return;
+    InnerTAVisitor::Visit(TA);
+  }
+
+  void VisitEnumType(const EnumType *T) {
+    const EnumDecl *ED = T->getDecl();
+    if (!ED->isScoped() && !ED->isFixed()) {
+      S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named)
+          << /* Unscoped enum requires fixed underlying type */ 2;
+      S.Diag(ED->getSourceRange().getBegin(), diag::note_entity_declared_at)
+          << ED;
+    }
+  }
+
+  void VisitRecordType(const RecordType *T) {
+    return VisitTagDecl(T->getDecl());
+  }
+
+  void VisitTagDecl(const TagDecl *Tag) {
+    bool UnnamedLambdaEnabled =
+        S.getASTContext().getLangOpts().SYCLUnnamedLambda;
+    if (!Tag->getDeclContext()->isTranslationUnit() &&
+        !isa<NamespaceDecl>(Tag->getDeclContext()) && !UnnamedLambdaEnabled) {
+      const bool KernelNameIsMissing = Tag->getName().empty();
+      if (KernelNameIsMissing) {
+        S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named)
+            << /* kernel name is missing */ 0;
+      } else {
+        if (Tag->isCompleteDefinition())
+          S.Diag(KernelInvocationFuncLoc,
+                 diag::err_sycl_kernel_incorrectly_named)
+              << /* kernel name is not globally-visible */ 1;
+        else
+          S.Diag(KernelInvocationFuncLoc, diag::warn_sycl_implicit_decl);
+
+        S.Diag(Tag->getSourceRange().getBegin(), diag::note_previous_decl)
+            << Tag->getName();
+      }
+    }
+  }
+
+  void VisitTypeTemplateArgument(const TemplateArgument &TA) {
+    QualType T = TA.getAsType();
+    if (const auto *ET = T->getAs<EnumType>())
+      VisitEnumType(ET);
+    else
+      Visit(T);
+  }
+
+  void VisitIntegralTemplateArgument(const TemplateArgument &TA) {
+    QualType T = TA.getIntegralType();
+    if (const EnumType *ET = T->getAs<EnumType>())
+      VisitEnumType(ET);
+  }
+
+  void VisitTemplateTemplateArgument(const TemplateArgument &TA) {
+    TemplateDecl *TD = TA.getAsTemplate().getAsTemplateDecl();
+    TemplateParameterList *TemplateParams = TD->getTemplateParameters();
+    for (NamedDecl *P : *TemplateParams) {
+      if (NonTypeTemplateParmDecl *TemplateParam =
+              dyn_cast<NonTypeTemplateParmDecl>(P))
+        if (const EnumType *ET = TemplateParam->getType()->getAs<EnumType>())
+          VisitEnumType(ET);
+    }
+  }
+};
+
 void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
                                ArrayRef<const Expr *> Args) {
   const CXXRecordDecl *KernelObj = getKernelObjectType(KernelFunc);
+  QualType KernelNameType =
+      calculateKernelNameType(getASTContext(), KernelFunc);
   if (!KernelObj) {
     Diag(Args[0]->getExprLoc(), diag::err_sycl_kernel_not_function_object);
     KernelFunc->setInvalidDecl();
@@ -2511,6 +2615,10 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
     return;
 
   KernelObjVisitor Visitor{*this};
+  SYCLKernelNameTypeVisitor KernelTypeVisitor(*this, Args[0]->getExprLoc());
+  // Emit diagnostics for SYCL device kernels only
+  if (LangOpts.SYCLIsDevice)
+    KernelTypeVisitor.Visit(KernelNameType);
   DiagnosingSYCLKernel = true;
   Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker,
                            ArgsSizeChecker);
@@ -2856,18 +2964,6 @@ static void emitWithoutAnonNamespaces(llvm::raw_ostream &OS, StringRef Source) {
   OS << Source;
 }
 
-static bool checkEnumTemplateParameter(const EnumDecl *ED,
-                                       DiagnosticsEngine &Diag,
-                                       SourceLocation KernelLocation) {
-  if (!ED->isScoped() && !ED->isFixed()) {
-    Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named) << 2;
-    Diag.Report(ED->getSourceRange().getBegin(), diag::note_entity_declared_at)
-        << ED;
-    return true;
-  }
-  return false;
-}
-
 // Emits a forward declaration
 void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
                                         SourceLocation KernelLocation) {
@@ -2880,32 +2976,6 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
     auto *NS = dyn_cast_or_null<NamespaceDecl>(DC);
 
     if (!NS) {
-      if (!DC->isTranslationUnit()) {
-        const TagDecl *TD = isa<ClassTemplateDecl>(D)
-                                ? cast<ClassTemplateDecl>(D)->getTemplatedDecl()
-                                : dyn_cast<TagDecl>(D);
-
-        if (TD && !UnnamedLambdaSupport) {
-          // defined class constituting the kernel name is not globally
-          // accessible - contradicts the spec
-          const bool KernelNameIsMissing = TD->getName().empty();
-          if (KernelNameIsMissing) {
-            Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
-                << /* kernel name is missing */ 0;
-            // Don't emit note if kernel name was completely omitted
-          } else {
-            if (TD->isCompleteDefinition())
-              Diag.Report(KernelLocation,
-                          diag::err_sycl_kernel_incorrectly_named)
-                  << /* kernel name is not globally-visible */ 1;
-            else
-              Diag.Report(KernelLocation, diag::warn_sycl_implicit_decl);
-            Diag.Report(D->getSourceRange().getBegin(),
-                        diag::note_previous_decl)
-                << TD->getName();
-          }
-        }
-      }
       break;
     }
 
@@ -3025,7 +3095,6 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
         // Handle Kernel Name Type templated using enum type and value.
         if (const auto *ET = T->getAs<EnumType>()) {
           const EnumDecl *ED = ET->getDecl();
-          if (!checkEnumTemplateParameter(ED, Diag, KernelLocation))
             emitFwdDecl(O, ED, KernelLocation);
         } else if (Arg.getKind() == TemplateArgument::ArgKind::Type)
           emitForwardClassDecls(O, T, KernelLocation, Printed);
@@ -3085,7 +3154,6 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
             QualType T = TemplateParam->getType();
             if (const auto *ET = T->getAs<EnumType>()) {
               const EnumDecl *ED = ET->getDecl();
-              if (!checkEnumTemplateParameter(ED, Diag, KernelLocation))
                 emitFwdDecl(O, ED, KernelLocation);
             }
           }

clang/test/SemaSYCL/args-size-overflow.cpp

@@ -4,6 +4,7 @@
 // RUN: %clang_cc1 -fsycl -triple spir64_gen -Werror=sycl-strict -DERROR -fsycl-is-device -fsyntax-only -verify %s
 
 #include "Inputs/sycl.hpp"
+class Foo;
 
 template <typename Name, typename F>
 __attribute__((sycl_kernel)) void kernel(F KernelFunc) {
@@ -37,5 +38,5 @@ void use() {
 #if defined(GPU) || defined(ERROR)
   // expected-note@+2 {{in instantiation of function template specialization 'parallel_for<Foo}}
 #endif
-  parallel_for<class Foo>(L);
+  parallel_for<Foo>(L);
 }

clang/test/SemaSYCL/implicit_kernel_type.cpp

@@ -51,10 +51,11 @@ class myWrapper2;
 int main() {
   queue q;
 #ifndef __SYCL_UNNAMED_LAMBDA__
-  // expected-note@+1 {{InvalidKernelName1 declared here}}
+  // expected-note@+3 {{InvalidKernelName1 declared here}}
+  // expected-note@+4{{in instantiation of function template specialization}}
+  // expected-error@28 {{kernel needs to have a globally-visible name}}
   class InvalidKernelName1 {};
   q.submit([&](handler &h) {
-    // expected-error@+1 {{kernel needs to have a globally-visible name}}
     h.single_task<InvalidKernelName1>([]() {});
   });
 #endif

clang/test/SemaSYCL/kernelname-enum.cpp

@@ -1,13 +1,14 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -Wno-sycl-2017-compat -verify %s
 
+// expected-error@Inputs/sycl.hpp:220 2{{kernel name is invalid. Unscoped enum requires fixed underlying type}}
 #include "Inputs/sycl.hpp"
 
 enum unscoped_enum_int : int {
   val_1,
   val_2
 };
 
-// expected-note@+1 {{'unscoped_enum_no_type_set' declared here}}
+// expected-note@+1 2 {{'unscoped_enum_no_type_set' declared here}}
 enum unscoped_enum_no_type_set {
   val_3,
   val_4
@@ -29,13 +30,18 @@ class dummy_functor_1 {
   void operator()() const {}
 };
 
-// expected-error@+2 {{kernel name is invalid. Unscoped enum requires fixed underlying type}}
 template <unscoped_enum_no_type_set EnumType>
 class dummy_functor_2 {
 public:
   void operator()() const {}
 };
 
+template <template <unscoped_enum_no_type_set EnumType> class C>
+class templated_functor {
+public:
+  void operator()() const {}
+};
+
 template <scoped_enum_int EnumType>
 class dummy_functor_3 {
 public:
@@ -54,6 +60,7 @@ int main() {
   dummy_functor_2<val_3> f2;
   dummy_functor_3<scoped_enum_int::val_2> f3;
   dummy_functor_4<scoped_enum_no_type_set::val_4> f4;
+  templated_functor<dummy_functor_2> f5;
 
   cl::sycl::queue q;
 
@@ -62,9 +69,15 @@ int main() {
   });
 
   q.submit([&](cl::sycl::handler &cgh) {
+    // expected-note@+1{{in instantiation of function template specialization}}
     cgh.single_task(f2);
   });
 
+  q.submit([&](cl::sycl::handler &cgh) {
+    // expected-note@+1{{in instantiation of function template specialization}}
+    cgh.single_task(f5);
+  });
+
   q.submit([&](cl::sycl::handler &cgh) {
     cgh.single_task(f3);
   });

clang/test/SemaSYCL/unnamed-kernel.cpp

@@ -27,34 +27,38 @@ struct MyWrapper {
   void test() {
     cl::sycl::queue q;
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@+5 {{kernel needs to have a globally-visible name}}
-    // expected-note@+2 {{InvalidKernelName1 declared here}}
+    // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
+    // expected-note@+3 {{InvalidKernelName1 declared here}}
+    // expected-note@+4{{in instantiation of function template specialization}}
 #endif
     class InvalidKernelName1 {};
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<InvalidKernelName1>([] {});
     });
 
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@+5 {{kernel needs to have a globally-visible name}}
-    // expected-note@+2 {{InvalidKernelName2 declared here}}
+    // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
+    // expected-note@+3 {{InvalidKernelName2 declared here}}
+    // expected-note@+4{{in instantiation of function template specialization}}
 #endif
     class InvalidKernelName2 {};
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<namespace1::KernelName<InvalidKernelName2>>([] {});
     });
 
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@+4 {{kernel needs to have a globally-visible name}}
+    // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
     // expected-note@21 {{InvalidKernelName0 declared here}}
+    // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<InvalidKernelName0>([] {});
     });
 
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@+4 {{kernel needs to have a globally-visible name}}
+    // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
     // expected-note@22 {{InvalidKernelName3 declared here}}
+    // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<namespace1::KernelName<InvalidKernelName3>>([] {});
@@ -74,17 +78,19 @@ struct MyWrapper {
 
     using InvalidAlias = InvalidKernelName4;
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@+4 {{kernel needs to have a globally-visible name}}
+    // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
     // expected-note@23 {{InvalidKernelName4 declared here}}
+    // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<InvalidAlias>([] {});
     });
 
     using InvalidAlias1 = InvalidKernelName5;
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@+4 {{kernel needs to have a globally-visible name}}
+    // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
     // expected-note@24 {{InvalidKernelName5 declared here}}
+    // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<namespace1::KernelName<InvalidAlias1>>([] {});
@@ -95,7 +101,8 @@ struct MyWrapper {
 int main() {
   cl::sycl::queue q;
 #ifndef __SYCL_UNNAMED_LAMBDA__
-  // expected-error@+2 {{kernel name is missing}}
+  // expected-error@Inputs/sycl.hpp:220 {{kernel name is missing}}
+  // expected-note@+2{{in instantiation of function template specialization}}
 #endif
   q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });