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.
This commit is contained in:
Ben Sander
2016-04-06 16:44:31 -05:00
parent e2d19d7f7a
commit 97772d6363
6 changed files with 59 additions and 38 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/CONTRIBUTING.md
+2 -2
View File
@@ -55,8 +55,8 @@ When adding a new HIP feature, add a new unit test as well.
See [tests/README.md](README.md) for more information.
## Development Flow
The Unit testing environment automatically rebuilds libhip_hcc.a and the tests when a change it made to the HIP source, and this
is a great place to develop new features alongside the associated test.
It is recommended that developers set the flag HIP_BUILD_LOCAL=1 so that the unit testing environment automatically rebuilds libhip_hcc.a and the tests when a change it made to the HIP source.
Directed tests provide a great place to develop new features alongside the associated test.
For applications and benchmarks outside the directed test environment, developments should use a two-step development flow:
- #1. Compile, link, and install HCC. See [Installation](README.md#Installation) notes.
hipamd/include/hcc_detail/hip_hcc.h
+5 -1
View File
@@ -35,6 +35,10 @@ THE SOFTWARE.
//Use the new HCC accelerator_view::copy instead of am_copy
#define USE_AV_COPY 0
// Compile peer-to-peer support.
// 2= use upcoming HCC APIs.
#define USE_PEER_TO_PEER 1
//#define INLINE static inline
//---
@@ -523,7 +527,7 @@ public:
bool removePeer(ihipDevice_t *peer);
uint32_t peerCnt() const { return _peerCnt; };
uint32_t peerAgents() const { return _peerAgents; };
hsa_agent_t *peerAgents() const { return _peerAgents; };
private:
hipamd/include/hcc_detail/hip_runtime_api.h
+7 -4
View File
@@ -919,19 +919,22 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDev
*
* @param [in] peerDeviceId
*
* TODO:cudaErrorPeerAccessNotEnabled and cudaErrorInvalidDevice error not supported in HIP, return hipErrorUnknown
* Returns #hipSuccess, #hipErrorUnknown
* Returns #hipSuccess, #hipErrorPeerAccessNotEnabled
*/
hipError_t hipDeviceDisablePeerAccess (int peerDeviceId);
/**
* @brief Enable direct access from current device's virtual address space to memory allocations physically located on a peer device.
*
* Memory which already allocated on peer device will be mapped into the address space of the current device. In addition, all
* future memory allocations on peerDeviceId will be mapped into the address space of the current device when the memory is allocated.
* The peer memory remains accessible from the current device until a call to hipDeviceDisablePeerAccess or @hipDeviceReset.
*
*
* @param [in] peerDeviceId
* @param [in] flags
*
* TODO:cudaErrorInvalidDevice error not supported in HIP, return hipErrorUnknown
* Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, #hipErrorUnknown
* Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, #hipErrorPeerAccessAlreadyEnabled
*/
hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags);
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);
};
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);
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;
}