[llvm] Revert "[OpenMP][Offload] Handle `present/to/from` when a different entry did `alloc/delete`." (PR #184240)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Mar 2 13:50:41 PST 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-offload
Author: Abhinav Gaba (abhinavgaba)
<details>
<summary>Changes</summary>
Reverts llvm/llvm-project#<!-- -->165494
Some buildbots are not happy about CHECKs enforcing strict ordering of prints inside/after target regions.
---
Patch is 50.21 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/184240.diff
17 Files Affected:
- (modified) offload/include/OpenMP/Mapping.h (+15-103)
- (modified) offload/libomptarget/OpenMP/Mapping.cpp (+4-31)
- (modified) offload/libomptarget/interface.cpp (+7-10)
- (modified) offload/libomptarget/omptarget.cpp (+50-223)
- (removed) offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c (-48)
- (removed) offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c (-49)
- (removed) offload/test/mapping/map_ordering_tgt_alloc_from_to.c (-26)
- (removed) offload/test/mapping/map_ordering_tgt_alloc_present_tofrom.c (-26)
- (removed) offload/test/mapping/map_ordering_tgt_alloc_tofrom.c (-15)
- (removed) offload/test/mapping/map_ordering_tgt_data_alloc_from.c (-15)
- (removed) offload/test/mapping/map_ordering_tgt_data_alloc_to_from.c (-18)
- (removed) offload/test/mapping/map_ordering_tgt_data_alloc_tofrom.c (-18)
- (removed) offload/test/mapping/map_ordering_tgt_exit_data_always_always.c (-28)
- (removed) offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c (-20)
- (removed) offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c (-42)
- (removed) offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c (-42)
- (removed) offload/test/mapping/map_ordering_tgt_exit_data_from_mapper_overlap.c (-50)
``````````diff
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index e4024abf26690..d40e2b188076a 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -495,110 +495,20 @@ struct AttachMapInfo {
MapType(Type), Pointername(Name) {}
};
-/// Structure to track new allocations, ATTACH entries, DELETE entries and
-/// skipped FROM data transfer information for a given construct, across
-/// recursive calls (for handling mappers) to targetDataBegin/targetDataEnd.
-struct StateInfoTy {
- /// ATTACH map entries for deferred processing until all other maps are done.
+/// Structure to track ATTACH entries and new allocations across recursive calls
+/// (for handling mappers) to targetDataBegin for a given construct.
+struct AttachInfoTy {
+ /// ATTACH map entries for deferred processing.
llvm::SmallVector<AttachMapInfo> AttachEntries;
- /// Host pointers for which new memory was allocated.
/// Key: host pointer, Value: allocation size.
llvm::DenseMap<void *, int64_t> NewAllocations;
- /// Host pointers that had a FROM entry, but for which a data transfer was
- /// skipped due to the ref-count not being zero.
- /// Key: host pointer, Value: data size.
- llvm::DenseMap<void *, int64_t> SkippedFromEntries;
-
- /// Host pointers for which we have triggered a FROM transfer at some point
- /// during targetDataEnd. It's used to avoid duplicate transfers.
- /// Key: host pointer, Value: transferred size.
- llvm::DenseMap<void *, int64_t> TransferredFromEntries;
-
- /// Starting host address and size of entries whose ref-count went to zero.
- /// This includes entries released through explicit DELETE, or normal
- /// ref-count decrements. It's used to ensure transfers are performed for FROM
- /// entries whose ref-count is already zero when the entry is encountered.
- /// Key: host pointer, Value: size.
- llvm::DenseMap<void *, int64_t> ReleasedEntries;
-
- StateInfoTy() = default;
+ AttachInfoTy() = default;
// Delete copy constructor and copy assignment operator to prevent copying
- StateInfoTy(const StateInfoTy &) = delete;
- StateInfoTy &operator=(const StateInfoTy &) = delete;
-
-private:
- /// Helper to find an entry in \p EntryMap that contains the pointer.
- /// Returns the matching entry if found, otherwise std::nullopt.
- std::optional<std::pair<void *, int64_t>>
- findEntryForPtr(void *Ptr,
- const llvm::DenseMap<void *, int64_t> &EntryMap) const {
- for (const auto &Entry : EntryMap) {
- void *EntryBegin = Entry.first;
- int64_t EntrySize = Entry.second;
- if (Ptr >= EntryBegin &&
- Ptr < static_cast<void *>(static_cast<char *>(EntryBegin) +
- EntrySize)) {
- return Entry;
- }
- }
- return std::nullopt;
- }
-
-public:
- /// Check if a pointer falls within any of the newly allocated ranges.
- /// Returns the matching entry if found, otherwise std::nullopt.
- std::optional<std::pair<void *, int64_t>> wasNewlyAllocated(void *Ptr) const {
- return findEntryForPtr(Ptr, NewAllocations);
- }
-
- /// Check if a pointer range [Ptr, Ptr+Size) is fully contained within any
- /// previously completed FROM transfer.
- /// Returns the matching entry if found, otherwise std::nullopt.
- std::optional<std::pair<void *, int64_t>>
- wasTransferredFrom(void *Ptr, int64_t Size) const {
- uintptr_t CheckBegin = reinterpret_cast<uintptr_t>(Ptr);
- uintptr_t CheckEnd = CheckBegin + Size;
-
- for (const auto &Entry : TransferredFromEntries) {
- void *RangePtr = Entry.first;
- int64_t RangeSize = Entry.second;
- uintptr_t RangeBegin = reinterpret_cast<uintptr_t>(RangePtr);
- uintptr_t RangeEnd = RangeBegin + RangeSize;
-
- if (CheckBegin >= RangeBegin && CheckEnd <= RangeEnd) {
- return Entry;
- }
- }
- return std::nullopt;
- }
-
- /// Check if a pointer falls within any released entry's range.
- /// Returns the matching entry if found, otherwise std::nullopt.
- std::optional<std::pair<void *, int64_t>>
- wasPreviouslyReleased(void *Ptr) const {
- return findEntryForPtr(Ptr, ReleasedEntries);
- }
-
- /// Add a skipped FROM entry. Only updates the entry if this is a new pointer
- /// or if the new size is larger than the existing entry.
- void addSkippedFromEntry(void *Ptr, int64_t Size) {
- auto It = SkippedFromEntries.find(Ptr);
- if (It == SkippedFromEntries.end() || Size > It->second) {
- SkippedFromEntries[Ptr] = Size;
- }
- }
-
- /// Add a transferred FROM entry. Only updates the entry if this is a new
- /// pointer or if the new size is larger than the existing entry.
- void addTransferredFromEntry(void *Ptr, int64_t Size) {
- auto It = TransferredFromEntries.find(Ptr);
- if (It == TransferredFromEntries.end() || Size > It->second) {
- TransferredFromEntries[Ptr] = Size;
- }
- }
+ AttachInfoTy(const AttachInfoTy &) = delete;
+ AttachInfoTy &operator=(const AttachInfoTy &) = delete;
};
// Function pointer type for targetData* functions (targetDataBegin,
@@ -606,7 +516,7 @@ struct StateInfoTy {
typedef int (*TargetDataFuncPtrTy)(ident_t *, DeviceTy &, int32_t, void **,
void **, int64_t *, int64_t *,
map_var_info_t *, void **, AsyncInfoTy &,
- StateInfoTy *, bool);
+ AttachInfoTy *, bool);
void dumpTargetPointerMappings(const ident_t *Loc, DeviceTy &Device,
bool toStdOut = false);
@@ -615,22 +525,24 @@ int targetDataBegin(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,
- StateInfoTy *StateInfo = nullptr, bool FromMapper = false);
+ AttachInfoTy *AttachInfo = nullptr,
+ bool FromMapper = false);
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,
- StateInfoTy *StateInfo = nullptr, bool FromMapper = false);
+ AttachInfoTy *AttachInfo = nullptr, bool FromMapper = false);
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,
- StateInfoTy *StateInfo = nullptr, bool FromMapper = false);
+ AttachInfoTy *AttachInfo = nullptr,
+ bool FromMapper = false);
// Process deferred ATTACH map entries collected during targetDataBegin.
-int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
+int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
AsyncInfoTy &AsyncInfo);
struct MappingInfoTy {
@@ -671,7 +583,7 @@ struct MappingInfoTy {
bool HasFlagTo, bool HasFlagAlways, bool IsImplicit, bool UpdateRefCount,
bool HasCloseModifier, bool HasPresentModifier, bool HasHoldModifier,
AsyncInfoTy &AsyncInfo, HostDataToTargetTy *OwnedTPR = nullptr,
- bool ReleaseHDTTMap = true, StateInfoTy *StateInfo = nullptr);
+ bool ReleaseHDTTMap = true);
/// Return the target pointer for \p HstPtrBegin in \p HDTTMap. The accessor
/// ensures exclusive access to the HDTT map.
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index 1bb2e424bd083..b8edd7978951f 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -209,8 +209,7 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
int64_t TgtPadding, int64_t Size, map_var_info_t HstPtrName, bool HasFlagTo,
bool HasFlagAlways, bool IsImplicit, bool UpdateRefCount,
bool HasCloseModifier, bool HasPresentModifier, bool HasHoldModifier,
- AsyncInfoTy &AsyncInfo, HostDataToTargetTy *OwnedTPR, bool ReleaseHDTTMap,
- StateInfoTy *StateInfo) {
+ AsyncInfoTy &AsyncInfo, HostDataToTargetTy *OwnedTPR, bool ReleaseHDTTMap) {
LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size, OwnedTPR);
LR.TPR.Flags.IsPresent = true;
@@ -329,36 +328,10 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
if (ReleaseHDTTMap)
HDTTMap.destroy();
- // Lambda to check if this pointer was newly allocated on the current region.
- // This is needed to handle cases when the TO entry is encountered after an
- // alloc entry for the same pointer. In such cases, the ref-count is already
- // non-zero when TO is encountered, but we still need to do a transfer. e.g.
- //
- // struct S {
- // int *p;
- // };
- // #pragma omp declare mapper(id : S s) map(to: s.p, s.p[0 : 10])
- //
- // S s1;
- // ...
- // #pragma omp target map(alloc : s1.p[0 : 10]) map(mapper(id), to : s1)
- auto WasNewlyAllocatedForCurrentRegion = [&]() {
- if (!StateInfo)
- return false;
- bool WasNewlyAllocated =
- StateInfo->wasNewlyAllocated(HstPtrBegin).has_value();
- if (WasNewlyAllocated)
- ODBG(ODT_Mapping) << "HstPtrBegin " << HstPtrBegin
- << " was newly allocated for the current region";
- return WasNewlyAllocated;
- };
-
- // Even if this isn't a new entry, we still need to do a data-transfer if
- // the pointer was newly allocated on the current target region.
+ // If the target pointer is valid, and we need to transfer data, issue the
+ // data transfer.
if (LR.TPR.TargetPointer && !LR.TPR.Flags.IsHostPointer && HasFlagTo &&
- (LR.TPR.Flags.IsNewEntry || HasFlagAlways ||
- WasNewlyAllocatedForCurrentRegion()) &&
- Size != 0) {
+ (LR.TPR.Flags.IsNewEntry || HasFlagAlways) && Size != 0) {
// If we have something like:
// #pragma omp target map(to: s.myarr[0:10]) map(to: s.myarr[0:10])
diff --git a/offload/libomptarget/interface.cpp b/offload/libomptarget/interface.cpp
index 354a0e33c03fc..31c0fe281ac4f 100644
--- a/offload/libomptarget/interface.cpp
+++ b/offload/libomptarget/interface.cpp
@@ -168,22 +168,19 @@ targetData(ident_t *Loc, int64_t DeviceId, int32_t ArgNum, void **ArgsBase,
int Rc = OFFLOAD_SUCCESS;
- // Allocate StateInfo for targetDataBegin and targetDataEnd to track
- // allocations, pointer attachments and deferred transfers.
- // This is not needed for targetDataUpdate.
- std::unique_ptr<StateInfoTy> StateInfo;
- if (TargetDataFunction == targetDataBegin ||
- TargetDataFunction == targetDataEnd)
- StateInfo = std::make_unique<StateInfoTy>();
+ // Only allocate AttachInfo for targetDataBegin
+ std::unique_ptr<AttachInfoTy> AttachInfo;
+ if (TargetDataFunction == targetDataBegin)
+ AttachInfo = std::make_unique<AttachInfoTy>();
Rc = TargetDataFunction(Loc, *DeviceOrErr, ArgNum, ArgsBase, Args, ArgSizes,
ArgTypes, ArgNames, ArgMappers, AsyncInfo,
- StateInfo.get(), /*FromMapper=*/false);
+ AttachInfo.get(), /*FromMapper=*/false);
if (Rc == OFFLOAD_SUCCESS) {
// Process deferred ATTACH entries BEFORE synchronization
- if (StateInfo && !StateInfo->AttachEntries.empty())
- Rc = processAttachEntries(*DeviceOrErr, *StateInfo, AsyncInfo);
+ if (AttachInfo && !AttachInfo->AttachEntries.empty())
+ Rc = processAttachEntries(*DeviceOrErr, *AttachInfo, AsyncInfo);
if (Rc == OFFLOAD_SUCCESS)
Rc = AsyncInfo.synchronize();
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 344c388e794af..bd99edee5e1b3 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -298,7 +298,7 @@ int targetDataMapper(ident_t *Loc, DeviceTy &Device, void *ArgBase, void *Arg,
int64_t ArgSize, int64_t ArgType, map_var_info_t ArgNames,
void *ArgMapper, AsyncInfoTy &AsyncInfo,
TargetDataFuncPtrTy TargetDataFunction,
- StateInfoTy *StateInfo = nullptr) {
+ AttachInfoTy *AttachInfo = nullptr) {
ODBG(ODT_Interface) << "Calling the mapper function " << ArgMapper;
// The mapper function fills up Components.
@@ -329,7 +329,7 @@ int targetDataMapper(ident_t *Loc, DeviceTy &Device, void *ArgBase, void *Arg,
MapperArgsBase.data(), MapperArgs.data(),
MapperArgSizes.data(), MapperArgTypes.data(),
MapperArgNames.data(), /*arg_mappers*/ nullptr,
- AsyncInfo, StateInfo, /*FromMapper=*/true);
+ AsyncInfo, AttachInfo, /*FromMapper=*/true);
return Rc;
}
@@ -512,9 +512,9 @@ int targetDataBegin(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,
- StateInfoTy *StateInfo, bool FromMapper) {
- assert(StateInfo && "StateInfo must be available for targetDataBegin for "
- "handling ATTACH and TO/TOFROM map-types.");
+ AttachInfoTy *AttachInfo, bool FromMapper) {
+ assert(AttachInfo && "AttachInfo must be available for targetDataBegin for "
+ "handling ATTACH map-types.");
// process each input.
for (int32_t I = 0; I < ArgNum; ++I) {
// Ignore private variables and arrays - there is no mapping for them.
@@ -533,7 +533,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,
- targetDataBegin, StateInfo);
+ targetDataBegin, AttachInfo);
if (Rc != OFFLOAD_SUCCESS) {
REPORT() << "Call to targetDataBegin via targetDataMapper for custom "
@@ -560,7 +560,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// similar to firstprivate (PRIVATE | TO) entries by
// PrivateArgumentManager.
if (!IsCorrespondingPointerInit)
- StateInfo->AttachEntries.emplace_back(
+ AttachInfo->AttachEntries.emplace_back(
/*PointerBase=*/HstPtrBase, /*PointeeBegin=*/HstPtrBegin,
/*PointerSize=*/DataSize, /*MapType=*/ArgTypes[I],
/*PointeeName=*/HstPtrName);
@@ -637,7 +637,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// Track new allocation, for eventual use in attachment decision-making.
if (PointerTpr.Flags.IsNewEntry && !IsHostPtr)
- StateInfo->NewAllocations[HstPtrBase] = sizeof(void *);
+ AttachInfo->NewAllocations[HstPtrBase] = sizeof(void *);
ODBG(ODT_Mapping) << "There are " << sizeof(void *)
<< " bytes allocated at target address "
@@ -659,8 +659,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
auto TPR = Device.getMappingInfo().getTargetPointer(
HDTTMap, HstPtrBegin, HstPtrBase, TgtPadding, DataSize, HstPtrName,
HasFlagTo, HasFlagAlways, IsImplicit, UpdateRef, HasCloseModifier,
- HasPresentModifier, HasHoldModifier, AsyncInfo, PointerTpr.getEntry(),
- /*ReleaseHDTTMap=*/true, StateInfo);
+ HasPresentModifier, HasHoldModifier, AsyncInfo, PointerTpr.getEntry());
void *TgtPtrBegin = TPR.TargetPointer;
IsHostPtr = TPR.Flags.IsHostPointer;
// If data_size==0, then the argument could be a zero-length pointer to
@@ -671,26 +670,11 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
: "device failure or illegal mapping")
<< ").";
return OFFLOAD_FAIL;
- } else if (TgtPtrBegin && HasPresentModifier &&
- StateInfo->wasNewlyAllocated(HstPtrBegin).has_value()) {
- // For "PRESENT" entries, we may have cases like the following:
- // int *xp = &x[0];
- // map(alloc: x[:]) map(present, alloc: xp[1])
- // The "PRESENT" entry may be encountered after a previous entry
- // allocated new storage for the pointer.
- // To catch such cases, we need to look at any existing allocations
- // and error out if we have any matching the pointer.
- MESSAGE("device mapping required by 'present' map type modifier does not "
- "exist for host address " DPxMOD " (%" PRId64 " bytes)\n",
- DPxPTR(HstPtrBegin), DataSize);
- REPORT() << "Pointer " << HstPtrBegin
- << " was not present on the device upon entry to the region.";
- return OFFLOAD_FAIL;
}
- // Track new allocation, for eventual use in attachment/to decision-making.
+ // Track new allocation, for eventual use in attachment decision-making.
if (TPR.Flags.IsNewEntry && !IsHostPtr && TgtPtrBegin)
- StateInfo->NewAllocations[HstPtrBegin] = DataSize;
+ AttachInfo->NewAllocations[HstPtrBegin] = DataSize;
ODBG(ODT_Mapping) << "There are " << DataSize
<< " bytes allocated at target address " << TgtPtrBegin
@@ -810,24 +794,24 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
///
/// For this purpose, we insert a data_fence before the first
/// pointer-attachment, (3), to ensure that all pending transfers finish first.
-int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
+int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
AsyncInfoTy &AsyncInfo) {
// Report all tracked allocations from both main loop and ATTACH processing
- if (!StateInfo.NewAllocations.empty()) {
+ if (!AttachInfo.NewAllocations.empty()) {
ODBG_OS(ODT_Mapping, [&](llvm::raw_ostream &OS) {
- OS << "Tracked " << StateInfo.NewAllocations.size()
+ OS << "Tracked " << AttachInfo.NewAllocations.size()
<< " total new allocations:";
- for (const auto &Alloc : StateInfo.NewAllocations) {
+ for (const auto &Alloc : AttachInfo.NewAllocations) {
OS << " Host ptr: " << Alloc.first << ", Size: " << Alloc.second
<< " bytes";
}
});
}
- if (StateInfo.AttachEntries.empty())
+ if (AttachInfo.AttachEntries.empty())
return OFFLOAD_SUCCESS;
- ODBG(ODT_Mapping) << "Processing " << StateInfo.AttachEntries.size()
+ ODBG(ODT_Mapping) << "Processing " << AttachInfo.AttachEntries.size()
<< " deferred ATTACH map entries";
bool TreatAttachAutoAsAlways = MappingConfig::get().TreatAttachAutoAsAlways;
@@ -837,9 +821,9 @@ int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
int Ret = OFFLOAD_SUCCESS;
bool IsFirstPointerAttachment = true;
- for (size_t EntryIdx = 0; EntryIdx < StateInfo.AttachEntries.size();
+ for (size_t EntryIdx = 0; EntryIdx < AttachInfo.AttachEntries.size();
++EntryIdx) {
- const auto &AttachEntry = StateInfo.AttachEntries[EntryIdx];
+ const auto &AttachEntry = AttachInfo.AttachEntries[EntryIdx];
void **HstPtr = reinterpret_cast<void **>(AttachEntry.PointerBase);
@@ -860,11 +844,18 @@ int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
// Lambda to check if a pointer was newly allocated
auto WasNewlyAllocated = [&](void *Ptr, const char *PtrName) {
- bool WasNewlyAllocated = StateInfo.wasNewlyAllocated(Ptr).has_value();
+ bool IsNewlyAllocated =
+ llvm::any_of(AttachInfo.NewAllocations, [&](const auto &Alloc) {
+ void *AllocPtr = Alloc.first;
+ int64_t AllocSize = Alloc.second;
+ return Ptr >= AllocPtr &&
+ Ptr < reinterpret_cast<void *>(
+ reinterpret_cast<char *>(AllocPtr) + AllocSize);
+ });
...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/184240
More information about the llvm-commits
mailing list