diff --git a/projects/hip-tests/catch/perftests/dispatch/hipPerfDispatchSpeed.cc b/projects/hip-tests/catch/perftests/dispatch/hipPerfDispatchSpeed.cc index 66e3985600..897999d8b7 100644 --- a/projects/hip-tests/catch/perftests/dispatch/hipPerfDispatchSpeed.cc +++ b/projects/hip-tests/catch/perftests/dispatch/hipPerfDispatchSpeed.cc @@ -104,42 +104,24 @@ TEST_CASE("Perf_hipPerfDispatchSpeed") { for (; test <= numTests; test++) { int openTest = test % testListSize; bool sleep = false; - bool doWarmup = false; - if ((test / testListSize) % 2) { - doWarmup = true; - } if (test >= (testListSize * 2)) { sleep = true; } int threads = (bufSize_ / sizeof(float)); int threads_per_block = 64; int blocks = (threads/threads_per_block) + (threads % threads_per_block); - hipEvent_t start, stop; - // NULL stream check: - err = hipEventCreate(&start); - REQUIRE(err == hipSuccess); - err = hipEventCreate(&stop); + // warmup + hipLaunchKernelGGL(_dispatchSpeed, dim3(blocks), dim3(threads_per_block), + 0, hipStream_t(0), srcBuffer); + err = hipDeviceSynchronize(); REQUIRE(err == hipSuccess); - if (doWarmup) { - hipLaunchKernelGGL(_dispatchSpeed, dim3(blocks), dim3(threads_per_block), - 0, hipStream_t(0), srcBuffer); - err = hipDeviceSynchronize(); - REQUIRE(err == hipSuccess); - } - - // Do a warm up event record and sync - HIP_CHECK(hipEventRecord(start, NULL)); - HIP_CHECK(hipStreamSynchronize(0)); - - auto Start = std::chrono::high_resolution_clock::now(); + auto start = std::chrono::high_resolution_clock::now(); for (unsigned int i = 0; i < testList[openTest].iterations; i++) { - HIP_CHECK(hipEventRecord(start, NULL)); hipLaunchKernelGGL(_dispatchSpeed, dim3(blocks), dim3(threads_per_block), 0, hipStream_t(0), srcBuffer); - HIP_CHECK(hipEventRecord(stop, NULL)); if ((testList[openTest].flushEvery > 0) && (((i + 1) % testList[openTest].flushEvery) == 0)) { if (sleep) { @@ -147,7 +129,7 @@ TEST_CASE("Perf_hipPerfDispatchSpeed") { REQUIRE(err == hipSuccess); } else { do { - err = hipEventQuery(stop); + err = hipStreamQuery(NULL); } while (err == hipErrorNotReady); } } @@ -157,20 +139,17 @@ TEST_CASE("Perf_hipPerfDispatchSpeed") { REQUIRE(err == hipSuccess); } else { do { - err = hipEventQuery(stop); + err = hipStreamQuery(NULL); } while (err == hipErrorNotReady); } - auto Stop = std::chrono::high_resolution_clock::now(); - HIP_CHECK(hipEventDestroy(start)); - HIP_CHECK(hipEventDestroy(stop)); - double sec = std::chrono::duration(Stop - Start).count(); + auto stop = std::chrono::high_resolution_clock::now(); + double microSec = std::chrono::duration(stop - start).count(); // microseconds per launch - double perf = (1000000.f*sec/testList[openTest].iterations); + double perf = (microSec/testList[openTest].iterations); const char *waitType; const char *extraChar; const char *n; - const char *warmup; if (sleep) { waitType = "sleep"; extraChar = ""; @@ -180,20 +159,15 @@ TEST_CASE("Perf_hipPerfDispatchSpeed") { n = "n"; extraChar = " "; } - if (doWarmup) { - warmup = "warmup"; - } else { - warmup = ""; - } char buf[256]; if (testList[openTest].flushEvery > 0) { - SNPRINTF(buf, sizeof(buf), "HIPPerfDispatchSpeed[%3d] %7d dispatches %s%sing every %5d %6s (us/disp) %3f", + SNPRINTF(buf, sizeof(buf), "HIPPerfDispatchSpeed[%3d] %7d dispatches %s%sing every %5d (us/disp) %3f", test, testList[openTest].iterations, - waitType, n, testList[openTest].flushEvery, warmup, (float)perf); + waitType, n, testList[openTest].flushEvery, (float)perf); } else { - SNPRINTF(buf, sizeof(buf), "HIPPerfDispatchSpeed[%3d] %7d dispatches (%s%s) %6s (us/disp) %3f", + SNPRINTF(buf, sizeof(buf), "HIPPerfDispatchSpeed[%3d] %7d dispatches (%s%s) (us/disp) %3f", test, testList[openTest].iterations, - waitType, extraChar, warmup, (float)perf); + waitType, extraChar, (float)perf); } printf("%s\n", buf); } diff --git a/projects/hip-tests/perftests/dispatch/hipPerfDispatchSpeed.cpp b/projects/hip-tests/perftests/dispatch/hipPerfDispatchSpeed.cpp index 35db04874a..56a757a547 100644 --- a/projects/hip-tests/perftests/dispatch/hipPerfDispatchSpeed.cpp +++ b/projects/hip-tests/perftests/dispatch/hipPerfDispatchSpeed.cpp @@ -110,12 +110,7 @@ int main(int argc, char* argv[]) { { int openTest = test % testListSize; bool sleep = false; - bool doWarmup = false; - if ((test / testListSize) % 2) - { - doWarmup = true; - } if (test >= (testListSize * 2)) { sleep = true; @@ -124,24 +119,12 @@ int main(int argc, char* argv[]) { int threads = (bufSize_ / sizeof(float)); int threads_per_block = 64; int blocks = (threads/threads_per_block) + (threads % threads_per_block); - hipEvent_t start, stop; - // NULL stream check: - err = hipEventCreate(&start); - err = hipEventCreate(&stop); - - CHECK_RESULT(err != hipSuccess, "hipEventCreate failed"); - - // Do a warm up event record and sync - hipEventRecord(start, NULL); - hipStreamSynchronize(0); - - if (doWarmup) - { - hipLaunchKernelGGL(_dispatchSpeed, dim3(blocks), dim3(threads_per_block), 0, hipStream_t(0), srcBuffer); - err = hipDeviceSynchronize(); - CHECK_RESULT(err != hipSuccess, "hipDeviceSynchronize failed"); - } + // warmup + hipLaunchKernelGGL(_dispatchSpeed, dim3(blocks), dim3(threads_per_block), + 0, hipStream_t(0), srcBuffer); + err = hipDeviceSynchronize(); + CHECK_RESULT(err != hipSuccess, "hipDeviceSynchronize failed"); CPerfCounter timer; @@ -149,9 +132,8 @@ int main(int argc, char* argv[]) { timer.Start(); for (unsigned int i = 0; i < testList[openTest].iterations; i++) { - hipEventRecord(start, NULL); - hipLaunchKernelGGL(_dispatchSpeed, dim3(blocks), dim3(threads_per_block), 0, hipStream_t(0), srcBuffer); - hipEventRecord(stop, NULL); + hipLaunchKernelGGL(_dispatchSpeed, dim3(blocks), dim3(threads_per_block), + 0, hipStream_t(0), srcBuffer); if ((testList[openTest].flushEvery > 0) && (((i + 1) % testList[openTest].flushEvery) == 0)) @@ -164,7 +146,7 @@ int main(int argc, char* argv[]) { else { do { - err = hipEventQuery(stop); + err = hipStreamQuery(NULL); } while (err == hipErrorNotReady); } } @@ -177,13 +159,11 @@ int main(int argc, char* argv[]) { else { do { - err = hipEventQuery(stop); + err = hipStreamQuery(NULL); } while (err == hipErrorNotReady); } timer.Stop(); - hipEventDestroy(start); - hipEventDestroy(stop); double sec = timer.GetElapsedTime(); // microseconds per launch @@ -191,7 +171,6 @@ int main(int argc, char* argv[]) { const char *waitType; const char *extraChar; const char *n; - const char *warmup; if (sleep) { waitType = "sleep"; @@ -204,26 +183,21 @@ int main(int argc, char* argv[]) { n = "n"; extraChar = " "; } - if (doWarmup) - { - warmup = "warmup"; - } - else - { - warmup = ""; - } char buf[256]; if (testList[openTest].flushEvery > 0) { - SNPRINTF(buf, sizeof(buf), "HIPPerfDispatchSpeed[%3d] %7d dispatches %s%sing every %5d %6s (us/disp) %3f", test, testList[openTest].iterations, - waitType, n, testList[openTest].flushEvery, warmup, (float)perf); + SNPRINTF(buf, sizeof(buf), + "HIPPerfDispatchSpeed[%3d] %7d dispatches %s%sing every %5d (us/disp) %3f", + test, testList[openTest].iterations, + waitType, n, testList[openTest].flushEvery, (float)perf); } else { - SNPRINTF(buf, sizeof(buf), "HIPPerfDispatchSpeed[%3d] %7d dispatches (%s%s) %6s (us/disp) %3f", test, testList[openTest].iterations, - waitType, extraChar, warmup, (float)perf); + SNPRINTF(buf, sizeof(buf), + "HIPPerfDispatchSpeed[%3d] %7d dispatches (%s%s) (us/disp) %3f", + test, testList[openTest].iterations, waitType, extraChar, (float)perf); } printf("%s\n", buf); }