Skip to content

Commit 60d71a2

Browse files
alexey-bataevalexey-bataev
committed
[OPENMP50]Allow overlapping mapping in target constructs.
OpenMP 5.0 removed a lot of restriction for overlapped mapped items comparing to OpenMP 4.5. Patch restricts the checks for overlapped data mappings only for OpenMP 4.5 and less and reorders mapping of the arguments so, that present and alloc mappings are processed first and then all others. Differential Revision: https://reviews.llvm.org/D86119
1 parent 77dc203 commit 60d71a2

22 files changed

+975
-863
lines changed

‎clang/lib/CodeGen/CGOpenMPRuntime.cpp

Lines changed: 322 additions & 295 deletions
Large diffs are not rendered by default.

‎clang/lib/Sema/SemaOpenMP.cpp

Lines changed: 8 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -3564,9 +3564,11 @@ class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
35643564
!Stack->isLoopControlVariable(VD).first) {
35653565
if (!Stack->checkMappableExprComponentListsForDecl(
35663566
VD, /*CurrentRegionOnly=*/true,
3567-
[](OMPClauseMappableExprCommon::MappableExprComponentListRef
3568-
StackComponents,
3569-
OpenMPClauseKind) {
3567+
[this](OMPClauseMappableExprCommon::MappableExprComponentListRef
3568+
StackComponents,
3569+
OpenMPClauseKind) {
3570+
if (SemaRef.LangOpts.OpenMP >= 50)
3571+
return !StackComponents.empty();
35703572
// Variable is used if it has been marked as an array, array
35713573
// section, array shaping or the variable iself.
35723574
return StackComponents.size() == 1 ||
@@ -17579,7 +17581,9 @@ static bool checkMapConflicts(
1757917581
ERange, CKind, &EnclosingExpr,
1758017582
CurComponents](OMPClauseMappableExprCommon::MappableExprComponentListRef
1758117583
StackComponents,
17582-
OpenMPClauseKind) {
17584+
OpenMPClauseKind Kind) {
17585+
if (CKind == Kind && SemaRef.LangOpts.OpenMP >= 50)
17586+
return false;
1758317587
assert(!StackComponents.empty() &&
1758417588
"Map clause expression with no components!");
1758517589
assert(StackComponents.back().getAssociatedDeclaration() == VD &&
@@ -18133,8 +18137,6 @@ static void checkMappableExpressionList(
1813318137
DSAS, Type))
1813418138
continue;
1813518139

18136-
Type = I->getAssociatedDeclaration()->getType().getNonReferenceType();
18137-
1813818140
if (CKind == OMPC_map) {
1813918141
// target enter data
1814018142
// OpenMP [2.10.2, Restrictions, p. 99]

‎clang/test/OpenMP/declare_mapper_codegen.cpp

Lines changed: 133 additions & 156 deletions
Large diffs are not rendered by default.

‎clang/test/OpenMP/target_data_codegen.cpp

Lines changed: 43 additions & 37 deletions
Original file line numberDiff line numberDiff line change
@@ -561,8 +561,9 @@ struct S2 {
561561

562562
void test_close_modifier(int arg) {
563563
S2 *ps;
564-
// CK5: private unnamed_addr constant [5 x i64] [i64 1027, i64 0, i64 562949953421328, i64 16, i64 1043]
565-
#pragma omp target data map(close,tofrom: arg, ps->ps->ps->ps->s)
564+
// CK5: private unnamed_addr constant [5 x i64] [i64 1027, i64 0, i64 562949953421328, i64 16, i64 1043]
565+
#pragma omp target data map(close, tofrom \
566+
: arg, ps->ps->ps->ps->s)
566567
{
567568
++(arg);
568569
}
@@ -585,8 +586,9 @@ void test_close_modifier(int arg) {
585586
// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
586587
#ifdef CK6
587588
void test_close_modifier(int arg) {
588-
// CK6: private unnamed_addr constant [1 x i64] [i64 1027]
589-
#pragma omp target data map(close,tofrom: arg)
589+
// CK6: private unnamed_addr constant [1 x i64] [i64 1027]
590+
#pragma omp target data map(close, tofrom \
591+
: arg)
590592
{++arg;}
591593
}
592594
#endif
@@ -642,36 +644,39 @@ void test_present_modifier(int arg) {
642644

643645
// CK8: private unnamed_addr constant [11 x i64]
644646

645-
// ps1
646-
//
647-
// PRESENT=0x1000 = 0x1000
648-
// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
649-
// MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1000000001010
650-
// PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010
651-
// PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013
652-
//
653-
// CK8-SAME: {{^}} [i64 [[#0x1000]], i64 [[#0x1000000000003]],
654-
// CK8-SAME: {{^}} i64 [[#0x1000000001010]], i64 [[#0x1010]], i64 [[#0x1013]],
655-
656-
// arg
657-
//
658-
// PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003
659-
//
660-
// CK8-SAME: {{^}} i64 [[#0x1003]],
661-
662-
// ps2
663-
//
664-
// PRESENT=0x1000 = 0x1000
665-
// MEMBER_OF_7=0x7000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x7000000001003
666-
// MEMBER_OF_7=0x7000000000000 | PTR_AND_OBJ=0x10 = 0x7000000000010
667-
// PTR_AND_OBJ=0x10 = 0x10
668-
// PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13
669-
//
670-
// CK8-SAME: {{^}} i64 [[#0x1000]], i64 [[#0x7000000001003]],
671-
// CK8-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
672-
#pragma omp target data map(tofrom: ps1->s) \
673-
map(present,tofrom: arg, ps1->ps->ps->ps->s, ps2->s) \
674-
map(tofrom: ps2->ps->ps->ps->s)
647+
// ps1
648+
//
649+
// PRESENT=0x1000 = 0x1000
650+
// MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1000000001010
651+
// PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010
652+
// PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013
653+
// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
654+
//
655+
// CK8-SAME: {{^}} [i64 [[#0x1000]], i64 [[#0x1000000001010]],
656+
// CK8-SAME: {{^}} i64 [[#0x1010]], i64 [[#0x1013]], i64 [[#0x1000000000003]],
657+
658+
// arg
659+
//
660+
// PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003
661+
//
662+
// CK8-SAME: {{^}} i64 [[#0x1003]],
663+
664+
// ps2
665+
//
666+
// PRESENT=0x1000 = 0x1000
667+
// MEMBER_OF_7=0x7000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x7000000001003
668+
// MEMBER_OF_7=0x7000000000000 | PTR_AND_OBJ=0x10 = 0x7000000000010
669+
// PTR_AND_OBJ=0x10 = 0x10
670+
// PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13
671+
//
672+
// CK8-SAME: {{^}} i64 [[#0x1000]], i64 [[#0x7000000001003]],
673+
// CK8-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
674+
#pragma omp target data map(tofrom \
675+
: ps1->s) \
676+
map(present, tofrom \
677+
: arg, ps1->ps->ps->ps->s, ps2->s) \
678+
map(tofrom \
679+
: ps2->ps->ps->ps->s)
675680
{
676681
++(arg);
677682
}
@@ -694,9 +699,10 @@ void test_present_modifier(int arg) {
694699
// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
695700
#ifdef CK9
696701
void test_present_modifier(int arg) {
697-
// PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003
698-
// CK9: private unnamed_addr constant [1 x i64] [i64 [[#0x1003]]]
699-
#pragma omp target data map(present,tofrom: arg)
702+
// PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003
703+
// CK9: private unnamed_addr constant [1 x i64] [i64 [[#0x1003]]]
704+
#pragma omp target data map(present, tofrom \
705+
: arg)
700706
{++arg;}
701707
}
702708
#endif

‎clang/test/OpenMP/target_enter_data_codegen.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -267,7 +267,7 @@ double gc[100];
267267
// PRESENT=0x1000 | TO=0x1 = 0x1001
268268
// CK1A: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]]
269269

270-
// PRESENT=0x1000 | CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x1425
270+
// PRESENT=0x1000 | CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x1405
271271
// CK1A: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1405]]]
272272

273273
// CK1A-LABEL: _Z3fooi

‎clang/test/OpenMP/target_is_device_ptr_messages.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,8 @@
1-
// RUN: %clang_cc1 -std=c++11 -verify -fopenmp -ferror-limit 200 %s -Wuninitialized
1+
// RUN: %clang_cc1 -std=c++11 -verify=expected,omp45 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wuninitialized
2+
// RUN: %clang_cc1 -std=c++11 -verify=expected,omp50 -fopenmp -ferror-limit 200 %s -Wuninitialized
23

3-
// RUN: %clang_cc1 -std=c++11 -verify -fopenmp-simd -ferror-limit 200 %s -Wuninitialized
4+
// RUN: %clang_cc1 -std=c++11 -verify=expected,omp45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wuninitialized
5+
// RUN: %clang_cc1 -std=c++11 -verify=expected,omp50 -fopenmp-simd -ferror-limit 200 %s -Wuninitialized
46
struct ST {
57
int *a;
68
};

‎clang/test/OpenMP/target_map_codegen_29.cpp

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -73,9 +73,8 @@ typedef struct StructWithPtrTag : public Base {
7373
// CK30-DAG: [[S_BEGIN_BC]] = ptrtoint i8* [[S_BEGIN:%.+]] to i64
7474
// CK30-DAG: [[S_END_BC]] = ptrtoint i8* [[S_END:%.+]] to i64
7575
// CK30-DAG: [[S_BEGIN]] = bitcast [[STRUCT]]* [[S]] to i8*
76-
// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST:%.+]], i32 1
77-
// CK30-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_BC:%.+]], i{{64|32}} {{55|27}}
78-
// CK30-DAG: [[S_BC]] = bitcast [[STRUCT]]* [[S]] to i8*
76+
// CK30-DAG: [[S_END]] = bitcast [[STRUCT]]* [[REAL_S_END:%.+]] to i8*
77+
// CK30-DAG: [[REAL_S_END]] = getelementptr [[STRUCT]], [[STRUCT]]* [[S]], i32 1
7978

8079
// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 1
8180
// CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to [[STRUCT]]**
@@ -125,7 +124,9 @@ typedef struct StructWithPtrTag : public Base {
125124
// CK30-DAG: [[S_PTR1_BC]] = ptrtoint i8* [[S_PTR1:%.+]] to i64
126125
// CK30-DAG: [[S_END_BC]] = ptrtoint i8* [[S_END:%.+]] to i64
127126
// CK30-DAG: [[S_PTR1]] = bitcast i32** [[PTR2]] to i8*
128-
// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST]], i{{64|32}} 1
127+
// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST:%.+]], i{{64|32}} 1
128+
// CK30-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_BC:%.+]], i{{64|32}} {{55|27}}
129+
// CK30-DAG: [[S_BC]] = bitcast [[STRUCT]]* [[S]] to i8*
129130

130131
// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 4
131132
// CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to i32***

‎clang/test/OpenMP/target_map_codegen_31.cpp

Lines changed: 91 additions & 88 deletions
Original file line numberDiff line numberDiff line change
@@ -45,11 +45,11 @@ struct ST {
4545
// CK31A-USE: [[MTYPE00:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x1020]],
4646
// CK31A-NOUSE: [[MTYPE00:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x1000]],
4747
//
48-
// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
4948
// MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1000000001003
49+
// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
5050
// PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 | TO=0x1 = 0x1023
51-
// CK31A-USE-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]], i64 [[#0x1023]],
52-
// CK31A-NOUSE-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]], i64 [[#0x1003]],
51+
// CK31A-USE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]], i64 [[#0x1023]],
52+
// CK31A-NOUSE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]], i64 [[#0x1003]],
5353
//
5454
// PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
5555
// MEMBER_OF_5=0x5000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x5000000001003
@@ -77,91 +77,94 @@ void explicit_maps_single (int ii){
7777
struct ST st1;
7878
struct ST st2;
7979

80-
// Make sure the struct picks up present even if another element of the struct
81-
// doesn't have present.
82-
// Region 00
83-
// CK31A: [[ST1_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 0
84-
// CK31A: [[ST1_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 1
85-
// CK31A: [[ST2_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 0
86-
// CK31A: [[ST2_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 1
87-
// CK31A-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 7, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[7 x i{{.+}}]* [[MTYPE00]]{{.+}})
88-
// CK31A-DAG: [[GEPS]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0
89-
// CK31A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
90-
// CK31A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
91-
92-
// st1
93-
// CK31A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
94-
// CK31A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
95-
// CK31A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
96-
// CK31A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
97-
// CK31A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
98-
// CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP0]]
99-
// CK31A-DAG: store i32* [[ST1_I]], i32** [[CP0]]
100-
// CK31A-DAG: store i64 %{{.+}}, i64* [[S0]]
101-
102-
// st1.i
103-
// CK31A-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
104-
// CK31A-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
105-
// CK31A-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
106-
// CK31A-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
107-
// CK31A-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
108-
// CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP1]]
109-
// CK31A-DAG: store i32* [[ST1_I]], i32** [[CP1]]
110-
// CK31A-DAG: store i64 4, i64* [[S1]]
111-
112-
// st1.j
113-
// CK31A-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
114-
// CK31A-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
115-
// CK31A-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
116-
// CK31A-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]**
117-
// CK31A-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32**
118-
// CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP2]]
119-
// CK31A-DAG: store i32* [[ST1_J]], i32** [[CP2]]
120-
// CK31A-DAG: store i64 4, i64* [[S2]]
121-
122-
// a
123-
// CK31A-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
124-
// CK31A-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
125-
// CK31A-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3
126-
// CK31A-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to i32**
127-
// CK31A-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32**
128-
// CK31A-DAG: store i32* [[A]], i32** [[CBP3]]
129-
// CK31A-DAG: store i32* [[A]], i32** [[CP3]]
130-
// CK31A-DAG: store i64 4, i64* [[S3]]
131-
132-
// st2
133-
// CK31A-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 4
134-
// CK31A-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 4
135-
// CK31A-DAG: [[S4:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 4
136-
// CK31A-DAG: [[CBP4:%.+]] = bitcast i8** [[BP4]] to [[ST]]**
137-
// CK31A-DAG: [[CP4:%.+]] = bitcast i8** [[P4]] to i32**
138-
// CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP4]]
139-
// CK31A-DAG: store i32* [[ST2_I]], i32** [[CP4]]
140-
// CK31A-DAG: store i64 %{{.+}}, i64* [[S4]]
141-
142-
// st2.i
143-
// CK31A-DAG: [[BP5:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 5
144-
// CK31A-DAG: [[P5:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 5
145-
// CK31A-DAG: [[S5:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 5
146-
// CK31A-DAG: [[CBP5:%.+]] = bitcast i8** [[BP5]] to [[ST]]**
147-
// CK31A-DAG: [[CP5:%.+]] = bitcast i8** [[P5]] to i32**
148-
// CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP5]]
149-
// CK31A-DAG: store i32* [[ST2_I]], i32** [[CP5]]
150-
// CK31A-DAG: store i64 4, i64* [[S5]]
151-
152-
// st2.j
153-
// CK31A-DAG: [[BP6:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 6
154-
// CK31A-DAG: [[P6:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 6
155-
// CK31A-DAG: [[S6:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 6
156-
// CK31A-DAG: [[CBP6:%.+]] = bitcast i8** [[BP6]] to [[ST]]**
157-
// CK31A-DAG: [[CP6:%.+]] = bitcast i8** [[P6]] to i32**
158-
// CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP6]]
159-
// CK31A-DAG: store i32* [[ST2_J]], i32** [[CP6]]
160-
// CK31A-DAG: store i64 4, i64* [[S6]]
161-
162-
// CK31A-USE: call void [[CALL00:@.+]]([[ST]]* [[ST1]], i32* [[A]], [[ST]]* [[ST2]])
163-
// CK31A-NOUSE: call void [[CALL00:@.+]]()
164-
#pragma omp target map(tofrom: st1.i) map(present, tofrom: a, st1.j, st2.i) map(tofrom: st2.j)
80+
// Make sure the struct picks up present even if another element of the struct
81+
// doesn't have present.
82+
// Region 00
83+
// CK31A: [[ST1_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 1
84+
// CK31A: [[ST1_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 0
85+
// CK31A: [[ST2_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 0
86+
// CK31A: [[ST2_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 1
87+
// CK31A-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 7, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[7 x i{{.+}}]* [[MTYPE00]]{{.+}})
88+
// CK31A-DAG: [[GEPS]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0
89+
// CK31A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
90+
// CK31A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
91+
92+
// st1
93+
// CK31A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
94+
// CK31A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
95+
// CK31A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
96+
// CK31A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
97+
// CK31A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
98+
// CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP0]]
99+
// CK31A-DAG: store i32* [[ST1_I]], i32** [[CP0]]
100+
// CK31A-DAG: store i64 %{{.+}}, i64* [[S0]]
101+
102+
// st1.j
103+
// CK31A-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
104+
// CK31A-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
105+
// CK31A-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
106+
// CK31A-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
107+
// CK31A-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
108+
// CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP1]]
109+
// CK31A-DAG: store i32* [[ST1_J]], i32** [[CP1]]
110+
// CK31A-DAG: store i64 4, i64* [[S1]]
111+
112+
// st1.i
113+
// CK31A-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
114+
// CK31A-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
115+
// CK31A-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
116+
// CK31A-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]**
117+
// CK31A-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32**
118+
// CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP2]]
119+
// CK31A-DAG: store i32* [[ST1_I]], i32** [[CP2]]
120+
// CK31A-DAG: store i64 4, i64* [[S2]]
121+
122+
// a
123+
// CK31A-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
124+
// CK31A-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
125+
// CK31A-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3
126+
// CK31A-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to i32**
127+
// CK31A-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32**
128+
// CK31A-DAG: store i32* [[A]], i32** [[CBP3]]
129+
// CK31A-DAG: store i32* [[A]], i32** [[CP3]]
130+
// CK31A-DAG: store i64 4, i64* [[S3]]
131+
132+
// st2
133+
// CK31A-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 4
134+
// CK31A-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 4
135+
// CK31A-DAG: [[S4:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 4
136+
// CK31A-DAG: [[CBP4:%.+]] = bitcast i8** [[BP4]] to [[ST]]**
137+
// CK31A-DAG: [[CP4:%.+]] = bitcast i8** [[P4]] to i32**
138+
// CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP4]]
139+
// CK31A-DAG: store i32* [[ST2_I]], i32** [[CP4]]
140+
// CK31A-DAG: store i64 %{{.+}}, i64* [[S4]]
141+
142+
// st2.i
143+
// CK31A-DAG: [[BP5:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 5
144+
// CK31A-DAG: [[P5:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 5
145+
// CK31A-DAG: [[S5:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 5
146+
// CK31A-DAG: [[CBP5:%.+]] = bitcast i8** [[BP5]] to [[ST]]**
147+
// CK31A-DAG: [[CP5:%.+]] = bitcast i8** [[P5]] to i32**
148+
// CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP5]]
149+
// CK31A-DAG: store i32* [[ST2_I]], i32** [[CP5]]
150+
// CK31A-DAG: store i64 4, i64* [[S5]]
151+
152+
// st2.j
153+
// CK31A-DAG: [[BP6:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 6
154+
// CK31A-DAG: [[P6:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 6
155+
// CK31A-DAG: [[S6:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 6
156+
// CK31A-DAG: [[CBP6:%.+]] = bitcast i8** [[BP6]] to [[ST]]**
157+
// CK31A-DAG: [[CP6:%.+]] = bitcast i8** [[P6]] to i32**
158+
// CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP6]]
159+
// CK31A-DAG: store i32* [[ST2_J]], i32** [[CP6]]
160+
// CK31A-DAG: store i64 4, i64* [[S6]]
161+
162+
// CK31A-USE: call void [[CALL00:@.+]]([[ST]]* [[ST1]], i32* [[A]], [[ST]]* [[ST2]])
163+
// CK31A-NOUSE: call void [[CALL00:@.+]]()
164+
#pragma omp target map(tofrom \
165+
: st1.i) map(present, tofrom \
166+
: a, st1.j, st2.i) map(tofrom \
167+
: st2.j)
165168
{
166169
#ifdef USE
167170
st1.i++;

0 commit comments

Comments
 (0)