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

[Alexey Bataev via cfe-commits]

Author: abataev
Date: Tue Aug 15 07:34:04 2017
New Revision: 310920

URL: http://llvm.org/viewvc/llvm-project?rev=310920&view=rev
Log:
[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.

Added:
    cfe/trunk/test/OpenMP/nvptx_param_translate.c
Modified:
    cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp?rev=310920&r1=310919&r2=310920&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp Tue Aug 15 07:34:04 2017
@@ -2295,9 +2295,14 @@ void CGOpenMPRuntimeNVPTX::emitOutlinedF
     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()) {

Added: cfe/trunk/test/OpenMP/nvptx_param_translate.c
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_param_translate.c?rev=310920&view=auto
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_param_translate.c (added)
+++ cfe/trunk/test/OpenMP/nvptx_param_translate.c Tue Aug 15 07:34:04 2017
@@ -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;
+}