[llvm] [Offload]: Skip copying of unused kernel-mapped data (PR #124723)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Feb 13 14:48:28 PST 2025
https://github.com/pradt2 updated https://github.com/llvm/llvm-project/pull/124723
>From 6e42d525cf807d35a9522795d681d90eb9b30eb7 Mon Sep 17 00:00:00 2001
From: pradt2 <12902844+pradt2 at users.noreply.github.com>
Date: Tue, 28 Jan 2025 01:00:58 -0800
Subject: [PATCH 1/4] [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<int64_t[]> maskIgnorableMappings(int64_t DeviceId, int32_t ArgNum, int64_t *ArgTypes,
+ int64_t *ArgSizes, map_var_info_t *ArgNames) {
+ std::unique_ptr<int64_t[]> ArgTypesOverride = std::make_unique<int64_t[]>(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<int64_t[]> 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 at users.noreply.github.com>
Date: Wed, 5 Feb 2025 19:17:11 -0800
Subject: [PATCH 2/4] [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<int64_t[]> 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<int64_t[]> maskRedundantTransfers(DeviceTy &Device, int32_t ArgNum,
+ int64_t *ArgTypes, int64_t *ArgSizes,
+ map_var_info_t *ArgNames, void **ArgPtrs,
+ void **ArgMappers) {
std::unique_ptr<int64_t[]> ArgTypesOverride = std::make_unique<int64_t[]>(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<int64_t[]> 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 at users.noreply.github.com>
Date: Thu, 13 Feb 2025 14:36:21 -0800
Subject: [PATCH 3/4] [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<int64_t[]> maskRedundantTransfers(DeviceTy &Device, int32_t ArgNum,
int64_t *ArgTypes, int64_t *ArgSizes,
map_var_info_t *ArgNames, void **ArgPtrs,
void **ArgMappers) {
std::unique_ptr<int64_t[]> ArgTypesOverride = std::make_unique<int64_t[]>(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 at users.noreply.github.com>
Date: Thu, 13 Feb 2025 14:47:41 -0800
Subject: [PATCH 4/4] [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<int64_t[]> 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<int64_t[]> 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;
More information about the llvm-commits
mailing list