AI/rocm-systems
2
0
Tükrözés 0
Kód Hibajegyek Egyesítési kérések Actions Packages Projektek Kiadások Wiki Tevékenység

SWDEV-494221 - Remove unnecessary hipEventRecord

Forráskód böngészése 
- Removed unnecessary hipEventRecord and fixed time calculation in
  hipPerfDispatchSpeed.cc where it was off by a factor of 1,000.

Change-Id: If538e1d236cf0e6d3c69caf7af53c9095d812ad6


[ROCm/hip-tests commit: b1f9f86543]
This commit is contained in:
Michael Xie
2024-11-21 11:51:04 -08:00
committed by Rakesh Roy
szülő 415bd0878b
commit b2f415e866
2 fájl változott, egészen pontosan 30 új sor hozzáadva és 82 régi sor törölve
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/perftests/dispatch/hipPerfDispatchSpeed.cc
+14 -40
Fájl megtekintése
@@ -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<double, std::milli>(Stop - Start).count();
auto stop = std::chrono::high_resolution_clock::now();
double microSec = std::chrono::duration<double, std::micro>(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);
}
projects/hip-tests/perftests/dispatch/hipPerfDispatchSpeed.cpp
+16 -42
Fájl megtekintése
@@ -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);
}