[SYCL] Diagnose implicit declaration of kernel function type (#1450)

SYCL 1.2.1 specification requires an explicit forward declaration of
a kernel type name.  The current implementation implicitly forward
declares the name if the user fails to do so.  Diagnose this in strict
sycl modes, as such user code is not portable.

Signed-off-by: Premanand M Rao <premanand.m.rao@intel.com>

Author: premanandrao

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10719,6 +10719,11 @@ def warn_sycl_attibute_function_raw_ptr
               "to a function with a raw pointer "
               "%select{return type|parameter type}1">,
       InGroup<SyclStrict>, DefaultError;
+def warn_sycl_implicit_decl
+    : Warning<"SYCL 1.2.1 specification requires an explicit forward "
+              "declaration for a kernel type name; your program may not "
+              "be portable">,
+      InGroup<SyclStrict>, DefaultIgnore;
 def err_ivdep_duplicate_arg : Error<
   "duplicate argument to 'ivdep'. attribute requires one or both of a safelen "
   "and array">;

clang/lib/Sema/SemaSYCL.cpp

@@ -1503,7 +1503,7 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
                                 ? cast<ClassTemplateDecl>(D)->getTemplatedDecl()
                                 : dyn_cast<TagDecl>(D);
 
-        if (TD && TD->isCompleteDefinition() && !UnnamedLambdaSupport) {
+        if (TD && !UnnamedLambdaSupport) {
           // defined class constituting the kernel name is not globally
           // accessible - contradicts the spec
           const bool KernelNameIsMissing = TD->getName().empty();
@@ -1512,8 +1512,12 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
                 << /* kernel name is missing */ 0;
             // Don't emit note if kernel name was completely omitted
           } else {
-            Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
-                << /* kernel name is not globally-visible */ 1;
+            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();

clang/test/SemaSYCL/implicit_kernel_type.cpp

@@ -0,0 +1,53 @@
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s -Werror=sycl-strict -DERROR
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s -Wsycl-strict -DWARN
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -verify %s -Werror=sycl-strict
+#include <sycl.hpp>
+
+#ifdef __SYCL_UNNAMED_LAMBDA__
+// expected-no-diagnostics
+#endif
+
+using namespace cl::sycl;
+
+void function() {
+}
+
+// user-defined class
+struct myWrapper {
+};
+
+// user-declared class
+class myWrapper2;
+
+int main() {
+  cl::sycl::queue q;
+#ifndef __SYCL_UNNAMED_LAMBDA__
+  // expected-note@+1 {{InvalidKernelName1 declared here}}
+  class InvalidKernelName1 {};
+  q.submit([&](cl::sycl::handler &h) {
+    // expected-error@+1 {{kernel needs to have a globally-visible name}}
+    h.single_task<InvalidKernelName1>([]() {});
+  });
+#endif
+#if defined(WARN)
+  // expected-warning@+6 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
+  // expected-note@+5 {{fake_kernel declared here}}
+#elif defined(ERROR)
+  // expected-error@+3 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
+  // expected-note@+2 {{fake_kernel declared here}}
+#endif
+  cl::sycl::kernel_single_task<class fake_kernel>([]() { function(); });
+#if defined(WARN)
+  // expected-warning@+6 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
+  // expected-note@+5 {{fake_kernel2 declared here}}
+#elif defined(ERROR)
+  // expected-error@+3 {{SYCL 1.2.1 specification requires an explicit forward declaration for a kernel type name; your program may not be portable}}
+  // expected-note@+2 {{fake_kernel2 declared here}}
+#endif
+  cl::sycl::kernel_single_task<class fake_kernel2>([]() {
+    auto l = [](auto f) { f(); };
+  });
+  cl::sycl::kernel_single_task<class myWrapper>([]() { function(); });
+  cl::sycl::kernel_single_task<class myWrapper2>([]() { function(); });
+  return 0;
+}