diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h index 1dd8777b1c..0d24862764 100644 --- a/projects/hip/include/hcc_detail/hip_hcc.h +++ b/projects/hip/include/hcc_detail/hip_hcc.h @@ -38,7 +38,7 @@ THE SOFTWARE. // Compile peer-to-peer support. // >= 2 : use HCC hc:accelerator::get_is_peer // >= 3 : use hc::am_memtracker_update_peers(...) -#define USE_PEER_TO_PEER 2 +#define USE_PEER_TO_PEER 3 //#define INLINE static inline @@ -526,6 +526,7 @@ public: bool addPeer(ihipDevice_t *peer); bool removePeer(ihipDevice_t *peer); + void resetPeers(ihipDevice_t *thisDevice); uint32_t peerCnt() const { return _peerCnt; }; hsa_agent_t *peerAgents() const { return _peerAgents; }; diff --git a/projects/hip/include/hcc_detail/hip_runtime_api.h b/projects/hip/include/hcc_detail/hip_runtime_api.h index 2f70ae47c3..58df3a2068 100644 --- a/projects/hip/include/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hcc_detail/hip_runtime_api.h @@ -907,10 +907,33 @@ hipError_t hipMemGetInfo (size_t * free, size_t * total) ; * * Returns "1" in @p canAccessPeer if the specified @p device is capable * of directly accessing memory physically located on peerDevice , or "0" if not. + * + * Returns "0" in @p canAccessPeer if deviceId == peerDeviceId, and both are valid devices : a device is not a peer of itself. + * + * + * + * @returns #hipSuccess, + * @returns #hipErrorInvalidDevice if deviceId or peerDeviceId are not valid devices */ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, 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 + * + * Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, + * @returns #hipErrorPeerAccessAlreadyEnabled if peer access is already enabled for this device. + */ +hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags); + /** * @brief Disable direct access from current device's virtual address space to memory allocations physically located on a peer device. @@ -923,20 +946,6 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDev */ 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 - * - * Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, #hipErrorPeerAccessAlreadyEnabled - */ -hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags); /** * @brief Copies memory from one device to memory on another device. diff --git a/projects/hip/samples/1_Utils/hipInfo/hipInfo.cpp b/projects/hip/samples/1_Utils/hipInfo/hipInfo.cpp index 9151d5058e..581194f624 100644 --- a/projects/hip/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/projects/hip/samples/1_Utils/hipInfo/hipInfo.cpp @@ -127,7 +127,7 @@ void printDeviceProp (int deviceId) for (int i=0; ilocked_waitAllStreams(); diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index ddf6bb1691..a9194c266a 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -240,6 +240,15 @@ bool ihipDeviceCriticalBase_t::removePeer(ihipDevice_t *peer) } } + +template<> +void ihipDeviceCriticalBase_t::resetPeers(ihipDevice_t *thisDevice) +{ + _peers.clear(); + _peerCnt = 0; + addPeer(thisDevice); // peer-list always contains self agent. +} + //------------------------------------------------------------------------------------------------- //--- @@ -444,12 +453,18 @@ 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? + + +#if USE_PEER_TO_PEER>=2 + // This resest peer list to just me: + crit->resetPeers(this); + #endif // Reset and release all memory stored in the tracker: + // Reset will remove peer mapping so don't need to do this explicitly. am_memtracker_reset(_acc); + }; @@ -474,10 +489,13 @@ void ihipDevice_t::init(unsigned device_index, unsigned deviceCnt, hc::accelerat getProperties(&_props); + _criticalData.init(deviceCnt); + + locked_reset(); + _default_stream = new ihipStream_t(device_index, acc.get_default_view(), hipStreamDefault); locked_addStream(_default_stream); - - _criticalData.init(deviceCnt); + tprintf(DB_SYNC, "created device with default_stream=%p\n", _default_stream); diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 9e8c963433..a3409a3e2e 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -132,7 +132,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) hc::am_memtracker_update(*ptr, device->_device_index, 0); { LockedAccessor_DeviceCrit_t crit(device->criticalData()); - if (crit->peerCnt()) { + if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved: hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); if (hsa_status != HSA_STATUS_SUCCESS) { hip_status = hipErrorMemoryAllocation; @@ -173,8 +173,9 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) }else{ hc::am_memtracker_update(*ptr, device->_device_index, flags); { + // TODO - allow_access only works for device memory, need to change am_alloc to allocate host directly. LockedAccessor_DeviceCrit_t crit(device->criticalData()); - if (crit->peerCnt()) { + if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved: hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); if (hsa_status != HSA_STATUS_SUCCESS) { hip_status = hipErrorMemoryAllocation; diff --git a/projects/hip/src/hip_peer.cpp b/projects/hip/src/hip_peer.cpp index aeb9c7fd49..612cd6e309 100644 --- a/projects/hip/src/hip_peer.cpp +++ b/projects/hip/src/hip_peer.cpp @@ -17,6 +17,8 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +#include + #include "hip_runtime.h" #include "hcc_detail/hip_hcc.h" #include "hcc_detail/trace_helper.h" @@ -35,11 +37,15 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDe auto peerDevice = ihipGetDevice(peerDeviceId); if ((thisDevice != NULL) && (peerDevice != NULL)) { + if (deviceId == peerDeviceId) { + *canAccessPeer = 0; + } else { #if USE_PEER_TO_PEER>=2 - *canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc); + *canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc); #else - *canAccessPeer = 0; + *canAccessPeer = 0; #endif + } } else { *canAccessPeer = 0; @@ -69,15 +75,15 @@ hipError_t hipDeviceDisablePeerAccess (int peerDeviceId) #endif if (! canAccessPeer) { err = hipErrorInvalidDevice; // P2P not allowed between these devices. + } else if (thisDevice == peerDevice) { + err = hipErrorInvalidDevice; // Can't disable peer access to self. } else { - - LockedAccessor_DeviceCrit_t crit(thisDevice->criticalData()); bool changed = crit->removePeer(peerDevice); if (changed) { #if USE_PEER_TO_PEER>=3 // Update the peers for all memory already saved in the tracker: - am_memtracker_update_peers(device->_acc, crit->peerCnt(), crit->peerAgents()); + am_memtracker_update_peers(thisDevice->_acc, crit->peerCnt(), crit->peerAgents()); #endif } else { err = hipErrorPeerAccessNotEnabled; // never enabled P2P access. @@ -92,9 +98,6 @@ hipError_t hipDeviceDisablePeerAccess (int peerDeviceId) }; -/** - * @warning Need to update this function when RT supports P2P - */ //--- // Enable registering memory on peerDevice for direct access from the current device. hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags) @@ -113,7 +116,7 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags) bool isNewPeer = crit->addPeer(peerDevice); if (isNewPeer) { #if USE_PEER_TO_PEER>=3 - am_memtracker_update_peers(device->_acc, crit->peerCnt(), crit->peerAgents()); + am_memtracker_update_peers(thisDevice->_acc, crit->peerCnt(), crit->peerAgents()); #endif } else { err = hipErrorPeerAccessAlreadyEnabled; diff --git a/projects/hip/tests/src/CMakeLists.txt b/projects/hip/tests/src/CMakeLists.txt index 28ea1bf10b..a75a48076c 100644 --- a/projects/hip/tests/src/CMakeLists.txt +++ b/projects/hip/tests/src/CMakeLists.txt @@ -8,6 +8,10 @@ include_directories( ${PROJECT_SOURCE_DIR}/include ) set (HIP_Unit_Test_VERSION_MAJOR 1) set (HIP_Unit_Test_VERSION_MINOR 0) +if(NOT DEFINED HIP_MULTI_GPU) + set(HIP_MULTI_GPU 0 CACHE BOOL "Run tests requiring more than one GPU") +endif() + if(NOT DEFINED HIP_BUILD_LOCAL) if(NOT DEFINED ENV{HIP_BUILD_LOCAL}) set(HIP_BUILD_LOCAL 0 CACHE BOOL "Build HIP in local folder") @@ -218,6 +222,12 @@ make_test(hipFuncDeviceSynchronize " ") make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-serial" --tests 0x1) make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-pyramid" --tests 0x4) make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-nearzero" --tests 0x10) -make_test(hipPeerToPeer_simple " " ) + +if (${HIP_MULTI_GPU}) + make_test(hipPeerToPeer_simple ) # use current device for copy, this fails. + make_test(hipPeerToPeer_simple --memcpyWithPeer) + make_test(hipPeerToPeer_simple --mirrorPeers) # mirror mapping: test to ensure mirror doesn't destroy orig mapping. + +endif() make_hipify_test(specialFunc.cu ) diff --git a/projects/hip/tests/src/hipPeerToPeer_simple.cpp b/projects/hip/tests/src/hipPeerToPeer_simple.cpp index 4aaa6a452b..5bfb583f3f 100644 --- a/projects/hip/tests/src/hipPeerToPeer_simple.cpp +++ b/projects/hip/tests/src/hipPeerToPeer_simple.cpp @@ -1,4 +1,3 @@ - /* Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. @@ -52,9 +51,11 @@ void parseMyArguments(int argc, char *argv[]) }; -int main(int argc, char *argv[]) +//--- +// Test which enables peer2peer first, then allocates the memory. +void enablePeerFirst() { - parseMyArguments(argc, argv); + printf ("\n==testing: %s\n", __func__); int deviceCnt; @@ -74,6 +75,110 @@ int main(int argc, char *argv[]) assert(canAccessPeer); + HIPCHECK (hipSetDevice(currentDevice)); + HIPCHECK(hipDeviceReset()); + HIPCHECK (hipSetDevice(peerDevice)); + HIPCHECK(hipDeviceReset()); + + HIPCHECK(hipSetDevice(currentDevice)); + HIPCHECK(hipDeviceEnablePeerAccess(peerDevice, 0)); + + if (p_mirrorPeers) { + int canAccessPeer; + HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, peerDevice, currentDevice)); + assert(canAccessPeer); + + HIPCHECK(hipSetDevice(peerDevice)); + HIPCHECK(hipDeviceEnablePeerAccess(currentDevice, 0)); + } + + size_t Nbytes = N*sizeof(char); + + char *A_d0, *A_d1; + char *A_h; + + A_h = (char*)malloc(Nbytes); + + // allocate and initialize memory on device0 + HIPCHECK (hipSetDevice(currentDevice)); + HIPCHECK (hipMalloc(&A_d0, Nbytes) ); + HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) ); + + // allocate and initialize memory on peer device + HIPCHECK (hipSetDevice(peerDevice)); + HIPCHECK (hipMalloc(&A_d1, Nbytes) ); + HIPCHECK ( hipMemset(A_d1, 0x13, Nbytes) ); + + + + // Device0 push to device1, using P2P: + HIPCHECK (hipSetDevice(p_memcpyWithPeer ? peerDevice : currentDevice)); + HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault)); + + // Copy data back to host: + HIPCHECK (hipSetDevice(peerDevice)); + HIPCHECK (hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost)); + + // Check host data: + for (int i=0; i