From d9caf4be3ca091f1ec1a2fc829849ab3e9c9517c Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva@intel.com>
Date: Tue, 19 Nov 2024 02:51:46 -0800
Subject: [PATCH] [clang][SYCL] Print canonical types in free function shim
 functions

Somehow even if the outer type was canonical, TypePrinter manages to
skip nns printing in case of nested default template arguments.
See added test case for the example.
---
 clang/lib/Sema/SemaSYCL.cpp                   |  2 ++
 ...ee_function_default_template_arguments.cpp | 32 +++++++++++++++++++
 2 files changed, 34 insertions(+)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 767dde6512b83..a4cf8c20058f8 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -6510,6 +6510,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
     std::string ParmList;
     bool FirstParam = true;
     Policy.SuppressDefaultTemplateArgs = false;
+    Policy.PrintCanonicalTypes = true;
     for (ParmVarDecl *Param : K.SyclKernel->parameters()) {
       if (FirstParam)
         FirstParam = false;
@@ -6518,6 +6519,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
       ParmList += Param->getType().getCanonicalType().getAsString(Policy);
     }
     FunctionTemplateDecl *FTD = K.SyclKernel->getPrimaryTemplate();
+    Policy.PrintCanonicalTypes = false;
     Policy.SuppressDefinition = true;
     Policy.PolishForDeclaration = true;
     Policy.FullyQualifiedName = true;
diff --git a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
index 808f7b93d8112..62a121d218b8b 100644
--- a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
+++ b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
@@ -60,6 +60,24 @@ templated3(Arg<T, notatuple, a, ns1::hasDefaultArg<>, int, int>, T end) {
 
 template void templated3(Arg<int, notatuple, 3, ns1::hasDefaultArg<>, int, int>, int);
 
+
+namespace sycl {
+template <typename T> struct X {};
+template <> struct X<int> {};
+namespace detail {
+struct Y {};
+} // namespace detail
+template <> struct X<detail::Y> {};
+} // namespace sycl
+using namespace sycl;
+template <typename T, typename = X<detail::Y>> struct Arg1 { T val; };
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel",
+                                              2)]] void
+foo(Arg1<int> arg) {
+  arg.val = 42;
+}
+
 // CHECK: Forward declarations of kernel and its argument types:
 // CHECK-NEXT: namespace ns { 
 // CHECK-NEXT: struct notatuple;
@@ -98,3 +116,17 @@ template void templated3(Arg<int, notatuple, 3, ns1::hasDefaultArg<>, int, int>,
 // CHECK-NEXT: static constexpr auto __sycl_shim5() {
 // CHECK-NEXT:   return (void (*)(struct ns::Arg<int, struct ns::notatuple, 3, class ns::ns1::hasDefaultArg<struct ns::notatuple>, int, int>, int))templated3<int, 3>;
 // CHECK-NEXT: }
+
+// CHECK Forward declarations of kernel and its argument types:
+// CHECK: namespace sycl { namespace detail {
+// CHECK-NEXT: struct Y;
+// CHECK-NEXT: }}
+// CHECK-NEXT: namespace sycl {
+// CHECK-NEXT: template <typename T> struct X;
+// CHECK-NEXT: }
+// CHECK-NEXT: template <typename T, typename> struct Arg1;
+
+// CHECK: void foo(Arg1<int, sycl::X<sycl::detail::Y> > arg);
+// CHECK-NEXT: static constexpr auto __sycl_shim6() {
+// CHECK-NEXT:   return (void (*)(struct Arg1<int, struct sycl::X<struct sycl::detail::Y> >))foo;
+// CHECK-NEXT: }