[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

2 files changed, 24 insertions, 0 deletions

diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
index fca0a50ca16..3202121e8c8 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
@@ -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()) {
diff --git a/clang/test/OpenMP/nvptx_param_translate.c b/clang/test/OpenMP/nvptx_param_translate.c
new file mode 100644
index 00000000000..ec123ce3811
--- /dev/null
+++ b/clang/test/OpenMP/nvptx_param_translate.c
@@ -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;
+}