Empty kernel test enhancements [tools] (#1999)
* Initial commit
* Improvements-1
* Initial commit for PR
* Updates warning, run.sh, decoupled loops
* Forcing seq cst for CPU timimg
[ROCm/rccl commit: 85baa0d113]
Этот коммит содержится в:
Avinash
GitHub
@@ -55,6 +55,7 @@ THE SOFTWARE.
|
||||
#include <algorithm>
|
||||
#include <vector>
|
||||
#include <numeric>
|
||||
#include <stdatomic.h>
|
||||
|
||||
// Helper macro for catching HIP errors
|
||||
#define HIP_CALL(cmd) \
|
||||
@@ -85,13 +86,21 @@ int main(int argc, char **argv)
|
||||
int numIterations = (argc > 1 ? atoi(argv[1]) : 10);
|
||||
int gridSize = (argc > 2 ? atoi(argv[2]) : 1);
|
||||
int blockSize = (argc > 3 ? atoi(argv[3]) : 1);
|
||||
bool cpuTime = (argc > 4 ? atoi(argv[4]) != 0 : true);
|
||||
bool gpuTime = (argc > 5 ? atoi(argv[5]) != 0 : true);
|
||||
bool outerLoop = (argc > 6 ? atoi(argv[6]) != 0 : false);
|
||||
bool innerLoop = (argc > 7 ? atoi(argv[7]) != 0 : true);
|
||||
int numWarmups = 3;
|
||||
printf("Running %d iterations <<<%d,%d>>>\n", numIterations, gridSize, blockSize);
|
||||
printf("Running %d iterations <<<%d,%d>>> cpuTime:%d gpuTime:%d outerLoop:%d innerLoop:%d\n", numIterations, gridSize, blockSize, cpuTime , gpuTime, outerLoop,innerLoop);
|
||||
if(!innerLoop && !outerLoop) {printf("Warning: Atleast one of inner and outer loops should be enabled\n");return 0;}
|
||||
if(outerLoop && innerLoop && (cpuTime || gpuTime)){printf("Warning: given cpuTime %d, gpuTime %d, outerLoop %d may not be meaningful combination, recommended are 001,100,010,110\n",cpuTime , gpuTime, outerLoop);}
|
||||
|
||||
// Create events and stream
|
||||
hipEvent_t startEvent, stopEvent;
|
||||
hipEvent_t startEvent, stopEvent, startEventOuterLoop, stopEventOuterLoop;
|
||||
HIP_CALL(hipEventCreate(&startEvent));
|
||||
HIP_CALL(hipEventCreate(&stopEvent));
|
||||
HIP_CALL(hipEventCreate(&startEventOuterLoop));
|
||||
HIP_CALL(hipEventCreate(&stopEventOuterLoop));
|
||||
hipStream_t stream;
|
||||
HIP_CALL(hipStreamCreate(&stream));
|
||||
|
||||
@@ -108,50 +117,85 @@ int main(int argc, char **argv)
|
||||
// NOTE: Timing is done per-iteration, instead of batching multiple iterations
|
||||
double cpuSum = 0.0;
|
||||
double gpuSum = 0.0;
|
||||
auto cpuStart = std::chrono::high_resolution_clock::now();
|
||||
auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
|
||||
double cpuDeltaMsec;
|
||||
float gpuDeltaMsec;
|
||||
|
||||
auto cpuStartOuterloop = std::chrono::high_resolution_clock::now();
|
||||
auto cpuDeltaOuterloop = std::chrono::high_resolution_clock::now() - cpuStartOuterloop;
|
||||
double cpuDeltaMsecOuterLoop;
|
||||
float gpuDeltaMsecOuterLoop;
|
||||
if(outerLoop && gpuTime) {HIP_CALL(hipEventRecord(startEventOuterLoop, stream));}
|
||||
for (int iteration = 0; iteration < numIterations; iteration++)
|
||||
{
|
||||
// Start timing
|
||||
auto cpuStart = std::chrono::high_resolution_clock::now();
|
||||
HIP_CALL(hipEventRecord(startEvent, stream));
|
||||
if(innerLoop && cpuTime){ atomic_signal_fence(memory_order_seq_cst); cpuStart = std::chrono::high_resolution_clock::now();}
|
||||
if(innerLoop && gpuTime){HIP_CALL(hipEventRecord(startEvent, stream));}
|
||||
|
||||
// Launch kernel and wait for completion
|
||||
EmptyKernel<<<gridSize, blockSize, 0, stream>>>();
|
||||
HIP_CALL(hipEventRecord(stopEvent, stream));
|
||||
HIP_CALL(hipStreamSynchronize(stream));
|
||||
if(innerLoop && gpuTime){HIP_CALL(hipEventRecord(stopEvent, stream));}
|
||||
if(innerLoop){HIP_CALL(hipStreamSynchronize(stream));}
|
||||
|
||||
// Collect timing info
|
||||
auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
|
||||
double cpuDeltaMsec = std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() * 1000.0;
|
||||
float gpuDeltaMsec;
|
||||
HIP_CALL(hipEventElapsedTime(&gpuDeltaMsec, startEvent, stopEvent));
|
||||
if(innerLoop && cpuTime) {
|
||||
atomic_signal_fence(memory_order_seq_cst);
|
||||
cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
|
||||
cpuDeltaMsec = std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() * 1000.0;
|
||||
}
|
||||
if(innerLoop && gpuTime){HIP_CALL(hipEventElapsedTime(&gpuDeltaMsec, startEvent, stopEvent));}
|
||||
|
||||
// Report timing
|
||||
printf("Iteration %03d Kernel Launch Time (usec) %10.5f (CPU) %10.5f (GPU)\n", iteration, cpuDeltaMsec *1000.0, gpuDeltaMsec * 1000.0);
|
||||
allGpuDeltaMsec[iteration] = gpuDeltaMsec * 1000.0;
|
||||
allCpuDeltaMsec[iteration] = cpuDeltaMsec * 1000.0;
|
||||
cpuSum += cpuDeltaMsec * 1000.0;
|
||||
gpuSum += gpuDeltaMsec * 1000.0;
|
||||
if(innerLoop && (cpuTime || gpuTime)) {printf("Iteration %03d Kernel Launch Time (usec) %10.5f (CPU) %10.5f (GPU)\n", iteration, cpuDeltaMsec *1000.0, gpuDeltaMsec * 1000.0);}
|
||||
if(innerLoop && gpuTime){
|
||||
allGpuDeltaMsec[iteration] = gpuDeltaMsec * 1000.0;
|
||||
gpuSum += gpuDeltaMsec * 1000.0;
|
||||
}
|
||||
if(innerLoop && cpuTime){
|
||||
allCpuDeltaMsec[iteration] = cpuDeltaMsec * 1000.0;
|
||||
cpuSum += cpuDeltaMsec * 1000.0;
|
||||
}
|
||||
}
|
||||
if(outerLoop) {
|
||||
atomic_signal_fence(memory_order_seq_cst);
|
||||
cpuDeltaOuterloop = std::chrono::high_resolution_clock::now() - cpuStartOuterloop;
|
||||
if(gpuTime){
|
||||
HIP_CALL(hipEventRecord(stopEventOuterLoop, stream));
|
||||
HIP_CALL(hipStreamSynchronize(stream));
|
||||
HIP_CALL(hipEventElapsedTime(&gpuDeltaMsecOuterLoop, startEventOuterLoop, stopEventOuterLoop));
|
||||
}
|
||||
cpuDeltaMsecOuterLoop = std::chrono::duration_cast<std::chrono::duration<double>>(cpuDeltaOuterloop).count() * 1000.0;
|
||||
}
|
||||
printf("\n");
|
||||
|
||||
// Report averages
|
||||
double avgCpuUsec = cpuSum / numIterations;
|
||||
double avgGpuUsec = gpuSum / numIterations;
|
||||
auto minCpuUsec = std::min_element(std::begin(allCpuDeltaMsec), std::end(allCpuDeltaMsec));
|
||||
auto minGpuUsec = std::min_element(std::begin(allGpuDeltaMsec), std::end(allGpuDeltaMsec));
|
||||
auto maxCpuUsec = std::max_element(std::begin(allCpuDeltaMsec), std::end(allCpuDeltaMsec));
|
||||
auto maxGpuUsec = std::max_element(std::begin(allGpuDeltaMsec), std::end(allGpuDeltaMsec));
|
||||
auto varCpuUsec = calStdDev(allCpuDeltaMsec, avgCpuUsec);
|
||||
auto varGpuUsec = calStdDev(allGpuDeltaMsec, avgGpuUsec);
|
||||
if(innerLoop && (cpuTime || gpuTime)){
|
||||
double avgCpuUsec = cpuSum / numIterations;
|
||||
double avgGpuUsec = gpuSum / numIterations;
|
||||
auto minCpuUsec = std::min_element(std::begin(allCpuDeltaMsec), std::end(allCpuDeltaMsec));
|
||||
auto minGpuUsec = std::min_element(std::begin(allGpuDeltaMsec), std::end(allGpuDeltaMsec));
|
||||
auto maxCpuUsec = std::max_element(std::begin(allCpuDeltaMsec), std::end(allCpuDeltaMsec));
|
||||
auto maxGpuUsec = std::max_element(std::begin(allGpuDeltaMsec), std::end(allGpuDeltaMsec));
|
||||
auto varCpuUsec = calStdDev(allCpuDeltaMsec, avgCpuUsec);
|
||||
auto varGpuUsec = calStdDev(allGpuDeltaMsec, avgGpuUsec);
|
||||
|
||||
printf("Average Kernel Launch time (usec) %10.5f (CPU) %10.5f (GPU)\n", avgCpuUsec, avgGpuUsec);
|
||||
printf("Minimum Kernel Launch time (usec) %10.5f (CPU) %10.5f (GPU)\n", *minCpuUsec, *minGpuUsec);
|
||||
printf("Maximum Kernel Launch time (usec) %10.5f (CPU) %10.5f (GPU)\n", *maxCpuUsec, *maxGpuUsec);
|
||||
printf("Stddev Kernel Launch time (usec) %10.5f (CPU) %10.5f (GPU)\n", varCpuUsec, varGpuUsec);
|
||||
printf("Average Kernel Launch time (usec) %10.5f (CPU) %10.5f (GPU)\n", avgCpuUsec, avgGpuUsec);
|
||||
printf("Minimum Kernel Launch time (usec) %10.5f (CPU) %10.5f (GPU)\n", *minCpuUsec, *minGpuUsec);
|
||||
printf("Maximum Kernel Launch time (usec) %10.5f (CPU) %10.5f (GPU)\n", *maxCpuUsec, *maxGpuUsec);
|
||||
printf("Stddev Kernel Launch time (usec) %10.5f (CPU) %10.5f (GPU)\n", varCpuUsec, varGpuUsec);
|
||||
}
|
||||
if(outerLoop) {
|
||||
double avgCpuUsecOuterLoop = (cpuDeltaMsecOuterLoop* 1000.0)/numIterations;
|
||||
double avgGpuUsecOuterLoop = 0.0;
|
||||
if(gpuTime){ avgGpuUsecOuterLoop = (gpuDeltaMsecOuterLoop* 1000.0)/numIterations; }
|
||||
printf("Average-OL Kernel Launch time (usec) %10.5f (CPU) %10.5f (GPU)\n", avgCpuUsecOuterLoop, avgGpuUsecOuterLoop);
|
||||
}
|
||||
// Cleanup events and stream
|
||||
HIP_CALL(hipStreamDestroy(stream));
|
||||
HIP_CALL(hipEventDestroy(startEvent));
|
||||
HIP_CALL(hipEventDestroy(stopEvent));
|
||||
|
||||
HIP_CALL(hipEventDestroy(startEventOuterLoop));
|
||||
HIP_CALL(hipEventDestroy(stopEventOuterLoop));
|
||||
return 0;
|
||||
}
|
||||
|
||||
Исполняемый файл
+20
@@ -0,0 +1,20 @@
|
||||
#!/bin/bash
|
||||
make clean
|
||||
make
|
||||
echo -e "\n\n"
|
||||
|
||||
echo "# Test 1 : CPU Only"
|
||||
./EmptyKernelTest 10 1 1 1 0
|
||||
echo -e "\n\n"
|
||||
|
||||
echo "# Test 2 : GPU Only"
|
||||
./EmptyKernelTest 10 1 1 0 1
|
||||
echo -e "\n\n"
|
||||
|
||||
echo "# Test 3 : CPU and GPU"
|
||||
./EmptyKernelTest 10 1 1 1 1 0
|
||||
echo -e "\n\n"
|
||||
|
||||
echo "# Test 4 : Outer loop - CPU Only"
|
||||
./EmptyKernelTest 10 1 1 0 0 1
|
||||
echo -e "\n\n"
|
||||
Ссылка в новой задаче
Block a user