Add a ROCTX test

Change-Id: I90dab2f349ade6cbfc86f0399c8b9ca905de23b3
2022-05-20 23:10:05 -07:00
commit 1e7783af58
@@ -243,6 +243,7 @@ def gen_events_info(tracefile, trace_level, no_events_cnt, events2ignore, events
  # 1822810364769411:1822810364771941 116477:116477 hsa_agent_get_info(<agent 0x8990e0>, 17, 0x7ffeac015fec) = 0
  # tool_gpu_act_record
  # 3632773658039902:3632773658046462 0:0 hcCommandMarker:273
+  roctx_record = re.compile(r'\d+\s\d+:(\d)+\s(\d):\d+:\".*\"')

  with open(tracefile) as f:
    for line in f:
@@ -262,6 +263,10 @@ def gen_events_info(tracefile, trace_level, no_events_cnt, events2ignore, events
      if tool_record_match:
        event = tool_record_match.group(2)
        tid = int(tool_record_match.group(1))
+      roctx_record_match = roctx_record.match(line)
+      if roctx_record_match:
+        event = roctx_record_match.group(2)
+        tid = int(roctx_record_match.group(1))
      if event == '' or event == '(null)': #some traces has these null events
        continue

@@ -136,9 +136,9 @@ target_include_directories(roctracer
  PUBLIC
    ${HIP_INCLUDE_DIRECTORIES} ${HSA_RUNTIME_INCLUDE_DIRECTORIES}
    $<BUILD_INTERFACE:${CMAKE_CURRENT_BINARY_DIR}>
+    $<BUILD_INTERFACE:${PROJECT_SOURCE_DIR}/inc>
  PRIVATE
-    ${CMAKE_CURRENT_SOURCE_DIR}/roctracer ${CMAKE_CURRENT_SOURCE_DIR}
-    ${PROJECT_SOURCE_DIR}/inc)
+    ${CMAKE_CURRENT_SOURCE_DIR}/roctracer ${CMAKE_CURRENT_SOURCE_DIR})

 target_link_options(roctracer PRIVATE -Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/roctracer/exportmap -Wl,--no-undefined)
 target_link_libraries(roctracer PRIVATE hsa-runtime64::hsa-runtime64 Threads::Threads dl)
@@ -158,7 +158,8 @@ set_target_properties(roctx PROPERTIES
                      SOVERSION ${ROCTRACER_VERSION_MAJOR})

 target_include_directories(roctx
-  PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/roctracer ${CMAKE_CURRENT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/inc)
+  PUBLIC $<BUILD_INTERFACE:${PROJECT_SOURCE_DIR}/inc>
+  PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/roctracer ${CMAKE_CURRENT_SOURCE_DIR})

 target_link_options(roctx PRIVATE -Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/roctx/exportmap -Wl,--no-undefined)

@@ -58,12 +58,14 @@ int ROCTX_API roctxRangePop() {

 roctx_range_id_t ROCTX_API roctxRangeStartA(const char* message) {
  static std::atomic<roctx_range_id_t> start_stop_range_id(1);
+  auto id = start_stop_range_id++;

  roctx_api_data_t api_data{};
  api_data.args.roctxRangeStartA.message = message;
+  api_data.args.roctxRangeStartA.id = id;
  callbacks.Invoke(ROCTX_API_ID_roctxRangeStartA, &api_data);

-  return start_stop_range_id++;
+  return id;
 }

 void ROCTX_API roctxRangeStop(roctx_range_id_t rangeId) {
@@ -105,6 +105,12 @@ target_link_libraries(copy hsa-runtime64::hsa-runtime64 Threads::Threads dl)
 add_dependencies(copy hsaco_targets)
 add_dependencies(mytest copy)

+## Build the ROCTX test
+set_source_files_properties(app/roctx_test.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)
+hip_add_executable(roctx_test EXCLUDE_FROM_ALL app/roctx_test.cpp)
+target_link_libraries(roctx_test Threads::Threads roctx)
+add_dependencies(mytest roctx_test)
+
 ## Copy the golden traces and test scripts
 configure_file(run.sh ${PROJECT_BINARY_DIR} COPYONLY)
 execute_process(COMMAND ${CMAKE_COMMAND} -E create_symlink run.sh ${PROJECT_BINARY_DIR}/run_ci.sh)
@@ -0,0 +1,64 @@
+/* Copyright (c) 2022 Advanced Micro Devices, Inc.
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights
+ to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ copies of the Software, and to permit persons to whom the Software is
+ furnished to do so, subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE. */
+
+#include <hip/hip_runtime.h>
+
+#include "roctx.h"
+
+__global__ void kernel() {}
+
+int main(int argc, char* argv[]) {
+  hipSetDevice(0);
+
+  // Not in a roctx range.
+  kernel<<<1, 1>>>();
+
+  int ret = roctxRangePush("NestedRangeA");
+
+  // In a simple first level roctx range.
+  kernel<<<1, 1>>>();
+
+  if (roctxRangePop() != ret) return -1;
+
+  roctxRangePush("NestedRangeB");
+  roctxRangePush("NestedRangeC");
+  roctx_range_id_t id = roctxRangeStart("StartStopRangeA");
+
+  // In a nested roctx range.
+  kernel<<<1, 1>>>();
+
+  roctxRangePop();
+  roctxRangePop();
+
+  std::thread thread([id]() { roctxRangeStop(id); });
+  thread.join();
+
+  roctxRangePush("NestedRangeD");
+  roctxRangePush("NestedRangeE");
+  roctxRangePop();
+
+  // In a first level roctx range, but after a nested range.
+  kernel<<<1, 1>>>();
+
+  if (roctxRangePop() != 0) return -1;
+
+  hipDeviceSynchronize();
+  return 0;
+}
@@ -0,0 +1,18 @@
+ROCTracer (pid=993231):
+    rocTX-trace()
+0xce5450 agent cpu
+0xd1d520 agent gpu
+0xd1fe80 agent gpu
+628584618590744
+628584859661999 993231:993231 1:0:"NestedRangeA"
+628584859674021 993231:993231 2:0:""
+628584859674693 993231:993231 1:0:"NestedRangeB"
+628584859675344 993231:993231 1:0:"NestedRangeC"
+628584859676115 993231:993231 3:1:"StartStopRangeA"
+628584859678390 993231:993231 2:0:""
+628584859678921 993231:993231 2:0:""
+628584859755545 993231:993233 4:1:""
+628584859819756 993231:993231 1:0:"NestedRangeD"
+628584859820708 993231:993231 1:0:"NestedRangeE"
+628584859821219 993231:993231 2:0:""
+628584859824095 993231:993231 2:0:""
@@ -16,3 +16,4 @@ copy_hsa_input_trace --check-events .*
 hsa_co_trace --check-none
 code_obj_trace --check-none
 trace_buffer_trace --check-none
+roctx_test_trace --check-count .*
@@ -68,6 +68,12 @@ xeval_test() {
 }

 eval_test() {
+  bright=$(tput bold)
+  red=$(tput setaf 1)
+  green=$(tput setaf 2)
+  blue=$(tput setaf 4)
+  normal=$(tput sgr0)
+
  label=$1
  cmdline=$2
  test_name=$3
@@ -98,9 +104,9 @@ eval_test() {
      fi
    fi
    if [ $is_failed = 0 ] ; then
-      echo "$test_name: PASSED"
+      echo "${bright}${blue}$test_name: ${green}PASSED${normal}"
    else
-      echo "$test_name: FAILED"
+      echo "${bright}${blue}$test_name: ${red}FAILED${normal}"
      failed_tests="$failed_tests\n  $test_number: $test_name - \"$label\""
      test_status=$(($test_status + 1))
    fi
@@ -124,6 +130,10 @@ eval_test "standalone HIP MGPU test" "./test/MatrixTranspose_mgpu" MatrixTranspo
 # rocTracer/tool is loaded by HSA runtime
 export HSA_TOOLS_LIB="$ROCTRACER_LIB_PATH/libroctracer64.so $ROCTRACER_TOOL_PATH/libroctracer_tool.so"

+# ROCTX test
+export ROCTRACER_DOMAIN="roctx"
+eval_test "roctx test" ./test/roctx_test roctx_test_trace
+
 # SYS test
 export ROCTRACER_DOMAIN="sys:roctx"
 eval_test "tool SYS test" ./test/MatrixTranspose MatrixTranspose_sys_trace