[OPENMP]Fix overlapped mapping for dereferenced pointer members.

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. Differential Revision: https://reviews.llvm.org/D105562
2021-07-07 05:26:53 -07:00 · 2021-07-07 05:26:53 -07:00 · ab8989ab87
parent 4a3b055653
commit ab8989ab87
3 changed files with 86 additions and 68 deletions
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@ -8997,11 +8997,17 @@ public:
          // 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)
-                  ->getAssociatedExpression()
-                  ->getType()
-                  ->isPointerType())
+              (std::prev(It)->getAssociatedDeclaration() &&
+               std::prev(It)
+                   ->getAssociatedDeclaration()
+                   ->getType()
+                   ->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 =
@ -9061,7 +9067,7 @@ public:
            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;
                });
--- a/clang/test/OpenMP/target_map_codegen_29.cpp
+++ b/clang/test/OpenMP/target_map_codegen_29.cpp
@ -38,9 +38,9 @@

 // 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;
@ -55,18 +55,18 @@ typedef struct StructWithPtrTag : public Base {
  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:%.+]]
@ -76,78 +76,35 @@ typedef struct StructWithPtrTag : public Base {
 // 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: [[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 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: [[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]]*
--- a/openmp/libomptarget/test/mapping/target_pointers_members_map.cpp
+++ b/openmp/libomptarget/test/mapping/target_pointers_members_map.cpp
@ -0,0 +1,55 @@
+// 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]);
+}