[SYCL] Improve diagnostics for invalid SYCL kernel names. (#2730)

This patch modifies the error diagnostic to print the fully instantiated
kernel name to easily identify problematic types in kernel name. A note
diagnostic with the problematic type was also added.

Signed-off-by: Elizabeth Andrews <[email protected]>

Author: elizabethandrews

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -11066,12 +11066,14 @@ def err_ext_int_max_size : Error<"%select{signed|unsigned}0 _ExtInt of bit "
 def err_esimd_glob_cant_init : Error<
   "SYCL explicit SIMD does not permit private global variable to have an initializer">;
 
-def err_sycl_kernel_incorrectly_named : Error<
-  "kernel %select{name is missing"
-  "|needs to have a globally-visible name"
-  "|name is invalid. Unscoped enum requires fixed underlying type"
-  "|name cannot be a type in the \"std\" namespace"
+def err_sycl_kernel_incorrectly_named : Error<"%0 is an invalid kernel name type">;
+def note_invalid_type_in_sycl_kernel : Note<
+  "%select{%1 should be globally-visible"
+  "|unscoped enum %1 requires fixed underlying type"
+  "|type %1 cannot be in the \"std\" namespace"
+  "|kernel name is missing"
   "}0">;
+
 def err_sycl_kernel_not_function_object
     : Error<"kernel parameter must be a lambda or function object">;
 def err_sycl_restrict : Error<

clang/lib/Sema/SemaSYCL.cpp

@@ -2833,6 +2833,7 @@ class SYCLKernelNameTypeVisitor
       public ConstTemplateArgumentVisitor<SYCLKernelNameTypeVisitor> {
   Sema &S;
   SourceLocation KernelInvocationFuncLoc;
+  QualType KernelNameType;
   using InnerTypeVisitor = TypeVisitor<SYCLKernelNameTypeVisitor>;
   using InnerTemplArgVisitor =
       ConstTemplateArgumentVisitor<SYCLKernelNameTypeVisitor>;
@@ -2844,8 +2845,10 @@ class SYCLKernelNameTypeVisitor
   }
 
 public:
-  SYCLKernelNameTypeVisitor(Sema &S, SourceLocation KernelInvocationFuncLoc)
-      : S(S), KernelInvocationFuncLoc(KernelInvocationFuncLoc) {}
+  SYCLKernelNameTypeVisitor(Sema &S, SourceLocation KernelInvocationFuncLoc,
+                            QualType KernelNameType)
+      : S(S), KernelInvocationFuncLoc(KernelInvocationFuncLoc),
+        KernelNameType(KernelNameType) {}
 
   bool isValid() { return !IsInvalid; }
 
@@ -2856,7 +2859,9 @@ class SYCLKernelNameTypeVisitor
     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;
+            << KernelNameType;
+        S.Diag(KernelInvocationFuncLoc, diag::note_invalid_type_in_sycl_kernel)
+            << /* kernel name cannot be a type in the std namespace */ 2 << T;
         IsInvalid = true;
       }
       return;
@@ -2881,9 +2886,10 @@ class SYCLKernelNameTypeVisitor
     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;
+          << KernelNameType;
+      S.Diag(KernelInvocationFuncLoc, diag::note_invalid_type_in_sycl_kernel)
+          << /* Unscoped enum requires fixed underlying type */ 1
+          << QualType(ED->getTypeForDecl(), 0);
       IsInvalid = true;
     }
   }
@@ -2900,7 +2906,10 @@ class SYCLKernelNameTypeVisitor
       auto *NameSpace = dyn_cast_or_null<NamespaceDecl>(DeclCtx);
       if (NameSpace && NameSpace->isStdNamespace()) {
         S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named)
-            << /* kernel name cannot be a type in the std namespace */ 3;
+            << KernelNameType;
+        S.Diag(KernelInvocationFuncLoc, diag::note_invalid_type_in_sycl_kernel)
+            << /* kernel name cannot be a type in the std namespace */ 2
+            << QualType(Tag->getTypeForDecl(), 0);
         IsInvalid = true;
         return;
       }
@@ -2909,19 +2918,27 @@ class SYCLKernelNameTypeVisitor
         if (KernelNameIsMissing) {
           S.Diag(KernelInvocationFuncLoc,
                  diag::err_sycl_kernel_incorrectly_named)
-              << /* kernel name is missing */ 0;
+              << KernelNameType;
+          S.Diag(KernelInvocationFuncLoc,
+                 diag::note_invalid_type_in_sycl_kernel)
+              << /* kernel name is missing */ 3;
           IsInvalid = true;
           return;
         }
         if (Tag->isCompleteDefinition()) {
           S.Diag(KernelInvocationFuncLoc,
                  diag::err_sycl_kernel_incorrectly_named)
-              << /* kernel name is not globally-visible */ 1;
+              << KernelNameType;
+          S.Diag(KernelInvocationFuncLoc,
+                 diag::note_invalid_type_in_sycl_kernel)
+              << /* kernel name is not globally-visible */ 0
+              << QualType(Tag->getTypeForDecl(), 0);
           IsInvalid = true;
-        } else
+        } else {
           S.Diag(KernelInvocationFuncLoc, diag::warn_sycl_implicit_decl);
-        S.Diag(Tag->getSourceRange().getBegin(), diag::note_previous_decl)
-            << Tag->getName();
+          S.Diag(Tag->getSourceRange().getBegin(), diag::note_previous_decl)
+              << Tag->getName();
+        }
       }
     }
   }
@@ -2998,7 +3015,8 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
   SyclKernelArgsSizeChecker ArgsSizeChecker(*this, Args[0]->getExprLoc());
 
   KernelObjVisitor Visitor{*this};
-  SYCLKernelNameTypeVisitor KernelNameTypeVisitor(*this, Args[0]->getExprLoc());
+  SYCLKernelNameTypeVisitor KernelNameTypeVisitor(*this, Args[0]->getExprLoc(),
+                                                  KernelNameType);
 
   DiagnosingSYCLKernel = true;
 

clang/test/SemaSYCL/implicit_kernel_type.cpp

@@ -25,12 +25,12 @@ int main() {
   queue q;
 
 #if defined(WARN)
-  // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
-  // expected-note@+7 {{InvalidKernelName1 declared here}}
+  // expected-error@Inputs/sycl.hpp:220 {{'InvalidKernelName1' is an invalid kernel name type}}
+  // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName1' should be globally-visible}}
   // expected-note@+8 {{in instantiation of function template specialization}}
 #elif defined(ERROR)
-  // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
-  // expected-note@+3 {{InvalidKernelName1 declared here}}
+  // expected-error@Inputs/sycl.hpp:220 {{'InvalidKernelName1' is an invalid kernel name type}}
+  // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName1' should be globally-visible}}
   // expected-note@+4 {{in instantiation of function template specialization}}
 #endif
   class InvalidKernelName1 {};

clang/test/SemaSYCL/kernelname-enum.cpp

@@ -1,14 +1,12 @@
 // 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 2 {{'unscoped_enum_no_type_set' declared here}}
 enum unscoped_enum_no_type_set {
   val_3,
   val_4
@@ -69,11 +67,15 @@ int main() {
   });
 
   q.submit([&](cl::sycl::handler &cgh) {
+    // expected-error@Inputs/sycl.hpp:220 {{'dummy_functor_2<val_3>' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}}
     // expected-note@+1{{in instantiation of function template specialization}}
     cgh.single_task(f2);
   });
 
   q.submit([&](cl::sycl::handler &cgh) {
+    // expected-error@Inputs/sycl.hpp:220 {{'templated_functor<dummy_functor_2>' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{unscoped enum 'unscoped_enum_no_type_set' requires fixed underlying type}}
     // expected-note@+1{{in instantiation of function template specialization}}
     cgh.single_task(f5);
   });

clang/test/SemaSYCL/stdtypes_kernel_type.cpp

@@ -8,36 +8,55 @@ typedef long int ptrdiff_t;
 typedef decltype(nullptr) nullptr_t;
 class T;
 class U;
+class Foo;
 } // namespace std
 
 template <typename T>
 struct Templated_kernel_name;
 
+template <typename T>
+struct Templated_kernel_name2;
+
 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-error@Inputs/sycl.hpp:220 {{'nullptr_t' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{type 'nullptr_t' cannot be in the "std" namespace}}
     // expected-note@+1{{in instantiation of function template specialization}}
     h.single_task<std::nullptr_t>([=] {});
   });
   q.submit([&](handler &h) {
+    // expected-error@Inputs/sycl.hpp:220 {{'std::T' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{type 'std::T' cannot be in the "std" namespace}}
     // expected-note@+1{{in instantiation of function template specialization}}
     h.single_task<std::T>([=] {});
   });
   q.submit([&](handler &h) {
+    // expected-error@Inputs/sycl.hpp:220 {{'Templated_kernel_name<nullptr_t>' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{type 'nullptr_t' cannot be in the "std" namespace}}
     // expected-note@+1{{in instantiation of function template specialization}}
     h.single_task<Templated_kernel_name<std::nullptr_t>>([=] {});
   });
   q.submit([&](handler &h) {
+    // expected-error@Inputs/sycl.hpp:220 {{'Templated_kernel_name<std::U>' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{type 'std::U' cannot be in the "std" namespace}}
     // expected-note@+1{{in instantiation of function template specialization}}
     h.single_task<Templated_kernel_name<std::U>>([=] {});
   });
   q.submit([&](handler &cgh) {
+    // expected-error@Inputs/sycl.hpp:220 {{'Templated_kernel_name2<Templated_kernel_name<std::Foo>>' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220{{type 'std::Foo' cannot be in the "std" namespace}}
+    // expected-note@+1{{in instantiation of function template specialization}}
+    cgh.single_task<Templated_kernel_name2<Templated_kernel_name<std::Foo>>>([]() {});
+  });
+  q.submit([&](handler &cgh) {
+    // expected-error@Inputs/sycl.hpp:220 {{'TemplParamPack<int, float, nullptr_t, double>' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{type 'nullptr_t' cannot be in the "std" namespace}}
     // expected-note@+1{{in instantiation of function template specialization}}
     cgh.single_task<TemplParamPack<int, float, std::nullptr_t, double>>([]() {});
   });

clang/test/SemaSYCL/unnamed-kernel.cpp

@@ -16,6 +16,12 @@ typedef struct {
 } max_align_t;
 } // namespace std
 
+template <typename T>
+struct Templated_kernel_name;
+
+template <typename T>
+struct Templated_kernel_name2;
+
 struct MyWrapper {
 private:
   class InvalidKernelName0 {};
@@ -27,8 +33,8 @@ struct MyWrapper {
   void test() {
     cl::sycl::queue q;
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
-    // expected-note@+3 {{InvalidKernelName1 declared here}}
+    // expected-error@Inputs/sycl.hpp:220 {{'InvalidKernelName1' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName1' should be globally-visible}}
     // expected-note@+4{{in instantiation of function template specialization}}
 #endif
     class InvalidKernelName1 {};
@@ -37,8 +43,8 @@ struct MyWrapper {
     });
 
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
-    // expected-note@+3 {{InvalidKernelName2 declared here}}
+    // expected-error@Inputs/sycl.hpp:220 {{'namespace1::KernelName<InvalidKernelName2>' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName2' should be globally-visible}}
     // expected-note@+4{{in instantiation of function template specialization}}
 #endif
     class InvalidKernelName2 {};
@@ -47,17 +53,17 @@ struct MyWrapper {
     });
 
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
-    // expected-note@21 {{InvalidKernelName0 declared here}}
+    // expected-error@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName0' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName0' should be globally-visible}}
     // 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@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
-    // expected-note@22 {{InvalidKernelName3 declared here}}
+    // expected-error@Inputs/sycl.hpp:220 {{'namespace1::KernelName<MyWrapper::InvalidKernelName3>' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName3' should be globally-visible}}
     // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
@@ -70,7 +76,8 @@ struct MyWrapper {
     });
 
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@Inputs/sycl.hpp:220 {{kernel name cannot be a type in the "std" namespace}}
+    // expected-error@Inputs/sycl.hpp:220 {{'std::max_align_t' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{type 'std::max_align_t' cannot be in the "std" namespace}}
     // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
@@ -79,8 +86,8 @@ struct MyWrapper {
 
     using InvalidAlias = InvalidKernelName4;
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
-    // expected-note@23 {{InvalidKernelName4 declared here}}
+    // expected-error@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName4' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName4' should be globally-visible}}
     // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
@@ -89,20 +96,29 @@ struct MyWrapper {
 
     using InvalidAlias1 = InvalidKernelName5;
 #ifndef __SYCL_UNNAMED_LAMBDA__
-    // expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
-    // expected-note@24 {{InvalidKernelName5 declared here}}
+    // expected-error@Inputs/sycl.hpp:220 {{'namespace1::KernelName<MyWrapper::InvalidKernelName5>' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName5' should be globally-visible}}
     // expected-note@+3{{in instantiation of function template specialization}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<namespace1::KernelName<InvalidAlias1>>([] {});
     });
+#ifndef __SYCL_UNNAMED_LAMBDA__
+    // expected-error@Inputs/sycl.hpp:220 {{'Templated_kernel_name2<Templated_kernel_name<InvalidKernelName1>>' is an invalid kernel name type}}
+    // expected-note@Inputs/sycl.hpp:220 {{'InvalidKernelName1' should be globally-visible}}
+    // expected-note@+3{{in instantiation of function template specialization}}
+#endif
+    q.submit([&](cl::sycl::handler &h) {
+      h.single_task<Templated_kernel_name2<Templated_kernel_name<InvalidKernelName1>>>([] {});
+    });
   }
 };
 
 int main() {
   cl::sycl::queue q;
 #ifndef __SYCL_UNNAMED_LAMBDA__
-  // expected-error@Inputs/sycl.hpp:220 {{kernel name is missing}}
+  // expected-error-re@Inputs/sycl.hpp:220 {{'(lambda at {{.*}}unnamed-kernel.cpp{{.*}}' is an invalid kernel name type}}
+  // expected-note@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([] {}); });