From 39c18e5e5ff1685ee84e1e35e28fb0231834e7e5 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 5 Jun 2017 00:41:18 -0500 Subject: [PATCH] Fix HIP_SYNC_NULL_STREAM=0 mode. - 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 --- include/hip/hcc_detail/hip_runtime_api.h | 10 +- src/hip_event.cpp | 74 ++++----- src/hip_hcc.cpp | 56 ++++--- src/hip_hcc_internal.h | 10 +- src/hip_stream.cpp | 30 ++-- tests/src/runtimeApi/event/record_event.cpp | 141 ++++++++++++------ .../src/runtimeApi/stream/hipStreamSync2.cpp | 72 ++++++--- 7 files changed, 249 insertions(+), 144 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 25eac31ec6..fde38c8395 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -658,10 +658,12 @@ hipError_t hipStreamSynchronize(hipStream_t stream); * * This function inserts a wait operation into the specified stream. * All future work submitted to @p stream will wait until @p event reports completion before beginning execution. - * This function is host-asynchronous and the function may return before the wait has completed. + * + * This function only waits for commands in the current stream to complete. Notably,, this function does + * not impliciy wait for commands in the default stream to complete, even if the specified stream is + * created with hipStreamNonBlocking = 0. * * @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamSynchronize, hipStreamDestroy - * */ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags); @@ -766,10 +768,10 @@ hipError_t hipEventCreate(hipEvent_t* event); * the specified stream, after all previous * commands in that stream have completed executing. * - * If hipEventRecord() has been previously called aon event, then this call will overwrite any existing state in event. + * If hipEventRecord() has been previously called on this event, then this call will overwrite any existing state in event. * * If this function is called on a an event that is currently being recorded, results are undefined - either - * outstanding recording may save state into the event, and the order is not guaranteed. This shoul be avoided. + * outstanding recording may save state into the event, and the order is not guaranteed. * * @see hipEventCreate, hipEventCreateWithFlags, hipEventQuery, hipEventSynchronize, hipEventDestroy, hipEventElapsedTime * diff --git a/src/hip_event.cpp b/src/hip_event.cpp index 8ef652489a..ab1c43a00b 100644 --- a/src/hip_event.cpp +++ b/src/hip_event.cpp @@ -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 (_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); } diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 5d8846da1e..0cdc57eaab 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -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()); } } } diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h index b15d5a73e4..c3f8b72311 100644 --- a/src/hip_hcc_internal.h +++ b/src/hip_hcc_internal.h @@ -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; }; diff --git a/src/hip_stream.cpp b/src/hip_stream.cpp index 9f1228d6f7..40aade28b9 100644 --- a/src/hip_stream.cpp +++ b/src/hip_stream.cpp @@ -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); diff --git a/tests/src/runtimeApi/event/record_event.cpp b/tests/src/runtimeApi/event/record_event.cpp index 66027b1643..bd8a3ada8e 100644 --- a/tests/src/runtimeApi/event/record_event.cpp +++ b/tests/src/runtimeApi/event/record_event.cpp @@ -28,8 +28,8 @@ THE SOFTWARE. enum SyncMode { syncNone, - syncNullStream, - syncOtherStream, + syncStream, + syncStopEvent, }; @@ -37,19 +37,23 @@ const char *syncModeString(int syncMode) { switch (syncMode) { case syncNone: return "syncNone"; - case syncNullStream: - return "syncNullStream"; - case syncOtherStream: - return "syncOtherStream"; + case syncStream: + return "syncStream"; + case syncStopEvent: + return "syncStopEvent"; default: return "unknown"; }; }; -void test(int *C_d, int *C_h, int64_t numElements, SyncMode syncMode) +void test(unsigned testMask, int *C_d, int *C_h, int64_t numElements, hipStream_t stream, int waitStart, SyncMode syncMode) { - printf ("\ntest: syncMode=%s\n", syncModeString(syncMode)); + if (!(testMask & p_tests)) { + return; + } + printf ("\ntest 0x%3x: stream=%p waitStart=%d syncMode=%s\n", + testMask, stream, waitStart, syncModeString(syncMode)); size_t sizeBytes = numElements * sizeof(int); @@ -60,55 +64,95 @@ void test(int *C_d, int *C_h, int64_t numElements, SyncMode syncMode) C_h[i] = -1; // initialize } - hipStream_t stream = 0; + hipEvent_t neverCreated=0, neverRecorded, timingDisabled; + HIPCHECK(hipEventCreate(&neverRecorded)); + HIPCHECK(hipEventCreateWithFlags(&timingDisabled, hipEventDisableTiming)); - unsigned flags=0; - if (syncMode == syncOtherStream) { - HIPCHECK(hipStreamCreateWithFlags(&stream, flags)); - } - - hipEvent_t neverCreated=0; - hipEvent_t start, stop, neverRecorded; + hipEvent_t start, stop; HIPCHECK(hipEventCreate(&start)); HIPCHECK(hipEventCreate(&stop)); - HIPCHECK(hipEventCreate(&neverRecorded)); unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements); + HIPCHECK(hipEventRecord(timingDisabled, stream)); // sandwhich a kernel: HIPCHECK(hipEventRecord(start, stream)); hipLaunchKernelGGL(HipTest::addCountReverse , dim3(blocks), dim3(threadsPerBlock), 0, stream, C_d, C_h, numElements, count); HIPCHECK(hipEventRecord(stop, stream)); - HIPCHECK(hipStreamSynchronize(stream)); // wait for recording to finish... + + if (waitStart) { + HIPCHECK(hipEventSynchronize(start)); + } + + + hipError_t expectedStopError = hipSuccess; + + // How to wait for the events to finish: + switch (syncMode) { + case syncNone: + expectedStopError = hipErrorNotReady; + break; + case syncStream: + HIPCHECK(hipStreamSynchronize(stream)); // wait for recording to finish... + break; + case syncStopEvent: + HIPCHECK(hipEventSynchronize(stop)); + break; + default: + assert(0); + }; + float t; - HIPCHECK_API(hipEventElapsedTime(&t, neverCreated, stop), hipErrorInvalidResourceHandle); - HIPCHECK_API(hipEventElapsedTime(&t, start, neverCreated), hipErrorInvalidResourceHandle); - HIPCHECK_API(hipEventElapsedTime(&t, neverRecorded, stop), hipErrorInvalidResourceHandle); - HIPCHECK_API(hipEventElapsedTime(&t, start, neverRecorded), hipErrorInvalidResourceHandle); - - HIPCHECK(hipEventElapsedTime(&t, start, stop)); - assert (t>0.0f); - printf ("time=%6.2f\n", t); - - HIPCHECK(hipEventElapsedTime(&t, stop, start)); - assert (t<0.0f); - printf ("negtime=%6.2f\n", t); - - HIPCHECK(hipEventElapsedTime(&t, start, start)); - assert (t==0.0f); - HIPCHECK(hipEventElapsedTime(&t, stop, stop)); - assert (t==0.0f); - - - if (stream) { - HIPCHECK(hipStreamDestroy(stream)); + hipError_t e = hipEventElapsedTime(&t, start, start); + if ((e != hipSuccess) && (e != hipErrorNotReady)) { + failed ("start event not in expected state, was %d=%s\n", e, hipGetErrorName(e)); } + + if (e == hipSuccess) + assert (t==0.0f); + + + // stop usually ready unless we skipped the synchronization (syncNone) + HIPCHECK_API(hipEventElapsedTime(&t, stop, stop), expectedStopError); + if (e == hipSuccess) + assert (t==0.0f); + + + e = hipEventElapsedTime(&t, start, stop); + HIPCHECK_API(e, expectedStopError); + if (expectedStopError == hipSuccess) + assert (t>0.0f); + printf ("time=%6.2f error=%s\n", t, hipGetErrorName(e)); + + e = hipEventElapsedTime(&t, stop, start); + HIPCHECK_API(e, expectedStopError); + if (expectedStopError == hipSuccess) + assert (t<0.0f); + printf ("negtime=%6.2f error=%s\n", t, hipGetErrorName(e)); + + + + { + // Check some error conditions for incomplete events: + HIPCHECK_API(hipEventElapsedTime(&t, timingDisabled, stop), hipErrorInvalidResourceHandle); + HIPCHECK_API(hipEventElapsedTime(&t, start, timingDisabled), hipErrorInvalidResourceHandle); + + HIPCHECK_API(hipEventElapsedTime(&t, neverCreated, stop), hipErrorInvalidResourceHandle); + HIPCHECK_API(hipEventElapsedTime(&t, start, neverCreated), hipErrorInvalidResourceHandle); + + HIPCHECK_API(hipEventElapsedTime(&t, neverRecorded, stop), hipErrorInvalidResourceHandle); + HIPCHECK_API(hipEventElapsedTime(&t, start, neverRecorded), hipErrorInvalidResourceHandle); + } + HIPCHECK(hipEventDestroy(start)); HIPCHECK(hipEventDestroy(stop)); + // Clear out everything: + HIPCHECK(hipDeviceSynchronize()); + printf ("test: OK \n"); } @@ -125,15 +169,22 @@ void runTests(int64_t numElements) HIPCHECK(hipMalloc(&C_d, sizeBytes)); HIPCHECK(hipHostMalloc(&C_h, sizeBytes)); + hipStream_t stream; + HIPCHECK(hipStreamCreateWithFlags(&stream, 0x0)); - { - test (C_d, C_h, numElements, syncNone); - test (C_d, C_h, numElements, syncNullStream); - test (C_d, C_h, numElements, syncOtherStream); - //test (C_d, C_h, numElements, syncDevice); + //for (int waitStart=0; waitStart<2; waitStart++) { + for (int waitStart=1; waitStart>=0; waitStart--) { + unsigned W = waitStart ? 0x1000:0; + test (W | 0x01, C_d, C_h, numElements, 0 , waitStart, syncNone); + test (W | 0x02, C_d, C_h, numElements, stream, waitStart, syncNone); + test (W | 0x04, C_d, C_h, numElements, 0 , waitStart, syncStream); + test (W | 0x08, C_d, C_h, numElements, stream, waitStart, syncStream); + test (W | 0x10, C_d, C_h, numElements, 0, waitStart, syncStopEvent); + test (W | 0x20, C_d, C_h, numElements, stream, waitStart, syncStopEvent); } + HIPCHECK(hipStreamDestroy(stream)); HIPCHECK(hipFree(C_d)); HIPCHECK(hipHostFree(C_h)); } @@ -143,7 +194,7 @@ int main(int argc, char *argv[]) { HipTest::parseStandardArguments(argc, argv, true /*failOnUndefinedArg*/); - runTests(4000000); + runTests(80000000); passed(); } diff --git a/tests/src/runtimeApi/stream/hipStreamSync2.cpp b/tests/src/runtimeApi/stream/hipStreamSync2.cpp index b57e120dcc..c6a58ce7d4 100644 --- a/tests/src/runtimeApi/stream/hipStreamSync2.cpp +++ b/tests/src/runtimeApi/stream/hipStreamSync2.cpp @@ -56,9 +56,27 @@ const char *syncModeString(int syncMode) { }; -void test(int *C_d, int *C_h, int64_t numElements, SyncMode syncMode, bool expectMismatch) +void test(unsigned testMask, int *C_d, int *C_h, int64_t numElements, SyncMode syncMode, bool expectMismatch) { - printf ("\ntest: syncMode=%s expectMismatch=%d\n", syncModeString(syncMode), expectMismatch); + + // This test sends a long-running kernel to the null stream, then tests to see if the + // specified synchronization technique is effective. + // + // Some syncMode are not expected to correctly sync (for example "syncNone"). in these + // cases the test sets expectMismatch and the check logic below will attempt to ensure that + // the undesired synchronization did not occur - ie ensure the kernel is still running and did + // not yet update the stop event. This can be tricky since if the kernel runs fast enough it + // may complete before the check. To prevent this, the addCountReverse has a count parameter + // which causes it to loop repeatedly, and the results are checked in reverse order. + // + // Tests with expectMismatch=true should ensure the kernel finishes correctly. This results + // are checked and we test to make sure stop event has completed. + + if (!(testMask & p_tests)) { + return; + } + printf ("\ntest 0x%02x: syncMode=%s expectMismatch=%d\n", + testMask, syncModeString(syncMode), expectMismatch); size_t sizeBytes = numElements * sizeof(int); @@ -72,13 +90,15 @@ void test(int *C_d, int *C_h, int64_t numElements, SyncMode syncMode, bool expec hipStream_t otherStream = 0; unsigned flags = (syncMode == syncMarkerThenOtherNonBlockingStream) ? hipStreamNonBlocking : hipStreamDefault; HIPCHECK(hipStreamCreateWithFlags(&otherStream, flags)); - hipEvent_t e; - HIPCHECK(hipEventCreate(&e)); + hipEvent_t stop, otherStreamEvent; + HIPCHECK(hipEventCreate(&stop)); + HIPCHECK(hipEventCreate(&otherStreamEvent)); unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, numElements); // Launch kernel into null stream, should result in C_h == count. hipLaunchKernelGGL(HipTest::addCountReverse , dim3(blocks), dim3(threadsPerBlock), 0, 0 /*stream*/, C_d, C_h, numElements, count); + HIPCHECK(hipEventRecord(stop, 0/*default*/)); switch (syncMode) { case syncNone: @@ -92,7 +112,10 @@ void test(int *C_d, int *C_h, int64_t numElements, SyncMode syncMode, bool expec break; case syncMarkerThenOtherStream: case syncMarkerThenOtherNonBlockingStream: - HIPCHECK(hipEventRecord(e, otherStream)); // this may wait for NULL stream depending hipStreamNonBlocking flag above + + // this may wait for NULL stream depending hipStreamNonBlocking flag above + HIPCHECK(hipEventRecord(otherStreamEvent, otherStream)); + HIPCHECK(hipStreamSynchronize(otherStream)); break; case syncDevice: @@ -102,6 +125,14 @@ void test(int *C_d, int *C_h, int64_t numElements, SyncMode syncMode, bool expec assert(0); }; + hipError_t done = hipEventQuery(stop); + + if (expectMismatch) { + assert (done == hipErrorNotReady); + } else { + assert (done == hipSuccess); + } + int mismatches = 0; int expected = init0 + count; for (int i=0; i