Adding LaunchBench tool (#952)
This commit is contained in:
gilbertlee-amd
GitHub
@@ -0,0 +1,53 @@
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <iostream>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
#define HIP_CALL(cmd) \
|
||||
do { \
|
||||
hipError_t error = (cmd); \
|
||||
if (error != hipSuccess) \
|
||||
{ \
|
||||
std::cout << "Encountered HIP error (" << hipGetErrorString(error) \
|
||||
<< ") at line " << __LINE__ << " in file " << __FILE__ << "\n"; \
|
||||
exit(-1); \
|
||||
} \
|
||||
} while (0)
|
||||
|
||||
// Macro for collecting HW_REG_XCC_ID
|
||||
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
|
||||
#define GetXccId(val) \
|
||||
asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_XCC_ID)" : "=s" (val));
|
||||
#else
|
||||
#define GetXccId(val) \
|
||||
val = 0
|
||||
#endif
|
||||
|
||||
// Macro for collecting HW_REG_HW_ID
|
||||
#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__NVCC__)
|
||||
#define GetHwId(val) \
|
||||
val = 0
|
||||
#else
|
||||
#define GetHwId(val) \
|
||||
asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_HW_ID)" : "=s" (val));
|
||||
#endif
|
||||
@@ -0,0 +1,230 @@
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <cstdio>
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
#include "Common.hpp"
|
||||
#include <omp.h>
|
||||
#include <unistd.h>
|
||||
|
||||
struct SyncData
|
||||
{
|
||||
uint64_t cpuStart;
|
||||
uint64_t cpuStop;
|
||||
int32_t xccId;
|
||||
};
|
||||
|
||||
__global__ void SyncKernel(volatile uint64_t* cpuTime,
|
||||
volatile uint32_t* abortFlag,
|
||||
SyncData* syncData)
|
||||
{
|
||||
// Collect timestamp upon kernel entry
|
||||
uint64_t cpuStart = *cpuTime;
|
||||
|
||||
// Wait for abort flag to be modified
|
||||
while (*abortFlag == 0);
|
||||
|
||||
// Collect timestamps after abort flag
|
||||
uint64_t cpuStop = *cpuTime;
|
||||
|
||||
// Save timestamps
|
||||
syncData[blockIdx.x].cpuStart = cpuStart;
|
||||
syncData[blockIdx.x].cpuStop = cpuStop;
|
||||
GetXccId(syncData[blockIdx.x].xccId);
|
||||
}
|
||||
|
||||
void UpdateCpuTime(volatile uint64_t* cpuTimestamp, volatile bool& abortThread)
|
||||
{
|
||||
while (!abortThread)
|
||||
{
|
||||
*cpuTimestamp = std::chrono::steady_clock::now().time_since_epoch().count();
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
int main(int argc, char **argv)
|
||||
{
|
||||
int numBlocks = (argc > 1 ? atoi(argv[1]) : 4);
|
||||
int numIterations = (argc > 2 ? atoi(argv[2]) : 1);
|
||||
int numWarmups = (argc > 3 ? atoi(argv[3]) : 100);
|
||||
int numSleepUsec = (argc > 4 ? atoi(argv[4]) : 20);
|
||||
|
||||
int totalIterations = numWarmups + numIterations;
|
||||
|
||||
int numGpus;
|
||||
HIP_CALL(hipGetDeviceCount(&numGpus));
|
||||
printf("Running %d GPUs with %d block(s) each, %d timed iterations, %d untimed warmup iterations, sleeping for %d usec\n",
|
||||
numGpus, numBlocks, numIterations, numWarmups, numSleepUsec);
|
||||
for (int i = 0; i < numGpus; i++)
|
||||
{
|
||||
hipDeviceProp_t prop;
|
||||
HIP_CALL(hipGetDeviceProperties(&prop, i));
|
||||
printf("GPU %02d: %s\n", i, prop.gcnArchName);
|
||||
}
|
||||
|
||||
typedef typename std::ratio_multiply<std::chrono::steady_clock::period,std::mega>::type MicroSec;
|
||||
printf("std::chrono::steady_clock precision: %8.3f usec\n",
|
||||
static_cast<double>(MicroSec::num)/MicroSec::den);
|
||||
|
||||
|
||||
// Allocate pinned host memory for CPU timestamp / abort flag
|
||||
volatile uint64_t* cpuTimestamp;
|
||||
volatile uint32_t* abortFlag;
|
||||
HIP_CALL(hipHostMalloc((void**)&cpuTimestamp, sizeof(uint64_t)));
|
||||
HIP_CALL(hipHostMalloc((void**)&abortFlag, sizeof(uint32_t)));
|
||||
|
||||
// Allocate device memory for collecting timestamps
|
||||
std::vector<SyncData*> syncDataList(numGpus);
|
||||
std::vector<hipStream_t>streams(numGpus);
|
||||
for (int i = 0; i < numGpus; i++)
|
||||
{
|
||||
HIP_CALL(hipSetDevice(i));
|
||||
HIP_CALL(hipMalloc((void**)&syncDataList[i], numIterations * numBlocks * sizeof(SyncData)));
|
||||
HIP_CALL(hipStreamCreate(&streams[i]));
|
||||
}
|
||||
|
||||
// Start update thread
|
||||
// NOTE: NPKit usually runs 1 GPU per process which means 1 update thread per GPU
|
||||
// However in this case, only a single CPU update thread is used
|
||||
volatile bool abortThread = false;
|
||||
std::thread updateThread(UpdateCpuTime, cpuTimestamp, std::ref(abortThread));
|
||||
|
||||
// Launch one thread per GPU
|
||||
std::vector<uint64_t> cpuAbortTime(totalIterations);
|
||||
std::vector<std::vector<uint64_t>> cpuStartList(numGpus, std::vector<uint64_t>(totalIterations, 0));
|
||||
std::vector<std::vector<uint64_t>> cpuReturnList(numGpus, std::vector<uint64_t>(totalIterations, 0));
|
||||
std::vector<std::vector<uint64_t>> cpuStopList(numGpus, std::vector<uint64_t>(totalIterations, 0));
|
||||
|
||||
uint64_t cpuAbort;
|
||||
#pragma omp parallel num_threads(numGpus)
|
||||
{
|
||||
int deviceId = omp_get_thread_num();
|
||||
HIP_CALL(hipSetDevice(deviceId));
|
||||
|
||||
for (int iteration = 0; iteration < totalIterations; iteration++)
|
||||
{
|
||||
// Single thread resets abort flag
|
||||
#pragma omp single
|
||||
*abortFlag = 0;
|
||||
|
||||
// Prepare for this iteration
|
||||
SyncData* syncData = syncDataList[deviceId] + (iteration * numBlocks);
|
||||
|
||||
// Wait for all threads to arrive before launching all kernels
|
||||
#pragma omp barrier
|
||||
uint64_t cpuStart = std::chrono::steady_clock::now().time_since_epoch().count();
|
||||
SyncKernel<<<numBlocks, 1, 0, streams[deviceId]>>>(cpuTimestamp, abortFlag, syncData);
|
||||
uint64_t cpuReturn = std::chrono::steady_clock::now().time_since_epoch().count();
|
||||
|
||||
// Busy wait performs more accurately than usleep / sleep_for
|
||||
if (deviceId == 0)
|
||||
{
|
||||
while (std::chrono::steady_clock::now().time_since_epoch().count() - cpuStart < numSleepUsec * 1000);
|
||||
*abortFlag = 1;
|
||||
cpuAbort = std::chrono::steady_clock::now().time_since_epoch().count();
|
||||
}
|
||||
|
||||
// Wait for kernels to finish
|
||||
HIP_CALL(hipStreamSynchronize(streams[deviceId]));
|
||||
uint64_t cpuStop = std::chrono::steady_clock::now().time_since_epoch().count();
|
||||
|
||||
// Store values (after all timings to avoid false sharing)
|
||||
cpuStartList[deviceId][iteration] = cpuStart;
|
||||
cpuReturnList[deviceId][iteration] = cpuReturn;
|
||||
#pragma omp single
|
||||
cpuAbortTime[iteration] = cpuAbort;
|
||||
cpuStopList[deviceId][iteration] = cpuStop;
|
||||
|
||||
#pragma omp barrier
|
||||
}
|
||||
}
|
||||
|
||||
abortThread = true;
|
||||
updateThread.join();
|
||||
|
||||
for (int iteration = 1; iteration <= numIterations; iteration++)
|
||||
{
|
||||
// Ignore warmup iterations
|
||||
int iter = iteration + numWarmups - 1;
|
||||
printf("---------------------------------------------------------------------------------------------------\n");
|
||||
printf("Iteration %d: (All times in usec)\n", iteration);
|
||||
|
||||
uint64_t origin = cpuStartList[0][iter];
|
||||
for (int gpu = 0; gpu < numGpus; gpu++)
|
||||
{
|
||||
for (int block = 0; block < numBlocks; block++)
|
||||
{
|
||||
origin = std::min(origin, cpuStartList[gpu][iter]);
|
||||
origin = std::min(origin, syncDataList[gpu][iter * numBlocks + block].cpuStart);
|
||||
origin = std::min(origin, cpuAbortTime[iter]);
|
||||
origin = std::min(origin, syncDataList[gpu][iter * numBlocks + block].cpuStop);
|
||||
origin = std::min(origin, cpuStopList[gpu][iter]);
|
||||
}
|
||||
}
|
||||
|
||||
printf("| GPU | BLOCK | XCC | START(CPU) | RETURN(CPU)| START(GPU) | ABORT(CPU) | STOP (GPU) | STOP (CPU) |\n");
|
||||
|
||||
double minCpuStart, minGpuStart, minCpuReturn, minCpuAbort, minGpuStop, minCpuStop;
|
||||
double maxCpuStart, maxGpuStart, maxCpuReturn, maxCpuAbort, maxGpuStop, maxCpuStop;
|
||||
|
||||
for (int gpu = 0; gpu < numGpus; gpu++)
|
||||
{
|
||||
for (int block = 0; block < numBlocks; block++)
|
||||
{
|
||||
int xccId = syncDataList[gpu][iter * numBlocks + block].xccId;
|
||||
double cpuStart = (cpuStartList[gpu][iter] - origin) / 1000.0;
|
||||
double gpuStart = (syncDataList[gpu][iter * numBlocks + block].cpuStart - origin) / 1000.0;
|
||||
double cpuReturn = (cpuReturnList[gpu][iter] - origin) / 1000.0;
|
||||
double cpuAbort = (cpuAbortTime[iter] - origin) / 1000.0;
|
||||
double gpuStop = (syncDataList[gpu][iter * numBlocks + block].cpuStop - origin) / 1000.0;
|
||||
double cpuStop = (cpuStopList[gpu][iter] - origin) / 1000.0;
|
||||
|
||||
minCpuStart = ((gpu == 0 && block == 0) || (minCpuStart > cpuStart)) ? cpuStart : minCpuStart;
|
||||
maxCpuStart = ((gpu == 0 && block == 0) || (maxCpuStart < cpuStart)) ? cpuStart : maxCpuStart;
|
||||
minGpuStart = ((gpu == 0 && block == 0) || (minGpuStart > gpuStart)) ? gpuStart : minGpuStart;
|
||||
maxGpuStart = ((gpu == 0 && block == 0) || (maxGpuStart < gpuStart)) ? gpuStart : maxGpuStart;
|
||||
minCpuReturn = ((gpu == 0 && block == 0) || (minCpuReturn > cpuReturn)) ? cpuReturn : minCpuReturn;
|
||||
maxCpuReturn = ((gpu == 0 && block == 0) || (maxCpuReturn < cpuReturn)) ? cpuReturn : maxCpuReturn;
|
||||
minCpuAbort = ((gpu == 0 && block == 0) || (minCpuAbort > cpuAbort)) ? cpuAbort : minCpuAbort;
|
||||
maxCpuAbort = ((gpu == 0 && block == 0) || (maxCpuAbort < cpuAbort)) ? cpuAbort : maxCpuAbort;
|
||||
minGpuStop = ((gpu == 0 && block == 0) || (minGpuStop > gpuStop)) ? gpuStop : minGpuStop;
|
||||
maxGpuStop = ((gpu == 0 && block == 0) || (maxGpuStop < gpuStop)) ? gpuStop : maxGpuStop;
|
||||
minCpuStop = ((gpu == 0 && block == 0) || (minCpuStop > gpuStop)) ? gpuStop : minCpuStop;
|
||||
maxCpuStop = ((gpu == 0 && block == 0) || (maxCpuStop < gpuStop)) ? gpuStop : maxCpuStop;
|
||||
|
||||
printf("| %3d | %3d | %3d | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f |\n",
|
||||
gpu, block, xccId, cpuStart, cpuReturn, gpuStart, cpuAbort, gpuStop, cpuStop);
|
||||
}
|
||||
}
|
||||
printf("---------------------------------------------------------------------------------------------------\n");
|
||||
printf("| MIN | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f |\n",
|
||||
minCpuStart, minCpuReturn, minGpuStart, minCpuAbort, minGpuStop, minCpuStop);
|
||||
printf("| MAX | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f |\n",
|
||||
maxCpuStart, maxCpuReturn, maxGpuStart, maxCpuAbort, maxGpuStop, maxCpuStop);
|
||||
printf("| DIFF | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f | %10.3f |\n",
|
||||
maxCpuStart - minCpuStart, maxCpuReturn - minCpuReturn, maxGpuStart - minGpuStart,
|
||||
maxCpuAbort - minCpuAbort, maxGpuStop - minGpuStop, maxCpuStop - minCpuStop);
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1,11 @@
|
||||
ROCM_DIR = /opt/rocm
|
||||
HIPCC = $(ROCM_DIR)/bin/hipcc
|
||||
CCFLAGS = -lhsa-runtime64 -fopenmp
|
||||
|
||||
all: LaunchBench
|
||||
|
||||
LaunchBench: LaunchBench.cpp Common.hpp
|
||||
$(HIPCC) $(CCFLAGS) $< -o $@
|
||||
|
||||
clean:
|
||||
rm -f ./LaunchBench
|
||||
Reference in New Issue
Block a user