[OPENMP]Fix PR46012: declare target pointer cannot be accessed in tar…

…get region.

Summary:
Need to avoid an optimization for base pointer mapping for target data
directives.

Reviewers: jdoerfert, ye-luo

Subscribers: yaxunl, guansong, cfe-commits, sstefan1, caomhin

Tags: #clang

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

clang/lib/CodeGen/CGOpenMPRuntime.cpp

@@ -7379,6 +7379,9 @@ class MappableExprsHandler {
     //
     // map(p[1:24])
     // p, &p[1], 24*sizeof(float), TARGET_PARAM | TO | FROM
+    // for data directives
+    // p, p, sizeof(float*), TARGET_PARAM | TO | FROM
+    // p, &p[1], 24*sizeof(float), PTR_AND_OBJ | TO | FROM
     //
     // map(s)
     // &s, &s, sizeof(S2), TARGET_PARAM | TO | FROM
@@ -7557,9 +7560,15 @@ class MappableExprsHandler {
       if (Ty->isAnyPointerType() && std::next(I) != CE) {
         BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
 
-        // We do not need to generate individual map information for the
-        // pointer, it can be associated with the combined storage.
-        ++I;
+        // For non-data directives, we do not need to generate individual map
+        // information for the pointer, it can be associated with the combined
+        // storage.
+        if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
+            !CurDir.is<const OMPExecutableDirective *>() ||
+            !isOpenMPTargetDataManagementDirective(
+                CurDir.get<const OMPExecutableDirective *>()
+                    ->getDirectiveKind()))
+          ++I;
       }
     }
 
@@ -7633,6 +7642,7 @@ class MappableExprsHandler {
                 isa<MemberExpr>(Next->getAssociatedExpression()) ||
                 isa<ArraySubscriptExpr>(Next->getAssociatedExpression()) ||
                 isa<OMPArraySectionExpr>(Next->getAssociatedExpression()) ||
+                isa<OMPArrayShapingExpr>(Next->getAssociatedExpression()) ||
                 isa<UnaryOperator>(Next->getAssociatedExpression()) ||
                 isa<BinaryOperator>(Next->getAssociatedExpression())) &&
                "Unexpected expression");

clang/test/OpenMP/target_data_codegen.cpp

@@ -462,7 +462,7 @@ struct S2 {
 
 void test_close_modifier(int arg) {
   S2 *ps;
-  // CK5: private unnamed_addr constant [5 x i64] [i64 1059, i64 32, {{.*}}, i64 16, i64 1043]
+  // CK5: private unnamed_addr constant [6 x i64] [i64 1059, i64 32, i64 562949953422339, i64 562949953421328, i64 16, i64 1043]
   #pragma omp target data map(close,tofrom: arg, ps->ps->ps->ps->s)
   {
     ++(arg);

clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp

@@ -22,18 +22,18 @@
 double *g;
 
 // CK1: @g = global double*
-// CK1: [[MTYPE00:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE01:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE03:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE04:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE05:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE06:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE07:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE08:@.+]] = {{.*}}constant [2 x i64] [{{i64 35, i64 99|i64 99, i64 35}}]
-// CK1: [[MTYPE09:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 99]
-// CK1: [[MTYPE10:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 99]
-// CK1: [[MTYPE11:@.+]] = {{.*}}constant [2 x i64] [i64 96, i64 35]
-// CK1: [[MTYPE12:@.+]] = {{.*}}constant [2 x i64] [i64 96, i64 35]
+// CK1: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
+// CK1: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
+// CK1: [[MTYPE03:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
+// CK1: [[MTYPE04:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
+// CK1: [[MTYPE05:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
+// CK1: [[MTYPE06:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
+// CK1: [[MTYPE07:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
+// CK1: [[MTYPE08:@.+]] = {{.*}}constant [4 x i64] [i64 99, i64 19, i64 35, i64 19]
+// CK1: [[MTYPE09:@.+]] = {{.*}}constant [4 x i64] [i64 99, i64 19, i64 99, i64 19]
+// CK1: [[MTYPE10:@.+]] = {{.*}}constant [4 x i64] [i64 99, i64 19, i64 99, i64 19]
+// CK1: [[MTYPE11:@.+]] = {{.*}}constant [3 x i64] [i64 96, i64 35, i64 19]
+// CK1: [[MTYPE12:@.+]] = {{.*}}constant [3 x i64] [i64 96, i64 35, i64 19]
 
 // CK1-LABEL: @_Z3foo
 template<typename T>
@@ -42,7 +42,7 @@ void foo(float *&lr, T *&tr) {
   T *t;
 
   // CK1:     [[T:%.+]] = load double*, double** [[DECL:@g]],
-  // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to double**
   // CK1:     store double* [[T]], double** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE00]]
@@ -61,7 +61,7 @@ void foo(float *&lr, T *&tr) {
   ++g;
 
   // CK1:     [[T1:%.+]] = load float*, float** [[DECL:%.+]],
-  // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float**
   // CK1:     store float* [[T1]], float** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE01]]
@@ -92,7 +92,7 @@ void foo(float *&lr, T *&tr) {
   ++l;
 
   // CK1:     [[T1:%.+]] = load float*, float** [[DECL:%.+]],
-  // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float**
   // CK1:     store float* [[T1]], float** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE03]]
@@ -115,7 +115,7 @@ void foo(float *&lr, T *&tr) {
 
   // CK1:     [[BTHEN]]:
   // CK1:     [[T1:%.+]] = load float*, float** [[DECL:%.+]],
-  // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float**
   // CK1:     store float* [[T1]], float** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE04]]
@@ -152,7 +152,7 @@ void foo(float *&lr, T *&tr) {
 
   // CK1:     [[T2:%.+]] = load float**, float*** [[DECL:%.+]],
   // CK1:     [[T1:%.+]] = load float*, float** [[T2]],
-  // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float**
   // CK1:     store float* [[T1]], float** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE05]]
@@ -174,7 +174,7 @@ void foo(float *&lr, T *&tr) {
   ++lr;
 
   // CK1:     [[T1:%.+]] = load i32*, i32** [[DECL:%.+]],
-  // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to i32**
   // CK1:     store i32* [[T1]], i32** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE06]]
@@ -194,7 +194,7 @@ void foo(float *&lr, T *&tr) {
 
   // CK1:     [[T2:%.+]] = load i32**, i32*** [[DECL:%.+]],
   // CK1:     [[T1:%.+]] = load i32*, i32** [[T2]],
-  // CK1:     [[BP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to i32**
   // CK1:     store i32* [[T1]], i32** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE07]]
@@ -216,7 +216,7 @@ void foo(float *&lr, T *&tr) {
   ++tr;
 
   // CK1:     [[T1:%.+]] = load float*, float** [[DECL:%.+]],
-  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* %{{.+}}, i32 0, i32
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float**
   // CK1:     store float* [[T1]], float** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE08]]
@@ -280,7 +280,7 @@ void foo(float *&lr, T *&tr) {
   ++l; ++t;
 
   // CK1:     [[T1:%.+]] = load i32*, i32** [[DECL:%.+]],
-  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 0
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to i32**
   // CK1:     store i32* [[T1]], i32** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE11]]
@@ -300,7 +300,7 @@ void foo(float *&lr, T *&tr) {
 
   // CK1:     [[T2:%.+]] = load i32**, i32*** [[DECL:%.+]],
   // CK1:     [[T1:%.+]] = load i32*, i32** [[T2]],
-  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 0
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to i32**
   // CK1:     store i32* [[T1]], i32** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE12]]
@@ -348,7 +348,7 @@ void bar(float *&a, int *&b) {
 // CK2: [[ST:%.+]] = type { double*, double** }
 // CK2: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 32, i64 281474976710739]
 // CK2: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 32, i64 281474976710739]
-// CK2: [[MTYPE02:@.+]] = {{.*}}constant [3 x i64] [i64 35, i64 32, i64 562949953421392]
+// CK2: [[MTYPE02:@.+]] = {{.*}}constant [4 x i64] [i64 35, i64 19, i64 32, i64 844424930132048]
 // CK2: [[MTYPE03:@.+]] = {{.*}}constant [3 x i64] [i64 32, i64 281474976710739, i64 281474976710736]
 
 template <typename T>
@@ -404,7 +404,7 @@ struct ST {
     // CK2:     getelementptr inbounds double, double* [[TTTT]], i32 1
     b++;
 
-    // CK2:     [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 2
+    // CK2:     [[BP:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* %{{.+}}, i32 0, i32 3
     // CK2:     [[CBP:%.+]] = bitcast i8** [[BP]] to double***
     // CK2:     store double** [[RVAL:%.+]], double*** [[CBP]],
     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE02]]