From bee7955da607ca879f02abcd4f94a556d1afd885 Mon Sep 17 00:00:00 2001
From: Erich Keane <erich.keane@intel.com>
Date: Fri, 6 Aug 2021 08:36:44 -0700
Subject: [PATCH] [SYCL] Ensure int-header fwd decls work with references

We weren't properly descending down into reference types, so we ended up
not forward declaring any types that were referred to only as a
reference type.  This patch ensures we do so, just like we did with
pointer-type.
---
 clang/lib/Sema/SemaSYCL.cpp                   |  6 ++++
 .../int_header_reference_fwd_decl.cpp         | 33 +++++++++++++++++++
 2 files changed, 39 insertions(+)
 create mode 100644 clang/test/CodeGenSYCL/int_header_reference_fwd_decl.cpp

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>([]{});
+}