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