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.


[ROCm/hip commit: 813b063888]

projects/hip/CONTRIBUTING.md

@@ -55,8 +55,8 @@ When adding a new HIP feature, add a new unit test as well.
 See [tests/README.md](README.md) for more information.
 
 ## Development Flow
-The Unit testing environment automatically rebuilds libhip_hcc.a and the tests when a change it made to the HIP source, and this 
-is a great place to develop new features alongside the associated test.  
+It is recommended that developers set the flag HIP_BUILD_LOCAL=1 so that the unit testing environment automatically rebuilds libhip_hcc.a and the tests when a change it made to the HIP source. 
+Directed tests provide a great place to develop new features alongside the associated test.  
 
 For applications and benchmarks outside the directed test environment, developments should use a two-step development flow:
 - #1. Compile, link, and install HCC.  See [Installation](README.md#Installation) notes.

projects/hip/include/hcc_detail/hip_hcc.h

@@ -35,6 +35,10 @@ THE SOFTWARE.
 //Use the new HCC accelerator_view::copy instead of am_copy
 #define USE_AV_COPY 0
 
+// Compile peer-to-peer support.
+// 2= use upcoming HCC APIs.
+#define USE_PEER_TO_PEER 1
+
 //#define INLINE static inline
 
 //---
@@ -523,7 +527,7 @@ public:
     bool removePeer(ihipDevice_t *peer);
 
     uint32_t peerCnt() const { return _peerCnt; };
-    uint32_t peerAgents() const { return _peerAgents; };
+    hsa_agent_t *peerAgents() const { return _peerAgents; };
 
 
 private:

projects/hip/include/hcc_detail/hip_runtime_api.h

@@ -919,19 +919,22 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDev
  *
  * @param [in] peerDeviceId
  *
- * TODO:cudaErrorPeerAccessNotEnabled and cudaErrorInvalidDevice error not supported in HIP, return hipErrorUnknown
- * Returns #hipSuccess, #hipErrorUnknown
+ * Returns #hipSuccess, #hipErrorPeerAccessNotEnabled
  */
 hipError_t  hipDeviceDisablePeerAccess (int peerDeviceId);
 
 /**
  * @brief Enable direct access from current device's virtual address space to memory allocations physically located on a peer device.  
  *
+ * Memory which already allocated on peer device will be mapped into the address space of the current device.  In addition, all
+ * future memory allocations on peerDeviceId will be mapped into the address space of the current device when the memory is allocated.
+ * The peer memory remains accessible from the current device until a call to hipDeviceDisablePeerAccess or @hipDeviceReset.
+ *
+ *
  * @param [in] peerDeviceId
  * @param [in] flags
  *
- * TODO:cudaErrorInvalidDevice error not supported in HIP, return hipErrorUnknown
- * Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, #hipErrorUnknown
+ * Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, #hipErrorPeerAccessAlreadyEnabled
  */
 hipError_t  hipDeviceEnablePeerAccess (int  peerDeviceId, unsigned int flags);
 

projects/hip/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);
-
 };
 
 

projects/hip/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);

projects/hip/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;
     }