Skip to content

Commit d97c89d

Browse files
vladimirlazvladimirlaz
committed
[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]>
1 parent d37cadc commit d97c89d

File tree

9 files changed

+26
-21
lines changed

9 files changed

+26
-21
lines changed

‎clang/lib/CodeGen/BackendUtil.cpp

Lines changed: 15 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -704,9 +704,6 @@ void EmitAssemblyHelper::CreatePasses(legacy::PassManager &MPM,
704704
if (!CodeGenOpts.SampleProfileFile.empty())
705705
PMBuilder.PGOSampleUse = CodeGenOpts.SampleProfileFile;
706706

707-
if (LangOpts.SYCL)
708-
MPM.add(createASFixerPass());
709-
710707
PMBuilder.populateFunctionPassManager(FPM);
711708
PMBuilder.populateModulePassManager(MPM);
712709
}
@@ -842,10 +839,15 @@ void EmitAssemblyHelper::EmitAssembly(BackendAction Action,
842839
}
843840
break;
844841

845-
846842
case Backend_EmitSPIRV:
847-
if (LangOpts.SYCL)
843+
if (LangOpts.SYCL) {
848844
SPIRV::SPIRVNoDerefAttr = true;
845+
// TODO: this pass added to work around missing linkonce_odr in SPIR-V
846+
PerModulePasses.add(
847+
createAlwaysInlinerLegacyPass(true /*InsertLifetimeIntrinsics*/));
848+
PerModulePasses.add(createASFixerPass());
849+
PerModulePasses.add(createDeadStoreEliminationPass());
850+
}
849851
PerModulePasses.add(createSPIRVWriterPass(*OS));
850852

851853
break;
@@ -1086,6 +1088,14 @@ void EmitAssemblyHelper::EmitAssemblyWithNewPassManager(
10861088
break;
10871089

10881090
case Backend_EmitSPIRV:
1091+
if (LangOpts.SYCL) {
1092+
SPIRV::SPIRVNoDerefAttr = true;
1093+
// TODO: this pass added to work around missing linkonce_odr in SPIR-V
1094+
CodeGenPasses.add(
1095+
createAlwaysInlinerLegacyPass(true /*InsertLifetimeIntrinsics*/));
1096+
CodeGenPasses.add(createASFixerPass());
1097+
CodeGenPasses.add(createDeadStoreEliminationPass());
1098+
}
10891099
CodeGenPasses.add(createSPIRVWriterPass(*OS));
10901100
break;
10911101

‎clang/lib/Driver/ToolChains/Clang.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3500,6 +3500,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
35003500
llvm::Triple(llvm::sys::getProcessTriple()).normalize();
35013501
CmdArgs.push_back("-aux-triple");
35023502
CmdArgs.push_back(Args.MakeArgString(NormalizedTriple));
3503+
CmdArgs.push_back("-disable-llvm-passes");
35033504
}
35043505

35053506
if (IsOpenMPDevice) {

‎clang/lib/Frontend/CompilerInvocation.cpp

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -121,12 +121,10 @@ CompilerInvocationBase::~CompilerInvocationBase() = default;
121121
static unsigned getOptimizationLevel(ArgList &Args, InputKind IK,
122122
DiagnosticsEngine &Diags) {
123123
unsigned DefaultOpt = llvm::CodeGenOpt::None;
124-
if (IK.getLanguage() == InputKind::OpenCL && !Args.hasArg(OPT_cl_opt_disable))
124+
if (IK.getLanguage() == InputKind::OpenCL &&
125+
!Args.hasArg(OPT_cl_opt_disable) || Args.hasArg(OPT_fsycl_is_device))
125126
DefaultOpt = llvm::CodeGenOpt::Default;
126127

127-
if (Args.hasArg(OPT_emit_spirv))
128-
return 0; // LLVM-SPIRV translator expects not optimized IR
129-
130128
if (Arg *A = Args.getLastArg(options::OPT_O_Group)) {
131129
if (A->getOption().matches(options::OPT_O0))
132130
return llvm::CodeGenOpt::None;

‎clang/test/CodeGenSYCL/device-functions.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang -cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -emit-llvm -x c++ %s -o - | FileCheck %s
1+
// 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
22

33
template <typename T>
44
T bar(T arg);

‎clang/test/CodeGenSYCL/emit-kernel-in-virtual-func.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang -cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -emit-llvm -x c++ %s -o - | FileCheck %s
1+
// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -emit-llvm -x c++ %s -o - | FileCheck %s
22

33
template <typename name, typename Func>
44
__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
@@ -27,6 +27,4 @@ int main() {
2727

2828
// Ensure that the SPIR-Kernel function is actually emitted.
2929
// CHECK: define spir_kernel void @FF
30-
// CHECK: call spir_func void @_ZZN7DERIVEDIiE10initializeEvENKUlvE_clEv
31-
// CHECK: define linkonce_odr spir_func void @_ZZN7DERIVEDIiE10initializeEvENKUlvE_clEv
3230

‎clang/test/CodeGenSYCL/emit-kernel-in-virtual-func2.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang -cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -emit-llvm -x c++ %s -o - | FileCheck %s
1+
// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -emit-llvm -x c++ %s -o - | FileCheck %s
22

33
template <typename name, typename Func>
44
__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
@@ -32,6 +32,4 @@ int main() {
3232

3333
// Ensure that the SPIR-Kernel function is actually emitted.
3434
// CHECK: define spir_kernel void @PP
35-
// CHECK: call spir_func void @_ZZ2TTIiEvvENKUlvE_clEv
36-
// CHECK: define linkonce_odr spir_func void @_ZZ2TTIiEvvENKUlvE_clEv
3735

‎clang/test/CodeGenSYCL/kernel-with-id.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -emit-llvm %s -o - | FileCheck %s
1+
// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm %s -o - | FileCheck %s
22

33
namespace cl {
44
namespace sycl {
@@ -69,8 +69,8 @@ __attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
6969

7070
int main() {
7171
cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorA;
72-
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EE13__set_pointerEPU3AS1i(%"class.cl::sycl::accessor"* %1, i32 addrspace(1)* %2)
73-
// 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)
72+
// CHECK: call spir_func void @{{.*}}__set_pointer{{.*}}(%"class.cl::sycl::accessor"* %{{.*}}, i32 addrspace(1)* %{{.*}})
73+
// CHECK: call spir_func void @{{.*}}__set_range{{.*}}(%"class.cl::sycl::accessor"* %{{.*}}, %"struct.cl::sycl::range"* byval align 1 %{{.*}})
7474
kernel<class kernel_function>(
7575
[=]() {
7676
accessorA.use();

‎clang/test/CodeGenSYCL/spir-calling-conv.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -emit-llvm -x c++ %s -o - | FileCheck %s
1+
// 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
22

33
template <typename name, typename Func>
44
__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {

‎clang/test/CodeGenSYCL/spir-no-deref-attr.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang -cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -emit-spirv -x c++ %s -o %t.spv
1+
// 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
22
// RUN: llvm-spirv %t.spv -to-text -o %t.txt
33
// RUN: FileCheck < %t.txt %s --check-prefix=CHECK
44

0 commit comments

Comments
 (0)