AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Handle hsa_queue_destroy after finalization (#679)

Browse Source 
* Handle hsa_queue_destroy after finalization

- fixes issue where hsa_queue_destroy(...) is invoked after rocprofiler-sdk has finalized
- hsa::get_queue_controller() returns pointer
- if queue controller is a null pointer, skip invoking QueueController::destroy_queue

* Update HIP/HSA/marker update_table logging

* Update rocprofv3 tests

- remove HSA_TOOLS_LIB env variable
- remove setting ROCPROFILER_LOG_LEVEL env variable
- add timeouts to tests which are missing them

* Disable thread sanitizer deadlock detection

* Update CI workflow

- rename vega20-ubuntu job to core-ci
- enable navi32 in core-ci and sanitizers

* Update run-ci.py

- set gcovr html medium and high threshold

* Update lib/rocprofiler-sdk/hsa/queue_controller.cpp

- remove this capture from enable/disable serialization

* Update lib/rocprofiler-sdk/hsa/{hsa_barrier,profile_serializer}.*

- hsa_barrier::set_barrier accepts const-ref to queue map
- profile_serializer::enable and profile_serializer::disable accept const-ref to queue map

* Logging for HIP/HSA/marker/profile_serializer

* Logging for HIP/HSA/marker/queue_controller

* Improve test_retired_correlation_ids asserts

* Fix tests/counter-collection/validate.py

- scale expected SQ_WAVES counter value based on warp size of GPU

* Tweak github comment for code coverage

* Remove gcovr html high/medium threshold args

* Fix tests/counter-collection/validate.py

- round before casting to int in test_counter_values

* operator bool for profile_serializer

- only wait on CV if profile_serializer is used

* Logging updates (profile_serializer + code_object)

* Update counter-collection validate.py

* QueueController does not wait on CV if finalizing/finalized

* Update CI workflow

- remove navi32 from core job

* Improve HIP/HSA/marker tracing get_functor/functor

- remove lambda wrapper around functor

* Update lib/rocprofiler-sdk/hsa/queue_controller.cpp

- do not acquire cvmutex lock during finalization

* Update lib/rocprofiler-sdk/hsa/hsa_barrier.*

- move ctor and dtor to implementation
- skip signal store screlease and destroy if already finalized

* Update CI workflow

- remove navi32 runners

* bwelton fixes for hangs

* CMake improvements + simplified demangle

- remove amd-comgr from common target (and thus removed from roctx DT_NEEDED)

---------

Co-authored-by: Benjamin Welton <bewelton@amd.com>
This commit is contained in:
Jonathan R. Madsen
2024-03-21 17:52:15 -05:00
committed by GitHub
parent 78939e705a
commit 2f9b1767e9
48 changed files with 494 additions and 498 deletions
Download Patch File Download Diff File Expand all files Collapse all files
.github/workflows/continuous_integration.yml
+20 -77
View File
@@ -24,28 +24,16 @@ env:
PATH: "/usr/bin:$PATH"
jobs:
# Changed job name from mi200-ubuntu to vega20-ubuntu
# TODO: Change it back when re-enabling on mi200
vega20-ubuntu:
core:
# See: https://docs.github.com/en/free-pro-team@latest/actions/learn-github-actions/managing-complex-workflows#using-a-build-matrix
strategy:
fail-fast: true
max-parallel: 4
matrix:
include:
# Run job on vega20 instead of mi200 as mi200 is unstable, as per ammar's instructions.
# TODO: Change it back when re-enabling on mi200
# - os: 'ubuntu-22.04'
# runner: 'renderD131'
# device: '/renderD131'
# build-type: 'RelWithDebInfo'
# ci-flags: '--linter clang-tidy'
# name-tag: ''
- os: 'ubuntu-22.04'
runner: 'vega20'
build-type: 'RelWithDebInfo'
ci-flags: '--linter clang-tidy'
name-tag: ''
os: ['ubuntu-22.04']
runner: ['vega20']
build-type: ['RelWithDebInfo']
ci-flags: ['--linter clang-tidy']
name-tag: ['']
runs-on: ${{ matrix.runner }}
@@ -53,16 +41,6 @@ jobs:
env:
GIT_DISCOVERY_ACROSS_FILESYSTEM: 1
# TODO: Uncomment this when re-enabling tests on the mi200 as it contains --memory and --cpus flag for the mi200. Remove these 2 options when running on vega20.
# vega20 machine only has 24 cpus available.
# container:
# # TODO(aelwazir): Adding multiple os distros
# image: evuedsoacr.azurecr.io/ml-lang-comp-us/rocm-tools/rocm-rocprofiler-sdk-staging-base:latest
# options: --memory=128g --cpus=32 --ipc=host --device=/dev/kfd --group-add video --cap-add=SYS_PTRACE --cap-add CAP_SYS_ADMIN --security-opt seccomp=unconfined
# credentials:
# username: ${{ secrets.DOCKER_USERNAME }}
# password: ${{ secrets.DOCKER_PASSWORD }}
container:
# TODO(aelwazir): Adding multiple os distros
image: evuedsoacr.azurecr.io/ml-lang-comp-us/rocm-tools/rocm-rocprofiler-sdk-staging-base:latest
@@ -99,8 +77,6 @@ jobs:
- name: Configure, Build, and Test
timeout-minutes: 30
shell: bash
# Replaced 'mi200' with '${{ matrix.runner }}' when disabling jobs on mi200 and running it on vega20.
# TODO: Change it back when re-enabling on mi200
run:
python3 ./source/scripts/run-ci.py -B build
--name ${{ github.repository }}-${{ github.ref_name }}-${{ matrix.runner }}-${{ matrix.os }}${{ matrix.name-tag }}
@@ -112,7 +88,7 @@ jobs:
-DROCPROFILER_DEP_ROCMCORE=ON
-DROCPROFILER_BUILD_DOCS=ON
-DCMAKE_BUILD_TYPE=${{ matrix.build-type }}
-DCMAKE_INSTALL_PREFIX=/opt/rocprofiler/v2
-DCMAKE_INSTALL_PREFIX=/opt/rocprofiler-sdk
-DCPACK_GENERATOR='DEB;RPM;TGZ'
-DCPACK_PACKAGING_INSTALL_PREFIX="$(realpath /opt/rocm)"
-DPython3_EXECUTABLE=$(which python3)
@@ -131,9 +107,9 @@ jobs:
timeout-minutes: 10
shell: bash
run: |
CMAKE_PREFIX_PATH=/opt/rocprofiler/v2 cmake -B build-samples samples
CMAKE_PREFIX_PATH=/opt/rocprofiler/v2 cmake -B build-tests tests
export LD_LIBRARY_PATH=/opt/rocprofiler/v2/lib:${LD_LIBRARY_PATH}
CMAKE_PREFIX_PATH=/opt/rocprofiler-sdk cmake -B build-samples samples
CMAKE_PREFIX_PATH=/opt/rocprofiler-sdk cmake -B build-tests tests
export LD_LIBRARY_PATH=/opt/rocprofiler-sdk/lib:${LD_LIBRARY_PATH}
cmake --build build-samples --target all --parallel 8
cmake --build build-tests --target all --parallel 8
ctest --test-dir build-samples --output-on-failure
@@ -171,7 +147,6 @@ jobs:
code-coverage:
strategy:
fail-fast: true
max-parallel: 4
matrix:
# TODO: Change it back when re-enabling on mi200
include:
@@ -263,8 +238,6 @@ jobs:
- name: Configure, Build, and Test (Total Code Coverage)
timeout-minutes: 30
shell: bash
# Replaced 'mi200' with '${{ matrix.runner }}' when disabling jobs on mi200 and running it on vega20.
# TODO: Change it back when re-enabling on mi200
run:
python3 ./source/scripts/run-ci.py -B build
--name ${{ github.repository }}-${{ github.ref_name }}-${{ matrix.runner }}-${{ matrix.os }}-codecov
@@ -337,18 +310,19 @@ jobs:
echo -e "\n.codecov:"
ls -la .codecov
get-md-contents() { cat .codecov/${1}.png.md .codecov/${1}.md; }
get-base-md-contents() { cat .codecov/${1}.png.md; }
get-full-md-contents() { cat .codecov/${1}.png.md .codecov/${1}.md; }
cat << EOF > .codecov/report.md
# Code Coverage Report
## Tests Only
$(get-md-contents tests)
$(get-base-md-contents tests)
## Samples Only
$(get-md-contents samples)
$(get-base-md-contents samples)
## Tests + Samples
$(get-md-contents all)
$(get-full-md-contents all)
EOF
@@ -405,42 +379,11 @@ jobs:
strategy:
fail-fast: false
matrix:
include:
# Run job on vega20 instead of mi200 as mi200 is unstable, as per ammar's instructions.
# TODO: Change it back when re-enabling on mi200
- os: 'ubuntu-22.04'
runner: 'vega20'
build-type: 'RelWithDebInfo'
ci-flags: ''
sanitizer: 'AddressSanitizer'
- os: 'ubuntu-22.04'
runner: 'vega20'
build-type: 'RelWithDebInfo'
ci-flags: ''
sanitizer: 'ThreadSanitizer'
- os: 'ubuntu-22.04'
runner: 'vega20'
build-type: 'RelWithDebInfo'
ci-flags: ''
sanitizer: 'LeakSanitizer'
# - os: 'ubuntu-22.04'
# runner: 'renderD131'
# device: '/renderD131'
# build-type: 'RelWithDebInfo'
# ci-flags: ''
# sanitizer: 'AddressSanitizer'
# - os: 'ubuntu-22.04'
# runner: 'renderD131'
# device: '/renderD131'
# build-type: 'RelWithDebInfo'
# ci-flags: ''
# sanitizer: 'ThreadSanitizer'
# - os: 'ubuntu-22.04'
# runner: 'renderD131'
# device: '/renderD131'
# build-type: 'RelWithDebInfo'
# ci-flags: ''
# sanitizer: 'LeakSanitizer'
os: ['ubuntu-22.04']
runner: ['vega20']
build-type: ['RelWithDebInfo']
sanitizer: ['AddressSanitizer', 'ThreadSanitizer', 'LeakSanitizer']
ci-flags: ['']
runs-on: ${{ matrix.runner }}
cmake/rocprofiler_config_interfaces.cmake
+30 -8
View File
@@ -79,6 +79,14 @@ foreach(_LIB dl rt)
endif()
endforeach()
# ----------------------------------------------------------------------------------------#
#
# atomic library
#
# ----------------------------------------------------------------------------------------#
target_link_libraries(rocprofiler-atomic INTERFACE atomic)
# ----------------------------------------------------------------------------------------#
#
# filesystem library
@@ -183,6 +191,28 @@ find_library(
target_link_libraries(rocprofiler-hsa-aql INTERFACE ${hsa-amd-aqlprofile64_library})
# ----------------------------------------------------------------------------------------#
#
# HSAKMT
#
# ----------------------------------------------------------------------------------------#
find_package(
hsakmt
REQUIRED
CONFIG
HINTS
${rocm_version_DIR}
${ROCM_PATH}
PATHS
${rocm_version_DIR}
${ROCM_PATH}
PATH_SUFFIXES
lib/cmake/hsakmt)
target_link_libraries(rocprofiler-hsakmt INTERFACE hsakmt::hsakmt)
rocprofiler_config_nolink_target(rocprofiler-hsakmt-nolink hsakmt::hsakmt)
# ----------------------------------------------------------------------------------------#
#
# drm
@@ -218,11 +248,3 @@ find_library(
target_include_directories(rocprofiler-drm SYSTEM INTERFACE ${drm_INCLUDE_DIR}
${xf86drm_INCLUDE_DIR})
target_link_libraries(rocprofiler-drm INTERFACE ${drm_LIBRARY} ${drm_amdgpu_LIBRARY})
# ----------------------------------------------------------------------------------------#
#
# interface targets which emulate another interface target but do not link to the library.
# E.g. rocprofiler-hip-nolink has the include directories, compile definitions, and
# compile options of rocprofiler-hip but does not link to the HIP runtime library
#
# ----------------------------------------------------------------------------------------#
cmake/rocprofiler_interfaces.cmake
+20 -4
View File
@@ -36,26 +36,42 @@ rocprofiler_add_interface_library(rocprofiler-stack-protector
rocprofiler_add_interface_library(rocprofiler-memcheck INTERFACE INTERNAL)
#
# interfaces for libraries
# interfaces for libraries (general)
#
rocprofiler_add_interface_library(rocprofiler-dl
"Build flags for dynamic linking library" INTERNAL)
rocprofiler_add_interface_library(rocprofiler-rt "Build flags for runtime library"
INTERNAL)
rocprofiler_add_interface_library(rocprofiler-hip "HIP library" INTERNAL)
rocprofiler_add_interface_library(rocprofiler-hsa-runtime "HSA runtime library" INTERNAL)
rocprofiler_add_interface_library(rocprofiler-amd-comgr "AMD comgr library" INTERNAL)
rocprofiler_add_interface_library(rocprofiler-atomic "atomic library" INTERNAL)
rocprofiler_add_interface_library(rocprofiler-gtest "Google Test library" INTERNAL)
rocprofiler_add_interface_library(rocprofiler-glog "Google Log library" INTERNAL)
rocprofiler_add_interface_library(rocprofiler-fmt "C++ format string library" INTERNAL)
rocprofiler_add_interface_library(rocprofiler-cxx-filesystem "C++ filesystem library"
INTERNAL)
rocprofiler_add_interface_library(rocprofiler-ptl "Parallel Tasking Library" INTERNAL)
#
# interface for libraries (ROCm-specific)
#
rocprofiler_add_interface_library(rocprofiler-hip "HIP library" INTERNAL)
rocprofiler_add_interface_library(rocprofiler-hsa-runtime "HSA runtime library" INTERNAL)
rocprofiler_add_interface_library(rocprofiler-amd-comgr "AMD comgr library" INTERNAL)
rocprofiler_add_interface_library(rocprofiler-hsa-aql "AQL library" INTERNAL)
rocprofiler_add_interface_library(rocprofiler-hsakmt "HSAKMT library for AMD KFD support"
INTERNAL)
rocprofiler_add_interface_library(rocprofiler-drm "drm (amdgpu) library" INTERNAL)
#
# "nolink" interface targets emulate another interface target but do not link to the
# library. E.g. rocprofiler-hip-nolink has the include directories, compile definitions,
# and compile options of rocprofiler-hip but does not link to the HIP runtime library
#
rocprofiler_add_interface_library(
rocprofiler-hip-nolink "rocprofiler-hip without linking to HIP library" IMPORTED)
rocprofiler_add_interface_library(
rocprofiler-hsa-runtime-nolink
"rocprofiler-hsa-runtime without linking to HSA library" IMPORTED)
rocprofiler_add_interface_library(
rocprofiler-hsakmt-nolink "rocprofiler-hsakmt without linking to HSAKMT library"
IMPORTED)
source/lib/common/CMakeLists.txt
+3 -5
View File
@@ -38,11 +38,9 @@ target_link_libraries(
$<BUILD_INTERFACE:rocprofiler::rocprofiler-glog>
$<BUILD_INTERFACE:rocprofiler::rocprofiler-fmt>
$<BUILD_INTERFACE:rocprofiler::rocprofiler-dl>
$<BUILD_INTERFACE:rocprofiler::rocprofiler-hip-nolink>
$<BUILD_INTERFACE:rocprofiler::rocprofiler-amd-comgr>
$<BUILD_INTERFACE:rocprofiler::rocprofiler-hsa-runtime-nolink>
$<BUILD_INTERFACE:rocprofiler::rocprofiler-ptl>
$<BUILD_INTERFACE:rocprofiler::rocprofiler-hsa-aql>
$<BUILD_INTERFACE:rocprofiler::rocprofiler-drm>)
$<BUILD_INTERFACE:rocprofiler::rocprofiler-atomic>
$<BUILD_INTERFACE:rocprofiler::rocprofiler-hsakmt-nolink>)
set_target_properties(rocprofiler-common-library PROPERTIES OUTPUT_NAME
rocprofiler-common)
source/lib/common/demangle.cpp
+7 -48
View File
@@ -22,10 +22,11 @@
#include "lib/common/demangle.hpp"
#include <amd_comgr/amd_comgr.h>
#include <glog/logging.h>
#include <cxxabi.h>
#include <cstdarg>
#include <cstdint>
#include <cstdio>
#include <cstring>
#include <iomanip>
@@ -34,18 +35,6 @@
#include <sstream>
#include <string>
#define amd_comgr_(call) \
do \
{ \
if(amd_comgr_status_t status = amd_comgr_##call; status != AMD_COMGR_STATUS_SUCCESS) \
{ \
const char* reason = ""; \
amd_comgr_status_string(status, &reason); \
fprintf(stderr, #call " failed: %s\n", reason); \
abort(); \
} \
} while(false)
namespace rocprofiler
{
namespace common
@@ -53,7 +42,6 @@ namespace common
std::string
cxa_demangle(std::string_view _mangled_name, int* _status)
{
constexpr size_t buffer_len = 4096;
// return the mangled since there is no buffer
if(_mangled_name.empty())
{
@@ -91,26 +79,14 @@ cxa_demangle(std::string_view _mangled_name, int* _status)
}
case -1:
{
char _msg[buffer_len];
::memset(_msg, '\0', buffer_len * sizeof(char));
::snprintf(_msg,
buffer_len,
"memory allocation failure occurred demangling %s",
_demangled_name.c_str());
::perror(_msg);
PLOG(ERROR) << "memory allocation failure occurred demangling " << _demangled_name;
break;
}
case -2: break;
case -3:
{
char _msg[buffer_len];
::memset(_msg, '\0', buffer_len * sizeof(char));
::snprintf(_msg,
buffer_len,
"Invalid argument in: (\"%s\", nullptr, nullptr, %p)",
_demangled_name.c_str(),
(void*) _status);
::perror(_msg);
PLOG(ERROR) << "Invalid argument in: (\"" << _demangled_name << "\", nullptr, nullptr, "
<< _status << ")";
break;
}
default: break;
@@ -130,27 +106,10 @@ cxx_demangle(std::string_view symbol)
{
int _status = 0;
auto demangled_str = cxa_demangle(symbol, &_status);
if(_status == 0)
{
return demangled_str;
}
amd_comgr_data_t mangled_data;
amd_comgr_(create_data(AMD_COMGR_DATA_KIND_BYTES, &mangled_data));
amd_comgr_(set_data(mangled_data, symbol.size(), symbol.data()));
if(_status == 0) return demangled_str;
amd_comgr_data_t demangled_data;
amd_comgr_(demangle_symbol_name(mangled_data, &demangled_data));
size_t demangled_size = 0;
amd_comgr_(get_data(demangled_data, &demangled_size, nullptr));
demangled_str.resize(demangled_size);
amd_comgr_(get_data(demangled_data, &demangled_size, demangled_str.data()));
amd_comgr_(release_data(mangled_data));
amd_comgr_(release_data(demangled_data));
return demangled_str;
return std::string{symbol};
}
// The function extracts the kernel name from
source/lib/rocprofiler-sdk-tool/CMakeLists.txt
+3 -7
View File
@@ -11,13 +11,9 @@ add_subdirectory(plugins)
target_link_libraries(
rocprofiler-sdk-tool
PRIVATE rocprofiler::rocprofiler-shared-library
rocprofiler::rocprofiler-headers
rocprofiler::rocprofiler-build-flags
rocprofiler::rocprofiler-memcheck
rocprofiler::rocprofiler-common-library
rocprofiler::rocprofiler-cxx-filesystem
atomic)
PRIVATE rocprofiler::rocprofiler-shared-library rocprofiler::rocprofiler-headers
rocprofiler::rocprofiler-build-flags rocprofiler::rocprofiler-memcheck
rocprofiler::rocprofiler-common-library)
set_target_properties(
rocprofiler-sdk-tool
source/lib/rocprofiler-sdk/CMakeLists.txt
+4 -3
View File
@@ -50,9 +50,10 @@ target_link_libraries(
PRIVATE rocprofiler::rocprofiler-build-flags
rocprofiler::rocprofiler-memcheck
rocprofiler::rocprofiler-common-library
rocprofiler::rocprofiler-cxx-filesystem
rocprofiler::rocprofiler-dl
rocprofiler::rocprofiler-amd-comgr)
rocprofiler::rocprofiler-amd-comgr
rocprofiler::rocprofiler-hsa-aql
rocprofiler::rocprofiler-drm
rocprofiler::rocprofiler-hsakmt-nolink)
target_compile_definitions(rocprofiler-object-library PRIVATE rocprofiler_EXPORTS=1)
source/lib/rocprofiler-sdk/aql/tests/aql_test.cpp
+12 -10
View File
@@ -61,11 +61,12 @@ CoreApiTable&
get_api_table()
{
static auto _v = []() {
auto val = CoreApiTable{};
val.hsa_iterate_agents_fn = hsa_iterate_agents;
val.hsa_agent_get_info_fn = hsa_agent_get_info;
val.hsa_queue_create_fn = hsa_queue_create;
val.hsa_queue_destroy_fn = hsa_queue_destroy;
auto val = CoreApiTable{};
val.hsa_iterate_agents_fn = hsa_iterate_agents;
val.hsa_agent_get_info_fn = hsa_agent_get_info;
val.hsa_queue_create_fn = hsa_queue_create;
val.hsa_queue_destroy_fn = hsa_queue_destroy;
val.hsa_signal_wait_relaxed_fn = hsa_signal_wait_relaxed;
return val;
}();
return _v;
@@ -102,7 +103,8 @@ test_init()
table.amd_ext_ = &get_ext_table();
table.core_ = &get_api_table();
agent::construct_agent_cache(&table);
hsa::get_queue_controller().init(get_api_table(), get_ext_table());
ASSERT_TRUE(hsa::get_queue_controller() != nullptr);
hsa::get_queue_controller()->init(get_api_table(), get_ext_table());
}
} // namespace rocprofiler
@@ -113,7 +115,7 @@ TEST(aql_profile, construct_packets)
{
ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS);
rocprofiler::test_init();
auto agents = rocprofiler::hsa::get_queue_controller().get_supported_agents();
auto agents = rocprofiler::hsa::get_queue_controller()->get_supported_agents();
ASSERT_GT(agents.size(), 0);
for(const auto& [_, agent] : agents)
{
@@ -129,7 +131,7 @@ TEST(aql_profile, too_many_counters)
{
ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS);
rocprofiler::test_init();
auto agents = rocprofiler::hsa::get_queue_controller().get_supported_agents();
auto agents = rocprofiler::hsa::get_queue_controller()->get_supported_agents();
ASSERT_GT(agents.size(), 0);
for(const auto& [_, agent] : agents)
{
@@ -157,7 +159,7 @@ TEST(aql_profile, packet_generation_single)
ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS);
rocprofiler::test_init();
auto agents = rocprofiler::hsa::get_queue_controller().get_supported_agents();
auto agents = rocprofiler::hsa::get_queue_controller()->get_supported_agents();
ASSERT_GT(agents.size(), 0);
for(const auto& [_, agent] : agents)
{
@@ -175,7 +177,7 @@ TEST(aql_profile, packet_generation_multi)
ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS);
rocprofiler::test_init();
auto agents = rocprofiler::hsa::get_queue_controller().get_supported_agents();
auto agents = rocprofiler::hsa::get_queue_controller()->get_supported_agents();
ASSERT_GT(agents.size(), 0);
for(const auto& [_, agent] : agents)
{
source/lib/rocprofiler-sdk/counters/core.cpp
+12 -10
View File
@@ -236,8 +236,8 @@ counter_callback_info::get_packet(std::unique_ptr<rocprofiler::hsa::AQLPacket>&
if(!ret_pkt)
{
// If we do not have a packet in the cache, create one.
ret_pkt =
profile->pkt_generator->construct_packet(hsa::get_queue_controller().get_ext_table());
ret_pkt = profile->pkt_generator->construct_packet(
CHECK_NOTNULL(hsa::get_queue_controller())->get_ext_table());
}
ret_pkt->before_krn_pkt.clear();
@@ -268,7 +268,7 @@ queue_cb(const context::context* ctx,
// and maybe adds barrier packets if the state is transitioning from serialized <->
// unserialized
auto maybe_add_serialization = [&](auto& gen_pkt) {
hsa::get_queue_controller().serializer().rlock([&](const auto& serializer) {
CHECK_NOTNULL(hsa::get_queue_controller())->serializer().rlock([&](const auto& serializer) {
for(auto& s_pkt : serializer.kernel_dispatch(queue))
{
gen_pkt->before_krn_pkt.push_back(s_pkt.ext_amd_aql_pm4);
@@ -392,8 +392,9 @@ completed_cb(const context::context* ctx,
if(!pkt) return;
hsa::get_queue_controller().serializer().wlock(
[&](auto& serializer) { serializer.kernel_completion_signal(session.queue); });
CHECK_NOTNULL(hsa::get_queue_controller())->serializer().wlock([&](auto& serializer) {
serializer.kernel_completion_signal(session.queue);
});
// We have no profile config, nothing to output.
if(!prof_config) return;
@@ -478,10 +479,10 @@ start_context(const context::context* ctx)
{
if(!ctx || !ctx->counter_collection) return;
auto& controller = hsa::get_queue_controller();
auto* controller = hsa::get_queue_controller();
bool already_enabled = true;
controller.enable_serialization();
CHECK_NOTNULL(controller)->enable_serialization();
ctx->counter_collection->enabled.wlock([&](auto& enabled) {
if(enabled) return;
already_enabled = false;
@@ -495,7 +496,7 @@ start_context(const context::context* ctx)
// Insert our callbacks into HSA Interceptor. This
// turns on counter instrumentation.
if(cb->queue_id != rocprofiler::hsa::ClientID{-1}) continue;
cb->queue_id = controller.add_callback(
cb->queue_id = controller->add_callback(
std::nullopt,
[=](const hsa::Queue& q,
const hsa::rocprofiler_packet& kern_pkt,
@@ -526,13 +527,14 @@ stop_context(const context::context* ctx)
{
if(!ctx || !ctx->counter_collection) return;
auto& controller = hsa::get_queue_controller();
auto* controller = hsa::get_queue_controller();
ctx->counter_collection->enabled.wlock([&](auto& enabled) {
if(!enabled) return;
enabled = false;
});
controller.disable_serialization();
if(controller) controller->disable_serialization();
}
bool
source/lib/rocprofiler-sdk/counters/dimensions.cpp
+5 -2
View File
@@ -56,7 +56,8 @@ getBlockDimensions(std::string_view agent, const Metric& metric)
std::vector<MetricDimension> ret;
for(const auto& [_, maybe_agent] : hsa::get_queue_controller().get_supported_agents())
for(const auto& [_, maybe_agent] :
CHECK_NOTNULL(hsa::get_queue_controller())->get_supported_agents())
{
if(maybe_agent.name() == agent)
{
@@ -106,7 +107,9 @@ get_dimension_cache()
* Fails if HSA is not loaded by retruning nothing. This should not remain after
* AQL is transistioned away from HSA.
*/
if(rocprofiler::hsa::get_queue_controller().get_supported_agents().empty())
if(CHECK_NOTNULL(rocprofiler::hsa::get_queue_controller())
->get_supported_agents()
.empty())
{
return {};
}
source/lib/rocprofiler-sdk/counters/tests/core.cpp
+24 -15
View File
@@ -90,11 +90,18 @@ CoreApiTable&
get_api_table()
{
static auto _v = []() {
auto val = CoreApiTable{};
val.hsa_iterate_agents_fn = hsa_iterate_agents;
val.hsa_agent_get_info_fn = hsa_agent_get_info;
val.hsa_queue_create_fn = hsa_queue_create;
val.hsa_queue_destroy_fn = hsa_queue_destroy;
auto val = CoreApiTable{};
val.hsa_iterate_agents_fn = hsa_iterate_agents;
val.hsa_agent_get_info_fn = hsa_agent_get_info;
val.hsa_queue_create_fn = hsa_queue_create;
val.hsa_queue_destroy_fn = hsa_queue_destroy;
val.hsa_signal_create_fn = hsa_signal_create;
val.hsa_signal_destroy_fn = hsa_signal_destroy;
val.hsa_signal_store_screlease_fn = hsa_signal_store_screlease;
val.hsa_signal_load_scacquire_fn = hsa_signal_load_scacquire;
val.hsa_signal_add_relaxed_fn = hsa_signal_add_relaxed;
val.hsa_signal_subtract_relaxed_fn = hsa_signal_subtract_relaxed;
val.hsa_signal_wait_relaxed_fn = hsa_signal_wait_relaxed;
return val;
}();
return _v;
@@ -131,7 +138,8 @@ test_init()
table.amd_ext_ = &get_ext_table();
table.core_ = &get_api_table();
agent::construct_agent_cache(&table);
hsa::get_queue_controller().init(get_api_table(), get_ext_table());
ASSERT_TRUE(hsa::get_queue_controller() != nullptr);
hsa::get_queue_controller()->init(get_api_table(), get_ext_table());
}
} // namespace
@@ -235,8 +243,8 @@ TEST(core, check_packet_generation)
{
ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS);
test_init();
auto agents = hsa::get_queue_controller().get_supported_agents();
ASSERT_TRUE(hsa::get_queue_controller() != nullptr);
auto agents = hsa::get_queue_controller()->get_supported_agents();
ASSERT_GT(agents.size(), 0);
for(const auto& [_, agent] : agents)
{
@@ -319,7 +327,7 @@ class FakeQueue : public Queue
{
public:
FakeQueue(const AgentCache& a, rocprofiler_queue_id_t id)
: Queue(a)
: Queue(a, get_api_table())
, _agent(a)
, _id(id)
{}
@@ -418,9 +426,10 @@ TEST(core, check_callbacks)
ctx.counter_collection = std::make_unique<rocprofiler::context::counter_collection_service>();
ctx.counter_collection->enabled.wlock([](auto& data) { data = true; });
auto agents = hsa::get_queue_controller().get_supported_agents();
ASSERT_TRUE(hsa::get_queue_controller() != nullptr);
auto agents = hsa::get_queue_controller()->get_supported_agents();
ASSERT_GT(agents.size(), 0);
hsa::get_queue_controller().disable_serialization();
hsa::get_queue_controller()->disable_serialization();
for(const auto& [_, agent] : agents)
{
@@ -540,7 +549,7 @@ TEST(core, destroy_counter_profile)
context::push_client(1);
ROCPROFILER_CALL(rocprofiler_create_context(&get_client_ctx()), "context creation failed");
auto agents = hsa::get_queue_controller().get_supported_agents();
auto agents = hsa::get_queue_controller()->get_supported_agents();
ASSERT_GT(agents.size(), 0);
for(const auto& [_, agent] : agents)
{
@@ -613,7 +622,7 @@ TEST(core, start_stop_buffered_ctx)
EXPECT_TRUE(found);
found = false;
hsa::get_queue_controller().iterate_callbacks([&](auto cid, const auto&) {
hsa::get_queue_controller()->iterate_callbacks([&](auto cid, const auto&) {
if(cid == ctx.counter_collection->callbacks.at(0)->queue_id)
{
found = true;
@@ -678,7 +687,7 @@ TEST(core, start_stop_callback_ctx)
EXPECT_TRUE(found);
found = false;
hsa::get_queue_controller().iterate_callbacks([&](auto cid, const auto&) {
hsa::get_queue_controller()->iterate_callbacks([&](auto cid, const auto&) {
if(cid == ctx.counter_collection->callbacks.at(0)->queue_id)
{
found = true;
@@ -707,7 +716,7 @@ TEST(core, public_api_iterate_agents)
registration::init_logging();
registration::set_init_status(-1);
context::push_client(1);
auto agents = hsa::get_queue_controller().get_supported_agents();
auto agents = hsa::get_queue_controller()->get_supported_agents();
for(const auto& [_, agent] : agents)
{
std::set<uint64_t> from_api;
source/lib/rocprofiler-sdk/counters/tests/dimension.cpp
+9 -7
View File
@@ -170,11 +170,12 @@ CoreApiTable&
get_api_table()
{
static auto _v = []() {
auto val = CoreApiTable{};
val.hsa_iterate_agents_fn = hsa_iterate_agents;
val.hsa_agent_get_info_fn = hsa_agent_get_info;
val.hsa_queue_create_fn = hsa_queue_create;
val.hsa_queue_destroy_fn = hsa_queue_destroy;
auto val = CoreApiTable{};
val.hsa_iterate_agents_fn = hsa_iterate_agents;
val.hsa_agent_get_info_fn = hsa_agent_get_info;
val.hsa_queue_create_fn = hsa_queue_create;
val.hsa_queue_destroy_fn = hsa_queue_destroy;
val.hsa_signal_wait_relaxed_fn = hsa_signal_wait_relaxed;
return val;
}();
return _v;
@@ -211,7 +212,8 @@ test_init()
table.amd_ext_ = &get_ext_table();
table.core_ = &get_api_table();
agent::construct_agent_cache(&table);
hsa::get_queue_controller().init(get_api_table(), get_ext_table());
ASSERT_TRUE(hsa::get_queue_controller() != nullptr);
hsa::get_queue_controller()->init(get_api_table(), get_ext_table());
}
} // namespace
@@ -221,7 +223,7 @@ TEST(dimension, block_dim_test)
ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS);
test_init();
auto agents = hsa::get_queue_controller().get_supported_agents();
auto agents = hsa::get_queue_controller()->get_supported_agents();
ASSERT_GT(agents.size(), 0);
for(const auto& [_, agent] : agents)
{
source/lib/rocprofiler-sdk/hip/defines.hpp
+2 -12
View File
@@ -88,14 +88,9 @@
template <typename RetT, typename... Args> \
static auto get_functor(RetT (*)(Args...)) \
{ \
if constexpr(std::is_void<RetT>::value) \
return [](Args... args) -> RetT { base_type::functor(args...); }; \
else \
return [](Args... args) -> RetT { return base_type::functor(args...); }; \
return &base_type::functor<RetT, Args...>; \
} \
\
static auto get_functor() { return get_functor(get_table_func()); } \
\
static std::vector<void*> as_arg_addr(callback_data_type) { return std::vector<void*>{}; } \
\
static std::vector<common::stringified_argument> as_arg_list(callback_data_type, int32_t) \
@@ -171,14 +166,9 @@
template <typename RetT, typename... Args> \
static auto get_functor(RetT (*)(Args...)) \
{ \
if constexpr(std::is_same<RetT, void>::value) \
return [](Args... args) -> RetT { base_type::functor(args...); }; \
else \
return [](Args... args) -> RetT { return base_type::functor(args...); }; \
return &base_type::functor<RetT, Args...>; \
} \
\
static auto get_functor() { return get_functor(get_table_func()); } \
\
static std::vector<void*> as_arg_addr(callback_data_type trace_data) \
{ \
return std::vector<void*>{ \
source/lib/rocprofiler-sdk/hip/hip.cpp
+11 -11
View File
@@ -28,6 +28,7 @@
#include "lib/rocprofiler-sdk/hip/details/ostream.hpp"
#include "lib/rocprofiler-sdk/hip/types.hpp"
#include "lib/rocprofiler-sdk/hip/utils.hpp"
#include "lib/rocprofiler-sdk/registration.hpp"
#include <rocprofiler-sdk/buffer.h>
#include <rocprofiler-sdk/callback_tracing.h>
@@ -230,14 +231,16 @@ populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx,
} // namespace
template <size_t TableIdx, size_t OpIdx>
template <typename... Args>
auto
hip_api_impl<TableIdx, OpIdx>::functor(Args&&... args)
template <typename RetT, typename... Args>
RetT
hip_api_impl<TableIdx, OpIdx>::functor(Args... args)
{
using info_type = hip_api_info<TableIdx, OpIdx>;
using callback_api_data_t = typename hip_domain_info<TableIdx>::callback_data_type;
using buffered_api_data_t = typename hip_domain_info<TableIdx>::buffered_data_type;
LOG_IF(INFO, registration::get_fini_status() != 0) << "Executing " << info_type::name;
auto thr_id = common::get_tid();
auto callback_contexts = std::vector<callback_context_data>{};
auto buffered_contexts = std::vector<buffered_context_data>{};
@@ -251,10 +254,10 @@ hip_api_impl<TableIdx, OpIdx>::functor(Args&&... args)
if(callback_contexts.empty() && buffered_contexts.empty())
{
auto _ret = exec(info_type::get_table_func(), std::forward<Args>(args)...);
if constexpr(!std::is_same<decltype(_ret), null_type>::value)
if constexpr(!std::is_void<RetT>::value)
return _ret;
else
return 0;
return;
}
constexpr auto ref_count = 2;
@@ -379,10 +382,7 @@ hip_api_impl<TableIdx, OpIdx>::functor(Args&&... args)
context::pop_latest_correlation_id(corr_id);
if constexpr(!std::is_same<decltype(_ret), null_type>::value)
return _ret;
else
return 0;
if constexpr(!std::is_void<RetT>::value) return _ret;
}
} // namespace hip
} // namespace rocprofiler
@@ -554,8 +554,6 @@ update_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
{
auto _info = hip_api_info<TableIdx, OpIdx>{};
LOG(INFO) << "updating table entry for " << _info.name;
// make sure we don't access a field that doesn't exist in input table
if(_info.offset() >= _orig->size) return;
@@ -564,6 +562,8 @@ update_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
_info.callback_domain_idx, _info.buffered_domain_idx, _info.operation_idx))
return;
LOG(INFO) << "updating table entry for " << _info.name;
// 1. get the sub-table containing the function pointer in original table
// 2. get reference to function pointer in sub-table in original table
// 3. update function pointer with wrapper
source/lib/rocprofiler-sdk/hip/hip.hpp
+2 -2
View File
@@ -74,8 +74,8 @@ struct hip_api_impl : hip_domain_info<TableIdx>
template <typename FuncT, typename... Args>
static auto exec(FuncT&&, Args&&... args);
template <typename... Args>
static auto functor(Args&&... args);
template <typename RetT, typename... Args>
static RetT functor(Args... args);
};
template <size_t TableIdx>
source/lib/rocprofiler-sdk/hsa/code_object.cpp
+9 -1
View File
@@ -255,7 +255,15 @@ accum_vgpr_count(std::string_view name, kernel_descriptor_t kernel_code)
1) *
(8 - arch_vgpr_count(name, kernel_code));
LOG(WARNING) << "Missing support for accum_vgpr_count for " << name;
bool emplaced = false;
{
static auto warned = std::unordered_set<std::string>{};
static auto mtx = std::mutex{};
auto lk = std::unique_lock<std::mutex>{mtx};
emplaced = warned.emplace(name).second;
}
LOG_IF(WARNING, emplaced) << "Missing support for accum_vgpr_count for " << name;
return 0;
}
source/lib/rocprofiler-sdk/hsa/defines.hpp
+2 -12
View File
@@ -123,14 +123,9 @@
template <typename RetT, typename... Args> \
static auto get_functor(RetT (*)(Args...)) \
{ \
if constexpr(std::is_void<RetT>::value) \
return [](Args... args) -> RetT { base_type::functor(args...); }; \
else \
return [](Args... args) -> RetT { return base_type::functor(args...); }; \
return &base_type::functor<RetT, Args...>; \
} \
\
static auto get_functor() { return get_functor(get_table_func()); } \
\
static std::vector<void*> as_arg_addr(rocprofiler_callback_tracing_hsa_api_data_t) \
{ \
return std::vector<void*>{}; \
@@ -203,14 +198,9 @@
template <typename RetT, typename... Args> \
static auto get_functor(RetT (*)(Args...)) \
{ \
if constexpr(std::is_void<RetT>::value) \
return [](Args... args) -> RetT { base_type::functor(args...); }; \
else \
return [](Args... args) -> RetT { return base_type::functor(args...); }; \
return &base_type::functor<RetT, Args...>; \
} \
\
static auto get_functor() { return get_functor(get_table_func()); } \
\
static std::vector<void*> as_arg_addr( \
rocprofiler_callback_tracing_hsa_api_data_t trace_data) \
{ \
source/lib/rocprofiler-sdk/hsa/hsa.cpp
+15 -17
View File
@@ -167,8 +167,7 @@ get_table()
.core_ = get_core_table(),
.amd_ext_ = get_amd_ext_table(),
.finalizer_ext_ = get_fini_ext_table(),
.image_ext_ = get_img_ext_table(),
.tools_ = nullptr};
.image_ext_ = get_img_ext_table()};
return tbl;
}
@@ -267,19 +266,21 @@ populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx,
} // namespace
template <size_t TableIdx, size_t OpIdx>
template <typename... Args>
auto
hsa_api_impl<TableIdx, OpIdx>::functor(Args&&... args)
template <typename RetT, typename... Args>
RetT
hsa_api_impl<TableIdx, OpIdx>::functor(Args... args)
{
using info_type = hsa_api_info<TableIdx, OpIdx>;
LOG_IF(INFO, registration::get_fini_status() != 0) << "Executing " << info_type::name;
if(registration::get_fini_status() != 0)
{
auto _ret = exec(info_type::get_table_func(), std::forward<Args>(args)...);
if constexpr(!std::is_same<decltype(_ret), null_type>::value)
[[maybe_unused]] auto _ret = exec(info_type::get_table_func(), std::forward<Args>(args)...);
if constexpr(!std::is_void<RetT>::value)
return _ret;
else
return HSA_STATUS_SUCCESS;
return;
}
auto thr_id = common::get_tid();
@@ -294,11 +295,11 @@ hsa_api_impl<TableIdx, OpIdx>::functor(Args&&... args)
if(callback_contexts.empty() && buffered_contexts.empty())
{
auto _ret = exec(info_type::get_table_func(), std::forward<Args>(args)...);
if constexpr(!std::is_same<decltype(_ret), null_type>::value)
[[maybe_unused]] auto _ret = exec(info_type::get_table_func(), std::forward<Args>(args)...);
if constexpr(!std::is_void<RetT>::value)
return _ret;
else
return HSA_STATUS_SUCCESS;
return;
}
constexpr auto ref_count = 2;
@@ -423,10 +424,7 @@ hsa_api_impl<TableIdx, OpIdx>::functor(Args&&... args)
context::pop_latest_correlation_id(corr_id);
if constexpr(!std::is_same<decltype(_ret), null_type>::value)
return _ret;
else
return HSA_STATUS_SUCCESS;
if constexpr(!std::is_void<RetT>::value) return _ret;
}
} // namespace hsa
} // namespace rocprofiler
@@ -601,8 +599,6 @@ update_table(const context::context_array_t& _contexts,
{
auto _info = hsa_api_info<TableIdx, OpIdx>{};
LOG(INFO) << "updating table entry for " << _info.name;
// make sure we don't access a field that doesn't exist in input table
if(_info.offset() >= _orig->version.minor_id) return;
@@ -614,6 +610,8 @@ update_table(const context::context_array_t& _contexts,
_info.operation_idx))
return;
LOG(INFO) << "updating table entry for " << _info.name;
// 1. get the sub-table containing the function pointer in original table
// 2. get reference to function pointer in sub-table in original table
// 3. update function pointer with wrapper
source/lib/rocprofiler-sdk/hsa/hsa.hpp
+2 -2
View File
@@ -96,8 +96,8 @@ struct hsa_api_impl
template <typename FuncT, typename... Args>
static auto exec(FuncT&&, Args&&... args);
template <typename... Args>
static auto functor(Args&&... args);
template <typename RetT, typename... Args>
static RetT functor(Args... args);
};
template <size_t TableIdx>
source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp
+21 -2
View File
@@ -21,17 +21,36 @@
// SOFTWARE.
#include "lib/rocprofiler-sdk/hsa/hsa_barrier.hpp"
#include "lib/rocprofiler-sdk/registration.hpp"
namespace rocprofiler
{
namespace hsa
{
hsa_barrier::hsa_barrier(std::function<void()>&& finished, CoreApiTable core_api)
: _barried_finished(std::move(finished))
, _core_api(core_api)
{
// Create the barrier signal
_core_api.hsa_signal_create_fn(0, 0, nullptr, &_barrier_signal);
}
hsa_barrier::~hsa_barrier()
{
// Destroy the barrier signal
if(registration::get_fini_status() < 1)
{
_core_api.hsa_signal_store_screlease_fn(_barrier_signal, 0);
_core_api.hsa_signal_destroy_fn(_barrier_signal);
}
}
void
hsa_barrier::set_barrier(queue_map_t& q)
hsa_barrier::set_barrier(const queue_map_t& q)
{
_core_api.hsa_signal_store_screlease_fn(_barrier_signal, 1);
_queue_waiting.wlock([&](auto& queue_waiting) {
for(auto& [_, queue] : q)
for(const auto& [_, queue] : q)
{
queue->lock_queue([ptr = queue.get(), &queue_waiting]() {
if(ptr->active_async_packets() > 0)
source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp
+3 -14
View File
@@ -44,22 +44,11 @@ class hsa_barrier
{
public:
using queue_map_t = std::unordered_map<hsa_queue_t*, std::unique_ptr<Queue>>;
hsa_barrier(std::function<void()>&& finished, CoreApiTable core_api)
: _barried_finished(std::move(finished))
, _core_api(core_api)
{
// Create the barrier signal
_core_api.hsa_signal_create_fn(0, 0, nullptr, &_barrier_signal);
}
~hsa_barrier()
{
// Destroy the barrier signal
_core_api.hsa_signal_store_screlease_fn(_barrier_signal, 0);
_core_api.hsa_signal_destroy_fn(_barrier_signal);
}
hsa_barrier(std::function<void()>&& finished, CoreApiTable core_api);
~hsa_barrier();
void set_barrier(queue_map_t& q);
void set_barrier(const queue_map_t& q);
std::optional<rocprofiler_packet> enqueue_packet(const Queue* queue);
bool register_completion(const Queue* queue);
source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp
+69 -28
View File
@@ -34,10 +34,11 @@ bool
profiler_serializer_ready_signal_handler(hsa_signal_value_t /* signal_value */, void* data)
{
auto* hsa_queue = static_cast<hsa_queue_t*>(data);
const auto* queue = get_queue_controller().get_queue(*hsa_queue);
const auto* queue = CHECK_NOTNULL(get_queue_controller())->get_queue(*hsa_queue);
CHECK(queue);
get_queue_controller().serializer().wlock(
[&](auto& serializer) { serializer.queue_ready(hsa_queue, *queue); });
CHECK_NOTNULL(get_queue_controller())->serializer().wlock([&](auto& serializer) {
serializer.queue_ready(hsa_queue, *queue);
});
return true;
}
@@ -63,8 +64,14 @@ void
profiler_serializer::add_queue(hsa_queue_t** hsa_queues, const Queue& queue)
{
hsa_signal_t signal = queue.ready_signal;
hsa_status_t status = get_queue_controller().get_ext_table().hsa_amd_signal_async_handler_fn(
signal, HSA_SIGNAL_CONDITION_EQ, -1, profiler_serializer_ready_signal_handler, *hsa_queues);
hsa_status_t status =
CHECK_NOTNULL(get_queue_controller())
->get_ext_table()
.hsa_amd_signal_async_handler_fn(signal,
HSA_SIGNAL_CONDITION_EQ,
-1,
profiler_serializer_ready_signal_handler,
*hsa_queues);
if(status != HSA_STATUS_SUCCESS) LOG(FATAL) << "hsa_amd_signal_async_handler failed";
}
@@ -96,16 +103,19 @@ profiler_serializer::kernel_completion_signal(const Queue& completed)
CHECK(_dispatch_queue);
_dispatch_queue = nullptr;
get_queue_controller().get_core_table().hsa_signal_store_screlease_fn(completed.block_signal,
1);
get_queue_controller().get_core_table().hsa_signal_store_screlease_fn(completed.ready_signal,
0);
CHECK_NOTNULL(get_queue_controller())
->get_core_table()
.hsa_signal_store_screlease_fn(completed.block_signal, 1);
CHECK_NOTNULL(get_queue_controller())
->get_core_table()
.hsa_signal_store_screlease_fn(completed.ready_signal, 0);
if(!_dispatch_ready.empty())
{
const auto* queue = _dispatch_ready.front();
_dispatch_ready.erase(_dispatch_ready.begin());
get_queue_controller().get_core_table().hsa_signal_store_screlease_fn(queue->block_signal,
0);
CHECK_NOTNULL(get_queue_controller())
->get_core_table()
.hsa_signal_store_screlease_fn(queue->block_signal, 0);
_dispatch_queue = queue;
}
}
@@ -114,20 +124,34 @@ void
profiler_serializer::queue_ready(hsa_queue_t* hsa_queue, const Queue& queue)
{
{
LOG(INFO) << "Obtaining queue mutex lock...";
std::lock_guard<std::mutex> cv_lock(queue.cv_mutex);
LOG(INFO) << "Queue mutex lock obtained";
if(queue.get_state() == queue_state::to_destroy)
{
get_queue_controller().set_queue_state(queue_state::done_destroy, hsa_queue);
get_queue_controller().get_core_table().hsa_signal_destroy_fn(queue.ready_signal);
LOG(INFO) << "Setting queue state to done_destroy...";
CHECK_NOTNULL(get_queue_controller())
->set_queue_state(queue_state::done_destroy, hsa_queue);
LOG(INFO) << "Destroying ready signal...";
CHECK_NOTNULL(get_queue_controller())
->get_core_table()
.hsa_signal_destroy_fn(queue.ready_signal);
LOG(INFO) << "Notifying queue condition variable...";
queue.cv_ready_signal.notify_one();
return;
}
}
get_queue_controller().get_core_table().hsa_signal_store_screlease_fn(queue.ready_signal, 1);
LOG(INFO) << "setting queue ready signal to 1...";
CHECK_NOTNULL(get_queue_controller())
->get_core_table()
.hsa_signal_store_screlease_fn(queue.ready_signal, 1);
if(_dispatch_queue == nullptr)
{
get_queue_controller().get_core_table().hsa_signal_store_screlease_fn(queue.block_signal,
0);
CHECK_NOTNULL(get_queue_controller())
->get_core_table()
.hsa_signal_store_screlease_fn(queue.block_signal, 0);
_dispatch_queue = &queue;
}
else
@@ -178,8 +202,10 @@ profiler_serializer::kernel_dispatch(const Queue& queue) const
}
void
profiler_serializer::destory_queue(hsa_queue_t* id, const Queue& queue)
profiler_serializer::destroy_queue(hsa_queue_t* id, const Queue& queue)
{
LOG(INFO) << "destroying queue...";
/*Deletes the queue to be destructed from the dispatch ready.*/
for(auto& barriers : _barrier)
{
@@ -206,42 +232,57 @@ profiler_serializer::destory_queue(hsa_queue_t* id, const Queue& queue)
return false;
}),
_dispatch_ready.end());
get_queue_controller().set_queue_state(queue_state::to_destroy, id);
get_queue_controller().get_core_table().hsa_signal_store_screlease_fn(queue.ready_signal, 0);
CHECK_NOTNULL(get_queue_controller())->set_queue_state(queue_state::to_destroy, id);
CHECK_NOTNULL(get_queue_controller())
->get_core_table()
.hsa_signal_store_screlease_fn(queue.ready_signal, 0);
LOG(INFO) << "queue destroyed";
}
// Enable the serializer
void
profiler_serializer::enable(queue_map_t& queues)
profiler_serializer::enable(const queue_map_t& queues)
{
if(_serializer_status == Status::ENABLED) return;
LOG(INFO) << "Enabling profiler serialization...";
_serializer_status = Status::ENABLED;
if(queues.empty()) return;
clear_complete_barriers(_barrier);
_barrier.emplace_back(
Status::DISABLED,
std::make_unique<hsa_barrier>([] {}, get_queue_controller().get_core_table()));
_barrier.emplace_back(Status::DISABLED,
std::make_unique<hsa_barrier>(
[] {}, CHECK_NOTNULL(get_queue_controller())->get_core_table()));
_serializer_status = Status::ENABLED;
_barrier.back().barrier->set_barrier(queues);
LOG(INFO) << "Profiler serialization enabled";
}
// Disable the serializer
void
profiler_serializer::disable(queue_map_t& queues)
profiler_serializer::disable(const queue_map_t& queues)
{
if(_serializer_status == Status::DISABLED) return;
LOG(INFO) << "Disabling profiler serialization...";
_serializer_status = Status::DISABLED;
if(queues.empty()) return;
clear_complete_barriers(_barrier);
_barrier.emplace_back(
Status::ENABLED,
std::make_unique<hsa_barrier>([] {}, get_queue_controller().get_core_table()));
_barrier.emplace_back(Status::ENABLED,
std::make_unique<hsa_barrier>(
[] {}, CHECK_NOTNULL(get_queue_controller())->get_core_table()));
_serializer_status = Status::DISABLED;
_barrier.back().barrier->set_barrier(queues);
LOG(INFO) << "Profiler serialization disabled";
}
} // namespace hsa
} // namespace rocprofiler
} // namespace rocprofiler
source/lib/rocprofiler-sdk/hsa/profile_serializer.hpp
+3 -3
View File
@@ -79,11 +79,11 @@ public:
void queue_ready(hsa_queue_t* hsa_queue, const Queue& queue);
// Enable the serializer
void enable(queue_map_t& queues);
void enable(const queue_map_t& queues);
// Disable the serializer
void disable(queue_map_t& queues);
void disable(const queue_map_t& queues);
void destory_queue(hsa_queue_t* id, const Queue& queue);
void destroy_queue(hsa_queue_t* id, const Queue& queue);
static void add_queue(hsa_queue_t** hsa_queues, const Queue& queue);
source/lib/rocprofiler-sdk/hsa/queue.cpp
+44 -44
View File
@@ -78,8 +78,6 @@ context_filter(const context::context* ctx)
bool
AsyncSignalHandler(hsa_signal_value_t /*signal_v*/, void* data)
{
// LOG(ERROR) << "signal value is " << signal_v;
if(!data) return true;
auto& queue_info_session = *static_cast<Queue::queue_info_session_t*>(data);
@@ -173,8 +171,9 @@ AsyncSignalHandler(hsa_signal_value_t /*signal_v*/, void* data)
if(queue_info_session.interrupt_signal.handle != 0u)
{
#if !defined(NDEBUG)
hsa::get_queue_controller()._debug_signals.wlock(
[&](auto& signals) { signals.erase(queue_info_session.interrupt_signal.handle); });
CHECK_NOTNULL(hsa::get_queue_controller())->_debug_signals.wlock([&](auto& signals) {
signals.erase(queue_info_session.interrupt_signal.handle);
});
#endif
hsa::get_core_table()->hsa_signal_store_screlease_fn(queue_info_session.interrupt_signal,
-1);
@@ -235,8 +234,7 @@ WriteInterceptor(const void* packets,
void* data,
hsa_amd_queue_intercept_packet_writer writer)
{
using context_array_t = Queue::context_array_t;
using context_array_t = Queue::context_array_t;
auto&& CreateBarrierPacket = [](hsa_signal_t* dependency_signal,
hsa_signal_t* completion_signal,
std::vector<rocprofiler_packet>& _packets) {
@@ -345,7 +343,7 @@ WriteInterceptor(const void* packets,
barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
// barrier.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE;
// barrier.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE;
// barrier.header |= 1 << HSA_PACKET_HEADER_BARRIER;
barrier.header |= 1 << HSA_PACKET_HEADER_BARRIER;
barrier.completion_signal = original_packet.completion_signal;
transformed_packets.emplace_back(barrier);
}
@@ -410,27 +408,11 @@ WriteInterceptor(const void* packets,
}
} // namespace
Queue::~Queue() { sync(); }
void
Queue::signal_async_handler(const hsa_signal_t& signal, Queue::queue_info_session_t* data) const
Queue::Queue(const AgentCache& agent, CoreApiTable table)
: _core_api(table)
, _agent(agent)
{
#if !defined(NDEBUG)
hsa::get_queue_controller()._debug_signals.wlock(
[&](auto& signals) { signals[signal.handle] = signal; });
#endif
hsa_status_t status = _ext_api.hsa_amd_signal_async_handler_fn(
signal, HSA_SIGNAL_CONDITION_EQ, -1, AsyncSignalHandler, static_cast<void*>(data));
LOG_IF(FATAL, status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK)
<< "Error: hsa_amd_signal_async_handler failed";
}
void
Queue::create_signal(uint32_t attribute, hsa_signal_t* signal) const
{
hsa_status_t status = _ext_api.hsa_amd_signal_create_fn(1, 0, nullptr, attribute, signal);
LOG_IF(FATAL, status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK)
<< "Error: hsa_amd_signal_create failed";
_core_api.hsa_signal_create_fn(0, 0, nullptr, &_active_kernels);
}
Queue::Queue(const AgentCache& agent,
@@ -470,30 +452,48 @@ Queue::Queue(const AgentCache& agent,
create_signal(0, &ready_signal);
create_signal(0, &block_signal);
create_signal(0, &_active_kernels);
_core_api.hsa_signal_store_screlease_fn(ready_signal, 0);
_core_api.hsa_signal_store_screlease_fn(_active_kernels, 0);
*queue = _intercept_queue;
}
Queue::~Queue()
{
sync();
_core_api.hsa_signal_destroy_fn(_active_kernels);
}
void
Queue::signal_async_handler(const hsa_signal_t& signal, Queue::queue_info_session_t* data) const
{
#if !defined(NDEBUG)
CHECK_NOTNULL(hsa::get_queue_controller())->_debug_signals.wlock([&](auto& signals) {
signals[signal.handle] = signal;
});
#endif
hsa_status_t status = _ext_api.hsa_amd_signal_async_handler_fn(
signal, HSA_SIGNAL_CONDITION_EQ, -1, AsyncSignalHandler, static_cast<void*>(data));
LOG_IF(FATAL, status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK)
<< "Error: hsa_amd_signal_async_handler failed";
}
void
Queue::create_signal(uint32_t attribute, hsa_signal_t* signal) const
{
hsa_status_t status = _ext_api.hsa_amd_signal_create_fn(1, 0, nullptr, attribute, signal);
LOG_IF(FATAL, status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK)
<< "Error: hsa_amd_signal_create failed";
}
void
Queue::sync() const
{
// Potentially replace with condition variable at some point
// but performance may not matter here.
constexpr auto max_wait_time = std::chrono::milliseconds{1000};
constexpr auto query_interval = std::chrono::milliseconds{10};
auto _orig_active = _active_async_packets.load(std::memory_order_relaxed);
auto _curr_active = _orig_active;
auto inactive = common::yield(
[this, &_curr_active]() {
return ((_curr_active = _active_async_packets.load(std::memory_order_relaxed)) == 0);
},
max_wait_time,
query_interval);
LOG_IF(WARNING, !inactive)
<< "rocprofiler-sdk Queue (instance=" << this << ") abandoned waiting for " << _orig_active
<< " async completion callbacks after " << max_wait_time.count() << " msecs. There were "
<< _curr_active << " async completion callbacks which were not delivered at that time.";
if(_active_kernels.handle != 0u)
{
_core_api.hsa_signal_wait_relaxed_fn(
_active_kernels, HSA_SIGNAL_CONDITION_EQ, 0, -1, HSA_WAIT_STATE_ACTIVE);
}
}
void
source/lib/rocprofiler-sdk/hsa/queue.hpp
+10 -9
View File
@@ -147,6 +147,8 @@ public:
const Queue::queue_info_session_t&,
inst_pkt_t&)>;
using callback_map_t = std::unordered_map<ClientID, std::pair<queue_cb_t, completed_cb_t>>;
Queue(const AgentCache& agent, CoreApiTable table);
Queue(const AgentCache& agent,
uint32_t size,
hsa_queue_type32_t type,
@@ -157,11 +159,6 @@ public:
CoreApiTable core_api,
AmdExtTable ext_api,
hsa_queue_t** queue);
Queue(const AgentCache& agent)
: _agent(agent)
{}
virtual ~Queue();
const hsa_queue_t* intercept_queue() const { return _intercept_queue; };
@@ -184,10 +181,13 @@ public:
// Tracks the number of in flight kernel executions we
// are waiting on. We cannot destroy Queue until all kernels
// have comleted.
void async_started() { _active_async_packets++; }
void async_complete() { _active_async_packets--; }
int64_t active_async_packets() const { return _active_async_packets; }
void sync() const;
void async_started() { _core_api.hsa_signal_add_relaxed_fn(_active_kernels, 1); }
void async_complete() { _core_api.hsa_signal_subtract_relaxed_fn(_active_kernels, 1); }
int64_t active_async_packets() const
{
return _core_api.hsa_signal_load_scacquire_fn(_active_kernels);
}
void sync() const;
void register_callback(ClientID id, queue_cb_t enqueue_cb, completed_cb_t complete_cb);
void remove_callback(ClientID id);
@@ -211,6 +211,7 @@ private:
hsa_queue_t* _intercept_queue = nullptr;
queue_state _state = queue_state::normal;
std::mutex _lock_queue;
hsa_signal_t _active_kernels = {.handle = 0};
};
inline rocprofiler_queue_id_t
source/lib/rocprofiler-sdk/hsa/queue_controller.cpp
+33 -24
View File
@@ -25,6 +25,7 @@
#include "lib/rocprofiler-sdk/agent.hpp"
#include "lib/rocprofiler-sdk/context/context.hpp"
#include "lib/rocprofiler-sdk/hsa/agent_cache.hpp"
#include "lib/rocprofiler-sdk/registration.hpp"
#include <rocprofiler-sdk/fwd.h>
@@ -47,7 +48,8 @@ create_queue(hsa_agent_t agent,
uint32_t group_segment_size,
hsa_queue_t** queue)
{
for(const auto& [_, agent_info] : get_queue_controller().get_supported_agents())
auto* controller = CHECK_NOTNULL(get_queue_controller());
for(const auto& [_, agent_info] : controller->get_supported_agents())
{
if(agent_info.get_hsa_agent().handle == agent.handle)
{
@@ -58,13 +60,13 @@ create_queue(hsa_agent_t agent,
data,
private_segment_size,
group_segment_size,
get_queue_controller().get_core_table(),
get_queue_controller().get_ext_table(),
controller->get_core_table(),
controller->get_ext_table(),
queue);
get_queue_controller().serializer().wlock(
controller->serializer().wlock(
[&](auto& serializer) { serializer.add_queue(queue, *new_queue); });
get_queue_controller().add_queue(*queue, std::move(new_queue));
controller->add_queue(*queue, std::move(new_queue));
return HSA_STATUS_SUCCESS;
}
@@ -76,7 +78,7 @@ create_queue(hsa_agent_t agent,
hsa_status_t
destroy_queue(hsa_queue_t* hsa_queue)
{
get_queue_controller().destroy_queue(hsa_queue);
if(get_queue_controller()) get_queue_controller()->destroy_queue(hsa_queue);
return HSA_STATUS_SUCCESS;
}
@@ -108,16 +110,20 @@ QueueController::add_queue(hsa_queue_t* id, std::unique_ptr<Queue> queue)
void
QueueController::destroy_queue(hsa_queue_t* id)
{
const auto* queue = get_queue_controller().get_queue(*id);
std::unique_lock<std::mutex> cvlock(queue->cv_mutex);
if(!id) return;
serializer().wlock([&](auto& serializer) { serializer.destory_queue(id, *queue); });
const auto* queue = get_queue(*id);
queue->cv_ready_signal.wait(
cvlock, [queue] { return queue->get_state() == queue_state::done_destroy; });
if(queue->block_signal.handle != 0)
get_queue_controller().get_core_table().hsa_signal_destroy_fn(queue->block_signal);
// return if queue does not exist
if(!queue) return;
LOG(INFO) << "destroying queue...";
queue->sync();
if(queue->block_signal.handle != 0) get_core_table().hsa_signal_destroy_fn(queue->block_signal);
_queues.wlock([&](auto& map) { map.erase(id); });
LOG(INFO) << "queue destroyed";
}
ClientID
@@ -232,18 +238,20 @@ QueueController::get_queue(const hsa_queue_t& _hsa_queue) const
void
QueueController::disable_serialization()
{
_queues.wlock([](queue_map_t& _queues_v) {
get_queue_controller().serializer().wlock(
[&](auto& serializer) { serializer.disable(_queues_v); });
_queues.rlock([](const queue_map_t& _queues_v) {
if(get_queue_controller())
get_queue_controller()->serializer().wlock(
[&](auto& serializer) { serializer.disable(_queues_v); });
});
}
void
QueueController::enable_serialization()
{
_queues.wlock([](queue_map_t& _queues_v) {
get_queue_controller().serializer().wlock(
[&](auto& serializer) { serializer.enable(_queues_v); });
_queues.rlock([](const queue_map_t& _queues_v) {
if(get_queue_controller())
get_queue_controller()->serializer().wlock(
[&](auto& serializer) { serializer.enable(_queues_v); });
});
}
@@ -272,7 +280,7 @@ QueueController::print_debug_signals() const
}
void
QueueController::set_queue_state(enum queue_state state, hsa_queue_t* hsa_queue)
QueueController::set_queue_state(queue_state state, hsa_queue_t* hsa_queue)
{
_queues.wlock([&](auto& map) { map[hsa_queue]->set_state(state); });
}
@@ -299,23 +307,24 @@ QueueController::iterate_callbacks(const callback_iterator_cb_t& cb) const
});
}
QueueController&
QueueController*
get_queue_controller()
{
static auto*& controller = common::static_object<QueueController>::construct();
return *(CHECK_NOTNULL(controller));
return controller;
}
void
queue_controller_init(HsaApiTable* table)
{
get_queue_controller().init(*table->core_, *table->amd_ext_);
CHECK_NOTNULL(get_queue_controller())->init(*table->core_, *table->amd_ext_);
}
void
queue_controller_fini()
{
get_queue_controller().iterate_queues([](const Queue* _queue) { _queue->sync(); });
if(get_queue_controller())
get_queue_controller()->iterate_queues([](const Queue* _queue) { _queue->sync(); });
}
} // namespace hsa
} // namespace rocprofiler
source/lib/rocprofiler-sdk/hsa/queue_controller.hpp
+1 -1
View File
@@ -112,7 +112,7 @@ private:
common::Synchronized<hsa::profiler_serializer> _profiler_serializer;
};
QueueController&
QueueController*
get_queue_controller();
void
source/lib/rocprofiler-sdk/marker/defines.hpp
+2 -12
View File
@@ -89,14 +89,9 @@
template <typename RetT, typename... Args> \
static auto get_functor(RetT (*)(Args...)) \
{ \
if constexpr(std::is_void<RetT>::value) \
return [](Args... args) -> RetT { base_type::functor(args...); }; \
else \
return [](Args... args) -> RetT { return base_type::functor(args...); }; \
return &base_type::functor<RetT, Args...>; \
} \
\
static auto get_functor() { return get_functor(get_table_func()); } \
\
static std::vector<void*> as_arg_addr(callback_data_type) { return std::vector<void*>{}; } \
\
static std::vector<common::stringified_argument> as_arg_list(callback_data_type, int32_t) \
@@ -174,14 +169,9 @@
template <typename RetT, typename... Args> \
static auto get_functor(RetT (*)(Args...)) \
{ \
if constexpr(std::is_same<RetT, void>::value) \
return [](Args... args) -> RetT { base_type::functor(args...); }; \
else \
return [](Args... args) -> RetT { return base_type::functor(args...); }; \
return &base_type::functor<RetT, Args...>; \
} \
\
static auto get_functor() { return get_functor(get_table_func()); } \
\
static std::vector<void*> as_arg_addr(callback_data_type trace_data) \
{ \
return std::vector<void*>{ \
source/lib/rocprofiler-sdk/marker/marker.cpp
+12 -12
View File
@@ -27,6 +27,7 @@
#include "lib/rocprofiler-sdk/buffer.hpp"
#include "lib/rocprofiler-sdk/context/context.hpp"
#include "lib/rocprofiler-sdk/marker/utils.hpp"
#include "lib/rocprofiler-sdk/registration.hpp"
#include <rocprofiler-sdk/buffer.h>
#include <rocprofiler-sdk/callback_tracing.h>
@@ -181,14 +182,16 @@ populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx,
} // namespace
template <size_t TableIdx, size_t OpIdx>
template <typename... Args>
auto
roctx_api_impl<TableIdx, OpIdx>::functor(Args&&... args)
template <typename RetT, typename... Args>
RetT
roctx_api_impl<TableIdx, OpIdx>::functor(Args... args)
{
using info_type = roctx_api_info<TableIdx, OpIdx>;
using callback_api_data_t = typename roctx_domain_info<TableIdx>::callback_data_type;
using buffered_api_data_t = typename roctx_domain_info<TableIdx>::buffer_data_type;
LOG_IF(INFO, registration::get_fini_status() != 0) << "Executing " << info_type::name;
auto thr_id = common::get_tid();
auto callback_contexts = std::vector<callback_context_data>{};
auto buffered_contexts = std::vector<buffered_context_data>{};
@@ -201,11 +204,11 @@ roctx_api_impl<TableIdx, OpIdx>::functor(Args&&... args)
if(callback_contexts.empty() && buffered_contexts.empty())
{
auto _ret = exec(info_type::get_table_func(), std::forward<Args>(args)...);
if constexpr(!std::is_same<decltype(_ret), null_type>::value)
[[maybe_unused]] auto _ret = exec(info_type::get_table_func(), std::forward<Args>(args)...);
if constexpr(!std::is_void<RetT>::value)
return _ret;
else
return 0;
return;
}
auto ref_count = 2;
@@ -330,10 +333,7 @@ roctx_api_impl<TableIdx, OpIdx>::functor(Args&&... args)
context::pop_latest_correlation_id(corr_id);
if constexpr(!std::is_same<decltype(_ret), null_type>::value)
return _ret;
else
return 0;
if constexpr(!std::is_void<RetT>::value) return _ret;
}
} // namespace marker
} // namespace rocprofiler
@@ -505,8 +505,6 @@ update_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
{
auto _info = roctx_api_info<TableIdx, OpIdx>{};
LOG(INFO) << "updating table entry for " << _info.name;
// make sure we don't access a field that doesn't exist in input table
if(_info.offset() >= _orig->size) return;
@@ -516,6 +514,8 @@ update_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
_info.callback_domain_idx, _info.buffered_domain_idx, _info.operation_idx))
return;
LOG(INFO) << "updating table entry for " << _info.name;
// 1. get the sub-table containing the function pointer in original table
// 2. get reference to function pointer in sub-table in original table
// 3. update function pointer with wrapper
source/lib/rocprofiler-sdk/marker/marker.hpp
+2 -2
View File
@@ -62,8 +62,8 @@ struct roctx_api_impl : roctx_domain_info<TableIdx>
template <typename FuncT, typename... Args>
static auto exec(FuncT&&, Args&&... args);
template <typename... Args>
static auto functor(Args&&... args);
template <typename RetT, typename... Args>
static RetT functor(Args... args);
};
template <size_t TableIdx>
source/lib/rocprofiler-sdk/tests/CMakeLists.txt
-1
View File
@@ -59,7 +59,6 @@ gtest_add_tests(
WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR})
set(rocprofiler-lib-tests-env
"HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler-shared-library>"
"LD_LIBRARY_PATH=${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_LIBDIR}:$ENV{LD_LIBRARY_PATH}")
set_tests_properties(
source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp
+62 -46
View File
@@ -14,36 +14,10 @@
#include "lib/rocprofiler-sdk/hsa/hsa_barrier.hpp"
#include "lib/rocprofiler-sdk/hsa/queue_controller.hpp"
#include "lib/rocprofiler-sdk/registration.hpp"
#include "rocprofiler-sdk/registration.h"
using namespace rocprofiler;
using namespace rocprofiler::hsa;
namespace rocprofiler
{
namespace hsa
{
class FakeQueue : public Queue
{
public:
FakeQueue(const AgentCache& a, rocprofiler_queue_id_t id)
: Queue(a)
, _agent(a)
, _id(id)
{}
virtual const AgentCache& get_agent() const override final { return _agent; };
virtual rocprofiler_queue_id_t get_id() const override final { return _id; };
~FakeQueue() {}
private:
const AgentCache& _agent;
rocprofiler_queue_id_t _id = {};
};
} // namespace hsa
} // namespace rocprofiler
namespace
{
AmdExtTable&
@@ -66,35 +40,76 @@ CoreApiTable&
get_api_table()
{
static auto _v = []() {
auto val = CoreApiTable{};
val.hsa_iterate_agents_fn = hsa_iterate_agents;
val.hsa_agent_get_info_fn = hsa_agent_get_info;
val.hsa_queue_create_fn = hsa_queue_create;
val.hsa_queue_destroy_fn = hsa_queue_destroy;
val.hsa_signal_create_fn = hsa_signal_create;
val.hsa_signal_destroy_fn = hsa_signal_destroy;
val.hsa_signal_store_screlease_fn = hsa_signal_store_screlease;
val.hsa_signal_load_scacquire_fn = hsa_signal_load_scacquire;
auto val = CoreApiTable{};
val.hsa_iterate_agents_fn = hsa_iterate_agents;
val.hsa_agent_get_info_fn = hsa_agent_get_info;
val.hsa_queue_create_fn = hsa_queue_create;
val.hsa_queue_destroy_fn = hsa_queue_destroy;
val.hsa_signal_create_fn = hsa_signal_create;
val.hsa_signal_destroy_fn = hsa_signal_destroy;
val.hsa_signal_store_screlease_fn = hsa_signal_store_screlease;
val.hsa_signal_load_scacquire_fn = hsa_signal_load_scacquire;
val.hsa_signal_add_relaxed_fn = hsa_signal_add_relaxed;
val.hsa_signal_subtract_relaxed_fn = hsa_signal_subtract_relaxed;
val.hsa_signal_wait_relaxed_fn = hsa_signal_wait_relaxed;
return val;
}();
return _v;
}
namespace rocprofiler
{
namespace hsa
{
class FakeQueue : public Queue
{
public:
FakeQueue(const AgentCache& a, rocprofiler_queue_id_t id)
: Queue(a, get_api_table())
, _agent(a)
, _id(id)
{}
virtual const AgentCache& get_agent() const override final { return _agent; };
virtual rocprofiler_queue_id_t get_id() const override final { return _id; };
~FakeQueue() {}
private:
const AgentCache& _agent;
rocprofiler_queue_id_t _id = {};
};
} // namespace hsa
} // namespace rocprofiler
QueueController::queue_map_t
create_queue_map(size_t count)
{
QueueController::queue_map_t ret;
auto agents = hsa::get_queue_controller().get_supported_agents();
// ensure test fails if null
EXPECT_TRUE(hsa::get_queue_controller() != nullptr);
// prevent segfault
if(!hsa::get_queue_controller()) return ret;
auto agents = hsa::get_queue_controller()->get_supported_agents();
for(size_t i = 0; i < count; i++)
{
auto& agent_cache = agents.begin()->second;
// Create queue
hsa_queue_t* queue;
hsa_queue_create(
agent_cache.get_hsa_agent(), 2048, HSA_QUEUE_TYPE_SINGLE, NULL, NULL, 0, 0, &queue);
ret[queue] = std::make_unique<FakeQueue>(agent_cache, rocprofiler_queue_id_t{.handle = i});
hsa_queue_create(agent_cache.get_hsa_agent(),
2048,
HSA_QUEUE_TYPE_SINGLE,
nullptr,
nullptr,
0,
0,
&queue);
ret[queue] = std::make_unique<rocprofiler::hsa::FakeQueue>(
agent_cache, rocprofiler_queue_id_t{.handle = i});
}
return ret;
@@ -171,8 +186,9 @@ test_init()
HsaApiTable table;
table.amd_ext_ = &get_ext_table();
table.core_ = &get_api_table();
rocprofiler::agent::construct_agent_cache(&table);
hsa::get_queue_controller().init(get_api_table(), get_ext_table());
agent::construct_agent_cache(&table);
ASSERT_TRUE(hsa::get_queue_controller() != nullptr);
hsa::get_queue_controller()->init(get_api_table(), get_ext_table());
}
} // namespace
@@ -191,7 +207,7 @@ TEST(hsa_barrier, no_block_single)
auto queues = create_queue_map(1);
// Immediate return of barrier due to no active async packets
rocprofiler::hsa::hsa_barrier barrier(finished_func, get_api_table());
hsa::hsa_barrier barrier(finished_func, get_api_table());
barrier.set_barrier(queues);
executed_handlers = 0;
ASSERT_TRUE(barrier.complete());
@@ -222,7 +238,7 @@ TEST(hsa_barrier, no_block_multi)
auto queues = create_queue_map(10);
// Immediate return of barrier due to no active async packets
rocprofiler::hsa::hsa_barrier barrier(finished_func, get_api_table());
hsa::hsa_barrier barrier(finished_func, get_api_table());
barrier.set_barrier(queues);
ASSERT_TRUE(barrier.complete());
should_execute_handler = true;
@@ -253,7 +269,7 @@ TEST(hsa_barrier, block_single)
auto queues = create_queue_map(1);
rocprofiler::hsa::hsa_barrier barrier(finished_func, get_api_table());
hsa::hsa_barrier barrier(finished_func, get_api_table());
// Simulate waiting on packets already in the queue to complete
for(auto& [_, queue] : queues)
@@ -308,7 +324,7 @@ TEST(hsa_barrier, block_multi)
auto queues = create_queue_map(10);
// Immediate return of barrier due to no active async packets
rocprofiler::hsa::hsa_barrier barrier(finished_func, get_api_table());
hsa::hsa_barrier barrier(finished_func, get_api_table());
// Simulate waiting on packets already in the queue to complete
for(auto& [_, queue] : queues)
@@ -355,4 +371,4 @@ TEST(hsa_barrier, block_multi)
registration::set_init_status(1);
registration::finalize();
}
}
source/scripts/run-ci.py
+1 -1
View File
@@ -80,7 +80,7 @@ def generate_custom(args, cmake_args, ctest_args):
os.environ["TSAN_OPTIONS"] = " ".join(
[
"history_size=5",
"second_deadlock_stack=1",
"detect_deadlocks=0",
f"suppressions={SOURCE_DIR}/source/scripts/thread-sanitizer-suppr.txt",
external_symbolizer_path,
os.environ.get("TSAN_OPTIONS", ""),
source/scripts/setup-sanitizer-env.sh
+1 -1
View File
@@ -27,7 +27,7 @@ fi
: ${ASAN_OPTIONS="detect_leaks=0 use_sigaltstack=0 suppressions=${SUPPR_DIR}/address-sanitizer-suppr.txt"}
: ${LSAN_OPTIONS="suppressions=${SUPPR_DIR}/leak-sanitizer-suppr.txt"}
: ${TSAN_OPTIONS="history_size=5 second_deadlock_stack=1 suppressions=${SUPPR_DIR}/thread-sanitizer-suppr.txt${EXTERNAL_SYMBOLIZER}"}
: ${TSAN_OPTIONS="history_size=5 detect_deadlocks=0 suppressions=${SUPPR_DIR}/thread-sanitizer-suppr.txt${EXTERNAL_SYMBOLIZER}"}
export ASAN_OPTIONS
export LSAN_OPTIONS
source/scripts/thread-sanitizer-suppr.txt
+2 -1
View File
@@ -10,6 +10,7 @@ race:libamdhip64.so
# data race arising from hsa runtime
race:libhsa-runtime64.so
# unlock of an unlocked mutex (or by a wrong thread)
mutex:libhsa-runtime64.so
@@ -25,4 +26,4 @@ race:tzset_internal
mutex:external/ptl/source/PTL/TaskGroup.hh
# lock order inversion that cannot happen
mutex:source/lib/common/synchronized.hpp
mutex:source/lib/common/synchronized.hpp
tests/async-copy-tracing/CMakeLists.txt
-1
View File
@@ -21,7 +21,6 @@ add_test(NAME test-async-copy-tracing-execute COMMAND $<TARGET_FILE:transpose>)
set(async-copy-tracing-env
"${PRELOAD_ENV}"
"HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>"
"ROCPROFILER_TOOL_OUTPUT_FILE=async-copy-tracing-test.json"
"LD_LIBRARY_PATH=$<TARGET_FILE_DIR:rocprofiler::rocprofiler-shared-library>:$ENV{LD_LIBRARY_PATH}"
)
tests/async-copy-tracing/validate.py
+2 -2
View File
@@ -224,12 +224,12 @@ def test_retired_correlation_ids(input_data):
for cid, itr in async_corr_ids.items():
assert cid in retired_corr_ids.keys()
ts = retired_corr_ids[cid]["timestamp"]
assert ts > itr["end_timestamp"]
assert (ts - itr["end_timestamp"]) > 0, f"correlation-id: {cid}, data: {itr}"
for cid, itr in api_corr_ids.items():
assert cid in retired_corr_ids.keys()
ts = retired_corr_ids[cid]["timestamp"]
assert ts > itr["end_timestamp"]
assert (ts - itr["end_timestamp"]) > 0, f"correlation-id: {cid}, data: {itr}"
assert len(api_corr_ids.keys()) == (len(retired_corr_ids.keys()))
tests/bin/multistream/multistream_app.cpp
+5 -5
View File
@@ -78,15 +78,15 @@ LaunchMultiStreamKernels()
HIP_ASSERT(hipMemcpy(d_x, x, N * sizeof(float), hipMemcpyHostToDevice));
HIP_ASSERT(hipMemcpy(d_y, y, N * sizeof(float), hipMemcpyHostToDevice));
// Launch kernel on 1M elements on the GPU
// Launch kernel on one or two warps on the GPU
int blockSize = 64;
// This Kernel will always be launched with one wave
int numBlocks = 1;
for(int i = 0; i < 100; i++)
{
for(size_t j = 0; j < hip_streams.size(); j++)
for(auto& hip_stream : hip_streams)
{
hipLaunchKernelGGL(add, numBlocks, blockSize, 0, hip_streams[j], N, d_x, d_y);
hipLaunchKernelGGL(add, numBlocks, blockSize, 0, hip_stream, N, d_x, d_y);
}
}
@@ -103,9 +103,9 @@ LaunchMultiStreamKernels()
delete[] x;
delete[] y;
for(size_t i = 0; i < hip_streams.size(); i++)
for(auto& hip_stream : hip_streams)
{
HIP_ASSERT(hipStreamDestroy(hip_streams[i]));
HIP_ASSERT(hipStreamDestroy(hip_stream));
}
}
tests/counter-collection/CMakeLists.txt
+7 -9
View File
@@ -17,19 +17,17 @@ else()
set(PRELOAD_ENV "LD_PRELOAD=$<TARGET_FILE:rocprofiler-sdk-json-tool>")
endif()
set(counter-collection-env
"${PRELOAD_ENV}" "ROCPROFILER_TOOL_OUTPUT_FILE=counter-collection-test.json"
"ROCPROFILER_TOOL_CONTEXTS=COUNTER_COLLECTION" "ROCPROF_COUNTERS=SQ_WAVES_sum")
add_test(NAME test-counter-collection-execute COMMAND $<TARGET_FILE:multistream>)
set_tests_properties(
test-counter-collection-execute
PROPERTIES
TIMEOUT
45
LABELS
"integration-tests"
ENVIRONMENT
"${PRELOAD_ENV};HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>;ROCPROFILER_TOOL_OUTPUT_FILE=counter-collection-test.json;ROCPROFILER_TOOL_CONTEXTS=COUNTER_COLLECTION;ROCPROF_COUNTERS=SQ_WAVES_sum"
FAIL_REGULAR_EXPRESSION
"${ROCPROFILER_DEFAULT_FAIL_REGEX}")
PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT
"${counter-collection-env}" FAIL_REGULAR_EXPRESSION
"${ROCPROFILER_DEFAULT_FAIL_REGEX}")
foreach(FILENAME validate.py pytest.ini conftest.py)
configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME}
tests/counter-collection/validate.py
+8 -2
View File
@@ -23,10 +23,16 @@ def test_data_structure(input_data):
def test_counter_values(input_data):
data = input_data
counter_info = {}
scaling_factor = 1
for itr in data["rocprofiler-sdk-json-tool"]["agents"]:
if itr["type"] == 2 and itr["wave_front_size"] > 0:
scaling_factor = 64 / itr["wave_front_size"]
break
for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["counter_collection"]:
value = itr["counter_value"]
assert value == 1
assert int(round(value, 0)) == int(round(1 * scaling_factor, 0))
if __name__ == "__main__":
tests/kernel-tracing/CMakeLists.txt
-1
View File
@@ -26,7 +26,6 @@ add_test(NAME test-kernel-tracing-execute COMMAND $<TARGET_FILE:reproducible-run
set(kernel-tracing-env
"${PRELOAD_ENV}"
"HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>"
"ROCPROFILER_TOOL_OUTPUT_FILE=kernel-tracing-test.json"
"LD_LIBRARY_PATH=$<TARGET_FILE_DIR:rocprofiler::rocprofiler-shared-library>:$ENV{LD_LIBRARY_PATH}"
)
tests/kernel-tracing/validate.py
+2 -2
View File
@@ -158,12 +158,12 @@ def test_retired_correlation_ids(input_data):
for cid, itr in async_corr_ids.items():
assert cid in retired_corr_ids.keys()
ts = retired_corr_ids[cid]["timestamp"]
assert ts > itr["end_timestamp"]
assert (ts - itr["end_timestamp"]) > 0, f"correlation-id: {cid}, data: {itr}"
for cid, itr in api_corr_ids.items():
assert cid in retired_corr_ids.keys()
ts = retired_corr_ids[cid]["timestamp"]
assert ts > itr["end_timestamp"]
assert (ts - itr["end_timestamp"]) > 0, f"correlation-id: {cid}, data: {itr}"
assert len(api_corr_ids.keys()) == (len(retired_corr_ids.keys()))
tests/rocprofv3/counter-collection/input1/CMakeLists.txt
+1 -4
View File
@@ -32,10 +32,7 @@ else()
set(LOG_LEVEL "info")
endif()
set(cc-env-pmc1
"${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}"
"ROCPROFILER_LOG_LEVEL=${LOG_LEVEL}"
"HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>")
set(cc-env-pmc1 "${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}")
set_tests_properties(
rocprofv3-test-counter-collection-pmc1-execute
tests/rocprofv3/counter-collection/input2/CMakeLists.txt
+1 -4
View File
@@ -32,10 +32,7 @@ else()
set(LOG_LEVEL "info")
endif()
set(cc-env-pmc2
"${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}"
"ROCPROFILER_LOG_LEVEL=${LOG_LEVEL}"
"HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>")
set(cc-env-pmc2 "${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}")
set_tests_properties(
rocprofv3-test-counter-collection-pmc2-execute
tests/rocprofv3/counter-collection/list_metrics/CMakeLists.txt
+1 -4
View File
@@ -33,10 +33,7 @@ else()
set(LOG_LEVEL "info")
endif()
set(cc-env-list-metrics
"${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}"
"ROCPROFILER_LOG_LEVEL=${LOG_LEVEL}"
"HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>")
set(cc-env-list-metrics "${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}")
set_tests_properties(
rocprofv3-test-list-metrics-execute
tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt
+3 -2
View File
@@ -21,12 +21,13 @@ add_test(
string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV
"${ROCPROFILER_MEMCHECK_PRELOAD_ENV}")
set(tracing-env "${PRELOAD_ENV}"
"HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>")
set(tracing-env "${PRELOAD_ENV}")
set_tests_properties(
rocprofv3-test-trace-hip-in-libraries-execute
PROPERTIES
TIMEOUT
45
LABELS
"integration-tests"
ENVIRONMENT
tests/rocprofv3/tracing-plus-cc/CMakeLists.txt
+1 -4
View File
@@ -32,10 +32,7 @@ else()
set(LOG_LEVEL "info")
endif()
set(cc-tracing-env
"${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}"
"ROCPROFILER_LOG_LEVEL=${LOG_LEVEL}"
"HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>")
set(cc-tracing-env "${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}")
set_tests_properties(
rocprofv3-test-tracing-plus-cc-execute
tests/rocprofv3/tracing/CMakeLists.txt
+5 -4
View File
@@ -26,14 +26,13 @@ else()
set(LOG_LEVEL "info")
endif()
set(tracing-env
"${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}"
"ROCPROFILER_LOG_LEVEL=${LOG_LEVEL}"
"HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>")
set(tracing-env "${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}")
set_tests_properties(
rocprofv3-test-trace-execute
PROPERTIES
TIMEOUT
45
LABELS
"integration-tests"
ENVIRONMENT
@@ -91,6 +90,8 @@ add_test(
set_tests_properties(
rocprofv3-test-systrace-execute
PROPERTIES
TIMEOUT
45
LABELS
"integration-tests"
ENVIRONMENT