Skip to content

[SYCL] Free function kernels bugfix #19535

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 3 commits into
base: sycl
Choose a base branch
from
Open

[SYCL] Free function kernels bugfix #19535

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
7eb843d
[SYCL] Free function kernels bugfix
AlexeySachkov Jun 24, 2025
8c32037
Apply clang-format
AlexeySachkov Jul 24, 2025
414cb2a
Expand SYCL E2E tests to cover changes
AlexeySachkov Jul 24, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
61 changes: 57 additions & 4 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,7 @@
#include "clang/AST/SYCLKernelInfo.h"
#include "clang/AST/StmtSYCL.h"
#include "clang/AST/TemplateArgumentVisitor.h"
#include "clang/AST/Type.h"
#include "clang/AST/TypeOrdering.h"
#include "clang/AST/TypeVisitor.h"
#include "clang/Analysis/CallGraph.h"
Expand Down Expand Up @@ -6710,23 +6711,75 @@ class FreeFunctionPrinter {
/// returned string Example:
/// \code
/// template <typename T1, typename T2>
/// void foo(T1 a, T2 b);
/// void foo(T1 a, int b, T2 c);
/// \endcode
/// returns string "T1 a, T2 b"
/// returns string "T1, int, T2"
std::string
getTemplatedParamList(const llvm::ArrayRef<clang::ParmVarDecl *> Parameters,
PrintingPolicy Policy) {
bool FirstParam = true;
llvm::SmallString<128> ParamList;
llvm::raw_svector_ostream ParmListOstream{ParamList};
Policy.SuppressTagKeyword = true;

for (ParmVarDecl *Param : Parameters) {
if (FirstParam)
FirstParam = false;
else
ParmListOstream << ", ";
ParmListOstream << Param->getType().getAsString(Policy);
ParmListOstream << " " << Param->getNameAsString();

// 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();

auto *ET = dyn_cast<ElaboratedType>(T.getTypePtr());
if (!ET) {
ParmListOstream << T.getAsString(Policy);
continue;
}

auto *TST = dyn_cast<TemplateSpecializationType>(
ET->getNamedType().getTypePtr());
auto *CTST = dyn_cast<TemplateSpecializationType>(CT.getTypePtr());
if (!TST || !CTST) {
ParmListOstream << T.getAsString(Policy);
continue;
}

TemplateName TN = TST->getTemplateName();
auto SpecArgs = TST->template_arguments();
auto DeclArgs = CTST->template_arguments();

TN.getAsTemplateDecl()->printQualifiedName(ParmListOstream);
ParmListOstream << "<";

for (size_t I = 0, E = std::max(DeclArgs.size(), SpecArgs.size()),
SE = SpecArgs.size();
I < E; ++I) {
if (I != 0)
ParmListOstream << ", ";
if (I < SE) // A specialized argument exists, use it
SpecArgs[I].print(Policy, ParmListOstream, false /* IncludeType */);
else // Print a canonical form of a default argument
DeclArgs[I].print(Policy, ParmListOstream, false /* IncludeType */);
}

ParmListOstream << ">";
}
return ParamList.str().str();
}
Expand Down
54 changes: 54 additions & 0 deletions clang/test/CodeGenSYCL/free-function-kernel-templated-arg-with-enum.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,54 @@
// 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:
// - free function kernel is a function template
// - its argument is templated and has non-type template parameter (with default
// value) that is an enumeration defined within a namespace

namespace ns {

enum class enum_A { A, B, C };

template<typename T, enum_A V = enum_A::B>
class feature_A {};

namespace nested {
enum class enum_B { A, B, C };

template<typename T, int V, enum_B V2 = enum_B::A, enum_A V3 = enum_A::C>
struct feature_B {};
}

inline namespace nested_inline {
namespace nested2 {
enum class enum_C { A, B, C };

template<int V = 42, enum_C V2 = enum_C::B>
struct feature_C {};
}
} // namespace nested_inline
} // namespace ns

template<typename T>
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
void templated_on_A(ns::feature_A<T> Arg) {}
template void templated_on_A(ns::feature_A<int>);

// CHECK: template <typename T> void templated_on_A(ns::feature_A<T, ns::enum_A::B>);

template<typename T, int V = 42>
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
void templated_on_B(ns::nested::feature_B<T, V> Arg) {}
template void templated_on_B(ns::nested::feature_B<int, 12>);

// CHECK: template <typename T, int V> void templated_on_B(ns::nested::feature_B<T, V, ns::nested::enum_B::A, ns::enum_A::C>);
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Note for reviewers: the exact thing which was missing is ns::nested:: for enum_B::A and ns:: for enum_A::C.


template<int V>
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
void templated_on_C(ns::nested2::feature_C<V> Arg) {}
template void templated_on_C(ns::nested2::feature_C<42>);

// CHECK: template <int V> void templated_on_C(ns::nested2::feature_C<V, ns::nested2::enum_C::B>);
48 changes: 24 additions & 24 deletions clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -267,17 +267,17 @@ namespace Testing::Tests {
// CHECK-NEXT: return (void (*)(struct ns::Arg<class ns::ns1::hasDefaultArg<struct ns::notatuple>, int, 12, struct ns::notatuple>))simple1;
// CHECK-NEXT: }

// CHECK: template <typename T> void templated(ns::Arg<T, float, 3, ns::notatuple> , T end);
// CHECK: template <typename T> void templated(ns::Arg<T, float, 3, ns::notatuple, <>>, T);
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Two changes here:

  • names of arguments in forward-declarations were dropped intentionally, because that's unnecessary information
  • <> is the way how empty type packs are printed. I'm not 100% sure that this is a legal C++ code, but I checked that generated integration header compiles with clang

// CHECK-NEXT: static constexpr auto __sycl_shim3() {
// CHECK-NEXT: return (void (*)(struct ns::Arg<int, float, 3, struct ns::notatuple>, int))templated<int>;
// CHECK-NEXT: }

// CHECK: template <typename T> void templated2(ns::Arg<T, ns::notatuple, 12, ns::notatuple> , T end);
// CHECK: template <typename T> void templated2(ns::Arg<T, ns::notatuple, 12, ns::notatuple, <>>, T);
// CHECK-NEXT: static constexpr auto __sycl_shim4() {
// CHECK-NEXT: return (void (*)(struct ns::Arg<int, struct ns::notatuple, 12, struct ns::notatuple>, int))templated2<int>;
// CHECK-NEXT: }

// CHECK: template <typename T, int a> void templated3(ns::Arg<T, ns::notatuple, a, ns::ns1::hasDefaultArg<ns::notatuple>, int, int> , T end);
// CHECK: template <typename T, int a> void templated3(ns::Arg<T, ns::notatuple, a, ns::ns1::hasDefaultArg<ns::notatuple>, int, int>, T);
// 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: }
Expand All @@ -297,7 +297,7 @@ namespace Testing::Tests {
// CHECK-NEXT: }

// CHECK: namespace TestNamespace {
// CHECK-NEXT: template <typename T> void templated(ns::Arg<T, float, 3, ns::notatuple> , T end);
// CHECK-NEXT: template <typename T> void templated(ns::Arg<T, float, 3, ns::notatuple, <>>, T);
// CHECK-NEXT: } // namespace TestNamespace

// CHECK: static constexpr auto __sycl_shim8() {
Expand All @@ -316,7 +316,7 @@ namespace Testing::Tests {

// CHECK: namespace TestNamespace {
// CHECK-NEXT: inline namespace _V1 {
// CHECK-NEXT: template <typename T, int a> void templated1(ns::Arg<T, float, a, ns::notatuple> , T end);
// CHECK-NEXT: template <typename T, int a> void templated1(ns::Arg<T, float, a, ns::notatuple, <>>, T);
// CHECK-NEXT: } // inline namespace _V1
// CHECK-NEXT: } // namespace TestNamespace
// CHECK: static constexpr auto __sycl_shim9() {
Expand All @@ -335,7 +335,7 @@ namespace Testing::Tests {

// CHECK: namespace TestNamespace {
// CHECK-NEXT: inline namespace _V2 {
// CHECK-NEXT: template <typename T, int a> void templated1(ns::Arg<T, T, a, ns::notatuple> , T end);
// CHECK-NEXT: template <typename T, int a> void templated1(ns::Arg<T, T, a, ns::notatuple, <>>, T);
// CHECK-NEXT: } // inline namespace _V2
// CHECK-NEXT: } // namespace TestNamespace
// CHECK: static constexpr auto __sycl_shim10() {
Expand All @@ -353,7 +353,7 @@ namespace Testing::Tests {
// CHECK-NEXT: }

// CHECK: namespace {
// CHECK-NEXT: template <typename T> void templated(T start, T end);
// CHECK-NEXT: template <typename T> void templated(T, T);
// CHECK-NEXT: } // namespace
// CHECK: static constexpr auto __sycl_shim11() {
// CHECK-NEXT: return (void (*)(float, float))templated<float>;
Expand All @@ -370,7 +370,7 @@ namespace Testing::Tests {
// CHECK-NEXT: }

// CHECK: struct TestStruct;
// CHECK: template <typename T> void templated(ns::Arg<T, float, 3, ns::notatuple> , T end);
// CHECK: template <typename T> void templated(ns::Arg<T, float, 3, ns::notatuple, <>>, T);
// CHECK-NEXT: static constexpr auto __sycl_shim12() {
// CHECK-NEXT: return (void (*)(struct ns::Arg<struct TestStruct, float, 3, struct ns::notatuple>, struct TestStruct))templated<struct TestStruct>;
// CHECK-NEXT:}
Expand All @@ -387,7 +387,7 @@ namespace Testing::Tests {

// CHECK: class BaseClass;
// CHECK: namespace {
// CHECK-NEXT: template <typename T> void templated(T start, T end);
// CHECK-NEXT: template <typename T> void templated(T, T);
// CHECK-NEXT: } // namespace
// CHECK: static constexpr auto __sycl_shim13() {
// CHECK-NEXT: return (void (*)(class BaseClass, class BaseClass))templated<class BaseClass>;
Expand All @@ -405,7 +405,7 @@ namespace Testing::Tests {

// CHECK: class ChildOne;
// CHECK: namespace {
// CHECK-NEXT: template <typename T> void templated(T start, T end);
// CHECK-NEXT: template <typename T> void templated(T, T);
// CHECK-NEXT: } // namespace
// CHECK: static constexpr auto __sycl_shim14() {
// CHECK-NEXT: return (void (*)(class ChildOne, class ChildOne))templated<class ChildOne>;
Expand All @@ -423,7 +423,7 @@ namespace Testing::Tests {

// CHECK: class ChildTwo;
// CHECK: namespace {
// CHECK-NEXT: template <typename T> void templated(T start, T end);
// CHECK-NEXT: template <typename T> void templated(T, T);
// CHECK-NEXT: } // namespace
// CHECK: static constexpr auto __sycl_shim15() {
// CHECK-NEXT: return (void (*)(class ChildTwo, class ChildTwo))templated<class ChildTwo>;
Expand All @@ -441,7 +441,7 @@ namespace Testing::Tests {

// CHECK: class ChildThree;
// CHECK: namespace {
// CHECK-NEXT: template <typename T> void templated(T start, T end);
// CHECK-NEXT: template <typename T> void templated(T, T);
// CHECK-NEXT: } // namespace
// CHECK: static constexpr auto __sycl_shim16() {
// CHECK-NEXT: return (void (*)(class ChildThree, class ChildThree))templated<class ChildThree>;
Expand All @@ -461,7 +461,7 @@ namespace Testing::Tests {
// CHECK-NEXT: template <int dim> struct id;
// CHECK-NEXT: }}
// CHECK: namespace {
// CHECK-NEXT: template <typename T> void templated(T start, T end);
// CHECK-NEXT: template <typename T> void templated(T, T);
// CHECK-NEXT: } // namespace
// CHECK: static constexpr auto __sycl_shim17() {
// CHECK-NEXT: return (void (*)(struct sycl::id<2>, struct sycl::id<2>))templated<struct sycl::id<2>>;
Expand All @@ -481,7 +481,7 @@ namespace Testing::Tests {
// CHECK-NEXT: template <int dim> struct range;
// CHECK-NEXT: }}
// CHECK: namespace {
// CHECK-NEXT: template <typename T> void templated(T start, T end);
// CHECK-NEXT: template <typename T> void templated(T, T);
// CHECK-NEXT: } // namespace
// CHECK: static constexpr auto __sycl_shim18() {
// CHECK-NEXT: return (void (*)(struct sycl::range<3>, struct sycl::range<3>))templated<struct sycl::range<3>>;
Expand All @@ -498,7 +498,7 @@ namespace Testing::Tests {
// CHECK-NEXT: }

// CHECK: namespace {
// CHECK-NEXT: template <typename T> void templated(T start, T end);
// CHECK-NEXT: template <typename T> void templated(T, T);
// CHECK-NEXT: } // namespace
// CHECK: static constexpr auto __sycl_shim19() {
// CHECK-NEXT: return (void (*)(int *, int *))templated<int *>;
Expand All @@ -515,7 +515,7 @@ namespace Testing::Tests {
// CHECK-NEXT: }

// CHECK: namespace {
// CHECK-NEXT: template <typename T> void templated(T start, T end);
// CHECK-NEXT: template <typename T> void templated(T, T);
// CHECK-NEXT: } // namespace
// CHECK: static constexpr auto __sycl_shim20() {
// CHECK-NEXT: return (void (*)(struct sycl::X<class ChildTwo>, struct sycl::X<class ChildTwo>))templated<struct sycl::X<class ChildTwo>>;
Expand All @@ -536,7 +536,7 @@ namespace Testing::Tests {
// CHECK-NEXT: }}}
// CHECK: namespace TestNamespace {
// CHECK-NEXT: inline namespace _V1 {
// CHECK-NEXT: template <typename T, int a> void templated1(ns::Arg<T, float, a, ns::notatuple> , T end);
// CHECK-NEXT: template <typename T, int a> void templated1(ns::Arg<T, float, a, ns::notatuple, <>>, T);
// CHECK-NEXT: } // inline namespace _V1
// CHECK-NEXT: } // namespace TestNamespace
// CHECK: static constexpr auto __sycl_shim21() {
Expand All @@ -553,7 +553,7 @@ namespace Testing::Tests {
// CHECK-NEXT: };
// CHECK-NEXT: }

// CHECK: template <typename ... Args> void variadic_templated(Args... args);
// CHECK: template <typename ... Args> void variadic_templated(Args...);
// CHECK-NEXT: static constexpr auto __sycl_shim22() {
// CHECK-NEXT: return (void (*)(int, float, char))variadic_templated<int, float, char>;
// CHECK-NEXT: }
Expand All @@ -568,7 +568,7 @@ namespace Testing::Tests {
// CHECK-NEXT: };
// CHECK-NEXT: }

// CHECK: template <typename ... Args> void variadic_templated(Args... args);
// CHECK: template <typename ... Args> void variadic_templated(Args...);
// CHECK-NEXT: static constexpr auto __sycl_shim23() {
// CHECK-NEXT: return (void (*)(int, float, char, int))variadic_templated<int, float, char, int>;
// CHECK-NEXT: }
Expand All @@ -583,7 +583,7 @@ namespace Testing::Tests {
// CHECK-NEXT: };
// CHECK-NEXT: }

// CHECK: template <typename ... Args> void variadic_templated(Args... args);
// CHECK: template <typename ... Args> void variadic_templated(Args...);
// CHECK-NEXT: static constexpr auto __sycl_shim24() {
// CHECK-NEXT: return (void (*)(float, float))variadic_templated<float, float>;
// CHECK-NEXT: }
Expand All @@ -598,7 +598,7 @@ namespace Testing::Tests {
// CHECK-NEXT: };
// CHECK-NEXT: }

// CHECK: template <typename T, typename ... Args> void variadic_templated1(T b, Args... args);
// CHECK: template <typename T, typename ... Args> void variadic_templated1(T, Args...);
// CHECK-NEXT: static constexpr auto __sycl_shim25() {
// CHECK-NEXT: return (void (*)(float, char, char))variadic_templated1<float, char, char>;
// CHECK-NEXT: }
Expand All @@ -613,7 +613,7 @@ namespace Testing::Tests {
// CHECK-NEXT: };
// CHECK-NEXT: }

// CHECK: template <typename T, typename ... Args> void variadic_templated1(T b, Args... args);
// CHECK: template <typename T, typename ... Args> void variadic_templated1(T, Args...);
// CHECK-NEXT: static constexpr auto __sycl_shim26() {
// CHECK-NEXT: return (void (*)(int, float, char))variadic_templated1<int, float, char>;
// CHECK-NEXT: }
Expand All @@ -630,7 +630,7 @@ namespace Testing::Tests {

// CHECK: namespace Testing {
// CHECK-NEXT: namespace Tests {
// CHECK-NEXT: template <typename T, typename ... Args> void variadic_templated(T b, Args... args);
// CHECK-NEXT: template <typename T, typename ... Args> void variadic_templated(T, Args...);
// CHECK-NEXT: } // namespace Tests
// CHECK-NEXT: } // namespace Testing
// CHECK: static constexpr auto __sycl_shim27() {
Expand All @@ -649,7 +649,7 @@ namespace Testing::Tests {

// CHECK: namespace Testing {
// CHECK-NEXT: namespace Tests {
// CHECK-NEXT: template <typename T, typename ... Args> void variadic_templated(T b, Args... args);
// CHECK-NEXT: template <typename T, typename ... Args> void variadic_templated(T, Args...);
// CHECK-NEXT: } // namespace Tests
// CHECK-NEXT: } // namespace Testing
// CHECK: static constexpr auto __sycl_shim28() {
Expand Down
Loading
Loading