diff --git a/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp b/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp index c2e5965e39..90e7112356 100644 --- a/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp +++ b/projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp @@ -23,7 +23,7 @@ THE SOFTWARE. // Also serves as a template for other tests. /* HIT_START - * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM vdi + * BUILD: %t %s ../test_common.cpp * TEST: %t EXCLUDE_HIP_PLATFORM hcc vdi * TEST: %t --memcpyWithPeer EXCLUDE_HIP_PLATFORM hcc vdi * TEST: %t --mirrorPeers EXCLUDE_HIP_PLATFORM hcc vdi diff --git a/projects/hip/vdi/hip_internal.hpp b/projects/hip/vdi/hip_internal.hpp index 129978dd68..a1cbe33a29 100755 --- a/projects/hip/vdi/hip_internal.hpp +++ b/projects/hip/vdi/hip_internal.hpp @@ -73,11 +73,14 @@ namespace hip { /// HIP Device class class Device { + amd::Monitor lock_{"Device lock"}; /// VDI context amd::Context* context_; /// Device's ID /// Store it here so we don't have to loop through the device list every time int deviceId_; + //Maintain list of user enabled peers + std::list userEnabledPeers; public: Device(amd::Context* ctx, int devId): context_(ctx), deviceId_(devId) { assert(ctx != nullptr); } ~Device() {} @@ -87,6 +90,25 @@ namespace hip { void retain() const { context_->retain(); } void release() const { context_->release(); } const std::vector& devices() const { return context_->devices(); } + hipError_t EnablePeerAccess(int peerDeviceId){ + amd::ScopedLock lock(lock_); + bool found = (std::find(userEnabledPeers.begin(), userEnabledPeers.end(), peerDeviceId) != userEnabledPeers.end()); + if (found) { + return hipErrorPeerAccessAlreadyEnabled; + } + userEnabledPeers.push_back(peerDeviceId); + return hipSuccess; + } + hipError_t DisablePeerAccess(int peerDeviceId) { + amd::ScopedLock lock(lock_); + bool found = (std::find(userEnabledPeers.begin(), userEnabledPeers.end(), peerDeviceId) != userEnabledPeers.end()); + if (found) { + userEnabledPeers.remove(peerDeviceId); + return hipSuccess; + } else { + return hipErrorPeerAccessNotEnabled; + } + } }; extern std::once_flag g_ihipInitialized; diff --git a/projects/hip/vdi/hip_peer.cpp b/projects/hip/vdi/hip_peer.cpp index 13cfd1cc7e..225361d525 100644 --- a/projects/hip/vdi/hip_peer.cpp +++ b/projects/hip/vdi/hip_peer.cpp @@ -48,48 +48,56 @@ hipError_t hipMemcpyPeerAsync(void* dst, hipCtx_t dstDevice, const void* src, hi HIP_RETURN(hipErrorNotSupported); } -hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId) { - HIP_INIT_API(hipDeviceCanAccessPeer, canAccessPeer, deviceId, peerDeviceId); - +hipError_t canAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId){ amd::Device* device = nullptr; amd::Device* peer_device = nullptr; - if (canAccessPeer == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - /* Peer cannot be self */ if (deviceId == peerDeviceId) { *canAccessPeer = 0; - return HIP_RETURN(hipSuccess); + HIP_RETURN(hipSuccess); } - /* Cannot exceed the max number of devices */ if (static_cast(deviceId) >= g_devices.size() || static_cast(peerDeviceId) >= g_devices.size()) { - return HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(hipErrorInvalidDevice); } - device = g_devices[deviceId]->devices()[0]; peer_device = g_devices[peerDeviceId]->devices()[0]; - *canAccessPeer = static_cast(std::find(device->p2pDevices_.begin(), device->p2pDevices_.end(), as_cl(peer_device)) != device->p2pDevices_.end()); + HIP_RETURN(hipSuccess); +} - return HIP_RETURN(hipSuccess); +hipError_t hipDeviceCanAccessPeer(int* canAccess, int deviceId, int peerDeviceId) { + HIP_INIT_API(hipDeviceCanAccessPeer, canAccess, deviceId, peerDeviceId); + HIP_RETURN(canAccessPeer(canAccess, deviceId, peerDeviceId)); } hipError_t hipDeviceDisablePeerAccess(int peerDeviceId) { HIP_INIT_API(hipDeviceDisablePeerAccess, peerDeviceId); - - HIP_RETURN(hipSuccess); + int deviceId = hip::getCurrentDevice()->deviceId(); + int canAccess = 0; + if ((hipSuccess != canAccessPeer(&canAccess, deviceId, peerDeviceId)) || (canAccess == 0)) { + HIP_RETURN(hipErrorInvalidDevice); + } + HIP_RETURN(hip::getCurrentDevice()->DisablePeerAccess(peerDeviceId)); } hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags) { HIP_INIT_API(hipDeviceEnablePeerAccess, peerDeviceId, flags); - - HIP_RETURN(hipSuccess); + int deviceId = hip::getCurrentDevice()->deviceId(); + int canAccess = 0; + if (flags != 0) { + HIP_RETURN(hipErrorInvalidValue); + } + if ((hipSuccess != canAccessPeer(&canAccess, deviceId, peerDeviceId)) || (canAccess == 0)) { + HIP_RETURN(hipErrorInvalidDevice); + } + HIP_RETURN(hip::getCurrentDevice()->EnablePeerAccess(peerDeviceId)); } hipError_t hipMemcpyPeer(void* dst, int dstDevice, const void* src, int srcDevice,