dnrops
/
clang
mirror of https://github.com/microsoft/clang.git
1
0
Fork 0
Code Issues Packages Projects 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.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@310920 91177308-0d34-0410-b5e6-96231b3b80d8
This commit is contained in:
Alexey Bataev 2017-08-15 14:34:04 +00:00
parent 8872ad1bd2
commit bcc9bb7226
2 changed files with 24 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
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
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;
}