From b2c23f63d5fbf8684ec035d5519817884884eef2 Mon Sep 17 00:00:00 2001
From: Evgeny <evgeny.shcherbakov@amd.com>
Date: Tue, 15 Oct 2019 11:15:48 -0500
Subject: [PATCH] conditional building of kfd wrapper

---
 inc/roctracer_ext.h                           |  3 ---
 src/CMakeLists.txt                            | 21 ++++++++++++-------
 test/MatrixTranspose/MatrixTranspose.cpp      |  4 ++++
 test/MatrixTranspose_test/MatrixTranspose.cpp | 10 +++++----
 test/run.sh                                   |  2 +-
 5 files changed, 24 insertions(+), 16 deletions(-)

diff --git a/inc/roctracer_ext.h b/inc/roctracer_ext.h
index f91927962b..d555055d0b 100644
--- a/inc/roctracer_ext.h
+++ b/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();
 
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index bf0fe6c06a..4966827737 100644
--- a/src/CMakeLists.txt
+++ b/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
diff --git a/test/MatrixTranspose/MatrixTranspose.cpp b/test/MatrixTranspose/MatrixTranspose.cpp
index 8a5ff54a90..d2ecfb8484 100644
--- a/test/MatrixTranspose/MatrixTranspose.cpp
+++ b/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;
diff --git a/test/MatrixTranspose_test/MatrixTranspose.cpp b/test/MatrixTranspose_test/MatrixTranspose.cpp
index 1a4978bab7..c2c50ebc41 100644
--- a/test/MatrixTranspose_test/MatrixTranspose.cpp
+++ b/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;
 }
diff --git a/test/run.sh b/test/run.sh
index e5170a1416..3b383da201 100755
--- a/test/run.sh
+++ b/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