diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 23bfd9eed4e30..3b03f317afdf4 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -10696,6 +10696,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">;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 3724984e71b66..2a4ed4c183f89 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1534,7 +1534,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();
@@ -1543,8 +1543,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();
diff --git a/clang/test/SemaSYCL/implicit_kernel_type.cpp b/clang/test/SemaSYCL/implicit_kernel_type.cpp
new file mode 100644
index 0000000000000..97a0ca90a2bc9
--- /dev/null
+++ b/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;
+}