00001 /***************************************************************************
00002 *cr
00003 *cr (C) Copyright 1995-2019 The Board of Trustees of the
00004 *cr University of Illinois
00005 *cr All Rights Reserved
00006 *cr
00007 ***************************************************************************/
00008
00009 /***************************************************************************
00010 * RCS INFORMATION:
00011 *
00012 * $RCSfile: ProfileHooks.h,v $
00013 * $Author: johns $ $Locker: $ $State: Exp $
00014 * $Revision: 1.23 $ $Date: 2020年05月26日 20:38:35 $
00015 *
00016 ***************************************************************************/
00026 #ifndef PROFILEHOOKS_H
00027 #define PROFILEHOOKS_H
00028
00029 #if defined(VMDNVTX)
00030
00031 #if 1
00032
00033 #define VMDUSEGETTID 1
00034 #include <unistd.h>
00035 #include <sys/types.h>
00036 #include <sys/syscall.h>
00037
00038 #ifndef gettid
00039
00040 #define gettid() syscall(SYS_gettid)
00041 #endif
00042 #else
00043
00044 #include <pthread.h>
00045 #endif
00046
00047 #include <cuda_runtime.h>
00048 #include <cuda_profiler_api.h>
00049
00050 #if CUDART_VERSION >= 10000
00051 #include <nvtx3/nvToolsExt.h> // CUDA >= 10 has NVTX V3+
00052 #else
00053 #error NVTXv3 requires CUDA 10.0 or greater
00054 //#include <nvToolsExt.h> // CUDA < 10 has NVTX V2
00055 #endif
00056
00057
00060 const uint32_t VMD_nvtx_colors[] = {
00061 0xff00ff00, // 0 green
00062 0xff0000ff, // 1 blue
00063 0xffffff00, // 2 yellow
00064 0xffff00ff, // 3 purple
00065 0xff00ffff, // 4 teal
00066 0xffff0000, // 5 red
00067 0xffffffff, // 6 white
00068 };
00069 const int VMD_nvtx_colors_len = sizeof(VMD_nvtx_colors)/sizeof(uint32_t);
00070
00071
00072 #define PROFILE_INITIALIZE() do { nvtxInitialize(NULL); } while(0) // terminate with semicolon
00073
00074 #define PROFILE_START() \
00075 do { \
00076 cudaProfilerStart(); \
00077 } while (0) // terminate with semicolon
00078
00079 #define PROFILE_STOP() \
00080 do { \
00081 cudaDeviceSynchronize(); \
00082 cudaProfilerStop(); \
00083 } while (0) // terminate with semicolon
00084
00085
00094 #if defined(VMDUSEGETTID)
00095
00096 // On Linux use gettid() to get current thread ID
00097 #define PROFILE_MAIN_THREAD() \
00098 do { \
00099 nvtxNameOsThread(gettid(), "Main VMD thread"); \
00100 } while (0) // terminate with semicolon
00101
00102 #define PROFILE_NAME_THREAD(name) \
00103 do { \
00104 nvtxNameOsThread(gettid(), name); \
00105 } while (0) // terminate with semicolon
00106
00107 #else
00108
00109 // Nn MacOS X or other platforms use pthread_threadid_np()
00110 #define PROFILE_MAIN_THREAD() \
00111 do { \
00112 __uint64_t tid;
00113 pthread_threadid_np(pthread_self(), &tid);
00114 nvtxNameOsThread(tid, "Main VMD thread"); \
00115 } while (0) // terminate with semicolon
00116
00117 #define PROFILE_NAME_THREAD(name) \
00118 do { \
00119 __uint64_t tid;
00120 pthread_threadid_np(pthread_self(), &tid);
00121 nvtxNameOsThread(gettid(), name); \
00122 } while (0) // terminate with semicolon
00123
00124 #endif
00125
00126
00127 #define PROFILE_MARK(name,cid) \
00128 do { \
00129 /* create an ASCII event marker */ \
00130 /* nvtxMarkA(name); */ \
00131 int color_id = cid; \
00132 color_id = color_id % VMD_nvtx_colors_len; \
00133 nvtxEventAttributes_t eventAttrib = {0}; \
00134 eventAttrib.version = NVTX_VERSION; \
00135 eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE; \
00136 eventAttrib.colorType = NVTX_COLOR_ARGB; \
00137 eventAttrib.color = VMD_nvtx_colors[color_id]; \
00138 eventAttrib.messageType = NVTX_MESSAGE_TYPE_ASCII; \
00139 eventAttrib.message.ascii = name; \
00140 nvtxMarkEx(&eventAttrib); \
00141 } while(0) // terminate with semicolon
00142
00143 // start recording an event
00144 #define PROFILE_PUSH_RANGE(name,cid) \
00145 do { \
00146 int color_id = cid; \
00147 color_id = color_id % VMD_nvtx_colors_len; \
00148 nvtxEventAttributes_t eventAttrib = {0}; \
00149 eventAttrib.version = NVTX_VERSION; \
00150 eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE; \
00151 eventAttrib.colorType = NVTX_COLOR_ARGB; \
00152 eventAttrib.color = VMD_nvtx_colors[color_id]; \
00153 eventAttrib.messageType = NVTX_MESSAGE_TYPE_ASCII; \
00154 eventAttrib.message.ascii = name; \
00155 nvtxRangePushEx(&eventAttrib); \
00156 } while(0) // must terminate with semi-colon
00157
00158 // stop recording an event
00159 #define PROFILE_POP_RANGE(empty) \
00160 do { \
00161 nvtxRangePop(); \
00162 } while(0) // terminate with semicolon
00163
00164 // embed event recording in class to automatically pop when destroyed
00165 class VMD_NVTX_Tracer {
00166 public:
00167 VMD_NVTX_Tracer(const char *name, int cid = 0) { PROFILE_PUSH_RANGE(name, cid); }
00168 ~VMD_NVTX_Tracer() { PROFILE_POP_RANGE(); }
00169 };
00170
00171 // include cid as part of the name
00172 // call RANGE at beginning of function to push event recording
00173 // destructor is automatically called on return to pop event recording
00174 #define PROFILE_RANGE(name,cid) \
00175 VMD_NVTX_Tracer vmd_nvtx_tracer##cid(name,cid)
00176 // must terminate with semi-colon
00177
00178
00179
00180 #else
00181
00182 //
00183 // Otherwise the NVTX profiling macros become no-ops.
00184 //
00185 #define PROFILE_INITIALIZE() do { } while(0) // terminate with semicolon
00186 #define PROFILE_START() do { } while(0) // terminate with semicolon
00187 #define PROFILE_STOP() do { } while(0) // terminate with semicolon
00188 #define PROFILE_MAIN_THREAD() do { } while(0) // terminate with semicolon
00189 #define PROFILE_NAME_THREAD(name) do { } while(0) // terminate with semicolon
00190 #define PROFILE_MARK(name,cid) do { } while(0) // terminate with semicolon
00191 #define PROFILE_PUSH_RANGE(name,cid) do { } while(0) // terminate with semicolon
00192 #define PROFILE_POP_RANGE() do { } while(0) // terminate with semicolon
00193 #define PROFILE_RANGE(namd,cid) do { } while(0) // terminate with semicolon
00194
00195 #endif
00196
00197 #endif