ProfileHooks.h

Go to the documentation of this file.

/*
 * TachyonOptiX.cu - OptiX host-side RT engine implementation
 *
 * (C) Copyright 2013-2022 John E. Stone
 * SPDX-License-Identifier: BSD-3-Clause
 *
 * $Id: ProfileHooks.h,v 1.3 2022/03/11 00:45:37 johns Exp $
 *
 */

#ifndef PROFILEHOOKS_H
#define PROFILEHOOKS_H

#if defined(WKFNVTX)

#if 1
#define WKFUSEGETTID 1
#include <unistd.h>
#include <sys/types.h>
#include <sys/syscall.h>

#ifndef gettid
#define gettid() syscall(SYS_gettid)
#endif
#else 
#include <pthread.h>
#endif

#include <cuda_runtime.h>
#include <cuda_profiler_api.h>

// We really only want to use CUDA >= 10.0 w/ NVTX V3+,
// which bypasses the need to include/ship/link against 
// any additional libs.
#if CUDART_VERSION >= 10000
#include <nvtx3/nvToolsExt.h>  // CUDA >= 10 has NVTX V3+
#else
#error NVTXv3 requires CUDA 10.0 or greater
#include <nvToolsExt.h>        // CUDA < 10 has NVTX V2
#endif


const uint32_t WKF_nvtx_colors[] = {
  0xff629f57, 
  0xffee8c40, 
  0xff507ba4, 
  0xffecc65c, 
  0xffac7c9f, 
  0xff7cb7b2, 
  0xffdb565c, 
  0xffb9b0ac, 
  0xffffffff, 
};

const int WKF_nvtx_colors_len = sizeof(WKF_nvtx_colors)/sizeof(uint32_t);

#define PROFILE_GREEN  0
#define PROFILE_ORANGE 1
#define PROFILE_BLUE   2
#define PROFILE_YELLOW 3
#define PROFILE_PURPLE 4
#define PROFILE_TEAL   5
#define PROFILE_RED    6
#define PROFILE_GRAY   7
#define PROFILE_WHITE  8

#define PROFILE_INITIALIZE() do { nvtxInitialize(NULL); } while(0) // terminate with semicolon

#define PROFILE_START() \
  do { \
    cudaProfilerStart(); \
  } while (0) // terminate with semicolon

#define PROFILE_STOP() \
  do { \
    cudaDeviceSynchronize(); \
    cudaProfilerStop(); \
  } while (0) // terminate with semicolon


#if defined(WKFUSEGETTID)

#define PROFILE_MAIN_THREAD() \
  do { \
    /* On Linux use gettid() to get current thread ID */ \
    nvtxNameOsThread(gettid(), "Main thread"); \
  } while (0) // terminate with semicolon

#define PROFILE_NAME_THREAD(name) \
  do { \
    nvtxNameOsThread(gettid(), name); \
  } while (0) // terminate with semicolon

#else

#define PROFILE_MAIN_THREAD() \
  do { \
    /* On MacOS X or other platforms use pthread_threadid_np() */ \
    __uint64_t tid;
    pthread_threadid_np(pthread_self(), &tid);
    nvtxNameOsThread(tid, "Main thread"); \
  } while (0) // terminate with semicolon

#define PROFILE_NAME_THREAD(name) \
  do { \
    __uint64_t tid;
    pthread_threadid_np(pthread_self(), &tid);
    nvtxNameOsThread(gettid(), name); \
  } while (0) // terminate with semicolon

#endif


#define PROFILE_MARK(name,cid) \
  do { \
    /* create an ASCII event marker */ \
    /* nvtxMarkA(name); */ \
    int color_id = cid; \
    color_id = color_id % WKF_nvtx_colors_len; \
    nvtxEventAttributes_t eventAttrib = {0}; \
    eventAttrib.version = NVTX_VERSION; \
    eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE; \
    eventAttrib.colorType = NVTX_COLOR_ARGB; \
    eventAttrib.color = WKF_nvtx_colors[color_id]; \
    eventAttrib.messageType = NVTX_MESSAGE_TYPE_ASCII; \
    eventAttrib.message.ascii = name; \
    nvtxMarkEx(&eventAttrib); \
  } while(0) // terminate with semicolon


#define PROFILE_PUSH_RANGE(name,cid) \
  do { \
    int color_id = cid; \
    color_id = color_id % WKF_nvtx_colors_len; \
    nvtxEventAttributes_t eventAttrib = {0}; \
    eventAttrib.version = NVTX_VERSION; \
    eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE; \
    eventAttrib.colorType = NVTX_COLOR_ARGB; \
    eventAttrib.color = WKF_nvtx_colors[color_id]; \
    eventAttrib.messageType = NVTX_MESSAGE_TYPE_ASCII; \
    eventAttrib.message.ascii = name; \
    nvtxRangePushEx(&eventAttrib); \
  } while(0)  // must terminate with semi-colon


#define PROFILE_POP_RANGE(empty) \
  do { \
    nvtxRangePop(); \
  } while(0) // terminate with semicolon 


// embed event recording in class to automatically pop when destroyed
class WKF_NVTX_Tracer {
  public:
    WKF_NVTX_Tracer(const char *name, int cid = 0) { PROFILE_PUSH_RANGE(name, cid); }
    ~WKF_NVTX_Tracer() { PROFILE_POP_RANGE(); }
};

#define PROFILE_RANGE(name,cid) \
  /* include cid as part of the name */ \
  /* call RANGE at beginning of function to push event recording */ \
  /* destructor is automatically called on return to pop event recording */ \
  WKF_NVTX_Tracer wkf_nvtx_tracer##cid(name,cid)
  // must terminate with semi-colon

#if defined(WKFNVTX_SYNCPRETTY)
#define PROFILE_STREAM_SYNC_PRETTY(stream) \
  do { \
    /* Add a CUDA stream sync call, but only for the benefit of */ \
    /* profile trace clarity, so that it can be disabled on demand */ \
    cudaStreamSynchronize(stream); \
  } while(0) // terminate with semicolon

#else
#define PROFILE_STREAM_SYNC_PRETTY(stream) do { } while(0) // term w/ semicolon 
#endif


#else

//
// If NVTX isn't enabled, then the profiling macros become no-ops.
// We add inline documentation here since Doxygen sees this branch by default.
//

#define PROFILE_INITIALIZE()         do { } while(0) // terminate with semicolon

#define PROFILE_START()              do { } while(0) // terminate with semicolon

#define PROFILE_STOP()               do { } while(0) // terminate with semicolon

#define PROFILE_MAIN_THREAD()        do { } while(0) // terminate with semicolon

#define PROFILE_NAME_THREAD(name)    do { } while(0) // terminate with semicolon

#define PROFILE_MARK(name,cid)       do { } while(0) // terminate with semicolon

#define PROFILE_PUSH_RANGE(name,cid) do { } while(0) // terminate with semicolon

#define PROFILE_POP_RANGE()          do { } while(0) // terminate with semicolon

#define PROFILE_RANGE(name,cid)      do { } while(0) // terminate with semicolon

#define PROFILE_STREAM_SYNC_PRETTY(stream) do { } while(0) // term w/ semicolon 
#endif

#endif