Skip to content

Commit ce020c9

Browse files
kbobrovskbobrovs
authored
[SYCL] Fix processing of spec consts referenced twice (#1524)
Signed-off-by: Konstantin S Bobrovsky <[email protected]>
1 parent e4b530b commit ce020c9

File tree

5 files changed

+66
-11
lines changed

5 files changed

+66
-11
lines changed

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

Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,27 @@
1+
; This test checks that the tool does not crash and removes the unused spec
2+
; constant global symbol when it is referenced more than once.
3+
4+
; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \
5+
; RUN: | FileCheck %s
6+
7+
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
8+
target triple = "spir64-unknown-unknown-sycldevice"
9+
10+
%"sycl::experimental::spec_constant" = type { i8 }
11+
12+
@SCSymID = private unnamed_addr constant [10 x i8] c"SpecConst\00", align 1
13+
; CHECK-NOT: @SCSymID
14+
15+
declare dso_local spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 addrspace(4)*)
16+
17+
; Function Attrs: norecurse
18+
define weak_odr dso_local spir_kernel void @Kernel() {
19+
%1 = call spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([10 x i8], [10 x i8]* @SCSymID, i64 0, i64 0) to i8 addrspace(4)*))
20+
ret void
21+
}
22+
23+
; Function Attrs: norecurse
24+
define dso_local spir_func float @foo_float(%"sycl::experimental::spec_constant" addrspace(4)* nocapture readnone dereferenceable(1) %0) local_unnamed_addr #3 {
25+
%2 = tail call spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([10 x i8], [10 x i8]* @SCSymID, i64 0, i64 0) to i8 addrspace(4)*))
26+
ret float %2
27+
}

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,7 @@ declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1
2626
declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1
2727

2828
; Function Attrs: norecurse
29-
define linkonce_odr dso_local spir_func zeroext i1 @FOO(%"UserSpecConstIDType" addrspace(4)* %0) comdat align 2 {
29+
define spir_func zeroext i1 @FOO(%"UserSpecConstIDType" addrspace(4)* %0) comdat align 2 {
3030
%2 = alloca %"UserSpecConstIDType" addrspace(4)*, align 8
3131
%3 = alloca i8 addrspace(4)*, align 8
3232
store %"UserSpecConstIDType" addrspace(4)* %0, %"UserSpecConstIDType" addrspace(4)** %2, align 8, !tbaa !8

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

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,26 @@
1+
; This test checks that the post-link tool works correctly when both
2+
; device code splitting and specialization constant processing are
3+
; requested.
4+
;
5+
; RUN: sycl-post-link -split=kernel -spec-const=rt -S %s -o %t.files.table
6+
; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefixes CHECK0,CHECK
7+
; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefixes CHECK1,CHECK
8+
9+
@SCSymID = private unnamed_addr constant [10 x i8] c"SpecConst\00", align 1
10+
; CHECK-NOT: @SCSymID
11+
12+
declare dso_local spir_func zeroext i1 @_Z27__sycl_getSpecConstantValueIbET_PKc(i8 addrspace(4)*)
13+
14+
define dso_local spir_kernel void @KERNEL_AAA() {
15+
%1 = call spir_func zeroext i1 @_Z27__sycl_getSpecConstantValueIbET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([10 x i8], [10 x i8]* @SCSymID, i64 0, i64 0) to i8 addrspace(4)*))
16+
; CHECK0: %{{[0-9]+}} = call i1 @_Z20__spirv_SpecConstantib(i32 0, i1 false), !SYCL_SPEC_CONST_SYM_ID ![[MD_ID:[0-9]+]]
17+
ret void
18+
}
19+
20+
define dso_local spir_kernel void @KERNEL_BBB() {
21+
%1 = call spir_func zeroext i1 @_Z27__sycl_getSpecConstantValueIbET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([10 x i8], [10 x i8]* @SCSymID, i64 0, i64 0) to i8 addrspace(4)*))
22+
; CHECK1: %{{[0-9]+}} = call i1 @_Z20__spirv_SpecConstantib(i32 0, i1 false), !SYCL_SPEC_CONST_SYM_ID ![[MD_ID:[0-9]+]]
23+
ret void
24+
}
25+
26+
; CHECK: ![[MD_ID]] = !{!"SpecConst", i32 0}

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

Lines changed: 8 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -40,19 +40,18 @@ static void AssertRelease(bool Cond, const char *Msg) {
4040

4141
StringRef getStringLiteralArg(const CallInst *CI, unsigned ArgNo,
4242
SmallVectorImpl<Instruction *> &DelInsts,
43-
GlobalVariable *&DelGlob) {
43+
GlobalVariable *&SymGlob) {
4444
Value *V = CI->getArgOperand(ArgNo)->stripPointerCasts();
4545

4646
if (auto *L = dyn_cast<LoadInst>(V)) {
4747
// Must be a
4848
// vvvvvvvvvvvvvvvvvvvv
49-
// @.str = private unnamed_addr constant[18 x i8]
50-
// c"_ZTS11MyBoolConst\00", align 1
49+
// @.str = private unnamed_addr constant[10 x i8] c"SpecConst\00", align 1
5150
// ...
5251
// %TName = alloca i8 addrspace(4)*, align 8
5352
// ...
5453
// store i8 addrspace(4)* addrspacecast(
55-
// i8* getelementptr inbounds([18 x i8], [18 x i8] * @.str, i32 0, i32 0)
54+
// i8* getelementptr inbounds([10 x i8], [10 x i8] * @.str, i32 0, i32 0)
5655
// to i8 addrspace(4)*), i8 addrspace(4)** %TName, align 8, !tbaa !10
5756
// %1 = load i8 addrspace(4)*, i8 addrspace(4)** %TName, align 8, !tbaa !10
5857
// %call = call spir_func zeroext
@@ -96,7 +95,7 @@ StringRef getStringLiteralArg(const CallInst *CI, unsigned ArgNo,
9695
V = Store->getValueOperand()->stripPointerCasts();
9796
}
9897
const Constant *Init = cast<GlobalVariable>(V)->getInitializer();
99-
DelGlob = cast<GlobalVariable>(V);
98+
SymGlob = cast<GlobalVariable>(V);
10099
StringRef Res = cast<ConstantDataArray>(Init)->getAsString();
101100
if (Res.size() > 0 && Res[Res.size() - 1] == '\0')
102101
Res = Res.substr(0, Res.size() - 1);
@@ -214,8 +213,8 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
214213
// code can't use this intrinsic directly.
215214
SmallVector<Instruction *, 3> DelInsts;
216215
DelInsts.push_back(CI);
217-
GlobalVariable *DelGlob = nullptr;
218-
StringRef SymID = getStringLiteralArg(CI, 0, DelInsts, DelGlob);
216+
GlobalVariable *SymGlob = nullptr;
217+
StringRef SymID = getStringLiteralArg(CI, 0, DelInsts, SymGlob);
219218
Type *SCTy = CI->getType();
220219

221220
if (SetValAtRT) {
@@ -262,9 +261,8 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
262261
I->removeFromParent();
263262
I->deleteValue();
264263
}
265-
DelGlob->replaceAllUsesWith(ConstantPointerNull::get(DelGlob->getType()));
266-
DelGlob->removeFromParent();
267-
DelGlob->deleteValue();
264+
// Don't delete SymGlob here, as it may be referenced from multiple
265+
// functions if __sycl_getSpecConstantValue is inlined.
268266
}
269267
}
270268
return IRModified ? PreservedAnalyses::none() : PreservedAnalyses::all();

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

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -31,6 +31,7 @@
3131
#include "llvm/Support/SystemUtils.h"
3232
#include "llvm/Support/WithColor.h"
3333
#include "llvm/Transforms/IPO.h"
34+
#include "llvm/Transforms/IPO/GlobalDCE.h"
3435
#include "llvm/Transforms/Utils/Cloning.h"
3536
#include <memory>
3637

@@ -427,6 +428,9 @@ int main(int argc, char **argv) {
427428
// Register required analysis
428429
MAM.registerPass([&] { return PassInstrumentationAnalysis(); });
429430
RunSpecConst.addPass(SCP);
431+
if (!DoSplit)
432+
// This pass deletes unreachable globals. Code splitter runs it later.
433+
RunSpecConst.addPass(GlobalDCEPass());
430434
PreservedAnalyses Res = RunSpecConst.run(*MPtr, MAM);
431435
SpecConstsMet = !Res.areAllPreserved();
432436
}

0 commit comments

Comments
 (0)