Checkpoint initial peer2peer implementation.

hipamd/include/hcc_detail/hip_hcc.h

@@ -494,11 +494,23 @@ struct ihipEvent_t {
 // will lock the mutex on construction and unlock on destruction.
 //
 // MUTEX_TYPE is template argument so can easily convert to FakeMutex for performance or stress testing.
-template <typename MUTEX_TYPE>
+template <class MUTEX_TYPE>
 class ihipDeviceCriticalBase_t : LockedBase<MUTEX_TYPE>
 {
 public:
-    ihipDeviceCriticalBase_t() :  _stream_id(0) {};
+    ihipDeviceCriticalBase_t() :  _stream_id(0), _peerAgents(nullptr) {};
+   
+    void init(unsigned deviceCnt) {
+        assert(_peerAgents == nullptr);
+        _peerAgents = new hsa_agent_t[deviceCnt];
+    };
+
+    ~ihipDeviceCriticalBase_t()  {
+        if (_peerAgents != nullptr) {
+            delete _peerAgents;
+            _peerAgents = nullptr;
+        }
+    }
     friend class LockedAccessor<ihipDeviceCriticalBase_t>;
 
     std::list<ihipStream_t*> &streams() { return _streams; };
@@ -507,10 +519,19 @@ public:
     // "Allocate" a stream ID:
     ihipStream_t::SeqNum_t  incStreamId() { return _stream_id++; };
 
+    void recomputePeerAgents();
+    void addPeer(ihipDevice_t *peer);
+    void removePeer(ihipDevice_t *peer);
+
 
 private:
     std::list<ihipStream_t*> _streams;   // streams associated with this device.
     ihipStream_t::SeqNum_t   _stream_id;
+
+    // These reflect the currently Enabled set of peers for this GPU:
+    std::list<ihipDevice_t*>  _peers;     // list of enabled peer devices.
+    uint32_t                  _peerCnt;     // number of enabled peers
+    hsa_agent_t              *_peerAgents;  // efficient packed array of enabled agents (to use for allocations.)
 };
 
 // Note Mutex selected based on DeviceMutex
@@ -530,7 +551,7 @@ class ihipDevice_t
 {
 public: // Functions:
     ihipDevice_t() {}; // note: calls constructor for _criticalData 
-    void init(unsigned device_index, hc::accelerator &acc, unsigned flags);
+    void init(unsigned device_index, unsigned deviceCnt, hc::accelerator &acc, unsigned flags);
     ~ihipDevice_t();
 
     void locked_addStream(ihipStream_t *s);

hipamd/include/hcc_detail/hip_runtime_api.h

@@ -908,12 +908,12 @@ hipError_t hipMemGetInfo  (size_t * free, size_t * total)   ;
  * Returns "1" in @p canAccessPeer if the specified @p device is capable
  * of directly accessing memory physically located on peerDevice , or "0" if not.
  */
-hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int  device, int  peerDevice );
+hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDeviceId);
 
 
 
 /**
- * @brief Disables registering memory on peerDevice for direct access from the current device.
+ * @brief Disable registering memory on peerDevice for direct access from the current device.
  *
  * If there are any allocations on peerDevice which were registered in the current device using hipPeerRegister() then these allocations will be automatically unregistered.
  * Returns hipErrorPeerAccessNotEnabled if direct access to memory on peerDevice has not yet been enabled from the current device.
@@ -922,10 +922,10 @@ hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int  device, int  peerDe
  * TODO:cudaErrorPeerAccessNotEnabled and cudaErrorInvalidDevice error not supported in HIP, return hipErrorUnknown
  * Returns #hipSuccess, #hipErrorUnknown
  */
-hipError_t  hipDeviceDisablePeerAccess ( int  peerDevice );
+hipError_t  hipDeviceDisablePeerAccess (int peerDeviceId);
 
 /**
- * @brief Enables registering memory on peerDevice for direct access from the current device.
+ * @brief Enable registering memory on peerDevice for direct access from the current device.
  *
  * @param [in] peerDevice
  * @param [in] flags
@@ -933,7 +933,7 @@ hipError_t  hipDeviceDisablePeerAccess ( int  peerDevice );
  * TODO:cudaErrorInvalidDevice error not supported in HIP, return hipErrorUnknown
  * Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, #hipErrorUnknown
  */
-hipError_t  hipDeviceEnablePeerAccess ( int  peerDevice, unsigned int  flags );
+hipError_t  hipDeviceEnablePeerAccess (int  peerDeviceId, unsigned int flags);
 
 /**
  * @brief Copies memory from one device to memory on another device.

hipamd/src/hip_hcc.cpp

@@ -198,6 +198,40 @@ void ihipStream_t::locked_wait(bool assertQueueEmpty)
 };
 
 
+// Recompute the peercnt and the packed _peerAgents whenever a peer is added or deleted.
+// The packed _peerAgents can efficiently be used on each memory allocation.
+template<> 
+void ihipDeviceCriticalBase_t<DeviceMutex>::recomputePeerAgents()
+{
+    _peerCnt = 0;
+    std::for_each (_peers.begin(), _peers.end(), [this](ihipDevice_t* device) {
+        _peerAgents[_peerCnt++] = device->_hsa_agent;
+    });
+}
+
+
+template<>
+void ihipDeviceCriticalBase_t<DeviceMutex>::addPeer(ihipDevice_t *peer) 
+{
+    auto match = std::find(_peers.begin(), _peers.end(), peer);
+    if (match != std::end(_peers)) {
+        _peers.push_back(peer);
+        recomputePeerAgents();
+        return;
+    }
+
+    // If we get here - peer was already on list, silently ignore.
+}
+
+
+template<>
+void ihipDeviceCriticalBase_t<DeviceMutex>::removePeer(ihipDevice_t *peer) 
+{
+    _peers.remove(peer);
+    recomputePeerAgents();
+}
+
+//-------------------------------------------------------------------------------------------------
 
 //---
 ihipDevice_t * ihipStream_t::getDevice() const
@@ -408,7 +442,7 @@ void ihipDevice_t::locked_reset()
 
 
 //---
-void ihipDevice_t::init(unsigned device_index, hc::accelerator &acc, unsigned flags)
+void ihipDevice_t::init(unsigned device_index, unsigned deviceCnt, hc::accelerator &acc, unsigned flags)
 {
     _device_index = device_index;
     _device_flags = flags;
@@ -431,6 +465,8 @@ void ihipDevice_t::init(unsigned device_index, hc::accelerator &acc, unsigned fl
     _default_stream = new ihipStream_t(device_index, acc.get_default_view(), hipStreamDefault);
     locked_addStream(_default_stream);
 
+    _criticalData.init(deviceCnt);
+
     tprintf(DB_SYNC, "created device with default_stream=%p\n", _default_stream);
 
     hsa_region_t *pinnedHostRegion;
@@ -905,7 +941,7 @@ void ihipInit()
                 //If device is not in visible devices list, ignore
                 continue;
             }
-            g_devices[g_deviceCnt].init(g_deviceCnt, accs[i], hipDeviceMapHost);
+            g_devices[g_deviceCnt].init(g_deviceCnt, deviceCnt, accs[i], hipDeviceMapHost);
             g_deviceCnt++;
         }
     }

hipamd/src/hip_peer.cpp

@@ -20,22 +20,44 @@ THE SOFTWARE.
 #include "hip_runtime.h"
 #include "hcc_detail/hip_hcc.h"
 #include "hcc_detail/trace_helper.h"
+#define USE_PEER_TO_PEER 1
 
 /**
  * @warning HCC returns 0 in *canAccessPeer ; Need to update this function when RT supports P2P
  */
 //---
-hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int  device, int  peerDevice )
+hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int  deviceId, int peerDeviceId)
 {
-    HIP_INIT_API(canAccessPeer, device, peerDevice);
+    HIP_INIT_API(canAccessPeer, deviceId, peerDeviceId);
 
+    hipError_t err = hipSuccess;
+
+#if USE_PEER_TO_PEER
+    auto device = ihipGetDevice(deviceId);
+    auto peerDevice = ihipGetDevice(peerDeviceId);
+
+    if ((device != NULL) && (peerDevice != NULL)) {
+#if USE_PEER_TO_PEER==2
+        *canAccessPeer = peerDevice->_acc.is_peer(device->_acc);
+#else
+        *canAccessPeer = 0;
+#endif
+
+    } else {
+        *canAccessPeer = false;
+        err = hipErrorInvalidDevice;
+    }
+
+
+#else
     *canAccessPeer = false;
-    return ihipLogStatus(hipSuccess);
+#endif
+    return ihipLogStatus(err);
 }
 
 
 /**
- * @warning Need to update this function when RT supports P2P
+ * warning Need to update this function when RT supports P2P
  */
 //---
 hipError_t  hipDeviceDisablePeerAccess ( int  peerDevice )
@@ -51,11 +73,26 @@ hipError_t  hipDeviceDisablePeerAccess ( int  peerDevice )
  * @warning Need to update this function when RT supports P2P
  */
 //---
-hipError_t  hipDeviceEnablePeerAccess ( int  peerDevice, unsigned int  flags )
+ // Enable registering memory on peerDevice for direct access from the current device.
+hipError_t  hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags)
 {
     std::call_once(hip_initialized, ihipInit);
-    // TODO-p2p
+
-    return ihipLogStatus(hipSuccess);
+    hipError_t err = hipSuccess;
+#if USE_PEER_TO_PEER
+    if (flags != 0) {
+        err = hipErrorInvalidValue;
+    } else {
+        auto peerDevice = ihipGetDevice(peerDeviceId);
+        if (peerDevice != NULL) {
+
+        } else {
+            err = hipErrorInvalidDevice;
+        }
+    }
+#endif
+
+    return ihipLogStatus(err);
 }