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 <Tony.Tye@amd.com> * 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 <Tony.Tye@amd.com>
Este commit está contenido en:
Jonathan R. Madsen
cometido por
GitHub
GitHub
padre
14373c57be
commit
4f2dc896d3
@@ -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)
|
||||
|
||||
@@ -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<call_stack_t*>(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<call_stack_t*>(tool_data)->emplace_back(
|
||||
source_location{__FUNCTION__,
|
||||
|
||||
@@ -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
|
||||
$<TARGET_NAME_IF_EXISTS:rocprofiler::samples-build-flags>)
|
||||
|
||||
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
|
||||
$<TARGET_NAME_IF_EXISTS:rocprofiler::samples-build-flags>)
|
||||
|
||||
add_test(NAME intercept-table COMMAND $<TARGET_FILE:intercept-table>)
|
||||
|
||||
set_tests_properties(
|
||||
intercept-table
|
||||
PROPERTIES
|
||||
TIMEOUT
|
||||
45
|
||||
LABELS
|
||||
"samples"
|
||||
ENVIRONMENT
|
||||
"${ROCPROFILER_MEMCHECK_PRELOAD_ENV};HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler>"
|
||||
)
|
||||
@@ -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 <rocprofiler/registration.h>
|
||||
#include <rocprofiler/rocprofiler.h>
|
||||
|
||||
#include <cassert>
|
||||
#include <chrono>
|
||||
#include <cstddef>
|
||||
#include <cstdint>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <filesystem>
|
||||
#include <fstream>
|
||||
#include <functional>
|
||||
#include <iostream>
|
||||
#include <map>
|
||||
#include <mutex>
|
||||
#include <ratio>
|
||||
#include <string>
|
||||
#include <string_view>
|
||||
#include <vector>
|
||||
|
||||
#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<source_location>;
|
||||
using callback_kind_names_t = std::map<rocprofiler_service_callback_tracing_kind_t, const char*>;
|
||||
using callback_kind_operation_names_t =
|
||||
std::map<rocprofiler_service_callback_tracing_kind_t, std::map<uint32_t, const char*>>;
|
||||
using wrap_count_t = std::pair<source_location, size_t>;
|
||||
|
||||
rocprofiler_client_id_t* client_id = nullptr;
|
||||
std::map<size_t, wrap_count_t> 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<void(std::ostream*&)>{};
|
||||
|
||||
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<call_stack_t*>(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 <size_t Idx, typename RetT, typename... Args>
|
||||
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<RetT>::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<int>(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<std::vector<client::source_location>*>(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<HsaApiTable*>(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<uint32_t, 3>{};
|
||||
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<uint32_t, 3>{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<client::source_location>{};
|
||||
|
||||
// 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<void*>(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<void*>(client_tool_data)};
|
||||
|
||||
// return pointer to configure data
|
||||
return &cfg;
|
||||
}
|
||||
@@ -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
|
||||
@@ -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 <chrono>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <iostream>
|
||||
#include <mutex>
|
||||
#include <random>
|
||||
#include <stdexcept>
|
||||
|
||||
#define HIP_API_CALL(CALL) \
|
||||
{ \
|
||||
hipError_t error_ = (CALL); \
|
||||
if(error_ != hipSuccess) \
|
||||
{ \
|
||||
auto _hip_api_print_lk = auto_lock_t{print_lock}; \
|
||||
fprintf(stderr, \
|
||||
"%s:%d :: HIP error : %s\n", \
|
||||
__FILE__, \
|
||||
__LINE__, \
|
||||
hipGetErrorString(error_)); \
|
||||
throw std::runtime_error("hip_api_call"); \
|
||||
} \
|
||||
}
|
||||
|
||||
namespace
|
||||
{
|
||||
using auto_lock_t = std::unique_lock<std::mutex>;
|
||||
auto print_lock = std::mutex{};
|
||||
size_t 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<std::thread> _threads{};
|
||||
std::vector<hipStream_t> _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<int> _dist{0, 1000};
|
||||
|
||||
size_t size = sizeof(int) * M * N;
|
||||
int* inp_matrix = new int[size];
|
||||
int* out_matrix = new int[size];
|
||||
for(size_t i = 0; i < M * N; i++)
|
||||
{
|
||||
inp_matrix[i] = _dist(_engine);
|
||||
out_matrix[i] = 0;
|
||||
}
|
||||
int* in = nullptr;
|
||||
int* out = nullptr;
|
||||
|
||||
HIP_API_CALL(hipMalloc(&in, size));
|
||||
HIP_API_CALL(hipMalloc(&out, size));
|
||||
HIP_API_CALL(hipMemsetAsync(in, 0, size, stream));
|
||||
HIP_API_CALL(hipMemsetAsync(out, 0, size, stream));
|
||||
HIP_API_CALL(hipMemcpyAsync(in, inp_matrix, size, hipMemcpyHostToDevice, stream));
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
|
||||
dim3 grid(M / 32, N / 32, 1);
|
||||
dim3 block(32, 32, 1); // transpose_a
|
||||
|
||||
auto t1 = std::chrono::high_resolution_clock::now();
|
||||
for(size_t i = 0; i < nitr; ++i)
|
||||
{
|
||||
transpose_a<<<grid, block, 0, stream>>>(in, out, M, N);
|
||||
check_hip_error();
|
||||
if(i % nsync == (nsync - 1)) HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
}
|
||||
auto t2 = std::chrono::high_resolution_clock::now();
|
||||
HIP_API_CALL(hipStreamSynchronize(stream));
|
||||
HIP_API_CALL(hipMemcpyAsync(out_matrix, out, size, hipMemcpyDeviceToHost, stream));
|
||||
double time = std::chrono::duration_cast<std::chrono::duration<double>>(t2 - t1).count();
|
||||
float GB = (float) size * nitr * 2 / (1 << 30);
|
||||
|
||||
print_lock.lock();
|
||||
std::cout << "[" << 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
|
||||
Referencia en una nueva incidencia
Block a user