From 45439cca14fc26790b5d768ef2c93fcdfa44a9af Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Mon, 28 Mar 2022 22:04:58 +0300
Subject: [PATCH 1/2] [SYCL] Fix device code outlining for static local
 variables

When static variable was declared in host code and used from device code
that forced emission of parent host function. It could cause emission of
the code that is not valid for device (but still valid for host). This
change makes sure that doesn't happen anymore.
---
 clang/lib/CodeGen/CGDecl.cpp                  |  5 ++
 clang/lib/CodeGen/CGExpr.cpp                  |  8 +++
 .../CodeGenSYCL/static-vars-in-host-code.cpp  | 50 +++++++++++++++++++
 3 files changed, 63 insertions(+)
 create mode 100644 clang/test/CodeGenSYCL/static-vars-in-host-code.cpp

diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp
index 89435c95348e2..278e872cd4f89 100644
--- a/clang/lib/CodeGen/CGDecl.cpp
+++ b/clang/lib/CodeGen/CGDecl.cpp
@@ -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());
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index a79612feae938..d1642681824bf 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -2841,6 +2841,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));
 
diff --git a/clang/test/CodeGenSYCL/static-vars-in-host-code.cpp b/clang/test/CodeGenSYCL/static-vars-in-host-code.cpp
new file mode 100644
index 0000000000000..9ac165b71cb60
--- /dev/null
+++ b/clang/test/CodeGenSYCL/static-vars-in-host-code.cpp
@@ -0,0 +1,50 @@
+// 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;
+}

From 9b45f593a8308b76a91401e86875e6d5a59feb8d Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Tue, 29 Mar 2022 11:19:26 +0300
Subject: [PATCH 2/2] Make clang-format happy

---
 clang/test/CodeGenSYCL/static-vars-in-host-code.cpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/clang/test/CodeGenSYCL/static-vars-in-host-code.cpp b/clang/test/CodeGenSYCL/static-vars-in-host-code.cpp
index 9ac165b71cb60..e4c56dedca0ed 100644
--- a/clang/test/CodeGenSYCL/static-vars-in-host-code.cpp
+++ b/clang/test/CodeGenSYCL/static-vars-in-host-code.cpp
@@ -17,7 +17,6 @@
 // CHECK: define{{.*}}@_Z3fooi(i32 noundef %In)
 // CHECK-NOT: define{{.*}}@main()
 
-
 template <class T> struct S {
   static const T Member = 1;
   int Parrots = 38;