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 all 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
42 changes: 42 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,44 @@ 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() {\n";
O << "#ifndef NDEBUG\n";
O << " return \""
<< std::string(PLoc.getFilename())
.substr(std::string(PLoc.getFilename()).find_last_of("/\\") + 1);
O << "\";\n";
O << "#else\n";
O << " return \"\";\n";
O << "#endif\n";
O << " }\n";
O << " __SYCL_DLL_LOCAL\n";
O << " static constexpr const char* getFunctionName() {\n";
O << "#ifndef NDEBUG\n";
O << " return \"";
SYCLKernelNameTypePrinter Printer(O, Policy);
Printer.Visit(K.NameType);
O << "\";\n";
O << "#else\n";
O << " return \"\";\n";
O << "#endif\n";
O << " }\n";
O << " __SYCL_DLL_LOCAL\n";
O << " static constexpr unsigned getLineNumber() {\n";
O << "#ifndef NDEBUG\n";
O << " return " << PLoc.getLine() << ";\n";
O << "#else\n";
O << " return 0;\n";
O << "#endif\n";
O << " }\n";
O << " __SYCL_DLL_LOCAL\n";
O << " static constexpr unsigned getColumnNumber() {\n";
O << "#ifndef NDEBUG\n";
O << " return " << PLoc.getColumn() << ";\n";
O << "#else\n";
O << " return 0;\n";
O << "#endif\n";
O << " }\n";
O << "};\n";
CurStart += N;
}
Expand Down
308 changes: 308 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,308 @@
// 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

#include "Inputs/sycl.hpp"

// Check that meaningful information is returned when NDEBUG is not defined
// and empty strings and 0s are emitted when it is.
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;
}
// 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() {
// CHECK: #ifndef NDEBUG
// CHECK: return "code_location.cpp";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr const char* getFunctionName() {
// CHECK: #ifndef NDEBUG
// CHECK: return "";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr unsigned getLineNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 11;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr unsigned getColumnNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 54;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: };

// CHECK: template <> struct KernelInfo<KernelName> {
// CHECK: static constexpr const char* getFileName() {
// CHECK: #ifndef NDEBUG
// CHECK: return "code_location.cpp";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: __SYCL_DLL_LOCAL
// CHECK: static constexpr const char* getFunctionName() {
// CHECK: #ifndef NDEBUG
// CHECK: return "KernelName";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: __SYCL_DLL_LOCAL
// CHECK: static constexpr unsigned getLineNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 12;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: __SYCL_DLL_LOCAL
// CHECK: static constexpr unsigned getColumnNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 72;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: };

// Check that the right name and location is returned when
// lambda and kernel name are defined on different lines
class KernelName2;
int test2() {
cl::sycl::queue q;
q.submit([&](cl::sycl::handler &h) { h.single_task<KernelName2>(
[] { int i = 2; }); });
return 0;
}
// CHECK: template <> struct KernelInfo<::KernelName2> {
// CHECK: static constexpr const char* getFileName() {
// CHECK: #ifndef NDEBUG
// CHECK: return "code_location.cpp";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr const char* getFunctionName() {
// CHECK: #ifndef NDEBUG
// CHECK: return "::KernelName2";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr unsigned getLineNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 86;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr unsigned getColumnNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 44;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: };

// Check that fully qualified name is returned
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() {
// CHECK: #ifndef NDEBUG
// CHECK: return "code_location.cpp";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr const char* getFunctionName() {
// CHECK: #ifndef NDEBUG
// CHECK: return "::KernelName3<::KernelName2>";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr unsigned getLineNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 125;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr unsigned getColumnNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 44;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: };

// Check that the location information returned is that of l4
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() {
// CHECK: #ifndef NDEBUG
// CHECK: return "code_location.cpp";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr const char* getFunctionName() {
// CHECK: #ifndef NDEBUG
// CHECK: return "KernelName4";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr unsigned getLineNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 160;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr unsigned getColumnNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 11;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: };

// Check that fully qualified name is returned when unnamed lambda
// kernel is enclosed in a namespace
namespace NS {
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;
}
} // namespace NS
// CHECK: template <> struct KernelInfoData<'_', 'Z', 'T', 'S', 'Z', 'Z', 'N', '2', 'N', 'S', '5', 't', 'e', 's', 't', '5', 'E', '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', '3', '_', 'E', 'U', 'l', 'v', 'E', '_'> {
// CHECK: static constexpr const char* getFileName() {
// CHECK: #ifndef NDEBUG
// CHECK: return "code_location.cpp";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr const char* getFunctionName() {
// CHECK: #ifndef NDEBUG
// CHECK: return "NS::";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr unsigned getLineNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 202;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr unsigned getColumnNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 54;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: };
// CHECK: template <> struct KernelInfo<NS::KernelName5> {
// CHECK: static constexpr const char* getFileName() {
// CHECK: #ifndef NDEBUG
// CHECK: return "code_location.cpp";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr const char* getFunctionName() {
// CHECK: #ifndef NDEBUG
// CHECK: return "NS::KernelName5";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr unsigned getLineNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 203;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr unsigned getColumnNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 73;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: };

// Check that the location information returned is that of the Functor
struct Functor {
void operator()() const {
}
};
int test6() {
Functor F;
cl::sycl::queue q;
q.submit([=](cl::sycl::handler &h) { h.single_task<class KernelName6>(F); });
return 0;
}
// CHECK: template <> struct KernelInfo<KernelName6> {
// CHECK: static constexpr const char* getFileName() {
// CHECK: #ifndef NDEBUG
// CHECK: return "code_location.cpp";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr const char* getFunctionName() {
// CHECK: #ifndef NDEBUG
// CHECK: return "KernelName6";
// CHECK: #else
// CHECK: return "";
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr unsigned getLineNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 269;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: static constexpr unsigned getColumnNumber() {
// CHECK: #ifndef NDEBUG
// CHECK: return 8;
// CHECK: #else
// CHECK: return 0;
// CHECK: #endif
// CHECK: }
// CHECK: };
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