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.
Bu işleme şunda yer alıyor:
Ben Sander
@@ -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.
|
||||
|
||||
@@ -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 "<<hip-api: hipLaunchKernel '%s' gridDim:[%d.%d.%d] groupDim:[%d.%d.%d] groupMem:+%d stream=%p\n" KNRM, \
|
||||
#_kernelName, lp.gridDim.z, lp.gridDim.y, lp.gridDim.x, lp.groupDim.z, lp.groupDim.y, lp.groupDim.x, lp.groupMemBytes, (void*)(_stream));\
|
||||
}\
|
||||
_kernelName (lp, __VA_ARGS__);\
|
||||
@@ -524,7 +528,7 @@ do {\
|
||||
lp.cf = &cf; \
|
||||
hipStream_t trueStream = (ihipPreLaunchKernel(_stream, &lp.av)); \
|
||||
if (HIP_TRACE_API) {\
|
||||
fprintf(stderr, "hiptrace1: launch '%s' gridDim:[%d.%d.%d] groupDim:[%d.%d.%d] groupMem:+%d stream=%p\n", \
|
||||
fprintf(stderr, "==hip-api: launch '%s' gridDim:[%d.%d.%d] groupDim:[%d.%d.%d] groupMem:+%d stream=%p\n", \
|
||||
#_kernelName, lp.gridDim.z, lp.gridDim.y, lp.gridDim.x, lp.groupDim.z, lp.groupDim.y, lp.groupDim.x, lp.groupMemBytes, (void*)(_stream));\
|
||||
}\
|
||||
_kernelName (lp, __VA_ARGS__);\
|
||||
|
||||
@@ -2,6 +2,14 @@
|
||||
#include <iomanip>
|
||||
#include <string>
|
||||
|
||||
//---
|
||||
// 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 <typename T>
|
||||
@@ -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 <typename T, typename... Args>
|
||||
std::string ToString(T first, Args... args) {
|
||||
return ToString(first) + ", " + ToString(args...) ;
|
||||
}
|
||||
|
||||
|
||||
|
||||
@@ -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
|
||||
|
||||
Yeni konuda referans
Bir kullanıcı engelle