[SYCL] Improve error diagnostic for invalid SYCL kernel name (#4867)

According to SYCL 2020 spec:
The kernel name must be forward declarable at namespace scope (including global namespace scope) and may not be forward declared other than at namespace scope.

The use of nested struct-case (where the name IS globally visible but just not forward-declarable) in a SYCL kernel name currently generates invalid error message:

error: '{{.*}}' should be globally visible

This patch modifies the error diagnostic message when the kernel name is declared at non-namespace scope (i.e. Inside a function or class/struct).

Signed-off-by: Soumi Manna soumi.manna@intel.com

Author: smanna12

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -11451,7 +11451,8 @@ def err_invalid_std_type_in_sycl_kernel : Error<"%0 is an invalid kernel name, "
                                           "%q1 is declared in the 'std' namespace ">;
 
 def err_sycl_kernel_incorrectly_named : Error<
-  "%select{%1 should be globally visible"
+  "%select{%1 is invalid; kernel name should be forward declarable "
+  "at namespace scope"
   "|unscoped enum %1 requires fixed underlying type"
   "|unnamed type %1 is invalid; provide a kernel name, or use "
   "'-fsycl-unnamed-lambda' to enable unnamed kernel lambdas"

clang/lib/Sema/SemaSYCL.cpp

@@ -3508,13 +3508,12 @@ class SYCLKernelNameTypeVisitor
           }
           // Check if the declaration is completely defined within a
           // function or class/struct.
-
           if (Tag->isCompleteDefinition()) {
             S.Diag(KernelInvocationFuncLoc,
                    diag::err_sycl_kernel_incorrectly_named)
-                << /* kernel name should be globally visible */ 0
-                << KernelNameType;
-
+                << /* kernel name should be forward declarable at namespace
+                      scope */
+                0 << KernelNameType;
             IsInvalid = true;
           } else {
             S.Diag(KernelInvocationFuncLoc, diag::warn_sycl_implicit_decl);

clang/test/SemaSYCL/implicit_kernel_type.cpp

@@ -22,14 +22,9 @@ class myWrapper2;
 int main() {
   queue q;
 
-#if defined(WARN)
-  // expected-error@#KernelSingleTask {{'InvalidKernelName1' should be globally visible}}
-  // expected-note@+7 {{in instantiation of function template specialization}}
-#elif defined(ERROR)
-  // expected-error@#KernelSingleTask {{'InvalidKernelName1' should be globally visible}}
-  // expected-note@+4 {{in instantiation of function template specialization}}
-#endif
   class InvalidKernelName1 {};
+  // expected-error@#KernelSingleTask {{'InvalidKernelName1' is invalid; kernel name should be forward declarable at namespace scope}}
+  // expected-note@+2 {{in instantiation of function template specialization}}
   q.submit([&](handler &h) {
     h.single_task<InvalidKernelName1>([]() {});
   });

clang/test/SemaSYCL/nested-anon-and-std-ns.cpp

@@ -37,7 +37,7 @@ struct MyWrapper {
       h.single_task<ValidNS::StructinValidNS>([] {});
     });
 
-    // expected-error@#KernelSingleTask {{'ParentStruct::ChildStruct' should be globally visible}}
+    // expected-error@#KernelSingleTask {{'ParentStruct::ChildStruct' is invalid; kernel name should be forward declarable at namespace scope}}
     // expected-note@+2{{in instantiation of function template specialization}}
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<ParentStruct::ChildStruct>([] {});

clang/test/SemaSYCL/unnamed-kernel.cpp

@@ -32,27 +32,27 @@ struct MyWrapper {
 public:
   void test() {
     cl::sycl::queue q;
-    // expected-error@#KernelSingleTask {{'InvalidKernelName1' should be globally visible}}
+    // expected-error@#KernelSingleTask {{'InvalidKernelName1' is invalid; kernel name should be forward declarable at namespace scope}}
     // expected-note@+3{{in instantiation of function template specialization}}
     class InvalidKernelName1 {};
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<InvalidKernelName1>([] {});
     });
 
-    // expected-error@#KernelSingleTask {{'namespace1::KernelName<InvalidKernelName2>' should be globally visible}}
+    // expected-error@#KernelSingleTask {{'namespace1::KernelName<InvalidKernelName2>' is invalid; kernel name should be forward declarable at namespace scope}}
     // expected-note@+3{{in instantiation of function template specialization}}
     class InvalidKernelName2 {};
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<namespace1::KernelName<InvalidKernelName2>>([] {});
     });
 
-    // expected-error@#KernelSingleTask {{'MyWrapper::InvalidKernelName0' should be globally visible}}
+    // expected-error@#KernelSingleTask {{'MyWrapper::InvalidKernelName0' is invalid; kernel name should be forward declarable at namespace scope}}
     // expected-note@+2{{in instantiation of function template specialization}}
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<InvalidKernelName0>([] {});
     });
 
-    // expected-error@#KernelSingleTask {{'namespace1::KernelName<MyWrapper::InvalidKernelName3>' should be globally visible}}
+    // expected-error@#KernelSingleTask {{'namespace1::KernelName<MyWrapper::InvalidKernelName3>' is invalid; kernel name should be forward declarable at namespace scope}}
     // expected-note@+2{{in instantiation of function template specialization}}
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<namespace1::KernelName<InvalidKernelName3>>([] {});
@@ -70,19 +70,19 @@ struct MyWrapper {
     });
 
     using InvalidAlias = InvalidKernelName4;
-    // expected-error@#KernelSingleTask {{'MyWrapper::InvalidKernelName4' should be globally visible}}
+    // expected-error@#KernelSingleTask {{'MyWrapper::InvalidKernelName4' is invalid; kernel name should be forward declarable at namespace scope}}
     // expected-note@+2{{in instantiation of function template specialization}}
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<InvalidAlias>([] {});
     });
 
     using InvalidAlias1 = InvalidKernelName5;
-    // expected-error@#KernelSingleTask {{'namespace1::KernelName<MyWrapper::InvalidKernelName5>' should be globally visible}}
+    // expected-error@#KernelSingleTask {{'namespace1::KernelName<MyWrapper::InvalidKernelName5>' is invalid; kernel name should be forward declarable at namespace scope}}
     // expected-note@+2{{in instantiation of function template specialization}}
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<namespace1::KernelName<InvalidAlias1>>([] {});
     });
-    // expected-error@#KernelSingleTask {{'Templated_kernel_name2<Templated_kernel_name<InvalidKernelName1>>' should be globally visible}}
+    // expected-error@#KernelSingleTask {{'Templated_kernel_name2<Templated_kernel_name<InvalidKernelName1>>' is invalid; kernel name should be forward declarable at namespace scope}}
     // expected-note@+2{{in instantiation of function template specialization}}
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<Templated_kernel_name2<Templated_kernel_name<InvalidKernelName1>>>([] {});