diff --git a/hipamd/include/hcc_detail/hip_hcc.h b/hipamd/include/hcc_detail/hip_hcc.h index 4b40a0ac5e..99cc46c380 100644 --- a/hipamd/include/hcc_detail/hip_hcc.h +++ b/hipamd/include/hcc_detail/hip_hcc.h @@ -531,6 +531,9 @@ public: bool removePeer(ihipDevice_t *peer); void resetPeers(ihipDevice_t *thisDevice); + + void addStream(ihipStream_t *stream); + uint32_t peerCnt() const { return _peerCnt; }; hsa_agent_t *peerAgents() const { return _peerAgents; }; diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index 90f5afa929..44c9bb2807 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -249,6 +249,14 @@ void ihipDeviceCriticalBase_t::resetPeers(ihipDevice_t *thisDevice) addPeer(thisDevice); // peer-list always contains self agent. } + +template<> +void ihipDeviceCriticalBase_t::addStream(ihipStream_t *stream) +{ + _streams.push_back(stream); + stream->_id = incStreamId(); +} + //------------------------------------------------------------------------------------------------- //--- @@ -450,9 +458,29 @@ void ihipDevice_t::locked_reset() // Obtain mutex access to the device critical data, release by destructor LockedAccessor_DeviceCrit_t crit(_criticalData); + + //--- + //Wait for pending activity to complete? TODO - check if this is required behavior: + tprintf(DB_SYNC, "locked_reset waiting for activity to complete.\n"); + // Reset and remove streams: + // Delete all created streams including the default one. + for (auto streamI=crit->const_streams().begin(); streamI!=crit->const_streams().end(); streamI++) { + ihipStream_t *stream = *streamI; + (*streamI)->locked_wait(); + tprintf(DB_SYNC, " delete stream=%p\n", stream); + + delete stream; + } + // Clear the list. crit->streams().clear(); + + // Create a fresh default stream and add it: + _default_stream = new ihipStream_t(_device_index, _acc.get_default_view(), hipStreamDefault); + crit->addStream(_default_stream); + + // This resest peer list to just me: crit->resetPeers(this); @@ -488,9 +516,6 @@ void ihipDevice_t::init(unsigned device_index, unsigned deviceCnt, hc::accelerat locked_reset(); - _default_stream = new ihipStream_t(device_index, acc.get_default_view(), hipStreamDefault); - locked_addStream(_default_stream); - tprintf(DB_SYNC, "created device with default_stream=%p\n", _default_stream); @@ -772,8 +797,7 @@ void ihipDevice_t::locked_addStream(ihipStream_t *s) { LockedAccessor_DeviceCrit_t crit(_criticalData); - crit->streams().push_back(s); - s->_id = crit->incStreamId(); + crit->addStream(s); } //---