Author: Alexey Bataev Date: 2021-07-09T12:51:26-07:00 New Revision: ab8989ab8710c693e83edbccf221746c897c835f
URL: https://github.com/llvm/llvm-project/commit/ab8989ab8710c693e83edbccf221746c897c835f DIFF: https://github.com/llvm/llvm-project/commit/ab8989ab8710c693e83edbccf221746c897c835f.diff LOG: [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 Added: openmp/libomptarget/test/mapping/target_pointers_members_map.cpp Modified: clang/lib/CodeGen/CGOpenMPRuntime.cpp clang/test/OpenMP/target_map_codegen_29.cpp Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index c2ef95cb1d28..8b0462988345 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -8997,11 +8997,17 @@ class MappableExprsHandler { // 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 @@ class MappableExprsHandler { 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; }); diff --git a/clang/test/OpenMP/target_map_codegen_29.cpp b/clang/test/OpenMP/target_map_codegen_29.cpp index 360a44812197..2be0e2534d6d 100644 --- 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: [[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]]* diff --git a/openmp/libomptarget/test/mapping/target_pointers_members_map.cpp b/openmp/libomptarget/test/mapping/target_pointers_members_map.cpp new file mode 100644 index 000000000000..6de639a0e16e --- /dev/null +++ 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]); +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
