[llvm] [Offload] Unlock mapping entry before deleting it in disassociatePtr (PR #223319)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Sep 14 23:26:19 PDT 2026
https://github.com/StevenYangCC updated https://github.com/llvm/llvm-project/pull/223319
>From aa06372ffa5547a45bf900cabc392eef2aa86f33 Mon Sep 17 00:00:00 2001
From: "chengcang.yang" <yangchengcang at gmail.com>
Date: Mon, 14 Sep 2026 15:26:52 +0800
Subject: [PATCH] [Offload] Unlock mapping entry before deleting it in
disassociatePtr
HostDataToTargetTy stores its mutex as a member. omp_target_disassociate_ptr
deleted the mapping entry while a lock_guard still held that mutex, so the
guard's destructor unlocked freed memory.
Release the per-entry lock, erase the mapping, then delete the entry. Add a
lit test that cycles associate/disassociate on the same host pointer.
---
offload/libompaccsupport/Mapping.cpp | 51 ++++++-----
.../omp_target_associate_disassociate_ptr.c | 85 +++++++++++++++++++
2 files changed, 114 insertions(+), 22 deletions(-)
create mode 100644 offload/test/api/omp_target_associate_disassociate_ptr.c
diff --git a/offload/libompaccsupport/Mapping.cpp b/offload/libompaccsupport/Mapping.cpp
index 1bb2e424bd083..a83d71305a179 100644
--- a/offload/libompaccsupport/Mapping.cpp
+++ b/offload/libompaccsupport/Mapping.cpp
@@ -107,32 +107,39 @@ int MappingInfoTy::disassociatePtr(void *HstPtrBegin) {
REPORT() << "Association not found";
return OFFLOAD_FAIL;
}
- // Mapping exists
- HostDataToTargetTy &HDTT = *It->HDTT;
- std::lock_guard<HostDataToTargetTy> LG(HDTT);
-
- if (HDTT.getHoldRefCount()) {
- // This is based on OpenACC 3.1, sec 3.2.33 "acc_unmap_data", L3656-3657:
- // "It is an error to call acc_unmap_data if the structured reference
- // count for the pointer is not zero."
- REPORT() << "Trying to disassociate a pointer with a non-zero "
- << "hold reference count";
- return OFFLOAD_FAIL;
- }
- if (HDTT.isDynRefCountInf()) {
+ // Mapping exists. The per-entry mutex is a member of HostDataToTargetTy, so
+ // it must be released before the entry is destroyed.
+ HostDataToTargetTy *Entry = It->HDTT;
+ void *Event = nullptr;
+ {
+ std::lock_guard<HostDataToTargetTy> LG(*Entry);
+
+ if (Entry->getHoldRefCount()) {
+ // This is based on OpenACC 3.1, sec 3.2.33 "acc_unmap_data", L3656-3657:
+ // "It is an error to call acc_unmap_data if the structured reference
+ // count for the pointer is not zero."
+ REPORT() << "Trying to disassociate a pointer with a non-zero "
+ << "hold reference count";
+ return OFFLOAD_FAIL;
+ }
+
+ if (!Entry->isDynRefCountInf()) {
+ REPORT() << "Trying to disassociate a pointer which was not mapped via "
+ << "omp_target_associate_ptr";
+ return OFFLOAD_FAIL;
+ }
+
ODBG(ODT_Mapping) << "Association found, removing it";
- void *Event = HDTT.getEvent();
- delete &HDTT;
- if (Event)
- Device.destroyEvent(Event);
- HDTTMap->erase(It);
- return Device.notifyDataUnmapped(HstPtrBegin);
+ Event = Entry->getEvent();
}
- REPORT() << "Trying to disassociate a pointer which was not mapped via "
- << "omp_target_associate_ptr";
- return OFFLOAD_FAIL;
+ HDTTMap->erase(It);
+ if (Event)
+ Device.destroyEvent(Event);
+ int Ret = Device.notifyDataUnmapped(HstPtrBegin);
+ delete Entry;
+ return Ret;
}
LookupResult MappingInfoTy::lookupMapping(HDTTMapAccessorTy &HDTTMap,
diff --git a/offload/test/api/omp_target_associate_disassociate_ptr.c b/offload/test/api/omp_target_associate_disassociate_ptr.c
new file mode 100644
index 0000000000000..cac01d39049a9
--- /dev/null
+++ b/offload/test/api/omp_target_associate_disassociate_ptr.c
@@ -0,0 +1,85 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// omp_target_disassociate_ptr must unlock the host-to-device map entry before
+// destroying it. Repeating associate/disassociate on the same host pointer
+// would otherwise use a destroyed mutex.
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ int Dev = omp_get_default_device();
+ int HostVal = 42;
+ int *DevPtr = (int *)omp_target_alloc(sizeof(int), Dev);
+ if (!DevPtr) {
+ printf("omp_target_alloc failed\n");
+ return 1;
+ }
+
+ // CHECK: present before associate: 0
+ printf("present before associate: %d\n",
+ omp_target_is_present(&HostVal, Dev));
+
+ int Rc = omp_target_associate_ptr(&HostVal, DevPtr, sizeof(int), 0, Dev);
+ // CHECK: associate: 0
+ printf("associate: %d\n", Rc);
+
+ // CHECK: present after associate: 1
+ printf("present after associate: %d\n", omp_target_is_present(&HostVal, Dev));
+ // CHECK: mapped matches: 1
+ printf("mapped matches: %d\n",
+ omp_get_mapped_ptr(&HostVal, Dev) == (void *)DevPtr);
+
+ Rc = omp_target_disassociate_ptr(&HostVal, Dev);
+ // CHECK: disassociate: 0
+ printf("disassociate: %d\n", Rc);
+
+ // CHECK: present after disassociate: 0
+ printf("present after disassociate: %d\n",
+ omp_target_is_present(&HostVal, Dev));
+ // CHECK: mapped after disassociate is null: 1
+ printf("mapped after disassociate is null: %d\n",
+ omp_get_mapped_ptr(&HostVal, Dev) == NULL);
+
+ for (int I = 0; I < 8; ++I) {
+ if (omp_target_associate_ptr(&HostVal, DevPtr, sizeof(int), 0, Dev)) {
+ printf("repeated associate failed at %d\n", I);
+ omp_target_free(DevPtr, Dev);
+ return 1;
+ }
+ if (omp_target_disassociate_ptr(&HostVal, Dev)) {
+ printf("repeated disassociate failed at %d\n", I);
+ omp_target_free(DevPtr, Dev);
+ return 1;
+ }
+ }
+ // CHECK: repeated associate/disassociate: ok
+ printf("repeated associate/disassociate: ok\n");
+
+ Rc = omp_target_disassociate_ptr(&HostVal, Dev);
+ // CHECK: disassociate missing: 1
+ printf("disassociate missing: %d\n", Rc != 0);
+
+ // Device storage is independent of the host association.
+ int In = 7, Out = 0;
+ if (omp_target_memcpy(DevPtr, &In, sizeof(int), 0, 0, Dev,
+ omp_get_initial_device()) ||
+ omp_target_memcpy(&Out, DevPtr, sizeof(int), 0, 0,
+ omp_get_initial_device(), Dev)) {
+ printf("omp_target_memcpy failed\n");
+ omp_target_free(DevPtr, Dev);
+ return 1;
+ }
+ // CHECK: device memory after disassociate: 7
+ printf("device memory after disassociate: %d\n", Out);
+
+ int MappedByEnter = 0;
+#pragma omp target enter data map(alloc : MappedByEnter)
+ Rc = omp_target_disassociate_ptr(&MappedByEnter, Dev);
+ // CHECK: disassociate of mapped data: 1
+ printf("disassociate of mapped data: %d\n", Rc != 0);
+#pragma omp target exit data map(delete : MappedByEnter)
+
+ omp_target_free(DevPtr, Dev);
+ return 0;
+}
More information about the llvm-commits
mailing list