maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[OPENMP] Fix compiler crash on argument translate for NVPTX. 

When translating arguments for NVPTX target it is not taken into account
that function may have variable number of arguments. Patch fixes this
problem.

llvm-svn: 310920
This commit is contained in:
Alexey Bataev 2017-08-15 14:34:04 +00:00
parent 16e6dd3cd6
commit 07ed94a7c7
2 changed files with 24 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
View File

@ -2295,9 +2295,14 @@ void CGOpenMPRuntimeNVPTX::emitOutlinedFunctionCall(
CodeGenFunction &CGF, SourceLocation Loc, llvm::Value *OutlinedFn,
ArrayRef<llvm::Value *> Args) const {
SmallVector<llvm::Value *, 4> TargetArgs;
TargetArgs.reserve(Args.size());
auto *FnType =
cast<llvm::FunctionType>(OutlinedFn->getType()->getPointerElementType());
for (unsigned I = 0, E = Args.size(); I < E; ++I) {
if (FnType->isVarArg() && FnType->getNumParams() <= I) {
TargetArgs.append(std::next(Args.begin(), I), Args.end());
break;
}
llvm::Type *TargetType = FnType->getParamType(I);
llvm::Value *NativeArg = Args[I];
if (!TargetType->isPointerTy()) {

19
clang/test/OpenMP/nvptx_param_translate.c Normal file
View File

@ -0,0 +1,19 @@
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
// expected-no-diagnostics
// CHECK: [[MAP_FN:%.+]] = load void (i8*, ...)*, void (i8*, ...)** %
// CHECK: call void (i8*, ...) [[MAP_FN]](i8* %
int main() {
double a, b;
#pragma omp target map(tofrom \
: a) map(to \
: b)
{
#pragma omp taskgroup
#pragma omp task shared(a)
a = b;
}
return 0;
}