diff --git a/llvm/test/tools/sycl-post-link/sc_sym_two_refs.ll b/llvm/test/tools/sycl-post-link/sc_sym_two_refs.ll new file mode 100644 index 000000000000..5eadf6d40bc2 --- /dev/null +++ b/llvm/test/tools/sycl-post-link/sc_sym_two_refs.ll @@ -0,0 +1,27 @@ +; This test checks that the tool does not crash and removes the unused spec +; constant global symbol when it is referenced more than once. + +; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \ +; RUN: | FileCheck %s + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" +target triple = "spir64-unknown-unknown-sycldevice" + +%"sycl::experimental::spec_constant" = type { i8 } + +@SCSymID = private unnamed_addr constant [10 x i8] c"SpecConst\00", align 1 +; CHECK-NOT: @SCSymID + +declare dso_local spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 addrspace(4)*) + +; Function Attrs: norecurse +define weak_odr dso_local spir_kernel void @Kernel() { + %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)*)) + ret void +} + +; Function Attrs: norecurse +define dso_local spir_func float @foo_float(%"sycl::experimental::spec_constant" addrspace(4)* nocapture readnone dereferenceable(1) %0) local_unnamed_addr #3 { + %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)*)) + ret float %2 +} diff --git a/llvm/test/tools/sycl-post-link/spec_const_O0.ll b/llvm/test/tools/sycl-post-link/spec_const_O0.ll index 363b7170e8c0..65a63cf709c9 100644 --- a/llvm/test/tools/sycl-post-link/spec_const_O0.ll +++ b/llvm/test/tools/sycl-post-link/spec_const_O0.ll @@ -26,7 +26,7 @@ declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1 declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1 ; Function Attrs: norecurse -define linkonce_odr dso_local spir_func zeroext i1 @FOO(%"UserSpecConstIDType" addrspace(4)* %0) comdat align 2 { +define spir_func zeroext i1 @FOO(%"UserSpecConstIDType" addrspace(4)* %0) comdat align 2 { %2 = alloca %"UserSpecConstIDType" addrspace(4)*, align 8 %3 = alloca i8 addrspace(4)*, align 8 store %"UserSpecConstIDType" addrspace(4)* %0, %"UserSpecConstIDType" addrspace(4)** %2, align 8, !tbaa !8 diff --git a/llvm/test/tools/sycl-post-link/spec_const_and_split.ll b/llvm/test/tools/sycl-post-link/spec_const_and_split.ll new file mode 100644 index 000000000000..2df5ad205a51 --- /dev/null +++ b/llvm/test/tools/sycl-post-link/spec_const_and_split.ll @@ -0,0 +1,26 @@ +; This test checks that the post-link tool works correctly when both +; device code splitting and specialization constant processing are +; requested. +; +; RUN: sycl-post-link -split=kernel -spec-const=rt -S %s -o %t.files.table +; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefixes CHECK0,CHECK +; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefixes CHECK1,CHECK + +@SCSymID = private unnamed_addr constant [10 x i8] c"SpecConst\00", align 1 +; CHECK-NOT: @SCSymID + +declare dso_local spir_func zeroext i1 @_Z27__sycl_getSpecConstantValueIbET_PKc(i8 addrspace(4)*) + +define dso_local spir_kernel void @KERNEL_AAA() { + %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)*)) +; CHECK0: %{{[0-9]+}} = call i1 @_Z20__spirv_SpecConstantib(i32 0, i1 false), !SYCL_SPEC_CONST_SYM_ID ![[MD_ID:[0-9]+]] + ret void +} + +define dso_local spir_kernel void @KERNEL_BBB() { + %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)*)) +; CHECK1: %{{[0-9]+}} = call i1 @_Z20__spirv_SpecConstantib(i32 0, i1 false), !SYCL_SPEC_CONST_SYM_ID ![[MD_ID:[0-9]+]] + ret void +} + +; CHECK: ![[MD_ID]] = !{!"SpecConst", i32 0} diff --git a/llvm/tools/sycl-post-link/SpecConstants.cpp b/llvm/tools/sycl-post-link/SpecConstants.cpp index 02218cab5de2..3cdead43da15 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.cpp +++ b/llvm/tools/sycl-post-link/SpecConstants.cpp @@ -40,19 +40,18 @@ static void AssertRelease(bool Cond, const char *Msg) { StringRef getStringLiteralArg(const CallInst *CI, unsigned ArgNo, SmallVectorImpl &DelInsts, - GlobalVariable *&DelGlob) { + GlobalVariable *&SymGlob) { Value *V = CI->getArgOperand(ArgNo)->stripPointerCasts(); if (auto *L = dyn_cast(V)) { // Must be a // vvvvvvvvvvvvvvvvvvvv - // @.str = private unnamed_addr constant[18 x i8] - // c"_ZTS11MyBoolConst\00", align 1 + // @.str = private unnamed_addr constant[10 x i8] c"SpecConst\00", align 1 // ... // %TName = alloca i8 addrspace(4)*, align 8 // ... // store i8 addrspace(4)* addrspacecast( - // i8* getelementptr inbounds([18 x i8], [18 x i8] * @.str, i32 0, i32 0) + // i8* getelementptr inbounds([10 x i8], [10 x i8] * @.str, i32 0, i32 0) // to i8 addrspace(4)*), i8 addrspace(4)** %TName, align 8, !tbaa !10 // %1 = load i8 addrspace(4)*, i8 addrspace(4)** %TName, align 8, !tbaa !10 // %call = call spir_func zeroext @@ -96,7 +95,7 @@ StringRef getStringLiteralArg(const CallInst *CI, unsigned ArgNo, V = Store->getValueOperand()->stripPointerCasts(); } const Constant *Init = cast(V)->getInitializer(); - DelGlob = cast(V); + SymGlob = cast(V); StringRef Res = cast(Init)->getAsString(); if (Res.size() > 0 && Res[Res.size() - 1] == '\0') Res = Res.substr(0, Res.size() - 1); @@ -214,8 +213,8 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, // code can't use this intrinsic directly. SmallVector DelInsts; DelInsts.push_back(CI); - GlobalVariable *DelGlob = nullptr; - StringRef SymID = getStringLiteralArg(CI, 0, DelInsts, DelGlob); + GlobalVariable *SymGlob = nullptr; + StringRef SymID = getStringLiteralArg(CI, 0, DelInsts, SymGlob); Type *SCTy = CI->getType(); if (SetValAtRT) { @@ -262,9 +261,8 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, I->removeFromParent(); I->deleteValue(); } - DelGlob->replaceAllUsesWith(ConstantPointerNull::get(DelGlob->getType())); - DelGlob->removeFromParent(); - DelGlob->deleteValue(); + // Don't delete SymGlob here, as it may be referenced from multiple + // functions if __sycl_getSpecConstantValue is inlined. } } return IRModified ? PreservedAnalyses::none() : PreservedAnalyses::all(); diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 2d1d01ce7801..66c842ef285e 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -31,6 +31,7 @@ #include "llvm/Support/SystemUtils.h" #include "llvm/Support/WithColor.h" #include "llvm/Transforms/IPO.h" +#include "llvm/Transforms/IPO/GlobalDCE.h" #include "llvm/Transforms/Utils/Cloning.h" #include @@ -427,6 +428,9 @@ int main(int argc, char **argv) { // Register required analysis MAM.registerPass([&] { return PassInstrumentationAnalysis(); }); RunSpecConst.addPass(SCP); + if (!DoSplit) + // This pass deletes unreachable globals. Code splitter runs it later. + RunSpecConst.addPass(GlobalDCEPass()); PreservedAnalyses Res = RunSpecConst.run(*MPtr, MAM); SpecConstsMet = !Res.areAllPreserved(); }