Skip to content

SYCL free function namespace support #17585

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

Merged
merged 26 commits into from
Apr 9, 2025
Merged

SYCL free function namespace support #17585

Show file tree
Hide file tree
Changes from 9 commits
Commits
Show all changes
26 commits
Select commit Hold shift + click to select a range
b305e94
[SYCL] add e2e test for namespace
dklochkov-emb Mar 7, 2025
5ff7b2c
[SYCL] implement namespace support for free function
dklochkov-emb Mar 18, 2025
12df05e
[SYCL] fix test issues after adding namespace support
dklochkov-emb Mar 19, 2025
bd15ef7
Merge remote-tracking branch 'upstream/sycl' into sycl-free-function-…
dklochkov-emb Mar 19, 2025
65ef84f
[SYCL] do not use macros in tests
dklochkov-emb Mar 19, 2025
117e97a
[SYCL] fix tests
dklochkov-emb Mar 21, 2025
c53312a
[SYCL] exclude cuda from free function tests
dklochkov-emb Mar 25, 2025
190ac32
[SYCL] update unsupported list of tests
dklochkov-emb Mar 25, 2025
0858e70
[SYCL] fix typo
dklochkov-emb Mar 26, 2025
32113db
[SYCL] rework free functions to use a separate entity
dklochkov-emb Mar 27, 2025
2183658
[SYCL] Update code style
dklochkov-emb Mar 28, 2025
5dd5894
[SYCL] update formating
dklochkov-emb Mar 28, 2025
2499963
[SYCL] fix includes in SemaSYCL
dklochkov-emb Mar 28, 2025
a5edf00
[SYCL] include missing header
dklochkov-emb Mar 28, 2025
2bb7c21
[SYCL] add unit tests of free function namespace support
dklochkov-emb Mar 28, 2025
b3bd8ae
[SYCL] fix includes
dklochkov-emb Mar 28, 2025
c755d26
Merge remote-tracking branch 'upstream/sycl' into sycl-free-function-…
dklochkov-emb Mar 28, 2025
64906f8
[SYCL] fix post merge issue
dklochkov-emb Mar 28, 2025
c976900
Merge remote-tracking branch 'upstream/sycl' into sycl-free-function-…
dklochkov-emb Mar 31, 2025
4cebf72
[SYCL] check two functions with the same namespace
dklochkov-emb Apr 1, 2025
157b39a
[SYCL][E2E] fix formatting
dklochkov-emb Apr 1, 2025
e3ff53a
[SYCL][E2E] update free function integration header test
dklochkov-emb Apr 1, 2025
bdb3967
[SYCL][E2E] do not update no-unsupported test
dklochkov-emb Apr 1, 2025
cae876a
[SYCL] add new tests to cover free fuction namespace support
dklochkov-emb Apr 4, 2025
afb8b59
[SYCL] remove unused variable
dklochkov-emb Apr 4, 2025
09c7786
[SYCL][E2E] add new tests for free functions namespace support
dklochkov-emb Apr 4, 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
165 changes: 105 additions & 60 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,7 +27,6 @@
#include "clang/Basic/Diagnostic.h"
#include "clang/Basic/TargetInfo.h"
#include "clang/Basic/Version.h"
#include "clang/AST/SYCLKernelInfo.h"
#include "clang/Sema/Attr.h"
#include "clang/Sema/Initialization.h"
#include "clang/Sema/ParsedAttr.h"
Expand Down Expand Up @@ -6388,6 +6387,91 @@ static void EmitPragmaDiagnosticPop(raw_ostream &O) {
O << "\n";
}

template <typename BeforeFn, typename AfterFn>
static void PrintNSHelper(BeforeFn Before, AfterFn After, raw_ostream &OS,
const DeclContext *DC) {
if (DC->isTranslationUnit())
return;

const auto *CurDecl = cast<Decl>(DC);
// Ensure we are in the canonical version, so that we know we have the 'full'
// name of the thing.
CurDecl = CurDecl->getCanonicalDecl();

// We are intentionally skipping linkage decls and record decls. Namespaces
// can appear in a linkage decl, but not a record decl, so we don't have to
// worry about the names getting messed up from that. We handle record decls
// later when printing the name of the thing.
const auto *NS = dyn_cast<NamespaceDecl>(CurDecl);
if (NS)
Before(OS, NS);

if (const DeclContext *NewDC = CurDecl->getDeclContext())
PrintNSHelper(Before, After, OS, NewDC);

if (NS)
After(OS, NS);
}

static void PrintNamespaces(raw_ostream &OS, const DeclContext *DC,
bool isPrintNamesOnly = false) {
PrintNSHelper([](raw_ostream &OS, const NamespaceDecl *NS) {},
[isPrintNamesOnly](raw_ostream &OS, const NamespaceDecl *NS) {
if (!isPrintNamesOnly) {
if (NS->isInline())
OS << "inline ";
OS << "namespace ";
}
if (!NS->isAnonymousNamespace()) {
OS << NS->getName();
if (isPrintNamesOnly)
OS << "::";
else
OS << " ";
}
if (!isPrintNamesOnly) {
OS << "{\n";
}
},
OS, DC);
}

static void PrintNSClosingBraces(raw_ostream &OS, const DeclContext *DC) {
PrintNSHelper(
[](raw_ostream &OS, const NamespaceDecl *NS) {
OS << "} // ";
if (NS->isInline())
OS << "inline ";

OS << "namespace ";
if (!NS->isAnonymousNamespace())
OS << NS->getName();

OS << '\n';
},
[](raw_ostream &OS, const NamespaceDecl *NS) {}, OS, DC);
}

static bool insertFreeFunctionDeclaration(const PrintingPolicy &Policy,
Copy link
Contributor

@Fznamznon Fznamznon Mar 27, 2025

Choose a reason for hiding this comment

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

Maybe a documentation comment here makes sense? The function is named insertFreeFunctionDeclaration yet it returns true if a namespace was inserted. Normally functions that are named like that return success status in a bool flag. Alternatively the bool return value can be converted to a reference parameter with a good name.

Copy link
Contributor Author

@dklochkov-emb dklochkov-emb Mar 28, 2025

Choose a reason for hiding this comment

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

Reworked to have a separate class to emit free function and shim for that. Added comments for each method.

const FunctionDecl *FD,
const std::string& Args,
raw_ostream &O) {
const auto *DC = FD->getDeclContext();
bool NSInserted{false};
if (DC) {
if (isa<NamespaceDecl>(DC)) {
PrintNamespaces(O, FD);
NSInserted = true;
}
O << FD->getReturnType().getAsString() << " ";
Copy link
Contributor

@Fznamznon Fznamznon Mar 27, 2025

Choose a reason for hiding this comment

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

I wonder why doing FD->print() does not suffice anymore?

Copy link
Contributor Author

@dklochkov-emb dklochkov-emb Mar 28, 2025

Choose a reason for hiding this comment

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

FD->print() works fine for free function declared outside of namespace. If it is declared in it, print method will emit:
void NS1::NS2::some_func(...)
for function
namespace NS1::NS2{ void some_func(...){} }

The form
void NS1::NS2::some_func(...){}
works fine but forward declaration
void NS1::NS2::some_func(...) does not work. That is why it is needed to emit namespace first into intermediate header for function declaration.
P.S.
I reviewed Policy to use FD->print() but did not find.

O << FD->getNameAsString() << "(" << Args << ");";
if (NSInserted) {
PrintNSClosingBraces(O, FD);
}
}
return NSInserted;
}

void SYCLIntegrationHeader::emit(raw_ostream &O) {
O << "// This is auto-generated SYCL integration header.\n";
O << "\n";
Expand Down Expand Up @@ -6676,16 +6760,26 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
if (K.SyclKernel->getLanguageLinkage() == CLanguageLinkage)
O << "extern \"C\" ";
std::string ParmList;
std::string ParmListWithNames;
bool FirstParam = true;
Policy.SuppressDefaultTemplateArgs = false;
Policy.PrintCanonicalTypes = true;
llvm::raw_string_ostream ParmListWithNamesOstream{ParmListWithNames};
for (ParmVarDecl *Param : K.SyclKernel->parameters()) {
if (FirstParam)
FirstParam = false;
else
{
ParmList += ", ";
ParmListWithNamesOstream << ", ";
}
Policy.SuppressTagKeyword = true;
Param->getType().print(ParmListWithNamesOstream, Policy);
Policy.SuppressTagKeyword = false;
ParmListWithNamesOstream << " " << Param->getNameAsString();
Comment on lines +6840 to +6843
Copy link
Contributor

@Fznamznon Fznamznon Mar 27, 2025

Choose a reason for hiding this comment

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

Could you please elaborate what this particular addition is trying to achieve, why the previous code did not suffice and how does it relate to namespace printing?

Copy link
Contributor Author

@dklochkov-emb dklochkov-emb Mar 28, 2025

Choose a reason for hiding this comment

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

In general, ParamList contains only parameter types, i.e. for the function
void some_func(float a, float* b)
ParamList contains {float, float*}
It was added new list to have additional list of parameters with names to pass already existed tests for free function which checked generated header.
flag
Policy.SuppressTagKeyword = true;
forces printing without type tags, i.e. without words class and struct.

ParmList += Param->getType().getCanonicalType().getAsString(Policy);
}
ParmListWithNamesOstream.flush();
FunctionTemplateDecl *FTD = K.SyclKernel->getPrimaryTemplate();
Policy.PrintCanonicalTypes = false;
Policy.SuppressDefinition = true;
Expand Down Expand Up @@ -6719,17 +6813,23 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
// template arguments that match default template arguments while printing
// template-ids, even if the source code doesn't reference them.
Policy.EnforceDefaultTemplateArgs = true;
bool NSInserted{false};
if (FTD) {
FTD->print(O, Policy);
O << ";\n";
} else {
K.SyclKernel->print(O, Policy);
NSInserted = insertFreeFunctionDeclaration(Policy, K.SyclKernel, ParmListWithNames, O);
O << "\n";
}
O << ";\n";

// Generate a shim function that returns the address of the free function.
O << "static constexpr auto __sycl_shim" << ShimCounter << "() {\n";
O << " return (void (*)(" << ParmList << "))"
<< K.SyclKernel->getIdentifier()->getName().data();
O << " return (void (*)(" << ParmList << "))";
if (NSInserted) {
PrintNamespaces(O, K.SyclKernel, true);
}

O << K.SyclKernel->getIdentifier()->getName().data();
if (FTD) {
const TemplateArgumentList *TAL =
K.SyclKernel->getTemplateSpecializationArgs();
Expand Down Expand Up @@ -6898,61 +6998,6 @@ bool SYCLIntegrationFooter::emit(StringRef IntHeaderName) {
return emit(Out);
}

template <typename BeforeFn, typename AfterFn>
static void PrintNSHelper(BeforeFn Before, AfterFn After, raw_ostream &OS,
const DeclContext *DC) {
if (DC->isTranslationUnit())
return;

const auto *CurDecl = cast<Decl>(DC);
// Ensure we are in the canonical version, so that we know we have the 'full'
// name of the thing.
CurDecl = CurDecl->getCanonicalDecl();

// We are intentionally skipping linkage decls and record decls. Namespaces
// can appear in a linkage decl, but not a record decl, so we don't have to
// worry about the names getting messed up from that. We handle record decls
// later when printing the name of the thing.
const auto *NS = dyn_cast<NamespaceDecl>(CurDecl);
if (NS)
Before(OS, NS);

if (const DeclContext *NewDC = CurDecl->getDeclContext())
PrintNSHelper(Before, After, OS, NewDC);

if (NS)
After(OS, NS);
}

static void PrintNamespaces(raw_ostream &OS, const DeclContext *DC) {
PrintNSHelper([](raw_ostream &OS, const NamespaceDecl *NS) {},
[](raw_ostream &OS, const NamespaceDecl *NS) {
if (NS->isInline())
OS << "inline ";
OS << "namespace ";
if (!NS->isAnonymousNamespace())
OS << NS->getName() << " ";
OS << "{\n";
},
OS, DC);
}

static void PrintNSClosingBraces(raw_ostream &OS, const DeclContext *DC) {
PrintNSHelper(
[](raw_ostream &OS, const NamespaceDecl *NS) {
OS << "} // ";
if (NS->isInline())
OS << "inline ";

OS << "namespace ";
if (!NS->isAnonymousNamespace())
OS << NS->getName();

OS << '\n';
},
[](raw_ostream &OS, const NamespaceDecl *NS) {}, OS, DC);
}

static std::string EmitShim(raw_ostream &OS, unsigned &ShimCounter,
const std::string &LastShim,
const NamespaceDecl *AnonNS) {
Expand Down
11 changes: 6 additions & 5 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 @@ -86,18 +86,19 @@ foo(Arg1<int> arg) {
// CHECK-NEXT: template <typename T, typename, int a, typename, typename ...TS> struct Arg;
// CHECK-NEXT: }

// CHECK: void ns::simple(ns::Arg<char, int, 12, ns::notatuple>);
// CHECK-NEXT: static constexpr auto __sycl_shim1() {
// CHECK-NEXT: return (void (*)(struct ns::Arg<char, int, 12, struct ns::notatuple>))simple;
// CHECK: namespace ns {
// CHECK-NEXT: void simple(ns::Arg<char, int, 12, ns::notatuple> );} // namespace ns
// CHECK: static constexpr auto __sycl_shim1() {
// CHECK-NEXT: return (void (*)(struct ns::Arg<char, int, 12, struct ns::notatuple>))ns::simple;
// CHECK-NEXT: }

// CHECK: Forward declarations of kernel and its argument types:
// CHECK: namespace ns {
// CHECK: namespace ns1 {
// CHECK-NEXT: template <typename A> class hasDefaultArg;
// CHECK-NEXT: }
// CHECK-NEXT: }}
Copy link
Contributor

@Fznamznon Fznamznon Mar 27, 2025

Choose a reason for hiding this comment

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

Can we add more FE tests? Like with various combinations of namespaces around the free function kernel declaration? With inline namespace and not. Can we also test that codegen and semantic analysis is ok for free function kernels defined in a (maybe nested) namespace?

Copy link
Contributor Author

@dklochkov-emb dklochkov-emb Mar 28, 2025

Choose a reason for hiding this comment

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

I added new e2e tests to check any possible namespaces: nested, anonymous, inline etc. Is it enough or add in these tests too? New tests do not check header directly but if something is emitted wrong, they will fail.

Copy link
Contributor

@Fznamznon Fznamznon Mar 28, 2025

Choose a reason for hiding this comment

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

SYCL compiler is complicated and has a lot of components. If we only have a e2e test and it fails suddenly (for example after a pulldown), it may take a while to identify which component now has a problem.
This is one of reasons why we normally check each component separately with unit tests and everything together in e2e tests. FE-only tests are "unit" tests in this scenario. They will help more quickly to identify that the problem is in FE. They will also help people to fix any FE problems without needing to have sycl rt and device. So, I still encourage to add FE-only tests.

Copy link
Contributor Author

@dklochkov-emb dklochkov-emb Mar 28, 2025

Choose a reason for hiding this comment

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

Good point, I did not see that these tests are units. Added new checks.

Copy link
Contributor

@Fznamznon Fznamznon Mar 28, 2025

Choose a reason for hiding this comment

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

Well they are "unit" because clang is enormous itself and has its own unit tests but in terms of SYCL compiler we can consider them as unit tests.


// CHECK: void simple1(ns::Arg<ns::ns1::hasDefaultArg<ns::notatuple>, int, 12, ns::notatuple>);
// CHECK: void simple1(ns::Arg<ns::ns1::hasDefaultArg<ns::notatuple>, int, 12, ns::notatuple> );
// CHECK-NEXT: static constexpr auto __sycl_shim2() {
// CHECK-NEXT: return (void (*)(struct ns::Arg<class ns::ns1::hasDefaultArg<struct ns::notatuple>, int, 12, struct ns::notatuple>))simple1;
// CHECK-NEXT: }
Expand Down
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/free_function_int_header.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -163,7 +163,7 @@ void ff_8(sycl::work_group_memory<int>) {

// CHECK: Definition of _Z18__sycl_kernel_ff_2Piii as a free function kernel
// CHECK: Forward declarations of kernel and its argument types:
// CHECK: void ff_2(int *ptr, int start, int end);
// CHECK: void ff_2(int * ptr, int start, int end);
// CHECK-NEXT: static constexpr auto __sycl_shim1() {
// CHECK-NEXT: return (void (*)(int *, int, int))ff_2;
// CHECK-NEXT: }
Expand All @@ -180,7 +180,7 @@ void ff_8(sycl::work_group_memory<int>) {

// CHECK: Definition of _Z18__sycl_kernel_ff_2Piiii as a free function kernel
// CHECK: Forward declarations of kernel and its argument types:
// CHECK: void ff_2(int *ptr, int start, int end, int value);
// CHECK: void ff_2(int * ptr, int start, int end, int value);
// CHECK-NEXT: static constexpr auto __sycl_shim2() {
// CHECK-NEXT: return (void (*)(int *, int, int, int))ff_2;
// CHECK-NEXT: }
Expand Down Expand Up @@ -309,7 +309,7 @@ void ff_8(sycl::work_group_memory<int>) {
// CHECK: Forward declarations of kernel and its argument types:
// CHECK: template <typename DataT> class work_group_memory;

// CHECK: void ff_8(sycl::work_group_memory<int>);
// CHECK: void ff_8(sycl::work_group_memory<int> );
// CHECK-NEXT: static constexpr auto __sycl_shim9() {
// CHECK-NEXT: return (void (*)(class sycl::work_group_memory<int>))ff_8;
// CHECK-NEXT: }
Expand Down
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -52,7 +52,7 @@ int main(){
// CHECK-RTC-NOT: free_function_single_kernel
// CHECK-RTC-NOT: free_function_nd_range

// CHECK-NORTC: void free_function_single(int *ptr, int start, int end);
// CHECK-NORTC: void free_function_single(int * ptr, int start, int end);
// CHECK-NORTC: static constexpr auto __sycl_shim[[#FIRST:]]()
// CHECK-NORTC-NEXT: return (void (*)(int *, int, int))free_function_single;

Expand All @@ -63,7 +63,7 @@ int main(){
// CHECK-NORTC-NEXT: static constexpr bool value = true;


// CHECK-NORTC: void free_function_nd_range(int *ptr, int start, int end);
// CHECK-NORTC: void free_function_nd_range(int * ptr, int start, int end);
// CHECK-NORTC: static constexpr auto __sycl_shim[[#SECOND:]]() {
// CHECK-NORTC-NEXT: return (void (*)(int *, int, int))free_function_nd_range;

Expand Down
Loading
Loading