AI/rocm-systems
2
0
Förgrening 0
Kod Ärenden Pull-förfrågningar Actions Packages Projekt Släpp Wiki Aktiviteter

Fix HIP_SYNC_NULL_STREAM=0 mode.

Bläddra i källkod 
- Fix null-stream sync
- hipStreamDestroy of null stream returns hipErrorInvalidResourceHandle
- Update documentation.
- Add tests for null stream sync, hipEventElapsedTime.
- Rename internal enum hipEventStatusRecorded to hipEventStatusComplete
- refactor hipStreamWaitEvent to streamline control-flow
This commit is contained in:
Ben Sander
2017-06-05 00:41:18 -05:00
förälder 8ce6d17983
incheckning 39c18e5e5f
7 ändrade filer med 249 tillägg och 144 borttagningar
Download Patch File Download Diff File Expand all files Collapse all files
src/hip_event.cpp
+39 -35
Visa fil
@@ -53,15 +53,12 @@ void ihipEvent_t::attachToCompletionFuture(const hc::completion_future *cf,
void ihipEvent_t::setTimestamp()
void ihipEvent_t::refereshEventStatus()
{
bool isReady0 = _marker.is_ready();
bool isReady1;
int val = 0;
if (_state == hipEventStatusRecorded) {
// already recorded, done:
return;
} else {
if (_state == hipEventStatusRecording) {
// TODO - use completion-future functions to obtain ticks and timestamps:
hsa_signal_t *sig = static_cast<hsa_signal_t*> (_marker.get_native_handle());
isReady1 = _marker.is_ready();
@@ -78,12 +75,12 @@ void ihipEvent_t::setTimestamp()
_timestamp = 0;
}
_state = hipEventStatusRecorded;
_state = hipEventStatusComplete;
}
}
}
if (_state != hipEventStatusRecorded) {
if (_state != hipEventStatusComplete) {
//printf (" not ready isReady0=%d val=%d isReady1=%d\n", isReady0, val, isReady1);
}
}
@@ -103,12 +100,10 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags)
const unsigned releaseFlags = (hipEventReleaseToDevice | hipEventReleaseToSystem);
const bool illegalFlags = (flags & ~supportedFlags) || // can't set any unsupported flags.
(flags & releaseFlags) == releaseFlags; // can't set both
(flags & releaseFlags) == releaseFlags; // can't set both release flags
if (!illegalFlags) {
ihipEvent_t *eh = new ihipEvent_t(flags);
*event = eh;
*event = new ihipEvent_t(flags);
} else {
e = hipErrorInvalidValue;
}
@@ -148,7 +143,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
ctx->locked_syncDefaultStream(true, true);
event->_timestamp = hc::get_system_ticks();
event->_state = hipEventStatusRecorded;
event->_state = hipEventStatusComplete;
return ihipLogStatus(hipSuccess);
} else {
event->_state = hipEventStatusRecording;
@@ -209,41 +204,50 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop)
{
HIP_INIT_API(ms, start, stop);
start->setTimestamp();
stop->setTimestamp();
hipError_t status = hipSuccess;
*ms = 0.0f;
if (start && stop) {
// refresh status:
if ((start->_state == hipEventStatusRecorded) && (stop->_state == hipEventStatusRecorded)) {
// Common case, we have good information for both events.
if ((start == nullptr) ||
(start->_flags & hipEventDisableTiming) ||
(start->_state == hipEventStatusUnitialized) || (start->_state == hipEventStatusCreated) ||
(stop == nullptr) ||
(stop->_flags & hipEventDisableTiming) ||
( stop->_state == hipEventStatusUnitialized) || ( stop->_state == hipEventStatusCreated)) {
int64_t tickDiff = (stop->timestamp() - start->timestamp());
// Both events must be at least recorded else return hipErrorInvalidResourceHandle
uint64_t freqHz;
hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &freqHz);
if (freqHz) {
*ms = ((double)(tickDiff) / (double)(freqHz)) * 1000.0f;
status = hipSuccess;
} else {
* ms = 0.0f;
status = hipErrorInvalidValue;
}
status = hipErrorInvalidResourceHandle;
} else {
// Refresh status, if still recording...
start->refereshEventStatus();
stop->refereshEventStatus();
if ((start->_state == hipEventStatusComplete) && (stop->_state == hipEventStatusComplete)) {
// Common case, we have good information for both events.
int64_t tickDiff = (stop->timestamp() - start->timestamp());
uint64_t freqHz;
hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &freqHz);
if (freqHz) {
*ms = ((double)(tickDiff) / (double)(freqHz)) * 1000.0f;
status = hipSuccess;
} else {
* ms = 0.0f;
status = hipErrorInvalidValue;
}
} else if ((start->_state == hipEventStatusRecording) ||
(stop->_state == hipEventStatusRecording)) {
status = hipErrorNotReady;
} else if ((start->_state == hipEventStatusUnitialized) ||
(stop->_state == hipEventStatusUnitialized)) {
status = hipErrorInvalidResourceHandle;
} else {
assert(0);
}
} else {
status = hipErrorInvalidResourceHandle;
}
}
return ihipLogStatus(status);
}
src/hip_hcc.cpp
+35 -21
Visa fil
@@ -92,7 +92,8 @@ int HIP_COHERENT_HOST_ALLOC = 0;
// USE_ HIP_SYNC_HOST_ALLOC
int HIP_SYNC_HOST_ALLOC = 1;
// Sync on host between
// Chicken bit to sync on host to implement null stream.
// If 0, null stream synchronization is performed on the GPU
int HIP_SYNC_NULL_STREAM = 1;
// HIP needs to change some behavior based on HCC_OPT_FLUSH :
@@ -987,11 +988,17 @@ std::string ihipCtx_t::toString() const
// Implement "default" stream syncronization
// This waits for all other streams to drain before continuing.
// This called for submissions that are sent to the null/default stream. This routine ensures
// that this new command waits for activity in the other streams to complete before proceeding.
//
// HIP_SYNC_NULL_STREAM=0 does all dependency resolutiokn on the GPU
// HIP_SYNC_NULL_STREAM=1 s legacy non-optimal mode which conservatively waits on host.
//
// If waitOnSelf is set, this additionally waits for the default stream to empty.
// In new HIP_SYNC_NULL_STREAM=0 mode, this enqueues a marker which causes the default stream to wait for other
// activity, but doesn't actually block the host. If host blocking is desired, the caller should set syncHost.
//
// syncToHost causes host to wait for the stream to finish.
// Note HIP_SYNC_NULL_STREAM=1 path always sync to Host.
void ihipCtx_t::locked_syncDefaultStream(bool waitOnSelf, bool syncHost)
{
@@ -1005,34 +1012,36 @@ void ihipCtx_t::locked_syncDefaultStream(bool waitOnSelf, bool syncHost)
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.
// And - don't wait for the NULL stream, unless waitOnSelf specified.
bool waitThisStream = (!(stream->_flags & hipStreamNonBlocking)) &&
(waitOnSelf || (stream != _defaultStream));
if (HIP_SYNC_NULL_STREAM) {
// Don't wait for streams that have "opted-out" of syncing with NULL stream.
// And - don't wait for the NULL stream
if (!(stream->_flags & hipStreamNonBlocking)) {
if (waitOnSelf || (stream != _defaultStream)) {
stream->locked_wait();
}
if (waitThisStream) {
stream->locked_wait();
}
} else {
if (!(stream->_flags & hipStreamNonBlocking) && (stream != _defaultStream)) {
if (waitThisStream) {
LockedAccessor_StreamCrit_t streamCrit(stream->_criticalData);
// The last marker will provide appropriate visibility:
if (!streamCrit->_av.get_is_empty()) {
depOps.push_back(streamCrit->_av.create_marker(hc::accelerator_scope));
tprintf(DB_SYNC, " push marker to wait for stream=%s\n", ToString(stream).c_str());
} else {
tprintf(DB_SYNC, " skipped stream=%s since it is empty\n", ToString(stream).c_str());
}
}
}
}
// Enqueue a barrier to wait on all the barriers we sent above:
if (!HIP_SYNC_NULL_STREAM && !depOps.empty()) {
LockedAccessor_StreamCrit_t defaultStreamCrit(_defaultStream->_criticalData);
tprintf(DB_SYNC, " null-stream wait on %zu non-empty streams\n", depOps.size());
tprintf(DB_SYNC, " null-stream wait on %zu non-empty streams. sync_host=%d\n", depOps.size(), syncHost);
hc::completion_future defaultCf = defaultStreamCrit->_av.create_blocking_marker(depOps.begin(), depOps.end(), hc::accelerator_scope);
if (syncHost) {
defaultCf.wait(); // TODO - account for active or blocking here.
@@ -1374,6 +1383,7 @@ void ihipInit()
hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
{
if (stream == hipStreamNull ) {
// Submitting to NULL stream, call locked_syncDefaultStream to wait for all other streams:
ihipCtx_t *ctx = ihipGetTlsDefaultCtx();
tprintf(DB_SYNC, "ihipSyncAndResolveStream %s wait on default stream\n", ToString(stream).c_str());
@@ -1382,34 +1392,38 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
#endif
return ctx->_defaultStream;
} else {
// All streams have to wait for legacy default stream to be empty:
// Submitting to a "normal" stream, just wait for null stream:
if (!(stream->_flags & hipStreamNonBlocking)) {
if (HIP_SYNC_NULL_STREAM) {
tprintf(DB_SYNC, "ihipSyncAndResolveStream %s wait on default stream\n", ToString(stream).c_str());
tprintf(DB_SYNC, "ihipSyncAndResolveStream %s host-wait on default stream\n", ToString(stream).c_str());
stream->getCtx()->_defaultStream->locked_wait();
} else {
ihipStream_t *defaultStream = stream->getCtx()->_defaultStream;
tprintf(DB_SYNC, "%s marker wait default stream\n", ToString(stream).c_str());
bool needMarker = false;
bool needGatherMarker = false; // used to gather together other markers.
hc::completion_future dcf;
{
LockedAccessor_StreamCrit_t defaultStreamCrit(defaultStream->criticalData());
// TODO - could call create_blocking_marker(queue)
// TODO - could call create_blocking_marker(queue) or uses existing marker.
if (!defaultStreamCrit->_av.get_is_empty()) {
needMarker = true;
needGatherMarker = true;
// TODO - add "none_scope".
tprintf(DB_SYNC, " %s adding marker to default %s for dependency\n",
ToString(stream).c_str(), ToString(defaultStream).c_str());
dcf = defaultStreamCrit->_av.create_marker(hc::accelerator_scope);
} else {
tprintf(DB_SYNC, " %s skipping marker since default stream is empty\n", ToString(stream).c_str());
}
}
if (needMarker) {
if (needGatherMarker) {
// ensure any commands sent to this stream wait on the NULL stream before continuing
LockedAccessor_StreamCrit_t thisStreamCrit(stream->criticalData());
// TODO - could be "noret" version of create_blocking_marker
thisStreamCrit->_av.create_blocking_marker(dcf, hc::accelerator_scope);
tprintf(DB_SYNC, " %s adding marker to wait for freshly recorded default-stream marker \n",
ToString(stream).c_str());
}
}
}
src/hip_hcc_internal.h
+5 -5
Visa fil
@@ -586,10 +586,10 @@ private: // Data
//----
// Internal event structure:
enum hipEventStatus_t {
hipEventStatusUnitialized = 0, // event is unutilized, must be "Created" before use.
hipEventStatusCreated = 1,
hipEventStatusRecording = 2, // event has been enqueued to record something.
hipEventStatusRecorded = 3, // event has been recorded - timestamps are valid.
hipEventStatusUnitialized = 0, // event is uninitialized, must be "Created" before use.
hipEventStatusCreated = 1, // event created, but not yet Recorded
hipEventStatusRecording = 2, // event has been recorded into a stream but not completed yet.
hipEventStatusComplete = 3, // event has been recorded - timestamps are valid.
} ;
// TODO - rename to ihip type of some kind
@@ -604,7 +604,7 @@ class ihipEvent_t {
public:
ihipEvent_t(unsigned flags);
void attachToCompletionFuture(const hc::completion_future *cf, hipStream_t stream, ihipEventType_t eventType);
void setTimestamp();
void refereshEventStatus();
uint64_t timestamp() const { return _timestamp; } ;
ihipEventType_t type() const { return _type; };
src/hip_stream.cpp
+13 -17
Visa fil
@@ -93,20 +93,17 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
} else if (event->_state != hipEventStatusUnitialized) {
bool fastWait = false;
if (stream != hipStreamNull) {
// This will user create_blocking_marker to wait on the specified queue.
stream->locked_waitEvent(event);
fastWait = true; // don't use the slow host-side synchronization.
}
if (!fastWait) {
} else {
// 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.
// TODO-HIP_SYNC_NULL_STREAM
stream->locked_wait();
e = hipSuccess;
}
} // else event not recorded, return immediately and don't create marker.
@@ -150,6 +147,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream)
ihipCtx_t *ctx = ihipGetTlsDefaultCtx();
ctx->locked_syncDefaultStream(true/*waitOnSelf*/, true/*syncToHost*/);
} else {
// note this does not synchornize with the NULL stream:
stream->locked_wait();
e = hipSuccess;
}
@@ -171,20 +169,18 @@ hipError_t hipStreamDestroy(hipStream_t stream)
//--- Drain the stream:
if (stream == NULL) {
ihipCtx_t *ctx = ihipGetTlsDefaultCtx();
ctx->locked_syncDefaultStream(true/*waitOnSelf*/, true /*syncToHost*/);
e = hipErrorInvalidResourceHandle; // TODO - review - what happens if try to destroy null stream
} else {
stream->locked_wait();
e = hipSuccess;
}
ihipCtx_t *ctx = stream->getCtx();
ihipCtx_t *ctx = stream->getCtx();
if (ctx) {
ctx->locked_removeStream(stream);
delete stream;
} else {
e = hipErrorInvalidResourceHandle;
if (ctx) {
ctx->locked_removeStream(stream);
delete stream;
} else {
e = hipErrorInvalidResourceHandle;
}
}
return ihipLogStatus(e);