[OpenMP] Add support for the 'private pointer' flag to signal variables captured in target regions and used in first-private clauses.
Summary: If a variable is implicitly mapped (doesn't show in a map clause), the runtime library has to be informed if the corresponding capture shows up in first-private clause, so that the storage previously allocated in the device is used. This patch adds the support for that.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: caomhin, cfe-commits
Differential Revision: http://reviews.llvm.org/D20112
llvm-svn: 270870
diff --git a/clang/test/OpenMP/target_map_codegen.cpp b/clang/test/OpenMP/target_map_codegen.cpp
index 2f6cf3c..61f5ece 100644
--- a/clang/test/OpenMP/target_map_codegen.cpp
+++ b/clang/test/OpenMP/target_map_codegen.cpp
@@ -4281,8 +4281,17 @@
// CK27: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] zeroinitializer
// CK27: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
-// CK27-LABEL: zero_size_section_maps
-void zero_size_section_maps (int ii){
+// CK27: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] zeroinitializer
+// CK27: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
+
+// CK27: [[SIZE07:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK27: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] [i32 288]
+
+// CK27: [[SIZE09:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 40]
+// CK27: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 161]
+
+// CK27-LABEL: zero_size_section_and_private_maps
+void zero_size_section_and_private_maps (int ii){
// Map of a pointer.
int *pa;
@@ -4367,12 +4376,99 @@
{
pa[50]++;
}
+
+ int *pvtPtr;
+ int pvtScl;
+ int pvtArr[10];
+
+ // Region 04
+ // CK27: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 0, i8** null, i8** null, i{{64|32}}* null, i32* null)
+ // CK27: call void [[CALL04:@.+]]()
+ #pragma omp target private(pvtPtr)
+ {
+ pvtPtr[5]++;
+ }
+
+ // Region 05
+ // CK27-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}})
+ // CK27-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK27-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK27-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+ // CK27-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+ // CK27-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8*
+ // CK27-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8*
+
+ // CK27: call void [[CALL05:@.+]](i32* {{[^,]+}})
+ #pragma omp target firstprivate(pvtPtr)
+ {
+ pvtPtr[5]++;
+ }
+
+ // Region 06
+ // CK27: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 0, i8** null, i8** null, i{{64|32}}* null, i32* null)
+ // CK27: call void [[CALL06:@.+]]()
+ #pragma omp target private(pvtScl)
+ {
+ pvtScl++;
+ }
+
+ // Region 07
+ // CK27-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZE07]]{{.+}}, {{.+}}[[MTYPE07]]{{.+}})
+ // CK27-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
+ // CK27-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
+ // CK27-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
+ // CK27-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
+ // CK27-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
+ // CK27-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
+ // CK27-DAG: [[VALBP]] = inttoptr i[[Z]] [[VAL:%.+]] to i8*
+ // CK27-DAG: [[VALP]] = inttoptr i[[Z]] [[VAL:%.+]] to i8*
+ // CK27-DAG: [[VAL]] = load i[[Z]], i[[Z]]* [[ADDR:%.+]],
+ // CK27-64-DAG: [[CADDR:%.+]] = bitcast i[[Z]]* [[ADDR]] to i32*
+ // CK27-64-DAG: store i32 {{.+}}, i32* [[CADDR]],
+
+ // CK27: call void [[CALL07:@.+]](i[[Z]] [[VAL]])
+ #pragma omp target firstprivate(pvtScl)
+ {
+ pvtScl++;
+ }
+
+ // Region 08
+ // CK27: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 0, i8** null, i8** null, i{{64|32}}* null, i32* null)
+ // CK27: call void [[CALL08:@.+]]()
+ #pragma omp target private(pvtArr)
+ {
+ pvtArr[5]++;
+ }
+
+ // Region 09
+ // CK27-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}})
+ // CK27-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK27-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK27-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+ // CK27-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+ // CK27-DAG: [[CBPVAL0]] = bitcast [10 x i32]* [[VAR0:%.+]] to i8*
+ // CK27-DAG: [[CPVAL0]] = bitcast [10 x i32]* [[VAR0]] to i8*
+
+ // CK27: call void [[CALL09:@.+]]([10 x i32]* {{[^,]+}})
+ #pragma omp target firstprivate(pvtArr)
+ {
+ pvtArr[5]++;
+ }
}
// CK27: define {{.+}}[[CALL00]]
// CK27: define {{.+}}[[CALL01]]
// CK27: define {{.+}}[[CALL02]]
// CK27: define {{.+}}[[CALL03]]
-
+// CK27: define {{.+}}[[CALL04]]
+// CK27: define {{.+}}[[CALL05]]
+// CK27: define {{.+}}[[CALL06]]
+// CK27: define {{.+}}[[CALL07]]
#endif
#endif