From 6e42d525cf807d35a9522795d681d90eb9b30eb7 Mon Sep 17 00:00:00 2001 From: pradt2 <12902844+pradt2@users.noreply.github.com> Date: Tue, 28 Jan 2025 01:00:58 -0800 Subject: [PATCH 1/7] [Offload]: Skip copying of unused kernel-mapped data --- offload/libomptarget/omptarget.cpp | 38 ++++++++++++++++++++++++++++-- 1 file changed, 36 insertions(+), 2 deletions(-) diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp index 5b25d955dd320..729669caa07bc 100644 --- a/offload/libomptarget/omptarget.cpp +++ b/offload/libomptarget/omptarget.cpp @@ -1197,6 +1197,35 @@ class PrivateArgumentManagerTy { } }; +static std::unique_ptr maskIgnorableMappings(int64_t DeviceId, int32_t ArgNum, int64_t *ArgTypes, + int64_t *ArgSizes, map_var_info_t *ArgNames) { + std::unique_ptr ArgTypesOverride = std::make_unique(ArgNum); + + for (int32_t I = 0; I < ArgNum; ++I) { + bool IsTargetParam = ArgTypes[I] & OMP_TGT_MAPTYPE_TARGET_PARAM; + + bool IsMapTo = ArgTypes[I] & OMP_TGT_MAPTYPE_TO; + if (IsTargetParam || !IsMapTo) { + ArgTypesOverride[I] = ArgTypes[I]; + continue; + } + + bool IsMapFrom = ArgTypes[I] & OMP_TGT_MAPTYPE_FROM; + const char *Type = IsMapFrom ? "tofrom" : "to"; + + // Optimisation: A 'to' or 'tofrom' mapping is not + // used by the kernel. Change its type such that + // no new mapping is created, but any existing + // mapping has its counter decremented. + INFO(OMP_INFOTYPE_ALL, DeviceId, "%s(%s)[%" PRId64 "] %s\n", Type, + getNameFromMapping(ArgNames[I]).c_str(), ArgSizes[I], "is not used and will not be copied"); + + ArgTypesOverride[I] = ArgTypes[I] & ~(OMP_TGT_MAPTYPE_TO | OMP_TGT_MAPTYPE_FROM); + } + + return ArgTypesOverride; +} + /// Process data before launching the kernel, including calling targetDataBegin /// to map and transfer data to target device, transferring (first-)private /// variables. @@ -1417,11 +1446,16 @@ int target(ident_t *Loc, DeviceTy &Device, void *HostPtr, int NumClangLaunchArgs = KernelArgs.NumArgs; int Ret = OFFLOAD_SUCCESS; + + std::unique_ptr ArgTypesOverride = + maskIgnorableMappings(DeviceId, NumClangLaunchArgs, KernelArgs.ArgTypes, + KernelArgs.ArgSizes, KernelArgs.ArgNames); + if (NumClangLaunchArgs) { // Process data, such as data mapping, before launching the kernel Ret = processDataBefore(Loc, DeviceId, HostPtr, NumClangLaunchArgs, KernelArgs.ArgBasePtrs, KernelArgs.ArgPtrs, - KernelArgs.ArgSizes, KernelArgs.ArgTypes, + KernelArgs.ArgSizes, ArgTypesOverride.get(), KernelArgs.ArgNames, KernelArgs.ArgMappers, TgtArgs, TgtOffsets, PrivateArgumentManager, AsyncInfo); if (Ret != OFFLOAD_SUCCESS) { @@ -1473,7 +1507,7 @@ int target(ident_t *Loc, DeviceTy &Device, void *HostPtr, // variables Ret = processDataAfter(Loc, DeviceId, HostPtr, NumClangLaunchArgs, KernelArgs.ArgBasePtrs, KernelArgs.ArgPtrs, - KernelArgs.ArgSizes, KernelArgs.ArgTypes, + KernelArgs.ArgSizes, ArgTypesOverride.get(), KernelArgs.ArgNames, KernelArgs.ArgMappers, PrivateArgumentManager, AsyncInfo); if (Ret != OFFLOAD_SUCCESS) { From 12561c445d04f13a769f0f8a2c0c01934a0a8062 Mon Sep 17 00:00:00 2001 From: pradt2 <12902844+pradt2@users.noreply.github.com> Date: Wed, 5 Feb 2025 19:17:11 -0800 Subject: [PATCH 2/7] [Offload]: Skip copying of unused kernel-mapped data --- offload/include/OpenMP/Mapping.h | 7 +++ offload/include/Shared/Debug.h | 2 + offload/libomptarget/omptarget.cpp | 66 +++++++++++++++++------ offload/test/mapping/skip_transfers.cpp | 72 +++++++++++++++++++++++++ 4 files changed, 132 insertions(+), 15 deletions(-) create mode 100644 offload/test/mapping/skip_transfers.cpp diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h index b9f5c16582931..1595e0671419b 100644 --- a/offload/include/OpenMP/Mapping.h +++ b/offload/include/OpenMP/Mapping.h @@ -386,6 +386,13 @@ struct LookupResult { LookupResult() : Flags({0, 0, 0}), TPR() {} TargetPointerResultTy TPR; + + bool isEmpty() const { + bool IsEmpty = Flags.IsContained == 0 + & Flags.ExtendsBefore == 0 + & Flags.ExtendsAfter == 0; + return IsEmpty; + } }; // This structure stores information of a mapped memory region. diff --git a/offload/include/Shared/Debug.h b/offload/include/Shared/Debug.h index 7c3db8dbf119f..ec8f4bbedaaeb 100644 --- a/offload/include/Shared/Debug.h +++ b/offload/include/Shared/Debug.h @@ -58,6 +58,8 @@ enum OpenMPInfoType : uint32_t { OMP_INFOTYPE_DATA_TRANSFER = 0x0020, // Print whenever data does not have a viable device counterpart. OMP_INFOTYPE_EMPTY_MAPPING = 0x0040, + // Print whenever data does not need to be transferred + OMP_INFOTYPE_REDUNDANT_TRANSFER = 0x0080, // Enable every flag. OMP_INFOTYPE_ALL = 0xffffffff, }; diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp index 729669caa07bc..6fe3def424cba 100644 --- a/offload/libomptarget/omptarget.cpp +++ b/offload/libomptarget/omptarget.cpp @@ -1197,30 +1197,66 @@ class PrivateArgumentManagerTy { } }; -static std::unique_ptr maskIgnorableMappings(int64_t DeviceId, int32_t ArgNum, int64_t *ArgTypes, - int64_t *ArgSizes, map_var_info_t *ArgNames) { +/// Try to find redundant mappings associated with a kernel launch, +/// and provide a masked version of the kernel argument types that +/// avoid redundant to data transfers between the host and device. +static std::unique_ptr maskRedundantTransfers(DeviceTy &Device, int32_t ArgNum, + int64_t *ArgTypes, int64_t *ArgSizes, + map_var_info_t *ArgNames, void **ArgPtrs, + void **ArgMappers) { std::unique_ptr ArgTypesOverride = std::make_unique(ArgNum); + MappingInfoTy &MappingInfo = Device.getMappingInfo(); + MappingInfoTy::HDTTMapAccessorTy HDTTMap = MappingInfo + .HostDataToTargetMap.getExclusiveAccessor(); + + int64_t UnusedArgs = 0; + for (int32_t I = 0; I < ArgNum; ++I) { - bool IsTargetParam = ArgTypes[I] & OMP_TGT_MAPTYPE_TARGET_PARAM; + tgt_map_type ArgType = (tgt_map_type) ArgTypes[I]; + + // Check for unused implicit mappings + bool IsArgUnused = ArgType == OMP_TGT_MAPTYPE_NONE; + + // Check for unused `map(buf[0:size])` mappings + IsArgUnused |= ArgType == OMP_TGT_MAPTYPE_FROM + || ArgType == OMP_TGT_MAPTYPE_TO + || ArgType == (OMP_TGT_MAPTYPE_FROM | OMP_TGT_MAPTYPE_TO); + + // Check for unused `map(wrapper.buf[0:size])` mappings + IsArgUnused |= UnusedArgs == ArgNum - 1 && ArgType & OMP_TGT_MAPTYPE_MEMBER_OF + && ((ArgType & ~OMP_TGT_MAPTYPE_MEMBER_OF) == OMP_TGT_MAPTYPE_PTR_AND_OBJ + || (ArgType & ~OMP_TGT_MAPTYPE_MEMBER_OF) == (OMP_TGT_MAPTYPE_PTR_AND_OBJ | OMP_TGT_MAPTYPE_TO) + || (ArgType & ~OMP_TGT_MAPTYPE_MEMBER_OF) == (OMP_TGT_MAPTYPE_PTR_AND_OBJ | OMP_TGT_MAPTYPE_FROM | OMP_TGT_MAPTYPE_TO)); - bool IsMapTo = ArgTypes[I] & OMP_TGT_MAPTYPE_TO; - if (IsTargetParam || !IsMapTo) { + bool IsExistingMapping = !MappingInfo.lookupMapping(HDTTMap, ArgPtrs[I], ArgSizes[I]).isEmpty(); + + bool IsCustomMapped = ArgMappers && ArgMappers[I]; + + if (IsExistingMapping | IsCustomMapped | !IsArgUnused) { ArgTypesOverride[I] = ArgTypes[I]; continue; } - bool IsMapFrom = ArgTypes[I] & OMP_TGT_MAPTYPE_FROM; - const char *Type = IsMapFrom ? "tofrom" : "to"; + const std::string Name = ArgNames && ArgNames[I] ? + getNameFromMapping(ArgNames[I]) : std::string("unknown"); + + bool IsArgFrom = ArgType & OMP_TGT_MAPTYPE_FROM; + bool IsArgTo = ArgType & OMP_TGT_MAPTYPE_TO; + + const char *Type = IsArgFrom && IsArgTo ? "tofrom" + : IsArgFrom ? "from" + : IsArgTo ? "to" + : "unknown"; - // Optimisation: A 'to' or 'tofrom' mapping is not - // used by the kernel. Change its type such that - // no new mapping is created, but any existing - // mapping has its counter decremented. - INFO(OMP_INFOTYPE_ALL, DeviceId, "%s(%s)[%" PRId64 "] %s\n", Type, - getNameFromMapping(ArgNames[I]).c_str(), ArgSizes[I], "is not used and will not be copied"); + // Optimisation: + // A new mapping is not used by the kernel. + // Change the type such that no data is transferred to and/or from the device. + INFO(OMP_INFOTYPE_REDUNDANT_TRANSFER, Device.DeviceID, "%s(%s)[%" PRId64 "] %s\n", Type, + Name.c_str(), ArgSizes[I], "is not used and will not be copied"); ArgTypesOverride[I] = ArgTypes[I] & ~(OMP_TGT_MAPTYPE_TO | OMP_TGT_MAPTYPE_FROM); + UnusedArgs++; } return ArgTypesOverride; @@ -1448,8 +1484,8 @@ int target(ident_t *Loc, DeviceTy &Device, void *HostPtr, int Ret = OFFLOAD_SUCCESS; std::unique_ptr ArgTypesOverride = - maskIgnorableMappings(DeviceId, NumClangLaunchArgs, KernelArgs.ArgTypes, - KernelArgs.ArgSizes, KernelArgs.ArgNames); + maskRedundantTransfers(Device, NumClangLaunchArgs, KernelArgs.ArgTypes, + KernelArgs.ArgSizes, KernelArgs.ArgNames, KernelArgs.ArgPtrs, KernelArgs.ArgMappers); if (NumClangLaunchArgs) { // Process data, such as data mapping, before launching the kernel diff --git a/offload/test/mapping/skip_transfers.cpp b/offload/test/mapping/skip_transfers.cpp new file mode 100644 index 0000000000000..ff0459f01f935 --- /dev/null +++ b/offload/test/mapping/skip_transfers.cpp @@ -0,0 +1,72 @@ +// clang-format off +// RUN: %libomptarget-compilexx-generic +// RUN: env LIBOMPTARGET_INFO=128 %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic + +// REQUIRES: gpu +// clang-format on + +int main() { + float DataStack = 0; + + // CHECK-NOT: omptarget device 0 info: from(unknown)[4] is not used and will not be copied + #pragma omp target map(from: DataStack) + { + DataStack = 1; + } + + // CHECK-NOT: omptarget device 0 info: to(unknown)[4] is not used and will not be copied + #pragma omp target map(always to: DataStack) + ; + + // CHECK: omptarget device 0 info: tofrom(unknown)[4] is not used and will not be copied + #pragma omp target map(tofrom: DataStack) + ; + + int Size = 16; + double *Data = new double[Size]; + + // CHECK-NOT: omptarget device 0 info: tofrom(unknown)[8] is not used and will not be copied + #pragma omp target map(tofrom: Data[0:1]) + { + Data[0] = 1; + } + + // CHECK-NOT: omptarget device 0 info: tofrom(unknown)[16] is not used and will not be copied + #pragma omp target map(always tofrom: Data[0:2]) + ; + + // CHECK: omptarget device 0 info: from(unknown)[24] is not used and will not be copied + #pragma omp target map(from: Data[0:3]) + ; + + // CHECK: omptarget device 0 info: to(unknown)[24] is not used and will not be copied + #pragma omp target map(to: Data[0:3]) + ; + + // CHECK: omptarget device 0 info: tofrom(unknown)[32] is not used and will not be copied + #pragma omp target map(tofrom: Data[0:4]) + ; + + // CHECK-NOT: omptarget device 0 info: to(unknown)[40] is not used and will not be copied + #pragma omp target map(to: Data[0:5]) + { + #pragma omp teams + Data[0] = 1; + } + + struct { + double *Data; + } Wrapper { .Data = Data }; + + // CHECK-NOT: omptarget device 0 info: tofrom(unknown)[48] is not used and will not be copied + #pragma omp target map(tofrom: Wrapper.Data[0:6]) + { + Wrapper.Data[0] = 1; + } + + // CHECK: omptarget device 0 info: unknown(unknown)[8] is not used and will not be copied + // CHECK: omptarget device 0 info: tofrom(unknown)[56] is not used and will not be copied + #pragma omp target map(tofrom: Wrapper.Data[0:7]) + ; +} From 02f78ec9831a6ec6063ffb912a750eb78acdc7cc Mon Sep 17 00:00:00 2001 From: pradt2 <12902844+pradt2@users.noreply.github.com> Date: Thu, 13 Feb 2025 14:36:21 -0800 Subject: [PATCH 3/7] [Offload]: Skip copying of unused kernel-mapped data --- offload/include/OpenMP/Mapping.h | 4 +- offload/libomptarget/omptarget.cpp | 69 +++++++++++++++++++----------- 2 files changed, 46 insertions(+), 27 deletions(-) diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h index 1595e0671419b..2b20be28391b9 100644 --- a/offload/include/OpenMP/Mapping.h +++ b/offload/include/OpenMP/Mapping.h @@ -389,8 +389,8 @@ struct LookupResult { bool isEmpty() const { bool IsEmpty = Flags.IsContained == 0 - & Flags.ExtendsBefore == 0 - & Flags.ExtendsAfter == 0; + && Flags.ExtendsBefore == 0 + && Flags.ExtendsAfter == 0; return IsEmpty; } }; diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp index 6fe3def424cba..8a8113c24e5b8 100644 --- a/offload/libomptarget/omptarget.cpp +++ b/offload/libomptarget/omptarget.cpp @@ -1197,66 +1197,85 @@ class PrivateArgumentManagerTy { } }; +/// Try to determine if kernel argument is unused. This method +/// takes a conservative approach, i.e. it may return false +/// negatives but it should never return a false positive. +static bool isArgUnused(tgt_map_type ArgType) { + bool IsArgUnused = ArgType == OMP_TGT_MAPTYPE_NONE + || ArgType == OMP_TGT_MAPTYPE_FROM + || ArgType == OMP_TGT_MAPTYPE_TO + || ArgType == (OMP_TGT_MAPTYPE_FROM | OMP_TGT_MAPTYPE_TO); + return IsArgUnused; +} + /// Try to find redundant mappings associated with a kernel launch, /// and provide a masked version of the kernel argument types that -/// avoid redundant to data transfers between the host and device. +/// avoid redundant data transfers between the host and the device. static std::unique_ptr maskRedundantTransfers(DeviceTy &Device, int32_t ArgNum, int64_t *ArgTypes, int64_t *ArgSizes, map_var_info_t *ArgNames, void **ArgPtrs, void **ArgMappers) { std::unique_ptr ArgTypesOverride = std::make_unique(ArgNum); - MappingInfoTy &MappingInfo = Device.getMappingInfo(); - MappingInfoTy::HDTTMapAccessorTy HDTTMap = MappingInfo - .HostDataToTargetMap.getExclusiveAccessor(); - - int64_t UnusedArgs = 0; + bool AllArgsUnused = true; for (int32_t I = 0; I < ArgNum; ++I) { + bool IsCustomMapped = ArgMappers && ArgMappers[I]; + + if (IsCustomMapped) { + ArgTypesOverride[I] = ArgTypes[I]; + AllArgsUnused = false; + continue; + } + tgt_map_type ArgType = (tgt_map_type) ArgTypes[I]; - // Check for unused implicit mappings - bool IsArgUnused = ArgType == OMP_TGT_MAPTYPE_NONE; + bool IsArgUnused = true; // Check for unused `map(buf[0:size])` mappings - IsArgUnused |= ArgType == OMP_TGT_MAPTYPE_FROM - || ArgType == OMP_TGT_MAPTYPE_TO - || ArgType == (OMP_TGT_MAPTYPE_FROM | OMP_TGT_MAPTYPE_TO); + IsArgUnused |= isArgUnused(ArgType); + + bool IsArgMemberPtr = ArgType & OMP_TGT_MAPTYPE_MEMBER_OF + && ArgType & OMP_TGT_MAPTYPE_PTR_AND_OBJ; + + tgt_map_type ArgTypeMemberPtrMasked = (tgt_map_type) (ArgType & ~(OMP_TGT_MAPTYPE_MEMBER_OF | OMP_TGT_MAPTYPE_PTR_AND_OBJ)); // Check for unused `map(wrapper.buf[0:size])` mappings - IsArgUnused |= UnusedArgs == ArgNum - 1 && ArgType & OMP_TGT_MAPTYPE_MEMBER_OF - && ((ArgType & ~OMP_TGT_MAPTYPE_MEMBER_OF) == OMP_TGT_MAPTYPE_PTR_AND_OBJ - || (ArgType & ~OMP_TGT_MAPTYPE_MEMBER_OF) == (OMP_TGT_MAPTYPE_PTR_AND_OBJ | OMP_TGT_MAPTYPE_TO) - || (ArgType & ~OMP_TGT_MAPTYPE_MEMBER_OF) == (OMP_TGT_MAPTYPE_PTR_AND_OBJ | OMP_TGT_MAPTYPE_FROM | OMP_TGT_MAPTYPE_TO)); + IsArgUnused |= AllArgsUnused && IsArgMemberPtr && isArgUnused(ArgTypeMemberPtrMasked); - bool IsExistingMapping = !MappingInfo.lookupMapping(HDTTMap, ArgPtrs[I], ArgSizes[I]).isEmpty(); + if (!IsArgUnused) { + ArgTypesOverride[I] = ArgTypes[I]; + AllArgsUnused = false; + continue; + } - bool IsCustomMapped = ArgMappers && ArgMappers[I]; + MappingInfoTy &MappingInfo = Device.getMappingInfo(); + MappingInfoTy::HDTTMapAccessorTy HDTTMap = MappingInfo + .HostDataToTargetMap.getExclusiveAccessor(); + + bool IsExistingMapping = !MappingInfo.lookupMapping(HDTTMap, ArgPtrs[I], ArgSizes[I]).isEmpty(); - if (IsExistingMapping | IsCustomMapped | !IsArgUnused) { + if (IsExistingMapping) { ArgTypesOverride[I] = ArgTypes[I]; + AllArgsUnused = false; continue; } - const std::string Name = ArgNames && ArgNames[I] ? + [[maybe_unused]] const std::string Name = ArgNames && ArgNames[I] ? getNameFromMapping(ArgNames[I]) : std::string("unknown"); bool IsArgFrom = ArgType & OMP_TGT_MAPTYPE_FROM; bool IsArgTo = ArgType & OMP_TGT_MAPTYPE_TO; - const char *Type = IsArgFrom && IsArgTo ? "tofrom" + [[maybe_unused]] const char *Type = IsArgFrom && IsArgTo ? "tofrom" : IsArgFrom ? "from" : IsArgTo ? "to" : "unknown"; - // Optimisation: - // A new mapping is not used by the kernel. - // Change the type such that no data is transferred to and/or from the device. INFO(OMP_INFOTYPE_REDUNDANT_TRANSFER, Device.DeviceID, "%s(%s)[%" PRId64 "] %s\n", Type, Name.c_str(), ArgSizes[I], "is not used and will not be copied"); - ArgTypesOverride[I] = ArgTypes[I] & ~(OMP_TGT_MAPTYPE_TO | OMP_TGT_MAPTYPE_FROM); - UnusedArgs++; + ArgTypesOverride[I] = ArgType & ~(OMP_TGT_MAPTYPE_TO | OMP_TGT_MAPTYPE_FROM); } return ArgTypesOverride; From 153b3f9798718ab45aca8723cfdf83e5a8fdf129 Mon Sep 17 00:00:00 2001 From: pradt2 <12902844+pradt2@users.noreply.github.com> Date: Thu, 13 Feb 2025 14:47:41 -0800 Subject: [PATCH 4/7] [Offload]: Skip copying of unused kernel-mapped data --- offload/include/OpenMP/Mapping.h | 7 ++-- offload/libomptarget/omptarget.cpp | 67 +++++++++++++++++------------- 2 files changed, 40 insertions(+), 34 deletions(-) diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h index 2b20be28391b9..77fc90711abce 100644 --- a/offload/include/OpenMP/Mapping.h +++ b/offload/include/OpenMP/Mapping.h @@ -388,10 +388,9 @@ struct LookupResult { TargetPointerResultTy TPR; bool isEmpty() const { - bool IsEmpty = Flags.IsContained == 0 - && Flags.ExtendsBefore == 0 - && Flags.ExtendsAfter == 0; - return IsEmpty; + bool IsEmpty = Flags.IsContained == 0 && Flags.ExtendsBefore == 0 && + Flags.ExtendsAfter == 0; + return IsEmpty; } }; diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp index 8a8113c24e5b8..e48b0ffb7abc6 100644 --- a/offload/libomptarget/omptarget.cpp +++ b/offload/libomptarget/omptarget.cpp @@ -1201,11 +1201,11 @@ class PrivateArgumentManagerTy { /// takes a conservative approach, i.e. it may return false /// negatives but it should never return a false positive. static bool isArgUnused(tgt_map_type ArgType) { - bool IsArgUnused = ArgType == OMP_TGT_MAPTYPE_NONE - || ArgType == OMP_TGT_MAPTYPE_FROM - || ArgType == OMP_TGT_MAPTYPE_TO - || ArgType == (OMP_TGT_MAPTYPE_FROM | OMP_TGT_MAPTYPE_TO); - return IsArgUnused; + bool IsArgUnused = ArgType == OMP_TGT_MAPTYPE_NONE || + ArgType == OMP_TGT_MAPTYPE_FROM || + ArgType == OMP_TGT_MAPTYPE_TO || + ArgType == (OMP_TGT_MAPTYPE_FROM | OMP_TGT_MAPTYPE_TO); + return IsArgUnused; } /// Try to find redundant mappings associated with a kernel launch, @@ -1220,13 +1220,13 @@ static std::unique_ptr maskRedundantTransfers(DeviceTy &Device, int32 bool AllArgsUnused = true; for (int32_t I = 0; I < ArgNum; ++I) { - bool IsCustomMapped = ArgMappers && ArgMappers[I]; + bool IsCustomMapped = ArgMappers && ArgMappers[I]; - if (IsCustomMapped) { - ArgTypesOverride[I] = ArgTypes[I]; - AllArgsUnused = false; - continue; - } + if (IsCustomMapped) { + ArgTypesOverride[I] = ArgTypes[I]; + AllArgsUnused = false; + continue; + } tgt_map_type ArgType = (tgt_map_type) ArgTypes[I]; @@ -1235,47 +1235,54 @@ static std::unique_ptr maskRedundantTransfers(DeviceTy &Device, int32 // Check for unused `map(buf[0:size])` mappings IsArgUnused |= isArgUnused(ArgType); - bool IsArgMemberPtr = ArgType & OMP_TGT_MAPTYPE_MEMBER_OF - && ArgType & OMP_TGT_MAPTYPE_PTR_AND_OBJ; + bool IsArgMemberPtr = ArgType & OMP_TGT_MAPTYPE_MEMBER_OF && + ArgType & OMP_TGT_MAPTYPE_PTR_AND_OBJ; - tgt_map_type ArgTypeMemberPtrMasked = (tgt_map_type) (ArgType & ~(OMP_TGT_MAPTYPE_MEMBER_OF | OMP_TGT_MAPTYPE_PTR_AND_OBJ)); + tgt_map_type ArgTypeMemberPtrMasked = + (tgt_map_type)(ArgType & ~(OMP_TGT_MAPTYPE_MEMBER_OF | + OMP_TGT_MAPTYPE_PTR_AND_OBJ)); // Check for unused `map(wrapper.buf[0:size])` mappings - IsArgUnused |= AllArgsUnused && IsArgMemberPtr && isArgUnused(ArgTypeMemberPtrMasked); + IsArgUnused |= AllArgsUnused && IsArgMemberPtr && + isArgUnused(ArgTypeMemberPtrMasked); if (!IsArgUnused) { - ArgTypesOverride[I] = ArgTypes[I]; - AllArgsUnused = false; - continue; + ArgTypesOverride[I] = ArgTypes[I]; + AllArgsUnused = false; + continue; } MappingInfoTy &MappingInfo = Device.getMappingInfo(); - MappingInfoTy::HDTTMapAccessorTy HDTTMap = MappingInfo - .HostDataToTargetMap.getExclusiveAccessor(); + MappingInfoTy::HDTTMapAccessorTy HDTTMap = + MappingInfo.HostDataToTargetMap.getExclusiveAccessor(); - bool IsExistingMapping = !MappingInfo.lookupMapping(HDTTMap, ArgPtrs[I], ArgSizes[I]).isEmpty(); + bool IsExistingMapping = + !MappingInfo.lookupMapping(HDTTMap, ArgPtrs[I], ArgSizes[I]) + .isEmpty(); if (IsExistingMapping) { - ArgTypesOverride[I] = ArgTypes[I]; - AllArgsUnused = false; - continue; + ArgTypesOverride[I] = ArgTypes[I]; + AllArgsUnused = false; + continue; } - [[maybe_unused]] const std::string Name = ArgNames && ArgNames[I] ? - getNameFromMapping(ArgNames[I]) : std::string("unknown"); + [[maybe_unused]] const std::string Name = + ArgNames && ArgNames[I] ? getNameFromMapping(ArgNames[I]) + : std::string("unknown"); bool IsArgFrom = ArgType & OMP_TGT_MAPTYPE_FROM; bool IsArgTo = ArgType & OMP_TGT_MAPTYPE_TO; [[maybe_unused]] const char *Type = IsArgFrom && IsArgTo ? "tofrom" - : IsArgFrom ? "from" - : IsArgTo ? "to" - : "unknown"; + : IsArgFrom ? "from" + : IsArgTo ? "to" + : "unknown"; INFO(OMP_INFOTYPE_REDUNDANT_TRANSFER, Device.DeviceID, "%s(%s)[%" PRId64 "] %s\n", Type, Name.c_str(), ArgSizes[I], "is not used and will not be copied"); - ArgTypesOverride[I] = ArgType & ~(OMP_TGT_MAPTYPE_TO | OMP_TGT_MAPTYPE_FROM); + ArgTypesOverride[I] = + ArgType & ~(OMP_TGT_MAPTYPE_TO | OMP_TGT_MAPTYPE_FROM); } return ArgTypesOverride; From de77e4bdb3b09edcb11490bfe1d5559d3fadce44 Mon Sep 17 00:00:00 2001 From: pradt2 <12902844+pradt2@users.noreply.github.com> Date: Thu, 13 Feb 2025 15:06:52 -0800 Subject: [PATCH 5/7] [Offload]: Skip copying of unused kernel-mapped data --- offload/libomptarget/omptarget.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp index e48b0ffb7abc6..4a09d1b09ef27 100644 --- a/offload/libomptarget/omptarget.cpp +++ b/offload/libomptarget/omptarget.cpp @@ -1230,7 +1230,7 @@ static std::unique_ptr maskRedundantTransfers(DeviceTy &Device, int32 tgt_map_type ArgType = (tgt_map_type) ArgTypes[I]; - bool IsArgUnused = true; + bool IsArgUnused = false; // Check for unused `map(buf[0:size])` mappings IsArgUnused |= isArgUnused(ArgType); From 050c5e3333e430d4833b86feb18bb6fbf28ceb93 Mon Sep 17 00:00:00 2001 From: pradt2 <12902844+pradt2@users.noreply.github.com> Date: Thu, 13 Feb 2025 15:33:37 -0800 Subject: [PATCH 6/7] [Offload]: Skip copying of unused kernel-mapped data --- offload/test/mapping/skip_transfers.cpp | 17 ++++++++++------- 1 file changed, 10 insertions(+), 7 deletions(-) diff --git a/offload/test/mapping/skip_transfers.cpp b/offload/test/mapping/skip_transfers.cpp index ff0459f01f935..93b615a0ab3d6 100644 --- a/offload/test/mapping/skip_transfers.cpp +++ b/offload/test/mapping/skip_transfers.cpp @@ -1,6 +1,6 @@ // clang-format off // RUN: %libomptarget-compilexx-generic -// RUN: env LIBOMPTARGET_INFO=128 %libomptarget-run-generic 2>&1 \ +// RUN: env LIBOMPTARGET_INFO=160 %libomptarget-run-generic 2>&1 \ // RUN: | %fcheck-generic // REQUIRES: gpu @@ -9,13 +9,13 @@ int main() { float DataStack = 0; - // CHECK-NOT: omptarget device 0 info: from(unknown)[4] is not used and will not be copied + // CHECK: omptarget device 0 info: Copying data from device to host, TgtPtr=0x{{.*}}, HstPtr=0x{{.*}}, Size=4, Name=unknown #pragma omp target map(from: DataStack) { DataStack = 1; } - // CHECK-NOT: omptarget device 0 info: to(unknown)[4] is not used and will not be copied + // CHECK: omptarget device 0 info: Copying data from host to device, HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=4, Name=unknown #pragma omp target map(always to: DataStack) ; @@ -26,13 +26,15 @@ int main() { int Size = 16; double *Data = new double[Size]; - // CHECK-NOT: omptarget device 0 info: tofrom(unknown)[8] is not used and will not be copied + // CHECK: omptarget device 0 info: Copying data from host to device, HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=8, Name=unknown + // CHECK: omptarget device 0 info: Copying data from device to host, TgtPtr=0x{{.*}}, HstPtr=0x{{.*}}, Size=8, Name=unknown #pragma omp target map(tofrom: Data[0:1]) { Data[0] = 1; } - // CHECK-NOT: omptarget device 0 info: tofrom(unknown)[16] is not used and will not be copied + // CHECK: omptarget device 0 info: Copying data from host to device, HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=16, Name=unknown + // CHECK: omptarget device 0 info: Copying data from device to host, TgtPtr=0x{{.*}}, HstPtr=0x{{.*}}, Size=16, Name=unknown #pragma omp target map(always tofrom: Data[0:2]) ; @@ -48,7 +50,7 @@ int main() { #pragma omp target map(tofrom: Data[0:4]) ; - // CHECK-NOT: omptarget device 0 info: to(unknown)[40] is not used and will not be copied + // CHECK: omptarget device 0 info: Copying data from host to device, HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=40, Name=unknown #pragma omp target map(to: Data[0:5]) { #pragma omp teams @@ -59,7 +61,8 @@ int main() { double *Data; } Wrapper { .Data = Data }; - // CHECK-NOT: omptarget device 0 info: tofrom(unknown)[48] is not used and will not be copied + // CHECK: omptarget device 0 info: Copying data from host to device, HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=48, Name=unknown + // CHECK: omptarget device 0 info: Copying data from device to host, TgtPtr=0x{{.*}}, HstPtr=0x{{.*}}, Size=48, Name=unknown #pragma omp target map(tofrom: Wrapper.Data[0:6]) { Wrapper.Data[0] = 1; From 99bd34b8297164f8d9e73372c3e3e573ad506eeb Mon Sep 17 00:00:00 2001 From: pradt2 <12902844+pradt2@users.noreply.github.com> Date: Thu, 13 Feb 2025 15:33:53 -0800 Subject: [PATCH 7/7] [Offload]: Skip copying of unused kernel-mapped data --- offload/test/mapping/skip_transfers.cpp | 109 +++++++++++++----------- 1 file changed, 61 insertions(+), 48 deletions(-) diff --git a/offload/test/mapping/skip_transfers.cpp b/offload/test/mapping/skip_transfers.cpp index 93b615a0ab3d6..6e5bc605613bc 100644 --- a/offload/test/mapping/skip_transfers.cpp +++ b/offload/test/mapping/skip_transfers.cpp @@ -9,67 +9,80 @@ int main() { float DataStack = 0; - // CHECK: omptarget device 0 info: Copying data from device to host, TgtPtr=0x{{.*}}, HstPtr=0x{{.*}}, Size=4, Name=unknown - #pragma omp target map(from: DataStack) - { - DataStack = 1; - } +// CHECK: omptarget device 0 info: Copying data from device to host, +// TgtPtr=0x{{.*}}, HstPtr=0x{{.*}}, Size=4, Name=unknown +#pragma omp target map(from : DataStack) + { + DataStack = 1; + } - // CHECK: omptarget device 0 info: Copying data from host to device, HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=4, Name=unknown - #pragma omp target map(always to: DataStack) - ; +// CHECK: omptarget device 0 info: Copying data from host to device, +// HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=4, Name=unknown +#pragma omp target map(always to : DataStack) + ; - // CHECK: omptarget device 0 info: tofrom(unknown)[4] is not used and will not be copied - #pragma omp target map(tofrom: DataStack) - ; +// CHECK: omptarget device 0 info: tofrom(unknown)[4] is not used and will not +// be copied +#pragma omp target map(tofrom : DataStack) + ; - int Size = 16; - double *Data = new double[Size]; + int Size = 16; + double *Data = new double[Size]; - // CHECK: omptarget device 0 info: Copying data from host to device, HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=8, Name=unknown - // CHECK: omptarget device 0 info: Copying data from device to host, TgtPtr=0x{{.*}}, HstPtr=0x{{.*}}, Size=8, Name=unknown - #pragma omp target map(tofrom: Data[0:1]) - { - Data[0] = 1; - } +// CHECK: omptarget device 0 info: Copying data from host to device, +// HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=8, Name=unknown CHECK: omptarget +// device 0 info: Copying data from device to host, TgtPtr=0x{{.*}}, +// HstPtr=0x{{.*}}, Size=8, Name=unknown +#pragma omp target map(tofrom : Data[0 : 1]) + { + Data[0] = 1; + } - // CHECK: omptarget device 0 info: Copying data from host to device, HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=16, Name=unknown - // CHECK: omptarget device 0 info: Copying data from device to host, TgtPtr=0x{{.*}}, HstPtr=0x{{.*}}, Size=16, Name=unknown - #pragma omp target map(always tofrom: Data[0:2]) - ; +// CHECK: omptarget device 0 info: Copying data from host to device, +// HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=16, Name=unknown CHECK: omptarget +// device 0 info: Copying data from device to host, TgtPtr=0x{{.*}}, +// HstPtr=0x{{.*}}, Size=16, Name=unknown +#pragma omp target map(always tofrom : Data[0 : 2]) + ; - // CHECK: omptarget device 0 info: from(unknown)[24] is not used and will not be copied - #pragma omp target map(from: Data[0:3]) - ; +// CHECK: omptarget device 0 info: from(unknown)[24] is not used and will not be +// copied +#pragma omp target map(from : Data[0 : 3]) + ; - // CHECK: omptarget device 0 info: to(unknown)[24] is not used and will not be copied - #pragma omp target map(to: Data[0:3]) - ; +// CHECK: omptarget device 0 info: to(unknown)[24] is not used and will not be +// copied +#pragma omp target map(to : Data[0 : 3]) + ; - // CHECK: omptarget device 0 info: tofrom(unknown)[32] is not used and will not be copied - #pragma omp target map(tofrom: Data[0:4]) - ; +// CHECK: omptarget device 0 info: tofrom(unknown)[32] is not used and will not +// be copied +#pragma omp target map(tofrom : Data[0 : 4]) + ; - // CHECK: omptarget device 0 info: Copying data from host to device, HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=40, Name=unknown - #pragma omp target map(to: Data[0:5]) - { - #pragma omp teams - Data[0] = 1; - } +// CHECK: omptarget device 0 info: Copying data from host to device, +// HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=40, Name=unknown +#pragma omp target map(to : Data[0 : 5]) + { +#pragma omp teams + Data[0] = 1; + } - struct { - double *Data; - } Wrapper { .Data = Data }; + struct { + double *Data; + } Wrapper{.Data = Data}; - // CHECK: omptarget device 0 info: Copying data from host to device, HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=48, Name=unknown - // CHECK: omptarget device 0 info: Copying data from device to host, TgtPtr=0x{{.*}}, HstPtr=0x{{.*}}, Size=48, Name=unknown - #pragma omp target map(tofrom: Wrapper.Data[0:6]) - { - Wrapper.Data[0] = 1; - } +// CHECK: omptarget device 0 info: Copying data from host to device, +// HstPtr=0x{{.*}}, TgtPtr=0x{{.*}}, Size=48, Name=unknown CHECK: omptarget +// device 0 info: Copying data from device to host, TgtPtr=0x{{.*}}, +// HstPtr=0x{{.*}}, Size=48, Name=unknown +#pragma omp target map(tofrom : Wrapper.Data[0 : 6]) + { + Wrapper.Data[0] = 1; + } // CHECK: omptarget device 0 info: unknown(unknown)[8] is not used and will not be copied // CHECK: omptarget device 0 info: tofrom(unknown)[56] is not used and will not be copied #pragma omp target map(tofrom: Wrapper.Data[0:7]) - ; + ; }