Skip to content

[SYCL] Fix spec constants support in integration header #2896

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 2 commits into from
Dec 16, 2020
Merged

[SYCL] Fix spec constants support in integration header #2896

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
db34864
[SYCL] Fix spec constants support in integration header
AlexeySachkov Dec 15, 2020
9b7086f
Fix clang-format & refactor
AlexeySachkov Dec 16, 2020
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
12 changes: 9 additions & 3 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3757,8 +3757,14 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
PrintingPolicy Policy(LO);
Policy.SuppressTypedefs = true;
Policy.SuppressUnwrittenScope = true;
SYCLFwdDeclEmitter FwdDeclEmitter(O, S.getLangOpts());

if (SpecConsts.size() > 0) {
O << "// Forward declarations of templated spec constant types:\n";
for (const auto &SC : SpecConsts)
FwdDeclEmitter.Visit(SC.first);
O << "\n";

// Remove duplicates.
std::sort(SpecConsts.begin(), SpecConsts.end(),
[](const SpecConstID &SC1, const SpecConstID &SC2) {
Expand All @@ -3772,10 +3778,12 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
// Here can do faster comparison of types.
return SC1.first == SC2.first;
});

Copy link
Contributor

Choose a reason for hiding this comment

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

Nit:

Suggested change

O << "// Specialization constants IDs:\n";
for (const auto &P : llvm::make_range(SpecConsts.begin(), End)) {
O << "template <> struct sycl::detail::SpecConstantInfo<";
O << P.first.getAsString(Policy);
SYCLKernelNameTypePrinter Printer(O, Policy);
Printer.Visit(P.first);
O << "> {\n";
O << " static constexpr const char* getName() {\n";
O << " return \"" << P.second << "\";\n";
Expand All @@ -3786,8 +3794,6 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {

if (!UnnamedLambdaSupport) {
O << "// Forward declarations of templated kernel function types:\n";

SYCLFwdDeclEmitter FwdDeclEmitter(O, S.getLangOpts());
for (const KernelDesc &K : KernelDescs)
FwdDeclEmitter.Visit(K.NameType);
}
Expand Down
32 changes: 29 additions & 3 deletions clang/test/CodeGenSYCL/int_header_spec_const.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,10 @@ class MyUInt32Const;
class MyFloatConst;
class MyDoubleConst;

namespace test {
class MySpecConstantWithinANamespace;
};

int main() {
// Create specialization constants.
cl::sycl::ONEAPI::experimental::spec_constant<bool, MyBoolConst> i1(false);
Expand All @@ -32,13 +36,31 @@ int main() {
cl::sycl::ONEAPI::experimental::spec_constant<unsigned int, MyUInt32Const> ui32(0);
cl::sycl::ONEAPI::experimental::spec_constant<float, MyFloatConst> f32(0);
cl::sycl::ONEAPI::experimental::spec_constant<double, MyDoubleConst> f64(0);
// Kernel name can be used as a spec constant name
cl::sycl::ONEAPI::experimental::spec_constant<int, SpecializedKernel> spec1(0);
// Spec constant name can be declared within a namespace
cl::sycl::ONEAPI::experimental::spec_constant<int, test::MySpecConstantWithinANamespace> spec2(0);

double val;
double *ptr = &val; // to avoid "unused" warnings

// CHECK: // Forward declarations of templated spec constant types:
// CHECK: class MyInt8Const;
// CHECK: class MyUInt8Const;
// CHECK: class MyInt16Const;
// CHECK: class MyUInt16Const;
// CHECK: class MyInt32Const;
// CHECK: class MyUInt32Const;
// CHECK: class MyFloatConst;
// CHECK: class MyDoubleConst;
// CHECK: class SpecializedKernel;
// CHECK: namespace test {
// CHECK: class MySpecConstantWithinANamespace;
// CHECK: }

cl::sycl::kernel_single_task<SpecializedKernel>([=]() {
*ptr = i1.get() +
// CHECK-DAG: template <> struct sycl::detail::SpecConstantInfo<class MyBoolConst> {
// CHECK-DAG: template <> struct sycl::detail::SpecConstantInfo<::MyBoolConst> {
// CHECK-DAG-NEXT: static constexpr const char* getName() {
// CHECK-DAG-NEXT: return "_ZTS11MyBoolConst";
// CHECK-DAG-NEXT: }
Expand All @@ -58,7 +80,11 @@ int main() {
// CHECK-DAG: return "_ZTS13MyUInt32Const";
f32.get() +
// CHECK-DAG: return "_ZTS12MyFloatConst";
f64.get();
// CHECK-DAG: return "_ZTS13MyDoubleConst";
f64.get() +
// CHECK-DAG: return "_ZTS13MyDoubleConst";
spec1.get() +
// CHECK-DAG: return "_ZTS17SpecializedKernel"
spec2.get();
// CHECK-DAG: return "_ZTSN4test30MySpecConstantWithinANamespaceE"
});
}