From e013d9993edaa84bfccc00e5d0fcd1b1a2fe72e1 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Mon, 28 Mar 2022 13:27:50 +0300 Subject: [PATCH 1/3] [SYCL] Emit integration header even if no kernels provided Since device_global variables can be declared in one file and used in another, it is possible to receive an input with device_global and no kernels. Make sure integration header and footer are emitted in this case. --- clang/include/clang/Sema/Sema.h | 6 ++- clang/lib/Sema/SemaSYCL.cpp | 5 ++- .../int_header_without_kernels.cpp | 38 +++++++++++++++++++ 3 files changed, 46 insertions(+), 3 deletions(-) create mode 100644 clang/test/CodeGenSYCL/int_header_without_kernels.cpp diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index e0116feddc39..d09705446b76 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -359,6 +359,8 @@ class SYCLIntegrationHeader { Itr->updateKernelNames(Name, StableName); } + void addDeviceGlobalMap() { NeedToEmitDeviceGlobalMap = true; } + private: // Kernel actual parameter descriptor. struct KernelParamDesc { @@ -433,6 +435,8 @@ class SYCLIntegrationHeader { llvm::SmallVector SpecConsts; Sema &S; + + bool NeedToEmitDeviceGlobalMap = false; }; class SYCLIntegrationFooter { @@ -440,14 +444,12 @@ class SYCLIntegrationFooter { SYCLIntegrationFooter(Sema &S) : S(S) {} bool emit(StringRef MainSrc); void addVarDecl(const VarDecl *VD); - bool isDeviceGlobalsEmitted() { return DeviceGlobalsEmitted; } private: bool emit(raw_ostream &O); Sema &S; llvm::SmallVector GlobalVars; void emitSpecIDName(raw_ostream &O, const VarDecl *VD); - bool DeviceGlobalsEmitted = false; }; /// Tracks expected type during expression parsing, for use in code completion. diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8174a24bbc27..49c9d39ddd73 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -4656,7 +4656,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { // whose sole purpose is to run its constructor before the application's // main() function. - if (S.getSyclIntegrationFooter().isDeviceGlobalsEmitted()) { + if (NeedToEmitDeviceGlobalMap) { O << "namespace {\n"; O << "class __sycl_device_global_registration {\n"; @@ -5020,6 +5020,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { llvm::SmallSet Visited; bool EmittedFirstSpecConstant = false; + bool DeviceGlobalsEmitted = false; // Used to uniquely name the 'shim's as we generate the names in each // anonymous namespace. @@ -5103,6 +5104,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { OS << "}\n"; OS << "} // namespace (unnamed)\n"; OS << "} // namespace sycl::detail\n"; + + S.getSyclIntegrationHeader().addDeviceGlobalMap(); } return true; } diff --git a/clang/test/CodeGenSYCL/int_header_without_kernels.cpp b/clang/test/CodeGenSYCL/int_header_without_kernels.cpp new file mode 100644 index 000000000000..6f9ab2403d75 --- /dev/null +++ b/clang/test/CodeGenSYCL/int_header_without_kernels.cpp @@ -0,0 +1,38 @@ +// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -fsycl-int-header=%t.header.h -fsycl-unique-prefix=THE_PREFIX %s -emit-llvm -o %t.ll +// RUN: FileCheck -input-file=%t.footer.h %s --check-prefix=CHECK-FOOTER +// RUN: FileCheck -input-file=%t.header.h %s --check-prefix=CHECK-HEADER + +// This test checks that integration header and footer are emitted correctly +// for device_global variables even without kernels. + +#include "sycl.hpp" + +using namespace cl::sycl::ext::oneapi; + +// CHECK-HEADER: __SYCL_INLINE_NAMESPACE(cl) { +// CHECK-HEADER-NEXT: namespace sycl { +// CHECK-HEADER-NEXT: namespace detail { +// CHECK-HEADER-NEXT: namespace { +// CHECK-HEADER-NEXT: class __sycl_device_global_registration { +// CHECK-HEADER-NEXT: public: +// CHECK-HEADER-NEXT: __sycl_device_global_registration() noexcept; +// CHECK-HEADER-NEXT: }; +// CHECK-HEADER-NEXT: __sycl_device_global_registration __sycl_device_global_registrar; +// CHECK-HEADER-NEXT: } // namespace +// CHECK-HEADER: } // namespace detail +// CHECK-HEADER: } // namespace sycl +// CHECK-HEADER: } // __SYCL_INLINE_NAMESPACE(cl) + +// CHECK-FOOTER: #include + +// CHECK-FOOTER: #include +// CHECK-FOOTER: namespace sycl::detail { +// CHECK-FOOTER-NEXT: namespace { +// CHECK-FOOTER-NEXT: __sycl_device_global_registration::__sycl_device_global_registration() noexcept { + +device_global Basic; +// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::Basic, "_Z5Basic"); + +// CHECK-FOOTER-NEXT: } +// CHECK-FOOTER-NEXT: } +// CHECK-FOOTER-NEXT: } From 7881a0fbdff52f440c477fdc7c6fe0ae651ef2d5 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Mon, 28 Mar 2022 22:01:10 +0300 Subject: [PATCH 2/3] Update clang/test/CodeGenSYCL/int_header_without_kernels.cpp Co-authored-by: premanandrao --- clang/test/CodeGenSYCL/int_header_without_kernels.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/CodeGenSYCL/int_header_without_kernels.cpp b/clang/test/CodeGenSYCL/int_header_without_kernels.cpp index 6f9ab2403d75..14fad42a87ab 100644 --- a/clang/test/CodeGenSYCL/int_header_without_kernels.cpp +++ b/clang/test/CodeGenSYCL/int_header_without_kernels.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -fsycl-int-header=%t.header.h -fsycl-unique-prefix=THE_PREFIX %s -emit-llvm -o %t.ll +// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -fsycl-int-header=%t.header.h %s -emit-llvm -o %t.ll // RUN: FileCheck -input-file=%t.footer.h %s --check-prefix=CHECK-FOOTER // RUN: FileCheck -input-file=%t.header.h %s --check-prefix=CHECK-HEADER From 9a5a3ff6d2e19853c5e85daa0f477d5e78dbf898 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 29 Mar 2022 14:18:17 +0300 Subject: [PATCH 3/3] Add code comments --- clang/include/clang/Sema/Sema.h | 11 +++++++++-- clang/lib/Sema/SemaSYCL.cpp | 4 ++-- 2 files changed, 11 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index d09705446b76..9d9e15553665 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -359,7 +359,12 @@ class SYCLIntegrationHeader { Itr->updateKernelNames(Name, StableName); } - void addDeviceGlobalMap() { NeedToEmitDeviceGlobalMap = true; } + /// Signals that emission of __sycl_device_global_registration type and + /// declaration of variable __sycl_device_global_registrar of this type in + /// integration header is required. + void addDeviceGlobalRegistration() { + NeedToEmitDeviceGlobalRegistration = true; + } private: // Kernel actual parameter descriptor. @@ -436,7 +441,9 @@ class SYCLIntegrationHeader { Sema &S; - bool NeedToEmitDeviceGlobalMap = false; + /// Keeps track of whether declaration of __sycl_device_global_registration + /// type and __sycl_device_global_registrar variable are required to emit. + bool NeedToEmitDeviceGlobalRegistration = false; }; class SYCLIntegrationFooter { diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 49c9d39ddd73..7ade4155c328 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -4656,7 +4656,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { // whose sole purpose is to run its constructor before the application's // main() function. - if (NeedToEmitDeviceGlobalMap) { + if (NeedToEmitDeviceGlobalRegistration) { O << "namespace {\n"; O << "class __sycl_device_global_registration {\n"; @@ -5105,7 +5105,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { OS << "} // namespace (unnamed)\n"; OS << "} // namespace sycl::detail\n"; - S.getSyclIntegrationHeader().addDeviceGlobalMap(); + S.getSyclIntegrationHeader().addDeviceGlobalRegistration(); } return true; }