[OPENMP50]Mapping of the subcomponents with the 'default' mappers.
If the mapped structure has data members, which have 'default' mappers,
need to map these members individually using their 'default' mappers.
Differential Revision: https://reviews.llvm.org/D92195
GitOrigin-RevId: 0caf736d7e1d16d1059553fc28dbac31f0b9f788
diff --git a/libomptarget/src/omptarget.cpp b/libomptarget/src/omptarget.cpp
index f273b66..6731439 100644
--- a/libomptarget/src/omptarget.cpp
+++ b/libomptarget/src/omptarget.cpp
@@ -269,10 +269,11 @@
MapperArgNames[I] = C.Name;
}
- int rc = target_data_function(
- loc, Device, MapperComponents.Components.size(), MapperArgsBase.data(),
- MapperArgs.data(), MapperArgSizes.data(), MapperArgTypes.data(),
- MapperArgNames.data(), /*arg_mappers*/ nullptr, AsyncInfo);
+ int rc = target_data_function(loc, Device, MapperComponents.Components.size(),
+ MapperArgsBase.data(), MapperArgs.data(),
+ MapperArgSizes.data(), MapperArgTypes.data(),
+ MapperArgNames.data(), /*arg_mappers*/ nullptr,
+ AsyncInfo, /*FromMapper=*/true);
return rc;
}
@@ -281,7 +282,8 @@
int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
void **args_base, void **args, int64_t *arg_sizes,
int64_t *arg_types, map_var_info_t *arg_names,
- void **arg_mappers, AsyncInfoTy &AsyncInfo) {
+ void **arg_mappers, AsyncInfoTy &AsyncInfo,
+ bool FromMapper) {
// process each input.
for (int32_t i = 0; i < arg_num; ++i) {
// Ignore private variables and arrays - there is no mapping for them.
@@ -379,7 +381,10 @@
Pointer_HstPtrBegin = HstPtrBase;
// modify current entry.
HstPtrBase = *(void **)HstPtrBase;
- UpdateRef = true; // subsequently update ref count of pointee
+ // No need to update pointee ref count for the first element of the
+ // subelement that comes from mapper.
+ UpdateRef =
+ (!FromMapper || i != 0); // subsequently update ref count of pointee
}
void *TgtPtrBegin = Device.getOrAllocTgtPtr(
@@ -483,7 +488,7 @@
int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
void **ArgBases, void **Args, int64_t *ArgSizes,
int64_t *ArgTypes, map_var_info_t *ArgNames,
- void **ArgMappers, AsyncInfoTy &AsyncInfo) {
+ void **ArgMappers, AsyncInfoTy &AsyncInfo, bool FromMapper) {
int Ret;
std::vector<DeallocTgtPtrInfo> DeallocTgtPtrs;
// process each input.
@@ -536,7 +541,8 @@
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);
+ (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ &&
+ (!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;
@@ -584,8 +590,13 @@
bool DelEntry = IsLast || ForceDelete;
- if ((ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) &&
- !(ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) {
+ // If the last element from the mapper (for end transfer args comes in
+ // reverse order), do not remove the partial entry, the parent struct still
+ // 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)) {
DelEntry = false; // protect parent struct from being deallocated
}
@@ -822,7 +833,7 @@
int targetDataUpdate(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
void **ArgsBase, void **Args, int64_t *ArgSizes,
int64_t *ArgTypes, map_var_info_t *ArgNames,
- void **ArgMappers, AsyncInfoTy &AsyncInfo) {
+ void **ArgMappers, AsyncInfoTy &AsyncInfo, bool) {
// process each input.
for (int32_t I = 0; I < ArgNum; ++I) {
if ((ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) ||
diff --git a/libomptarget/src/private.h b/libomptarget/src/private.h
index fc6997a..2eb0c81 100644
--- a/libomptarget/src/private.h
+++ b/libomptarget/src/private.h
@@ -23,17 +23,20 @@
extern int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
void **args_base, void **args, int64_t *arg_sizes,
int64_t *arg_types, map_var_info_t *arg_names,
- void **arg_mappers, AsyncInfoTy &AsyncInfo);
+ void **arg_mappers, AsyncInfoTy &AsyncInfo,
+ bool FromMapper = false);
extern int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
void **ArgBases, void **Args, int64_t *ArgSizes,
int64_t *ArgTypes, map_var_info_t *arg_names,
- void **ArgMappers, AsyncInfoTy &AsyncInfo);
+ void **ArgMappers, AsyncInfoTy &AsyncInfo,
+ bool FromMapper = false);
extern int targetDataUpdate(ident_t *loc, DeviceTy &Device, int32_t arg_num,
void **args_base, void **args, int64_t *arg_sizes,
int64_t *arg_types, map_var_info_t *arg_names,
- void **arg_mappers, AsyncInfoTy &AsyncInfo);
+ void **arg_mappers, AsyncInfoTy &AsyncInfo,
+ bool FromMapper = false);
extern int target(ident_t *loc, DeviceTy &Device, void *HostPtr, int32_t ArgNum,
void **ArgBases, void **Args, int64_t *ArgSizes,
@@ -76,7 +79,8 @@
// targetDataEnd and targetDataUpdate).
typedef int (*TargetDataFuncPtrTy)(ident_t *, DeviceTy &, int32_t, void **,
void **, int64_t *, int64_t *,
- map_var_info_t *, void **, AsyncInfoTy &);
+ map_var_info_t *, void **, AsyncInfoTy &,
+ bool);
// Implemented in libomp, they are called from within __tgt_* functions.
#ifdef __cplusplus
diff --git a/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp b/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp
new file mode 100644
index 0000000..ae2902a
--- /dev/null
+++ b/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp
@@ -0,0 +1,63 @@
+// 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;
+ D **spp = &sp;
+
+ printf("%d %d %d %4.5f %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.c.a,
+ spp[0][0].f.b[1], spp[0][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 : spp[0][0]) firstprivate(p)
+ {
+ printf("%d %d %d\n", spp[0][0].f.a, spp[0][0].f.c.a,
+ spp[0][0].f.b == reinterpret_cast<void *>(p) ? 1 : 0);
+ // CHECK: 222 777 0
+ spp[0][0].e = 333;
+ spp[0][0].f.a = 444;
+ spp[0][0].f.c.a = 555;
+ spp[0][0].f.b[1] = 40;
+ }
+ printf("%d %d %d %4.5f %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.c.a,
+ spp[0][0].f.b[1], spp[0][0].f.b == &x[0] ? 1 : 0);
+ // CHECK: 333 222 777 40.00000 1
+}
