Skip to content

Commit 3f81007

Browse files
illuhadilluhad
committed
[SSCP] Fix inline assembly by checking CallBase::getCalledFunction() and add test case
1 parent 29fe4c1 commit 3f81007

File tree

7 files changed

+97
-32
lines changed

7 files changed

+97
-32
lines changed

‎src/compiler/llvm-to-backend/AddressSpaceInferencePass.cpp‎

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -132,8 +132,9 @@ llvm::PreservedAnalyses AddressSpaceInferencePass::run(llvm::Module &M,
132132
// so we cannot just make them use ASCastInst instead of AI now.
133133
forEachUseOfPointerValue(AI, [&](llvm::Value* U){
134134
if(auto* CB = llvm::dyn_cast<llvm::CallBase>(U)) {
135-
llvm::StringRef CalleeName = CB->getCalledFunction()->getName();
136-
if(llvmutils::starts_with(CalleeName,"llvm.lifetime")) {
135+
if (CB->getCalledFunction() &&
136+
llvmutils::starts_with(CB->getCalledFunction()->getName(), "llvm.lifetime")) {
137+
llvm::StringRef CalleeName = CB->getCalledFunction()->getName();
137138
InstsToRemove.push_back(CB);
138139

139140
llvm::Intrinsic::ID Id = llvmutils::starts_with(CalleeName, "llvm.lifetime.start")

‎src/compiler/llvm-to-backend/LLVMToBackend.cpp‎

Lines changed: 8 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -182,12 +182,14 @@ class InstructionCleanupPass : public llvm::PassInfoMixin<InstructionCleanupPass
182182
for(auto& BB : F) {
183183
for(auto& I : BB) {
184184
if(llvm::CallBase* CB = llvm::dyn_cast<llvm::CallBase>(&I)) {
185-
// these instructions can sometimes appear as a byproduct of some transformations
186-
// even without dynamic allocas, but they are generally unsupported on device
187-
// backends.
188-
if (llvmutils::starts_with(CB->getCalledFunction()->getName(), "llvm.stacksave") ||
189-
llvmutils::starts_with(CB->getCalledFunction()->getName(), "llvm.stackrestore"))
190-
CallsToRemove.push_back(CB);
185+
if(CB->getCalledFunction()) {
186+
// these instructions can sometimes appear as a byproduct of some transformations
187+
// even without dynamic allocas, but they are generally unsupported on device
188+
// backends.
189+
if (llvmutils::starts_with(CB->getCalledFunction()->getName(), "llvm.stacksave") ||
190+
llvmutils::starts_with(CB->getCalledFunction()->getName(), "llvm.stackrestore"))
191+
CallsToRemove.push_back(CB);
192+
}
191193
}
192194
}
193195
}

‎src/compiler/llvm-to-backend/spirv/LLVMToSpirv.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -258,8 +258,8 @@ bool LLVMToSpirvTranslator::toBackendFlavor(llvm::Module &M, PassHandler& PH) {
258258
// llvm-spirv translator does not like llvm.lifetime.start/end operate on generic
259259
// pointers.
260260
auto* CalledF = CB->getCalledFunction();
261-
if (llvmutils::starts_with(CalledF->getName(), "llvm.lifetime.start") ||
262-
llvmutils::starts_with(CalledF->getName(), "llvm.lifetime.end")) {
261+
if (CalledF && (llvmutils::starts_with(CalledF->getName(), "llvm.lifetime.start") ||
262+
llvmutils::starts_with(CalledF->getName(), "llvm.lifetime.end"))) {
263263
if(CB->getNumOperands() > 1 && CB->getArgOperand(1)->getType()->isPointerTy())
264264
if (CB->getArgOperand(1)->getType()->getPointerAddressSpace() ==
265265
ASMap[AddressSpace::Generic])

‎src/compiler/reflection/FunctionNameExtractionPass.cpp‎

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -47,9 +47,10 @@ bool isReflectionAnnotatedFunction(llvm::Function* F, const utils::ProcessFuncti
4747
bool isAnyUserReflectionAnnotatedFunction(llvm::Function* F, const utils::ProcessFunctionAnnotationPass& PFA) {
4848
for(auto* U : F->users()) {
4949
if(auto* CB = llvm::dyn_cast<llvm::CallBase>(U)) {
50-
if(CB->getCalledFunction() != F) {
50+
auto* CalledF = CB->getCalledFunction();
51+
if(CalledF && (CalledF != F)) {
5152

52-
if(isReflectionAnnotatedFunction(CB->getCalledFunction(), PFA))
53+
if(isReflectionAnnotatedFunction(CalledF, PFA))
5354
return true;
5455

5556
}

‎src/compiler/sscp/KernelOutliningPass.cpp‎

Lines changed: 17 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -161,21 +161,23 @@ class FunctionArgPointeeTypeInfererence {
161161
// Follow address space casts, we don't care about pointer address spaces
162162
rankUsers(Current, Scores, CurrentScore);
163163
} else if(auto CI = llvm::dyn_cast<llvm::CallBase>(Current)) {
164-
// Ugh, the value is forwarded as an argument into some other function, need
165-
// to continue looking there...
166-
167-
// First, check if we have any interesting allocas in the called function
168-
scanAllocas(CI->getCalledFunction(), Scores);
169-
170-
// Next, follow the argument that was passed in there
171-
for (int i = 0; i < CI->getCalledFunction()->getFunctionType()->getNumParams(); ++i) {
172-
if(CI->getArgOperand(i) == Parent) {
173-
auto Arg = CI->getCalledFunction()->getArg(i);
174-
// Never, ever take into account the callee argument. This should never happen,
175-
// but if it does, it will go terribly because we will take into account users of functions,
176-
// not arguments anymore.
177-
if(!llvm::isa<llvm::Function>(Arg))
178-
rankUsers(Arg, Scores, CurrentScore);
164+
if(CI->getCalledFunction()) {
165+
// Ugh, the value is forwarded as an argument into some other function, need
166+
// to continue looking there...
167+
168+
// First, check if we have any interesting allocas in the called function
169+
scanAllocas(CI->getCalledFunction(), Scores);
170+
171+
// Next, follow the argument that was passed in there
172+
for (int i = 0; i < CI->getCalledFunction()->getFunctionType()->getNumParams(); ++i) {
173+
if (CI->getArgOperand(i) == Parent) {
174+
auto Arg = CI->getCalledFunction()->getArg(i);
175+
// Never, ever take into account the callee argument. This should never happen,
176+
// but if it does, it will go terribly because we will take into account users of
177+
// functions, not arguments anymore.
178+
if (!llvm::isa<llvm::Function>(Arg))
179+
rankUsers(Arg, Scores, CurrentScore);
180+
}
179181
}
180182
}
181183
}

‎src/compiler/stdpar/SyncElision.cpp‎

Lines changed: 8 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -79,11 +79,14 @@ void identifyStoresPotentiallyForStdparArgHandling(
7979
return true;
8080
}
8181
} else if (auto *CB = llvm::dyn_cast<llvm::CallBase>(Current)) {
82-
if (StdparFunctions.contains(CB->getCalledFunction())) {
83-
Users.push_back(Current);
84-
return true;
85-
} else if(llvmutils::starts_with(CB->getCalledFunction()->getName(), "llvm.lifetime")) {
86-
return true;
82+
auto* Callee = CB->getCalledFunction();
83+
if(Callee) {
84+
if (StdparFunctions.contains(Callee)) {
85+
Users.push_back(Current);
86+
return true;
87+
} else if (llvmutils::starts_with(Callee->getName(), "llvm.lifetime")) {
88+
return true;
89+
}
8790
}
8891
}
8992

‎tests/compiler/sscp/inline_assembly.cpp‎

Lines changed: 56 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,56 @@
1+
// RUN: %acpp %s -o %t --acpp-targets=generic
2+
// RUN: %t | FileCheck %s
3+
// RUN: %acpp %s -o %t --acpp-targets=generic -O3
4+
// RUN: %t | FileCheck %s
5+
// RUN: %acpp %s -o %t --acpp-targets=generic -g
6+
// RUN: %t | FileCheck %s
7+
8+
#include <iostream>
9+
10+
#include <sycl/sycl.hpp>
11+
#include "common.hpp"
12+
13+
unsigned run_cpuid() {
14+
#ifdef __x86_64__
15+
unsigned a = 0x1, b, c, d;
16+
asm volatile("cpuid"
17+
: "=a"(a), "=b"(b), "=c"(c), "=d"(d)
18+
: "a"(a), "b"(b), "c"(c), "d"(d));
19+
return a;
20+
#else
21+
return 0;
22+
#endif
23+
}
24+
25+
int main()
26+
{
27+
sycl::queue q = get_queue();
28+
29+
unsigned* data = sycl::malloc_shared<unsigned>(4, q);
30+
for(int i = 0; i < 4; ++i)
31+
data[i] = 0;
32+
33+
q.parallel_for(sycl::range{1024}, [=](auto idx) {
34+
__acpp_if_target_sscp(
35+
sycl::AdaptiveCpp_jit::compile_if(
36+
__acpp_sscp_jit_reflect_compiler_backend() ==
37+
sycl::AdaptiveCpp_jit::compiler_backend::host,
38+
[&]() { *data = run_cpuid(); });
39+
40+
);
41+
}).wait();
42+
43+
// CHECK: 1
44+
#ifdef __x86_64__
45+
if(q.get_device().get_backend() != sycl::backend::omp) {
46+
std::cout << 1 << std::endl;
47+
} else {
48+
unsigned result = *data;
49+
std::cout << (result > 1) << std::endl;
50+
}
51+
#else
52+
std::cout << 1 << std::endl;
53+
#endif
54+
55+
sycl::free(data,q);
56+
}

0 commit comments

Comments
 (0)