Code object, kernel dispatch, and memory copy tracing (#177)

* Update samples/api_buffered_tracing - external correlation id - support ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH * Update lib/rocprofiler/context.cpp - update alternative get_active_contexts paradigm * Update lib/rocprofiler/external_correlation.cpp - inherit correlation id from main thread * Update lib/rocprofiler/hsa/queue.* - typedef changes - rocprofiler_packet union - modify Queue::queue_info_session_t - use rocprofiler_packet - add thread id - add kernel id - add correlation id - out of line definitions - AsyncSignalHandler function update - handle kernel dispatch tracing - Move CreateBarrierPacket and AddVendorSpecificPacket to lambdas - handle contexts * Update lib/rocprofiler/hsa/hsa.cpp - remove unnecessary log function - use new get_active_contexts paradigm - use new correlation id updates * Update AgentCache and kernel dispatch record - include const rocprofiler_agent_t* in rocprofiler_buffer_tracing_kernel_dispatch_record_t - AgentCache::get_rocp_agent returns const pointer * Replace ROCPROFILER_SERVICE_ with ROCPROFILER_ * source formatting * Code Object Tracing - include/rocprofiler/callback_tracing.h - remove rocprofiler_callback_tracing_code_object_unload_data_t - remove rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t - include/rocprofiler/fwd.h - remove ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT_UNLOAD - remove ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_UNREGISTER - lib/common/utility.hpp - assert_public_api_struct_properties() - init_public_api_struct(...) - lib/rocprofiler/registration.cpp - invoke hsa::code_object_init - lib/rocprofiler/hsa/CMakeLists.txt - compile code_object code - lib/rocprofiler/hsa/code_object.{hpp,cpp} - tracing code object load/unload - lib/rocprofiler/hsa/queue.cpp - get_kernel_id * Update lib/rocprofiler/hsa/hsa.cpp - fix should_wrap_functor logic (which was not handling callback_tracer + buffered_tracer properly) * Update lib/rocprofiler/hsa/queue.cpp - fix rocprofiler_buffer_tracing_kernel_dispatch_record_t construction * Update samples/api_buffered_tracing/client.cpp - print kernel names * Move samples/apps to tests/apps * Update lib/rocprofiler/hsa/code_object.cpp - ensure unload callbacks when application is exiting - support user data in between load/unload callbacks * Update lib/rocprofiler/hsa/queue.{hpp,cpp} - store contexts and external correlation ids in queue_info_session - reduce signal_limiter to 96 to fix hangs - fix support for kernel tracing and async memory copies * Add lib/common/scope_destructor.hpp - similar to static_cleanup_wrapper but different * Update include/rocprofiler/buffer_tracing.h - update rocprofiler_buffer_tracing_memory_copy_record_t - remove operation: user can figure that out from correlation id - add kernel id - add rocprofiler agent id * Update include/rocprofiler/callback_tracing.h - fix data type of load_delta field in code object - remove rocp_agent from kernel_symbol_register_data_t (known via code_object_id) * Add samples/code_object_tracing - sample demonstrating code object tracing * Update samples - minor tweak to print_call_stack * Update lib/rocprofiler/hsa/code_object.cpp - flip ordering of unload callbacks for code object unloading and kernel symbol deregistering * clang-tidy fixes * Update lib/rocprofiler/hsa/code_object.cpp - fix heap-use-after-free issue with code object * Update include/rocprofiler/external_correlation.h - update documentation to include info about default value of external correlation value * Use common::container::small_vector for contexts - small_vector<const context*> is an ideal data structure for array of active contexts * Update context handling for code object unload - code object unload is only called for contexts which received the load callback * Update samples - improve ROCPROFILER_CALL macro to include status string - api_buffered_tracing handles ROCPROFILER_STATUS_ERROR_BUFFER_BUSY * Code object shutdown - ensure code object callbacks are invoked prior to finalizing * Update lib/common (memory allocators) - added lib/common/memory folder with allocators * Add lib/rocprofiler/allocator.* - rocprofiler::allocator::static_data_allocator - special allocator for static data which finalizes before any data gets destroyed - rocprofiler::allocator::unique_static_ptr_t - unique_ptr that uses static data deleter (ensure finalize is called) * Update lib/rocprofiler/buffer.cpp - flush checks fini status - use unique_static_ptr_t * Update lib/rocprofiler/internal_threading.* - change meaning of thread_pool_t and task_group_t - improve finalization to prevent data races and heap-use-after-free * Update lib/rocprofiler/registration.cpp - use static_data_allocator for client_library vector * Update lib/rocprofiler/context/context.* - use allocator::unique_static_ptr_t * Update lib/rocprofiler/allocator.cpp - avoid deadlock in deleter<static_data>::operator() * Update lib/rocprofiler/registration.cpp - avoid deadlock in rocprofiler::registration::finalize() * Update lib/rocprofiler/hsa/code_object.cpp - suppress duplicate reporting of code-object/kernel-symbol load/unload * Update leak sanitizer suppressions - __new_exitfn (via stdlib/cxa_atexit.c leaks [ROCm/rocprofiler-sdk commit: 3082288a25]
2023-11-13 22:30:15 -06:00
@@ -0,0 +1,60 @@
+#
+#
+#
+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-code-object-tracing 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(code-object-tracing-client SHARED)
+target_sources(code-object-tracing-client PRIVATE client.cpp)
+target_link_libraries(
+    code-object-tracing-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(code-object-tracing)
+target_sources(code-object-tracing PRIVATE main.cpp)
+target_link_libraries(
+    code-object-tracing PRIVATE code-object-tracing-client Threads::Threads
+                                $<TARGET_NAME_IF_EXISTS:rocprofiler::samples-build-flags>)
+
+add_test(NAME code-object-tracing COMMAND $<TARGET_FILE:code-object-tracing>)
+
+set_tests_properties(
+    code-object-tracing
+    PROPERTIES
+        TIMEOUT
+        45
+        LABELS
+        "samples"
+        ENVIRONMENT
+        "${ROCPROFILER_MEMCHECK_PRELOAD_ENV};HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler>"
+        FAIL_REGULAR_EXPRESSION
+        "threw an exception")
@@ -0,0 +1,405 @@
+// 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
+#include <regex>
+#ifdef NDEBUG
+#    undef NDEBUG
+#endif
+
+/**
+ * @file samples/code_object_tracing/client.cpp
+ *
+ * @brief Example rocprofiler client (tool)
+ */
+
+#include <rocprofiler/buffer.h>
+#include <rocprofiler/callback_tracing.h>
+#include <rocprofiler/fwd.h>
+#include <rocprofiler/registration.h>
+#include <rocprofiler/rocprofiler.h>
+
+#include <cxxabi.h>
+#include <atomic>
+#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 <string>
+#include <string_view>
+#include <thread>
+#include <vector>
+
+#define ROCPROFILER_CALL(result, msg)                                                              \
+    {                                                                                              \
+        rocprofiler_status_t CHECKSTATUS = result;                                                 \
+        if(CHECKSTATUS != ROCPROFILER_STATUS_SUCCESS)                                              \
+        {                                                                                          \
+            std::string status_msg = rocprofiler_get_status_string(CHECKSTATUS);                   \
+            std::cerr << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg            \
+                      << " failed with error code " << CHECKSTATUS << ": " << status_msg           \
+                      << std::endl;                                                                \
+            std::stringstream errmsg{};                                                            \
+            errmsg << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg " failure ("  \
+                   << status_msg << ")";                                                           \
+            throw std::runtime_error(errmsg.str());                                                \
+        }                                                                                          \
+    }
+
+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 code_obj_load_data_t = rocprofiler_callback_tracing_code_object_load_data_t;
+using kernel_symbol_data_t = rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t;
+using kernel_symbol_map_t  = std::unordered_map<rocprofiler_kernel_id_t, kernel_symbol_data_t>;
+
+rocprofiler_client_id_t*      client_id        = nullptr;
+rocprofiler_client_finalize_t client_fini_func = nullptr;
+rocprofiler_context_id_t      client_ctx       = {};
+kernel_symbol_map_t           client_kernels   = {};
+
+std::string
+cxa_demangle(std::string_view _mangled_name, int* _status)
+{
+    constexpr size_t buffer_len = 4096;
+    // return the mangled since there is no buffer
+    if(_mangled_name.empty())
+    {
+        *_status = -2;
+        return std::string{};
+    }
+
+    auto _demangled_name = std::string{_mangled_name};
+
+    // PARAMETERS to __cxa_demangle
+    //  mangled_name:
+    //      A NULL-terminated character string containing the name to be demangled.
+    //  buffer:
+    //      A region of memory, allocated with malloc, of *length bytes, into which the
+    //      demangled name is stored. If output_buffer is not long enough, it is expanded
+    //      using realloc. output_buffer may instead be NULL; in that case, the demangled
+    //      name is placed in a region of memory allocated with malloc.
+    //  _buflen:
+    //      If length is non-NULL, the length of the buffer containing the demangled name
+    //      is placed in *length.
+    //  status:
+    //      *status is set to one of the following values
+    size_t _demang_len = 0;
+    char*  _demang = abi::__cxa_demangle(_demangled_name.c_str(), nullptr, &_demang_len, _status);
+    switch(*_status)
+    {
+        //  0 : The demangling operation succeeded.
+        // -1 : A memory allocation failure occurred.
+        // -2 : mangled_name is not a valid name under the C++ ABI mangling rules.
+        // -3 : One of the arguments is invalid.
+        case 0:
+        {
+            if(_demang) _demangled_name = std::string{_demang};
+            break;
+        }
+        case -1:
+        {
+            char _msg[buffer_len];
+            ::memset(_msg, '\0', buffer_len * sizeof(char));
+            ::snprintf(_msg,
+                       buffer_len,
+                       "memory allocation failure occurred demangling %s",
+                       _demangled_name.c_str());
+            ::perror(_msg);
+            break;
+        }
+        case -2: break;
+        case -3:
+        {
+            char _msg[buffer_len];
+            ::memset(_msg, '\0', buffer_len * sizeof(char));
+            ::snprintf(_msg,
+                       buffer_len,
+                       "Invalid argument in: (\"%s\", nullptr, nullptr, %p)",
+                       _demangled_name.c_str(),
+                       (void*) _status);
+            ::perror(_msg);
+            break;
+        }
+        default: break;
+    };
+
+    // if it "demangled" but the length is zero, set the status to -2
+    if(_demang_len == 0 && *_status == 0) *_status = -2;
+
+    // free allocated buffer
+    ::free(_demang);
+    return _demangled_name;
+}
+
+void
+print_call_stack(const call_stack_t& _call_stack)
+{
+    namespace fs = ::std::filesystem;
+
+    auto ofname = std::string{"code_object_trace.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;
+    for(const auto& itr : _call_stack)
+    {
+        *ofs << std::left << std::setw(2) << ++n << "/" << std::setw(2) << _call_stack.size()
+             << " [" << 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);
+}
+
+template <typename Tp>
+std::string
+as_hex(Tp _v, size_t _width = 16)
+{
+    auto _ss = std::stringstream{};
+    _ss.fill('0');
+    _ss << "0x" << std::hex << std::setw(_width) << _v;
+    return _ss.str();
+}
+
+void
+tool_tracing_callback(rocprofiler_callback_tracing_record_t record,
+                      rocprofiler_user_data_t*              user_data,
+                      void*                                 callback_data)
+{
+    if(record.kind == ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT &&
+       record.operation == ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT_LOAD)
+    {
+        auto* data         = static_cast<code_obj_load_data_t*>(record.payload);
+        auto* call_stack_v = static_cast<call_stack_t*>(callback_data);
+        auto  info         = std::stringstream{};
+
+        if(record.phase == ROCPROFILER_CALLBACK_PHASE_LOAD)
+        {
+            info << "code object load :: ";
+        }
+        else if(record.phase == ROCPROFILER_CALLBACK_PHASE_UNLOAD)
+        {
+            info << "code object unload :: ";
+        }
+
+        info << "code_object_id=" << data->code_object_id
+             << ", rocp_agent=" << data->rocp_agent.handle << ", uri=" << data->uri
+             << ", load_base=" << as_hex(data->load_base) << ", load_size=" << data->load_size
+             << ", load_delta=" << as_hex(data->load_delta);
+        if(data->storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_FILE)
+            info << ", storage_file_descr=" << data->storage_file;
+        else if(data->storage_type == ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_MEMORY)
+            info << ", storage_memory_base=" << as_hex(data->memory_base)
+                 << ", storage_memory_size=" << data->memory_size;
+
+        call_stack_v->emplace_back(source_location{__FUNCTION__, __FILE__, __LINE__, info.str()});
+    }
+    if(record.kind == ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT &&
+       record.operation == ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER)
+    {
+        auto* data         = static_cast<kernel_symbol_data_t*>(record.payload);
+        auto* call_stack_v = static_cast<call_stack_t*>(callback_data);
+        auto  info         = std::stringstream{};
+
+        if(record.phase == ROCPROFILER_CALLBACK_PHASE_LOAD)
+        {
+            info << "kernel symbol load :: ";
+            client_kernels.emplace(data->kernel_id, *data);
+        }
+        else if(record.phase == ROCPROFILER_CALLBACK_PHASE_UNLOAD)
+        {
+            info << "kernel symbol unload :: ";
+            client_kernels.erase(data->kernel_id);
+        }
+
+        auto kernel_name     = std::regex_replace(data->kernel_name, std::regex{"(\\.kd)$"}, "");
+        int  demangle_status = 0;
+        kernel_name          = cxa_demangle(kernel_name, &demangle_status);
+
+        info << "code_object_id=" << data->code_object_id << ", kernel_id=" << data->kernel_id
+             << ", kernel_object=" << as_hex(data->kernel_object)
+             << ", kernarg_segment_size=" << data->kernarg_segment_size
+             << ", kernarg_segment_alignment=" << data->kernarg_segment_alignment
+             << ", group_segment_size=" << data->group_segment_size
+             << ", private_segment_size=" << data->private_segment_size
+             << ", kernel_name=" << kernel_name;
+
+        call_stack_v->emplace_back(source_location{__FUNCTION__, __FILE__, __LINE__, info.str()});
+    }
+
+    (void) user_data;
+}
+
+int
+tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data)
+{
+    assert(tool_data != nullptr);
+
+    auto* call_stack_v = static_cast<call_stack_t*>(tool_data);
+
+    call_stack_v->emplace_back(source_location{__FUNCTION__, __FILE__, __LINE__, ""});
+
+    client_fini_func = fini_func;
+
+    ROCPROFILER_CALL(rocprofiler_create_context(&client_ctx), "context creation");
+
+    ROCPROFILER_CALL(
+        rocprofiler_configure_callback_tracing_service(client_ctx,
+                                                       ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT,
+                                                       nullptr,
+                                                       0,
+                                                       tool_tracing_callback,
+                                                       tool_data),
+        "code object tracing service configure");
+
+    int valid_ctx = 0;
+    ROCPROFILER_CALL(rocprofiler_context_is_valid(client_ctx, &valid_ctx),
+                     "context validity check");
+    if(valid_ctx == 0)
+    {
+        // notify rocprofiler that initialization failed
+        // and all the contexts, buffers, etc. created
+        // should be ignored
+        return -1;
+    }
+
+    ROCPROFILER_CALL(rocprofiler_start_context(client_ctx), "context start");
+
+    // no errors
+    return 0;
+}
+
+void
+tool_fini(void* tool_data)
+{
+    assert(tool_data != nullptr);
+
+    auto* _call_stack = static_cast<call_stack_t*>(tool_data);
+    _call_stack->emplace_back(source_location{__FUNCTION__, __FILE__, __LINE__, ""});
+
+    print_call_stack(*_call_stack);
+
+    delete _call_stack;
+}
+
+void
+setup()
+{
+    if(int status = 0;
+       rocprofiler_is_initialized(&status) == ROCPROFILER_STATUS_SUCCESS && status == 0)
+    {
+        ROCPROFILER_CALL(rocprofiler_force_configure(&rocprofiler_configure),
+                         "force configuration");
+    }
+}
+}  // namespace
+
+// force configuration when library is loaded
+bool cfg_on_load = (client::setup(), true);
+}  // 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;
+
+    auto* client_tool_data = new std::vector<client::source_location>{};
+
+    client_tool_data->emplace_back(
+        client::source_location{__FUNCTION__, __FILE__, __LINE__, info.str()});
+
+    // create configure data
+    static auto cfg =
+        rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t),
+                                            &client::tool_init,
+                                            &client::tool_fini,
+                                            static_cast<void*>(client_tool_data)};
+
+    // return pointer to configure data
+    return &cfg;
+}
@@ -0,0 +1,247 @@
+// 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(const 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(const int* in, int* out, int M, int N)
+{
+    __shared__ int tile[shared_mem_tile_dim][shared_mem_tile_dim];
+
+    int idx = (blockIdx.y * blockDim.y + threadIdx.y) * M + blockIdx.x * blockDim.x + threadIdx.x;
+    tile[threadIdx.y][threadIdx.x] = in[idx];
+    __syncthreads();
+    idx      = (blockIdx.x * blockDim.x + threadIdx.y) * N + blockIdx.y * blockDim.y + threadIdx.x;
+    out[idx] = tile[threadIdx.x][threadIdx.y];
+}
+
+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
+
+    print_lock.lock();
+    printf("[%i][%i] grid=(%i,%i,%i), block=(%i,%i,%i)\n",
+           rank,
+           tid,
+           grid.x,
+           grid.y,
+           grid.z,
+           block.x,
+           block.y,
+           block.z);
+    print_lock.unlock();
+
+    auto t1 = std::chrono::high_resolution_clock::now();
+    for(size_t i = 0; i < nitr; ++i)
+    {
+        transpose_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