Add initial hipProfileStart/Stop

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

Change-Id: I2579824d2dd7863ea23874d34f0dabb3cb305d3e


[ROCm/hip commit: bb58f4f6fc]

projects/hip/include/hip/hcc_detail/hip_runtime_api.h

@@ -1737,11 +1737,24 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
  *
  *  @warning The cudaProfilerInitialize API format for "configFile" is not supported.
  *
- *  On AMD platforms, hipProfilerStart and hipProfilerStop require installation of AMD's GPU
- *  perf counter API and defining GPU_PERF
  */
 
 
+// TODO - expand descriptions:
+/**
+ * @brief Start recording of profiling information
+ * @warning : hipProfilerStart API is under development.
+ */
+hipError_t hipProfilerStart();
+
+
+/**
+ * @brief Stop recording of profiling information.
+ * @warning : hipProfilerStop API is under development.
+ */
+hipError_t hipProfilerStop();
+
+
 /**
  * @}
  */

projects/hip/include/hip/nvcc_detail/hip_runtime_api.h

@@ -643,6 +643,18 @@ inline static hipError_t hipMemcpyPeerAsync ( void* dst, int  dstDevice, const v
     return hipCUDAErrorTohipError(cudaMemcpyPeerAsync(dst, dstDevice, src, srcDevice, count, stream));
 }
 
+// Profile APIs:
+inline hipError_t hipProfilerStart() 
+{ 
+    return hipCUDAErrorTohipError(cudaProfileStart());
+}
+
+inline hipError_t hipProfilerStop() 
+{ 
+    return hipCUDAErrorTohipError(cudaProfileStop());
+}
+
+
 inline static hipError_t hipSetDeviceFlags (unsigned int flags)
 {
     return hipCUDAErrorTohipError(cudaSetDeviceFlags(flags));
@@ -802,6 +814,8 @@ inline static hipError_t hipModuleLaunchKernel(hipFunction_t f,
                     sharedMemBytes, stream, kernelParams, extra));
 }
 
+
+
 #ifdef __cplusplus
 }
 #endif

projects/hip/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)
 

projects/hip/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;
 }

projects/hip/src/hip_hcc.cpp

@@ -63,7 +63,9 @@ int HIP_LAUNCH_BLOCKING = 0;
 int HIP_PRINT_ENV = 0;
 int HIP_TRACE_API= 0;
 std::string HIP_TRACE_API_COLOR("green");
-int HIP_PROFILE_API= 0;
+int HIP_PROFILE_API= 0;S
+
+// TODO - DB_START/STOP need more testing.
 std::string HIP_DB_START_API;
 std::string HIP_DB_STOP_API;
 int HIP_DB= 0;
@@ -1891,6 +1893,31 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
     }
 }
 
+//-------------------------------------------------------------------------------------------------
+//-------------------------------------------------------------------------------------------------
+//Profiler, really these should live elsewhere:
+hipError_t hipProfilerStart()
+{
+    HIP_INIT_API();
+#if COMPILE_HIP_ATP_MARKER
+    amdtResumeProfiling(AMDT_ALL_PROFILING);
+#endif
+
+    return ihipLogStatus(hipSuccess);
+};
+
+
+hipError_t hipProfilerStop()
+{
+    HIP_INIT_API();
+#if COMPILE_HIP_ATP_MARKER
+    amdtStopProfiling(AMDT_ALL_PROFILING);
+#endif
+
+    return ihipLogStatus(hipSuccess);
+};
+
+
 //-------------------------------------------------------------------------------------------------
 //-------------------------------------------------------------------------------------------------
 // HCC-specific accessor functions: