From 4f2dc896d3a2a9ad3bbd31c8afad0c390ba066ab Mon Sep 17 00:00:00 2001 From: "Jonathan R. Madsen" Date: Thu, 2 Nov 2023 19:10:10 -0500 Subject: [PATCH] Support Tool Intercept API Tables (#165) * Update include/rocprofiler - intercept_table.h header - generic rocprofiler_runtime_library_t - rocprofiler_internal_thread_library_t is not typedef for rocprofiler_runtime_library_t - rocprofiler_at_runtime_api_registration * Update lib/rocprofiler - minor tweaks to context.cpp - check if none context early - disallow stop_context when finalizing - add intercept_table.hpp and intercept_table.cpp - implement rocprofiler_at_runtime_api_registration - implement notify_runtime_api_registration - update registration.cpp - invoke notify_runtime_api_registration - tweak to fini status when invoking client finalizer * Update lib/rocprofiler/tests - add tests for intercept table * Add samples/intercept_table - demonstrate how to install custom API function wrappers instead of relying on HSA callback tracing * Update lib/rocprofiler/tests/intercept_table.cpp - remove _SERVICE from ROCPROFILER_SERVICE_ * Update include/rocprofiler/intercept_table.h - Update doxygen comments * Update lib/rocprofiler/intercept_table.cpp - return error config locked if already initialized * Update lib/rocprofiler/intercept_table.cpp - remove unnecessary alias * Apply suggestions from code review Co-authored-by: Tony Tye * Update doxygen comments - clarify when rocprofiler_at_runtime_api_registration can be invoked * Use rocprofiler_runtime_library_t for intercept table and internal threading - remove rocprofiler_intercept_library_t alias to rocprofiler_runtime_library_t - remove rocprofiler_internal_thread_library_t alias to rocprofiler_runtime_library_t - move around documentation with regard to rocprofiler_runtime_library_t enumeration - added some extra doxygen documentation to internal threading functions --------- Co-authored-by: Tony Tye --- samples/CMakeLists.txt | 1 + samples/api_buffered_tracing/client.cpp | 4 +- samples/intercept_table/CMakeLists.txt | 59 ++ samples/intercept_table/client.cpp | 305 +++++++++ samples/intercept_table/client.hpp | 38 ++ samples/intercept_table/main.cpp | 235 +++++++ source/include/rocprofiler/CMakeLists.txt | 1 + source/include/rocprofiler/fwd.h | 14 + source/include/rocprofiler/intercept_table.h | 223 +++++++ .../include/rocprofiler/internal_threading.h | 58 +- source/include/rocprofiler/rocprofiler.h | 7 +- source/lib/rocprofiler/CMakeLists.txt | 3 +- source/lib/rocprofiler/context.cpp | 26 +- source/lib/rocprofiler/intercept_table.cpp | 184 +++++ source/lib/rocprofiler/intercept_table.hpp | 40 ++ source/lib/rocprofiler/internal_threading.cpp | 22 +- source/lib/rocprofiler/internal_threading.hpp | 4 +- source/lib/rocprofiler/registration.cpp | 9 + source/lib/rocprofiler/tests/CMakeLists.txt | 3 +- .../tests/external_correlation.cpp | 4 +- .../lib/rocprofiler/tests/intercept_table.cpp | 631 ++++++++++++++++++ source/lib/rocprofiler/tests/registration.cpp | 4 +- 22 files changed, 1823 insertions(+), 52 deletions(-) create mode 100644 samples/intercept_table/CMakeLists.txt create mode 100644 samples/intercept_table/client.cpp create mode 100644 samples/intercept_table/client.hpp create mode 100644 samples/intercept_table/main.cpp create mode 100644 source/include/rocprofiler/intercept_table.h create mode 100644 source/lib/rocprofiler/intercept_table.cpp create mode 100644 source/lib/rocprofiler/intercept_table.hpp create mode 100644 source/lib/rocprofiler/tests/intercept_table.cpp diff --git a/samples/CMakeLists.txt b/samples/CMakeLists.txt index 513bea9db4..4884d85270 100644 --- a/samples/CMakeLists.txt +++ b/samples/CMakeLists.txt @@ -19,3 +19,4 @@ add_subdirectory(pc_sampling) add_subdirectory(api_callback_tracing) add_subdirectory(api_buffered_tracing) add_subdirectory(counter_collection) +add_subdirectory(intercept_table) diff --git a/samples/api_buffered_tracing/client.cpp b/samples/api_buffered_tracing/client.cpp index 58a721edd0..64f8fc1007 100644 --- a/samples/api_buffered_tracing/client.cpp +++ b/samples/api_buffered_tracing/client.cpp @@ -245,7 +245,7 @@ tool_tracing_callback(rocprofiler_context_id_t context, } void -thread_precreate(rocprofiler_internal_thread_library_t lib, void* tool_data) +thread_precreate(rocprofiler_runtime_library_t lib, void* tool_data) { static_cast(tool_data)->emplace_back( source_location{__FUNCTION__, @@ -256,7 +256,7 @@ thread_precreate(rocprofiler_internal_thread_library_t lib, void* tool_data) } void -thread_postcreate(rocprofiler_internal_thread_library_t lib, void* tool_data) +thread_postcreate(rocprofiler_runtime_library_t lib, void* tool_data) { static_cast(tool_data)->emplace_back( source_location{__FUNCTION__, diff --git a/samples/intercept_table/CMakeLists.txt b/samples/intercept_table/CMakeLists.txt new file mode 100644 index 0000000000..a365ea5d8e --- /dev/null +++ b/samples/intercept_table/CMakeLists.txt @@ -0,0 +1,59 @@ +# +# +# +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-samples-intercept-table 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() + +if(NOT TARGET rocprofiler::rocprofiler) + find_package(rocprofiler REQUIRED) +endif() + +add_library(intercept-table-client SHARED) +target_sources(intercept-table-client PRIVATE client.cpp client.hpp) +target_link_libraries( + intercept-table-client + PRIVATE rocprofiler::rocprofiler + $) + +set_source_files_properties(main.cpp PROPERTIES LANGUAGE HIP) +find_package(Threads REQUIRED) + +add_executable(intercept-table) +target_sources(intercept-table PRIVATE main.cpp) +target_link_libraries( + intercept-table PRIVATE intercept-table-client Threads::Threads + $) + +add_test(NAME intercept-table COMMAND $) + +set_tests_properties( + intercept-table + PROPERTIES + TIMEOUT + 45 + LABELS + "samples" + ENVIRONMENT + "${ROCPROFILER_MEMCHECK_PRELOAD_ENV};HSA_TOOLS_LIB=$" + ) diff --git a/samples/intercept_table/client.cpp b/samples/intercept_table/client.cpp new file mode 100644 index 0000000000..96f317b843 --- /dev/null +++ b/samples/intercept_table/client.cpp @@ -0,0 +1,305 @@ +// MIT License +// +// Copyright (c) 2023 ROCm Developer Tools +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. +// +// undefine NDEBUG so asserts are implemented +#ifdef NDEBUG +# undef NDEBUG +#endif + +/** + * @file samples/intercept_table/client.cpp + * + * @brief Example rocprofiler client (tool) + */ + +#include "client.hpp" + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#define ROCPROFILER_CALL(result, msg) \ + { \ + rocprofiler_status_t CHECKSTATUS = result; \ + if(CHECKSTATUS != ROCPROFILER_STATUS_SUCCESS) \ + { \ + std::cerr << #result << " failed with error code " << CHECKSTATUS << std::endl; \ + throw std::runtime_error(#result " failure"); \ + } \ + } + +namespace client +{ +namespace +{ +struct source_location +{ + std::string function = {}; + std::string file = {}; + uint32_t line = 0; + std::string context = {}; +}; + +using call_stack_t = std::vector; +using callback_kind_names_t = std::map; +using callback_kind_operation_names_t = + std::map>; +using wrap_count_t = std::pair; + +rocprofiler_client_id_t* client_id = nullptr; +std::map client_wrap_data = {}; + +void +print_call_stack(const call_stack_t& _call_stack) +{ + namespace fs = ::std::filesystem; + + auto ofname = std::string{"intercept_table.log"}; + if(auto* eofname = getenv("ROCPROFILER_SAMPLE_OUTPUT_FILE")) ofname = eofname; + + std::ostream* ofs = nullptr; + auto cleanup = std::function{}; + + if(ofname == "stdout") + ofs = &std::cout; + else if(ofname == "stderr") + ofs = &std::cerr; + else + { + ofs = new std::ofstream{ofname}; + if(ofs && *ofs) + cleanup = [](std::ostream*& _os) { delete _os; }; + else + { + std::cerr << "Error outputting to " << ofname << ". Redirecting to stderr...\n"; + ofname = "stderr"; + ofs = &std::cerr; + } + } + + std::cout << "Outputting collected data to " << ofname << "...\n" << std::flush; + + size_t n = 0; + *ofs << std::left; + for(const auto& itr : _call_stack) + { + *ofs << std::setw(2) << ++n << "/" << std::setw(2) << _call_stack.size() << " "; + *ofs << "[" << fs::path{itr.file}.filename() << ":" << itr.line << "] " << std::setw(20) + << itr.function; + if(!itr.context.empty()) *ofs << " :: " << itr.context; + *ofs << "\n"; + } + + *ofs << std::flush; + + if(cleanup) cleanup(ofs); +} + +void +tool_fini(void* tool_data) +{ + assert(tool_data != nullptr); + + auto* _call_stack = static_cast(tool_data); + + size_t wrapped_count = 0; + for(const auto& itr : client_wrap_data) + { + auto src_loc = itr.second.first; + src_loc.context += "call_count=" + std::to_string(itr.second.second); + _call_stack->emplace_back(std::move(src_loc)); + wrapped_count += itr.second.second; + } + + _call_stack->emplace_back(source_location{__FUNCTION__, __FILE__, __LINE__, ""}); + + print_call_stack(*_call_stack); + + delete _call_stack; + + if(wrapped_count == 0) + { + throw std::runtime_error{"intercept_table sample did not wrap HSA API table"}; + } +} + +template +auto +generate_wrapper(const char* name, RetT (*func)(Args...)) +{ + using functor_type = RetT (*)(Args...); + + client_wrap_data.emplace(Idx, wrap_count_t{source_location{name, __FILE__, __LINE__, ""}, 0}); + + static functor_type underlying_func = func; + static functor_type wrapped_func = [](Args... args) -> RetT { + client_wrap_data.at(Idx).second += 1; + if(underlying_func) return underlying_func(args...); + if constexpr(!std::is_void::value) return RetT{}; + }; + + return wrapped_func; +} + +#define GENERATE_WRAPPER(TABLE, FUNC) \ + TABLE->FUNC##_fn = generate_wrapper<__COUNTER__>(#FUNC, TABLE->FUNC##_fn) + +void +api_registration_callback(rocprofiler_runtime_library_t type, + uint64_t lib_version, + uint64_t lib_instance, + void** tables, + uint64_t num_tables, + void* user_data) +{ + if(type != ROCPROFILER_HSA_LIBRARY) + throw std::runtime_error{"unexpected library type: " + + std::to_string(static_cast(type))}; + if(lib_instance != 0) throw std::runtime_error{"multiple instances of HSA runtime library"}; + if(num_tables != 1) throw std::runtime_error{"expected only one table of type HsaApiTable"}; + + auto* call_stack = static_cast*>(user_data); + + uint32_t major = lib_version / 10000; + uint32_t minor = (lib_version % 10000) / 100; + uint32_t patch = lib_version % 100; + + auto info = std::stringstream{}; + info << client_id->name << " is using HSA runtime v" << major << "." << minor << "." << patch; + + std::clog << info.str() << "\n" << std::flush; + + call_stack->emplace_back(client::source_location{__FUNCTION__, __FILE__, __LINE__, info.str()}); + + auto* hsa_api_table = static_cast(tables[0]); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_agent_get_info); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_agent_iterate_isas); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_code_object_reader_create_from_memory); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_executable_create_alt); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_executable_freeze); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_executable_get_symbol_by_name); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_executable_iterate_symbols); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_executable_load_agent_code_object); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_executable_symbol_get_info); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_isa_get_info_alt); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_iterate_agents); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_queue_add_write_index_screlease); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_queue_create); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_queue_load_read_index_relaxed); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_queue_load_read_index_scacquire); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_signal_create); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_signal_destroy); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_signal_load_relaxed); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_signal_silent_store_relaxed); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_signal_store_screlease); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_signal_wait_scacquire); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_system_get_info); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_system_get_major_extension_table); +} +} // namespace + +void +setup() +{} + +void +shutdown() +{} +} // namespace client + +extern "C" rocprofiler_tool_configure_result_t* +rocprofiler_configure(uint32_t version, + const char* runtime_version, + uint32_t priority, + rocprofiler_client_id_t* id) +{ + // only activate if main tool + if(priority > 0) return nullptr; + + // set the client name + id->name = "ExampleTool"; + + // store client info + client::client_id = id; + + // compute major/minor/patch version info + uint32_t major = version / 10000; + uint32_t minor = (version % 10000) / 100; + uint32_t patch = version % 100; + + // generate info string + auto info = std::stringstream{}; + info << id->name << " is using rocprofiler v" << major << "." << minor << "." << patch << " (" + << runtime_version << ")"; + + std::clog << info.str() << std::endl; + + // demonstration of alternative way to get the version info + { + auto version_info = std::array{}; + ROCPROFILER_CALL( + rocprofiler_get_version(&version_info.at(0), &version_info.at(1), &version_info.at(2)), + "failed to get version info"); + + if(std::array{major, minor, patch} != version_info) + { + throw std::runtime_error{"version info mismatch"}; + } + } + + // data passed around all the callbacks + auto* client_tool_data = new std::vector{}; + + // add first entry + client_tool_data->emplace_back( + client::source_location{__FUNCTION__, __FILE__, __LINE__, info.str()}); + + rocprofiler_at_runtime_api_registration(client::api_registration_callback, + ROCPROFILER_HSA_LIBRARY, + static_cast(client_tool_data)); + + // create configure data + static auto cfg = + rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), + nullptr, + &client::tool_fini, + static_cast(client_tool_data)}; + + // return pointer to configure data + return &cfg; +} diff --git a/samples/intercept_table/client.hpp b/samples/intercept_table/client.hpp new file mode 100644 index 0000000000..976f5afccc --- /dev/null +++ b/samples/intercept_table/client.hpp @@ -0,0 +1,38 @@ +// MIT License +// +// Copyright (c) 2023 ROCm Developer Tools +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#ifdef intercept_table_client_EXPORTS +# define CLIENT_API __attribute__((visibility("default"))) +#else +# define CLIENT_API +#endif + +namespace client +{ +void +setup() CLIENT_API; + +void +shutdown() CLIENT_API; +} // namespace client diff --git a/samples/intercept_table/main.cpp b/samples/intercept_table/main.cpp new file mode 100644 index 0000000000..8f63b3be27 --- /dev/null +++ b/samples/intercept_table/main.cpp @@ -0,0 +1,235 @@ +// MIT License +// +// Copyright (c) 2023 ROCm Developer Tools +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include "hip/hip_runtime.h" + +#include +#include +#include +#include +#include +#include +#include + +#define HIP_API_CALL(CALL) \ + { \ + hipError_t error_ = (CALL); \ + if(error_ != hipSuccess) \ + { \ + auto _hip_api_print_lk = auto_lock_t{print_lock}; \ + fprintf(stderr, \ + "%s:%d :: HIP error : %s\n", \ + __FILE__, \ + __LINE__, \ + hipGetErrorString(error_)); \ + throw std::runtime_error("hip_api_call"); \ + } \ + } + +namespace +{ +using auto_lock_t = std::unique_lock; +auto print_lock = std::mutex{}; +size_t nthreads = 2; +size_t nitr = 500; +size_t nsync = 10; +constexpr unsigned shared_mem_tile_dim = 32; + +void +check_hip_error(void); + +void +verify(int* in, int* out, int M, int N); +} // namespace + +__global__ void +transpose_a(int* in, int* out, int M, int N); + +void +run(int rank, int tid, hipStream_t stream, int argc, char** argv); + +int +main(int argc, char** argv) +{ + int rank = 0; + for(int i = 1; i < argc; ++i) + { + auto _arg = std::string{argv[i]}; + if(_arg == "?" || _arg == "-h" || _arg == "--help") + { + fprintf(stderr, + "usage: transpose [NUM_THREADS (%zu)] [NUM_ITERATION (%zu)] " + "[SYNC_EVERY_N_ITERATIONS (%zu)]\n", + nthreads, + nitr, + nsync); + exit(EXIT_SUCCESS); + } + } + if(argc > 1) nthreads = atoll(argv[1]); + if(argc > 2) nitr = atoll(argv[2]); + if(argc > 3) nsync = atoll(argv[3]); + + printf("[transpose] Number of threads: %zu\n", nthreads); + printf("[transpose] Number of iterations: %zu\n", nitr); + printf("[transpose] Syncing every %zu iterations\n", nsync); + + // this is a temporary workaround in omnitrace when HIP + MPI is enabled + int ndevice = 0; + int devid = rank; + HIP_API_CALL(hipGetDeviceCount(&ndevice)); + printf("[transpose] Number of devices found: %i\n", ndevice); + if(ndevice > 0) + { + devid = rank % ndevice; + HIP_API_CALL(hipSetDevice(devid)); + printf("[transpose] Rank %i assigned to device %i\n", rank, devid); + } + if(rank == devid && rank < ndevice) + { + std::vector _threads{}; + std::vector _streams(nthreads); + for(size_t i = 0; i < nthreads; ++i) + HIP_API_CALL(hipStreamCreate(&_streams.at(i))); + for(size_t i = 1; i < nthreads; ++i) + _threads.emplace_back(run, rank, i, _streams.at(i), argc, argv); + run(rank, 0, _streams.at(0), argc, argv); + for(auto& itr : _threads) + itr.join(); + for(size_t i = 0; i < nthreads; ++i) + HIP_API_CALL(hipStreamDestroy(_streams.at(i))); + } + HIP_API_CALL(hipDeviceSynchronize()); + HIP_API_CALL(hipDeviceReset()); + + return 0; +} + +__global__ void +transpose_a(int* in, int* out, int M, int N) +{ + __shared__ int tile[shared_mem_tile_dim][shared_mem_tile_dim]; + + int idx = (blockIdx.y * blockDim.y + threadIdx.y) * M + blockIdx.x * blockDim.x + threadIdx.x; + tile[threadIdx.y][threadIdx.x] = in[idx]; + __syncthreads(); + idx = (blockIdx.x * blockDim.x + threadIdx.y) * N + blockIdx.y * blockDim.y + threadIdx.x; + out[idx] = tile[threadIdx.x][threadIdx.y]; +} + +void +run(int rank, int tid, hipStream_t stream, int argc, char** argv) +{ + unsigned int M = 4960 * 2; + unsigned int N = 4960 * 2; + if(argc > 2) nitr = atoll(argv[2]); + if(argc > 3) nsync = atoll(argv[3]); + + auto_lock_t _lk{print_lock}; + std::cout << "[" << rank << "][" << tid << "] M: " << M << " N: " << N << std::endl; + _lk.unlock(); + + std::default_random_engine _engine{std::random_device{}() * (rank + 1) * (tid + 1)}; + std::uniform_int_distribution _dist{0, 1000}; + + size_t size = sizeof(int) * M * N; + int* inp_matrix = new int[size]; + int* out_matrix = new int[size]; + for(size_t i = 0; i < M * N; i++) + { + inp_matrix[i] = _dist(_engine); + out_matrix[i] = 0; + } + int* in = nullptr; + int* out = nullptr; + + HIP_API_CALL(hipMalloc(&in, size)); + HIP_API_CALL(hipMalloc(&out, size)); + HIP_API_CALL(hipMemsetAsync(in, 0, size, stream)); + HIP_API_CALL(hipMemsetAsync(out, 0, size, stream)); + HIP_API_CALL(hipMemcpyAsync(in, inp_matrix, size, hipMemcpyHostToDevice, stream)); + HIP_API_CALL(hipStreamSynchronize(stream)); + + dim3 grid(M / 32, N / 32, 1); + dim3 block(32, 32, 1); // transpose_a + + auto t1 = std::chrono::high_resolution_clock::now(); + for(size_t i = 0; i < nitr; ++i) + { + transpose_a<<>>(in, out, M, N); + check_hip_error(); + if(i % nsync == (nsync - 1)) HIP_API_CALL(hipStreamSynchronize(stream)); + } + auto t2 = std::chrono::high_resolution_clock::now(); + HIP_API_CALL(hipStreamSynchronize(stream)); + HIP_API_CALL(hipMemcpyAsync(out_matrix, out, size, hipMemcpyDeviceToHost, stream)); + double time = std::chrono::duration_cast>(t2 - t1).count(); + float GB = (float) size * nitr * 2 / (1 << 30); + + print_lock.lock(); + std::cout << "[" << rank << "][" << tid << "] Runtime of transpose is " << time << " sec\n" + << "The average performance of transpose is " << GB / time << " GBytes/sec" + << std::endl; + print_lock.unlock(); + + HIP_API_CALL(hipStreamSynchronize(stream)); + + // cpu_transpose(matrix, out_matrix, M, N); + verify(inp_matrix, out_matrix, M, N); + + HIP_API_CALL(hipFree(in)); + HIP_API_CALL(hipFree(out)); + + delete[] inp_matrix; + delete[] out_matrix; +} + +namespace +{ +void +check_hip_error(void) +{ + hipError_t err = hipGetLastError(); + if(err != hipSuccess) + { + auto_lock_t _lk{print_lock}; + std::cerr << "Error: " << hipGetErrorString(err) << std::endl; + throw std::runtime_error("hip_api_call"); + } +} + +void +verify(int* in, int* out, int M, int N) +{ + for(int i = 0; i < 10; i++) + { + int row = rand() % M; + int col = rand() % N; + if(in[row * N + col] != out[col * M + row]) + { + auto_lock_t _lk{print_lock}; + std::cout << "mismatch: " << row << ", " << col << " : " << in[row * N + col] << " | " + << out[col * M + row] << "\n"; + } + } +} +} // namespace diff --git a/source/include/rocprofiler/CMakeLists.txt b/source/include/rocprofiler/CMakeLists.txt index 48c343d1e0..5bceeb96ab 100644 --- a/source/include/rocprofiler/CMakeLists.txt +++ b/source/include/rocprofiler/CMakeLists.txt @@ -24,6 +24,7 @@ set(ROCPROFILER_HEADER_FILES fwd.h hip.h hsa.h + intercept_table.h internal_threading.h marker.h pc_sampling.h diff --git a/source/include/rocprofiler/fwd.h b/source/include/rocprofiler/fwd.h index 2be3c03b6e..e770397a79 100644 --- a/source/include/rocprofiler/fwd.h +++ b/source/include/rocprofiler/fwd.h @@ -209,6 +209,20 @@ typedef enum // NOLINT(performance-enum-size) ROCPROFILER_BUFFER_POLICY_LAST, } rocprofiler_buffer_policy_t; +/** + * @brief Enumeration for specifying runtime libraries supported by rocprofiler. This enumeration is + * used for intercept tables and thread creation callbacks. @see INTERCEPT_TABLE and @see + * INTERNAL_THREADING. + */ +typedef enum +{ + ROCPROFILER_LIBRARY = (1 << 0), + ROCPROFILER_HSA_LIBRARY = (1 << 1), + ROCPROFILER_HIP_LIBRARY = (1 << 2), + ROCPROFILER_MARKER_LIBRARY = (1 << 3), + ROCPROFILER_LIBRARY_LAST = ROCPROFILER_MARKER_LIBRARY, +} rocprofiler_runtime_library_t; + //--------------------------------------------------------------------------------------// // // ALIASES diff --git a/source/include/rocprofiler/intercept_table.h b/source/include/rocprofiler/intercept_table.h new file mode 100644 index 0000000000..f0019d3bc0 --- /dev/null +++ b/source/include/rocprofiler/intercept_table.h @@ -0,0 +1,223 @@ +// MIT License +// +// Copyright (c) 2023 ROCm Developer Tools +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#include +#include + +ROCPROFILER_EXTERN_C_INIT + +/** + * @defgroup INTERCEPT_TABLE Intercept table for runtime libraries + * @brief Enable tools to wrap the runtime API function calls of HIP, HSA, and ROCTx before and + * after the "real" implementation is called. + * + * When an application invokes the public function from the HIP, HSA, and ROCTx libraries, these + * functions invoke a function pointer which, when properly chained, allow tools to wrap these + * function calls to collect information. When this capability is used alongside the rocprofiler API + * tracing, tools will wrap the rocprofiler wrappers of the API function, e.g. if the tool installs + * a wrapper around the `hsa_init` function called `tool_hsa_init`, and rocprofiler installs a + * wrapper around the `hsa_init` function called `rocp_hsa_init`, and within the HSA runtime + * library, the "real" implementation of the `hsa_init` invokes a function called `real_hsa_init`, + * the invocation chain (starting from within the user application) will be: `` -> + * `hsa_init` -> `tool_hsa_init` -> `rocp_hsa_init` -> `real_hsa_init`. The return sequence will be + * the inverse of invocation chain: `real_hsa_init` -> `rocp_hsa_init` -> `tool_hsa_init` -> + * ``. Thus, it is important for tools that use this feature to (A) call the next + * function in the chain and (B) properly handle the return value. + * + * @{ + */ + +/** + * @brief Callback type when a new runtime library is loaded. @see + * rocprofiler_at_runtime_api_registration + * @param [in] type Type of API table + * @param [in] lib_version Major, minor, and patch version of library encoded into single number + * similar to @ref ROCPROFILER_VERSION + * @param [in] lib_instance The number of times this runtime library has been registered previously + * @param [in] tables An array of pointers to the API tables + * @param [in] num_tables The size of the array of pointers to the API tables + * @param [in] user_data The pointer to the data provided to @ref + * rocprofiler_at_runtime_api_registration + */ +typedef void (*rocprofiler_intercept_library_cb_t)(rocprofiler_runtime_library_t type, + uint64_t lib_version, + uint64_t lib_instance, + void** tables, + uint64_t num_tables, + void* user_data); + +/** + * @brief Invoke this function to receive callbacks when a ROCm library registers its API + * intercept table with rocprofiler. Use the @ref rocprofiler_runtime_library_t enumeration for + * specifying which raw API tables the tool would like to have access to. E.g. including @ref + * ROCPROFILER_HSA_LIBRARY in the @ref rocprofiler_at_runtime_api_registration function call + * communicates to rocprofiler that, when rocprofiler receives a `HsaApiTable` instance, the tool + * would like rocprofiler to provide it access too. + * + * When the HIP, HSA, and ROCTx libraries are initialized (either explicitly or on the first + * invocation of one of their public API functions), these runtimes will provide a table of function + * pointers to the rocprofiler library via the rocprofiler-register library if the + * `rocprofiler_configure` symbol is visible in the application's symbol table. The vast majority of + * tools will want to use the @ref CALLBACK_TRACING_SERVICE to trace these runtime APIs, however, + * some tools may want or require installing their own intercept functions in lieu of receiving + * these callbacks and those tools should use the @ref rocprofiler_at_runtime_api_registration to + * install their intercept functions. There are no restrictions to where or how early this function + * can be invoked but it will return ::ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED if it is + * invoked after rocprofiler has requested all the tool configurations. Thus, it is highly + * recommended to invoke this function within the @ref rocprofiler_configure function or the + * callback passed to the @ref rocprofiler_force_configure function -- the reason for this + * recommendation is that if @ref rocprofiler_at_runtime_api_registration is invoked in one of these + * locations, rocprofiler can guarantee that the tool will be passed the API table because, at the + * first instance of a runtime registering it's API table, rocprofiler will ensure that, in the case + * of the former, rocprofiler will invoke all of the @ref rocprofiler_configure symbols that are + * visible before checking the list of tools which want to receive the API tables and, in the case + * of the latter, @ref rocprofiler_force_configure will fail with error code @ref + * ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED if a runtime has already been registered (and, + * therefore, already scanned and invoked the visible @ref rocprofiler_configure symbols and + * completed the tool initialization). If @ref rocprofiler_at_runtime_api_registration is invoked + * outside of these recommended places, even if it is done before the `main` function starts (e.g. + * in a library init/constructor function), it is possible that another library, such as ROCm-aware + * MPI, caused the HIP and HSA runtime libraries to be initialized when that library was loaded. In + * this aforementioned scenario, if the ROCm-aware MPI library library init/constructor function + * runs before your library init/constructor function, rocprofiler will have already processed the + * API table and will not provide the API table to the tool due to the fact that the API may already + * be in use and, thus, any modifications to the table might result in thread-safety violations or + * more disastrous consequences. + * + * @param [in] callback Callback to tool invoked when a runtime registers their API table with + * rocprofiler + * @param [in] libs Bitwise-or of libraries, e.g. `ROCPROFILER_HSA_LIBRARY | ROCPROFILER_HIP_LIBRARY + * | ROCPROFILER_MARKER_LIBRARY` means the callbacks will be invoked whenever the HSA, HIP, and + * ROCTx libraries register the intercept table. + * @param [in] data Data to provide to callback(s) + * @return ::rocprofiler_status_t + * @retval ::ROCPROFILER_STATUS_SUCCESS Callback was registered for specified runtime(s) + * @retval ::ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED rocprofiler has already initialized + * @retval ::ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT this error code is returned if + * `ROCPROFILER_LIBRARY` is included in bitwise-or of the libs + * @retval ::ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED this error code is returned if one of the + * specified libraries does not have support for API intercept tables (which should not be the case + * by the time this code is publicly released) + * + * @code{.cpp} + * namespace + * { + * // this function generates a wrapper around the original function that + * // prints out the function name and then invokes the original function + * template + * auto + * generate_wrapper(const char* name, RetT (*func)(Args...)) + * { + * using functor_type = RetT (*)(Args...); + * + * // save function name, "real function" + * static const auto* func_name = name; + * static functor_type underlying_func = func; + * static functor_type wrapped_func = [](Args... args) -> RetT { + * std::cout << "Wrapping " << func_name << "..." << std::endl; + * if(underlying_func) return underlying_func(args...); + * if constexpr(!std::is_void::value) return RetT{}; + * }; + * + * return wrapped_func; + * } + * + * + * // this macro installs the wrapper in place of the original function + * #define GENERATE_WRAPPER(TABLE, FUNC) \ + * TABLE->FUNC##_fn = generate_wrapper<__COUNTER__>(#FUNC, TABLE->FUNC##_fn) + * + * + * // this is the function that gets called when the HSA runtime + * // intercept table is registered with rocprofiler + * void + * api_registration_callback(rocprofiler_runtime_library_t type, + * uint64_t lib_version, + * uint64_t lib_instance, + * void** tables, + * uint64_t num_tables, + * void* user_data) + * { + * if(type != ROCPROFILER_HSA_LIBRARY) + * throw std::runtime_error{"unexpected library type: " + + * std::to_string(static_cast(type))}; + * if(lib_instance != 0) throw std::runtime_error{"multiple instances of HSA runtime library"}; + * if(num_tables != 1) throw std::runtime_error{"expected only one table of type HsaApiTable"}; + * + * auto* hsa_api_table = static_cast(tables[0]); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_agent_get_info); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_agent_iterate_isas); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_code_object_reader_create_from_memory); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_executable_create_alt); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_executable_freeze); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_executable_get_symbol_by_name); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_executable_iterate_symbols); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_executable_load_agent_code_object); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_executable_symbol_get_info); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_isa_get_info_alt); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_iterate_agents); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_queue_add_write_index_screlease); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_queue_create); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_queue_load_read_index_relaxed); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_queue_load_read_index_scacquire); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_signal_create); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_signal_destroy); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_signal_load_relaxed); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_signal_silent_store_relaxed); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_signal_store_screlease); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_signal_wait_scacquire); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_system_get_info); + * GENERATE_WRAPPER(hsa_api_table->core_, hsa_system_get_major_extension_table); + * } + * } // namespace + * + * + * 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 = "ExampleTool"; + * + * // specify that we only want to intercept the HSA library + * rocprofiler_at_runtime_api_registration(api_registration_callback, + * ROCPROFILER_HSA_LIBRARY, nullptr); + * + * return nullptr; + * } + * @endcode + * + * @example intercept_table/client.cpp + * Example demonstrating @ref rocprofiler_at_runtime_api_registration usage + */ +rocprofiler_status_t +rocprofiler_at_runtime_api_registration(rocprofiler_intercept_library_cb_t callback, + int libs, + void* data) ROCPROFILER_API; + +/** @} */ + +ROCPROFILER_EXTERN_C_FINI diff --git a/source/include/rocprofiler/internal_threading.h b/source/include/rocprofiler/internal_threading.h index 4c9d373e04..154fc65d32 100644 --- a/source/include/rocprofiler/internal_threading.h +++ b/source/include/rocprofiler/internal_threading.h @@ -32,44 +32,37 @@ ROCPROFILER_EXTERN_C_INIT * @brief Callbacks before and after threads created internally by libraries * * @{ + * @example api_buffered_tracing/client.cpp + * Example demonstrating @ref BUFFER_TRACING_SERVICE that includes usage of @ref + * rocprofiler_at_internal_thread_create, @ref rocprofiler_create_callback_thread, and @ref + * rocprofiler_assign_callback_thread. */ -/** - * @brief Enumeration for specifying which libraries you want callbacks before and after the library - * creates an internal thread. These callbacks will be invoked on the thread that is about to create - * the new thread (not on the newly created thread). In thread-aware tools that wrap pthread_create, - * this can be used to disable the wrapper before the pthread_create invocation and re-enable the - * wrapper afterwards. In many cases, tools will want to ignore the thread(s) created by rocprofiler - * since these threads do not exist in the normal application execution, whereas the internal - * threads for HSA, HIP, etc. are created in normal application execution; however, the HIP, HSA, - * etc. internal threads are typically background threads which just monitor kernel completion and - * are unlikely to contribute to any performance issues. - */ -typedef enum -{ - ROCPROFILER_LIBRARY = (1 << 0), - ROCPROFILER_HSA_LIBRARY = (1 << 1), - ROCPROFILER_HIP_LIBRARY = (1 << 2), - ROCPROFILER_MARKER_LIBRARY = (1 << 3), - ROCPROFILER_LIBRARY_LAST = ROCPROFILER_MARKER_LIBRARY, -} rocprofiler_internal_thread_library_t; - /** * @brief Callback type before and after internal thread creation. @see * rocprofiler_at_internal_thread_create * */ -typedef void (*rocprofiler_internal_thread_library_cb_t)(rocprofiler_internal_thread_library_t, - void*); +typedef void (*rocprofiler_internal_thread_library_cb_t)(rocprofiler_runtime_library_t, void*); /** * @brief Invoke this function to receive callbacks before and after the creation of an internal * thread by a library which as invoked on the thread which is creating the internal thread(s). - * Please note that the postcreate callback is guaranteed to be invoked after the underlying - * system call to create a new thread but it does not guarantee that the new thread has been - * started. Please note, that once these callbacks are registered, they cannot be removed so the - * caller is responsible for ignoring these callbacks if they want to ignore them beyond a certain - * point in the application. + * + * Use the @ref rocprofiler_runtime_library_t enumeration for specifying which libraries you want + * callbacks before and after the library creates an internal thread. These callbacks will be + * invoked on the thread that is about to create the new thread (not on the newly created thread). + * In thread-aware tools that wrap pthread_create, this can be used to disable the wrapper before + * the pthread_create invocation and re-enable the wrapper afterwards. In many cases, tools will + * want to ignore the thread(s) created by rocprofiler since these threads do not exist in the + * normal application execution, whereas the internal threads for HSA, HIP, etc. are created in + * normal application execution; however, the HIP, HSA, etc. internal threads are typically + * background threads which just monitor kernel completion and are unlikely to contribute to any + * performance issues. Please note that the postcreate callback is guaranteed to be invoked after + * the underlying system call to create a new thread but it does not guarantee that the new thread + * has been started. Please note, that once these callbacks are registered, they cannot be removed + * so the caller is responsible for ignoring these callbacks if they want to ignore them beyond a + * certain point in the application. * * @param [in] precreate Callback invoked immediately before a new internal thread is created * @param [in] postcreate Callback invoked immediately after a new internal thread is created @@ -77,6 +70,9 @@ typedef void (*rocprofiler_internal_thread_library_cb_t)(rocprofiler_internal_th * means the callbacks will be invoked whenever rocprofiler and/or the marker library create * internal threads but not when the HSA or HIP libraries create internal threads. * @param [in] data Data shared between callbacks + * @return ::rocprofiler_status_t + * @retval ::ROCPROFILER_STATUS_SUCCESS There are currently no conditions which result in any other + * value, even if internal threads have already been created */ rocprofiler_status_t ROCPROFILER_API rocprofiler_at_internal_thread_create(rocprofiler_internal_thread_library_cb_t precreate, @@ -100,6 +96,8 @@ typedef struct * * @param [in] cb_thread_id User-provided pointer to a @ref rocprofiler_callback_thread_t * @return ::rocprofiler_status_t + * @retval ::ROCPROFILER_STATUS_SUCCESS Successful thread creation + * @retval ::ROCPROFILER_STATUS_ERROR Failed to create thread */ rocprofiler_status_t ROCPROFILER_API rocprofiler_create_callback_thread(rocprofiler_callback_thread_t* cb_thread_id) @@ -113,6 +111,12 @@ rocprofiler_create_callback_thread(rocprofiler_callback_thread_t* cb_thread_id) * @param [in] buffer_id Buffer identifier * @param [in] cb_thread_id Callback thread identifier via @ref rocprofiler_create_callback_thread * @return ::rocprofiler_status_t + * @retval ::ROCPROFILER_STATUS_SUCCESS Successful assignment of the delivery thread for the given + * buffer + * @retval ::ROCPROFILER_STATUS_ERROR_THREAD_NOT_FOUND Thread identifier did not match any of the + * threads created by rocprofiler + * @retval ::ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND Buffer identifier did not match any of the + * buffers registered with rocprofiler */ rocprofiler_status_t ROCPROFILER_API rocprofiler_assign_callback_thread(rocprofiler_buffer_id_t buffer_id, diff --git a/source/include/rocprofiler/rocprofiler.h b/source/include/rocprofiler/rocprofiler.h index 7fec25a63d..07c84ac976 100644 --- a/source/include/rocprofiler/rocprofiler.h +++ b/source/include/rocprofiler/rocprofiler.h @@ -65,7 +65,7 @@ ROCPROFILER_EXTERN_C_FINI /** @} */ #include "rocprofiler/agent.h" -#include "rocprofiler/agent_profile.h" +// #include "rocprofiler/agent_profile.h" #include "rocprofiler/buffer.h" #include "rocprofiler/buffer_tracing.h" #include "rocprofiler/callback_tracing.h" @@ -75,11 +75,12 @@ ROCPROFILER_EXTERN_C_FINI #include "rocprofiler/external_correlation.h" #include "rocprofiler/hip.h" #include "rocprofiler/hsa.h" +#include "rocprofiler/intercept_table.h" #include "rocprofiler/internal_threading.h" -#include "rocprofiler/marker.h" +// #include "rocprofiler/marker.h" #include "rocprofiler/pc_sampling.h" #include "rocprofiler/profile_config.h" -#include "rocprofiler/spm.h" +// #include "rocprofiler/spm.h" ROCPROFILER_EXTERN_C_INIT diff --git a/source/lib/rocprofiler/CMakeLists.txt b/source/lib/rocprofiler/CMakeLists.txt index 24e2caf220..96e7dd5d5a 100644 --- a/source/lib/rocprofiler/CMakeLists.txt +++ b/source/lib/rocprofiler/CMakeLists.txt @@ -4,7 +4,7 @@ rocprofiler_activate_clang_tidy() set(ROCPROFILER_LIB_HEADERS agent.hpp buffer.hpp external_correlation.hpp - internal_threading.hpp registration.hpp) + intercept_table.hpp internal_threading.hpp registration.hpp) set(ROCPROFILER_LIB_SOURCES agent.cpp buffer.cpp @@ -14,6 +14,7 @@ set(ROCPROFILER_LIB_SOURCES counters.cpp dispatch_profile.cpp external_correlation.cpp + intercept_table.cpp internal_threading.cpp pc_sampling.cpp profile_config.cpp diff --git a/source/lib/rocprofiler/context.cpp b/source/lib/rocprofiler/context.cpp index 2c93fd0668..f1db18da77 100644 --- a/source/lib/rocprofiler/context.cpp +++ b/source/lib/rocprofiler/context.cpp @@ -21,6 +21,7 @@ // SOFTWARE. #include +#include #include #include "lib/rocprofiler/context/context.hpp" @@ -31,12 +32,17 @@ #include #include +namespace +{ +constexpr auto rocprofiler_context_none = ROCPROFILER_CONTEXT_NONE; +} + extern "C" { rocprofiler_status_t rocprofiler_create_context(rocprofiler_context_id_t* context_id) { // always set to none first - *context_id = ROCPROFILER_CONTEXT_NONE; + *context_id = rocprofiler_context_none; if(rocprofiler::registration::get_init_status() > -1) return ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED; @@ -50,12 +56,22 @@ rocprofiler_create_context(rocprofiler_context_id_t* context_id) rocprofiler_status_t rocprofiler_start_context(rocprofiler_context_id_t context_id) { + if(context_id.handle == rocprofiler_context_none.handle) + return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; + + // if currently finalizing or finalized, don't allow starting a context + if(rocprofiler::registration::get_fini_status() != 0) + return ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED; + return rocprofiler::context::start_context(context_id); } rocprofiler_status_t rocprofiler_stop_context(rocprofiler_context_id_t context_id) { + if(context_id.handle == rocprofiler_context_none.handle) + return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; + return rocprofiler::context::stop_context(context_id); } @@ -63,6 +79,10 @@ rocprofiler_status_t rocprofiler_context_is_active(rocprofiler_context_id_t context_id, int* status) { *status = 0; + + if(context_id.handle == rocprofiler_context_none.handle) + return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; + for(const auto& itr : rocprofiler::context::get_active_contexts()) { const auto* cfg = itr.load(std::memory_order_relaxed); @@ -79,6 +99,10 @@ rocprofiler_status_t rocprofiler_context_is_valid(rocprofiler_context_id_t context_id, int* status) { *status = 0; + + if(context_id.handle == rocprofiler_context_none.handle) + return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; + for(const auto& itr : rocprofiler::context::get_registered_contexts()) { if(itr && itr->context_idx == context_id.handle) diff --git a/source/lib/rocprofiler/intercept_table.cpp b/source/lib/rocprofiler/intercept_table.cpp new file mode 100644 index 0000000000..0e31c08769 --- /dev/null +++ b/source/lib/rocprofiler/intercept_table.cpp @@ -0,0 +1,184 @@ +// MIT License +// +// Copyright (c) 2023 ROCm Developer Tools +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include +#include +#include + +#include "lib/common/container/stable_vector.hpp" +#include "lib/common/utility.hpp" +#include "lib/rocprofiler/buffer.hpp" +#include "lib/rocprofiler/context/context.hpp" +#include "lib/rocprofiler/internal_threading.hpp" +#include "lib/rocprofiler/registration.hpp" + +#include + +#include +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace intercept_table +{ +namespace +{ +template +using library_sequence_t = std::integer_sequence; + +// this is used to loop over the different libraries +constexpr auto intercept_library_seq = library_sequence_t{}; + +// check that intercept_library_seq is up to date +static_assert((1 << (intercept_library_seq.size())) == ROCPROFILER_LIBRARY_LAST, + "Update intercept_library_seq to include new libraries"); + +// data structure holding list of callbacks +template +struct intercept +{ + static constexpr auto value = LibT; + + std::vector callbacks = {}; + std::vector user_data = {}; + std::mutex mutex = {}; +}; + +// static accessor for intercept instance +template +auto& +get_intercept() +{ + static auto _v = intercept{}; + return _v; +} + +// adds callbacks to intercept instance(s) +template +void +update_intercepts(rocprofiler_intercept_library_cb_t cb, + int libs, + void* data, + library_sequence_t) +{ + auto update = [cb, libs, data](auto& notifier) { + if(libs == 0 || ((libs & notifier.value) == notifier.value)) + { + notifier.mutex.lock(); + notifier.callbacks.emplace_back(cb); + notifier.user_data.emplace_back(data); + notifier.mutex.unlock(); + } + }; + + (update(get_intercept()), ...); +} + +template +auto +get_void_array(std::tuple data, std::index_sequence) +{ + constexpr auto size = sizeof...(Idx); + return std::array{static_cast(std::get(data))...}; +}; + +// invokes creation notifiers +template +void +execute_intercepts(rocprofiler_runtime_library_t lib, + uint64_t lib_version, + uint64_t lib_instance, + std::tuple tables, + std::integer_sequence) +{ + auto execute = [lib, lib_version, lib_instance, tables](auto& notifier) { + if(((lib & notifier.value) == notifier.value)) + { + constexpr uint64_t num_tables = sizeof...(ApiTableT); + auto tables_v = get_void_array(tables, std::make_index_sequence{}); + + notifier.mutex.lock(); + for(size_t i = 0; i < notifier.callbacks.size(); ++i) + { + auto itr = notifier.callbacks.at(i); + if(itr) + itr(notifier.value, + lib_version, + lib_instance, + tables_v.data(), + num_tables, + notifier.user_data.at(i)); + } + notifier.mutex.unlock(); + } + }; + + (execute(get_intercept()), ...); +} +} // namespace + +template +void +notify_runtime_api_registration(rocprofiler_runtime_library_t lib, + uint64_t lib_version, + uint64_t lib_instance, + std::tuple tables) +{ + execute_intercepts(lib, lib_version, lib_instance, tables, intercept_library_seq); +} + +// template instantiation for HsaApiTable +template void notify_runtime_api_registration(rocprofiler_runtime_library_t, + uint64_t, + uint64_t, + std::tuple); +} // namespace intercept_table +} // namespace rocprofiler + +extern "C" { +rocprofiler_status_t +rocprofiler_at_runtime_api_registration(rocprofiler_intercept_library_cb_t callback, + int libs, + void* data) +{ + // if this function is invoked after initialization, we cannot guarantee that the runtime + // intercept API has not already be registered and returned to the runtime. + if(rocprofiler::registration::get_init_status() > 0) + return ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED; + + if((libs & ROCPROFILER_LIBRARY) == ROCPROFILER_LIBRARY) + return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; + else if((libs & ROCPROFILER_HIP_LIBRARY) == ROCPROFILER_HIP_LIBRARY) + return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; + else if((libs & ROCPROFILER_MARKER_LIBRARY) == ROCPROFILER_MARKER_LIBRARY) + return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; + + rocprofiler::intercept_table::update_intercepts( + callback, libs, data, rocprofiler::intercept_table::intercept_library_seq); + return ROCPROFILER_STATUS_SUCCESS; +} +} diff --git a/source/lib/rocprofiler/intercept_table.hpp b/source/lib/rocprofiler/intercept_table.hpp new file mode 100644 index 0000000000..4619ed95e6 --- /dev/null +++ b/source/lib/rocprofiler/intercept_table.hpp @@ -0,0 +1,40 @@ +// MIT License +// +// Copyright (c) 2023 ROCm Developer Tools +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#include + +#include + +namespace rocprofiler +{ +namespace intercept_table +{ +template +void +notify_runtime_api_registration(rocprofiler_runtime_library_t lib, + uint64_t lib_version, + uint64_t lib_instance, + std::tuple tables); +} +} // namespace rocprofiler diff --git a/source/lib/rocprofiler/internal_threading.cpp b/source/lib/rocprofiler/internal_threading.cpp index 16df79415e..16148bf304 100644 --- a/source/lib/rocprofiler/internal_threading.cpp +++ b/source/lib/rocprofiler/internal_threading.cpp @@ -43,9 +43,9 @@ namespace internal_threading { namespace { -template -using library_sequence_t = std::integer_sequence; -using creation_notifier_cb_t = void (*)(rocprofiler_internal_thread_library_t, void*); +template +using library_sequence_t = std::integer_sequence; +using creation_notifier_cb_t = void (*)(rocprofiler_runtime_library_t, void*); using thread_pool_config_t = PTL::ThreadPool::Config; // this is used to loop over the different libraries @@ -66,7 +66,7 @@ enum class notifier_stage }; // data structure holding list of callbacks -template +template struct creation_notifier { static constexpr auto value = LibT; @@ -78,7 +78,7 @@ struct creation_notifier }; // static accessor for creation_notifier instance -template +template auto& get_creation_notifier() { @@ -87,7 +87,7 @@ get_creation_notifier() } // adds callbacks to creation_notifier instance(s) -template +template void update_creation_notifiers(creation_notifier_cb_t pre, creation_notifier_cb_t post, @@ -110,10 +110,10 @@ update_creation_notifiers(creation_notifier_cb_t pre, } // invokes creation notifiers -template +template void -execute_creation_notifiers(rocprofiler_internal_thread_library_t libs, - std::integer_sequence) +execute_creation_notifiers(rocprofiler_runtime_library_t libs, + std::integer_sequence) { auto execute = [libs](auto& notifier) { if(((libs & notifier.value) == notifier.value)) @@ -189,13 +189,13 @@ finalize() } void -notify_pre_internal_thread_create(rocprofiler_internal_thread_library_t libs) +notify_pre_internal_thread_create(rocprofiler_runtime_library_t libs) { execute_creation_notifiers(libs, creation_notifier_library_seq); } void -notify_post_internal_thread_create(rocprofiler_internal_thread_library_t libs) +notify_post_internal_thread_create(rocprofiler_runtime_library_t libs) { execute_creation_notifiers(libs, creation_notifier_library_seq); } diff --git a/source/lib/rocprofiler/internal_threading.hpp b/source/lib/rocprofiler/internal_threading.hpp index 02417444f2..04dc0b3340 100644 --- a/source/lib/rocprofiler/internal_threading.hpp +++ b/source/lib/rocprofiler/internal_threading.hpp @@ -54,8 +54,8 @@ using task_group_vec_t = rocprofiler::common::static_cleanup_wrapper< std::vector, std::function&)>>; -void notify_pre_internal_thread_create(rocprofiler_internal_thread_library_t); -void notify_post_internal_thread_create(rocprofiler_internal_thread_library_t); +void notify_pre_internal_thread_create(rocprofiler_runtime_library_t); +void notify_post_internal_thread_create(rocprofiler_runtime_library_t); // initialize the default thread pool void diff --git a/source/lib/rocprofiler/registration.cpp b/source/lib/rocprofiler/registration.cpp index 76b79d4981..d6ab592ff9 100644 --- a/source/lib/rocprofiler/registration.cpp +++ b/source/lib/rocprofiler/registration.cpp @@ -26,6 +26,7 @@ #include "lib/rocprofiler/hsa/hsa.hpp" #include "lib/rocprofiler/hsa/queue.hpp" #include "lib/rocprofiler/hsa/queue_controller.hpp" +#include "lib/rocprofiler/intercept_table.hpp" #include "lib/rocprofiler/internal_threading.hpp" #include @@ -369,7 +370,10 @@ invoke_client_finalizer(rocprofiler_client_id_t client_id) { if(itr.configure_result && itr.configure_result->finalize) { + auto _fini_status = get_fini_status(); + if(_fini_status == 0) set_fini_status(-1); itr.configure_result->finalize(itr.configure_result->tool_data); + if(_fini_status == 0) set_fini_status(_fini_status); // set to nullptr so finalize only gets called once itr.configure_result->finalize = nullptr; } @@ -553,7 +557,12 @@ rocprofiler_set_api_table(const char* name, auto& saved_hsa_api_table = rocprofiler::hsa::get_table(); ::copyTables(hsa_api_table, &saved_hsa_api_table); + // install rocprofiler API wrappers rocprofiler::hsa::update_table(hsa_api_table); + + // allow tools to install API wrappers + rocprofiler::intercept_table::notify_runtime_api_registration( + ROCPROFILER_HSA_LIBRARY, lib_version, lib_instance, std::make_tuple(hsa_api_table)); } else if(std::string_view{name} == "roctx") { diff --git a/source/lib/rocprofiler/tests/CMakeLists.txt b/source/lib/rocprofiler/tests/CMakeLists.txt index 13e7d38980..d4f05f8dd3 100644 --- a/source/lib/rocprofiler/tests/CMakeLists.txt +++ b/source/lib/rocprofiler/tests/CMakeLists.txt @@ -34,7 +34,8 @@ set_tests_properties(${lib_TESTS} PROPERTIES TIMEOUT 45 LABELS "unittests") # # -------------------------------------------------------------------------------------- # -set(rocprofiler_shared_lib_sources external_correlation.cpp registration.cpp status.cpp) +set(rocprofiler_shared_lib_sources external_correlation.cpp intercept_table.cpp + registration.cpp status.cpp) add_executable(rocprofiler-lib-tests-shared) target_sources(rocprofiler-lib-tests-shared PRIVATE ${rocprofiler_shared_lib_sources}) diff --git a/source/lib/rocprofiler/tests/external_correlation.cpp b/source/lib/rocprofiler/tests/external_correlation.cpp index 5869607c70..bbf2e8c64f 100644 --- a/source/lib/rocprofiler/tests/external_correlation.cpp +++ b/source/lib/rocprofiler/tests/external_correlation.cpp @@ -208,14 +208,14 @@ tool_tracing_buffered(rocprofiler_context_id_t context, } void -thread_precreate(rocprofiler_internal_thread_library_t /*lib*/, void* tool_data) +thread_precreate(rocprofiler_runtime_library_t /*lib*/, void* tool_data) { auto* cb_data = static_cast(tool_data); cb_data->client_workflow_count++; } void -thread_postcreate(rocprofiler_internal_thread_library_t /*lib*/, void* tool_data) +thread_postcreate(rocprofiler_runtime_library_t /*lib*/, void* tool_data) { auto* cb_data = static_cast(tool_data); cb_data->client_workflow_count++; diff --git a/source/lib/rocprofiler/tests/intercept_table.cpp b/source/lib/rocprofiler/tests/intercept_table.cpp new file mode 100644 index 0000000000..97e21d76ea --- /dev/null +++ b/source/lib/rocprofiler/tests/intercept_table.cpp @@ -0,0 +1,631 @@ +// MIT License +// +// Copyright (c) 2023 ROCm Developer Tools +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include +#include +#include +#include + +#include "lib/common/environment.hpp" +#include "lib/common/units.hpp" +#include "lib/common/utility.hpp" + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#define ROCPROFILER_CALL(ARG, MSG) \ + { \ + auto _status = (ARG); \ + EXPECT_EQ(_status, ROCPROFILER_STATUS_SUCCESS) \ + << MSG << " (" << _status << "=" << rocprofiler_get_status_string(_status) \ + << ") :: " << #ARG; \ + } + +#define ROCPROFILER_CALL_EXPECT(ARG, MSG, STATUS) \ + { \ + auto _status = (ARG); \ + EXPECT_EQ(_status, STATUS) << MSG << " (" << _status << "=" \ + << rocprofiler_get_status_string(_status) << ") :: " << #ARG; \ + } + +namespace +{ +struct callback_data +{ + using callback_count_data_t = std::pair; + using callback_count_map_t = std::map; + + rocprofiler_client_id_t* client_id = nullptr; + rocprofiler_client_finalize_t client_fini_func = nullptr; + rocprofiler_context_id_t client_ctx = {}; + rocprofiler_buffer_id_t client_buffer = {}; + rocprofiler_callback_thread_t client_thread = {}; + uint64_t client_workflow_count = {}; + uint64_t client_elapsed = {}; + int64_t current_depth = 0; + int64_t max_depth = 0; + callback_count_map_t client_callback_count = {}; +}; + +struct agent_data +{ + uint64_t agent_count = 0; + std::vector agents = {}; +}; + +using callback_kind_names_t = std::map; +using callback_kind_operation_names_t = + std::map>; +using wrap_count_t = std::pair; + +struct callback_name_info +{ + callback_kind_names_t kind_names = {}; + callback_kind_operation_names_t operation_names = {}; +}; + +auto& +get_preincrement() +{ + static size_t _v = 0; + return _v; +} + +auto& +get_postincrement() +{ + static size_t _v = 0; + return _v; +} + +auto& +get_client_callback_data() +{ + static auto _v = callback_data{}; + return _v; +} + +auto +get_callback_tracing_names() +{ + auto cb_name_info = callback_name_info{}; + // + // callback for each kind operation + // + static auto tracing_kind_operation_cb = + [](rocprofiler_service_callback_tracing_kind_t kindv, uint32_t operation, void* data_v) { + auto* name_info_v = static_cast(data_v); + + if(kindv == ROCPROFILER_CALLBACK_TRACING_HSA_API) + { + const char* name = nullptr; + ROCPROFILER_CALL(rocprofiler_query_callback_tracing_kind_operation_name( + kindv, operation, &name, nullptr), + "query callback tracing kind operation name"); + if(name) name_info_v->operation_names[kindv][operation] = name; + } + return 0; + }; + + // + // callback for each callback kind (i.e. domain) + // + static auto tracing_kind_cb = [](rocprofiler_service_callback_tracing_kind_t kind, void* data) { + // store the callback kind name + auto* name_info_v = static_cast(data); + const char* name = nullptr; + ROCPROFILER_CALL(rocprofiler_query_callback_tracing_kind_name(kind, &name, nullptr), + "query callback tracing kind operation name"); + if(name) name_info_v->kind_names[kind] = name; + + if(kind == ROCPROFILER_CALLBACK_TRACING_HSA_API) + { + ROCPROFILER_CALL(rocprofiler_iterate_callback_tracing_kind_operations( + kind, tracing_kind_operation_cb, static_cast(data)), + "iterating callback tracing kind operations"); + } + return 0; + }; + + ROCPROFILER_CALL(rocprofiler_iterate_callback_tracing_kinds(tracing_kind_cb, + static_cast(&cb_name_info)), + "iterating callback tracing kinds"); + + return cb_name_info; +} + +void +tool_tracing_callback(rocprofiler_callback_tracing_record_t record, + rocprofiler_user_data_t* /*user_data*/, + void* client_data) +{ + auto* cb_data = static_cast(client_data); + + static auto name_map = get_callback_tracing_names(); + if(record.phase == ROCPROFILER_CALLBACK_PHASE_ENTER) + { + cb_data->client_callback_count[name_map.operation_names[record.kind][record.operation]] + .second += get_preincrement(); + } + else if(record.phase == ROCPROFILER_CALLBACK_PHASE_EXIT) + { + cb_data->client_callback_count[name_map.operation_names[record.kind][record.operation]] + .second += get_postincrement(); + } +} + +template +auto +generate_wrapper(const char* name, RetT (*func)(Args...)) +{ + using functor_type = RetT (*)(Args...); + + static functor_type underlying_func = func; + static auto func_name = std::string_view{name}; + + get_client_callback_data().client_callback_count.emplace( + std::string_view{name}, callback_data::callback_count_data_t{0, 0}); + + static functor_type wrapped_func = [](Args... args) -> RetT { + get_client_callback_data().client_callback_count.at(func_name).first += + get_preincrement() + get_postincrement(); + if(underlying_func) return underlying_func(args...); + if constexpr(!std::is_void::value) return RetT{}; + }; + + return wrapped_func; +} + +#define GENERATE_WRAPPER(TABLE, FUNC) \ + TABLE->FUNC##_fn = generate_wrapper<__COUNTER__>(#FUNC, TABLE->FUNC##_fn) + +void +api_registration_callback(rocprofiler_runtime_library_t type, + uint64_t lib_version, + uint64_t lib_instance, + void** tables, + uint64_t num_tables, + void* client_data) +{ + auto* cb_data = static_cast(client_data); + + cb_data->client_workflow_count++; + + EXPECT_EQ(type, ROCPROFILER_HSA_LIBRARY) << "unexpected library type: " << type; + EXPECT_EQ(lib_instance, 0) << "multiple instances of HSA runtime library"; + EXPECT_EQ(num_tables, 1) << "expected only one table of type HsaApiTable"; + EXPECT_GT(lib_version, 0) << "expected library version > 0"; + + auto* hsa_api_table = static_cast(tables[0]); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_init); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_agent_get_info); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_iterate_agents); + GENERATE_WRAPPER(hsa_api_table->core_, hsa_shut_down); +} +} // namespace + +TEST(rocprofiler_lib, intercept_table_and_callback_tracing) +{ + // test layering of tool interception of API table on top of rocprofiler API tracing. + // This test enables both rocprofiler API tracing and a tool installing it's own + // wrappers via the HsaApiTable. With both active, one should see the same + // number of calls to "hsa_init", "hsa_iterate_agents", "hsa_agent_get_info", and + // "hsa_shut_down" + + using init_func_t = int (*)(rocprofiler_client_finalize_t, void*); + using fini_func_t = void (*)(void*); + using hsa_iterate_agents_cb_t = hsa_status_t (*)(hsa_agent_t, void*); + + get_preincrement() = 1; + get_postincrement() = 0; + + static init_func_t tool_init = [](rocprofiler_client_finalize_t fini_func, + void* client_data) -> int { + auto* cb_data = static_cast(client_data); + + cb_data->client_workflow_count++; + cb_data->client_fini_func = fini_func; + + ROCPROFILER_CALL(rocprofiler_create_context(&cb_data->client_ctx), + "failed to create context"); + + auto operations = std::vector{ROCPROFILER_HSA_API_ID_hsa_init, + ROCPROFILER_HSA_API_ID_hsa_iterate_agents, + ROCPROFILER_HSA_API_ID_hsa_agent_get_info, + ROCPROFILER_HSA_API_ID_hsa_shut_down}; + + ROCPROFILER_CALL( + rocprofiler_configure_callback_tracing_service(cb_data->client_ctx, + ROCPROFILER_CALLBACK_TRACING_HSA_API, + operations.data(), + operations.size(), + tool_tracing_callback, + client_data), + "callback tracing service failed to configure"); + + int valid_ctx = 0; + ROCPROFILER_CALL(rocprofiler_context_is_valid(cb_data->client_ctx, &valid_ctx), + "failure checking context validity"); + + EXPECT_EQ(valid_ctx, 1); + + ROCPROFILER_CALL(rocprofiler_start_context(cb_data->client_ctx), + "rocprofiler context start failed"); + + // no errors + return 0; + }; + + static fini_func_t tool_fini = [](void* client_data) -> void { + auto* cb_data = static_cast(client_data); + ROCPROFILER_CALL(rocprofiler_stop_context(cb_data->client_ctx), + "rocprofiler context stop failed"); + + static_cast(client_data)->client_workflow_count++; + }; + + static auto& cb_data = get_client_callback_data(); + + static auto cfg_result = + rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), + tool_init, + tool_fini, + static_cast(&cb_data)}; + + static rocprofiler_configure_func_t rocp_init = + [](uint32_t version, + const char* runtime_version, + uint32_t prio, + rocprofiler_client_id_t* client_id) -> rocprofiler_tool_configure_result_t* { + auto expected_version = ROCPROFILER_VERSION; + EXPECT_EQ(expected_version, version); + EXPECT_EQ(std::string_view{runtime_version}, std::string_view{ROCPROFILER_VERSION_STRING}); + EXPECT_EQ(prio, 0); + EXPECT_EQ(client_id->name, nullptr); + cb_data.client_id = client_id; + cb_data.client_id->name = ::testing::UnitTest::GetInstance()->current_test_info()->name(); + + ROCPROFILER_CALL_EXPECT( + rocprofiler_at_runtime_api_registration(api_registration_callback, + ROCPROFILER_LIBRARY | ROCPROFILER_HSA_LIBRARY | + ROCPROFILER_HIP_LIBRARY | + ROCPROFILER_MARKER_LIBRARY, + static_cast(&cb_data)), + "function should return invalid argument if ROCPROFILER_LIBRARY included", + ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT); + + using init_list_t = std::initializer_list; + for(auto itr : init_list_t{ + (ROCPROFILER_HSA_LIBRARY | ROCPROFILER_HIP_LIBRARY | ROCPROFILER_MARKER_LIBRARY), + (ROCPROFILER_HSA_LIBRARY | ROCPROFILER_HIP_LIBRARY), + (ROCPROFILER_HIP_LIBRARY), + (ROCPROFILER_HSA_LIBRARY | ROCPROFILER_MARKER_LIBRARY), + (ROCPROFILER_MARKER_LIBRARY)}) + { + ROCPROFILER_CALL_EXPECT( + rocprofiler_at_runtime_api_registration( + api_registration_callback, itr, static_cast(&cb_data)), + "test should be updated if new (non-HSA) intercept table is supported", + ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED); + } + + ROCPROFILER_CALL_EXPECT(rocprofiler_at_runtime_api_registration( + api_registration_callback, + ROCPROFILER_HSA_LIBRARY | ROCPROFILER_HIP_LIBRARY, + static_cast(&cb_data)), + "test should be updated if HIP intercept table is supported", + ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED); + + ROCPROFILER_CALL_EXPECT( + rocprofiler_at_runtime_api_registration( + api_registration_callback, ROCPROFILER_HIP_LIBRARY, static_cast(&cb_data)), + "test should be updated if HIP intercept table is supported", + ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED); + + ROCPROFILER_CALL_EXPECT(rocprofiler_at_runtime_api_registration( + api_registration_callback, + ROCPROFILER_HSA_LIBRARY | ROCPROFILER_MARKER_LIBRARY, + static_cast(&cb_data)), + "test should be updated if ROCTx intercept table is supported", + ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED); + + ROCPROFILER_CALL( + rocprofiler_at_runtime_api_registration( + api_registration_callback, ROCPROFILER_HSA_LIBRARY, static_cast(&cb_data)), + "HSA API intercept table registration failed"); + + return &cfg_result; + }; + + EXPECT_EQ(rocprofiler_force_configure(rocp_init), ROCPROFILER_STATUS_SUCCESS); + + hsa_iterate_agents_cb_t agent_cb = [](hsa_agent_t agent, void* data) { + static_cast(data)->agent_count++; + + auto status = HSA_STATUS_SUCCESS; + auto agent_type = hsa_device_type_t{}; + if((status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &agent_type)) == + HSA_STATUS_SUCCESS) + static_cast(data)->agents.emplace_back(agent_type); + + return status; + }; + + hsa_init(); + auto _agent_data = agent_data{}; + hsa_status_t itr_status = hsa_iterate_agents(agent_cb, static_cast(&_agent_data)); + hsa_shut_down(); + + EXPECT_EQ(itr_status, HSA_STATUS_SUCCESS); + EXPECT_GT(_agent_data.agent_count, 0); + EXPECT_EQ(_agent_data.agent_count, _agent_data.agents.size()); + + ASSERT_NE(cb_data.client_id, nullptr); + ASSERT_NE(cb_data.client_fini_func, nullptr); + + cb_data.client_fini_func(*cb_data.client_id); + + EXPECT_EQ(cb_data.client_workflow_count, 3); + + for(auto itr : cb_data.client_callback_count) + { + EXPECT_EQ(itr.second.first, itr.second.second) + << "mismatched wrap counts for " << itr.first + << " (lhs=tool_wrapper, rhs=rocprofiler_wrapper)"; + if(itr.first != "hsa_init") + { + EXPECT_GT(itr.second.first, 0) << itr.first << " not wrapped"; + } + else + { + EXPECT_EQ(itr.second.first, 0) << itr.first + << " was wrapped. If hsa runtime has been updated to " + "include first call to hsa_init, update this test"; + } + } + + auto get_count = [](std::string_view func_name) { + // we already checked that first == second so we can just check first here + return cb_data.client_callback_count.at(func_name).first; + }; + + EXPECT_EQ(get_count("hsa_init"), 0); + EXPECT_EQ(get_count("hsa_iterate_agents"), 1); + EXPECT_EQ(get_count("hsa_agent_get_info"), _agent_data.agent_count); + EXPECT_EQ(get_count("hsa_shut_down"), 1); +} + +TEST(rocprofiler_lib, intercept_table_and_callback_tracing_disable_context) +{ + // test layering of tool interception of API table on top of rocprofiler API tracing. + // Similar to intercept_table_and_callback_tracing test except on the + // first call to hsa_iterate_agents the context is disabled. As a result, + // one should only see the rocprofiler API tracing for hsa_iterate_agents + // and not for hsa_agent_get_info or hsa_shut_down. + + using init_func_t = int (*)(rocprofiler_client_finalize_t, void*); + using fini_func_t = void (*)(void*); + using hsa_iterate_agents_cb_t = hsa_status_t (*)(hsa_agent_t, void*); + + get_preincrement() = 1; + get_postincrement() = 1; + + static init_func_t tool_init = [](rocprofiler_client_finalize_t fini_func, + void* client_data) -> int { + auto* cb_data = static_cast(client_data); + + cb_data->client_workflow_count++; + cb_data->client_fini_func = fini_func; + + ROCPROFILER_CALL(rocprofiler_create_context(&cb_data->client_ctx), + "failed to create context"); + + auto operations = std::vector{ROCPROFILER_HSA_API_ID_hsa_init, + ROCPROFILER_HSA_API_ID_hsa_iterate_agents, + ROCPROFILER_HSA_API_ID_hsa_agent_get_info, + ROCPROFILER_HSA_API_ID_hsa_shut_down}; + + ROCPROFILER_CALL( + rocprofiler_configure_callback_tracing_service(cb_data->client_ctx, + ROCPROFILER_CALLBACK_TRACING_HSA_API, + operations.data(), + operations.size(), + tool_tracing_callback, + client_data), + "callback tracing service failed to configure"); + + int valid_ctx = 0; + ROCPROFILER_CALL(rocprofiler_context_is_valid(cb_data->client_ctx, &valid_ctx), + "failure checking context validity"); + + EXPECT_EQ(valid_ctx, 1); + + ROCPROFILER_CALL(rocprofiler_start_context(cb_data->client_ctx), + "rocprofiler context start failed"); + + // no errors + return 0; + }; + + static fini_func_t tool_fini = [](void* client_data) -> void { + auto* cb_data = static_cast(client_data); + ROCPROFILER_CALL_EXPECT(rocprofiler_stop_context(cb_data->client_ctx), + "rocprofiler context stop", + ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND); + + ROCPROFILER_CALL_EXPECT(rocprofiler_start_context(cb_data->client_ctx), + "rocprofiler context start", + ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED); + + static_cast(client_data)->client_workflow_count++; + }; + + static auto& cb_data = get_client_callback_data(); + cb_data = callback_data{}; + + static auto cfg_result = + rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), + tool_init, + tool_fini, + static_cast(&cb_data)}; + + static rocprofiler_configure_func_t rocp_init = + [](uint32_t version, + const char* runtime_version, + uint32_t prio, + rocprofiler_client_id_t* client_id) -> rocprofiler_tool_configure_result_t* { + auto expected_version = ROCPROFILER_VERSION; + EXPECT_EQ(expected_version, version); + EXPECT_EQ(std::string_view{runtime_version}, std::string_view{ROCPROFILER_VERSION_STRING}); + EXPECT_EQ(prio, 0); + EXPECT_EQ(client_id->name, nullptr); + cb_data.client_id = client_id; + cb_data.client_id->name = ::testing::UnitTest::GetInstance()->current_test_info()->name(); + + ROCPROFILER_CALL_EXPECT( + rocprofiler_at_runtime_api_registration(api_registration_callback, + ROCPROFILER_LIBRARY | ROCPROFILER_HSA_LIBRARY | + ROCPROFILER_HIP_LIBRARY | + ROCPROFILER_MARKER_LIBRARY, + static_cast(&cb_data)), + "function should return invalid argument if ROCPROFILER_LIBRARY included", + ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT); + + using init_list_t = std::initializer_list; + for(auto itr : init_list_t{ + (ROCPROFILER_HSA_LIBRARY | ROCPROFILER_HIP_LIBRARY | ROCPROFILER_MARKER_LIBRARY), + (ROCPROFILER_HSA_LIBRARY | ROCPROFILER_HIP_LIBRARY), + (ROCPROFILER_HIP_LIBRARY), + (ROCPROFILER_HSA_LIBRARY | ROCPROFILER_MARKER_LIBRARY), + (ROCPROFILER_MARKER_LIBRARY)}) + { + ROCPROFILER_CALL_EXPECT( + rocprofiler_at_runtime_api_registration( + api_registration_callback, itr, static_cast(&cb_data)), + "test should be updated if new (non-HSA) intercept table is supported", + ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED); + } + + ROCPROFILER_CALL_EXPECT(rocprofiler_at_runtime_api_registration( + api_registration_callback, + ROCPROFILER_HSA_LIBRARY | ROCPROFILER_HIP_LIBRARY, + static_cast(&cb_data)), + "test should be updated if HIP intercept table is supported", + ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED); + + ROCPROFILER_CALL_EXPECT( + rocprofiler_at_runtime_api_registration( + api_registration_callback, ROCPROFILER_HIP_LIBRARY, static_cast(&cb_data)), + "test should be updated if HIP intercept table is supported", + ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED); + + ROCPROFILER_CALL_EXPECT(rocprofiler_at_runtime_api_registration( + api_registration_callback, + ROCPROFILER_HSA_LIBRARY | ROCPROFILER_MARKER_LIBRARY, + static_cast(&cb_data)), + "test should be updated if ROCTx intercept table is supported", + ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED); + + ROCPROFILER_CALL( + rocprofiler_at_runtime_api_registration( + api_registration_callback, ROCPROFILER_HSA_LIBRARY, static_cast(&cb_data)), + "HSA API intercept table registration failed"); + + return &cfg_result; + }; + + EXPECT_EQ(rocprofiler_force_configure(rocp_init), ROCPROFILER_STATUS_SUCCESS); + + hsa_iterate_agents_cb_t agent_cb = [](hsa_agent_t agent, void* data) { + auto* _data = static_cast*>(data); + _data->first->agent_count++; + + if(int _is_active = 0; + rocprofiler_context_is_active(_data->second->client_ctx, &_is_active) == + ROCPROFILER_STATUS_SUCCESS && + _is_active != 0) + { + ROCPROFILER_CALL(rocprofiler_stop_context(_data->second->client_ctx), + "rocprofiler context stop failed"); + } + + auto status = HSA_STATUS_SUCCESS; + auto agent_type = hsa_device_type_t{}; + if((status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &agent_type)) == + HSA_STATUS_SUCCESS) + _data->first->agents.emplace_back(agent_type); + + return status; + }; + + hsa_init(); + auto _agent_data = agent_data{}; + auto _pair_data = std::make_pair(&_agent_data, &cb_data); + hsa_status_t itr_status = hsa_iterate_agents(agent_cb, static_cast(&_pair_data)); + hsa_shut_down(); + + EXPECT_EQ(itr_status, HSA_STATUS_SUCCESS); + EXPECT_GT(_agent_data.agent_count, 0); + EXPECT_EQ(_agent_data.agent_count, _agent_data.agents.size()); + + ASSERT_NE(cb_data.client_id, nullptr); + ASSERT_NE(cb_data.client_fini_func, nullptr); + + cb_data.client_fini_func(*cb_data.client_id); + + EXPECT_EQ(cb_data.client_workflow_count, 3); + + auto get_tool_count = [](std::string_view func_name) { + // we already checked that first == second so we can just check first here + return cb_data.client_callback_count.at(func_name).first; + }; + + auto get_rocp_count = [](std::string_view func_name) { + // we already checked that first == second so we can just check first here + return cb_data.client_callback_count.at(func_name).second; + }; + + EXPECT_EQ(get_tool_count("hsa_init"), 0); + EXPECT_EQ(get_tool_count("hsa_iterate_agents"), 2); + EXPECT_EQ(get_tool_count("hsa_agent_get_info"), 2 * _agent_data.agent_count); + EXPECT_EQ(get_tool_count("hsa_shut_down"), 2); + + EXPECT_EQ(get_rocp_count("hsa_init"), 0); + EXPECT_EQ(get_rocp_count("hsa_iterate_agents"), 2) + << "if equal to 1, then ENTER phase was invoked but EXIT phase was not (incorrect)"; + EXPECT_EQ(get_rocp_count("hsa_agent_get_info"), 0); + EXPECT_EQ(get_rocp_count("hsa_shut_down"), 0); +} diff --git a/source/lib/rocprofiler/tests/registration.cpp b/source/lib/rocprofiler/tests/registration.cpp index 3b1d6ca1c4..919ec60dd5 100644 --- a/source/lib/rocprofiler/tests/registration.cpp +++ b/source/lib/rocprofiler/tests/registration.cpp @@ -337,14 +337,14 @@ tool_tracing_buffered(rocprofiler_context_id_t context, } void -thread_precreate(rocprofiler_internal_thread_library_t /*lib*/, void* tool_data) +thread_precreate(rocprofiler_runtime_library_t /*lib*/, void* tool_data) { auto* cb_data = static_cast(tool_data); cb_data->client_workflow_count++; } void -thread_postcreate(rocprofiler_internal_thread_library_t /*lib*/, void* tool_data) +thread_postcreate(rocprofiler_runtime_library_t /*lib*/, void* tool_data) { auto* cb_data = static_cast(tool_data); cb_data->client_workflow_count++;