AI/rocm-systems
2
0
Fork 0
Code Issues Pull-Requests Actions Pakete Projekte Releases Wiki Aktivität

SWDEV-516846: Fix serialization services conflicts and ATT counter streaming (#230)

Quellcode durchsuchen 
* Update TT API

* Rework serialization

* update att_core

* Fix tests

* Fix tool

* Formatting

* Fix perfcounter

* Formatting

* Rename agent TT

* Format

* Workaround for codeQL alert

* Tidy fix

* Fix compiler error

* Tidy

* Fix some tests

* Fixing some tests

* formatting

* Fixing ATT serialization

* Format

* Fix test commandline

* Fixing init order

* Format

* Tidy fixes

* Removing unused sample

* Fix tests and schema

* Added ATT + PMC test

* Fix mode

* Fix file mode

* Review comments

* Fix typo

* Review comments

* Review comments

* Fix missing id inc after review comment

* Review comments

* Suggested Fixes

* Testing changes

* Test fix

* Build fixes

* Minor build fix

---------

Co-authored-by: Giovanni Baraldi <gbaraldi@amd.com>
Co-authored-by: Benjamin Welton <bewelton@amd.com>
Co-authored-by: Welton, Benjamin <Benjamin.Welton@amd.com>
Dieser Commit ist enthalten in:
Baraldi, Giovanni
2025-03-15 02:11:10 +01:00
committet von GitHub
Ursprung 914923f688
Commit 821918a512
48 geänderte Dateien mit 763 neuen und 1525 gelöschten Zeilen
Patch-Datei herunterladen Vergleichs-Datei herunterladen Alle Dateien ausklappen Alle Dateien einklappen
samples/advanced_thread_trace/CMakeLists.txt
-59
Datei anzeigen
@@ -1,59 +0,0 @@
#
#
#
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-advanced-thread-trace 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)
find_package(
amd_comgr
REQUIRED
CONFIG
HINTS
${rocm_version_DIR}
${ROCM_PATH}
PATHS
${rocm_version_DIR}
${ROCM_PATH}
PATH_SUFFIXES
lib/cmake/amd_comgr)
set_source_files_properties(main.cpp PROPERTIES LANGUAGE HIP)
set_source_files_properties(main.cpp PROPERTIES COMPILE_FLAGS "-g")
add_executable(advanced-thread-trace)
target_sources(advanced-thread-trace PRIVATE main.cpp client.cpp)
target_link_libraries(
advanced-thread-trace
PRIVATE rocprofiler-sdk::rocprofiler-sdk amd_comgr
rocprofiler-sdk::samples-common-library rocprofiler-sdk::samples-build-flags)
rocprofiler_samples_get_preload_env(PRELOAD_ENV)
add_test(NAME advanced-thread-trace COMMAND $<TARGET_FILE:advanced-thread-trace>)
set_tests_properties(
advanced-thread-trace
PROPERTIES TIMEOUT 45 LABELS "samples" ENVIRONMENT "${PRELOAD_ENV}"
FAIL_REGULAR_EXPRESSION "threw an exception")
samples/advanced_thread_trace/client.cpp
-541
Datei anzeigen
@@ -1,541 +0,0 @@
// MIT License
//
// Copyright (c) 2023-2025 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/advanced_thread_trace/client.cpp
*
* @brief Example rocprofiler client (tool)
*/
#include <rocprofiler-sdk/amd_detail/thread_trace.h>
#include <rocprofiler-sdk/buffer.h>
#include <rocprofiler-sdk/callback_tracing.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/registration.h>
#include <rocprofiler-sdk/rocprofiler.h>
#include <rocprofiler-sdk/cxx/codeobj/code_printing.hpp>
#include <shared_mutex>
#include "common/defines.hpp"
#include "common/filesystem.hpp"
#include <cxxabi.h>
#include <atomic>
#include <cassert>
#include <chrono>
#include <cstddef>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <fstream>
#include <functional>
#include <iomanip>
#include <iostream>
#include <map>
#include <mutex>
#include <regex>
#include <string>
#include <string_view>
#include <thread>
#include <unordered_map>
#include <vector>
#define OUTPUT_OFSTREAM "advanced_thread_trace.log"
#define TARGET_CU 1
#define SIMD_SELECT 0x3
#define BUFFER_SIZE 0x6000000
#define SE_MASK 0x11
constexpr bool COPY_MEMORY_CODEOBJ = true;
#define C_API_BEGIN \
try \
{
#define C_API_END \
} \
catch(std::exception & e) \
{ \
std::cerr << "Error in " << __FILE__ << ':' << __LINE__ << ' ' << e.what() << std::endl; \
} \
catch(...) { std::cerr << "Error in " << __FILE__ << ':' << __LINE__ << std::endl; }
struct pcinfo_t
{
uint64_t marker_id;
uint64_t addr;
};
bool
operator==(const pcinfo_t& a, const pcinfo_t& b)
{
return a.addr == b.addr && a.marker_id == b.marker_id;
};
bool
operator<(const pcinfo_t& a, const pcinfo_t& b)
{
if(a.marker_id == b.marker_id) return a.addr < b.addr;
return a.marker_id < b.marker_id;
};
namespace client
{
using code_obj_load_data_t = rocprofiler_callback_tracing_code_object_load_data_t;
using kernel_symbol_data_t = rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t;
using Instruction = rocprofiler::sdk::codeobj::disassembly::Instruction;
using CodeobjAddressTranslate = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate;
using SymbolInfo = rocprofiler::sdk::codeobj::disassembly::SymbolInfo;
rocprofiler_client_id_t* client_id = nullptr;
rocprofiler_context_id_t client_ctx = {0};
struct isa_map_elem_t
{
std::atomic<size_t> hitcount{0};
std::atomic<size_t> latency{0};
std::unique_ptr<Instruction> code_line{nullptr};
};
struct ToolData
{
ToolData()
{
try
{
output_file.open(OUTPUT_OFSTREAM);
} catch(...)
{}
if(output_file.is_open())
std::cout << "Writing code-object-isa-decode log to: " << OUTPUT_OFSTREAM << std::endl;
else
std::cout << "Could not open log file: " << OUTPUT_OFSTREAM << ", writing to stdout\n";
};
std::shared_mutex isa_map_mut;
std::mutex output_mut;
CodeobjAddressTranslate codeobjTranslate;
std::map<pcinfo_t, std::unique_ptr<isa_map_elem_t>> isa_map;
std::unordered_map<uint64_t, SymbolInfo> kernels_in_codeobj = {};
std::unordered_map<uint64_t, std::string> kernel_id_to_kernel_name = {};
int num_waves = 0;
std::ostream& output()
{
if(output_file.is_open())
return output_file;
else
return std::cout;
}
std::stringstream printKernel(uint64_t vaddr)
{
std::stringstream ss;
try
{
ss << '\n' << std::hex;
SymbolInfo& info = kernels_in_codeobj.at(vaddr);
ss << std::hex << "Found: " << info.name << " at addr: 0x" << vaddr << " with offset 0x"
<< info.faddr << " vaddr 0x" << info.vaddr << std::dec << '\n';
} catch(std::exception& e)
{
ss << e.what() << '\n';
}
return ss;
}
private:
std::ofstream output_file;
};
struct source_location
{
std::string function = {};
std::string file = {};
uint32_t line = 0;
std::string context = {};
};
struct trace_data_t
{
int64_t id;
uint8_t* data;
uint64_t size;
};
auto* tool = new ToolData{};
void
tool_codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record,
rocprofiler_user_data_t* /* user_data */,
void* /* callback_data */)
{
C_API_BEGIN
if(record.kind != ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT) return;
if(record.phase != ROCPROFILER_CALLBACK_PHASE_LOAD) return;
if(record.operation == ROCPROFILER_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER)
{
std::unique_lock<std::shared_mutex> lg(tool->isa_map_mut);
auto* data = static_cast<kernel_symbol_data_t*>(record.payload);
tool->kernel_id_to_kernel_name.emplace(data->kernel_id, data->kernel_name);
}
if(record.operation != ROCPROFILER_CODE_OBJECT_LOAD) return;
std::unique_lock<std::shared_mutex> lg(tool->isa_map_mut);
auto* data = static_cast<code_obj_load_data_t*>(record.payload);
if(std::string_view(data->uri).find("file:///") == 0)
{
tool->codeobjTranslate.addDecoder(
data->uri, data->code_object_id, data->load_delta, data->load_size);
auto symbolmap = tool->codeobjTranslate.getSymbolMap(data->code_object_id);
for(auto& [vaddr, symbol] : symbolmap)
tool->kernels_in_codeobj[vaddr] = symbol;
}
else if(COPY_MEMORY_CODEOBJ)
{
tool->codeobjTranslate.addDecoder(reinterpret_cast<const void*>(data->memory_base),
data->memory_size,
data->code_object_id,
data->load_delta,
data->load_size);
auto symbolmap = tool->codeobjTranslate.getSymbolMap(data->code_object_id);
for(auto& [vaddr, symbol] : symbolmap)
tool->kernels_in_codeobj[vaddr] = symbol;
}
C_API_END
}
rocprofiler_att_control_flags_t
dispatch_callback(rocprofiler_queue_id_t /* queue_id */,
const rocprofiler_agent_t* /* agent */,
rocprofiler_correlation_id_t /* correlation_id */,
rocprofiler_kernel_id_t kernel_id,
rocprofiler_dispatch_id_t /* dispatch_id */,
rocprofiler_user_data_t* /* userdata */,
void* /* userdata */)
{
C_API_BEGIN
std::shared_lock<std::shared_mutex> lg(tool->isa_map_mut);
static std::atomic<int> call_id{0};
static std::string_view desired_func_name = "transposeLds";
try
{
auto& kernel_name = tool->kernel_id_to_kernel_name.at(kernel_id);
if(kernel_name.find(desired_func_name) == std::string::npos)
return ROCPROFILER_ATT_CONTROL_NONE;
int id = call_id.fetch_add(1);
if(id == 1) return ROCPROFILER_ATT_CONTROL_START_AND_STOP;
} catch(...)
{
std::cerr << "Could not find kernel id: " << kernel_id << std::endl;
}
C_API_END
return ROCPROFILER_ATT_CONTROL_NONE;
}
void
get_trace_data(rocprofiler_att_parser_data_type_t type, void* att_data, void* userdata)
{
C_API_BEGIN
assert(userdata && "ISA callback passed null!");
std::shared_lock<std::shared_mutex> shared_lock(tool->isa_map_mut);
if(type == ROCPROFILER_ATT_PARSER_DATA_TYPE_OCCUPANCY) tool->num_waves++;
if(type != ROCPROFILER_ATT_PARSER_DATA_TYPE_ISA) return;
auto& event = *reinterpret_cast<rocprofiler_att_data_type_isa_t*>(att_data);
pcinfo_t pc{event.marker_id, event.offset};
auto it = tool->isa_map.find(pc);
if(it == tool->isa_map.end())
{
shared_lock.unlock();
{
std::unique_lock<std::shared_mutex> unique_lock(tool->isa_map_mut);
auto ptr = std::make_unique<isa_map_elem_t>();
try
{
ptr->code_line = tool->codeobjTranslate.get(pc.marker_id, pc.addr);
} catch(std::exception& e)
{
std::cerr << pc.marker_id << ":" << pc.addr << ' ' << e.what() << std::endl;
return;
} catch(...)
{
std::cerr << "Could not fetch: " << pc.marker_id << ':' << pc.addr << std::endl;
return;
}
it = tool->isa_map.emplace(pc, std::move(ptr)).first;
}
shared_lock.lock();
}
it->second->hitcount.fetch_add(event.hitcount, std::memory_order_relaxed);
it->second->latency.fetch_add(event.latency, std::memory_order_relaxed);
C_API_END
}
uint64_t
copy_trace_data(int* seid, uint8_t** buffer, uint64_t* buffer_size, void* userdata)
{
trace_data_t& data = *reinterpret_cast<trace_data_t*>(userdata);
*seid = data.id;
*buffer_size = data.size;
*buffer = data.data;
data.size = 0;
return *buffer_size;
}
rocprofiler_status_t
isa_callback(char* isa_instruction,
uint64_t* isa_memory_size,
uint64_t* isa_size,
uint64_t marker_id,
uint64_t offset,
void* userdata)
{
C_API_BEGIN
assert(userdata && "ISA callback passed null!");
std::unique_ptr<Instruction> instruction;
{
std::unique_lock<std::shared_mutex> unique_lock(tool->isa_map_mut);
instruction = tool->codeobjTranslate.get(marker_id, offset);
}
if(!instruction.get()) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT;
{
size_t tmp_isa_size = *isa_size;
*isa_size = instruction->inst.size();
if(*isa_size > tmp_isa_size) return ROCPROFILER_STATUS_ERROR_OUT_OF_RESOURCES;
}
memcpy(isa_instruction, instruction->inst.data(), *isa_size);
*isa_memory_size = instruction->size;
auto ptr = std::make_unique<isa_map_elem_t>();
ptr->code_line = std::move(instruction);
tool->isa_map.emplace(pcinfo_t{marker_id, offset}, std::move(ptr));
return ROCPROFILER_STATUS_SUCCESS;
C_API_END
return ROCPROFILER_STATUS_ERROR;
}
void
shader_data_callback(int64_t se_id,
void* se_data,
size_t data_size,
rocprofiler_user_data_t /* userdata */)
{
C_API_BEGIN
{
std::unique_lock<std::mutex> lk(tool->output_mut);
tool->output() << "SE ID: " << se_id << " with size " << data_size << std::hex << '\n';
}
trace_data_t data{.id = se_id, .data = (uint8_t*) se_data, .size = data_size};
auto status = rocprofiler_att_parse_data(copy_trace_data, get_trace_data, isa_callback, &data);
if(status != ROCPROFILER_STATUS_SUCCESS)
std::cerr << "shader_data_callback failed with status " << status << std::endl;
C_API_END
}
int
tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
{
(void) fini_func;
ROCPROFILER_CALL(rocprofiler_create_context(&client_ctx), "context creation");
ROCPROFILER_CALL(
rocprofiler_configure_callback_tracing_service(client_ctx,
ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT,
nullptr,
0,
tool_codeobj_tracing_callback,
tool_data),
"code object tracing service configure");
std::vector<rocprofiler_att_parameter_t> parameters = {
{ROCPROFILER_ATT_PARAMETER_TARGET_CU, {TARGET_CU}},
{ROCPROFILER_ATT_PARAMETER_SIMD_SELECT, {SIMD_SELECT}},
{ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE, {BUFFER_SIZE}},
{ROCPROFILER_ATT_PARAMETER_SHADER_ENGINE_MASK, {SE_MASK}}};
ROCPROFILER_CALL(rocprofiler_configure_dispatch_thread_trace_service(client_ctx,
parameters.data(),
parameters.size(),
dispatch_callback,
shader_data_callback,
tool_data),
"thread trace service configure");
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;
}
ROCPROFILER_CALL(rocprofiler_start_context(client_ctx), "context start");
// no errors
return 0;
}
void
tool_fini(void* /* data */)
{
std::unique_lock<std::shared_mutex> isa_lk(client::tool->isa_map_mut);
std::unique_lock<std::mutex> out_lk(client::tool->output_mut);
// Find largest instruction
size_t max_inst_size = 0;
for(auto& [addr, lines] : client::tool->isa_map)
if(lines.get()) max_inst_size = std::max(max_inst_size, lines->code_line->inst.size());
std::string empty_space;
empty_space.resize(max_inst_size, ' ');
size_t vmc_latency = 0;
size_t lgk_latency = 0;
size_t scalar_latency = 0;
size_t vector_latency = 0;
size_t other_latency = 0;
size_t scalar_exec = 0;
size_t vector_exec = 0;
size_t other_exec = 0;
for(auto& [addr, line] : client::tool->isa_map)
if(line.get())
{
size_t hitcount = line->hitcount.load(std::memory_order_relaxed);
size_t latency = line->latency.load(std::memory_order_relaxed);
auto& code_line = line->code_line->inst;
client::tool->output() << std::hex << "0x" << addr.addr << std::dec << ' ' << code_line
<< empty_space.substr(0, max_inst_size - code_line.size())
<< " Hit: " << hitcount << " - Latency: " << latency << '\n';
if(code_line.find("s_waitcnt") == 0)
{
other_exec += hitcount;
if(code_line.find("lgkmcnt") != std::string::npos)
lgk_latency += latency;
else
vmc_latency += latency;
}
else if(code_line.find("v_") == 0)
{
vector_exec += hitcount;
vector_latency += latency;
}
else if(code_line.find("s_") == 0)
{
scalar_exec += hitcount;
scalar_latency += latency;
}
else
{
other_exec += hitcount;
other_latency += latency;
}
}
size_t total_exec = vector_exec + scalar_exec + other_exec;
size_t memory_latency = vmc_latency + lgk_latency;
size_t total_latency = memory_latency + vector_latency + scalar_latency + other_latency;
float vmc_fraction = 100 * vmc_latency / float(total_latency);
float lgk_fraction = 100 * lgk_latency / float(total_latency);
client::tool->output() << "Total executed instructions: " << total_exec << '\n'
<< "Total executed vector instructions: " << vector_exec
<< " with average " << vector_latency / float(vector_exec)
<< " cycles.\n"
<< "Total executed scalar instructions: " << scalar_exec
<< " with average " << scalar_latency / float(scalar_exec)
<< " cycles.\n"
<< "Vector memory ops occupied: " << vmc_fraction << "% of cycles.\n"
<< "Scalar and LDS memory ops occupied: " << lgk_fraction
<< "% of cycles.\n"
<< "Num waves created: " << (client::tool->num_waves / 2) << std::endl;
}
} // 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)
{
// set the client name
id->name = "Adv_Thread_Trace_Sample";
// 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 << " (priority=" << priority << ") is using rocprofiler-sdk v" << major << "."
<< minor << "." << patch << " (" << runtime_version << ")";
std::clog << info.str() << std::endl;
// create configure data
static auto cfg =
rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t),
&client::tool_init,
&client::tool_fini,
nullptr};
// return pointer to configure data
return &cfg;
}
samples/advanced_thread_trace/main.cpp
-245
Datei anzeigen
@@ -1,245 +0,0 @@
// MIT License
//
// Copyright (c) 2024-2025 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 <cstdio>
#include <cstdlib>
#include <iomanip>
#include <iostream>
#include <mutex>
#include <random>
#include <stdexcept>
#include "transpose_kernels.hpp"
#define PRINT_ALIGN 36
namespace
{
using lock_guard_t = std::lock_guard<std::mutex>;
auto print_lock = std::mutex{};
} // namespace
enum TransposeType
{
TRANSPOSE_NAIVE,
TRANSPOSE_INPLACE_LDS,
TRANSPOSE_NO_BANK_CONFLICTS
};
class ITranspose
{
public:
virtual void run(TransposeType ttype, int numThreadsY, int num_iter) = 0;
virtual ~ITranspose(){};
};
template <typename T>
class Transpose : public ITranspose
{
public:
Transpose(int dev, size_t _M)
: devID(dev)
, M(_M)
, databytes(_M * _M * sizeof(T))
{
HIP_API_CALL(hipSetDevice(devID));
HIP_API_CALL(hipStreamCreate(&stream));
std::default_random_engine _engine{std::random_device{}() * rand()};
std::uniform_int_distribution<int> _dist{0, 1000};
inp_matrix = new T[M * M];
out_matrix = new T[M * M];
for(size_t i = 0; i < M * M; i++)
inp_matrix[i] = static_cast<T>(_dist(_engine));
memset(out_matrix, 0, databytes);
HIP_API_CALL(hipMalloc(&in, databytes));
HIP_API_CALL(hipMalloc(&out, databytes));
HIP_API_CALL(hipMemsetAsync(in, 0, databytes, stream));
HIP_API_CALL(hipMemsetAsync(out, 0, databytes, stream));
HIP_API_CALL(hipMemcpyAsync(in, inp_matrix, databytes, hipMemcpyDefault, stream));
HIP_API_CALL(hipEventCreate(&start));
HIP_API_CALL(hipEventCreate(&stop));
}
void run(TransposeType ttype, int numThreadsY, int num_iter) override
{
HIP_API_CALL(hipSetDevice(devID));
dim3 grid(M / TILE_DIM, M / TILE_DIM, 1);
dim3 block(TILE_DIM, numThreadsY, 1);
auto Kernel = transposeNaive<T>;
std::string KernelName = "transposeNaive";
if(ttype == TransposeType::TRANSPOSE_NO_BANK_CONFLICTS)
{
Kernel = transposeLdsNoBankConflicts<T>;
KernelName = "transposeLdsNoBankConflicts";
}
else if(ttype == TransposeType::TRANSPOSE_INPLACE_LDS)
{
Kernel = transposeLdsSwapInplace<T>;
KernelName = "transposeLdsSwapInplace";
}
{
std::string functypeid = __PRETTY_FUNCTION__;
auto it_beg = functypeid.rfind("[T = ");
auto it_end = functypeid.rfind(']');
if(it_beg != std::string::npos) it_beg += std::string("[T = ").size();
if(it_beg < it_end && it_end != std::string::npos)
KernelName += '<' + functypeid.substr(it_beg, it_end - it_beg) + '>';
}
HIP_API_CALL(hipStreamSynchronize(stream));
HIP_API_CALL(hipEventRecord(start, stream));
for(int i = 0; i < num_iter; i++)
{
Kernel<<<grid, block, 0, stream>>>(out, in, M);
HIP_API_CALL(hipGetLastError());
}
HIP_API_CALL(hipEventRecord(stop, stream));
HIP_API_CALL(hipMemcpyAsync(out_matrix, out, databytes, hipMemcpyDefault, stream));
HIP_API_CALL(hipEventSynchronize(stop));
float time;
HIP_API_CALL(hipEventElapsedTime(&time, start, stop));
float GB = databytes * num_iter * 2 / float(1 << 30);
{
lock_guard_t _lk{print_lock};
std::cout << "The average performance of " << std::setw(38) << KernelName << " : "
<< (1000 * GB / time) << " GB/s" << std::endl;
}
verify();
}
void verify() const
{
HIP_API_CALL(hipStreamSynchronize(stream));
for(int i = 0; i < 10; i++)
{
int row = rand() % M;
int col = rand() % M;
if(inp_matrix[row * M + col] != out_matrix[col * M + row])
{
lock_guard_t _lk{print_lock};
std::cout << "mismatch: " << row << ", " << col << " : "
<< inp_matrix[row * M + col] << " | " << out_matrix[col * M + row]
<< std::endl;
}
}
}
virtual ~Transpose()
{
HIP_API_CALL(hipSetDevice(devID));
HIP_API_CALL(hipEventDestroy(start));
HIP_API_CALL(hipEventDestroy(stop));
HIP_API_CALL(hipFree(in));
HIP_API_CALL(hipFree(out));
HIP_API_CALL(hipStreamDestroy(stream));
delete[] inp_matrix;
delete[] out_matrix;
}
const int devID;
const size_t M;
const size_t databytes;
hipStream_t stream;
hipEvent_t start, stop;
T* inp_matrix = nullptr;
T* out_matrix = nullptr;
T* in = nullptr;
T* out = nullptr;
};
int
main(int argc, char** argv)
{
int deviceId = 0;
int blockDimY = 8;
int num_iter = 1;
int mat_size = 8192;
for(int i = 1; i < argc; ++i)
{
auto _arg = std::string{argv[i]};
if(_arg == "?" || _arg == "-h" || _arg == "--help")
{
std::cout << "usage: transpose "
<< "[MatrixSize (" << mat_size << ")] "
<< "[numIter (" << num_iter << ")] "
<< "[blockDimY (" << blockDimY << ")] "
<< "[DEVICE_ID (" << deviceId << ")] " << std::endl;
exit(EXIT_SUCCESS);
}
}
if(argc > 1) mat_size = atoll(argv[1]);
if(argc > 2) num_iter = atoll(argv[2]);
if(argc > 3) blockDimY = atoll(argv[3]);
if(argc > 4) deviceId = atoll(argv[4]);
printf("[transpose] Matrix size: %d, device ID: %d, num iter: %d, blockDimY: %d\n",
mat_size,
deviceId,
num_iter,
blockDimY);
int ndevice = 0;
HIP_API_CALL(hipGetDeviceCount(&ndevice));
printf("[transpose] Number of devices found: %i\n", ndevice);
assert(ndevice > 0);
if(deviceId >= ndevice) exit(EXIT_FAILURE);
{
std::vector<std::unique_ptr<ITranspose>> kernels;
kernels.push_back(std::make_unique<Transpose<int>>(deviceId, mat_size));
kernels.push_back(std::make_unique<Transpose<float>>(deviceId, mat_size));
kernels.push_back(std::make_unique<Transpose<double>>(deviceId, mat_size));
for(auto& kernel : kernels)
{
kernel->run(TransposeType::TRANSPOSE_NAIVE, blockDimY, num_iter);
kernel->run(TransposeType::TRANSPOSE_INPLACE_LDS, blockDimY, num_iter);
kernel->run(TransposeType::TRANSPOSE_NO_BANK_CONFLICTS, blockDimY, num_iter);
}
}
HIP_API_CALL(hipDeviceSynchronize());
return 0;
}
samples/advanced_thread_trace/transpose_kernels.hpp
-106
Datei anzeigen
@@ -1,106 +0,0 @@
// MIT License
//
// Copyright (c) 2024-2025 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.
#pragma once
#include "hip/hip_runtime.h"
#define HIP_API_CALL(CALL) \
{ \
hipError_t error_ = (CALL); \
if(error_ != hipSuccess) \
{ \
lock_guard_t _hip_api_print_lk{print_lock}; \
fprintf(stderr, \
"%s:%d :: HIP error : %s\n", \
__FILE__, \
__LINE__, \
hipGetErrorString(error_)); \
exit(EXIT_FAILURE); \
} \
}
#define TILE_DIM 64
template <typename T>
__global__ void
transposeNaive(T* odata, const T* idata, size_t size)
{
size_t idx = blockIdx.x * TILE_DIM + threadIdx.x;
size_t block_posy = blockIdx.y * TILE_DIM;
for(size_t idy = threadIdx.y; idy < TILE_DIM; idy += blockDim.y)
odata[size * idx + block_posy + idy] = idata[idx + (block_posy + idy) * size];
}
template <typename T>
__global__ void
transposeLdsNoBankConflicts(T* odata, const T* idata, size_t size)
{
__shared__ T tile[TILE_DIM][TILE_DIM + 1];
size_t idx_in = blockIdx.x * TILE_DIM + threadIdx.x;
size_t idy_in = blockIdx.y * TILE_DIM + threadIdx.y;
size_t index_in = idx_in + idy_in * size;
size_t idx_out = blockIdx.y * TILE_DIM + threadIdx.x;
size_t idy_out = blockIdx.x * TILE_DIM + threadIdx.y;
size_t index_out = idx_out + idy_out * size;
for(size_t y = 0; y < TILE_DIM; y += blockDim.y)
tile[threadIdx.y + y][threadIdx.x] = idata[index_in + y * size];
__syncthreads();
for(size_t y = 0; y < TILE_DIM; y += blockDim.y)
odata[index_out + y * size] = tile[threadIdx.x][threadIdx.y + y];
}
// Generates more interesting ISA
template <typename T>
__global__ void
transposeLdsSwapInplace(T* odata, const T* idata, size_t size)
{
__shared__ T tile[TILE_DIM][TILE_DIM];
const size_t idx_in = blockIdx.x * TILE_DIM + threadIdx.x;
for(size_t idy = threadIdx.y; idy < TILE_DIM; idy += blockDim.y)
tile[idy][threadIdx.x] = idata[idx_in + (idy + blockIdx.y * TILE_DIM) * size];
__syncthreads();
for(size_t idy = threadIdx.y; idy < TILE_DIM; idy += blockDim.y)
if(idy < threadIdx.x)
{
T temp = tile[idy][threadIdx.x];
tile[idy][threadIdx.x] = tile[threadIdx.x][idy];
tile[threadIdx.x][idy] = temp;
}
__syncthreads();
const size_t idx_out = blockIdx.y * TILE_DIM + threadIdx.x;
for(size_t idy = threadIdx.y; idy < TILE_DIM; idy += blockDim.y)
odata[(blockIdx.x * TILE_DIM + idy) * size + idx_out] = tile[idy][threadIdx.x];
}
source/bin/rocprofv3.py
+36 -12
Datei anzeigen
@@ -90,7 +90,7 @@ def strtobool(val):
def search_path(path_list):
supported_option = []
lib_att_pattern = r"libatt_decoder_(trace|summary|debug|testing)\.so"
lib_att_pattern = r"libatt_decoder_(trace|debug|testing1|testing2)\.so"
file_list = []
for path in path_list:
@@ -100,7 +100,7 @@ def search_path(path_list):
for itr in file_list:
_match = re.match(lib_att_pattern, itr)
if _match:
lst = re.findall("trace|debug|summary|testing", itr)
lst = re.findall("trace|debug|testing1|testing2", itr)
supported_option.extend(lst)
return set(supported_option)
@@ -682,6 +682,20 @@ For MPI applications (or other job launchers such as SLURM), place rocprofv3 ins
choices=set(choice_list),
)
att_options.add_argument(
"--att-perfcounters",
help="Set performance counters, and optionally their mask",
default=None,
type=str.upper,
)
att_options.add_argument(
"--att-perfcounter-ctrl",
help="Integer in [0,32] range specifying collection period.",
default=None,
type=int,
)
add_parser_bool_argument(
att_options,
"--att-serialize-all",
@@ -1305,15 +1319,13 @@ def run(app_args, args, **kwargs):
f"{type(num_str)} is not supported. {num_str} should be of type integer or string."
)
if args.pmc or (
if (
args.pc_sampling_beta_enabled
or args.pc_sampling_unit
or args.pc_sampling_method
or args.pc_sampling_interval
):
fatal_error(
"Advanced thread trace cannot be enabled with counter collection or pc sampling"
)
fatal_error("Advanced thread trace cannot be enabled with pc sampling")
if not args.att_parse:
fatal_error("provide the parser choice")
@@ -1358,12 +1370,24 @@ def run(app_args, args, **kwargs):
":".join(args.att_library_path),
overwrite=True,
)
if args.att_percounters:
update_env(
"ROCPROF_ATT_PARAM_PERFCOUNTERS",
" ".join(args.att_perfcounters),
overwrite=True,
)
if args.att_perfcounters:
if args.pmc:
fatal_error("ATT perfcounters cannot be enabled with PMC")
else:
update_env(
"ROCPROF_ATT_PARAM_PERFCOUNTERS",
args.att_perfcounters,
overwrite=True,
)
if args.att_perfcounter_ctrl:
if args.pmc:
fatal_error("ATT perfcounters cannot be enabled with PMC")
else:
update_env(
"ROCPROF_ATT_PARAM_PERFCOUNTER_CTRL",
args.att_perfcounter_ctrl,
overwrite=True,
)
if args.log_level in ("info", "trace", "env"):
log_config(app_env)
source/docs/rocprofv3-schema.json
+17 -12
Datei anzeigen
@@ -569,22 +569,27 @@
"description": "Comments matching assembly instructions from pc_sample_instructions array. If debug symbols are available, comments provide instructions to source-line mapping. Otherwise, a comment is an empty string."
},
"att_filenames" : {
"type": "object",
"properties": {
"key": {
"type": "integer",
"description": "Dispatch id."
},
"value": {
"type": "array",
"description": "An array of ATT filenames."
}
"type": "array",
"items": {
"type": "string",
"description": "Filename of saved .att data in the form name_shader_engine_seId_dispatchId.att"
}
},
"code_object_snapshot_filenames": {
"type": "array",
"description": "Code object snapshot filename"
"items": {
"type": "object",
"properties": {
"key": {
"type": "integer",
"description": "Code object load id."
},
"value": {
"type": "string",
"description": "Saved code object filename"
}
}
}
}
}
},
source/include/rocprofiler-sdk/amd_detail/thread_trace_agent.h
+2 -2
Datei anzeigen
@@ -57,11 +57,11 @@ ROCPROFILER_EXTERN_C_INIT
* @retval ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT for invalid rocprofiler_att_parameter_t
*/
rocprofiler_status_t
rocprofiler_configure_agent_thread_trace_service(
rocprofiler_configure_device_thread_trace_service(
rocprofiler_context_id_t context_id,
rocprofiler_agent_id_t agent_id,
rocprofiler_att_parameter_t* parameters,
size_t num_parameters,
rocprofiler_agent_id_t agent_id,
rocprofiler_att_shader_data_callback_t shader_callback,
rocprofiler_user_data_t callback_userdata) ROCPROFILER_API;
source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h
+4 -2
Datei anzeigen
@@ -66,8 +66,9 @@ typedef rocprofiler_att_control_flags_t (*rocprofiler_att_dispatch_callback_t)(
/**
* @brief Enables the advanced thread trace service for dispatch-based tracing.
* The tool has an option to enable/disable thread trace on every dispatch callback.
* This service enables kernel serialization.
* @param [in] context_id context_id.
* This service serializes all traced kernels, and optionally all non-traced kernels.
* @param [in] context_id id of the context used for start/stop thread_trace.
* @param [in] agent_id rocprofiler_agent_id_t to configure thread trace.
* @param [in] parameters List of ATT-specific parameters.
* @param [in] num_parameters Number of parameters. Zero is allowed.
* @param [in] dispatch_callback Control fn which decides when ATT starts/stop collecting.
@@ -84,6 +85,7 @@ typedef rocprofiler_att_control_flags_t (*rocprofiler_att_dispatch_callback_t)(
rocprofiler_status_t
rocprofiler_configure_dispatch_thread_trace_service(
rocprofiler_context_id_t context_id,
rocprofiler_agent_id_t agent_id,
rocprofiler_att_parameter_t* parameters,
size_t num_parameters,
rocprofiler_att_dispatch_callback_t dispatch_callback,
source/lib/output/generateJSON.cpp
+3 -2
Datei anzeigen
@@ -125,9 +125,10 @@ write_json(json_output& json_ar,
auto att_filenames = tool_metadata.get_att_filenames();
auto code_object_snapshot_filenames = std::vector<std::string>{};
code_object_snapshot_filenames.reserve(code_object_load_info.size());
code_object_snapshot_filenames.resize(code_object_load_info.size());
for(const auto& info : code_object_load_info)
code_object_snapshot_filenames.emplace_back(fs::path(info.name).filename());
code_object_snapshot_filenames.at(info.id) = fs::path(info.name).filename();
json_ar.setNextName("strings");
json_ar.startNode();
source/lib/output/metadata.cpp
+16 -1
Datei anzeigen
@@ -23,6 +23,7 @@
#include "metadata.hpp"
#include "lib/common/filesystem.hpp"
#include "lib/common/logging.hpp"
#include "lib/common/string_entry.hpp"
#include "lib/output/agent_info.hpp"
#include "lib/output/host_symbol_info.hpp"
@@ -31,6 +32,7 @@
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/rocprofiler.h>
#include <cstdint>
#include <rocprofiler-sdk/cxx/details/tokenize.hpp>
#include <fmt/core.h>
@@ -246,7 +248,20 @@ metadata::get_code_object_load_info() const
return _info;
});
return _data;
uint64_t _sz = 0;
for(const auto& itr : _data)
_sz = std::max(_sz, itr.id);
ROCP_WARNING_IF((_sz + 1) - _data.size() > 1000) << fmt::format(
"Spares index detected for code object load info: {} < {}", _data.size(), _sz);
auto _code_obj_data = std::vector<rocprofiler::att_wrapper::CodeobjLoadInfo>{};
_code_obj_data.resize(_sz + 1, rocprofiler::att_wrapper::CodeobjLoadInfo{});
// index by the code object id
for(auto& itr : _data)
_code_obj_data.at(itr.id) = itr;
return _code_obj_data;
}
std::vector<std::string>
source/lib/rocprofiler-sdk-att/CMakeLists.txt
+1
Datei anzeigen
@@ -12,6 +12,7 @@ set(ATT_TOOL_SOURCE_FILES
filenames.cpp
occupancy.cpp
wstates.cpp
perfcounter.cpp
profile_interface.cpp
dl.cpp)
source/lib/rocprofiler-sdk-att/att_decoder.h
+11
Datei anzeigen
@@ -124,6 +124,17 @@ typedef struct
att_wave_instruction_t* instructions_array;
} att_wave_data_t;
typedef struct att_perfevent_t
{
int64_t time;
uint16_t events0;
uint16_t events1;
uint16_t events2;
uint16_t events3;
uint8_t CU;
uint8_t bank;
} att_perfevent_t;
typedef rocprofiler_att_decoder_status_t (*rocprofiler_att_decoder_isa_callback_t)(
char* instruction,
uint64_t* memory_size,
source/lib/rocprofiler-sdk-att/att_lib_wrapper.cpp
+7 -4
Datei anzeigen
@@ -48,13 +48,13 @@ get_lib_names()
std::vector<std::pair<tool_att_capability_t, const char*>> lib_names = {
{tool_att_capability_t::ATT_CAPABILITIES_DEBUG, "libatt_decoder_debug.so"},
{tool_att_capability_t::ATT_CAPABILITIES_TRACE, "libatt_decoder_trace.so"},
{tool_att_capability_t::ATT_CAPABILITIES_SUMMARY, "libatt_decoder_summary.so"},
{tool_att_capability_t::ATT_CAPABILITIES_TESTING, "libatt_decoder_testing.so"},
{tool_att_capability_t::ATT_CAPABILITIES_TESTING1, "libatt_decoder_testing1.so"},
{tool_att_capability_t::ATT_CAPABILITIES_TESTING2, "libatt_decoder_testing2.so"},
};
return lib_names;
}
ATTFileMgr::ATTFileMgr(Fspath _dir, std::shared_ptr<DL> _dl)
ATTFileMgr::ATTFileMgr(Fspath _dir, std::shared_ptr<DL> _dl, std::vector<std::string> _counters)
: dir(std::move(_dir))
, dl(std::move(_dl))
{
@@ -62,8 +62,10 @@ ATTFileMgr::ATTFileMgr(Fspath _dir, std::shared_ptr<DL> _dl)
table = std::make_shared<AddressTable>();
codefile = std::make_shared<CodeFile>(dir, table);
filenames = std::make_shared<FilenameMgr>(dir);
for(size_t i = 0; i < ATT_WAVE_STATE_LAST; i++)
wstates.at(i) = std::make_shared<WstatesFile>(i, dir);
filenames->perfcounters = std::move(_counters);
}
ATTFileMgr::~ATTFileMgr() { OccupancyFile::OccupancyFile(dir, table, occupancy); }
@@ -124,6 +126,7 @@ ATTDecoder::parse(const Fspath& input_dir,
const Fspath& output_dir,
const std::vector<std::string>& att_files,
const std::vector<CodeobjLoadInfo>& codeobj_files,
const std::vector<std::string>& counters_names,
const std::string& output_formats)
{
auto& formats = GlobalDefs::get().output_formats;
@@ -132,7 +135,7 @@ ATTDecoder::parse(const Fspath& input_dir,
return std::tolower(c);
});
ATTFileMgr mgr(output_dir, dl);
ATTFileMgr mgr(output_dir, dl, counters_names);
for(const auto& file : codeobj_files)
{
source/lib/rocprofiler-sdk-att/att_lib_wrapper.hpp
+5 -4
Datei anzeigen
@@ -47,9 +47,9 @@ struct CodeobjLoadInfo
enum tool_att_capability_t
{
ATT_CAPABILITIES_TESTING = 0, // used for code coverage testing
ATT_CAPABILITIES_SUMMARY, // used for CSV output only
ATT_CAPABILITIES_TRACE, // used for all outputs
ATT_CAPABILITIES_TESTING1 = 0, // used for code coverage testing
ATT_CAPABILITIES_TESTING2, // used for code coverage testing
ATT_CAPABILITIES_TRACE, // used for all outputs
ATT_CAPABILITIES_DEBUG,
ATT_CAPABILITIES_LAST = ATT_CAPABILITIES_DEBUG,
};
@@ -77,6 +77,7 @@ public:
const Fspath& output_dir,
const std::vector<std::string>& att_files,
const std::vector<CodeobjLoadInfo>& codeobj_files,
const std::vector<std::string>& counters_names,
const std::string& output_formats);
bool valid() const;
@@ -90,7 +91,7 @@ class ATTFileMgr
using AddressTable = rocprofiler::sdk::codeobj::disassembly::CodeobjAddressTranslate;
public:
ATTFileMgr(Fspath _dir, std::shared_ptr<DL> _dl);
ATTFileMgr(Fspath _dir, std::shared_ptr<DL> _dl, std::vector<std::string> _counters);
~ATTFileMgr();
void parseShader(int se_id, const std::vector<char>& data);
source/lib/rocprofiler-sdk-att/filenames.cpp
+6 -5
Datei anzeigen
@@ -50,12 +50,13 @@ FilenameMgr::~FilenameMgr()
[to_string(coord.id)] = {data.name, data.begin, data.end};
}
nlohmann::json metadata = {{"global_begin_time", 0},
{"gfxv", (gfxip > 9) ? "navi" : "vega"},
{"gfxip", gfxip},
{"version", TOOL_VERSION}};
const nlohmann::json metadata = {{"global_begin_time", 0},
{"gfxv", (gfxip > 9) ? "navi" : "vega"},
{"gfxip", gfxip},
{"version", TOOL_VERSION},
{"counter_names", perfcounters},
{"wave_filenames", namelist}};
metadata["wave_filenames"] = namelist;
OutputFile(filename) << metadata;
}
source/lib/rocprofiler-sdk-att/filenames.hpp
+1
Datei anzeigen
@@ -72,6 +72,7 @@ public:
Fspath dir{};
Fspath filename{};
std::map<Coord, WaveName> streams{};
std::vector<std::string> perfcounters{};
int gfxip = 9;
};
source/lib/rocprofiler-sdk-att/perfcounter.cpp
+62
Datei anzeigen
@@ -0,0 +1,62 @@
// MIT License
//
// Copyright (c) 2024-2025 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 "perfcounter.hpp"
#include "outputfile.hpp"
#include "wave.hpp"
#include <fstream>
#include <iostream>
#include <nlohmann/json.hpp>
#include <sstream>
namespace rocprofiler
{
namespace att_wrapper
{
void
PerfcounterFile(WaveConfig& config, const att_perfevent_t* events, size_t event_count)
{
nlohmann::json data;
for(size_t i = 0; i < event_count; i++)
{
const auto& event = events[i];
nlohmann::json json_event;
json_event.push_back(event.time);
json_event.push_back(event.events0);
json_event.push_back(event.events1);
json_event.push_back(event.events2);
json_event.push_back(event.events3);
json_event.push_back(event.CU);
json_event.push_back(event.bank);
data.push_back(json_event);
}
const nlohmann::json json{{"data", data}};
std::string filename = "se" + std::to_string(config.shader_engine) + "_perfcounter.json";
OutputFile(config.filemgr->dir / filename) << json;
}
} // namespace att_wrapper
} // namespace rocprofiler
source/lib/rocprofiler-sdk-att/perfcounter.hpp
+34
Datei anzeigen
@@ -0,0 +1,34 @@
// MIT License
//
// Copyright (c) 2024-2025 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.
#pragma once
#include "att_lib_wrapper.hpp"
namespace rocprofiler
{
namespace att_wrapper
{
void
PerfcounterFile(class WaveConfig& config, const att_perfevent_t* events, size_t event_count);
} // namespace att_wrapper
} // namespace rocprofiler
source/lib/rocprofiler-sdk-att/profile_interface.cpp
+9 -6
Datei anzeigen
@@ -26,12 +26,14 @@
#endif
#include "profile_interface.hpp"
#include "att_decoder.h"
#include "dl.hpp"
#include "perfcounter.hpp"
#include <cxxabi.h>
#include <cstring>
#include <fstream>
#include <shared_mutex>
#include "att_decoder.h"
#include "dl.hpp"
namespace rocprofiler
{
@@ -64,11 +66,8 @@ get_trace_data(rocprofiler_att_decoder_record_type_t trace_id,
auto* infos = (rocprofiler_att_decoder_info_t*) trace_events;
for(size_t i = 0; i < trace_size; i++)
ROCP_WARNING << tool.dl->att_info_fn(infos[i]);
return ROCPROFILER_ATT_DECODER_STATUS_SUCCESS;
}
if(trace_id == ROCPROFILER_ATT_DECODER_TYPE_GFXIP)
else if(trace_id == ROCPROFILER_ATT_DECODER_TYPE_GFXIP)
{
tool.config.filemgr->gfxip = reinterpret_cast<size_t>(trace_events);
}
@@ -78,6 +77,10 @@ get_trace_data(rocprofiler_att_decoder_record_type_t trace_id,
tool.config.occupancy.push_back(
reinterpret_cast<const att_occupancy_info_v2_t*>(trace_events)[i]);
}
else if(trace_id == ROCPROFILER_ATT_DECODER_TYPE_PERFEVENT)
{
PerfcounterFile(tool.config, reinterpret_cast<att_perfevent_t*>(trace_events), trace_size);
}
if(trace_id != ROCPROFILER_ATT_DECODER_TYPE_WAVE) return ROCPROFILER_ATT_DECODER_STATUS_SUCCESS;
source/lib/rocprofiler-sdk-att/tests/CMakeLists.txt
+9 -9
Datei anzeigen
@@ -24,17 +24,17 @@ target_link_libraries(
GTest::gtest
GTest::gtest_main)
add_library(att_decoder_testing SHARED)
add_library(rocprofiler-sdk::att-decoder-testing ALIAS att_decoder_testing)
target_sources(att_decoder_testing PRIVATE dummy_decoder.cpp)
set_target_properties(att_decoder_testing PROPERTIES VERSION ${PROJECT_VERSION}
SOVERSION ${PROJECT_VERSION_MAJOR})
add_library(att_decoder_testing1 SHARED)
add_library(rocprofiler-sdk::att-decoder-testing1 ALIAS att_decoder_testing1)
target_sources(att_decoder_testing1 PRIVATE dummy_decoder.cpp)
set_target_properties(att_decoder_testing1 PROPERTIES VERSION ${PROJECT_VERSION}
SOVERSION ${PROJECT_VERSION_MAJOR})
add_library(att_decoder_summary SHARED)
add_library(rocprofiler-sdk::att-decoder-summary ALIAS att_decoder_summary)
target_sources(att_decoder_summary PRIVATE dummy_decoder.cpp)
add_library(att_decoder_testing2 SHARED)
add_library(rocprofiler-sdk::att-decoder-testing2 ALIAS att_decoder_testing2)
target_sources(att_decoder_testing2 PRIVATE dummy_decoder.cpp)
set_target_properties(
att_decoder_summary
att_decoder_testing2
PROPERTIES LIBRARY_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR}/lib/att
VERSION ${PROJECT_VERSION}
SOVERSION ${PROJECT_VERSION_MAJOR})
source/lib/rocprofiler-sdk-att/tests/att_decoder_test.cpp
+3 -3
Datei anzeigen
@@ -37,7 +37,7 @@ class ATTDecoderTest : public ATTDecoder
{
public:
ATTDecoderTest()
: ATTDecoder(rocprofiler::att_wrapper::tool_att_capability_t::ATT_CAPABILITIES_TESTING)
: ATTDecoder(rocprofiler::att_wrapper::tool_att_capability_t::ATT_CAPABILITIES_TESTING1)
{
rocprofiler::att_wrapper::OutputFile::Enabled() = false;
GlobalDefs::get().output_formats = "json,csv";
@@ -46,7 +46,7 @@ public:
void test_parse()
{
ATTFileMgr mgr("out/", dl);
ATTFileMgr mgr("out/", dl, {});
auto append_isa = [&](const char* line) {
// matches addresses in dummy_decoder.cpp
@@ -107,7 +107,7 @@ TEST(att_decoder_test, warn_failures)
ATTDecoderTest decoder;
ROCP_FATAL_IF(!decoder.valid()) << "Failed to initialize decoder library!";
decoder.parse(".", ".", att_files, codeobjs, "csv,json");
decoder.parse(".", ".", att_files, codeobjs, {}, "csv,json");
}
TEST(att_decoder_test, code_write)
source/lib/rocprofiler-sdk-att/tests/standalone_tool_main.cpp
+5 -4
Datei anzeigen
@@ -53,7 +53,7 @@ main(int argc, char** argv)
FLAGS_colorlogtostderr = true;
});
auto cap = rocprofiler::att_wrapper::tool_att_capability_t::ATT_CAPABILITIES_SUMMARY;
auto cap = rocprofiler::att_wrapper::tool_att_capability_t::ATT_CAPABILITIES_TRACE;
{
auto query = rocprofiler::att_wrapper::query_att_decode_capability();
ROCP_FATAL_IF(query.empty()) << "No decoder capability available!";
@@ -98,7 +98,7 @@ main(int argc, char** argv)
std::unordered_map<int, std::vector<std::string>> all_runs;
for(auto& file : sdk_json["strings"]["att_files"])
for(auto& file : sdk_json["strings"]["att_filenames"])
{
try
{
@@ -119,7 +119,7 @@ main(int argc, char** argv)
std::vector<rocprofiler::att_wrapper::CodeobjLoadInfo> codeobj_files{};
std::map<size_t, std::string> snapshot_files{};
for(auto elem : sdk_json["strings"]["code_object_snapshot_files"])
for(auto elem : sdk_json["strings"]["code_object_snapshot_filenames"])
snapshot_files[elem["key"]] = elem["value"];
for(auto& codeobj : sdk_json["code_objects"])
@@ -158,7 +158,8 @@ main(int argc, char** argv)
std::string run_name = input_path.filename().c_str();
std::string ui_name = run_name.substr(0, run_name.find(".json"));
auto output_dir = output_path / ("ui_output_" + ui_name + std::to_string(run_number));
decoder.parse(input_path.parent_path(), output_dir, att_filenames, codeobj_files, formats);
decoder.parse(
input_path.parent_path(), output_dir, att_filenames, codeobj_files, {}, formats);
}
ROCP_INFO << "Finalizing ATT Tool";
source/lib/rocprofiler-sdk-tool/config.cpp
+4 -12
Datei anzeigen
@@ -113,13 +113,13 @@ has_counter_format(std::string const& str)
}
// validate kernel names
std::unordered_set<uint32_t>
std::unordered_set<size_t>
get_kernel_filter_range(const std::string& kernel_filter)
{
if(kernel_filter.empty()) return {};
auto delim = rocprofiler::sdk::parse::tokenize(kernel_filter, "[], ");
auto range_set = std::unordered_set<uint32_t>{};
auto range_set = std::unordered_set<size_t>{};
for(const auto& itr : delim)
{
if(itr.find('-') != std::string::npos)
@@ -129,8 +129,8 @@ get_kernel_filter_range(const std::string& kernel_filter)
ROCP_FATAL_IF(drange.size() != 2)
<< "bad range format for '" << itr << "'. Expected [A-B] where A and B are numbers";
uint32_t start_range = std::stoul(drange.front());
uint32_t end_range = std::stoul(drange.back());
size_t start_range = std::stoul(drange.front());
size_t end_range = std::stoul(drange.back());
for(auto i = start_range; i <= end_range; i++)
range_set.emplace(i);
}
@@ -180,7 +180,6 @@ parse_att_counters(std::string line)
};
// regex to check if string is of the form "counter_name:simd_mask"
std::regex pattern(R"([a-zA-Z0-9_]+(:0x[0-9a-fA-F]+)?)");
std::set<std::string> unique_counters;
auto input_ss = std::stringstream{line};
@@ -190,13 +189,6 @@ parse_att_counters(std::string line)
input_ss >> counter;
if(counter.empty()) break;
// check if the counter string matches the pattern
if(!std::regex_match(counter, pattern))
{
ROCP_FATAL << "Invalid counter format for ATT: " << counter
<< ". Expected format : Counter_name:optional_simd_mask(hexadecimal)";
}
// Consider only those counters where combination of counter name and simd mask is unique
if(unique_counters.insert(counter).second == false) continue;
source/lib/rocprofiler-sdk-tool/config.hpp
+2 -1
Datei anzeigen
@@ -128,6 +128,7 @@ struct config : output_config
uint64_t att_param_buffer_size = get_env<uint64_t>("ROCPROF_ATT_PARAM_BUFFER_SIZE", 0x6000000);
uint64_t att_param_simd_select = get_env<uint64_t>("ROCPROF_ATT_PARAM_SIMD_SELECT", 0xF);
uint64_t att_param_target_cu = get_env<uint64_t>("ROCPROF_ATT_PARAM_TARGET_CU", 1);
uint64_t att_param_perf_ctrl = get_env<uint64_t>("ROCPROF_ATT_PARAM_PERFCOUNTER_CTRL", 0);
std::string kernel_filter_include = get_env("ROCPROF_KERNEL_FILTER_INCLUDE_REGEX", ".*");
std::string kernel_filter_exclude = get_env("ROCPROF_KERNEL_FILTER_EXCLUDE_REGEX", "");
@@ -135,7 +136,7 @@ struct config : output_config
std::string pc_sampling_unit = get_env("ROCPROF_PC_SAMPLING_UNIT", "none");
std::string extra_counters_contents = get_env("ROCPROF_EXTRA_COUNTERS_CONTENTS", "");
std::unordered_set<uint32_t> kernel_filter_range = {};
std::unordered_set<size_t> kernel_filter_range = {};
std::vector<std::set<std::string>> counters = {};
std::string att_capability = get_env("ROCPROF_ATT_CAPABILITY", "");
std::vector<att_perfcounter> att_param_perfcounters = {};
source/lib/rocprofiler-sdk-tool/tool.cpp
+135 -121
Datei anzeigen
@@ -193,17 +193,16 @@ as_pointer()
}
using targeted_kernels_map_t =
std::unordered_map<rocprofiler_kernel_id_t, std::unordered_set<uint32_t>>;
std::unordered_map<rocprofiler_kernel_id_t, std::unordered_set<size_t>>;
using counter_dimension_info_map_t =
std::unordered_map<uint64_t, std::vector<rocprofiler_record_dimension_info_t>>;
using agent_info_map_t = std::unordered_map<rocprofiler_agent_id_t, rocprofiler_agent_t>;
using kernel_iteration_t = std::unordered_map<rocprofiler_kernel_id_t, uint32_t>;
using kernel_iteration_t = std::unordered_map<rocprofiler_kernel_id_t, size_t>;
using kernel_rename_map_t = std::unordered_map<uint64_t, uint64_t>;
using kernel_rename_stack_t = std::stack<uint64_t>;
auto* tool_metadata = as_pointer<tool::metadata>(tool::metadata::inprocess{});
auto target_kernels = common::Synchronized<targeted_kernels_map_t>{};
auto kernel_iteration = common::Synchronized<kernel_iteration_t, true>{};
auto* tool_metadata = as_pointer<tool::metadata>(tool::metadata::inprocess{});
auto target_kernels = common::Synchronized<targeted_kernels_map_t>{};
std::mutex att_shader_data;
thread_local auto thread_dispatch_rename = as_pointer<kernel_rename_stack_t>();
@@ -239,13 +238,13 @@ add_kernel_rename_and_stream_display_pairs(kernel_rename_and_stream_display_pair
}
bool
add_kernel_target(uint64_t _kern_id, const std::unordered_set<uint32_t>& range)
add_kernel_target(uint64_t _kern_id, const std::unordered_set<size_t>& range)
{
return target_kernels
.wlock(
[](targeted_kernels_map_t& _targets_v,
uint64_t _kern_id_v,
const std::unordered_set<uint32_t>& _range) {
[](targeted_kernels_map_t& _targets_v,
uint64_t _kern_id_v,
const std::unordered_set<size_t>& _range) {
return _targets_v.emplace(_kern_id_v, _range);
},
_kern_id,
@@ -254,10 +253,11 @@ add_kernel_target(uint64_t _kern_id, const std::unordered_set<uint32_t>& range)
}
bool
is_targeted_kernel(uint64_t _kern_id)
is_targeted_kernel(uint64_t _kern_id,
common::Synchronized<kernel_iteration_t, true>& _kernel_iteration)
{
const std::unordered_set<uint32_t>* range = target_kernels.rlock(
[](const auto& _targets_v, uint64_t _kern_id_v) -> const std::unordered_set<uint32_t>* {
const std::unordered_set<size_t>* range = target_kernels.rlock(
[](const auto& _targets_v, uint64_t _kern_id_v) -> const std::unordered_set<size_t>* {
if(_targets_v.find(_kern_id_v) != _targets_v.end()) return &_targets_v.at(_kern_id_v);
return nullptr;
},
@@ -265,21 +265,28 @@ is_targeted_kernel(uint64_t _kern_id)
if(range)
{
return kernel_iteration.rlock(
[](const auto& _kernel_iter,
uint64_t _kernel_id,
const std::unordered_set<uint32_t>& _range) {
auto itr = _kernel_iter.at(_kernel_id);
_kernel_iteration.wlock(
[](auto& _kernel_iter, rocprofiler_kernel_id_t _kernel_id) {
auto itr = _kernel_iter.find(_kernel_id);
if(itr == _kernel_iter.end())
_kernel_iter.emplace(_kernel_id, 1);
else
itr->second++;
},
_kern_id);
return _kernel_iteration.rlock(
[](const auto& _kernel_iter,
uint64_t _kernel_id,
const std::unordered_set<size_t>& _range) {
auto itr = _kernel_iter.at(_kernel_id);
// If the iteration range is not given then all iterations of the kernel is profiled
if(_range.empty())
{
if(!tool::get_config().advanced_thread_trace)
return true;
else
{
if(itr == 1) return true;
}
else if(itr == 1)
return true;
}
else if(_range.find(itr) != _range.end())
return true;
@@ -1129,32 +1136,33 @@ get_instruction_index(rocprofiler_pc_t pc)
} // namespace
std::vector<rocprofiler_att_parameter_t>
get_att_perfcounter_params(std::vector<rocprofiler::tool::att_perfcounter>& att_perf_counters)
get_att_perfcounter_params(rocprofiler_agent_id_t agent,
std::vector<rocprofiler::tool::att_perfcounter>& att_perf_counters)
{
std::vector<rocprofiler_att_parameter_t> _data;
std::vector<rocprofiler_att_parameter_t> _data{};
if(att_perf_counters.empty()) return _data;
static const auto gpu_agents = get_gpu_agents();
static const auto gpu_agents_counter_info = get_agent_counter_info();
static const auto agent_counter_info = get_agent_counter_info();
for(const auto& [agent_, tool_counter_info_] : gpu_agents_counter_info)
for(const auto& att_perf_counter : att_perf_counters)
{
for(const auto& counter_info_ : tool_counter_info_)
{
if(std::string_view(counter_info_.block) != "SQ") continue;
bool counter_found = false;
for(const auto& att_perf_counter : att_perf_counters)
{
if(std::string_view(counter_info_.name) == att_perf_counter.counter_name)
{
auto param = rocprofiler_att_parameter_t{};
param.type = ROCPROFILER_ATT_PARAMETER_PERFCOUNTER,
param.counter_id = counter_info_.id,
param.simd_mask = att_perf_counter.simd_mask;
_data.emplace_back(param);
}
}
for(const auto& counter_info_ : agent_counter_info.at(agent))
{
if(std::string_view(counter_info_.name) != att_perf_counter.counter_name) continue;
auto param = rocprofiler_att_parameter_t{};
param.type = ROCPROFILER_ATT_PARAMETER_PERFCOUNTER;
param.counter_id = counter_info_.id;
param.simd_mask = att_perf_counter.simd_mask;
_data.emplace_back(param);
counter_found = true;
break;
}
ROCP_WARNING_IF(!counter_found)
<< "Agent " << agent.handle << " counter not found: " << att_perf_counter.counter_name;
}
return _data;
@@ -1232,23 +1240,12 @@ att_dispatch_callback(rocprofiler_agent_id_t /* agent_id */,
void* /*userdata_config*/,
rocprofiler_user_data_t* userdata_shader)
{
userdata_shader->value = dispatch_id;
kernel_iteration.wlock(
[](auto& _kernel_iter, rocprofiler_kernel_id_t _kernel_id) {
auto itr = _kernel_iter.find(_kernel_id);
if(itr == _kernel_iter.end())
_kernel_iter.emplace(_kernel_id, 1);
else
{
itr->second++;
}
},
kernel_id);
static auto kernel_iteration = common::Synchronized<kernel_iteration_t, true>{};
if(is_targeted_kernel(kernel_id))
{
userdata_shader->value = dispatch_id;
if(is_targeted_kernel(kernel_id, kernel_iteration))
return ROCPROFILER_ATT_CONTROL_START_AND_STOP;
}
return ROCPROFILER_ATT_CONTROL_NONE;
}
@@ -1258,22 +1255,12 @@ dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data,
rocprofiler_user_data_t* user_data,
void* /*callback_data_args*/)
{
static auto kernel_iteration = common::Synchronized<kernel_iteration_t, true>{};
auto kernel_id = dispatch_data.dispatch_info.kernel_id;
auto agent_id = dispatch_data.dispatch_info.agent_id;
kernel_iteration.wlock(
[](auto& _kernel_iter, rocprofiler_kernel_id_t _kernel_id) {
auto itr = _kernel_iter.find(_kernel_id);
if(itr == _kernel_iter.end())
_kernel_iter.emplace(_kernel_id, 1);
else
{
itr->second++;
}
},
kernel_id);
if(!is_targeted_kernel(kernel_id))
if(!is_targeted_kernel(kernel_id, kernel_iteration))
{
return;
}
@@ -1555,34 +1542,50 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
if(tool::get_config().advanced_thread_trace)
{
auto parameters = std::vector<rocprofiler_att_parameter_t>{};
auto global_parameters = std::vector<rocprofiler_att_parameter_t>{};
uint64_t target_cu = tool::get_config().att_param_target_cu;
uint64_t simd_select = tool::get_config().att_param_simd_select;
uint64_t buffer_sz = tool::get_config().att_param_buffer_size;
uint64_t shader_mask = tool::get_config().att_param_shader_engine_mask;
uint64_t perfcounter_ctrl = tool::get_config().att_param_perf_ctrl;
auto& att_perf = tool::get_config().att_param_perfcounters;
bool att_serialize_all = tool::get_config().att_serialize_all;
auto att_perf_params = get_att_perfcounter_params(att_perf);
parameters.insert(parameters.end(), att_perf_params.begin(), att_perf_params.end());
// TODO: att params could be different for different devices. How to support?
// Input file schema might also need to change to support multiple ATT params
parameters.push_back({ROCPROFILER_ATT_PARAMETER_TARGET_CU, {target_cu}});
parameters.push_back({ROCPROFILER_ATT_PARAMETER_SIMD_SELECT, {simd_select}});
parameters.push_back({ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE, {buffer_sz}});
parameters.push_back({ROCPROFILER_ATT_PARAMETER_SHADER_ENGINE_MASK, {shader_mask}});
parameters.push_back(
global_parameters.push_back({ROCPROFILER_ATT_PARAMETER_TARGET_CU, {target_cu}});
global_parameters.push_back({ROCPROFILER_ATT_PARAMETER_SIMD_SELECT, {simd_select}});
global_parameters.push_back({ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE, {buffer_sz}});
global_parameters.push_back({ROCPROFILER_ATT_PARAMETER_SHADER_ENGINE_MASK, {shader_mask}});
global_parameters.push_back(
{ROCPROFILER_ATT_PARAMETER_SERIALIZE_ALL, {static_cast<uint64_t>(att_serialize_all)}});
ROCPROFILER_CALL(
rocprofiler_configure_dispatch_thread_trace_service(get_client_ctx(),
parameters.data(),
parameters.size(),
att_dispatch_callback,
att_shader_data_callback,
tool_data),
"thread trace service configure");
if(perfcounter_ctrl != 0 && !att_perf.empty())
{
global_parameters.push_back(
{ROCPROFILER_ATT_PARAMETER_PERFCOUNTERS_CTRL, {perfcounter_ctrl}});
}
else if(perfcounter_ctrl != 0 || !att_perf.empty())
{
ROCP_FATAL << "ATT Perf requires setting both perfcounter_ctrl and perfcounter list!";
}
for(auto& [id, agent] : tool_metadata->agents_map)
{
if(agent.type != ROCPROFILER_AGENT_TYPE_GPU) continue;
auto agent_params = global_parameters;
for(auto& counter : get_att_perfcounter_params(id, att_perf))
agent_params.push_back(counter);
ROCPROFILER_CALL(
rocprofiler_configure_dispatch_thread_trace_service(get_client_ctx(),
id,
agent_params.data(),
agent_params.size(),
att_dispatch_callback,
att_shader_data_callback,
tool_data),
"thread trace service configure");
}
}
if(tool::get_config().hip_runtime_api_trace || tool::get_config().hip_compiler_api_trace)
@@ -1933,40 +1936,6 @@ tool_fini(void* /*tool_data*/)
tool::generate_csv(tool::get_config(), *tool_metadata, contributions);
}
if(tool::get_config().advanced_thread_trace)
{
const std::unordered_map<std::string_view, rocprofiler::att_wrapper::tool_att_capability_t>
tool_att_capability_map = {
{"testing", rocprofiler::att_wrapper::ATT_CAPABILITIES_TESTING},
{"summary", rocprofiler::att_wrapper::ATT_CAPABILITIES_SUMMARY},
{"trace", rocprofiler::att_wrapper::ATT_CAPABILITIES_TRACE},
{"debug", rocprofiler::att_wrapper::ATT_CAPABILITIES_DEBUG}};
ROCP_FATAL_IF(tool::get_config().att_capability.empty())
<< "Provide the decoder parser method as input";
auto att_capability_value = tool_att_capability_map.at(tool::get_config().att_capability);
auto decoder = rocprofiler::att_wrapper::ATTDecoder(att_capability_value);
ROCP_FATAL_IF(!decoder.valid()) << "Decoder library not found at ROCPROF_ATT_LIBRARY_PATH";
auto codeobj = tool_metadata->get_code_object_load_info();
auto output_path = tool::format_path(tool::get_config().output_path);
for(auto& [dispatch_id, att_filename_data] : tool_metadata->att_filenames)
{
std::string formats = "json,csv";
// if(tool::get_config().json_output) formats += "json,";
// if(tool::get_config().csv_output) formats += "csv,";
std::stringstream ui_name;
ui_name << fmt::format("ui_output_agent_{}_dispatch_{}",
std::to_string(att_filename_data.first.handle),
dispatch_id);
auto out_path = fmt::format("{}/{}", output_path, ui_name.str());
std::string in_path = ".";
decoder.parse(in_path, out_path, att_filename_data.second, codeobj, formats);
}
}
if(tool::get_config().json_output && num_output > 0)
{
auto json_ar = tool::open_json(tool::get_config());
@@ -2044,6 +2013,51 @@ tool_fini(void* /*tool_data*/)
{
tool::generate_stats(tool::get_config(), *tool_metadata, contributions);
}
if(tool::get_config().advanced_thread_trace)
{
const std::unordered_map<std::string_view, rocprofiler::att_wrapper::tool_att_capability_t>
tool_att_capability_map = {
{"testing1", rocprofiler::att_wrapper::ATT_CAPABILITIES_TESTING1},
{"testing2", rocprofiler::att_wrapper::ATT_CAPABILITIES_TESTING2},
{"trace", rocprofiler::att_wrapper::ATT_CAPABILITIES_TRACE},
{"debug", rocprofiler::att_wrapper::ATT_CAPABILITIES_DEBUG}};
ROCP_FATAL_IF(tool::get_config().att_capability.empty())
<< "Provide the decoder parser method as input";
auto att_capability_value = tool_att_capability_map.at(tool::get_config().att_capability);
auto decoder = rocprofiler::att_wrapper::ATTDecoder(att_capability_value);
ROCP_FATAL_IF(!decoder.valid()) << "Decoder library not found at ROCPROF_ATT_LIBRARY_PATH";
auto codeobj = tool_metadata->get_code_object_load_info();
auto output_path = tool::format_path(tool::get_config().output_path);
std::vector<std::string> perf{};
for(auto& counter : tool::get_config().att_param_perfcounters)
{
std::stringstream ss;
ss << counter.counter_name;
if(counter.simd_mask != 0xF) ss << ':' << std::hex << counter.simd_mask;
perf.emplace_back(ss.str());
}
for(auto& [dispatch_id, att_filename_data] : tool_metadata->att_filenames)
{
std::string formats = "json,csv";
auto ui_name = std::stringstream{};
ui_name << fmt::format("ui_output_agent_{}_dispatch_{}",
std::to_string(att_filename_data.first.handle),
dispatch_id);
auto out_path = fmt::format("{}/{}", output_path, ui_name.str());
auto in_path = std::string(".");
decoder.parse(in_path, out_path, att_filename_data.second, codeobj, perf, formats);
}
}
auto destroy_output = [](auto& _buffered_output_v) { _buffered_output_v.destroy(); };
destroy_output(kernel_dispatch_with_stream_output);
source/lib/rocprofiler-sdk/context/context.cpp
+2 -2
Datei anzeigen
@@ -339,7 +339,7 @@ start_context(rocprofiler_context_id_t context_id)
auto status = ROCPROFILER_STATUS_SUCCESS;
if(cfg->counter_collection) rocprofiler::counters::start_context(cfg);
if(cfg->agent_thread_trace) cfg->agent_thread_trace->start_context();
if(cfg->device_thread_trace) cfg->device_thread_trace->start_context();
if(cfg->dispatch_thread_trace) cfg->dispatch_thread_trace->start_context();
if(cfg->device_counter_collection) status = rocprofiler::counters::start_agent_ctx(cfg);
#if ROCPROFILER_SDK_HSA_PC_SAMPLING > 0
@@ -374,7 +374,7 @@ stop_context(rocprofiler_context_id_t idx)
rocprofiler::counters::stop_context(const_cast<context*>(_expected));
}
if(_expected->agent_thread_trace) _expected->agent_thread_trace->stop_context();
if(_expected->device_thread_trace) _expected->device_thread_trace->stop_context();
if(_expected->dispatch_thread_trace)
_expected->dispatch_thread_trace->stop_context();
source/lib/rocprofiler-sdk/context/context.hpp
+1 -1
Datei anzeigen
@@ -131,7 +131,7 @@ struct context
std::unique_ptr<pc_sampling_service> pc_sampler = {};
std::unique_ptr<thread_trace::DispatchThreadTracer> dispatch_thread_trace = {};
std::unique_ptr<thread_trace::AgentThreadTracer> agent_thread_trace = {};
std::unique_ptr<thread_trace::DeviceThreadTracer> device_thread_trace = {};
template <typename KindT>
bool is_tracing(KindT _kind) const;
source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp
+12 -43
Datei anzeigen
@@ -46,7 +46,7 @@ namespace counters
*
* We return an AQLPacket containing the start/stop/read packets for injection.
*/
std::unique_ptr<rocprofiler::hsa::AQLPacket>
hsa::Queue::pkt_and_serialize_t
queue_cb(const context::context* ctx,
const std::shared_ptr<counter_callback_info>& info,
const hsa::Queue& queue,
@@ -59,43 +59,25 @@ queue_cb(const context::context* ctx,
{
CHECK(info && ctx);
// Maybe adds serialization packets to the AQLPacket (if serializer is enabled)
// and maybe adds barrier packets if the state is transitioning from serialized <->
// unserialized
auto maybe_add_serialization = [&](auto& gen_pkt) {
CHECK_NOTNULL(hsa::get_queue_controller())
->serializer(&queue)
.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);
}
});
};
// Packet generated when no instrumentation is performed. May contain serialization
// packets/barrier packets (and can be empty).
auto no_instrumentation = [&]() {
auto ret_pkt = std::make_unique<rocprofiler::hsa::EmptyAQLPacket>();
info->packet_return_map.wlock([&](auto& data) { data.emplace(ret_pkt.get(), nullptr); });
// If we have a counter collection context but it is not enabled, we still might need
// to add barrier packets to transition from serialized -> unserialized execution. This
// transition is coordinated by the serializer.
maybe_add_serialization(ret_pkt);
info->packet_return_map.wlock([&](auto& data) { data.emplace(ret_pkt.get(), nullptr); });
return ret_pkt;
};
if(!ctx || !ctx->counter_collection) return nullptr;
if(!ctx || !ctx->counter_collection) return {nullptr, false};
bool is_enabled = false;
ctx->counter_collection->enabled.rlock(
[&](const auto& collect_ctx) { is_enabled = collect_ctx; });
if(!is_enabled || !info->user_cb)
{
return no_instrumentation();
}
if(!is_enabled || !info->user_cb) return {no_instrumentation(), true};
auto _corr_id_v =
rocprofiler_correlation_id_t{.internal = 0, .external = context::null_user_data};
@@ -133,10 +115,7 @@ queue_cb(const context::context* ctx,
info->user_cb(dispatch_data, &req_profile, user_data, info->callback_args);
if(req_profile.handle == 0)
{
return no_instrumentation();
}
if(req_profile.handle == 0) return {no_instrumentation(), true};
auto prof_config = get_controller().get_profile_cfg(req_profile);
CHECK(prof_config);
@@ -145,18 +124,15 @@ queue_cb(const context::context* ctx,
auto status = info->get_packet(ret_pkt, prof_config);
CHECK_EQ(status, ROCPROFILER_STATUS_SUCCESS) << rocprofiler_get_status_string(status);
maybe_add_serialization(ret_pkt);
if(ret_pkt->empty)
if(!ret_pkt->empty)
{
return ret_pkt;
ret_pkt->populate_before();
ret_pkt->populate_after();
for(auto& aql_pkt : ret_pkt->after_krn_pkt)
aql_pkt.completion_signal.handle = 0;
}
ret_pkt->populate_before();
ret_pkt->populate_after();
for(auto& aql_pkt : ret_pkt->after_krn_pkt)
aql_pkt.completion_signal.handle = 0;
return ret_pkt;
return {std::move(ret_pkt), true};
}
/**
@@ -169,7 +145,6 @@ completed_cb(const context::context* ctx,
inst_pkt_t& pkts,
kernel_dispatch::profiling_time dispatch_time)
{
auto& session = *ptr_session;
CHECK(info && ctx);
std::shared_ptr<profile_config> prof_config;
@@ -189,14 +164,8 @@ completed_cb(const context::context* ctx,
}
});
if(!pkt) return;
CHECK_NOTNULL(hsa::get_queue_controller())
->serializer(&session.queue)
.wlock([&](auto& serializer) { serializer.kernel_completion_signal(session.queue); });
// We have no profile config, nothing to output.
if(!prof_config) return;
if(!pkt || !prof_config) return;
completed_cb_params_t params{info, ptr_session, dispatch_time, prof_config, std::move(pkt)};
process_callback_data(std::move(params));
source/lib/rocprofiler-sdk/counters/dispatch_handlers.hpp
+1 -1
Datei anzeigen
@@ -34,7 +34,7 @@ using ClientID = int64_t;
using inst_pkt_t = common::container::
small_vector<std::pair<std::unique_ptr<rocprofiler::hsa::AQLPacket>, ClientID>, 4>;
std::unique_ptr<rocprofiler::hsa::AQLPacket>
hsa::Queue::pkt_and_serialize_t
queue_cb(const context::context* ctx,
const std::shared_ptr<counter_callback_info>& info,
const hsa::Queue& queue,
source/lib/rocprofiler-sdk/counters/tests/core.cpp
+3 -3
Datei anzeigen
@@ -444,8 +444,8 @@ TEST(core, check_callbacks)
extern_ids,
&corr_id);
ASSERT_TRUE(ret_pkt) << fmt::format("Expected a packet to be generated for - {}",
metric.name());
ASSERT_TRUE(ret_pkt.pkt)
<< fmt::format("Expected a packet to be generated for - {}", metric.name());
/**
* Create the buffer and run test
@@ -469,7 +469,7 @@ TEST(core, check_callbacks)
counters::inst_pkt_t pkts;
pkts.emplace_back(
std::make_pair(std::move(ret_pkt), static_cast<counters::ClientID>(0)));
std::make_pair(std::move(ret_pkt.pkt), static_cast<counters::ClientID>(0)));
completed_cb(&ctx, cb_info, sess, pkts, kernel_dispatch::profiling_time{});
rocprofiler_flush_buffer(opt_buff_id);
rocprofiler_destroy_buffer(opt_buff_id);
source/lib/rocprofiler-sdk/hsa/queue.cpp
+35 -12
Datei anzeigen
@@ -137,6 +137,15 @@ AsyncSignalHandler(hsa_signal_value_t /*signal_v*/, void* data)
}
});
if(queue_info_session.is_serialized)
{
CHECK_NOTNULL(hsa::get_queue_controller())
->serializer(&queue_info_session.queue)
.wlock([&](auto& serializer) {
serializer.kernel_completion_signal(queue_info_session.queue);
});
}
// Delete signals and packets, signal we have completed.
if(queue_info_session.interrupt_signal.handle != 0u)
{
@@ -367,24 +376,36 @@ WriteInterceptor(const void* packets,
// completed_cb_t)
auto inst_pkt = inst_pkt_t{};
// True if any service (ATT,SPM,CC) requests this dispatch to be serialized
bool bRequest_Serialize = false;
// Signal callbacks that a kernel_pkt is being enqueued
queue.signal_callback([&](const auto& map) {
for(const auto& [client_id, cb_pair] : map)
{
if(auto maybe_pkt = cb_pair.first(queue,
kernel_pkt,
kernel_id,
dispatch_id,
&user_data,
tracing_data_v.external_correlation_ids,
corr_id))
{
inst_pkt.push_back(std::make_pair(std::move(maybe_pkt), client_id));
}
auto [packet, bSerial] = cb_pair.first(queue,
kernel_pkt,
kernel_id,
dispatch_id,
&user_data,
tracing_data_v.external_correlation_ids,
corr_id);
bRequest_Serialize |= bSerial;
if(packet) inst_pkt.push_back(std::make_pair(std::move(packet), client_id));
}
});
bool inserted_before = false;
if(bRequest_Serialize)
{
inserted_before = true;
CHECK_NOTNULL(hsa::get_queue_controller())
->serializer(&queue)
.rlock([&](const auto& serializer) {
for(auto& s_pkt : serializer.kernel_dispatch(queue))
transformed_packets.emplace_back(s_pkt.ext_amd_aql_pm4);
});
}
for(const auto& pkt_injection : inst_pkt)
{
for(const auto& pkt : pkt_injection.first->before_krn_pkt)
@@ -460,7 +481,8 @@ WriteInterceptor(const void* packets,
.correlation_id = corr_id,
.kernel_pkt = kernel_pkt,
.callback_record = callback_record,
.tracing_data = tracing_data_v};
.tracing_data = tracing_data_v,
.is_serialized = bRequest_Serialize};
auto shared = std::make_shared<Queue::queue_info_session_t>(std::move(info_session));
@@ -521,7 +543,8 @@ Queue::Queue(const AgentCache& agent,
<< "Could not setup intercept profiler";
if(!context::get_registered_contexts([](const context::context* ctx) {
return (ctx->counter_collection || ctx->device_counter_collection);
return (ctx->counter_collection || ctx->device_counter_collection ||
ctx->dispatch_thread_trace || ctx->device_thread_trace);
}).empty())
{
CHECK(_agent.cpu_pool().handle != 0);
source/lib/rocprofiler-sdk/hsa/queue.hpp
+17 -11
Datei anzeigen
@@ -73,17 +73,23 @@ public:
using callback_t = void (*)(hsa_status_t status, hsa_queue_t* source, void* data);
using queue_info_session_t = queue_info_session;
// Function prototype used to notify consumers that a kernel has been
// enqueued. An AQL packet can be returned that will be injected into
// the queue.
using queue_cb_t = std::function<std::unique_ptr<AQLPacket>(
const Queue&,
const rocprofiler_packet&,
rocprofiler_kernel_id_t,
rocprofiler_dispatch_id_t,
rocprofiler_user_data_t*,
const queue_info_session_t::external_corr_id_map_t&,
const context::correlation_id*)>;
struct pkt_and_serialize_t
{
std::unique_ptr<AQLPacket> pkt{nullptr};
bool request_serialize{false};
};
// Function prototype used to notify consumers that a kernel has been enqueued.
// Pair first: An AQL packet can be returned that will be injected into the queue.
// Pair second: Boolean flag indicating the dispatch needs to be serialized.
using queue_cb_t =
std::function<pkt_and_serialize_t(const Queue&,
const rocprofiler_packet&,
rocprofiler_kernel_id_t,
rocprofiler_dispatch_id_t,
rocprofiler_user_data_t*,
const queue_info_session_t::external_corr_id_map_t&,
const context::correlation_id*)>;
// Signals the completion of the kernel packet.
using completed_cb_t = std::function<void(const Queue&,
const rocprofiler_packet&,
source/lib/rocprofiler-sdk/hsa/queue_controller.cpp
+1 -17
Datei anzeigen
@@ -148,12 +148,6 @@ constexpr rocprofiler_agent_t default_agent =
void
QueueController::add_queue(hsa_queue_t* id, std::unique_ptr<Queue> queue)
{
for(const auto& itr : context::get_registered_contexts())
{
if(auto* trace = itr->dispatch_thread_trace.get())
trace->resource_init(queue->get_agent(), get_core_table(), get_ext_table());
}
CHECK(queue);
_callback_cache.wlock([&](auto& callbacks) {
_queues.wlock([&](auto& map) {
@@ -176,16 +170,6 @@ QueueController::destroy_queue(hsa_queue_t* id)
{
if(!id) return;
for(const auto& itr : context::get_registered_contexts())
{
if(!itr->dispatch_thread_trace) continue;
_queues.wlock([&](auto& map) {
if(map.find(id) != map.end())
itr->dispatch_thread_trace->resource_deinit(map.at(id)->get_agent());
});
}
const auto* queue = get_queue(*id);
// return if queue does not exist
@@ -469,7 +453,7 @@ enable_queue_intercept()
itr->is_tracing(ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY);
if(itr->counter_collection || itr->pc_sampler || has_kernel_tracing ||
has_scratch_reporting || itr->device_counter_collection || itr->agent_thread_trace ||
has_scratch_reporting || itr->device_counter_collection || itr->device_thread_trace ||
itr->dispatch_thread_trace)
return true;
}
source/lib/rocprofiler-sdk/hsa/queue_info_session.hpp
+1
Datei anzeigen
@@ -64,6 +64,7 @@ struct queue_info_session
rocprofiler_packet kernel_pkt = {};
callback_record_t callback_record = {};
tracing::tracing_data tracing_data = {};
bool is_serialized = false;
};
} // namespace hsa
} // namespace rocprofiler
source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp
+3 -1
Datei anzeigen
@@ -347,7 +347,9 @@ pc_sampling_service_finish_configuration(context::pc_sampling_service* service)
rocprofiler_dispatch_id_t /*dispatch_id*/,
rocprofiler_user_data_t*,
const rocprofiler::hsa::Queue::queue_info_session_t::external_corr_id_map_t&,
const context::correlation_id*) { return nullptr; },
const context::correlation_id*) {
return rocprofiler::hsa::Queue::pkt_and_serialize_t{};
},
// Completion CB
[](const rocprofiler::hsa::Queue& q,
rocprofiler::hsa::rocprofiler_packet kern_pkt,
source/lib/rocprofiler-sdk/registration.cpp
+1 -1
Datei anzeigen
@@ -890,6 +890,7 @@ rocprofiler_set_api_table(const char* name,
// need to construct agent mappings before initializing the queue controller
rocprofiler::agent::construct_agent_cache(hsa_api_table);
rocprofiler::thread_trace::initialize(hsa_api_table);
rocprofiler::hsa::queue_controller_init(hsa_api_table);
// Process agent ctx's that were started prior to HSA init
rocprofiler::counters::device_counting_service_hsa_registration();
@@ -898,7 +899,6 @@ rocprofiler_set_api_table(const char* name,
rocprofiler::hsa::memory_allocation_init(hsa_api_table->core_, lib_instance);
rocprofiler::hsa::memory_allocation_init(hsa_api_table->amd_ext_, lib_instance);
rocprofiler::code_object::initialize(hsa_api_table);
rocprofiler::thread_trace::initialize(hsa_api_table);
#if ROCPROFILER_SDK_HSA_PC_SAMPLING > 0
if(runtime_pc_sampling_table)
rocprofiler::pc_sampling::code_object::initialize(hsa_api_table);
source/lib/rocprofiler-sdk/thread_trace/att_core.cpp
+100 -148
Datei anzeigen
@@ -26,6 +26,7 @@
#include "lib/common/container/stable_vector.hpp"
#include "lib/common/utility.hpp"
#include "lib/rocprofiler-sdk/agent.hpp"
#include "lib/rocprofiler-sdk/buffer.hpp"
#include "lib/rocprofiler-sdk/context/context.hpp"
#include "lib/rocprofiler-sdk/hsa/queue_controller.hpp"
@@ -68,20 +69,6 @@ struct cbdata_t
common::Synchronized<std::optional<int64_t>> client;
CoreApiTable&
get_core()
{
static CoreApiTable api{};
return api;
}
AmdExtTable&
get_ext()
{
static AmdExtTable api{};
return api;
}
bool
thread_trace_parameter_pack::are_params_valid() const
{
@@ -110,14 +97,16 @@ class Signal
public:
Signal(hsa_ext_amd_aql_pm4_packet_t* packet)
{
get_ext().hsa_amd_signal_create_fn(0, 0, nullptr, 0, &signal);
auto& core = *hsa::get_core_table();
auto& ext = *hsa::get_amd_ext_table();
ext.hsa_amd_signal_create_fn(0, 0, nullptr, 0, &signal);
packet->completion_signal = signal;
get_core().hsa_signal_store_screlease_fn(signal, 1);
core.hsa_signal_store_screlease_fn(signal, 1);
}
~Signal()
{
WaitOn();
get_core().hsa_signal_destroy_fn(signal);
hsa::get_core_table()->hsa_signal_destroy_fn(signal);
}
Signal(Signal& other) = delete;
Signal(const Signal& other) = delete;
@@ -126,7 +115,7 @@ public:
void WaitOn() const
{
auto* wait_fn = get_core().hsa_signal_wait_scacquire_fn;
auto wait_fn = hsa::get_core_table()->hsa_signal_wait_scacquire_fn;
while(wait_fn(signal, HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX, HSA_WAIT_STATE_BLOCKED) != 0)
{}
}
@@ -136,10 +125,12 @@ public:
};
std::unique_ptr<Signal>
ThreadTracerQueue::Submit(hsa_ext_amd_aql_pm4_packet_t* packet, bool bWait)
ThreadTracerQueue::Submit(hsa_ext_amd_aql_pm4_packet_t* packet, bool bWait) const
{
auto* core = hsa::get_core_table();
std::unique_ptr<Signal> signal{};
const uint64_t write_idx = add_write_index_relaxed_fn(queue, 1);
const uint64_t write_idx = core->hsa_queue_add_write_index_relaxed_fn(queue, 1);
size_t index = (write_idx % queue->size) * sizeof(hsa_ext_amd_aql_pm4_packet_t);
// NOLINTNEXTLINE(performance-no-int-to-ptr)
@@ -154,40 +145,44 @@ ThreadTracerQueue::Submit(hsa_ext_amd_aql_pm4_packet_t* packet, bool bWait)
auto* header = reinterpret_cast<std::atomic<uint32_t>*>(queue_slot);
header->store(slot_data[0], std::memory_order_release);
signal_store_screlease_fn(queue->doorbell_signal, write_idx);
core->hsa_signal_store_screlease_fn(queue->doorbell_signal, write_idx);
return signal;
}
ThreadTracerQueue::ThreadTracerQueue(thread_trace_parameter_pack _params,
const hsa::AgentCache& cache,
const CoreApiTable& coreapi,
const AmdExtTable& ext)
rocprofiler_agent_id_t cache)
: params(std::move(_params))
, agent_id(cache.get_rocp_agent()->id)
, agent_id(cache)
{
factory = std::make_unique<aql::ThreadTraceAQLPacketFactory>(cache, this->params, coreapi, ext);
ROCP_TRACE << "Constructing ATT instance for agent " << agent_id.handle;
auto* core = hsa::get_core_table();
auto* ext = hsa::get_amd_ext_table();
factory = std::make_unique<aql::ThreadTraceAQLPacketFactory>(
*rocprofiler::agent::get_agent_cache(rocprofiler::agent::get_agent(agent_id)),
this->params,
*core,
*ext);
control_packet = factory->construct_control_packet();
auto status = coreapi.hsa_queue_create_fn(cache.get_hsa_agent(),
QUEUE_SIZE,
HSA_QUEUE_TYPE_SINGLE,
nullptr,
nullptr,
UINT32_MAX,
UINT32_MAX,
&this->queue);
auto hsa_agent = rocprofiler::agent::get_hsa_agent(agent_id);
CHECK(hsa_agent.has_value());
auto status = core->hsa_queue_create_fn(*hsa_agent,
QUEUE_SIZE,
HSA_QUEUE_TYPE_SINGLE,
nullptr,
nullptr,
UINT32_MAX,
UINT32_MAX,
&this->queue);
if(status != HSA_STATUS_SUCCESS)
{
ROCP_ERROR << "Failed to create thread trace async queue";
this->queue = nullptr;
}
queue_destroy_fn = coreapi.hsa_queue_destroy_fn;
signal_store_screlease_fn = coreapi.hsa_signal_store_screlease_fn;
add_write_index_relaxed_fn = coreapi.hsa_queue_add_write_index_relaxed_fn;
load_read_index_relaxed_fn = coreapi.hsa_queue_load_read_index_relaxed_fn;
codeobj_reg = std::make_unique<code_object::CodeobjCallbackRegistry>(
[this](rocprofiler_agent_id_t agent, uint64_t codeobj_id, uint64_t addr, uint64_t size) {
if(agent == this->agent_id) this->load_codeobj(codeobj_id, addr, size);
@@ -199,14 +194,15 @@ ThreadTracerQueue::ThreadTracerQueue(thread_trace_parameter_pack _params,
ThreadTracerQueue::~ThreadTracerQueue()
{
ROCP_TRACE << "Destroying ATT Queue...";
std::unique_lock<std::mutex> lk(trace_resources_mut);
if(active_traces.load() < 1)
{
if(queue_destroy_fn) queue_destroy_fn(this->queue);
hsa::get_core_table()->hsa_queue_destroy_fn(this->queue);
return;
}
ROCP_WARNING << "Thread tracer being destroyed with thread trace active";
ROCP_CI_LOG(WARNING) << "Thread tracer being destroyed with thread trace active";
control_packet->clear();
control_packet->populate_after();
@@ -253,7 +249,10 @@ ThreadTracerQueue::iterate_data(aqlprofile_handle_t handle, rocprofiler_user_dat
cb_dt.userdata = &data;
auto status = aqlprofile_att_iterate_data(handle, thread_trace_callback, &cb_dt);
CHECK_HSA(status, "Failed to iterate ATT data");
if(status == HSA_STATUS_ERROR_OUT_OF_RESOURCES)
ROCP_WARNING << "Thread trace buffer full!";
else
CHECK_HSA(status, "Failed to iterate ATT data");
active_traces.fetch_sub(1);
}
@@ -284,42 +283,36 @@ ThreadTracerQueue::unload_codeobj(code_object_id_t id)
}
void
DispatchThreadTracer::resource_init(const hsa::AgentCache& cache,
const CoreApiTable& coreapi,
const AmdExtTable& ext)
DispatchThreadTracer::resource_init()
{
auto agent = cache.get_hsa_agent();
std::unique_lock<std::shared_mutex> lk(agents_map_mut);
auto rocp_agents = rocprofiler::agent::get_agents();
auto agent_it = agents.find(agent);
if(agent_it != agents.end())
auto lk = std::unique_lock{agents_map_mut};
for(const auto* rocp_agent : rocp_agents)
{
agent_it->second->active_queues.fetch_add(1);
return;
}
auto it = params.find(rocp_agent->id);
if(it == params.end()) continue;
auto new_tracer = std::make_unique<ThreadTracerQueue>(this->params, cache, coreapi, ext);
agents.emplace(agent, std::move(new_tracer));
auto cache = rocprofiler::agent::get_hsa_agent(rocp_agent);
CHECK(cache.has_value());
agents[*cache] = std::make_unique<ThreadTracerQueue>(it->second, rocp_agent->id);
}
}
void
DispatchThreadTracer::resource_deinit(const hsa::AgentCache& cache)
DispatchThreadTracer::resource_deinit()
{
std::unique_lock<std::shared_mutex> lk(agents_map_mut);
auto agent_it = agents.find(cache.get_hsa_agent());
if(agent_it == agents.end()) return;
if(agent_it->second->active_queues.fetch_sub(1) > 1) return;
agents.erase(cache.get_hsa_agent());
ROCP_TRACE << "Clearing agents";
auto lk = std::unique_lock{agents_map_mut};
agents.clear();
}
/**
* Callback we get from HSA interceptor when a kernel packet is being enqueued.
* We return an AQLPacket containing the start/stop/read packets for injection.
*/
std::unique_ptr<hsa::AQLPacket>
hsa::Queue::pkt_and_serialize_t
DispatchThreadTracer::pre_kernel_call(const hsa::Queue& queue,
rocprofiler_kernel_id_t kernel_id,
rocprofiler_dispatch_id_t dispatch_id,
@@ -332,73 +325,36 @@ DispatchThreadTracer::pre_kernel_call(const hsa::Queue& queue,
if(corr_id) rocprof_corr_id.internal = corr_id->internal;
// TODO: Get external
// Maybe adds serialization packets to the AQLPacket (if serializer is enabled)
// and maybe adds barrier packets if the state is transitioning from serialized <->
// unserialized
auto maybe_add_serialization = [&](auto& gen_pkt) {
CHECK_NOTNULL(hsa::get_queue_controller())
->serializer(&queue)
.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);
});
};
auto control_flags = params.dispatch_cb_fn(queue.get_agent().get_rocp_agent()->id,
queue.get_id(),
rocprof_corr_id,
kernel_id,
dispatch_id,
params.callback_userdata.ptr,
user_data);
if(control_flags == ROCPROFILER_ATT_CONTROL_NONE)
{
auto empty = std::make_unique<hsa::EmptyAQLPacket>();
if(params.bSerialize) maybe_add_serialization(empty);
return empty;
}
std::shared_lock<std::shared_mutex> lk(agents_map_mut);
auto it = agents.find(queue.get_agent().get_hsa_agent());
assert(it != agents.end() && it->second != nullptr);
auto packet = it->second->get_control(true);
if(it == agents.end()) return {nullptr, false};
auto& agent = *CHECK_NOTNULL(it->second);
const auto& parameters = agent.params;
auto control_flags = parameters.dispatch_cb_fn(queue.get_agent().get_rocp_agent()->id,
queue.get_id(),
rocprof_corr_id,
kernel_id,
dispatch_id,
parameters.callback_userdata.ptr,
user_data);
if(control_flags == ROCPROFILER_ATT_CONTROL_NONE) return {nullptr, parameters.bSerialize};
auto packet = agent.get_control(true);
post_move_data.fetch_add(1);
maybe_add_serialization(packet);
packet->populate_before();
packet->populate_after();
return packet;
return {std::move(packet), true};
}
class SignalSerializerExit
{
public:
SignalSerializerExit(const hsa::Queue::queue_info_session_t& _session)
: session(_session)
{}
~SignalSerializerExit()
{
auto* controller = hsa::get_queue_controller();
if(!controller) return;
controller->serializer(&session.queue).wlock([&](auto& serializer) {
serializer.kernel_completion_signal(session.queue);
});
}
const hsa::Queue::queue_info_session_t& session;
};
void
DispatchThreadTracer::post_kernel_call(DispatchThreadTracer::inst_pkt_t& aql,
const hsa::Queue::queue_info_session_t& session)
{
std::unique_ptr<SignalSerializerExit> signal{nullptr};
if(params.bSerialize) signal = std::make_unique<SignalSerializerExit>(session);
if(post_move_data.load() < 1) return;
for(auto& aql_pkt : aql)
@@ -414,8 +370,6 @@ DispatchThreadTracer::post_kernel_call(DispatchThreadTracer::inst_pkt_t& a
auto it = agents.find(pkt->GetAgent());
if(it != agents.end() && it->second != nullptr)
it->second->iterate_data(pkt->GetHandle(), session.user_data);
if(!signal) std::make_unique<SignalSerializerExit>(session);
}
}
@@ -456,19 +410,22 @@ DispatchThreadTracer::start_context()
void
DispatchThreadTracer::stop_context() // NOLINT(readability-convert-member-functions-to-static)
{
CHECK_NOTNULL(hsa::get_queue_controller())->disable_serialization();
auto* controller = hsa::get_queue_controller();
if(!controller) return;
client.wlock([&](auto& client_id) {
if(!client_id) return;
// Remove our callbacks from HSA's queue controller
CHECK_NOTNULL(hsa::get_queue_controller())->remove_callback(*client_id);
controller->remove_callback(*client_id);
client_id = std::nullopt;
});
controller->disable_serialization();
}
void
AgentThreadTracer::resource_init(const CoreApiTable& coreapi, const AmdExtTable& ext)
DeviceThreadTracer::resource_init()
{
auto rocp_agents = rocprofiler::agent::get_agents();
@@ -476,41 +433,36 @@ AgentThreadTracer::resource_init(const CoreApiTable& coreapi, const AmdExtTable&
for(const auto* rocp_agent : rocp_agents)
{
auto id = rocp_agent->id;
const auto* cache = rocprofiler::agent::get_agent_cache(rocp_agent);
auto it = params.find(CHECK_NOTNULL(rocp_agent)->id);
if(it == params.end()) continue;
if(tracers.find(id) != tracers.end())
ROCP_WARNING << "Agent configured twice: " << id.handle;
else if(params.find(id) == params.end())
ROCP_INFO << "Skipping agent " << id.handle;
else if(cache == nullptr)
ROCP_WARNING << "Invalid agent id: " << id.handle;
else
tracers[id] = std::make_unique<ThreadTracerQueue>(params.at(id), *cache, coreapi, ext);
agents[it->first] = std::make_unique<ThreadTracerQueue>(it->second, rocp_agent->id);
}
}
void
AgentThreadTracer::resource_deinit()
DeviceThreadTracer::resource_deinit()
{
ROCP_TRACE << "Clearing agents";
std::unique_lock<std::mutex> lk(agent_mut);
tracers.clear();
agents.clear();
}
void
AgentThreadTracer::start_context()
DeviceThreadTracer::start_context()
{
ROCP_TRACE << "Start device context";
std::unique_lock<std::mutex> lk(agent_mut);
if(tracers.empty())
if(agents.empty())
{
ROCP_FATAL << "Thread trace context not present for agent!";
ROCP_WARNING << "Thread trace context not present for agent!";
return;
}
std::vector<std::unique_ptr<Signal>> wait_list{};
for(auto& [_, tracer] : tracers)
for(auto& [_, tracer] : agents)
{
auto packet = tracer->get_control(true);
packet->populate_before();
@@ -521,20 +473,20 @@ AgentThreadTracer::start_context()
}
void
AgentThreadTracer::stop_context()
DeviceThreadTracer::stop_context()
{
using wait_t = std::tuple<ThreadTracerQueue*, aqlprofile_handle_t, std::unique_ptr<Signal>>;
std::unique_lock<std::mutex> lk(agent_mut);
if(tracers.empty())
if(agents.empty())
{
ROCP_FATAL << "Thread trace context not present for agent!";
ROCP_WARNING << "Thread trace context not present for agent!";
return;
}
std::vector<wait_t> wait_list{};
for(auto& [_, tracer] : tracers)
for(auto& [_, tracer] : agents)
{
auto packet = tracer->get_control(false);
packet->populate_after();
@@ -553,25 +505,25 @@ AgentThreadTracer::stop_context()
void
initialize(HsaApiTable* table)
{
assert(table->core_ && table->amd_ext_);
get_core() = *table->core_;
get_ext() = *table->amd_ext_;
ROCP_FATAL_IF(!table->core_ || !table->amd_ext_);
code_object::initialize(table);
for(auto& ctx : context::get_registered_contexts())
{
if(ctx->agent_thread_trace)
ctx->agent_thread_trace->resource_init(*table->core_, *table->amd_ext_);
if(ctx->device_thread_trace) ctx->device_thread_trace->resource_init();
if(ctx->dispatch_thread_trace) ctx->dispatch_thread_trace->resource_init();
}
}
void
finalize()
{
ROCP_TRACE << "Finalize called";
for(auto& ctx : context::get_registered_contexts())
{
if(ctx->agent_thread_trace) ctx->agent_thread_trace->resource_deinit();
if(ctx->device_thread_trace) ctx->device_thread_trace->resource_deinit();
if(ctx->dispatch_thread_trace) ctx->dispatch_thread_trace->resource_deinit();
}
code_object::finalize();
source/lib/rocprofiler-sdk/thread_trace/att_core.hpp
+25 -31
Datei anzeigen
@@ -25,6 +25,7 @@
#include "lib/rocprofiler-sdk/context/correlation_id.hpp"
#include "lib/rocprofiler-sdk/hsa/agent_cache.hpp"
#include "lib/rocprofiler-sdk/hsa/aql_packet.hpp"
#include "lib/rocprofiler-sdk/hsa/queue.hpp"
#include "lib/rocprofiler-sdk/hsa/queue_info_session.hpp"
#include "lib/rocprofiler-sdk/thread_trace/code_object.hpp"
@@ -72,11 +73,9 @@ struct thread_trace_parameter_pack
// GFX9 Only
std::vector<std::pair<uint32_t, uint32_t>> perfcounters;
static constexpr size_t DEFAULT_SIMD = 0x7;
static constexpr size_t DEFAULT_PERFCOUNTER_SIMD_MASK = 0xF;
static constexpr size_t DEFAULT_SE_MASK = 0x21;
static constexpr size_t DEFAULT_BUFFER_SIZE = 0x8000000;
static constexpr size_t PERFCOUNTER_SIMD_MASK_SHIFT = 28;
static constexpr size_t DEFAULT_SIMD = 0xF;
static constexpr size_t DEFAULT_SE_MASK = 0x1;
static constexpr size_t DEFAULT_BUFFER_SIZE = 0x8000000;
bool are_params_valid() const;
};
@@ -86,10 +85,7 @@ class ThreadTracerQueue
using code_object_id_t = uint64_t;
public:
ThreadTracerQueue(thread_trace_parameter_pack _params,
const hsa::AgentCache&,
const CoreApiTable&,
const AmdExtTable&);
ThreadTracerQueue(thread_trace_parameter_pack _params, rocprofiler_agent_id_t);
virtual ~ThreadTracerQueue();
void load_codeobj(code_object_id_t id, uint64_t addr, uint64_t size);
@@ -98,17 +94,17 @@ public:
std::unique_ptr<hsa::TraceControlAQLPacket> get_control(bool bStart);
void iterate_data(aqlprofile_handle_t handle, rocprofiler_user_data_t data);
hsa_queue_t* queue = nullptr;
hsa_queue_t* queue{nullptr};
std::mutex trace_resources_mut;
thread_trace_parameter_pack params;
std::atomic<int> active_traces{0};
std::atomic<int> active_queues{1};
std::unique_ptr<hsa::TraceControlAQLPacket> control_packet;
std::unique_ptr<aql::ThreadTraceAQLPacketFactory> factory;
[[nodiscard]] std::unique_ptr<class Signal> Submit(hsa_ext_amd_aql_pm4_packet_t* packet,
bool bWait);
bool bWait) const;
template <typename VecType>
[[nodiscard]] std::unique_ptr<class Signal> SubmitAndSignalLast(VecType vec)
@@ -125,11 +121,6 @@ private:
std::unique_ptr<code_object::CodeobjCallbackRegistry> codeobj_reg{nullptr};
rocprofiler_agent_id_t agent_id;
decltype(hsa_queue_load_read_index_relaxed)* load_read_index_relaxed_fn{nullptr};
decltype(hsa_queue_add_write_index_relaxed)* add_write_index_relaxed_fn{nullptr};
decltype(hsa_signal_store_screlease)* signal_store_screlease_fn{nullptr};
decltype(hsa_queue_destroy)* queue_destroy_fn{nullptr};
};
class DispatchThreadTracer
@@ -139,17 +130,21 @@ class DispatchThreadTracer
using inst_pkt_t = common::container::small_vector<std::pair<AQLPacketPtr, int64_t>, 4>;
public:
DispatchThreadTracer(thread_trace_parameter_pack _params)
: params(std::move(_params))
{}
DispatchThreadTracer() = default;
~DispatchThreadTracer() = default;
void start_context();
void stop_context();
void resource_init(const hsa::AgentCache&, const CoreApiTable&, const AmdExtTable&);
void resource_deinit(const hsa::AgentCache&);
void resource_init();
void resource_deinit();
std::unique_ptr<hsa::AQLPacket> pre_kernel_call(const hsa::Queue& queue,
void add_agent(rocprofiler_agent_id_t agent, thread_trace_parameter_pack pack)
{
auto lk = std::unique_lock{agents_map_mut};
params[agent] = std::move(pack);
}
hsa::Queue::pkt_and_serialize_t pre_kernel_call(const hsa::Queue& queue,
uint64_t kernel_id,
rocprofiler_dispatch_id_t dispatch_id,
rocprofiler_user_data_t* user_data,
@@ -157,23 +152,22 @@ public:
void post_kernel_call(inst_pkt_t& aql, const hsa::queue_info_session& session);
std::unordered_map<hsa_agent_t, std::unique_ptr<ThreadTracerQueue>> agents{};
std::unordered_map<hsa_agent_t, std::unique_ptr<ThreadTracerQueue>> agents{};
std::unordered_map<rocprofiler_agent_id_t, thread_trace_parameter_pack> params{};
std::shared_mutex agents_map_mut{};
std::atomic<int> post_move_data{0};
thread_trace_parameter_pack params{};
};
class AgentThreadTracer
class DeviceThreadTracer
{
public:
AgentThreadTracer() = default;
~AgentThreadTracer() = default;
DeviceThreadTracer() = default;
~DeviceThreadTracer() = default;
void start_context();
void stop_context();
void resource_init(const CoreApiTable&, const AmdExtTable&);
void resource_init();
void resource_deinit();
void add_agent(rocprofiler_agent_id_t id, thread_trace_parameter_pack _params)
@@ -187,7 +181,7 @@ public:
return params.find(id) != params.end();
}
std::map<rocprofiler_agent_id_t, std::unique_ptr<ThreadTracerQueue>> tracers{};
std::map<rocprofiler_agent_id_t, std::unique_ptr<ThreadTracerQueue>> agents{};
std::map<rocprofiler_agent_id_t, thread_trace_parameter_pack> params{};
std::mutex agent_mut;
source/lib/rocprofiler-sdk/thread_trace/att_service.cpp
+15 -9
Datei anzeigen
@@ -31,25 +31,30 @@
#include "rocprofiler-sdk/amd_detail/thread_trace.h"
using DispatchThreadTracer = rocprofiler::thread_trace::DispatchThreadTracer;
using AgentThreadTracer = rocprofiler::thread_trace::AgentThreadTracer;
using DeviceThreadTracer = rocprofiler::thread_trace::DeviceThreadTracer;
extern "C" {
rocprofiler_status_t
rocprofiler_configure_dispatch_thread_trace_service(
rocprofiler_context_id_t context_id,
rocprofiler_agent_id_t agent_id,
rocprofiler_att_parameter_t* parameters,
size_t num_parameters,
rocprofiler_att_dispatch_callback_t dispatch_callback,
rocprofiler_att_shader_data_callback_t shader_callback,
void* callback_userdata)
{
ROCP_TRACE << "Configuring Dispatch ATT for agent " << agent_id.handle;
if(rocprofiler::registration::get_init_status() > -1)
return ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED;
auto* ctx = rocprofiler::context::get_mutable_registered_context(context_id);
if(!ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND;
if(ctx->dispatch_thread_trace) return ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED;
if(ctx->agent_thread_trace) return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID;
if(ctx->device_thread_trace) return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID;
if(!ctx->dispatch_thread_trace)
ctx->dispatch_thread_trace = std::make_unique<DispatchThreadTracer>();
auto pack = rocprofiler::thread_trace::thread_trace_parameter_pack{};
@@ -92,19 +97,21 @@ rocprofiler_configure_dispatch_thread_trace_service(
if(!pack.are_params_valid()) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT;
ctx->dispatch_thread_trace = std::make_unique<DispatchThreadTracer>(pack);
ctx->dispatch_thread_trace->add_agent(agent_id, pack);
return ROCPROFILER_STATUS_SUCCESS;
}
rocprofiler_status_t
rocprofiler_configure_agent_thread_trace_service(
rocprofiler_configure_device_thread_trace_service(
rocprofiler_context_id_t context_id,
rocprofiler_agent_id_t agent_id,
rocprofiler_att_parameter_t* parameters,
size_t num_parameters,
rocprofiler_agent_id_t agent,
rocprofiler_att_shader_data_callback_t shader_callback,
rocprofiler_user_data_t userdata)
{
ROCP_TRACE << "Configuring Device ATT for agent " << agent_id.handle;
if(rocprofiler::registration::get_init_status() > -1)
return ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED;
@@ -112,7 +119,7 @@ rocprofiler_configure_agent_thread_trace_service(
if(!ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND;
if(ctx->dispatch_thread_trace) return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID;
if(!ctx->agent_thread_trace) ctx->agent_thread_trace = std::make_unique<AgentThreadTracer>();
if(!ctx->device_thread_trace) ctx->device_thread_trace = std::make_unique<DeviceThreadTracer>();
auto pack = rocprofiler::thread_trace::thread_trace_parameter_pack{};
@@ -154,8 +161,7 @@ rocprofiler_configure_agent_thread_trace_service(
if(!pack.are_params_valid()) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT;
assert(ctx->agent_thread_trace);
ctx->agent_thread_trace->add_agent(agent, pack);
ctx->device_thread_trace->add_agent(agent_id, pack);
return ROCPROFILER_STATUS_SUCCESS;
}
}
source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp
+65 -53
Datei anzeigen
@@ -27,6 +27,7 @@
#include "lib/rocprofiler-sdk/counters/metrics.hpp"
#include "lib/rocprofiler-sdk/counters/tests/hsa_tables.hpp"
#include "lib/rocprofiler-sdk/hsa/agent_cache.hpp"
#include "lib/rocprofiler-sdk/hsa/hsa.hpp"
#include "lib/rocprofiler-sdk/hsa/queue.hpp"
#include "lib/rocprofiler-sdk/hsa/queue_controller.hpp"
#include "lib/rocprofiler-sdk/registration.hpp"
@@ -64,11 +65,13 @@ test_init()
HsaApiTable table;
table.amd_ext_ = &get_ext_table();
table.core_ = &get_api_table();
rocprofiler::hsa::copy_table(table.core_, 0);
rocprofiler::hsa::copy_table(table.amd_ext_, 0);
agent::construct_agent_cache(&table);
hsa::get_queue_controller()->init(get_api_table(), get_ext_table());
return true;
};
[[maybe_unused]] static bool run_ince = init();
[[maybe_unused]] static bool run_once = init();
}
} // namespace rocprofiler
@@ -85,6 +88,7 @@ TEST(thread_trace, resource_creation)
auto agents = hsa::get_queue_controller()->get_supported_agents();
ASSERT_GT(agents.size(), 0);
for(const auto& [_, agent] : agents)
{
auto params = thread_trace::thread_trace_parameter_pack{};
@@ -101,14 +105,12 @@ TEST(thread_trace, resource_creation)
{
thread_trace::thread_trace_parameter_pack params{};
thread_trace::DispatchThreadTracer tracer(std::move(params));
thread_trace::DispatchThreadTracer tracer{};
for(const auto& [_, agent] : agents)
{
// Init twice to simulate two queues
tracer.resource_init(agent, get_api_table(), get_ext_table());
tracer.resource_init(agent, get_api_table(), get_ext_table());
}
tracer.add_agent(agent.get_rocp_agent()->id, params);
tracer.resource_init();
for(auto& [_, agenttracer] : tracer.agents)
{
@@ -117,18 +119,13 @@ TEST(thread_trace, resource_creation)
agenttracer->unload_codeobj(1);
}
for(const auto& [_, agent] : agents)
{
// Deinit twice to remove both queues
tracer.resource_deinit(agent);
tracer.resource_deinit(agent);
}
tracer.resource_deinit();
}
hsa_shut_down();
}
TEST(thread_trace, configure_test)
{
ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS);
test_init();
registration::init_logging();
@@ -143,29 +140,35 @@ TEST(thread_trace, configure_test)
params.push_back({ROCPROFILER_ATT_PARAMETER_BUFFER_SIZE, {0x1000000}});
params.push_back({ROCPROFILER_ATT_PARAMETER_SIMD_SELECT, {0xF}});
rocprofiler_configure_dispatch_thread_trace_service(
ctx,
params.data(),
params.size(),
[](rocprofiler_agent_id_t,
rocprofiler_queue_id_t,
rocprofiler_correlation_id_t,
rocprofiler_kernel_id_t,
rocprofiler_dispatch_id_t,
void*,
rocprofiler_user_data_t*) { return ROCPROFILER_ATT_CONTROL_NONE; },
[](rocprofiler_agent_id_t, int64_t, void*, size_t, rocprofiler_user_data_t) {},
nullptr);
auto agents = hsa::get_queue_controller()->get_supported_agents();
ASSERT_GT(agents.size(), 0);
for(auto& [_, agent] : agents)
{
rocprofiler_configure_dispatch_thread_trace_service(
ctx,
agent.get_rocp_agent()->id,
params.data(),
params.size(),
[](rocprofiler_agent_id_t,
rocprofiler_queue_id_t,
rocprofiler_correlation_id_t,
rocprofiler_kernel_id_t,
rocprofiler_dispatch_id_t,
void*,
rocprofiler_user_data_t*) { return ROCPROFILER_ATT_CONTROL_NONE; },
[](rocprofiler_agent_id_t, int64_t, void*, size_t, rocprofiler_user_data_t) {},
nullptr);
}
ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS);
ROCPROFILER_CALL(rocprofiler_start_context(ctx), "context start failed");
ROCPROFILER_CALL(rocprofiler_stop_context(ctx), "context stop failed");
context::pop_client(1);
hsa_shut_down();
}
TEST(thread_trace, perfcounters_configure_test)
{
ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS);
test_init();
registration::init_logging();
@@ -194,36 +197,45 @@ TEST(thread_trace, perfcounters_configure_test)
expected.insert({std::atoi(metric.event().c_str()), simd_mask});
}
rocprofiler_configure_dispatch_thread_trace_service(
ctx,
params.data(),
params.size(),
[](rocprofiler_agent_id_t,
rocprofiler_queue_id_t,
rocprofiler_correlation_id_t,
rocprofiler_kernel_id_t,
rocprofiler_dispatch_id_t,
void*,
rocprofiler_user_data_t*) { return ROCPROFILER_ATT_CONTROL_NONE; },
[](rocprofiler_agent_id_t, int64_t, void*, size_t, rocprofiler_user_data_t) {},
nullptr);
auto agents = hsa::get_queue_controller()->get_supported_agents();
ASSERT_GT(agents.size(), 0);
for(auto& [_, agent] : agents)
{
rocprofiler_configure_dispatch_thread_trace_service(
ctx,
agent.get_rocp_agent()->id,
params.data(),
params.size(),
[](rocprofiler_agent_id_t,
rocprofiler_queue_id_t,
rocprofiler_correlation_id_t,
rocprofiler_kernel_id_t,
rocprofiler_dispatch_id_t,
void*,
rocprofiler_user_data_t*) { return ROCPROFILER_ATT_CONTROL_NONE; },
[](rocprofiler_agent_id_t, int64_t, void*, size_t, rocprofiler_user_data_t) {},
nullptr);
}
auto* context = rocprofiler::context::get_mutable_registered_context(ctx);
auto* tracer = context->dispatch_thread_trace.get();
ASSERT_NE(tracer, nullptr);
ASSERT_EQ(tracer->params.perfcounter_ctrl, 1);
ASSERT_EQ(tracer->params.perfcounters.size(), 3);
for(const auto& param : tracer->params.perfcounters)
EXPECT_TRUE(expected.find(param) != expected.end())
<< "valid AQLprofile mask not generated for perfcounters";
for(auto& [id, agent] : tracer->agents)
{
ASSERT_EQ(agent->params.perfcounter_ctrl, 1);
ASSERT_EQ(agent->params.perfcounters.size(), 3);
for(const auto& param : agent->params.perfcounters)
EXPECT_TRUE(expected.find(param) != expected.end())
<< "valid AQLprofile mask not generated for perfcounters";
}
context::pop_client(1);
hsa_shut_down();
}
TEST(thread_trace, perfcounters_aql_options_test)
{
hsa_init();
ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS);
test_init();
registration::init_logging();
@@ -243,12 +255,11 @@ TEST(thread_trace, perfcounters_aql_options_test)
_params.perfcounters.push_back({std::atoi(metric.event().c_str()), simd_mask});
_params.perfcounter_ctrl = 2;
auto new_tracer = std::make_unique<thread_trace::ThreadTracerQueue>(
_params, begin(agents)->second, get_api_table(), get_ext_table());
_params, begin(agents)->second.get_rocp_agent()->id);
ASSERT_EQ(new_tracer->factory->aql_params.size(),
sqtt_default_num_options + perf_counters.size());
context::pop_client(1);
hsa_shut_down();
}
rocprofiler_status_t
@@ -281,11 +292,11 @@ query_available_agents(rocprofiler_agent_version_t /* version */,
params.push_back(att_param);
}
rocprofiler_configure_agent_thread_trace_service(
rocprofiler_configure_device_thread_trace_service(
*reinterpret_cast<rocprofiler_context_id_t*>(ctx_ptr),
agent->id,
params.data(),
params.size(),
agent->id,
[](rocprofiler_agent_id_t, int64_t, void*, size_t, rocprofiler_user_data_t) {},
rocprofiler_user_data_t{});
}
@@ -294,6 +305,7 @@ query_available_agents(rocprofiler_agent_version_t /* version */,
TEST(thread_trace, agent_configure_test)
{
ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS);
test_init();
registration::init_logging();
tests/rocprofv3/CMakeLists.txt
+1 -1
Datei anzeigen
@@ -38,7 +38,7 @@ add_subdirectory(pc-sampling)
add_subdirectory(collection-period)
add_subdirectory(rocdecode-trace)
add_subdirectory(rocjpeg-trace)
if(TARGET att_decoder_testing)
if(TARGET att_decoder_testing1)
add_subdirectory(advanced-thread-trace)
endif()
add_subdirectory(hip-stream-display)
tests/rocprofv3/advanced-thread-trace/CMakeLists.txt
+26 -8
Datei anzeigen
@@ -18,12 +18,14 @@ find_package(rocprofiler-sdk REQUIRED)
# hsa multiqueue dependency test
message("Add rocprofv3-test-hsa-multiqueue-att-cmd-env-ld-lib-path-execute")
add_test(
NAME rocprofv3-test-hsa-multiqueue-att-cmd-env-ld-lib-path-execute
COMMAND
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --log-level env --advanced-thread-trace
1 --att-target-cu 1 --att-shader-engine-mask 0x11 --kernel-include-regex copyD
--att-buffer-size 0x6000000 --att-simd-select 0x3 --att-parse testing
--att-buffer-size 0x6000000 --att-simd-select 0x3 --att-parse testing1
--att-serialize-all 1 -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/cmd_input -o
out --output-format json ${PRELOAD_ARGS} --
$<TARGET_FILE:hsa_code_object_testapp>)
@@ -38,7 +40,7 @@ add_test(
COMMAND
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --log-level env --advanced-thread-trace
1 --att-target-cu 1 --att-shader-engine-mask 0x11 --kernel-include-regex copyD
--att-buffer-size 0x6000000 --att-simd-select 0x3 --att-parse testing
--att-buffer-size 0x6000000 --att-simd-select 0x3 --att-parse testing1
--att-serialize-all 1 -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/cmd_input -o
out --output-format json ${PRELOAD_ARGS} --
$<TARGET_FILE:hsa_code_object_testapp>)
@@ -96,8 +98,8 @@ set_tests_properties(
"rocprofv3-test-hsa-multiqueue-att-json-execute" FAIL_REGULAR_EXPRESSION
"AssertionError")
if(TARGET rocprofiler-sdk::att-decoder-testing AND TARGET
rocprofiler-sdk::att-decoder-summary)
if(TARGET rocprofiler-sdk::att-decoder-testing1 AND TARGET
rocprofiler-sdk::att-decoder-testing2)
set(MISSING_TEST_DECODER_LIBS OFF)
else()
set(MISSING_TEST_DECODER_LIBS ON)
@@ -144,7 +146,7 @@ add_test(
NAME rocprofv3-test-att-library-path-cmd-line
COMMAND
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --att --att-library-path
${CMAKE_LIBRARY_OUTPUT_DIRECTORY}/att --att-parse summary --log-level env --echo
${CMAKE_LIBRARY_OUTPUT_DIRECTORY}/att --att-parse testing2 --log-level env --echo
-- sleep 0)
set_tests_properties(
@@ -156,7 +158,7 @@ add_test(
NAME rocprofv3-test-att-library-path-cmd-line-will-fail
COMMAND
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --att --att-library-path
${CMAKE_LIBRARY_OUTPUT_DIRECTORY}/att --att-parse testing --log-level env --echo
${CMAKE_LIBRARY_OUTPUT_DIRECTORY}/att --att-parse testing1 --log-level env --echo
-- sleep 0)
set_tests_properties(
@@ -168,7 +170,7 @@ set_tests_properties(
# Uses ROCPROF_ATT_LIBRARY_PATH to specify ATT library path
#
add_test(NAME rocprofv3-test-att-library-path-env-var
COMMAND $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --att --att-parse summary
COMMAND $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --att --att-parse testing2
--log-level env --echo -- sleep 0)
set_tests_properties(
@@ -178,7 +180,7 @@ set_tests_properties(
"${MISSING_TEST_DECODER_LIBS}")
add_test(NAME rocprofv3-test-att-library-path-env-var-will-fail
COMMAND $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --att --att-parse testing
COMMAND $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --att --att-parse testing1
--log-level env --echo -- sleep 0)
set_tests_properties(
@@ -193,3 +195,19 @@ set_tests_properties(
ON
DISABLED
"${MISSING_TEST_DECODER_LIBS}")
#
# Uses ATT and Counter Collection at the same time
#
add_test(
NAME rocprofv3-test-hsa-multiqueue-att-plus-pmc-execute
COMMAND
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> --log-level env --pmc SQ_WAVES
--advanced-thread-trace --att-parse testing1 -d
${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace/cmd_input -o out --output-format json
${PRELOAD_ARGS} -- $<TARGET_FILE:vector-ops>)
set_tests_properties(
rocprofv3-test-hsa-multiqueue-att-plus-pmc-execute
PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT
LD_LIBRARY_PATH=${CMAKE_LIBRARY_OUTPUT_DIRECTORY}:$ENV{LD_LIBRARY_PATH})
tests/rocprofv3/advanced-thread-trace/att_input.json
+3 -2
Datei anzeigen
@@ -9,12 +9,13 @@
],
"truncate_kernels": true,
"advanced_thread_trace": true,
"att_parse": "testing",
"att_parse": "testing1",
"att_target_cu": 1,
"att_shader_engine_mask": "0x11",
"att_simd_select": "0x3",
"att_buffer_size": "0x6000000",
"att_perfcounters": "SQ_WAVES:0x1 SQ_INSTS_VALU:0x3 SQ_INSTS_SALU:0xF"
"att_perfcounters": "SQ_WAVES:0x1 SQ_INSTS_VALU:0x3 SQ_INSTS_SALU",
"att_perfcounter_ctrl": 3
}
]
}
tests/rocprofv3/advanced-thread-trace/att_input.yml.in
+6 -6
Datei anzeigen
@@ -1,30 +1,30 @@
jobs:
- advanced_thread_trace: True
att_parse: summary
att_parse: testing2
att_library_path:
- @LIBRARY_OUTPUT_DIR@/att
- advanced_thread_trace: True
att_parse: summary
att_parse: testing2
att_library_path:
- @LIBRARY_OUTPUT_DIR@/att
- @LIBRARY_OUTPUT_DIR@
- advanced_thread_trace: True
att_parse: testing
att_parse: testing1
att_library_path:
- @LIBRARY_OUTPUT_DIR@/att
- @LIBRARY_OUTPUT_DIR@
- advanced_thread_trace: True
att_parse: testing
att_parse: testing1
- advanced_thread_trace: True
att_parse: testing
att_parse: testing1
att_library_path:
- @LIBRARY_OUTPUT_DIR@
- advanced_thread_trace: True
att_parse: testing
att_parse: testing1
att_library_path:
- @LIBRARY_OUTPUT_DIR@
tests/rocprofv3/advanced-thread-trace/validate.py
+9 -8
Datei anzeigen
@@ -19,11 +19,14 @@ def test_code_object_memory(code_object_file_path, json_data, output_path):
data = json_data["rocprofiler-sdk-tool"]
tool_memory_load = data["strings"]["code_object_snapshot_filenames"]
gfx_pattern = "gfx[a-z0-9]+"
match = re.search(gfx_pattern, tool_memory_load[0])
match = re.search(gfx_pattern, tool_memory_load[1])
assert match != None
gpu_name = match.group(0)
tool_memory_load_1 = open(os.path.join(output_path, tool_memory_load[0]), "rb")
tool_memory_load_2 = open(os.path.join(output_path, tool_memory_load[1]), "rb")
read_bytes = lambda filename: open(os.path.join(output_path, filename), "rb").read()
# Loads all saved code objects
tool_memory = [read_bytes(saved) for saved in tool_memory_load[1:]]
found = False
for hsa_file in code_object_file_path["hsa_memory_load"]:
@@ -33,11 +36,9 @@ def test_code_object_memory(code_object_file_path, json_data, output_path):
if gpu == gpu_name:
found = True
hsa_memory_load = open(hsa_file, "rb")
hsa_memory_fs = hsa_memory_load.read()
tool_memory_fs_1 = tool_memory_load_1.read()
tool_memory_fs_2 = tool_memory_load_2.read()
assert hsa_memory_fs == tool_memory_fs_2 or hsa_memory_fs == tool_memory_fs_1
hsa_memory_bytes = open(hsa_file, "rb").read()
# Checks if hsa_file is one of the saved code objects
assert any([hsa_memory_bytes == fs for fs in tool_memory])
break
assert found == True
tests/thread-trace/agent.cpp
+6 -6
Datei anzeigen
@@ -123,12 +123,12 @@ query_available_agents(rocprofiler_agent_version_t /* version */,
parameters.push_back({ROCPROFILER_ATT_PARAMETER_SERIALIZE_ALL, 0});
ROCPROFILER_CALL(
rocprofiler_configure_agent_thread_trace_service(agent_ctx,
parameters.data(),
parameters.size(),
agent->id,
Callbacks::shader_data_callback,
user),
rocprofiler_configure_device_thread_trace_service(agent_ctx,
agent->id,
parameters.data(),
parameters.size(),
Callbacks::shader_data_callback,
user),
"thread trace service configure");
}
return ROCPROFILER_STATUS_SUCCESS;
tests/thread-trace/multi_dispatch.cpp
+29 -7
Datei anzeigen
@@ -73,14 +73,36 @@ tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data)
std::vector<rocprofiler_att_parameter_t> params{};
params.push_back({ROCPROFILER_ATT_PARAMETER_SERIALIZE_ALL, 1});
std::vector<rocprofiler_agent_id_t> agents{};
ROCPROFILER_CALL(
rocprofiler_configure_dispatch_thread_trace_service(client_ctx,
params.data(),
params.size(),
dispatch_callback,
Callbacks::shader_data_callback,
tool_data),
"thread trace service configure");
rocprofiler_query_available_agents(
ROCPROFILER_AGENT_INFO_VERSION_0,
[](rocprofiler_agent_version_t, const void** _agents, size_t _num_agents, void* _data) {
auto* agent_v = static_cast<std::vector<rocprofiler_agent_id_t>*>(_data);
for(size_t i = 0; i < _num_agents; ++i)
{
auto* agent = static_cast<const rocprofiler_agent_v0_t*>(_agents[i]);
if(agent->type == ROCPROFILER_AGENT_TYPE_GPU) agent_v->emplace_back(agent->id);
}
return ROCPROFILER_STATUS_SUCCESS;
},
sizeof(rocprofiler_agent_v0_t),
&agents),
"Failed to iterate agents");
for(auto id : agents)
{
ROCPROFILER_CALL(
rocprofiler_configure_dispatch_thread_trace_service(client_ctx,
id,
params.data(),
params.size(),
dispatch_callback,
Callbacks::shader_data_callback,
tool_data),
"thread trace service configure");
}
int valid_ctx = 0;
ROCPROFILER_CALL(rocprofiler_context_is_valid(client_ctx, &valid_ctx),
tests/thread-trace/single_dispatch.cpp
+29 -3
Datei anzeigen
@@ -83,10 +83,36 @@ tool_init(rocprofiler_client_finalize_t /* fini_func */, void* tool_data)
tool_data),
"code object tracing service configure");
std::vector<rocprofiler_agent_id_t> agents{};
ROCPROFILER_CALL(
rocprofiler_configure_dispatch_thread_trace_service(
client_ctx, nullptr, 0, dispatch_callback, Callbacks::shader_data_callback, tool_data),
"thread trace service configure");
rocprofiler_query_available_agents(
ROCPROFILER_AGENT_INFO_VERSION_0,
[](rocprofiler_agent_version_t, const void** _agents, size_t _num_agents, void* _data) {
auto* agent_v = static_cast<std::vector<rocprofiler_agent_id_t>*>(_data);
for(size_t i = 0; i < _num_agents; ++i)
{
auto* agent = static_cast<const rocprofiler_agent_v0_t*>(_agents[i]);
if(agent->type == ROCPROFILER_AGENT_TYPE_GPU) agent_v->emplace_back(agent->id);
}
return ROCPROFILER_STATUS_SUCCESS;
},
sizeof(rocprofiler_agent_v0_t),
&agents),
"Failed to iterate agents");
for(auto id : agents)
{
ROCPROFILER_CALL(
rocprofiler_configure_dispatch_thread_trace_service(client_ctx,
id,
nullptr,
0,
dispatch_callback,
Callbacks::shader_data_callback,
tool_data),
"thread trace service configure");
}
int valid_ctx = 0;
ROCPROFILER_CALL(rocprofiler_context_is_valid(client_ctx, &valid_ctx),