diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp index 6c59bc1cf38a8..eb9e5abe00b46 100644 --- a/openmp/libomptarget/src/omptarget.cpp +++ b/openmp/libomptarget/src/omptarget.cpp @@ -585,11 +585,38 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, void **ArgMappers, AsyncInfoTy &AsyncInfo, bool FromMapper) { TIMESCOPE_WITH_IDENT(Loc); + + // Initialize new map type with old type: + SmallVector NewArgTypes(ArgTypes, ArgTypes + ArgNum); + + // Try to prevent mapping a struct multiple times in the same construct. + // Mapping the struct more than once will potentially overwrite previously + // mapped information. + for (int32_t I = 0; I < ArgNum; ++I) { + if (NewArgTypes[I] < 0) + continue; + if ((NewArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) || + (NewArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE)) + continue; + for (int32_t J = I + 1; J < ArgNum; ++J) { + if (Args[I] == ArgsBase[I] && Args[I] == Args[J] && + ArgsBase[I] == ArgsBase[J] && ArgSizes[I] == ArgSizes[J] && + ArgSizes[I] > 0 && NewArgTypes[J] >= 0) { + NewArgTypes[I] |= ArgTypes[J]; + NewArgTypes[J] = -1; + } + } + } + // process each input. for (int32_t I = 0; I < ArgNum; ++I) { + int64_t ArgType = NewArgTypes[I]; + if (ArgType < 0) + continue; + // Ignore private variables and arrays - there is no mapping for them. - if ((ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) || - (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE)) + if ((ArgType & OMP_TGT_MAPTYPE_LITERAL) || + (ArgType & OMP_TGT_MAPTYPE_PRIVATE)) continue; if (ArgMappers && ArgMappers[I]) { @@ -600,7 +627,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, map_var_info_t ArgName = (!ArgNames) ? nullptr : ArgNames[I]; int Rc = targetDataMapper(Loc, Device, ArgsBase[I], Args[I], ArgSizes[I], - ArgTypes[I], ArgName, ArgMappers[I], AsyncInfo, + ArgType, ArgName, ArgMappers[I], AsyncInfo, targetDataBegin); if (Rc != OFFLOAD_SUCCESS) { @@ -623,8 +650,8 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, // is a combined entry. int64_t TgtPadding = 0; const int NextI = I + 1; - if (getParentIndex(ArgTypes[I]) < 0 && NextI < ArgNum && - getParentIndex(ArgTypes[NextI]) == I) { + if (getParentIndex(ArgType) < 0 && NextI < ArgNum && + NewArgTypes[NextI] >= 0 && getParentIndex(NewArgTypes[NextI]) == I) { int64_t Alignment = getPartialStructRequiredAlignment(HstPtrBase); TgtPadding = (int64_t)HstPtrBegin % Alignment; if (TgtPadding) { @@ -638,23 +665,23 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, void *PointerHstPtrBegin, *PointerTgtPtrBegin; TargetPointerResultTy PointerTpr; bool IsHostPtr = false; - bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT; + bool IsImplicit = ArgType & OMP_TGT_MAPTYPE_IMPLICIT; // Force the creation of a device side copy of the data when: // a close map modifier was associated with a map that contained a to. - bool HasCloseModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_CLOSE; - bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT; - bool HasHoldModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_OMPX_HOLD; + bool HasCloseModifier = ArgType & OMP_TGT_MAPTYPE_CLOSE; + bool HasPresentModifier = ArgType & OMP_TGT_MAPTYPE_PRESENT; + bool HasHoldModifier = ArgType & OMP_TGT_MAPTYPE_OMPX_HOLD; // UpdateRef is based on MEMBER_OF instead of TARGET_PARAM because if we // have reached this point via __tgt_target_data_begin and not __tgt_target // 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 = - !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) && !(FromMapper && I == 0); + !(ArgType & OMP_TGT_MAPTYPE_MEMBER_OF) && !(FromMapper && I == 0); DeviceTy::HDTTMapAccessorTy HDTTMap = Device.HostDataToTargetMap.getExclusiveAccessor(); - if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) { + if (ArgType & OMP_TGT_MAPTYPE_PTR_AND_OBJ) { DP("Has a pointer entry: \n"); // Base is address of pointer. // @@ -696,8 +723,8 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, (!FromMapper || I != 0); // subsequently update ref count of pointee } - const bool HasFlagTo = ArgTypes[I] & OMP_TGT_MAPTYPE_TO; - const bool HasFlagAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS; + const bool HasFlagTo = ArgType & OMP_TGT_MAPTYPE_TO; + const bool HasFlagAlways = ArgType & OMP_TGT_MAPTYPE_ALWAYS; // Note that HDTTMap will be released in getTargetPointer. auto TPR = Device.getTargetPointer( HDTTMap, HstPtrBegin, HstPtrBase, TgtPadding, DataSize, HstPtrName, @@ -717,14 +744,14 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, " - is%s new\n", DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsNewEntry ? "" : " not")); - if (ArgTypes[I] & OMP_TGT_MAPTYPE_RETURN_PARAM) { + if (ArgType & OMP_TGT_MAPTYPE_RETURN_PARAM) { uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase; void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta); DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase)); ArgsBase[I] = TgtPtrBase; } - if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) { + if (ArgType & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) { uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase; void *ExpectedTgtPtrBase = (void *)((uint64_t)TgtPtrBegin - Delta); @@ -752,8 +779,8 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, } // Check if variable can be used on the device: - bool IsStructMember = ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF; - if (getInfoLevel() & OMP_INFOTYPE_EMPTY_MAPPING && ArgTypes[I] != 0 && + bool IsStructMember = ArgType & OMP_TGT_MAPTYPE_MEMBER_OF; + if (getInfoLevel() & OMP_INFOTYPE_EMPTY_MAPPING && ArgType != 0 && !IsStructMember && !IsImplicit && !TPR.isPresent() && !TPR.isContained() && !TPR.isHostPointer()) INFO(OMP_INFOTYPE_EMPTY_MAPPING, Device.DeviceID, diff --git a/openmp/libomptarget/test/offloading/struct_mapping_with_pointers.cpp b/openmp/libomptarget/test/offloading/struct_mapping_with_pointers.cpp new file mode 100644 index 0000000000000..4c9d5c7a23e81 --- /dev/null +++ b/openmp/libomptarget/test/offloading/struct_mapping_with_pointers.cpp @@ -0,0 +1,114 @@ +// clang-format off +// RUN: %libomptarget-compilexx-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic +// clang-format on + +#include +#include + +struct Descriptor { + int *datum; + long int x; + int *more_datum; + int xi; + int val_datum, val_more_datum; + long int arr[1][30]; + int val_arr; +}; + +int main() { + Descriptor dat = Descriptor(); + dat.datum = (int *)malloc(sizeof(int) * 10); + dat.more_datum = (int *)malloc(sizeof(int) * 20); + dat.xi = 3; + dat.arr[0][0] = 1; + + dat.datum[7] = 7; + dat.more_datum[17] = 17; + + /// The struct is mapped with type 0x0 when the pointer fields are mapped. + /// The struct is also map explicitely by the user. The second mapping by + /// the user must not overwrite the mapping set up for the pointer fields + /// when mapping the struct happens after the mapping of the pointers. + + // clang-format off + // CHECK: Libomptarget --> Entry 0: Base=[[DAT_HST_PTR_BASE:0x.*]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x0, Name=unknown + // CHECK: Libomptarget --> Entry 1: Base=[[DAT_HST_PTR_BASE]], Begin=[[DATUM_HST_PTR_BASE:0x.*]], Size=40, Type=0x1000000000011, Name=unknown + // CHECK: Libomptarget --> Entry 2: Base=[[MORE_DATUM_HST_PTR_BASE:0x.*]], Begin=[[MORE_DATUM_HST_PTR_BEGIN:0x.*]], Size=80, Type=0x1000000000011, Name=unknown + // CHECK: Libomptarget --> Entry 3: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x1000000000001, Name=unknown + // clang-format on + + /// The struct will be mapped in the same order as the above entries. + + /// First argument is the struct itself and it will be mapped once. + + // clang-format off + // CHECK: Libomptarget --> Looking up mapping(HstPtrBegin=[[DAT_HST_PTR_BASE]], Size=288)... + // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 288 with host pointer [[DAT_HST_PTR_BASE]]. + // CHECK: Libomptarget --> Creating new map entry with HstPtrBase=[[DAT_HST_PTR_BASE]], HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtAllocBegin=[[DAT_DEVICE_PTR_BASE:0x.*]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=1, HoldRefCount=0, Name=unknown + // CHECK: Libomptarget --> Moving 288 bytes (hst:[[DAT_HST_PTR_BASE]]) -> (tgt:[[DAT_DEVICE_PTR_BASE]]) + // clang-format on + + /// Second argument is dat.datum: + // clang-format off + // CHECK: Libomptarget --> Looking up mapping(HstPtrBegin=[[DATUM_HST_PTR_BASE]], Size=40)... + // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 40 with host pointer [[DATUM_HST_PTR_BASE]]. + // CHECK: Libomptarget --> Creating new map entry with HstPtrBase=[[DATUM_HST_PTR_BASE]], HstPtrBegin=[[DATUM_HST_PTR_BASE]], TgtAllocBegin=[[DATUM_DEVICE_PTR_BASE:0x.*]], TgtPtrBegin=[[DATUM_DEVICE_PTR_BASE]], Size=40, DynRefCount=1, HoldRefCount=0, Name=unknown + // CHECK: Libomptarget --> Moving 40 bytes (hst:[[DATUM_HST_PTR_BASE]]) -> (tgt:[[DATUM_DEVICE_PTR_BASE]]) + // clang-format on + + /// Third argument is dat.more_datum: + // clang-format off + // CHECK: Libomptarget --> Looking up mapping(HstPtrBegin=[[MORE_DATUM_HST_PTR_BEGIN]], Size=80)... + // CHECK: PluginInterface --> MemoryManagerTy::allocate: size 80 with host pointer [[MORE_DATUM_HST_PTR_BEGIN]]. + // CHECK: Libomptarget --> Creating new map entry with HstPtrBase=[[MORE_DATUM_HST_PTR_BEGIN]], HstPtrBegin=[[MORE_DATUM_HST_PTR_BEGIN]], TgtAllocBegin=[[MORE_DATUM_DEVICE_PTR_BEGIN:0x.*]], TgtPtrBegin=[[MORE_DATUM_DEVICE_PTR_BEGIN]], Size=80, DynRefCount=1, HoldRefCount=0, Name=unknown + // CHECK: Libomptarget --> Moving 80 bytes (hst:[[MORE_DATUM_HST_PTR_BEGIN]]) -> (tgt:[[MORE_DATUM_DEVICE_PTR_BEGIN]]) + // clang-format on + +#pragma omp target enter data map(to : dat.datum[ : 10]) \ + map(to : dat.more_datum[ : 20]) map(to : dat) + + /// Checks induced by having a target region: + // clang-format off + // CHECK: Libomptarget --> Entry 0: Base=[[DAT_HST_PTR_BASE]], Begin=[[DAT_HST_PTR_BASE]], Size=288, Type=0x223, Name=unknown + // CHECK: Libomptarget --> Mapping exists (implicit) with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=2 (incremented), HoldRefCount=0, Name=unknown + // CHECK: Libomptarget --> Obtained target argument [[DAT_DEVICE_PTR_BASE]] from host pointer [[DAT_HST_PTR_BASE]] + // clang-format on + +#pragma omp target + { + dat.xi = 4; + dat.datum[7]++; + dat.more_datum[17]++; + dat.val_datum = dat.datum[7]; + dat.val_more_datum = dat.more_datum[17]; + dat.datum[dat.arr[0][0]] = dat.xi; + dat.val_arr = dat.datum[dat.arr[0][0]]; + } + + /// Post-target region checks: + // clang-format off + // CHECK: Libomptarget --> Mapping exists with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=1 (decremented), HoldRefCount=0 + // clang-format on + +#pragma omp target exit data map(from : dat) + + /// Target data end checks: + // clang-format off + // CHECK: Libomptarget --> Mapping exists with HstPtrBegin=[[DAT_HST_PTR_BASE]], TgtPtrBegin=[[DAT_DEVICE_PTR_BASE]], Size=288, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0 + // CHECK: Libomptarget --> Moving 288 bytes (tgt:[[DAT_DEVICE_PTR_BASE]]) -> (hst:[[DAT_HST_PTR_BASE]]) + // clang-format on + + // CHECK: dat.xi = 4 + // CHECK: dat.val_datum = 8 + // CHECK: dat.val_more_datum = 18 + // CHECK: dat.datum[dat.arr[0][0]] = 0 + // CHECK: dat.val_arr = 4 + + printf("dat.xi = %d\n", dat.xi); + printf("dat.val_datum = %d\n", dat.val_datum); + printf("dat.val_more_datum = %d\n", dat.val_more_datum); + printf("dat.datum[dat.arr[0][0]] = %d\n", dat.datum[dat.arr[0][0]]); + printf("dat.val_arr = %d\n", dat.val_arr); + + return 0; +}