[OPENMP]Fix overlapped mapping for dereferenced pointer members.

[ABataev]

Summary

If the base is used in a map clause and later we have a memberexpr with
this base, and the member is a pointer, and this pointer is dereferenced
anyhow (subscript, array section, dereference, etc.), such components
should be considered as overlapped, otherwise it may lead to incorrect
size computations, since we try to map a pointee as a part of the whole
struct, which is not true for the pointer members.

[jyu2]

This will eliminate some IRS and it works okay with me. Thanks! But I don't know enough about this domain, please wait for some time to see if others may command on it.

clang/test/OpenMP/target_map_codegen_29.cpp

86

So instead generate map for individual pointer field, just generate one for each base with whole struct size. Make sense.

[ABataev]

Rebase + added runtime test intended to be fixed by the patch.

[abhinavgaba]

Thanks for the fix, Alexey.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

for (const MapData &L : DeclComponentLists) {
         if (CI == CE || SI == SE) {
           // Ignore it if it is the same component.
           if (CI == CE && SI == SE)
             continue;
           const auto It = (SI == SE) ? CI : SI;
           // If one component is a pointer and another one is a kind of
           // dereference of this pointer (array subscript, section, dereference,
           // etc.), it is not an overlapping.
+          // Same, if one component is a base and another component is a
+          // dereferenced pointer memberexpr with the same base.
           if (!isa<MemberExpr>(It->getAssociatedExpression()) ||
+              (std::prev(It)->getAssociatedDeclaration() &&
                std::prev(It)
-                  ->getAssociatedExpression()
+                   ->getAssociatedDeclaration()
                    ->getType()
-                  ->isPointerType())
+                   ->isPointerType()) ||
+              (It->getAssociatedDeclaration() &&
+               It->getAssociatedDeclaration()->getType()->isPointerType() &&
+               std::next(It) != CE && std::next(It) != SE))
             continue;
           const MapData &BaseData = CI == CE ? L : L1;
           OMPClauseMappableExprCommon::MappableExprComponentListRef SubData =
               SI == SE ? Components : Components1;
           auto &OverlappedElements = OverlappedData.FindAndConstruct(&BaseData);
           OverlappedElements.getSecond().push_back(SubData);
         }
       }
for (auto &Pair : OverlappedData) {
             // List with less elements is less than list with more elements.
             if (CI == CE || SI == SE)
               return CI == CE;
 
             const auto *FD1 = cast<FieldDecl>(CI->getAssociatedDeclaration());
             const auto *FD2 = cast<FieldDecl>(SI->getAssociatedDeclaration());
             if (FD1->getParent() == FD2->getParent())
               return FD1->getFieldIndex() < FD2->getFieldIndex();
-            const auto It =
+            const auto *It =
                 llvm::find_if(Layout, [FD1, FD2](const FieldDecl *FD) {
                   return FD == FD1 || FD == FD2;
                 });
             return *It == FD1;
           });
     }
 
     // Associated with a capture, because the mapping flags depend on it.

clang/test/OpenMP/target_map_codegen_29.cpp

 // SIMD-ONLY30-NOT: {{__kmpc|__tgt}}
 #ifdef CK30
 
 // CK30-DAG: [[BASE:%.+]] = type { i32*, i32, i32* }
 // CK30-DAG: [[STRUCT:%.+]] = type { [[BASE]], i32*, i32*, i32, i32* }
 
 // CK30-LABEL: @.__omp_offloading_{{.*}}map_with_deep_copy{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // The first element: 0x20 - OMP_MAP_TARGET_PARAM
-// 2-4: 0x1000000000003 - OMP_MAP_MEMBER_OF(0) | OMP_MAP_TO | OMP_MAP_FROM - copies all the data in structs excluding deep-copied elements (from &s to &s.ptrBase1, from &s.ptr to &s.ptr1, from &s.ptr1 to end of s).
-// 5-6: 0x1000000000013 - OMP_MAP_MEMBER_OF(0) | OMP_MAP_PTR_AND_OBJ | OMP_MAP_TO | OMP_MAP_FROM - deep copy of the pointers + pointee.
-// CK30: [[MTYPE00:@.+]] = private {{.*}}constant [6 x i64] [i64 32, i64 281474976710659, i64 281474976710659, i64 281474976710659, i64 281474976710675, i64 281474976710675]
+// 2: 0x1000000000003 - OMP_MAP_MEMBER_OF(0) | OMP_MAP_TO | OMP_MAP_FROM - copies all the data in structs excluding deep-copied elements (from &s to end of s).
+// 3-4: 0x1000000000013 - OMP_MAP_MEMBER_OF(0) | OMP_MAP_PTR_AND_OBJ | OMP_MAP_TO | OMP_MAP_FROM - deep copy of the pointers + pointee.
+// CK30: [[MTYPE00:@.+]] = private {{.*}}constant [4 x i64] [i64 32, i64 281474976710659, i64 281474976710675, i64 281474976710675]
 
 typedef struct {
   int *ptrBase;
   int valBase;
   int *ptrBase1;
 } Base;
 
 typedef struct StructWithPtrTag : public Base {
   int *ptr;
   int *ptr2;
   int val;
   int *ptr1;
 } StructWithPtr;
 
-// CK30-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @.__omp_offloading_{{.*}}map_with_deep_copy{{.*}}_l{{[0-9]+}}.region_id, i32 6, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MTYPE00]], i32 0, i32 0), i8** null, i8** null)
-// CK30-DAG: [[GEPS]] = getelementptr inbounds [6 x i{{64|32}}], [6 x i64]* [[SIZES:%.+]], i32 0, i32 0
-// CK30-DAG: [[GEPP]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS:%.+]], i32 0, i32 0
-// CK30-DAG: [[GEPBP]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES:%.+]], i32 0, i32 0
+// CK30-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @.__omp_offloading_{{.*}}map_with_deep_copy{{.*}}_l{{[0-9]+}}.region_id, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], i64* getelementptr inbounds ([4 x i64], [4 x i64]* [[MTYPE00]], i32 0, i32 0), i8** null, i8** null)
+// CK30-DAG: [[GEPS]] = getelementptr inbounds [4 x i{{64|32}}], [4 x i64]* [[SIZES:%.+]], i32 0, i32 0
+// CK30-DAG: [[GEPP]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[PTRS:%.+]], i32 0, i32 0
+// CK30-DAG: [[GEPBP]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BASES:%.+]], i32 0, i32 0
 
-// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 0
+// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BASES]], i32 0, i32 0
 // CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to [[STRUCT]]**
 // CK30-DAG: store [[STRUCT]]* [[S:%.+]], [[STRUCT]]** [[BC]],
-// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 0
+// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[PTRS]], i32 0, i32 0
 // CK30-DAG: [[BC:%.+]] = bitcast i8** [[PTR]] to [[STRUCT]]**
 // CK30-DAG: store [[STRUCT]]* [[S]], [[STRUCT]]** [[BC]],
-// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i{{64|32}}], [6 x i64]* [[SIZES]], i32 0, i32 0
+// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [4 x i{{64|32}}], [4 x i64]* [[SIZES]], i32 0, i32 0
 // CK30-DAG: store i64 [[S_ALLOC_SIZE:%.+]], i64* [[SIZE]],
 // CK30-DAG: [[S_ALLOC_SIZE]] = sdiv exact i64 [[DIFF:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
 // CK30-DAG: [[DIFF]] = sub i64 [[S_END_BC:%.+]], [[S_BEGIN_BC:%.+]]
 // CK30-DAG: [[S_BEGIN_BC]] = ptrtoint i8* [[S_BEGIN:%.+]] to i64
 // CK30-DAG: [[S_END_BC]] = ptrtoint i8* [[S_END:%.+]] to i64
 // CK30-DAG: [[S_BEGIN]] = bitcast [[STRUCT]]* [[S]] to i8*
 // CK30-DAG: [[S_END]] = bitcast [[STRUCT]]* [[REAL_S_END:%.+]] to i8*
 // CK30-DAG: [[REAL_S_END]] = getelementptr [[STRUCT]], [[STRUCT]]* [[S]], i32 1
 
-// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 1
+// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BASES]], i32 0, i32 1
 // CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to [[STRUCT]]**
 // CK30-DAG: store [[STRUCT]]* [[S]], [[STRUCT]]** [[BC]],
-// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 1
+// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[PTRS]], i32 0, i32 1
 // CK30-DAG: [[BC:%.+]] = bitcast i8** [[PTR]] to [[STRUCT]]**
 // CK30-DAG: store [[STRUCT]]* [[S]], [[STRUCT]]** [[BC]],
-// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[SIZES]], i32 0, i32 1
-// CK30-DAG: store i64 [[SIZE1:%.+]], i64* [[SIZE]],
-// CK30-DAG: [[SIZE1]] = sdiv exact i64 [[DIFF:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
-// CK30-DAG: [[DIFF]] = sub i64 [[S_PTRBASE1_BC:%.+]], [[S_BEGIN_BC:%.+]]
-// CK30-DAG: [[S_BEGIN_BC]] = ptrtoint i8* [[S_BEGIN:%.+]] to i64
-// CK30-DAG: [[S_PTRBASE1_BC]] = ptrtoint i8* [[S_PTRBASE1:%.+]] to i64
-// CK30-DAG: [[S_PTRBASE1]] = bitcast i32** [[S_PTRBASE1_REF:%.+]] to i8*
-// CK30-DAG: [[S_BEGIN]] = bitcast [[STRUCT]]* [[S]] to i8*
-// CK30-DAG: [[S_PTRBASE1_REF]] = getelementptr inbounds [[BASE]], [[BASE]]* [[BASE_ADDR:%.+]], i32 0, i32 2
-// CK30-DAG: [[BASE_ADDR]] = bitcast [[STRUCT]]* [[S]] to [[BASE]]*
-
-// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 2
-// CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to [[STRUCT]]**
-// CK30-DAG: store [[STRUCT]]* [[S]], [[STRUCT]]** [[BC]],
-// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 2
-// CK30-DAG: [[BC:%.+]] = bitcast i8** [[PTR]] to i32***
-// CK30-DAG: store i32** [[PTR1:%.+]], i32*** [[BC]],
-// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[SIZES]], i32 0, i32 2
-// CK30-DAG: store i64 [[SIZE2:%.+]], i64* [[SIZE]],
-// CK30-DAG: [[PTR1]] = getelementptr i32*, i32** [[S_PTRBASE1_REF]], i{{64|32}} 1
-// CK30-DAG: [[SIZE2]] = sdiv exact i64 [[DIFF:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
-// CK30-DAG: [[DIFF]] = sub i64 [[S_PTR1_BC:%.+]], [[S_PTRBASE1_BC:%.+]]
-// CK30-DAG: [[S_PTR1_BC]] = ptrtoint i8* [[S_PTR1:%.+]] to i64
-// CK30-DAG: [[S_PTRBASE1_BC]] = ptrtoint i8* [[S_PTRBASE1:%.+]] to i64
-// CK30-DAG: [[S_PTR1]] = bitcast i32** [[S_PTR1_REF:%.+]] to i8*
-// CK30-DAG: [[S_PTRBASE1]] = bitcast i32** [[PTR1]] to i8*
-// CK30-DAG: [[S_PTR1_REF]] = getelementptr inbounds [[STRUCT]], [[STRUCT]]* [[S]], i32 0, i32 4
-
-// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 3
-// CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to [[STRUCT]]**
-// CK30-DAG: store [[STRUCT]]* [[S]], [[STRUCT]]** [[BC]],
-// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 3
-// CK30-DAG: [[BC:%.+]] = bitcast i8** [[PTR]] to i32***
-// CK30-DAG: store i32** [[PTR2:%.+]], i32*** [[BC]],
-// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[SIZES]], i32 0, i32 3
-// CK30-DAG: store i64 [[SIZE3:%.+]], i64* [[SIZE]],
-// CK30-DAG: [[PTR2]] = getelementptr i32*, i32** [[S_PTR1_REF]], i{{64|32}} 1
-// CK30-DAG: [[SIZE3]] = sdiv exact i64 [[DIFF:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
-// CK30-DAG: [[DIFF]] = sub i64 [[S_END_BC:%.+]], [[S_PTR1_BC:%.+]]
-// CK30-DAG: [[S_PTR1_BC]] = ptrtoint i8* [[S_PTR1:%.+]] to i64
-// CK30-DAG: [[S_END_BC]] = ptrtoint i8* [[S_END:%.+]] to i64
-// CK30-DAG: [[S_PTR1]] = bitcast i32** [[PTR2]] to i8*
-// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST:%.+]], i{{64|32}} 1
-// CK30-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_BC:%.+]], i{{64|32}} {{55|27}}
-// CK30-DAG: [[S_BC]] = bitcast [[STRUCT]]* [[S]] to i8*
+// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [4 x i64], [4 x i64]* [[SIZES]], i32 0, i32 1
+// CK30-DAG: store i64 {{56|28}}, i64* [[SIZE]],
 
-// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 4
+// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BASES]], i32 0, i32 2
 // CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to i32***
 // CK30-DAG: store i32** [[S_PTR1:%.+]], i32*** [[BC]],
-// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 4
+// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[PTRS]], i32 0, i32 2
 // CK30-DAG: [[BC:%.+]] = bitcast i8** [[PTR]] to i32**
 // CK30-DAG: store i32* [[S_PTR1_BEGIN:%.+]], i32** [[BC]],
-// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[SIZES]], i32 0, i32 4
+// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [4 x i64], [4 x i64]* [[SIZES]], i32 0, i32 2
 // CK30-DAG: store i64 4, i64* [[SIZE]],
 // CK30-DAG: [[S_PTR1]] = getelementptr inbounds [[STRUCT]], [[STRUCT]]* [[S]], i32 0, i32 4
 // CK30-DAG: [[S_PTR1_BEGIN]] = getelementptr inbounds i32, i32* [[S_PTR1_BEGIN_REF:%.+]], i{{64|32}} 0
 // CK30-DAG: [[S_PTR1_BEGIN_REF]] = load i32*, i32** [[S_PTR1:%.+]],
 // CK30-DAG: [[S_PTR1]] = getelementptr inbounds [[STRUCT]], [[STRUCT]]* [[S]], i32 0, i32 4
 
-// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 5
+// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BASES]], i32 0, i32 3
 // CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to i32***
 // CK30-DAG: store i32** [[S_PTRBASE1:%.+]], i32*** [[BC]],
-// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 5
+// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[PTRS]], i32 0, i32 3
 // CK30-DAG: [[BC:%.+]] = bitcast i8** [[PTR]] to i32**
 // CK30-DAG: store i32* [[S_PTRBASE1_BEGIN:%.+]], i32** [[BC]],
-// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i{{64|32}}], [6 x i{{64|32}}]* [[SIZES]], i32 0, i32 5
+// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [4 x i{{64|32}}], [4 x i{{64|32}}]* [[SIZES]], i32 0, i32 3
 // CK30-DAG: store i{{64|32}} 4, i{{64|32}}* [[SIZE]],
 // CK30-DAG: [[S_PTRBASE1]] = getelementptr inbounds [[BASE]], [[BASE]]* [[S_BASE:%.+]], i32 0, i32 2
 // CK30-DAG: [[S_BASE]] = bitcast [[STRUCT]]* [[S]] to [[BASE]]*
 // CK30-DAG: [[S_PTRBASE1_BEGIN]] = getelementptr inbounds i32, i32* [[S_PTRBASE1_BEGIN_REF:%.+]], i{{64|32}} 0
 // CK30-DAG: [[S_PTRBASE1_BEGIN_REF]] = load i32*, i32** [[S_PTRBASE1:%.+]],
 // CK30-DAG: [[S_PTRBASE1]] = getelementptr inbounds [[BASE]], [[BASE]]* [[S_BASE:%.+]], i32 0, i32 2
 // CK30-DAG: [[S_BASE]] = bitcast [[STRUCT]]* [[S]] to [[BASE]]*
 void map_with_deep_copy() {

openmp/libomptarget/test/mapping/target_pointers_members_map.cpp

+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <cstdio>
+#include <cstdlib>
+
+typedef struct {
+  short *a;
+  long d1, d2;
+} DV_A;
+
+typedef struct {
+  DV_A b;
+  long d3;
+} C;
+
+typedef struct {
+  C *c;
+  long d4, d5;
+} DV_B;
+
+int main() {
+
+  short arr1[10] = {10, 11, 12, 13, 14, 15, 16, 17, 18, 19};
+  short arr2[10] = {20, 31, 22, 23, 24, 25, 26, 27, 28, 29};
+
+  C c1[2];
+  c1[0].b.a = (short *)arr1;
+  c1[1].b.a = (short *)arr2;
+  c1[0].b.d1 = 111;
+
+  DV_B dvb1;
+  dvb1.c = (C *)&c1;
+
+  // CHECK: 10 111
+  printf("%d %ld %p %p %p %p\n", dvb1.c[0].b.a[0], dvb1.c[0].b.d1, &dvb1,
+         &dvb1.c[0], &dvb1.c[0].b, &dvb1.c[0].b.a[0]);
+#pragma omp target map(to                                                      \
+                       : dvb1, dvb1.c [0:2])                                   \
+    map(tofrom                                                                 \
+        : dvb1.c[0].b.a [0:10], dvb1.c[1].b.a [0:10])
+  {
+    // CHECK: 10 111
+    printf("%d %ld %p %p %p %p\n", dvb1.c[0].b.a[0], dvb1.c[0].b.d1, &dvb1,
+           &dvb1.c[0], &dvb1.c[0].b, &dvb1.c[0].b.a[0]);
+    dvb1.c[0].b.a[0] = 333;
+    dvb1.c[0].b.d1 = 444;
+  }
+  // CHECK: 333 111
+  printf("%d %ld %p %p %p %p\n", dvb1.c[0].b.a[0], dvb1.c[0].b.d1, &dvb1,
+         &dvb1.c[0], &dvb1.c[0].b, &dvb1.c[0].b.a[0]);
+}