diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index 20c6f42451..ddf6bb1691 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -214,7 +214,8 @@ template<> bool ihipDeviceCriticalBase_t::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::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; diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 4cbec08ba0..2a4ec205e2 100644 --- a/hipamd/src/hip_memory.cpp +++ b/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; + } } } } diff --git a/hipamd/src/hip_peer.cpp b/hipamd/src/hip_peer.cpp index 2e52f8e586..ad9cf14634 100644 --- a/hipamd/src/hip_peer.cpp +++ b/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); }