diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 1bcf48f2b06e2..ea7d983d77bd2 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -4459,6 +4459,12 @@ class SYCLFwdDeclEmitter
     InnerTypeVisitor::Visit(T.getTypePtr());
   }
 
+  void VisitReferenceType(const ReferenceType *RT) {
+    // Our forward declarations don't care about references, so we should just
+    // ignore the reference and continue on.
+    Visit(RT->getPointeeType());
+  }
+
   void Visit(const TemplateArgument &TA) {
     if (TA.isNull())
       return;
diff --git a/clang/test/CodeGenSYCL/int_header_reference_fwd_decl.cpp b/clang/test/CodeGenSYCL/int_header_reference_fwd_decl.cpp
new file mode 100644
index 0000000000000..42804a74e2c09
--- /dev/null
+++ b/clang/test/CodeGenSYCL/int_header_reference_fwd_decl.cpp
@@ -0,0 +1,33 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -fsycl-int-header=%t.h %s -o %t.out
+// RUN: FileCheck -input-file=%t.h %s
+
+// This test validates that we handle reference types properly in forward
+// declarations, the bug was that the PassedAsRefType didn't get forward
+// declared, so it resulted in a compile error during host compile. SO, we make
+// sure that we properly forward declare the type (and any template children!).
+
+#include <sycl.hpp>
+
+// CHECK: // Forward declarations of templated kernel function types:
+// CHECK-NEXT: namespace WrapsType {
+// CHECK-NEXT: struct InsidePassedAsRef;
+// CHECK-NEXT: }
+// CHECK: namespace WrapsType {
+// CHECK-NEXT: template <typename T> struct PassedAsRef;
+// CHECK-NEXT: }
+// CHECK: template <typename T> class Wrapper;
+
+namespace WrapsType {
+  struct InsidePassedAsRef{};
+  template<typename T>
+  struct PassedAsRef{};
+}
+
+template<typename T>
+class Wrapper{};
+
+void foo() {
+  using namespace WrapsType;
+  using KernelName = Wrapper<PassedAsRef<InsidePassedAsRef>&>;
+  sycl::kernel_single_task<KernelName>([]{});
+}