P2p checkpoint.

- set USE_PEER_TO_PEER=3 (requires HCC "am_memtracker_update_peers")
- when enabling peer, turn it on for previously allocated memory.
- hipDeviceCanAccessPeer is no longer self-ware (self does not qualify
  as a peer)
- device peerlist always includes self, so when we call allow_access
  we never remove self access.
- hipDeviceReset() removes old peer mappings.


[ROCm/hip commit: 83f0de7806]

projects/hip/src/hip_device.cpp

@@ -174,6 +174,7 @@ hipError_t hipDeviceReset(void)
     if (device) {
         //---
         //Wait for pending activity to complete?  TODO - check if this is required behavior:
+        //TODO, also we have small window between wait and reset.
         
         device->locked_waitAllStreams();
 

projects/hip/src/hip_hcc.cpp

@@ -240,6 +240,15 @@ bool ihipDeviceCriticalBase_t<DeviceMutex>::removePeer(ihipDevice_t *peer)
     }
 }
 
+
+template<>
+void ihipDeviceCriticalBase_t<DeviceMutex>::resetPeers(ihipDevice_t *thisDevice)
+{
+    _peers.clear();
+    _peerCnt = 0;
+    addPeer(thisDevice); // peer-list always contains self agent.
+}
+
 //-------------------------------------------------------------------------------------------------
 
 //---
@@ -444,12 +453,18 @@ void ihipDevice_t::locked_reset()
     // Reset and remove streams:
     crit->streams().clear();
 
-#if USE_PEER_TO_PEER==2
-    // remove peer mappings to this device?  Call removePeer on all other devices?
+
+
+#if USE_PEER_TO_PEER>=2
+    // This resest peer list to just me:
+    crit->resetPeers(this);
+
 #endif
 
     // Reset and release all memory stored in the tracker:
+    // Reset will remove peer mapping so don't need to do this explicitly.
     am_memtracker_reset(_acc);
+
 };
 
 
@@ -474,10 +489,13 @@ void ihipDevice_t::init(unsigned device_index, unsigned deviceCnt, hc::accelerat
 
     getProperties(&_props);
 
+    _criticalData.init(deviceCnt);
+
+    locked_reset();
+
     _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);
 

projects/hip/src/hip_memory.cpp

@@ -132,7 +132,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
             hc::am_memtracker_update(*ptr, device->_device_index, 0);
             {
                 LockedAccessor_DeviceCrit_t crit(device->criticalData());
-                if (crit->peerCnt()) {
+                if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved:
                     hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
                     if (hsa_status != HSA_STATUS_SUCCESS) {
                         hip_status = hipErrorMemoryAllocation; 
@@ -173,8 +173,9 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
             }else{
                 hc::am_memtracker_update(*ptr, device->_device_index, flags);
                 {
+                    // TODO - allow_access only works for device memory, need to change am_alloc to allocate host directly.
                     LockedAccessor_DeviceCrit_t crit(device->criticalData());
-                    if (crit->peerCnt()) {
+                    if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved:
                         hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
                         if (hsa_status != HSA_STATUS_SUCCESS) {
                             hip_status = hipErrorMemoryAllocation; 

projects/hip/src/hip_peer.cpp

@@ -17,6 +17,8 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
+#include <hc_am.hpp>
+
 #include "hip_runtime.h"
 #include "hcc_detail/hip_hcc.h"
 #include "hcc_detail/trace_helper.h"
@@ -35,11 +37,15 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int  deviceId, int peerDe
     auto peerDevice = ihipGetDevice(peerDeviceId);
 
     if ((thisDevice != NULL) && (peerDevice != NULL)) {
+        if (deviceId == peerDeviceId) {
+            *canAccessPeer = 0;
+        } else {
 #if USE_PEER_TO_PEER>=2
-        *canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc);
+            *canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc);
 #else
-        *canAccessPeer = 0;
+            *canAccessPeer = 0;
 #endif
+        }
 
     } else {
         *canAccessPeer = 0;
@@ -69,15 +75,15 @@ hipError_t hipDeviceDisablePeerAccess (int peerDeviceId)
 #endif
         if (! canAccessPeer) {
             err = hipErrorInvalidDevice;  // P2P not allowed between these devices.
+        } else if (thisDevice == peerDevice)  {
+            err = hipErrorInvalidDevice;  // Can't disable peer access to self.
         } else {
-
-
             LockedAccessor_DeviceCrit_t crit(thisDevice->criticalData());
             bool changed = crit->removePeer(peerDevice);
             if (changed) {
 #if USE_PEER_TO_PEER>=3
                 // Update the peers for all memory already saved in the tracker:
-                am_memtracker_update_peers(device->_acc, crit->peerCnt(), crit->peerAgents());
+                am_memtracker_update_peers(thisDevice->_acc, crit->peerCnt(), crit->peerAgents());
 #endif
             } else {
                 err = hipErrorPeerAccessNotEnabled; // never enabled P2P access.
@@ -92,9 +98,6 @@ hipError_t hipDeviceDisablePeerAccess (int peerDeviceId)
 };
 
 
-/**
- * @warning Need to update this function when RT supports P2P
- */
 //---
  // Enable registering memory on peerDevice for direct access from the current device.
 hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags)
@@ -113,7 +116,7 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags)
             bool isNewPeer = crit->addPeer(peerDevice);
             if (isNewPeer) {
 #if USE_PEER_TO_PEER>=3
-                am_memtracker_update_peers(device->_acc, crit->peerCnt(), crit->peerAgents());
+                am_memtracker_update_peers(thisDevice->_acc, crit->peerCnt(), crit->peerAgents());
 #endif
             } else {
                 err = hipErrorPeerAccessAlreadyEnabled;