From 06e2dbd128177cf87ca2693c2cc5b49e4af69acd Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 6 Apr 2016 16:44:31 -0500 Subject: [PATCH] For P2P, use the peer list when allocating Device memory or pinned host. Each new allocation is automatically mapped into the address space of all enabled peers. [ROCm/hip commit: 813b0638884e484e2e7a2a28c349932c29e37593] --- projects/hip/CONTRIBUTING.md | 4 +- projects/hip/include/hcc_detail/hip_hcc.h | 6 ++- .../hip/include/hcc_detail/hip_runtime_api.h | 11 +++-- projects/hip/src/hip_hcc.cpp | 5 +- projects/hip/src/hip_memory.cpp | 47 ++++++++++--------- projects/hip/src/hip_peer.cpp | 24 +++++++--- 6 files changed, 59 insertions(+), 38 deletions(-) diff --git a/projects/hip/CONTRIBUTING.md b/projects/hip/CONTRIBUTING.md index b593cadd64..4d5050ca89 100644 --- a/projects/hip/CONTRIBUTING.md +++ b/projects/hip/CONTRIBUTING.md @@ -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. diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h index 05fa2818b6..23c66759f4 100644 --- a/projects/hip/include/hcc_detail/hip_hcc.h +++ b/projects/hip/include/hcc_detail/hip_hcc.h @@ -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: diff --git a/projects/hip/include/hcc_detail/hip_runtime_api.h b/projects/hip/include/hcc_detail/hip_runtime_api.h index f33895cd98..2f70ae47c3 100644 --- a/projects/hip/include/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hcc_detail/hip_runtime_api.h @@ -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); diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 0cae125509..e3239c4119 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -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); - }; diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index b78eb7af6e..dd7f374cba 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -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); diff --git a/projects/hip/src/hip_peer.cpp b/projects/hip/src/hip_peer.cpp index de86688ede..eaac47d82f 100644 --- a/projects/hip/src/hip_peer.cpp +++ b/projects/hip/src/hip_peer.cpp @@ -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; }