Skip to content

[SYCL] Emit integration header even if no kernels provided #5899

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
bader merged 4 commits into from
Mar 30, 2022
Merged

[SYCL] Emit integration header even if no kernels provided #5899

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
e013d99
[SYCL] Emit integration header even if no kernels provided
Fznamznon Mar 28, 2022
7881a0f
Update clang/test/CodeGenSYCL/int_header_without_kernels.cpp
Fznamznon Mar 28, 2022
9a5a3ff
Add code comments
Fznamznon Mar 29, 2022
8aff4cd
Merge branch 'int-header-no-kernels' of https://github.com/Fznamznon/…
Fznamznon Mar 29, 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
6 changes: 4 additions & 2 deletions clang/include/clang/Sema/Sema.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -359,6 +359,8 @@ class SYCLIntegrationHeader {
Itr->updateKernelNames(Name, StableName);
}

void addDeviceGlobalMap() { NeedToEmitDeviceGlobalMap = true; }
Comment thread
smanna12 marked this conversation as resolved.
Outdated

private:
// Kernel actual parameter descriptor.
struct KernelParamDesc {
Expand Down Expand Up @@ -433,21 +435,21 @@ class SYCLIntegrationHeader {
llvm::SmallVector<SpecConstID, 4> SpecConsts;

Sema &S;

bool NeedToEmitDeviceGlobalMap = false;
Comment thread
smanna12 marked this conversation as resolved.
Outdated
};

class SYCLIntegrationFooter {
public:
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<const VarDecl *> GlobalVars;
void emitSpecIDName(raw_ostream &O, const VarDecl *VD);
bool DeviceGlobalsEmitted = false;
};

/// Tracks expected type during expression parsing, for use in code completion.
Expand Down
5 changes: 4 additions & 1 deletion clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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";
Expand Down Expand Up @@ -5020,6 +5020,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {

llvm::SmallSet<const VarDecl *, 8> Visited;
bool EmittedFirstSpecConstant = false;
bool DeviceGlobalsEmitted = false;

// Used to uniquely name the 'shim's as we generate the names in each
// anonymous namespace.
Expand Down Expand Up @@ -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;
}
Expand Down
38 changes: 38 additions & 0 deletions clang/test/CodeGenSYCL/int_header_without_kernels.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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 %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 <CL/sycl/detail/defines_elementary.hpp>

// CHECK-FOOTER: #include <CL/sycl/detail/device_global_map.hpp>
// CHECK-FOOTER: namespace sycl::detail {
// CHECK-FOOTER-NEXT: namespace {
// CHECK-FOOTER-NEXT: __sycl_device_global_registration::__sycl_device_global_registration() noexcept {

device_global<int> Basic;
// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::Basic, "_Z5Basic");

// CHECK-FOOTER-NEXT: }
// CHECK-FOOTER-NEXT: }
// CHECK-FOOTER-NEXT: }