From 2f9b1767e9a8c35a176482c7f45b7c960df678e8 Mon Sep 17 00:00:00 2001 From: "Jonathan R. Madsen" Date: Thu, 21 Mar 2024 17:52:15 -0500 Subject: [PATCH] Handle hsa_queue_destroy after finalization (#679) * 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 --- .github/workflows/continuous_integration.yml | 97 ++++------------ cmake/rocprofiler_config_interfaces.cmake | 38 ++++-- cmake/rocprofiler_interfaces.cmake | 24 +++- source/lib/common/CMakeLists.txt | 8 +- source/lib/common/demangle.cpp | 55 ++------- .../lib/rocprofiler-sdk-tool/CMakeLists.txt | 10 +- source/lib/rocprofiler-sdk/CMakeLists.txt | 7 +- .../rocprofiler-sdk/aql/tests/aql_test.cpp | 22 ++-- source/lib/rocprofiler-sdk/counters/core.cpp | 22 ++-- .../rocprofiler-sdk/counters/dimensions.cpp | 7 +- .../rocprofiler-sdk/counters/tests/core.cpp | 39 ++++--- .../counters/tests/dimension.cpp | 16 +-- source/lib/rocprofiler-sdk/hip/defines.hpp | 14 +-- source/lib/rocprofiler-sdk/hip/hip.cpp | 22 ++-- source/lib/rocprofiler-sdk/hip/hip.hpp | 4 +- .../lib/rocprofiler-sdk/hsa/code_object.cpp | 10 +- source/lib/rocprofiler-sdk/hsa/defines.hpp | 14 +-- source/lib/rocprofiler-sdk/hsa/hsa.cpp | 32 +++--- source/lib/rocprofiler-sdk/hsa/hsa.hpp | 4 +- .../lib/rocprofiler-sdk/hsa/hsa_barrier.cpp | 23 +++- .../lib/rocprofiler-sdk/hsa/hsa_barrier.hpp | 17 +-- .../hsa/profile_serializer.cpp | 97 +++++++++++----- .../hsa/profile_serializer.hpp | 6 +- source/lib/rocprofiler-sdk/hsa/queue.cpp | 88 +++++++------- source/lib/rocprofiler-sdk/hsa/queue.hpp | 19 +-- .../rocprofiler-sdk/hsa/queue_controller.cpp | 57 +++++---- .../rocprofiler-sdk/hsa/queue_controller.hpp | 2 +- source/lib/rocprofiler-sdk/marker/defines.hpp | 14 +-- source/lib/rocprofiler-sdk/marker/marker.cpp | 24 ++-- source/lib/rocprofiler-sdk/marker/marker.hpp | 4 +- .../lib/rocprofiler-sdk/tests/CMakeLists.txt | 1 - .../lib/rocprofiler-sdk/tests/hsa_barrier.cpp | 108 ++++++++++-------- source/scripts/run-ci.py | 2 +- source/scripts/setup-sanitizer-env.sh | 2 +- source/scripts/thread-sanitizer-suppr.txt | 3 +- tests/async-copy-tracing/CMakeLists.txt | 1 - tests/async-copy-tracing/validate.py | 4 +- tests/bin/multistream/multistream_app.cpp | 10 +- tests/counter-collection/CMakeLists.txt | 16 ++- tests/counter-collection/validate.py | 10 +- tests/kernel-tracing/CMakeLists.txt | 1 - tests/kernel-tracing/validate.py | 4 +- .../counter-collection/input1/CMakeLists.txt | 5 +- .../counter-collection/input2/CMakeLists.txt | 5 +- .../list_metrics/CMakeLists.txt | 5 +- .../tracing-hip-in-libraries/CMakeLists.txt | 5 +- .../rocprofv3/tracing-plus-cc/CMakeLists.txt | 5 +- tests/rocprofv3/tracing/CMakeLists.txt | 9 +- 48 files changed, 494 insertions(+), 498 deletions(-) diff --git a/.github/workflows/continuous_integration.yml b/.github/workflows/continuous_integration.yml index 4456c3b026..02c49167e7 100644 --- a/.github/workflows/continuous_integration.yml +++ b/.github/workflows/continuous_integration.yml @@ -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 }} diff --git a/cmake/rocprofiler_config_interfaces.cmake b/cmake/rocprofiler_config_interfaces.cmake index c0ad2260f6..87a1e7f276 100644 --- a/cmake/rocprofiler_config_interfaces.cmake +++ b/cmake/rocprofiler_config_interfaces.cmake @@ -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 -# -# ----------------------------------------------------------------------------------------# diff --git a/cmake/rocprofiler_interfaces.cmake b/cmake/rocprofiler_interfaces.cmake index 6f3c382598..07234c1553 100644 --- a/cmake/rocprofiler_interfaces.cmake +++ b/cmake/rocprofiler_interfaces.cmake @@ -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) diff --git a/source/lib/common/CMakeLists.txt b/source/lib/common/CMakeLists.txt index d9803be307..743ca98559 100644 --- a/source/lib/common/CMakeLists.txt +++ b/source/lib/common/CMakeLists.txt @@ -38,11 +38,9 @@ target_link_libraries( $ $ $ - $ - $ - $ $ - $ - $) + $ + $) + set_target_properties(rocprofiler-common-library PROPERTIES OUTPUT_NAME rocprofiler-common) diff --git a/source/lib/common/demangle.cpp b/source/lib/common/demangle.cpp index 0ee86d9818..cf316deb9f 100644 --- a/source/lib/common/demangle.cpp +++ b/source/lib/common/demangle.cpp @@ -22,10 +22,11 @@ #include "lib/common/demangle.hpp" -#include +#include #include #include +#include #include #include #include @@ -34,18 +35,6 @@ #include #include -#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 diff --git a/source/lib/rocprofiler-sdk-tool/CMakeLists.txt b/source/lib/rocprofiler-sdk-tool/CMakeLists.txt index ff5d0e5496..2b7d5b2229 100644 --- a/source/lib/rocprofiler-sdk-tool/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk-tool/CMakeLists.txt @@ -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 diff --git a/source/lib/rocprofiler-sdk/CMakeLists.txt b/source/lib/rocprofiler-sdk/CMakeLists.txt index ca1f39142b..9ae2c3aba8 100644 --- a/source/lib/rocprofiler-sdk/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk/CMakeLists.txt @@ -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) diff --git a/source/lib/rocprofiler-sdk/aql/tests/aql_test.cpp b/source/lib/rocprofiler-sdk/aql/tests/aql_test.cpp index 40af5da201..0aef048867 100644 --- a/source/lib/rocprofiler-sdk/aql/tests/aql_test.cpp +++ b/source/lib/rocprofiler-sdk/aql/tests/aql_test.cpp @@ -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) { diff --git a/source/lib/rocprofiler-sdk/counters/core.cpp b/source/lib/rocprofiler-sdk/counters/core.cpp index ba4c5c9870..a8b07cd0cc 100644 --- a/source/lib/rocprofiler-sdk/counters/core.cpp +++ b/source/lib/rocprofiler-sdk/counters/core.cpp @@ -236,8 +236,8 @@ counter_callback_info::get_packet(std::unique_ptr& 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 diff --git a/source/lib/rocprofiler-sdk/counters/dimensions.cpp b/source/lib/rocprofiler-sdk/counters/dimensions.cpp index 8d97863041..792182f8f5 100644 --- a/source/lib/rocprofiler-sdk/counters/dimensions.cpp +++ b/source/lib/rocprofiler-sdk/counters/dimensions.cpp @@ -56,7 +56,8 @@ getBlockDimensions(std::string_view agent, const Metric& metric) std::vector 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 {}; } diff --git a/source/lib/rocprofiler-sdk/counters/tests/core.cpp b/source/lib/rocprofiler-sdk/counters/tests/core.cpp index d63f3d8af4..d5ac27d88e 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/core.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/core.cpp @@ -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(); 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 from_api; diff --git a/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp b/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp index b50fed0600..647bf1ce47 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp @@ -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) { diff --git a/source/lib/rocprofiler-sdk/hip/defines.hpp b/source/lib/rocprofiler-sdk/hip/defines.hpp index 9ce3bd41b4..d7234955a0 100644 --- a/source/lib/rocprofiler-sdk/hip/defines.hpp +++ b/source/lib/rocprofiler-sdk/hip/defines.hpp @@ -88,14 +88,9 @@ template \ static auto get_functor(RetT (*)(Args...)) \ { \ - if constexpr(std::is_void::value) \ - return [](Args... args) -> RetT { base_type::functor(args...); }; \ - else \ - return [](Args... args) -> RetT { return base_type::functor(args...); }; \ + return &base_type::functor; \ } \ \ - static auto get_functor() { return get_functor(get_table_func()); } \ - \ static std::vector as_arg_addr(callback_data_type) { return std::vector{}; } \ \ static std::vector as_arg_list(callback_data_type, int32_t) \ @@ -171,14 +166,9 @@ template \ static auto get_functor(RetT (*)(Args...)) \ { \ - if constexpr(std::is_same::value) \ - return [](Args... args) -> RetT { base_type::functor(args...); }; \ - else \ - return [](Args... args) -> RetT { return base_type::functor(args...); }; \ + return &base_type::functor; \ } \ \ - static auto get_functor() { return get_functor(get_table_func()); } \ - \ static std::vector as_arg_addr(callback_data_type trace_data) \ { \ return std::vector{ \ diff --git a/source/lib/rocprofiler-sdk/hip/hip.cpp b/source/lib/rocprofiler-sdk/hip/hip.cpp index efbb0f66ee..ab61007167 100644 --- a/source/lib/rocprofiler-sdk/hip/hip.cpp +++ b/source/lib/rocprofiler-sdk/hip/hip.cpp @@ -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 #include @@ -230,14 +231,16 @@ populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx, } // namespace template -template -auto -hip_api_impl::functor(Args&&... args) +template +RetT +hip_api_impl::functor(Args... args) { using info_type = hip_api_info; using callback_api_data_t = typename hip_domain_info::callback_data_type; using buffered_api_data_t = typename hip_domain_info::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{}; auto buffered_contexts = std::vector{}; @@ -251,10 +254,10 @@ hip_api_impl::functor(Args&&... args) if(callback_contexts.empty() && buffered_contexts.empty()) { auto _ret = exec(info_type::get_table_func(), std::forward(args)...); - if constexpr(!std::is_same::value) + if constexpr(!std::is_void::value) return _ret; else - return 0; + return; } constexpr auto ref_count = 2; @@ -379,10 +382,7 @@ hip_api_impl::functor(Args&&... args) context::pop_latest_correlation_id(corr_id); - if constexpr(!std::is_same::value) - return _ret; - else - return 0; + if constexpr(!std::is_void::value) return _ret; } } // namespace hip } // namespace rocprofiler @@ -554,8 +554,6 @@ update_table(Tp* _orig, std::integral_constant) { auto _info = hip_api_info{}; - 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) _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 diff --git a/source/lib/rocprofiler-sdk/hip/hip.hpp b/source/lib/rocprofiler-sdk/hip/hip.hpp index 4d39385be5..e74cc59700 100644 --- a/source/lib/rocprofiler-sdk/hip/hip.hpp +++ b/source/lib/rocprofiler-sdk/hip/hip.hpp @@ -74,8 +74,8 @@ struct hip_api_impl : hip_domain_info template static auto exec(FuncT&&, Args&&... args); - template - static auto functor(Args&&... args); + template + static RetT functor(Args... args); }; template diff --git a/source/lib/rocprofiler-sdk/hsa/code_object.cpp b/source/lib/rocprofiler-sdk/hsa/code_object.cpp index 026e501826..c36cacfa49 100644 --- a/source/lib/rocprofiler-sdk/hsa/code_object.cpp +++ b/source/lib/rocprofiler-sdk/hsa/code_object.cpp @@ -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{}; + static auto mtx = std::mutex{}; + auto lk = std::unique_lock{mtx}; + emplaced = warned.emplace(name).second; + } + + LOG_IF(WARNING, emplaced) << "Missing support for accum_vgpr_count for " << name; return 0; } diff --git a/source/lib/rocprofiler-sdk/hsa/defines.hpp b/source/lib/rocprofiler-sdk/hsa/defines.hpp index d9b1d0568f..8ea93360e3 100644 --- a/source/lib/rocprofiler-sdk/hsa/defines.hpp +++ b/source/lib/rocprofiler-sdk/hsa/defines.hpp @@ -123,14 +123,9 @@ template \ static auto get_functor(RetT (*)(Args...)) \ { \ - if constexpr(std::is_void::value) \ - return [](Args... args) -> RetT { base_type::functor(args...); }; \ - else \ - return [](Args... args) -> RetT { return base_type::functor(args...); }; \ + return &base_type::functor; \ } \ \ - static auto get_functor() { return get_functor(get_table_func()); } \ - \ static std::vector as_arg_addr(rocprofiler_callback_tracing_hsa_api_data_t) \ { \ return std::vector{}; \ @@ -203,14 +198,9 @@ template \ static auto get_functor(RetT (*)(Args...)) \ { \ - if constexpr(std::is_void::value) \ - return [](Args... args) -> RetT { base_type::functor(args...); }; \ - else \ - return [](Args... args) -> RetT { return base_type::functor(args...); }; \ + return &base_type::functor; \ } \ \ - static auto get_functor() { return get_functor(get_table_func()); } \ - \ static std::vector as_arg_addr( \ rocprofiler_callback_tracing_hsa_api_data_t trace_data) \ { \ diff --git a/source/lib/rocprofiler-sdk/hsa/hsa.cpp b/source/lib/rocprofiler-sdk/hsa/hsa.cpp index b07a2df00c..1cc3e05c4e 100644 --- a/source/lib/rocprofiler-sdk/hsa/hsa.cpp +++ b/source/lib/rocprofiler-sdk/hsa/hsa.cpp @@ -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 -template -auto -hsa_api_impl::functor(Args&&... args) +template +RetT +hsa_api_impl::functor(Args... args) { using info_type = hsa_api_info; + 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)...); - if constexpr(!std::is_same::value) + [[maybe_unused]] auto _ret = exec(info_type::get_table_func(), std::forward(args)...); + if constexpr(!std::is_void::value) return _ret; else - return HSA_STATUS_SUCCESS; + return; } auto thr_id = common::get_tid(); @@ -294,11 +295,11 @@ hsa_api_impl::functor(Args&&... args) if(callback_contexts.empty() && buffered_contexts.empty()) { - auto _ret = exec(info_type::get_table_func(), std::forward(args)...); - if constexpr(!std::is_same::value) + [[maybe_unused]] auto _ret = exec(info_type::get_table_func(), std::forward(args)...); + if constexpr(!std::is_void::value) return _ret; else - return HSA_STATUS_SUCCESS; + return; } constexpr auto ref_count = 2; @@ -423,10 +424,7 @@ hsa_api_impl::functor(Args&&... args) context::pop_latest_correlation_id(corr_id); - if constexpr(!std::is_same::value) - return _ret; - else - return HSA_STATUS_SUCCESS; + if constexpr(!std::is_void::value) return _ret; } } // namespace hsa } // namespace rocprofiler @@ -601,8 +599,6 @@ update_table(const context::context_array_t& _contexts, { auto _info = hsa_api_info{}; - 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 diff --git a/source/lib/rocprofiler-sdk/hsa/hsa.hpp b/source/lib/rocprofiler-sdk/hsa/hsa.hpp index 0c5b08ee47..79aefa9a15 100644 --- a/source/lib/rocprofiler-sdk/hsa/hsa.hpp +++ b/source/lib/rocprofiler-sdk/hsa/hsa.hpp @@ -96,8 +96,8 @@ struct hsa_api_impl template static auto exec(FuncT&&, Args&&... args); - template - static auto functor(Args&&... args); + template + static RetT functor(Args... args); }; template diff --git a/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp b/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp index c3a214d79d..a359141445 100644 --- a/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp +++ b/source/lib/rocprofiler-sdk/hsa/hsa_barrier.cpp @@ -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&& 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) diff --git a/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp b/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp index 543f4c7ad0..b8114b700c 100644 --- a/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp +++ b/source/lib/rocprofiler-sdk/hsa/hsa_barrier.hpp @@ -44,22 +44,11 @@ class hsa_barrier { public: using queue_map_t = std::unordered_map>; - hsa_barrier(std::function&& 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&& finished, CoreApiTable core_api); + ~hsa_barrier(); - void set_barrier(queue_map_t& q); + void set_barrier(const queue_map_t& q); std::optional enqueue_packet(const Queue* queue); bool register_completion(const Queue* queue); diff --git a/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp b/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp index 91acf343d3..db9684e272 100644 --- a/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp +++ b/source/lib/rocprofiler-sdk/hsa/profile_serializer.cpp @@ -34,10 +34,11 @@ bool profiler_serializer_ready_signal_handler(hsa_signal_value_t /* signal_value */, void* data) { auto* hsa_queue = static_cast(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 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([] {}, get_queue_controller().get_core_table())); + _barrier.emplace_back(Status::DISABLED, + std::make_unique( + [] {}, 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([] {}, get_queue_controller().get_core_table())); + _barrier.emplace_back(Status::ENABLED, + std::make_unique( + [] {}, 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 \ No newline at end of file +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/hsa/profile_serializer.hpp b/source/lib/rocprofiler-sdk/hsa/profile_serializer.hpp index 5a71b540db..20bb2e8c11 100644 --- a/source/lib/rocprofiler-sdk/hsa/profile_serializer.hpp +++ b/source/lib/rocprofiler-sdk/hsa/profile_serializer.hpp @@ -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); diff --git a/source/lib/rocprofiler-sdk/hsa/queue.cpp b/source/lib/rocprofiler-sdk/hsa/queue.cpp index fe1201fdf6..190311f7d9 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -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(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& _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(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(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 diff --git a/source/lib/rocprofiler-sdk/hsa/queue.hpp b/source/lib/rocprofiler-sdk/hsa/queue.hpp index 78a1cb3a8a..183e979a84 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue.hpp +++ b/source/lib/rocprofiler-sdk/hsa/queue.hpp @@ -147,6 +147,8 @@ public: const Queue::queue_info_session_t&, inst_pkt_t&)>; using callback_map_t = std::unordered_map>; + + 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 diff --git a/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp b/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp index 7c65d6e9bd..e6e45e52fb 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp +++ b/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp @@ -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 @@ -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) void QueueController::destroy_queue(hsa_queue_t* id) { - const auto* queue = get_queue_controller().get_queue(*id); - std::unique_lock 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::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 diff --git a/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp b/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp index dcb853044d..59a14f5c6c 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp +++ b/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp @@ -112,7 +112,7 @@ private: common::Synchronized _profiler_serializer; }; -QueueController& +QueueController* get_queue_controller(); void diff --git a/source/lib/rocprofiler-sdk/marker/defines.hpp b/source/lib/rocprofiler-sdk/marker/defines.hpp index 4ba956fc2e..d93479b2ba 100644 --- a/source/lib/rocprofiler-sdk/marker/defines.hpp +++ b/source/lib/rocprofiler-sdk/marker/defines.hpp @@ -89,14 +89,9 @@ template \ static auto get_functor(RetT (*)(Args...)) \ { \ - if constexpr(std::is_void::value) \ - return [](Args... args) -> RetT { base_type::functor(args...); }; \ - else \ - return [](Args... args) -> RetT { return base_type::functor(args...); }; \ + return &base_type::functor; \ } \ \ - static auto get_functor() { return get_functor(get_table_func()); } \ - \ static std::vector as_arg_addr(callback_data_type) { return std::vector{}; } \ \ static std::vector as_arg_list(callback_data_type, int32_t) \ @@ -174,14 +169,9 @@ template \ static auto get_functor(RetT (*)(Args...)) \ { \ - if constexpr(std::is_same::value) \ - return [](Args... args) -> RetT { base_type::functor(args...); }; \ - else \ - return [](Args... args) -> RetT { return base_type::functor(args...); }; \ + return &base_type::functor; \ } \ \ - static auto get_functor() { return get_functor(get_table_func()); } \ - \ static std::vector as_arg_addr(callback_data_type trace_data) \ { \ return std::vector{ \ diff --git a/source/lib/rocprofiler-sdk/marker/marker.cpp b/source/lib/rocprofiler-sdk/marker/marker.cpp index 8466875aed..3de208cb87 100644 --- a/source/lib/rocprofiler-sdk/marker/marker.cpp +++ b/source/lib/rocprofiler-sdk/marker/marker.cpp @@ -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 #include @@ -181,14 +182,16 @@ populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx, } // namespace template -template -auto -roctx_api_impl::functor(Args&&... args) +template +RetT +roctx_api_impl::functor(Args... args) { using info_type = roctx_api_info; using callback_api_data_t = typename roctx_domain_info::callback_data_type; using buffered_api_data_t = typename roctx_domain_info::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{}; auto buffered_contexts = std::vector{}; @@ -201,11 +204,11 @@ roctx_api_impl::functor(Args&&... args) if(callback_contexts.empty() && buffered_contexts.empty()) { - auto _ret = exec(info_type::get_table_func(), std::forward(args)...); - if constexpr(!std::is_same::value) + [[maybe_unused]] auto _ret = exec(info_type::get_table_func(), std::forward(args)...); + if constexpr(!std::is_void::value) return _ret; else - return 0; + return; } auto ref_count = 2; @@ -330,10 +333,7 @@ roctx_api_impl::functor(Args&&... args) context::pop_latest_correlation_id(corr_id); - if constexpr(!std::is_same::value) - return _ret; - else - return 0; + if constexpr(!std::is_void::value) return _ret; } } // namespace marker } // namespace rocprofiler @@ -505,8 +505,6 @@ update_table(Tp* _orig, std::integral_constant) { auto _info = roctx_api_info{}; - 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) _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 diff --git a/source/lib/rocprofiler-sdk/marker/marker.hpp b/source/lib/rocprofiler-sdk/marker/marker.hpp index fefe606d32..5b1626101b 100644 --- a/source/lib/rocprofiler-sdk/marker/marker.hpp +++ b/source/lib/rocprofiler-sdk/marker/marker.hpp @@ -62,8 +62,8 @@ struct roctx_api_impl : roctx_domain_info template static auto exec(FuncT&&, Args&&... args); - template - static auto functor(Args&&... args); + template + static RetT functor(Args... args); }; template diff --git a/source/lib/rocprofiler-sdk/tests/CMakeLists.txt b/source/lib/rocprofiler-sdk/tests/CMakeLists.txt index e33377ff7f..5828bfcc0f 100644 --- a/source/lib/rocprofiler-sdk/tests/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk/tests/CMakeLists.txt @@ -59,7 +59,6 @@ gtest_add_tests( WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}) set(rocprofiler-lib-tests-env - "HSA_TOOLS_LIB=$" "LD_LIBRARY_PATH=${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_LIBDIR}:$ENV{LD_LIBRARY_PATH}") set_tests_properties( diff --git a/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp b/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp index f5aaa77de2..b567c5c7d2 100644 --- a/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp +++ b/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp @@ -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(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( + 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(); -} \ No newline at end of file +} diff --git a/source/scripts/run-ci.py b/source/scripts/run-ci.py index c966c72b71..fd212f582f 100755 --- a/source/scripts/run-ci.py +++ b/source/scripts/run-ci.py @@ -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", ""), diff --git a/source/scripts/setup-sanitizer-env.sh b/source/scripts/setup-sanitizer-env.sh index 7df9811542..9e576735aa 100755 --- a/source/scripts/setup-sanitizer-env.sh +++ b/source/scripts/setup-sanitizer-env.sh @@ -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 diff --git a/source/scripts/thread-sanitizer-suppr.txt b/source/scripts/thread-sanitizer-suppr.txt index 8c8f3f0d55..54d588d3cb 100644 --- a/source/scripts/thread-sanitizer-suppr.txt +++ b/source/scripts/thread-sanitizer-suppr.txt @@ -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 \ No newline at end of file +mutex:source/lib/common/synchronized.hpp diff --git a/tests/async-copy-tracing/CMakeLists.txt b/tests/async-copy-tracing/CMakeLists.txt index bfa8ab16d5..07f868e558 100644 --- a/tests/async-copy-tracing/CMakeLists.txt +++ b/tests/async-copy-tracing/CMakeLists.txt @@ -21,7 +21,6 @@ add_test(NAME test-async-copy-tracing-execute COMMAND $) set(async-copy-tracing-env "${PRELOAD_ENV}" - "HSA_TOOLS_LIB=$" "ROCPROFILER_TOOL_OUTPUT_FILE=async-copy-tracing-test.json" "LD_LIBRARY_PATH=$:$ENV{LD_LIBRARY_PATH}" ) diff --git a/tests/async-copy-tracing/validate.py b/tests/async-copy-tracing/validate.py index 0a85fc334d..5069519548 100644 --- a/tests/async-copy-tracing/validate.py +++ b/tests/async-copy-tracing/validate.py @@ -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())) diff --git a/tests/bin/multistream/multistream_app.cpp b/tests/bin/multistream/multistream_app.cpp index 35138a97b0..cdb5466f73 100644 --- a/tests/bin/multistream/multistream_app.cpp +++ b/tests/bin/multistream/multistream_app.cpp @@ -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)); } } diff --git a/tests/counter-collection/CMakeLists.txt b/tests/counter-collection/CMakeLists.txt index 9efd75dd67..75b758ed50 100644 --- a/tests/counter-collection/CMakeLists.txt +++ b/tests/counter-collection/CMakeLists.txt @@ -17,19 +17,17 @@ else() set(PRELOAD_ENV "LD_PRELOAD=$") 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 $) set_tests_properties( test-counter-collection-execute - PROPERTIES - TIMEOUT - 45 - LABELS - "integration-tests" - ENVIRONMENT - "${PRELOAD_ENV};HSA_TOOLS_LIB=$;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} diff --git a/tests/counter-collection/validate.py b/tests/counter-collection/validate.py index 42e9d6a9f8..a216cb15a5 100644 --- a/tests/counter-collection/validate.py +++ b/tests/counter-collection/validate.py @@ -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__": diff --git a/tests/kernel-tracing/CMakeLists.txt b/tests/kernel-tracing/CMakeLists.txt index 74e911fa67..5d2f7ea87f 100644 --- a/tests/kernel-tracing/CMakeLists.txt +++ b/tests/kernel-tracing/CMakeLists.txt @@ -26,7 +26,6 @@ add_test(NAME test-kernel-tracing-execute COMMAND $" "ROCPROFILER_TOOL_OUTPUT_FILE=kernel-tracing-test.json" "LD_LIBRARY_PATH=$:$ENV{LD_LIBRARY_PATH}" ) diff --git a/tests/kernel-tracing/validate.py b/tests/kernel-tracing/validate.py index 84808ff5bc..f5b985364a 100644 --- a/tests/kernel-tracing/validate.py +++ b/tests/kernel-tracing/validate.py @@ -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())) diff --git a/tests/rocprofv3/counter-collection/input1/CMakeLists.txt b/tests/rocprofv3/counter-collection/input1/CMakeLists.txt index 097f25587e..cdd28bba5e 100644 --- a/tests/rocprofv3/counter-collection/input1/CMakeLists.txt +++ b/tests/rocprofv3/counter-collection/input1/CMakeLists.txt @@ -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=$") +set(cc-env-pmc1 "${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}") set_tests_properties( rocprofv3-test-counter-collection-pmc1-execute diff --git a/tests/rocprofv3/counter-collection/input2/CMakeLists.txt b/tests/rocprofv3/counter-collection/input2/CMakeLists.txt index eefd04a20d..28c65ce1c5 100644 --- a/tests/rocprofv3/counter-collection/input2/CMakeLists.txt +++ b/tests/rocprofv3/counter-collection/input2/CMakeLists.txt @@ -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=$") +set(cc-env-pmc2 "${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}") set_tests_properties( rocprofv3-test-counter-collection-pmc2-execute diff --git a/tests/rocprofv3/counter-collection/list_metrics/CMakeLists.txt b/tests/rocprofv3/counter-collection/list_metrics/CMakeLists.txt index b660439d1c..929aaeeea8 100644 --- a/tests/rocprofv3/counter-collection/list_metrics/CMakeLists.txt +++ b/tests/rocprofv3/counter-collection/list_metrics/CMakeLists.txt @@ -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=$") +set(cc-env-list-metrics "${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}") set_tests_properties( rocprofv3-test-list-metrics-execute diff --git a/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt b/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt index 0579b18563..65b7f21bde 100644 --- a/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt +++ b/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt @@ -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=$") +set(tracing-env "${PRELOAD_ENV}") set_tests_properties( rocprofv3-test-trace-hip-in-libraries-execute PROPERTIES + TIMEOUT + 45 LABELS "integration-tests" ENVIRONMENT diff --git a/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt b/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt index 9ad9fe8ed2..e4e96f91f4 100644 --- a/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt +++ b/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt @@ -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=$") +set(cc-tracing-env "${PRELOAD_ENV}" "ROCPROF_LOG_LEVEL=${LOG_LEVEL}") set_tests_properties( rocprofv3-test-tracing-plus-cc-execute diff --git a/tests/rocprofv3/tracing/CMakeLists.txt b/tests/rocprofv3/tracing/CMakeLists.txt index 5c72cf70a1..3e6d58e15a 100644 --- a/tests/rocprofv3/tracing/CMakeLists.txt +++ b/tests/rocprofv3/tracing/CMakeLists.txt @@ -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=$") +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