From 91368967bb7b76175ce0c13f8e8723113ccbe9b3 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 27 Jul 2016 18:31:11 -0500 Subject: [PATCH] Cleanup sync code. Remove dead depFutures, enqueueBarrier call. Rename some parms to reflect usage. Add comments to better explain tricky parts of sync code. Change-Id: I763296421d9c2b3b58fc8cef5f010b12ab49553c [ROCm/hip commit: 02dd7a739923f0a55a81ed89d9459586e8b9461a] --- projects/hip/include/hcc_detail/hip_hcc.h | 15 +++++----- projects/hip/src/hip_hcc.cpp | 34 ++++++++++------------- 2 files changed, 22 insertions(+), 27 deletions(-) diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h index db51a497e9..405e2de67c 100644 --- a/projects/hip/include/hcc_detail/hip_hcc.h +++ b/projects/hip/include/hcc_detail/hip_hcc.h @@ -364,7 +364,6 @@ public: ihipStreamCriticalBase_t * mlock() { LockedBase::lock(); return this;}; - public: // Critical Data: ihipCommand_t _last_command_type; // type of the last command @@ -380,8 +379,10 @@ public: int _signalCursor; SIGSEQNUM _oldest_live_sig_id; // oldest live seq_id, anything < this can be allocated. std::deque _signalPool; // Pool of signals for use by this stream. - uint32_t _signalCnt; - uint32_t _kernelCnt; + uint32_t _signalCnt; // Count of inflight commands using signals from the signal pool. + // Each copy may use 1-2 signals depending on command transitions: + // 2 are required if a barrier packet is inserted. + uint32_t _kernelCnt; // Count of inflight kernels in this stream. Reset at ::wait(). SIGSEQNUM _stream_sig_id; // Monotonically increasing unique signal id. }; @@ -421,7 +422,7 @@ typedef uint64_t SeqNum_t ; // Non-threadsafe accessors - must be protected by high-level stream lock with accessor passed to function. - SIGSEQNUM lastCopySeqId (LockedAccessor_StreamCrit_t &crit) { return crit->_last_copy_signal ? crit->_last_copy_signal->_sig_id : 0; }; + SIGSEQNUM lastCopySeqId (LockedAccessor_StreamCrit_t &crit) const { return crit->_last_copy_signal ? crit->_last_copy_signal->_sig_id : 0; }; ihipSignal_t * allocSignal (LockedAccessor_StreamCrit_t &crit); @@ -441,14 +442,12 @@ private: // Critical Data. THis MUST be accessed through LockedAccessor_StreamCrit_t ihipStreamCritical_t _criticalData; - // Array of dependency completion_future. - std::vector _depFutures; - private: void enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal, ihipSignal_t *completionSignal); - void enqueueBarrier(hsa_queue_t* queue, hsa_signal_t *depSignal); void waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal); + + // The unsigned return is hipMemcpyKind unsigned resolveMemcpyDirection(bool srcTracked, bool dstTracked, bool srcInDeviceMem, bool dstInDeviceMem); void setAsyncCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent); diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 54a16a56d8..48d1eebf50 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -146,6 +146,9 @@ ihipStream_t::~ihipStream_t() //--- +//TODO - this function is dangerous since it does not propertly account +//for younger commands which may be depending on the signals we are reclaiming. +//Will fix when we move to HCC management of copy signals. void ihipStream_t::locked_reclaimSignals(SIGSEQNUM sigNum) { LockedAccessor_StreamCrit_t crit(_criticalData); @@ -180,18 +183,21 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty if (! assertQueueEmpty) { tprintf (DB_SYNC, "stream %p wait for queue-empty..\n", this); _av.wait(); - crit->_kernelCnt = 0; + } else { + assert (crit->_kernelCnt == 0); } + if (crit->_last_copy_signal) { tprintf (DB_SYNC, "stream %p wait for lastCopy:#%lu...\n", this, lastCopySeqId(crit) ); this->waitCopy(crit, crit->_last_copy_signal); } + crit->_kernelCnt = 0; + // Reset the stream to "empty" - next command will not set up an inpute dependency on any older signal. crit->_last_command_type = ihipCommandCopyH2D; crit->_last_copy_signal = NULL; - - _depFutures.clear(); + crit->_signalCnt = 0; } @@ -306,7 +312,6 @@ ihipSignal_t *ihipStream_t::allocSignal(LockedAccessor_StreamCrit_t &crit) crit->_signalCnt++; if(crit->_signalCnt == HIP_NUM_SIGNALS_PER_STREAM){ - crit->_signalCnt = 0; this->wait(crit); } @@ -470,14 +475,12 @@ void ihipStream_t::lockclose_postKernelCommand(hc::completion_future &kernelFutu //--- // Called whenever a copy command is set to the stream. // Examines the last command sent to this stream and returns a signal to wait on, if required. -int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType) +int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *copyCompletionSignal, hsa_signal_t *waitSignal, ihipCommand_t copyType) { int needSync = 0; waitSignal->handle = 0; - //_mutex.lock(); // will be unlocked in postCopyCommand - // If switching command types, we need to add a barrier packet to synchronize things. if (FORCE_SAMEDIR_COPY_DEP || (crit->_last_command_type != copyType)) { @@ -490,15 +493,6 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t hsa_signal_store_relaxed(depSignal->_hsa_signal,1); this->enqueueBarrier(static_cast(_av.get_hsa_queue()), NULL, depSignal); *waitSignal = depSignal->_hsa_signal; -// hsa_signal_t *hsaSignal = (static_cast (crit->_last_kernel_future.get_native_handle())); -/* if (hsaSignal) { - // Keep reference to the kernel future in order to keep the - // dependent signal alive. - _depFutures.push_back(crit->_last_kernel_future); - *waitSignal = * hsaSignal; - } else { - assert(0); // if NULL signal, and we return 1, hsa_amd_memory_copy_async will fail. Confirm this never happens. - }*/ } else if (crit->_last_copy_signal) { needSync = 1; tprintf (DB_SYNC, "stream %p switch %s to %s (async copy dep on other copy #%lu)\n", @@ -521,7 +515,7 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t crit->_last_command_type = copyType; } - crit->_last_copy_signal = lastCopy; + crit->_last_copy_signal = copyCompletionSignal; return needSync; } @@ -1611,7 +1605,8 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const hc::am_copy(dst, src, sizeBytes); #endif } - }else{ + } else { + // This is H2D copy, and source is pinned host memory : we can copy directly w/o using staging buffer. hsa_agent_t dstAgent = *(static_cast(dstPtrInfo._acc.get_hsa_agent())); hsa_agent_t srcAgent = *(static_cast(srcPtrInfo._acc.get_hsa_agent())); @@ -1661,7 +1656,8 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const hc::am_copy(dst, src, sizeBytes); #endif } - }else{ + } else { + // This is D2H copy, and destination is pinned host memory : we can copy directly w/o using staging buffer. hsa_agent_t dstAgent = *(static_cast(dstPtrInfo._acc.get_hsa_agent())); hsa_agent_t srcAgent = *(static_cast(srcPtrInfo._acc.get_hsa_agent()));