Add initial hipProfileStart/Stop

And modify sample to show how to use.
Still needs some work to understand interaction with CXL.

Change-Id: I2579824d2dd7863ea23874d34f0dabb3cb305d3e

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)
 

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<ITERATIONS; i++) {
 
+    if (i==startTriggerIteration) {
+      hipProfilerStart();
+    }
+    if (i==stopTriggerIteration) {
+      hipProfilerStop();
+    }
+
     float eventMs = 0.0f;
 
     hipEvent_t start, stop;
@@ -129,7 +140,16 @@ void runGPU(float *Matrix, float *TransposeMatrix,
 };
 
 
-int main() {
+int main(int argc, char *argv[]) {
+
+  if (argc >= 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;
 }