float finalizeEvents(hipEvent_t start, hipEvent_t stop){
CUDA_SAFE_CALL( hipGetLastError() );
CUDA_SAFE_CALL( hipEventRecord(stop, 0) );
CUDA_SAFE_CALL( hipEventSynchronize(stop) );
float kernel_time;
CUDA_SAFE_CALL( hipEventElapsedTime(&kernel_time, start, stop) );
CUDA_SAFE_CALL( hipEventDestroy(start) );
CUDA_SAFE_CALL( hipEventDestroy(stop) );
return kernel_time;
}
void test(unsigned testMask, int *C_d, int *C_h, int64_t numElements, SyncMode syncMode, bool 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);
int count =100;
int init0 = 0;
HIPCHECK(hipMemset(C_d, init0, sizeBytes));
for (int i=0; i<numElements; i++) {
C_h[i] = -1; // initialize
}
hipStream_t otherStream = 0;
unsigned flags = (syncMode == syncMarkerThenOtherNonBlockingStream) ? hipStreamNonBlocking : hipStreamDefault;
HIPCHECK(hipStreamCreateWithFlags(&otherStream, flags));
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*/,
static_cast<const int*>(C_d),
C_h,
numElements,
count);
HIPCHECK(hipEventRecord(stop, 0/*default*/));
switch (syncMode) {
case syncNone:
break;
case syncNullStream:
HIPCHECK(hipStreamSynchronize(0)); // wait on host for null stream:
break;
case syncOtherStream:
// Does this synchronize with the null stream?
HIPCHECK(hipStreamSynchronize(otherStream));
break;
case syncMarkerThenOtherStream:
case syncMarkerThenOtherNonBlockingStream:
// this may wait for NULL stream depending hipStreamNonBlocking flag above
HIPCHECK(hipEventRecord(otherStreamEvent, otherStream));
HIPCHECK(hipStreamSynchronize(otherStream));
break;
case syncDevice:
HIPCHECK(hipDeviceSynchronize());
break;
default:
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<numElements; i++) {
bool compareEqual = (C_h[i] == expected);
if (!compareEqual) {
mismatches ++;
if (!expectMismatch) {
printf ("C_h[%d] (%d) != %d\n", i, C_h[i], expected);
assert(C_h[i] == expected);
}
}
}
if (expectMismatch) {
assert (mismatches > 0);
}
HIPCHECK(hipStreamDestroy(otherStream));
HIPCHECK(hipEventDestroy(stop));
HIPCHECK(hipEventDestroy(otherStreamEvent));
HIPCHECK(hipDeviceSynchronize());
printf ("test: OK - %d mismatches (%6.2f%%)\n", mismatches, ((double)(mismatches)*100.0)/numElements);
}
