SWDEV-494221 - Remove unnecessary hipEventRecord

- Removed unnecessary hipEventRecord and fixed time calculation in
  hipPerfDispatchSpeed.cc where it was off by a factor of 1,000.

Change-Id: If538e1d236cf0e6d3c69caf7af53c9095d812ad6
Этот коммит содержится в:
Michael Xie
2024-11-21 11:51:04 -08:00
коммит произвёл Rakesh Roy
родитель 274fb5d974
Коммит b1f9f86543
2 изменённых файлов: 30 добавлений и 82 удалений
+16 -42
Просмотреть файл
@@ -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);
}