AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Cleanup sync code.

Browse Source 
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: 02dd7a7399]
This commit is contained in:
Ben Sander
2016-07-27 18:31:11 -05:00
parent 07c9fb8b38
commit 91368967bb
2 changed files with 22 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/include/hcc_detail/hip_hcc.h
+7 -8
View File
@@ -364,7 +364,6 @@ public:
ihipStreamCriticalBase_t<StreamMutex> * mlock() { LockedBase<MUTEX_TYPE>::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<ihipSignal_t> _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<hc::completion_future> _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);
projects/hip/src/hip_hcc.cpp
+15 -19
View File
@@ -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<hsa_queue_t*>(_av.get_hsa_queue()), NULL, depSignal);
*waitSignal = depSignal->_hsa_signal;
// hsa_signal_t *hsaSignal = (static_cast<hsa_signal_t*> (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<hsa_agent_t*>(dstPtrInfo._acc.get_hsa_agent()));
hsa_agent_t srcAgent = *(static_cast<hsa_agent_t*>(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<hsa_agent_t*>(dstPtrInfo._acc.get_hsa_agent()));
hsa_agent_t srcAgent = *(static_cast<hsa_agent_t*>(srcPtrInfo._acc.get_hsa_agent()));