[SYCL] Move diagnostic from integration header emission

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -14,6 +14,7 @@ let Component = "Sema" in {
 let CategoryName = "Semantic Issue" in {
 def note_previous_decl : Note<"%0 declared here">;
 def note_entity_declared_at : Note<"%0 declared here">;
+def note_nullptr_used : Note<"nullptr is a prvalue of type std::nullptr_t">;
 def note_callee_decl : Note<"%0 declared here">;
 def note_defined_here : Note<"%0 defined here">;
 

clang/lib/Sema/SemaSYCL.cpp

@@ -2834,10 +2834,15 @@ class SYCLKernelNameTypeVisitor
   Sema &S;
   SourceLocation KernelInvocationFuncLoc;
   using InnerTypeVisitor = TypeVisitor<SYCLKernelNameTypeVisitor>;
-  using InnerTAVisitor =
+  using InnerTemplArgVisitor =
       ConstTemplateArgumentVisitor<SYCLKernelNameTypeVisitor>;
   bool IsInvalid = false;
 
+  void VisitTemplateArgs(ArrayRef<TemplateArgument> Args) {
+    for (auto &A : Args)
+      Visit(A);
+  }
+
 public:
   SYCLKernelNameTypeVisitor(Sema &S, SourceLocation KernelInvocationFuncLoc)
       : S(S), KernelInvocationFuncLoc(KernelInvocationFuncLoc) {}
@@ -2848,15 +2853,20 @@ class SYCLKernelNameTypeVisitor
     if (T.isNull())
       return;
     const CXXRecordDecl *RD = T->getAsCXXRecordDecl();
-    if (!RD)
+    if (!RD) {
+      if (T->isNullPtrType()) {
+        S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named)
+            << /* kernel name cannot be a type in the std namespace */ 3;
+        S.Diag(KernelInvocationFuncLoc, diag::note_nullptr_used);
+        IsInvalid = true;
+      }
       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]);
-      }
+      ArrayRef<TemplateArgument> Args = TSD->getTemplateArgs().asArray();
+      VisitTemplateArgs(Args);
     } else {
       InnerTypeVisitor::Visit(T.getTypePtr());
     }
@@ -2865,7 +2875,7 @@ class SYCLKernelNameTypeVisitor
   void Visit(const TemplateArgument &TA) {
     if (TA.isNull())
       return;
-    InnerTAVisitor::Visit(TA);
+    InnerTemplArgVisitor::Visit(TA);
   }
 
   void VisitEnumType(const EnumType *T) {
@@ -2886,22 +2896,31 @@ class SYCLKernelNameTypeVisitor
   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) {
+    const DeclContext *DeclCtx = Tag->getDeclContext();
+    if (DeclCtx && !UnnamedLambdaEnabled) {
+      auto *NameSpace = dyn_cast_or_null<NamespaceDecl>(DeclCtx);
+      if (NameSpace && NameSpace->isStdNamespace()) {
         S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named)
-            << /* kernel name is missing */ 0;
+            << /* kernel name cannot be a type in the std namespace */ 3;
         IsInvalid = true;
-      } else {
+        return;
+      }
+      if (!DeclCtx->isTranslationUnit() && !isa<NamespaceDecl>(DeclCtx)) {
+        const bool KernelNameIsMissing = Tag->getName().empty();
+        if (KernelNameIsMissing) {
+          S.Diag(KernelInvocationFuncLoc,
+                 diag::err_sycl_kernel_incorrectly_named)
+              << /* kernel name is missing */ 0;
+          IsInvalid = true;
+          return;
+        }
         if (Tag->isCompleteDefinition()) {
           S.Diag(KernelInvocationFuncLoc,
                  diag::err_sycl_kernel_incorrectly_named)
               << /* kernel name is not globally-visible */ 1;
           IsInvalid = true;
         } else
           S.Diag(KernelInvocationFuncLoc, diag::warn_sycl_implicit_decl);
-
         S.Diag(Tag->getSourceRange().getBegin(), diag::note_previous_decl)
             << Tag->getName();
       }
@@ -2932,6 +2951,10 @@ class SYCLKernelNameTypeVisitor
           VisitEnumType(ET);
     }
   }
+
+  void VisitPackTemplateArgument(const TemplateArgument &TA) {
+    VisitTemplateArgs(TA.getPackAsArray());
+  }
 };
 
 void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
@@ -3337,12 +3360,6 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
       break;
     }
 
-    if (NS->isStdNamespace()) {
-      Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
-          << /* name cannot be a type in the std namespace */ 3;
-      return;
-    }
-
     ++NamespaceCnt;
     const StringRef NSInlinePrefix = NS->isInline() ? "inline " : "";
     NSStr.insert(
@@ -3426,9 +3443,6 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
   const CXXRecordDecl *RD = T->getAsCXXRecordDecl();
 
   if (!RD) {
-    if (T->isNullPtrType())
-      Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
-          << /* name cannot be a type in the std namespace */ 3;
 
     return;
   }

clang/test/CodeGenSYCL/stdtypes_kernel_type.cpp

@@ -1,29 +1,6 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -DCHECK_ERROR -verify %s
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s
-// RUN: FileCheck -input-file=%t.h %s
-//
-// CHECK: #include <CL/sycl/detail/defines_elementary.hpp>
-// CHECK-NEXT: #include <CL/sycl/detail/kernel_desc.hpp>
-//
-// CHECK: static constexpr
-// CHECK-NEXT: const char* const kernel_names[] = {
-// CHECK-NEXT:   "_ZTSm",
-// CHECK-NEXT:   "_ZTSl"
-// CHECK-NEXT: };
-//
-// CHECK: static constexpr
-// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
-// CHECK-NEXT:   //--- _ZTSm
-// CHECK-EMPTY:
-// CHECK-NEXT:   //--- _ZTSl
-// CHECK-EMPTY:
-// CHECK-NEXT: };
-
-// CHECK: template <> struct KernelInfo<unsigned long> {
-// CHECK: template <> struct KernelInfo<long> {
-
-void usage() {
-}
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -sycl-std=2020 -DCHECK_ERROR -verify %s
+
+#include "Inputs/sycl.hpp"
 
 namespace std {
 typedef long unsigned int size_t;
@@ -36,22 +13,44 @@ class U;
 template <typename T>
 struct Templated_kernel_name;
 
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
-  kernelFunc();
-}
+template <typename T, typename... Args> class TemplParamPack;
+
+using namespace cl::sycl;
+queue q;
 
 int main() {
 #ifdef CHECK_ERROR
-  kernel_single_task<std::nullptr_t>([=]() {});                        // expected-error {{kernel name cannot be a type in the "std" namespace}}
-  kernel_single_task<std::T>([=]() {});                                // expected-error {{kernel name cannot be a type in the "std" namespace}}
-  kernel_single_task<Templated_kernel_name<std::nullptr_t>>([=]() {}); // expected-error {{kernel name cannot be a type in the "std" namespace}}
-  kernel_single_task<Templated_kernel_name<std::U>>([=]() {});         // expected-error {{kernel name cannot be a type in the "std" namespace}}
+  // expected-error@Inputs/sycl.hpp:328 5 {{kernel name cannot be a type in the "std" namespace}}
+  // expected-note@Inputs/sycl.hpp:328 3 {{nullptr is a prvalue of type std::nullptr_t}}
+  q.submit([&](handler &h) {
+    // expected-note@+1{{in instantiation of function template specialization}}
+    h.single_task<std::nullptr_t>([=] {});
+  });
+  q.submit([&](handler &h) {
+    // expected-note@+1{{in instantiation of function template specialization}}
+    h.single_task<std::T>([=] {});
+  });
+  q.submit([&](handler &h) {
+    // expected-note@+1{{in instantiation of function template specialization}}
+    h.single_task<Templated_kernel_name<std::nullptr_t>>([=] {});
+  });
+  q.submit([&](handler &h) {
+    // expected-note@+1{{in instantiation of function template specialization}}
+    h.single_task<Templated_kernel_name<std::U>>([=] {});
+  });
+  q.submit([&](handler &cgh) {
+    // expected-note@+1{{in instantiation of function template specialization}}
+    cgh.single_task<TemplParamPack<int, float, std::nullptr_t, double>>([]() {});
+  });
 #endif
 
   // Although in the std namespace, these resolve to builtins such as `int` that are allowed in kernel names
-  kernel_single_task<std::size_t>([=]() {});
-  kernel_single_task<std::ptrdiff_t>([=]() {});
+  q.submit([&](handler &h) {
+    h.single_task<std::size_t>([=] {});
+  });
+  q.submit([&](handler &h) {
+    h.single_task<std::ptrdiff_t>([=] {});
+  });
 
   return 0;
 }

clang/test/CodeGenSYCL/template-template-parameter.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -sycl-std=2020 %s
 // RUN: FileCheck -input-file=%t.h %s
 
 #include "Inputs/sycl.hpp"

clang/test/SemaSYCL/unnamed-kernel.cpp

@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -Wno-sycl-2017-compat -verify %s
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -Wno-sycl-2017-compat -verify %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-unnamed-lambda -fsyntax-only -Wno-sycl-2017-compat -verify %s
 #include "Inputs/sycl.hpp"
 
 #ifdef __SYCL_UNNAMED_LAMBDA__
@@ -70,7 +70,8 @@ struct MyWrapper {
     });
 
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@+3 {{kernel name cannot be a type in the "std" namespace}}
+    // expected-error@Inputs/sycl.hpp:220 {{kernel name cannot be a type in the "std" namespace}}
+    // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<std::max_align_t>([] {});