commit f7c4d7b0b161be9c1528095c012504d15a6da045
author Jonas Hahnfeld <Hahnfeld@itc.rwth-aachen.de> Sat Jul 01 10:40:50 2017 +0000
committer Jonas Hahnfeld <Hahnfeld@itc.rwth-aachen.de> Sat Jul 01 10:40:50 2017 +0000
tree 0fc6a05477638000eeef43369f81aec85093eb80
parent 85c529c98807de59540b8fc682a570526b1b2ca0

[OpenMP] Fix mapping of scalars for combined directives

Combined directives like 'target parallel' have two captured statements.
Sema has to check the right one from the right direction.

Previously, Sema::IsOpenMPCapturedByRef would return false for mapped
scalars on combined directives. This results in a wrong signature of
the outlined function which triggers an assertion:
void llvm::CallInst::init(llvm::FunctionType *, llvm::Value *, ArrayRef<llvm::Value *>, ArrayRef<OperandBundleDef>, const llvm::Twine &): Assertion `(i >= FTy->getNumParams() || FTy->getParamType(i) == Args[i]->getType()) && "Calling a function with a bad signature!"' failed.

Fixes PR30975 (and PR31985). New function was taken from clang-ykt.

Differential Revision: https://reviews.llvm.org/D34888

llvm-svn: 306956

clang/test/OpenMP/target_map_codegen.cpp

diff --git a/clang/test/OpenMP/target_map_codegen.cpp b/clang/test/OpenMP/target_map_codegen.cpp
index 4933bf3..69e80bb 100644
--- a/clang/test/OpenMP/target_map_codegen.cpp
+++ b/clang/test/OpenMP/target_map_codegen.cpp
@@ -1056,6 +1056,9 @@
 // CK19: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
 // CK19: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
 
+// CK19: [[SIZE00n:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
+// CK19: [[MTYPE00n:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
+
 // CK19: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
 // CK19: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
 
@@ -1194,6 +1197,28 @@
     ++a;
   }
 
+  // Map of a scalar in nested region.
+  int b = a;
+
+  // Region 00n
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00n]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00n]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK19-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]]
+  // CK19-DAG: store i32* [[VAR0]], i32** [[CP0]]
+
+  // CK19: call void [[CALL00n:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(alloc:b)
+  #pragma omp parallel
+  {
+    ++b;
+  }
+
   // Map of an array.
   int arra[100];
 
@@ -2388,6 +2413,7 @@
 }
 
 // CK19: define {{.+}}[[CALL00]]
+// CK19: define {{.+}}[[CALL00n]]
 // CK19: define {{.+}}[[CALL01]]
 // CK19: define {{.+}}[[CALL02]]
 // CK19: define {{.+}}[[CALL03]]