[Openmp-commits] [openmp] [OpenMP][libomptarget] Add map checks when running under unified shared memory (PR #69005)
via Openmp-commits
openmp-commits at lists.llvm.org
Wed Nov 1 13:07:08 PDT 2023
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-openmp
Author: Gheorghe-Teodor Bercea (doru1004)
<details>
<summary>Changes</summary>
This patch enables the OpenMP 6.0 feature of map checks even when running under unified shared memory.
---
Patch is 54.64 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/69005.diff
11 Files Affected:
- (modified) openmp/libomptarget/include/device.h (+4)
- (modified) openmp/libomptarget/src/device.cpp (+73-4)
- (modified) openmp/libomptarget/src/omptarget.cpp (+87-69)
- (added) openmp/libomptarget/test/unified_shared_memory/unified_map_checks_arrays.cpp (+48)
- (added) openmp/libomptarget/test/unified_shared_memory/unified_map_checks_close_enter_exit.cpp (+222)
- (added) openmp/libomptarget/test/unified_shared_memory/unified_map_checks_close_modifier.cpp (+184)
- (added) openmp/libomptarget/test/unified_shared_memory/unified_map_checks_error.cpp (+36)
- (added) openmp/libomptarget/test/unified_shared_memory/unified_map_checks_no_target.cpp (+36)
- (added) openmp/libomptarget/test/unified_shared_memory/unified_map_checks_scalars.cpp (+94)
- (added) openmp/libomptarget/test/unified_shared_memory/unified_map_checks_shared_update.cpp (+137)
- (added) openmp/libomptarget/test/unified_shared_memory/zero_sized_array.cpp (+31)
``````````diff
diff --git a/openmp/libomptarget/include/device.h b/openmp/libomptarget/include/device.h
index cd76d88618be4ee..56a4f5ba4242c18 100644
--- a/openmp/libomptarget/include/device.h
+++ b/openmp/libomptarget/include/device.h
@@ -447,6 +447,10 @@ struct DeviceTy {
/// - Data allocation failed;
/// - The user tried to do an illegal mapping;
/// - Data transfer issue fails.
+ /// If unified shared memory is enabled the data will not be transferred to
+ /// the device and will be used from the host. Data will be added to the
+ /// mapping table to allow checks to happen even when in unified shared
+ /// memory.
TargetPointerResultTy getTargetPointer(
HDTTMapAccessorTy &HDTTMap, void *HstPtrBegin, void *HstPtrBase,
int64_t TgtPadding, int64_t Size, map_var_info_t HstPtrName,
diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp
index 8a2fe4620b39cbe..a8e93adcab5fc7c 100644
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -268,6 +268,27 @@ TargetPointerResultTy DeviceTy::getTargetPointer(
LR.TPR.getEntry()->holdRefCountToStr().c_str(), HoldRefCountAction,
(HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
LR.TPR.TargetPointer = (void *)Ptr;
+
+ // When the target pointer is retrieved again, then the condition for this
+ // branch can be true hence preventing the unified shared memory to be
+ // taken at all. This ensures that the IsHostPointer and IsPresent flags
+ // are correctly set even in that situation.
+ if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+ !HasCloseModifier && !LR.TPR.Flags.IsHostPointer) {
+ // This is a host pointer and is not present if the pointers match:
+ if (LR.TPR.getEntry()->TgtPtrBegin == LR.TPR.getEntry()->HstPtrBegin) {
+ LR.TPR.Flags.IsPresent = false;
+ LR.TPR.Flags.IsHostPointer = true;
+ }
+
+ // Catch the case where incoming HstPtrBegin is not consistent with the
+ // entry HstPtrBegin.
+ if (LR.TPR.Flags.IsHostPointer &&
+ ((uintptr_t)HstPtrBegin - LR.TPR.getEntry()->HstPtrBegin) != 0) {
+ assert(false &&
+ "Incoming HstPtrBegin different from entry HstPtrBegin");
+ }
+ }
} else if ((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && !IsImplicit) {
// Explicit extension of mapped data - not allowed.
MESSAGE("explicit extension not allowed: host address specified is " DPxMOD
@@ -289,13 +310,38 @@ TargetPointerResultTy DeviceTy::getTargetPointer(
// In addition to the mapping rules above, the close map modifier forces the
// mapping of the variable to the device.
if (Size) {
- DP("Return HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared "
- "memory\n",
- DPxPTR((uintptr_t)HstPtrBegin), Size);
- LR.TPR.Flags.IsPresent = false;
+ LR.TPR.Flags.IsNewEntry = true;
+ assert(TgtPadding == 0 && "TgtPadding must always be zero in USM mode");
+ uintptr_t TgtPtrBegin = (uintptr_t)HstPtrBegin + TgtPadding;
+ LR.TPR.setEntry(
+ HDTTMap
+ ->emplace(new HostDataToTargetTy(
+ (uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
+ (uintptr_t)HstPtrBegin + Size, (uintptr_t)HstPtrBegin,
+ TgtPtrBegin, HasHoldModifier, HstPtrName))
+ .first->HDTT);
+ INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID,
+ "Creating new map entry ONLY with HstPtrBase=" DPxMOD
+ ", HstPtrBegin=" DPxMOD ", TgtAllocBegin=" DPxMOD
+ ", TgtPtrBegin=" DPxMOD
+ ", Size=%ld, DynRefCount=%s, HoldRefCount=%s, Name=%s\n",
+ DPxPTR(HstPtrBase), DPxPTR(HstPtrBegin), DPxPTR(HstPtrBegin),
+ DPxPTR(TgtPtrBegin), Size,
+ LR.TPR.getEntry()->dynRefCountToStr().c_str(),
+ LR.TPR.getEntry()->holdRefCountToStr().c_str(),
+ (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
LR.TPR.Flags.IsHostPointer = true;
+
+ // The following assert should catch any case in which the pointers
+ // do not match to understand if this case can ever happen.
+ assert((uintptr_t)HstPtrBegin == TgtPtrBegin &&
+ "Pointers must always match");
+
+ // If the above assert is ever hit the following should be changed to =
+ // TgtPtrBegin
LR.TPR.TargetPointer = HstPtrBegin;
}
+ LR.TPR.Flags.IsPresent = false;
} else if (HasPresentModifier) {
DP("Mapping required by 'present' map type modifier does not exist for "
"HstPtrBegin=" DPxMOD ", Size=%" PRId64 "\n",
@@ -444,6 +490,29 @@ DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool UpdateRefCount,
LR.TPR.getEntry()->dynRefCountToStr().c_str(), DynRefCountAction,
LR.TPR.getEntry()->holdRefCountToStr().c_str(), HoldRefCountAction);
LR.TPR.TargetPointer = (void *)TP;
+
+ // If this entry is not marked as being host pointer (the way the
+ // implementation works today this is never true, mistake?) then we
+ // have to check if this is a host pointer or not. This is a host pointer
+ // if the host address matches the target address.
+ if ((PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) &&
+ !LR.TPR.Flags.IsHostPointer) {
+ // If addresses match it means that we are dealing with a host pointer
+ // which has to be marked as one and present flag reset:
+ if (LR.TPR.getEntry()->TgtPtrBegin == LR.TPR.getEntry()->HstPtrBegin) {
+ LR.TPR.Flags.IsPresent = false;
+ LR.TPR.Flags.IsHostPointer = true;
+ }
+
+ // We want to catch the case where (uintptr_t)HstPtrBegin and
+ // LR.TPR.getEntry()->HstPtrBegin are not the same when LR is a host
+ // pointer. This case should never happen.
+ if (LR.TPR.Flags.IsHostPointer &&
+ ((uintptr_t)HstPtrBegin - LR.TPR.getEntry()->HstPtrBegin) != 0) {
+ assert(false &&
+ "Incoming HstPtrBegin different from entry HstPtrBegin");
+ }
+ }
} else if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) {
// If the value isn't found in the mapping and unified shared memory
// is on then it means we have stumbled upon a value which we need to
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 65f2a49abc714ce..020b87e989f7230 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -788,7 +788,9 @@ postProcessingTargetDataEnd(DeviceTy *Device,
int Ret = OFFLOAD_SUCCESS;
for (auto &[HstPtrBegin, DataSize, ArgType, TPR] : EntriesInfo) {
- bool DelEntry = !TPR.isHostPointer();
+ // Delete entry from the mapping table even when we are dealing with a
+ // host pointer.
+ bool DelEntry = true;
// If the last element from the mapper (for end transfer args comes in
// reverse order), do not remove the partial entry, the parent struct still
@@ -846,10 +848,12 @@ postProcessingTargetDataEnd(DeviceTy *Device,
Ret = Device->eraseMapEntry(HDTTMap, Entry, DataSize);
// Entry is already remove from the map, we can unlock it now.
HDTTMap.destroy();
- Ret |= Device->deallocTgtPtrAndEntry(Entry, DataSize);
- if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Deallocating data from device failed.\n");
- break;
+ if (!TPR.Flags.IsHostPointer) {
+ Ret |= Device->deallocTgtPtrAndEntry(Entry, DataSize);
+ if (Ret != OFFLOAD_SUCCESS) {
+ REPORT("Deallocating data from device failed.\n");
+ break;
+ }
}
}
@@ -908,78 +912,92 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
Device.getTgtPtrBegin(HstPtrBegin, DataSize, UpdateRef, HasHoldModifier,
!IsImplicit, ForceDelete, /*FromDataEnd=*/true);
void *TgtPtrBegin = TPR.TargetPointer;
- if (!TPR.isPresent() && !TPR.isHostPointer() &&
- (DataSize || HasPresentModifier)) {
- DP("Mapping does not exist (%s)\n",
- (HasPresentModifier ? "'present' map type modifier" : "ignored"));
- if (HasPresentModifier) {
- // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 350 L10-13:
- // "If a map clause appears on a target, target data, target enter data
- // or target exit data construct with a present map-type-modifier then
- // on entry to the region if the corresponding list item does not appear
- // in the device data environment then an error occurs and the program
- // terminates."
- //
- // This should be an error upon entering an "omp target exit data". It
- // should not be an error upon exiting an "omp target data" or "omp
- // target". For "omp target data", Clang thus doesn't include present
- // modifiers for end calls. For "omp target", we have not found a valid
- // OpenMP program for which the error matters: it appears that, if a
- // program can guarantee that data is present at the beginning of an
- // "omp target" region so that there's no error there, that data is also
- // guaranteed to be present at the end.
- MESSAGE("device mapping required by 'present' map type modifier does "
- "not exist for host address " DPxMOD " (%" PRId64 " bytes)",
- DPxPTR(HstPtrBegin), DataSize);
- return OFFLOAD_FAIL;
- }
- } else {
- DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
- " - is%s last\n",
- DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsLast ? "" : " not"));
- }
-
- // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 351 L14-16:
- // "If the map clause appears on a target, target data, or target exit data
- // construct and a corresponding list item of the original list item is not
- // present in the device data environment on exit from the region then the
- // list item is ignored."
- if (!TPR.isPresent())
- continue;
- // Move data back to the host
- const bool HasAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
- const bool HasFrom = ArgTypes[I] & OMP_TGT_MAPTYPE_FROM;
- if (HasFrom && (HasAlways || TPR.Flags.IsLast) &&
- !TPR.Flags.IsHostPointer && DataSize != 0) {
- DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
- DataSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
-
- // Wait for any previous transfer if an event is present.
- if (void *Event = TPR.getEntry()->getEvent()) {
- if (Device.waitEvent(Event, AsyncInfo) != OFFLOAD_SUCCESS) {
- REPORT("Failed to wait for event " DPxMOD ".\n", DPxPTR(Event));
+ // Check if HstPtrBegin matches the State HstPtrBegin or if any HstPtrBegin
+ // values have been registered:
+ bool HostPointerMismatch = true;
+ if (TPR.getEntry())
+ HostPointerMismatch =
+ TPR.getEntry()->HstPtrBegin != (uintptr_t)HstPtrBegin;
+
+ if (!TPR.isHostPointer()) {
+ if (!TPR.isPresent() && (DataSize || HasPresentModifier)) {
+ DP("Mapping does not exist (%s)\n",
+ (HasPresentModifier ? "'present' map type modifier" : "ignored"));
+ if (HasPresentModifier) {
+ // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 350 L10-13:
+ // "If a map clause appears on a target, target data, target enter
+ // data or target exit data construct with a present map-type-modifier
+ // then on entry to the region if the corresponding list item does not
+ // appear in the device data environment then an error occurs and the
+ // program terminates."
+ //
+ // This should be an error upon entering an "omp target exit data". It
+ // should not be an error upon exiting an "omp target data" or "omp
+ // target". For "omp target data", Clang thus doesn't include present
+ // modifiers for end calls. For "omp target", we have not found a
+ // valid OpenMP program for which the error matters: it appears that,
+ // if a program can guarantee that data is present at the beginning of
+ // an "omp target" region so that there's no error there, that data is
+ // also guaranteed to be present at the end.
+ MESSAGE("device mapping required by 'present' map type modifier does "
+ "not exist for host address " DPxMOD " (%" PRId64 " bytes)",
+ DPxPTR(HstPtrBegin), DataSize);
return OFFLOAD_FAIL;
}
+ } else {
+ DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
+ " - is%s last\n",
+ DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsLast ? "" : " not"));
}
- Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, DataSize, AsyncInfo,
- TPR.getEntry());
- if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Copying data from device failed.\n");
- return OFFLOAD_FAIL;
- }
+ // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 351 L14-16:
+ // "If the map clause appears on a target, target data, or target exit
+ // data construct and a corresponding list item of the original list item
+ // is not present in the device data environment on exit from the region
+ // then the list item is ignored."
+ if (!TPR.isPresent())
+ continue;
- // As we are expecting to delete the entry the d2h copy might race
- // with another one that also tries to delete the entry. This happens
- // as the entry can be reused and the reuse might happen after the
- // copy-back was issued but before it completed. Since the reuse might
- // also copy-back a value we would race.
- if (TPR.Flags.IsLast) {
- if (TPR.getEntry()->addEventIfNecessary(Device, AsyncInfo) !=
- OFFLOAD_SUCCESS)
+ // Move data back to the host
+ const bool HasAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
+ const bool HasFrom = ArgTypes[I] & OMP_TGT_MAPTYPE_FROM;
+ if (HasFrom && (HasAlways || TPR.Flags.IsLast) && DataSize != 0) {
+ DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
+ DataSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
+
+ // Wait for any previous transfer if an event is present.
+ if (void *Event = TPR.getEntry()->getEvent()) {
+ if (Device.waitEvent(Event, AsyncInfo) != OFFLOAD_SUCCESS) {
+ REPORT("Failed to wait for event " DPxMOD ".\n", DPxPTR(Event));
+ return OFFLOAD_FAIL;
+ }
+ }
+
+ Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, DataSize, AsyncInfo,
+ TPR.getEntry());
+ if (Ret != OFFLOAD_SUCCESS) {
+ REPORT("Copying data from device failed.\n");
return OFFLOAD_FAIL;
+ }
+
+ // As we are expecting to delete the entry the d2h copy might race
+ // with another one that also tries to delete the entry. This happens
+ // as the entry can be reused and the reuse might happen after the
+ // copy-back was issued but before it completed. Since the reuse might
+ // also copy-back a value we would race.
+ if (TPR.Flags.IsLast) {
+ if (TPR.getEntry()->addEventIfNecessary(Device, AsyncInfo) !=
+ OFFLOAD_SUCCESS)
+ return OFFLOAD_FAIL;
+ }
}
+ } else {
+ // Some zero-sized arrays are not mapped or added to the mapping table so
+ // they do not need to be removed. These arrays are not part of the
+ // current entry.
+ if (DataSize == 0 && !TPR.isPresent() && HostPointerMismatch)
+ continue;
}
// Add pointer to the buffer for post-synchronize processing.
diff --git a/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_arrays.cpp b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_arrays.cpp
new file mode 100644
index 000000000000000..1c0257f9f246aa3
--- /dev/null
+++ b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_arrays.cpp
@@ -0,0 +1,48 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic && env HSA_XNACK=1 LIBOMPTARGET_INFO=-1 %libomptarget-run-generic 2>&1 | %fcheck-generic
+// clang-format on
+
+// REQUIRES: amdgcn-amd-amdhsa
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+#include <omp.h>
+
+#include <cassert>
+#include <iostream>
+
+#pragma omp requires unified_shared_memory
+
+int main(int argc, char *argv[]) {
+ int *v = (int *)malloc(sizeof(int) * 100);
+
+// clang-format off
+// CHECK: Entering OpenMP kernel at {{.*}} with 1 arguments:
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[V_HST_PTR_ADDR:0x.*]], HstPtrBegin=[[V_HST_PTR_ADDR]], TgtAllocBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=200, DynRefCount=1, HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=200, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks and 256 threads in Generic mode
+
+// CHECK: Mapping exists with HstPtrBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=200, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Removing map entry with HstPtrBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=200
+
+// CHECK: Entering OpenMP kernel at {{.*}} with 1 arguments:
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[V_HST_PTR_ADDR]], HstPtrBegin=[[V_HST_PTR_ADDR]], TgtAllocBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=280, DynRefCount=1, HoldRefCount=0
+// CHECK: Mapping exists with HstPtrBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=280, DynRefCount=1 (update suppressed), HoldRefCount=0
+// CHECK: Launching kernel __omp_offloading_{{.*}}_main_l{{.*}} with 1 blocks and 256 threads in Generic mode
+// CHECK: Mapping exists with HstPtrBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=280, DynRefCount=0 (decremented, delayed deletion), HoldRefCount=0
+// CHECK: Removing map entry with HstPtrBegin=[[V_HST_PTR_ADDR]], TgtPtrBegin=[[V_HST_PTR_ADDR]], Size=280
+// clang-format on
+#pragma omp target map(tofrom : v[ : 50])
+ { v[32] = 32; }
+
+#pragma omp target map(tofrom : v[ : 70])
+ { v[64] = 64; }
+
+ printf("v[32] = %d, v[64] = %d\n", v[32], v[64]);
+
+ free(v);
+
+ std::cout << "PASS\n";
+ return 0;
+}
+// CHECK: v[32] = 32, v[64] = 64
+// CHECK: PASS
diff --git a/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_close_enter_exit.cpp b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_close_enter_exit.cpp
new file mode 100644
index 000000000000000..8d3dd1f72200aa8
--- /dev/null
+++ b/openmp/libomptarget/test/unified_shared_memory/unified_map_checks_close_enter_exit.cpp
@@ -0,0 +1,222 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic && env HSA_XNACK=1 LIBOMPTARGET_INFO=-1 %libomptarget-run-generic 2>&1 | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+// REQUIRES: amdgcn-amd-amdhsa
+
+#include <omp.h>
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+#define N 1024
+
+int main(int argc, char *argv[]) {
+ void *host_alloc = nullptr, *device_alloc = nullptr;
+ int *a = (int *)malloc(N * sizeof(int));
+ int dev = omp_get_default_device();
+
+ // Init
+ for (int i = 0; i < N; ++i) {
+ a[i] = 10;
+ }
+ host_alloc = &a[0];
+
+ //
+ // map + target no close
+ //
+
+// clang-format off
+// CHECK: Entering OpenMP data region with being_mapper at {{.*}} with 2 arguments:
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[A_HST_PTR:0x.*]], HstPtrBegin=[[A_HST_PTR]], TgtAllocBegin=[[A_HST_PTR]], TgtPtrBegin=[[A_HST_PTR]], Size=4096, DynRefCount=1, HoldRefCount=0
+// CHECK: Creating new map entry ONLY with HstPtrBase=[[DEVICE_ALLOC_HST_PTR:0x.*]], HstPtrBegin=[[DEVICE_ALLOC_HST_PTR]], TgtAllocBegin=[[DEVICE_ALLOC_HST_PTR]], TgtPtrBegin=[[DEVICE_ALLOC_HST_PTR]], Size=8, DynRefCount=1, HoldRefCount=0
+// CHECK: OpenMP Host-Device pointer mappings after block
+// CHECK: Host Ptr
+// CHECK: [[A_HST_PTR]]
+// CHECK: [[DEVICE_ALLOC_HST_PTR]]
+// clang-format on
+#pragma omp target data map(tofrom : a[ : N]) map(tofrom : device_alloc)
+ {
+// clan...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/69005
More information about the Openmp-commits
mailing list