From db348644f105bad60de8539aab7f467c22b7db06 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 15 Dec 2020 19:16:52 +0300 Subject: [PATCH 1/2] [SYCL] Fix spec constants support in integration header Added emissions of forward-declarations of types used as names for specialization constants, which allows to use types declared within namespaces as specialization constant names. --- clang/lib/Sema/SemaSYCL.cpp | 10 ++++++- .../CodeGenSYCL/int_header_spec_const.cpp | 30 +++++++++++++++++-- 2 files changed, 37 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c15e58e9d28b1..85a91d213c410 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3759,6 +3759,12 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { Policy.SuppressUnwrittenScope = true; if (SpecConsts.size() > 0) { + O << "// Forward declarations of templated spec constant types:\n"; + SYCLFwdDeclEmitter FwdDeclEmitter(O, S.getLangOpts()); + 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) { @@ -3772,10 +3778,12 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { // Here can do faster comparison of types. return SC1.first == SC2.first; }); + 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"; diff --git a/clang/test/CodeGenSYCL/int_header_spec_const.cpp b/clang/test/CodeGenSYCL/int_header_spec_const.cpp index 39e3d2b05ad0f..bfeb573e70123 100644 --- a/clang/test/CodeGenSYCL/int_header_spec_const.cpp +++ b/clang/test/CodeGenSYCL/int_header_spec_const.cpp @@ -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 i1(false); @@ -32,13 +36,31 @@ int main() { cl::sycl::ONEAPI::experimental::spec_constant ui32(0); cl::sycl::ONEAPI::experimental::spec_constant f32(0); cl::sycl::ONEAPI::experimental::spec_constant f64(0); + // Kernel name can be used as a spec constant name + cl::sycl::ONEAPI::experimental::spec_constant spec1(0); + // Spec constant name can be declared within a namespace + cl::sycl::ONEAPI::experimental::spec_constant 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([=]() { *ptr = i1.get() + - // CHECK-DAG: template <> struct sycl::detail::SpecConstantInfo { + // CHECK-DAG: template <> struct sycl::detail::SpecConstantInfo<::MyBoolConst> { // CHECK-DAG-NEXT: static constexpr const char* getName() { // CHECK-DAG-NEXT: return "_ZTS11MyBoolConst"; // CHECK-DAG-NEXT: } @@ -59,6 +81,10 @@ int main() { f32.get() + // CHECK-DAG: return "_ZTS12MyFloatConst"; f64.get(); - // CHECK-DAG: return "_ZTS13MyDoubleConst"; + // CHECK-DAG: return "_ZTS13MyDoubleConst"; + spec1.get(); + // CHECK-DAG: return "_ZTS17SpecializedKernel" + spec2.get(); + // CHECK-DAG: return "_ZTSN4test30MySpecConstantWithinANamespaceE" }); } From 9b7086f7cef079b80ac5e137394f8d77d5d49c3e Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 16 Dec 2020 12:30:40 +0300 Subject: [PATCH 2/2] Fix clang-format & refactor Moved definition of `SYCLFwdDeclEmitter` to a wider scope to avoid getting duplicated forward-declarations if there is a spec constant and a kernel with the same name. --- clang/lib/Sema/SemaSYCL.cpp | 4 +--- clang/test/CodeGenSYCL/int_header_spec_const.cpp | 8 ++++---- 2 files changed, 5 insertions(+), 7 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 85a91d213c410..bde9632b19e7d 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3757,10 +3757,10 @@ 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"; - SYCLFwdDeclEmitter FwdDeclEmitter(O, S.getLangOpts()); for (const auto &SC : SpecConsts) FwdDeclEmitter.Visit(SC.first); O << "\n"; @@ -3794,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); } diff --git a/clang/test/CodeGenSYCL/int_header_spec_const.cpp b/clang/test/CodeGenSYCL/int_header_spec_const.cpp index bfeb573e70123..c9ca63bcecf64 100644 --- a/clang/test/CodeGenSYCL/int_header_spec_const.cpp +++ b/clang/test/CodeGenSYCL/int_header_spec_const.cpp @@ -19,7 +19,7 @@ class MyFloatConst; class MyDoubleConst; namespace test { - class MySpecConstantWithinANamespace; +class MySpecConstantWithinANamespace; }; int main() { @@ -80,11 +80,11 @@ int main() { // CHECK-DAG: return "_ZTS13MyUInt32Const"; f32.get() + // CHECK-DAG: return "_ZTS12MyFloatConst"; - f64.get(); + f64.get() + // CHECK-DAG: return "_ZTS13MyDoubleConst"; - spec1.get(); + spec1.get() + // CHECK-DAG: return "_ZTS17SpecializedKernel" spec2.get(); - // CHECK-DAG: return "_ZTSN4test30MySpecConstantWithinANamespaceE" + // CHECK-DAG: return "_ZTSN4test30MySpecConstantWithinANamespaceE" }); }