From: Alexey Bataev <a.bataev@hotmail.com>
Date: Tue, 15 Aug 2017 14:34:04 +0000 (+0000)
Subject: [OPENMP] Fix compiler crash on argument translate for NVPTX.
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=bcc9bb722691c90459b10f7b8016fb84d62d3039;p=clang

[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
---

diff --git a/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp b/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
index fca0a50ca1..3202121e8c 100644
--- a/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
+++ b/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/test/OpenMP/nvptx_param_translate.c b/test/OpenMP/nvptx_param_translate.c
new file mode 100644
index 0000000000..ec123ce381
--- /dev/null
+++ b/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;
+}