Revert "[OPENMP]Fix PR46824: Global declare target pointer cannot be …

…accessed in target region." This reverts commit 142d0d3 to investigate undefined behavior revealed by buildbots.
llvm · Jul 30, 2020 · b69357c · b69357c
1 parent 4e6176f
commit b69357c
Show file tree

Hide file tree

Showing 7 changed files with 212 additions and 255 deletions.
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7392,9 +7392,10 @@ class MappableExprsHandler {
     // &p, &p, sizeof(float*), TARGET_PARAM | TO | FROM
     //
     // map(p[1:24])
-    // &p, &p[1], 24*sizeof(float), TARGET_PARAM | TO | FROM | PTR_AND_OBJ
-    // in unified shared memory mode or for local pointers
     // 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
@@ -7529,7 +7530,6 @@ class MappableExprsHandler {
     // Track if the map information being generated is the first for a list of
     // components.
     bool IsExpressionFirstInfo = true;
-    bool FirstPointerInComplexData = false;
     Address BP = Address::invalid();
     const Expr *AssocExpr = I->getAssociatedExpression();
     const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr);
@@ -7572,16 +7572,17 @@ class MappableExprsHandler {
       QualType Ty =
           I->getAssociatedDeclaration()->getType().getNonReferenceType();
       if (Ty->isAnyPointerType() && std::next(I) != CE) {
-        // No need to generate individual map information for the pointer, it
-        // can be associated with the combined storage if shared memory mode is
-        // active or the base declaration is not global variable.
-        const auto *VD = dyn_cast<VarDecl>(I->getAssociatedDeclaration());
+        BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
+
+        // 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() ||
-            !VD || VD->hasLocalStorage())
-          BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
-        else
-          FirstPointerInComplexData = true;
-        ++I;
+            !CurDir.is<const OMPExecutableDirective *>() ||
+            !isOpenMPTargetDataManagementDirective(
+                CurDir.get<const OMPExecutableDirective *>()
+                    ->getDirectiveKind()))
+          ++I;
       }
     }
 
@@ -7616,19 +7617,8 @@ class MappableExprsHandler {
         EncounteredME = dyn_cast<MemberExpr>(I->getAssociatedExpression());
         // If we encounter a PTR_AND_OBJ entry from now on it should be marked
         // as MEMBER_OF the parent struct.
-        if (EncounteredME) {
+        if (EncounteredME)
           ShouldBeMemberOf = true;
-          // Do not emit as complex pointer if this is actually not array-like
-          // expression.
-          if (FirstPointerInComplexData) {
-            QualType Ty = std::prev(I)
-                              ->getAssociatedDeclaration()
-                              ->getType()
-                              .getNonReferenceType();
-            BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
-            FirstPointerInComplexData = false;
-          }
-        }
       }
 
       auto Next = std::next(I);
@@ -7770,8 +7760,7 @@ class MappableExprsHandler {
           // (there is a set of entries for each capture).
           OpenMPOffloadMappingFlags Flags =
               getMapTypeBits(MapType, MapModifiers, MotionModifiers, IsImplicit,
-                             !IsExpressionFirstInfo || RequiresReference ||
-                                 FirstPointerInComplexData,
+                             !IsExpressionFirstInfo || RequiresReference,
                              IsCaptureFirstInfo && !RequiresReference);
 
           if (!IsExpressionFirstInfo) {
@@ -7830,7 +7819,6 @@ class MappableExprsHandler {
 
         IsExpressionFirstInfo = false;
         IsCaptureFirstInfo = false;
-        FirstPointerInComplexData = false;
       }
     }
   }
@@ -8079,7 +8067,6 @@ class MappableExprsHandler {
     // emission of that entry until the whole struct has been processed.
     llvm::MapVector<const ValueDecl *, SmallVector<DeferredDevicePtrEntryTy, 4>>
         DeferredInfo;
-    MapCombinedInfoTy UseDevicePtrCombinedInfo;
 
     for (const auto *C :
          CurExecDir->getClausesOfKind<OMPUseDevicePtrClause>()) {
@@ -8099,24 +8086,13 @@ class MappableExprsHandler {
         // We potentially have map information for this declaration already.
         // Look for the first set of components that refer to it.
         if (It != Info.end()) {
-          auto *CI = llvm::find_if(It->second, [VD](const MapInfo &MI) {
-            return MI.Components.back().getAssociatedDeclaration() == VD;
-          });
+          auto CI = std::find_if(
+              It->second.begin(), It->second.end(), [VD](const MapInfo &MI) {
+                return MI.Components.back().getAssociatedDeclaration() == VD;
+              });
           // If we found a map entry, signal that the pointer has to be returned
           // and move on to the next declaration.
-          // Exclude cases where the base pointer is mapped as array subscript,
-          // array section or array shaping. The base address is passed as a
-          // pointer to base in this case and cannot be used as a base for
-          // use_device_ptr list item.
-          auto PrevCI = std::next(CI->Components.rbegin());
-          const auto *VarD = dyn_cast<VarDecl>(VD);
-          if (CI != It->second.end() &&
-              (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
-               isa<MemberExpr>(IE) ||
-               !VD->getType().getNonReferenceType()->isPointerType() ||
-               PrevCI == CI->Components.rend() ||
-               isa<MemberExpr>(PrevCI->getAssociatedExpression()) || !VarD ||
-               VarD->hasLocalStorage())) {
+          if (CI != It->second.end()) {
             CI->ReturnDevicePointer = true;
             continue;
           }
@@ -8139,13 +8115,13 @@ class MappableExprsHandler {
         } else {
           llvm::Value *Ptr =
               CGF.EmitLoadOfScalar(CGF.EmitLValue(IE), IE->getExprLoc());
-          UseDevicePtrCombinedInfo.BasePointers.emplace_back(Ptr, VD);
-          UseDevicePtrCombinedInfo.Pointers.push_back(Ptr);
-          UseDevicePtrCombinedInfo.Sizes.push_back(
+          CombinedInfo.BasePointers.emplace_back(Ptr, VD);
+          CombinedInfo.Pointers.push_back(Ptr);
+          CombinedInfo.Sizes.push_back(
               llvm::Constant::getNullValue(CGF.Int64Ty));
-          UseDevicePtrCombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM |
-                                                   OMP_MAP_TARGET_PARAM);
-          UseDevicePtrCombinedInfo.Mappers.push_back(nullptr);
+          CombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM |
+                                       OMP_MAP_TARGET_PARAM);
+          CombinedInfo.Mappers.push_back(nullptr);
         }
       }
     }
@@ -8297,8 +8273,6 @@ class MappableExprsHandler {
       // We need to append the results of this capture to what we already have.
       CombinedInfo.append(CurInfo);
     }
-    // Append data for use_device_ptr clauses.
-    CombinedInfo.append(UseDevicePtrCombinedInfo);
   }
 
   /// Generate all the base pointers, section pointers, sizes, map types, and

diff --git a/clang/test/OpenMP/target_data_codegen.cpp b/clang/test/OpenMP/target_data_codegen.cpp
@@ -555,7 +555,7 @@ struct S2 {
 
 void test_close_modifier(int arg) {
   S2 *ps;
-  // CK5: private unnamed_addr constant [5 x i64] [i64 1059, i64 32, i64 562949953421328, 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);
@@ -634,17 +634,20 @@ void test_present_modifier(int arg) {
   // Make sure the struct picks up present even if another element of the struct
   // doesn't have present.
 
-  // CK8: private unnamed_addr constant [11 x i64]
+  // CK8: private unnamed_addr constant [15 x i64]
 
   // ps1
   //
   // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
   // MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
+  // MEMBER_OF_1=0x1000000000000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1000000000013
+  // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1000000001003
   // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1000000001010
   // PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010
   // PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013
   //
   // CK8-SAME: {{^}} [i64 [[#0x1020]], i64 [[#0x1000000000003]],
+  // CK8-SAME: {{^}} i64 [[#0x1000000000013]], i64 [[#0x1000000001003]],
   // CK8-SAME: {{^}} i64 [[#0x1000000001010]], i64 [[#0x1010]], i64 [[#0x1013]],
 
   // arg
@@ -656,13 +659,16 @@ void test_present_modifier(int arg) {
   // ps2
   //
   // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
-  // MEMBER_OF_7=0x7000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x7000000001003
-  // MEMBER_OF_7=0x7000000000000 | PTR_AND_OBJ=0x10 = 0x7000000000010
+  // MEMBER_OF_9=0x9000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x9000000001003
+  // MEMBER_OF_9=0x9000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x9000000001013
+  // MEMBER_OF_9=0x9000000000000 | FROM=0x2 | TO=0x1 = 0x9000000000003
+  // MEMBER_OF_9=0x9000000000000 | PTR_AND_OBJ=0x10 = 0x9000000000010
   // PTR_AND_OBJ=0x10 = 0x10
   // PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13
   //
-  // CK8-SAME: {{^}} i64 [[#0x1020]], i64 [[#0x7000000001003]],
-  // CK8-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
+  // CK8-SAME: {{^}} i64 [[#0x1020]], i64 [[#0x9000000001003]],
+  // CK8-SAME: {{^}} i64 [[#0x9000000001013]], i64 [[#0x9000000000003]],
+  // CK8-SAME: {{^}} i64 [[#0x9000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
   #pragma omp target data map(tofrom: ps1->s) \
                           map(present,tofrom: arg, ps1->ps->ps->ps->s, ps2->s) \
                           map(tofrom: ps2->ps->ps->ps->s)

diff --git a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
@@ -22,18 +22,18 @@
 double *g;
 
 // CK1: @g = global double*
-// CK1: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96]
-// 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 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 35, i64 96]
-// CK1: [[MTYPE12:@.+]] = {{.*}}constant [2 x i64] [i64 35, i64 96]
+// 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 [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
+  // 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 0
+  // 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 1
+  // 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 1
+  // 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]]