For P2P, use the peer list when allocating Device memory or pinned host.

Each new allocation is automatically mapped into the address space of
all enabled peers.

src/hip_hcc.cpp

@@ -442,9 +442,12 @@ 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?
+#endif
+
     // Reset and release all memory stored in the tracker:
     am_memtracker_reset(_acc);
-
 };
 
 

src/hip_memory.cpp

@@ -130,6 +130,12 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
             hip_status = hipErrorMemoryAllocation;
         } else {
             hc::am_memtracker_update(*ptr, device->_device_index, 0);
+            {
+                LockedAccessor_DeviceCrit_t crit(device->criticalData());
+                if (crit->peerCnt()) {
+                    hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
+                }
+            }
         }
     } else {
         hip_status = hipErrorMemoryAllocation;
@@ -139,29 +145,6 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
 }
 
 
-hipError_t hipMallocHost(void** ptr, size_t sizeBytes)
-{
-    HIP_INIT_API(ptr, sizeBytes);
-
-    hipError_t  hip_status = hipSuccess;
-
-    const unsigned am_flags = amHostPinned;
-	auto device = ihipGetTlsDefaultDevice();
-
-    if (device) {
-        *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags);
-        if (sizeBytes && (*ptr == NULL)) {
-            hip_status = hipErrorMemoryAllocation;
-        } else {
-            hc::am_memtracker_update(*ptr, device->_device_index, 0);
-        }
-
-        tprintf (DB_MEM, "  %s: pinned ptr=%p\n", __func__, *ptr);
-    }
-
-    return ihipLogStatus(hip_status);
-}
-
 
 hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
 {
@@ -186,6 +169,12 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
                 hip_status = hipErrorMemoryAllocation;
             }else{
                 hc::am_memtracker_update(*ptr, device->_device_index, flags);
+                {
+                    LockedAccessor_DeviceCrit_t crit(device->criticalData());
+                    if (crit->peerCnt()) {
+                        hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
+                    }
+                }
             }
             tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
         }
@@ -194,6 +183,7 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
 }
 
 
+//---
 // TODO - remove me, this is deprecated.
 hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags)
 {
@@ -201,6 +191,15 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags)
 };
 
 
+//---
+// TODO - remove me, this is deprecated.
+hipError_t hipMallocHost(void** ptr, size_t sizeBytes)
+{
+    return hipHostMalloc(ptr, sizeBytes, 0);
+}
+
+
+//---
 hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr)
 {
     HIP_INIT_API(flagsPtr, hostPtr);
@@ -225,6 +224,8 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr)
 	return ihipLogStatus(hip_status);
 }
 
+
+//---
 hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags)
 {
     HIP_INIT_API(hostPtr, sizeBytes, flags);

src/hip_peer.cpp

@@ -20,7 +20,6 @@ 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
@@ -67,15 +66,26 @@ hipError_t hipDeviceDisablePeerAccess (int peerDeviceId)
     auto thisDevice = ihipGetTlsDefaultDevice();
     auto peerDevice = ihipGetDevice(peerDeviceId);
     if ((thisDevice != NULL) && (peerDevice != NULL)) {
-        LockedAccessor_DeviceCrit_t crit(thisDevice->criticalData());
-        bool changed = crit->removePeer(peerDevice);
-        if (changed) {
 #if USE_PEER_TO_PEER==2
-            am_memtracker_update_peers(device->_acc, crit->peerCnt(), crit->peerAgents());
+        bool canAccessPeer =  peerDevice->_acc.get_is_peer(device->_acc);
+#else
+        bool canAccessPeer = 0;
 #endif
+        if (! canAccessPeer) {
+            err = hipErrorInvalidDevice;  // P2P not allowed between these devices.
         } else {
-            err = hipErrorPeerAccessNotEnabled; // never enabled P2P access.
-        }
+
+
+            LockedAccessor_DeviceCrit_t crit(thisDevice->criticalData());
+            bool changed = crit->removePeer(peerDevice);
+            if (changed) {
+#if USE_PEER_TO_PEER==2
+                am_memtracker_update_peers(device->_acc, crit->peerCnt(), crit->peerAgents());
+#endif
+            } else {
+                err = hipErrorPeerAccessNotEnabled; // never enabled P2P access.
+            }
+        } 
     } else {
         err = hipErrorInvalidDevice;
     }