diff --git a/tools/EmptyKernelTest/EmptyKernelTest.cpp b/tools/EmptyKernelTest/EmptyKernelTest.cpp index 1ea33da4f6..f9bf31199d 100644 --- a/tools/EmptyKernelTest/EmptyKernelTest.cpp +++ b/tools/EmptyKernelTest/EmptyKernelTest.cpp @@ -55,6 +55,7 @@ THE SOFTWARE. #include #include #include +#include // 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<<>>(); - 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>(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>(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>(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; } diff --git a/tools/EmptyKernelTest/run.sh b/tools/EmptyKernelTest/run.sh new file mode 100755 index 0000000000..31e257bb6a --- /dev/null +++ b/tools/EmptyKernelTest/run.sh @@ -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"