fix bugs in P2P implementation

- addPeers polarity reversed, would never add.
- check allow_access return value, pipe error to hipMalloc.

hipamd/src/hip_hcc.cpp

@@ -214,7 +214,8 @@ template<>
 bool ihipDeviceCriticalBase_t<DeviceMutex>::addPeer(ihipDevice_t *peer) 
 {
     auto match = std::find(_peers.begin(), _peers.end(), peer);
-    if (match != std::end(_peers)) {
+    if (match == std::end(_peers)) {
+        // Not already a peer, let's update the list:
         _peers.push_back(peer);
         recomputePeerAgents();
         return true;
@@ -230,6 +231,7 @@ bool ihipDeviceCriticalBase_t<DeviceMutex>::removePeer(ihipDevice_t *peer)
 {
     auto match = std::find(_peers.begin(), _peers.end(), peer);
     if (match != std::end(_peers)) {
+        // Found a valid peer, let's remove it.
         _peers.remove(peer);
         recomputePeerAgents();
         return true;

hipamd/src/hip_memory.cpp

@@ -133,7 +133,10 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
             {
                 LockedAccessor_DeviceCrit_t crit(device->criticalData());
                 if (crit->peerCnt()) {
-                    hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
+                    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; 
+                    }
                 }
             }
         }
@@ -172,7 +175,10 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
                 {
                     LockedAccessor_DeviceCrit_t crit(device->criticalData());
                     if (crit->peerCnt()) {
-                        hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
+                        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; 
+                        }
                     }
                 }
             }

hipamd/src/hip_peer.cpp

@@ -31,26 +31,23 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int  deviceId, int peerDe
 
     hipError_t err = hipSuccess;
 
-#if USE_PEER_TO_PEER
     auto thisDevice = ihipGetDevice(deviceId);
     auto peerDevice = ihipGetDevice(peerDeviceId);
 
     if ((thisDevice != NULL) && (peerDevice != NULL)) {
 #if USE_PEER_TO_PEER>=2
         *canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc);
+        printf ("canAccessPeer=%d\n", *canAccessPeer);
 #else
         *canAccessPeer = 0;
 #endif
 
     } else {
-        *canAccessPeer = false;
+        *canAccessPeer = 0;
         err = hipErrorInvalidDevice;
     }
 
 
-#else
-    *canAccessPeer = false;
-#endif
     return ihipLogStatus(err);
 }