diff options
author | Samuel Antao <sfantao@us.ibm.com> | 2016-05-26 16:53:38 +0000 |
---|---|---|
committer | Samuel Antao <sfantao@us.ibm.com> | 2016-05-26 16:53:38 +0000 |
commit | 666dfeba374e6b79f68a5ae1007c92aea184a0f5 (patch) | |
tree | 8c713106e6d2b70e0b643faeeae1106cf8d62851 /test/OpenMP/target_map_codegen.cpp | |
parent | 42720cb0acda06dd20caf809a0c4d931ad249dcf (diff) |
[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
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@270870 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/OpenMP/target_map_codegen.cpp')
-rw-r--r-- | test/OpenMP/target_map_codegen.cpp | 102 |
1 files changed, 99 insertions, 3 deletions
diff --git a/test/OpenMP/target_map_codegen.cpp b/test/OpenMP/target_map_codegen.cpp index 2f6cf3ccea..61f5ecea75 100644 --- a/test/OpenMP/target_map_codegen.cpp +++ b/test/OpenMP/target_map_codegen.cpp @@ -4281,8 +4281,17 @@ int explicit_maps_with_private_class_members(){ // 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 @@ void zero_size_section_maps (int ii){ { 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 |