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;
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
