Merge pull request #12 from rkebichi/amd-master

Update roctx.cpp

[ROCm/roctracer commit: 3088073037]

projects/roctracer/src/roctx/roctx.cpp

@@ -28,6 +28,7 @@ THE SOFTWARE.
 #include "inc/ext/prof_protocol.h"
 #include "util/exception.h"
 #include "util/logger.h"
+#include <stack>
 
 #define PUBLIC_API __attribute__((visibility("default")))
 #define CONSTRUCTOR_API __attribute__((constructor))
@@ -61,6 +62,8 @@ THE SOFTWARE.
   (void)err;                                                                                       \
   return X;
 
+static thread_local std::stack<std::string> message_stack;
+
 static inline uint32_t GetPid() { return syscall(__NR_getpid); }
 static inline uint32_t GetTid() { return syscall(__NR_gettid); }
 
@@ -114,14 +117,31 @@ PUBLIC_API void roctxMarkA(const char* message) {
 
 PUBLIC_API int roctxRangePushA(const char* message) {
   API_METHOD_PREFIX
-  EXC_ABORT(ROCTX_STATUS_ERROR, "method is not implemented");
-  API_METHOD_SUFFIX
+  roctx_api_data_t api_data{};
+  api_data.args.roctxRangePushA.message = strdup(message);
+  activity_rtapi_callback_t api_callback_fun = NULL;
+  void* api_callback_arg = NULL;
+  roctx::cb_table.get(ROCTX_API_ID_roctxRangePushA, &api_callback_fun, &api_callback_arg);
+  if (api_callback_fun) api_callback_fun(ACTIVITY_DOMAIN_ROCTX, ROCTX_API_ID_roctxRangePushA, &api_data, api_callback_arg);
+  message_stack.push(strdup(message));
+  API_METHOD_CATCH(-1);
+  return 0;
 }
 
 PUBLIC_API int roctxRangePop() {
   API_METHOD_PREFIX
-  EXC_ABORT(ROCTX_STATUS_ERROR, "method is not implemented");
-  API_METHOD_SUFFIX
+  roctx_api_data_t api_data{};
+  activity_rtapi_callback_t api_callback_fun = NULL;
+  void* api_callback_arg = NULL;
+  roctx::cb_table.get(ROCTX_API_ID_roctxRangePop, &api_callback_fun, &api_callback_arg);
+  if (api_callback_fun) api_callback_fun(ACTIVITY_DOMAIN_ROCTX, ROCTX_API_ID_roctxRangePop, &api_data, api_callback_arg);
+  if (message_stack.empty()) {
+      EXC_ABORT(ROCTX_STATUS_ERROR, "Pop from empty stack!");
+  } else {
+      message_stack.pop();
+  }
+  API_METHOD_CATCH(-1)
+  return 0;
 }
 
 }  // extern "C"

projects/roctracer/test/MatrixTranspose_test/MatrixTranspose.cpp

@@ -28,6 +28,9 @@ THE SOFTWARE.
 // hip header file
 #include <hip/hip_runtime.h>
 
+// roctx header file
+#include <inc/roctx.h>
+
 #ifndef ITERATIONS
 # define ITERATIONS 100
 #endif
@@ -104,20 +107,30 @@ int main() {
 
         // Memory transfer from host to device
         hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice);
-    
+
         // correlation reagion33
         roctracer_activity_push_external_correlation_id(33);
 
+        roctxMarkA("before hipLaunchKernel");
+        roctxRangePushA("hipLaunchKernel");
+
         // Lauching kernel from host
         hipLaunchKernel(matrixTranspose, dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y),
                         dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix,
                         gpuMatrix, WIDTH);
-    
+
+        roctxMarkA("after hipLaunchKernel");
+
         // correlation reagion end
         roctracer_activity_pop_external_correlation_id(NULL);
 
         // Memory transfer from device to host
+        roctxRangePushA("hipMemcpy");
+
         hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost);
+
+        roctxRangePop(); // for "hipMemcpy"
+        roctxRangePop(); // for "hipLaunchKernel"
     
         // correlation reagion end
         roctracer_activity_pop_external_correlation_id();
@@ -165,6 +178,7 @@ int main() {
 #if 1
 #include <inc/roctracer_hip.h>
 #include <inc/roctracer_hcc.h>
+#include <inc/roctracer_roctx.h>
 
 // Macro to check ROC-tracer calls status
 #define ROCTRACER_CALL(call)                                                                       \
@@ -184,6 +198,13 @@ void api_callback(
     void* arg)
 {
   (void)arg;
+
+  if (domain == ACTIVITY_DOMAIN_ROCTX) {
+    const roctx_api_data_t* data = reinterpret_cast<const roctx_api_data_t*>(callback_data);
+    fprintf(stdout, "ROCTX: \"%s\"\n", data->args.message);
+    return;
+  }
+
   const hip_api_data_t* data = reinterpret_cast<const hip_api_data_t*>(callback_data);
   fprintf(stdout, "<%s id(%u)\tcorrelation_id(%lu) %s> ",
     roctracer_op_string(ACTIVITY_DOMAIN_HIP_API, cid, 0),