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

For P2P, use the peer list when allocating Device memory or pinned host.

Browse Source 
Each new allocation is automatically mapped into the address space of
all enabled peers.


[ROCm/clr commit: 97772d6363]
This commit is contained in:
Ben Sander
2016-04-06 16:44:31 -05:00
parent 7eb8c10215
commit 2890517a23
6 changed files with 59 additions and 38 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hip_hcc.cpp
+4 -1
View File
@@ -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/clr/hipamd/src/hip_memory.cpp
+24 -23
View File
@@ -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/clr/hipamd/src/hip_peer.cpp
+17 -7
View File
@@ -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;
}