AI/rocm-systems
2
0
フォーク 0
コード イシュー プルリクエスト Actions パッケージ プロジェクト リリース Wiki アクティビティ

Merge branch 'privatestaging' of https://github.com/AMDComputeLibraries/HIP-privatestaging into privatestaging

ソースを参照
このコミットが含まれているのは:
Aditya Atluri
2016-03-29 11:18:09 -05:00
親 61cd2bb399 a37a019b01
コミット 39af5d9caf
17個のファイルの変更、678行の追加、245行の削除
Patchファイルをダウンロード Diffファイルをダウンロード すべてのファイルを展開 すべてのファイルを折り畳む
src/hip_device.cpp
+5 -9
ファイルの表示
@@ -150,8 +150,7 @@ hipError_t hipDeviceSynchronize(void)
{
HIP_INIT_API();
ihipGetTlsDefaultDevice()->waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
ihipGetTlsDefaultDevice()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
return ihipLogStatus(hipSuccess);
}
@@ -174,15 +173,12 @@ hipError_t hipDeviceReset(void)
if (device) {
//---
//Wait for pending activity to complete?
//TODO - check if this is required behavior:
for (auto streamI=device->_streams.begin(); streamI!=device->_streams.end(); streamI++) {
ihipStream_t *stream = *streamI;
stream->wait();
}
//Wait for pending activity to complete? TODO - check if this is required behavior:
device->locked_waitAllStreams();
// Release device resources (streams and memory):
device->reset();
device->locked_reset();
}
return ihipLogStatus(hipSuccess);
src/hip_event.cpp
+27 -12
ファイルの表示
@@ -25,16 +25,13 @@ THE SOFTWARE.
//-------------------------------------------------------------------------------------------------
// Events
//---
/**
* @warning : flags must be 0.
*/
hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags)
{
// TODO - support hipEventDefault, hipEventBlockingSync, hipEventDisableTiming
std::call_once(hip_initialized, ihipInit);
hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags)
{
hipError_t e = hipSuccess;
// TODO - support hipEventDefault, hipEventBlockingSync, hipEventDisableTiming
if (flags == 0) {
ihipEvent_t *eh = event->_handle = new ihipEvent_t();
@@ -47,8 +44,25 @@ hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags)
e = hipErrorInvalidValue;
}
return e;
}
return ihipLogStatus(e);
/**
* @warning : flags must be 0.
*/
hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags)
{
HIP_INIT_API(event, flags);
return ihipLogStatus(ihipEventCreate(event, flags));
}
hipError_t hipEventCreate(hipEvent_t* event)
{
HIP_INIT_API(event);
return ihipLogStatus(ihipEventCreate(event, 0));
}
@@ -67,7 +81,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
// TODO-HCC fix this - is CUDA this conservative or still uses device timestamps?
// TODO-HCC can we use barrier or event marker to implement better solution?
ihipDevice_t *device = ihipGetTlsDefaultDevice();
device->syncDefaultStream(true);
device->locked_syncDefaultStream(true);
eh->_timestamp = hc::get_system_ticks();
eh->_state = hipEventStatusRecorded;
@@ -77,7 +91,8 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
// Clear timestamps
eh->_timestamp = 0;
eh->_marker = stream->_av.create_marker();
eh->_copy_seq_id = stream->lastCopySeqId();
eh->_copy_seq_id = stream->locked_lastCopySeqId();
return ihipLogStatus(hipSuccess);
}
@@ -117,7 +132,7 @@ hipError_t hipEventSynchronize(hipEvent_t event)
return ihipLogStatus(hipSuccess);
} else if (eh->_stream == NULL) {
ihipDevice_t *device = ihipGetTlsDefaultDevice();
device->syncDefaultStream(true);
device->locked_syncDefaultStream(true);
return ihipLogStatus(hipSuccess);
} else {
#if __hcc_workweek__ >= 16033
@@ -125,7 +140,7 @@ hipError_t hipEventSynchronize(hipEvent_t event)
#else
eh->_marker.wait();
#endif
eh->_stream->reclaimSignals_ts(eh->_copy_seq_id);
eh->_stream->locked_reclaimSignals(eh->_copy_seq_id);
return ihipLogStatus(hipSuccess);
}
src/hip_hcc.cpp
+155 -103
ファイルの表示
@@ -59,6 +59,7 @@ int HIP_LAUNCH_BLOCKING = 0;
int HIP_PRINT_ENV = 0;
int HIP_TRACE_API= 0;
int HIP_ATP_MARKER= 0;
int HIP_DB= 0;
int HIP_STAGING_SIZE = 64; /* size of staging buffers, in KB */
int HIP_STAGING_BUFFERS = 2; // TODO - remove, two buffers should be enough.
@@ -124,42 +125,36 @@ ihipSignal_t::~ihipSignal_t()
// ihipStream_t:
//=================================================================================================
//---
ihipStream_t::ihipStream_t(unsigned device_index, hc::accelerator_view av, SeqNum_t id, unsigned int flags) :
_id(id),
ihipStream_t::ihipStream_t(unsigned device_index, hc::accelerator_view av, unsigned int flags) :
_id(0), // will be set by add function.
_av(av),
_flags(flags),
_device_index(device_index),
_last_command_type(ihipCommandCopyH2D),
_last_copy_signal(NULL),
_signalCursor(0),
_stream_sig_id(0),
_oldest_live_sig_id(1)
_device_index(device_index)
{
tprintf(DB_SYNC, " streamCreate: stream=%p\n", this);
_signalPool.resize(HIP_STREAM_SIGNALS > 0 ? HIP_STREAM_SIGNALS : 1);
};
//---
ihipStream_t::~ihipStream_t()
{
_signalPool.clear();
{
}
//---
void ihipStream_t::reclaimSignals_ts(SIGSEQNUM sigNum)
void ihipStream_t::locked_reclaimSignals(SIGSEQNUM sigNum)
{
LockedAccessor_StreamCrit_t crit(_criticalData);
tprintf(DB_SIGNAL, "reclaim signal #%lu\n", sigNum);
// Mark all signals older and including this one as available for
_oldest_live_sig_id = sigNum+1;
// Mark all signals older and including this one as available for re-allocation.
crit->_oldest_live_sig_id = sigNum+1;
}
//---
void ihipStream_t::waitCopy(ihipSignal_t *signal)
void ihipStream_t::waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal)
{
hsa_signal_wait_acquire(signal->_hsa_signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE);
@@ -167,32 +162,43 @@ void ihipStream_t::waitCopy(ihipSignal_t *signal)
tprintf(DB_SIGNAL, "waitCopy reclaim signal #%lu\n", sigNum);
// Mark all signals older and including this one as available for reclaim
if (sigNum > _oldest_live_sig_id) {
_oldest_live_sig_id = sigNum+1; // TODO, +1 here seems dangerous.
if (sigNum > crit->_oldest_live_sig_id) {
crit->_oldest_live_sig_id = sigNum+1; // TODO, +1 here seems dangerous.
}
}
//---
//Wait for all kernel and data copy commands in this stream to complete.
void ihipStream_t::wait(bool assertQueueEmpty)
//This signature should be used in routines that already have locked the stream mutex
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();
}
if (_last_copy_signal) {
tprintf (DB_SYNC, "stream %p wait for lastCopy:#%lu...\n", this, _last_copy_signal ? _last_copy_signal->_sig_id: 0x0 );
this->waitCopy(_last_copy_signal);
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);
}
// Reset the stream to "empty" - next command will not set up an inpute dependency on any older signal.
_last_command_type = ihipCommandCopyH2D;
_last_copy_signal = NULL;
crit->_last_command_type = ihipCommandCopyH2D;
crit->_last_copy_signal = NULL;
}
//---
//Wait for all kernel and data copy commands in this stream to complete.
void ihipStream_t::locked_wait(bool assertQueueEmpty)
{
LockedAccessor_StreamCrit_t crit(_criticalData);
wait(crit, assertQueueEmpty);
};
//---
ihipDevice_t * ihipStream_t::getDevice() const
{
@@ -208,26 +214,26 @@ ihipDevice_t * ihipStream_t::getDevice() const
// Allocate a new signal from the signal pool.
// Returned signals have value of 0.
// Signals are intended for use in this stream and are always reclaimed "in-order".
ihipSignal_t *ihipStream_t::allocSignal()
ihipSignal_t *ihipStream_t::allocSignal(LockedAccessor_StreamCrit_t &crit)
{
int numToScan = _signalPool.size();
int numToScan = crit->_signalPool.size();
do {
auto thisCursor = _signalCursor;
if (++_signalCursor == _signalPool.size()) {
_signalCursor = 0;
auto thisCursor = crit->_signalCursor;
if (++crit->_signalCursor == crit->_signalPool.size()) {
crit->_signalCursor = 0;
}
if (_signalPool[thisCursor]._sig_id < _oldest_live_sig_id) {
SIGSEQNUM oldSigId = _signalPool[thisCursor]._sig_id;
_signalPool[thisCursor]._index = thisCursor;
_signalPool[thisCursor]._sig_id = ++_stream_sig_id; // allocate it.
if (crit->_signalPool[thisCursor]._sig_id < crit->_oldest_live_sig_id) {
SIGSEQNUM oldSigId = crit->_signalPool[thisCursor]._sig_id;
crit->_signalPool[thisCursor]._index = thisCursor;
crit->_signalPool[thisCursor]._sig_id = ++crit->_stream_sig_id; // allocate it.
tprintf(DB_SIGNAL, "allocatSignal #%lu at pos:%i (old sigId:%lu < oldest_live:%lu)\n",
_signalPool[thisCursor]._sig_id,
thisCursor, oldSigId, _oldest_live_sig_id);
crit->_signalPool[thisCursor]._sig_id,
thisCursor, oldSigId, crit->_oldest_live_sig_id);
return &_signalPool[thisCursor];
return &crit->_signalPool[thisCursor];
}
} while (--numToScan) ;
@@ -235,13 +241,13 @@ ihipSignal_t *ihipStream_t::allocSignal()
assert(numToScan == 0);
// Have to grow the pool:
_signalCursor = _signalPool.size(); // set to the beginning of the new entries:
if (_signalCursor > 10000) {
fprintf (stderr, "warning: signal pool size=%d, may indicate runaway number of inflight commands\n", _signalCursor);
crit->_signalCursor = crit->_signalPool.size(); // set to the beginning of the new entries:
if (crit->_signalCursor > 10000) {
fprintf (stderr, "warning: signal pool size=%d, may indicate runaway number of inflight commands\n", crit->_signalCursor);
}
_signalPool.resize(_signalPool.size() * 2);
tprintf (DB_SIGNAL, "grow signal pool to %zu entries, cursor=%d\n", _signalPool.size(), _signalCursor);
return allocSignal(); // try again,
crit->_signalPool.resize(crit->_signalPool.size() * 2);
tprintf (DB_SIGNAL, "grow signal pool to %zu entries, cursor=%d\n", crit->_signalPool.size(), crit->_signalCursor);
return allocSignal(crit); // try again,
// Should never reach here.
assert(0);
@@ -284,32 +290,32 @@ void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal)
//into the stream to mimic CUDA stream semantics. (some hardware uses separate
//queues for data commands and kernel commands, and no implicit ordering is provided).
//
bool ihipStream_t::preKernelCommand()
bool ihipStream_t::lockopen_preKernelCommand()
{
_mutex.lock(); // will be unlocked in postKernelCommand
LockedAccessor_StreamCrit_t crit(_criticalData, false/*no unlock at destruction*/);
bool addedSync = false;
// If switching command types, we need to add a barrier packet to synchronize things.
if (_last_command_type != ihipCommandKernel) {
if (_last_copy_signal) {
if (crit->_last_command_type != ihipCommandKernel) {
if (crit->_last_copy_signal) {
addedSync = true;
hsa_queue_t * q = (hsa_queue_t*)_av.get_hsa_queue();
if (HIP_DISABLE_HW_KERNEL_DEP == 0) {
this->enqueueBarrier(q, _last_copy_signal);
this->enqueueBarrier(q, crit->_last_copy_signal);
tprintf (DB_SYNC, "stream %p switch %s to %s (barrier pkt inserted with wait on #%lu)\n",
this, ihipCommandName[_last_command_type], ihipCommandName[ihipCommandKernel], _last_copy_signal->_sig_id)
this, ihipCommandName[crit->_last_command_type], ihipCommandName[ihipCommandKernel], crit->_last_copy_signal->_sig_id)
} else if (HIP_DISABLE_HW_KERNEL_DEP>0) {
tprintf (DB_SYNC, "stream %p switch %s to %s (HOST wait for previous...)\n",
this, ihipCommandName[_last_command_type], ihipCommandName[ihipCommandKernel]);
this->waitCopy(_last_copy_signal);
this, ihipCommandName[crit->_last_command_type], ihipCommandName[ihipCommandKernel]);
this->waitCopy(crit, crit->_last_copy_signal);
} else if (HIP_DISABLE_HW_KERNEL_DEP==-1) {
tprintf (DB_SYNC, "stream %p switch %s to %s (IGNORE dependency)\n",
this, ihipCommandName[_last_command_type], ihipCommandName[ihipCommandKernel]);
this, ihipCommandName[crit->_last_command_type], ihipCommandName[ihipCommandKernel]);
}
}
_last_command_type = ihipCommandKernel;
crit->_last_command_type = ihipCommandKernel;
}
return addedSync;
@@ -317,11 +323,13 @@ bool ihipStream_t::preKernelCommand()
//---
void ihipStream_t::postKernelCommand(hc::completion_future &kernelFuture)
// Must be called after kernel finishes, this releases the lock on the stream so other commands can submit.
void ihipStream_t::lockclose_postKernelCommand(hc::completion_future &kernelFuture)
{
_last_kernel_future = kernelFuture;
// We locked _criticalData in the lockopen_preKernelCommand() so OK to access here:
_criticalData._last_kernel_future = kernelFuture;
_mutex.unlock();
_criticalData.unlock(); // paired with lock from lockopen_preKernelCommand.
};
@@ -329,7 +337,7 @@ void ihipStream_t::postKernelCommand(hc::completion_future &kernelFuture)
//---
// 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(ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType)
int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType)
{
int needSync = 0;
@@ -338,22 +346,24 @@ int ihipStream_t::preCopyCommand(ihipSignal_t *lastCopy, hsa_signal_t *waitSigna
//_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 || (_last_command_type != copyType)) {
if (FORCE_SAMEDIR_COPY_DEP || (crit->_last_command_type != copyType)) {
if (_last_command_type == ihipCommandKernel) {
if (crit->_last_command_type == ihipCommandKernel) {
tprintf (DB_SYNC, "stream %p switch %s to %s (async copy dep on prev kernel)\n",
this, ihipCommandName[_last_command_type], ihipCommandName[copyType]);
this, ihipCommandName[crit->_last_command_type], ihipCommandName[copyType]);
needSync = 1;
hsa_signal_t *hsaSignal = (static_cast<hsa_signal_t*> (_last_kernel_future.get_native_handle()));
hsa_signal_t *hsaSignal = (static_cast<hsa_signal_t*> (crit->_last_kernel_future.get_native_handle()));
if (hsaSignal) {
*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 (_last_copy_signal) {
} 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",
this, ihipCommandName[_last_command_type], ihipCommandName[copyType], _last_copy_signal->_sig_id);
*waitSignal = _last_copy_signal->_hsa_signal;
this, ihipCommandName[crit->_last_command_type], ihipCommandName[copyType], crit->_last_copy_signal->_sig_id);
*waitSignal = crit->_last_copy_signal->_hsa_signal;
}
if (HIP_DISABLE_HW_COPY_DEP && needSync) {
@@ -368,10 +378,10 @@ int ihipStream_t::preCopyCommand(ihipSignal_t *lastCopy, hsa_signal_t *waitSigna
}
}
_last_command_type = copyType;
crit->_last_command_type = copyType;
}
_last_copy_signal = lastCopy;
crit->_last_copy_signal = lastCopy;
return needSync;
}
@@ -383,10 +393,13 @@ int ihipStream_t::preCopyCommand(ihipSignal_t *lastCopy, hsa_signal_t *waitSigna
//
//Reset the device - this is called from hipDeviceReset.
//Device may be reset multiple times, and may be reset after init.
void ihipDevice_t::reset()
void ihipDevice_t::locked_reset()
{
// Obtain mutex access to the device critical data, release by destructor
LockedAccessor_DeviceCrit_t crit(_criticalData);
// Reset and remove streams:
_streams.clear();
crit->streams().clear();
// Reset and release all memory stored in the tracker:
am_memtracker_reset(_acc);
@@ -395,13 +408,12 @@ void ihipDevice_t::reset()
//---
void ihipDevice_t::init(unsigned device_index, hc::accelerator acc, unsigned flags)
void ihipDevice_t::init(unsigned device_index, hc::accelerator &acc, unsigned flags)
{
_stream_id = 0;
_device_index = device_index;
_device_flags = flags;
_acc = acc;
hsa_agent_t *agent = static_cast<hsa_agent_t*> (acc.get_hsa_agent());
if (agent) {
int err = hsa_agent_get_info(*agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, &_compute_units);
@@ -416,10 +428,10 @@ void ihipDevice_t::init(unsigned device_index, hc::accelerator acc, unsigned fla
getProperties(&_props);
_default_stream = new ihipStream_t(device_index, acc.get_default_view(), _stream_id++, hipStreamDefault);
this->_streams.push_back(_default_stream);
tprintf(DB_SYNC, "created device with default_stream=%p\n", _default_stream);
_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);
hsa_region_t *pinnedHostRegion;
pinnedHostRegion = static_cast<hsa_region_t*>(_acc.get_hsa_am_system_region());
@@ -429,6 +441,8 @@ void ihipDevice_t::init(unsigned device_index, hc::accelerator acc, unsigned fla
};
ihipDevice_t::~ihipDevice_t()
{
if (_default_stream) {
@@ -670,11 +684,13 @@ hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop)
// Implement "default" stream syncronization
// This waits for all other streams to drain before continuing.
// If waitOnSelf is set, this additionally waits for the default stream to empty.
void ihipDevice_t::syncDefaultStream(bool waitOnSelf)
void ihipDevice_t::locked_syncDefaultStream(bool waitOnSelf)
{
LockedAccessor_DeviceCrit_t crit(_criticalData);
tprintf(DB_SYNC, "syncDefaultStream\n");
for (auto streamI=_streams.begin(); streamI!=_streams.end(); streamI++) {
for (auto streamI=crit->const_streams().begin(); streamI!=crit->const_streams().end(); streamI++) {
ihipStream_t *stream = *streamI;
// Don't wait for streams that have "opted-out" of syncing with NULL stream.
@@ -684,20 +700,39 @@ void ihipDevice_t::syncDefaultStream(bool waitOnSelf)
if (waitOnSelf || (stream != _default_stream)) {
// TODO-hcc - use blocking or active wait here?
// TODO-sync - cudaDeviceBlockingSync
stream->wait();
stream->locked_wait();
}
}
}
}
//---
void ihipDevice_t::locked_addStream(ihipStream_t *s)
{
LockedAccessor_DeviceCrit_t crit(_criticalData);
crit->streams().push_back(s);
s->_id = crit->incStreamId();
}
//---
void ihipDevice_t::locked_removeStream(ihipStream_t *s)
{
LockedAccessor_DeviceCrit_t crit(_criticalData);
crit->streams().remove(s);
}
//---
//Heavyweight synchronization that waits on all streams, ignoring hipStreamNonBlocking flag.
void ihipDevice_t::waitAllStreams()
void ihipDevice_t::locked_waitAllStreams()
{
LockedAccessor_DeviceCrit_t crit(_criticalData);
tprintf(DB_SYNC, "waitAllStream\n");
for (auto streamI=_streams.begin(); streamI!=_streams.end(); streamI++) {
(*streamI)->wait();
for (auto streamI=crit->const_streams().begin(); streamI!=crit->const_streams().end(); streamI++) {
(*streamI)->locked_wait();
}
}
@@ -808,12 +843,12 @@ void ihipInit()
READ_ENV_I(release, HIP_LAUNCH_BLOCKING, CUDA_LAUNCH_BLOCKING, "Make HIP APIs 'host-synchronous', so they block until any kernel launches or data copy commands complete. Alias: CUDA_LAUNCH_BLOCKING." );
READ_ENV_I(release, HIP_DB, 0, "Print various debug info. Bitmask, see hip_hcc.cpp for more information.");
if ((HIP_DB & DB_API) && (HIP_TRACE_API == 0)) {
// Set HIP_TRACE_API before we read it, so it is printed correctly.
// Set HIP_TRACE_API default before we read it, so it is printed correctly.
HIP_TRACE_API = 1;
}
READ_ENV_I(release, HIP_TRACE_API, 0, "Trace each HIP API call. Print function name and return code to stderr as program executes.");
READ_ENV_I(release, HIP_ATP_MARKER, 0, "Add HIP function begin/end to ATP file generated with CodeXL");
READ_ENV_I(release, HIP_STAGING_SIZE, 0, "Size of each staging buffer (in KB)" );
READ_ENV_I(release, HIP_STAGING_BUFFERS, 0, "Number of staging buffers to use in each direction. 0=use hsa_memory_copy.");
READ_ENV_I(release, HIP_PININPLACE, 0, "For unpinned transfers, pin the memory in-place in chunks before doing the copy. Under development.");
@@ -823,6 +858,8 @@ void ihipInit()
READ_ENV_I(release, HIP_DISABLE_HW_KERNEL_DEP, 0, "Disable HW dependencies before kernel commands - instead wait for dependency on host. -1 means ignore these dependencies. (debug mode)");
READ_ENV_I(release, HIP_DISABLE_HW_COPY_DEP, 0, "Disable HW dependencies before copy commands - instead wait for dependency on host. -1 means ifnore these dependencies (debug mode)");
// Some flags have both compile-time and runtime flags - generate a warning if user enables the runtime flag but the compile-time flag is disabled.
if (HIP_DB && !COMPILE_HIP_DB) {
fprintf (stderr, "warning: env var HIP_DB=0x%x but COMPILE_HIP_DB=0. (perhaps enable COMPILE_HIP_DB in src code before compiling?)", HIP_DB);
}
@@ -831,6 +868,10 @@ void ihipInit()
fprintf (stderr, "warning: env var HIP_TRACE_API=0x%x but COMPILE_HIP_TRACE_API=0. (perhaps enable COMPILE_HIP_DB in src code before compiling?)", HIP_DB);
}
if (HIP_ATP_MARKER && !COMPILE_HIP_ATP_MARKER) {
fprintf (stderr, "warning: env var HIP_ATP_MARKER=0x%x but COMPILE_HIP_ATP_MARKER=0. (perhaps enable COMPILE_HIP_DB in src code before compiling?)", HIP_ATP_MARKER);
}
/*
* Build a table of valid compute devices.
@@ -933,14 +974,14 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
ihipDevice_t *device = ihipGetTlsDefaultDevice();
#ifndef HIP_API_PER_THREAD_DEFAULT_STREAM
device->syncDefaultStream(false);
device->locked_syncDefaultStream(false);
#endif
return device->_default_stream;
} else {
// Have to wait for legacy default stream to be empty:
if (!(stream->_flags & hipStreamNonBlocking)) {
tprintf(DB_SYNC, "stream %p wait default stream\n", stream);
stream->getDevice()->_default_stream->wait();
stream->getDevice()->_default_stream->locked_wait();
}
return stream;
@@ -956,7 +997,8 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, hc::accelerator_view **av)
std::call_once(hip_initialized, ihipInit);
stream = ihipSyncAndResolveStream(stream);
stream->preKernelCommand();
stream->lockopen_preKernelCommand();
*av = &stream->_av;
@@ -968,7 +1010,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, hc::accelerator_view **av)
//Called after kernel finishes execution.
void ihipPostLaunchKernel(hipStream_t stream, hc::completion_future &kernelFuture)
{
stream->postKernelCommand(kernelFuture);
stream->lockclose_postKernelCommand(kernelFuture);
if (HIP_LAUNCH_BLOCKING) {
tprintf(DB_SYNC, " stream:%p LAUNCH_BLOCKING for kernel completion\n", stream);
}
@@ -1072,7 +1114,7 @@ void ihipStream_t::setCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_
}
void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind)
void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind)
{
ihipDevice_t *device = this->getDevice();
@@ -1098,7 +1140,7 @@ void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsign
hsa_signal_t depSignal;
if ((kind == hipMemcpyHostToDevice) && (!srcTracked)) {
int depSignalCnt = preCopyCommand(NULL, &depSignal, ihipCommandCopyH2D);
int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyH2D);
if (HIP_STAGING_BUFFERS) {
tprintf(DB_COPY1, "D2H && !dstTracked: staged copy H2D dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
@@ -1109,7 +1151,7 @@ void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsign
}
// The copy waits for inputs and then completes before returning so can reset queue to empty:
this->wait(true);
this->wait(crit, true);
} else {
// TODO - remove, slow path.
tprintf(DB_COPY1, "H2D && ! srcTracked: am_copy dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
@@ -1120,14 +1162,14 @@ void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsign
#endif
}
} else if ((kind == hipMemcpyDeviceToHost) && (!dstTracked)) {
int depSignalCnt = preCopyCommand(NULL, &depSignal, ihipCommandCopyD2H);
int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyD2H);
if (HIP_STAGING_BUFFERS) {
tprintf(DB_COPY1, "D2H && !dstTracked: staged copy D2H dst=%p src=%p sz=%zu\n", dst, src, sizeBytes);
//printf ("staged-copy- read dep signals\n");
device->_staging_buffer[1]->CopyDeviceToHost(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL);
// The copy waits for inputs and then completes before returning so can reset queue to empty:
this->wait(true);
// The copy completes before returning so can reset queue to empty:
this->wait(crit, true);
} else {
// TODO - remove, slow path.
@@ -1139,7 +1181,7 @@ void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsign
#endif
}
} else if (kind == hipMemcpyHostToHost) {
int depSignalCnt = preCopyCommand(NULL, &depSignal, ihipCommandCopyH2H);
int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, ihipCommandCopyH2H);
if (depSignalCnt) {
// host waits before doing host memory copy.
@@ -1154,10 +1196,10 @@ void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsign
hsa_agent_t srcAgent, dstAgent;
setCopyAgents(kind, &commandType, &srcAgent, &dstAgent);
int depSignalCnt = preCopyCommand(NULL, &depSignal, commandType);
int depSignalCnt = preCopyCommand(crit, NULL, &depSignal, commandType);
// Get a completion signal:
ihipSignal_t *ihipSignal = allocSignal();
ihipSignal_t *ihipSignal = allocSignal(crit);
hsa_signal_t copyCompleteSignal = ihipSignal->_hsa_signal;
hsa_signal_store_relaxed(copyCompleteSignal, 1);
@@ -1168,7 +1210,7 @@ void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsign
// This is sync copy, so let's wait for copy right here:
if (hsa_status == HSA_STATUS_SUCCESS) {
waitCopy(ihipSignal); // wait for copy, and return to pool.
waitCopy(crit, ihipSignal); // wait for copy, and return to pool.
} else {
throw ihipException(hipErrorInvalidValue);
}
@@ -1176,10 +1218,19 @@ void ihipStream_t::copySync(void* dst, const void* src, size_t sizeBytes, unsign
}
// Sync copy that acquires lock:
void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind)
{
LockedAccessor_StreamCrit_t crit (_criticalData);
copySync(crit, dst, src, sizeBytes, kind);
}
void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind)
{
LockedAccessor_StreamCrit_t crit(_criticalData);
ihipDevice_t *device = this->getDevice();
if (device == NULL) {
@@ -1193,7 +1244,7 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig
/* As this is a CPU op, we need to wait until all
the commands in current stream are finished.
*/
this->wait();
this->wait(crit);
memcpy(dst, src, sizeBytes);
@@ -1220,8 +1271,7 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig
}
ihipSignal_t *ihip_signal = allocSignal();
ihipSignal_t *ihip_signal = allocSignal(crit);
hsa_signal_store_relaxed(ihip_signal->_hsa_signal, 1);
@@ -1232,7 +1282,7 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig
setCopyAgents(kind, &commandType, &srcAgent, &dstAgent);
hsa_signal_t depSignal;
int depSignalCnt = preCopyCommand(ihip_signal, &depSignal, commandType);
int depSignalCnt = preCopyCommand(crit, ihip_signal, &depSignal, commandType);
tprintf (DB_SYNC, " copy-async, waitFor=%lu completion=#%lu(%lu)\n", depSignalCnt? depSignal.handle:0x0, ihip_signal->_sig_id, ihip_signal->_hsa_signal.handle);
@@ -1242,7 +1292,7 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig
if (hsa_status == HSA_STATUS_SUCCESS) {
if (HIP_LAUNCH_BLOCKING) {
tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpyAsync(%zu)\n", sizeBytes);
this->wait();
this->wait(crit);
}
} else {
// This path can be hit if src or dst point to unpinned host memory.
@@ -1250,7 +1300,7 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig
throw ihipException(hipErrorInvalidValue);
}
} else {
copySync(dst, src, sizeBytes, kind);
copySync(crit, dst, src, sizeBytes, kind);
}
}
}
@@ -1302,6 +1352,8 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **a
// TODO - describe naming convention. ihip _. No accessors. No early returns from functions. Set status to success at top, only set error codes in implementation. No tabs.
// Caps convention _ or camelCase
// if { }
// Should use ihip* data structures inside code rather than app-facing hip. For example, use ihipDevice_t (rather than hipDevice_t), ihipStream_t (rather than hipStream_t).
// locked_
// TODO - describe MT strategy
//
//// TODO - add identifier numbers for streams and devices to help with debugging.
src/hip_memory.cpp
+5 -8
ファイルの表示
@@ -307,17 +307,14 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
hipError_t e = hipSuccess;
try {
stream->copySync(dst, src, sizeBytes, kind);
stream->locked_copySync(dst, src, sizeBytes, kind);
}
catch (ihipException ex) {
e = ex._code;
}
if (HIP_LAUNCH_BLOCKING) {
tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpy\n");
stream->wait();
}
return ihipLogStatus(e);
}
@@ -366,9 +363,9 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s
hipError_t e = hipSuccess;
stream = ihipSyncAndResolveStream(stream);
stream->preKernelCommand();
if (stream) {
stream->lockopen_preKernelCommand();
hc::completion_future cf ;
@@ -392,7 +389,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t s
}
}
stream->postKernelCommand(cf);
stream->lockclose_postKernelCommand(cf);
if (HIP_LAUNCH_BLOCKING) {
@@ -459,7 +456,7 @@ hipError_t hipFree(void* ptr)
hipError_t hipStatus = hipErrorInvalidDevicePointer;
// Synchronize to ensure all work has finished.
ihipGetTlsDefaultDevice()->waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
ihipGetTlsDefaultDevice()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish.
if (ptr) {
hc::accelerator acc;
src/hip_stream.cpp
+34 -17
ファイルの表示
@@ -28,10 +28,8 @@ THE SOFTWARE.
//
//---
hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int flags)
{
std::call_once(hip_initialized, ihipInit);
ihipDevice_t *device = ihipGetTlsDefaultDevice();
hc::accelerator acc = device->_acc;
@@ -41,12 +39,32 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
//Note this is an execute_in_order queue, so all kernels submitted will atuomatically wait for prev to complete:
//This matches CUDA stream behavior:
auto istream = new ihipStream_t(device->_device_index, acc.create_view(), device->_stream_id, flags);
device->_streams.push_back(istream);
auto istream = new ihipStream_t(device->_device_index, acc.create_view(), flags);
device->locked_addStream(istream);
*stream = istream;
tprintf(DB_SYNC, "hipStreamCreate, stream=%p\n", *stream);
return ihipLogStatus(hipSuccess);
return hipSuccess;
}
//---
hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
{
HIP_INIT_API(stream, flags);
return ihipLogStatus(ihipStreamCreate(stream, flags));
}
//---
hipError_t hipStreamCreate(hipStream_t *stream)
{
HIP_INIT_API(stream);
return ihipLogStatus(ihipStreamCreate(stream, hipStreamDefault));
}
@@ -56,8 +74,7 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
*/
hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags)
{
std::call_once(hip_initialized, ihipInit);
HIP_INIT_API(stream, event, flags);
hipError_t e = hipSuccess;
@@ -65,7 +82,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
// TODO-hcc Convert to use create_blocking_marker(...) functionality.
// Currently we have a super-conservative version of this - block on host, and drain the queue.
// This should create a barrier packet in the target queue.
stream->wait();
stream->locked_wait();
e = hipSuccess;
}
@@ -76,15 +93,15 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
//---
hipError_t hipStreamSynchronize(hipStream_t stream)
{
std::call_once(hip_initialized, ihipInit);
HIP_INIT_API(stream);
hipError_t e = hipSuccess;
if (stream == NULL) {
ihipDevice_t *device = ihipGetTlsDefaultDevice();
device->syncDefaultStream(true/*waitOnSelf*/);
device->locked_syncDefaultStream(true/*waitOnSelf*/);
} else {
stream->wait();
stream->locked_wait();
e = hipSuccess;
}
@@ -99,23 +116,23 @@ hipError_t hipStreamSynchronize(hipStream_t stream)
*/
hipError_t hipStreamDestroy(hipStream_t stream)
{
std::call_once(hip_initialized, ihipInit);
HIP_INIT_API(stream);
hipError_t e = hipSuccess;
//--- Drain the stream:
if (stream == NULL) {
ihipDevice_t *device = ihipGetTlsDefaultDevice();
device->syncDefaultStream(true/*waitOnSelf*/);
device->locked_syncDefaultStream(true/*waitOnSelf*/);
} else {
stream->wait();
stream->locked_wait();
e = hipSuccess;
}
ihipDevice_t *device = stream->getDevice();
if (device) {
device->_streams.remove(stream);
device->locked_removeStream(stream);
delete stream;
} else {
e = hipErrorInvalidResourceHandle;
@@ -128,7 +145,7 @@ hipError_t hipStreamDestroy(hipStream_t stream)
//---
hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags)
{
std::call_once(hip_initialized, ihipInit);
HIP_INIT_API(stream, flags);
if (flags == NULL) {
return ihipLogStatus(hipErrorInvalidValue);