Merge branch 'amd-master' into rkebichi-patch-2
This commit is contained in:
eshcherb
GitHub
@@ -8,6 +8,7 @@ b
|
||||
build
|
||||
inc/hsa_prof_str.h
|
||||
inc/kfd_prof_str.h
|
||||
inc/kfd_ostream_ops.h
|
||||
test/hsa
|
||||
test/MatrixTranspose/MatrixTranspose
|
||||
test/MatrixTranspose_test/MatrixTranspose
|
||||
|
||||
+33
-9
@@ -28,6 +28,9 @@ cmake_minimum_required ( VERSION 3.5.0 )
|
||||
## Verbose output.
|
||||
set ( CMAKE_VERBOSE_MAKEFILE TRUE CACHE BOOL "Verbose Output" FORCE )
|
||||
|
||||
# Install prefix
|
||||
set(CMAKE_INSTALL_PREFIX "/opt/rocm" CACHE PATH "Install path prefix, prepended onto install directories")
|
||||
|
||||
## Set module name and project name.
|
||||
set ( ROCTRACER_NAME "roctracer" )
|
||||
set ( ROCTRACER_TARGET "${ROCTRACER_NAME}64" )
|
||||
@@ -43,17 +46,25 @@ include ( env )
|
||||
|
||||
## Setup the package version.
|
||||
get_version ( "1.0.0" )
|
||||
message ( "-- LIB-VERSION: ${VERSION_MAJOR}.${VERSION_MINOR}.${VERSION_PATCH}" )
|
||||
|
||||
set ( BUILD_VERSION_MAJOR ${VERSION_MAJOR} )
|
||||
set ( BUILD_VERSION_MINOR ${VERSION_MINOR} )
|
||||
set ( BUILD_VERSION_PATCH ${VERSION_PATCH} )
|
||||
set ( LIB_VERSION_STRING "${BUILD_VERSION_MAJOR}.${BUILD_VERSION_MINOR}.${BUILD_VERSION_PATCH}" )
|
||||
|
||||
if ( DEFINED VERSION_BUILD AND NOT ${VERSION_BUILD} STREQUAL "" )
|
||||
message ( "VERSION BUILD DEFINED ${VERSION_BUILD}" )
|
||||
set ( BUILD_VERSION_PATCH "${BUILD_VERSION_PATCH}-${VERSION_BUILD}" )
|
||||
endif ()
|
||||
set ( BUILD_VERSION_STRING "${BUILD_VERSION_MAJOR}.${BUILD_VERSION_MINOR}.${BUILD_VERSION_PATCH}" )
|
||||
|
||||
set ( LIB_VERSION_MAJOR ${BUILD_VERSION_MAJOR} )
|
||||
set ( LIB_VERSION_MINOR ${BUILD_VERSION_MINOR} )
|
||||
if (DEFINED ENV{ROCM_LIBPATCH_VERSION})
|
||||
set (LIB_VERSION_PATCH $ENV{ROCM_LIBPATCH_VERSION} )
|
||||
else ()
|
||||
set (LIB_VERSION_PATCH ${BUILD_VERSION_PATCH} )
|
||||
endif()
|
||||
set ( LIB_VERSION_STRING "${LIB_VERSION_MAJOR}.${LIB_VERSION_MINOR}.${LIB_VERSION_PATCH}" )
|
||||
message ( "-- LIB-VERSION: ${LIB_VERSION_MAJOR}.${LIB_VERSION_MINOR}.${LIB_VERSION_PATCH}" )
|
||||
|
||||
## Set target and root/lib/test directory
|
||||
set ( TARGET_NAME "${ROCTRACER_TARGET}" )
|
||||
@@ -66,7 +77,7 @@ include ( ${LIB_DIR}/CMakeLists.txt )
|
||||
|
||||
## Set the VERSION and SOVERSION values
|
||||
set_property ( TARGET ${TARGET_NAME} PROPERTY VERSION "${LIB_VERSION_STRING}" )
|
||||
set_property ( TARGET ${TARGET_NAME} PROPERTY SOVERSION "${BUILD_VERSION_MAJOR}" )
|
||||
set_property ( TARGET ${TARGET_NAME} PROPERTY SOVERSION "${LIB_VERSION_MAJOR}" )
|
||||
|
||||
# If the library is a release, strip the target library
|
||||
if ( "${CMAKE_BUILD_TYPE}" STREQUAL release )
|
||||
@@ -84,6 +95,10 @@ add_custom_target ( inc-link ALL WORKING_DIRECTORY ${PROJECT_BINARY_DIR}
|
||||
COMMAND ${CMAKE_COMMAND} -E create_symlink ../${ROCTRACER_NAME}/include inc-link )
|
||||
add_custom_target ( so-link ALL WORKING_DIRECTORY ${PROJECT_BINARY_DIR}
|
||||
COMMAND ${CMAKE_COMMAND} -E create_symlink ../${ROCTRACER_NAME}/lib/${ROCTRACER_LIBRARY}.so so-link )
|
||||
add_custom_target ( so-major-link ALL WORKING_DIRECTORY ${PROJECT_BINARY_DIR}
|
||||
COMMAND ${CMAKE_COMMAND} -E create_symlink ../${ROCTRACER_NAME}/lib/${ROCTRACER_LIBRARY}.so.${LIB_VERSION_MAJOR} so-major-link )
|
||||
add_custom_target ( so-patch-link ALL WORKING_DIRECTORY ${PROJECT_BINARY_DIR}
|
||||
COMMAND ${CMAKE_COMMAND} -E create_symlink ../${ROCTRACER_NAME}/lib/${ROCTRACER_LIBRARY}.so.${LIB_VERSION_STRING} so-patch-link )
|
||||
|
||||
## Install information
|
||||
install ( TARGETS ${ROCTRACER_TARGET} LIBRARY DESTINATION lib )
|
||||
@@ -95,24 +110,33 @@ install ( FILES ${CMAKE_CURRENT_SOURCE_DIR}/inc/ext/prof_protocol.h DESTINATION
|
||||
install ( FILES ${CMAKE_CURRENT_SOURCE_DIR}/inc/ext/hsa_rt_utils.hpp DESTINATION include/ext )
|
||||
install ( FILES ${PROJECT_BINARY_DIR}/inc-link DESTINATION ../include RENAME ${ROCTRACER_NAME} )
|
||||
install ( FILES ${PROJECT_BINARY_DIR}/so-link DESTINATION ../lib RENAME ${ROCTRACER_LIBRARY}.so )
|
||||
install ( FILES ${PROJECT_BINARY_DIR}/so-major-link DESTINATION ../lib RENAME ${ROCTRACER_LIBRARY}.so.${LIB_VERSION_MAJOR} )
|
||||
install ( FILES ${PROJECT_BINARY_DIR}/so-patch-link DESTINATION ../lib RENAME ${ROCTRACER_LIBRARY}.so.${LIB_VERSION_STRING} )
|
||||
install ( FILES ${PROJECT_BINARY_DIR}/test/libtracer_tool.so DESTINATION tool )
|
||||
|
||||
## rocTX
|
||||
set ( ROCTX_TARGET "roctx64" )
|
||||
set ( ROCTX_LIBRARY "lib${ROCTX_TARGET}" )
|
||||
|
||||
## Set the VERSION and SOVERSION values
|
||||
set_property ( TARGET ${ROCTX_TARGET} PROPERTY VERSION "${LIB_VERSION_STRING}" )
|
||||
set_property ( TARGET ${ROCTX_TARGET} PROPERTY SOVERSION "${LIB_VERSION_MAJOR}" )
|
||||
|
||||
add_custom_target ( so-roctx-link ALL WORKING_DIRECTORY ${PROJECT_BINARY_DIR}
|
||||
COMMAND ${CMAKE_COMMAND} -E create_symlink ../${ROCTRACER_NAME}/lib/${ROCTX_LIBRARY}.so so-roctx-link )
|
||||
add_custom_target ( so-roctx-major-link ALL WORKING_DIRECTORY ${PROJECT_BINARY_DIR}
|
||||
COMMAND ${CMAKE_COMMAND} -E create_symlink ../${ROCTRACER_NAME}/lib/${ROCTX_LIBRARY}.so.${LIB_VERSION_MAJOR} so-roctx-major-link )
|
||||
add_custom_target ( so-roctx-patch-link ALL WORKING_DIRECTORY ${PROJECT_BINARY_DIR}
|
||||
COMMAND ${CMAKE_COMMAND} -E create_symlink ../${ROCTRACER_NAME}/lib/${ROCTX_LIBRARY}.so.${LIB_VERSION_STRING} so-roctx-patch-link )
|
||||
|
||||
install ( TARGETS "roctx64" LIBRARY DESTINATION lib )
|
||||
install ( FILES ${CMAKE_CURRENT_SOURCE_DIR}/inc/roctx.h DESTINATION include )
|
||||
install ( FILES ${CMAKE_CURRENT_SOURCE_DIR}/inc/roctracer_roctx.h DESTINATION include )
|
||||
install ( FILES ${PROJECT_BINARY_DIR}/so-roctx-link DESTINATION ../lib RENAME ${ROCTX_LIBRARY}.so )
|
||||
install ( FILES ${PROJECT_BINARY_DIR}/so-roctx-major-link DESTINATION ../lib RENAME ${ROCTX_LIBRARY}.so.${LIB_VERSION_MAJOR} )
|
||||
install ( FILES ${PROJECT_BINARY_DIR}/so-roctx-patch-link DESTINATION ../lib RENAME ${ROCTX_LIBRARY}.so.${LIB_VERSION_STRING} )
|
||||
|
||||
## KFD wrapper
|
||||
if ( DEFINED KFD_WRAPPER )
|
||||
install ( TARGETS "kfdwrapper64" LIBRARY DESTINATION lib )
|
||||
endif ()
|
||||
install ( TARGETS "kfdwrapper64" LIBRARY DESTINATION lib )
|
||||
|
||||
## Packaging directives
|
||||
set ( CPACK_GENERATOR "DEB" "RPM" "TGZ" )
|
||||
@@ -129,7 +153,7 @@ set ( CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE" )
|
||||
set ( CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "${CMAKE_CURRENT_SOURCE_DIR}/DEBIAN/postinst;${CMAKE_CURRENT_SOURCE_DIR}/DEBIAN/prerm" )
|
||||
|
||||
## RPM package specific variables
|
||||
set ( CPACK_RPM_PRE_INSTALL_SCRIPT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/RPM/rpm_post" )
|
||||
set ( CPACK_RPM_POST_INSTALL_SCRIPT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/RPM/rpm_post" )
|
||||
set ( CPACK_RPM_POST_UNINSTALL_SCRIPT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/RPM/rpm_postun" )
|
||||
|
||||
include ( CPack )
|
||||
|
||||
+4
-1
@@ -3,7 +3,10 @@
|
||||
set -e
|
||||
|
||||
do_ldconfig() {
|
||||
echo /opt/rocm/roctracer/lib > /etc/ld.so.conf.d/libroctracer64.conf && ldconfig
|
||||
INSTALL_PATH=/opt/rocm/roctracer
|
||||
if [ -e "${INSTALL_PATH}" ] ; then
|
||||
echo /opt/rocm/roctracer/lib > /etc/ld.so.conf.d/libroctracer64.conf && ldconfig
|
||||
fi
|
||||
}
|
||||
|
||||
case "$1" in
|
||||
|
||||
@@ -1,4 +1,5 @@
|
||||
Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
|
||||
[MITx11 License]
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
|
||||
@@ -1,52 +1,75 @@
|
||||
# ROC-tracer
|
||||
```
|
||||
ROC-tracer library, Runtimes Generic Callback/Activity APIs.
|
||||
ROC-tracer library: Runtimes Generic Callback/Activity APIs.
|
||||
The goal of the implementation is to provide a generic independent from
|
||||
specific runtime profiler to trace API and asyncronous activity.
|
||||
|
||||
The API provides functionality for registering the runtimes API callbacks and
|
||||
asyncronous activity records pool support.
|
||||
|
||||
ROC-TX librray: code annotation evemts API
|
||||
Includes basic API: roctxMark, roctxRangePush, roctxRangePop.
|
||||
```
|
||||
|
||||
## Usage
|
||||
```
|
||||
rocTracer API:
|
||||
To use the rocTracer API you need the API header and to link your application with roctracer .so librray:
|
||||
- the API header: /opt/rocm/roctracer/include/roctracer.h
|
||||
- the .so library: /opt/rocm/lib/libroctracer64.so
|
||||
|
||||
rocTX API:
|
||||
To use the rocTX API you need the API header and to link your application with roctx .so librray:
|
||||
- the API header: /opt/rocm/roctracer/include/roctx.h
|
||||
- the .so library: /opt/rocm/lib/libroctx64.so
|
||||
```
|
||||
|
||||
## The library source tree
|
||||
```
|
||||
- inc/roctracer.h - Library public API
|
||||
- doc - documentation
|
||||
- inc/roctracer.h - rocTacer library public API header
|
||||
- inc/roctx.h - rocTX library puiblic API header
|
||||
- src - Library sources
|
||||
- core - Library API sources
|
||||
- util - Library utils sources
|
||||
- core - rocTracer library API sources
|
||||
- roctx - rocTX library API sources
|
||||
- util - library utils sources
|
||||
- test - test suit
|
||||
- MatrixTranspose - test based on HIP MatrixTranspose sample
|
||||
```
|
||||
|
||||
## Documentation
|
||||
```
|
||||
- API description: inc/roctracer.h
|
||||
- Code example: test/MatrixTranspose_test/MatrixTranspose.cpp
|
||||
```
|
||||
- API description:
|
||||
- ['roctracer' / 'rocTX' profiling C API specification](doc/roctracer_spec.md)
|
||||
- Code examples:
|
||||
- [test/MatrixTranspose_test/MatrixTranspose.cpp](test/MatrixTranspose_test/MatrixTranspose.cpp)
|
||||
- [test/MatrixTranspose/MatrixTranspose.cpp](test/MatrixTranspose/MatrixTranspose.cpp)
|
||||
|
||||
## To build and run test
|
||||
```
|
||||
- ROCm-2.3 or higher is required
|
||||
|
||||
cd <your path>
|
||||
- ROCm is required
|
||||
|
||||
- Python modules requirements: CppHeaderParser, argparse.
|
||||
To install:
|
||||
sudo pip install CppHeaderParser argparse
|
||||
|
||||
- CLone development branch of roctracer:
|
||||
git clone -b amd-master https://github.com/ROCm-Developer-Tools/roctracer
|
||||
|
||||
- Set environment:
|
||||
export CMAKE_PREFIX_PATH=/opt/rocm
|
||||
- To use custom HIP/HCC versions:
|
||||
export HIP_PATH=/opt/rocm/hip
|
||||
export HCC_HOME=/opt/rocm/hcc
|
||||
export CMAKE_PREFIX_PATH=/opt/rocm
|
||||
|
||||
- Build ROCtracer
|
||||
- To build roctracer library:
|
||||
export CMAKE_BUILD_TYPE=<debug|release> # release by default
|
||||
cd <your path>/roctracer && mkdir build && cd build && cmake -DCMAKE_INSTALL_PREFIX=/opt/rocm .. && make -j <nproc>
|
||||
|
||||
- To build and run test
|
||||
- To build and run test:
|
||||
make mytest
|
||||
run.sh
|
||||
|
||||
- To install
|
||||
- To install:
|
||||
make install
|
||||
or
|
||||
make package && dpkg -i *.deb
|
||||
|
||||
+4
-1
@@ -1 +1,4 @@
|
||||
echo /opt/rocm/roctracer/lib > /etc/ld.so.conf.d/libroctracer64.conf && ldconfig
|
||||
INSTALL_PATH=/opt/rocm/roctracer
|
||||
if [ -e "${INSTALL_PATH}" ] ; then
|
||||
echo /opt/rocm/roctracer/lib > /etc/ld.so.conf.d/libroctracer64.conf && ldconfig
|
||||
fi
|
||||
|
||||
@@ -1,7 +1,8 @@
|
||||
#!/bin/bash -x
|
||||
SRC_DIR=`dirname $0`
|
||||
COMPONENT="roctracer"
|
||||
ROCM_PATH="/opt/rocm"
|
||||
ROCM_PATH="${ROCM_PATH:=/opt/rocm}"
|
||||
LD_RUNPATH_FLAG=" -Wl,--enable-new-dtags -Wl,--rpath,$ROCM_PATH/lib:$ROCM_PATH/lib64"
|
||||
|
||||
fatal() {
|
||||
echo "$1"
|
||||
@@ -19,9 +20,9 @@ if [ -z "$PACKAGE_ROOT" ] ; then PACKAGE_ROOT=$ROCM_PATH; fi
|
||||
if [ -z "$PACKAGE_PREFIX" ] ; then PACKAGE_PREFIX="$ROCM_PATH/$COMPONENT"; fi
|
||||
if [ -z "$PREFIX_PATH" ] ; then PREFIX_PATH=$PACKAGE_ROOT; fi
|
||||
if [ -n "$HIP_VDI" ] ; then HIP_VDI_OPT="-DHIP_VDI=1"; fi
|
||||
if [ -n "$ROCM_RPATH" ] ; then LD_RUNPATH_FLAG=" -Wl,--enable-new-dtags -Wl,--rpath,${ROCM_RPATH}"; fi
|
||||
|
||||
ROCTRACER_ROOT=$(cd $ROCTRACER_ROOT && echo $PWD)
|
||||
MAKE_OPTS="-j 8 -C $BUILD_DIR"
|
||||
|
||||
mkdir -p $BUILD_DIR
|
||||
pushd $BUILD_DIR
|
||||
@@ -33,6 +34,7 @@ cmake \
|
||||
-DCMAKE_INSTALL_PREFIX=$PACKAGE_ROOT \
|
||||
-DCPACK_PACKAGING_INSTALL_PREFIX=$PACKAGE_PREFIX \
|
||||
-DCPACK_GENERATOR="DEB;RPM" \
|
||||
-DCMAKE_SHARED_LINKER_FLAGS="$LD_RUNPATH_FLAG" \
|
||||
$HIP_VDI_OPT \
|
||||
$ROCTRACER_ROOT
|
||||
make
|
||||
|
||||
@@ -50,11 +50,6 @@ else()
|
||||
set ( HIP_VDI 0 )
|
||||
endif()
|
||||
|
||||
## Enable KFD wrapper
|
||||
if ( DEFINED KFD_WRAPPER )
|
||||
add_definitions ( -DKFD_WRAPPER=${KFD_WRAPPER} )
|
||||
endif()
|
||||
|
||||
## Enable HIP/HCC local build
|
||||
if ( DEFINED LOCAL_BUILD )
|
||||
add_definitions ( -DLOCAL_BUILD=${LOCAL_BUILD} )
|
||||
@@ -123,6 +118,7 @@ endif ()
|
||||
find_library ( HSA_KMT_LIB "libhsakmt.so" )
|
||||
get_filename_component ( HSA_KMT_LIB_PATH ${HSA_KMT_LIB} DIRECTORY )
|
||||
set ( HSA_KMT_INC_PATH "${HSA_KMT_LIB_PATH}/../include" )
|
||||
set ( ROCM_INC_PATH ${HSA_KMT_INC_PATH} )
|
||||
|
||||
## Basic Tool Chain Information
|
||||
message ( "----------------NBIT: ${NBIT}" )
|
||||
@@ -136,6 +132,5 @@ message ( "-------------HCC-Inc: ${HCC_INC_DIR}" )
|
||||
message ( "-------------HIP-Inc: ${HIP_INC_DIR}" )
|
||||
message ( "-------------KFD-Inc: ${HSA_KMT_INC_PATH}" )
|
||||
message ( "-------------HIP-VDI: ${HIP_VDI}" )
|
||||
message ( "---------KFD_WRAPPER: ${KFD_WRAPPER}" )
|
||||
message ( "-----CMAKE_CXX_FLAGS: ${CMAKE_CXX_FLAGS}" )
|
||||
message ( "---CMAKE_PREFIX_PATH: ${CMAKE_PREFIX_PATH}" )
|
||||
|
||||
@@ -0,0 +1,748 @@
|
||||
# ROC Tracer / ROC-TX Libraries Specification
|
||||
```
|
||||
ROC Tracer API version 2
|
||||
ROC-TX API version 1
|
||||
|
||||
- The rocTracer API is agnostic to specific runtime and may trace
|
||||
the runtime API calls and asynchronous GPU activity.
|
||||
- The rocTX API is provided for application code annotation.
|
||||
```
|
||||
## 1. High level overview
|
||||
```
|
||||
The goal of the implementation is to provide a runtime independent API
|
||||
for tracing of runtime calls and asynchronous activity, like GPU kernel
|
||||
dispatches and memory moves. The tracing includes callback API for
|
||||
runtime API tracing and activity API for asynchronous activity records
|
||||
logging.
|
||||
|
||||
Depending on particular runtime intercepting mechanism, the rocTracer
|
||||
library can be dynamically linked, dynamically loaded by the runtime as
|
||||
a plugin or some API wrapper can be loaded using LD_PRELOAD.
|
||||
The library has a C API.
|
||||
|
||||
The rocTracer library is an API that intercepts runtime API calls and
|
||||
traces asynchronous activity. The activity tracing results are recorded
|
||||
in a ring buffer.
|
||||
|
||||
The rocTX contains application code instrumentation API to support high
|
||||
level correlation of runtime API/activity events. The API includes mark
|
||||
and nested ranges.
|
||||
```
|
||||
## 2. General API
|
||||
### 2.1. Description
|
||||
```
|
||||
The library supports method for getting the error number and error string
|
||||
of the last failed library API call. It allows to check the conformance
|
||||
of used library API header and the library binary, the version macros and
|
||||
API methods can be used.
|
||||
|
||||
Returning the error and error string methods:
|
||||
• roctracer_status_t – error code enumeration
|
||||
• roctracer_error_string – method for returning the error string
|
||||
|
||||
Library version:
|
||||
• ROCTRACER_VERSION_MAJOR – API major version macro
|
||||
• ROCTRACER_VERSION_MINOR – API minor version macro
|
||||
• roctracer_version_major – library major version
|
||||
• roctracer_version_minor – library minor version
|
||||
```
|
||||
### 2.2. Error codes and error string methods
|
||||
```
|
||||
Error code enumeration:
|
||||
typedef enum {
|
||||
ROCTRACER_STATUS_SUCCESS = 0,
|
||||
ROCTRACER_STATUS_ERROR = 1,
|
||||
ROCTRACER_STATUS_UNINIT = 2,
|
||||
ROCTRACER_STATUS_BREAK = 3,
|
||||
ROCTRACER_STATUS_BAD_DOMAIN = 4,
|
||||
ROCTRACER_STATUS_BAD_PARAMETER = 5,
|
||||
ROCTRACER_STATUS_HIP_API_ERR = 6,
|
||||
ROCTRACER_STATUS_HCC_OPS_ERR = 7,
|
||||
ROCTRACER_STATUS_ROCTX_ERR = 8,
|
||||
} roctracer_status_t;
|
||||
|
||||
Return error string:
|
||||
const char* roctracer_error_string();
|
||||
```
|
||||
### 2.3. Library version
|
||||
```
|
||||
The library provides major and minor versions. Major version is for
|
||||
incompatible API changes and minor version for bug fixes.
|
||||
|
||||
API version macros defined in the library API header ‘roctracer.h’:
|
||||
ROCTRACER_VERSION_MAJOR
|
||||
ROCTRACER_VERSION_MINOR
|
||||
|
||||
Methods to check library major and minor venison:
|
||||
uint32_t roctracer_major_version();
|
||||
uint32_t roctracer_minor_version();
|
||||
```
|
||||
## 3. Frontend API
|
||||
### 3.1. Description
|
||||
```
|
||||
The rocTracer provides support for runtime API callbacks and activity
|
||||
records logging. The APIs of different runtimes at different levels
|
||||
are considered as different API domains with assigned domain IDs. For
|
||||
example, language level and driver level. The API callbacks provide
|
||||
the API calls arguments and are called on two phases on “enter” and
|
||||
on “exit”. The activity records are logged to the ring buffer and can
|
||||
be associated with the respective API calls using the correlation ID.
|
||||
Activity API can be used to enable collecting of the records with
|
||||
timestamping data for API calls and asynchronous activity like the
|
||||
kernel submits, memory copies and barriers
|
||||
|
||||
Tracing domains:
|
||||
• roctracer_domain_t – runtime API domains, HIP, HSA, etc…
|
||||
• roctracer_op_string – Return Op string by given domain and
|
||||
activity Op code
|
||||
• roctracer_op_code – Return Op code and kind by given string
|
||||
|
||||
Callback API:
|
||||
• roctracer_rtapi_callback_t – runtime API callback type
|
||||
• roctracer_enable_op_callback – enable runtime API callback
|
||||
by domain and Op code
|
||||
• roctracer_enable_domain_callback – enable runtime API callback
|
||||
by domain for all Ops
|
||||
• roctracer_enable_callback – enable runtime API callback for
|
||||
all domains, all Ops
|
||||
• roctracer_disable_op_callback – disable runtime API callback
|
||||
by domain and Op code
|
||||
• roctracer_enable_op_callback – enable runtime API callback
|
||||
by domain for all Ops
|
||||
• roctracer_enable_op_callback – enable runtime API callback for
|
||||
all domains, all Ops
|
||||
|
||||
Activity API:
|
||||
• roctracer_record_t – activity record
|
||||
• roctracer_pool_t – records pool type
|
||||
• roctracer_allocator_t – tracer allocator type
|
||||
• roctracer_buffer_callback_t – pool callback type
|
||||
• roctracer_open_pool[_expl] – create records pool
|
||||
• roctracer_close_pool[_expl] – close records pool
|
||||
• roctracer_default_pool[_expl] – get/set default pool
|
||||
• roctracer_properties_t – tracer properties
|
||||
• roctracer_enable_op_activity[_expl] – enable activity records logging
|
||||
• roctracer_enable_domain_activity[_expl] – enable activity records logging
|
||||
• roctracer_enable_activity[_expl] – enable activity records logging
|
||||
• roctracer_disable_op_activity – disable activity records logging
|
||||
• roctracer_disable_domain_activity – disable activity records logging
|
||||
• roctracer_disable_activity – disable activity records logging
|
||||
• roctracer_flush_activity[_expl] – disable activity records logging
|
||||
• roctracer_next_record – return next record
|
||||
• roctracer_get_timestamp – return correlated GPU/CPU system timestamp
|
||||
|
||||
External correlation ID API:
|
||||
• roctracer_activity_push_external_correlation_id - push an external
|
||||
correlation id for the calling thread
|
||||
• roctracer_activity_pop_external_correlation_id - pop an external
|
||||
correlation id for the calling thread
|
||||
|
||||
Tracing control API:
|
||||
• roctracer_start – tracing start
|
||||
• roctracer_stop – tracer stop
|
||||
|
||||
```
|
||||
### 3.2. Tracing Domains
|
||||
```
|
||||
Various tracing domains are supported. Each domain is assigned with
|
||||
a domain ID. The domains include HSA, HIP, and HCC runtime levels.
|
||||
|
||||
Traced API domains:
|
||||
typedef enum {
|
||||
ACTIVITY_DOMAIN_HSA_API = 0, // HSA API domain
|
||||
ACTIVITY_DOMAIN_HSA_OPS = 1, // HSA async activity domain
|
||||
ACTIVITY_DOMAIN_HIP_API = 2, // HIP API domain
|
||||
ACTIVITY_DOMAIN_HIP_OPS = 3, // HIP async activity domain
|
||||
ACTIVITY_DOMAIN_KFD_API = 4, // KFD API domain
|
||||
ACTIVITY_DOMAIN_EXT_API = 5, // External ID domain
|
||||
ACTIVITY_DOMAIN_ROCTX = 6, // ROCTX domain
|
||||
ACTIVITY_DOMAIN_NUMBER = 7
|
||||
} activity_domain_t;
|
||||
|
||||
Return name by given domain and Op code:
|
||||
const char* roctracer_op_string( // NULL returned on error and error number is set
|
||||
uint32_t domain, // tracing domain
|
||||
uint32_t op, // activity op code
|
||||
uint32_t kind); // activity kind
|
||||
Return Op code and kind by given string:
|
||||
roctracer_status_t roctracer_op_code(
|
||||
uint32_t domain, // tracing domain
|
||||
const char* str, // [in] op string
|
||||
uint32_t* op, // [out] op code
|
||||
uint32_t* kind); // [out] op kind code if not NULL
|
||||
```
|
||||
### 3.3. Callback API
|
||||
```
|
||||
The tracer provides support for runtime API callbacks and activity records
|
||||
logging. The API callbacks provide the API calls arguments and are called
|
||||
on two phases on “enter”, on “exit”.
|
||||
|
||||
API phase passed to the callbacks:
|
||||
typedef enum {
|
||||
ROCTRACER_API_PHASE_ENTER,
|
||||
ROCTRACER_API_PHASE_EXIT,
|
||||
} roctracer_api_phase_t;
|
||||
|
||||
Runtime API callback type:
|
||||
typedef void (*roctracer_rtapi_callback_t)(
|
||||
uint32_t domain, // runtime API domain
|
||||
uint32_t cid, // API call ID
|
||||
const void* data, // [in] callback data with correlation id and the call
|
||||
// arguments
|
||||
void* arg); // [in/out] user passed data
|
||||
|
||||
Enable runtime API callbacks:
|
||||
roctracer_status_t roctracer_enable_op_callback(
|
||||
activity_domain_t domain, // tracing domain
|
||||
uint32_t op, // API call ID
|
||||
activity_rtapi_callback_t callback, // callback function pointer
|
||||
void* arg); // [in/out] callback arg
|
||||
|
||||
roctracer_status_t roctracer_enable_domain_callback(
|
||||
activity_domain_t domain, // tracing domain
|
||||
activity_rtapi_callback_t callback, // callback function pointer
|
||||
void* arg); // [in/out] callback arg
|
||||
|
||||
|
||||
roctracer_status_t roctracer_enable_callback(
|
||||
activity_rtapi_callback_t callback, // callback function pointer
|
||||
void* arg); // [in/out] callback arg
|
||||
|
||||
Disable runtime API callbacks:
|
||||
roctracer_status_t roctracer_disable_op_callback(
|
||||
activity_domain_t domain, // tracing domain
|
||||
uint32_t op); // API call ID
|
||||
|
||||
roctracer_status_t roctracer_disable_domain_callback(
|
||||
activity_domain_t domain); // tracing domain
|
||||
|
||||
roctracer_status_t roctracer_disable_callback();
|
||||
```
|
||||
### 3.4 Activity API
|
||||
```
|
||||
The activity records are asynchronously logged to the pool and can be
|
||||
associated with the respective API callbacks using the correlation ID.
|
||||
Activity API can be used to enable collecting the records with
|
||||
timestamp data for API calls and GPU activity like kernel submits,
|
||||
memory copies, and barriers.
|
||||
|
||||
// Correlation id
|
||||
typedef uint64_t activity_correlation_id_t;
|
||||
|
||||
Activity record type:
|
||||
|
||||
// Activity record type
|
||||
struct activity_record_t {
|
||||
uint32_t domain; // activity domain id
|
||||
activity_kind_t kind; // activity kind
|
||||
activity_op_t op; // activity op
|
||||
activity_correlation_id_t correlation_id; // activity ID
|
||||
uint64_t begin_ns; // host begin timestamp
|
||||
uint64_t end_ns; // host end timestamp
|
||||
union {
|
||||
struct {
|
||||
int device_id; // device id
|
||||
uint64_t queue_id; // queue id
|
||||
};
|
||||
struct {
|
||||
uint32_t process_id; // device id
|
||||
uint32_t thread_id; // thread id
|
||||
};
|
||||
struct {
|
||||
activity_correlation_id_t external_id; // external correlation id
|
||||
};
|
||||
};
|
||||
size_t bytes; // data size bytes
|
||||
};
|
||||
|
||||
Return next record:
|
||||
static inline int roctracer_next_record(
|
||||
const activity_record_t* record, // [in] record ptr
|
||||
const activity_record_t** next); // [out] next record ptr
|
||||
|
||||
Tracer allocator type:
|
||||
typedef void (*roctracer_allocator_t)(
|
||||
char** ptr, // memory pointer
|
||||
size_t size, // memory size
|
||||
void* arg); // allocator arg
|
||||
|
||||
Pool callback type:
|
||||
typedef void (*roctracer_buffer_callback_t)(
|
||||
const char* begin, // [in] available buffered trace records
|
||||
const char* end, // [in] end of buffered trace records
|
||||
void* arg); // [in/out] callback arg
|
||||
|
||||
Tracer properties:
|
||||
typedef struct {
|
||||
uint32_t mode; // roctracer mode
|
||||
size_t buffer_size; // buffer size
|
||||
// power of 2
|
||||
roctracer_allocator_t alloc_fun; // memory allocator
|
||||
// function pointer
|
||||
void* alloc_arg; // memory allocator
|
||||
// function pointer
|
||||
roctracer_buffer_callback_t buffer_callback_fun; // tracer record
|
||||
// callback function
|
||||
void* buffer_callback_arg; // tracer record
|
||||
// callback arg
|
||||
} roctracer_properties_t;
|
||||
|
||||
Tracer memory pool handle type:
|
||||
typedef void roctracer_pool_t;
|
||||
|
||||
Create tracer memory pool:
|
||||
roctracer_status_t roctracer_open_pool(
|
||||
const roctracer_properties_t* properties); // tracer pool properties
|
||||
|
||||
roctracer_status_t roctracer_open_pool_expl(
|
||||
const roctracer_properties_t* properties, // tracer pool properties
|
||||
roctracer_pool_t** pool); // [out] returns tracer pool if
|
||||
// not NULL, otherwise sets the
|
||||
// default one if it is not set
|
||||
// yet; otherwise the error is
|
||||
// generated
|
||||
|
||||
Close tracer memory pool:
|
||||
roctracer_status_t roctracer_close_pool();
|
||||
|
||||
roctracer_status_t roctracer_close_pool_expl(
|
||||
roctracer_pool_t* pool); // memory pool, NULL means default pool
|
||||
|
||||
Return current default pool. Set new default pool if the argument is not NULL:
|
||||
roctracer_pool_t* roctracer_default_pool();
|
||||
|
||||
roctracer_pool_t* roctracer_default_pool_expl(
|
||||
roctracer_pool_t* pool); // new default pool if not NULL
|
||||
```
|
||||
Enable activity records logging:
|
||||
```
|
||||
roctracer_status_t roctracer_enable_op_activity(
|
||||
activity_domain_t domain, // tracing domain
|
||||
uint32_t op); // activity op ID
|
||||
|
||||
roctracer_status_t roctracer_enable_op_activity_expl(
|
||||
activity_domain_t domain, // tracing domain
|
||||
uint32_t op, // activity op ID
|
||||
roctracer_pool_t* pool); // memory pool, NULL means default pool
|
||||
|
||||
roctracer_status_t roctracer_enable_domain_activity(
|
||||
activity_domain_t domain); // tracing domain
|
||||
|
||||
roctracer_status_t roctracer_enable_domain_activity_expl(
|
||||
activity_domain_t domain, // tracing domain
|
||||
roctracer_pool_t* pool); // memory pool, NULL means default pool
|
||||
|
||||
roctracer_status_t roctracer_enable_activity();
|
||||
|
||||
roctracer_status_t roctracer_enable_activity_expl(
|
||||
roctracer_pool_t* pool); // memory pool, NULL means default pool
|
||||
|
||||
Disable activity records logging:
|
||||
roctracer_status_t roctracer_disable_op_activity(
|
||||
activity_domain_t domain, // tracing domain
|
||||
uint32_t op); // activity op ID
|
||||
|
||||
roctracer_status_t roctracer_disable_domain_activity(
|
||||
activity_domain_t domain); // tracing domain
|
||||
|
||||
roctracer_status_t roctracer_disable_activity();
|
||||
|
||||
Flush available activity records:
|
||||
roctracer_status_t roctracer_flush_activity();
|
||||
|
||||
roctracer_status_t roctracer_flush_activity_expl(
|
||||
roctracer_pool_t* pool); // memory pool, NULL means default pool
|
||||
|
||||
Return correlated GPU/CPU system timestamp:
|
||||
roctracer_status_t roctracer_get_timestamp(
|
||||
uint64_t* timestamp); // [out] return timestamp
|
||||
```
|
||||
External correlation ID API
|
||||
```
|
||||
The API provides activity records to associate rocTracer correlation IDs with
|
||||
IDs provided by external APIs. The external ID records are identified by
|
||||
ACTIVITY_DOMAIN_EXT_API domain value.
|
||||
Using the ‘push’ method an external ID is pushed to a per CPU thread stack and
|
||||
the ‘pop’ method can be used to remove the last pushed ID.
|
||||
An external ID record is inserted before any generated rocTracer activity record
|
||||
if the same CPU external ID stack is non-empty.
|
||||
|
||||
Notifies that the calling thread is entering an external API region.
|
||||
Push an external correlation id for the calling thread.
|
||||
roctracer_status_t roctracer_activity_push_external_correlation_id(
|
||||
activity_correlation_id_t id); // external correlation id
|
||||
|
||||
Notifies that the calling thread is leaving an external API region.
|
||||
Pop an external correlation id for the calling thread.
|
||||
roctracer_status_t roctracer_activity_pop_external_correlation_id(
|
||||
activity_correlation_id_t* last_id); // returns the last external correlation id
|
||||
// if not NULL
|
||||
```
|
||||
Tracing control API
|
||||
```
|
||||
Tracing start:
|
||||
void roctracer_start();
|
||||
|
||||
Tracing stop:
|
||||
void roctracer_stop();
|
||||
```
|
||||
## 4. rocTracer Usage Code Examples
|
||||
### 4.1. HIP API and HCC ops, GPU Activity Tracing
|
||||
```
|
||||
#include <inc/roctracer_hip.h>
|
||||
#include <inc/roctracer_hcc.h>
|
||||
|
||||
// HIP API callback function
|
||||
void hip_api_callback(
|
||||
uint32_t domain,
|
||||
uint32_t cid,
|
||||
const void* callback_data,
|
||||
void* arg)
|
||||
{
|
||||
(void)arg;
|
||||
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_id_string(ACTIVITY_DOMAIN_HIP_API, cid),
|
||||
cid,
|
||||
data->correlation_id,
|
||||
(data->phase == ACTIVITY_API_PHASE_ENTER) ? "on-enter" : "on-exit");
|
||||
<some code . . .>
|
||||
}
|
||||
|
||||
// Activity tracing callback
|
||||
void activity_callback(const char* begin, const char* end, void* arg) {
|
||||
const roctracer_record_t* record = reinterpret_cast<const
|
||||
roctracer_record_t*>(begin);
|
||||
const roctracer_record_t* end_record = reinterpret_cast<const
|
||||
roctracer_record_t*>(end);
|
||||
fprintf(stdout, "\tActivity records:\n");
|
||||
while (record < end_record) {
|
||||
const char * name = roctracer_op_string(record->domain,
|
||||
record->activity_id, 0);
|
||||
fprintf(stdout, "\t%s\tcorrelation_id(%lu) time_ns(%lu:%lu)
|
||||
device_id(%d) stream_id(%lu)\n",
|
||||
name,
|
||||
record->correlation_id,
|
||||
record->begin_ns,
|
||||
record->end_ns,
|
||||
record->device_id,
|
||||
record->stream_id
|
||||
);
|
||||
<some code . . .>
|
||||
ROCTRACER_CALL(roctracer_next_record(record, &record));
|
||||
}
|
||||
}
|
||||
|
||||
int main() {
|
||||
// Allocating tracing pool
|
||||
roctracer_properties_t properties{};
|
||||
properties.buffer_size = 12;
|
||||
properties.buffer_callback_fun = activity_callback;
|
||||
ROCTRACER_CALL(roctracer_open_pool(&properties));
|
||||
|
||||
// Enable HIP API callbacks. HIP_API_ID_ANY can be used to trace all HIP
|
||||
// API calls.
|
||||
ROCTRACER_CALL(roctracer_enable_op_callback(ACTIVITY_DOMAIN_HIP_API,
|
||||
HIP_API_ID_hipModuleLaunchKernel,
|
||||
hip_api_callback, NULL));
|
||||
ROCTRACER_CALL(roctracer_enable_op_acticity(ACTIVITY_DOMAIN_HIP_API,
|
||||
HIP_API_ID_hipModuleLaunchKernel));
|
||||
// Enable HIP kernel dispatch activity tracing
|
||||
ROCTRACER_CALL(roctracer_enable_op_activity(ACTIVITY_DOMAIN_HCC_OPS,
|
||||
hc::HSA_OP_ID_DISPATCH));
|
||||
|
||||
<test code>
|
||||
|
||||
// Disable tracing and closing the pool
|
||||
ROCTRACER_CALL(roctracer_disable_callback());
|
||||
ROCTRACER_CALL(roctracer_disable_activity());
|
||||
ROCTRACER_CALL(roctracer_close_pool());
|
||||
}
|
||||
```
|
||||
### 4.2. MatrixTranspose HIP sample with all APIs/activity tracing enabled
|
||||
```
|
||||
This shows a MatrixTranspose HIP sample with enabled tracing of
|
||||
all HIP API and all GPU asynchronous activity.
|
||||
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
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 <iostream>
|
||||
|
||||
// hip header file
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
#ifndef ITERATIONS
|
||||
# define ITERATIONS 100
|
||||
#endif
|
||||
#define WIDTH 1024
|
||||
|
||||
|
||||
#define NUM (WIDTH * WIDTH)
|
||||
|
||||
#define THREADS_PER_BLOCK_X 4
|
||||
#define THREADS_PER_BLOCK_Y 4
|
||||
#define THREADS_PER_BLOCK_Z 1
|
||||
|
||||
// Device (Kernel) function, it must be void
|
||||
// hipLaunchParm provides the execution configuration
|
||||
__global__ void matrixTranspose(hipLaunchParm lp, float* out, float* in,
|
||||
const int width) {
|
||||
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
|
||||
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
|
||||
|
||||
out[y * width + x] = in[x * width + y];
|
||||
}
|
||||
|
||||
// CPU implementation of matrix transpose
|
||||
void matrixTransposeCPUReference(float* output, float* input, const unsigned
|
||||
int width) {
|
||||
for (unsigned int j = 0; j < width; j++) {
|
||||
for (unsigned int i = 0; i < width; i++) {
|
||||
output[i * width + j] = input[j * width + i];
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
int iterations = ITERATIONS;
|
||||
void start_tracing();
|
||||
void stop_tracing();
|
||||
|
||||
int main() {
|
||||
float* Matrix;
|
||||
float* TransposeMatrix;
|
||||
float* cpuTransposeMatrix;
|
||||
|
||||
float* gpuMatrix;
|
||||
float* gpuTransposeMatrix;
|
||||
|
||||
hipDeviceProp_t devProp;
|
||||
hipGetDeviceProperties(&devProp, 0);
|
||||
|
||||
std::cout << "Device name " << devProp.name << std::endl;
|
||||
|
||||
int i;
|
||||
int errors;
|
||||
|
||||
while (iterations-- > 0) {
|
||||
start_tracing();
|
||||
|
||||
Matrix = (float*)malloc(NUM * sizeof(float));
|
||||
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
|
||||
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
|
||||
|
||||
// initialize the input data
|
||||
for (i = 0; i < NUM; i++) {
|
||||
Matrix[i] = (float)i * 10.0f;
|
||||
}
|
||||
|
||||
// allocate the memory on the device side
|
||||
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
|
||||
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
|
||||
|
||||
// Memory transfer from host to device
|
||||
hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float),
|
||||
hipMemcpyHostToDevice);
|
||||
|
||||
// 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);
|
||||
|
||||
// Memory transfer from device to host
|
||||
hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float),
|
||||
hipMemcpyDeviceToHost);
|
||||
|
||||
// CPU MatrixTranspose computation
|
||||
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
|
||||
|
||||
// verify the results
|
||||
errors = 0;
|
||||
double eps = 1.0E-6;
|
||||
for (i = 0; i < NUM; i++) {
|
||||
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) {
|
||||
errors++;
|
||||
}
|
||||
}
|
||||
if (errors != 0) {
|
||||
printf("FAILED: %d errors\n", errors);
|
||||
} else {
|
||||
printf("PASSED!\n");
|
||||
}
|
||||
|
||||
// free the resources on device side
|
||||
hipFree(gpuMatrix);
|
||||
hipFree(gpuTransposeMatrix);
|
||||
|
||||
// free the resources on host side
|
||||
free(Matrix);
|
||||
free(TransposeMatrix);
|
||||
free(cpuTransposeMatrix);
|
||||
|
||||
stop_tracing();
|
||||
}
|
||||
|
||||
return errors;
|
||||
}
|
||||
|
||||
/////////////////////////////////////////////////////////////////////////////
|
||||
// HIP/HCC Callbacks/Activity tracing
|
||||
/////////////////////////////////////////////////////////////////////////////
|
||||
#include <inc/roctracer_hip.h>
|
||||
#include <inc/roctracer_hcc.h>
|
||||
|
||||
// Macro to check ROC-tracer calls status
|
||||
#define ROCTRACER_CALL(call) \
|
||||
do { \
|
||||
int err = call; \
|
||||
if (err != 0) { \
|
||||
std::cerr << roctracer_error_string() << std::endl << std::flush; \
|
||||
abort(); \
|
||||
} \
|
||||
} while (0)
|
||||
|
||||
// HIP API callback function
|
||||
void hip_api_callback(
|
||||
uint32_t domain,
|
||||
uint32_t cid,
|
||||
const void* callback_data,
|
||||
void* arg)
|
||||
{
|
||||
(void)arg;
|
||||
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),
|
||||
cid,
|
||||
data->correlation_id,
|
||||
(data->phase == ACTIVITY_API_PHASE_ENTER) ? "on-enter" : "on-exit");
|
||||
if (data->phase == ACTIVITY_API_PHASE_ENTER) {
|
||||
switch (cid) {
|
||||
case HIP_API_ID_hipMemcpy:
|
||||
fprintf(stdout, "dst(%p) src(%p) size(0x%x) kind(%u)",
|
||||
data->args.hipMemcpy.dst,
|
||||
data->args.hipMemcpy.src,
|
||||
(uint32_t)(data->args.hipMemcpy.sizeBytes),
|
||||
(uint32_t)(data->args.hipMemcpy.kind));
|
||||
break;
|
||||
case HIP_API_ID_hipMalloc:
|
||||
fprintf(stdout, "ptr(%p) size(0x%x)",
|
||||
data->args.hipMalloc.ptr,
|
||||
(uint32_t)(data->args.hipMalloc.size));
|
||||
break;
|
||||
case HIP_API_ID_hipFree:
|
||||
fprintf(stdout, "ptr(%p),
|
||||
data->args.hipFree.ptr);
|
||||
break;
|
||||
case HIP_API_ID_hipModuleLaunchKernel:
|
||||
fprintf(stdout, "kernel(\"%s\") stream(%p)",
|
||||
hipKernelNameRef(data->args.hipModuleLaunchKernel.f),
|
||||
data->args.hipModuleLaunchKernel.stream);
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
} else {
|
||||
switch (cid) {
|
||||
case HIP_API_ID_hipMalloc:
|
||||
fprintf(stdout, "*ptr(0x%p)",
|
||||
*(data->args.hipMalloc.ptr));
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
}
|
||||
fprintf(stdout, "\n"); fflush(stdout);
|
||||
}
|
||||
|
||||
// Activity tracing callback
|
||||
// hipMalloc id(3) correlation_id(1):
|
||||
// begin_ns(1525888652762640464) end_ns(1525888652762877067)
|
||||
void activity_callback(const char* begin, const char* end, void* arg) {
|
||||
const roctracer_record_t* record = reinterpret_cast
|
||||
<const roctracer_record_t*>(begin);
|
||||
const roctracer_record_t* end_record = reinterpret_cast
|
||||
<const roctracer_record_t*>(end);
|
||||
fprintf(stdout, "\tActivity records:\n"); fflush(stdout);
|
||||
while (record < end_record) {
|
||||
const char * name = roctracer_op_string(record->domain,
|
||||
record->activity_id, 0);
|
||||
fprintf(stdout, "\t%s\tcorrelation_id(%lu) time_ns(%lu:%lu) \
|
||||
device_id(%d) stream_id(%lu)",
|
||||
name,
|
||||
record->correlation_id,
|
||||
record->begin_ns,
|
||||
record->end_ns,
|
||||
record->device_id,
|
||||
record->stream_id
|
||||
);
|
||||
if (record->kind == hc::HSA_OP_ID_COPY)
|
||||
fprintf(stdout, " bytes(0x%zx)", record->bytes);
|
||||
fprintf(stdout, "\n");
|
||||
fflush(stdout);
|
||||
ROCTRACER_CALL(roctracer_next_record(record, &record));
|
||||
}
|
||||
}
|
||||
|
||||
// Start tracing routine
|
||||
void start_tracing() {
|
||||
std::cout << "# START #############################" << std::endl
|
||||
<< std::flush;
|
||||
// Allocating tracing pool
|
||||
roctracer_properties_t properties{};
|
||||
properties.buffer_size = 0x1000;
|
||||
properties.buffer_callback_fun = activity_callback;
|
||||
ROCTRACER_CALL(roctracer_open_pool(&properties));
|
||||
// Enable API callbacks, all domains
|
||||
ROCTRACER_CALL(roctracer_enable_callback(hip_api_callback, NULL));
|
||||
// Enable activity tracing, all domains
|
||||
ROCTRACER_CALL(roctracer_enable_activity());
|
||||
}
|
||||
|
||||
// Stop tracing routine
|
||||
void stop_tracing() {
|
||||
ROCTRACER_CALL(roctracer_disable_api_callback());
|
||||
ROCTRACER_CALL(roctracer_disable_api_activity());
|
||||
ROCTRACER_CALL(roctracer_close_pool());
|
||||
std::cout << "# STOP #############################" << std::endl
|
||||
<< std::flush;
|
||||
}
|
||||
/////////////////////////////////////////////////////////////////////////////
|
||||
```
|
||||
## 5. rocTX application code annotation API
|
||||
```
|
||||
Basic annotation API: markers and nested ranges.
|
||||
// A marker created by given ASCII massage
|
||||
void roctxMark(const char* message);
|
||||
|
||||
// Returns the 0 based level of a nested range being started by given message associated to this range.
|
||||
// A negative value is returned on the error.
|
||||
int roctxRangePush(const char* message);
|
||||
|
||||
// Marks the end of a nested range.
|
||||
// Returns the 0 based level the range.
|
||||
// A negative value is returned on the error.
|
||||
int roctxRangePop();
|
||||
```
|
||||
@@ -29,9 +29,10 @@ THE SOFTWARE.
|
||||
typedef enum {
|
||||
ACTIVITY_DOMAIN_HSA_API = 0, // HSA API domain
|
||||
ACTIVITY_DOMAIN_HSA_OPS = 1, // HSA async activity domain
|
||||
ACTIVITY_DOMAIN_HCC_OPS = 2, // HCC async activity domain
|
||||
ACTIVITY_DOMAIN_HIP_OPS = 2, // HIP async activity domain
|
||||
ACTIVITY_DOMAIN_HCC_OPS = ACTIVITY_DOMAIN_HIP_OPS, // HCC async activity domain
|
||||
ACTIVITY_DOMAIN_HIP_VDI = ACTIVITY_DOMAIN_HIP_OPS, // HIP VDI async activity domain
|
||||
ACTIVITY_DOMAIN_HIP_API = 3, // HIP API domain
|
||||
ACTIVITY_DOMAIN_HIP_VDI = ACTIVITY_DOMAIN_HCC_OPS, // HIP VDI domain
|
||||
ACTIVITY_DOMAIN_KFD_API = 4, // KFD API domain
|
||||
ACTIVITY_DOMAIN_EXT_API = 5, // External ID domain
|
||||
ACTIVITY_DOMAIN_ROCTX = 6, // ROCTX domain
|
||||
@@ -60,7 +61,7 @@ typedef enum {
|
||||
typedef uint64_t activity_correlation_id_t;
|
||||
|
||||
// Activity record type
|
||||
struct activity_record_t {
|
||||
typedef struct activity_record_s {
|
||||
uint32_t domain; // activity domain id
|
||||
activity_kind_t kind; // activity kind
|
||||
activity_op_t op; // activity op
|
||||
@@ -81,7 +82,7 @@ struct activity_record_t {
|
||||
};
|
||||
};
|
||||
size_t bytes; // data size bytes
|
||||
};
|
||||
} activity_record_t;
|
||||
|
||||
// Activity sync calback type
|
||||
typedef void* (*activity_sync_callback_t)(uint32_t cid, activity_record_t* record, const void* data, void* arg);
|
||||
|
||||
+59
-26
@@ -38,10 +38,13 @@ THE SOFTWARE.
|
||||
|
||||
#include <stdint.h>
|
||||
#include <stddef.h>
|
||||
#ifndef __cplusplus
|
||||
#include <stdbool.h>
|
||||
#endif
|
||||
|
||||
#include "ext/prof_protocol.h"
|
||||
|
||||
#define ROCTRACER_VERSION_MAJOR 1
|
||||
#define ROCTRACER_VERSION_MAJOR 2
|
||||
#define ROCTRACER_VERSION_MINOR 0
|
||||
|
||||
#ifdef __cplusplus
|
||||
@@ -89,7 +92,7 @@ roctracer_status_t roctracer_op_code(
|
||||
uint32_t domain, // tracing domain
|
||||
const char* str, // [in] op string
|
||||
uint32_t* op, // [out] op code
|
||||
uint32_t* kind = NULL); // [out] op kind code
|
||||
uint32_t* kind); // [out] op kind code if not NULL
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// Callback API
|
||||
@@ -137,7 +140,7 @@ roctracer_status_t roctracer_disable_callback();
|
||||
typedef activity_record_t roctracer_record_t;
|
||||
|
||||
// Return next record
|
||||
static inline int roctracer_next_record(
|
||||
static inline roctracer_status_t roctracer_next_record(
|
||||
const activity_record_t* record, // [in] record ptr
|
||||
const activity_record_t** next) // [out] next record ptr
|
||||
{
|
||||
@@ -172,31 +175,59 @@ typedef void roctracer_pool_t;
|
||||
|
||||
// Create tracer memory pool
|
||||
// The first invocation sets the default pool
|
||||
roctracer_status_t roctracer_open_pool(
|
||||
roctracer_status_t roctracer_open_pool_expl(
|
||||
const roctracer_properties_t* properties, // tracer pool properties
|
||||
roctracer_pool_t** pool = NULL); // [out] returns tracer pool if not NULL,
|
||||
roctracer_pool_t** pool); // [out] returns tracer pool if not NULL,
|
||||
// otherwise sets the default one if it is not set yet
|
||||
static inline roctracer_status_t roctracer_open_pool(
|
||||
const roctracer_properties_t* properties) // tracer pool properties
|
||||
{
|
||||
return roctracer_open_pool_expl(properties, NULL);
|
||||
}
|
||||
// otherwise the error is generated
|
||||
|
||||
// Close tracer memory pool
|
||||
roctracer_status_t roctracer_close_pool(
|
||||
roctracer_pool_t* pool = NULL); // [in] memory pool, NULL is a default one
|
||||
roctracer_status_t roctracer_close_pool_expl(
|
||||
roctracer_pool_t* pool); // [in] memory pool, NULL is a default one
|
||||
static inline roctracer_status_t roctracer_close_pool()
|
||||
{
|
||||
return roctracer_close_pool_expl(NULL);
|
||||
}
|
||||
|
||||
// Return current default pool
|
||||
// Set new default pool if the argument is not NULL
|
||||
roctracer_pool_t* roctracer_default_pool(
|
||||
roctracer_pool_t* pool = NULL); // [in] new default pool if not NULL
|
||||
roctracer_pool_t* roctracer_default_pool_expl(
|
||||
roctracer_pool_t* pool); // [in] new default pool if not NULL
|
||||
static inline roctracer_pool_t* roctracer_default_pool()
|
||||
{
|
||||
return roctracer_default_pool_expl(NULL);
|
||||
}
|
||||
|
||||
// Enable activity records logging
|
||||
roctracer_status_t roctracer_enable_op_activity(
|
||||
roctracer_status_t roctracer_enable_op_activity_expl(
|
||||
activity_domain_t domain, // tracing domain
|
||||
uint32_t op, // activity op ID
|
||||
roctracer_pool_t* pool = NULL); // memory pool, NULL is a default one
|
||||
roctracer_status_t roctracer_enable_domain_activity(
|
||||
roctracer_pool_t* pool); // memory pool, NULL is a default one
|
||||
static inline roctracer_status_t roctracer_enable_op_activity(
|
||||
activity_domain_t domain, // tracing domain
|
||||
roctracer_pool_t* pool = NULL); // memory pool, NULL is a default one
|
||||
roctracer_status_t roctracer_enable_activity(
|
||||
roctracer_pool_t* pool = NULL); // memory pool, NULL is a default one
|
||||
uint32_t op) // activity op ID
|
||||
{
|
||||
return roctracer_enable_op_activity_expl(domain, op, NULL);
|
||||
}
|
||||
roctracer_status_t roctracer_enable_domain_activity_expl(
|
||||
activity_domain_t domain, // tracing domain
|
||||
roctracer_pool_t* pool); // memory pool, NULL is a default one
|
||||
static inline roctracer_status_t roctracer_enable_domain_activity(
|
||||
activity_domain_t domain) // tracing domain
|
||||
{
|
||||
return roctracer_enable_domain_activity_expl(domain, NULL);
|
||||
}
|
||||
roctracer_status_t roctracer_enable_activity_expl(
|
||||
roctracer_pool_t* pool); // memory pool, NULL is a default one
|
||||
static inline roctracer_status_t roctracer_enable_activity()
|
||||
{
|
||||
return roctracer_enable_activity_expl(NULL);
|
||||
}
|
||||
|
||||
// Disable activity records logging
|
||||
roctracer_status_t roctracer_disable_op_activity(
|
||||
@@ -207,24 +238,26 @@ roctracer_status_t roctracer_disable_domain_activity(
|
||||
roctracer_status_t roctracer_disable_activity();
|
||||
|
||||
// Flush available activity records
|
||||
roctracer_status_t roctracer_flush_activity(
|
||||
roctracer_pool_t* pool = NULL); // memory pool, NULL is a default one
|
||||
roctracer_status_t roctracer_flush_activity_expl(
|
||||
roctracer_pool_t* pool); // memory pool, NULL is a default one
|
||||
static inline roctracer_status_t roctracer_flush_activity()
|
||||
{
|
||||
return roctracer_flush_activity_expl(NULL);
|
||||
}
|
||||
|
||||
// Get system timestamp
|
||||
roctracer_status_t roctracer_get_timestamp(
|
||||
uint64_t* timestamp); // [out] return timestamp
|
||||
|
||||
// Load/Unload methods
|
||||
bool roctracer_load();
|
||||
void roctracer_unload();
|
||||
|
||||
// Set properties
|
||||
roctracer_status_t roctracer_set_properties(
|
||||
roctracer_domain_t domain, // tracing domain
|
||||
void* propertes); // tracing properties
|
||||
|
||||
struct HsaApiTable;
|
||||
bool roctracer_load(
|
||||
HsaApiTable* table,
|
||||
uint64_t runtime_version,
|
||||
uint64_t failed_tool_count,
|
||||
const char* const* failed_tool_names);
|
||||
|
||||
void roctracer_unload(bool destruct);
|
||||
|
||||
#ifdef __cplusplus
|
||||
} // extern "C" block
|
||||
#endif // __cplusplus
|
||||
|
||||
+3
-3
@@ -46,7 +46,7 @@ extern "C" {
|
||||
#endif // __cplusplus
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// Application annotatin API
|
||||
// Application annotation API
|
||||
|
||||
// Tracing start API
|
||||
void roctracer_start();
|
||||
@@ -63,8 +63,8 @@ roctracer_status_t roctracer_activity_push_external_correlation_id(activity_corr
|
||||
|
||||
// Notifies that the calling thread is leaving an external API region.
|
||||
// Pop an external correlation id for the calling thread.
|
||||
// 'lastId' returns the last external correlation
|
||||
roctracer_status_t roctracer_activity_pop_external_correlation_id(activity_correlation_id_t* last_id = NULL);
|
||||
// 'lastId' returns the last external correlation if not NULL
|
||||
roctracer_status_t roctracer_activity_pop_external_correlation_id(activity_correlation_id_t* last_id);
|
||||
|
||||
#ifdef __cplusplus
|
||||
} // extern "C" block
|
||||
|
||||
+10
-14
@@ -23,26 +23,22 @@ THE SOFTWARE.
|
||||
#ifndef INC_ROCTRACER_HCC_H_
|
||||
#define INC_ROCTRACER_HCC_H_
|
||||
|
||||
#if HIP_VDI
|
||||
#define HIP_OP_ID_NUMBER 3
|
||||
#define HIP_OP_ID_COPY 1
|
||||
enum {
|
||||
HIP_OP_ID_DISPATCH = 0,
|
||||
HIP_OP_ID_COPY = 1,
|
||||
HIP_OP_ID_BARRIER = 2,
|
||||
HIP_OP_ID_NUMBER = 3
|
||||
};
|
||||
|
||||
#ifdef __cplusplus
|
||||
extern "C" {
|
||||
#endif
|
||||
typedef void (hipInitAsyncActivityCallback_t)(void* id_callback, void* op_callback, void* arg);
|
||||
typedef bool (hipEnableAsyncActivityCallback_t)(unsigned op, bool enable);
|
||||
typedef const char* (hipGetOpName_t)(unsigned op);
|
||||
#ifdef __cplusplus
|
||||
}
|
||||
#else // !HIP_VDI
|
||||
#if LOCAL_BUILD
|
||||
#include <hc_prof_runtime.h>
|
||||
#else
|
||||
#include <hcc/hc_prof_runtime.h>
|
||||
#endif
|
||||
#define HIP_OP_ID_NUMBER hc::HSA_OP_ID_NUMBER
|
||||
#define HIP_OP_ID_COPY hc::HSA_OP_ID_COPY
|
||||
typedef decltype(Kalmar::CLAMP::InitActivityCallback) hipInitAsyncActivityCallback_t;
|
||||
typedef decltype(Kalmar::CLAMP::EnableActivityCallback) hipEnableAsyncActivityCallback_t;
|
||||
typedef decltype(Kalmar::CLAMP::GetCmdName) hipGetOpName_t;
|
||||
#endif // !HIP_VDI
|
||||
|
||||
#include "roctracer.h"
|
||||
|
||||
|
||||
+1
-1
@@ -33,7 +33,7 @@ extern "C" {
|
||||
#endif // __cplusplus
|
||||
|
||||
// Traced calls ID enumeration
|
||||
typedef hip_api_id_t roctracer_hip_api_cid_t;
|
||||
typedef enum hip_api_id_t roctracer_hip_api_cid_t;
|
||||
|
||||
#ifdef __cplusplus
|
||||
} // extern "C" block
|
||||
|
||||
+8
-4
@@ -22,16 +22,16 @@ THE SOFTWARE.
|
||||
|
||||
#ifndef INC_ROCTRACER_HSA_H_
|
||||
#define INC_ROCTRACER_HSA_H_
|
||||
#include <iostream>
|
||||
#include <mutex>
|
||||
|
||||
#include <hsa.h>
|
||||
#include <hsa_api_trace.h>
|
||||
#include <hsa_ext_amd.h>
|
||||
|
||||
#include "cb_table.h"
|
||||
#include "roctracer.h"
|
||||
|
||||
#ifdef __cplusplus
|
||||
#include <iostream>
|
||||
#include <hsa_api_trace.h>
|
||||
|
||||
namespace roctracer {
|
||||
namespace hsa_support {
|
||||
enum {
|
||||
@@ -129,6 +129,10 @@ struct output_streamer<hsa_queue_t**> {
|
||||
inline static std::ostream& put(std::ostream& out, hsa_queue_t** v) { out << "<queue " << *v << ">"; return out; }
|
||||
};
|
||||
};};
|
||||
#else // !__cplusplus
|
||||
typedef void* hsa_amd_queue_intercept_handler;
|
||||
typedef void* hsa_amd_runtime_queue_notifier;
|
||||
#endif //! __cplusplus
|
||||
|
||||
#include "inc/hsa_prof_str.h"
|
||||
#endif // INC_ROCTRACER_HSA_H_
|
||||
|
||||
+2
-2
@@ -23,11 +23,11 @@ THE SOFTWARE.
|
||||
/////////////////////////////////////////////////////////////////////////////
|
||||
#ifndef INC_ROCTRACER_KFD_H_
|
||||
#define INC_ROCTRACER_KFD_H_
|
||||
#include <iostream>
|
||||
|
||||
#include "roctracer.h"
|
||||
#include "hsakmt.h"
|
||||
#ifdef __cplusplus
|
||||
#include "inc/kfd_ostream_ops.h"
|
||||
#endif
|
||||
#include "inc/kfd_prof_str.h"
|
||||
|
||||
#endif // INC_ROCTRACER_KFD_H_
|
||||
|
||||
@@ -33,8 +33,6 @@ THE SOFTWARE.
|
||||
#ifndef INC_ROCTRACER_ROCTX_H_
|
||||
#define INC_ROCTRACER_ROCTX_H_
|
||||
|
||||
#include "cb_table.h"
|
||||
|
||||
// ROC-TX API ID enumeration
|
||||
enum roctx_api_id_t {
|
||||
ROCTX_API_ID_roctxMarkA = 0,
|
||||
@@ -45,7 +43,7 @@ enum roctx_api_id_t {
|
||||
};
|
||||
|
||||
// ROCTX callbacks data type
|
||||
struct roctx_api_data_t {
|
||||
typedef struct roctx_api_data_s {
|
||||
union {
|
||||
const char* message;
|
||||
struct {
|
||||
@@ -58,14 +56,15 @@ struct roctx_api_data_t {
|
||||
const char* message;
|
||||
} roctxRangePop;
|
||||
} args;
|
||||
};
|
||||
} roctx_api_data_t;
|
||||
|
||||
#ifdef __cplusplus
|
||||
#include "cb_table.h"
|
||||
namespace roctx {
|
||||
|
||||
// ROCTX callbacks table type
|
||||
typedef roctracer::CbTable<ROCTX_API_ID_NUMBER> cb_table_t;
|
||||
|
||||
} // namespace roctx
|
||||
#endif
|
||||
|
||||
#ifdef __cplusplus
|
||||
extern "C" {
|
||||
|
||||
Regular → Executable
@@ -440,6 +440,7 @@ class API_DescrParser:
|
||||
# generate stream operator
|
||||
def gen_out_stream(self, n, name, call, struct):
|
||||
if n == -1:
|
||||
self.content += '#ifdef __cplusplus\n'
|
||||
self.content += 'typedef std::pair<uint32_t, hsa_api_data_t> hsa_api_data_pair_t;\n'
|
||||
self.content += 'inline std::ostream& operator<< (std::ostream& out, const hsa_api_data_pair_t& data_pair) {\n'
|
||||
self.content += ' const uint32_t cid = data_pair.first;\n'
|
||||
@@ -483,6 +484,7 @@ class API_DescrParser:
|
||||
self.content += ' }\n'
|
||||
self.content += ' return out;\n'
|
||||
self.content += '}\n'
|
||||
self.content += '#endif\n'
|
||||
|
||||
#############################################################
|
||||
# main
|
||||
|
||||
+7
-6
@@ -284,7 +284,6 @@ class API_DescrParser:
|
||||
self.content_h += '#include <string.h>\n'
|
||||
self.content_h += '#include \"roctracer_kfd.h\"\n'
|
||||
self.content_h += '#include \"hsakmt.h\"\n'
|
||||
self.content_h += '#include \"cb_table.h\"\n'
|
||||
|
||||
self.content_h += '#define PUBLIC_API __attribute__((visibility(\"default\")))\n'
|
||||
|
||||
@@ -293,6 +292,7 @@ class API_DescrParser:
|
||||
|
||||
self.content_h += '\n'
|
||||
self.content_h += '#if PROF_API_IMPL\n'
|
||||
self.content_h += '#include \"cb_table.h\"\n'
|
||||
self.content_h += 'namespace roctracer {\n'
|
||||
self.content_h += 'namespace kfd_support {\n'
|
||||
|
||||
@@ -372,7 +372,7 @@ class API_DescrParser:
|
||||
# generate API args structure
|
||||
def gen_arg_struct(self, n, name, call, struct):
|
||||
if n == -1:
|
||||
self.content_h += 'struct kfd_api_data_t {\n'
|
||||
self.content_h += 'typedef struct kfd_api_data_s {\n'
|
||||
self.content_h += ' uint64_t correlation_id;\n'
|
||||
self.content_h += ' uint32_t phase;\n'
|
||||
if len(self.api_rettypes) != 0:
|
||||
@@ -394,7 +394,7 @@ class API_DescrParser:
|
||||
self.content_h += ' } ' + call + ';\n'
|
||||
else:
|
||||
self.content_h += ' } args;\n'
|
||||
self.content_h += '};\n'
|
||||
self.content_h += '} kfd_api_data_t;\n'
|
||||
|
||||
# generate API callbacks
|
||||
def gen_callbacks(self, n, name, call, struct):
|
||||
@@ -406,8 +406,7 @@ class API_DescrParser:
|
||||
call_id = self.api_id[call];
|
||||
ret_type = struct['ret']
|
||||
self.content_h += ret_type + ' ' + call + '_callback(' + struct['args'] + ') {\n' # 'static ' +
|
||||
if call == 'hsaKmtOpenKFD':
|
||||
self.content_h += ' if (' + name + '_table == NULL) intercept_KFDApiTable();\n'
|
||||
self.content_h += ' if (' + name + '_table == NULL) intercept_KFDApiTable();\n'
|
||||
self.content_h += ' kfd_api_data_t api_data{};\n'
|
||||
for var in struct['alst']:
|
||||
self.content_h += ' api_data.args.' + call + '.' + var.replace("[]","") + ' = ' + var.replace("[]","") + ';\n'
|
||||
@@ -477,6 +476,7 @@ class API_DescrParser:
|
||||
# generate stream operator
|
||||
def gen_out_stream(self, n, name, call, struct):
|
||||
if n == -1:
|
||||
self.content_h += '#ifdef __cplusplus\n'
|
||||
self.content_h += 'typedef std::pair<uint32_t, kfd_api_data_t> kfd_api_data_pair_t;\n'
|
||||
self.content_h += 'inline std::ostream& operator<< (std::ostream& out, const kfd_api_data_pair_t& data_pair) {\n'
|
||||
self.content_h += ' const uint32_t cid = data_pair.first;\n'
|
||||
@@ -510,6 +510,7 @@ class API_DescrParser:
|
||||
self.content_h += ' }\n'
|
||||
self.content_h += ' return out;\n'
|
||||
self.content_h += '}\n'
|
||||
self.content_h += '#endif\n'
|
||||
self.content_cpp += 'inline std::ostream& operator<< (std::ostream& out, const HsaMemFlags& v) { out << "HsaMemFlags"; return out; }\n'
|
||||
|
||||
# generate PUBLIC_API for all API fcts
|
||||
@@ -525,7 +526,7 @@ class API_DescrParser:
|
||||
self.content_cpp += ' return true;\n';
|
||||
self.content_cpp += '}\n\n';
|
||||
|
||||
if call != '-':
|
||||
if call != '-' and call != 'hsaKmtCloseKFD' and call != 'hsaKmtOpenKFD':
|
||||
self.content_cpp += 'PUBLIC_API ' + struct['ret'] + " " + call + '(' + struct['args'] + ') { return roctracer::kfd_support::' + call + '_callback('
|
||||
for i in range(0,len(struct['alst'])):
|
||||
if i == (len(struct['alst'])-1):
|
||||
|
||||
+10
-12
@@ -17,18 +17,16 @@ target_link_libraries( ${TARGET_LIB} PRIVATE ${HSA_RUNTIME_LIB} c stdc++ )
|
||||
execute_process ( COMMAND sh -xc "${ROOT_DIR}/script/hsaap.py ${ROOT_DIR} ${HSA_RUNTIME_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
|
||||
)
|
||||
execute_process ( COMMAND sh -xc "gcc -E ${HSA_KMT_INC_PATH}/hsakmttypes.h > ${PROJECT_BINARY_DIR}/hsakmttypes_pp.h" )
|
||||
execute_process ( COMMAND sh -xc "${ROOT_DIR}/script/gen_ostream_ops.py -in ${PROJECT_BINARY_DIR}/hsakmttypes_pp.h -out ${ROOT_DIR}/inc/kfd_ostream_ops.h" )
|
||||
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 ( KFD_LIB "kfdwrapper64" )
|
||||
set ( KFD_LIB_SRC
|
||||
${LIB_DIR}/kfd/kfd_wrapper.cpp
|
||||
)
|
||||
execute_process ( COMMAND sh -xc "gcc -E ${HSA_KMT_INC_PATH}/hsakmttypes.h > ${PROJECT_BINARY_DIR}/hsakmttypes_pp.h" )
|
||||
execute_process ( COMMAND sh -xc "${ROOT_DIR}/script/gen_ostream_ops.py -in ${PROJECT_BINARY_DIR}/hsakmttypes_pp.h -out ${ROOT_DIR}/inc/kfd_ostream_ops.h" )
|
||||
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}" )
|
||||
|
||||
set ( ROCTX_LIB "roctx64" )
|
||||
set ( ROCTX_LIB_SRC
|
||||
|
||||
+15
-7
@@ -46,9 +46,9 @@ class BaseLoader : public T {
|
||||
|
||||
private:
|
||||
BaseLoader() {
|
||||
const int flags = RTLD_LAZY;
|
||||
const int flags = (to_load_ == true) ? RTLD_LAZY : RTLD_LAZY|RTLD_NOLOAD;
|
||||
handle_ = dlopen(lib_name_, flags);
|
||||
if (handle_ == NULL) {
|
||||
if ((to_check_ == true) && (handle_ == NULL)) {
|
||||
fprintf(stderr, "roctracer: Loading '%s' failed, %s\n", lib_name_, dlerror());
|
||||
abort();
|
||||
}
|
||||
@@ -61,6 +61,9 @@ class BaseLoader : public T {
|
||||
if (handle_ != NULL) dlclose(handle_);
|
||||
}
|
||||
|
||||
static bool to_load_;
|
||||
static bool to_check_;
|
||||
|
||||
static mutex_t mutex_;
|
||||
static const char* lib_name_;
|
||||
static std::atomic<loader_t*> instance_;
|
||||
@@ -110,9 +113,9 @@ class HccApi {
|
||||
protected:
|
||||
void init(Loader* loader) {
|
||||
#if HIP_VDI
|
||||
InitActivityCallback = loader->GetFun<hipInitAsyncActivityCallback_t>("InitActivityCallback");
|
||||
EnableActivityCallback = loader->GetFun<hipEnableAsyncActivityCallback_t>("EnableActivityCallback");
|
||||
GetOpName = loader->GetFun<hipGetOpName_t>("GetCmdName");
|
||||
InitActivityCallback = loader->GetFun<hipInitAsyncActivityCallback_t>("hipInitActivityCallback");
|
||||
EnableActivityCallback = loader->GetFun<hipEnableAsyncActivityCallback_t>("hipEnableActivityCallback");
|
||||
GetOpName = loader->GetFun<hipGetOpName_t>("hipGetCmdName");
|
||||
#else
|
||||
InitActivityCallback = loader->GetFun<hipInitAsyncActivityCallback_t>("InitActivityCallbackImpl");
|
||||
EnableActivityCallback = loader->GetFun<hipEnableAsyncActivityCallback_t>("EnableActivityCallbackImpl");
|
||||
@@ -171,9 +174,14 @@ typedef BaseLoader<RocTxApi> RocTxLoader;
|
||||
#define LOADER_INSTANTIATE() \
|
||||
template<class T> typename roctracer::BaseLoader<T>::mutex_t roctracer::BaseLoader<T>::mutex_; \
|
||||
template<class T> std::atomic<roctracer::BaseLoader<T>*> roctracer::BaseLoader<T>::instance_{}; \
|
||||
template<class T> bool roctracer::BaseLoader<T>::to_load_ = false; \
|
||||
template<class T> bool roctracer::BaseLoader<T>::to_check_ = true; \
|
||||
template<> const char* roctracer::HipLoader::lib_name_ = "libhip_hcc.so"; \
|
||||
template<> const char* roctracer::HccLoader::lib_name_ = "libmcwamp_hsa.so"; \
|
||||
template<> bool roctracer::HipLoader::to_check_ = false; \
|
||||
template<> const char* roctracer::HccLoader::lib_name_ = "libmcwamp.so"; \
|
||||
template<> bool roctracer::HccLoader::to_check_ = false; \
|
||||
template<> const char* roctracer::KfdLoader::lib_name_ = "libkfdwrapper64.so"; \
|
||||
template<> const char* roctracer::RocTxLoader::lib_name_ = "libroctx64.so";
|
||||
template<> const char* roctracer::RocTxLoader::lib_name_ = "libroctx64.so"; \
|
||||
template<> bool roctracer::RocTxLoader::to_load_ = true;
|
||||
|
||||
#endif // SRC_CORE_LOADER_H_
|
||||
|
||||
+96
-59
@@ -27,9 +27,7 @@ THE SOFTWARE.
|
||||
#include "inc/roctracer_roctx.h"
|
||||
#define PROF_API_IMPL 1
|
||||
#include "inc/roctracer_hsa.h"
|
||||
#ifdef KFD_WRAPPER
|
||||
#include "inc/roctracer_kfd.h"
|
||||
#endif
|
||||
|
||||
#include <dirent.h>
|
||||
#include <pthread.h>
|
||||
@@ -87,9 +85,13 @@ THE SOFTWARE.
|
||||
(void)err; \
|
||||
return X;
|
||||
|
||||
#ifndef onload_debug
|
||||
#define onload_debug false
|
||||
#endif
|
||||
#define ONLOAD_TRACE(str) \
|
||||
if (getenv("ROCP_ONLOAD_TRACE")) do { \
|
||||
std::cout << "PID(" << GetPid() << "): TRACER_LIB::" << __FUNCTION__ << " " << str << std::endl << std::flush; \
|
||||
} while(0);
|
||||
#define ONLOAD_TRACE_BEG() ONLOAD_TRACE("begin")
|
||||
#define ONLOAD_TRACE_END() ONLOAD_TRACE("end")
|
||||
|
||||
|
||||
static inline uint32_t GetPid() { return syscall(__NR_getpid); }
|
||||
|
||||
@@ -174,7 +176,7 @@ decltype(hsa_amd_memory_async_copy_rect)* hsa_amd_memory_async_copy_rect_fn;
|
||||
|
||||
typedef decltype(roctracer_enable_op_callback)* roctracer_enable_op_callback_t;
|
||||
typedef decltype(roctracer_disable_op_callback)* roctracer_disable_op_callback_t;
|
||||
typedef decltype(roctracer_enable_op_activity)* roctracer_enable_op_activity_t;
|
||||
typedef decltype(roctracer_enable_op_activity_expl)* roctracer_enable_op_activity_t;
|
||||
typedef decltype(roctracer_disable_op_activity)* roctracer_disable_op_activity_t;
|
||||
|
||||
struct cb_journal_data_t {
|
||||
@@ -251,18 +253,16 @@ class GlobalCounter {
|
||||
public:
|
||||
typedef std::mutex mutex_t;
|
||||
typedef uint64_t counter_t;
|
||||
typedef std::atomic<counter_t> atomic_counter_t;
|
||||
|
||||
static counter_t Increment() {
|
||||
std::lock_guard<mutex_t> lock(mutex_);
|
||||
return ++counter_;
|
||||
}
|
||||
static counter_t Increment() { return counter_.fetch_add(1, std::memory_order_relaxed); }
|
||||
|
||||
private:
|
||||
static mutex_t mutex_;
|
||||
static counter_t counter_;
|
||||
static atomic_counter_t counter_;
|
||||
};
|
||||
GlobalCounter::mutex_t GlobalCounter::mutex_;
|
||||
GlobalCounter::counter_t GlobalCounter::counter_ = 0;
|
||||
GlobalCounter::atomic_counter_t GlobalCounter::counter_{1};
|
||||
|
||||
// Records storage
|
||||
struct roctracer_api_data_t {
|
||||
@@ -284,6 +284,7 @@ typedef std::map<activity_correlation_id_t, activity_correlation_id_t> correlati
|
||||
typedef std::mutex correlation_id_mutex_t;
|
||||
correlation_id_map_t* correlation_id_map = NULL;
|
||||
correlation_id_mutex_t correlation_id_mutex;
|
||||
bool correlation_id_wait = true;
|
||||
|
||||
static thread_local std::stack<activity_correlation_id_t> external_id_stack;
|
||||
|
||||
@@ -296,6 +297,7 @@ static inline void CorrelationIdRegistr(const activity_correlation_id_t& correla
|
||||
|
||||
static inline activity_correlation_id_t CorrelationIdLookup(const activity_correlation_id_t& correlation_id) {
|
||||
auto it = correlation_id_map->find(correlation_id);
|
||||
if (correlation_id_wait) while (it == correlation_id_map->end()) it = correlation_id_map->find(correlation_id);
|
||||
if (it == correlation_id_map->end()) EXC_ABORT(ROCTRACER_STATUS_ERROR, "HCC activity id lookup failed(" << correlation_id << ")");
|
||||
return it->second;
|
||||
}
|
||||
@@ -330,6 +332,7 @@ void* HIP_SyncActivityCallback(
|
||||
data = &(top.data.hip);
|
||||
data_ptr = const_cast<hip_api_data_t*>(data);
|
||||
data_ptr->phase = phase;
|
||||
data_ptr->correlation_id = 0;
|
||||
}
|
||||
|
||||
// Filing record info
|
||||
@@ -523,14 +526,29 @@ hsa_status_t hsa_amd_memory_async_copy_rect_interceptor(
|
||||
return status;
|
||||
}
|
||||
|
||||
// Logger routines and primitives
|
||||
util::Logger::mutex_t util::Logger::mutex_;
|
||||
std::atomic<util::Logger*> util::Logger::instance_{};
|
||||
|
||||
// Memory pool routines and primitives
|
||||
MemoryPool* memory_pool = NULL;
|
||||
typedef std::recursive_mutex memory_pool_mutex_t;
|
||||
memory_pool_mutex_t memory_pool_mutex;
|
||||
|
||||
// Stop sttaus routines and primitives
|
||||
unsigned stop_status_value = 0;
|
||||
typedef std::mutex stop_status_mutex_t;
|
||||
stop_status_mutex_t stop_status_mutex;
|
||||
unsigned set_stopped(unsigned val) {
|
||||
std::lock_guard<stop_status_mutex_t> lock(stop_status_mutex);
|
||||
const unsigned ret = (stop_status_value ^ val);
|
||||
stop_status_value = val;
|
||||
return ret;
|
||||
}
|
||||
} // namespace roctracer
|
||||
|
||||
LOADER_INSTANTIATE();
|
||||
TRACE_BUFFER_INSTANTIATE();
|
||||
|
||||
///////////////////////////////////////////////////////////////////////////////////////////////////
|
||||
// Public library methods
|
||||
@@ -567,12 +585,10 @@ PUBLIC_API const char* roctracer_op_string(
|
||||
return roctracer::HipLoader::Instance().ApiName(op);
|
||||
break;
|
||||
}
|
||||
#if KFD_WRAPPER
|
||||
case ACTIVITY_DOMAIN_KFD_API: {
|
||||
return roctracer::kfd_support::GetApiName(op);
|
||||
break;
|
||||
}
|
||||
#endif
|
||||
default:
|
||||
EXC_RAISING(ROCTRACER_STATUS_BAD_DOMAIN, "invalid domain ID(" << domain << ")");
|
||||
}
|
||||
@@ -593,13 +609,11 @@ PUBLIC_API roctracer_status_t roctracer_op_code(
|
||||
if (kind != NULL) *kind = 0;
|
||||
break;
|
||||
}
|
||||
#ifdef KFD_WRAPPER
|
||||
case ACTIVITY_DOMAIN_KFD_API: {
|
||||
*op = roctracer::kfd_support::GetApiCode(str);
|
||||
if (kind != NULL) *kind = 0;
|
||||
break;
|
||||
}
|
||||
#endif
|
||||
default:
|
||||
EXC_RAISING(ROCTRACER_STATUS_BAD_DOMAIN, "limited domain ID(" << domain << ")");
|
||||
}
|
||||
@@ -612,9 +626,7 @@ static inline uint32_t get_op_num(const uint32_t& domain) {
|
||||
case ACTIVITY_DOMAIN_HSA_API: return HSA_API_ID_NUMBER;
|
||||
case ACTIVITY_DOMAIN_HCC_OPS: return HIP_OP_ID_NUMBER;
|
||||
case ACTIVITY_DOMAIN_HIP_API: return HIP_API_ID_NUMBER;
|
||||
#ifdef KFD_WRAPPER
|
||||
case ACTIVITY_DOMAIN_KFD_API: return KFD_API_ID_NUMBER;
|
||||
#endif
|
||||
case ACTIVITY_DOMAIN_EXT_API: return 0;
|
||||
case ACTIVITY_DOMAIN_ROCTX: return ROCTX_API_ID_NUMBER;
|
||||
default:
|
||||
@@ -631,13 +643,11 @@ static roctracer_status_t roctracer_enable_callback_fun(
|
||||
void* user_data)
|
||||
{
|
||||
switch (domain) {
|
||||
#ifdef KFD_WRAPPER
|
||||
case ACTIVITY_DOMAIN_KFD_API: {
|
||||
const bool succ = roctracer::KfdLoader::Instance().RegisterApiCallback(op, (void*)callback, user_data);
|
||||
if (succ == false) EXC_RAISING(ROCTRACER_STATUS_ERROR, "KFD RegisterApiCallback error");
|
||||
break;
|
||||
}
|
||||
#endif
|
||||
case ACTIVITY_DOMAIN_HSA_OPS: break;
|
||||
case ACTIVITY_DOMAIN_HSA_API: {
|
||||
roctracer::hsa_support::cb_table.set(op, callback, user_data);
|
||||
@@ -645,6 +655,8 @@ static roctracer_status_t roctracer_enable_callback_fun(
|
||||
}
|
||||
case ACTIVITY_DOMAIN_HCC_OPS: break;
|
||||
case ACTIVITY_DOMAIN_HIP_API: {
|
||||
if (roctracer::HipLoader::Instance().Enabled() == false) break;
|
||||
|
||||
hipError_t hip_err = roctracer::HipLoader::Instance().RegisterApiCallback(op, (void*)callback, user_data);
|
||||
if (hip_err != hipSuccess) HIP_EXC_RAISING(ROCTRACER_STATUS_HIP_API_ERR, "hipRegisterApiCallback(" << op << ") error(" << hip_err << ")");
|
||||
break;
|
||||
@@ -712,17 +724,17 @@ static roctracer_status_t roctracer_disable_callback_fun(
|
||||
uint32_t op)
|
||||
{
|
||||
switch (domain) {
|
||||
#ifdef KFD_WRAPPER
|
||||
case ACTIVITY_DOMAIN_KFD_API: {
|
||||
const bool succ = roctracer::KfdLoader::Instance().RemoveApiCallback(op);
|
||||
if (succ == false) EXC_RAISING(ROCTRACER_STATUS_ERROR, "KFD RemoveApiCallback error");
|
||||
break;
|
||||
}
|
||||
#endif
|
||||
case ACTIVITY_DOMAIN_HSA_OPS: break;
|
||||
case ACTIVITY_DOMAIN_HSA_API: break;
|
||||
case ACTIVITY_DOMAIN_HCC_OPS: break;
|
||||
case ACTIVITY_DOMAIN_HIP_API: {
|
||||
if (roctracer::HipLoader::Instance().Enabled() == false) break;
|
||||
|
||||
hipError_t hip_err = roctracer::HipLoader::Instance().RemoveApiCallback(op);
|
||||
if (hip_err != hipSuccess) HIP_EXC_RAISING(ROCTRACER_STATUS_HIP_API_ERR, "hipRemoveApiCallback error(" << hip_err << ")");
|
||||
break;
|
||||
@@ -777,7 +789,7 @@ PUBLIC_API roctracer_status_t roctracer_disable_callback()
|
||||
}
|
||||
|
||||
// Return default pool and set new one if parameter pool is not NULL.
|
||||
PUBLIC_API roctracer_pool_t* roctracer_default_pool(roctracer_pool_t* pool) {
|
||||
PUBLIC_API roctracer_pool_t* roctracer_default_pool_expl(roctracer_pool_t* pool) {
|
||||
std::lock_guard<roctracer::memory_pool_mutex_t> lock(roctracer::memory_pool_mutex);
|
||||
roctracer_pool_t* p = reinterpret_cast<roctracer_pool_t*>(roctracer::memory_pool);
|
||||
if (pool != NULL) roctracer::memory_pool = reinterpret_cast<roctracer::MemoryPool*>(pool);
|
||||
@@ -785,7 +797,7 @@ PUBLIC_API roctracer_pool_t* roctracer_default_pool(roctracer_pool_t* pool) {
|
||||
}
|
||||
|
||||
// Open memory pool
|
||||
PUBLIC_API roctracer_status_t roctracer_open_pool(
|
||||
PUBLIC_API roctracer_status_t roctracer_open_pool_expl(
|
||||
const roctracer_properties_t* properties,
|
||||
roctracer_pool_t** pool)
|
||||
{
|
||||
@@ -802,7 +814,7 @@ PUBLIC_API roctracer_status_t roctracer_open_pool(
|
||||
}
|
||||
|
||||
// Close memory pool
|
||||
PUBLIC_API roctracer_status_t roctracer_close_pool(roctracer_pool_t* pool) {
|
||||
PUBLIC_API roctracer_status_t roctracer_close_pool_expl(roctracer_pool_t* pool) {
|
||||
API_METHOD_PREFIX
|
||||
std::lock_guard<roctracer::memory_pool_mutex_t> lock(roctracer::memory_pool_mutex);
|
||||
roctracer_pool_t* ptr = (pool == NULL) ? roctracer_default_pool() : pool;
|
||||
@@ -828,7 +840,18 @@ static roctracer_status_t roctracer_enable_activity_fun(
|
||||
case ACTIVITY_DOMAIN_HSA_API: break;
|
||||
case ACTIVITY_DOMAIN_KFD_API: break;
|
||||
case ACTIVITY_DOMAIN_HCC_OPS: {
|
||||
if (roctracer::HccLoader::GetRef() == NULL) {
|
||||
const bool init_phase = (roctracer::HccLoader::GetRef() == NULL);
|
||||
if (roctracer::HccLoader::Instance().Enabled() == false) break;
|
||||
|
||||
if (init_phase == true) {
|
||||
if (getenv("ROCP_HCC_CORRID_WAIT") != NULL) {
|
||||
roctracer::correlation_id_wait = true;
|
||||
fprintf(stdout, "roctracer: HCC correlation ID wait enabled\n"); fflush(stdout);
|
||||
}
|
||||
if (getenv("ROCP_HCC_CORRID_NOWAIT") != NULL) {
|
||||
roctracer::correlation_id_wait = false;
|
||||
fprintf(stdout, "roctracer: HCC correlation ID wait disabled\n"); fflush(stdout);
|
||||
}
|
||||
roctracer::HccLoader::Instance().InitActivityCallback((void*)roctracer::HCC_ActivityIdCallback,
|
||||
(void*)roctracer::HCC_AsyncActivityCallback,
|
||||
(void*)pool);
|
||||
@@ -838,6 +861,8 @@ static roctracer_status_t roctracer_enable_activity_fun(
|
||||
break;
|
||||
}
|
||||
case ACTIVITY_DOMAIN_HIP_API: {
|
||||
if (roctracer::HipLoader::Instance().Enabled() == false) break;
|
||||
|
||||
const hipError_t hip_err = roctracer::HipLoader::Instance().RegisterActivityCallback(op, (void*)roctracer::HIP_SyncActivityCallback, (void*)pool);
|
||||
if (hip_err != hipSuccess) HIP_EXC_RAISING(ROCTRACER_STATUS_HIP_API_ERR, "hipRegisterActivityCallback error(" << hip_err << ")");
|
||||
break;
|
||||
@@ -858,7 +883,7 @@ static void roctracer_enable_activity_impl(
|
||||
roctracer_enable_activity_fun((roctracer_domain_t)domain, op, pool);
|
||||
}
|
||||
|
||||
PUBLIC_API roctracer_status_t roctracer_enable_op_activity(
|
||||
PUBLIC_API roctracer_status_t roctracer_enable_op_activity_expl(
|
||||
roctracer_domain_t domain,
|
||||
uint32_t op,
|
||||
roctracer_pool_t* pool)
|
||||
@@ -868,7 +893,7 @@ PUBLIC_API roctracer_status_t roctracer_enable_op_activity(
|
||||
API_METHOD_SUFFIX
|
||||
}
|
||||
|
||||
PUBLIC_API roctracer_status_t roctracer_enable_domain_activity(
|
||||
PUBLIC_API roctracer_status_t roctracer_enable_domain_activity_expl(
|
||||
roctracer_domain_t domain,
|
||||
roctracer_pool_t* pool)
|
||||
{
|
||||
@@ -878,7 +903,7 @@ PUBLIC_API roctracer_status_t roctracer_enable_domain_activity(
|
||||
API_METHOD_SUFFIX
|
||||
}
|
||||
|
||||
PUBLIC_API roctracer_status_t roctracer_enable_activity(
|
||||
PUBLIC_API roctracer_status_t roctracer_enable_activity_expl(
|
||||
roctracer_pool_t* pool)
|
||||
{
|
||||
API_METHOD_PREFIX
|
||||
@@ -903,11 +928,15 @@ static roctracer_status_t roctracer_disable_activity_fun(
|
||||
case ACTIVITY_DOMAIN_HSA_API: break;
|
||||
case ACTIVITY_DOMAIN_KFD_API: break;
|
||||
case ACTIVITY_DOMAIN_HCC_OPS: {
|
||||
if (roctracer::HccLoader::Instance().Enabled() == false) break;
|
||||
|
||||
const bool succ = roctracer::HccLoader::Instance().EnableActivityCallback(op, false);
|
||||
if (succ == false) HCC_EXC_RAISING(ROCTRACER_STATUS_HCC_OPS_ERR, "HCC::EnableActivityCallback(NULL) error domain(" << domain << ") op(" << op << ")");
|
||||
break;
|
||||
}
|
||||
case ACTIVITY_DOMAIN_HIP_API: {
|
||||
if (roctracer::HipLoader::Instance().Enabled() == false) break;
|
||||
|
||||
const hipError_t hip_err = roctracer::HipLoader::Instance().RemoveActivityCallback(op);
|
||||
if (hip_err != hipSuccess) HIP_EXC_RAISING(ROCTRACER_STATUS_HIP_API_ERR, "hipRemoveActivityCallback error(" << hip_err << ")");
|
||||
break;
|
||||
@@ -956,11 +985,12 @@ PUBLIC_API roctracer_status_t roctracer_disable_activity()
|
||||
}
|
||||
|
||||
// Flush available activity records
|
||||
PUBLIC_API roctracer_status_t roctracer_flush_activity(roctracer_pool_t* pool) {
|
||||
PUBLIC_API roctracer_status_t roctracer_flush_activity_expl(roctracer_pool_t* pool) {
|
||||
API_METHOD_PREFIX
|
||||
if (pool == NULL) pool = roctracer_default_pool();
|
||||
roctracer::MemoryPool* memory_pool = reinterpret_cast<roctracer::MemoryPool*>(pool);
|
||||
memory_pool->Flush();
|
||||
roctracer::TraceBufferBase::FlushAll();
|
||||
API_METHOD_SUFFIX
|
||||
}
|
||||
|
||||
@@ -1001,16 +1031,26 @@ PUBLIC_API void roctracer_mark(const char* str) {
|
||||
|
||||
// Start API
|
||||
PUBLIC_API void roctracer_start() {
|
||||
if (roctracer::ext_support::roctracer_start_cb) roctracer::ext_support::roctracer_start_cb();
|
||||
roctracer::cb_journal->foreach(roctracer::cb_en_functor_t(roctracer_enable_callback_fun));
|
||||
roctracer::act_journal->foreach(roctracer::act_en_functor_t(roctracer_enable_activity_fun));
|
||||
if (roctracer::set_stopped(0)) {
|
||||
if (roctracer::ext_support::roctracer_start_cb) roctracer::ext_support::roctracer_start_cb();
|
||||
roctracer::cb_journal->foreach(roctracer::cb_en_functor_t(roctracer_enable_callback_fun));
|
||||
roctracer::act_journal->foreach(roctracer::act_en_functor_t(roctracer_enable_activity_fun));
|
||||
}
|
||||
}
|
||||
|
||||
// Stop API
|
||||
PUBLIC_API void roctracer_stop() {
|
||||
roctracer::cb_journal->foreach(roctracer::cb_dis_functor_t(roctracer_disable_callback_fun));
|
||||
roctracer::act_journal->foreach(roctracer::act_dis_functor_t(roctracer_disable_activity_fun));
|
||||
if (roctracer::ext_support::roctracer_stop_cb) roctracer::ext_support::roctracer_stop_cb();
|
||||
if (roctracer::set_stopped(1)) {
|
||||
roctracer::cb_journal->foreach(roctracer::cb_dis_functor_t(roctracer_disable_callback_fun));
|
||||
roctracer::act_journal->foreach(roctracer::act_dis_functor_t(roctracer_disable_activity_fun));
|
||||
if (roctracer::ext_support::roctracer_stop_cb) roctracer::ext_support::roctracer_stop_cb();
|
||||
}
|
||||
}
|
||||
|
||||
PUBLIC_API roctracer_status_t roctracer_get_timestamp(uint64_t* timestamp) {
|
||||
API_METHOD_PREFIX
|
||||
*timestamp = util::HsaRsrcFactory::Instance().TimestampNs();
|
||||
API_METHOD_SUFFIX
|
||||
}
|
||||
|
||||
// Set properties
|
||||
@@ -1021,6 +1061,8 @@ PUBLIC_API roctracer_status_t roctracer_set_properties(
|
||||
API_METHOD_PREFIX
|
||||
switch (domain) {
|
||||
case ACTIVITY_DOMAIN_HSA_OPS: {
|
||||
roctracer::trace_buffer.StartWorkerThread();
|
||||
|
||||
// HSA OPS properties
|
||||
roctracer::hsa_ops_properties_t* ops_properties = reinterpret_cast<roctracer::hsa_ops_properties_t*>(properties);
|
||||
HsaApiTable* table = reinterpret_cast<HsaApiTable*>(ops_properties->table);
|
||||
@@ -1046,12 +1088,10 @@ PUBLIC_API roctracer_status_t roctracer_set_properties(
|
||||
|
||||
break;
|
||||
}
|
||||
#ifdef KFD_WRAPPER
|
||||
case ACTIVITY_DOMAIN_KFD_API: {
|
||||
roctracer::kfd_support::intercept_KFDApiTable();
|
||||
break;
|
||||
}
|
||||
#endif
|
||||
case ACTIVITY_DOMAIN_HSA_API: {
|
||||
// HSA API properties
|
||||
HsaApiTable* table = reinterpret_cast<HsaApiTable*>(properties);
|
||||
@@ -1062,7 +1102,7 @@ PUBLIC_API roctracer_status_t roctracer_set_properties(
|
||||
}
|
||||
case ACTIVITY_DOMAIN_HCC_OPS:
|
||||
case ACTIVITY_DOMAIN_HIP_API: {
|
||||
#ifdef HIP_VDI
|
||||
#if HIP_VDI
|
||||
const char* hip_lib_name = "libamdhip64.so";
|
||||
roctracer::HccLoader::SetLibName(hip_lib_name);
|
||||
roctracer::HipLoader::SetLibName(hip_lib_name);
|
||||
@@ -1082,58 +1122,55 @@ PUBLIC_API roctracer_status_t roctracer_set_properties(
|
||||
API_METHOD_SUFFIX
|
||||
}
|
||||
|
||||
// HSA-runtime tool on-load method
|
||||
PUBLIC_API bool roctracer_load(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
|
||||
const char* const* failed_tool_names) {
|
||||
if (onload_debug) { printf("LIB roctracer_load\n"); fflush(stdout); }
|
||||
PUBLIC_API bool roctracer_load() {
|
||||
static bool is_loaded = false;
|
||||
ONLOAD_TRACE("begin, loaded(" << is_loaded << ")");
|
||||
|
||||
if (is_loaded) return true;
|
||||
is_loaded = true;
|
||||
|
||||
if (onload_debug) { printf("LIB roctracer_load end\n"); fflush(stdout); }
|
||||
ONLOAD_TRACE_END();
|
||||
return true;
|
||||
}
|
||||
|
||||
PUBLIC_API void roctracer_unload(bool destruct) {
|
||||
PUBLIC_API void roctracer_unload() {
|
||||
static bool is_unloaded = false;
|
||||
ONLOAD_TRACE("begin, unloaded(" << is_unloaded << ")");
|
||||
|
||||
if (onload_debug) { printf("LIB roctracer_unload (%d, %d)\n", (int)destruct, (int)is_unloaded); fflush(stdout); }
|
||||
if (destruct == false) return;
|
||||
if (is_unloaded == true) return;
|
||||
is_unloaded = true;
|
||||
|
||||
roctracer::trace_buffer.Flush();
|
||||
roctracer::close_output_file(roctracer::kernel_file_handle);
|
||||
if (onload_debug) { printf("LIB roctracer_unload end\n"); fflush(stdout); }
|
||||
ONLOAD_TRACE_END();
|
||||
}
|
||||
|
||||
// HSA-runtime tool on-load/unload methods
|
||||
PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
|
||||
const char* const* failed_tool_names) {
|
||||
if (onload_debug) { printf("LIB OnLoad\n"); fflush(stdout); }
|
||||
const bool ret = roctracer_load(table, runtime_version, failed_tool_count, failed_tool_names);
|
||||
if (onload_debug) { printf("LIB OnLoad end\n"); fflush(stdout); }
|
||||
ONLOAD_TRACE_BEG();
|
||||
const bool ret = roctracer_load();
|
||||
ONLOAD_TRACE_END();
|
||||
return ret;
|
||||
}
|
||||
PUBLIC_API void OnUnload() {
|
||||
if (onload_debug) { printf("LIB OnUnload\n"); fflush(stdout); }
|
||||
roctracer_unload(false);
|
||||
if (onload_debug) { printf("LIB OnUnload end\n"); fflush(stdout); }
|
||||
ONLOAD_TRACE("done");
|
||||
}
|
||||
|
||||
CONSTRUCTOR_API void constructor() {
|
||||
if (onload_debug) { printf("LIB constructor\n"); fflush(stdout); }
|
||||
ONLOAD_TRACE_BEG();
|
||||
roctracer::util::Logger::Create();
|
||||
if (roctracer::cb_journal == NULL) roctracer::cb_journal = new roctracer::CbJournal;
|
||||
if (roctracer::act_journal == NULL) roctracer::act_journal = new roctracer::ActJournal;
|
||||
if (onload_debug) { printf("LIB constructor end\n"); fflush(stdout); }
|
||||
ONLOAD_TRACE_END();
|
||||
}
|
||||
|
||||
DESTRUCTOR_API void destructor() {
|
||||
if (onload_debug) { printf("LIB destructor\n"); fflush(stdout); }
|
||||
roctracer_unload(true);
|
||||
ONLOAD_TRACE_BEG();
|
||||
roctracer_unload();
|
||||
util::HsaRsrcFactory::Destroy();
|
||||
roctracer::util::Logger::Destroy();
|
||||
if (onload_debug) { printf("LIB destructor end\n"); fflush(stdout); }
|
||||
ONLOAD_TRACE_END();
|
||||
}
|
||||
|
||||
} // extern "C"
|
||||
|
||||
+94
-22
@@ -2,12 +2,23 @@
|
||||
#define SRC_CORE_TRACE_BUFFER_H_
|
||||
|
||||
#include <atomic>
|
||||
#include <iostream>
|
||||
#include <list>
|
||||
#include <mutex>
|
||||
#include <sstream>
|
||||
|
||||
#include <pthread.h>
|
||||
#include <string.h>
|
||||
#include <unistd.h>
|
||||
|
||||
#define FATAL(stream) \
|
||||
do { \
|
||||
std::ostringstream oss; \
|
||||
oss << __FUNCTION__ << "(), " << stream; \
|
||||
std::cout << oss.str() << std::endl; \
|
||||
abort(); \
|
||||
} while (0)
|
||||
|
||||
#define PTHREAD_CALL(call) \
|
||||
do { \
|
||||
int err = call; \
|
||||
@@ -53,8 +64,55 @@ struct trace_entry_t {
|
||||
};
|
||||
};
|
||||
|
||||
template <class T>
|
||||
struct push_element_fun {
|
||||
T* const elem_;
|
||||
void fun(T* node) { if (node->next_elem_ == NULL) node->next_elem_ = elem_; }
|
||||
push_element_fun(T* elem) : elem_(elem) {}
|
||||
};
|
||||
|
||||
template <class T>
|
||||
struct call_element_fun {
|
||||
void (T::*fptr_)();
|
||||
void fun(T* node) { (node->*fptr_)(); }
|
||||
call_element_fun(void (T::*f)()) : fptr_(f) {}
|
||||
};
|
||||
|
||||
struct TraceBufferBase {
|
||||
typedef std::mutex mutex_t;
|
||||
|
||||
virtual void StartWorkerThread() = 0;
|
||||
virtual void Flush() = 0;
|
||||
|
||||
static void StartWorkerThreadAll() { foreach(call_element_fun<TraceBufferBase>(&TraceBufferBase::StartWorkerThread)); }
|
||||
static void FlushAll() { foreach(call_element_fun<TraceBufferBase>(&TraceBufferBase::Flush)); }
|
||||
|
||||
static void Push(TraceBufferBase* elem) {
|
||||
if (head_elem_ == NULL) head_elem_ = elem;
|
||||
else foreach(push_element_fun<TraceBufferBase>(elem));
|
||||
}
|
||||
|
||||
TraceBufferBase() : next_elem_(NULL) {}
|
||||
|
||||
template<class F>
|
||||
static void foreach(const F& f_in) {
|
||||
std::lock_guard<mutex_t> lck(mutex_);
|
||||
F f = f_in;
|
||||
TraceBufferBase* p = head_elem_;
|
||||
while (p != NULL) {
|
||||
TraceBufferBase* next = p->next_elem_;
|
||||
f.fun(p);
|
||||
p = next;
|
||||
}
|
||||
}
|
||||
|
||||
TraceBufferBase* next_elem_;
|
||||
static TraceBufferBase* head_elem_;
|
||||
static mutex_t mutex_;
|
||||
};
|
||||
|
||||
template <typename Entry>
|
||||
class TraceBuffer {
|
||||
class TraceBuffer : protected TraceBufferBase {
|
||||
public:
|
||||
typedef void (*callback_t)(Entry*);
|
||||
typedef TraceBuffer<Entry> Obj;
|
||||
@@ -67,7 +125,8 @@ class TraceBuffer {
|
||||
};
|
||||
|
||||
TraceBuffer(const char* name, uint32_t size, flush_prm_t* flush_prm_arr, uint32_t flush_prm_count) :
|
||||
is_flushed_(false)
|
||||
is_flushed_(false),
|
||||
work_thread_started_(false)
|
||||
{
|
||||
name_ = strdup(name);
|
||||
size_ = size;
|
||||
@@ -80,31 +139,43 @@ class TraceBuffer {
|
||||
flush_prm_arr_ = flush_prm_arr;
|
||||
flush_prm_count_ = flush_prm_count;
|
||||
|
||||
PTHREAD_CALL(pthread_mutex_init(&work_mutex_, NULL));
|
||||
PTHREAD_CALL(pthread_cond_init(&work_cond_, NULL));
|
||||
PTHREAD_CALL(pthread_create(&work_thread_, NULL, allocate_worker, this));
|
||||
TraceBufferBase::Push(this);
|
||||
}
|
||||
|
||||
~TraceBuffer() {
|
||||
PTHREAD_CALL(pthread_cancel(work_thread_));
|
||||
void *res;
|
||||
PTHREAD_CALL(pthread_join(work_thread_, &res));
|
||||
if (res != PTHREAD_CANCELED) abort_run("~TraceBuffer: consumer thread wasn't stopped correctly");
|
||||
|
||||
StopWorkerThread();
|
||||
Flush();
|
||||
}
|
||||
|
||||
void StartWorkerThread() {
|
||||
std::lock_guard<mutex_t> lck(mutex_);
|
||||
if (work_thread_started_ == false) {
|
||||
PTHREAD_CALL(pthread_mutex_init(&work_mutex_, NULL));
|
||||
PTHREAD_CALL(pthread_cond_init(&work_cond_, NULL));
|
||||
PTHREAD_CALL(pthread_create(&work_thread_, NULL, allocate_worker, this));
|
||||
work_thread_started_ = true;
|
||||
}
|
||||
}
|
||||
|
||||
void StopWorkerThread() {
|
||||
std::lock_guard<mutex_t> lck(mutex_);
|
||||
if (work_thread_started_ == true) {
|
||||
PTHREAD_CALL(pthread_cancel(work_thread_));
|
||||
void *res;
|
||||
PTHREAD_CALL(pthread_join(work_thread_, &res));
|
||||
if (res != PTHREAD_CANCELED) FATAL("consumer thread wasn't stopped correctly");
|
||||
work_thread_started_ = false;
|
||||
}
|
||||
}
|
||||
|
||||
Entry* GetEntry() {
|
||||
const pointer_t pointer = read_pointer_.fetch_add(1);
|
||||
if (pointer >= end_pointer_) wrap_buffer(pointer);
|
||||
if (pointer >= end_pointer_) abort_run("pointer >= end_pointer_ after buffer wrap");
|
||||
if (pointer >= end_pointer_) FATAL("pointer >= end_pointer_ after buffer wrap");
|
||||
return data_ + (pointer + size_ - end_pointer_);
|
||||
}
|
||||
|
||||
void Flush() {
|
||||
flush_buf();
|
||||
}
|
||||
void Flush() { flush_buf(); }
|
||||
|
||||
private:
|
||||
void flush_buf() {
|
||||
@@ -134,7 +205,7 @@ class TraceBuffer {
|
||||
|
||||
inline Entry* allocate_fun() {
|
||||
Entry* ptr = (Entry*) malloc(size_ * sizeof(Entry));
|
||||
if (ptr == NULL) abort_run("TraceBuffer::allocate_fun: calloc failed");
|
||||
if (ptr == NULL) FATAL("malloc failed");
|
||||
//memset(ptr, 0, size_ * sizeof(Entry));
|
||||
return ptr;
|
||||
}
|
||||
@@ -156,24 +227,20 @@ class TraceBuffer {
|
||||
|
||||
void wrap_buffer(const pointer_t pointer) {
|
||||
std::lock_guard<mutex_t> lck(mutex_);
|
||||
if (work_thread_started_ == false) FATAL("worker thread is not started");
|
||||
|
||||
PTHREAD_CALL(pthread_mutex_lock(&work_mutex_));
|
||||
if (pointer >= end_pointer_) {
|
||||
data_ = next_;
|
||||
next_ = NULL;
|
||||
PTHREAD_CALL(pthread_cond_signal(&work_cond_));
|
||||
end_pointer_ += size_;
|
||||
if (end_pointer_ == 0) abort_run("TraceBuffer::wrap_buffer: pointer overflow");
|
||||
if (end_pointer_ == 0) FATAL("pointer overflow");
|
||||
buf_list_.push_back(data_);
|
||||
}
|
||||
PTHREAD_CALL(pthread_mutex_unlock(&work_mutex_));
|
||||
}
|
||||
|
||||
void abort_run(const char* str) {
|
||||
fprintf(stderr, "%s\n", str);
|
||||
fflush(stderr);
|
||||
abort();
|
||||
}
|
||||
|
||||
const char* name_;
|
||||
uint32_t size_;
|
||||
Entry* data_;
|
||||
@@ -189,9 +256,14 @@ class TraceBuffer {
|
||||
pthread_t work_thread_;
|
||||
pthread_mutex_t work_mutex_;
|
||||
pthread_cond_t work_cond_;
|
||||
bool work_thread_started_;
|
||||
|
||||
mutex_t mutex_;
|
||||
};
|
||||
} // namespace roctracer
|
||||
|
||||
#define TRACE_BUFFER_INSTANTIATE() \
|
||||
roctracer::TraceBufferBase* roctracer::TraceBufferBase::head_elem_ = NULL; \
|
||||
roctracer::TraceBufferBase::mutex_t roctracer::TraceBufferBase::mutex_;
|
||||
|
||||
#endif // SRC_CORE_TRACE_BUFFER_H_
|
||||
|
||||
+7
-2
@@ -30,11 +30,16 @@ set ( HSA_TEST_DIR "${TEST_DIR}/hsa/test" )
|
||||
set ( RUN_SCRIPT "${TEST_DIR}/run.sh" )
|
||||
|
||||
## build HIP tests
|
||||
set ( TEST_CFLAGS HIP_VDI=${HIP_VDI} ROCM_INC_PATH=${ROCM_INC_PATH} )
|
||||
add_custom_target( mytest
|
||||
COMMAND make -C "${TEST_DIR}/MatrixTranspose"
|
||||
COMMAND sh -xc "cp ${TEST_DIR}/MatrixTranspose/MatrixTranspose ${PROJECT_BINARY_DIR}/test"
|
||||
COMMAND HIP_VDI=${HIP_VDI} KFD_WRAPPER=${KFD_WRAPPER} make -C "${TEST_DIR}/MatrixTranspose_test"
|
||||
COMMAND ${TEST_CFLAGS} make -C "${TEST_DIR}/MatrixTranspose_test"
|
||||
COMMAND sh -xc "cp ${TEST_DIR}/MatrixTranspose_test/MatrixTranspose ${PROJECT_BINARY_DIR}/test/MatrixTranspose_test"
|
||||
COMMAND MGPU_TEST=1 ${TEST_CFLAGS} make -C "${TEST_DIR}/MatrixTranspose_test"
|
||||
COMMAND sh -xc "cp ${TEST_DIR}/MatrixTranspose_test/MatrixTranspose ${PROJECT_BINARY_DIR}/test/MatrixTranspose_mgpu"
|
||||
COMMAND C_TEST=1 ${TEST_CFLAGS} make -C "${TEST_DIR}/MatrixTranspose_test"
|
||||
COMMAND sh -xc "cp ${TEST_DIR}/MatrixTranspose_test/MatrixTranspose ${PROJECT_BINARY_DIR}/test/MatrixTranspose_ctest"
|
||||
)
|
||||
|
||||
## Util sources
|
||||
@@ -49,7 +54,7 @@ target_link_libraries ( ${TEST_LIB} ${ROCTRACER_TARGET} ${HSA_RUNTIME_LIB} c std
|
||||
|
||||
## Build HSA test
|
||||
execute_process ( COMMAND sh -xc "if [ ! -e ${TEST_DIR}/hsa ] ; then git clone https://github.com/ROCmSoftwarePlatform/hsa-class.git ${TEST_DIR}/hsa; fi" )
|
||||
execute_process ( COMMAND sh -xc "if [ -e ${TEST_DIR}/hsa ] ; then cd ${TEST_DIR}/hsa && git fetch origin && git checkout 7defb6d; fi" )
|
||||
execute_process ( COMMAND sh -xc "if [ -e ${TEST_DIR}/hsa ] ; then cd ${TEST_DIR}/hsa && git fetch origin && git checkout a3aabb5; fi" )
|
||||
set ( TEST_DIR ${HSA_TEST_DIR} )
|
||||
add_subdirectory ( ${TEST_DIR} ${PROJECT_BINARY_DIR}/test/hsa )
|
||||
|
||||
|
||||
@@ -1,7 +1,6 @@
|
||||
ROOT_PATH = ../..
|
||||
LIB_PATH = $(ROOT_PATH)/build
|
||||
ROC_LIBS = -L$(LIB_PATH) -lroctracer64
|
||||
export LD_LIBRARY_PATH=$(LIB_PATH)
|
||||
ROC_LIBS = -Wl,--rpath,${LIB_PATH} $(LIB_PATH)/libroctracer64.so $(LIB_PATH)/libroctx64.so
|
||||
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
|
||||
@@ -23,8 +23,10 @@ THE SOFTWARE.
|
||||
#include <iostream>
|
||||
|
||||
// hip header file
|
||||
#include "hip/hip_runtime.h"
|
||||
#include <hip/hip_runtime.h>
|
||||
#include "roctracer_ext.h"
|
||||
// roctx header file
|
||||
#include <inc/roctx.h>
|
||||
|
||||
|
||||
#define WIDTH 1024
|
||||
@@ -94,15 +96,23 @@ int main() {
|
||||
hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice);
|
||||
|
||||
roctracer_mark("before HIP LaunchKernel");
|
||||
roctxMark("before hipLaunchKernel");
|
||||
roctxRangePush("hipLaunchKernel");
|
||||
// Lauching kernel from host
|
||||
hipLaunchKernelGGL(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);
|
||||
roctracer_mark("after HIP LaunchKernel");
|
||||
roctxMark("after hipLaunchKernel");
|
||||
|
||||
// Memory transfer from device to host
|
||||
roctxRangePush("hipMemcpy");
|
||||
|
||||
hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost);
|
||||
|
||||
roctxRangePop(); // for "hipMemcpy"
|
||||
roctxRangePop(); // for "hipLaunchKernel"
|
||||
|
||||
// CPU MatrixTranspose computation
|
||||
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
|
||||
|
||||
|
||||
@@ -1,8 +1,7 @@
|
||||
ROOT_PATH = ../..
|
||||
LIB_PATH = $(ROOT_PATH)/build
|
||||
ROC_LIBS = -L$(LIB_PATH) -lroctracer64 -lroctx64
|
||||
HSA_KMT_INC_PATH ?= /opt/rocm/include
|
||||
export LD_LIBRARY_PATH=$(LIB_PATH)
|
||||
ROC_LIBS = -Wl,--rpath,${LIB_PATH} $(LIB_PATH)/libroctracer64.so $(LIB_PATH)/libroctx64.so
|
||||
ROCM_INC_PATH ?= /opt/rocm/include
|
||||
HIP_VDI ?= 0
|
||||
ITERATIONS ?= 100
|
||||
|
||||
@@ -12,27 +11,37 @@ ifeq (,$(HIP_PATH))
|
||||
endif
|
||||
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
|
||||
TARGET=hcc
|
||||
|
||||
SOURCES = MatrixTranspose.cpp
|
||||
OBJECTS = $(SOURCES:.cpp=.o)
|
||||
|
||||
EXECUTABLE=./MatrixTranspose
|
||||
OBJECTS = MatrixTranspose.o
|
||||
FLAGS =-g -I$(ROOT_PATH) -I$(ROOT_PATH)/inc -I${ROCM_INC_PATH}/hsa -I${ROCM_INC_PATH} -DLOCAL_BUILD=1 -DHIP_VDI=${HIP_VDI} -DITERATIONS=$(ITERATIONS) -DAMD_INTERNAL_BUILD=1
|
||||
|
||||
ifeq ($(C_TEST), 1)
|
||||
COMP=gcc
|
||||
SOURCES = MatrixTranspose.c
|
||||
FLAGS += -DHIP_TEST=0 -D__HIP_PLATFORM_HCC__=1 -I${ROCM_INC_PATH}/hcc
|
||||
else
|
||||
COMP=$(HIPCC)
|
||||
SOURCES = MatrixTranspose.cpp
|
||||
FLAGS += -DHIP_TEST=1
|
||||
endif
|
||||
ifeq ($(MGPU_TEST), 1)
|
||||
FLAGS += -DMGPU_TEST=1
|
||||
endif
|
||||
|
||||
.PHONY: test
|
||||
|
||||
|
||||
all: clean $(EXECUTABLE)
|
||||
|
||||
CXXFLAGS =-g -I$(ROOT_PATH) -I$(ROOT_PATH)/inc -I${HSA_KMT_INC_PATH} -DLOCAL_BUILD=1 -DHIP_VDI=${HIP_VDI} -DITERATIONS=$(ITERATIONS)
|
||||
CXX=$(HIPCC)
|
||||
$(OBJECTS): $(SOURCES)
|
||||
$(COMP) $(FLAGS) -c -o $@ $<
|
||||
|
||||
$(EXECUTABLE): $(OBJECTS)
|
||||
$(HIPCC) $(OBJECTS) -o $@ $(ROC_LIBS)
|
||||
|
||||
test: $(EXECUTABLE)
|
||||
$(EXECUTABLE)
|
||||
LD_PRELOAD=$(LIB_PATH)/libkfdwrapper64.so $(EXECUTABLE)
|
||||
|
||||
clean:
|
||||
rm -f $(EXECUTABLE)
|
||||
|
||||
@@ -0,0 +1 @@
|
||||
MatrixTranspose.cpp
|
||||
@@ -20,34 +20,39 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <iostream>
|
||||
#include <stdio.h>
|
||||
|
||||
// roctracer extension API
|
||||
#include <inc/roctracer_ext.h>
|
||||
|
||||
// hip header file
|
||||
#include <hip/hip_runtime.h>
|
||||
#ifdef __cplusplus
|
||||
#include <cstdlib>
|
||||
using namespace std;
|
||||
#else
|
||||
#include <stdlib.h>
|
||||
#endif
|
||||
|
||||
// roctx header file
|
||||
#include <inc/roctx.h>
|
||||
// roctracer extension API
|
||||
#include <inc/roctracer_ext.h>
|
||||
|
||||
// kfd header file
|
||||
#ifdef KFD_WRAPPER
|
||||
#include <inc/roctracer_kfd.h>
|
||||
#if HIP_TEST
|
||||
// hip header file
|
||||
#include <hip/hip_runtime.h>
|
||||
// Macro to call HIP API
|
||||
#define HIP_CALL(call) do { call; } while(0)
|
||||
#else
|
||||
#define HIP_CALL(call) do {} while(0)
|
||||
#endif
|
||||
|
||||
#ifndef ITERATIONS
|
||||
# define ITERATIONS 101
|
||||
#endif
|
||||
#define WIDTH 1024
|
||||
|
||||
|
||||
#define NUM (WIDTH * WIDTH)
|
||||
|
||||
#define THREADS_PER_BLOCK_X 4
|
||||
#define THREADS_PER_BLOCK_Y 4
|
||||
#define THREADS_PER_BLOCK_Z 1
|
||||
|
||||
#if HIP_TEST
|
||||
// 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;
|
||||
@@ -55,6 +60,7 @@ __global__ void matrixTranspose(float* out, float* in, const int width) {
|
||||
|
||||
out[y * width + x] = in[x * width + y];
|
||||
}
|
||||
#endif
|
||||
|
||||
// CPU implementation of matrix transpose
|
||||
void matrixTransposeCPUReference(float* output, float* input, const unsigned int width) {
|
||||
@@ -78,19 +84,33 @@ int main() {
|
||||
float* gpuMatrix;
|
||||
float* gpuTransposeMatrix;
|
||||
|
||||
hipDeviceProp_t devProp;
|
||||
hipGetDeviceProperties(&devProp, 0);
|
||||
|
||||
std::cout << "Device name " << devProp.name << std::endl;
|
||||
|
||||
int i;
|
||||
int errors;
|
||||
|
||||
init_tracing();
|
||||
|
||||
#if HIP_TEST
|
||||
int gpuCount = 1;
|
||||
#if MGPU_TEST
|
||||
hipGetDeviceCount(&gpuCount);
|
||||
printf("Number of GPUs: %d\n", gpuCount);
|
||||
#endif
|
||||
iterations *= gpuCount;
|
||||
#endif
|
||||
|
||||
while (iterations-- > 0) {
|
||||
start_tracing();
|
||||
|
||||
#if HIP_TEST
|
||||
// set GPU
|
||||
const int devIndex = iterations % gpuCount;
|
||||
hipSetDevice(devIndex);
|
||||
|
||||
hipDeviceProp_t devProp;
|
||||
HIP_CALL(hipGetDeviceProperties(&devProp, 0));
|
||||
printf("Device %d name: %s\n", devIndex, devProp.name);
|
||||
#endif
|
||||
|
||||
Matrix = (float*)malloc(NUM * sizeof(float));
|
||||
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
|
||||
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
|
||||
@@ -101,8 +121,8 @@ int main() {
|
||||
}
|
||||
|
||||
// allocate the memory on the device side
|
||||
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
|
||||
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
|
||||
HIP_CALL(hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)));
|
||||
HIP_CALL(hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)));
|
||||
|
||||
// correlation reagion32
|
||||
roctracer_activity_push_external_correlation_id(31);
|
||||
@@ -110,7 +130,7 @@ int main() {
|
||||
roctracer_activity_push_external_correlation_id(32);
|
||||
|
||||
// Memory transfer from host to device
|
||||
hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice);
|
||||
HIP_CALL(hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice));
|
||||
|
||||
// correlation reagion33
|
||||
roctracer_activity_push_external_correlation_id(33);
|
||||
@@ -119,9 +139,9 @@ int main() {
|
||||
roctxRangePush("hipLaunchKernel");
|
||||
|
||||
// Lauching kernel from host
|
||||
hipLaunchKernelGGL(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);
|
||||
HIP_CALL(hipLaunchKernelGGL(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));
|
||||
|
||||
roctxMark("after hipLaunchKernel");
|
||||
|
||||
@@ -131,39 +151,40 @@ int main() {
|
||||
// Memory transfer from device to host
|
||||
roctxRangePush("hipMemcpy");
|
||||
|
||||
hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost);
|
||||
HIP_CALL(hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost));
|
||||
|
||||
roctxRangePop(); // for "hipMemcpy"
|
||||
roctxRangePop(); // for "hipLaunchKernel"
|
||||
|
||||
// correlation reagion end
|
||||
roctracer_activity_pop_external_correlation_id();
|
||||
roctracer_activity_pop_external_correlation_id(NULL);
|
||||
|
||||
// CPU MatrixTranspose computation
|
||||
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
|
||||
HIP_CALL(matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH));
|
||||
|
||||
// verify the results
|
||||
errors = 0;
|
||||
double eps = 1.0E-6;
|
||||
for (i = 0; i < NUM; i++) {
|
||||
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) {
|
||||
if (abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) {
|
||||
errors++;
|
||||
}
|
||||
}
|
||||
if (errors != 0) {
|
||||
if ((HIP_TEST != 0) && (errors != 0)) {
|
||||
printf("FAILED: %d errors\n", errors);
|
||||
} else {
|
||||
errors = 0;
|
||||
printf("PASSED!\n");
|
||||
}
|
||||
|
||||
// free the resources on device side
|
||||
hipFree(gpuMatrix);
|
||||
hipFree(gpuTransposeMatrix);
|
||||
HIP_CALL(hipFree(gpuMatrix));
|
||||
HIP_CALL(hipFree(gpuTransposeMatrix));
|
||||
|
||||
// correlation reagion end
|
||||
roctracer_activity_pop_external_correlation_id();
|
||||
roctracer_activity_pop_external_correlation_id(NULL);
|
||||
// correlation reagion end
|
||||
roctracer_activity_pop_external_correlation_id();
|
||||
roctracer_activity_pop_external_correlation_id(NULL);
|
||||
|
||||
// free the resources on host side
|
||||
free(Matrix);
|
||||
@@ -182,6 +203,8 @@ int main() {
|
||||
#if 1
|
||||
#include <inc/roctracer_hip.h>
|
||||
#include <inc/roctracer_hcc.h>
|
||||
#include <inc/roctracer_hsa.h>
|
||||
#include <inc/roctracer_kfd.h>
|
||||
#include <inc/roctracer_roctx.h>
|
||||
|
||||
// Macro to check ROC-tracer calls status
|
||||
@@ -189,7 +212,7 @@ int main() {
|
||||
do { \
|
||||
int err = call; \
|
||||
if (err != 0) { \
|
||||
std::cerr << roctracer_error_string() << std::endl << std::flush; \
|
||||
fprintf(stderr, "%s\n", roctracer_error_string()); \
|
||||
abort(); \
|
||||
} \
|
||||
} while (0)
|
||||
@@ -204,22 +227,20 @@ void api_callback(
|
||||
(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);
|
||||
const roctx_api_data_t* data = (const roctx_api_data_t*)(callback_data);
|
||||
fprintf(stdout, "<rocTX \"%s\">\n", data->args.message);
|
||||
return;
|
||||
}
|
||||
#ifdef KFD_WRAPPER
|
||||
if (domain == ACTIVITY_DOMAIN_KFD_API) {
|
||||
const kfd_api_data_t* data = reinterpret_cast<const kfd_api_data_t*>(callback_data);
|
||||
fprintf(stdout, "KFD: <%s id(%u)\tcorrelation_id(%lu) %s> \n",
|
||||
const kfd_api_data_t* data = (const kfd_api_data_t*)(callback_data);
|
||||
fprintf(stdout, "<%s id(%u)\tcorrelation_id(%lu) %s> \n",
|
||||
roctracer_op_string(ACTIVITY_DOMAIN_KFD_API, cid, 0),
|
||||
cid,
|
||||
data->correlation_id,
|
||||
(data->phase == ACTIVITY_API_PHASE_ENTER) ? "on-enter" : "on-exit");
|
||||
return;
|
||||
}
|
||||
#endif
|
||||
const hip_api_data_t* data = reinterpret_cast<const hip_api_data_t*>(callback_data);
|
||||
const hip_api_data_t* data = (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),
|
||||
cid,
|
||||
@@ -267,8 +288,8 @@ void api_callback(
|
||||
// Activity tracing callback
|
||||
// hipMalloc id(3) correlation_id(1): begin_ns(1525888652762640464) end_ns(1525888652762877067)
|
||||
void activity_callback(const char* begin, const char* end, void* arg) {
|
||||
const roctracer_record_t* record = reinterpret_cast<const roctracer_record_t*>(begin);
|
||||
const roctracer_record_t* end_record = reinterpret_cast<const roctracer_record_t*>(end);
|
||||
const roctracer_record_t* record = (const roctracer_record_t*)(begin);
|
||||
const roctracer_record_t* end_record = (const roctracer_record_t*)(end);
|
||||
fprintf(stdout, "\tActivity records:\n"); fflush(stdout);
|
||||
while (record < end_record) {
|
||||
const char * name = roctracer_op_string(record->domain, record->op, record->kind);
|
||||
@@ -278,7 +299,7 @@ void activity_callback(const char* begin, const char* end, void* arg) {
|
||||
record->begin_ns,
|
||||
record->end_ns
|
||||
);
|
||||
if (record->domain == ACTIVITY_DOMAIN_HIP_API or record->domain == ACTIVITY_DOMAIN_KFD_API) {
|
||||
if ((record->domain == ACTIVITY_DOMAIN_HIP_API) || (record->domain == ACTIVITY_DOMAIN_KFD_API)) {
|
||||
fprintf(stdout, " process_id(%u) thread_id(%u)",
|
||||
record->process_id,
|
||||
record->thread_id
|
||||
@@ -305,11 +326,12 @@ void activity_callback(const char* begin, const char* end, void* arg) {
|
||||
|
||||
// Init tracing routine
|
||||
void init_tracing() {
|
||||
std::cout << "# INIT #############################" << std::endl << std::flush;
|
||||
printf("# INIT #############################\n");
|
||||
// roctracer properties
|
||||
roctracer_set_properties(ACTIVITY_DOMAIN_HIP_API, NULL);
|
||||
// Allocating tracing pool
|
||||
roctracer_properties_t properties{};
|
||||
roctracer_properties_t properties;
|
||||
memset(&properties, 0, sizeof(roctracer_properties_t));
|
||||
properties.buffer_size = 0x1000;
|
||||
properties.buffer_callback_fun = activity_callback;
|
||||
ROCTRACER_CALL(roctracer_open_pool(&properties));
|
||||
@@ -318,11 +340,16 @@ void init_tracing() {
|
||||
// Enable HIP activity tracing
|
||||
ROCTRACER_CALL(roctracer_enable_domain_activity(ACTIVITY_DOMAIN_HIP_API));
|
||||
ROCTRACER_CALL(roctracer_enable_domain_activity(ACTIVITY_DOMAIN_HCC_OPS));
|
||||
// Enable KFD API tracing
|
||||
ROCTRACER_CALL(roctracer_enable_domain_callback(ACTIVITY_DOMAIN_KFD_API, api_callback, NULL));
|
||||
ROCTRACER_CALL(roctracer_enable_domain_activity(ACTIVITY_DOMAIN_KFD_API));
|
||||
// Enable rocTX
|
||||
ROCTRACER_CALL(roctracer_enable_domain_callback(ACTIVITY_DOMAIN_ROCTX, api_callback, NULL));
|
||||
}
|
||||
|
||||
// Start tracing routine
|
||||
void start_tracing() {
|
||||
std::cout << "# START (" << iterations << ") #############################" << std::endl << std::flush;
|
||||
printf("# START (%d) #############################\n", iterations);
|
||||
// Start
|
||||
if ((iterations & 1) == 1) roctracer_start();
|
||||
else roctracer_stop();
|
||||
@@ -333,8 +360,9 @@ void stop_tracing() {
|
||||
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_disable_domain_activity(ACTIVITY_DOMAIN_KFD_API));
|
||||
ROCTRACER_CALL(roctracer_flush_activity());
|
||||
std::cout << "# STOP #############################" << std::endl << std::flush;
|
||||
printf("# STOP #############################\n");
|
||||
}
|
||||
#else
|
||||
void init_tracing() {}
|
||||
|
||||
+12
-6
@@ -65,17 +65,23 @@ eval_test() {
|
||||
|
||||
# Standalone test
|
||||
# rocTrecer is used explicitely by test
|
||||
eval_test "standalone HIP test" "./test/MatrixTranspose_test"
|
||||
eval_test "standalone C test" "LD_PRELOAD=libkfdwrapper64.so ./test/MatrixTranspose_ctest"
|
||||
eval_test "standalone HIP test" "LD_PRELOAD=libkfdwrapper64.so ./test/MatrixTranspose_test"
|
||||
eval_test "standalone HIP MGPU test" "LD_PRELOAD=libkfdwrapper64.so ./test/MatrixTranspose_mgpu"
|
||||
|
||||
# Tool test
|
||||
# rocTracer/tool is loaded by HSA runtime
|
||||
export HSA_TOOLS_LIB="test/libtracer_tool.so"
|
||||
export ROCTRACER_DOMAIN="hip"
|
||||
|
||||
# HIP test
|
||||
eval_test "tool HIP test" ./test/MatrixTranspose
|
||||
# with trace sampling control <delay:length:rate>
|
||||
eval_test "tool HIP period test" "ROCP_CTRL_RATE=10:100000:1000000 ./test/MatrixTranspose"
|
||||
# SYS test
|
||||
export ROCTRACER_DOMAIN="sys:roctx"
|
||||
eval_test "tool SYS test" ./test/MatrixTranspose
|
||||
export ROCTRACER_DOMAIN="sys:hsa:roctx"
|
||||
eval_test "tool SYS/HSA test" ./test/MatrixTranspose
|
||||
# Tracing control <delay:length:rate>
|
||||
export ROCTRACER_DOMAIN="hip"
|
||||
eval_test "tool period test" "ROCP_CTRL_RATE=10:100000:1000000 ./test/MatrixTranspose"
|
||||
eval_test "tool flushing test" "ROCP_FLUSH_RATE=100000 ./test/MatrixTranspose"
|
||||
|
||||
# HSA test
|
||||
export ROCTRACER_DOMAIN="hsa"
|
||||
|
||||
+124
-96
@@ -37,9 +37,7 @@ THE SOFTWARE.
|
||||
#include <inc/roctracer_hsa.h>
|
||||
#include <inc/roctracer_hip.h>
|
||||
#include <inc/roctracer_hcc.h>
|
||||
#ifdef KFD_WRAPPER
|
||||
#include <inc/roctracer_kfd.h>
|
||||
#endif
|
||||
#include <inc/ext/hsa_rt_utils.hpp>
|
||||
#include <src/core/loader.h>
|
||||
#include <src/core/trace_buffer.h>
|
||||
@@ -59,9 +57,12 @@ THE SOFTWARE.
|
||||
} \
|
||||
} while (0)
|
||||
|
||||
#ifndef onload_debug
|
||||
#define onload_debug false
|
||||
#endif
|
||||
#define ONLOAD_TRACE(str) \
|
||||
if (getenv("ROCP_ONLOAD_TRACE")) do { \
|
||||
std::cout << "PID(" << GetPid() << "): TRACER_TOOL::" << __FUNCTION__ << " " << str << std::endl << std::flush; \
|
||||
} while(0);
|
||||
#define ONLOAD_TRACE_BEG() ONLOAD_TRACE("begin")
|
||||
#define ONLOAD_TRACE_END() ONLOAD_TRACE("end")
|
||||
|
||||
typedef hsa_rt_utils::Timer::timestamp_t timestamp_t;
|
||||
hsa_rt_utils::Timer* timer = NULL;
|
||||
@@ -76,6 +77,7 @@ bool trace_hip_activity = false;
|
||||
bool trace_kfd = false;
|
||||
|
||||
LOADER_INSTANTIATE();
|
||||
TRACE_BUFFER_INSTANTIATE();
|
||||
|
||||
// Global output file handle
|
||||
FILE* roctx_file_handle = NULL;
|
||||
@@ -121,19 +123,19 @@ void* control_thr_fun(void*) {
|
||||
const uint32_t len_us = control_len_us % 1000000;
|
||||
const uint32_t dist_sec = control_dist_us / 1000000;
|
||||
const uint32_t dist_us = control_dist_us % 1000000;
|
||||
bool start = true;
|
||||
bool to_start = true;
|
||||
|
||||
sleep(delay_sec);
|
||||
usleep(delay_us);
|
||||
|
||||
while (1) {
|
||||
if (start) {
|
||||
start = false;
|
||||
if (to_start) {
|
||||
to_start = false;
|
||||
roctracer_start();
|
||||
sleep(len_sec);
|
||||
usleep(len_us);
|
||||
} else {
|
||||
start = true;
|
||||
to_start = true;
|
||||
roctracer_stop();
|
||||
sleep(dist_sec);
|
||||
usleep(dist_us);
|
||||
@@ -141,6 +143,20 @@ void* control_thr_fun(void*) {
|
||||
}
|
||||
}
|
||||
|
||||
// Flushing control thread
|
||||
uint32_t control_flush_us = 0;
|
||||
void* flush_thr_fun(void*) {
|
||||
const uint32_t dist_sec = control_flush_us / 1000000;
|
||||
const uint32_t dist_us = control_flush_us % 1000000;
|
||||
|
||||
while (1) {
|
||||
sleep(dist_sec);
|
||||
usleep(dist_us);
|
||||
roctracer_flush_activity();
|
||||
roctracer::TraceBufferBase::FlushAll();
|
||||
}
|
||||
}
|
||||
|
||||
///////////////////////////////////////////////////////////////////////////////////////////////////////
|
||||
// rocTX annotation tracing
|
||||
|
||||
@@ -187,22 +203,17 @@ void roctx_api_callback(
|
||||
roctx_callback_fun(domain, cid, GetTid(), data->args.message);
|
||||
}
|
||||
|
||||
// Start/Stop callbacks
|
||||
void roctx_range_stack_callback(const roctx_range_data_t* data, void* arg) {
|
||||
const bool* is_stop_ptr = (bool*)arg;
|
||||
const uint32_t cid = (*is_stop_ptr == true) ? ROCTX_API_ID_roctxRangePop : ROCTX_API_ID_roctxRangePushA;
|
||||
const char* message = (*is_stop_ptr == true) ? NULL : data->message;
|
||||
roctx_callback_fun(ACTIVITY_DOMAIN_ROCTX, cid, data->tid, message);
|
||||
// rocTX Start/Stop callbacks
|
||||
void roctx_range_start_callback(const roctx_range_data_t* data, void* arg) {
|
||||
roctx_callback_fun(ACTIVITY_DOMAIN_ROCTX, ROCTX_API_ID_roctxRangePushA, data->tid, data->message);
|
||||
}
|
||||
void stop_callback() {
|
||||
bool is_stop = true;
|
||||
roctracer::RocTxLoader::Instance().RangeStackIterate(roctx_range_stack_callback, (void*)&is_stop);
|
||||
}
|
||||
void start_callback() {
|
||||
bool is_stop = false;
|
||||
roctracer::RocTxLoader::Instance().RangeStackIterate(roctx_range_stack_callback, (void*)&is_stop);
|
||||
void roctx_range_stop_callback(const roctx_range_data_t* data, void* arg) {
|
||||
roctx_callback_fun(ACTIVITY_DOMAIN_ROCTX, ROCTX_API_ID_roctxRangePop, data->tid, NULL);
|
||||
}
|
||||
void start_callback() { roctracer::RocTxLoader::Instance().RangeStackIterate(roctx_range_start_callback, NULL); }
|
||||
void stop_callback() { roctracer::RocTxLoader::Instance().RangeStackIterate(roctx_range_stop_callback, NULL); }
|
||||
|
||||
// rocTX buffer flush function
|
||||
void roctx_flush_cb(roctx_trace_entry_t* entry) {
|
||||
std::ostringstream os;
|
||||
os << entry->timestamp << " " << entry->pid << ":" << entry->tid << " " << entry->cid;
|
||||
@@ -270,6 +281,9 @@ void hsa_activity_callback(
|
||||
index++;
|
||||
}
|
||||
|
||||
///////////////////////////////////////////////////////////////////////////////////////////////////////
|
||||
// HIP API tracing
|
||||
|
||||
struct hip_api_trace_entry_t {
|
||||
uint32_t valid;
|
||||
uint32_t type;
|
||||
@@ -284,9 +298,6 @@ struct hip_api_trace_entry_t {
|
||||
void* ptr;
|
||||
};
|
||||
|
||||
///////////////////////////////////////////////////////////////////////////////////////////////////////
|
||||
// HIP API tracing
|
||||
|
||||
void hip_api_flush_cb(hip_api_trace_entry_t* entry);
|
||||
roctracer::TraceBuffer<hip_api_trace_entry_t>::flush_prm_t hip_flush_prm[1] = {{0, hip_api_flush_cb}};
|
||||
roctracer::TraceBuffer<hip_api_trace_entry_t> hip_api_trace_buffer("HIP", 0x200000, hip_flush_prm, 1);
|
||||
@@ -437,7 +448,7 @@ void hcc_activity_callback(const char* begin, const char* end, void* arg) {
|
||||
// KFD API tracing
|
||||
|
||||
// KFD API callback function
|
||||
#ifdef KFD_WRAPPER
|
||||
static thread_local bool in_kfd_api_callback = false;
|
||||
void kfd_api_callback(
|
||||
uint32_t domain,
|
||||
uint32_t cid,
|
||||
@@ -445,6 +456,8 @@ void kfd_api_callback(
|
||||
void* arg)
|
||||
{
|
||||
(void)arg;
|
||||
if (in_kfd_api_callback) return;
|
||||
in_kfd_api_callback = true;
|
||||
const kfd_api_data_t* data = reinterpret_cast<const kfd_api_data_t*>(callback_data);
|
||||
if (data->phase == ACTIVITY_API_PHASE_ENTER) {
|
||||
kfd_begin_timestamp = timer->timestamp_fn_ns();
|
||||
@@ -454,8 +467,8 @@ void kfd_api_callback(
|
||||
os << kfd_begin_timestamp << ":" << end_timestamp << " " << GetPid() << ":" << GetTid() << " " << kfd_api_data_pair_t(cid, *data);
|
||||
fprintf(kfd_api_file_handle, "%s\n", os.str().c_str());
|
||||
}
|
||||
in_kfd_api_callback = false;
|
||||
}
|
||||
#endif
|
||||
|
||||
///////////////////////////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
@@ -526,10 +539,56 @@ void close_output_file(FILE* file_handle) {
|
||||
if ((file_handle != NULL) && (file_handle != stdout)) fclose(file_handle);
|
||||
}
|
||||
|
||||
// tool unload method
|
||||
void tool_unload() {
|
||||
static bool is_unloaded = false;
|
||||
ONLOAD_TRACE("begin, unloaded(" << is_unloaded << ")");
|
||||
|
||||
if (is_unloaded == true) return;
|
||||
is_unloaded = true;
|
||||
|
||||
roctracer_unload();
|
||||
|
||||
if (trace_roctx) {
|
||||
ROCTRACER_CALL(roctracer_disable_domain_callback(ACTIVITY_DOMAIN_ROCTX));
|
||||
|
||||
roctx_trace_buffer.Flush();
|
||||
close_output_file(roctx_file_handle);
|
||||
}
|
||||
if (trace_hsa_api) {
|
||||
ROCTRACER_CALL(roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HSA_API));
|
||||
|
||||
hsa_api_trace_buffer.Flush();
|
||||
close_output_file(hsa_api_file_handle);
|
||||
}
|
||||
if (trace_hsa_activity) {
|
||||
ROCTRACER_CALL(roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HSA_OPS));
|
||||
|
||||
close_output_file(hsa_async_copy_file_handle);
|
||||
}
|
||||
if (trace_hip_api || trace_hip_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());
|
||||
ROCTRACER_CALL(roctracer_close_pool());
|
||||
|
||||
hip_api_trace_buffer.Flush();
|
||||
close_output_file(hip_api_file_handle);
|
||||
close_output_file(hcc_activity_file_handle);
|
||||
}
|
||||
|
||||
if (trace_kfd) {
|
||||
ROCTRACER_CALL(roctracer_disable_domain_callback(ACTIVITY_DOMAIN_KFD_API));
|
||||
fclose(kfd_api_file_handle);
|
||||
}
|
||||
ONLOAD_TRACE_END();
|
||||
}
|
||||
|
||||
// HSA-runtime tool on-load method
|
||||
extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
|
||||
const char* const* failed_tool_names) {
|
||||
if (onload_debug) { printf("TOOL OnLoad\n"); fflush(stdout); }
|
||||
const char* const* failed_tool_names) {
|
||||
ONLOAD_TRACE_BEG();
|
||||
timer = new hsa_rt_utils::Timer(table->core_->hsa_system_get_info_fn);
|
||||
|
||||
// Output file
|
||||
@@ -660,7 +719,7 @@ extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
|
||||
for (unsigned i = 0; i < hsa_api_vec.size(); ++i) {
|
||||
uint32_t cid = HSA_API_ID_NUMBER;
|
||||
const char* api = hsa_api_vec[i].c_str();
|
||||
ROCTRACER_CALL(roctracer_op_code(ACTIVITY_DOMAIN_HSA_API, api, &cid));
|
||||
ROCTRACER_CALL(roctracer_op_code(ACTIVITY_DOMAIN_HSA_API, api, &cid, NULL));
|
||||
ROCTRACER_CALL(roctracer_enable_op_callback(ACTIVITY_DOMAIN_HSA_API, cid, hsa_api_callback, NULL));
|
||||
printf(" %s", api);
|
||||
}
|
||||
@@ -714,11 +773,9 @@ extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
|
||||
uint32_t ctrl_delay = 0;
|
||||
uint32_t ctrl_len = 0;
|
||||
uint32_t ctrl_rate = 0;
|
||||
int ret = sscanf(ctrl_str, "%d:%d:%d", &ctrl_delay, &ctrl_len, &ctrl_rate);
|
||||
if (ret != 3) {
|
||||
fprintf(stderr, "ROCTracer: control rate value invalid 'delay:length:rate': '%s'\n", ctrl_str);
|
||||
abort();
|
||||
}
|
||||
|
||||
sscanf(ctrl_str, "%d:%d:%d", &ctrl_delay, &ctrl_len, &ctrl_rate);
|
||||
|
||||
if (ctrl_len > ctrl_rate) {
|
||||
fprintf(stderr, "ROCTracer: control length value (%u) > rate value (%u)\n", ctrl_len, ctrl_rate);
|
||||
abort();
|
||||
@@ -727,18 +784,36 @@ extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
|
||||
control_len_us = ctrl_len;
|
||||
control_delay_us = ctrl_delay;
|
||||
|
||||
fprintf(stdout, "ROCTracer: trace control: delay(%uus), length(%uus), rate(%uus)\n", ctrl_delay, ctrl_len, ctrl_rate); fflush(stdout);
|
||||
|
||||
roctracer_stop();
|
||||
|
||||
if (ctrl_delay != UINT32_MAX) {
|
||||
fprintf(stdout, "ROCTracer: trace control: delay(%uus), length(%uus), rate(%uus)\n", ctrl_delay, ctrl_len, ctrl_rate); fflush(stdout);
|
||||
pthread_t thread;
|
||||
pthread_attr_t attr;
|
||||
int err = pthread_attr_init(&attr);
|
||||
if (err) { errno = err; perror("pthread_attr_init"); abort(); }
|
||||
err = pthread_create(&thread, &attr, control_thr_fun, NULL);
|
||||
} else {
|
||||
fprintf(stdout, "ROCTracer: trace start disabled\n"); fflush(stdout);
|
||||
}
|
||||
}
|
||||
|
||||
const char* flush_str = getenv("ROCP_FLUSH_RATE");
|
||||
if (flush_str != NULL) {
|
||||
sscanf(flush_str, "%d", &control_flush_us);
|
||||
if (control_flush_us == 0) {
|
||||
fprintf(stderr, "ROCTracer: control flush rate bad value\n");
|
||||
abort();
|
||||
}
|
||||
|
||||
fprintf(stdout, "ROCTracer: trace control flush rate(%uus)\n", control_flush_us); fflush(stdout);
|
||||
pthread_t thread;
|
||||
pthread_attr_t attr;
|
||||
int err = pthread_attr_init(&attr);
|
||||
if (err) { errno = err; perror("pthread_attr_init"); abort(); }
|
||||
err = pthread_create(&thread, &attr, control_thr_fun, NULL);
|
||||
err = pthread_create(&thread, &attr, flush_thr_fun, NULL);
|
||||
}
|
||||
|
||||
#ifdef KFD_WRAPPER
|
||||
// Enable KFD API callbacks/activity
|
||||
if (trace_kfd) {
|
||||
kfd_api_file_handle = open_output_file(output_prefix, "kfd_api_trace.txt");
|
||||
@@ -750,7 +825,7 @@ extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
|
||||
for (unsigned i = 0; i < kfd_api_vec.size(); ++i) {
|
||||
uint32_t cid = KFD_API_ID_NUMBER;
|
||||
const char* api = kfd_api_vec[i].c_str();
|
||||
ROCTRACER_CALL(roctracer_op_code(ACTIVITY_DOMAIN_KFD_API, api, &cid));
|
||||
ROCTRACER_CALL(roctracer_op_code(ACTIVITY_DOMAIN_KFD_API, api, &cid, NULL));
|
||||
ROCTRACER_CALL(roctracer_enable_op_callback(ACTIVITY_DOMAIN_KFD_API, cid, kfd_api_callback, NULL));
|
||||
printf(" %s", api);
|
||||
}
|
||||
@@ -759,70 +834,23 @@ extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
|
||||
}
|
||||
printf(")\n");
|
||||
}
|
||||
#endif
|
||||
|
||||
if (onload_debug) { printf("TOOL OnLoad end\n"); fflush(stdout); }
|
||||
return roctracer_load(table, runtime_version, failed_tool_count, failed_tool_names);
|
||||
}
|
||||
|
||||
// tool unload method
|
||||
void tool_unload(bool destruct) {
|
||||
static bool is_unloaded = false;
|
||||
|
||||
if (onload_debug) { printf("TOOL tool_unload (%d, %d)\n", (int)destruct, (int)is_unloaded); fflush(stdout); }
|
||||
if (destruct == false) return;
|
||||
if (is_unloaded == true) return;
|
||||
is_unloaded = true;
|
||||
roctracer_unload(destruct);
|
||||
|
||||
if (trace_roctx) {
|
||||
ROCTRACER_CALL(roctracer_disable_domain_callback(ACTIVITY_DOMAIN_ROCTX));
|
||||
|
||||
roctx_trace_buffer.Flush();
|
||||
close_output_file(roctx_file_handle);
|
||||
}
|
||||
if (trace_hsa_api) {
|
||||
ROCTRACER_CALL(roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HSA_API));
|
||||
|
||||
hsa_api_trace_buffer.Flush();
|
||||
close_output_file(hsa_api_file_handle);
|
||||
}
|
||||
if (trace_hsa_activity) {
|
||||
ROCTRACER_CALL(roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HSA_OPS));
|
||||
|
||||
close_output_file(hsa_async_copy_file_handle);
|
||||
}
|
||||
if (trace_hip_api || trace_hip_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());
|
||||
ROCTRACER_CALL(roctracer_close_pool());
|
||||
|
||||
hip_api_trace_buffer.Flush();
|
||||
close_output_file(hip_api_file_handle);
|
||||
close_output_file(hcc_activity_file_handle);
|
||||
}
|
||||
|
||||
if (trace_kfd) {
|
||||
ROCTRACER_CALL(roctracer_disable_domain_callback(ACTIVITY_DOMAIN_KFD_API));
|
||||
fclose(kfd_api_file_handle);
|
||||
}
|
||||
if (onload_debug) { printf("TOOL tool_unload end\n"); fflush(stdout); }
|
||||
roctracer::TraceBufferBase::StartWorkerThreadAll();
|
||||
const bool ret = roctracer_load();
|
||||
ONLOAD_TRACE_END();
|
||||
return ret;
|
||||
}
|
||||
|
||||
// HSA-runtime on-unload method
|
||||
extern "C" PUBLIC_API void OnUnload() {
|
||||
if (onload_debug) { printf("TOOL OnUnload\n"); fflush(stdout); }
|
||||
tool_unload(false);
|
||||
if (onload_debug) { printf("TOOL OnUnload end\n"); fflush(stdout); }
|
||||
ONLOAD_TRACE("");
|
||||
}
|
||||
|
||||
extern "C" CONSTRUCTOR_API void constructor() {
|
||||
if (onload_debug) { printf("TOOL constructor ...end\n"); fflush(stdout); }
|
||||
ONLOAD_TRACE("")
|
||||
}
|
||||
extern "C" DESTRUCTOR_API void destructor() {
|
||||
if (onload_debug) { printf("TOOL destructor\n"); fflush(stdout); }
|
||||
tool_unload(true);
|
||||
if (onload_debug) { printf("TOOL destructor end\n"); fflush(stdout); }
|
||||
ONLOAD_TRACE_BEG();
|
||||
tool_unload();
|
||||
ONLOAD_TRACE_END();
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user