Update Enable/Disable peers to match cuda behaviour

Change-Id: I67194ccf77a0019368579ff7d95b7790fcf228f3


[ROCm/hip commit: bdb3a4b393]

projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp

@@ -23,7 +23,7 @@ THE SOFTWARE.
 // Also serves as a template for other tests.
 
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp  EXCLUDE_HIP_PLATFORM vdi
+ * BUILD: %t %s ../test_common.cpp
  * TEST: %t EXCLUDE_HIP_PLATFORM hcc vdi
  * TEST: %t --memcpyWithPeer EXCLUDE_HIP_PLATFORM hcc vdi
  * TEST: %t --mirrorPeers EXCLUDE_HIP_PLATFORM hcc vdi

projects/hip/vdi/hip_internal.hpp

@@ -73,11 +73,14 @@ namespace hip {
 
   /// HIP Device class
   class Device {
+    amd::Monitor lock_{"Device lock"};
     /// VDI context
     amd::Context* context_;
     /// Device's ID
     /// Store it here so we don't have to loop through the device list every time
     int deviceId_;
+    //Maintain list of user enabled peers
+    std::list<int> userEnabledPeers;
   public:
     Device(amd::Context* ctx, int devId): context_(ctx), deviceId_(devId) { assert(ctx != nullptr); }
     ~Device() {}
@@ -87,6 +90,25 @@ namespace hip {
     void retain() const { context_->retain(); }
     void release() const { context_->release(); }
     const std::vector<amd::Device*>& devices() const { return context_->devices(); }
+    hipError_t EnablePeerAccess(int peerDeviceId){
+      amd::ScopedLock lock(lock_);
+      bool found = (std::find(userEnabledPeers.begin(), userEnabledPeers.end(), peerDeviceId) != userEnabledPeers.end());
+      if (found) {
+        return hipErrorPeerAccessAlreadyEnabled;
+      }
+      userEnabledPeers.push_back(peerDeviceId);
+      return hipSuccess;
+    }
+    hipError_t DisablePeerAccess(int peerDeviceId) {
+      amd::ScopedLock lock(lock_);
+      bool found = (std::find(userEnabledPeers.begin(), userEnabledPeers.end(), peerDeviceId) != userEnabledPeers.end());
+      if (found) {
+        userEnabledPeers.remove(peerDeviceId);
+        return hipSuccess;
+      } else {
+        return hipErrorPeerAccessNotEnabled;
+      }
+    }
   };
 
   extern std::once_flag g_ihipInitialized;

projects/hip/vdi/hip_peer.cpp

@@ -48,48 +48,56 @@ hipError_t hipMemcpyPeerAsync(void* dst, hipCtx_t dstDevice, const void* src, hi
   HIP_RETURN(hipErrorNotSupported);
 }
 
-hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId) {
-  HIP_INIT_API(hipDeviceCanAccessPeer, canAccessPeer, deviceId, peerDeviceId);
-
+hipError_t canAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId){
   amd::Device* device = nullptr;
   amd::Device* peer_device = nullptr;
-
   if (canAccessPeer == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
-
   /* Peer cannot be self */
   if (deviceId == peerDeviceId) {
     *canAccessPeer = 0;
-    return HIP_RETURN(hipSuccess);
+    HIP_RETURN(hipSuccess);
   }
-
   /* Cannot exceed the max number of devices */
   if (static_cast<size_t>(deviceId) >= g_devices.size()
        || static_cast<size_t>(peerDeviceId) >= g_devices.size()) {
-    return HIP_RETURN(hipErrorInvalidValue);
+    HIP_RETURN(hipErrorInvalidDevice);
   }
-
   device = g_devices[deviceId]->devices()[0];
   peer_device = g_devices[peerDeviceId]->devices()[0];
-
   *canAccessPeer = static_cast<int>(std::find(device->p2pDevices_.begin(),
                                               device->p2pDevices_.end(), as_cl(peer_device))
                                               != device->p2pDevices_.end());
+  HIP_RETURN(hipSuccess);
+}
 
-  return HIP_RETURN(hipSuccess);
+hipError_t hipDeviceCanAccessPeer(int* canAccess, int deviceId, int peerDeviceId) {
+  HIP_INIT_API(hipDeviceCanAccessPeer, canAccess, deviceId, peerDeviceId);
+  HIP_RETURN(canAccessPeer(canAccess, deviceId, peerDeviceId));
 }
 
 hipError_t hipDeviceDisablePeerAccess(int peerDeviceId) {
   HIP_INIT_API(hipDeviceDisablePeerAccess, peerDeviceId);
-
-  HIP_RETURN(hipSuccess);
+  int deviceId = hip::getCurrentDevice()->deviceId();
+  int canAccess = 0;
+  if ((hipSuccess != canAccessPeer(&canAccess, deviceId, peerDeviceId)) || (canAccess == 0)) {
+    HIP_RETURN(hipErrorInvalidDevice);
+  }
+  HIP_RETURN(hip::getCurrentDevice()->DisablePeerAccess(peerDeviceId));
 }
 
 hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags) {
   HIP_INIT_API(hipDeviceEnablePeerAccess, peerDeviceId, flags);
-
-  HIP_RETURN(hipSuccess);
+  int deviceId = hip::getCurrentDevice()->deviceId();
+  int canAccess = 0;
+  if (flags != 0) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  if ((hipSuccess != canAccessPeer(&canAccess, deviceId, peerDeviceId)) || (canAccess == 0)) {
+    HIP_RETURN(hipErrorInvalidDevice);
+  }
+  HIP_RETURN(hip::getCurrentDevice()->EnablePeerAccess(peerDeviceId));
 }
 
 hipError_t hipMemcpyPeer(void* dst, int dstDevice, const void* src, int srcDevice,