From fa8deac1add34aeb7351536082faadd6b8a197d0 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 23 Mar 2016 02:57:52 -0500 Subject: [PATCH] Improve trace API - Validate compile-time disables. - Add README.md section explain how to install/use CodeXL tracing - Add code docs on trace_helper.h - fix color on hipLaunchKernel to green. --- hipamd/README.md | 44 +++++++++++++++++++++++- hipamd/include/hcc_detail/hip_runtime.h | 8 +++-- hipamd/include/hcc_detail/trace_helper.h | 27 +++++++-------- hipamd/src/hip_hcc.cpp | 41 ++++++++++++---------- 4 files changed, 84 insertions(+), 36 deletions(-) diff --git a/hipamd/README.md b/hipamd/README.md index 8f7b659eb8..acb7c9e4a8 100644 --- a/hipamd/README.md +++ b/hipamd/README.md @@ -45,9 +45,51 @@ Verify your can find hipconfig (one of the hip tools in bin dir): ``` > hipconfig -pn /home/me/HIP +``` + +### HCC Options + +#### Compiling CodeXL markers for HIP Functions +HIP can generate markers at function begin/end which are displayed on the CodeXL timeline view. To do this, you need to install CodeXL, tell HIP +where the CodeXL install directory lives, and enable HIP to generate the markers: + +1. Install CodeXL +See [CodeXL Download](http://developer.amd.com/tools-and-sdks/opencl-zone/codexl/?webSyncID=9d9c2cb9-3d73-5e65-268a-c7b06428e5e0&sessionGUID=29beacd0-d654-ddc6-a3e2-b9e6c0b0cc77) for the installation file. +Also this [blog](http://gpuopen.com/getting-up-to-speed-with-the-codexl-gpu-profiler-and-radeon-open-compute/) provides more information and tips for using CodeXL. In addition to installing the CodeXL profiling +and visualization tools, CodeXL also comes with an SDK that allow applications to add markers to the timeline viewer. We'll be linking HIP against this library. + +2. Set CODEXL_PATH +``` +# set to your code-xl installation location: +export CODEXL_PATH=/opt/AMD/CodeXL +``` + +3. Enable in source code. +In src/hip_hcc.cpp, enable the define +``` +#define COMPILE_TRACE_MARKER 1 +``` -### Using HIP with the AMD Native-GCN compiler. +Then recompile the target application, run with profiler enabled to generate ATP file or trace log. +``` +# Use profiler to generate timeline view: +$ $CODEXL_PATH/CodeXLGpuProfiler -A -o ./myHipApp +... +Session output path: /home/me/HIP-privatestaging/tests/b1/mytrace.atp +``` + +You can also print the HIP function strings to stderr using HIP_TRACE_API environment variable. This can be useful for tracing application flow. Also can be combined with the more detailed debug information provided +by the HIP_DB switch. For example: +``` +# Trace to stderr showing begin/end of each function (with arguments) + intermediate debug trace during the execution of each function. +$ HIP_TRACE_API=1 HIP_DB=0x2 ./myHipApp +``` + +Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors. + + +#### Using HIP with the AMD Native-GCN compiler. AMD recently released a direct-to-GCN-ISA target. This compiler generates GCN ISA directly from LLVM, without going through an intermediate compiler IR such as HSAIL or PTX. The native GCN target is included with upstream LLVM, and has also been integrated with HCC compiler and can be used to compiler HIP programs for AMD. diff --git a/hipamd/include/hcc_detail/hip_runtime.h b/hipamd/include/hcc_detail/hip_runtime.h index b54e40ef65..b9a6981fee 100644 --- a/hipamd/include/hcc_detail/hip_runtime.h +++ b/hipamd/include/hcc_detail/hip_runtime.h @@ -485,6 +485,10 @@ __device__ inline float __dsqrt_rz(double x) {return hc::fast_math::sqrt(x); }; hipStream_t ihipPreLaunchKernel(hipStream_t stream, hc::accelerator_view **av); void ihipPostLaunchKernel(hipStream_t stream, hc::completion_future &cf); +// TODO - move to common header file. +#define KNRM "\x1B[0m" +#define KGRN "\x1B[32m" + #if not defined(DISABLE_GRID_LAUNCH) #define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \ do {\ @@ -500,7 +504,7 @@ do {\ lp.cf = &cf; \ hipStream_t trueStream = (ihipPreLaunchKernel(_stream, &lp.av)); \ if (HIP_TRACE_API) {\ - fprintf(stderr, "==hip-api: launch '%s' gridDim:[%d.%d.%d] groupDim:[%d.%d.%d] groupMem:+%d stream=%p\n", \ + fprintf(stderr, KGRN "< #include +//--- +// Helper functions to convert HIP function arguments into strings. +// Handles POD data types as well as enumerations (ie hipMemcpyKind). +// The implementation uses C++11 variadic templates and template specialization. +// The hipMemcpyKind example below is a good example that shows how to implement conversion for a new HSA type. + + +// Handy macro to convert an enumeration to a stringified version of same: #define CASE_STR(x) case x: return #x; @@ -15,7 +23,7 @@ std::string ToHexString(T v) }; - +//--- // Template overloads for ToString to handle various types: // Note these use C++11 variadic templates template @@ -26,17 +34,6 @@ std::string ToString(T v) { }; -#if 0 -template <> -std::string ToString(void* v) { - std::ostringstream ss; - //ss << "0x" << std::setw(16) << std::setfill('0') << std::hex << v; - ss << "0x" << std::hex << v; - return ss.str(); -}; -#endif - - template <> std::string ToString(hipMemcpyKind v) { switch(v) { @@ -61,11 +58,11 @@ std::string ToString() { return (""); } - /// + +//--- // C++11 variadic template - peels off first argument, converts to string, and calls itself again to peel the next arg. +// Strings are automatically separated by comma+space. template std::string ToString(T first, Args... args) { return ToString(first) + ", " + ToString(args...) ; } - - diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index 33f520dde0..dc9960f24f 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -48,6 +48,9 @@ THE SOFTWARE. #include "hcc_detail/staging_buffer.h" +// TODO, re-org header order. +extern const char *ihipErrorString(hipError_t hip_error); +#include "hcc_detail/trace_helper.h" #define INLINE static inline @@ -79,6 +82,18 @@ int HIP_DISABLE_HW_KERNEL_DEP = 1; int HIP_DISABLE_HW_COPY_DEP = 1; +// Color defs for debug messages: +#define KNRM "\x1B[0m" +#define KRED "\x1B[31m" +#define KGRN "\x1B[32m" +#define KYEL "\x1B[33m" +#define KBLU "\x1B[34m" +#define KMAG "\x1B[35m" +#define KCYN "\x1B[36m" +#define KWHT "\x1B[37m" + +#define API_COLOR KGRN + #define HIP_HCC @@ -111,7 +126,7 @@ int HIP_DISABLE_HW_COPY_DEP = 1; // Compile code that generate #ifndef COMPILE_TRACE_MARKER -#define COMPILE_TRACE_MARKER 1 +#define COMPILE_TRACE_MARKER 0 #endif @@ -119,15 +134,11 @@ int HIP_DISABLE_HW_COPY_DEP = 1; #define ONE_OBJECT_FILE 1 -// TODO, re-org header order. -extern const char *ihipErrorString(hipError_t hip_error); - // Compile support for trace markers that are displayed on CodeXL GUI at start/stop of each function boundary. // TODO - currently we print the trace message at the beginning. if we waited, we could also include return codes, and any values returned // through ptr-to-args (ie the pointers allocated by hipMalloc). #if COMPILE_TRACE_MARKER #include "AMDTActivityLogger.h" -#include "hcc_detail/trace_helper.h" #define SCOPED_MARKER(markerName,group,userString) amdtScopedMarker(markerName, group, userString) #else // Swallow scoped markers: @@ -146,25 +157,19 @@ extern const char *ihipErrorString(hipError_t hip_error); } #else // Swallow API_TRACE -#define API_TRACE() +#define API_TRACE(...) #endif + + +// This macro should be called at the beginning of every HIP API. +// It initialies the hip runtime (exactly once), and +// generate trace to stderr or to ATP file. #define HIP_INIT_API(...) \ std::call_once(hip_initialized, ihipInit);\ API_TRACE(__VA_ARGS__); -// Color defs for debug messages: -#define KNRM "\x1B[0m" -#define KRED "\x1B[31m" -#define KGRN "\x1B[32m" -#define KYEL "\x1B[33m" -#define KBLU "\x1B[34m" -#define KMAG "\x1B[35m" -#define KCYN "\x1B[36m" -#define KWHT "\x1B[37m" - -#define API_COLOR KGRN //--- @@ -1098,7 +1103,7 @@ void ihipReadEnv_I(int *var_ptr, const char *var_name1, const char *var_name2, c //It is called with C++11 call_once, which provided thread-safety. void ihipInit() { -#ifdef COMPILE_TRACE_MARKER +#if COMPILE_TRACE_MARKER amdtInitializeActivityLogger(); amdtScopedMarker("ihipInit", "HIP", NULL); #endif