[Libomptarget] Change device free routines to accept the allocation kind

[jhuber6]

Summary

Previous support for device memory allocators used a single free
routine and did not provide the original kind of the allocation. This is
problematic as some of these memory types required different handling.
Previously this was worked around using a map in runtime to record the
original kind of each pointer. Instead, this patch introduces new free
routines similar to the existing allocation routines. This allows us to
avoid a map traversal every time we free a device pointer.

The only interfaces defined by the standard are omp_target_alloc and
omp_target_free, these do not take a kind as omp_alloc does. The
standard dictates the following:

"The omp_target_alloc routine returns a device pointer that references
the device address of a storage location of size bytes. The storage
location is dynamically allocated in the device data environment of the
device specified by device_num."

Which suggests that these routines only allocate the default device
memory for the kind. So this has been changed to reflect this. This
change is somewhat breaking if users were using omp_target_free as
previously shown in the tests.

[JonChesterfield]

I think changing the API of __tgt_rtl_data_delete is good. I like the map lookup disappearing from the cuda plugin. It'll mean mixing plugins and libomptargets from different toolchains will tend to segfault. They're not certain to be semantically interchangeable even if we keep the function name and type the same though.

@jdoerfert / @ronlieb / @RaviNarayanaswamy any of you aware of a requirement that forces slower/clumsy/backwards-compatible here, aka renaming the entry point to __tgt_rtl_data_delete_v2 or similar?

In particular changing the llvm_ prefixed function is definitely OK, anyone using llvm_ prefixed openmp functions signed up for semantics changing en route to standardisation.

Patch itself looks good to me, thanks for writing it.

[JonChesterfield]

openmp/libomptarget/include/device.h

410

Might be better to avoid the default argument here in favour of explicit at call site since that'll help pair up allocation kind when reading the implementation. Probably as a follow up patch to this one since the verbosity would detract from the review.

[jhuber6]

openmp/libomptarget/include/device.h

410

The alloc above does the same default argument so it made sense to just copy it.

[tianshilei1992]

I think we'd better keep __tgt_rtl_data_delete unchanged, and add another interface. AFAIK we require libomptarget to be aligned with clang, but don't force the plugins.

[jhuber6]

In D133053#3764082, @tianshilei1992 wrote:

I think we'd better keep __tgt_rtl_data_delete unchanged, and add another interface. AFAIK we require libomptarget to be aligned with clang, but don't force the plugins.

I figured it was okay to change since I don't recall if we actually provide backwards compatibility for the plugins.

[tianshilei1992]

In D133053#3764130, @jhuber6 wrote:

In D133053#3764082, @tianshilei1992 wrote:

I think we'd better keep __tgt_rtl_data_delete unchanged, and add another interface. AFAIK we require libomptarget to be aligned with clang, but don't force the plugins.

I figured it was okay to change since I don't recall if we actually provide backwards compatibility for the plugins.

Is it stated somewhere? I had a lot of trouble about compatibility before when I tried to change the interface. I'm okay with breaking the backward compatibility, as long as we already tell users. :-)

[tianshilei1992]

FWIW, if necessary, we can say LLVM 15 is the last version for backward compatibility.

[jhuber6]

In D133053#3764147, @tianshilei1992 wrote:

Is it stated somewhere? I had a lot of trouble about compatibility before when I tried to change the interface. I'm okay with breaking the backward compatibility, as long as we already tell users. :-)

It's not too difficult to just make a new interface and have the old one call the new one with the default. Probably better safe than sorry so I may as well add it if it's a problem.

[JonChesterfield]

In D133053#3764152, @jhuber6 wrote:

In D133053#3764147, @tianshilei1992 wrote:

Is it stated somewhere? I had a lot of trouble about compatibility before when I tried to change the interface. I'm okay with breaking the backward compatibility, as long as we already tell users. :-)

It's not too difficult to just make a new interface and have the old one call the new one with the default. Probably better safe than sorry so I may as well add it if it's a problem.

It's simple to do but means we're stuck with all the cruft in the cuda plugin as we won't know which form of delete will be called.

It'll be a really easy patch for downstream plugins to adopt (Intel has one iiuc?). The other use case is end users who have their plugin that they use with unrelated clang who I'm not sure exist. In either case, these toolchains will have already manually changed their libomptarget to get it to load their plugin already as we don't have any form of automatic discovery in place.

I think we should change the calling convention and remove all the map lookup stuff from the cuda plugin unless we have a compelling reason not to, where miscellaneous concern about slightly breaking unknown users that are reliant on internal interfaces doesn't seem compelling.

[tianshilei1992]

In D133053#3764170, @JonChesterfield wrote:

In D133053#3764152, @jhuber6 wrote:

In D133053#3764147, @tianshilei1992 wrote:

Is it stated somewhere? I had a lot of trouble about compatibility before when I tried to change the interface. I'm okay with breaking the backward compatibility, as long as we already tell users. :-)

It's not too difficult to just make a new interface and have the old one call the new one with the default. Probably better safe than sorry so I may as well add it if it's a problem.

It's simple to do but means we're stuck with all the cruft in the cuda plugin as we won't know which form of delete will be called.

It'll be a really easy patch for downstream plugins to adopt (Intel has one iiuc?). The other use case is end users who have their plugin that they use with unrelated clang who I'm not sure exist. In either case, these toolchains will have already manually changed their libomptarget to get it to load their plugin already as we don't have any form of automatic discovery in place.

I think we should change the calling convention and remove all the map lookup stuff from the cuda plugin unless we have a compelling reason not to, where miscellaneous concern about slightly breaking unknown users that are reliant on internal interfaces doesn't seem compelling.

You can still use _v2 or whatever in libomptarget. For libomptarget, there is no old interface then. Just keep the old one and all the logic for compatibility in the plugin. There is no confusion.

[jhuber6]

In D133053#3764266, @tianshilei1992 wrote:

You can still use _v2 or whatever in libomptarget. For libomptarget, there is no old interface then. Just keep the old one and all the logic for compatibility in the plugin. There is no confusion.

The potential change here is users who used omp_target_free to free non-device pointers as this worked before. We technically export the plugin routines as they need to be visible to libomptarget but I don't know if we expect any users to call them directly.

[tianshilei1992]

In D133053#3764270, @jhuber6 wrote:

In D133053#3764266, @tianshilei1992 wrote:

You can still use _v2 or whatever in libomptarget. For libomptarget, there is no old interface then. Just keep the old one and all the logic for compatibility in the plugin. There is no confusion.

The potential change here is users who used omp_target_free to free non-device pointers as this worked before. We technically export the plugin routines as they need to be visible to libomptarget but I don't know if we expect any users to call them directly.

The "user" here is old libomptarget.

[jhuber6]

In D133053#3764276, @tianshilei1992 wrote:

In D133053#3764270, @jhuber6 wrote:

In D133053#3764266, @tianshilei1992 wrote:

You can still use _v2 or whatever in libomptarget. For libomptarget, there is no old interface then. Just keep the old one and all the logic for compatibility in the plugin. There is no confusion.

The potential change here is users who used omp_target_free to free non-device pointers as this worked before. We technically export the plugin routines as they need to be visible to libomptarget but I don't know if we expect any users to call them directly.

The "user" here is old libomptarget.

It would be an incredibly weird setup to have libomptarget load a plugin from a different install right? We install them to the same directory so it would mean that the user manually edited libomptarget after the install.

[tianshilei1992]

In D133053#3764277, @jhuber6 wrote:

In D133053#3764276, @tianshilei1992 wrote:

In D133053#3764270, @jhuber6 wrote:

In D133053#3764266, @tianshilei1992 wrote:

You can still use _v2 or whatever in libomptarget. For libomptarget, there is no old interface then. Just keep the old one and all the logic for compatibility in the plugin. There is no confusion.

The potential change here is users who used omp_target_free to free non-device pointers as this worked before. We technically export the plugin routines as they need to be visible to libomptarget but I don't know if we expect any users to call them directly.

The "user" here is old libomptarget.

It would be an incredibly weird setup to have libomptarget load a plugin from a different install right? We install them to the same directory so it would mean that the user manually edited libomptarget after the install.

We don't force the plugins to be loaded from the same dir as libomptarget right?

Again, like I said, we can claim LLVM 15 is the last version to guarantee backward compatibility for all the components we require. For later version all components have to be from the same build. We just need to make a consensus in the meeting, document it, and do whatever we want to change the APIs in the future.

[JonChesterfield]

In D133053#3764266, @tianshilei1992 wrote:

You can still use _v2 or whatever in libomptarget. For libomptarget, there is no old interface then. Just keep the old one and all the logic for compatibility in the plugin. There is no confusion.

Then we'd have to keep the spurious hashtable inserts in alloc, and consider fixing the pointer leak in the current implementation (the table this deletes grows without bound). I.e. to keep things working for people that we don't have reason to believe exist, the cuda plugin must remain slower than it could be.

I can make my peace with that, but it looks like discarding this worthwhile patch, not applying it and taking the extra complexity without the benefit.

[jhuber6]

In D133053#3764298, @tianshilei1992 wrote:

In D133053#3764277, @jhuber6 wrote:

In D133053#3764276, @tianshilei1992 wrote:

In D133053#3764270, @jhuber6 wrote:

In D133053#3764266, @tianshilei1992 wrote:

You can still use _v2 or whatever in libomptarget. For libomptarget, there is no old interface then. Just keep the old one and all the logic for compatibility in the plugin. There is no confusion.

The potential change here is users who used omp_target_free to free non-device pointers as this worked before. We technically export the plugin routines as they need to be visible to libomptarget but I don't know if we expect any users to call them directly.

The "user" here is old libomptarget.

It would be an incredibly weird setup to have libomptarget load a plugin from a different install right? We install them to the same directory so it would mean that the user manually edited libomptarget after the install.

We don't force the plugins to be loaded from the same dir as libomptarget right?

Again, like I said, we can claim LLVM 15 is the last version to guarantee backward compatibility for all the components we require. For later version all components have to be from the same build. We just need to make a consensus in the meeting, document it, and do whatever we want to change the APIs in the future.

When we load the plugin we just pass the name like libomptarget.rtl.cuda.so so it goes to the linker's search path which is going to search using the default path. I suppose since we use -rpath it's possible these could get mixed up.

[jhuber6]

In D133053#3764298, @tianshilei1992 wrote

Again, like I said, we can claim LLVM 15 is the last version to guarantee backward compatibility for all the components we require. For later version all components have to be from the same build. We just need to make a consensus in the meeting, document it, and do whatever we want to change the APIs in the future.

This would be nice considering a lot of the legacy we carry around. The fact that LLVM 15 released with proper versioned .so objects, e.g. libomptarget.so.15 helps us to suggest that. It would definitely need to be discussed.

[tianshilei1992]

In D133053#3764314, @JonChesterfield wrote:

In D133053#3764266, @tianshilei1992 wrote:

You can still use _v2 or whatever in libomptarget. For libomptarget, there is no old interface then. Just keep the old one and all the logic for compatibility in the plugin. There is no confusion.

Then we'd have to keep the spurious hashtable inserts in alloc, and consider fixing the pointer leak in the current implementation (the table this deletes grows without bound). I.e. to keep things working for people that we don't have reason to believe exist, the cuda plugin must remain slower than it could be.

FWIW, we don't have to fix issues for old interface. That's the whole point of deprecating the old interface.

I can make my peace with that, but it looks like discarding this worthwhile patch, not applying it and taking the extra complexity without the benefit.

Of course we will not discard this patch. We just need to find the right way to do it. As long as we make agreement on compatibility, we are good to go. That's a single cut, which can save a lot future arguments.

[jhuber6]

We discussed how the plugin interfaces don't need backwards compatibility in the meeting today. @tianshilei1992 is this good to land?

[tianshilei1992]

Yup, let's go!

openmp/libomptarget/include/device.h

struct DeviceTy {
   /// allocated target data will be associated with. If it is unknown, the
   /// default value of \p HstPtr is nullptr. Note: this function doesn't do
   /// pointer association. Actually, all the __tgt_rtl_data_alloc
   /// implementations ignore \p HstPtr. \p Kind dictates what allocator should
   /// be used (host, shared, device).
   void *allocData(int64_t Size, void *HstPtr = nullptr,
                   int32_t Kind = TARGET_ALLOC_DEFAULT);
   /// Deallocates memory which \p TgtPtrBegin points at and returns
-  /// OFFLOAD_SUCCESS/OFFLOAD_FAIL when succeeds/fails.
-  int32_t deleteData(void *TgtPtrBegin);
+  /// OFFLOAD_SUCCESS/OFFLOAD_FAIL when succeeds/fails. p Kind dictates what
+  /// allocator should be used (host, shared, device).
+  int32_t deleteData(void *TgtPtrBegin, int32_t Kind = TARGET_ALLOC_DEFAULT);
 
   // Data transfer. When AsyncInfo is nullptr, the transfer will be
   // synchronous.
   // Copy data from host to device
   int32_t submitData(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size,
                      AsyncInfoTy &AsyncInfo);
   // Copy data from device back to host
   int32_t retrieveData(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size,

openmp/libomptarget/include/omptarget.h

 int omp_target_disassociate_ptr(const void *HostPtr, int DeviceNum);
 
 /// Explicit target memory allocators
 /// Using the llvm_ prefix until they become part of the OpenMP standard.
 void *llvm_omp_target_alloc_device(size_t Size, int DeviceNum);
 void *llvm_omp_target_alloc_host(size_t Size, int DeviceNum);
 void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
 
+/// Explicit target memory deallocators
+/// Using the llvm_ prefix until they become part of the OpenMP standard.
+void llvm_omp_target_free_device(void *DevicePtr, int DeviceNum);
+void llvm_omp_target_free_host(void *DevicePtr, int DeviceNum);
+void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
+
 /// Dummy target so we have a symbol for generating host fallback.
 void *llvm_omp_target_dynamic_shared_alloc();
 
 /// add the clauses of the requires directives in a given file
 void __tgt_register_requires(int64_t Flags);
 
 /// adds a target shared library to the target execution image
 void __tgt_register_lib(__tgt_bin_desc *Desc);

openmp/libomptarget/include/omptargetplugin.h

int32_t __tgt_rtl_data_exchange(int32_t SrcID, void *SrcPtr, int32_t DstID,
                                 void *DstPtr, int64_t Size);
 
 // Asynchronous version of __tgt_rtl_data_exchange
 int32_t __tgt_rtl_data_exchange_async(int32_t SrcID, void *SrcPtr,
                                       int32_t DesID, void *DstPtr, int64_t Size,
                                       __tgt_async_info *AsyncInfo);
 
 // De-allocate the data referenced by target ptr on the device. In case of
-// success, return zero. Otherwise, return an error code.
-int32_t __tgt_rtl_data_delete(int32_t ID, void *TargetPtr);
+// success, return zero. Otherwise, return an error code. Kind dictates what
+// allocator to use (e.g. shared, host, device).
+int32_t __tgt_rtl_data_delete(int32_t ID, void *TargetPtr, int32_t Kind);
 
 // Transfer control to the offloaded entry Entry on the target device.
 // Args and Offsets are arrays of NumArgs size of target addresses and
 // offsets. An offset should be added to the target address before passing it
 // to the outlined function on device side. If AsyncInfo is nullptr, it is
 // synchronous; otherwise it is asynchronous. However, AsyncInfo may be
 // ignored on some platforms, like x86_64. In that case, it is synchronous. In
 // case of success, return zero. Otherwise, return an error code.

openmp/libomptarget/include/rtl.h

struct RTLInfoTy {
   typedef int32_t(data_submit_async_ty)(int32_t, void *, void *, int64_t,
                                         __tgt_async_info *);
   typedef int32_t(data_retrieve_ty)(int32_t, void *, void *, int64_t);
   typedef int32_t(data_retrieve_async_ty)(int32_t, void *, void *, int64_t,
                                           __tgt_async_info *);
   typedef int32_t(data_exchange_ty)(int32_t, void *, int32_t, void *, int64_t);
   typedef int32_t(data_exchange_async_ty)(int32_t, void *, int32_t, void *,
                                           int64_t, __tgt_async_info *);
-  typedef int32_t(data_delete_ty)(int32_t, void *);
+  typedef int32_t(data_delete_ty)(int32_t, void *, int32_t);
   typedef int32_t(run_region_ty)(int32_t, void *, void **, ptrdiff_t *,
                                  int32_t);
   typedef int32_t(run_region_async_ty)(int32_t, void *, void **, ptrdiff_t *,
                                        int32_t, __tgt_async_info *);
   typedef int32_t(run_team_region_ty)(int32_t, void *, void **, ptrdiff_t *,
                                       int32_t, int32_t, int32_t, uint64_t);
   typedef int32_t(run_team_region_async_ty)(int32_t, void *, void **,
                                             ptrdiff_t *, int32_t, int32_t,

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

int32_t __tgt_rtl_data_retrieve_async(int DeviceId, void *HstPtr, void *TgtPtr,
                                       int64_t Size,
                                       __tgt_async_info *AsyncInfo) {
   assert(AsyncInfo && "AsyncInfo is nullptr");
   assert(DeviceId < DeviceInfo().NumberOfDevices && "Device ID too large");
   initAsyncInfo(AsyncInfo);
   return dataRetrieve(DeviceId, HstPtr, TgtPtr, Size, AsyncInfo);
 }
 
-int32_t __tgt_rtl_data_delete(int DeviceId, void *TgtPtr) {
+int32_t __tgt_rtl_data_delete(int DeviceId, void *TgtPtr, int32_t) {
   assert(DeviceId < DeviceInfo().NumberOfDevices && "Device ID too large");
   // HSA can free pointers allocated from different types of memory pool.
   hsa_status_t Err;
   DP("Tgt free data (tgt:%016llx).\n", (long long unsigned)(Elf64_Addr)TgtPtr);
   Err = core::Runtime::Memfree(TgtPtr);
   if (Err != HSA_STATUS_SUCCESS) {
     DP("Error when freeing CUDA memory\n");
     return OFFLOAD_FAIL;

openmp/libomptarget/plugins/common/MemoryManager/MemoryManager.h

 
 /// Base class of per-device allocator.
 class DeviceAllocatorTy {
 public:
   virtual ~DeviceAllocatorTy() = default;
 
   /// Allocate a memory of size \p Size . \p HstPtr is used to assist the
   /// allocation.
-  virtual void *allocate(size_t Size, void *HstPtr, TargetAllocTy Kind) = 0;
+  virtual void *allocate(size_t Size, void *HstPtr,
+                         TargetAllocTy Kind = TARGET_ALLOC_DEFAULT) = 0;
 
   /// Delete the pointer \p TgtPtr on the device
-  virtual int free(void *TgtPtr) = 0;
+  virtual int free(void *TgtPtr, TargetAllocTy Kind = TARGET_ALLOC_DEFAULT) = 0;
 };
 
 /// Class of memory manager. The memory manager is per-device by using
 /// per-device allocator. Therefore, each plugin using memory manager should
 /// have an allocator for each device.
 class MemoryManagerTy {
   static constexpr const size_t BucketSize[] = {
       0,       1U << 2, 1U << 3,  1U << 4,  1U << 5,  1U << 6, 1U << 7,

openmp/libomptarget/plugins/cuda/src/rtl.cpp

 #include <algorithm>
 #include <cassert>
 #include <cstddef>
 #include <cuda.h>
 #include <list>
 #include <memory>
 #include <mutex>
 #include <string>
-#include <unordered_map>
 #include <vector>
 
 #include "Debug.h"
 #include "DeviceEnvironment.h"
 #include "omptarget.h"
 #include "omptargetplugin.h"
 
 #define TARGET_NAME CUDA
class DeviceRTLTy {
 
   enum class PeerAccessState : uint8_t { Unkown, Yes, No };
   std::vector<std::vector<PeerAccessState>> PeerAccessMatrix;
   std::mutex PeerAccessMatrixLock;
 
   /// A class responsible for interacting with device native runtime library to
   /// allocate and free memory.
   class CUDADeviceAllocatorTy : public DeviceAllocatorTy {
-    std::unordered_map<void *, TargetAllocTy> HostPinnedAllocs;
-
   public:
     void *allocate(size_t Size, void *, TargetAllocTy Kind) override {
       if (Size == 0)
         return nullptr;
 
       void *MemAlloc = nullptr;
       CUresult Err;
       switch (Kind) {
       case TARGET_ALLOC_DEFAULT:
       case TARGET_ALLOC_DEVICE:
         CUdeviceptr DevicePtr;
         Err = cuMemAlloc(&DevicePtr, Size);
         MemAlloc = (void *)DevicePtr;
         if (!checkResult(Err, "Error returned from cuMemAlloc\n"))
           return nullptr;
         break;
       case TARGET_ALLOC_HOST:
         void *HostPtr;
         Err = cuMemAllocHost(&HostPtr, Size);
         MemAlloc = HostPtr;
         if (!checkResult(Err, "Error returned from cuMemAllocHost\n"))
           return nullptr;
-        HostPinnedAllocs[MemAlloc] = Kind;
         break;
       case TARGET_ALLOC_SHARED:
         CUdeviceptr SharedPtr;
         Err = cuMemAllocManaged(&SharedPtr, Size, CU_MEM_ATTACH_GLOBAL);
         MemAlloc = (void *)SharedPtr;
         if (!checkResult(Err, "Error returned from cuMemAllocManaged\n"))
           return nullptr;
         break;
       }
 
       return MemAlloc;
     }
 
-    int free(void *TgtPtr) override {
+    int free(void *TgtPtr, TargetAllocTy Kind) override {
       CUresult Err;
       // Host pinned memory must be freed differently.
-      TargetAllocTy Kind =
-          (HostPinnedAllocs.find(TgtPtr) == HostPinnedAllocs.end())
-              ? TARGET_ALLOC_DEFAULT
-              : TARGET_ALLOC_HOST;
       switch (Kind) {
       case TARGET_ALLOC_DEFAULT:
       case TARGET_ALLOC_DEVICE:
       case TARGET_ALLOC_SHARED:
         Err = cuMemFree((CUdeviceptr)TgtPtr);
         if (!checkResult(Err, "Error returned from cuMemFree\n"))
           return OFFLOAD_FAIL;
         break;
if (SrcDevId != DstDevId) {
         return memcpyDtoD(SrcPtr, DstPtr, Size, Stream);
       }
       }
     }
 
     return memcpyDtoD(SrcPtr, DstPtr, Size, Stream);
   }
 
-  int dataDelete(const int DeviceId, void *TgtPtr) {
+  int dataDelete(const int DeviceId, void *TgtPtr, TargetAllocTy Kind) {
+    switch (Kind) {
+    case TARGET_ALLOC_DEFAULT:
+    case TARGET_ALLOC_DEVICE:
       if (UseMemoryManager)
         return MemoryManagers[DeviceId]->free(TgtPtr);
+      else
+        return DeviceAllocators[DeviceId].free(TgtPtr, Kind);
+    case TARGET_ALLOC_HOST:
+    case TARGET_ALLOC_SHARED:
+      return DeviceAllocators[DeviceId].free(TgtPtr, Kind);
+    }
+
+    REPORT("Invalid target data allocation kind or requested allocator not "
+           "implemented yet\n");
 
-    return DeviceAllocators[DeviceId].free(TgtPtr);
+    return OFFLOAD_FAIL;
   }
 
   int runTargetTeamRegion(const int DeviceId, void *TgtEntryPtr, void **TgtArgs,
                           ptrdiff_t *TgtOffsets, const int ArgNum,
                           const int TeamNum, const int ThreadLimit,
                           const unsigned int LoopTripCount,
                           __tgt_async_info *AsyncInfo) const {
     // All args are references.
int32_t __tgt_rtl_data_exchange(int32_t SrcDevId, void *SrcPtr,
   const int32_t Rc = __tgt_rtl_data_exchange_async(SrcDevId, SrcPtr, DstDevId,
                                                    DstPtr, Size, &AsyncInfo);
   if (Rc != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
   return __tgt_rtl_synchronize(SrcDevId, &AsyncInfo);
 }
 
-int32_t __tgt_rtl_data_delete(int32_t DeviceId, void *TgtPtr) {
+int32_t __tgt_rtl_data_delete(int32_t DeviceId, void *TgtPtr, int32_t Kind) {
   assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");
 
   if (DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
-  return DeviceRTL.dataDelete(DeviceId, TgtPtr);
+  return DeviceRTL.dataDelete(DeviceId, TgtPtr, (TargetAllocTy)Kind);
 }
 
 int32_t __tgt_rtl_run_target_team_region(int32_t DeviceId, void *TgtEntryPtr,
                                          void **TgtArgs, ptrdiff_t *TgtOffsets,
                                          int32_t ArgNum, int32_t TeamNum,
                                          int32_t ThreadLimit,
                                          uint64_t LoopTripcount) {
   assert(DeviceRTL.isValidDeviceId(DeviceId) && "device_id is invalid");

openmp/libomptarget/plugins/generic-elf-64bit/src/rtl.cpp

 }
 
 int32_t __tgt_rtl_data_retrieve(int32_t DeviceId, void *HstPtr, void *TgtPtr,
                                 int64_t Size) {
   memcpy(HstPtr, TgtPtr, Size);
   return OFFLOAD_SUCCESS;
 }
 
-int32_t __tgt_rtl_data_delete(int32_t DeviceId, void *TgtPtr) {
+int32_t __tgt_rtl_data_delete(int32_t DeviceId, void *TgtPtr, int32_t) {
   free(TgtPtr);
   return OFFLOAD_SUCCESS;
 }
 
 int32_t __tgt_rtl_run_target_team_region(int32_t DeviceId, void *TgtEntryPtr,
                                          void **TgtArgs, ptrdiff_t *TgtOffsets,
                                          int32_t ArgNum, int32_t TeamNum,
                                          int32_t ThreadLimit,

openmp/libomptarget/plugins/remote/src/rtl.cpp

int32_t __tgt_rtl_data_submit(int32_t DeviceId, void *TgtPtr, void *HstPtr,
   return Manager->dataSubmit(DeviceId, TgtPtr, HstPtr, Size);
 }
 
 int32_t __tgt_rtl_data_retrieve(int32_t DeviceId, void *HstPtr, void *TgtPtr,
                                 int64_t Size) {
   return Manager->dataRetrieve(DeviceId, HstPtr, TgtPtr, Size);
 }
 
-int32_t __tgt_rtl_data_delete(int32_t DeviceId, void *TgtPtr) {
+int32_t __tgt_rtl_data_delete(int32_t DeviceId, void *TgtPtr, int32_t) {
   return Manager->dataDelete(DeviceId, TgtPtr);
 }
 
 int32_t __tgt_rtl_data_exchange(int32_t SrcDevId, void *SrcPtr,
                                 int32_t DstDevId, void *DstPtr, int64_t Size) {
   return Manager->dataExchange(SrcDevId, SrcPtr, DstDevId, DstPtr, Size);
 }
 

openmp/libomptarget/plugins/ve/src/rtl.cpp

if (ret != 0) {
     DP("veo_read_mem() failed with error code %d\n", ret);
     return OFFLOAD_FAIL;
   }
   return OFFLOAD_SUCCESS;
 }
 
 // De-allocate the data referenced by target ptr on the device. In case of
 // success, return zero. Otherwise, return an error code.
-int32_t __tgt_rtl_data_delete(int32_t ID, void *TargetPtr) {
+int32_t __tgt_rtl_data_delete(int32_t ID, void *TargetPtr, int32_t) {
   int ret = veo_free_mem(DeviceInfo.ProcHandles[ID], (uint64_t)TargetPtr);
 
   if (ret != 0) {
     DP("veo_free_mem() failed with error code %d\n", ret);
     return OFFLOAD_FAIL;
   }
   return OFFLOAD_SUCCESS;
 }

openmp/libomptarget/src/api.cpp

 EXTERN void *llvm_omp_target_alloc_host(size_t Size, int DeviceNum) {
   return targetAllocExplicit(Size, DeviceNum, TARGET_ALLOC_HOST, __func__);
 }
 
 EXTERN void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum) {
   return targetAllocExplicit(Size, DeviceNum, TARGET_ALLOC_SHARED, __func__);
 }
 
-EXTERN void *llvm_omp_target_dynamic_shared_alloc() { return nullptr; }
-EXTERN void *llvm_omp_get_dynamic_shared() { return nullptr; }
-
-EXTERN void omp_target_free(void *DevicePtr, int DeviceNum) {
-  TIMESCOPE();
-  DP("Call to omp_target_free for device %d and address " DPxMOD "\n",
-     DeviceNum, DPxPTR(DevicePtr));
-
-  if (!DevicePtr) {
-    DP("Call to omp_target_free with NULL ptr\n");
-    return;
-  }
+EXTERN void omp_target_free(void *Ptr, int DeviceNum) {
+  return targetFreeExplicit(Ptr, DeviceNum, TARGET_ALLOC_DEFAULT, __func__);
 
-  if (DeviceNum == omp_get_initial_device()) {
-    free(DevicePtr);
-    DP("omp_target_free deallocated host ptr\n");
-    return;
-  }
+EXTERN void llvm_omp_target_free_device(void *Ptr, int DeviceNum) {
+  return targetFreeExplicit(Ptr, DeviceNum, TARGET_ALLOC_DEVICE, __func__);
 
-  if (!deviceIsReady(DeviceNum)) {
-    DP("omp_target_free returns, nothing to do\n");
-    return;
-  }
+EXTERN void llvm_omp_target_free_host(void *Ptr, int DeviceNum) {
+  return targetFreeExplicit(Ptr, DeviceNum, TARGET_ALLOC_HOST, __func__);
 
-  PM->Devices[DeviceNum]->deleteData(DevicePtr);
-  DP("omp_target_free deallocated device ptr\n");
+EXTERN void llvm_omp_target_free_shared(void *Ptre, int DeviceNum) {
+  return targetFreeExplicit(Ptre, DeviceNum, TARGET_ALLOC_SHARED, __func__);
 }
 
+EXTERN void *llvm_omp_target_dynamic_shared_alloc() { return nullptr; }
+EXTERN void *llvm_omp_get_dynamic_shared() { return nullptr; }
+
 EXTERN int omp_target_is_present(const void *Ptr, int DeviceNum) {
   TIMESCOPE();
   DP("Call to omp_target_is_present for device %d and address " DPxMOD "\n",
      DeviceNum, DPxPTR(Ptr));
 
   if (!Ptr) {
     DP("Call to omp_target_is_present with NULL ptr, returning false\n");
     return false;

openmp/libomptarget/src/device.cpp

__tgt_target_table *DeviceTy::loadBinary(void *Img) {
   std::lock_guard<decltype(RTL->Mtx)> LG(RTL->Mtx);
   return RTL->load_binary(RTLDeviceID, Img);
 }
 
 void *DeviceTy::allocData(int64_t Size, void *HstPtr, int32_t Kind) {
   return RTL->data_alloc(RTLDeviceID, Size, HstPtr, Kind);
 }
 
-int32_t DeviceTy::deleteData(void *TgtPtrBegin) {
-  return RTL->data_delete(RTLDeviceID, TgtPtrBegin);
+int32_t DeviceTy::deleteData(void *TgtPtrBegin, int32_t Kind) {
+  return RTL->data_delete(RTLDeviceID, TgtPtrBegin, Kind);
 }
 
 // Submit data to device
 int32_t DeviceTy::submitData(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size,
                              AsyncInfoTy &AsyncInfo) {
   if (getInfoLevel() & OMP_INFOTYPE_DATA_TRANSFER) {
     HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor();
     LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size);

openmp/libomptarget/src/exports

global:
     omp_target_is_present;
     omp_target_memcpy;
     omp_target_memcpy_rect;
     omp_target_associate_ptr;
     omp_target_disassociate_ptr;
     llvm_omp_target_alloc_host;
     llvm_omp_target_alloc_shared;
     llvm_omp_target_alloc_device;
+    llvm_omp_target_free_host;
+    llvm_omp_target_free_shared;
+    llvm_omp_target_free_device;
     llvm_omp_target_dynamic_shared_alloc;
     __tgt_set_info_flag;
     __tgt_print_device_info;
     omp_get_interop_ptr;
     omp_get_interop_str;
     omp_get_interop_int;
     omp_get_interop_name;
     omp_get_interop_type_desc;
     omp_get_interop_rc_desc;
     __tgt_interop_init;
     __tgt_interop_use;
     __tgt_interop_destroy;
   local:
     *;
 };
 

openmp/libomptarget/src/omptarget.cpp

void *targetAllocExplicit(size_t Size, int DeviceNum, int Kind,
   }
 
   DeviceTy &Device = *PM->Devices[DeviceNum];
   Rc = Device.allocData(Size, nullptr, Kind);
   DP("%s returns device ptr " DPxMOD "\n", Name, DPxPTR(Rc));
   return Rc;
 }
 
+void targetFreeExplicit(void *DevicePtr, int DeviceNum, int Kind,
+                        const char *Name) {
+  TIMESCOPE();
+  DP("Call to %s for device %d and address " DPxMOD "\n", Name, DeviceNum,
+     DPxPTR(DevicePtr));
+
+  if (!DevicePtr) {
+    DP("Call to %s with NULL ptr\n", Name);
+    return;
+  }
+
+  if (DeviceNum == omp_get_initial_device()) {
+    free(DevicePtr);
+    DP("%s deallocated host ptr\n", Name);
+    return;
+  }
+
+  if (!deviceIsReady(DeviceNum)) {
+    DP("%s returns, nothing to do\n", Name);
+    return;
+  }
+
+  PM->Devices[DeviceNum]->deleteData(DevicePtr, Kind);
+  DP("omp_target_free deallocated device ptr\n");
+}
+
 /// Call the user-defined mapper function followed by the appropriate
 // targetData* function (targetData{Begin,End,Update}).
 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) {
   TIMESCOPE_WITH_IDENT(Loc);
   DP("Calling the mapper function " DPxMOD "\n", DPxPTR(ArgMapper));

openmp/libomptarget/src/private.h

extern int target(ident_t *Loc, DeviceTy &Device, void *HostPtr, int32_t ArgNum,
                   void **ArgMappers, int32_t TeamNum, int32_t ThreadLimit,
                   uint64_t Tripcount, int IsTeamConstruct,
                   AsyncInfoTy &AsyncInfo);
 
 extern void handleTargetOutcome(bool Success, ident_t *Loc);
 extern bool checkDeviceAndCtors(int64_t &DeviceID, ident_t *Loc);
 extern void *targetAllocExplicit(size_t Size, int DeviceNum, int Kind,
                                  const char *Name);
+extern void targetFreeExplicit(void *DevicePtr, int DeviceNum, int Kind,
+                               const char *Name);
 
 // This structure stores information of a mapped memory region.
 struct MapComponentInfoTy {
   void *Base;
   void *Begin;
   int64_t Size;
   int64_t Type;
   void *Name;

openmp/libomptarget/test/api/omp_device_managed_memory.c

 // RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
 // REQUIRES: nvptx64-nvidia-cuda
 
 #include <omp.h>
 #include <stdio.h>
 
 void *llvm_omp_target_alloc_shared(size_t, int);
+void llvm_omp_target_free_shared(void *, int);
 
 int main() {
   const int N = 64;
   const int device = omp_get_default_device();
 
   int *shared_ptr = llvm_omp_target_alloc_shared(N * sizeof(int), device);
 
-#pragma omp target teams distribute parallel for device(device) \
-            is_device_ptr(shared_ptr)
+#pragma omp target teams distribute parallel for device(device)                \
   for (int i = 0; i < N; ++i) {
     shared_ptr[i] = 1;
   }
 
   int sum = 0;
   for (int i = 0; i < N; ++i)
     sum += shared_ptr[i];
 
-  omp_target_free(shared_ptr, device);
+  llvm_omp_target_free_shared(shared_ptr, device);
   // CHECK: PASS
   if (sum == N)
-    printf ("PASS\n");
+    printf("PASS\n");
 }

openmp/libomptarget/test/api/omp_host_pinned_memory.c

 // RUN: %libomptarget-compile-run-and-check-generic
 
 #include <omp.h>
 #include <stdio.h>
 
 // Allocate pinned memory on the host
 void *llvm_omp_target_alloc_host(size_t, int);
+void llvm_omp_target_free_host(void *, int);
 
 int main() {
   const int N = 64;
   const int device = omp_get_default_device();
   const int host = omp_get_initial_device();
 
   int *hst_ptr = llvm_omp_target_alloc_host(N * sizeof(int), device);
 
   for (int i = 0; i < N; ++i)
     hst_ptr[i] = 2;
 
 #pragma omp target teams distribute parallel for device(device) \
            map(tofrom:hst_ptr[0 : N])
   for (int i = 0; i < N; ++i)
     hst_ptr[i] -= 1;
 
   int sum = 0;
   for (int i = 0; i < N; ++i)
     sum += hst_ptr[i];
 
-  omp_target_free(hst_ptr, device);
+  llvm_omp_target_free_host(hst_ptr, device);
   // CHECK: PASS
   if (sum == N)
     printf ("PASS\n");
 }

openmp/runtime/src/kmp_alloc.cpp

 static void **mk_hbw_hugetlb;
 static void **mk_hbw_preferred_hugetlb;
 static void **mk_dax_kmem;
 static void **mk_dax_kmem_all;
 static void **mk_dax_kmem_preferred;
 static void *(*kmp_target_alloc_host)(size_t size, int device);
 static void *(*kmp_target_alloc_shared)(size_t size, int device);
 static void *(*kmp_target_alloc_device)(size_t size, int device);
-static void *(*kmp_target_free)(void *ptr, int device);
+static void *(*kmp_target_free_host)(void *ptr, int device);
+static void *(*kmp_target_free_shared)(void *ptr, int device);
+static void *(*kmp_target_free_device)(void *ptr, int device);
 static bool __kmp_target_mem_available;
 #define KMP_IS_TARGET_MEM_SPACE(MS)                                            \
   (MS == llvm_omp_target_host_mem_space ||                                     \
    MS == llvm_omp_target_shared_mem_space ||                                   \
    MS == llvm_omp_target_device_mem_space)
 #define KMP_IS_TARGET_MEM_ALLOC(MA)                                            \
   (MA == llvm_omp_target_host_mem_alloc ||                                     \
    MA == llvm_omp_target_shared_mem_alloc ||                                   \

 }
 
 void __kmp_init_target_mem() {
   *(void **)(&kmp_target_alloc_host) = KMP_DLSYM("llvm_omp_target_alloc_host");
   *(void **)(&kmp_target_alloc_shared) =
       KMP_DLSYM("llvm_omp_target_alloc_shared");
   *(void **)(&kmp_target_alloc_device) =
       KMP_DLSYM("llvm_omp_target_alloc_device");
-  *(void **)(&kmp_target_free) = KMP_DLSYM("omp_target_free");
-  __kmp_target_mem_available = kmp_target_alloc_host &&
-                               kmp_target_alloc_shared &&
-                               kmp_target_alloc_device && kmp_target_free;
+  *(void **)(&kmp_target_free_host) = KMP_DLSYM("llvm_omp_target_free_host");
+  *(void **)(&kmp_target_free_shared) =
+      KMP_DLSYM("llvm_omp_target_free_shared");
+  *(void **)(&kmp_target_free_device) =
+      KMP_DLSYM("llvm_omp_target_free_device");
+  __kmp_target_mem_available =
+      kmp_target_alloc_host && kmp_target_alloc_shared &&
+      kmp_target_alloc_device && kmp_target_free_host &&
+      kmp_target_free_shared && kmp_target_free_device;
 }
 
 omp_allocator_handle_t __kmpc_init_allocator(int gtid, omp_memspace_handle_t ms,
                                              int ntraits,
                                              omp_alloctrait_t traits[]) {
   // OpenMP 5.0 only allows predefined memspaces
   KMP_DEBUG_ASSERT(ms == omp_default_mem_space || ms == omp_low_lat_mem_space ||
                    ms == omp_large_cap_mem_space || ms == omp_const_mem_space ||
if (ptr == NULL)
     return;
 
   kmp_allocator_t *al;
   omp_allocator_handle_t oal;
   al = RCAST(kmp_allocator_t *, CCAST(omp_allocator_handle_t, allocator));
   kmp_mem_desc_t desc;
   kmp_uintptr_t addr_align; // address to return to caller
   kmp_uintptr_t addr_descr; // address of memory block descriptor
-  if (KMP_IS_TARGET_MEM_ALLOC(allocator) ||
+  if (__kmp_target_mem_available && (KMP_IS_TARGET_MEM_ALLOC(allocator) ||
                                      (allocator > kmp_max_mem_alloc &&
-       KMP_IS_TARGET_MEM_SPACE(al->memspace))) {
-    KMP_DEBUG_ASSERT(kmp_target_free);
+                                      KMP_IS_TARGET_MEM_SPACE(al->memspace)))) {
     kmp_int32 device =
         __kmp_threads[gtid]->th.th_current_task->td_icvs.default_device;
-    kmp_target_free(ptr, device);
+    if (allocator == llvm_omp_target_host_mem_alloc) {
+      kmp_target_free_host(ptr, device);
+    } else if (allocator == llvm_omp_target_shared_mem_alloc) {
+      kmp_target_free_shared(ptr, device);
+    } else if (allocator == llvm_omp_target_device_mem_alloc) {
+      kmp_target_free_device(ptr, device);
+    }
     return;
   }
 
   addr_align = (kmp_uintptr_t)ptr;
   addr_descr = addr_align - sizeof(kmp_mem_desc_t);
   desc = *((kmp_mem_desc_t *)addr_descr); // read descriptor
 
   KMP_DEBUG_ASSERT(desc.ptr_align == ptr);