diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 8c575eedc0..b62d0c4957 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/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(); + + /** * @} */ diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index 1436008dd1..f4a9fd8e1b 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/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 diff --git a/projects/hip/samples/2_Cookbook/2_Profiler/Makefile b/projects/hip/samples/2_Cookbook/2_Profiler/Makefile index 4b9a063f38..db2d008182 100644 --- a/projects/hip/samples/2_Cookbook/2_Profiler/Makefile +++ b/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) diff --git a/projects/hip/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp b/projects/hip/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp index b6a6b141d2..3747bb4ec5 100644 --- a/projects/hip/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp +++ b/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= 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; } diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index e0a979b1e9..7d79f0a930 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/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: