[OPENMP]Use the real pointer value as base, not indexed value.

[ABataev]

Summary

After fix for PR48174 the base pointer for pointer-based
array-sections/array-subscripts will be emitted as &ptr[idx], but
actually it should be just ptr, i.e. the address stored in the ponter
to point correctly to the beginning of the array. Currently it may lead
to a crash in the runtime.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

for (; I != CE; ++I) {
         // The pointer becomes the base for the next element.
         if (Next != CE)
           BP = LB;
 
         IsExpressionFirstInfo = false;
         IsCaptureFirstInfo = false;
         FirstPointerInComplexData = false;
       } else if (FirstPointerInComplexData) {
-        BP = CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
-                 .getAddress(CGF);
+        QualType Ty = Components.rbegin()
+                          ->getAssociatedDeclaration()
+                          ->getType()
+                          .getNonReferenceType();
+        BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
         FirstPointerInComplexData = false;
       }
     }
 
     if (!IsNonContiguous)
       return;
 
     const ASTContext &Context = CGF.getContext();

clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp

 MyObject *objects;
 #pragma omp end declare target
 
 // CHECK-DAG: [[SIZES0:@.+]] = private unnamed_addr constant [1 x i64] [i64 {{8|4}}]
 // CHECK-DAG: [[MAPS0:@.+]] = private unnamed_addr constant [1 x i64] [i64 49]
 // CHECK-DAG: [[MAPS1:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 281474976710673]
 // CHECK: @main
 int main(void) {
+// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
+// CHECK: [[BPTRC0:%.+]] = bitcast i8** [[BPTR0]] to %struct.MyObject***
+// CHECK: store %struct.MyObject** @objects, %struct.MyObject*** [[BPTRC0]],
 // CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[SIZES0]], i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[MAPS0]], i32 0, i32 0), i8** null, i8** null)
-#pragma omp target enter data map(to : objects [0:1])
+#pragma omp target enter data map(to : objects [1:1])
+// CHECK: [[OBJ:%.+]] = load %struct.MyObject*, %struct.MyObject** @objects,
+// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+// CHECK: [[BPTRC0:%.+]] = bitcast i8** [[BPTR0]] to %struct.MyObject**
+// CHECK: store %struct.MyObject* [[OBJ]], %struct.MyObject** [[BPTRC0]],
 // CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 2, i8** %{{.+}}, i8** %{{.+}}, i64* %{{.+}}, i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPS1]], i32 0, i32 0), i8** null, i8** null)
-#pragma omp target enter data map(to : objects[0].arr [0:1])
+#pragma omp target enter data map(to : objects[1].arr [0:1])
 
   return 0;
 }
 #endif