Remove LastDeviceWritingToMemObj

LastEventWritingToMemObj is never an interop event, so always has an
associated queue. So using an extra LastDeviceWritingToMemObj is not
necessary.

Author: Hugh Delaney

source/adapters/hip/enqueue.cpp

@@ -218,8 +218,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferRead(
     // last queue to write to the MemBuffer, meaning we must perform the copy
     // from a different device
     if (hBuffer->LastEventWritingToMemObj &&
-        hBuffer->LastDeviceWritingToMemObj != hQueue->getDevice()) {
-      Device = hBuffer->LastDeviceWritingToMemObj;
+        hBuffer->LastEventWritingToMemObj->getQueue()->getDevice() !=
+            hQueue->getDevice()) {
+      Device = hBuffer->LastEventWritingToMemObj->getQueue()->getDevice();
       ScopedContext Active(Device);
       HIPStream = hipStream_t{0}; // Default stream for different device
       // We may have to wait for an event on another queue if it is the last
@@ -367,8 +368,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
         // if it has been written to
         if (phEvent && (MemArg.AccessFlags &
                         (UR_MEM_FLAG_READ_WRITE | UR_MEM_FLAG_WRITE_ONLY))) {
-          MemArg.Mem->setLastEventWritingToMemObj(RetImplEvent.get(),
-                                                  hQueue->getDevice());
+          MemArg.Mem->setLastEventWritingToMemObj(RetImplEvent.get());
         }
       }
       // We can release the MemoryMigrationMutexes now
@@ -585,8 +585,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferReadRect(
     // last queue to write to the MemBuffer, meaning we must perform the copy
     // from a different device
     if (hBuffer->LastEventWritingToMemObj &&
-        hBuffer->LastDeviceWritingToMemObj != hQueue->getDevice()) {
-      Device = hBuffer->LastDeviceWritingToMemObj;
+        hBuffer->LastEventWritingToMemObj->getQueue()->getDevice() !=
+            hQueue->getDevice()) {
+      Device = hBuffer->LastEventWritingToMemObj->getQueue()->getDevice();
       ScopedContext Active(Device);
       HIPStream = hipStream_t{0}; // Default stream for different device
       // We may have to wait for an event on another queue if it is the last
@@ -1018,8 +1019,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemImageRead(
     // last queue to write to the MemBuffer, meaning we must perform the copy
     // from a different device
     if (hImage->LastEventWritingToMemObj &&
-        hImage->LastDeviceWritingToMemObj != hQueue->getDevice()) {
-      Device = hImage->LastDeviceWritingToMemObj;
+        hImage->LastEventWritingToMemObj->getQueue()->getDevice() !=
+            hQueue->getDevice()) {
+      Device = hImage->LastEventWritingToMemObj->getQueue()->getDevice();
       ScopedContext Active(Device);
       HIPStream = hipStream_t{0}; // Default stream for different device
       // We may have to wait for an event on another queue if it is the last

source/adapters/hip/memory.cpp

@@ -525,10 +525,12 @@ inline ur_result_t migrateBufferToDevice(ur_mem_handle_t Mem,
       UR_CHECK_ERROR(
           hipMemcpyHtoD(Buffer.getPtr(hDevice), Buffer.HostPtr, Buffer.Size));
     }
-  } else if (Mem->LastDeviceWritingToMemObj != hDevice) {
-    UR_CHECK_ERROR(hipMemcpyDtoD(Buffer.getPtr(hDevice),
-                                 Buffer.getPtr(Mem->LastDeviceWritingToMemObj),
-                                 Buffer.Size));
+  } else if (Mem->LastEventWritingToMemObj->getQueue()->getDevice() !=
+             hDevice) {
+    UR_CHECK_ERROR(hipMemcpyDtoD(
+        Buffer.getPtr(hDevice),
+        Buffer.getPtr(Mem->LastEventWritingToMemObj->getQueue()->getDevice()),
+        Buffer.Size));
   }
   return UR_RESULT_SUCCESS;
 }
@@ -576,19 +578,24 @@ inline ur_result_t migrateImageToDevice(ur_mem_handle_t Mem,
       CpyDesc3D.srcHost = Image.HostPtr;
       UR_CHECK_ERROR(hipDrvMemcpy3D(&CpyDesc3D));
     }
-  } else if (Mem->LastDeviceWritingToMemObj != hDevice) {
+  } else if (Mem->LastEventWritingToMemObj->getQueue()->getDevice() !=
+             hDevice) {
     if (Image.ImageDesc.type == UR_MEM_TYPE_IMAGE1D) {
       // FIXME: 1D memcpy from DtoD going through the host.
       UR_CHECK_ERROR(hipMemcpyAtoH(
-          Image.HostPtr, Image.getArray(Mem->LastDeviceWritingToMemObj),
+          Image.HostPtr,
+          Image.getArray(
+              Mem->LastEventWritingToMemObj->getQueue()->getDevice()),
           0 /*srcOffset*/, ImageSizeBytes));
       UR_CHECK_ERROR(
           hipMemcpyHtoA(ImageArray, 0, Image.HostPtr, ImageSizeBytes));
     } else if (Image.ImageDesc.type == UR_MEM_TYPE_IMAGE2D) {
-      CpyDesc2D.srcArray = Image.getArray(Mem->LastDeviceWritingToMemObj);
+      CpyDesc2D.srcArray = Image.getArray(
+          Mem->LastEventWritingToMemObj->getQueue()->getDevice());
       UR_CHECK_ERROR(hipMemcpyParam2D(&CpyDesc2D));
     } else if (Image.ImageDesc.type == UR_MEM_TYPE_IMAGE3D) {
-      CpyDesc3D.srcArray = Image.getArray(Mem->LastDeviceWritingToMemObj);
+      CpyDesc3D.srcArray = Image.getArray(
+          Mem->LastEventWritingToMemObj->getQueue()->getDevice());
       UR_CHECK_ERROR(hipDrvMemcpy3D(&CpyDesc3D));
     }
   }

source/adapters/hip/memory.hpp

@@ -393,10 +393,6 @@ struct ur_mem_handle_t_ {
   // We should wait on this event prior to migrating memory across allocations
   // in this ur_mem_handle_t_
   ur_event_handle_t LastEventWritingToMemObj{nullptr};
-  // Since the event may not contain device info (if using interop, which
-  // doesn't take a queue) we should use this member var to keep track of which
-  // device has most recent view of data
-  ur_device_handle_t LastDeviceWritingToMemObj{nullptr};
 
   // Enumerates all possible types of accesses.
   enum access_mode_t { unknown, read_write, read_only, write_only };
@@ -491,23 +487,18 @@ struct ur_mem_handle_t_ {
 
   uint32_t getReferenceCount() const noexcept { return RefCount; }
 
-  void setLastEventWritingToMemObj(ur_event_handle_t NewEvent,
-                                   ur_device_handle_t RecentDevice) {
+  void setLastEventWritingToMemObj(ur_event_handle_t NewEvent) {
     assert(NewEvent && "Invalid event!");
     // This entry point should only ever be called when using multi device ctx
     assert(Context->Devices.size() > 1);
-    urEventRetain(NewEvent);
-    urDeviceRetain(RecentDevice);
     if (LastEventWritingToMemObj != nullptr) {
       urEventRelease(LastEventWritingToMemObj);
     }
-    if (LastDeviceWritingToMemObj != nullptr) {
-      urDeviceRelease(LastDeviceWritingToMemObj);
-    }
+    urEventRetain(NewEvent);
     LastEventWritingToMemObj = NewEvent;
     for (const auto &Device : Context->getDevices()) {
       HaveMigratedToDeviceSinceLastWrite[Device->getIndex()] =
-          Device == RecentDevice;
+          Device == NewEvent->getQueue()->getDevice();
     }
   }
 };