[clang] [llvm] [NFC][Clang][OpenMP] Add a util to implicitly map attach-ptr-exprs on `target`, when applicable. (PR #161294)
Abhinav Gaba via llvm-commits
llvm-commits at lists.llvm.org
Mon Oct 6 12:40:46 PDT 2025
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/161294
>From f0cd9cdd7cacbe219b6c73711feeef804c41fa71 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 29 Sep 2025 14:29:12 -0700
Subject: [PATCH 1/4] Reuse the same comparator object for the same MEH.
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 31 ++++++++++++++++-----------
1 file changed, 18 insertions(+), 13 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index c90e1a487daf9..9e64ea8ad1c2e 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -6804,12 +6804,13 @@ class MappableExprsHandler {
/// they were computed by collectAttachPtrExprInfo(), if they are semantically
/// different.
struct AttachPtrExprComparator {
- const MappableExprsHandler *Handler = nullptr;
+ const MappableExprsHandler &Handler;
// Cache of previous equality comparison results.
mutable llvm::DenseMap<std::pair<const Expr *, const Expr *>, bool>
CachedEqualityComparisons;
- AttachPtrExprComparator(const MappableExprsHandler *H) : Handler(H) {}
+ AttachPtrExprComparator(const MappableExprsHandler &H) : Handler(H) {}
+ AttachPtrExprComparator() = delete;
// Return true iff LHS is "less than" RHS.
bool operator()(const Expr *LHS, const Expr *RHS) const {
@@ -6817,15 +6818,15 @@ class MappableExprsHandler {
return false;
// First, compare by complexity (depth)
- const auto ItLHS = Handler->AttachPtrComponentDepthMap.find(LHS);
- const auto ItRHS = Handler->AttachPtrComponentDepthMap.find(RHS);
+ const auto ItLHS = Handler.AttachPtrComponentDepthMap.find(LHS);
+ const auto ItRHS = Handler.AttachPtrComponentDepthMap.find(RHS);
std::optional<size_t> DepthLHS =
- (ItLHS != Handler->AttachPtrComponentDepthMap.end()) ? ItLHS->second
- : std::nullopt;
+ (ItLHS != Handler.AttachPtrComponentDepthMap.end()) ? ItLHS->second
+ : std::nullopt;
std::optional<size_t> DepthRHS =
- (ItRHS != Handler->AttachPtrComponentDepthMap.end()) ? ItRHS->second
- : std::nullopt;
+ (ItRHS != Handler.AttachPtrComponentDepthMap.end()) ? ItRHS->second
+ : std::nullopt;
// std::nullopt (no attach pointer) has lowest complexity
if (!DepthLHS.has_value() && !DepthRHS.has_value()) {
@@ -6873,8 +6874,8 @@ class MappableExprsHandler {
/// Returns true iff LHS was computed before RHS by
/// collectAttachPtrExprInfo().
bool wasComputedBefore(const Expr *LHS, const Expr *RHS) const {
- const size_t &OrderLHS = Handler->AttachPtrComputationOrderMap.at(LHS);
- const size_t &OrderRHS = Handler->AttachPtrComputationOrderMap.at(RHS);
+ const size_t &OrderLHS = Handler.AttachPtrComputationOrderMap.at(LHS);
+ const size_t &OrderRHS = Handler.AttachPtrComputationOrderMap.at(RHS);
return OrderLHS < OrderRHS;
}
@@ -6893,7 +6894,7 @@ class MappableExprsHandler {
if (!LHS || !RHS)
return false;
- ASTContext &Ctx = Handler->CGF.getContext();
+ ASTContext &Ctx = Handler.CGF.getContext();
// Strip away parentheses and no-op casts to get to the core expression
LHS = LHS->IgnoreParenNoopCasts(Ctx);
RHS = RHS->IgnoreParenNoopCasts(Ctx);
@@ -7242,6 +7243,10 @@ class MappableExprsHandler {
llvm::DenseMap<const Expr *, size_t> AttachPtrComputationOrderMap = {
{nullptr, 0}};
+ /// An instance of attach-ptr-expr comparator that can be used throughout the
+ /// lifetime of this handler.
+ AttachPtrExprComparator AttachPtrComparator;
+
llvm::Value *getExprTypeSize(const Expr *E) const {
QualType ExprTy = E->getType().getCanonicalType();
@@ -8959,7 +8964,7 @@ class MappableExprsHandler {
public:
MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF)
- : CurDir(&Dir), CGF(CGF) {
+ : CurDir(&Dir), CGF(CGF), AttachPtrComparator(*this) {
// Extract firstprivate clause information.
for (const auto *C : Dir.getClausesOfKind<OMPFirstprivateClause>())
for (const auto *D : C->varlist())
@@ -9005,7 +9010,7 @@ class MappableExprsHandler {
/// Constructor for the declare mapper directive.
MappableExprsHandler(const OMPDeclareMapperDecl &Dir, CodeGenFunction &CGF)
- : CurDir(&Dir), CGF(CGF) {}
+ : CurDir(&Dir), CGF(CGF), AttachPtrComparator(*this) {}
/// Generate code for the combined entry if we have a partially mapped struct
/// and take care of the mapping flags of the arguments corresponding to
>From 08e051ad3767bcd7944e5406623ad64b43ec2de3 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 29 Sep 2025 15:58:20 -0700
Subject: [PATCH 2/4] [NFC][Clang][OpenMP] Add a util to implicitly map
`attach-ptr-exprs` on `target, when applicable.
On a target construct, if there's an implicit map on a struct, or that of
this[:], and an explicit map with a member of that struct/class as the
base-pointer, we need to make sure that base-pointer is implicitly mapped,
to make sure we don't map the full struct/class. For example:
```cpp
struct S {
int dummy[10000];
int *p;
void f1() {
#pragma omp target map(p[0:1])
(void)this;
}
}; S s;
void f2() {
#pragma omp target map(s.p[0:10])
(void)s;
}
```
Only `this-p` and `s.p` should be mapped in the two cases above. If we
were to implicitly map the full struct `s`, or `this[0:1]`, it would map
the `dummy` field as well.
This was pulled out of #153683 to make that PR smaller. it also pulls
out one other utility, and an NFC changes to the AttachPtrExpr
comparator from that PR.
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 183 ++++++++++++++++++
.../llvm/Frontend/OpenMP/OMPConstants.h | 3 +
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 6 +
3 files changed, 192 insertions(+)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 9e64ea8ad1c2e..0aee3c081af7e 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7414,6 +7414,39 @@ class MappableExprsHandler {
return ConstLength.getSExtValue() != 1;
}
+ /// Emit an attach entry into \p CombinedInfo, using the information from \p
+ /// AttachInfo. For example, for a map of form `int *p; ... map(p[1:10])`,
+ /// an attach entry has the following form:
+ /// &p, &p[1], sizeof(void*), ATTACH
+ void emitAttachEntry(CodeGenFunction &CGF, MapCombinedInfoTy &CombinedInfo,
+ const AttachInfoTy &AttachInfo) const {
+ assert(AttachInfo.isValid() &&
+ "Expected valid attach pointer/pointee information!");
+
+ // Size is the size of the pointer itself - use pointer size, not BaseDecl
+ // size
+ llvm::Value *PointerSize = CGF.Builder.CreateIntCast(
+ llvm::ConstantInt::get(
+ CGF.CGM.SizeTy,
+ CGF.getContext().getTypeSize(
+ CGF.getContext().getPointerType(CGF.getContext().VoidTy)) /
+ 8),
+ CGF.Int64Ty, /*isSigned=*/true);
+
+ CombinedInfo.Exprs.emplace_back(AttachInfo.AttachPtrDecl,
+ AttachInfo.AttachMapExpr);
+ CombinedInfo.BasePointers.push_back(
+ AttachInfo.AttachPtrAddr.emitRawPointer(CGF));
+ CombinedInfo.DevicePtrDecls.push_back(nullptr);
+ CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
+ CombinedInfo.Pointers.push_back(
+ AttachInfo.AttachPteeAddr.emitRawPointer(CGF));
+ CombinedInfo.Sizes.push_back(PointerSize);
+ CombinedInfo.Types.push_back(OpenMPOffloadMappingFlags::OMP_MAP_ATTACH);
+ CombinedInfo.Mappers.push_back(nullptr);
+ CombinedInfo.NonContigInfo.Dims.push_back(1);
+ }
+
/// A helper class to copy structures with overlapped elements, i.e. those
/// which have mappings of both "s" and "s.mem". Consecutive elements that
/// are not explicitly copied have mapping nodes synthesized for them,
@@ -9252,6 +9285,156 @@ class MappableExprsHandler {
}
}
+ /// On a target construct, if there's an implicit map on a struct, or that of
+ /// this[:], and an explicit map with a member of that struct/class as the
+ /// base-pointer, we need to make sure that base-pointer is implicitly mapped,
+ /// to make sure we don't map the full struct/class. For example:
+ ///
+ /// \code
+ /// struct S {
+ /// int dummy[10000];
+ /// int *p;
+ /// void f1() {
+ /// #pragma omp target map(p[0:1])
+ /// (void)this;
+ /// }
+ /// }; S s;
+ ///
+ /// void f2() {
+ /// #pragma omp target map(s.p[0:10])
+ /// (void)s;
+ /// }
+ /// \endcode
+ ///
+ /// Only `this-p` and `s.p` should be mapped in the two cases above.
+ //
+ // OpenMP 6.0: 7.9.6 map clause, pg 285
+ // If a list item with an implicitly determined data-mapping attribute does
+ // not have any corresponding storage in the device data environment prior to
+ // a task encountering the construct associated with the map clause, and one
+ // or more contiguous parts of the original storage are either list items or
+ // base pointers to list items that are explicitly mapped on the construct,
+ // only those parts of the original storage will have corresponding storage in
+ // the device data environment as a result of the map clauses on the
+ // construct.
+ void addImplicitMapForAttachPtrBaseIfMemberOfCapturedVD(
+ const ValueDecl *CapturedVD, MapDataArrayTy &DeclComponentLists,
+ SmallVectorImpl<
+ SmallVector<OMPClauseMappableExprCommon::MappableComponent, 8>>
+ &ComponentVectorStorage) const {
+ bool IsThisCapture = CapturedVD == nullptr;
+
+ for (const auto &ComponentsAndAttachPtr : AttachPtrExprMap) {
+ OMPClauseMappableExprCommon::MappableExprComponentListRef
+ ComponentsWithAttachPtr = ComponentsAndAttachPtr.first;
+ const Expr *AttachPtrExpr = ComponentsAndAttachPtr.second;
+ if (!AttachPtrExpr)
+ continue;
+
+ const auto *ME = dyn_cast<MemberExpr>(AttachPtrExpr);
+ if (!ME)
+ continue;
+
+ const Expr *Base = ME->getBase()->IgnoreParenImpCasts();
+
+ // If we are handling a "this" capture, then we are looking for
+ // attach-ptrs of form `this->p`, either explicitly or implicitly.
+ if (IsThisCapture && !ME->isImplicitCXXThis() && !isa<CXXThisExpr>(Base))
+ continue;
+
+ if (!IsThisCapture && (!isa<DeclRefExpr>(Base) ||
+ cast<DeclRefExpr>(Base)->getDecl() != CapturedVD))
+ continue;
+
+ // For non-this captures, we are looking for attach-ptrs of form
+ // `s.p`.
+ // For non-this captures, we are looking for attach-ptrs like `s.p`.
+ if (!IsThisCapture && (ME->isArrow() || !isa<DeclRefExpr>(Base) ||
+ cast<DeclRefExpr>(Base)->getDecl() != CapturedVD))
+ continue;
+
+ // Check if we have an existing map on either:
+ // this[:], s, this->p, or s.p, in which case, we don't need to add
+ // an implicit one for the attach-ptr s.p/this->p.
+ bool FoundExistingMap = false;
+ for (const MapData &ExistingL : DeclComponentLists) {
+ OMPClauseMappableExprCommon::MappableExprComponentListRef
+ ExistingComponents = std::get<0>(ExistingL);
+
+ if (ExistingComponents.empty())
+ continue;
+
+ // First check if we have a map like map(this->p) or map(s.p).
+ const auto &FirstComponent = ExistingComponents.front();
+ const Expr *FirstExpr = FirstComponent.getAssociatedExpression();
+
+ if (!FirstExpr)
+ continue;
+
+ // First check if we have a map like map(this->p) or map(s.p).
+ if (AttachPtrComparator.areEqual(FirstExpr, AttachPtrExpr)) {
+ FoundExistingMap = true;
+ break;
+ }
+
+ // Check if we have a map like this[0:1]
+ if (IsThisCapture) {
+ if (const auto *OASE = dyn_cast<ArraySectionExpr>(FirstExpr)) {
+ if (isa<CXXThisExpr>(OASE->getBase()->IgnoreParenImpCasts())) {
+ FoundExistingMap = true;
+ break;
+ }
+ }
+ continue;
+ }
+
+ // When the attach-ptr is something like `s.p`, check if
+ // `s` itself is mapped explicitly.
+ if (const auto *DRE = dyn_cast<DeclRefExpr>(FirstExpr)) {
+ if (DRE->getDecl() == CapturedVD) {
+ FoundExistingMap = true;
+ break;
+ }
+ }
+ }
+
+ if (FoundExistingMap)
+ continue;
+
+ // If no base map is found, we need to create an implicit map for the
+ // attach-pointer expr.
+
+ ComponentVectorStorage.emplace_back();
+ auto &AttachPtrComponents = ComponentVectorStorage.back();
+
+ static const OpenMPMapModifierKind Unknown = OMPC_MAP_MODIFIER_unknown;
+ bool SeenAttachPtrComponent = false;
+ // For creating a map on the attach-ptr `s.p/this->p`, we copy all
+ // components from the component-list which has `s.p/this->p`
+ // as the attach-ptr, starting from the component which matches
+ // `s.p/this->p`. This way, we'll have component-lists of
+ // `s.p` -> `s`, and `this->p` -> `this`.
+ for (size_t i = 0; i < ComponentsWithAttachPtr.size(); ++i) {
+ const auto &Component = ComponentsWithAttachPtr[i];
+ const Expr *ComponentExpr = Component.getAssociatedExpression();
+
+ if (!SeenAttachPtrComponent && ComponentExpr != AttachPtrExpr)
+ continue;
+ SeenAttachPtrComponent = true;
+
+ AttachPtrComponents.emplace_back(Component.getAssociatedExpression(),
+ Component.getAssociatedDeclaration(),
+ Component.isNonContiguous());
+ }
+ assert(!AttachPtrComponents.empty() &&
+ "Could not populate component-lists for mapping attach-ptr");
+
+ DeclComponentLists.emplace_back(
+ AttachPtrComponents, OMPC_MAP_tofrom, Unknown,
+ /*IsImplicit=*/true, /*mapper=*/nullptr, AttachPtrExpr);
+ }
+ }
+
/// For a capture that has an associated clause, generate the base pointers,
/// section pointers, sizes, map types, and mappers (all included in
/// \a CurCaptureVarInfo).
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h b/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h
index 6e1bce12af8e4..7bec7e0c6736d 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h
@@ -239,6 +239,9 @@ enum class OpenMPOffloadMappingFlags : uint64_t {
// dynamic.
// This is an OpenMP extension for the sake of OpenACC support.
OMP_MAP_OMPX_HOLD = 0x2000,
+ // Attach pointer and pointee, after processing all other maps.
+ // Applicable to map-entering directives. Does not change ref-count.
+ OMP_MAP_ATTACH = 0x4000,
/// Signal that the runtime library should use args as an array of
/// descriptor_dim pointers and use args_size as dims. Used when we have
/// non-contiguous list items in target update directive
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 9b67465faab0b..181f923d8c470 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -10358,6 +10358,12 @@ void OpenMPIRBuilder::setCorrectMemberOfFlag(
omp::OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF))
return;
+ // Entries with ATTACH are not members-of anything. They are handled
+ // separately by the runtime after other maps have been handled.
+ if (static_cast<std::underlying_type_t<omp::OpenMPOffloadMappingFlags>>(
+ Flags & omp::OpenMPOffloadMappingFlags::OMP_MAP_ATTACH))
+ return;
+
// Reset the placeholder value to prepare the flag for the assignment of the
// proper MEMBER_OF value.
Flags &= ~omp::OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF;
>From 117a718b9385411a0aa78361f8644da91f15c0d8 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Tue, 30 Sep 2025 10:15:26 -0700
Subject: [PATCH 3/4] Use getTypeSizeInChars
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 7 +++----
1 file changed, 3 insertions(+), 4 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 0aee3c081af7e..9c664cd4a1485 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7427,10 +7427,9 @@ class MappableExprsHandler {
// size
llvm::Value *PointerSize = CGF.Builder.CreateIntCast(
llvm::ConstantInt::get(
- CGF.CGM.SizeTy,
- CGF.getContext().getTypeSize(
- CGF.getContext().getPointerType(CGF.getContext().VoidTy)) /
- 8),
+ CGF.CGM.SizeTy, CGF.getContext()
+ .getTypeSizeInChars(CGF.getContext().VoidPtrTy)
+ .getQuantity()),
CGF.Int64Ty, /*isSigned=*/true);
CombinedInfo.Exprs.emplace_back(AttachInfo.AttachPtrDecl,
>From 01c9fb108608dbdc0e3f0355c02d97923d704111 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 2 Oct 2025 23:11:19 -0700
Subject: [PATCH 4/4] Remove emitAttachEntry, and changes to
OMPIRBuilder/OMPConstants.h
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 32 -------------------
.../llvm/Frontend/OpenMP/OMPConstants.h | 3 --
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 6 ----
3 files changed, 41 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 9c664cd4a1485..160dd584bc3ca 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7414,38 +7414,6 @@ class MappableExprsHandler {
return ConstLength.getSExtValue() != 1;
}
- /// Emit an attach entry into \p CombinedInfo, using the information from \p
- /// AttachInfo. For example, for a map of form `int *p; ... map(p[1:10])`,
- /// an attach entry has the following form:
- /// &p, &p[1], sizeof(void*), ATTACH
- void emitAttachEntry(CodeGenFunction &CGF, MapCombinedInfoTy &CombinedInfo,
- const AttachInfoTy &AttachInfo) const {
- assert(AttachInfo.isValid() &&
- "Expected valid attach pointer/pointee information!");
-
- // Size is the size of the pointer itself - use pointer size, not BaseDecl
- // size
- llvm::Value *PointerSize = CGF.Builder.CreateIntCast(
- llvm::ConstantInt::get(
- CGF.CGM.SizeTy, CGF.getContext()
- .getTypeSizeInChars(CGF.getContext().VoidPtrTy)
- .getQuantity()),
- CGF.Int64Ty, /*isSigned=*/true);
-
- CombinedInfo.Exprs.emplace_back(AttachInfo.AttachPtrDecl,
- AttachInfo.AttachMapExpr);
- CombinedInfo.BasePointers.push_back(
- AttachInfo.AttachPtrAddr.emitRawPointer(CGF));
- CombinedInfo.DevicePtrDecls.push_back(nullptr);
- CombinedInfo.DevicePointers.push_back(DeviceInfoTy::None);
- CombinedInfo.Pointers.push_back(
- AttachInfo.AttachPteeAddr.emitRawPointer(CGF));
- CombinedInfo.Sizes.push_back(PointerSize);
- CombinedInfo.Types.push_back(OpenMPOffloadMappingFlags::OMP_MAP_ATTACH);
- CombinedInfo.Mappers.push_back(nullptr);
- CombinedInfo.NonContigInfo.Dims.push_back(1);
- }
-
/// A helper class to copy structures with overlapped elements, i.e. those
/// which have mappings of both "s" and "s.mem". Consecutive elements that
/// are not explicitly copied have mapping nodes synthesized for them,
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h b/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h
index 7bec7e0c6736d..6e1bce12af8e4 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h
@@ -239,9 +239,6 @@ enum class OpenMPOffloadMappingFlags : uint64_t {
// dynamic.
// This is an OpenMP extension for the sake of OpenACC support.
OMP_MAP_OMPX_HOLD = 0x2000,
- // Attach pointer and pointee, after processing all other maps.
- // Applicable to map-entering directives. Does not change ref-count.
- OMP_MAP_ATTACH = 0x4000,
/// Signal that the runtime library should use args as an array of
/// descriptor_dim pointers and use args_size as dims. Used when we have
/// non-contiguous list items in target update directive
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 181f923d8c470..9b67465faab0b 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -10358,12 +10358,6 @@ void OpenMPIRBuilder::setCorrectMemberOfFlag(
omp::OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF))
return;
- // Entries with ATTACH are not members-of anything. They are handled
- // separately by the runtime after other maps have been handled.
- if (static_cast<std::underlying_type_t<omp::OpenMPOffloadMappingFlags>>(
- Flags & omp::OpenMPOffloadMappingFlags::OMP_MAP_ATTACH))
- return;
-
// Reset the placeholder value to prepare the flag for the assignment of the
// proper MEMBER_OF value.
Flags &= ~omp::OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF;
More information about the llvm-commits
mailing list