ABataev created this revision. ABataev added reviewers: jdoerfert, abhinavgaba, jyu2, mikerice. Herald added subscribers: guansong, yaxunl. ABataev requested review of this revision. Herald added subscribers: openmp-commits, sstefan1. Herald added projects: clang, OpenMP.
The implicitly generated mappings for allocation/deallocation in mappers runtime should be mapped as implicit, also no need to clear member_of flag to avoid ref counter increment. Also, the ref counter should not be incremented for the very first element that comes from the mapper function. Repository: rG LLVM Github Monorepo https://reviews.llvm.org/D100673 Files: clang/lib/CodeGen/CGOpenMPRuntime.cpp clang/test/OpenMP/declare_mapper_codegen.cpp openmp/libomptarget/src/omptarget.cpp openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array_subscript.cpp openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_ptr_subscript.cpp openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_var.cpp
Index: openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_var.cpp =================================================================== --- /dev/null +++ openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_var.cpp @@ -0,0 +1,62 @@ +// 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 { + int a; + double *b; +} C1; +#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2]) + +typedef struct { + int a; + double *b; + C1 c; +} C; +#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2]) + +typedef struct { + int e; + C f; + int h; +} D; + +int main() { + constexpr int N = 10; + D s; + s.e = 111; + s.f.a = 222; + s.f.c.a = 777; + double x[2]; + double x1[2]; + x[1] = 20; + s.f.b = &x[0]; + s.f.c.b = &x1[0]; + s.h = N; + + printf("%d %d %d %4.5f %d\n", s.e, s.f.a, s.f.c.a, s.f.b[1], + s.f.b == &x[0] ? 1 : 0); + // CHECK: 111 222 777 20.00000 1 + + __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]); + +#pragma omp target map(tofrom : s) firstprivate(p) + { + printf("%d %d %d\n", s.f.a, s.f.c.a, + s.f.b == reinterpret_cast<void *>(p) ? 1 : 0); + // CHECK: 222 777 0 + s.e = 333; + s.f.a = 444; + s.f.c.a = 555; + s.f.b[1] = 40; + } + + printf("%d %d %d %4.5f %d\n", s.e, s.f.a, s.f.c.a, s.f.b[1], + s.f.b == &x[0] ? 1 : 0); + // CHECK: 333 222 777 40.00000 1 +} Index: openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_ptr_subscript.cpp =================================================================== --- /dev/null +++ openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_ptr_subscript.cpp @@ -0,0 +1,62 @@ +// 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 { + int a; + double *b; +} C1; +#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2]) + +typedef struct { + int a; + double *b; + C1 c; +} C; +#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2]) + +typedef struct { + int e; + C f; + int h; +} D; + +int main() { + constexpr int N = 10; + D s; + s.e = 111; + s.f.a = 222; + s.f.c.a = 777; + double x[2]; + double x1[2]; + x[1] = 20; + s.f.b = &x[0]; + s.f.c.b = &x1[0]; + s.h = N; + + D *sp = &s; + + printf("%d %d %d %4.5f %d\n", sp[0].e, sp[0].f.a, sp[0].f.c.a, sp[0].f.b[1], + sp[0].f.b == &x[0] ? 1 : 0); + // CHECK: 111 222 777 20.00000 1 + + __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]); +#pragma omp target map(tofrom : sp[0]) firstprivate(p) + { + printf("%d %d %d\n", sp[0].f.a, sp[0].f.c.a, + sp[0].f.b == reinterpret_cast<void *>(p) ? 1 : 0); + // CHECK: 222 777 0 + sp[0].e = 333; + sp[0].f.a = 444; + sp[0].f.c.a = 555; + sp[0].f.b[1] = 40; + } + printf("%d %d %d %4.5f %d\n", sp[0].e, sp[0].f.a, sp[0].f.c.a, sp[0].f.b[1], + sp[0].f.b == &x[0] ? 1 : 0); + // CHECK: 333 222 777 40.00000 1 +} Index: openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array_subscript.cpp =================================================================== --- /dev/null +++ openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array_subscript.cpp @@ -0,0 +1,60 @@ +// 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 { + int a; + double *b; +} C1; +#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2]) + +typedef struct { + int a; + double *b; + C1 c; +} C; +#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2]) + +typedef struct { + int e; + C f; + int h; +} D; + +int main() { + constexpr int N = 10; + D sa[10]; + sa[1].e = 111; + sa[1].f.a = 222; + sa[1].f.c.a = 777; + double x[2]; + double x1[2]; + x[1] = 20; + sa[1].f.b = &x[0]; + sa[1].f.c.b = &x1[0]; + sa[1].h = N; + + printf("%d %d %d %4.5f %d\n", sa[1].e, sa[1].f.a, sa[1].f.c.a, sa[1].f.b[1], + sa[1].f.b == &x[0] ? 1 : 0); + // CHECK: 111 222 777 20.00000 1 + + __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]); +#pragma omp target map(tofrom : sa[1]) firstprivate(p) + { + printf("%d %d %d\n", sa[1].f.a, sa[1].f.c.a, + sa[1].f.b == reinterpret_cast<void *>(p) ? 1 : 0); + // CHECK: 222 777 0 + sa[1].e = 333; + sa[1].f.a = 444; + sa[1].f.c.a = 555; + sa[1].f.b[1] = 40; + } + printf("%d %d %d %4.5f %d\n", sa[1].e, sa[1].f.a, sa[1].f.c.a, sa[1].f.b[1], + sa[1].f.b == &x[0] ? 1 : 0); + // CHECK: 333 222 777 40.00000 1 +} Index: openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp =================================================================== --- /dev/null +++ openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers_array.cpp @@ -0,0 +1,70 @@ +// 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 + +// XFAIL: clang + +#include <cstdio> +#include <cstdlib> + +typedef struct { + int a; + double *b; +} C1; +#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2]) + +typedef struct { + int a; + double *b; + C1 c; +} C; +#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2]) + +typedef struct { + int e; + C f; + int h; +} D; + +int main() { + constexpr int N = 10; + D sa[2]; + double x[2], y[2]; + double x1[2], y1[2]; + y[1] = x[1] = 20; + + sa[0].e = 111; + sa[0].f.a = 222; + sa[0].f.c.a = 777; + sa[0].f.b = &x[0]; + sa[0].f.c.b = &x1[0]; + sa[0].h = N; + + sa[1].e = 111; + sa[1].f.a = 222; + sa[1].f.c.a = 777; + sa[1].f.b = &y[0]; + sa[1].f.c.b = &y1[0]; + sa[1].h = N; + + printf("%d %d %d %4.5f %d\n", sa[1].e, sa[1].f.a, sa[1].f.c.a, sa[1].f.b[1], + sa[1].f.b == &x[0] ? 1 : 0); + // CHECK: 111 222 777 20.00000 1 + + __intptr_t p = reinterpret_cast<__intptr_t>(&y[0]); +#pragma omp target map(tofrom : sa) firstprivate(p) + { + printf("%d %d %d\n", sa[1].f.a, sa[1].f.c.a, + sa[1].f.b == reinterpret_cast<void *>(p) ? 1 : 0); + // CHECK: 222 777 0 + sa[1].e = 333; + sa[1].f.a = 444; + sa[1].f.c.a = 555; + sa[1].f.b[1] = 40; + } + printf("%d %d %d %4.5f %d\n", sa[1].e, sa[1].f.a, sa[1].f.c.a, sa[1].f.b[1], + sa[1].f.b == &x[0] ? 1 : 0); + // CHECK: 333 222 777 40.00000 1 +} Index: openmp/libomptarget/src/omptarget.cpp =================================================================== --- openmp/libomptarget/src/omptarget.cpp +++ openmp/libomptarget/src/omptarget.cpp @@ -472,7 +472,8 @@ // then no argument is marked as TARGET_PARAM ("omp target data map" is not // associated with a target region, so there are no target parameters). This // may be considered a hack, we could revise the scheme in the future. - bool UpdateRef = !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF); + bool UpdateRef = + !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) && !(FromMapper && i == 0); if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) { DP("Has a pointer entry: \n"); // Base is address of pointer. @@ -665,8 +666,7 @@ bool IsLast, IsHostPtr; bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT; bool UpdateRef = !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) || - (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && - (!FromMapper || I != ArgNum - 1)); + (!FromMapper || I != ArgNum - 1); bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE; bool HasCloseModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_CLOSE; bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT; @@ -719,8 +719,7 @@ // exists. if (((ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) && !(ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) || - (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && FromMapper && - I == ArgNum - 1)) { + (FromMapper && I == ArgNum - 1)) { DelEntry = false; // protect parent struct from being deallocated } Index: clang/test/OpenMP/declare_mapper_codegen.cpp =================================================================== --- clang/test/OpenMP/declare_mapper_codegen.cpp +++ clang/test/OpenMP/declare_mapper_codegen.cpp @@ -118,8 +118,11 @@ // CK0: [[INIT]] // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8 -// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652 -// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}}) + +// Remove movement mappings and mark as implicit +// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 +// CK0-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512 +// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}}) // CK0: br label %[[LHEAD:[^,]+]] // CK0: [[LHEAD]] @@ -228,8 +231,11 @@ // CK0: [[EVALDEL]] // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8 -// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652 -// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}}) + +// Remove movement mappings and mark as implicit +// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 +// CK0-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512 +// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}}) // CK0: br label %[[DONE]] // CK0: [[DONE]] // CK0: ret void @@ -672,8 +678,11 @@ // CK1: [[INITEVALDEL]] // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4 -// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652 -// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}}) + +// Remove movement mappings and mark as implicit +// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 +// CK1-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512 +// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}}) // CK1: br label %[[LHEAD:[^,]+]] // CK1: [[LHEAD]] @@ -718,8 +727,11 @@ // CK1: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 // CK1: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]] // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4 -// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652 -// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}}) + +// Remove movement mappings and mark as implicit +// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 +// CK1-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512 +// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}}) // CK1: br label %[[DONE]] // CK1: [[DONE]] // CK1: ret void @@ -793,8 +805,11 @@ // CK2: [[INITEVALDEL]] // CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 -// CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652 -// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}}) + +// Remove movement mappings and mark as implicit +// CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 +// CK2-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512 +// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}}) // CK2: br label %[[LHEAD:[^,]+]] // CK2: [[LHEAD]] @@ -841,8 +856,11 @@ // CK2: br i1 [[CMP1]], label %[[EVALDEL:[^,]+]], label %[[DONE]] // CK2: [[EVALDEL]] // CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 -// CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652 -// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}}) + +// Remove movement mappings and mark as implicit +// CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 +// CK2-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512 +// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}}) // CK2: br label %[[DONE]] // CK2: [[DONE]] // CK2: ret void @@ -998,8 +1016,11 @@ // CK4: [[INITEVALDEL]] // CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 // CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8 -// CK4-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652 -// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}}) + +// Remove movement mappings and mark as implicit +// CK4-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 +// CK4-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512 +// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}}) // CK4: br label %[[LHEAD:[^,]+]] // CK4: [[LHEAD]] @@ -1108,8 +1129,11 @@ // CK4: [[EVALDEL]] // CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 // CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8 -// CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652 -// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}}) + +// Remove movement mappings and mark as implicit +// CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 +// CK4-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512 +// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}}) // CK4: br label %[[DONE]] // CK4: [[DONE]] // CK4: ret void Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -10005,8 +10005,10 @@ llvm::Value *MapTypeArg = MapperCGF.Builder.CreateAnd( MapType, MapperCGF.Builder.getInt64(~(MappableExprsHandler::OMP_MAP_TO | - MappableExprsHandler::OMP_MAP_FROM | - MappableExprsHandler::OMP_MAP_MEMBER_OF))); + MappableExprsHandler::OMP_MAP_FROM))); + MapTypeArg = MapperCGF.Builder.CreateOr( + MapTypeArg, + MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_IMPLICIT)); // Call the runtime API __tgt_push_mapper_component to fill up the runtime // data structure.
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits