AI/rocm-systems
2
0
Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

Update bandwidth and latency calculations

Bladeren bron 
- Refined bandwidth and latency calculations for improved accuracy
This commit is contained in:
avinashkethineedi
2025-02-17 06:18:46 +00:00
bovenliggende 40bd8a38a0
commit c155636da4
26 gewijzigde bestanden met toevoegingen van 283 en 270 verwijderingen
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files
tests/functional_tests/tester.cpp
+59 -108
Bestand weergeven
@@ -58,14 +58,22 @@ Tester::Tester(TesterArguments args) : args(args) {
_shmem_context = args.shmem_context;
CHECK_HIP(hipGetDevice(&device_id));
CHECK_HIP(hipGetDeviceProperties(&deviceProps, device_id));
num_warps = (args.wg_size - 1) / deviceProps.warpSize + 1;
wf_size = deviceProps.warpSize;
num_warps = (args.wg_size - 1) / wf_size + 1;
CHECK_HIP(hipStreamCreate(&stream));
CHECK_HIP(hipEventCreate(&start_event));
CHECK_HIP(hipEventCreate(&stop_event));
CHECK_HIP(hipMalloc((void**)&timer, sizeof(uint64_t) * args.num_wgs));
CHECK_HIP(hipDeviceGetAttribute(&wall_clk_rate,
hipDeviceAttributeWallClockRate, device_id));
num_timers = args.num_wgs;
CHECK_HIP(hipMalloc((void**)&timer, sizeof(long long int) * num_timers));
CHECK_HIP(hipMalloc((void**)&start_time, sizeof(long long int) * num_timers));
CHECK_HIP(hipMalloc((void**)&end_time, sizeof(long long int) * num_timers));
}
Tester::~Tester() {
CHECK_HIP(hipFree(end_time));
CHECK_HIP(hipFree(start_time));
CHECK_HIP(hipFree(timer));
CHECK_HIP(hipEventDestroy(stop_event));
CHECK_HIP(hipEventDestroy(start_event));
@@ -75,11 +83,6 @@ Tester::~Tester() {
std::vector<Tester*> Tester::create(TesterArguments args) {
int rank = args.myid;
std::vector<Tester*> testers;
hipDeviceProp_t deviceProps;
int device_id, numWarps;
CHECK_HIP(hipGetDevice(&device_id));
CHECK_HIP(hipGetDeviceProperties(&deviceProps, device_id));
numWarps = (args.wg_size - 1) / deviceProps.warpSize + 1;
if (rank == 0) std::cout << "### Creating Test: ";
@@ -352,35 +355,23 @@ std::vector<Tester*> Tester::create(TesterArguments args) {
testers.push_back(new ShmemPtrTester(args));
return testers;
case WGGetTestType:
if (rank == 0) {
if (args.num_wgs > 1)
std::cout << "Tiled Blocking WG level Gets ###" << std::endl;
else std::cout << "Blocking WG level Gets ###" << std::endl;
}
if (rank == 0)
std::cout << "Blocking WG level Gets ###" << std::endl;
testers.push_back(new ExtendedPrimitiveTester(args));
return testers;
case WGGetNBITestType:
if (rank == 0) {
if (args.num_wgs > 1)
std::cout << "Tiled Non-Blocking WG level Gets ###" << std::endl;
else std::cout << "Non-Blocking WG level Gets ###" << std::endl;
}
if (rank == 0)
std::cout << "Non-Blocking WG level Gets ###" << std::endl;
testers.push_back(new ExtendedPrimitiveTester(args));
return testers;
case WGPutTestType:
if (rank == 0) {
if (args.num_wgs > 1)
std::cout << "Tiled Blocking WG level Puts ###" << std::endl;
else std::cout << "Blocking WG level Puts ###" << std::endl;
}
if (rank == 0)
std::cout << "Blocking WG level Puts ###" << std::endl;
testers.push_back(new ExtendedPrimitiveTester(args));
return testers;
case WGPutNBITestType:
if (rank == 0) {
if (args.num_wgs > 1)
std::cout << "Tiled Non-Blocking WG level Puts ###" << std::endl;
else std::cout << "Non-Blocking WG level Puts ###" << std::endl;
}
if (rank == 0)
std::cout << "Non-Blocking WG level Puts ###" << std::endl;
testers.push_back(new ExtendedPrimitiveTester(args));
return testers;
case PutNBIMRTestType:
@@ -389,35 +380,23 @@ std::vector<Tester*> Tester::create(TesterArguments args) {
testers.push_back(new PrimitiveMRTester(args));
return testers;
case WAVEGetTestType:
if (rank == 0) {
if (args.num_wgs > 1 || numWarps > 1)
std::cout << "Tiled Blocking WAVE level Gets ###" << std::endl;
else std::cout << "Blocking WAVE level Gets ###" << std::endl;
}
if (rank == 0)
std::cout << "Blocking WAVE level Gets ###" << std::endl;
testers.push_back(new WaveLevelPrimitiveTester(args));
return testers;
case WAVEGetNBITestType:
if (rank == 0) {
if (args.num_wgs > 1 || numWarps > 1)
std::cout << "Tiled Non-Blocking WAVE level Gets ###" << std::endl;
else std::cout << "Non-Blocking WAVE level Gets ###" << std::endl;
}
if (rank == 0)
std::cout << "Non-Blocking WAVE level Gets ###" << std::endl;
testers.push_back(new WaveLevelPrimitiveTester(args));
return testers;
case WAVEPutTestType:
if (rank == 0) {
if (args.num_wgs > 1 || numWarps > 1)
std::cout << "Tiled Blocking WAVE level Puts ###" << std::endl;
else std::cout << "Blocking WAVE level Puts ###" << std::endl;
}
if (rank == 0)
std::cout << "Blocking WAVE level Puts ###" << std::endl;
testers.push_back(new WaveLevelPrimitiveTester(args));
return testers;
case WAVEPutNBITestType:
if (rank == 0) {
if (args.num_wgs > 1 || numWarps > 1)
std::cout << "Tiled Non-Blocking WAVE level Puts ###" << std::endl;
else std::cout << "Non-Blocking WAVE level Puts ###" << std::endl;
}
if (rank == 0)
std::cout << "Non-Blocking WAVE level Puts ###" << std::endl;
testers.push_back(new WaveLevelPrimitiveTester(args));
return testers;
case PutSignalTestType:
@@ -501,11 +480,6 @@ void Tester::execute() {
* rocshmem pes.
*/
if (peLaunchesKernel()) {
/**
* TODO:
* Verify that this timer type is actually uint64_t on the
* device side.
*/
memset(timer, 0, sizeof(uint64_t) * args.num_wgs);
const dim3 blockSize(args.wg_size, 1, 1);
@@ -521,9 +495,6 @@ void Tester::execute() {
if (err != hipSuccess) {
printf("error = %d \n", err);
}
// rocshmem_dump_stats();
// rocshmem_reset_stats();
}
barrier();
@@ -533,32 +504,10 @@ void Tester::execute() {
// data validation
verifyResults(size);
/**
* Adjust size for *_wg and *_wave functions
*/
uint64_t size_ = size;
TestType type = (TestType)args.algorithm;
switch (type) {
case WAVEGetTestType:
case WAVEGetNBITestType:
case WAVEPutTestType:
case WAVEPutNBITestType:
size_ *= (args.num_wgs * num_warps);
break;
case WGGetTestType:
case WGGetNBITestType:
case WGPutTestType:
case WGPutNBITestType:
size_ *= args.num_wgs;
break;
default:
break;
}
barrier();
if (_type != TeamCtxInfraTestType) {
print(size_);
print(size);
}
}
}
@@ -589,30 +538,39 @@ void Tester::print(uint64_t size) {
return;
}
uint64_t timer_avg = timerAvgInMicroseconds();
double latency_avg = static_cast<double>(timer_avg) / num_timed_msgs;
/**
* Calculate total amount of data transfered
*/
uint64_t total_size = size * num_timed_msgs;
double timer_avg = timerAvgInMicroseconds();
double latency_avg = timer_avg / num_timed_msgs;
double avg_msg_rate = num_timed_msgs / (timer_avg / 1e6);
float total_kern_time_ms;
CHECK_HIP(hipEventElapsedTime(&total_kern_time_ms, start_event, stop_event));
float total_kern_time_s = total_kern_time_ms / 1000;
double time_us = gpuCyclesToMicroseconds(max_end_time - min_start_time);
double time_s = time_us / 1e6;
double bandwidth_avg_gbs =
num_msgs * size * bw_factor / total_kern_time_s / pow(2, 30);
static_cast<double>(total_size * bw_factor) / time_s / pow(2, 30);
int field_width = 20;
int float_precision = 2;
if (_print_header) {
printf("%-*s%*s%*s%*s",
10, "# Size (B)",
printf("%-*s%-*s%*s%*s%*s",
15, "# Size (B)",
15, "# of timed Msgs",
field_width, "Latency (us)",
field_width, "Bandwidth (GB/s)",
field_width + 1, "Msg Rate (Msg/s)\n");
_print_header = 0;
}
printf("%-*lu%*.*f%*.*f%*.*f\n",
10, size,
printf("%-*lu%-*d%*.*f%*.*f%*.*f\n",
15, size,
15, num_timed_msgs,
field_width, float_precision, latency_avg,
field_width, float_precision, bandwidth_avg_gbs,
field_width, float_precision, avg_msg_rate);
@@ -634,33 +592,26 @@ void Tester::barrier() {
flush_hdp();
}
uint64_t Tester::gpuCyclesToMicroseconds(uint64_t cycles) {
/**
* The dGPU asm core timer runs at 27MHz. This is different from the
* core clock returned by HIP. For an APU, this is different and might
* need adjusting.
*/
uint64_t gpu_frequency_MHz = 27;
/**
* hipDeviceGetAttribute(&gpu_frequency_khz,
* hipDeviceAttributeClockRate,
* 0);
*/
return cycles / gpu_frequency_MHz;
double Tester::gpuCyclesToMicroseconds(long long int cycles) {
return static_cast<double>(cycles) /
(static_cast<double>(wall_clk_rate) * 1e-3);
}
uint64_t Tester::timerAvgInMicroseconds() {
uint64_t sum = 0;
double Tester::timerAvgInMicroseconds() {
double sum = 0;
min_start_time = LLONG_MAX;
max_end_time = 0;
/**
* TODO: (bpotter/avinash) Modify the calcuation for the Tiled version of
* puts and gets at wavefront level
*/
for (uint64_t i = 0; i < args.num_wgs; i++) {
for (uint32_t i = 0; i < num_timers; i++) {
timer[i] = end_time[i] - start_time[i];
sum += gpuCyclesToMicroseconds(timer[i]);
min_start_time = (start_time[i] < min_start_time)
? start_time[i]
: min_start_time;
max_end_time = (end_time[i] > max_end_time)
? end_time[i]
: max_end_time;
}
return sum / args.num_wgs;
return sum / num_timers;
}