PC Sampling Support (#715)
* cmake formatting (cmake-format) (#188) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * source formatting (clang-format v11) (#189) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: design of the pc sampling data struct; guarding parts of code that uses ROCr marker packets * source formatting (clang-format v11) (#191) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * cmake formatting (cmake-format) (#192) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: shadow variable fix * pcs: fix for compiler errors reported by CI/CD * source formatting (clang-format v11) (#193) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: docs fix; samples uses rocprofiler::rocprofiler library * cmake formatting (cmake-format) (#195) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: client in samples folder fixed * pcs: client requires rocprofiler package as dependency * pcs: client uses single context * source formatting (clang-format v11) (#196) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: client using single buffer; no buffer destroy in client * pcs: client::setup explicitly called from the example * pcs: rocprofiler_pc_sample_record_t updated * pcs: fixed init of external correlation id * source formatting (clang-format v11) (#198) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: remove outdated files; update CMakeLists * cmake formatting (cmake-format) (#212) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: using rocprofiler_agent_id_t * pcs: Removing trailing whitespaces Co-authored-by: Jonathan R. Madsen <jrmadsen@users.noreply.github.com> * source formatting (clang-format v11) (#214) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: mapping agent_id to the agent * source formatting (clang-format v11) (#215) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: const while iterating over agents * source formatting (clang-format v11) (#216) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: calling get_buffer instead of get_buffers * pcs: workgroup typo * pcs: documentation for the public PC sampling API * pcs: queue_cb_t signature adaptation * pcs: mocks removed * pcs: updating HsaApiTable with HSA/ROCr PC sampling API * pcs: querying available PC sampling configs through IOCTL * pcs: create the PCS session in IOCTL * pcs: first actual PC samples delivered to the rocprofiler's client :) * pcs: works with marker packet too * pcs: using HSA table to call pc sampling related functions * pcs: using ioctl instead of kfd in naming * pcs: configuration service test fixed * pcs: sample processing test fixed * pcs: marker packet macro wrapper removed * pcs: marker packet is part of the rocprofiler_packet union * pcs: one fixme added * pcs: client that uses pc-sampling and code obj tracing * pcs: client that supprts PC sampling and code obj tracing refactored * pcs: show more info for each PC sample * pcs: hex output for the samples that do not belong to the matmul kernel * pcs: querying avail configuration happens immediately before configuring * pcs: hsa_ven_amd_pcs_create_from_id renamed * pcs: using hsa_stop; accessing a buffer by id from parser * pcs: includes reworked, tests returned to life * pcs: rocrofiler dir removed as outdated * cmake formatting (cmake-format) (#271) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * source formatting (clang-format v11) (#272) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: some warnings fixed * source formatting (clang-format v11) (#273) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * cmake formatting (cmake-format) (#274) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: show MI200 relevant information in the sample * pcs: queue cb fixed; rocr.h include fixed * source formatting (clang-format v11) (#296) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: getting hsa_agent and the doorbell_id from hsa_queue * source formatting (clang-format v11) (#297) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: correlation ID logic fixed * source formatting (clang-format v11) (#303) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: pure pc sampling example fixed * source formatting (clang-format v11) (#307) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * cmake formatting (cmake-format) (#308) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: interval value if the PC sampling is already configured * pcs: ROCPROFILER_STATUS_ERROR_PC_SAMPLING_ALREADY_CONFIGURED New status code if another process configured PC sampling service with different configuration. Samples are extended to consider this case and retry if it happens. * pcs: hsa_amd_queue_get_info mocked in tests * source formatting (clang-format v11) (#328) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs (tests): query configs after configuring service * source formatting (clang-format v11) (#329) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: sample checks workgroup_id_* and wave_id * source formatting (clang-format v11) (#330) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs samples: running samples on the device 0 * pcs: kfd_ioctl updated * pcs: ioctl config struct changed fields names * pcs: status when PC sampling is configured by another process is renamed * pcs: HSA PC sampling API table fixed * pcs: tmp hack to be able to use HSA pc sampling table * source formatting (clang-format v11) (#443) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs service use CIDs generated by HIP API tracing service * source formatting (clang-format v11) (#455) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * cmake formatting (cmake-format) (#456) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: CID manager * pcs: explicit flush with no delivered data executes retirement logic * source formatting (clang-format v11) (#464) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: rocprofiler_query_pc_sampling_agent_configurations docs update * source formatting (clang-format v11) (#465) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: rocprofiler_configure_pc_sampling_service docs update * pcs: explicit sync introduced in PCSCIDManager * pcs: new logic for retiring CIDs in PC sampling service documented * pcs: queue interception cb signature updated * source formatting (clang-format v11) (#471) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: if no agents supports PC sampling, fail gracefully * elaborating when KFD returns EBUSY and EEXIST * pcs: the second PC sampling examples fails gracefully * code samples use only single kernel for now * pcs: CID manager refactored * source formatting (clang-format v11) (#481) Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> * pcs: ioctl update * source formatting (clang-format v11) (#531) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * pcs:code sample to test PC sampling applied on concurrent kernels * source formatting (clang-format v11) (#533) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * pcs: pc sampling strest test included * cmake formatting (cmake-format) (#539) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * source formatting (clang-format v11) (#540) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * pcs: standalone benchmark * cmake formatting (cmake-format) (#555) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * pcs: glance in external correlation IDs * source formatting (clang-format v11) (#557) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * another change in ioctl interface * pcs: update queue interceptor callbacks and samples accroding to the agent 0 version * source formatting (clang-format v11) (#611) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * pcs: avoid running problematic PC sampling test * pcs: guarding tests not to fail on architectures not supporting PC sampling * source formatting (clang-format v11) (#617) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * pcs: check IOCTL version prior to each KFD call * pcs: ioctl refactoring * pcs: PC sampling service increases the ref_count of the correlation ID of the kernel dispatch * cmake formatting (cmake-format) (#631) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * source formatting (clang-format v11) (#632) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * pcs: PC sampling service provides external correlation IDs * source formatting (clang-format v11) (#644) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * pcs: use rocprofiler_dim3_t for workgrou_ip * source formatting (clang-format v11) (#645) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * pcs: minor fixes * pcs: updating the documentation for the pc sampling API functions * pcs: api table and queue controller fix * pcs: don't generate marker packets for the agent if PC sampling is not configured on it * pcs: multi-GPU and single-GPU clients * source formatting (clang-format v11) (#700) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * pcs: warning and errors fixed * source formatting (clang-format v11) (#702) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * pcs: clang compiler errors and warnings fixed * source formatting (clang-format v11) (#716) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * pcs: const reference in cid manager * source formatting (clang-format v11) (#717) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * pcs: const & func in manager explicit * pcs: test to cover creating PC sampling service of agent that does not exist * pcs: generate marker packets if service is active * source formatting (clang-format v11) (#719) Co-authored-by: vlaindic <139573562+vlaindic@users.noreply.github.com> * pcs: refactoring hsa_adapter; use the correlation_id->thread_idx * Update source/lib/rocprofiler-sdk/pc_sampling/cid_manager.cpp * Update source/lib/rocprofiler-sdk/pc_sampling/cid_manager.cpp * Update source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp * Update source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp * Update source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp * Update source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp * Update source/lib/rocprofiler-sdk/pc_sampling/utils.cpp * Update utils.cpp * moving pc-sampling tests and samples to pc-sampling label * Format fix * pcs: use configured instead of active service * Update source/lib/rocprofiler-sdk/pc_sampling/service.cpp * pcs: ensure configuring PC sampling on the HSA level is called only once * pcs: minor fix * Update CMakeLists.txt * Update CMakeLists.txt * Update CMakeLists.txt * Update CMakeLists.txt * pcs: refactoring IOCTL integration * Update source/lib/rocprofiler-sdk/pc_sampling/tests/CMakeLists.txt Co-authored-by: Ammar ELWazir <ammar.elwazir@amd.com> * Update source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> * Update source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter_types.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> * Update source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> * Update source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter_types.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> * Update source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> * pcs: reverting back what bot doubled * Update source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter_types.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> * pcs: retesting the bot * Update source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter_types.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> * pcs: why bot fails on this IOCTL status * pcs: why failing on <vector> * Update source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> * pcs: returning commits removed by bot * pcs: formatting locally * pcs: clients are flushing buffers inside the tool_fini * pcs: sync function in public API * pcs: sync prior to unloading the code object * pcs: sync function requires context * pcs: client uses CID retirement service * pcs: test for flusing internal ROCr buffers * pcs: source formatting * Update source/lib/rocprofiler-sdk/pc_sampling/tests/CMakeLists.txt Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> * pcs: code samples refactoring * pcs: public API header refactored * pcs: rocprofiler_buffer_flush drains internal PC sampling buffers too * pcs: remove unnecessary functions * pcs: do not call hsa's copytables * pcs: include reordering * pcs: using ROCP_ERROR inside PC sampling implementation * pcs: pc_sampling sample uses ostream instean of printfs * pcs: pc_sampling_codeobj tracing using ostream instead of prints * pcs: registering once for interceptor callbacks * pcs: do not generate internal CIDs if not in debug mode * pcs: rebasing fixed; missing external correlation IDs * pcs: code formatting * enable kernel tracing service to receive external correlation IDs * pcs: using ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_KERNEL * pcs: polishing parser * formatting * updating parser to use workgroup_id * kfd_ioctl.h extracted in details folder * refactoring * pcs: preparing to generate code object information * flush internal buffers prior to unloading code object * pcs: generating marker records * pcs: wrap code_object's shutdown function * ROCR_VISIBLE_DEVICES and HIP_VISISBLE_DEVICES unsupported at the moment * documenting the ignorance of ROCR/HIP_VISIBLE_DEVICES * pcs: separate structs for code object loading/unloading markers * pcs: inst_pkt_t changed the namespace * pcs: removing wrapper around the shutdown function * pcs: size in record field * pcs: documentation refactoring + typdefs * renaming PCSAgentConfig to PCSAgentSession * pcs: service does not keep a pointer to the context * pcs: static assertions related to the versioning * pcs: rocprofiler_pc_sampling_configuration_t size field * pcs: report API unimplemented unleass explicitly enabled * pcs: skip tests if KFD does not support PC sampling * pcs: if ROCr hides some devices, no PC samples will be delivered for it * pcs: hip error check after kernel launch * formatting * removing PCS info from agent.h * fix based on review * Update continuous integration workflow - use mi200 runner for code coverage (supports PC sampling) - split sanitizer jobs across navi3, vega20, and mi300 * Updating pc sampling test labels * ROCP_PC_SAMPLING_ENABLED env in CI * ROCP_PC_SAMPLING_ENABLED for all CI mi200 jobs * Rearrange sanitizer assignments * fixes according to review * removed unused functions * pcs: rocprofiler_agent_id_t instead of handle as a key in map * Update source/lib/rocprofiler-sdk/context/context.hpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> * removing drm_fd from the agent.h * pcs: removing one sample due to complexity * pcs: refactoring sample * simplifying sample * new lines * Improve queue_control enable intercepter logic * Update lib/rocprofiler-sdk/hsa/types.hpp - handle amd_ext size for HSA 1.12.0 * ROCP_PC_SAMPLING_ENABLED -> ROCPROFILER_PC_SAMPLING_BETA_ENABLED * Update hsa_adapter.cpp - anonymous namespace + remove debug * parser update * Apply suggestions from code review --------- Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Co-authored-by: vlaindic <vlaindic@users.noreply.github.com> Co-authored-by: vlaindic <vladimir.indic@amd.com> Co-authored-by: vlaindic <vlaindic@amd.com> Co-authored-by: Vladimir Indic <139573562+vlaindic@users.noreply.github.com> Co-authored-by: Jonathan R. Madsen <jrmadsen@users.noreply.github.com> Co-authored-by: gobhardw <gopesh.bhardwaj@amd.com> Co-authored-by: Jonathan R. Madsen <jonathanrmadsen@gmail.com>
This commit is contained in:
Ammar ELWazir
GitHub
@@ -0,0 +1,63 @@
|
||||
#
|
||||
#
|
||||
#
|
||||
cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
|
||||
|
||||
if(NOT CMAKE_HIP_COMPILER)
|
||||
find_program(
|
||||
amdclangpp_EXECUTABLE
|
||||
NAMES amdclang++
|
||||
HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
|
||||
PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
|
||||
PATH_SUFFIXES bin llvm/bin NO_CACHE)
|
||||
mark_as_advanced(amdclangpp_EXECUTABLE)
|
||||
|
||||
if(amdclangpp_EXECUTABLE)
|
||||
set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}")
|
||||
endif()
|
||||
endif()
|
||||
|
||||
project(rocprofiler-sdk-samples-pc-sampling LANGUAGES CXX HIP)
|
||||
|
||||
foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO)
|
||||
if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "")
|
||||
set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}")
|
||||
endif()
|
||||
endforeach()
|
||||
|
||||
find_package(rocprofiler-sdk REQUIRED)
|
||||
|
||||
add_library(pc-sampling-client SHARED)
|
||||
target_sources(pc-sampling-client PRIVATE client.cpp pcs.hpp pcs.cpp utils.hpp utils.cpp)
|
||||
target_link_libraries(
|
||||
pc-sampling-client
|
||||
PRIVATE rocprofiler-sdk::rocprofiler-sdk rocprofiler::samples-build-flags
|
||||
rocprofiler::samples-common-library)
|
||||
|
||||
set_source_files_properties(main.cpp PROPERTIES LANGUAGE HIP)
|
||||
find_package(Threads REQUIRED)
|
||||
|
||||
add_executable(pc-sampling)
|
||||
target_sources(pc-sampling PRIVATE main.cpp)
|
||||
target_link_libraries(pc-sampling PRIVATE pc-sampling-client Threads::Threads
|
||||
rocprofiler::samples-build-flags)
|
||||
|
||||
rocprofiler_samples_get_preload_env(PRELOAD_ENV pc-sampling-client)
|
||||
rocprofiler_samples_get_ld_library_path_env(LIBRARY_PATH_ENV)
|
||||
|
||||
set(pc-sampling-env ${PRELOAD_ENV} ${LIBRARY_PATH_ENV})
|
||||
|
||||
add_test(NAME pc-sampling COMMAND $<TARGET_FILE:pc-sampling>)
|
||||
|
||||
set_tests_properties(
|
||||
pc-sampling
|
||||
PROPERTIES TIMEOUT
|
||||
45
|
||||
LABELS
|
||||
"samples;pc-sampling"
|
||||
ENVIRONMENT
|
||||
"${pc-sampling-env}"
|
||||
FAIL_REGULAR_EXPRESSION
|
||||
"${ROCPROFILER_DEFAULT_FAIL_REGEX}"
|
||||
SKIP_REGULAR_EXPRESSION
|
||||
"PC sampling unavailable")
|
||||
@@ -0,0 +1,230 @@
|
||||
// MIT License
|
||||
//
|
||||
// Copyright (c) 2023 ROCm Developer Tools
|
||||
//
|
||||
// Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
// of this software and associated documentation files (the "Software"), to deal
|
||||
// in the Software without restriction, including without limitation the rights
|
||||
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
// copies of the Software, and to permit persons to whom the Software is
|
||||
// furnished to do so, subject to the following conditions:
|
||||
//
|
||||
// The above copyright notice and this permission notice shall be included in all
|
||||
// copies or substantial portions of the Software.
|
||||
//
|
||||
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
|
||||
// SOFTWARE.
|
||||
|
||||
// undefine NDEBUG so asserts are implemented
|
||||
#ifdef NDEBUG
|
||||
# undef NDEBUG
|
||||
#endif
|
||||
|
||||
/**
|
||||
* @file samples/pc_sampling_library/client.cpp
|
||||
*
|
||||
* @brief Example rocprofiler client (tool)
|
||||
*/
|
||||
|
||||
#include "pcs.hpp"
|
||||
#include "utils.hpp"
|
||||
|
||||
#include <rocprofiler-sdk/buffer.h>
|
||||
#include <rocprofiler-sdk/fwd.h>
|
||||
#include <rocprofiler-sdk/internal_threading.h>
|
||||
#include <rocprofiler-sdk/registration.h>
|
||||
#include <rocprofiler-sdk/rocprofiler.h>
|
||||
|
||||
#include "common/defines.hpp"
|
||||
|
||||
#include <cassert>
|
||||
#include <chrono>
|
||||
#include <cstddef>
|
||||
#include <cstdint>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <filesystem>
|
||||
#include <fstream>
|
||||
#include <iostream>
|
||||
#include <memory>
|
||||
#include <regex>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
namespace client
|
||||
{
|
||||
namespace
|
||||
{
|
||||
rocprofiler_client_id_t* client_id = nullptr;
|
||||
rocprofiler_client_finalize_t client_fini_func = nullptr;
|
||||
rocprofiler_context_id_t client_ctx;
|
||||
|
||||
int
|
||||
tool_init(rocprofiler_client_finalize_t fini_func, void* /*tool_data*/)
|
||||
{
|
||||
client_fini_func = fini_func;
|
||||
|
||||
client::pcs::find_all_gpu_agents_supporting_pc_sampling();
|
||||
|
||||
if(client::pcs::gpu_agents.empty())
|
||||
{
|
||||
*utils::get_output_stream() << "No availabe gpu agents supporting PC sampling" << std::endl;
|
||||
// Exit with no error if none of the GPUs support PC sampling.
|
||||
exit(0);
|
||||
}
|
||||
|
||||
// The relations assumed:
|
||||
// - One context for all gpu agents
|
||||
// - a buffer per agent
|
||||
// - a callback thread per buffer
|
||||
// - a pc sampling service per agent/buffer
|
||||
|
||||
ROCPROFILER_CHECK(rocprofiler_create_context(&client_ctx));
|
||||
|
||||
for(auto& gpu_agent : pcs::gpu_agents)
|
||||
{
|
||||
// creating a buffer that will hold pc sampling information
|
||||
rocprofiler_buffer_policy_t drop_buffer_action = ROCPROFILER_BUFFER_POLICY_LOSSLESS;
|
||||
auto buffer_id = rocprofiler_buffer_id_t{};
|
||||
ROCPROFILER_CHECK(rocprofiler_create_buffer(client_ctx,
|
||||
client::pcs::BUFFER_SIZE_BYTES,
|
||||
client::pcs::WATERMARK,
|
||||
drop_buffer_action,
|
||||
client::pcs::rocprofiler_pc_sampling_callback,
|
||||
nullptr,
|
||||
&buffer_id));
|
||||
|
||||
client::pcs::configure_pc_sampling_prefer_stochastic(
|
||||
gpu_agent.get(), client_ctx, buffer_id);
|
||||
|
||||
// One helper thread per GPU agent's buffer.
|
||||
auto client_agent_thread = rocprofiler_callback_thread_t{};
|
||||
ROCPROFILER_CHECK(rocprofiler_create_callback_thread(&client_agent_thread));
|
||||
|
||||
ROCPROFILER_CHECK(rocprofiler_assign_callback_thread(buffer_id, client_agent_thread));
|
||||
|
||||
client::pcs::buffer_ids.emplace_back(buffer_id);
|
||||
}
|
||||
|
||||
int valid_ctx = 0;
|
||||
ROCPROFILER_CHECK(rocprofiler_context_is_valid(client_ctx, &valid_ctx));
|
||||
if(valid_ctx == 0)
|
||||
{
|
||||
// notify rocprofiler that initialization failed
|
||||
// and all the contexts, buffers, etc. created
|
||||
// should be ignored
|
||||
return -1;
|
||||
}
|
||||
|
||||
// Start PC sampling
|
||||
ROCPROFILER_CHECK(rocprofiler_start_context(client_ctx));
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
void
|
||||
tool_fini(void* /*tool_data*/)
|
||||
{
|
||||
if(client_id)
|
||||
{
|
||||
// Assert the context is inactive.
|
||||
int state = -1;
|
||||
ROCPROFILER_CHECK(rocprofiler_context_is_active(client_ctx, &state))
|
||||
assert(state == 0);
|
||||
|
||||
// No need to stop the context, since it has been stopped implicitly by the rocprofiler-SDK.
|
||||
for(size_t i = 0; i < client::pcs::buffer_ids.size(); i++)
|
||||
{
|
||||
// Flush the buffer explicitly
|
||||
ROCPROFILER_CHECK(rocprofiler_flush_buffer(client::pcs::buffer_ids.at(i)));
|
||||
// Destroying the buffer
|
||||
rocprofiler_status_t status = rocprofiler_destroy_buffer(client::pcs::buffer_ids.at(i));
|
||||
if(status == ROCPROFILER_STATUS_ERROR_BUFFER_BUSY)
|
||||
{
|
||||
*utils::get_output_stream()
|
||||
<< "The buffer is busy, so we cannot destroy it at the moment." << std::endl;
|
||||
}
|
||||
else
|
||||
{
|
||||
ROCPROFILER_CHECK(status);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
||||
// forward declaration
|
||||
void
|
||||
setup();
|
||||
|
||||
void
|
||||
setup()
|
||||
{
|
||||
if(int status = 0;
|
||||
rocprofiler_is_initialized(&status) == ROCPROFILER_STATUS_SUCCESS && status == 0)
|
||||
{
|
||||
ROCPROFILER_CHECK(rocprofiler_force_configure(&rocprofiler_configure));
|
||||
}
|
||||
}
|
||||
|
||||
void
|
||||
shutdown()
|
||||
{}
|
||||
|
||||
} // namespace client
|
||||
|
||||
extern "C" rocprofiler_tool_configure_result_t*
|
||||
rocprofiler_configure(uint32_t version,
|
||||
const char* runtime_version,
|
||||
uint32_t priority,
|
||||
rocprofiler_client_id_t* id)
|
||||
{
|
||||
// only activate if main tool
|
||||
if(priority > 0) return nullptr;
|
||||
|
||||
// set the client name
|
||||
id->name = "PCSamplingExampleTool";
|
||||
|
||||
// store client info
|
||||
client::client_id = id;
|
||||
|
||||
// compute major/minor/patch version info
|
||||
uint32_t major = version / 10000;
|
||||
uint32_t minor = (version % 10000) / 100;
|
||||
uint32_t patch = version % 100;
|
||||
|
||||
// generate info string
|
||||
auto info = std::stringstream{};
|
||||
info << id->name << " is using rocprofiler v" << major << "." << minor << "." << patch << " ("
|
||||
<< runtime_version << ")";
|
||||
|
||||
std::clog << info.str() << std::endl;
|
||||
|
||||
std::ostream* output_stream = nullptr;
|
||||
std::string filename = "pc_sampling.log";
|
||||
if(auto* outfile = getenv("ROCPROFILER_SAMPLE_OUTPUT_FILE"); outfile) filename = outfile;
|
||||
if(filename == "stdout")
|
||||
output_stream = &std::cout;
|
||||
else if(filename == "stderr")
|
||||
output_stream = &std::cerr;
|
||||
else
|
||||
output_stream = new std::ofstream{filename};
|
||||
|
||||
client::utils::get_output_stream() = output_stream;
|
||||
|
||||
// create configure data
|
||||
static auto cfg =
|
||||
rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t),
|
||||
&client::tool_init,
|
||||
&client::tool_fini,
|
||||
static_cast<void*>(output_stream)};
|
||||
|
||||
// return pointer to configure data
|
||||
return &cfg;
|
||||
}
|
||||
@@ -0,0 +1,412 @@
|
||||
// MIT License
|
||||
//
|
||||
// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
//
|
||||
// Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
// of this software and associated documentation files (the "Software"), to deal
|
||||
// in the Software without restriction, including without limitation the rights
|
||||
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
// copies of the Software, and to permit persons to whom the Software is
|
||||
// furnished to do so, subject to the following conditions:
|
||||
//
|
||||
// The above copyright notice and this permission notice shall be included in
|
||||
// all copies or substantial portions of the Software.
|
||||
//
|
||||
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
// THE SOFTWARE.
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
#include <chrono>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <iostream>
|
||||
#include <mutex>
|
||||
#include <random>
|
||||
#include <sstream>
|
||||
#include <stdexcept>
|
||||
|
||||
#define HIP_API_CALL(CALL) \
|
||||
{ \
|
||||
hipError_t error_ = (CALL); \
|
||||
if(error_ != hipSuccess) \
|
||||
{ \
|
||||
auto _hip_api_print_lk = auto_lock_t{print_lock}; \
|
||||
fprintf(stderr, \
|
||||
"%s:%d :: HIP error : %s\n", \
|
||||
__FILE__, \
|
||||
__LINE__, \
|
||||
hipGetErrorString(error_)); \
|
||||
throw std::runtime_error("hip_api_call"); \
|
||||
} \
|
||||
}
|
||||
|
||||
namespace
|
||||
{
|
||||
using auto_lock_t = std::unique_lock<std::mutex>;
|
||||
auto print_lock = std::mutex{};
|
||||
size_t nthread_per_device = 2;
|
||||
size_t nitr = 500;
|
||||
size_t nsync = 10;
|
||||
constexpr unsigned shared_mem_tile_dim = 32;
|
||||
|
||||
void
|
||||
check_hip_error(void);
|
||||
|
||||
void
|
||||
verify(int* in, int* out, int M, int N);
|
||||
} // namespace
|
||||
|
||||
__global__ void
|
||||
transpose(const int* in, int* out, int M, int N);
|
||||
|
||||
void
|
||||
run(int rank, int tid, int devid, int argc, char** argv);
|
||||
|
||||
void
|
||||
run_transpose(int rank, int tid, hipStream_t stream, int argc, char** argv);
|
||||
|
||||
void
|
||||
run_migrate(int rank, int tid, hipStream_t stream, int, char** argv);
|
||||
|
||||
void
|
||||
run_scratch(int rank, int tid, hipStream_t stream, int argc, char** argv);
|
||||
|
||||
int
|
||||
main(int argc, char** argv)
|
||||
{
|
||||
auto* exe_name = ::basename(argv[0]);
|
||||
|
||||
int rank = 0;
|
||||
for(int i = 1; i < argc; ++i)
|
||||
{
|
||||
auto _arg = std::string{argv[i]};
|
||||
if(_arg == "?" || _arg == "-h" || _arg == "--help")
|
||||
{
|
||||
fprintf(stderr,
|
||||
"usage: %s [NUM_THREADS_PER_DEVICE (%zu)] [NUM_ITERATION (%zu)] "
|
||||
"[SYNC_EVERY_N_ITERATIONS (%zu)]\n",
|
||||
exe_name,
|
||||
nthread_per_device,
|
||||
nitr,
|
||||
nsync);
|
||||
exit(EXIT_SUCCESS);
|
||||
}
|
||||
}
|
||||
if(argc > 1) nthread_per_device = atoll(argv[1]);
|
||||
if(argc > 2) nitr = atoll(argv[2]);
|
||||
if(argc > 3) nsync = atoll(argv[3]);
|
||||
|
||||
int ndevice = 0;
|
||||
HIP_API_CALL(hipGetDeviceCount(&ndevice));
|
||||
|
||||
auto nthreads = (ndevice * nthread_per_device);
|
||||
|
||||
printf("[%s] Number of devices found: %i\n", exe_name, ndevice);
|
||||
printf("[%s] Number of threads (per device): %zu\n", exe_name, nthread_per_device);
|
||||
printf("[%s] Number of threads (total): %zu\n", exe_name, nthreads);
|
||||
printf("[%s] Number of iterations: %zu\n", exe_name, nitr);
|
||||
printf("[%s] Syncing every %zu iterations\n", exe_name, nsync);
|
||||
|
||||
{
|
||||
auto _threads = std::vector<std::thread>{};
|
||||
for(size_t i = 0; i < nthreads; ++i)
|
||||
_threads.emplace_back(run, rank, i, i % ndevice, argc, argv);
|
||||
for(auto& itr : _threads)
|
||||
itr.join();
|
||||
}
|
||||
|
||||
HIP_API_CALL(hipDeviceSynchronize());
|
||||
HIP_API_CALL(hipDeviceReset());
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
__global__ void
|
||||
transpose(const int* in, int* out, int M, int N)
|
||||
{
|
||||
__shared__ int tile[shared_mem_tile_dim][shared_mem_tile_dim];
|
||||
|
||||
int idx = (blockIdx.y * blockDim.y + threadIdx.y) * M + blockIdx.x * blockDim.x + threadIdx.x;
|
||||
tile[threadIdx.y][threadIdx.x] = in[idx];
|
||||
__syncthreads();
|
||||
idx = (blockIdx.x * blockDim.x + threadIdx.y) * N + blockIdx.y * blockDim.y + threadIdx.x;
|
||||
out[idx] = tile[threadIdx.x][threadIdx.y];
|
||||
}
|
||||
|
||||
template <typename Tp>
|
||||
__global__ void
|
||||
test_page_migrate(Tp* data, Tp val)
|
||||
{
|
||||
int idx = (blockIdx.x * blockDim.x) + threadIdx.x;
|
||||
data[idx] += val;
|
||||
}
|
||||
|
||||
__global__ void
|
||||
test_kern_large(uint64_t* output)
|
||||
{
|
||||
uint64_t result = 0;
|
||||
int test[4000];
|
||||
memset(test, 5, 4000);
|
||||
for(int& i : test)
|
||||
{
|
||||
i = i + 7;
|
||||
*output += i;
|
||||
result += i;
|
||||
}
|
||||
*output ^= result;
|
||||
*output ^= result;
|
||||
}
|
||||
|
||||
__global__ void
|
||||
test_kern_medium(uint64_t* output)
|
||||
{
|
||||
uint64_t result = 0;
|
||||
int test[175];
|
||||
memset(test, 5, 175);
|
||||
for(int& i : test)
|
||||
{
|
||||
i = i + 7;
|
||||
*output += i;
|
||||
result += i;
|
||||
}
|
||||
*output ^= result;
|
||||
*output ^= result;
|
||||
}
|
||||
|
||||
__global__ void
|
||||
test_kern_small(uint64_t* output)
|
||||
{
|
||||
uint64_t result = 0;
|
||||
int test[2];
|
||||
for(int& i : test)
|
||||
{
|
||||
i = i + 7;
|
||||
*output += i;
|
||||
result += i;
|
||||
}
|
||||
*output ^= result;
|
||||
*output ^= result;
|
||||
}
|
||||
|
||||
void
|
||||
run(int rank, int tid, int devid, int argc, char** argv)
|
||||
{
|
||||
auto* stream = hipStream_t{};
|
||||
HIP_API_CALL(hipSetDevice(devid));
|
||||
HIP_API_CALL(hipStreamCreate(&stream));
|
||||
|
||||
run_migrate(rank, tid, stream, argc, argv);
|
||||
run_scratch(rank, tid, stream, argc, argv);
|
||||
run_transpose(rank, tid, stream, argc, argv);
|
||||
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
HIP_API_CALL(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
void
|
||||
run_transpose(int rank, int tid, hipStream_t stream, int argc, char** argv)
|
||||
{
|
||||
auto* exe_name = ::basename(argv[0]);
|
||||
|
||||
unsigned int M = 4960 * 2;
|
||||
unsigned int N = 4960 * 2;
|
||||
if(argc > 2) nitr = atoll(argv[2]);
|
||||
if(argc > 3) nsync = atoll(argv[3]);
|
||||
|
||||
auto_lock_t _lk{print_lock};
|
||||
std::cout << "[" << exe_name << "][transpose][" << rank << "][" << tid << "] M: " << M
|
||||
<< " N: " << N << std::endl;
|
||||
_lk.unlock();
|
||||
|
||||
std::default_random_engine _engine{std::random_device{}() * (rank + 1) * (tid + 1)};
|
||||
std::uniform_int_distribution<int> _dist{0, 1000};
|
||||
|
||||
size_t size = sizeof(int) * M * N;
|
||||
int* inp_matrix = new int[size];
|
||||
int* out_matrix = new int[size];
|
||||
for(size_t i = 0; i < M * N; i++)
|
||||
{
|
||||
inp_matrix[i] = _dist(_engine);
|
||||
out_matrix[i] = 0;
|
||||
}
|
||||
int* in = nullptr;
|
||||
int* out = nullptr;
|
||||
|
||||
HIP_API_CALL(hipMalloc(&in, size));
|
||||
HIP_API_CALL(hipMalloc(&out, size));
|
||||
HIP_API_CALL(hipMemsetAsync(in, 0, size, stream));
|
||||
HIP_API_CALL(hipMemsetAsync(out, 0, size, stream));
|
||||
HIP_API_CALL(hipMemcpyAsync(in, inp_matrix, size, hipMemcpyHostToDevice, stream));
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
|
||||
dim3 grid(M / 32, N / 32, 1);
|
||||
dim3 block(32, 32, 1); // transpose
|
||||
|
||||
print_lock.lock();
|
||||
printf("[%s][transpose][%i][%i] grid=(%i,%i,%i), block=(%i,%i,%i)\n",
|
||||
exe_name,
|
||||
rank,
|
||||
tid,
|
||||
grid.x,
|
||||
grid.y,
|
||||
grid.z,
|
||||
block.x,
|
||||
block.y,
|
||||
block.z);
|
||||
print_lock.unlock();
|
||||
|
||||
auto t1 = std::chrono::high_resolution_clock::now();
|
||||
for(size_t i = 0; i < nitr; ++i)
|
||||
{
|
||||
transpose<<<grid, block, 0, stream>>>(in, out, M, N);
|
||||
check_hip_error();
|
||||
if(i % nsync == (nsync - 1)) HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
}
|
||||
auto t2 = std::chrono::high_resolution_clock::now();
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
HIP_API_CALL(hipMemcpyAsync(out_matrix, out, size, hipMemcpyDeviceToHost, stream));
|
||||
double time = std::chrono::duration_cast<std::chrono::duration<double>>(t2 - t1).count();
|
||||
float GB = (float) size * nitr * 2 / (1 << 30);
|
||||
|
||||
print_lock.lock();
|
||||
std::cout << "[" << exe_name << "][transpose][" << rank << "][" << tid
|
||||
<< "] Runtime of transpose is " << time << " sec\n";
|
||||
std::cout << "[" << exe_name << "][transpose][" << rank << "][" << tid
|
||||
<< "] The average performance of transpose is " << GB / time << " GBytes/sec"
|
||||
<< std::endl;
|
||||
print_lock.unlock();
|
||||
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
|
||||
// cpu_transpose(matrix, out_matrix, M, N);
|
||||
verify(inp_matrix, out_matrix, M, N);
|
||||
|
||||
HIP_API_CALL(hipFree(in));
|
||||
HIP_API_CALL(hipFree(out));
|
||||
|
||||
delete[] inp_matrix;
|
||||
delete[] out_matrix;
|
||||
}
|
||||
|
||||
void
|
||||
run_scratch(int rank, int tid, hipStream_t stream, int, char** argv)
|
||||
{
|
||||
auto t1 = std::chrono::high_resolution_clock::now();
|
||||
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
|
||||
const auto* exe_name = ::basename(argv[0]);
|
||||
|
||||
uint64_t* data_ptr = nullptr;
|
||||
HIP_API_CALL(hipHostMalloc(&data_ptr, sizeof(uint64_t), 0));
|
||||
*data_ptr = 0;
|
||||
|
||||
test_kern_small<<<1000, 1, 0, stream>>>(data_ptr);
|
||||
test_kern_medium<<<1000, 1, 0, stream>>>(data_ptr);
|
||||
test_kern_small<<<1000, 1, 0, stream>>>(data_ptr);
|
||||
test_kern_large<<<1100, 1, 0, stream>>>(data_ptr);
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
|
||||
test_kern_small<<<1000, 1, 0, stream>>>(data_ptr);
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
|
||||
test_kern_medium<<<1000, 1, 0, stream>>>(data_ptr);
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
|
||||
test_kern_small<<<1000, 1, 0, stream>>>(data_ptr);
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
|
||||
test_kern_large<<<1100, 1, 0, stream>>>(data_ptr);
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
|
||||
auto t2 = std::chrono::high_resolution_clock::now();
|
||||
double time = std::chrono::duration_cast<std::chrono::duration<double>>(t2 - t1).count();
|
||||
|
||||
print_lock.lock();
|
||||
std::cout << "[" << exe_name << "][scratch][" << rank << "][" << tid
|
||||
<< "] Runtime of scratch is " << time << " sec\n";
|
||||
print_lock.unlock();
|
||||
}
|
||||
|
||||
void
|
||||
run_migrate(int rank, int tid, hipStream_t stream, int, char** argv)
|
||||
{
|
||||
using data_type = uint64_t;
|
||||
constexpr data_type init_v = 1;
|
||||
constexpr data_type incr_v = 1;
|
||||
|
||||
auto t1 = std::chrono::high_resolution_clock::now();
|
||||
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
|
||||
const auto* exe_name = ::basename(argv[0]);
|
||||
auto page_data = std::vector<data_type>(1024, 0);
|
||||
|
||||
HIP_API_CALL(hipHostRegister(
|
||||
page_data.data(), page_data.size() * sizeof(data_type), hipHostRegisterDefault));
|
||||
|
||||
for(auto& itr : page_data)
|
||||
itr = init_v;
|
||||
|
||||
test_page_migrate<<<1, 1024, 0, stream>>>(page_data.data(), incr_v);
|
||||
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
|
||||
for(auto& itr : page_data)
|
||||
{
|
||||
auto diff = (itr - incr_v);
|
||||
if(diff != init_v)
|
||||
{
|
||||
auto msg = std::stringstream{};
|
||||
msg << "invalid diff: " << diff << ". expected: " << init_v;
|
||||
throw std::runtime_error{msg.str()};
|
||||
}
|
||||
}
|
||||
|
||||
HIP_API_CALL(hipHostUnregister(page_data.data()));
|
||||
|
||||
auto t2 = std::chrono::high_resolution_clock::now();
|
||||
double time = std::chrono::duration_cast<std::chrono::duration<double>>(t2 - t1).count();
|
||||
|
||||
print_lock.lock();
|
||||
std::cout << "[" << exe_name << "][migrate][" << rank << "][" << tid
|
||||
<< "] Runtime of migrate is " << time << " sec\n";
|
||||
print_lock.unlock();
|
||||
}
|
||||
|
||||
namespace
|
||||
{
|
||||
void
|
||||
check_hip_error(void)
|
||||
{
|
||||
hipError_t err = hipGetLastError();
|
||||
if(err != hipSuccess)
|
||||
{
|
||||
auto_lock_t _lk{print_lock};
|
||||
std::cerr << "Error: " << hipGetErrorString(err) << std::endl;
|
||||
throw std::runtime_error("hip_api_call");
|
||||
}
|
||||
}
|
||||
|
||||
void
|
||||
verify(int* in, int* out, int M, int N)
|
||||
{
|
||||
for(int i = 0; i < 10; i++)
|
||||
{
|
||||
int row = rand() % M;
|
||||
int col = rand() % N;
|
||||
if(in[row * N + col] != out[col * M + row])
|
||||
{
|
||||
auto_lock_t _lk{print_lock};
|
||||
std::cout << "mismatch: " << row << ", " << col << " : " << in[row * N + col] << " | "
|
||||
<< out[col * M + row] << "\n";
|
||||
}
|
||||
}
|
||||
}
|
||||
} // namespace
|
||||
@@ -0,0 +1,317 @@
|
||||
// MIT License
|
||||
//
|
||||
// Copyright (c) 2023 ROCm Developer Tools
|
||||
//
|
||||
// Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
// of this software and associated documentation files (the "Software"), to deal
|
||||
// in the Software without restriction, including without limitation the rights
|
||||
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
// copies of the Software, and to permit persons to whom the Software is
|
||||
// furnished to do so, subject to the following conditions:
|
||||
//
|
||||
// The above copyright notice and this permission notice shall be included in all
|
||||
// copies or substantial portions of the Software.
|
||||
//
|
||||
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
|
||||
// SOFTWARE.
|
||||
|
||||
// undefine NDEBUG so asserts are implemented
|
||||
#ifdef NDEBUG
|
||||
# undef NDEBUG
|
||||
#endif
|
||||
|
||||
#include "pcs.hpp"
|
||||
#include "utils.hpp"
|
||||
|
||||
#include "common/defines.hpp"
|
||||
|
||||
#include <cassert>
|
||||
#include <cstdio>
|
||||
#include <iomanip>
|
||||
#include <memory>
|
||||
#include <sstream>
|
||||
|
||||
namespace client
|
||||
{
|
||||
namespace pcs
|
||||
{
|
||||
tool_agent_info_vec_t gpu_agents;
|
||||
std::vector<rocprofiler_buffer_id_t> buffer_ids;
|
||||
|
||||
rocprofiler_status_t
|
||||
find_all_gpu_agents_supporting_pc_sampling_impl(rocprofiler_agent_version_t version,
|
||||
const void** agents,
|
||||
size_t num_agents,
|
||||
void* user_data)
|
||||
{
|
||||
assert(version == ROCPROFILER_AGENT_INFO_VERSION_0);
|
||||
// user_data represent the pointer to the array where gpu_agent will be stored
|
||||
if(!user_data) return ROCPROFILER_STATUS_ERROR;
|
||||
|
||||
std::stringstream ss;
|
||||
|
||||
auto* _out_agents = static_cast<tool_agent_info_vec_t*>(user_data);
|
||||
auto* _agents = reinterpret_cast<const rocprofiler_agent_t**>(agents);
|
||||
for(size_t i = 0; i < num_agents; i++)
|
||||
{
|
||||
if(_agents[i]->type == ROCPROFILER_AGENT_TYPE_GPU)
|
||||
{
|
||||
// Instantiate the tool_agent_info.
|
||||
// Store pointer to the rocprofiler_agent_t and instatiate a vector of
|
||||
// available configurations.
|
||||
// Move the ownership to the _out_agents
|
||||
auto tool_gpu_agent = std::make_unique<tool_agent_info>();
|
||||
tool_gpu_agent->agent_id = _agents[i]->id;
|
||||
tool_gpu_agent->avail_configs = std::make_unique<avail_configs_vec_t>();
|
||||
tool_gpu_agent->agent = _agents[i];
|
||||
// Check if the GPU agent supports PC sampling. If so, add it to the
|
||||
// output list `_out_agents`.
|
||||
if(query_avail_configs_for_agent(tool_gpu_agent.get()))
|
||||
_out_agents->push_back(std::move(tool_gpu_agent));
|
||||
}
|
||||
|
||||
ss << "[" << __FUNCTION__ << "] " << _agents[i]->name << " :: "
|
||||
<< "id=" << _agents[i]->id.handle << ", "
|
||||
<< "type=" << _agents[i]->type << "\n";
|
||||
}
|
||||
|
||||
*utils::get_output_stream() << ss.str() << std::endl;
|
||||
|
||||
return ROCPROFILER_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
void
|
||||
find_all_gpu_agents_supporting_pc_sampling()
|
||||
{
|
||||
// This function returns the all gpu agents supporting some kind of PC sampling
|
||||
ROCPROFILER_CHECK(
|
||||
rocprofiler_query_available_agents(ROCPROFILER_AGENT_INFO_VERSION_0,
|
||||
&find_all_gpu_agents_supporting_pc_sampling_impl,
|
||||
sizeof(rocprofiler_agent_t),
|
||||
static_cast<void*>(&gpu_agents)));
|
||||
}
|
||||
|
||||
/**
|
||||
* @brief The function queries available PC sampling configurations.
|
||||
* If there is at least one available configuration, it returns true.
|
||||
* Otherwise, this function returns false to indicate the agent does
|
||||
* not support PC sampling.
|
||||
*/
|
||||
bool
|
||||
query_avail_configs_for_agent(tool_agent_info* agent_info)
|
||||
{
|
||||
// Clear the available configurations vector
|
||||
agent_info->avail_configs->clear();
|
||||
|
||||
auto cb = [](const rocprofiler_pc_sampling_configuration_t* configs,
|
||||
size_t num_config,
|
||||
void* user_data) {
|
||||
auto* avail_configs = static_cast<avail_configs_vec_t*>(user_data);
|
||||
for(size_t i = 0; i < num_config; i++)
|
||||
{
|
||||
avail_configs->emplace_back(configs[i]);
|
||||
}
|
||||
return ROCPROFILER_STATUS_SUCCESS;
|
||||
};
|
||||
|
||||
auto status = rocprofiler_query_pc_sampling_agent_configurations(
|
||||
agent_info->agent_id, cb, agent_info->avail_configs.get());
|
||||
|
||||
std::stringstream ss;
|
||||
|
||||
if(status != ROCPROFILER_STATUS_SUCCESS)
|
||||
{
|
||||
// The query operation failed, so consider the PC sampling is unsupported at the agent.
|
||||
// This can happen if the PC sampling service is invoked within the ROCgdb.
|
||||
ss << "Querying PC sampling capabilities failed with status: " << status << std::endl;
|
||||
*utils::get_output_stream() << ss.str() << std::endl;
|
||||
return false;
|
||||
}
|
||||
else if(agent_info->avail_configs->size() == 0)
|
||||
{
|
||||
// No available configuration at the moment, so mark the PC sampling as unsupported.
|
||||
return false;
|
||||
}
|
||||
|
||||
ss << "The agent with the id: " << agent_info->agent_id.handle << " supports the "
|
||||
<< agent_info->avail_configs->size() << " configurations: " << std::endl;
|
||||
size_t ind = 0;
|
||||
for(auto& cfg : *agent_info->avail_configs)
|
||||
{
|
||||
ss << "(" << ++ind << ".) "
|
||||
<< "method: " << cfg.method << ", "
|
||||
<< "unit: " << cfg.unit << ", "
|
||||
<< "min_interval: " << cfg.min_interval << ", "
|
||||
<< "max_interval: " << cfg.max_interval << ", "
|
||||
<< "flags: " << std::hex << cfg.flags << std::dec << std::endl;
|
||||
}
|
||||
|
||||
*utils::get_output_stream() << ss.str() << std::flush;
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
void
|
||||
configure_pc_sampling_prefer_stochastic(tool_agent_info* agent_info,
|
||||
rocprofiler_context_id_t context_id,
|
||||
rocprofiler_buffer_id_t buffer_id)
|
||||
{
|
||||
int failures = 10;
|
||||
size_t interval = 0;
|
||||
do
|
||||
{
|
||||
// Update the list of available configurations
|
||||
auto success = query_avail_configs_for_agent(agent_info);
|
||||
if(!success)
|
||||
{
|
||||
// An error occured while querying PC sampling capabilities,
|
||||
// so avoid trying configuring PC sampling service.
|
||||
// Instead return false to indicated a failure.
|
||||
ROCPROFILER_CHECK(ROCPROFILER_STATUS_ERROR);
|
||||
}
|
||||
|
||||
const rocprofiler_pc_sampling_configuration_t* first_host_trap_config = nullptr;
|
||||
const rocprofiler_pc_sampling_configuration_t* first_stochastic_config = nullptr;
|
||||
// Search until encountering on the stochastic configuration, if any.
|
||||
// Otherwise, use the host trap config
|
||||
for(auto const& cfg : *agent_info->avail_configs)
|
||||
{
|
||||
if(cfg.method == ROCPROFILER_PC_SAMPLING_METHOD_STOCHASTIC)
|
||||
{
|
||||
first_stochastic_config = &cfg;
|
||||
break;
|
||||
}
|
||||
else if(!first_host_trap_config &&
|
||||
cfg.method == ROCPROFILER_PC_SAMPLING_METHOD_HOST_TRAP)
|
||||
{
|
||||
first_host_trap_config = &cfg;
|
||||
}
|
||||
}
|
||||
|
||||
// Check if the stochastic config is found. Use host trap config otherwise.
|
||||
const rocprofiler_pc_sampling_configuration_t* picked_cfg =
|
||||
(first_stochastic_config != nullptr) ? first_stochastic_config : first_host_trap_config;
|
||||
|
||||
if(picked_cfg->min_interval == picked_cfg->max_interval)
|
||||
{
|
||||
// Another process already configured PC sampling, so use the intreval it set up.
|
||||
interval = picked_cfg->min_interval;
|
||||
}
|
||||
else
|
||||
{
|
||||
interval = 10000;
|
||||
}
|
||||
|
||||
auto status = rocprofiler_configure_pc_sampling_service(context_id,
|
||||
agent_info->agent_id,
|
||||
picked_cfg->method,
|
||||
picked_cfg->unit,
|
||||
interval,
|
||||
buffer_id);
|
||||
if(status == ROCPROFILER_STATUS_SUCCESS)
|
||||
{
|
||||
*utils::get_output_stream()
|
||||
<< ">>> We chose PC sampling interval: " << interval
|
||||
<< " on the agent: " << agent_info->agent->id.handle << std::endl;
|
||||
return;
|
||||
}
|
||||
else if(status != ROCPROFILER_STATUS_ERROR_NOT_AVAILABLE)
|
||||
{
|
||||
ROCPROFILER_CHECK(status);
|
||||
}
|
||||
// status == ROCPROFILER_STATUS_ERROR_NOT_AVAILABLE
|
||||
// means another process P2 already configured PC sampling.
|
||||
// Query available configurations again and receive the configurations picked by P2.
|
||||
// However, if P2 destroys PC sampling service after query function finished,
|
||||
// but before the `rocprofiler_configure_pc_sampling_service` is called,
|
||||
// then the `rocprofiler_configure_pc_sampling_service` will fail again.
|
||||
// The process P1 executing this loop can spin wait (starve) if it is unlucky enough
|
||||
// to always be interuppted by some other process P2 that creates/destroys
|
||||
// PC sampling service on the same device while P1 is executing the code
|
||||
// after the `query_avail_configs_for_agent` and
|
||||
// before the `rocprofiler_configure_pc_sampling_service`.
|
||||
// This should happen very rarely, but just to be sure, we introduce a counter `failures`
|
||||
// that will allow certain amount of failures to process P1.
|
||||
} while(--failures);
|
||||
|
||||
// The process failed too many times configuring PC sampling,
|
||||
// report this to user;
|
||||
ROCPROFILER_CHECK(ROCPROFILER_STATUS_ERROR);
|
||||
}
|
||||
|
||||
void
|
||||
rocprofiler_pc_sampling_callback(rocprofiler_context_id_t /*context_id*/,
|
||||
rocprofiler_buffer_id_t /*buffer_id*/,
|
||||
rocprofiler_record_header_t** headers,
|
||||
size_t num_headers,
|
||||
void* /*data*/,
|
||||
uint64_t drop_count)
|
||||
{
|
||||
std::stringstream ss;
|
||||
ss << "The number of delivered samples is: " << num_headers << ", "
|
||||
<< "while the number of dropped samples is: " << drop_count << std::endl;
|
||||
|
||||
for(size_t i = 0; i < num_headers; i++)
|
||||
{
|
||||
auto* cur_header = headers[i];
|
||||
|
||||
if(cur_header == nullptr)
|
||||
{
|
||||
throw std::runtime_error{
|
||||
"rocprofiler provided a null pointer to header. this should never happen"};
|
||||
}
|
||||
else if(cur_header->hash !=
|
||||
rocprofiler_record_header_compute_hash(cur_header->category, cur_header->kind))
|
||||
{
|
||||
throw std::runtime_error{"rocprofiler_record_header_t (category | kind) != hash"};
|
||||
}
|
||||
else if(cur_header->category == ROCPROFILER_BUFFER_CATEGORY_PC_SAMPLING)
|
||||
{
|
||||
if(cur_header->kind == ROCPROFILER_PC_SAMPLING_RECORD_SAMPLE)
|
||||
{
|
||||
auto* pc_sample =
|
||||
static_cast<rocprofiler_pc_sampling_record_t*>(cur_header->payload);
|
||||
|
||||
ss << "pc: " << std::hex << pc_sample->pc << ", "
|
||||
<< "timestamp: " << std::dec << pc_sample->timestamp << ", "
|
||||
<< "exec: " << std::hex << std::setw(16) << pc_sample->exec_mask << ", "
|
||||
<< "workgroup_id_(x=" << std::dec << std::setw(5) << pc_sample->workgroup_id.x
|
||||
<< ", "
|
||||
<< "y=" << std::setw(5) << pc_sample->workgroup_id.y << ", "
|
||||
<< "z=" << std::setw(5) << pc_sample->workgroup_id.z << "), "
|
||||
<< "wave_id: " << std::setw(2) << static_cast<unsigned int>(pc_sample->wave_id)
|
||||
<< ", "
|
||||
<< "cu_id: " << pc_sample->hw_id << ", "
|
||||
<< "correlation: {internal=" << std::setw(7)
|
||||
<< pc_sample->correlation_id.internal << ", "
|
||||
<< "external=" << std::setw(5) << pc_sample->correlation_id.external.value << "}"
|
||||
<< std::endl;
|
||||
}
|
||||
else if(cur_header->kind == ROCPROFILER_PC_SAMPLING_RECORD_CODE_OBJECT_LOAD_MARKER)
|
||||
{
|
||||
auto* marker = static_cast<rocprofiler_pc_sampling_code_object_load_marker_t*>(
|
||||
cur_header->payload);
|
||||
ss << "code object loading: " << marker->code_object_id << std::endl;
|
||||
}
|
||||
else if(cur_header->kind == ROCPROFILER_PC_SAMPLING_RECORD_CODE_OBJECT_UNLOAD_MARKER)
|
||||
{
|
||||
auto* marker = static_cast<rocprofiler_pc_sampling_code_object_unload_marker_t*>(
|
||||
cur_header->payload);
|
||||
ss << "code object unloading: " << marker->code_object_id << std::endl;
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
throw std::runtime_error{"unexpected rocprofiler_record_header_t category + kind"};
|
||||
}
|
||||
}
|
||||
|
||||
*utils::get_output_stream() << ss.str() << std::endl;
|
||||
}
|
||||
} // namespace pcs
|
||||
} // namespace client
|
||||
@@ -0,0 +1,79 @@
|
||||
// MIT License
|
||||
//
|
||||
// Copyright (c) 2023 ROCm Developer Tools
|
||||
//
|
||||
// Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
// of this software and associated documentation files (the "Software"), to deal
|
||||
// in the Software without restriction, including without limitation the rights
|
||||
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
// copies of the Software, and to permit persons to whom the Software is
|
||||
// furnished to do so, subject to the following conditions:
|
||||
//
|
||||
// The above copyright notice and this permission notice shall be included in all
|
||||
// copies or substantial portions of the Software.
|
||||
//
|
||||
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
|
||||
// SOFTWARE.
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <rocprofiler-sdk/fwd.h>
|
||||
#include <rocprofiler-sdk/rocprofiler.h>
|
||||
|
||||
#include <vector>
|
||||
|
||||
namespace client
|
||||
{
|
||||
namespace pcs
|
||||
{
|
||||
constexpr size_t BUFFER_SIZE_BYTES = 8192;
|
||||
constexpr size_t WATERMARK = (BUFFER_SIZE_BYTES / 4);
|
||||
|
||||
struct tool_agent_info;
|
||||
using avail_configs_vec_t = std::vector<rocprofiler_pc_sampling_configuration_t>;
|
||||
using tool_agent_info_vec_t = std::vector<std::unique_ptr<tool_agent_info>>;
|
||||
|
||||
struct tool_agent_info
|
||||
{
|
||||
rocprofiler_agent_id_t agent_id;
|
||||
std::unique_ptr<avail_configs_vec_t> avail_configs;
|
||||
const rocprofiler_agent_t* agent;
|
||||
};
|
||||
|
||||
// GPU agents supporting some kind of PC sampling.
|
||||
// Note that for some of these agent, the corresponding context might be invalid,
|
||||
// meaning we were not able to enable PC sampling service.
|
||||
// Check the `tool_init` for more information.
|
||||
extern tool_agent_info_vec_t gpu_agents;
|
||||
// Ids of the buffers used as containers for PC sampling records
|
||||
extern std::vector<rocprofiler_buffer_id_t> buffer_ids;
|
||||
|
||||
void
|
||||
find_all_gpu_agents_supporting_pc_sampling();
|
||||
|
||||
/**
|
||||
* @brief The return value indicates if the agent supports PC sampling.
|
||||
* Check the implementation for more info.
|
||||
*/
|
||||
bool
|
||||
query_avail_configs_for_agent(tool_agent_info* agent_info);
|
||||
|
||||
void
|
||||
configure_pc_sampling_prefer_stochastic(tool_agent_info* agent_info,
|
||||
rocprofiler_context_id_t context_id,
|
||||
rocprofiler_buffer_id_t buffer_id);
|
||||
|
||||
void
|
||||
rocprofiler_pc_sampling_callback(rocprofiler_context_id_t context_id,
|
||||
rocprofiler_buffer_id_t buffer_id,
|
||||
rocprofiler_record_header_t** headers,
|
||||
size_t num_headers,
|
||||
void* data,
|
||||
uint64_t drop_count);
|
||||
} // namespace pcs
|
||||
} // namespace client
|
||||
@@ -0,0 +1,37 @@
|
||||
// MIT License
|
||||
//
|
||||
// Copyright (c) 2023 ROCm Developer Tools
|
||||
//
|
||||
// Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
// of this software and associated documentation files (the "Software"), to deal
|
||||
// in the Software without restriction, including without limitation the rights
|
||||
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
// copies of the Software, and to permit persons to whom the Software is
|
||||
// furnished to do so, subject to the following conditions:
|
||||
//
|
||||
// The above copyright notice and this permission notice shall be included in all
|
||||
// copies or substantial portions of the Software.
|
||||
//
|
||||
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
|
||||
// SOFTWARE.
|
||||
|
||||
#include "utils.hpp"
|
||||
|
||||
namespace client
|
||||
{
|
||||
namespace utils
|
||||
{
|
||||
std::ostream*&
|
||||
get_output_stream()
|
||||
{
|
||||
// The output strea is initially unitialized
|
||||
static std::ostream* _v = nullptr;
|
||||
return _v;
|
||||
}
|
||||
} // namespace utils
|
||||
} // namespace client
|
||||
@@ -0,0 +1,36 @@
|
||||
// MIT License
|
||||
//
|
||||
// Copyright (c) 2023 ROCm Developer Tools
|
||||
//
|
||||
// Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
// of this software and associated documentation files (the "Software"), to deal
|
||||
// in the Software without restriction, including without limitation the rights
|
||||
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
// copies of the Software, and to permit persons to whom the Software is
|
||||
// furnished to do so, subject to the following conditions:
|
||||
//
|
||||
// The above copyright notice and this permission notice shall be included in all
|
||||
// copies or substantial portions of the Software.
|
||||
//
|
||||
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
|
||||
// SOFTWARE.
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <rocprofiler-sdk/fwd.h>
|
||||
|
||||
#include <iostream>
|
||||
|
||||
namespace client
|
||||
{
|
||||
namespace utils
|
||||
{
|
||||
std::ostream*&
|
||||
get_output_stream();
|
||||
}
|
||||
} // namespace client
|
||||
Reference in New Issue
Block a user