Skip to content

[SYCL] Add clang support for code_location in KernelInfo #5335

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 23 commits into from
Feb 4, 2022
Merged

[SYCL] Add clang support for code_location in KernelInfo #5335

Show file tree
Hide file tree
Changes from 20 commits
Commits
Show all changes
23 commits
Select commit Hold shift + click to select a range
b4f196f
[SYCL] Add clang support for KernelInfo code_location() method
schittir Jan 6, 2022
4fada17
Fix NDEBUG typo
schittir Jan 19, 2022
3bc31c7
Populate code_location members
schittir Jan 19, 2022
457b97b
Populate fields
schittir Jan 20, 2022
5ed0c13
Make changes per Erich's comments
schittir Jan 25, 2022
114d8cd
Populate "FuncionName" and change the test
schittir Jan 27, 2022
7be77e7
Delete int_header2.cpp
schittir Jan 27, 2022
622d27a
Update FunctionName and test
schittir Jan 27, 2022
a04d561
Fix format
schittir Jan 27, 2022
d0007da
Tidying up
schittir Jan 27, 2022
e7e29b9
Remove extra space
schittir Jan 27, 2022
08a8d15
Remove temporary variables and path information from filename
schittir Jan 28, 2022
c848a82
Fix test and add more cases
schittir Jan 28, 2022
7201467
Remove comments
schittir Jan 28, 2022
338bff0
Fix format
schittir Jan 29, 2022
99242c8
Add NDEBUG cases
schittir Jan 30, 2022
c067405
Remove NDEBUG from kernel_desc.hpp
schittir Jan 31, 2022
4c59038
Fix quotes
schittir Jan 31, 2022
2b10cbf
Emit fully qualified name
schittir Feb 2, 2022
14124ff
Add an -fsycl-is-host RUN line with the int header as FileCheck input
schittir Feb 2, 2022
16d6af3
Emit NDEBUG case to integration header
schittir Feb 2, 2022
190a44d
Undo accidental deletion of __SYCL_DLL_LOCAL
schittir Feb 2, 2022
76a77f4
Add Functor test case
schittir Feb 3, 2022
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
34 changes: 34 additions & 0 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4684,6 +4684,10 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {

for (const KernelDesc &K : KernelDescs) {
const size_t N = K.Params.size();
PresumedLoc PLoc = S.Context.getSourceManager().getPresumedLoc(
S.Context.getSourceManager()
.getExpansionRange(K.KernelLocation)
.getEnd());
if (K.IsUnnamedKernel) {
O << "template <> struct KernelInfoData<";
OutputStableNameInChars(O, K.StableName);
Expand All @@ -4707,6 +4711,36 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
O << " __SYCL_DLL_LOCAL\n";
O << " static constexpr bool isESIMD() { return " << K.IsESIMDKernel
<< "; }\n";
O << " __SYCL_DLL_LOCAL\n";
O << " static constexpr const char* getFileName() { return \""
#ifndef NDEBUG
<< std::string(PLoc.getFilename())
.substr(std::string(PLoc.getFilename()).find_last_of("/\\") + 1)
#endif
<< "\"; }\n";
O << " __SYCL_DLL_LOCAL\n";
O << " static constexpr const char* getFunctionName() { return \"";
#ifndef NDEBUG
SYCLKernelNameTypePrinter Printer(O, Policy);
Printer.Visit(K.NameType);
#endif
O << "\"; }\n";
O << " __SYCL_DLL_LOCAL\n";
O << " static constexpr unsigned getLineNumber() { return "
#ifndef NDEBUG
<< PLoc.getLine()
#else
<< 0
#endif
<< "; }\n";
O << " __SYCL_DLL_LOCAL\n";
O << " static constexpr unsigned getColumnNumber() { return "
#ifndef NDEBUG
<< PLoc.getColumn()
#else
<< 0
#endif
<< "; }\n";
O << "};\n";
CurStart += N;
}
Expand Down
101 changes: 101 additions & 0 deletions clang/test/CodeGenSYCL/code_location.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,101 @@
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem -sycl-std=2020 -fsycl-int-header=%t.h %s -o %t.out
// RUN: FileCheck -input-file=%t.h %s
// RUN: %clang_cc1 -fsycl-is-host -sycl-std=2020 %s | FileCheck -input-file=%t.h %s
// : %clang_cc1 -fsycl-is-device -internal-isystem -sycl-std=2020 -DNDEBUG -fsycl-int-header=%t2.h %s -o %t2.out
// : FileCheck -input-file=%t2.h %s
#include "Inputs/sycl.hpp"
#ifndef NDEBUG
int test1() {
cl::sycl::queue q;
q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });
q.submit([&](cl::sycl::handler &h) { h.single_task<class KernelName>([]() {}); });
return 0;
}
// Specializations of KernelInfo for kernel function types:
// CHECK: template <> struct KernelInfoData<'_', 'Z', 'T', 'S', 'Z', 'Z', '5', 't', 'e', 's', 't', '1', 'v', 'E', 'N', 'K', 'U', 'l', 'R', 'N', '2', 'c', 'l', '4', 's', 'y', 'c', 'l', '7', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '_', 'c', 'l', 'E', 'S', '2', '_', 'E', 'U', 'l', 'v', 'E', '_'> {
// CHECK: static constexpr const char* getFileName() { return "code_location.cpp"; }
// CHECK: static constexpr const char* getFunctionName() { return ""; }
// CHECK: static constexpr unsigned getLineNumber() { return 10; }
// CHECK: static constexpr unsigned getColumnNumber() { return 54; }
//};

// CHECK: template <> struct KernelInfo<KernelName> {
// CHECK: static constexpr const char* getFileName() { return "code_location.cpp"; }
// CHECK: static constexpr const char* getFunctionName() { return "KernelName"; }
// CHECK: static constexpr unsigned getLineNumber() { return 11; }
// CHECK: static constexpr unsigned getColumnNumber() { return 72; }
//};

class KernelName2;
int test2() {
cl::sycl::queue q;
q.submit([&](cl::sycl::handler &h) { h.single_task(
[] { int i = 2; }); });
q.submit([&](cl::sycl::handler &h) { h.single_task<KernelName2>(
[] { int i = 2; }); });
return 0;
}

// CHECK: template <> struct KernelInfoData<'_', 'Z', 'T', 'S', 'Z', 'Z', '5', 't', 'e', 's', 't', '2', 'v', 'E', 'N', 'K', 'U', 'l', 'R', 'N', '2', 'c', 'l', '4', 's', 'y', 'c', 'l', '7', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '_', 'c', 'l', 'E', 'S', '2', '_', 'E', 'U', 'l', 'v', 'E', '_'> {
// CHECK: static constexpr const char* getFileName() { return "code_location.cpp"; }
// CHECK: static constexpr const char* getFunctionName() { return ""; }
// CHECK: static constexpr unsigned getLineNumber() { return 33; }
// CHECK: static constexpr unsigned getColumnNumber() { return 44; }
//};
// CHECK: template <> struct KernelInfo<::KernelName2> {
// CHECK: static constexpr const char* getFileName() { return "code_location.cpp"; }
// CHECK: static constexpr const char* getFunctionName() { return "::KernelName2"; }
// CHECK: static constexpr unsigned getLineNumber() { return 35; }
// CHECK: static constexpr unsigned getColumnNumber() { return 44; }
//};

template <typename T> class KernelName3;
int test3() {
cl::sycl::queue q;
q.submit([&](cl::sycl::handler &h) { h.single_task<KernelName3<KernelName2>>(
[] { int i = 3; }); });
return 0;
}

// CHECK: template <> struct KernelInfo<::KernelName3<::KernelName2>> {
// CHECK: static constexpr const char* getFileName() { return "code_location.cpp"; }
// CHECK: static constexpr const char* getFunctionName() { return "::KernelName3<::KernelName2>"; }
// CHECK: static constexpr unsigned getLineNumber() { return 56; }
// CHECK: static constexpr unsigned getColumnNumber() { return 44; }
//};

auto l4 = []() { return 4; };
int test4() {
cl::sycl::queue q;
q.submit([=](cl::sycl::handler &h) { h.single_task<class KernelName4>(l4); });
return 0;
}

// CHECK: template <> struct KernelInfo<KernelName4> {
// CHECK: static constexpr const char* getFileName() { return "code_location.cpp"; }
// CHECK: static constexpr const char* getFunctionName() { return "KernelName4"; }
// CHECK: static constexpr unsigned getLineNumber() { return 67; }
// CHECK: static constexpr unsigned getColumnNumber() { return 11; }
//};

#else
int test5() {
cl::sycl::queue q;
q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });
q.submit([&](cl::sycl::handler &h) { h.single_task<class KernelName5>([]() {}); });
return 0;
}

// HECK: template <> struct KernelInfoData<'_', 'Z', 'T', 'S', 'Z', 'Z', '5', 't', 'e', 's', 't', '5', 'v', 'E', 'N', 'K', 'U', 'l', 'R', 'N', '2', 'c', 'l', '4', 's', 'y', 'c', 'l', '7', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '_', 'c', 'l', 'E', 'S', '2', '_', 'E', 'U', 'l', 'v', 'E', '_'> {
// HECK: static constexpr const char* getFileName() { return ; }
// HECK: static constexpr const char* getFunctionName() { return ; }
// HECK: static constexpr unsigned getLineNumber() { return 0; }
// HECK: static constexpr unsigned getColumnNumber() { return 0; }
//};
// HECK: template <> struct KernelInfo<KernelName5> {
// HECK: static constexpr const char* getFileName() { return ; }
// HECK: static constexpr const char* getFunctionName() { return ; }
// HECK: static constexpr unsigned getLineNumber() { return 0; }
// HECK: static constexpr unsigned getColumnNumber() { return 0; }
//};
#endif
12 changes: 12 additions & 0 deletions sycl/include/CL/sycl/detail/kernel_desc.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -75,6 +75,10 @@ template <class KernelNameType> struct KernelInfo {
}
static constexpr const char *getName() { return ""; }
static constexpr bool isESIMD() { return 0; }
static constexpr const char *getFileName() { return ""; }
static constexpr const char *getFunctionName() { return ""; }
static constexpr unsigned getLineNumber() { return 0; }
static constexpr unsigned getColumnNumber() { return 0; }
};
#else
template <char...> struct KernelInfoData {
Expand All @@ -85,6 +89,10 @@ template <char...> struct KernelInfoData {
}
static constexpr const char *getName() { return ""; }
static constexpr bool isESIMD() { return 0; }
static constexpr const char *getFileName() { return ""; }
static constexpr const char *getFunctionName() { return ""; }
static constexpr unsigned getLineNumber() { return 0; }
static constexpr unsigned getColumnNumber() { return 0; }
};

// C++14 like index_sequence and make_index_sequence
Expand Down Expand Up @@ -123,6 +131,10 @@ template <class KernelNameType> struct KernelInfo {
}
static constexpr const char *getName() { return SubKernelInfo::getName(); }
static constexpr bool isESIMD() { return SubKernelInfo::isESIMD(); }
static constexpr const char *getFileName() { return ""; }
static constexpr const char *getFunctionName() { return ""; }
static constexpr unsigned getLineNumber() { return 0; }
static constexpr unsigned getColumnNumber() { return 0; }
};
#endif //__SYCL_UNNAMED_LAMBDA__

Expand Down