AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

P2P checkpoint.

Browse Source 
Maintain enabled peer tables for each device.
This commit is contained in:
Ben Sander
2016-04-06 15:49:32 -05:00
parent 69f2469cbb
commit f2aa470f7f
6 changed files with 96 additions and 52 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/hip_peer.cpp
+43 -17
View File
@@ -26,7 +26,7 @@ THE SOFTWARE.
* @warning HCC returns 0 in *canAccessPeer ; Need to update this function when RT supports P2P
*/
//---
hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int deviceId, int peerDeviceId)
hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDeviceId)
{
HIP_INIT_API(canAccessPeer, deviceId, peerDeviceId);
@@ -38,7 +38,7 @@ hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int deviceId, int peerD
if ((device != NULL) && (peerDevice != NULL)) {
#if USE_PEER_TO_PEER==2
*canAccessPeer = peerDevice->_acc.is_peer(device->_acc);
*canAccessPeer = peerDevice->_acc.get_is_peer(device->_acc);
#else
*canAccessPeer = 0;
#endif
@@ -56,16 +56,32 @@ hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int deviceId, int peerD
}
/**
* warning Need to update this function when RT supports P2P
*/
//---
hipError_t hipDeviceDisablePeerAccess ( int peerDevice )
hipError_t hipDeviceDisablePeerAccess (int peerDeviceId)
{
HIP_INIT_API(peerDevice);
HIP_INIT_API(peerDeviceId);
// TODO-p2p
return ihipLogStatus(hipSuccess);
hipError_t err = hipSuccess;
#if USE_PEER_TO_PEER
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());
#endif
} else {
err = hipErrorPeerAccessNotEnabled; // never enabled P2P access.
}
} else {
err = hipErrorInvalidDevice;
}
#endif
return ihipLogStatus(err);
};
@@ -74,18 +90,27 @@ hipError_t hipDeviceDisablePeerAccess ( int peerDevice )
*/
//---
// Enable registering memory on peerDevice for direct access from the current device.
hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags)
hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags)
{
std::call_once(hip_initialized, ihipInit);
HIP_INIT_API(peerDeviceId, flags);
hipError_t err = hipSuccess;
#if USE_PEER_TO_PEER
if (flags != 0) {
err = hipErrorInvalidValue;
} else {
auto thisDevice = ihipGetTlsDefaultDevice();
auto peerDevice = ihipGetDevice(peerDeviceId);
if (peerDevice != NULL) {
if ((thisDevice != NULL) && (peerDevice != NULL)) {
LockedAccessor_DeviceCrit_t crit(thisDevice->criticalData());
bool isNewPeer = crit->addPeer(peerDevice);
if (isNewPeer) {
#if USE_PEER_TO_PEER==2
am_memtracker_update_peers(device->_acc, crit->peerCnt(), crit->peerAgents());
#endif
} else {
err = hipErrorPeerAccessAlreadyEnabled;
}
} else {
err = hipErrorInvalidDevice;
}
@@ -97,9 +122,10 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags)
//---
hipError_t hipMemcpyPeer ( void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes )
hipError_t hipMemcpyPeer (void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes)
{
std::call_once(hip_initialized, ihipInit);
HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes);
// HCC has a unified memory architecture so device specifiers are not required.
return hipMemcpy(dst, src, sizeBytes, hipMemcpyDefault);
};
@@ -109,9 +135,9 @@ hipError_t hipMemcpyPeer ( void* dst, int dstDevice, const void* src, int srcD
* @bug This function uses a synchronous copy
*/
//---
hipError_t hipMemcpyPeerAsync ( void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream )
hipError_t hipMemcpyPeerAsync (void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream)
{
std::call_once(hip_initialized, ihipInit);
HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes, stream);
// HCC has a unified memory architecture so device specifiers are not required.
return hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream);
};