diff --git a/samples/2_Cookbook/2_Profiler/Makefile b/samples/2_Cookbook/2_Profiler/Makefile index 4b9a063f38..db2d008182 100644 --- a/samples/2_Cookbook/2_Profiler/Makefile +++ b/samples/2_Cookbook/2_Profiler/Makefile @@ -35,6 +35,13 @@ profile: $(EXECUTABLE) $(HIPPROFILER_POST_CMD) +# Pass option to control start and stop iterations for profiling - see MatrixTranspose.cpp for implementation: +# Note we start profiler in --startdisabled mode - no timing collected until app enabled it via hipProfilerStart() +profile_trigger: $(EXECUTABLE) + $(HIPPROFILER) $(PROFILER_OPT) --startdisabled $(EXECUTABLE) 3 6 + $(HIPPROFILER_POST_CMD) + + run: $(EXECUTABLE) $(EXECUTABLE) diff --git a/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp b/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp index b6a6b141d2..3747bb4ec5 100644 --- a/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp +++ b/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp @@ -36,6 +36,10 @@ THE SOFTWARE. #define ITERATIONS 10 +// Cmdline parms to control start and stop triggers +int startTriggerIteration=-1; +int stopTriggerIteration=-1; + // Device (Kernel) function, it must be void // hipLaunchParm provides the execution configuration __global__ void matrixTranspose(hipLaunchParm lp, @@ -74,6 +78,13 @@ void runGPU(float *Matrix, float *TransposeMatrix, for (int i=0; i= 2) { + startTriggerIteration = atoi(argv[1]); + printf ("info : will start tracing at iteration:%d\n", startTriggerIteration); + } + if (argc >= 3) { + stopTriggerIteration = atoi(argv[2]); + printf ("info : will stop tracing at iteration:%d\n", stopTriggerIteration); + } float* Matrix; float* TransposeMatrix; @@ -166,6 +186,8 @@ int main() { // allocate the memory on the device side hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + + // FYI, the scoped-marker will be destroyed here when the scope exits, and will record its "end" timestamp. } runGPU(Matrix, TransposeMatrix, gpuMatrix, gpuTransposeMatrix); @@ -204,7 +226,8 @@ int main() { free(TransposeMatrix); free(cpuTransposeMatrix); - HIP_END_MARKER(); + // This ends the last marker started in this thread, in this case "Check&TearDown" + HIP_END_MARKER(); return errors; }