Skip to content

Commit 1aed422

Browse files
maarquitos14maarquitos14
authored
[SYCL][ESIMD] Run all passes with O0 opt level (#19554)
In #19411, we simply went for running all passes regardless of the optimization level. However, the change made it quite difficult to test a lot of the ESIMD functionalities, and some tests were even useless after the change. #19411 --and it's follow up #19453, is now reverted, and we're taking a new approach: we still run all passes with `-O0`, but we add a new `-force-disable-opt` option to `sycl-post-link` so we can still test for ESIMD functionalities and keep the tests valid.
1 parent 29dee43 commit 1aed422

15 files changed

+490
-207
lines changed

‎llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp

Lines changed: 20 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -30,25 +30,30 @@ using namespace llvm::module_split;
3030

3131
namespace {
3232

33-
ModulePassManager buildESIMDLoweringPipeline(bool OptLevelO0, bool SplitESIMD) {
33+
ModulePassManager buildESIMDLoweringPipeline(bool ForceDisableESIMDOpt,
34+
bool SplitESIMD) {
3435
ModulePassManager MPM;
3536
MPM.addPass(SYCLLowerESIMDPass(!SplitESIMD));
3637

37-
FunctionPassManager FPM;
38-
FPM.addPass(SROAPass(SROAOptions::ModifyCFG));
39-
MPM.addPass(createModuleToFunctionPassAdaptor(std::move(FPM)));
38+
if (!ForceDisableESIMDOpt) {
39+
FunctionPassManager FPM;
40+
FPM.addPass(SROAPass(SROAOptions::ModifyCFG));
41+
MPM.addPass(createModuleToFunctionPassAdaptor(std::move(FPM)));
42+
}
4043
MPM.addPass(ESIMDOptimizeVecArgCallConvPass{});
4144
FunctionPassManager MainFPM;
4245
MainFPM.addPass(ESIMDLowerLoadStorePass{});
4346

44-
MainFPM.addPass(SROAPass(SROAOptions::ModifyCFG));
45-
MainFPM.addPass(EarlyCSEPass(true));
46-
MainFPM.addPass(InstCombinePass{});
47-
MainFPM.addPass(DCEPass{});
48-
MainFPM.addPass(SROAPass(SROAOptions::ModifyCFG));
49-
MainFPM.addPass(EarlyCSEPass(true));
50-
MainFPM.addPass(InstCombinePass{});
51-
MainFPM.addPass(DCEPass{});
47+
if (!ForceDisableESIMDOpt) {
48+
MainFPM.addPass(SROAPass(SROAOptions::ModifyCFG));
49+
MainFPM.addPass(EarlyCSEPass(true));
50+
MainFPM.addPass(InstCombinePass{});
51+
MainFPM.addPass(DCEPass{});
52+
MainFPM.addPass(SROAPass(SROAOptions::ModifyCFG));
53+
MainFPM.addPass(EarlyCSEPass(true));
54+
MainFPM.addPass(InstCombinePass{});
55+
MainFPM.addPass(DCEPass{});
56+
}
5257
MPM.addPass(ESIMDLowerSLMReservationCalls{});
5358
MPM.addPass(createModuleToFunctionPassAdaptor(std::move(MainFPM)));
5459
MPM.addPass(GenXSPIRVWriterAdaptor(/*RewriteTypes=*/true,
@@ -60,7 +65,7 @@ ModulePassManager buildESIMDLoweringPipeline(bool OptLevelO0, bool SplitESIMD) {
6065

6166
// When ESIMD code was separated from the regular SYCL code,
6267
// we can safely process ESIMD part.
63-
bool sycl::lowerESIMDConstructs(ModuleDesc &MD, bool OptLevelO0,
68+
bool sycl::lowerESIMDConstructs(ModuleDesc &MD, bool ForceDisableESIMDOpt,
6469
bool SplitESIMD) {
6570
// TODO: support options like -debug-pass, -print-[before|after], and others
6671
LoopAnalysisManager LAM;
@@ -77,7 +82,8 @@ bool sycl::lowerESIMDConstructs(ModuleDesc &MD, bool OptLevelO0,
7782

7883
std::vector<std::string> Names;
7984
MD.saveEntryPointNames(Names);
80-
ModulePassManager MPM = buildESIMDLoweringPipeline(OptLevelO0, SplitESIMD);
85+
ModulePassManager MPM =
86+
buildESIMDLoweringPipeline(ForceDisableESIMDOpt, SplitESIMD);
8187
PreservedAnalyses Res = MPM.run(MD.getModule(), MAM);
8288

8389
// GenXSPIRVWriterAdaptor pass replaced some functions with "rewritten"

‎llvm/test/tools/sycl-post-link/sycl-esimd/basic-esimd-lower.ll

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -14,6 +14,10 @@
1414
; RUN: sycl-post-link -properties -split-esimd -lower-esimd -O2 -S < %s -o %t.table
1515
; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK-O2
1616

17+
; -O0 lowering, requires `-force-disable-esimd-opt` to disable all optimizations.
18+
; RUN: sycl-post-link -properties -split-esimd -lower-esimd -O0 -force-disable-esimd-opt -S < %s -o %t.table
19+
; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK-O0
20+
1721
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
1822
target triple = "spir64-unknown-linux"
1923

@@ -50,6 +54,15 @@ attributes #0 = { "sycl-module-id"="a.cpp" }
5054
; CHECK-NO-LOWERING: ret void
5155
; CHECK-NO-LOWERING: }
5256

57+
; With -O0, we only lower ESIMD code, but no other optimizations
58+
; CHECK-O0: define dso_local spir_kernel void @ESIMD_kernel() #{{[0-9]}} !sycl_explicit_simd !{{[0-9]}} !intel_reqd_sub_group_size !{{[0-9]}} {
59+
; CHECK-O0: entry:
60+
; CHECK-O0: %0 = load <3 x i64>, {{.*}} addrspacecast {{.*}} @__spirv_BuiltInGlobalInvocationId
61+
; CHECK-O0: %1 = extractelement <3 x i64> %0, i64 0
62+
; CHECK-O0: call void @llvm.genx.barrier()
63+
; CHECK-O0: ret void
64+
; CHECK-O0: }
65+
5366
; With -O2, unused call was optimized away
5467
; CHECK-O2: define dso_local spir_kernel void @ESIMD_kernel()
5568
; CHECK-O2: entry:

‎llvm/test/tools/sycl-post-link/sycl-esimd/sycl-post-link-test.ll

Lines changed: 14 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,5 @@
1-
; RUN: sycl-post-link -properties -split-esimd -lower-esimd -O0 -S < %s -o %t.table
1+
; -O0 lowering, requires `-force-disable-esimd-opt` to disable all optimizations.
2+
; RUN: sycl-post-link -properties -split-esimd -lower-esimd -O0 -force-disable-esimd-opt -S < %s -o %t.table
23
; RUN: FileCheck %s -input-file=%t_esimd_0.ll
34
; This test checks that IR code below can be successfully processed by
45
; sycl-post-link. In this IR no extractelement instruction and no casting are used
@@ -20,8 +21,10 @@ entry:
2021
store i32 %add.i, ptr addrspace(1) %_arg_DoNotOptimize32, align 4
2122
ret void
2223
}
23-
; CHECK: store i64 0, ptr addrspace(1) %_arg_DoNotOptimize, align 8
24-
; CHECK: store i32 3, ptr addrspace(1) %_arg_DoNotOptimize32, align 4
24+
; CHECK: %conv.i = zext i32 0 to i64
25+
; CHECK: store i64 %conv.i, ptr addrspace(1) %_arg_DoNotOptimize, align 8
26+
; CHECK: %add.i = add i32 0, 3
27+
; CHECK: store i32 %add.i, ptr addrspace(1) %_arg_DoNotOptimize32, align 4
2528

2629
; Function Attrs: convergent norecurse
2730
define dso_local spir_kernel void @kernel_SubgroupSize(ptr addrspace(1) noundef align 8 %_arg_DoNotOptimize, ptr addrspace(1) noundef align 4 %_arg_DoNotOptimize32)#0 !sycl_explicit_simd !3{
@@ -33,8 +36,10 @@ entry:
3336
store i32 %add.i, ptr addrspace(1) %_arg_DoNotOptimize32, align 4
3437
ret void
3538
}
36-
; CHECK: store i64 1, ptr addrspace(1) %_arg_DoNotOptimize, align 8
37-
; CHECK: store i32 8, ptr addrspace(1) %_arg_DoNotOptimize32, align 4
39+
; CHECK: %conv.i = zext i32 1 to i64
40+
; CHECK: store i64 %conv.i, ptr addrspace(1) %_arg_DoNotOptimize, align 8
41+
; CHECK: %add.i = add i32 1, 7
42+
; CHECK: store i32 %add.i, ptr addrspace(1) %_arg_DoNotOptimize32, align 4
3843

3944
; Function Attrs: convergent norecurse
4045
define dso_local spir_kernel void @kernel_SubgroupMaxSize(ptr addrspace(1) noundef align 8 %_arg_DoNotOptimize, ptr addrspace(1) noundef align 4 %_arg_DoNotOptimize32) #0 !sycl_explicit_simd !3 {
@@ -46,8 +51,10 @@ entry:
4651
store i32 %add.i, ptr addrspace(1) %_arg_DoNotOptimize32, align 4
4752
ret void
4853
}
49-
; CHECK: store i64 1, ptr addrspace(1) %_arg_DoNotOptimize, align 8
50-
; CHECK: store i32 10, ptr addrspace(1) %_arg_DoNotOptimize32, align 4
54+
; CHECK: %conv.i = zext i32 1 to i64
55+
; CHECK: store i64 %conv.i, ptr addrspace(1) %_arg_DoNotOptimize, align 8
56+
; CHECK: %add.i = add i32 1, 9
57+
; CHECK: store i32 %add.i, ptr addrspace(1) %_arg_DoNotOptimize32, align 4
5158

5259
attributes #0 = { "sycl-module-id"="a.cpp" }
5360

‎llvm/tools/sycl-post-link/sycl-post-link.cpp

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -177,6 +177,10 @@ cl::opt<bool> OptLevelO3("O3",
177177
cl::desc("Optimization level 3. Similar to clang -O3"),
178178
cl::cat(PostLinkCat));
179179

180+
cl::opt<bool> ForceDisableESIMDOpt("force-disable-esimd-opt", cl::Hidden,
181+
cl::desc("Force no optimizations."),
182+
cl::cat(PostLinkCat));
183+
180184
cl::opt<module_split::IRSplitMode> SplitMode(
181185
"split", cl::desc("split input module"), cl::Optional,
182186
cl::init(module_split::SPLIT_NONE),
@@ -523,7 +527,8 @@ handleESIMD(module_split::ModuleDesc &&MDesc, bool &Modified,
523527
for (auto &MD : Result) {
524528
DUMP_ENTRY_POINTS(MD.entries(), MD.Name.c_str(), 3);
525529
if (LowerEsimd && MD.isESIMD())
526-
Modified |= sycl::lowerESIMDConstructs(MD, OptLevelO0, SplitEsimd);
530+
Modified |=
531+
sycl::lowerESIMDConstructs(MD, ForceDisableESIMDOpt, SplitEsimd);
527532
}
528533

529534
if (!SplitEsimd && Result.size() > 1) {

‎sycl/test/check_device_code/esimd/fp16_converts.cpp

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -5,7 +5,9 @@
55

66
// Checks that lowerESIMD pass builds proper vc-intrinsics
77
// RUN: %clangxx -O2 -fsycl -c -fsycl-device-only -Xclang -emit-llvm %s -o %t
8-
// RUN: sycl-post-link -properties -split-esimd -lower-esimd -O0 -S %t -o %t.table
8+
// -O0 lowering, requires `-force-disable-esimd-opt` to disable all
9+
// optimizations.
10+
// RUN: sycl-post-link -properties -split-esimd -lower-esimd -O0 -force-disable-esimd-opt -S %t -o %t.table
911
// RUN: FileCheck %s -input-file=%t_esimd_0.ll
1012

1113
#include <sycl/ext/intel/esimd.hpp>
@@ -34,7 +36,9 @@ __attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
3436
SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void bf16_vector() {
3537
simd<float, 8> F32 = 0;
3638
simd<bfloat16, 8> BF16 = F32;
39+
// CHECK: call <8 x half> @llvm.genx.bf.cvt.v8f16.v8f32(<8 x float> {{[^)]+}})
3740
simd<float, 8> F32_conv = BF16;
41+
// CHECK: call <8 x float> @llvm.genx.bf.cvt.v8f32.v8f16(<8 x half> {{[^)]+}})
3842
}
3943

4044
SYCL_ESIMD_FUNCTION SYCL_EXTERNAL void bf16_scalar() {

0 commit comments

Comments
 (0)