[SYCL] Remove noinline attribute from SPIRV

This change gets rid of noinline attribute by switching to O2 optimization level with llvm passes disabled when --sycl flag is used. Signed-off-by: Vladimir Lazarev <[email protected]>
intel · Jan 22, 2019 · d97c89d · d97c89d
1 parent d37cadc
commit d97c89d
Show file tree

Hide file tree

Showing 9 changed files with 26 additions and 21 deletions.
diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp
@@ -704,9 +704,6 @@ void EmitAssemblyHelper::CreatePasses(legacy::PassManager &MPM,
   if (!CodeGenOpts.SampleProfileFile.empty())
     PMBuilder.PGOSampleUse = CodeGenOpts.SampleProfileFile;
 
-  if (LangOpts.SYCL)
-    MPM.add(createASFixerPass());
-
   PMBuilder.populateFunctionPassManager(FPM);
   PMBuilder.populateModulePassManager(MPM);
 }
@@ -842,10 +839,15 @@ void EmitAssemblyHelper::EmitAssembly(BackendAction Action,
     }
     break;
 
-
   case Backend_EmitSPIRV:
-    if (LangOpts.SYCL)
+    if (LangOpts.SYCL) {
       SPIRV::SPIRVNoDerefAttr = true;
+      // TODO: this pass added to work around missing linkonce_odr in SPIR-V
+      PerModulePasses.add(
+          createAlwaysInlinerLegacyPass(true /*InsertLifetimeIntrinsics*/));
+      PerModulePasses.add(createASFixerPass());
+      PerModulePasses.add(createDeadStoreEliminationPass());
+    }
     PerModulePasses.add(createSPIRVWriterPass(*OS));
 
     break;
@@ -1086,6 +1088,14 @@ void EmitAssemblyHelper::EmitAssemblyWithNewPassManager(
     break;
 
   case Backend_EmitSPIRV:
+    if (LangOpts.SYCL) {
+      SPIRV::SPIRVNoDerefAttr = true;
+      // TODO: this pass added to work around missing linkonce_odr in SPIR-V
+      CodeGenPasses.add(
+          createAlwaysInlinerLegacyPass(true /*InsertLifetimeIntrinsics*/));
+      CodeGenPasses.add(createASFixerPass());
+      CodeGenPasses.add(createDeadStoreEliminationPass());
+    }
     CodeGenPasses.add(createSPIRVWriterPass(*OS));
     break;
 

diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -3500,6 +3500,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
         llvm::Triple(llvm::sys::getProcessTriple()).normalize();
     CmdArgs.push_back("-aux-triple");
     CmdArgs.push_back(Args.MakeArgString(NormalizedTriple));
+    CmdArgs.push_back("-disable-llvm-passes");
   }
 
   if (IsOpenMPDevice) {

diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -121,12 +121,10 @@ CompilerInvocationBase::~CompilerInvocationBase() = default;
 static unsigned getOptimizationLevel(ArgList &Args, InputKind IK,
                                      DiagnosticsEngine &Diags) {
   unsigned DefaultOpt = llvm::CodeGenOpt::None;
-  if (IK.getLanguage() == InputKind::OpenCL && !Args.hasArg(OPT_cl_opt_disable))
+  if (IK.getLanguage() == InputKind::OpenCL &&
+      !Args.hasArg(OPT_cl_opt_disable) || Args.hasArg(OPT_fsycl_is_device))
     DefaultOpt = llvm::CodeGenOpt::Default;
 
-  if (Args.hasArg(OPT_emit_spirv))
-    return 0; // LLVM-SPIRV translator expects not optimized IR
-
   if (Arg *A = Args.getLastArg(options::OPT_O_Group)) {
     if (A->getOption().matches(options::OPT_O0))
       return llvm::CodeGenOpt::None;

diff --git a/clang/test/CodeGenSYCL/device-functions.cpp b/clang/test/CodeGenSYCL/device-functions.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang -cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -emit-llvm -x c++ %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm -x c++ %s -o - | FileCheck %s
 
 template <typename T>
 T bar(T arg);

diff --git a/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func.cpp b/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang -cc1 -triple spir64-unknown-linux-sycldevice  -std=c++11 -fsycl-is-device -emit-llvm -x c++ %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice  -std=c++11 -fsycl-is-device -emit-llvm -x c++ %s -o - | FileCheck %s
 
 template <typename name, typename Func>
 __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
@@ -27,6 +27,4 @@ int main() {
 
 // Ensure that the SPIR-Kernel function is actually emitted.
 // CHECK: define spir_kernel void @FF
-// CHECK: call spir_func void @_ZZN7DERIVEDIiE10initializeEvENKUlvE_clEv
-// CHECK: define linkonce_odr spir_func void @_ZZN7DERIVEDIiE10initializeEvENKUlvE_clEv
 
diff --git a/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func2.cpp b/clang/test/CodeGenSYCL/emit-kernel-in-virtual-func2.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang -cc1 -triple spir64-unknown-linux-sycldevice  -std=c++11 -fsycl-is-device -emit-llvm -x c++ %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice  -std=c++11 -fsycl-is-device -emit-llvm -x c++ %s -o - | FileCheck %s
 
 template <typename name, typename Func>
 __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
@@ -32,6 +32,4 @@ int main() {
 
 // Ensure that the SPIR-Kernel function is actually emitted.
 // CHECK: define spir_kernel void @PP
-// CHECK: call spir_func void @_ZZ2TTIiEvvENKUlvE_clEv
-// CHECK: define linkonce_odr spir_func void @_ZZ2TTIiEvvENKUlvE_clEv
 
diff --git a/clang/test/CodeGenSYCL/kernel-with-id.cpp b/clang/test/CodeGenSYCL/kernel-with-id.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm %s -o - | FileCheck %s
 
 namespace cl {
 namespace sycl {
@@ -69,8 +69,8 @@ __attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
 
 int main() {
   cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorA;
-// CHECK: call spir_func void @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EE13__set_pointerEPU3AS1i(%"class.cl::sycl::accessor"* %1, i32 addrspace(1)* %2)
-// CHECK: call spir_func void @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EE11__set_rangeENS0_5rangeILi1EEE(%"class.cl::sycl::accessor"* %3, %"struct.cl::sycl::range"* byval align 1 %agg.tmp)
+// CHECK: call spir_func void @{{.*}}__set_pointer{{.*}}(%"class.cl::sycl::accessor"* %{{.*}}, i32 addrspace(1)* %{{.*}})
+// CHECK: call spir_func void @{{.*}}__set_range{{.*}}(%"class.cl::sycl::accessor"* %{{.*}}, %"struct.cl::sycl::range"* byval align 1 %{{.*}})
     kernel<class kernel_function>(
       [=]() {
         accessorA.use();

diff --git a/clang/test/CodeGenSYCL/spir-calling-conv.cpp b/clang/test/CodeGenSYCL/spir-calling-conv.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -emit-llvm -x c++ %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm -x c++ %s -o - | FileCheck %s
 
 template <typename name, typename Func>
 __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {

diff --git a/clang/test/CodeGenSYCL/spir-no-deref-attr.cpp b/clang/test/CodeGenSYCL/spir-no-deref-attr.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang -cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -emit-spirv -x c++ %s -o %t.spv
+// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-spirv -x c++ %s -o %t.spv
 // RUN: llvm-spirv %t.spv -to-text -o %t.txt
 // RUN: FileCheck < %t.txt %s --check-prefix=CHECK