[SYCL] Move diagnostic from integration header emission (#2644)

This patch moves the diagnostic for the kernel name (kernel name
cannot be a type in the "std" namespace) to the point in which the
kernel is called, rather than during integration header emission.

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,19 @@ 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;
+        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 +2874,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 +2895,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 +2950,10 @@ class SYCLKernelNameTypeVisitor
           VisitEnumType(ET);
     }
   }
+
+  void VisitPackTemplateArgument(const TemplateArgument &TA) {
+    VisitTemplateArgs(TA.getPackAsArray());
+  }
 };
 
 void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
@@ -3337,12 +3359,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 +3442,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/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/stdtypes_kernel_type.cpp

@@ -0,0 +1,55 @@
+// 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;
+typedef long int ptrdiff_t;
+typedef decltype(nullptr) nullptr_t;
+class T;
+class U;
+} // namespace std
+
+template <typename T>
+struct Templated_kernel_name;
+
+template <typename T, typename... Args> class TemplParamPack;
+
+using namespace cl::sycl;
+queue q;
+
+int main() {
+#ifdef CHECK_ERROR
+  // expected-error@Inputs/sycl.hpp:220 5 {{kernel name cannot be a type in the "std" namespace}}
+  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
+  q.submit([&](handler &h) {
+    h.single_task<std::size_t>([=] {});
+  });
+  q.submit([&](handler &h) {
+    h.single_task<std::ptrdiff_t>([=] {});
+  });
+
+  return 0;
+}

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>([] {});