From 549ca8fc684b2e0efcbde983b1d1b8be78bd6d80 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Thu, 11 Dec 2025 05:43:04 -0800
Subject: [PATCH 1/6] [SYCL][clang] Fix more free-function kernel integration
 header cases

This commit fixes a number of known issues when the integration header
generates prototypes of the free-function kernels. These changes focus
on the additional removal of aliasing and proper handling of templated
template arguments.

This commit also adds disabled test cases for a known issue with
unresolved nested templated type aliases. These are cases for future
work.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Co-authored-by: Sachkov, Alexey <alexey.sachkov@intel.com>
---
 clang/include/clang/AST/TypeBase.h            |   3 +
 clang/lib/AST/Type.cpp                        |  35 +++
 clang/lib/Sema/SemaSYCL.cpp                   | 279 ++++++++++++++----
 .../free-function-kernel-type-alias-arg.cpp   |  80 ++++-
 4 files changed, 335 insertions(+), 62 deletions(-)

diff --git a/clang/include/clang/AST/TypeBase.h b/clang/include/clang/AST/TypeBase.h
index c6bf312268647..899b9939b33d2 100644
--- a/clang/include/clang/AST/TypeBase.h
+++ b/clang/include/clang/AST/TypeBase.h
@@ -2956,6 +2956,9 @@ class alignas(TypeAlignment) Type : public ExtQualsTypeCommonBase {
     return TST;
   }
 
+  const TemplateSpecializationType *
+  getAsTemplateSpecializationTypeWithoutAliases(const ASTContext &Ctx) const;
+
   /// Member-template getAsAdjusted<specific type>. Look through specific kinds
   /// of sugar (parens, attributes, etc) for an instance of \<specific type>.
   /// This is used when you need to walk over sugar nodes that represent some
diff --git a/clang/lib/AST/Type.cpp b/clang/lib/AST/Type.cpp
index 8a0b11bfcedac..76c2b170835f2 100644
--- a/clang/lib/AST/Type.cpp
+++ b/clang/lib/AST/Type.cpp
@@ -1925,6 +1925,41 @@ Type::getAsNonAliasTemplateSpecializationType() const {
   return TST;
 }
 
+const TemplateSpecializationType *
+Type::getAsTemplateSpecializationTypeWithoutAliases(
+    const ASTContext &Ctx) const {
+  const TemplateSpecializationType *TST =
+      getAsNonAliasTemplateSpecializationType();
+  if (!TST)
+    return TST;
+
+  // Ensure the template arguments of the template specialization type are
+  // without aliases.
+  SmallVector<TemplateArgument, 4> ArgsWithoutAliases;
+  ArgsWithoutAliases.reserve(TST->template_arguments().size());
+  for (const TemplateArgument &TA : TST->template_arguments()) {
+    if (TA.getKind() == TemplateArgument::ArgKind::Type) {
+      QualType TAQTy = TA.getAsType();
+      const Type *TATy = TAQTy->getUnqualifiedDesugaredType();
+      if (isa<TemplateSpecializationType>(TATy))
+        TATy = TATy->getAsTemplateSpecializationTypeWithoutAliases(Ctx);
+      ArgsWithoutAliases.emplace_back(QualType(TATy, TAQTy.getCVRQualifiers()));
+    } else if (TA.getKind() == TemplateArgument::ArgKind::Template) {
+      TemplateName TN = TA.getAsTemplate();
+      while (std::optional<TemplateName> DesugaredTN =
+                 TN.desugar(/*IgnoreDeduced=*/false))
+        TN = *DesugaredTN;
+      ArgsWithoutAliases.emplace_back(TN);
+    } else {
+      ArgsWithoutAliases.push_back(TA);
+    }
+  }
+  return Ctx
+      .getTemplateSpecializationType(TST->getKeyword(), TST->getTemplateName(),
+                                     ArgsWithoutAliases, {}, QualType{})
+      ->getAs<TemplateSpecializationType>();
+}
+
 NestedNameSpecifier Type::getPrefix() const {
   switch (getTypeClass()) {
   case Type::DependentName:
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 639e3a2d16cc7..923cb2a7ccd6c 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -6626,6 +6626,195 @@ static void PrintNSClosingBraces(raw_ostream &OS, const DeclContext *DC) {
       [](raw_ostream &, const NamespaceDecl *) {}, OS, DC);
 }
 
+/// Dedicated visitor which helps with printing of kernel arguments in forward
+/// declarations of free function kernels which are declared as function
+/// templates.
+///
+/// Based on:
+/// \code
+/// template <typename T1, typename T2>
+/// void foo(T1 a, int b, T2 c);
+/// \endcode
+///
+/// It prints into the output stream "T1, int, T2".
+///
+/// The main complexity (which motivates addition of such visitor) comes from
+/// the fact that there could be type aliases and default template arguments.
+/// For example:
+/// \code
+/// template<typename T>
+/// void kernel(sycl::accessor<T, 1>);
+/// template void kernel(sycl::accessor<int, 1>);
+/// \endcode
+/// sycl::accessor has many template arguments which have default values. If
+/// we iterate over non-canonicalized argument type, we don't get those default
+/// values and we don't get necessary namespace qualifiers for all the template
+/// arguments. If we iterate over canonicalized argument type, then all
+/// references to T will be replaced with something like type-argument-X-Y.
+/// What this visitor does is it iterates over both in sync, picking the right
+/// values from one or another.
+///
+/// Moral of the story: drop integration header ASAP (but that is blocked
+/// by support for 3rd-party host compilers, which is important).
+class FreeFunctionTemplateKernelArgsPrinter
+    : public ConstTemplateArgumentVisitor<FreeFunctionTemplateKernelArgsPrinter,
+                                          void, ArrayRef<TemplateArgument>> {
+  raw_ostream &O;
+  PrintingPolicy &Policy;
+  ASTContext &Context;
+
+  using Base =
+      ConstTemplateArgumentVisitor<FreeFunctionTemplateKernelArgsPrinter, void,
+                                   ArrayRef<TemplateArgument>>;
+
+  void PrintTemplateDeclName(const TemplateDecl *TD,
+                             ArrayRef<TemplateArgument> SpecArgs) {}
+
+public:
+  FreeFunctionTemplateKernelArgsPrinter(raw_ostream &O, PrintingPolicy &Policy,
+                                        ASTContext &Context)
+      : O(O), Policy(Policy), Context(Context) {}
+
+  void Visit(const TemplateSpecializationType *T,
+             const TemplateSpecializationType *CT) {
+    ArrayRef<TemplateArgument> SpecArgs = T->template_arguments();
+    ArrayRef<TemplateArgument> DeclArgs = CT->template_arguments();
+
+    const TemplateDecl *TD = CT->getTemplateName().getAsTemplateDecl();
+    if (!TD->getIdentifier())
+      TD = T->getTemplateName().getAsTemplateDecl();
+    TD->printQualifiedName(O);
+
+    O << "<";
+    for (size_t I = 0, E = std::max(DeclArgs.size(), SpecArgs.size()),
+                SE = SpecArgs.size();
+         I < E; ++I) {
+      if (I != 0)
+        O << ", ";
+      // If we have a specialized argument, use it. Otherwise fallback to a
+      // default argument.
+      // We pass specialized arguments in case there are references to them
+      // from other types.
+      // FIXME: passing SpecArgs here is incorrect. It refers to template
+      // arguments of a single function argument, but DeclArgs contain
+      // references (in form of depth-index) to template arguments of the
+      // function itself which results in incorrect integration header being
+      // produced.
+      Base::Visit(I < SE ? SpecArgs[I] : DeclArgs[I], SpecArgs);
+    }
+    O << ">";
+  }
+
+  // Internal version of the function above that is used when template argument
+  // is a template by itself
+  void Visit(const TemplateSpecializationType *T,
+             ArrayRef<TemplateArgument> SpecArgs) {
+    const TemplateDecl *TD = T->getTemplateName().getAsTemplateDecl();
+    const auto *TTPD = dyn_cast<TemplateTemplateParmDecl>(TD);
+    if (TTPD && !TTPD->getIdentifier())
+      SpecArgs[TTPD->getIndex()].print(Policy, O, /* IncludeType = */ false);
+    else
+      TD->printQualifiedName(O);
+    O << "<";
+    ArrayRef<const TemplateArgument> DeclArgs = T->template_arguments();
+    for (size_t I = 0, E = DeclArgs.size(); I < E; ++I) {
+      if (I != 0)
+        O << ", ";
+      Base::Visit(DeclArgs[I], SpecArgs);
+    }
+    O << ">";
+  }
+
+  void VisitNullTemplateArgument(const TemplateArgument &,
+                                 ArrayRef<TemplateArgument>) {
+    llvm_unreachable("If template argument has not been deduced, then we can't "
+                     "forward-declare it, something went wrong");
+  }
+
+  void VisitTypeTemplateArgument(const TemplateArgument &Arg,
+                                 ArrayRef<TemplateArgument> SpecArgs) {
+    // If we reference an existing template argument without a known identifier,
+    // print it instead.
+    const auto *TPT = dyn_cast<TemplateTypeParmType>(Arg.getAsType());
+    if (TPT && !TPT->getIdentifier()) {
+      SpecArgs[TPT->getIndex()].print(Policy, O, /* IncludeType = */ false);
+      return;
+    }
+
+    const auto *TST = dyn_cast<TemplateSpecializationType>(Arg.getAsType());
+    if (TST && Arg.isInstantiationDependent()) {
+      // This is an instantiation dependent template specialization, meaning
+      // that some of its arguments reference template arguments of the free
+      // function kernel itself.
+      Visit(TST, SpecArgs);
+      return;
+    }
+
+    Arg.print(Policy, O, /* IncludeType = */ false);
+  }
+
+  void VisitDeclarationTemplateArgument(const TemplateArgument &,
+                                        ArrayRef<TemplateArgument>) {
+    llvm_unreachable("Free function kernels cannot have non-type template "
+                     "arguments which are pointers or references");
+  }
+
+  void VisitNullPtrTemplateArgument(const TemplateArgument &,
+                                    ArrayRef<TemplateArgument>) {
+    llvm_unreachable("Free function kernels cannot have non-type template "
+                     "arguments which are pointers or references");
+  }
+
+  void VisitIntegralTemplateArgument(const TemplateArgument &Arg,
+                                     ArrayRef<TemplateArgument>) {
+    Arg.print(Policy, O, /* IncludeType = */ false);
+  }
+
+  void VisitStructuralValueTemplateArgument(const TemplateArgument &Arg,
+                                            ArrayRef<TemplateArgument>) {
+    Arg.print(Policy, O, /* IncludeType = */ false);
+  }
+
+  void VisitTemplateTemplateArgument(const TemplateArgument &Arg,
+                                     ArrayRef<TemplateArgument>) {
+    Arg.print(Policy, O, /* IncludeType = */ false);
+  }
+
+  void VisitTemplateExpansionTemplateArgument(const TemplateArgument &Arg,
+                                              ArrayRef<TemplateArgument>) {
+    // Likely does not work similar to the one above
+    Arg.print(Policy, O, /* IncludeType = */ false);
+  }
+
+  void VisitExpressionTemplateArgument(const TemplateArgument &Arg,
+                                       ArrayRef<TemplateArgument>) {
+    Expr *E = Arg.getAsExpr();
+    assert(E && "Failed to get an Expr for an Expression template arg?");
+
+    if (Arg.isInstantiationDependent() ||
+        E->getType().getTypePtr()->isScopedEnumeralType()) {
+      // Scoped enumerations can't be implicitly cast from integers, so
+      // we don't need to evaluate them.
+      // If expression is instantiation-dependent, then we can't evaluate it
+      // either, let's fallback to default printing mechanism.
+      Arg.print(Policy, O, /* IncludeType = */ false);
+      return;
+    }
+
+    Expr::EvalResult Res;
+    [[maybe_unused]] bool Success =
+        Arg.getAsExpr()->EvaluateAsConstantExpr(Res, Context);
+    assert(Success && "invalid non-type template argument?");
+    assert(!Res.Val.isAbsent() && "couldn't read the evaulation result?");
+    Res.Val.printPretty(O, Policy, Arg.getAsExpr()->getType(), &Context);
+  }
+
+  void VisitPackTemplateArgument(const TemplateArgument &Arg,
+                                 ArrayRef<TemplateArgument>) {
+    Arg.print(Policy, O, /* IncludeType = */ false);
+  }
+};
+
 class FreeFunctionPrinter {
   raw_ostream &O;
   PrintingPolicy &Policy;
@@ -6789,7 +6978,10 @@ class FreeFunctionPrinter {
     llvm::raw_svector_ostream ParmListOstream{ParamList};
     Policy.SuppressTagKeyword = true;
 
-    for (ParmVarDecl *Param : Parameters) {
+    FreeFunctionTemplateKernelArgsPrinter Printer(ParmListOstream, Policy,
+                                                  Context);
+
+    for (const ParmVarDecl *Param : Parameters) {
       if (FirstParam)
         FirstParam = false;
       else
@@ -6822,53 +7014,11 @@ class FreeFunctionPrinter {
       }
 
       const TemplateSpecializationType *TSTAsNonAlias =
-          TST->getAsNonAliasTemplateSpecializationType();
+          TST->getAsTemplateSpecializationTypeWithoutAliases(Context);
       if (TSTAsNonAlias)
         TST = TSTAsNonAlias;
 
-      TemplateName CTN = CTST->getTemplateName();
-      CTN.getAsTemplateDecl()->printQualifiedName(ParmListOstream);
-      ParmListOstream << "<";
-
-      ArrayRef<TemplateArgument> SpecArgs = TST->template_arguments();
-      ArrayRef<TemplateArgument> DeclArgs = CTST->template_arguments();
-
-      auto TemplateArgPrinter = [&](const TemplateArgument &Arg) {
-        if (Arg.getKind() != TemplateArgument::ArgKind::Expression ||
-            Arg.isInstantiationDependent()) {
-          Arg.print(Policy, ParmListOstream, /* IncludeType = */ false);
-          return;
-        }
-
-        Expr *E = Arg.getAsExpr();
-        assert(E && "Failed to get an Expr for an Expression template arg?");
-        if (E->getType().getTypePtr()->isScopedEnumeralType()) {
-          // Scoped enumerations can't be implicitly cast from integers, so
-          // we don't need to evaluate them.
-          Arg.print(Policy, ParmListOstream, /* IncludeType = */ false);
-          return;
-        }
-
-        Expr::EvalResult Res;
-        [[maybe_unused]] bool Success =
-            Arg.getAsExpr()->EvaluateAsConstantExpr(Res, Context);
-        assert(Success && "invalid non-type template argument?");
-        assert(!Res.Val.isAbsent() && "couldn't read the evaulation result?");
-        Res.Val.printPretty(ParmListOstream, Policy, Arg.getAsExpr()->getType(),
-                            &Context);
-      };
-
-      for (size_t I = 0, E = std::max(DeclArgs.size(), SpecArgs.size()),
-                  SE = SpecArgs.size();
-           I < E; ++I) {
-        if (I != 0)
-          ParmListOstream << ", ";
-        // If we have a specialized argument, use it. Otherwise fallback to a
-        // default argument.
-        TemplateArgPrinter(I < SE ? SpecArgs[I] : DeclArgs[I]);
-      }
-
-      ParmListOstream << ">";
+      Printer.Visit(TST, CTST);
     }
     return ParamList.str().str();
   }
@@ -6886,26 +7036,39 @@ class FreeFunctionPrinter {
   std::string getTemplateParameters(const clang::TemplateParameterList *TPL) {
     std::string TemplateParams{"template <"};
     bool FirstParam{true};
-    for (NamedDecl *Param : *TPL) {
+    for (const NamedDecl *Param : *TPL) {
       if (!FirstParam)
         TemplateParams += ", ";
       FirstParam = false;
-      if (const auto *TemplateParam = dyn_cast<TemplateTypeParmDecl>(Param)) {
-        TemplateParams +=
-            TemplateParam->wasDeclaredWithTypename() ? "typename " : "class ";
-        if (TemplateParam->isParameterPack())
-          TemplateParams += "... ";
-        TemplateParams += TemplateParam->getNameAsString();
-      } else if (const auto *NonTypeParam =
-                     dyn_cast<NonTypeTemplateParmDecl>(Param)) {
-        TemplateParams += NonTypeParam->getType().getAsString();
-        TemplateParams += " ";
-        TemplateParams += NonTypeParam->getNameAsString();
-      }
+      TemplateParams += getTemplateParameter(Param);
     }
     TemplateParams += "> ";
     return TemplateParams;
   }
+
+  /// Helper method to get text representation of a template parameter.
+  /// \param Param The template parameter.
+  std::string getTemplateParameter(const NamedDecl *Param) {
+    auto GetTypenameOrClass = [](const auto *Param) {
+      return Param->wasDeclaredWithTypename() ? "typename " : "class ";
+    };
+    if (const auto *TemplateParam = dyn_cast<TemplateTypeParmDecl>(Param)) {
+      std::string TemplateParamStr = GetTypenameOrClass(TemplateParam);
+      if (TemplateParam->isParameterPack())
+        TemplateParamStr += "... ";
+      TemplateParamStr += TemplateParam->getNameAsString();
+      return TemplateParamStr;
+    } else if (const auto *NonTypeParam =
+                   dyn_cast<NonTypeTemplateParmDecl>(Param)) {
+      return NonTypeParam->getType().getAsString() + " " +
+             NonTypeParam->getNameAsString();
+    } else if (const auto *TTParam =
+                   dyn_cast<TemplateTemplateParmDecl>(Param)) {
+      return getTemplateParameters(TTParam->getTemplateParameters()) + " " +
+             GetTypenameOrClass(TTParam) + TTParam->getNameAsString();
+    }
+    return "";
+  }
 };
 
 void SYCLIntegrationHeader::emit(raw_ostream &O) {
diff --git a/clang/test/CodeGenSYCL/free-function-kernel-type-alias-arg.cpp b/clang/test/CodeGenSYCL/free-function-kernel-type-alias-arg.cpp
index 5d6ea216d7d38..d8717eaaab44f 100644
--- a/clang/test/CodeGenSYCL/free-function-kernel-type-alias-arg.cpp
+++ b/clang/test/CodeGenSYCL/free-function-kernel-type-alias-arg.cpp
@@ -14,6 +14,9 @@ typedef int IntTypedef;
 template <typename T>
 struct Foo {};
 
+template <typename T>
+using FooUsing = Foo<T>;
+
 using FooIntUsing = Foo<int>;
 typedef Foo<int> FooIntTypedef;
 
@@ -29,15 +32,35 @@ using BarUsing2 = Bar<Foo<T2>, T1>;
 template <typename T1>
 using BarUsingBarUsing2 = BarUsing2<T1, int>;
 
+template <typename T1>
+using BarUsingFooIntUsing = Bar<FooIntUsing, T1>;
+
+template <typename T1>
+using BarUsingBarUsingFooIntUsing = BarUsingFooIntUsing<T1>;
+
 class Baz {
 public:
   using type = BarUsing<double>;
 };
 
+template <template <typename> class T1>
+class Foz {};
+
+template <template <typename> class T1>
+using FozUsing = Foz<T1>;
+
+using FozFooUsing = Foz<Foo>;
+
 template <typename T1, typename T2, typename T3 = BarUsing2<T1, T2>,
           typename T4 = BarUsing<T2>>
 struct AliasAsDefaultArg {};
 
+template <typename T>
+struct List {};
+
+template <typename T>
+struct Array {};
+
 } // namespace ns
 
 [[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
@@ -81,18 +104,67 @@ template void bar_using_bar_using2(ns::BarUsingBarUsing2<int>);
 
 // CHECK: template <typename T1> void bar_using_bar_using2(ns::Bar<ns::Foo<int>, T1>);
 
+template<typename T>
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void bar_using_foo_int_using(ns::BarUsingFooIntUsing<T> Arg) {}
+template void bar_using_foo_int_using(ns::BarUsingFooIntUsing<float>);
+
+// CHECK: template <typename T> void bar_using_foo_int_using(ns::Bar<ns::Foo<int>, T>);
+
+template<typename T>
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void bar_using_bar_using_foo_int_using(ns::BarUsingBarUsingFooIntUsing<T> Arg) {}
+template void bar_using_bar_using_foo_int_using(ns::BarUsingBarUsingFooIntUsing<float>);
+
+// CHECK: template <typename T> void bar_using_bar_using_foo_int_using(ns::Bar<ns::Foo<int>, T>);
+
 [[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
 void baz_type(ns::Baz::type Arg) {}
 
 // CHECK: void baz_type(ns::Bar<double, float> Arg);
 
-#if 0
-// This test case fails, but it is added here in advance to add a record of a
-// known bug.
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void foz_using(ns::FozUsing<ns::Foo> Arg) {}
+
+// CHECK: void foz_using(ns::Foz<ns::Foo> Arg);
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void foz_foo_using(ns::FozFooUsing Arg) {}
+
+// CHECK: void foz_foo_using(ns::Foz<ns::Foo> Arg);
+
+template<template <typename> class T1, typename T2>
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void bar_using2_foz(ns::BarUsing2<ns::Foz<T1>, T2> Arg) {}
+template void bar_using2_foz(ns::BarUsing2<ns::Foz<ns::Foo>, float>);
+
+// CHECK: template <template <typename > class T1, typename T2> void bar_using2_foz(ns::Bar<ns::Foo<T2>, ns::Foz<T1>>)
+
 template<typename T1, typename T2>
 [[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
 void alias_as_default_template_arg(ns::AliasAsDefaultArg<T1, T2> Arg) {}
 template void alias_as_default_template_arg(ns::AliasAsDefaultArg<int, float>);
 
-// CHECK-DISABLED: template <typename T1, typename T2> void alias_as_default_template_arg(ns::AliasAsDefaultArg<T1, T2, ns::Bar<ns::Foo<T2>, T1>, ns::Bar<T2, float>>);
+// CHECK: template <typename T1, typename T2> void alias_as_default_template_arg(ns::AliasAsDefaultArg<T1, T2, ns::Bar<ns::Foo<T2>, T1>, ns::Bar<T2, float>>);
+
+template<typename T, template <typename> typename Container>
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void template_template_arg(Container<T> Arg) {}
+template void template_template_arg(ns::List<int>);
+
+// CHECK: template <typename T, template <typename > typename Container> void template_template_arg(Container<T>);
+
+// These test cases fail, but they are added here in advance to add a record of
+// known bugs.
+#if 0
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void foz_foo_using(ns::Foz<ns::FooUsing> Arg) {}
+
+// CHECK-DISABLED: void foz_foo_using(ns::Foz<ns::Foo> Arg);
+
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void foz_using_foo_using(ns::FozUsing<ns::FooUsing> Arg) {}
+
+// CHECK-DISABLED: void foz_using_foo_using(ns::Foz<ns::Foo> Arg);
 #endif

From bea8af016fb4071925e9e78ceeb911df39e7bdb0 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Sun, 14 Dec 2025 22:51:10 -0800
Subject: [PATCH 2/6] Remove getAsTemplateSpecializationTypeWithoutAliases

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 clang/include/clang/AST/TypeBase.h |  3 ---
 clang/lib/AST/Type.cpp             | 35 ------------------------------
 clang/lib/Sema/SemaSYCL.cpp        |  2 +-
 3 files changed, 1 insertion(+), 39 deletions(-)

diff --git a/clang/include/clang/AST/TypeBase.h b/clang/include/clang/AST/TypeBase.h
index 899b9939b33d2..c6bf312268647 100644
--- a/clang/include/clang/AST/TypeBase.h
+++ b/clang/include/clang/AST/TypeBase.h
@@ -2956,9 +2956,6 @@ class alignas(TypeAlignment) Type : public ExtQualsTypeCommonBase {
     return TST;
   }
 
-  const TemplateSpecializationType *
-  getAsTemplateSpecializationTypeWithoutAliases(const ASTContext &Ctx) const;
-
   /// Member-template getAsAdjusted<specific type>. Look through specific kinds
   /// of sugar (parens, attributes, etc) for an instance of \<specific type>.
   /// This is used when you need to walk over sugar nodes that represent some
diff --git a/clang/lib/AST/Type.cpp b/clang/lib/AST/Type.cpp
index 76c2b170835f2..8a0b11bfcedac 100644
--- a/clang/lib/AST/Type.cpp
+++ b/clang/lib/AST/Type.cpp
@@ -1925,41 +1925,6 @@ Type::getAsNonAliasTemplateSpecializationType() const {
   return TST;
 }
 
-const TemplateSpecializationType *
-Type::getAsTemplateSpecializationTypeWithoutAliases(
-    const ASTContext &Ctx) const {
-  const TemplateSpecializationType *TST =
-      getAsNonAliasTemplateSpecializationType();
-  if (!TST)
-    return TST;
-
-  // Ensure the template arguments of the template specialization type are
-  // without aliases.
-  SmallVector<TemplateArgument, 4> ArgsWithoutAliases;
-  ArgsWithoutAliases.reserve(TST->template_arguments().size());
-  for (const TemplateArgument &TA : TST->template_arguments()) {
-    if (TA.getKind() == TemplateArgument::ArgKind::Type) {
-      QualType TAQTy = TA.getAsType();
-      const Type *TATy = TAQTy->getUnqualifiedDesugaredType();
-      if (isa<TemplateSpecializationType>(TATy))
-        TATy = TATy->getAsTemplateSpecializationTypeWithoutAliases(Ctx);
-      ArgsWithoutAliases.emplace_back(QualType(TATy, TAQTy.getCVRQualifiers()));
-    } else if (TA.getKind() == TemplateArgument::ArgKind::Template) {
-      TemplateName TN = TA.getAsTemplate();
-      while (std::optional<TemplateName> DesugaredTN =
-                 TN.desugar(/*IgnoreDeduced=*/false))
-        TN = *DesugaredTN;
-      ArgsWithoutAliases.emplace_back(TN);
-    } else {
-      ArgsWithoutAliases.push_back(TA);
-    }
-  }
-  return Ctx
-      .getTemplateSpecializationType(TST->getKeyword(), TST->getTemplateName(),
-                                     ArgsWithoutAliases, {}, QualType{})
-      ->getAs<TemplateSpecializationType>();
-}
-
 NestedNameSpecifier Type::getPrefix() const {
   switch (getTypeClass()) {
   case Type::DependentName:
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 923cb2a7ccd6c..d99dfef4c8327 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -7014,7 +7014,7 @@ class FreeFunctionPrinter {
       }
 
       const TemplateSpecializationType *TSTAsNonAlias =
-          TST->getAsTemplateSpecializationTypeWithoutAliases(Context);
+          TST->getAsNonAliasTemplateSpecializationType();
       if (TSTAsNonAlias)
         TST = TSTAsNonAlias;
 

From 3c7f9728d468a4290f3d3d96763b495c27c5ae1a Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Sun, 14 Dec 2025 23:37:03 -0800
Subject: [PATCH 3/6] Integrate desugaring

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp | 49 +++++++++++++++++++++++++++----------
 1 file changed, 36 insertions(+), 13 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 93454c46202b5..30a826bc32f65 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -6654,8 +6654,29 @@ class FreeFunctionTemplateKernelArgsPrinter
       ConstTemplateArgumentVisitor<FreeFunctionTemplateKernelArgsPrinter, void,
                                    ArrayRef<TemplateArgument>>;
 
-  void PrintTemplateDeclName(const TemplateDecl *TD,
-                             ArrayRef<TemplateArgument> SpecArgs) {}
+  // Desugars a template argument. This helps avoid aliases.
+  static TemplateArgument DesugarTemplateArgument(const TemplateArgument &Arg) {
+    switch (Arg.getKind()) {
+    case TemplateArgument::ArgKind::Type: {
+      QualType ArgTy = Arg.getAsType();
+      return {QualType(ArgTy->getUnqualifiedDesugaredType(),
+                       ArgTy.getCVRQualifiers())};
+    }
+    case TemplateArgument::ArgKind::Template: {
+      TemplateName TN = Arg.getAsTemplate();
+      while (std::optional<TemplateName> DesugaredTN =
+                 TN.desugar(/*IgnoreDeduced=*/false))
+        TN = *DesugaredTN;
+      return {TN};
+    }
+    default:
+      return Arg;
+    }
+  }
+
+  void PrintDesugared(const TemplateArgument &Arg) {
+    DesugarTemplateArgument(Arg).print(Policy, O, /*IncludeType=*/false);
+  }
 
 public:
   FreeFunctionTemplateKernelArgsPrinter(raw_ostream &O, PrintingPolicy &Policy,
@@ -6699,7 +6720,7 @@ class FreeFunctionTemplateKernelArgsPrinter
     const TemplateDecl *TD = T->getTemplateName().getAsTemplateDecl();
     const auto *TTPD = dyn_cast<TemplateTemplateParmDecl>(TD);
     if (TTPD && !TTPD->getIdentifier())
-      SpecArgs[TTPD->getIndex()].print(Policy, O, /* IncludeType = */ false);
+      PrintDesugared(SpecArgs[TTPD->getIndex()]);
     else
       TD->printQualifiedName(O);
     O << "<";
@@ -6720,15 +6741,17 @@ class FreeFunctionTemplateKernelArgsPrinter
 
   void VisitTypeTemplateArgument(const TemplateArgument &Arg,
                                  ArrayRef<TemplateArgument> SpecArgs) {
+    TemplateArgument DesugaredArg = DesugarTemplateArgument(Arg);
     // If we reference an existing template argument without a known identifier,
     // print it instead.
-    const auto *TPT = dyn_cast<TemplateTypeParmType>(Arg.getAsType());
+    const auto *TPT = dyn_cast<TemplateTypeParmType>(DesugaredArg.getAsType());
     if (TPT && !TPT->getIdentifier()) {
-      SpecArgs[TPT->getIndex()].print(Policy, O, /* IncludeType = */ false);
+      PrintDesugared(SpecArgs[TPT->getIndex()]);
       return;
     }
 
-    const auto *TST = dyn_cast<TemplateSpecializationType>(Arg.getAsType());
+    const auto *TST =
+        dyn_cast<TemplateSpecializationType>(DesugaredArg.getAsType());
     if (TST && Arg.isInstantiationDependent()) {
       // This is an instantiation dependent template specialization, meaning
       // that some of its arguments reference template arguments of the free
@@ -6737,7 +6760,7 @@ class FreeFunctionTemplateKernelArgsPrinter
       return;
     }
 
-    Arg.print(Policy, O, /* IncludeType = */ false);
+    DesugaredArg.print(Policy, O, /* IncludeType = */ false);
   }
 
   void VisitDeclarationTemplateArgument(const TemplateArgument &,
@@ -6754,23 +6777,23 @@ class FreeFunctionTemplateKernelArgsPrinter
 
   void VisitIntegralTemplateArgument(const TemplateArgument &Arg,
                                      ArrayRef<TemplateArgument>) {
-    Arg.print(Policy, O, /* IncludeType = */ false);
+    PrintDesugared(Arg);
   }
 
   void VisitStructuralValueTemplateArgument(const TemplateArgument &Arg,
                                             ArrayRef<TemplateArgument>) {
-    Arg.print(Policy, O, /* IncludeType = */ false);
+    PrintDesugared(Arg);
   }
 
   void VisitTemplateTemplateArgument(const TemplateArgument &Arg,
                                      ArrayRef<TemplateArgument>) {
-    Arg.print(Policy, O, /* IncludeType = */ false);
+    PrintDesugared(Arg);
   }
 
   void VisitTemplateExpansionTemplateArgument(const TemplateArgument &Arg,
                                               ArrayRef<TemplateArgument>) {
     // Likely does not work similar to the one above
-    Arg.print(Policy, O, /* IncludeType = */ false);
+    PrintDesugared(Arg);
   }
 
   void VisitExpressionTemplateArgument(const TemplateArgument &Arg,
@@ -6784,7 +6807,7 @@ class FreeFunctionTemplateKernelArgsPrinter
       // we don't need to evaluate them.
       // If expression is instantiation-dependent, then we can't evaluate it
       // either, let's fallback to default printing mechanism.
-      Arg.print(Policy, O, /* IncludeType = */ false);
+      PrintDesugared(Arg);
       return;
     }
 
@@ -6798,7 +6821,7 @@ class FreeFunctionTemplateKernelArgsPrinter
 
   void VisitPackTemplateArgument(const TemplateArgument &Arg,
                                  ArrayRef<TemplateArgument>) {
-    Arg.print(Policy, O, /* IncludeType = */ false);
+    PrintDesugared(Arg);
   }
 };
 

From abbb457f04f7b28c6d238cd63b1adf152f6c0513 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Mon, 15 Dec 2025 03:30:25 -0800
Subject: [PATCH 4/6] Desugar aliases in independent free function arguments

This commit fixes the desugaring of independent aliased argument types
in dependent free function kernels.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp                   |  2 +-
 .../free-function-kernel-type-alias-arg.cpp   | 26 ++++++++++++++++---
 2 files changed, 24 insertions(+), 4 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 30a826bc32f65..7c7cf96a782b5 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -7019,7 +7019,7 @@ class FreeFunctionPrinter {
       const auto *TST = dyn_cast<TemplateSpecializationType>(T.getTypePtr());
       const auto *CTST = dyn_cast<TemplateSpecializationType>(CT.getTypePtr());
       if (!TST || !CTST) {
-        ParmListOstream << T.getAsString(Policy);
+        ParmListOstream << T.getDesugaredType(Context).getAsString(Policy);
         continue;
       }
 
diff --git a/clang/test/CodeGenSYCL/free-function-kernel-type-alias-arg.cpp b/clang/test/CodeGenSYCL/free-function-kernel-type-alias-arg.cpp
index d8717eaaab44f..aa7db827743c3 100644
--- a/clang/test/CodeGenSYCL/free-function-kernel-type-alias-arg.cpp
+++ b/clang/test/CodeGenSYCL/free-function-kernel-type-alias-arg.cpp
@@ -123,16 +123,37 @@ void baz_type(ns::Baz::type Arg) {}
 
 // CHECK: void baz_type(ns::Bar<double, float> Arg);
 
+template <int N>
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void baz_type_with_unrelated_template(ns::Baz::type Arg) {}
+template void baz_type_with_unrelated_template<1>(ns::Baz::type);
+
+// CHECK: template <int N> void baz_type_with_unrelated_template(ns::Bar<double, float>);
+
 [[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
 void foz_using(ns::FozUsing<ns::Foo> Arg) {}
 
 // CHECK: void foz_using(ns::Foz<ns::Foo> Arg);
 
+template <typename T>
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void foz_using_with_unrelated_template(ns::FozUsing<ns::Foo> Arg) {}
+template void foz_using_with_unrelated_template<int>(ns::FozUsing<ns::Foo>);
+
+// CHECK: template <typename T> void foz_using_with_unrelated_template(ns::Foz<ns::Foo>);
+
 [[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
 void foz_foo_using(ns::FozFooUsing Arg) {}
 
 // CHECK: void foz_foo_using(ns::Foz<ns::Foo> Arg);
 
+template <int N>
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void foz_foo_using_with_unrelated_template(ns::FozFooUsing Arg) {}
+template void foz_foo_using_with_unrelated_template<2>(ns::FozFooUsing);
+
+// CHECK: template <int N> void foz_foo_using_with_unrelated_template(ns::Foz<ns::Foo>);
+
 template<template <typename> class T1, typename T2>
 [[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
 void bar_using2_foz(ns::BarUsing2<ns::Foz<T1>, T2> Arg) {}
@@ -154,10 +175,9 @@ template void template_template_arg(ns::List<int>);
 
 // CHECK: template <typename T, template <typename > typename Container> void template_template_arg(Container<T>);
 
-// These test cases fail, but they are added here in advance to add a record of
-// known bugs.
 #if 0
-
+// This test case fails, but it is added here in advance to add a record of a
+// known bug.
 [[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
 void foz_foo_using(ns::Foz<ns::FooUsing> Arg) {}
 

From 9bdcbac11c36938d91b18c7b8444c9ea64f9e2d8 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Tue, 16 Dec 2025 02:39:45 -0800
Subject: [PATCH 5/6] Address comments

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp                   | 78 +++++++++----------
 ...free-function-kernel-pack-template-arg.cpp | 14 ++++
 2 files changed, 52 insertions(+), 40 deletions(-)
 create mode 100644 clang/test/CodeGenSYCL/free-function-kernel-pack-template-arg.cpp

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 7c7cf96a782b5..3b05f34730fb4 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -6683,14 +6683,44 @@ class FreeFunctionTemplateKernelArgsPrinter
                                         ASTContext &Context)
       : O(O), Policy(Policy), Context(Context) {}
 
-  void Visit(const TemplateSpecializationType *T,
-             const TemplateSpecializationType *CT) {
-    ArrayRef<TemplateArgument> SpecArgs = T->template_arguments();
-    ArrayRef<TemplateArgument> DeclArgs = CT->template_arguments();
+  void Visit(const ParmVarDecl *Param) {
+    // There are cases when we can't directly use neither the original
+    // argument type, nor its canonical version. An example would be:
+    // template<typename T>
+    // void kernel(sycl::accessor<T, 1>);
+    // template void kernel(sycl::accessor<int, 1>);
+    // Accessor has multiple non-type template arguments with default values
+    // and non-qualified type will not include necessary namespaces for all
+    // of them. Qualified type will have that information, but all references
+    // to T will be replaced to something like type-argument-0
+    // What we do instead is we iterate template arguments of both versions
+    // of a type in sync and take elements from one or another to get the best
+    // of both: proper references to template arguments of a kernel itself and
+    // fully-qualified names for enumerations.
+    //
+    // Moral of the story: drop integration header ASAP (but that is blocked
+    // by support for 3rd-party host compilers, which is important).
+    QualType T = Param->getType();
+    QualType CT = T.getCanonicalType();
+
+    const auto *TST = dyn_cast<TemplateSpecializationType>(T.getTypePtr());
+    const auto *CTST = dyn_cast<TemplateSpecializationType>(CT.getTypePtr());
+    if (!TST || !CTST) {
+      O << T.getDesugaredType(Context).getAsString(Policy);
+      return;
+    }
+
+    const TemplateSpecializationType *TSTAsNonAlias =
+        TST->getAsNonAliasTemplateSpecializationType();
+    if (TSTAsNonAlias)
+      TST = TSTAsNonAlias;
 
-    const TemplateDecl *TD = CT->getTemplateName().getAsTemplateDecl();
+    ArrayRef<TemplateArgument> SpecArgs = TST->template_arguments();
+    ArrayRef<TemplateArgument> DeclArgs = CTST->template_arguments();
+
+    const TemplateDecl *TD = CTST->getTemplateName().getAsTemplateDecl();
     if (!TD->getIdentifier())
-      TD = T->getTemplateName().getAsTemplateDecl();
+      TD = TST->getTemplateName().getAsTemplateDecl();
     TD->printQualifiedName(O);
 
     O << "<";
@@ -6792,7 +6822,6 @@ class FreeFunctionTemplateKernelArgsPrinter
 
   void VisitTemplateExpansionTemplateArgument(const TemplateArgument &Arg,
                                               ArrayRef<TemplateArgument>) {
-    // Likely does not work similar to the one above
     PrintDesugared(Arg);
   }
 
@@ -6802,7 +6831,7 @@ class FreeFunctionTemplateKernelArgsPrinter
     assert(E && "Failed to get an Expr for an Expression template arg?");
 
     if (Arg.isInstantiationDependent() ||
-        E->getType().getTypePtr()->isScopedEnumeralType()) {
+        E->getType()->isScopedEnumeralType()) {
       // Scoped enumerations can't be implicitly cast from integers, so
       // we don't need to evaluate them.
       // If expression is instantiation-dependent, then we can't evaluate it
@@ -6997,38 +7026,7 @@ class FreeFunctionPrinter {
       else
         ParmListOstream << ", ";
 
-      // There are cases when we can't directly use neither the original
-      // argument type, nor its canonical version. An example would be:
-      // template<typename T>
-      // void kernel(sycl::accessor<T, 1>);
-      // template void kernel(sycl::accessor<int, 1>);
-      // Accessor has multiple non-type template arguments with default values
-      // and non-qualified type will not include necessary namespaces for all
-      // of them. Qualified type will have that information, but all references
-      // to T will be replaced to something like type-argument-0
-      // What we do instead is we iterate template arguments of both versions
-      // of a type in sync and take elements from one or another to get the best
-      // of both: proper references to template arguments of a kernel itself and
-      // fully-qualified names for enumerations.
-      //
-      // Moral of the story: drop integration header ASAP (but that is blocked
-      // by support for 3rd-party host compilers, which is important).
-      QualType T = Param->getType();
-      QualType CT = T.getCanonicalType();
-
-      const auto *TST = dyn_cast<TemplateSpecializationType>(T.getTypePtr());
-      const auto *CTST = dyn_cast<TemplateSpecializationType>(CT.getTypePtr());
-      if (!TST || !CTST) {
-        ParmListOstream << T.getDesugaredType(Context).getAsString(Policy);
-        continue;
-      }
-
-      const TemplateSpecializationType *TSTAsNonAlias =
-          TST->getAsNonAliasTemplateSpecializationType();
-      if (TSTAsNonAlias)
-        TST = TSTAsNonAlias;
-
-      Printer.Visit(TST, CTST);
+      Printer.Visit(Param);
     }
     return ParamList.str().str();
   }
diff --git a/clang/test/CodeGenSYCL/free-function-kernel-pack-template-arg.cpp b/clang/test/CodeGenSYCL/free-function-kernel-pack-template-arg.cpp
new file mode 100644
index 0000000000000..90faccbd8807a
--- /dev/null
+++ b/clang/test/CodeGenSYCL/free-function-kernel-pack-template-arg.cpp
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -sycl-std=2020 -fsycl-int-header=%t.h %s
+// RUN: FileCheck -input-file=%t.h %s
+//
+// The purpose of this test is to ensure that forward declarations of free
+// function kernels are emitted properly.
+// However, this test checks a specific scenario:
+// - parameter packs are emitted correctly.
+
+template<typename... T>
+[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
+void parameter_pack(T... Args) {}
+template void parameter_pack(int Arg1, float Arg2);
+
+// CHECK: template <typename ... T> void parameter_pack(T...);

From b4a108db3c07b2af9a380029889aa0836e3541cc Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Tue, 16 Dec 2025 03:30:32 -0800
Subject: [PATCH 6/6] Assert and comment

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp | 7 +++++++
 1 file changed, 7 insertions(+)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 3b05f34730fb4..cf34c5aaeaf31 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -6641,6 +6641,11 @@ static void PrintNSClosingBraces(raw_ostream &OS, const DeclContext *DC) {
 /// What this visitor does is it iterates over both in sync, picking the right
 /// values from one or another.
 ///
+/// The template argument visitor functions take an additional
+/// ArrayRef<TemplateArgument> argument corresponding to the template arguments
+/// of the outermost template. This is used by some of these functions for
+/// mapping dependent template arguments.
+///
 /// Moral of the story: drop integration header ASAP (but that is blocked
 /// by support for 3rd-party host compilers, which is important).
 class FreeFunctionTemplateKernelArgsPrinter
@@ -6721,6 +6726,8 @@ class FreeFunctionTemplateKernelArgsPrinter
     const TemplateDecl *TD = CTST->getTemplateName().getAsTemplateDecl();
     if (!TD->getIdentifier())
       TD = TST->getTemplateName().getAsTemplateDecl();
+    assert(TD->getIdentifier() &&
+           "Either the type or the canonical type should have an identifier.");
     TD->printQualifiedName(O);
 
     O << "<";