diff --git a/script/check_trace.py b/script/check_trace.py index 025a690339..28772e3f25 100644 --- a/script/check_trace.py +++ b/script/check_trace.py @@ -243,6 +243,7 @@ def gen_events_info(tracefile, trace_level, no_events_cnt, events2ignore, events # 1822810364769411:1822810364771941 116477:116477 hsa_agent_get_info(, 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 diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 404df0b125..26f6c26888 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -136,9 +136,9 @@ target_include_directories(roctracer PUBLIC ${HIP_INCLUDE_DIRECTORIES} ${HSA_RUNTIME_INCLUDE_DIRECTORIES} $ + $ 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 $ + 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) diff --git a/src/roctx/roctx.cpp b/src/roctx/roctx.cpp index 988c2a55b9..bb27b5850a 100644 --- a/src/roctx/roctx.cpp +++ b/src/roctx/roctx.cpp @@ -58,12 +58,14 @@ int ROCTX_API roctxRangePop() { roctx_range_id_t ROCTX_API roctxRangeStartA(const char* message) { static std::atomic 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) { diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index daf4a97954..43915f3dbf 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -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) diff --git a/test/app/roctx_test.cpp b/test/app/roctx_test.cpp new file mode 100644 index 0000000000..ef863bdafc --- /dev/null +++ b/test/app/roctx_test.cpp @@ -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 + +#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; +} diff --git a/test/golden_traces/roctx_test_trace.txt b/test/golden_traces/roctx_test_trace.txt new file mode 100644 index 0000000000..87668cabe6 --- /dev/null +++ b/test/golden_traces/roctx_test_trace.txt @@ -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:"" diff --git a/test/golden_traces/tests_trace_cmp_levels.txt b/test/golden_traces/tests_trace_cmp_levels.txt index 111bfa5a5d..eab5b72648 100644 --- a/test/golden_traces/tests_trace_cmp_levels.txt +++ b/test/golden_traces/tests_trace_cmp_levels.txt @@ -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 .* diff --git a/test/run.sh b/test/run.sh index 059638ba90..5d1e26dda5 100755 --- a/test/run.sh +++ b/test/run.sh @@ -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