conditional building of kfd wrapper

inc/roctracer_ext.h

@@ -41,9 +41,6 @@ extern "C" {
 ////////////////////////////////////////////////////////////////////////////////
 // Application annotatin API
 
-// Mark API
-void roctracer_mark(const char* str);
-
 // Tracing start API
 void roctracer_start();
 

src/CMakeLists.txt

@@ -12,16 +12,21 @@ set ( LIB_SRC
 add_library ( ${TARGET_LIB} SHARED ${LIB_SRC} )
 target_include_directories ( ${TARGET_LIB} PRIVATE ${LIB_DIR} ${ROOT_DIR} ${ROOT_DIR}/inc ${HSA_RUNTIME_INC_PATH} ${HSA_RUNTIME_HSA_INC_PATH} ${HIP_INC_DIR} ${HCC_INC_DIR} ${HSA_KMT_INC_PATH} )
 target_link_libraries( ${TARGET_LIB} PRIVATE ${HSA_RUNTIME_LIB} c stdc++ )
+
+# Generating HSA tracing primitives
 execute_process ( COMMAND sh -xc "${ROOT_DIR}/script/hsaap.py ${ROOT_DIR} ${HSA_RUNTIME_INC_PATH}" )
 
-set ( KFD_LIB "kfdwrapper64" )
-set ( KFD_LIB_SRC
-  ${LIB_DIR}/kfd/kfd_wrapper.cpp
-)
-add_library ( ${KFD_LIB} SHARED ${KFD_LIB_SRC} )
-target_include_directories ( ${KFD_LIB} PRIVATE ${LIB_DIR} ${ROOT_DIR} ${ROOT_DIR}/inc ${HSA_RUNTIME_INC_PATH} ${HSA_RUNTIME_HSA_INC_PATH} ${HSA_KMT_INC_PATH} )
-target_link_libraries( ${KFD_LIB} PRIVATE c stdc++ )
-execute_process ( COMMAND sh -xc "${ROOT_DIR}/script/kfdap.py ${ROOT_DIR} ${HSA_KMT_INC_PATH}" )
+# Generating KFD/Thunk tracing primitives
+if ( DEFINED KFD_WRAPPER )
+  set ( KFD_LIB "kfdwrapper64" )
+  set ( KFD_LIB_SRC
+    ${LIB_DIR}/kfd/kfd_wrapper.cpp
+  )
+  add_library ( ${KFD_LIB} SHARED ${KFD_LIB_SRC} )
+  target_include_directories ( ${KFD_LIB} PRIVATE ${LIB_DIR} ${ROOT_DIR} ${ROOT_DIR}/inc ${HSA_RUNTIME_INC_PATH} ${HSA_RUNTIME_HSA_INC_PATH} ${HSA_KMT_INC_PATH} )
+  target_link_libraries( ${KFD_LIB} PRIVATE c stdc++ )
+  execute_process ( COMMAND sh -xc "${ROOT_DIR}/script/kfdap.py ${ROOT_DIR} ${HSA_KMT_INC_PATH}" )
+endif()
 
 set ( ROCTX_LIB "roctx64" )
 set ( ROCTX_LIB_SRC

test/MatrixTranspose/MatrixTranspose.cpp

@@ -36,6 +36,10 @@ THE SOFTWARE.
 #define THREADS_PER_BLOCK_Y 4
 #define THREADS_PER_BLOCK_Z 1
 
+// Mark API
+extern "C"
+void roctracer_mark(const char* str);
+
 // Device (Kernel) function, it must be void
 __global__ void matrixTranspose(float* out, float* in, const int width) {
     int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;

test/MatrixTranspose_test/MatrixTranspose.cpp

@@ -297,9 +297,10 @@ void init_tracing() {
   properties.buffer_callback_fun = activity_callback;
   ROCTRACER_CALL(roctracer_open_pool(&properties));
   // Enable HIP API callbacks
-  ROCTRACER_CALL(roctracer_enable_callback(api_callback, NULL));
+  ROCTRACER_CALL(roctracer_enable_domain_callback(ACTIVITY_DOMAIN_HIP_API, api_callback, NULL));
   // Enable HIP activity tracing
-  ROCTRACER_CALL(roctracer_enable_activity());
+  ROCTRACER_CALL(roctracer_enable_domain_activity(ACTIVITY_DOMAIN_HIP_API));
+  ROCTRACER_CALL(roctracer_enable_domain_activity(ACTIVITY_DOMAIN_HCC_OPS));
 }
 
 // Start tracing routine
@@ -312,8 +313,9 @@ void start_tracing() {
 
 // Stop tracing routine
 void stop_tracing() {
-  ROCTRACER_CALL(roctracer_disable_callback());
-  ROCTRACER_CALL(roctracer_disable_activity());
+  ROCTRACER_CALL(roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API));
+  ROCTRACER_CALL(roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HIP_API));
+  ROCTRACER_CALL(roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HCC_OPS));
   ROCTRACER_CALL(roctracer_flush_activity());
   std::cout << "# STOP  #############################" << std::endl << std::flush;
 }

test/run.sh

@@ -59,7 +59,7 @@ eval_test() {
 
 # Standalone test
 # rocTrecer is used explicitely by test
-eval_test "standalone HIP test" "LD_PRELOAD=libkfdwrapper64.so ./test/MatrixTranspose_test"
+eval_test "standalone HIP test" "./test/MatrixTranspose_test"
 
 # Tool test
 # rocTracer/tool is loaded by HSA runtime