Skip to content

[SYCL] Fix device code outlining for static local variables #5915

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.

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 3 commits into from
Apr 1, 2022
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
5 changes: 5 additions & 0 deletions clang/lib/CodeGen/CGDecl.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -296,6 +296,11 @@ llvm::Constant *CodeGenModule::getOrCreateStaticVarDecl(

setStaticLocalDeclAddress(&D, Addr);

// Do not force emission of the parent funtion since it can be a host function
// that contains illegal code for SYCL device.
if (getLangOpts().SYCLIsDevice)
return Addr;

// Ensure that the static local gets initialized by making sure the parent
// function gets emitted eventually.
const Decl *DC = cast<Decl>(D.getDeclContext());
Expand Down
8 changes: 8 additions & 0 deletions clang/lib/CodeGen/CGExpr.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -2840,6 +2840,14 @@ LValue CodeGenFunction::EmitDeclRefLValue(const DeclRefExpr *E) {
} else if (VD->isStaticLocal()) {
llvm::Constant *var = CGM.getOrCreateStaticVarDecl(
*VD, CGM.getLLVMLinkageVarDefinition(VD, /*IsConstant=*/false));

// Force completion of static variable for SYCL since if it wasn't emitted
// already that means it is defined in host code and its parent function
// won't be emitted.
if (getLangOpts().SYCLIsDevice)
EmitStaticVarDecl(
*VD, CGM.getLLVMLinkageVarDefinition(VD, /*IsConstant=*/false));

addr = Address(
var, ConvertTypeForMem(VD->getType()), getContext().getDeclAlign(VD));

Expand Down
49 changes: 49 additions & 0 deletions clang/test/CodeGenSYCL/static-vars-in-host-code.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,49 @@
// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s

// This test checks that static variables defined in host code and used in
// device code do not force emission of their parent host functions.

#include "sycl.hpp"

// CHECK-NOT: class.cl::sycl::queue

// CHECK: @_ZZ4mainE3Loc = internal addrspace(1) constant i32 42, align 4
// CHECK: @_ZZ4mainE6Struct = internal addrspace(1) constant %struct.S { i32 38 }, align 4
// CHECK: @_ZL4Glob = internal addrspace(1) constant i64 100500, align 8
// CHECK: @_ZZZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_ENKUlvE_clEvE6InKern = internal addrspace(1) constant i32 2022, align 4
// CHECK: @_ZN1SIiE6MemberE = available_externally addrspace(1) constant i32 1, align 4
// CHECK: @_ZZ3fooiE5InFoo = internal addrspace(1) constant i32 300, align 4

// CHECK: define{{.*}}@_Z3fooi(i32 noundef %In)
// CHECK-NOT: define{{.*}}@main()

template <class T> struct S {
static const T Member = 1;
int Parrots = 38;
};

static constexpr unsigned long Glob = 100500;
int foo(const int In) {
static constexpr int InFoo = 300;
return InFoo + In;
}

int main() {
sycl::queue q;
static constexpr int Loc = 42;
static const S<int> Struct;
q.submit([&](sycl::handler &cgh) {
cgh.single_task<class TheKernel>([=]() {
(void)Loc;
(void)Struct;

// Make sure other use cases with statics are not broken by the change.
(void)Glob;
static const int InKern = 2022;
foo(Loc);
(void)S<int>::Member;
});
});

return 0;
}