api_buffered_tracing/client.cpp#
|
Rocprofiler SDK Developer API 0.4.0
ROCm Profiling API and tools
|
api_buffered_tracing/client.cpp
Example demonstrating Asynchronous Tracing Service that includes usage of rocprofiler_at_internal_thread_create, rocprofiler_create_callback_thread, and rocprofiler_assign_callback_thread.
// 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.
//
// undefine NDEBUG so asserts are implemented
#ifdef NDEBUG
# undef NDEBUG
#endif
/**
* @file samples/api_buffered_tracing/client.cpp
*
* @brief Example rocprofiler client (tool)
*/
#include "client.hpp"
#include <rocprofiler-sdk/buffer.h>
#include <rocprofiler-sdk/buffer_tracing.h>
#include <rocprofiler-sdk/callback_tracing.h>
#include <rocprofiler-sdk/external_correlation.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/internal_threading.h>
#include <rocprofiler-sdk/registration.h>
#include <rocprofiler-sdk/rocprofiler.h>
#include "common/call_stack.hpp"
#include "common/defines.hpp"
#include "common/filesystem.hpp"
#include "common/name_info.hpp"
#include <atomic>
#include <cassert>
#include <chrono>
#include <cmath>
#include <cstddef>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <fstream>
#include <functional>
#include <iomanip>
#include <iostream>
#include <map>
#include <mutex>
#include <sstream>
#include <string>
#include <string_view>
#include <thread>
#include <unordered_set>
#include <vector>
namespace client
{
namespace
{
using common::buffer_name_info;
using common::call_stack_t;
using common::source_location;
using kernel_symbol_data_t = rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t;
using kernel_symbol_map_t = std::unordered_map<rocprofiler_kernel_id_t, kernel_symbol_data_t>;
rocprofiler_context_id_t client_ctx = {};
rocprofiler_buffer_id_t client_buffer = {};
buffer_name_info client_name_info = {};
kernel_symbol_map_t client_kernels = {};
void
print_call_stack(const call_stack_t& _call_stack)
{
common::print_call_stack("api_buffered_trace.log", _call_stack);
}
void
tool_code_object_callback(rocprofiler_callback_tracing_record_t record,
rocprofiler_user_data_t* user_data,
void* callback_data)
{
{
{
// flush the buffer to ensure that any lookups for the client kernel names for the code
// object are completed
ROCPROFILER_CALL(flush_status, "buffer flush");
}
}
{
{
client_kernels.emplace(data->kernel_id, *data);
}
{
client_kernels.erase(data->kernel_id);
}
}
(void) user_data;
(void) callback_data;
}
void
tool_tracing_callback(rocprofiler_context_id_t context,
rocprofiler_buffer_id_t buffer_id,
rocprofiler_record_header_t** headers,
size_t num_headers,
void* user_data,
uint64_t drop_count)
{
assert(user_data != nullptr);
assert(drop_count == 0 && "drop count should be zero for lossless policy");
if(num_headers == 0)
throw std::runtime_error{
"rocprofiler invoked a buffer callback with no headers. this should never happen"};
else if(headers == nullptr)
throw std::runtime_error{"rocprofiler invoked a buffer callback with a null pointer to the "
"array of headers. this should never happen"};
for(size_t i = 0; i < num_headers; ++i)
{
auto* header = headers[i];
auto kind_name = std::string{};
{
const char* _name = nullptr;
"query buffer tracing kind name");
if(_name)
{
static size_t len = 15;
kind_name = std::string{_name};
len = std::max(len, kind_name.length());
kind_name.resize(len, ' ');
kind_name += " :: ";
}
}
if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING &&
(header->kind == ROCPROFILER_BUFFER_TRACING_HSA_CORE_API ||
header->kind == ROCPROFILER_BUFFER_TRACING_HSA_AMD_EXT_API ||
header->kind == ROCPROFILER_BUFFER_TRACING_HSA_IMAGE_EXT_API ||
header->kind == ROCPROFILER_BUFFER_TRACING_HSA_FINALIZE_EXT_API))
{
auto* record =
auto info = std::stringstream{};
<< ", start=" << record->start_timestamp << ", stop=" << record->end_timestamp
if(record->start_timestamp > record->end_timestamp)
{
auto msg = std::stringstream{};
msg << "hsa api: start > end (" << record->start_timestamp << " > "
<< record->end_timestamp
<< "). diff = " << (record->start_timestamp - record->end_timestamp);
std::cerr << "threw an exception " << msg.str() << "\n" << std::flush;
// throw std::runtime_error{msg.str()};
}
static_cast<call_stack_t*>(user_data)->emplace_back(
source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()});
}
header->kind == ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API)
{
auto* record =
auto info = std::stringstream{};
<< ", buffer_id=" << buffer_id.handle
<< ", start=" << record->start_timestamp << ", stop=" << record->end_timestamp
if(record->start_timestamp > record->end_timestamp)
{
auto msg = std::stringstream{};
msg << "hip api: start > end (" << record->start_timestamp << " > "
<< record->end_timestamp
<< "). diff = " << (record->start_timestamp - record->end_timestamp);
std::cerr << "threw an exception " << msg.str() << "\n" << std::flush;
// throw std::runtime_error{msg.str()};
}
static_cast<call_stack_t*>(user_data)->emplace_back(
source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()});
}
header->kind == ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH)
{
auto* record =
auto info = std::stringstream{};
<< ", buffer_id=" << buffer_id.handle
<< ", agent_id=" << record->dispatch_info.agent_id.handle
<< ", queue_id=" << record->dispatch_info.queue_id.handle
<< ", kernel_id=" << record->dispatch_info.kernel_id
<< ", kernel=" << client_kernels.at(record->dispatch_info.kernel_id).kernel_name
<< ", start=" << record->start_timestamp << ", stop=" << record->end_timestamp
<< ", private_segment_size=" << record->dispatch_info.private_segment_size
<< ", group_segment_size=" << record->dispatch_info.group_segment_size
<< ", workgroup_size=(" << record->dispatch_info.workgroup_size.x << ","
<< record->dispatch_info.workgroup_size.y << ","
<< record->dispatch_info.workgroup_size.z << "), grid_size=("
<< record->dispatch_info.grid_size.x << "," << record->dispatch_info.grid_size.y
<< "," << record->dispatch_info.grid_size.z << ")";
if(record->start_timestamp > record->end_timestamp)
throw std::runtime_error("kernel dispatch: start > end");
static_cast<call_stack_t*>(user_data)->emplace_back(
source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()});
}
header->kind == ROCPROFILER_BUFFER_TRACING_MEMORY_COPY)
{
auto* record =
auto info = std::stringstream{};
<< ", buffer_id=" << buffer_id.handle
<< ", src_agent_id=" << record->src_agent_id.handle
<< ", dst_agent_id=" << record->dst_agent_id.handle
<< ", stop=" << record->end_timestamp
if(record->start_timestamp > record->end_timestamp)
throw std::runtime_error("memory copy: start > end");
static_cast<call_stack_t*>(user_data)->emplace_back(
source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()});
}
header->kind == ROCPROFILER_BUFFER_TRACING_PAGE_MIGRATION)
{
auto* record =
auto info = std::stringstream{};
<< ", pid=" << record->pid << ", start=" << record->start_timestamp
<< ", stop=" << record->end_timestamp
switch(record->operation)
{
{
info << ", page_fault=(" << record->page_fault.read_fault << ", "
<< record->page_fault.migrated << ", " << record->page_fault.node_id
<< ", " << std::hex << "0x" << record->page_fault.address << ")";
break;
}
{
info << ", page_migrate=(" << std::hex << "0x"
<< record->page_migrate.start_addr << ", 0x"
<< record->page_migrate.end_addr << ", " << std::dec
<< record->page_migrate.from_node << ", " << record->page_migrate.to_node
<< ", " << record->page_migrate.prefetch_node << ", "
<< record->page_migrate.preferred_node << ", "
<< record->page_migrate.trigger << ")";
break;
}
{
info << ", queue_suspend=(" << record->queue_suspend.rescheduled << ", "
<< record->queue_suspend.node_id << ", " << record->queue_suspend.trigger
<< ")";
break;
}
{
info << ", unmap_from_gpu=(" << record->unmap_from_gpu.node_id << std::hex
<< ", 0x" << record->unmap_from_gpu.start_addr << ", 0x"
<< record->unmap_from_gpu.end_addr << ", " << std::dec
<< record->unmap_from_gpu.trigger << ")";
break;
}
{
throw std::runtime_error{"unexpected page migration value"};
break;
}
}
if(record->start_timestamp > record->end_timestamp)
throw std::runtime_error("page migration: start > end");
static_cast<call_stack_t*>(user_data)->emplace_back(
source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()});
}
header->kind == ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY)
{
auto* record =
auto info = std::stringstream{};
auto _elapsed =
std::chrono::duration_cast<std::chrono::duration<double, std::micro>>(
std::chrono::nanoseconds{record->end_timestamp - record->start_timestamp})
.count();
<< ", buffer_id=" << buffer_id.handle
<< ", agent_id=" << record->agent_id.handle
<< ", elapsed=" << std::setprecision(3) << std::fixed << _elapsed
<< " usec, flags=" << record->flags
static_cast<call_stack_t*>(user_data)->emplace_back(
source_location{__FUNCTION__, __FILE__, __LINE__, kind_name + info.str()});
}
else
{
auto _msg = std::stringstream{};
_msg << "unexpected rocprofiler_record_header_t category + kind: (" << header->category
throw std::runtime_error{_msg.str()};
}
}
}
void
thread_precreate(rocprofiler_runtime_library_t lib, void* tool_data)
{
static_cast<call_stack_t*>(tool_data)->emplace_back(
source_location{__FUNCTION__,
__FILE__,
__LINE__,
std::string{"internal thread about to be created by rocprofiler (lib="} +
std::to_string(lib) + ")"});
}
void
thread_postcreate(rocprofiler_runtime_library_t lib, void* tool_data)
{
static_cast<call_stack_t*>(tool_data)->emplace_back(
source_location{__FUNCTION__,
__FILE__,
__LINE__,
std::string{"internal thread was created by rocprofiler (lib="} +
std::to_string(lib) + ")"});
}
int
tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
{
assert(tool_data != nullptr);
auto* call_stack_v = static_cast<call_stack_t*>(tool_data);
call_stack_v->emplace_back(source_location{__FUNCTION__, __FILE__, __LINE__, ""});
client_name_info = common::get_buffer_tracing_names();
for(const auto& itr : client_name_info)
{
auto name_idx = std::stringstream{};
name_idx << " [" << std::setw(3) << itr.value << "]";
call_stack_v->emplace_back(
source_location{"rocprofiler_buffer_tracing_kind_names " + name_idx.str(),
__FILE__,
__LINE__,
std::string{itr.name}});
for(auto [didx, ditr] : itr.items())
{
auto operation_idx = std::stringstream{};
operation_idx << " [" << std::setw(3) << didx << "]";
call_stack_v->emplace_back(source_location{
"rocprofiler_buffer_tracing_kind_operation_names" + operation_idx.str(),
__FILE__,
__LINE__,
std::string{"- "} + std::string{*ditr}});
}
}
client_fini_func = fini_func;
auto code_object_ops = std::vector<rocprofiler_tracing_operation_t>{
ROCPROFILER_CALL(
rocprofiler_configure_callback_tracing_service(client_ctx,
code_object_ops.data(),
code_object_ops.size(),
tool_code_object_callback,
nullptr),
"code object tracing service configure");
constexpr auto buffer_size_bytes = 4096;
constexpr auto buffer_watermark_bytes = buffer_size_bytes - (buffer_size_bytes / 8);
ROCPROFILER_CALL(rocprofiler_create_buffer(client_ctx,
buffer_size_bytes,
buffer_watermark_bytes,
tool_tracing_callback,
tool_data,
&client_buffer),
"buffer creation");
for(auto itr :
{
ROCPROFILER_CALL(rocprofiler_configure_buffer_tracing_service(
client_ctx, itr, nullptr, 0, client_buffer),
"buffer tracing service configure");
}
ROCPROFILER_CALL(
client_ctx, ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API, nullptr, 0, client_buffer),
"buffer tracing service configure");
ROCPROFILER_CALL(
client_ctx, ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, nullptr, 0, client_buffer),
"buffer tracing service for kernel dispatch configure");
ROCPROFILER_CALL(
client_ctx, ROCPROFILER_BUFFER_TRACING_MEMORY_COPY, nullptr, 0, client_buffer),
"buffer tracing service for memory copy configure");
// May have incompatible kernel so only emit a warning here
ROCPROFILER_WARN(rocprofiler_configure_buffer_tracing_service(
client_ctx, ROCPROFILER_BUFFER_TRACING_PAGE_MIGRATION, nullptr, 0, client_buffer));
ROCPROFILER_CALL(
client_ctx, ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY, nullptr, 0, client_buffer),
"buffer tracing service for page migration configure");
ROCPROFILER_CALL(rocprofiler_create_callback_thread(&client_thread),
"creating callback thread");
ROCPROFILER_CALL(rocprofiler_assign_callback_thread(client_buffer, client_thread),
"assignment of thread for buffer");
int valid_ctx = 0;
ROCPROFILER_CALL(rocprofiler_context_is_valid(client_ctx, &valid_ctx),
"context validity check");
if(valid_ctx == 0)
{
// notify rocprofiler that initialization failed
// and all the contexts, buffers, etc. created
// should be ignored
return -1;
}
// no errors
return 0;
}
void
tool_fini(void* tool_data)
{
assert(tool_data != nullptr);
auto* _call_stack = static_cast<call_stack_t*>(tool_data);
_call_stack->emplace_back(source_location{__FUNCTION__, __FILE__, __LINE__, ""});
print_call_stack(*_call_stack);
delete _call_stack;
}
} // namespace
void
setup()
{
if(int status = 0;
{
"force configuration");
}
}
void
shutdown()
{
if(client_id)
{
ROCPROFILER_CALL(rocprofiler_flush_buffer(client_buffer), "buffer flush");
client_fini_func(*client_id);
}
}
void
start()
{
ROCPROFILER_CALL(rocprofiler_start_context(client_ctx), "context start");
}
void
identify(uint64_t val)
{
auto _tid = rocprofiler_thread_id_t{};
rocprofiler_get_thread_id(&_tid);
rocprofiler_user_data_t user_data = {};
user_data.value = val;
rocprofiler_push_external_correlation_id(client_ctx, _tid, user_data);
}
void
stop()
{
}
} // namespace client
rocprofiler_configure(uint32_t version,
const char* runtime_version,
uint32_t priority,
{
// set the client name
id->name = "ExampleTool";
// 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{};
<< minor << "." << patch << " (" << runtime_version << ")";
std::clog << info.str() << std::endl;
auto* client_tool_data = new std::vector<client::source_location>{};
client_tool_data->emplace_back(
client::source_location{__FUNCTION__, __FILE__, __LINE__, info.str()});
ROCPROFILER_CALL(rocprofiler_at_internal_thread_create(
client::thread_precreate,
client::thread_postcreate,
static_cast<void*>(client_tool_data)),
"registration for thread creation notifications");
// create configure data
static auto cfg =
&client::tool_init,
&client::tool_fini,
static_cast<void*>(client_tool_data)};
// return pointer to configure data
return &cfg;
}
rocprofiler_callback_tracing_kind_t kind
Definition fwd.h:581
rocprofiler_thread_id_t thread_id
Definition fwd.h:579
rocprofiler_correlation_id_t correlation_id
Definition fwd.h:580
uint64_t rocprofiler_thread_id_t
Thread ID. Value will be equivalent to syscall(__NR_gettid)
Definition fwd.h:439
rocprofiler_runtime_library_t
Enumeration for specifying runtime libraries supported by rocprofiler. This enumeration is used for t...
Definition fwd.h:348
@ ROCPROFILER_PAGE_MIGRATION_QUEUE_SUSPEND
Definition fwd.h:231
@ ROCPROFILER_PAGE_MIGRATION_UNMAP_FROM_GPU
Definition fwd.h:232
@ ROCPROFILER_PAGE_MIGRATION_PAGE_MIGRATE
Definition fwd.h:229
@ ROCPROFILER_STATUS_ERROR_BUFFER_BUSY
buffer operation failed because it currently busy handling another request (e.g. flushing)
Definition fwd.h:71
@ ROCPROFILER_CALLBACK_PHASE_UNLOAD
Callback invoked prior to code object unloading.
Definition fwd.h:146
@ ROCPROFILER_BUFFER_TRACING_HSA_CORE_API
Definition fwd.h:181
@ ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY
Buffer scratch memory reclaimation info.
Definition fwd.h:194
@ ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API
Definition fwd.h:186
@ ROCPROFILER_BUFFER_TRACING_HSA_IMAGE_EXT_API
Definition fwd.h:183
@ ROCPROFILER_BUFFER_TRACING_HSA_AMD_EXT_API
Definition fwd.h:182
@ ROCPROFILER_BUFFER_TRACING_HSA_FINALIZE_EXT_API
Definition fwd.h:184
@ ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT
Definition fwd.h:168
Generic record with type identifier(s) and a pointer to data. This data type is used with buffered da...
Definition fwd.h:617
rocprofiler_status_t rocprofiler_create_buffer(rocprofiler_context_id_t context, unsigned long size, unsigned long watermark, rocprofiler_buffer_policy_t policy, rocprofiler_buffer_tracing_cb_t callback, void *callback_data, rocprofiler_buffer_id_t *buffer_id)
Create buffer.
rocprofiler_status_t rocprofiler_flush_buffer(rocprofiler_buffer_id_t buffer_id)
Flush buffer.
rocprofiler_status_t rocprofiler_configure_buffer_tracing_service(rocprofiler_context_id_t context_id, rocprofiler_buffer_tracing_kind_t kind, rocprofiler_tracing_operation_t *operations, unsigned long operations_count, rocprofiler_buffer_id_t buffer_id)
Configure Buffer Tracing Service.
rocprofiler_status_t rocprofiler_query_buffer_tracing_kind_name(rocprofiler_buffer_tracing_kind_t kind, const char **name, uint64_t *name_len)
Query the name of the buffer tracing kind. The name retrieved from this function is a string literal ...
ROCProfiler Buffer HIP API Tracer Record.
Definition buffer_tracing.h:111
ROCProfiler Buffer HSA API Tracer Record.
Definition buffer_tracing.h:87
ROCProfiler Buffer Kernel Dispatch Tracer Record.
Definition buffer_tracing.h:177
ROCProfiler Buffer Memory Copy Tracer Record.
Definition buffer_tracing.h:154
ROCProfiler Buffer Page Migration Tracer Record.
Definition buffer_tracing.h:231
ROCProfiler Buffer Scratch Memory Tracer Record.
Definition buffer_tracing.h:256
rocprofiler_status_t rocprofiler_configure_callback_tracing_service(rocprofiler_context_id_t context_id, rocprofiler_callback_tracing_kind_t kind, rocprofiler_tracing_operation_t *operations, unsigned long operations_count, rocprofiler_callback_tracing_cb_t callback, void *callback_args)
Configure Callback Tracing Service. The callback tracing service provides two synchronous callbacks a...
ROCProfiler Code Object Kernel Symbol Tracer Callback Record.
Definition callback_tracing.h:138
rocprofiler_status_t rocprofiler_start_context(rocprofiler_context_id_t context_id)
Start context.
rocprofiler_status_t rocprofiler_create_context(rocprofiler_context_id_t *context_id)
Create context.
rocprofiler_status_t rocprofiler_context_is_valid(rocprofiler_context_id_t context_id, int *status)
Query whether the context is valid.
rocprofiler_status_t rocprofiler_stop_context(rocprofiler_context_id_t context_id)
Stop context.
rocprofiler_status_t rocprofiler_push_external_correlation_id(rocprofiler_context_id_t context, rocprofiler_thread_id_t tid, rocprofiler_user_data_t external_correlation_id)
Push default value for external field in rocprofiler_correlation_id_t onto stack.
rocprofiler_status_t rocprofiler_assign_callback_thread(rocprofiler_buffer_id_t buffer_id, rocprofiler_callback_thread_t cb_thread_id)
By default, all buffered results are delivered on the same thread. Using rocprofiler_create_callback_...
rocprofiler_status_t rocprofiler_at_internal_thread_create(rocprofiler_internal_thread_library_cb_t precreate, rocprofiler_internal_thread_library_cb_t postcreate, int libs, void *data)
Invoke this function to receive callbacks before and after the creation of an internal thread by a li...
rocprofiler_status_t rocprofiler_create_callback_thread(rocprofiler_callback_thread_t *cb_thread_id)
Create a handle to a unique thread (created by rocprofiler) which, when associated with a particular ...
opaque handle to an internal thread identifier which delivers callbacks for buffers
Definition internal_threading.h:88
rocprofiler_status_t rocprofiler_get_thread_id(rocprofiler_thread_id_t *tid)
Get the identifier value of the current thread that is used by rocprofiler.
void(* rocprofiler_client_finalize_t)(rocprofiler_client_id_t)
Prototype for the function pointer provided to tool in rocprofiler_tool_initialize_t....
Definition registration.h:61
rocprofiler_tool_configure_result_t * rocprofiler_configure(uint32_t version, const char *runtime_version, uint32_t priority, rocprofiler_client_id_t *client_id)
This is the special function that tools define to enable rocprofiler support. The tool should return ...
rocprofiler_status_t rocprofiler_is_initialized(int *status)
Query whether rocprofiler has already scanned the binary for all the instances of rocprofiler_configu...
rocprofiler_status_t rocprofiler_force_configure(rocprofiler_configure_func_t configure_func)
Function for explicitly registering a configuration with rocprofiler. This can be invoked before any ...
A client refers to an individual or entity engaged in the configuration of ROCprofiler services....
Definition registration.h:48
Data structure containing a initialization, finalization, and data.
Definition registration.h:96
Generated by
1.9.8