From db91890f53276dc0067a3581372c4598e86471dc Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 6 Apr 2016 11:28:23 -0500 Subject: [PATCH] Checkpoint initial peer2peer implementation. --- hipamd/include/hcc_detail/hip_hcc.h | 27 +++++++++-- hipamd/include/hcc_detail/hip_runtime_api.h | 10 ++-- hipamd/src/hip_hcc.cpp | 40 +++++++++++++++- hipamd/src/hip_peer.cpp | 51 ++++++++++++++++++--- 4 files changed, 111 insertions(+), 17 deletions(-) diff --git a/hipamd/include/hcc_detail/hip_hcc.h b/hipamd/include/hcc_detail/hip_hcc.h index 42dd3b6df5..5c8aaf5659 100644 --- a/hipamd/include/hcc_detail/hip_hcc.h +++ b/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/hipamd/include/hcc_detail/hip_runtime_api.h b/hipamd/include/hcc_detail/hip_runtime_api.h index b5fdb312a4..56b3560427 100644 --- a/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/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/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index 6c80b84e0e..bbf6538817 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/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/hipamd/src/hip_peer.cpp b/hipamd/src/hip_peer.cpp index d4de6053de..c4a8a1bcf3 100644 --- a/hipamd/src/hip_peer.cpp +++ b/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); }