From d624c57d3922c6e4a2d2a21d778f56e636cfdc7a Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 1 Apr 2016 16:45:42 -0500 Subject: [PATCH 01/15] update proposed release notes [ROCm/clr commit: e180fa5b72b51e3c870b41a609f01b06bd022c8a] --- projects/clr/hipamd/RELEASE.md | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/RELEASE.md b/projects/clr/hipamd/RELEASE.md index 8041657b8e..3fdc529571 100644 --- a/projects/clr/hipamd/RELEASE.md +++ b/projects/clr/hipamd/RELEASE.md @@ -18,9 +18,14 @@ Next: - Move to CMake. - Split source into multiple modular .cpp and .h files. - Create static library and link. -- Deprecate hipDeviceGetProp, replace with hipGetDeviceProp + - Set HIP_PATH to install. +- Make hipDevice and hipStream thread-safe. + - Prefered hipStream usage is still to create new streams for each new thread, but it works even if you don;t. +- Improve automated platform detection: If AMD GPU is installed and detected by driver, default HIP_PLATFORM to hcc. +- HIP_TRACE_API now prints arguments to the HIP function (in addition to name of function). +- Deprecate hipDeviceGetProp (Replace with hipGetDeviceProp) - Deprecate hipMallocHost (Replace with hipHostMalloc) -- Deprecate hipFreeHost (Replace with hipHostFree). +- Deprecate hipFreeHost (Replace with hipHostFree) ## Revision History: From e195900a6ea22435a3c118355510571928dab8f7 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 6 Apr 2016 11:28:23 -0500 Subject: [PATCH 02/15] Checkpoint initial peer2peer implementation. [ROCm/clr commit: db91890f53276dc0067a3581372c4598e86471dc] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 27 ++++++++-- .../include/hcc_detail/hip_runtime_api.h | 10 ++-- projects/clr/hipamd/src/hip_hcc.cpp | 40 ++++++++++++++- projects/clr/hipamd/src/hip_peer.cpp | 51 ++++++++++++++++--- 4 files changed, 111 insertions(+), 17 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 42dd3b6df5..5c8aaf5659 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -494,11 +494,23 @@ struct ihipEvent_t { // will lock the mutex on construction and unlock on destruction. // // MUTEX_TYPE is template argument so can easily convert to FakeMutex for performance or stress testing. -template +template class ihipDeviceCriticalBase_t : LockedBase { public: - ihipDeviceCriticalBase_t() : _stream_id(0) {}; + ihipDeviceCriticalBase_t() : _stream_id(0), _peerAgents(nullptr) {}; + + void init(unsigned deviceCnt) { + assert(_peerAgents == nullptr); + _peerAgents = new hsa_agent_t[deviceCnt]; + }; + + ~ihipDeviceCriticalBase_t() { + if (_peerAgents != nullptr) { + delete _peerAgents; + _peerAgents = nullptr; + } + } friend class LockedAccessor; std::list &streams() { return _streams; }; @@ -507,10 +519,19 @@ public: // "Allocate" a stream ID: ihipStream_t::SeqNum_t incStreamId() { return _stream_id++; }; + void recomputePeerAgents(); + void addPeer(ihipDevice_t *peer); + void removePeer(ihipDevice_t *peer); + private: std::list _streams; // streams associated with this device. ihipStream_t::SeqNum_t _stream_id; + + // These reflect the currently Enabled set of peers for this GPU: + std::list _peers; // list of enabled peer devices. + uint32_t _peerCnt; // number of enabled peers + hsa_agent_t *_peerAgents; // efficient packed array of enabled agents (to use for allocations.) }; // Note Mutex selected based on DeviceMutex @@ -530,7 +551,7 @@ class ihipDevice_t { public: // Functions: ihipDevice_t() {}; // note: calls constructor for _criticalData - void init(unsigned device_index, hc::accelerator &acc, unsigned flags); + void init(unsigned device_index, unsigned deviceCnt, hc::accelerator &acc, unsigned flags); ~ihipDevice_t(); void locked_addStream(ihipStream_t *s); diff --git a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h index b5fdb312a4..56b3560427 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h @@ -908,12 +908,12 @@ 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. */ -hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int device, int peerDevice ); +hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDeviceId); /** - * @brief Disables registering memory on peerDevice for direct access from the current device. + * @brief Disable registering memory on peerDevice for direct access from the current device. * * If there are any allocations on peerDevice which were registered in the current device using hipPeerRegister() then these allocations will be automatically unregistered. * Returns hipErrorPeerAccessNotEnabled if direct access to memory on peerDevice has not yet been enabled from the current device. @@ -922,10 +922,10 @@ hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int device, int peerDe * TODO:cudaErrorPeerAccessNotEnabled and cudaErrorInvalidDevice error not supported in HIP, return hipErrorUnknown * Returns #hipSuccess, #hipErrorUnknown */ -hipError_t hipDeviceDisablePeerAccess ( int peerDevice ); +hipError_t hipDeviceDisablePeerAccess (int peerDeviceId); /** - * @brief Enables registering memory on peerDevice for direct access from the current device. + * @brief Enable registering memory on peerDevice for direct access from the current device. * * @param [in] peerDevice * @param [in] flags @@ -933,7 +933,7 @@ hipError_t hipDeviceDisablePeerAccess ( int peerDevice ); * TODO:cudaErrorInvalidDevice error not supported in HIP, return hipErrorUnknown * Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, #hipErrorUnknown */ -hipError_t hipDeviceEnablePeerAccess ( int peerDevice, unsigned int flags ); +hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags); /** * @brief Copies memory from one device to memory on another device. diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 6c80b84e0e..bbf6538817 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -198,6 +198,40 @@ void ihipStream_t::locked_wait(bool assertQueueEmpty) }; +// Recompute the peercnt and the packed _peerAgents whenever a peer is added or deleted. +// The packed _peerAgents can efficiently be used on each memory allocation. +template<> +void ihipDeviceCriticalBase_t::recomputePeerAgents() +{ + _peerCnt = 0; + std::for_each (_peers.begin(), _peers.end(), [this](ihipDevice_t* device) { + _peerAgents[_peerCnt++] = device->_hsa_agent; + }); +} + + +template<> +void ihipDeviceCriticalBase_t::addPeer(ihipDevice_t *peer) +{ + auto match = std::find(_peers.begin(), _peers.end(), peer); + if (match != std::end(_peers)) { + _peers.push_back(peer); + recomputePeerAgents(); + return; + } + + // If we get here - peer was already on list, silently ignore. +} + + +template<> +void ihipDeviceCriticalBase_t::removePeer(ihipDevice_t *peer) +{ + _peers.remove(peer); + recomputePeerAgents(); +} + +//------------------------------------------------------------------------------------------------- //--- ihipDevice_t * ihipStream_t::getDevice() const @@ -408,7 +442,7 @@ void ihipDevice_t::locked_reset() //--- -void ihipDevice_t::init(unsigned device_index, hc::accelerator &acc, unsigned flags) +void ihipDevice_t::init(unsigned device_index, unsigned deviceCnt, hc::accelerator &acc, unsigned flags) { _device_index = device_index; _device_flags = flags; @@ -431,6 +465,8 @@ void ihipDevice_t::init(unsigned device_index, hc::accelerator &acc, unsigned fl _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); hsa_region_t *pinnedHostRegion; @@ -905,7 +941,7 @@ void ihipInit() //If device is not in visible devices list, ignore continue; } - g_devices[g_deviceCnt].init(g_deviceCnt, accs[i], hipDeviceMapHost); + g_devices[g_deviceCnt].init(g_deviceCnt, deviceCnt, accs[i], hipDeviceMapHost); g_deviceCnt++; } } diff --git a/projects/clr/hipamd/src/hip_peer.cpp b/projects/clr/hipamd/src/hip_peer.cpp index d4de6053de..c4a8a1bcf3 100644 --- a/projects/clr/hipamd/src/hip_peer.cpp +++ b/projects/clr/hipamd/src/hip_peer.cpp @@ -20,22 +20,44 @@ 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 */ //--- -hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int device, int peerDevice ) +hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int deviceId, int peerDeviceId) { - HIP_INIT_API(canAccessPeer, device, peerDevice); + HIP_INIT_API(canAccessPeer, deviceId, peerDeviceId); + hipError_t err = hipSuccess; + +#if USE_PEER_TO_PEER + auto device = ihipGetDevice(deviceId); + auto peerDevice = ihipGetDevice(peerDeviceId); + + if ((device != NULL) && (peerDevice != NULL)) { +#if USE_PEER_TO_PEER==2 + *canAccessPeer = peerDevice->_acc.is_peer(device->_acc); +#else + *canAccessPeer = 0; +#endif + + } else { + *canAccessPeer = false; + err = hipErrorInvalidDevice; + } + + +#else *canAccessPeer = false; - return ihipLogStatus(hipSuccess); +#endif + return ihipLogStatus(err); } /** - * @warning Need to update this function when RT supports P2P + * warning Need to update this function when RT supports P2P */ //--- hipError_t hipDeviceDisablePeerAccess ( int peerDevice ) @@ -51,11 +73,26 @@ hipError_t hipDeviceDisablePeerAccess ( int peerDevice ) * @warning Need to update this function when RT supports P2P */ //--- -hipError_t hipDeviceEnablePeerAccess ( int peerDevice, unsigned int flags ) + // Enable registering memory on peerDevice for direct access from the current device. +hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags) { std::call_once(hip_initialized, ihipInit); - // TODO-p2p - return ihipLogStatus(hipSuccess); + + hipError_t err = hipSuccess; +#if USE_PEER_TO_PEER + if (flags != 0) { + err = hipErrorInvalidValue; + } else { + auto peerDevice = ihipGetDevice(peerDeviceId); + if (peerDevice != NULL) { + + } else { + err = hipErrorInvalidDevice; + } + } +#endif + + return ihipLogStatus(err); } From 6ccb0a5054bf13ebebcd8633b34c3fb3ac70922e Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 6 Apr 2016 15:49:32 -0500 Subject: [PATCH 03/15] P2P checkpoint. Maintain enabled peer tables for each device. [ROCm/clr commit: 6a182ce788bf28904d9456282718ca48903c1dc6] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 12 +++- .../include/hcc_detail/hip_runtime_api.h | 18 +++--- projects/clr/hipamd/include/hip_runtime_api.h | 4 +- projects/clr/hipamd/src/hip_error.cpp | 2 +- projects/clr/hipamd/src/hip_hcc.cpp | 52 +++++++++------- projects/clr/hipamd/src/hip_peer.cpp | 60 +++++++++++++------ 6 files changed, 96 insertions(+), 52 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 5c8aaf5659..05fa2818b6 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -519,9 +519,11 @@ public: // "Allocate" a stream ID: ihipStream_t::SeqNum_t incStreamId() { return _stream_id++; }; - void recomputePeerAgents(); - void addPeer(ihipDevice_t *peer); - void removePeer(ihipDevice_t *peer); + bool addPeer(ihipDevice_t *peer); + bool removePeer(ihipDevice_t *peer); + + uint32_t peerCnt() const { return _peerCnt; }; + uint32_t peerAgents() const { return _peerAgents; }; private: @@ -532,6 +534,8 @@ private: std::list _peers; // list of enabled peer devices. uint32_t _peerCnt; // number of enabled peers hsa_agent_t *_peerAgents; // efficient packed array of enabled agents (to use for allocations.) +private: + void recomputePeerAgents(); }; // Note Mutex selected based on DeviceMutex @@ -560,6 +564,8 @@ public: // Functions: void locked_waitAllStreams(); void locked_syncDefaultStream(bool waitOnSelf); + ihipDeviceCritical_t &criticalData() { return _criticalData; }; // TODO, move private. Fix P2P. + public: // Data, set at initialization: unsigned _device_index; // index into g_devices. diff --git a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h index 56b3560427..f33895cd98 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h @@ -913,21 +913,21 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDev /** - * @brief Disable registering memory on peerDevice for direct access from the current device. + * @brief Disable direct access from current device's virtual address space to memory allocations physically located on a peer device. * - * If there are any allocations on peerDevice which were registered in the current device using hipPeerRegister() then these allocations will be automatically unregistered. * Returns hipErrorPeerAccessNotEnabled if direct access to memory on peerDevice has not yet been enabled from the current device. * - * @param [in] peerDevice + * @param [in] peerDeviceId + * * TODO:cudaErrorPeerAccessNotEnabled and cudaErrorInvalidDevice error not supported in HIP, return hipErrorUnknown * Returns #hipSuccess, #hipErrorUnknown */ hipError_t hipDeviceDisablePeerAccess (int peerDeviceId); /** - * @brief Enable registering memory on peerDevice for direct access from the current device. + * @brief Enable direct access from current device's virtual address space to memory allocations physically located on a peer device. * - * @param [in] peerDevice + * @param [in] peerDeviceId * @param [in] flags * * TODO:cudaErrorInvalidDevice error not supported in HIP, return hipErrorUnknown @@ -939,14 +939,14 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags); * @brief Copies memory from one device to memory on another device. * * @param [out] dst - Destination device pointer. - * @param [in] dstDevice - Destination device + * @param [in] dstDeviceId - Destination device * @param [in] src - Source device pointer - * @param [in] srcDevice - Source device + * @param [in] srcDeviceId - Source device * @param [in] sizeBytes - Size of memory copy in bytes * * Returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDevice */ -hipError_t hipMemcpyPeer ( void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes ); +hipError_t hipMemcpyPeer (void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t sizeBytes); /** * @brief Copies memory from one device to memory on another device. @@ -961,7 +961,7 @@ hipError_t hipMemcpyPeer ( void* dst, int dstDevice, const void* src, int srcD * Returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDevice */ #if __cplusplus -hipError_t hipMemcpyPeerAsync ( void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream=0 ); +hipError_t hipMemcpyPeerAsync ( void* dst, int dstDeviceId, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream=0 ); #else hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream); #endif diff --git a/projects/clr/hipamd/include/hip_runtime_api.h b/projects/clr/hipamd/include/hip_runtime_api.h index bef881d066..ca49ab5d13 100644 --- a/projects/clr/hipamd/include/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip_runtime_api.h @@ -138,7 +138,7 @@ typedef struct hipPointerAttribute_t { * @enum * @ingroup Enumerations */ -// Developer note - when updating these, update the hipErrorName and hipErrorString functions +// Developer note - when updating these, update the hipErrorName and hipErrorString functions in NVCC and HCC paths typedef enum hipError_t { hipSuccess = 0 ///< Successful completion. ,hipErrorMemoryAllocation ///< Memory allocation error. @@ -155,6 +155,8 @@ typedef enum hipError_t { ,hipErrorNoDevice ///< Call to hipGetDeviceCount returned 0 devices ,hipErrorNotReady ///< Indicates that asynchronous operations enqueued earlier are not ready. This is not actually an error, but is used to distinguish from hipSuccess (which indicates completion). APIs that return this error include hipEventQuery and hipStreamQuery. ,hipErrorUnknown ///< Unknown error. + ,hipErrorPeerAccessNotEnabled ///< Peer access was never enabled from the current device. + ,hipErrorPeerAccessAlreadyEnabled ///< Peer access was already enabled from the current device. ,hipErrorRuntimeMemory ///< HSA runtime memory call returned error. Typically not seen in production systems. ,hipErrorRuntimeOther ///< HSA runtime call other than memory returned error. Typically not seen in production systems. ,hipErrorTbd ///< Marker that more error codes are needed. diff --git a/projects/clr/hipamd/src/hip_error.cpp b/projects/clr/hipamd/src/hip_error.cpp index 655ab3b8c7..7c723b1aa2 100644 --- a/projects/clr/hipamd/src/hip_error.cpp +++ b/projects/clr/hipamd/src/hip_error.cpp @@ -65,7 +65,7 @@ const char *hipGetErrorName(hipError_t hip_error) //--- const char *hipGetErrorString(hipError_t hip_error) { - std::call_once(hip_initialized, ihipInit); + HIP_INIT_API(hip_error); // TODO - return a message explaining the error. // TODO - This should be set up to return the same string reported in the the doxygen comments, somehow. diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index bbf6538817..992dbca724 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -211,24 +211,31 @@ void ihipDeviceCriticalBase_t::recomputePeerAgents() template<> -void ihipDeviceCriticalBase_t::addPeer(ihipDevice_t *peer) +bool ihipDeviceCriticalBase_t::addPeer(ihipDevice_t *peer) { auto match = std::find(_peers.begin(), _peers.end(), peer); if (match != std::end(_peers)) { _peers.push_back(peer); recomputePeerAgents(); - return; + return true; } // If we get here - peer was already on list, silently ignore. + return false; } template<> -void ihipDeviceCriticalBase_t::removePeer(ihipDevice_t *peer) +bool ihipDeviceCriticalBase_t::removePeer(ihipDevice_t *peer) { - _peers.remove(peer); - recomputePeerAgents(); + auto match = std::find(_peers.begin(), _peers.end(), peer); + if (match != std::end(_peers)) { + _peers.remove(peer); + recomputePeerAgents(); + return true; + } else { + return false; + } } //------------------------------------------------------------------------------------------------- @@ -1073,22 +1080,25 @@ void ihipPostLaunchKernel(hipStream_t stream, hc::completion_future &kernelFutur const char *ihipErrorString(hipError_t hip_error) { switch (hip_error) { - case hipSuccess : return "hipSuccess"; - case hipErrorMemoryAllocation : return "hipErrorMemoryAllocation"; - case hipErrorMemoryFree : return "hipErrorMemoryFree"; - case hipErrorUnknownSymbol : return "hipErrorUnknownSymbol"; - case hipErrorOutOfResources : return "hipErrorOutOfResources"; - case hipErrorInvalidValue : return "hipErrorInvalidValue"; - case hipErrorInvalidResourceHandle : return "hipErrorInvalidResourceHandle"; - case hipErrorInvalidDevice : return "hipErrorInvalidDevice"; - case hipErrorInvalidMemcpyDirection : return "hipErrorInvalidMemcpyDirection"; - case hipErrorNoDevice : return "hipErrorNoDevice"; - case hipErrorNotReady : return "hipErrorNotReady"; - case hipErrorRuntimeMemory : return "hipErrorRuntimeMemory"; - case hipErrorRuntimeOther : return "hipErrorRuntimeOther"; - case hipErrorUnknown : return "hipErrorUnknown"; - case hipErrorTbd : return "hipErrorTbd"; - default : return "hipErrorUnknown"; + case hipSuccess : return "hipSuccess"; + case hipErrorMemoryAllocation : return "hipErrorMemoryAllocation"; + case hipErrorMemoryFree : return "hipErrorMemoryFree"; + case hipErrorUnknownSymbol : return "hipErrorUnknownSymbol"; + case hipErrorOutOfResources : return "hipErrorOutOfResources"; + case hipErrorInvalidValue : return "hipErrorInvalidValue"; + case hipErrorInvalidResourceHandle : return "hipErrorInvalidResourceHandle"; + case hipErrorInvalidDevice : return "hipErrorInvalidDevice"; + case hipErrorInvalidMemcpyDirection : return "hipErrorInvalidMemcpyDirection"; + case hipErrorNoDevice : return "hipErrorNoDevice"; + case hipErrorNotReady : return "hipErrorNotReady"; + case hipErrorPeerAccessNotEnabled : return "hipErrorPeerAccessNotEnabled"; + case hipErrorPeerAccessAlreadyEnabled : return "hipErrorPeerAccessAlreadyEnabled"; + + case hipErrorRuntimeMemory : return "hipErrorRuntimeMemory"; + case hipErrorRuntimeOther : return "hipErrorRuntimeOther"; + case hipErrorUnknown : return "hipErrorUnknown"; + case hipErrorTbd : return "hipErrorTbd"; + default : return "hipErrorUnknown"; }; }; diff --git a/projects/clr/hipamd/src/hip_peer.cpp b/projects/clr/hipamd/src/hip_peer.cpp index c4a8a1bcf3..de86688ede 100644 --- a/projects/clr/hipamd/src/hip_peer.cpp +++ b/projects/clr/hipamd/src/hip_peer.cpp @@ -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); }; From e2879330512fa9d938082b2f58cbc68285df4a60 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 6 Apr 2016 16:44:31 -0500 Subject: [PATCH 04/15] 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/clr commit: 288682ccb324bfe3b0605735b149be36d659906f] --- projects/clr/hipamd/CONTRIBUTING.md | 4 +- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 6 ++- .../include/hcc_detail/hip_runtime_api.h | 11 +++-- projects/clr/hipamd/src/hip_hcc.cpp | 5 +- projects/clr/hipamd/src/hip_memory.cpp | 47 ++++++++++--------- projects/clr/hipamd/src/hip_peer.cpp | 24 +++++++--- 6 files changed, 59 insertions(+), 38 deletions(-) diff --git a/projects/clr/hipamd/CONTRIBUTING.md b/projects/clr/hipamd/CONTRIBUTING.md index b593cadd64..4d5050ca89 100644 --- a/projects/clr/hipamd/CONTRIBUTING.md +++ b/projects/clr/hipamd/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/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 05fa2818b6..23c66759f4 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/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/clr/hipamd/include/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h index f33895cd98..2f70ae47c3 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/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/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 992dbca724..20c6f42451 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/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/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 4d75a1d7a1..4cbec08ba0 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/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/clr/hipamd/src/hip_peer.cpp b/projects/clr/hipamd/src/hip_peer.cpp index de86688ede..eaac47d82f 100644 --- a/projects/clr/hipamd/src/hip_peer.cpp +++ b/projects/clr/hipamd/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; } From 038d8055bc0a3e07618a263f465f81e6a14e14ed Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 7 Apr 2016 14:28:42 -0500 Subject: [PATCH 05/15] Use HCC get_is_pool, add USE_PEER_TO_PEER level 2, 3. [ROCm/clr commit: bae62e02dc767ffa437a0e27cbe1f837f03a4d50] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 5 +++-- projects/clr/hipamd/src/hip_peer.cpp | 17 +++++++++-------- 2 files changed, 12 insertions(+), 10 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 23c66759f4..1dd8777b1c 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -36,8 +36,9 @@ THE SOFTWARE. #define USE_AV_COPY 0 // Compile peer-to-peer support. -// 2= use upcoming HCC APIs. -#define USE_PEER_TO_PEER 1 +// >= 2 : use HCC hc:accelerator::get_is_peer +// >= 3 : use hc::am_memtracker_update_peers(...) +#define USE_PEER_TO_PEER 2 //#define INLINE static inline diff --git a/projects/clr/hipamd/src/hip_peer.cpp b/projects/clr/hipamd/src/hip_peer.cpp index eaac47d82f..2e52f8e586 100644 --- a/projects/clr/hipamd/src/hip_peer.cpp +++ b/projects/clr/hipamd/src/hip_peer.cpp @@ -32,12 +32,12 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDe hipError_t err = hipSuccess; #if USE_PEER_TO_PEER - auto device = ihipGetDevice(deviceId); + auto thisDevice = ihipGetDevice(deviceId); auto peerDevice = ihipGetDevice(peerDeviceId); - if ((device != NULL) && (peerDevice != NULL)) { -#if USE_PEER_TO_PEER==2 - *canAccessPeer = peerDevice->_acc.get_is_peer(device->_acc); + if ((thisDevice != NULL) && (peerDevice != NULL)) { +#if USE_PEER_TO_PEER>=2 + *canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc); #else *canAccessPeer = 0; #endif @@ -66,8 +66,8 @@ hipError_t hipDeviceDisablePeerAccess (int peerDeviceId) auto thisDevice = ihipGetTlsDefaultDevice(); auto peerDevice = ihipGetDevice(peerDeviceId); if ((thisDevice != NULL) && (peerDevice != NULL)) { -#if USE_PEER_TO_PEER==2 - bool canAccessPeer = peerDevice->_acc.get_is_peer(device->_acc); +#if USE_PEER_TO_PEER>=2 + bool canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc); #else bool canAccessPeer = 0; #endif @@ -79,7 +79,8 @@ hipError_t hipDeviceDisablePeerAccess (int peerDeviceId) LockedAccessor_DeviceCrit_t crit(thisDevice->criticalData()); bool changed = crit->removePeer(peerDevice); if (changed) { -#if USE_PEER_TO_PEER==2 +#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()); #endif } else { @@ -115,7 +116,7 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags) LockedAccessor_DeviceCrit_t crit(thisDevice->criticalData()); bool isNewPeer = crit->addPeer(peerDevice); if (isNewPeer) { -#if USE_PEER_TO_PEER==2 +#if USE_PEER_TO_PEER>=3 am_memtracker_update_peers(device->_acc, crit->peerCnt(), crit->peerAgents()); #endif } else { From 10d8142b58b4bffbf1c71be11434255d53d79b1b Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 7 Apr 2016 15:51:08 -0500 Subject: [PATCH 06/15] Add simple P2P test [ROCm/clr commit: 5b37639efbd16f157c91660853d7ce965462f6d6] --- projects/clr/hipamd/tests/src/CMakeLists.txt | 2 + .../hipamd/tests/src/hipPeerToPeer_simple.cpp | 87 +++++++++++++++++++ 2 files changed, 89 insertions(+) create mode 100644 projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp diff --git a/projects/clr/hipamd/tests/src/CMakeLists.txt b/projects/clr/hipamd/tests/src/CMakeLists.txt index a0413eaad8..28ea1bf10b 100644 --- a/projects/clr/hipamd/tests/src/CMakeLists.txt +++ b/projects/clr/hipamd/tests/src/CMakeLists.txt @@ -170,6 +170,7 @@ make_hip_executable (hipFuncSetDeviceFlags hipFuncSetDeviceFlags.cpp) make_hip_executable (hipFuncGetDevice hipFuncGetDevice.cpp) make_hip_executable (hipFuncSetDevice hipFuncSetDevice.cpp) make_hip_executable (hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp) +make_hip_executable (hipPeerToPeer_simple hipPeerToPeer_simple.cpp) make_hip_executable (hipMultiThreadDevice hipMultiThreadDevice.cpp) @@ -217,5 +218,6 @@ 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 " " ) make_hipify_test(specialFunc.cu ) diff --git a/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp b/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp new file mode 100644 index 0000000000..79401d28b0 --- /dev/null +++ b/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp @@ -0,0 +1,87 @@ + +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +// Simple test for memset. +// Also serves as a template for other tests. + +#include "hip_runtime.h" +#include "test_common.h" + + +int main(int argc, char *argv[]) +{ + + HipTest::parseStandardArguments(argc, argv, true); + + int deviceCnt; + + HIPCHECK(hipGetDeviceCount(&deviceCnt)); + HIPCHECK(hipSetDevice(p_gpuDevice)); + + int peerDevice = ((p_gpuDevice + 1) % deviceCnt); + + printf ("N=%zu device=%d peerDevice=%d (%d devices total)\n", N, p_gpuDevice, peerDevice, deviceCnt); + + // Must be on a multi-gpu system: + assert (p_gpuDevice != peerDevice); + + int canAccessPeer; + HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, p_gpuDevice, peerDevice)); + printf ("dev#%d canAccessPeer:#%d=%d\n", p_gpuDevice, peerDevice, canAccessPeer); + + assert(canAccessPeer); + + HIPCHECK(hipDeviceEnablePeerAccess(peerDevice, 0)); + + size_t Nbytes = N*sizeof(float); + + float *A_d0, *A_d1; + float *A_h; + + A_h = (float*)malloc(Nbytes); + HIPCHECK (hipSetDevice(peerDevice)); + HIPCHECK (hipMalloc(&A_d1, Nbytes) ); + + HIPCHECK (hipSetDevice(p_gpuDevice)); + HIPCHECK (hipMalloc(&A_d0, Nbytes) ); + + + // Set memory on first device. + HIPCHECK (hipSetDevice(p_gpuDevice)); + HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) ); + + // Device0 push to device1, using P2P: + HIPCHECK ( hipMemcpy(A_d0, A_d1, Nbytes, hipMemcpyDefault)); + + // Copy data back to host: + HIPCHECK ( hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost)); + + // Check host data: + for (int i=0; i Date: Fri, 8 Apr 2016 02:15:46 -0500 Subject: [PATCH 07/15] Print peers in hipConfig. Also include peer APIs in vim hilighting. [ROCm/clr commit: b1b1db28253f14df632b4ae5414eda22cc86614b] --- .../hipamd/samples/1_Utils/hipInfo/hipInfo.cpp | 17 +++++++++++++++++ projects/clr/hipamd/util/vim/hip.vim | 6 ++++++ 2 files changed, 23 insertions(+) diff --git a/projects/clr/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp b/projects/clr/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp index 146d17e015..581194f624 100644 --- a/projects/clr/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/projects/clr/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp @@ -120,8 +120,25 @@ void printDeviceProp (int deviceId) cout << setw(w1) << "arch.hasSurfaceFuncs: " << props.arch.hasSurfaceFuncs << endl; cout << setw(w1) << "arch.has3dGrid: " << props.arch.has3dGrid << endl; cout << setw(w1) << "arch.hasDynamicParallelism: " << props.arch.hasDynamicParallelism << endl; + + int deviceCnt; + hipGetDeviceCount(&deviceCnt); + cout << setw(w1) << "peers: "; + for (int i=0; i Date: Fri, 8 Apr 2016 02:17:29 -0500 Subject: [PATCH 08/15] Use HIP_PATH if set else use relative ../... [ROCm/clr commit: 90d986ae472ed61d67befc6adf238af23259f026] --- projects/clr/hipamd/samples/0_Intro/bit_extract/Makefile | 2 +- projects/clr/hipamd/samples/0_Intro/square/Makefile | 2 +- projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/Makefile | 2 +- projects/clr/hipamd/samples/1_Utils/hipDispatchLatency/Makefile | 2 +- projects/clr/hipamd/samples/1_Utils/hipInfo/Makefile | 2 +- 5 files changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/clr/hipamd/samples/0_Intro/bit_extract/Makefile b/projects/clr/hipamd/samples/0_Intro/bit_extract/Makefile index 39fb5cf8c6..cdf793363b 100644 --- a/projects/clr/hipamd/samples/0_Intro/bit_extract/Makefile +++ b/projects/clr/hipamd/samples/0_Intro/bit_extract/Makefile @@ -1,6 +1,6 @@ #Dependencies : [MYHIP]/bin must be in user's path. -HIP_PATH=../../.. +HIP_PATH=?../../.. HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform) HIPCC=$(HIP_PATH)/bin/hipcc diff --git a/projects/clr/hipamd/samples/0_Intro/square/Makefile b/projects/clr/hipamd/samples/0_Intro/square/Makefile index 98ee0be4f6..817c556b26 100644 --- a/projects/clr/hipamd/samples/0_Intro/square/Makefile +++ b/projects/clr/hipamd/samples/0_Intro/square/Makefile @@ -1,4 +1,4 @@ -HIP_PATH=../../.. +HIP_PATH?=../../.. HIPCC=$(HIP_PATH)/bin/hipcc all: square.hip.out diff --git a/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/Makefile b/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/Makefile index a713379d8d..4599cacba2 100644 --- a/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/Makefile +++ b/projects/clr/hipamd/samples/1_Utils/hipBusBandwidth/Makefile @@ -1,4 +1,4 @@ -HIP_PATH=../../.. +HIP_PATH?=../../.. HIPCC=$(HIP_PATH)/bin/hipcc EXE=hipBusBandwidth diff --git a/projects/clr/hipamd/samples/1_Utils/hipDispatchLatency/Makefile b/projects/clr/hipamd/samples/1_Utils/hipDispatchLatency/Makefile index 9b2d558114..87e707923d 100644 --- a/projects/clr/hipamd/samples/1_Utils/hipDispatchLatency/Makefile +++ b/projects/clr/hipamd/samples/1_Utils/hipDispatchLatency/Makefile @@ -1,4 +1,4 @@ -HIP_PATH=../../.. +HIP_PATH?=../../.. HIPCC=$(HIP_PATH)/bin/hipcc EXE=hipDispatchLatency diff --git a/projects/clr/hipamd/samples/1_Utils/hipInfo/Makefile b/projects/clr/hipamd/samples/1_Utils/hipInfo/Makefile index f38f157bcb..d69067388e 100644 --- a/projects/clr/hipamd/samples/1_Utils/hipInfo/Makefile +++ b/projects/clr/hipamd/samples/1_Utils/hipInfo/Makefile @@ -1,4 +1,4 @@ -HIP_PATH=../../.. +HIP_PATH?=../../.. HIPCC=$(HIP_PATH)/bin/hipcc EXE=hipInfo From db2e56794553352b21c0b68057d9089cbab8e495 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 8 Apr 2016 09:22:24 -0500 Subject: [PATCH 09/15] Fix simple P2P test, had device ptrs swapped [ROCm/clr commit: 2b45fc35cbf33d5198a552b9f8c72375629b6dd2] --- projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp b/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp index 79401d28b0..ca2501438d 100644 --- a/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp +++ b/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp @@ -52,12 +52,12 @@ int main(int argc, char *argv[]) HIPCHECK(hipDeviceEnablePeerAccess(peerDevice, 0)); - size_t Nbytes = N*sizeof(float); + size_t Nbytes = N*sizeof(char); - float *A_d0, *A_d1; - float *A_h; + char *A_d0, *A_d1; + char *A_h; - A_h = (float*)malloc(Nbytes); + A_h = (char*)malloc(Nbytes); HIPCHECK (hipSetDevice(peerDevice)); HIPCHECK (hipMalloc(&A_d1, Nbytes) ); @@ -70,7 +70,7 @@ int main(int argc, char *argv[]) HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) ); // Device0 push to device1, using P2P: - HIPCHECK ( hipMemcpy(A_d0, A_d1, Nbytes, hipMemcpyDefault)); + HIPCHECK ( hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault)); // Copy data back to host: HIPCHECK ( hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost)); From b6071f5eca0f36fb4e050f5d35acab3220387fb3 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 9 Apr 2016 04:10:57 -0500 Subject: [PATCH 10/15] fix bugs in P2P implementation - addPeers polarity reversed, would never add. - check allow_access return value, pipe error to hipMalloc. [ROCm/clr commit: 7ca06d2fb164f856bda1218340f62c5132e8e2ae] --- projects/clr/hipamd/src/hip_hcc.cpp | 4 +++- projects/clr/hipamd/src/hip_memory.cpp | 10 ++++++++-- projects/clr/hipamd/src/hip_peer.cpp | 7 ++----- 3 files changed, 13 insertions(+), 8 deletions(-) diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 20c6f42451..ddf6bb1691 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -214,7 +214,8 @@ template<> bool ihipDeviceCriticalBase_t::addPeer(ihipDevice_t *peer) { auto match = std::find(_peers.begin(), _peers.end(), peer); - if (match != std::end(_peers)) { + if (match == std::end(_peers)) { + // Not already a peer, let's update the list: _peers.push_back(peer); recomputePeerAgents(); return true; @@ -230,6 +231,7 @@ bool ihipDeviceCriticalBase_t::removePeer(ihipDevice_t *peer) { auto match = std::find(_peers.begin(), _peers.end(), peer); if (match != std::end(_peers)) { + // Found a valid peer, let's remove it. _peers.remove(peer); recomputePeerAgents(); return true; diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 4cbec08ba0..2a4ec205e2 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -133,7 +133,10 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) { LockedAccessor_DeviceCrit_t crit(device->criticalData()); if (crit->peerCnt()) { - hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); + 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; + } } } } @@ -172,7 +175,10 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { LockedAccessor_DeviceCrit_t crit(device->criticalData()); if (crit->peerCnt()) { - hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); + 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/clr/hipamd/src/hip_peer.cpp b/projects/clr/hipamd/src/hip_peer.cpp index 2e52f8e586..ad9cf14634 100644 --- a/projects/clr/hipamd/src/hip_peer.cpp +++ b/projects/clr/hipamd/src/hip_peer.cpp @@ -31,26 +31,23 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDe hipError_t err = hipSuccess; -#if USE_PEER_TO_PEER auto thisDevice = ihipGetDevice(deviceId); auto peerDevice = ihipGetDevice(peerDeviceId); if ((thisDevice != NULL) && (peerDevice != NULL)) { #if USE_PEER_TO_PEER>=2 *canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc); + printf ("canAccessPeer=%d\n", *canAccessPeer); #else *canAccessPeer = 0; #endif } else { - *canAccessPeer = false; + *canAccessPeer = 0; err = hipErrorInvalidDevice; } -#else - *canAccessPeer = false; -#endif return ihipLogStatus(err); } From 7dbdd61088ae147d19e6652d6052dd63d02864f8 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 9 Apr 2016 04:35:06 -0500 Subject: [PATCH 11/15] Remove stray debug msgs, hipInfo don't print self as peer. [ROCm/clr commit: 715940672c5b1932c7f44224bc39c87d73684274] --- projects/clr/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp | 2 +- projects/clr/hipamd/src/hip_memory.cpp | 2 -- projects/clr/hipamd/src/hip_peer.cpp | 1 - 3 files changed, 1 insertion(+), 4 deletions(-) diff --git a/projects/clr/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp b/projects/clr/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp index 581194f624..9151d5058e 100644 --- a/projects/clr/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/projects/clr/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp @@ -127,7 +127,7 @@ void printDeviceProp (int deviceId) for (int i=0; i_acc, &deviceMemSize, &hostMemSize, &userMemSize); - printf ("deviceMemSize=%zu\n", deviceMemSize); - *free = hipDevice->_props.totalGlobalMem - deviceMemSize; } diff --git a/projects/clr/hipamd/src/hip_peer.cpp b/projects/clr/hipamd/src/hip_peer.cpp index ad9cf14634..aeb9c7fd49 100644 --- a/projects/clr/hipamd/src/hip_peer.cpp +++ b/projects/clr/hipamd/src/hip_peer.cpp @@ -37,7 +37,6 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDe if ((thisDevice != NULL) && (peerDevice != NULL)) { #if USE_PEER_TO_PEER>=2 *canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc); - printf ("canAccessPeer=%d\n", *canAccessPeer); #else *canAccessPeer = 0; #endif From dbdd0e0037327324de4edc3c733fbe3a50b415f3 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 9 Apr 2016 04:47:07 -0500 Subject: [PATCH 12/15] Improve P2P test. Add option to select which device does the copy. [ROCm/clr commit: 30fce5b3d832a058299ef6ed07e60ed9d715ba8f] --- .../hipamd/tests/src/hipPeerToPeer_simple.cpp | 41 ++++++++++++++----- 1 file changed, 30 insertions(+), 11 deletions(-) diff --git a/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp b/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp index ca2501438d..aa305f92fa 100644 --- a/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp +++ b/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp @@ -26,11 +26,26 @@ THE SOFTWARE. #include "hip_runtime.h" #include "test_common.h" +bool p_memcpyWithPeer = false; + +void parseMyArguments(int argc, char *argv[]) +{ + int more_argc = HipTest::parseStandardArguments(argc, argv, false); + // parse args for this test: + for (int i = 1; i < more_argc; i++) { + const char *arg = argv[i]; + + if (!strcmp(arg, "--memcpyWithPeer")) { + p_memcpyWithPeer = true; + } else { + failed("Bad argument '%s'", arg); + } + }; +}; int main(int argc, char *argv[]) { - - HipTest::parseStandardArguments(argc, argv, true); + parseMyArguments(argc, argv); int deviceCnt; @@ -58,27 +73,31 @@ int main(int argc, char *argv[]) char *A_h; A_h = (char*)malloc(Nbytes); - HIPCHECK (hipSetDevice(peerDevice)); - HIPCHECK (hipMalloc(&A_d1, Nbytes) ); + // allocate and initialize memory on device0 HIPCHECK (hipSetDevice(p_gpuDevice)); HIPCHECK (hipMalloc(&A_d0, Nbytes) ); - - - // Set memory on first device. - HIPCHECK (hipSetDevice(p_gpuDevice)); 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 ( hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault)); + HIPCHECK (hipSetDevice(p_memcpyWithPeer ? peerDevice : p_gpuDevice)); + HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault)); // Copy data back to host: - HIPCHECK ( hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost)); + HIPCHECK (hipSetDevice(peerDevice)); + HIPCHECK (hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost)); // Check host data: for (int i=0; i Date: Sat, 9 Apr 2016 05:03:08 -0500 Subject: [PATCH 13/15] Improve P2P test. - Ensure proper device is set before each command. - Add command line switches: +bool p_memcpyWithPeer = false; // use the peer device for the P2P copy +bool p_mirrorPeers = false; // in addition to mapping current to peer space, map peer to current space. +int p_peerDevice = -1; // explicly specify which peer to use, else use p_gpuDevice + 1. [ROCm/clr commit: e18284c7cf3585d197b48dc9f74f0cfd6b6c209e] --- .../hipamd/tests/src/hipPeerToPeer_simple.cpp | 37 ++++++++++++++----- projects/clr/hipamd/tests/src/test_common.cpp | 2 +- 2 files changed, 29 insertions(+), 10 deletions(-) diff --git a/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp b/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp index aa305f92fa..4aaa6a452b 100644 --- a/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp +++ b/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp @@ -26,7 +26,9 @@ THE SOFTWARE. #include "hip_runtime.h" #include "test_common.h" -bool p_memcpyWithPeer = false; +bool p_memcpyWithPeer = false; // use the peer device for the P2P copy +bool p_mirrorPeers = false; // in addition to mapping current to peer space, map peer to current space. +int p_peerDevice = -1; // explicly specify which peer to use, else use p_gpuDevice + 1. void parseMyArguments(int argc, char *argv[]) { @@ -37,12 +39,19 @@ void parseMyArguments(int argc, char *argv[]) if (!strcmp(arg, "--memcpyWithPeer")) { p_memcpyWithPeer = true; + } else if (!strcmp(arg, "--mirrorPeers")) { + p_mirrorPeers = true; + } else if (!strcmp(arg, "--peerDevice")) { + if (++i >= argc || !HipTest::parseInt(argv[i], &p_peerDevice)) { + failed("Bad peerDevice argument"); + } } else { failed("Bad argument '%s'", arg); } }; }; + int main(int argc, char *argv[]) { parseMyArguments(argc, argv); @@ -50,23 +59,33 @@ int main(int argc, char *argv[]) int deviceCnt; HIPCHECK(hipGetDeviceCount(&deviceCnt)); - HIPCHECK(hipSetDevice(p_gpuDevice)); - int peerDevice = ((p_gpuDevice + 1) % deviceCnt); + int currentDevice = p_gpuDevice; + int peerDevice = (p_peerDevice == -1) ? ((currentDevice + 1) % deviceCnt) : p_peerDevice; - printf ("N=%zu device=%d peerDevice=%d (%d devices total)\n", N, p_gpuDevice, peerDevice, deviceCnt); + printf ("N=%zu device=%d peerDevice=%d (%d devices total)\n", N, currentDevice, peerDevice, deviceCnt); // Must be on a multi-gpu system: - assert (p_gpuDevice != peerDevice); + assert (currentDevice != peerDevice); int canAccessPeer; - HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, p_gpuDevice, peerDevice)); - printf ("dev#%d canAccessPeer:#%d=%d\n", p_gpuDevice, peerDevice, canAccessPeer); + HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, currentDevice, peerDevice)); + printf ("dev#%d canAccessPeer:#%d=%d\n", currentDevice, peerDevice, canAccessPeer); assert(canAccessPeer); + 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; @@ -75,7 +94,7 @@ int main(int argc, char *argv[]) A_h = (char*)malloc(Nbytes); // allocate and initialize memory on device0 - HIPCHECK (hipSetDevice(p_gpuDevice)); + HIPCHECK (hipSetDevice(currentDevice)); HIPCHECK (hipMalloc(&A_d0, Nbytes) ); HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) ); @@ -87,7 +106,7 @@ int main(int argc, char *argv[]) // Device0 push to device1, using P2P: - HIPCHECK (hipSetDevice(p_memcpyWithPeer ? peerDevice : p_gpuDevice)); + HIPCHECK (hipSetDevice(p_memcpyWithPeer ? peerDevice : currentDevice)); HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault)); // Copy data back to host: diff --git a/projects/clr/hipamd/tests/src/test_common.cpp b/projects/clr/hipamd/tests/src/test_common.cpp index 332c2856d3..35e3d6d3f2 100644 --- a/projects/clr/hipamd/tests/src/test_common.cpp +++ b/projects/clr/hipamd/tests/src/test_common.cpp @@ -111,7 +111,7 @@ int parseStandardArguments(int argc, char *argv[], bool failOnUndefinedArg) failed("Bad iterations argument"); } - } else if (!strcmp(arg, "--gpu") || (!strcmp(arg, "-g"))) { + } else if (!strcmp(arg, "--gpu") || (!strcmp(arg, "-gpuDevice")) || (!strcmp(arg, "-g"))) { if (++i >= argc || !HipTest::parseInt(argv[i], &p_gpuDevice)) { failed("Bad gpuDevice argument"); } From 9bf70214f357ad1dd1e37aa9da8e29fa5c7cd5bf Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 11 Apr 2016 12:52:18 -0500 Subject: [PATCH 14/15] P2p checkpoint. - set USE_PEER_TO_PEER=3 (requires HCC "am_memtracker_update_peers") - when enabling peer, turn it on for previously allocated memory. - hipDeviceCanAccessPeer is no longer self-ware (self does not qualify as a peer) - device peerlist always includes self, so when we call allow_access we never remove self access. - hipDeviceReset() removes old peer mappings. [ROCm/clr commit: 5af4c901c63e53fa8051cd38735ba8b6c6e2c4d8] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 3 +- .../include/hcc_detail/hip_runtime_api.h | 37 +++-- .../samples/1_Utils/hipInfo/hipInfo.cpp | 2 +- projects/clr/hipamd/src/hip_device.cpp | 1 + projects/clr/hipamd/src/hip_hcc.cpp | 26 ++- projects/clr/hipamd/src/hip_memory.cpp | 5 +- projects/clr/hipamd/src/hip_peer.cpp | 21 +-- projects/clr/hipamd/tests/src/CMakeLists.txt | 12 +- .../hipamd/tests/src/hipPeerToPeer_simple.cpp | 149 +++++++++++++++--- 9 files changed, 202 insertions(+), 54 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 1dd8777b1c..0d24862764 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/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/clr/hipamd/include/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h index 2f70ae47c3..58df3a2068 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/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/clr/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp b/projects/clr/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp index 9151d5058e..581194f624 100644 --- a/projects/clr/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/projects/clr/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp @@ -127,7 +127,7 @@ void printDeviceProp (int deviceId) for (int i=0; ilocked_waitAllStreams(); diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index ddf6bb1691..a9194c266a 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/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/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 9e8c963433..a3409a3e2e 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/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/clr/hipamd/src/hip_peer.cpp b/projects/clr/hipamd/src/hip_peer.cpp index aeb9c7fd49..612cd6e309 100644 --- a/projects/clr/hipamd/src/hip_peer.cpp +++ b/projects/clr/hipamd/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/clr/hipamd/tests/src/CMakeLists.txt b/projects/clr/hipamd/tests/src/CMakeLists.txt index 28ea1bf10b..a75a48076c 100644 --- a/projects/clr/hipamd/tests/src/CMakeLists.txt +++ b/projects/clr/hipamd/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/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp b/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp index 4aaa6a452b..5bfb583f3f 100644 --- a/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp +++ b/projects/clr/hipamd/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 Date: Mon, 11 Apr 2016 13:46:53 -0500 Subject: [PATCH 15/15] add simple negative P2P tests [ROCm/clr commit: 2f0c589ffe2c0dc4add18c7cecb0ac5c02aef125] --- projects/clr/hipamd/tests/src/CMakeLists.txt | 2 +- .../hipamd/tests/src/hipMultiThreadDevice.cpp | 4 +- .../hipamd/tests/src/hipPeerToPeer_simple.cpp | 162 ++++++++++-------- 3 files changed, 97 insertions(+), 71 deletions(-) diff --git a/projects/clr/hipamd/tests/src/CMakeLists.txt b/projects/clr/hipamd/tests/src/CMakeLists.txt index a75a48076c..c13f90a8fe 100644 --- a/projects/clr/hipamd/tests/src/CMakeLists.txt +++ b/projects/clr/hipamd/tests/src/CMakeLists.txt @@ -224,7 +224,7 @@ make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-pyramid" --tests 0x4 make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-nearzero" --tests 0x10) if (${HIP_MULTI_GPU}) - make_test(hipPeerToPeer_simple ) # use current device for copy, this fails. + 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. diff --git a/projects/clr/hipamd/tests/src/hipMultiThreadDevice.cpp b/projects/clr/hipamd/tests/src/hipMultiThreadDevice.cpp index a1f64aceb3..d9afda59d0 100644 --- a/projects/clr/hipamd/tests/src/hipMultiThreadDevice.cpp +++ b/projects/clr/hipamd/tests/src/hipMultiThreadDevice.cpp @@ -116,12 +116,12 @@ int main(int argc, char *argv[]) /*disable, this takess a while and if the next one works then no need to run serial*/ if (1 && (p_tests & 0x2)) { printf ("\ntest 0x2 : serialized multiThread_pyramid(1) \n"); - multiThread_pyramid(true, 10); + multiThread_pyramid(true, 3); } if (p_tests & 0x4) { printf ("\ntest 0x4 : parallel multiThread_pyramid(1) \n"); - multiThread_pyramid(false, 10); + multiThread_pyramid(false, 3); } //if (p_tests & 0x8) { diff --git a/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp b/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp index 5bfb583f3f..2c0dd95b36 100644 --- a/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp +++ b/projects/clr/hipamd/tests/src/hipPeerToPeer_simple.cpp @@ -29,6 +29,10 @@ bool p_memcpyWithPeer = false; // use the peer device for the P2P copy bool p_mirrorPeers = false; // in addition to mapping current to peer space, map peer to current space. int p_peerDevice = -1; // explicly specify which peer to use, else use p_gpuDevice + 1. + +int g_currentDevice; +int g_peerDevice; + void parseMyArguments(int argc, char *argv[]) { int more_argc = HipTest::parseStandardArguments(argc, argv, false); @@ -51,45 +55,50 @@ void parseMyArguments(int argc, char *argv[]) }; +// Sets globals g_currentDevice, g_peerDevice +void setupPeerTests() +{ + int deviceCnt; + + HIPCHECK(hipGetDeviceCount(&deviceCnt)); + + g_currentDevice = p_gpuDevice; + g_peerDevice = (p_peerDevice == -1) ? ((g_currentDevice + 1) % deviceCnt) : p_peerDevice; + + printf ("N=%zu device=%d peerDevice=%d (%d devices total)\n", N, g_currentDevice, g_peerDevice, deviceCnt); + + // Must be on a multi-gpu system: + assert (g_currentDevice != g_peerDevice); + + int canAccessPeer; + HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, g_currentDevice, g_peerDevice)); + printf ("dev#%d canAccessPeer:#%d=%d\n", g_currentDevice, g_peerDevice, canAccessPeer); + + assert(canAccessPeer); + + HIPCHECK (hipSetDevice(g_currentDevice)); + HIPCHECK(hipDeviceReset()); + HIPCHECK (hipSetDevice(g_peerDevice)); + HIPCHECK(hipDeviceReset()); +} + //--- // Test which enables peer2peer first, then allocates the memory. void enablePeerFirst() { printf ("\n==testing: %s\n", __func__); - int deviceCnt; - HIPCHECK(hipGetDeviceCount(&deviceCnt)); - - int currentDevice = p_gpuDevice; - int peerDevice = (p_peerDevice == -1) ? ((currentDevice + 1) % deviceCnt) : p_peerDevice; - - printf ("N=%zu device=%d peerDevice=%d (%d devices total)\n", N, currentDevice, peerDevice, deviceCnt); - - // Must be on a multi-gpu system: - assert (currentDevice != peerDevice); - - int canAccessPeer; - HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, currentDevice, peerDevice)); - printf ("dev#%d canAccessPeer:#%d=%d\n", currentDevice, peerDevice, canAccessPeer); - - assert(canAccessPeer); - - HIPCHECK (hipSetDevice(currentDevice)); - HIPCHECK(hipDeviceReset()); - HIPCHECK (hipSetDevice(peerDevice)); - HIPCHECK(hipDeviceReset()); - - HIPCHECK(hipSetDevice(currentDevice)); - HIPCHECK(hipDeviceEnablePeerAccess(peerDevice, 0)); + HIPCHECK(hipSetDevice(g_currentDevice)); + HIPCHECK(hipDeviceEnablePeerAccess(g_peerDevice, 0)); if (p_mirrorPeers) { int canAccessPeer; - HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, peerDevice, currentDevice)); + HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, g_peerDevice, g_currentDevice)); assert(canAccessPeer); - HIPCHECK(hipSetDevice(peerDevice)); - HIPCHECK(hipDeviceEnablePeerAccess(currentDevice, 0)); + HIPCHECK(hipSetDevice(g_peerDevice)); + HIPCHECK(hipDeviceEnablePeerAccess(g_currentDevice, 0)); } size_t Nbytes = N*sizeof(char); @@ -100,23 +109,23 @@ void enablePeerFirst() A_h = (char*)malloc(Nbytes); // allocate and initialize memory on device0 - HIPCHECK (hipSetDevice(currentDevice)); + HIPCHECK (hipSetDevice(g_currentDevice)); HIPCHECK (hipMalloc(&A_d0, Nbytes) ); HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) ); // allocate and initialize memory on peer device - HIPCHECK (hipSetDevice(peerDevice)); + HIPCHECK (hipSetDevice(g_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 (hipSetDevice(p_memcpyWithPeer ? g_peerDevice : g_currentDevice)); HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault)); // Copy data back to host: - HIPCHECK (hipSetDevice(peerDevice)); + HIPCHECK (hipSetDevice(g_peerDevice)); HIPCHECK (hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost)); // Check host data: @@ -128,35 +137,14 @@ void enablePeerFirst() } - //--- - // Test which allocated memory first, then enables peer2peer. - // Enabling peer needs to scan all allocated memory and enable peer access. - void allocMemoryFirst() - { - printf ("\n==testing: %s\n", __func__); - int deviceCnt; - - HIPCHECK(hipGetDeviceCount(&deviceCnt)); - - int currentDevice = p_gpuDevice; - int peerDevice = (p_peerDevice == -1) ? ((currentDevice + 1) % deviceCnt) : p_peerDevice; - - printf ("N=%zu device=%d peerDevice=%d (%d devices total)\n", N, currentDevice, peerDevice, deviceCnt); - - // Must be on a multi-gpu system: - assert (currentDevice != peerDevice); - - int canAccessPeer; - HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, currentDevice, peerDevice)); - printf ("dev#%d canAccessPeer:#%d=%d\n", currentDevice, peerDevice, canAccessPeer); - - assert(canAccessPeer); - - HIPCHECK (hipSetDevice(currentDevice)); - HIPCHECK(hipDeviceReset()); - HIPCHECK (hipSetDevice(peerDevice)); - HIPCHECK(hipDeviceReset()); +//--- +// Test which allocated memory first, then enables peer2peer. +// Enabling peer needs to scan all allocated memory and enable peer access. +void allocMemoryFirst() +{ + printf ("\n==testing: %s\n", __func__); + setupPeerTests(); size_t Nbytes = N*sizeof(char); @@ -167,39 +155,39 @@ void enablePeerFirst() //--- // allocate and initialize memory on device0 - HIPCHECK (hipSetDevice(currentDevice)); + HIPCHECK (hipSetDevice(g_currentDevice)); HIPCHECK (hipMalloc(&A_d0, Nbytes) ); HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) ); // allocate and initialize memory on peer device - HIPCHECK (hipSetDevice(peerDevice)); + HIPCHECK (hipSetDevice(g_peerDevice)); HIPCHECK (hipMalloc(&A_d1, Nbytes) ); HIPCHECK ( hipMemset(A_d1, 0x13, Nbytes) ); //--- //Enable peer access, for memory already allocated: - HIPCHECK(hipSetDevice(currentDevice)); - HIPCHECK(hipDeviceEnablePeerAccess(peerDevice, 0)); + HIPCHECK(hipSetDevice(g_currentDevice)); + HIPCHECK(hipDeviceEnablePeerAccess(g_peerDevice, 0)); if (p_mirrorPeers) { int canAccessPeer; - HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, peerDevice, currentDevice)); + HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, g_peerDevice, g_currentDevice)); assert(canAccessPeer); - HIPCHECK(hipSetDevice(peerDevice)); - HIPCHECK(hipDeviceEnablePeerAccess(currentDevice, 0)); + HIPCHECK(hipSetDevice(g_peerDevice)); + HIPCHECK(hipDeviceEnablePeerAccess(g_currentDevice, 0)); } //--- // Copies to test functionality: // Device0 push to device1, using P2P: - HIPCHECK (hipSetDevice(p_memcpyWithPeer ? peerDevice : currentDevice)); + HIPCHECK (hipSetDevice(p_memcpyWithPeer ? g_peerDevice : g_currentDevice)); HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault)); // Copy data back to host: - HIPCHECK (hipSetDevice(peerDevice)); + HIPCHECK (hipSetDevice(g_peerDevice)); HIPCHECK (hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost)); @@ -212,6 +200,40 @@ void enablePeerFirst() } } +void simpleNegative() +{ + printf ("\n==testing: %s\n", __func__); + + setupPeerTests(); + + int deviceId; + HIPCHECK (hipGetDevice(&deviceId)); + + //--- + //-- self is not a peer + int canAccessPeer; + hipError_t e = hipDeviceCanAccessPeer(&canAccessPeer, deviceId, deviceId); + HIPASSERT( e == hipSuccess); // no error returned, it doesn't hurt to ask. + HIPASSERT (canAccessPeer == 0); // but self is not a peer. + + e = hipSuccess; + //--- + // Enable same device twice in a row: + HIPCHECK(hipSetDevice(g_currentDevice)); + HIPCHECK(hipDeviceEnablePeerAccess(g_peerDevice, 0)); + e =(hipDeviceEnablePeerAccess(g_peerDevice, 0)); + HIPASSERT (e == hipErrorPeerAccessAlreadyEnabled); + + //--- + // try disabling twice in a row + HIPCHECK(hipDeviceDisablePeerAccess(g_peerDevice)); + e =(hipDeviceDisablePeerAccess(g_peerDevice)); + HIPASSERT (e == hipErrorPeerAccessNotEnabled); + + + // More tests here: +} + int main(int argc, char *argv[]) @@ -226,5 +248,9 @@ int main(int argc, char *argv[]) allocMemoryFirst(); } + if (p_tests & 0x4) { + simpleNegative(); + } + passed(); }